CUDA warp蝴蝶洗牌求和Release模式下步长偏移错误问题

架构老林

2026-8-26

CUDA Release模式下warp洗牌求和结果异常问题

编辑说明： 我已将该问题作为bug提交至NVIDIA官方issue追踪系统，对应工单ID为3711214。

我在编写数值仿真程序时遇到问题：程序在Release模式下会输出存在细微偏差的错误结果，但Debug模式下运行结果符合预期。原程序使用curand实现随机采样，我已将问题简化为确定性更高的MVCE（最小可复现示例）：仅启动1个block、1个warp（共32个thread）的kernel，每个thread执行逻辑如下：

执行带循环的计算任务，该循环极易产生warp-divergent（warp分支发散），尤其在循环末期，部分thread会先于其他thread完成任务退出循环。
调用同步接口对齐所有thread的执行状态。
尝试通过warp内butterfly-shuffle（蝴蝶洗牌）操作与同warp的thread交换数据，最终计算得到全warp的求和结果。
[MVCE中省略该步骤] 0号thread将求和结果写回global memory，以便后续拷贝至host端。

最小可复现代码如下：

#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模式

洗牌操作运行完全符合预期。初始阶段每个thread持有独立的cIters值：

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
// ...

完成第一次offset为1的__shfl_xor_sync操作后，每两个相邻thread的cIters值保持一致：

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

完成offset为2的__shfl_xor_sync操作后，每4个thread为一组的cIters值保持一致：

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内所有thread的cIters值统一为4673。

Release模式

切换至Release模式编译运行后，结果出现隐蔽错误：进入洗牌逻辑的初始数值与Debug模式完全一致，第一轮offset为1的洗牌输出也与Debug版本匹配（相邻thread对数值一致），但执行offset为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

经核对，该错误输出与Debug模式下将第二轮洗牌的offset从2错误改为1时的输出完全一致，对应错误代码逻辑如下：

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错误改为1
    cIters += __shfl_xor_sync(0xffffffff, cIters, 1, 32); // 偏移量从2错误改为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);

两种模式下的完整输出差异稳定可复现。

测试环境

硬件：GA103核心RTX 3080Ti（移动版），运行在厂商推荐时钟频率，配备16GB VRAM。该设备运行其他CUDA程序未出现数据损坏问题（已通过primegrid-CUDA测试，任务结果经双重校验确认正确）
CUDA版本：11.0
主机编译器：MSVC 14.29.30133
Debug模式完整编译命令：

"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"

Release模式完整编译命令：

"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()调用（包括现有同步位置、各次洗牌操作之间），尽管洗牌操作本身自带warp同步语义，理论上无需额外同步
将compute capability修改为8.0以匹配当前显卡架构
强制GPU运行在基础时钟频率排除硬件稳定性问题
反转洗牌执行顺序（按16/8/4/2/1的offset顺序执行）
替换为__shfl_down_sync接口，使用相同步长模式实现warp求和。

经测试，若让每个thread将私有值写入global memory后在host端CPU求和，可得到正确结果。

临时规避方案测试

将所有洗牌操作替换为__shfl_sync调用、手动计算目标lane ID可解决问题；但仅替换出现错误的offset=2的__shfl_xor_sync为__shfl_sync无法修复问题；仅替换原本运行正常的第一次offset=1的__shfl_xor_sync为__shfl_sync反而可以修复问题。上述规避方案仅在MVCE中验证，尚未在完整程序中测试有效性，对应代码如下：

// 可意外实现正常运行的修改版本
    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）导致修复方案鲁棒性不足，待解答的核心疑问为：

现有代码中是否存在未定义行为？
该问题是否为已知的CUDA编译器bug？

内容的提问来源于stack exchange，提问作者nanofarad

CUDA warp蝴蝶洗牌求和Release模式下步长偏移错误问题

不同编译模式下的运行表现 #

Debug模式 #

Release模式 #

测试环境 #

已尝试的排查方案 #