printf() in my CUDA kernel doesn't result produce any output

Question

I have added some printf() statements in my CUDA program

__device__ __global__ void Kernel(float *, float * ,int );
void DeviceFunc(float *temp_h , int numvar , float *temp1_h)
{ .....
    //Kernel call
    printf("calling kernel
");
    Kernel<<<dimGrid , dimBlock>>>(a_d , b_d , numvar);
    printf("kernel called
");
  ....
}

int main(int argc , char **argv)
{   ....
    printf("beforeDeviceFunc

");
    DeviceFunc(a_h , numvar , b_h); //Showing the data
    printf("after DeviceFunc

");
    ....
}

Also in the Kernel.cu, I wrote:

#include<cuda.h>
#include <stdio.h>
__device__ __global__ void Kernel(float *a_d , float *b_d ,int size)
{
    int idx = threadIdx.x ;
    int idy = threadIdx.y ;
    //Allocating memory in the share memory of the device
    __shared__ float temp[16][16];

    //Copying the data to the shared memory
    temp[idy][idx] = a_d[(idy * (size+1)) + idx] ;
    printf("idx=%d, idy=%d, size=%d", idx, idy, size);
    ....
}

Then I compile using -arch=sm_20 like this:

nvcc -c -arch sm_20 main.cu
nvcc -c -arch sm_20 Kernel.cu
nvcc -arch sm_20 main.o Kernel.o -o main

Now when I run the program, I see:

beforeDeviceFunc

calling kernel
kernel called
after DeviceFunc

So the printf() inside the kernel is not printed. How can I fix that?

Answer 1

printf() output is only displayed if the kernel finishes successfully, so check the return codes of all CUDA function calls and make sure no errors are reported.

Furthermore printf() output is only displayed at certain points in the program. Appendix B.32.2 of the Programming Guide lists these as

• Kernel launch via <<<>>> or cuLaunchKernel() (at the start of the launch, and if the CUDA_LAUNCH_BLOCKING environment variable is set to 1, at the end of the launch as well),
• Synchronization via cudaDeviceSynchronize(), cuCtxSynchronize(), cudaStreamSynchronize(), cuStreamSynchronize(), cudaEventSynchronize(), or cuEventSynchronize(),
• Memory copies via any blocking version of cudaMemcpy*() or cuMemcpy*(),
• Module loading/unloading via cuModuleLoad() or cuModuleUnload(),
• Context destruction via cudaDeviceReset() or cuCtxDestroy().
• Prior to executing a stream callback added by cudaStreamAddCallback() or cuStreamAddCallback().

To check this is your problem, put the following code after your kernel invocation:

{
    cudaError_t cudaerr = cudaDeviceSynchronize();
    if (cudaerr != cudaSuccess)
        printf("kernel launch failed with error "%s".
",
               cudaGetErrorString(cudaerr));
}

You should then see either the output of your kernel or an error message.

More conveniently, cuda-memcheck will automatically check all return codes for you if you run your executable under it. While you should always check for errors anyway, this comes handy when resolving concrete issues.