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

Question

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

如果不是那么这意味着 block 中的一些线程可能在 block 中的其他线程进入它之前就已经退出了 __device__ 函数(在没有任何显式同步的情况下)？

任何指向相关信息/引用的指针都将不胜感激。

Answer 1

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