NVidia CUDA: синхронизация и shared memory

lexa - 02/Мар/2007 15:19
Экспериментально выяснилось, что содержимое shared memory не сохраняется между запусками кода на G80. Всякий раз оно инициализировано, причем значения разные, то 10 (float), то 125.

Плакала идея синхронизации между мультипроцессорами путем завершения kernel. Нет, синхронизироваться можно, конечно, но если хочется сохранить результат, то надо писать в глобальную память.

Comments

Виктор (not verified) - 02/Мар/2007 20:38

Время жизни shared переменных - до завершения CTA. Даже если бы GRF сохранял своё значение, появилось бы множество технических и логических трудностей: например, как связать переменные из первого и второго ядра? Скорее всего, пришлось бы требовать, чтобы __shared__ переменные были объявлены абсолютно одинаково и в таком же порядке. Что делать, в тех случаях, когда это не так? Даже если общий размер, типы, и порядок объявления совпадает, присутствует неопределённость: CTA из второго ядра не смогут автоматически узнать (только вручную), от какого CTA ему достался фрейм GRF.

lexa - 02/Мар/2007 21:06

Ну вот я бы не отказался от сайд-эффекта (документированного) - ваши байты на мультипроцессоре остались в том же порядке.

С другой стороны, действительно неизвестно, какой thread block жил на этой shared memory, но можно выйти из положения или явно (писать blockid туда же) или неявно - задавая ту же конфигурацию блоков-тредов и предполагая (лучше тоже документированно), что их распределение такое же.

Остается еще вопрос "что делать, когда блоков/тредов очень очень много", ну да это мы обсуждаем в личной переписке.