单个 Warp 的线程之间的原子竞争（并发）和一个块中不同 Warp 的线程之间哪个更好？我认为，当您访问共享内存时，一个经线的线程相互竞争比不同经线的线程少。而访问全局内存则相反，一个block的不同warp的线程竞争比单个warp的线程少，不是吗？
我需要它知道如何更好地解决竞争（并发）以及如何更好地分离存储：在单个 warp 中的线程之间或 warp 之间。
顺便说一句，可以说团队__syncthreads();同步它在一个块中的扭曲而不是一个扭曲的线程？
【问题讨论】：
标签： cuda synchronization gpgpu
如果块中的大量线程对相同的值执行原子更新，则性能会很差，因为这些线程必须全部序列化。在这种情况下，通常最好让每个线程将其结果写入单独的位置，然后在单独的内核中处理这些值。
如果 warp 中的每个线程对相同的值执行原子更新，则 warp 中的所有线程都在同一时钟周期内执行更新，因此它们必须在原子更新点都被序列化。这可能意味着 warp 被安排了 32 次以得到所有线程的服务（非常糟糕）。
另一方面，如果块中每个 warp 中的单个线程对相同的值执行原子更新，则影响会较小，因为成对的 warp（两个 warp 调度程序在每个时钟处理的两个 warp ) 在时间上偏移（一个时钟周期），因为它们在处理管道中移动。所以你最终只得到了两个原子更新（两个warp中的一个），在一个周期内发布，需要立即序列化。
所以，在第二种情况下，情况更好，但仍然存在问题。原因是，根据共享值的位置，您仍然可以在 SM 之间进行序列化，这可能非常慢，因为每个线程可能必须等待更新一直到全局内存，或者至少要等待L2，然后返回。可以重构算法，使块内的线程对共享内存 (L1) 中的值执行原子更新，然后让每个块中的一个线程对全局内存 (L2) 中的值执行原子更新)。
原子操作可以成为救命稻草，但它们往往会被 CUDA 新手过度使用。通常最好使用带有并行缩减或并行流压缩算法的单独步骤（请参阅thrust::copy_if）。
【讨论】：