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

This all depends on which version of the CUDA libraries you're using. I believe printf directly from kernel was added somewhere around 2.1. If you want to printf from hostside, you'll need to use cudaMemcpy, yes. — Tyler Jandreau
– Tyler Jandreau, Commented May 5, 2014 at 13:50

Robert Crovella · Accepted Answer · 2014-05-05 13:56:24Z

8

"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).

answered May 5, 2014 at 13:56

Robert Crovella

154k12 gold badges255 silver badges300 bronze badges

Sign up to request clarification or add additional context in comments.

Comments

Vitality · Accepted Answer · 2014-05-13 12:45:36Z

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

}

Collectives™ on Stack Overflow

Using CUDA printf outside the kernel to print device variables

2 Answers 2

Comments

Comments

Your Answer

Hot Network Questions

Collectives™ on Stack Overflow

2 Answers 2

Comments

Comments

Your Answer

Sign up or log in

Post as a guest

Related