对于__syncthreads()的理解

CUDA官方文档的原文是：“__syncthreads()is allowed in conditional code but only if the conditional evaluates identically across the entire thread block, otherwise the code execution is likely to hang or produce unintended side effects.”

对于evaluates identically我的理解是那些先到达__syncthreads()的线程等待其他可以到达该点的线程，而不是等待块内所有其他线程[1]。

一篇文章里提到“__syncthreads() is a barrier that waits for all warps to either call __syncthreads() or terminate”[2]，他提到__syncthreads()中同步的粒度在实际操作中不是block而是warps(as a barrier for all warps in the same block)，一个warp到达__syncthreads()需要等待其他warps到达__syncthreads()或者终止才能继续运行。

参考

[1] http://www.cnblogs.com/dwdxdy/p/3215136.html

[2] Demystifying GPU Microarchitecture through Microbenchmarking