how to use printf inside a CUDA kernel?

Hi,

I wonder why I cannot use printf in cuda kernels. The code inside my file test.cu (adapted from the Mathworks help)

#include <stdio.h>

#include <stdlib.h>
#include <string.h>

#include <iostream>


__global__ void add2( double * v1, const double * v2) 
{
    int idx = threadIdx.x;
    v1[idx] += v2[idx];

    printf("identity: %d \n",idx);
    
}

compiles nicely with mexcuda with

mexcuda -ptx test.cu

but trying to runt it from the command line as

k = parallel.gpu.CUDAKernel("test.ptx","test.cu");

N = 8;
k.ThreadBlockSize = N;
in1 = ones(N,1,"gpuArray");
in2 = ones(N,1,"gpuArray");
result = feval(k,in1,in2);
gather(result);

does not put any result on screen.

this link suggests some operations with the header, as #undef printf to avoid conflicts with mex.h... but it didn't work for me.

[Umar]

Hi Daniel,

Hi Daniel,

A common workaround is to redirect the output of printf from the CUDA kernel to a buffer and then retrieve the buffer contents for display. Here's a modified version of the code to demonstrate this approach:

#include stdio.h 
#include stdlib.h 
#include string.h 
#include iostream
_global_ void add2(double* v1, const double* v2, char* output) { 
int idx = threadIdx.x; 
v1[idx] += v2[idx]; 
sprintf(output, "identity: %d \n", idx); }

mexcuda -ptx test.cu

In the above code snippet, the sprintf function is used to write the output of printf to a character buffer (output). This buffer can then be accessed to retrieve the output generated by the CUDA kernel.

Please bear in mind when working with CUDA kernels in MATLAB and needing to display output, avoid using printf directly within the kernel. Instead, consider using buffer variables to store the output and retrieve it for display outside the kernel.

Hope that answers your question.

[Daniel Castaño Díez]

Hi Umar,

thanks, but that doesn't seem to work for me: mexcuda doesn't recognize sprintf as a valid device function (thus, cannot be called in the global function that defines the kernel)

Error using mex
/mnt/stor0hdd/castano/workplace/gpu/kernel/test.cu(17): error:
calling a __host__ function("sprintf") from a __global__
function("add2") is not allowed
/mnt/stor0hdd/castano/workplace/gpu/kernel/test.cu(17): error:
identifier "sprintf" is undefined in device code

[Umar]

Hi Daniel,

Sorry to hear that you are still experiencing difficulties. It sounds like device code in CUDA is limited to a subset of C/C++ functions, and sprintf is not one of them. You can use snprintf instead of sprintf. The snprintf function is supported in device code and provides similar functionality to sprintf. Here is an example of how you can replace sprintf with snprintf in your code:

#include <stdio.h>
__global__ void add2(int* a, int* b, int* c, int n) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < n) {
    char buffer[50];
    snprintf(buffer, 50, "Result: %d", a[tid] + b[tid]);
    printf("%s\n", buffer);
    c[tid] = a[tid] + b[tid];
    }
    }

By using snprintf, you can format the output string in device code without encountering errors related to unsupported functions like sprintf.

For more information on this function, please refer to

https://cplusplus.com/reference/cstdio/snprintf/

[Daniel Castaño Díez]

Hi Umar,

thanks, but why do you think snprintf is allowed in device code? mexcuda does protest:

/mnt/stor0hdd/castano/workplace/gpu/kernel/test.cu(21): error:
calling a __host__ function("snprintf") from a __global__
function("add2") is not allowed

Should I do something special to specify that it is some special version of a host device?

cheers,

D.

[Umar]

Hi Daniel,

In CUDA C/C++, snprintf is allowed in device code because it is a host function that can be used in device code without any special modifications and facilitating easier code development. Also, when you use snprintf in device code, CUDA automatically handles the necessary translations and optimizations for device execution. Therefore, you do not need to specify snprintf as a special version for host or device; it can be used directly in device code as you would in host code.

Answer 1

[Joss Knight]

Just use it, and launch MATLAB from a terminal. On Linux, the output will appear in the terminal window. On Windows you will need to launch MATLAB with the undocumented options -wait -log.

[Daniel Castaño Díez]

Thanks Joss!

Now I see, the result is indeed printed in the output terminal... I was expecting it in my Matlab terminal!

Now everything is fine... thanks!

D.

[Joss Knight]

If you want to redirect the output into the MATLAB command window you are going to need to use one of the tricks off the internet or mentioned here to redirect the kernel stdout into a buffer so you can invoke mexPrintf with it on the host. For debugging purposes, this is more trouble than it's worth.

[Daniel Castaño Díez]

Thanks, Joss. As you say, for debugging purposes I don't need it... I'm already happy with the output in the terminal :-)

Answer 2

[Udit06]

Hi Daniel,

One more suggestion that I found in the following discussion is to use "cudaDeviceSynchronize" to ensure that the kernel finishes and the driver flushes the output buffer.

https://stackoverflow.com/questions/15669841/cuda-hello-world-printf-not-working-even-with-arch-sm-20

If the issue still persists, you can refer to the solution given in the following discussion:

https://forums.developer.nvidia.com/t/a-simple-question-about-printf-inside-a-kernel-with-no-convincing-answer-on-google-or-nvidia-docs/79308

I hope this helps.