Использование CUDA __syncthreads() в варпе

Question

Если абсолютно необходимо, чтобы все потоки в блоке находились в одной и той же точке кода, требуется ли нам функция __syncthreads, если количество запускаемых потоков равно количеству потоков в деформации?

Примечание: никаких дополнительных потоков или блоков, только один перекос для ядра.

Пример кода:

shared _voltatile_ sdata[16];

int index = some_number_between_0_and_15;
sdata[tid] = some_number;
output[tid] = x ^ y ^ z ^ sdata[index];

Answer 1

Обновлено с дополнительной информацией об использовании volatile

Предположительно, вы хотите, чтобы все потоки находились в одной точке, поскольку они читают данные, записанные другими потоками, в общую память. Если вы запускаете одну деформацию (в каждом блоке), то вы знаете, что все потоки выполняются вместе. На первый взгляд это означает, что вы можете опустить __syncthreads(), практика, известная как "деформация синхронного программирования". Однако есть несколько вещей, на которые стоит обратить внимание.

• Помните, что компилятор будет предполагать, что он может оптимизировать, если семантика внутри потока остается верной, включая задержку хранения в памяти, где данные могут храниться в регистрах. __syncthreads() действует как барьер для этого и, следовательно, гарантирует, что данные будут записаны в общую память, прежде чем другие потоки прочитают данные. С помощью volatile заставляет компилятор выполнять запись в память, а не хранить ее в регистрах, однако это сопряжено с некоторыми рисками и является скорее хаком (то есть я не знаю, как это повлияет в будущем)
  • Технически, вы всегда должны использовать __syncthreads() соответствовать модели программирования CUDA
• Размер основы равен и всегда был равен 32, но вы можете:
  • Во время компиляции используйте специальную переменную warpSize в коде устройства (задокументировано в Руководстве по программированию CUDA, в разделе "встроенные переменные", раздел B.4 в версии 4.1)
  • Во время выполнения используйте поле warpSize структуры cudaDeviceProp (задокументировано в Справочном руководстве CUDA)

Обратите внимание, что некоторые из образцов SDK (особенно сокращение и сканирование) используют эту синхронизирующую технологию.

Answer 2

Тебе еще нужно __syncthreads() даже если деформации выполняются параллельно. Фактическое выполнение в аппаратном обеспечении может не быть параллельным, потому что число ядер в SM (потоковом мультипроцессоре) может быть меньше 32. Например, архитектура GT200 имеет 8 ядер в каждом SM, поэтому вы никогда не сможете быть уверены, что все потоки находятся в та же точка в коде.