CUDA中printf打印size_t的正确方式及异常原因探究

问题重现 #

用户编写的CUDA验证代码如下：

#include <stdio.h>

__global__ void testKernel(int *data, int n){
    size_t tid = threadIdx.x + blockIdx.x * blockDim.x;
    if(tid >= n) return;
    int v = data[tid];
    if(tid < 8) printf("%d: %d\n", tid, v);
}

int main(){
    int n = 128;
    auto h_data = new int[n];
    for(int i = 0; i < n; i++) h_data[i] = i;
    
    int *d_data;
    cudaMalloc((void**)&d_data, sizeof(int) * n);
    cudaMemcpy(d_data, h_data, sizeof(int) * n, cudaMemcpyHostToDevice);
    
    testKernel<<<4, 32>>>(d_data, n);
    cudaDeviceSynchronize();
    
    cudaFree(d_data);
    delete[] h_data;
    return 0;
}

运行后输出的v值全部为0，将printf格式改为printf("%lu: %d\n", tid, v)后输出恢复正常，但CPU端用%d打印size_t未出现异常，需解释该现象并给出CUDA中打印size_t的正确方法。

原因分析 #

• 设备端参数传递错位：在CUDA支持的64位GPU架构中，size_t是64位无符号整数类型，而%d是对应32位有符号整数的格式符。当调用printf("%d: %d\n", tid, v)时，printf会错误地解析参数：64位的tid会被拆分为两个32位单元读取，导致原本的v参数位置被挤占，最终读取到错误的数值（表现为0）。
• CPU端无异常的本质是未定义行为：CPU端如果是32位系统，size_t本身就是32位，和%d格式符匹配；即使是64位系统，当tid的数值小于2^31时，64位值的低32位和整体数值一致，且部分编译器的参数传递机制可能刚好掩盖了类型不匹配的问题，但这属于未定义行为，并非安全的写法。

CUDA中打印size_t的正确方法 #

1. 使用标准格式符%zu：这是C语言标准中专门为size_t类型定义的格式符，能自动适配size_t的位数（32位或64位），在CUDA设备端和CPU端都能正确工作，是最推荐的方式：

   printf("%zu: %d\n", tid, v);
   

2. 针对64位架构使用%lu：如果明确目标是64位GPU，也可以用%lu（无符号长整数格式符），因为64位系统中size_t和unsigned long长度一致，但%zu的通用性更强。

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