cuda - 为什么我的 CUDA warp shuffle sum 在一个洗牌步骤中使用了错误的偏移量？

Question

编辑: 我已将此作为错误提交到 https://developer.nvidia.com/nvidia_bug/3711214 .

我正在编写一个数值模拟程序，该程序在 Release模式下给出了微妙的错误结果，但在 Debug模式下却给出了看似正确的结果。原始程序使用 curand 进行随机采样，但我已将其简化为更简单且更具确定性的 MVCE，它启动 1 个 block * 1 个扭曲(32 个线程)的单个内核，其中每个线程:

• 使用可能会变得扭曲发散的循环执行计算，尤其是在接近尾声时，因为一些线程先于其他线程完成任务。
• 将线程重新同步。
• 尝试与 warp 中的其他线程对数据进行蝶式洗牌以获得单个总和。
• [在 MVCE 中不需要]线程 0 会将总和写回全局内存，以便可以将其复制到主机

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>


__global__ void test_kernel()
{

    int cSteps = 0;
    int cIters = 0;
    float pos = 0;

    //curandState localState = state[threadIdx.x];

    while (true) {
        float rn = threadIdx.x * 0.01 + 0.001;
        pos += rn;
        cSteps++;
        if (pos > 1.0f) {
            pos = 0;
            cIters++;
            if (cSteps > 1024) {
                break;
            }
        }
    }

    printf(" 0: Th %d cI %d\n", threadIdx.x, cIters);
    __syncthreads();
    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 1, 32);
    cIters += __shfl_xor_sync(0xffffffff, cIters, 1, 32);

    printf(" 1: Th %2d cI %d\n", threadIdx.x, cIters);
    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 2, 32);
    cIters += __shfl_xor_sync(0xffffffff, cIters, 2, 32);

    printf(" 2: Th %2d cI %d\n", threadIdx.x, cIters);

    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 4, 32);
    cIters += __shfl_xor_sync(0xffffffff, cIters, 4, 32);

    printf(" 4: Th %2d cI %d\n", threadIdx.x, cIters);

    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 8, 32);
    cIters += __shfl_xor_sync(0xffffffff, cIters, 8, 32);

    printf(" 8: Th %2d cI %d\n", threadIdx.x, cIters);

    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 16, 32);
    cIters += __shfl_xor_sync(0xffffffff, cIters, 16, 32);

    printf("16: Th %2d cI %d\n", threadIdx.x, cIters);
}

int main()
{
    test_kernel <<<1, 32>>> ();
    return 0;
}

在 Debug模式下，随机播放按预期工作。我看到每个线程都以自己的值开始:

 0: Th 0 cI 2
 0: Th 1 cI 12
 0: Th 2 cI 22
 0: Th 3 cI 32
 0: Th 4 cI 41
// ...

在第一次洗牌 xor 1 后，每对线程都同意相同的数字:

 1: Th  0 cI 14
 1: Th  1 cI 14
 1: Th  2 cI 54
 1: Th  3 cI 54

在 shuffle xor 2 之后，每组四个线程同意:

 2: Th  0 cI 68
 2: Th  1 cI 68
 2: Th  2 cI 68
 2: Th  3 cI 68
 2: Th  4 cI 223
 2: Th  5 cI 223
 2: Th  6 cI 223
 2: Th  7 cI 223

等等。在最后一次洗牌之后，warp 中的所有线程都同意相同的值 (4673)。

一旦我启用 Release模式，我得到的结果是微妙的垃圾。进入随机播放的值是相同的，第一轮随机播放后的值与调试版本一致(并且与之前一样在每一对内一致)。一旦我执行 shuffle xor 2，结果就会崩溃:

 2: Th  0 cI 28
 2: Th  1 cI 28
 2: Th  2 cI 108
 2: Th  3 cI 108
 2: Th  4 cI 186
 2: Th  5 cI 186
 2: Th  6 cI 260
 2: Th  7 cI 260

事实上，如果洗牌序列被这个特定损坏的序列替换，这就是调试构建(和手工检查)将产生的确切输出:

    printf(" 0: Th %d cI %d\n", threadIdx.x, cIters);
    __syncthreads();
    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 1, 32);
    cIters += __shfl_xor_sync(0xffffffff, cIters, 1, 32);

    printf(" 1: Th %2d cI %d\n", threadIdx.x, cIters);
    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 1, 32); // 2 changed to 1
    cIters += __shfl_xor_sync(0xffffffff, cIters, 1, 32); // 2 changed to 1

    printf(" 2: Th %2d cI %d\n", threadIdx.x, cIters);

    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 4, 32);
    cIters += __shfl_xor_sync(0xffffffff, cIters, 4, 32);

    printf(" 4: Th %2d cI %d\n", threadIdx.x, cIters);

    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 8, 32);
    cIters += __shfl_xor_sync(0xffffffff, cIters, 8, 32);

    printf(" 8: Th %2d cI %d\n", threadIdx.x, cIters);

    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 16, 32);
    cIters += __shfl_xor_sync(0xffffffff, cIters, 16, 32);

输出的完整差异是 here .

软硬件环境如下:

• GA103 3080Ti(移动)，采用制造商推荐的时钟，16 G VRAM。机器似乎没有与其他 Cuda 程序损坏(使用 primegrid-CUDA 测试并根据双重检查验证任务)

• CUDA 11.0

• MVSC 主机编译器 14.29.30133

• 完整调试命令行如下:

  "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\bin\nvcc.exe" -gencode=arch=compute_52,code=\"sm_52,compute_52\" --use-local-env -ccbin "C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\bin\HostX86\x64" -x cu   -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\include"  -G   --keep-dir x64\Debug -maxrregcount=0  --machine 64 --compile -cudart static  -g   -DWIN32 -DWIN64 -D_DEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /Od /Fdx64\Debug\vc142.pdb /FS /Zi /RTC1 /MDd " -o x64\Debug\kernel.cu.obj "C:\Users\[username]\source\repos\BugRepro\BugRepro\kernel.cu"
  

• 完整发布命令行如下:

  C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\bin\nvcc.exe" -gencode=arch=compute_52,code=\"sm_52,compute_52\" --use-local-env -ccbin "C:\Program Files (x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.29.30133\bin\HostX86\x64" -x cu   -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0\include"     --keep-dir x64\Release -maxrregcount=0  --machine 64 --compile -cudart static     -DWIN32 -DWIN64 -DNDEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /O2 /Fdx64\Release\vc142.pdb /FS /Zi  /MD " -o x64\Release\kernel.cu.obj "C:\Users\[username]\source\repos\BugRepro\BugRepro\kernel.cu"
  

我尝试过但没有解决的事情:

• 添加/删除 syncthreads 调用(其中显示一个，并且在 shuffle 调用之间)，即使它们不是必需的，因为每个 shuffle 都是同步的
• 将计算能力更改为 8.0 以更好地匹配我的卡
• 在 GPU 上强制使用基本时钟
• 以相反的顺序洗牌 (16/8/4/2/1)
• 使用 __shfl_down_sync 而不是 xor，具有相同的偏移模式。

让每个线程写入全局内存，然后在主机 CPU 上求和确实会产生正确的结果。

用调用 __shfl_sync 和手动计算的车道 ID 替换所有随机播放。 只是用__shfl_sync 替换损坏的shuffle xor 2 不会。用 __shfl_sync 替换 just 第一个 shuffle xor 1(它工作正常)确实似乎修复了它。 (这两个解决方法适用于我的 MVCE；我还没有机会评估它们是否适用于整个程序)

    // unexpectedly working
    int id = threadIdx.x;
    printf(" 0: Th %d cI %d\n", threadIdx.x, cIters);
    __syncthreads();
    cSteps += __shfl_sync(0xffffffff, cSteps, id ^ 1, 32);
    cIters += __shfl_sync(0xffffffff, cIters, id ^ 1, 32);

    printf(" 1: Th %2d cI %d\n", threadIdx.x, cIters);
    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 2, 32);
    cIters += __shfl_xor_sync(0xffffffff, cIters, 2, 32);

    printf(" 2: Th %2d cI %d\n", threadIdx.x, cIters);

    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 4, 32);
    cIters += __shfl_xor_sync(0xffffffff, cIters, 4, 32);

    printf(" 4: Th %2d cI %d\n", threadIdx.x, cIters);

    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 8, 32);
    cIters += __shfl_xor_sync(0xffffffff, cIters, 8, 32);

    printf(" 8: Th %2d cI %d\n", threadIdx.x, cIters);

    cSteps += __shfl_xor_sync(0xffffffff, cSteps, 16, 32);
    cIters += __shfl_xor_sync(0xffffffff, cIters, 16, 32);

    printf("16: Th %2d cI %d\n", threadIdx.x, cIters);

即使我有解决方法，我担心我仍然会在某处遇到未定义的行为，而且我的修复可能很脆弱。

任何人都可以阐明这一点吗？我的程序中确实有UB吗？这是一个已知的编译器错误吗？

Answer 1

根据 CUDA 工程团队，这已被确认为编译器错误。正如他们的来信所证实的那样，修复即将推出:

The fix is targeting a future major CUDA release after CUDA 11. The JIT fix will possibly be a little earlier in a Driver branch after latest R515 online.

编辑:似乎没有在 516.94 Game Ready 驱动程序中修复。它似乎在 522.25 和 Cuda 11.8 中得到修复。

他们还确认关闭优化可以解决问题；他们不会对任何在优化仍在进行中可靠工作的解决方法发表评论。

以下解决方法适用于我的硬件和编译器，但 YMMV:

• 使用 __shfl_sync 而不是 shfl_add_sync 或 shfl_xor_sync
• __reduce_add_sync