如果块中有相当数量的线程对同一个值执行原子更新,那么性能会很差,因为这些线程都必须被序列化。在这种情况下,让每个线程将其结果写入单独的位置,然后在单独的内核中处理这些值通常会更好。
如果warp中的每个线程都执行原子更新为相同的值,则warp中的所有线程都会在相同的时钟周期内执行更新,因此它们必须在原子更新点处全部序列化。这可能意味着该warp预定了32次以获得所有服务的线程(非常糟糕)。另一方面,如果块中的每个变形中的单个线程执行原子更新为相同的值,则该影响将会较低,这是因为经线对(在两个经线在每个时钟处处理的两个经线调度程序)在时间上偏移(一个时钟周期),因为它们在处理管道中移动。所以最终只有两个原子更新(两个经线中的每一个),在一个周期内发布并需要立即被序列化。
因此,在第二种情况下,情况较好,但仍存在问题。原因在于,取决于共享值的位置,您仍然可以在SM之间进行序列化,并且这可能非常缓慢,因为每个线程可能必须等待更新才能全部移出到全局内存,或者至少L2,然后回来。有可能以这样的方式重构算法,即块内的线程对共享内存(L1)中的值执行原子更新,然后每个块中的一个线程对全局内存中的值执行原子更新(L2 )。
原子操作可以是完整的救生员,但他们往往被CUDA新手过度使用。通常使用单独步骤并行简化或并行流压缩算法更好(请参阅thrust::copy_if
)。