为什么无warp divergence的CUDA kernel性能比有发散的更低？

问题背景 #

编写CUDA kernel研究warp divergence行为的代码如下：

#include <cuda_runtime.h>
#include <stdio.h>
#include "util.h"
#include <chrono>

__global__ void wardUp(float *c)
{
    float a = 0.0;
    float b = 0.0;
    int idx = threadIdx.x + blockIdx.x*blockDim.x;
    if ((idx/warpSize)%2 == 0){
        a = 100.0f;
    }
    else{
        b = 200.0f;
    }
    c[idx] = a+b;
}

__global__ void kernel1(float *c)
{
    float a = 0.0;
    float b = 0.0;
    int idx = threadIdx.x + blockIdx.x*blockDim.x;
    if ((idx/warpSize)%2 == 0){
        a = 100.0f;
    }
    else{
        b = 200.0f;
    }
    c[idx] = a+b;
}

__global__ void kernel2(float *c)
{
    float a = 0.0;
    float b = 0.0;
    int idx = threadIdx.x + blockIdx.x*blockDim.x;
    if (idx%2 == 0){
        a = 100.0f;
    }
    else{
        b = 200.0f;
    }
    c[idx] = a+b;
}

int main(int argc, char **argv)
{
    initDevice(0);
    int size = 64;
    int blocksize = 64;
    int nBytes = sizeof(float)*size;
    float *a_d;
    CHECK(cudaMalloc((float**)&a_d, nBytes));
    dim3 block(blocksize, 1);
    dim3 grid((blocksize-1)/block.x+1, 1);

    wardUp<<<grid, block>>>(a_d);

    float elapsed = 0;
    cudaEvent_t start1, stop1;
    CHECK(cudaEventCreate(&start1));
    CHECK(cudaEventCreate(&stop1));
    CHECK(cudaEventRecord(start1, 0));
    kernel1<<<grid, block>>>(a_d);
    CHECK(cudaEventRecord(stop1, 0));
    CHECK(cudaEventSynchronize(stop1));
    CHECK(cudaEventElapsedTime(&elapsed, start1, stop1));
    printf("kernel1 take:%2f ms\n", elapsed);

    float elapsed_1 = 0;
    cudaEvent_t start2, stop2;
    CHECK(cudaEventCreate(&start2));
    CHECK(cudaEventCreate(&stop2));
    CHECK(cudaEventRecord(start2, 0));
    kernel2<<<grid, block>>>(a_d);
    CHECK(cudaEventRecord(stop2, 0));
    CHECK(cudaEventSynchronize(stop2));
    CHECK(cudaEventElapsedTime(&elapsed_1, start2, stop2));
    printf("kernel2 take:%2f ms\n", elapsed_1);
    
    cudaFree(a_d);
    cudaEventDestroy(start1);
    cudaEventDestroy(stop1);
    cudaEventDestroy(start2);
    cudaEventDestroy(stop2);
    return 0;
}

观测现象 #

按照常规认知，kernel1不存在warp发散问题：if分支按warp粒度划分，0-31号线程同属一个warp走同一分支；kernel2存在warp发散问题，同一warp内奇偶线程走不同分支，理论上kernel1性能应该远优于kernel2。但实测结果相反：

使用设备：0: NVIDIA GeForce RTX 2080 Ti
kernel1耗时：0.008864 ms
kernel2耗时：0.006752 ms

即使用cudaEventRecord做高精度耗时统计，kernel1运行速度依然慢于kernel2。

原因分析 #

• kernel计算量过小，warp发散的性能差异被噪声覆盖
  测试仅启动1个block共64个线程，kernel本身执行时间仅几微秒，这个量级的耗时会被kernel启动开销、GPU调度波动完全覆盖，warp发散带来的性能损失根本无法体现。
• 分支条件的计算开销差异成为主导
  kernel1的分支判断(idx/warpSize)%2需要先做整数除法再做取模运算，kernel2的分支判断idx%2可以直接通过读取寄存器最低位完成，计算开销远低于kernel1的判断逻辑。在整体计算量极小的前提下，分支条件本身的开销差就决定了最终耗时。
• 编译器优化消除了warp发散的影响
  对于逻辑极其简单的kernel，nvcc会做激进优化：
  1. kernel2的奇偶分支可以直接优化为谓词执行，不需要分支跳转，两个分支的指令可以并行发射，所谓的warp发散两倍耗时损失在该场景下完全不存在。
  2. 两个kernel的最终输出a+b恒等于300，编译器甚至可以直接消除分支逻辑，直接给c[idx]赋值300，此时分支逻辑差异完全不影响执行耗时，耗时差完全来自判断条件本身的计算开销。

正确复现warp发散性能差异的方法 #

如果要观测到warp发散的性能损失，需要做如下调整：

• 扩大总线程规模到百万级别，让kernel实际执行时间远高于启动开销，至少达到毫秒级。
• 增大分支内的计算量，比如在分支中加入数十次浮点运算，避免编译器直接优化掉分支逻辑。
• 调整分支逻辑让输出结果无法被编译器提前推导，避免分支被直接消除。

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