CUDA:如何使用barrier.sync

Question

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

1. 它说有16个“屏障逻辑资源”，您可以指定参数"a“使用哪个屏障。什么是障碍逻辑资源？
2. 我有一段来自外部来源的代码，我知道这是可行的。但是，我无法理解"asm“中使用的语法以及”内存“的功能。我假设"name“代替"%0”，"numThreads“代替"%1"，但是”内存“是什么?冒号在做什么？ __device__ __forceinline__ void namedBarrierSync(int名称，int numThreads) {bar.sync %0，%1；：：“r”(名称)，"r"(numThreads)：“内存”)；}
3. 在一个由256个线程组成的块中，我只希望线程64 ~ 127同步。这在barrier.sync函数中是可能的吗？(例如，我有一个1块的网格，256个线程块。我们将块分成三个条件分支。线程0~ 63进入kernel1，线程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. “屏障逻辑资源”是同步线程块中的线程/翘曲(可能是原子计数器等)所必需的硬件。您不需要知道实际的硬件实现来编程，只需知道它们有16个可用实例就足够了。
2. 正如Robert在Nvidia论坛的交叉柱中指出的那样，内联PTX的文档位于https://docs.nvidia.com/cuda/inline-ptx-assembly/index.html。
3. 命名屏障和线程计数为64的barrier.sync同步到达指定屏障的前两个偏差(计算能力最高可达6.x)或到达指定屏障的前64个线程(对于计算能力7.0以后)。
4. 您的测试只启动一个线程(分配给它的共享内存为256字节)，这使得同步指令的测试毫无意义。您希望以test<<<1, 256>>>(128);的形式启动测试内核。