CUDA动态并行和全局内存同步-Java 学习之路

我无法弄清楚以下内容 .

如果我启动一个内核并考虑例如 0 块中的线程 0 ，在 __syncthreads() 调用之后，所有其他块中的所有其他线程是否会看到块 0 中的线程 0 对全局内存所做的更改？

我的猜测是否定的 . 实际上，在CUDA C编程指南的synchronization functions部分中，声明：

void __syncthreads（）;等待直到线程块中的所有线程都到达此点，并且__syncthreads（）之前由这些线程进行的所有全局和共享内存访问对块中的所有线程都是可见的 .

但是，在谈论动态并行性中的global memory consistency时，CUDA C编程指南指出：

仅在第二次__syncthreads（）调用之后，这些修改才可用于父网格的其他线程 .

那么，当涉及动态并行性时， __syncthreads() 是否也可以跨块进行更改？

谢谢

1 回答

4
__syncthreads() 执行的唯一操作是您在CUDA C编程指南中描述的引用 . 除了在多个内核启动中划分内核执行的天真方法之外，CUDA中没有办法跨块进行同步，并且在性能方面存在所有缺点 . 因此，你自己猜到的第一个问题的答案是否定的 .

在帖子的第二部分中，您指的是CUDA C编程指南的一个特定示例，即
```
__global__ void child_launch(int *data) {
    data[threadIdx.x] = data[threadIdx.x]+1;
}

__global__ void parent_launch(int *data) { 
    data[threadIdx.x] = threadIdx.x;

    __syncthreads();

    if (threadIdx.x == 0) {
        child_launch<<< 1, 256 >>>(data);
        cudaDeviceSynchronize();
    }

    __syncthreads();
}

void host_launch(int *data) {
    parent_launch<<< 1, 256 >>>(data);
}
```
这里， parent_launch 内核的所有 256 个线程在 data 中写入了一些东西 . 之后，线程 0 调用 child_launch . 需要第一个 __syncthreads() 来确保在子内核调用之前完成所有内存写入 . 在这一点上引用指南：

由于第一次__syncthreads（）调用，子进程将看到数据[0] = 0，数据[1] = 1，...，数据[255] = 255（没有__syncthreads（）调用，只有数据[0 ]将保证被孩子看到） .

关于第二个 __syncthreads() ，指南解释了这一点

当子网格返回时，线程0保证看到其子网格中的线程所做的修改 . 只有在第二次__syncthreads（）调用之后，这些修改才可用于父网格的其他线程 .

在该特定示例中，第二个 __syncthreads() 是冗余的，因为由于内核终止而存在隐式同步，但是在子内核启动之后必须执行其他操作时需要第二个 __syncthreads() .

最后，关于你在帖子中引用的句子：

仅在第二次__syncthreads（）调用之后，这些修改才可用于父网格的其他线程

请注意，在具体示例中， host_launch 函数只启动了一个线程块 . 这可能有点误导了你 .

在NVIDIA论坛上有一个有趣的讨论（可能甚至不止一个），涉及跨越块的线程同步

Synchronize all blocks in CUDA
回复于 2024-04-30T21:41:16+08:00

CUDA动态并行和全局内存同步

1 回答

相关问题