2012-08-07 47 views
1

哪一个更好,单个Warp的线程之间或不同Warps线程之间的原子竞争(并发)?我认为,当你访问共享内存的时候,一个warp的线程彼此竞争的时候会比不同warp的线程少。而相反地访问全局内存,最好是一个块的不同线程的线程比单个线程的线程竞争更少,不是吗?哪个更好,原子之间的竞争:单个Warp的线程还是不同的Warps线程?

我需要它来知道如何更好地解决竞争(并发)和更好地分隔存储:单线程之间或线程之间的线程。

顺便说一下,可以说团队__ syncthreads();同步它在一个块中变形而不是一个变形的线程?

回答

2

如果块中有相当数量的线程对同一个值执行原子更新,那么性能会很差,因为这些线程都必须被序列化。在这种情况下,让每个线程将其结果写入单独的位置,然后在单独的内核中处理这些值通常会更好。

如果warp中的每个线程都执行原子更新为相同的值,则warp中的所有线程都会在相同的时钟周期内执行更新,因此它们必须在原子更新点处全部序列化。这可能意味着该warp预定了32次以获得所有服务的线程(非常糟糕)。另一方面,如果块中的每个变形中的单个线程执行原子更新为相同的值,则该影响将会较低,这是因为经线对(在两个经线在每个时钟处处理的两个经线调度程序)在时间上偏移(一个时钟周期),因为它们在处理管道中移动。所以最终只有两个原子更新(两个经线中的每一个),在一个周期内发布并需要立即被序列化。

因此,在第二种情况下,情况较好,但仍存在问题。原因在于,取决于共享值的位置,您仍然可以在SM之间进行序列化,并且这可能非常缓慢,因为每个线程可能必须等待更新才能全部移出到全局内存,或者至少L2,然后回来。有可能以这样的方式重构算法,即块内的线程对共享内存(L1)中的值执行原子更新,然后每个块中的一个线程对全局内存中的值执行原子更新(L2 )。

原子操作可以是完整的救生员,但他们往往被CUDA新手过度使用。通常使用单独步骤并行简化或并行流压缩算法更好(请参阅thrust::copy_if)。

相关问题