use*_*543 2 cuda thread-synchronization
我正在尝试在 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 变量的可见值。
我该如何解决？
我不相信你的内核将能够做你想要做的事情，因为发散分支while(t<tid)导致经线的所有线程无限循环并且永远不会到达 line ++t。
如果您已经了解线程、块和经线，请滚动到重要内容的“重要部分”：
（我还没有使用开普勒架构的经验。如果不使用费米，其中一些数字可能会有所不同。）
需要解释一些术语才能理解下一节： 以下术语与逻辑（如软件结构中的逻辑）线程有关：
以下术语与物理（与硬件架构相关的物理）线程相关：
内核由一个或多个流式多处理器 (SM) 执行。Fermi 系列（GeForce 400 和 GeForce 500 系列）的典型中高端 GeForce 卡在单个 GPU 上具有 8-16 个 SM [ Fermi 白皮书]。每个 SM 由 32 个 CUDA Cores（核心）组成。线程由 warp 调度器调度执行，每个 SM 有两个以锁步方式工作的 warp 调度器单元。warp 调度器可以调度的最小单元称为 warp，它由截至撰写本文时发布的所有 CUDA 硬件上的 32 个线程组成。每个 SM 上一次只能执行一个扭曲。
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 可以有 48 个扭曲（1536 个线程），每个 SM 一次处于活动状态，因为这些线程使用很少的本地和共享内存来同时容纳所有线程。线程之间的上下文切换很快，因为寄存器被分配给线程，因此不需要在线程切换之间保存和恢复寄存器和共享内存。结果是实际上需要过度分配硬件，因为它会通过让warp调度程序在发生停顿时切换当前活动的warp来隐藏内核内部的内存停顿。
线程扭曲是在同一流多处理器(SM)上执行的硬件线程组。Warp 的线程可以比作线程之间共享一个公共程序计数器，因此所有线程都必须执行同一行程序代码。如果代码有一些分支语句，例如if ... then ... elsewarp 必须首先执行进入第一个块的线程，而 warp 的其他线程等待，接下来进入下一个块的线程将在其他线程等待时执行，依此类推。由于这种行为，应尽可能避免在 GPU 代码中使用条件语句。当一个经线的线程遵循不同的执行线时，它被称为具有发散线程。虽然 CUDA 内核中的条件块应该保持在最低限度，但有时可以重新排序语句，以便相同经线的所有线程仅遵循if ... then ... else块中的单个执行路径并减轻此限制。
该while和for语句分支语句，所以它不局限于if。