cuda - CUDA __device__ 函数的开头或结尾是否存在隐式屏障同步？

Question

如果 CUDA 内核调用 __device__ 函数，在进入或退出 __device__ 函数的块中的所有线程之间是否存在任何隐式同步？

如果不是，那么这意味着块中的某些线程可能在块中的其他线程甚至进入它之前退出 __device__ 函数（在没有任何显式同步的情况下）？

任何指向相关信息/参考的指针将不胜感激。

Answer 1

是的，正如 talonmies 指出的那样，经线中只有隐式同步 btw 线程。当内核启动时，硬件会查看任何经线（可能是第一个经线）并为其执行第一条指令，然后切换到另一个经线。有可能一些warp在另一个warp甚至执行内核的第一条指令之前就退休了，因为warp在退出时不需要相互等待

Answer 2

您应该__syncthreads()在内核代码中使用函数。调用该__device__函数后，添加__syncthreads();行以放置障碍并同步线程。