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

Question

如果一个CUDA内核调用了一个__device__函数，那么在进入或退出__device__函数时，所有线程之间是否存在任何隐式同步？

如果没有，那么这意味着在没有任何显式同步的情况下，一些线程可能已经退出了__device__函数，而其他线程甚至还没有进入它？

如果有相关信息或参考资料，请提供。

- Paul Caheny

---

在同一束线程中，线程之间只有隐式的同步。否则，需要使用显式同步原语。 - talonmies

Answer 1

是的，正如talonmies所指出的那样，在warp中只有隐式同步。当内核启动时，硬件会窥视任何一个warp（可能是第一个），并为其执行第一条指令，然后切换到另一个warp。由于warp在退出时不需要等待彼此，因此有可能某些warp在另一个warp甚至没有执行内核的第一条指令之前就被退役了。

- user1545642

Answer 2

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

- Bardia