Using CUDA printf outside the kernel to print device variables

Question

What is the best way to print device variables in CUDA outside of the kernel? Do I have to do a cudaMemcpy to the host and then print the resulting values? When I try to use printf on pointers created using cudaMalloc, the program crashes. It seems that most of the attention focuses on printing inside the kernel, not in regular code.

Thanks, Eric

Answer 1

"When I try to use printf on pointers created using cudaMalloc, the program crashes"

If you have this:

int *d_data, *h_data;
cudaMalloc(&d_data, DSIZE);

You cannot do this:

printf(" %d ", *d_data);

as this requires dereferencing a device pointer (d_data) in host code which is normally illegal in CUDA.

Instead you can do this:

h_data = (int *)malloc(DSIZE);
cudaMemcpy(h_data, d_data, DSIZE, cudaMemcpyDeviceToHost);
printf(" %d ", *h_data);

You can also investigate Unified Memory which is new in CUDA 6, and see if it will serve your purposes.

And, as mentioned in the comments, devices of cc2.0 or greater support printf from the kernel, which operates on device data (only).

Answer 2

An approach alternative to what suggested by Robert Crovella is to wrap the device pointer into a thrust::device_ptr by thrust::device_pointer_cast. This way is slightly more immediate when you need to access only very few elements of the device array. See the example below:

#include <thrust\device_vector.h>

void main() {

    const int N = 10;

    int *h_data = (int*)malloc(N*sizeof(int));
    for (int i=0; i<N; i++) h_data[i] = 3;

    int *d_data; cudaMalloc(&d_data, N*sizeof(int));    

    cudaMemcpy(d_data,h_data,N*sizeof(int),cudaMemcpyHostToDevice);

    // --- Alternative approach
    thrust::device_ptr<int> dev_ptr_key     = thrust::device_pointer_cast(d_data);
    int i = 4; printf("Element number %d is equal to %d\n",i,(int)*(dev_ptr_key+i));

    getchar();

}