关于cuda中修改flag数组的问题

我正在研究 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]，但它陷入了无限循环。为什么会出现这种情况？

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

1. 假设至少有 1 个线程，则前进进度将在至少 1 个（可发布、非退休）线程中交付。
2. 将遵守执行屏障语义

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

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

Longer:

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

为什么要做出这样的区分呢？因为对所提供的代码进行相对较小的更改（实际上并不能解决正确性问题）可能会导致编译器生成不挂起的代码。缺乏仔细的治疗可能会导致人们得出这样的结论：因为它没有悬挂，所以它一定没问题。关键是代码是否挂起与它是否正确是不同的。我已经向自己证明了这一点。但是我不想提供该代码。正确的做法是设计正确的代码。请参阅下面我的尝试。

在我们深入研究 SASS 之前，我想指出代码中的另一个缺陷。 CUDA 编译器可以自由地将任何全局数据“优化”到寄存器中，同时保持单线程语义正确性。编译器大多只考虑单个线程，因此这可能会给依赖线程间通信的程序员带来麻烦（正如此代码所示）。为了正确性，在此代码中，线程 x 修改的数据必须（最终）对线程 x-1 可见。 CUDA 编程模型不保证这种线程间可见性，编译器通常也不强制执行。为了正确性，有必要通知编译器使该数据可见，并命令加载和存储来实现这一点。有多种方法可以实现这一点。我会建议将数据标记为volatile https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#volatile-qualifier为了简单起见，尽管可以通过执行障碍来做到这一点（例如__syncthreads(), __syncwarp()）那也内置内存屏障 https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#synchronization-functions。无论选择哪种方法来强制执行线程间数据可见性，如果没有它，代码就会被破坏，而与任何其他考虑无关。

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

现在，我们观察到的是挂起。通过条件测试的线程和未通过条件测试的线程不会同时由 warp 调度程序调度。它必须安排一组（例如通过）或另一组（例如失败）。扭曲调度程序必须重复做出同样的决定。如果我们观察到挂起，唯一可能的结论是，未通过条件测试的线程被重复调度（选择发出），而通过条件测试的线程未得到调度。

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

（旁白：请注意，我们可能还观察到，warp 调度程序选择传递线程而不是失败线程来调度/发出。在这种情况下，因为这些传递线程最终在此实现中退出/退休，我预计这会导致在不挂起的代码中。传递的线程最终将全部退出，并且通过本答案顶部的第 1 项，warp 调度程序将被迫开始调度失败的线程。不挂在这里将是一个同样有效和可能的观察，就此处概述的扭曲调度特征而言。但基于该结果得出的任何正确性结论仍然是错误的。）

那么，延伸这个想法，人们可能会问“有没有一种合法的方式来实现这种模式？”我建议我们现在知道，如果我们要使这项工作成功，我们可能需要执行障碍。我们来选择一下__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 if wflag为零（投票操作后）。我认为这种认识没有任何有意义的差异。

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