how to use printf inside a CUDA kernel?

2024 6 28
2 个回答
回答已采纳
更新时间：2024 7 11
51 次查看（30 天）

0 个投票

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.

5 个评论
显示 3更早的评论 隐藏 3更早的评论

Umar
Umar 2024-6-28
编辑：Walter Roberson 2024-6-28
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.
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
Umar 2024-7-2
编辑：Walter Roberson 2024-7-2
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/
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
Umar 2024-7-4
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.

请先登录，再进行评论。

请先登录，再回答此问题。

 采纳的回答

Joss Knight
Joss Knight 2024-7-6
编辑：Joss Knight 2024-7-6

0 个投票

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.

3 个评论
显示 1更早的评论 隐藏 1更早的评论

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
Joss Knight 2024-7-8
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.
Thanks, Joss. As you say, for debugging purposes I don't need it... I'm already happy with the output in the terminal :-)

请先登录，再进行评论。

更多回答（1 个）

Udit06
Udit06 2024-7-1

0 个投票

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.

类别

在 帮助中心 和 File Exchange 中查找有关 Get Started with GPU Coder 的更多信息

产品

版本

R2024a

标签

Community Treasure Hunt

Find the treasures in MATLAB Central and discover how the community can help you!

Start Hunting!

Translated by