北京
分享至
一段CUDA的并行前缀和内核代码，乍看之下像是会发生数据竞争——多个线程似乎在同一时间访问同一块共享内存。但仔细分析,这件事并没有真的发生危险。
关键在于两个因素被妥善结合了起来：__syncthreads()这个栅栏同步，以及线程索引上的条件约束。__syncthreads()强制同一个线程块里的所有线程都到达调用点之后，才继续往下执行。在这道栅栏之前的所有共享内存写入，会变得对块内所有线程可见；栅栏之后的读取，用的都是安全的状态。
而线程索引条件则精确划定了哪些线程做加载、哪些做累加、哪些只旁观。这样一来，即使多个线程写着似乎是同一处地址，时间上也已经被栅栏切成了“前一批写入、全体同步、后一批读取”的节奏，不会出现同时读写的冲突。
所以这段内核表面上像在挑战并行安全，实际上是通过对同步和索引条件的分析，被证明是干净的并行实现。看上去冲突的原因，只是因为线性阅读代码的直觉与栅栏分割出的执行序没有对上。
特别声明：以上内容(如有图片或视频亦包括在内)为自媒体平台“网易号”用户上传并发布，本平台仅提供信息存储服务。
Notice: The content above (including the pictures and videos if any) is uploaded and posted by a user of NetEase Hao, which is a social media platform and only provides information storage services.
