concurrency - cuda修改flag数组的问题

Question

我正在研究 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;
}

内核应该用 data 完成执行预期持有 [1,1,1,1,1,1,1,1,1,1]，但它陷入无限循环。为什么会发生这种情况？

Answer 1

TL;DR:代码被检查破坏了。 CUDA 线程模型 不保证任何特定线程的前进进度 除非符合以下条件:

假设至少有 1 个线程，将在至少 1 个(已发布的、未退休的)线程中交付前进进度。

将遵守执行屏障语义

CUDA 编程模型未定义将为第 1 项选择哪个或哪些线程。除非程序员使用执行障碍进行显式控制，否则 CUDA 线程模型可以随意调度单个线程，直到该线程退休或遇到显式执行障碍为止。

由于提供的代码没有执行障碍，CUDA 工作调度程序(就 CUDA 语义而言)可以自由调度，例如线程 0，没有其他线程。如果我们将这个概念应用到所提供的代码中，很明显线程 0 如果单独运行，将呈现出无限循环。

更长:

这恰好是观察到的行为，尽管如果是我，我不会将两者联系起来。挂起的原因(根据我尝试描述的方式)不是“为了正确性，此代码依赖于 CUDA 编程模型未提供的保证”，尽管我相信这是一个真实的陈述。要了解挂起的原因，我建议有必要在查看 SASS(机器汇编代码)的情况下检查低级机器行为。我真的没有能力把这个话题穷尽，所以我将对此提出一个有限的看法。

为什么要划出这个区别？因为对所提供代码的相对较小的更改实际上并未解决正确性问题，可能会导致编译器生成不会挂起的代码。不小心处理可能会导致人们得出结论，因为它没有挂，所以它一定没问题。关键是代码是否挂起与它是否正确是不同的。我已经向自己证明了这一点。但是，我不想提供该代码。正确的做法是设计正确的代码。请参阅下文，了解我在这方面的尝试。

在我们深入研究 SASS 之前，我想指出代码中的另一个缺陷。 CUDA 编译器可以自由地将任何全局数据“优化”到寄存器中，同时维护 单线程语义正确性。编译器在 View 中大多只有一个线程，因此这可能会绊倒依赖线程间通信的程序员(如此代码所示)。为了正确起见，在此代码中，线程 x 修改的数据必须(最终)对线程 x-1 可见。 CUDA 编程模型不保证这种线程间可见性，编译器通常也不强制执行它。为了正确性，有必要通知编译器使这些数据可见，并排序加载和存储以实现这一点。有多种方法可以实现这一点。我会建议 marking the data with volatile 为简单起见，虽然可以使用执行障碍(例如 __syncthreads() 、 __syncwarp() )来做到这一点，但也可以使用 have memory barriers built-in .无论选择哪种方法来强制执行线程间数据可见性，如果没有它，代码就会被破坏，与任何其他考虑因素无关。

因此，在深入研究 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 指令上的 LD 是我们 volatile 修饰的结果)，条件测试在第 0050 行和条件分支在 0058 行执行。非零值，则它将继续执行第 0060 行(并最终执行原子操作并退出)。如果没有，它将返回到第 0040 行重复加载和测试。

现在，我们观察到的是挂起。通过和未通过条件测试的线程不会被 warp 调度器同时调度。它必须安排一组(例如通过)或另一组(例如失败)。经纱调度器必须反复做出同样的决定。如果我们观察到挂起，唯一可能的结论是条件测试失败的线程被重复调度(选择发布)，而通过条件测试的线程没有被调度。

这是合法的，根据 CUDA 编程模型和这个代码设计，任何关于传递线程应该“最终”被调度的结论都是无效的结论。保证传递的线程得到调度的唯一方法是将 warp 调度程序交给一个情况，以便它没有其他选择可用，这与本答案顶部的原则 1 保持一致。

(旁白:请注意，我们可能还观察到经线调度程序选择了传递线程而不是失败线程来调度/发出。在这种情况下，因为这些传递线程最终会在此实现中退出/退休，我希望这会导致在不会挂起的代码中。通过的线程最终将全部退休，并且经线调度程序将被此答案顶部的第 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) 重铸上面的代码。循环，并在循环内发出 break如果 wflag为零(投票操作后)。我认为这种认识没有任何有意义的差异。

我仍然没有声明此代码或我发布的任何其他代码的正确性。任何使用我发布的任何代码的人都需要自担风险。我只是声称我试图解决我在原始帖子中发现的缺陷，并提供一些解释。我并不是说我的代码没有缺陷，或者它适用于任何特定目的。使用(或不使用)风险由您自行承担。