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 }

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

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

我该如何解决？

4

3 回答 3

7

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

长解释

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

（我还没有使用 Kepler 架构的经验。如果不使用 Fermi，其中一些数字可能会有所不同。）

需要解释一些术语以理解下一部分：以下术语与逻辑（软件构造中的逻辑）线程有关：

  • thread – 单个执行线程。
  • 块——一组执行相同内核的多个线程。
  • 网格——一组块。

以下术语与物理（与硬件架构相关的物理）线程有关：

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

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

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

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

几个线程块可以依次组合成一个网格。网格是块的 3 维数组。最大块大小与可用的硬件资源相关，而网格可以是（几乎）任意大小。网格内的块只能通过全局内存共享数据，全局内存是延迟最高的 GPU 内存。

Fermi GPU 可以在每个 SM 上同时激活 48 个 warp（1536 个线程），因为线程使用足够少的本地和共享内存来同时容纳所有线程。线程之间的上下文切换很快，因为寄存器被分配给线程，因此不需要在线程切换之间保存和恢复寄存器和共享内存。结果是实际上需要过度分配硬件，因为它将通过让warp调度程序在发生停顿时切换当前活动的warp来隐藏内核内的内存停顿。

重要的部分

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

whileand语句是分支语句，因此for不限于if.

于 2012-11-15T12:23:41.637 回答
2

当您使用多个线程启动内核时，您会遇到一个无限循环，因为while(t<tid);对于任何idx大于零的线程来说都是一个无限循环。

此时，您的问题与线程同步无关，而是与您实现的循环有关。

于 2012-11-15T12:43:03.023 回答
1

如果您要做的是让一系列线程串行执行，那么您就是在滥用 CUDA。

它也行不通，因为第一个线程之后的任何线程都不会收到更新的 t - 您必须调用 __syncthreads() 以刷新共享变量，但只有在所有线程都执行相同的事情时才能这样做- 即不等待。

于 2012-11-15T12:29:46.793 回答