CUDA: как использовать барьер.sync

Question

Я прочитал https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-bar, в котором подробно описана функция синхронизации PTX.

1. В нем говорится, что существует 16 «барьерных логических ресурсов», и вы можете указать, какой барьер использоватьс параметром "а".Что такое барьерный логический ресурс?

2. У меня есть фрагмент кода из внешнего источника, который, я знаю, работает.Тем не менее, я не могу понять синтаксис, используемый внутри «asm» и что делает «память».Я предполагаю, что «имя» заменяет «% 0», а «numThreads» заменяет «% 1», но что такое «память» и что делают двоеточия?

   __device__ __forceinline__ void namedBarrierSync(int name, int numThreads) {
   asm volatile("bar.sync %0, %1;" : : "r"(name), "r"(numThreads) : "memory");}
   

3. Вблок из 256 потоков, я только хочу синхронизировать потоки 64 ~ 127.Это возможно с функцией barrier.sync?(например, скажем, у меня есть сетка из 1 блока, блока из 256 потоков. Мы разбиваем блок на 3 условных ветви: st-потоки 0 ~ 63 идут в kernel1, потоки 64 ~ 127 идут в kernel 2, а потоки 128 ~ 255перейдите к ядру 3. Я хочу, чтобы потоки в ядре 2. синхронизировались только между собой. Поэтому, если я использую определенную выше функцию "namedBarrierSync": "namedBarrierSync (1, 64)". Затем она синхронизирует только потоки 64 ~ 127 или потоки.0 ~ 63?

4. Я протестировал с помощью приведенного ниже кода (предположим, что gpuAssert - это функция проверки ошибок, определенная где-то в файле).

Вот код:

__global__ void test(int num_threads) 
{
    if (threadIdx.x >= 64 && threadIdx.x < 128) 
    {
        namedBarrierSync(0, num_threads) ;
    }
    __syncthreads();
}

int main(void) 
{
    test<<<1, 1, 256>>>(128);
    gpuAssert(cudaDeviceSynchronize(), __FILE__, __LINE_);
    printf("complete\n");
    return 1;
}

Answer 1

1. «Барьерный логический ресурс» - это оборудование, необходимое для синхронизации потоков / деформаций в блоке потоков (возможно, атомных счетчиков и т. Д.).Вам не нужно знать фактическую реализацию аппаратного обеспечения для их программирования, достаточно знать, что их доступно 16 экземпляров.
2. Как указал Роберт Кровелла в вашем кросс-посте на форуме Nvidia документация для встроенного PTX находится по адресу https://docs.nvidia.com/cuda/inline-ptx-assembly/index.html.
3. barrier.sync с именованным барьером и счетчиком потоков, равным 64, синхронизирует первые две деформации, поступающие на указанный барьер (для вычислительных возможностей до6.x) или первые 64 потока, прибывающие к названному барьеру (для вычислительных возможностей 7.0 и выше).
4. Ваш тест запускает только один поток (с выделенным ему 256 байтов разделяемой памяти), который выполняет тестыиз инструкции по синхронизацииВместо этого вы хотите запустить тестовое ядро ​​как test<<<1, 256>>>(128);.