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