CUDA kernel printf of (int) -1 gives wrong output with %d specifier

Question

When telling my computer in C

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

I do expect, and usually get, too, a '-1' response. However, on my Tesla M2090 Nvidia Card spoken to from my Ubuntu-based Cuda 5.0 this innocent demo program

/** 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();
}

compiled and called via

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

yields the output:

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

The last number actually being some kind of random number (two subseqent calls return different results) as if something was amiss in memory allocation. The (int) before the -1 is already trial and error. The original program did not have it.

And here is the question: Does anybody see, why the -1 is not written, and if so, could you please tell me why? Thanks a lot, Markus.

Answer 1

If you check appendix B.32.1 "Format Specifiers" of the CUDA C Programming Guide, you will find that the z modifier in %zd is not supported. You will have to cast to unsigned long and use %lu as format specifier:

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