CUDA：如何使用 barrier.sync

Posted 2023-02-16

【中文标题】CUDA：如何使用 barrier.sync

【英文标题】：CUDA: how to use barrier.sync

【发布时间】：2018-12-07 02:40:40

【问题描述】：

我已经阅读了https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-bar关于PTX同步功能的详细信息。

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

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

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

在一个有 256 个线程的块中，我只希望线程 64 ~ 127 同步。这可能与barrier.sync 功能？ （例如，假设我有 1 个块的网格，256 个线程的块。我们将块分成 3 个条件分支 s.t. 线程 0 ~ 63 进入内核 1，线程 64 ~ 127 进入内核 2，线程 128 ~ 255进入内核 3。我希望内核 2 中的线程只在它们之间同步。所以如果我使用上面定义的“namedBarrierSync”函数：“namedBarrierSync(1, 64)”。那么它只同步线程 64 ~ 127，还是线程0~63？

我已经用下面的代码进行了测试（假设 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;

【参考方案1】：

“屏障逻辑资源”是同步线程块中的线程/扭曲所必需的硬件（可能是原子计数器等）。您无需了解实际的硬件实现即可对其进行编程，知道有 16 个可用实例就足够了。 正如 Robert Crovella 在 Nvidia 论坛上的 cross-post 中指出的那样，内联 PTX 的文档位于 https://docs.nvidia.com/cuda/inline-ptx-assembly/index.html。 barrier.sync 具有命名屏障和 64 线程计数同步到达命名屏障的前两个扭曲（计算能力高达 6.x）或到达命名屏障的前 64 个线程（计算能力 7.0 及以上）。 您的测试只启动了一个线程（分配了 256 字节的共享内存），这使得同步指令的测试没有实际意义。您希望以 test<<<1, 256>>>(128); 的身份启动测试内核。