我正在研究 GPU 编程，并且有一个关于在线程中修改全局数组的问题。
__device__ float data[10] = {0,0,0,0,0,0,0,0,0,1};
__global__ void gradually_set_global_data() {
while (1) {
if (data[threadIdx.x + 1]) {
atomicAdd(&data[threadIdx.x], data[threadIdx.x + 1]);
break;
}
}
}
int main() {
gradually_set_global_data<<<1, 9>>>();
cudaDeviceReset();
return 0;
}
内核应该在预期保持 [1,1,1,1,1,1,1,1,1,1] 的情况下完成执行data，但它陷入了无限循环。为什么会这样？
TL;DR：代码被检查破坏。CUDA 线程模型不保证任何特定线程的前进进度，除非符合以下条件：
CUDA 编程模型未定义为第 1 项选择的线程。除非程序员使用执行屏障进行显式控制，否则 CUDA 线程模型可以随意调度单个线程，直到该线程退出或遇到显式执行屏障。
由于提供的代码没有执行障碍，CUDA 工作调度程序（相对于 CUDA 语义）可以自由调度，例如线程 0，而没有其他线程。如果我们将这个概念应用到提供的代码中，很明显线程 0，如果它自己运行，将表现出无限循环。
更长：
这恰好是观察到的行为，尽管如果是我，我不会将两者联系起来。挂起的原因（根据我试图描述的方式）不是“为了正确性，这段代码依赖于 CUDA 编程模型不提供的保证”，尽管我相信这是一个真实的陈述。要了解挂起的原因，我建议有必要在查看 SASS（机器汇编代码）的情况下检查低级机器行为。我真的没有能力穷尽这个话题，所以我将对此进行有限的看法。
为什么要做出这种区分？因为对所提供代码的相对较小的更改，实际上并没有解决正确性问题，可能会导致编译器生成不会挂起的代码。缺乏仔细的处理可能会导致人们得出结论，因为它没有挂起，所以它一定没问题。关键是代码是否挂起与它是否正确是不同的。我已经向自己证明了这一点。但是，我不想提供该代码。正确的做法是设计正确的代码。请参阅下文了解我的尝试。
在我们深入研究 SASS 之前，我想指出代码中的另一个缺陷。CUDA 编译器可以自由地将任何全局数据“优化”到寄存器中，同时保持单线程语义正确性。编译器主要有一个单一的线程，因此这可能会使依赖于线程间通信的程序员绊倒（就像这段代码一样）。为了正确起见，在此代码中，线程 x 修改的数据必须（最终）对线程 x-1 可见。CUDA 编程模型不保证这种线程间可见性，编译器通常也不强制执行它。为了正确起见，有必要通知编译器使这些数据可见，并订购加载和存储以实现这一点。有多种方法可以做到这一点。我会建议为简单起见，用 标记数据volatile，尽管可以使用也具有内置内存屏障的__syncthreads()执行屏障（例如）来执行此操作。无论选择哪种方法来强制执行线程间数据可见性，如果没有它，代码就会被破坏，与任何其他考虑无关。 __syncwarp()
因此，在深入研究 SASS 之前，我建议对所提供的代码以及其后的 SASS 进行以下修改：
$ cat t1691.cu
__device__ volatile float data[10] = {0,0,0,0,0,0,0,0,0,1};
__global__ void gradually_set_global_data() {
while (1) {
if (data[threadIdx.x + 1]) {
atomicAdd((float *)&data[threadIdx.x], data[threadIdx.x + 1]);
break;
}
}
}
int main() {
gradually_set_global_data<<<1, 9>>>();
cudaDeviceReset();
return 0;
}
$ nvcc -o t1691 t1691.cu
$ cuobjdump -sass ./t1691
Fatbin elf code:
================
arch = sm_30
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
code for sm_30
Fatbin elf code:
================
arch = sm_30
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
code for sm_30
Function : _Z25gradually_set_global_datav
.headerflags @"EF_CUDA_SM30 EF_CUDA_PTX_SM(EF_CUDA_SM30)"
/* 0x22f2c04272004307 */
/*0008*/ MOV R1, c[0x0][0x44]; /* 0x2800400110005de4 */
/*0010*/ S2R R0, SR_TID.X; /* 0x2c00000084001c04 */
/*0018*/ MOV32I R3, 0x0; /* 0x180000000000dde2 */
/*0020*/ SSY 0x68; /* 0x6000000100001c07 */
/*0028*/ IMAD R2.CC, R0, 0x4, R3; /* 0x2007c00010009ca3 */
/*0030*/ MOV32I R3, 0x0; /* 0x180000000000dde2 */
/*0038*/ IMAD.U32.U32.HI.X R3, R0, 0x4, R3; /* 0x2086c0001000dc43 */
/* 0x22f043f2f2e2c3f7 */
/*0048*/ LD.E.CV R0, [R2+0x4]; /* 0x8400000010201f85 */
/*0050*/ FSETP.NEU.AND P0, PT, R0, RZ, PT; /* 0x268e0000fc01dc00 */
/*0058*/ @!P0 BRA 0x40; /* 0x4003ffff800021e7 */
/*0060*/ NOP.S; /* 0x4000000000001df4 */
/*0068*/ LD.E.CV R4, [R2+0x4]; /* 0x8400000010211f85 */
/*0070*/ RED.E.ADD.F32.FTZ.RN [R2], R4; /* 0x2c00000000211e05 */
/*0078*/ EXIT; /* 0x8000000000001de7 */
/*0080*/ BRA 0x80; /* 0x4003ffffe0001de7 */
/*0088*/ NOP; /* 0x4000000000001de4 */
/*0090*/ NOP; /* 0x4000000000001de4 */
/*0098*/ NOP; /* 0x4000000000001de4 */
/*00a0*/ NOP; /* 0x4000000000001de4 */
/*00a8*/ NOP; /* 0x4000000000001de4 */
/*00b0*/ NOP; /* 0x4000000000001de4 */
/*00b8*/ NOP; /* 0x4000000000001de4 */
.........................................
Fatbin ptx code:
================
arch = sm_30
code version = [6,4]
producer = <unknown>
host = linux
compile_size = 64bit
compressed
$
根据我在 cc3.5 和 cc7.0 设备上的测试，上述代码仍然挂起，因此我们没有通过这些更改来修改其观察到的行为。（注意上面的 SASS 代码是针对 cc3.0 的，用 CUDA 10.1.243 编译）。
代码将表现出扭曲发散行为，而 IMO 这对于理解挂起至关重要，因此我们将重点关注 SASS 代码的条件区域：
/*0038*/ IMAD.U32.U32.HI.X R3, R0, 0x4, R3; /* 0x2086c0001000dc43 */
/* 0x22f043f2f2e2c3f7 */
/*0048*/ LD.E.CV R0, [R2+0x4]; /* 0x8400000010201f85 */
/*0050*/ FSETP.NEU.AND P0, PT, R0, RZ, PT; /* 0x268e0000fc01dc00 */
/*0058*/ @!P0 BRA 0x40; /* 0x4003ffff800021e7 */
/*0060*/ NOP.S; /* 0x4000000000001df4 */
/*0068*/ LD.E.CV R4, [R2+0x4]; /* 0x8400000010211f85 */
/*0070*/ RED.E.ADD.F32.FTZ.RN [R2], R4; /* 0x2c00000000211e05 */
/*0078*/ EXIT; /* 0x8000000000001de7 */
到 0038 行，所有的设置工作已经完成。在 0048 行，线程正在__device__ data从全局内存中加载它的值（指令.CV上的是我们修饰的结果），在 0050 行执行条件测试，在 0058 行执行条件分支。非零值，那么它将继续执行第 0060 行（并最终执行原子操作并退出）。如果没有，它会回到 0040 行重复加载和测试。LDvolatile
现在，我们观察到的是挂起。通过条件测试的线程和未通过条件测试的线程不会由 warp 调度程序同时调度。它必须安排一组（例如通过）或另一组（例如失败）。warp 调度器必须反复做出同样的决定。如果我们观察到挂起，唯一可能的结论是，条件测试失败的线程被重复调度（选择发出），而通过条件测试的线程没有被调度。
这是合法的，根据 CUDA 编程模型和此代码设计，任何“最终”应该安排通过的线程的结论都是无效的结论。保证传递的线程被调度的唯一方法是向 warp 调度程序提供一个没有其他选择的情况，这与此答案顶部的原则 1 保持一致。
（旁白：请注意，我们可能还观察到，warp 调度程序选择通过的线程而不是失败的线程来调度/发出。在这种情况下，因为那些通过的线程最终在这种实现中退出/退出，我希望这会导致在不挂起的代码中。通过的线程最终将全部退休，并且该答案顶部的第 1 项将强制扭曲调度程序开始调度失败的线程。不挂在这里将是同样有效的和可能的观察，在这里概述了扭曲调度特征。但基于该结果的任何正确性结论仍然是错误的。）
那么，扩展这个想法，有人可能会问“有没有一种合法的方式来实现这种模式？” 我建议我们现在知道，如果我们要完成这项工作，我们可能需要执行障碍。让我们选择__syncwarp()。对于那个屏障，屏障的合法使用通常要求我们有一个完整的经线（或多个经线）。所以我们需要重新编写代码以允许一个完整的经线处于活动状态，但只有所需的线程（总共 9 个）在做“工作”。
以下是实现这一目标的一种可能方法。我敢肯定还有其他方法。根据我的测试，此代码不会挂在 cc3.5 或 cc7.0 设备上：
__device__ volatile float data[10] = {0,0,0,0,0,0,0,0,0,1};
__global__ void gradually_set_global_data(int sz) {
int tflag = (threadIdx.x < sz) ? 1:0; // choose the needed threads to do the "work"
unsigned wflag = 1; // initially, the entire warp is marked active
while (wflag) { // run the entire warp, or exit the entire warp
if (tflag) // if this thread still needs to do its "work"
if (data[threadIdx.x + 1]) {
atomicAdd((float *)&data[threadIdx.x], data[threadIdx.x + 1]);
tflag = 0; // the work for this thread is completed
}
__syncwarp();
wflag = __ballot_sync(0xFFFFFFFFU, tflag); //deactivate warp when all threads done
}
}
int main() {
gradually_set_global_data<<<1, 32>>>(9);
cudaDeviceReset();
return 0;
}
请注意，如果我们想要更接近提供的代码，可以使用while(1)循环重铸上述代码，并在循环内部发出breakifwflag为零（在投票操作之后）。我认为这种认识没有任何有意义的差异。
我仍然没有声明此代码或我发布的任何其他代码的正确性。使用我发布的任何代码的任何人都需要自担风险。我只是声称我试图解决我在原始帖子中发现的缺陷，并提供一些解释。我并不是说我的代码没有缺陷，或者它适用于任何特定目的。使用（或不使用）风险自负。