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);

Answer 2

The basic problem seems to be that the %zd format specifier isn't being honoured by the device printf (I am not sure, off the top of my head whether it is supported).

EDIT:

The documentation for printf in CUDA 5 says this:

As for standard printf(), format specifiers take the form: %[flags][width][.precision][size]type

The following fields are supported (see widely-available documentation for a complete description of all behaviors):

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

Note that CUDA’s printf()will accept any combination of flag, width, precision, size and type, whether or not overall they form a valid format specifier. In other words, “%hd” will be accepted and printf will expect a double-precision variable in the corresponding location in the argument list.

So the %zd format specifier for size_t isn't supported. Modifying your kernel like this:

__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);
}

works fine.

Also note you have a pretty major mistake in your host code, although that won't have any effect on the manifest behaviour in your example. You are only allocating and copying one byte for dev_ptr which is obviously incorrect. It should look something like:

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

to transfer the full contents of data from host to device.