我不知道以下内容。
如果我启动一个内核和考虑,例如,线程0
块0
,一个后__syncthreads()
来电,都会在所有其它模块看到线程全局内存所做的更改其他线程0
块0
?
我的猜测不是。实际上,在《CUDA C编程指南》的同步功能部分中,指出:
void __syncthreads();
等待直到线程块中的所有线程都到达这个点,并且这些线程之前的所有全局和共享内存访问对该块中的所有线程__syncthreads()
都可见。
但是,在谈论动态并行性中的全局内存一致性时,《 CUDA C编程指南》指出:
仅在第二次调用之后,这些修改才对父网格的其他线程可用
__syncthreads()
。
__syncthreads()
当涉及动态并行性时,还可以使更改跨块可用吗?
谢谢
唯一执行的操作__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);
}
在这里,内核的所有256
线程都在中parent_launch
写入了一些内容data
。之后,线程0
调用child_launch
。首先__syncthreads()
需要确保在子内核调用之前完成所有内存写入。在这一点上引用指南:
由于有第一个
__syncthreads()
电话,孩子会看到data[0]=0
,,data[1]=1
...data[255]=255
(没有__syncthreads()
电话,只能data[0]
保证孩子会看到它)。
关于第二个__syncthreads()
,《指南》解释说,
当子网格返回时,
0
保证线程看到其子网格中的线程所做的修改。仅在第二次__syncthreads()
调用之后,这些修改才对父网格的其他线程可用。
在该特定示例中,第二个__syncthreads()
是多余的,因为由于内核终止而存在隐式同步,但是__syncthreads()
当在子内核启动之后必须执行其他操作时,第二个就变得很有用。
最后,关于您在帖子中引用的句子:
只有在第二次
__syncthreads()
调用之后,这些修改才可用于父网格的其他线程
请注意,在特定示例中,该host_launch
函数仅启动一个线程块。这也许在某种程度上误导了您。
NVIDIA论坛上有一个有趣的讨论(可能甚至超过一个),涉及跨块的线程同步,标题为
本文收集自互联网,转载请注明来源。
如有侵权,请联系 [email protected] 删除。
我来说两句