c++11 - CUDA:如何使用barrier.sync

Question

我已阅读https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-bar其中详细介绍了PTX同步功能。

1. 它说有16个“屏障逻辑资源”，你可以通过参数“a”指定使用哪个屏障。什么是屏障逻辑资源？

2. 我有一段来自外部的代码，我知道它可以工作。但是，我无法理解“asm”内部使用的语法以及“内存”的作用。我假设“name”替换“%0”，“numThreads”替换“%1”，但是“内存”是什么以及冒号在做什么？

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

3. 在 256 个线程的 block 中，我只希望线程 64 ~ 127 进行同步。这可以通过 barrier.sync 实现吗 功能？ (举个例子，假设我有一个 1 block 的网格，256 个线程的 block 。我们将 block 分成 3 个条件分支，线程 0 ~ 63 进入内核 1，线程 64 ~ 127 进入内核 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. “屏障逻辑资源”是同步线程 block 中的线程/线程束所需的硬件(可能是原子计数器等)。您不需要知道实际的硬件实现来对它们进行编程，知道有 16 个可用实例就足够了。
2. 正如 Robert Crovella 在您的 cross-post 中指出的那样在 Nvidia 论坛上，内联 PTX 的文档位于 https://docs.nvidia.com/cuda/inline-ptx-assembly/index.html .
3. barrier.sync命名屏障和线程数为 64 会同步到达命名屏障的前两个 warp(对于高达 6.x 的计算能力)或到达命名屏障的前 64 个线程(对于计算能力 7.0 及以上)。
4. 您的测试仅启动一个线程(分配有 256 字节的共享内存)，这使得同步指令的测试毫无意义。您想要将测试内核启动为 test<<<1, 256>>>(128);相反。