【问题标题】：CUDA: Thread synchronization in the same blockCUDA：同一块中的线程同步
【发布时间】：2012-11-15 21:17:58
【问题描述】：

我正在尝试在 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

【解决方案1】：

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

详细解释

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

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

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

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

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

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

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

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

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

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

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

重要的部分

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

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

【讨论】：

我认为线程在块中是独立的。函数 __syncthreads();是用于不同块的线程之间的同步吗？
warp 中的每个线程都试图以同步的方式执行代码 - cuda 的全部目的是在同一时间做很多次大致相同的事情。
一个warp的所有线程都执行相同的指令，或者等待，它们不能并行执行内核的不同部分
好的。块和经线有什么区别？例如，如果我调用 gpu_f>> 那么块有两个扭曲（32 个线程）？还是？
是的，gpu_f<<<1,64>>> 应该使用两个经纱执行。块是线程的逻辑分组，warp 是在芯片中实现 GPU 时来自设计选择的低级硬件分组。我更新了我的答案，介绍了块、扭曲和线程。

【解决方案2】：

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

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

【讨论】：

【解决方案3】：

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

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

【讨论】：