100 DAYS OF GPU Course preview
Day 64Module 7
in-technical-review

printf, assert and when printf lies

You put a printf in a kernel, the program runs, exits zero, and nothing prints. The kernel launched and every thread executed the printf. The missing output comes from how CUDA buffers and flushes device output.

This page covers three ways a device printf line can go missing. It also shows how assert names the thread with bad data and then invalidates the CUDA context. Neither call is a correctness check.

Where a printf line waits

A printf in device code does not write straight to your terminal. It appends a record containing the format string's address and packed arguments to a fixed-size device buffer. The host formats the record later.

You can read and set the buffer size through cudaDeviceGetLimit and cudaDeviceSetLimit with cudaLimitPrintfFifoSize, and the default is 1 megabyte (https://docs.nvidia.com/cuda/archive/12.6.3/cuda-c-programming-guide/index.html#formatted-output , section 7.35.3, checked 2026-09-01).

Two buffer properties explain most device printf behavior. First, it is a ring: "It is circular and if more output is produced during kernel execution than can fit in the buffer, older output is overwritten." Your first thousand lines are exactly the ones a full buffer throws away.

Second, it is emptied only at a short list of flush points, quoted from the same section: a kernel launch, synchronization via cudaDeviceSynchronize() and the stream and event variants, any blocking memory copy, module load and unload, and context destruction. Between flush points the lines exist only as records on the device.

The guide states a separate exit rule: "Note that the buffer is not flushed automatically when the program exits." A kernel that prints, followed by return 0;, produces a program whose records remain in the ring when the process ends. The kernel can run without error while the terminal receives no lines.

Diagram: a printf line's route to the screen has a waiting room. A kernel on the left, the printf ring buffer as a strip in the middle, the terminal on the right, arrows left to right. Band 1, the kernel runs and 32 records land in the strip; the terminal is empty. Caption "32 records buffered, 0 lines on screen." Band 2, a cudaDeviceSynchronize() bar crosses the strip and the records drain to the terminal. Caption "one flush point, all 32 lines at once." Band 3, the program exits with no bar; the strip still holds the records, greyed out. Caption "no flush point before exit: 32 written, 0 delivered." Alt text: "A kernel's printf lines land in a fixed circular buffer, not the screen. One sync point delivers all thirty-two at once, and exiting without one delivers zero."

The habit that printf is a terminal

Host print output often appears when a line runs, in program order. Device output waits for a flush point, and CUDA does not promise an order among threads. Each thread of a warp executes the printf "per-thread, and in the context of the calling thread", and the guide's own five-thread example prints in the order 2, 1, 4, 0, 3.

Day 1 measured the same non-promise one level up: across 20 runs, block 0 printed first 16 times and block 1 four times. Day 21 therefore checks ordering by writing to a device array instead of printing.

Using printf to inspect a race can change the race. The guide states: "Internally printf() uses a shared data structure and so it is possible that calling printf() might change the order of execution of threads." A thread that prints takes a longer, serialized path than one that does not.

Day 27's fence-free reduction produced the right answer on all 100 runs despite its race. Adding printf changes the timing that determines whether the race appears. Use day 62's racecheck under Compute Sanitizer, and the tool for inspecting a live thread is day 63's cuda-gdb, which stops the thread instead of racing it.

One file, four builds

Full program in code/day64-printf/printf_assert.cu. It is one source file the README builds four ways: the default demonstration, a -DDAY64_LOSE_OUTPUT=1 build carrying the deliberate exit-without-sync bug, a -DNDEBUG build where the assert vanishes, and a -DDAY64_POISON=0 build with clean data.

The FIFO is shrunk before anything prints. The runtime reference is strict: setting cudaLimitPrintfFifoSize "must not be performed after launching any kernel that uses the printf() device system call - in such case cudaErrorInvalidValue will be returned" (https://docs.nvidia.com/cuda/archive/12.6.3/cuda-runtime-api/group__CUDART__DEVICE.html , checked 2026-09-01). So part 0 reads the default, shrinks the ring to 4 KiB, and reads it back:

    size_t fifoDefault = 0;
    CUDA_CHECK(cudaDeviceGetLimit(&fifoDefault, cudaLimitPrintfFifoSize));
    std::printf("printf FIFO default: %zu bytes\n", fifoDefault);
    CUDA_CHECK(cudaDeviceSetLimit(cudaLimitPrintfFifoSize, kFifoBytes));
    size_t fifoNow = 0;
    CUDA_CHECK(cudaDeviceGetLimit(&fifoNow, cudaLimitPrintfFifoSize));
    std::printf("printf FIFO now:     %zu bytes\n", fifoNow);

Overflow comes from one thread, so the order is provable. With many threads you cannot say which records were oldest. One thread's loop has one order, so "older output is overwritten" becomes checkable: survivors must be one contiguous run ending at the last line.

__global__ void printPastFifo(int lines) {
    for (int k = 0; k < lines; ++k) {
        printf("fifo line %04d of %d\n", k, lines);
    }
}

The one assert in the file is the specimen, not the harness. The host fills 32 floats with their own indices and poisons element 5 on purpose. Exactly one thread's assert fires, and its message names that thread, which is the thing printf cannot do without you printing from all 32 and reading the pile:

__global__ void assertEachElement(const float* __restrict__ in, size_t n) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (i < n) {
        const float want = static_cast<float>(i);
        // NDEBUG deletes the next line; the -DNDEBUG build proves it.
        assert(in[i] == want);  // deliberate device-side assert (day 64)
    }
}

The guide's assertion section states the cost: "Any subsequent host-side synchronization calls made for the same device will return cudaErrorAssert." (https://docs.nvidia.com/cuda/archive/12.6.3/cuda-c-programming-guide/index.html#assertion , section 7.32, checked 2026-09-01.) After the assert, the program expects its fresh cudaMalloc and cudaFree(d_in) calls to report the same error.

This is the sticky error class from day 6. Every gate is a branch returning EXIT_FAILURE; on the poisoned build, a clean context fails the gate.

The file ends with an experiment prompted by two NVIDIA documents. The guide's assertion section says no more commands reach the device "until cudaDeviceReset() is called to reinitialize the device". The runtime reference's entry for cudaErrorAssert (710) says "the process must be terminated and relaunched" (https://docs.nvidia.com/cuda/archive/12.6.3/cuda-runtime-api/group__CUDART__TYPES.html , checked 2026-09-01).

The program calls reset, tries one more allocation, and reports the result. It does not gate that observation, just as day 27 does not gate its racy rows.

The shipped source keeps the deliberate no-sync variant behind a switch. The fix is visible in the same transcript: part 1 syncs after launching this same kernel and gets its 32 lines; the tail below never syncs and gets none.

    // deliberate bug (day 64): launch a printing kernel, then exit with
    // no sync. The 12.6 guide, section 7.35.2: "the buffer is not flushed
    // automatically when the program exits." These 32 lines are dropped.
    // The fix is the cudaDeviceSynchronize() part 1 runs on this same
    // kernel: this transcript holds 32 `device: thread` lines, not 64.
    std::printf("part 3: launching, then exiting with no sync\n");
    std::fflush(stdout);
    printPerThread<<<1, kProbeThreads>>>();
    // deliberately unchecked (day 64): even cudaGetLastError() is not a
    // flush point, but the point lands harder with nothing at all here.
    return EXIT_SUCCESS;

How many of the 2048 lines survive the 4 KiB ring is not derivable from the string lengths, because the ring stores argument records and metadata, not formatted text. The count is a driver property. The shape is not, and the shape is what gets predicted.

Results

Run on the project's Tesla T4 (driver 595.84, CUDA 12.6 V12.6.85) on 2026-09-01. Four builds, five transcripts, in code/day64-printf/evidence/. Two predictions held, two died, and one of the deaths changed the program.

CUDA 13.0 rebuilt and ran all four variants successfully. The FIFO limits, sync-dependent flushing, 32-line lose-output result, sticky assert statuses, reset result and -DNDEBUG behavior all matched. Survivor counts changed to 9,554, 5,554, 7,545 and 9,534 for the default, lose-output, -DNDEBUG and fixed builds, respectively, reinforcing that those counts are observations, not constants.

Died before anything else could run: the FIFO does not shrink to what you ask for. cudaDeviceGetLimit reported a default of 1,310,720 bytes. The program asked for 4,096 but got 262,144.

The driver rounds this limit up to its own granularity, and 4 KiB is far below it, so the original gate (fifoNow != kFifoBytes is a failure) refused to let the program past part 0. The gate now checks that the driver honoured a request at all, and part 2 sizes its line count from what was actually granted rather than from what was asked. Ask for a limit, then read it back, is the rule this cost us.

Held: the buffered lines appear at the sync, not at the launch. In every transcript the 32 device: thread lines sit after host: launch returned, no sync yet and before host: sync returned, the buffer is flushed. The host marker is pinned by an fflush, so the ordering is the device buffer's, not stdio's.

Held: exiting without a sync drops the lines. The -DDAY64_LOSE_OUTPUT=1 build launches printPerThread twice and its transcript holds 32 device: thread lines, not 64. Same kernel, same build, one cudaDeviceSynchronize apart.

Died: the survivors are not one contiguous tail. The prediction said a full FIFO leaves one run of the newest records, with line 0 gone. What 32,768 lines into a 262,144-byte buffer actually left was 6,083 survivors in 11 separate contiguous runs, starting at fifo line 0000 and ending at fifo line 32767. The reason is that nothing waits: the driver drains the FIFO while the kernel is still filling it, so records are consumed and overwritten at the same time.

The survivors are the records that had not drained when each wrap occurred. The counts are not stable either, 6,083 / 8,278 / 8,177 / 7,982 across the four builds of the same program. "The oldest records are overwritten" is true of a full buffer at an instant, and useless as a prediction about which lines you will see.

Held exactly: the assert fires once and poisons everything after it.

printf_assert.cu:121: void assertEachElement(const float *, unsigned long):
  block: [0,0,0], thread: [5,0,0] Assertion `in[i] == want` failed.
  cudaGetLastError at launch     cudaSuccess              no error
  cudaDeviceSynchronize          cudaErrorAssert          device-side assert triggered
  a fresh cudaMalloc             cudaErrorAssert          device-side assert triggered
  cudaFree(d_in)                 cudaErrorAssert          device-side assert triggered

The launch itself reports success. The failure surfaces at the sync, and after that a cudaMalloc that has nothing to do with the assert reports the assert too. This is day 6's sticky-context rule with a different cause.

The tiebreaker, settled. The programming guide says the device comes back after cudaDeviceReset(); the runtime reference says the process must be terminated and relaunched. On this card the reference wins:

  cudaDeviceReset                cudaSuccess              no error
  cudaMalloc after reset         cudaErrorDevicesUnavailable CUDA-capable device(s) is/are busy or unavailable

The reset returns success and the device still refuses to allocate. A fired device assert ends the process, whatever the reset call reports.

Held: -DNDEBUG deletes the check. Same poisoned element, same kernel, and every status is cudaSuccess. The bug is still there and nothing in the program's own output says so, which is why no gate in this course is an assert.

Run it yourself

This day's program fits Compiler Explorer: one file, no headers beyond the runtime, and work that fits the 20 second compile and 20 second run caps. Paste the file, then switch builds by typing -DDAY64_LOSE_OUTPUT=1, -DNDEBUG or -DDAY64_POISON=0 into the compiler options box.

One difference from a shell: CE shows stdout and stderr in separate panes, so the assert message will not sit between the printf lines. The local command below targets sm_75; replace it with the target for your GPU.

nvcc -std=c++17 -O3 -arch=sm_75 -o printf_assert printf_assert.cu

The README has all four build lines and the grep loop that counts survivors out of the transcripts.

Exercise

Run all four builds and collect three facts: the survivor count and first surviving line number from part 2, the device: thread line count in the lose-output transcript, and which document won the reset experiment. Then take the quiz.

Time: 20 to 35 minutes. Submit: the three facts, plus one sentence on why the first lines were the ones the overflow destroyed.

Check: three questions marked in the browser from content/quizzes/day64.toml, with an explanation on every option. The program also gates itself: every expected status is a real branch returning EXIT_FAILURE, the poisoned build expecting cudaErrorAssert from four calls, so a run that exits 0 has checked its own story.

Hint 1

Which operation flushes records from the fixed-size device buffer to the terminal? What happens to the oldest buffered record when the ring fills?

Hint 2

The ring stores what printf was asked to print, not the printed text: a format string address and the arguments. So which build knob changes the survivor count, and why can you not compute that count from the 21 visible characters of each line?

Solution

cudaDeviceSynchronize() is the flush point here. When the ring fills, a new record overwrites the oldest buffered record.

The survivor count depends on each record's argument bytes and metadata, which the visible text does not reveal. Only cudaLimitPrintfFifoSize changes the buffer limit. Device printf records can arrive late, out of order, overwritten, or not at all because they remain buffered until a flush point.

Pitfalls

Your kernel prints nothing and the program exits 0. There is no flush point between the launch and the exit, and "the buffer is not flushed automatically when the program exits." Put a cudaDeviceSynchronize() after the launch, which the error checking discipline from day 6 already had you doing. Day 6 covers what else that sync catches.

The first lines of your output are missing. The ring wrapped and overwrote the oldest records. Raise cudaLimitPrintfFifoSize with cudaDeviceSetLimit, and do it before the first printing kernel launches: after one, the set fails with cudaErrorInvalidValue.

You added a printf and the bug went away. The printf serializes threads through a shared structure and changes the timing your race depends on. Day 14 measured a race that hides at one block size, day 27 one that hid for 100 straight runs; a printf that "fixes" your kernel has diagnosed it. Confirm with day 62's racecheck.

%zu prints garbage from device code. Device printf supports only h, l and ll as size modifiers and will accept an invalid combination without complaint, leaving the result host-dependent. Cast to unsigned long long and print with %llu. This program's kernels print only int and unsigned for that reason; the host side keeps %zu.

You gated a test with assert and CI stayed green. Release builds define NDEBUG, which deletes the assert at the preprocessor, and the guide itself recommends disabling assertions in production code. Every gate in this course is a branch returning EXIT_FAILURE; the -DNDEBUG build of this very program is the demonstration. Day 66 builds the full testing story.

After one assert, every later call reports a device-side assert. cudaErrorAssert (710) is in day 6's sticky class: the reference says the device "cannot be used again". Read the first failure, not the twentieth, and treat the rest as repeated errors. Whether cudaDeviceReset() revives the device is this lesson's open experiment; the process relaunch always works.

Go deeper

Next

Day 65 diagnoses the hangs, where the program prints nothing because it is still running: barriers half a warp never reaches, host and device each waiting on the other. Day 66 then takes the lesson under this lesson, that neither printf nor assert is a test, and builds the day 5 harness out into property tests with sanitizer runs in CI.