CUDA如何实现printf()？

Question

我在我的 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;
}

上述代码的输出为：

0: 0
1: 0
2: 0
3: 0
4: 0
5: 0
6: 0
7: 0

显然，

v

的值是错误的。然后我将

printf()

修改为以下内容：

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

输出变得正确，如下：

0: 0
1: 1
2: 2
3: 3
4: 4
5: 5
6: 6
7: 7

我认为原因是我使用了错误的格式来打印

size_t

类型，但我不知道这是怎么发生的。

此外，我在CPU上尝试了相同的打印格式，如下：

for(int i = 0; i < 8; i++){
    size_t x = i;
    printf("%d: %d\n", x, h_data[i]);
}
/*
0: 0
1: 1
2: 2
3: 3
4: 4
5: 5
6: 6
7: 7
*/

虽然我使用了错误的打印格式，但结果没有显示任何错误。

Answer 1

在您的具体示例中，

tid

具有类型

size_t

，因此无法使用修饰符

%d

打印它。对于 printf 的主机版本，

size_t

的适当限定符是

%zu

，但 CUDA 不支持它。相反，您可以使用修饰符

%llu

。

以下示例：

__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("%llu: %d\n", tid, v);
}

我可以在本地确认它有效：

$ nvcc -o test test.cu
$ ./test
0: 0
1: 1
2: 2
3: 3
4: 4
5: 5
6: 6
7: 7