为何同步Warp投票(__ballot_sync)返回不一致结果？

在CUDA开发中，发现本应收敛的__ballot_sync函数返回了不一致结果，可通过以下代码复现：

__global__ void testKernel(uint32_t* testBuffer)
{
    if (!threadIdx.x)
        atomicAdd(&testBuffer[0], 1);

    printf("%u: %u \n", threadIdx.x, __ballot_sync(__activemask(), 1));
}

int main()
{
    uint32_t* tBuffer;
    cudaMalloc(&tBuffer, 1024 * sizeof(uint32_t));
    
    testKernel<<<1,32>>>(tBuffer);

    cudaFree(tBuffer);
    return 0;
}

执行后输出如下：

1     : 4294967294
 . . .
 31    : 4294967294
 0     : 1

我对SIMT编程模型的理解是：

• warp内的线程以逻辑同步而非物理锁步的方式执行；
• 当同一warp内的线程发生发散时，其中一个分支的线程会执行，另一分支的线程会被掩码。

基于以上两点，我预期代码中线程1-31会等待线程0完成atomicAdd操作后重新收敛，但实际线程1-31和线程0各自参与了独立的投票操作。

为何线程在原子操作后未能重新收敛？

相关注意事项 #

• 仅原子操作会引发该发散，执行非原子加法甚至循环10000次都不会导致发散；
• 在原子操作后添加__syncthreads可强制线程重新收敛；
• 禁用调试信息生成(-G)时不会出现该发散。

编译调试信息 #

-gencode=arch=compute_75,code="sm_75,compute_75"
--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\v12.0\include" -I"C:\Users\x\Desktop\VCPCKG\vcpkg\installed\x64-windows\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.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 /FS /Zi /RTC1 /MDd " -Xcompiler "/Fdx64\Debug\vc142.pdb" -o C:\Users\x\source\repos\WarpDivergence\x64\Debug\kernel.cu.obj "C:\Users\x\source\repos\WarpDivergence\kernel.cu"

运行硬件 #

RTX 2080 Super

原因分析 #

核心原因是调试模式(-G)下CUDA编译器会禁用warp级收敛优化。

调试模式的核心目标是保留代码的原始执行路径，方便单步调试，因此编译器不会对线程执行流程做合并优化。atomicAdd属于内存操作，在调试模式下，编译器会为执行该操作的线程0生成独立的执行分支，线程1-31则直接跳过分支执行后续的__ballot_sync，导致warp内线程没有收敛，各自生成投票结果。

在非调试模式下，编译器会自动进行warp级优化，确保分支执行完成后所有线程重新同步，因此不会出现这个问题。而__syncthreads作为块级同步原语，会强制整个线程块的所有线程到达同步点后再继续执行，所以能手动修复收敛问题。普通的非原子操作或循环不存在内存访问的强依赖，编译器在调试模式下也不会拆分执行路径，因此不会引发发散。

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