两个逻辑等价的空CUDA内核为何存在显著性能差异？

这个差异的核心原因是nvcc对包含设备端printf的代码优化策略不同——哪怕分支永远不会执行，printf绑定的隐性运行时开销也没被完全消除，具体拆解如下：

1. 设备端printf的特殊隐性开销 #

CUDA的printf不是普通的主机端打印函数：它需要设备端维护输出缓冲区，还涉及主机-设备间的同步、数据传递逻辑。哪怕你的printf放在永远不会进入的分支里，在CUDA 7.5的-O3优化下，编译器也没能彻底移除和printf绑定的运行时初始化/注册代码。这些冗余代码会被内核携带，导致内核启动和执行时产生额外开销。

而__empty2里的空语句就简单多了：编译器一眼就能识别这是完全无意义的死代码，直接把整个分支连同相关逻辑彻底删掉，最终内核几乎是空的，执行时间就是CUDA内核启动的最小调度开销。

2. CUDA 7.5的死代码消除局限性 #

较老版本的nvcc（比如你用的7.5）在死代码分析能力上不如新版本。对于包含printf的死分支，编译器无法彻底推导出来printf永远不会被调用，因此保留了一些必要的运行时支持代码。而空语句的死分支属于最容易被优化的场景，自然被完全清除。

3. 内核执行的开销细节 #

• 对于__empty1：虽然分支不执行，但内核中仍然保留了与printf相关的设备端运行时代码，这些代码会占用SM的少量资源，并且在启动时需要初始化输出相关的机制——哪怕最终没有任何输出，这些步骤也会消耗时间。
• 对于__empty2：被优化后内核本质上是一个空函数，执行时几乎没有任何实际操作，只需要完成内核启动的基本调度流程，所以耗时只有0.02ms左右。

补充验证建议 #

你可以做两个小测试来验证这个结论：

• 把__empty1里的printf换成一个普通的空设备函数（比如__device__ void dummy() {}），再编译运行，看看两个内核的耗时是否会接近。
• 升级到较新的CUDA版本（比如11.x）重新编译，新版本的死代码消除能力更强，大概率会把__empty1里的死分支连同printf相关代码一起优化掉，两个内核的耗时会趋于一致。

你的测试代码 #

#include <cstdio> 
#include <time.h> 
#include <sys/time.h> 
#include <cuda.h> 

inline double wtime(){ 
    double time[2]; 
    struct timeval time1; 
    gettimeofday(&time1, NULL); 
    time[0]=time1.tv_sec; 
    time[1]=time1.tv_usec; 
    return time[0]+time[1]*1.0e-6; 
} 

__global__ void __empty1(bool flag){if(flag){ printf("hh\n");}} 
__global__ void __empty2(bool flag){if(flag){ ; }} 

int main(){ 
    cudaDeviceSynchronize(); 
    double s = wtime(); 
    __empty1<<<256,256>>>(false); 
    cudaDeviceSynchronize(); 
    printf("empty1: %.3f\n", 1000*(wtime()-s)); 

    cudaDeviceSynchronize(); 
    s = wtime(); 
    __empty2<<<256,256>>>(false); 
    cudaDeviceSynchronize(); 
    printf("empty2: %.3f\n", 1000*(wtime()-s)); 
    return 0; 
}

编译命令 #

nvcc -O3 -arch=sm_35 ./main.cu

运行结果 #

empty1: 1.075
empty2: 0.019

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