问CUDA:同一块中的线程同步
EN

Stack Overflow用户

提问于 2012-11-15 12:18:36

回答 3查看 2.2K关注 0票数 2

我正在尝试写CUDA程序，但我有问题，在同一块线程之间的同步。

以下是模型情况：

 10 __global__ void gpu_test_sync()
 11 {
 12     __shared__ int t;
 13     int tid = threadIdx.x;
 14
 15     t = 0;
 16     __threadfence();
 17     __syncthreads();
 18
 19     // for(int i=0; i<1000000 && t<tid; i++); // with fuse
 20     while(t<tid);
 21
 22     t++;
 23     __threadfence();
 24 }
 25
 26 void f_cpu()
 27 {
 28     printf("TEST ... ");
 29     int blocks = 1;
 30     int threads = 2;
 31     gpu_test_sync<<< blocks , threads >>>();
 32     printf("OK\n");
 33 }

如果线程数= 1，则一切正常。如果线程数> 1，则无限循环。

为什么？函数__threadfence()；应使t变量的值对其他线程可见。

我该怎么解决它呢？

cuda

thread-synchronization

回答 3

Stack Overflow用户

回答已采纳

发布于 2012-11-15 12:23:42

我不相信你的内核能够完成你想要做的事情，因为while(t<tid)中的分支会导致warp的所有线程无限循环，并且永远不会到达++t行。

冗长的解释

如果您已经了解线程、块和扭曲，请滚动到“重要部分”查找重要内容：

(到目前为止，我还没有使用开普勒架构的经验。如果不使用费米，这些数字中的一些可能会不同。)

为了理解下一节，需要解释一些术语:以下术语与逻辑(与软件构造中的逻辑相同)线程相关：

网格线程- execution.

block的单线程-一组执行相同内核的多线程。
-一组块。

以下术语与物理(与硬件架构相关的物理线程)相关：

内核-单个计算内核，一个内核一次只能运行一条指令。
warp -在硬件上并行执行的一组线程，一个warp由当前代CUDA硬件上的32个线程组成。

内核由一个或多个流式多处理器(SM)执行。费米系列(GeForce 400和GeForce 500系列)的典型中高端GPU卡在单个GeForce [Fermi whitepaper]上具有8-16个SMs。每个SM由32个CUDA核心(核心)组成。线程由warp调度器调度执行，每个SM都有两个以锁步方式工作的warp调度器单元。warp调度器可以调度的最小单元称为warp，它由在编写本文时发布的所有CUDA硬件上的32个线程组成。在每个SM上一次只能执行一次warp。

CUDA中的线程比CPU线程轻得多，上下文切换更便宜，warp中的所有线程都执行相同的指令，或者在warp中的其他线程执行指令时必须等待，这称为单指令多线程(SIMT)，并且类似于传统的CPU单指令多数据(SIMD)指令，如SSE、AVX、NEON、Al- tivec等，这在使用条件语句时会产生后果，如下面所述。

为了解决需要超过32个线程才能解决的问题，CUDA线程被安排到由软件开发人员定义的称为块和网格大小的逻辑组中。块是线程的三维集合，块中的每个线程都有自己的单独的三维标识号，以允许开发人员区分内核代码中的线程。单个块中的线程可以通过共享内存共享数据，这减少了全局内存的负载。共享内存的延迟比全局内存低得多，但这是一种有限的资源，用户可以在(每块) 16 kB共享内存和48 kB L1缓存或48 kB共享内存和16 kB L1缓存之间进行选择。

几个线程块可以依次分组到一个网格中。栅格是块的三维数组。最大块大小与可用的硬件资源有关，而网格可以是(几乎)任意大小。网格中的块只能通过全局内存共享数据，全局内存是具有最高延迟的GPU内存。

一个Fermi GPU可以在每个SM中一次有48个warps (1536个线程)处于活动状态，假设线程使用的本地和共享内存很少，可以同时容纳所有线程。线程之间的上下文切换是快速的，因为寄存器被分配给线程，因此不需要保存和恢复寄存器以及线程切换之间的共享存储器。结果是，实际上需要过度分配硬件，因为它将通过让warp调度器在任何时候发生停顿时切换当前活动的warp来隐藏内核内的内存停顿。

重要的部分

线程扭曲是在同一流式多处理器(SM)上执行的一组硬件线程。可以将warp的线程比作在线程之间共享公共程序计数器，因此所有线程必须执行相同的程序代码行。如果代码有一些分支语句，比如if ... then ... else，则warp必须首先执行进入第一个块的线程，而warp的其他线程在等待，接下来进入下一个块的线程将执行，而其他线程则等待，依此类推。由于这种行为，应尽可能避免在GPU代码中使用条件语句。当一个warp的线程遵循不同的执行线时，它被称为具有不同的线程。虽然条件块应该在CUDA内核中保持最小，但有时可以对语句进行重新排序，以便相同warp的所有线程在if ... then ... else块中只遵循一条执行路径，从而减轻这一限制。

while和for语句是分支语句，因此并不局限于if。

票数 7