(int) -1 的 CUDA 内核 printf 使用 %d 说明符给出错误的输出

Question

当用 C 告诉我的电脑时

printf("%d",(int)-1);

我确实期望并且通常也会得到“-1”响应。然而，在我的 Tesla M2090 Nvidia Card 上，我的基于 Ubuntu 的 Cuda 5.0 与这个无辜的演示程序对话

/** cheese.cu */
#include <iostream>
#include <cuda.h>
#include <cstdio>

using namespace std;

template<typename T> struct SpacePtr
  { size_t size; T* ptr; };

 __global__ void f(SpacePtr<int>* sp)
{
  printf("This _is_ a 'minus one' or ain't it: %d\n",(int)-1);
  // Note: All appears to work fine with %lu instead of %zd
  // Still: How does the size value affect the -1?
  printf("On DEV: Size: %zd, Minus One: %d\n",sp->size,(int)-1);
}

int main()
{
  SpacePtr<int> data; data.ptr = 0; data.size = 168;
  SpacePtr<int>* devPtr = 0;
  cudaMalloc(&devPtr,1);
  cudaMemcpy(devPtr,&data,sizeof(SpacePtr<int>),cudaMemcpyHostToDevice);
  f<<<1,1,0,0>>>(devPtr);
  cudaError_t err = cudaGetLastError();
  cout << "The last error code was " << err << " (" <<
    cudaGetErrorString(err) << ")" << endl;
  cudaDeviceSynchronize();
}

编译和调用方式

nvcc -arch=sm_20 cheese.cu && ./a.out

产生输出：

The last error code was 0 (no error)
This _is_ a 'minus one' or ain't it: -1
On DEV: Size: 168, Minus One: 10005640

最后一个数字实际上是某种随机数（两个后续调用返回不同的结果），就好像内存分配有问题一样。 -1 之前的 (int) 已经是反复试验。原程序没有。

问题是：有人知道为什么没有写出 -1 吗？如果是，请告诉我为什么？非常感谢，马库斯。

Answer 1

如果你查看CUDA C Programming Guide的appendix B.32.1 "Format Specifiers"，你会发现%zd中的z修饰符不被支持。您必须转换为 unsigned long 并使用 %lu 作为格式说明符：

printf("On DEV: Size: %lu, Minus One: %d\n",(unsigned long)(sp->size), (int)-1);

Answer 2

基本问题似乎是设备 printf 不支持 %zd 格式说明符（我不确定它是否受支持）。

编辑：

CUDA 5 中 printf 的文档是这样说的：

对于标准 printf()，格式说明符采用以下形式： %[flags][width][.precision][size]type

支持以下字段（请参阅广泛可用的文档 有关所有行为的完整描述）：

Flags: ‘#’ ‘ ‘ ‘0’ ‘+’ ‘-‘
Width: ‘*’ ‘0-9’
Precision: ‘0-9’
Size: ‘h’ ‘l’ ‘ll’
Type: ‘%cdiouxXpeEfgGaAs’

请注意，CUDA 的 printf() 将接受标志、宽度、 精度、尺寸和类型，无论它们总体上是否构成有效 格式说明符。换句话说，“%hd”将被接受并 printf 将期望在相应位置有一个双精度变量 在参数列表中。

因此，不支持 size_t 的 %zd 格式说明符。像这样修改你的内核：

__global__ void f(SpacePtr<int>* sp)
{
    const int minus_one = -1;
    printf("This _is_ a 'minus one' or ain't it: %d\n", minus_one);
    printf("On DEV: Size: %d, Minus One: %d\n",int(sp->size), minus_one);
}

工作正常。

另请注意，您的主机代码中有一个相当大的错误，尽管这不会对您示例中的清单行为产生任何影响。您只是为dev_ptr 分配和复制一个字节，这显然是不正确的。它应该类似于：

cudaMalloc((void **)&devPtr, sizeof(data));
cudaMemcpy(devPtr, &data, sizeof(data), cudaMemcpyHostToDevice);

将数据的全部内容从主机传输到设备。