# No printf(".") output from the kernel

**URL:** <https://forums.developer.nvidia.com/t/no-printf-output-from-the-kernel/246115>\
**Category:** CUDA Programming and Performance\
**Created:** [March 15, 2023, 12:01am UTC](https://forums.developer.nvidia.com/t/no-printf-output-from-the-kernel/246115 "2023-03-15T00:01:40Z")\
**Posts on this page:** 7\
**Page:** 1

<div class="post-metadata">

**Author:** ![tiputa](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tiputa](https://forums.developer.nvidia.com/u/tiputa)\
**Post date:** [March 15, 2023, 12:01am UTC](https://forums.developer.nvidia.com/t/no-printf-output-from-the-kernel/246115/1 "2023-03-15T00:01:40Z")

</div>

I am running 496 threads in the 1 block. Each thread is running in a loop where n is about 100E9.

I would like to see the progress by adding print command:

```auto
if (threadIdx.x == 0 && n % 1000000ull == 0) {
  printf(".");
}

```

It seems that there is not print output to the console until the program finishes. Is there any way how to force the print command to output data immediately? Another option would be to use keyboard shortcut to close the program and let him flush the buffer to the console.

---

<div class="post-metadata">

**Author:** ![rs277](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@rs277](https://forums.developer.nvidia.com/u/rs277)\
**Post date:** [March 15, 2023, 12:10am UTC](https://forums.developer.nvidia.com/t/no-printf-output-from-the-kernel/246115/2 "2023-03-15T00:10:24Z")

</div>

printf buffer flushing occurs as per conditions outlined [here.](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#limitations)

---

<div class="post-metadata">

**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\
**Post date:** [March 15, 2023, 12:19am UTC](https://forums.developer.nvidia.com/t/no-printf-output-from-the-kernel/246115/3 "2023-03-15T00:19:53Z")

</div>

If it is a kernel progress report you want, [this](https://stackoverflow.com/questions/20345702/how-can-i-check-the-progress-of-matrix-multiplication/20381924#20381924) may also be of interest.

---

<div class="post-metadata">

**Author:** ![tiputa](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tiputa](https://forums.developer.nvidia.com/u/tiputa)\
**Post date:** [March 15, 2023, 7:10am UTC](https://forums.developer.nvidia.com/t/no-printf-output-from-the-kernel/246115/4 "2023-03-15T07:10:35Z")

</div>

[7.33.2. Limitations](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#limitations) states:  
The output buffer for `printf()` is set to a fixed size before kernel launch (see [Associated Host-Side API](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#associated-host-side-api)). It is circular and if more output is produced during kernel execution than can fit in the buffer, older output is overwritten. **It is flushed only when one of these actions is performed** :

- 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`.

What to concretely use in my case?

---

<div class="post-metadata">

**Author:** ![njuffa](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@njuffa](https://forums.developer.nvidia.com/u/njuffa)\
**Post date:** [March 15, 2023, 7:38am UTC](https://forums.developer.nvidia.com/t/no-printf-output-from-the-kernel/246115/5 "2023-03-15T07:38:35Z")

</div>

What that list tells you is that you _cannot_ use printing dots via device-side `printf()` for a progress report, because the device-side printf-buffer is _not_ continuously copied over to the host while a kernel is in progress. Specific synchronizing activity initiated by the host (enumerated above) is required for the buffer to be copied from device to host.

If you desire a progress report from a running kernel, follow the link to a worked example of how to accomplish that in @Robert_Crovella’s post above.

---

<div class="post-metadata">

**Author:** ![tiputa](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tiputa](https://forums.developer.nvidia.com/u/tiputa)\
**Post date:** [March 15, 2023, 8:18am UTC](https://forums.developer.nvidia.com/t/no-printf-output-from-the-kernel/246115/6 "2023-03-15T08:18:08Z")

</div>

Unfortunately, the example doesn’t work (VS22). It shows just 100% when the program finished:

> kernel starting  
> progress check finished  
> h\_data = 10  
> matrix multiply kernel starting  
> Progress:  
> percent complete = 100.0
> 
> matrix multiply finished. elapsed time = 6877.415039 milliseconds

Programm:

```auto
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <time.h>

#define TIME_INC 100000000
#define INCS 10
#define USE_PROGRESS 1
#define MAT_DIMX 4000
#define MAT_DIMY MAT_DIMX

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__ , __LINE__ ); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

__global__ void mykernel(volatile int* data) {

    unsigned long time;
    for (int i = 0; i < INCS; i++) {
        atomicAdd((int*)data, 1);
        __threadfence_system();
        time = clock64();
        while ((clock64() - time) < TIME_INC) {};
    }
    printf("progress check finished\n");
}

__global__ void matmult(float* a, float* b, float* c, unsigned int rowA, unsigned int colA, unsigned int colB, volatile int* progress) {
    unsigned int row = threadIdx.x + blockDim.x * blockIdx.x;
    unsigned int col = threadIdx.y + blockDim.y * blockIdx.y;
    if ((row < rowA) && (col < colB)) {
        float temp = 0.0f;
        for (unsigned int k = 0; k < colA; k++)
            temp += a[(row * colA) + k] * b[(k * colB) + col];
        c[(row * colB) + col] = temp;
#if USE_PROGRESS
        if (!(threadIdx.x || threadIdx.y)) {
            atomicAdd((int*)progress, 1);
            __threadfence_system();
        }
#endif
    }
}

int main() {
    // simple test to demonstrate reading progress data from kernel
    volatile int* d_data, * h_data;
    cudaSetDeviceFlags(cudaDeviceMapHost);
    cudaCheckErrors("cudaSetDeviceFlags error");
    cudaHostAlloc((void**)&h_data, sizeof(int), cudaHostAllocMapped);
    cudaCheckErrors("cudaHostAlloc error");
    cudaHostGetDevicePointer((int**)&d_data, (int*)h_data, 0);
    cudaCheckErrors("cudaHostGetDevicePointer error");
    *h_data = 0;
    printf("kernel starting\n");
    mykernel << <1, 1 >> > (d_data);
    cudaCheckErrors("kernel fail");
    int value = 0;
    do {
        int value1 = *h_data;
        if (value1 > value) {
            printf("h_data = %d\n", value1);
            value = value1;
        }
    } while (value < (INCS - 1));
    cudaDeviceSynchronize();
    cudaCheckErrors("kernel fail 2");

    // now try matrix multiply with progress

    float* h_c, * d_a, * d_b, * d_c;
    h_c = (float*)malloc(MAT_DIMX * MAT_DIMY * sizeof(float));
    if (h_c == NULL) { printf("malloc fail\n"); return 1; }
    cudaMalloc((void**)&d_a, MAT_DIMX * MAT_DIMY * sizeof(float));
    cudaCheckErrors("cudaMalloc a fail");
    cudaMalloc((void**)&d_b, MAT_DIMX * MAT_DIMY * sizeof(float));
    cudaCheckErrors("cudaMalloc b fail");
    cudaMalloc((void**)&d_c, MAT_DIMX * MAT_DIMY * sizeof(float));
    cudaCheckErrors("cudaMalloc c fail");

    for (int i = 0; i < MAT_DIMX * MAT_DIMY; i++) h_c[i] = rand() / (float)RAND_MAX;
    cudaMemcpy(d_a, h_c, MAT_DIMX * MAT_DIMY * sizeof(float), cudaMemcpyHostToDevice);
    cudaCheckErrors("cudaMemcpy a fail");
    cudaMemcpy(d_b, h_c, MAT_DIMX * MAT_DIMY * sizeof(float), cudaMemcpyHostToDevice);
    cudaCheckErrors("cudaMemcpy b fail");

    cudaEvent_t start, stop;
    cudaEventCreate(&start); cudaEventCreate(&stop);
    *h_data = 0;
    dim3 block(16, 16);
    dim3 grid(((MAT_DIMX + block.x - 1) / block.x), ((MAT_DIMY + block.y - 1) / block.y));
    printf("matrix multiply kernel starting\n");
    cudaEventRecord(start);
    matmult << <grid, block >> > (d_a, d_b, d_c, MAT_DIMY, MAT_DIMX, MAT_DIMX, d_data);
    cudaEventRecord(stop);
#if USE_PROGRESS
    unsigned int num_blocks = grid.x * grid.y;
    float my_progress = 0.0f;
    value = 0;
    printf("Progress:\n");
    do {
        cudaEventQuery(stop); // may help WDDM scenario
        int value1 = *h_data;
        float kern_progress = (float)value1 / (float)num_blocks;
        if ((kern_progress - my_progress) > 0.1f) {
            printf("percent complete = %2.1f\n", (kern_progress * 100));
            my_progress = kern_progress;
        }
    } while (my_progress < 0.9f);
    printf("\n");
#endif
    cudaEventSynchronize(stop);
    cudaCheckErrors("event sync fail");
    float et;
    cudaEventElapsedTime(&et, start, stop);
    cudaCheckErrors("event elapsed time fail");
    cudaDeviceSynchronize();
    cudaCheckErrors("mat mult kernel fail");
    printf("matrix multiply finished. elapsed time = %f milliseconds\n", et);

    return 0;
}

```

---

<div class="post-metadata">

**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\
**Post date:** [March 15, 2023, 2:28pm UTC](https://forums.developer.nvidia.com/t/no-printf-output-from-the-kernel/246115/7 "2023-03-15T14:28:28Z")

</div>

try enabling [windows hardware scheduling](https://www.howtogeek.com/756935/how-to-enable-hardware-accelerated-gpu-scheduling-in-windows-11/#:~:text=On%20Windows%2010%2C%20head%20into,Hardware%2DAccelerated%20GPU%20Scheduling.%22).