100 DAYS OF GPU Course preview
Day 6Module 1
in-technical-review

How to check for errors in CUDA

PyTorch ships this text inside its own CUDA error messages, from get_cuda_async_error_suffix in https://github.com/pytorch/pytorch/blob/main/c10/cuda/CUDAMiscFunctions.cpp (checked 2026-08-30):

CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.
For debugging consider passing CUDA_LAUNCH_BLOCKING=1

PyTorch warns that its stack trace may point to the wrong line because CUDA launches run after the host moves on. A kernel launch returns no status value.

The next CUDA runtime call often reports the fault. That call is usually a valid copy.

Day 5 introduced CUDA_CHECK. This page explains what the macro does, when an error appears late, and why some kernel errors make every later CUDA call fail.

What a call does with its status, and what a launch does without one

Every function in the CUDA runtime API returns a cudaError_t, and the reference says what each one gives you on failure. cudaMalloc is typical: "cudaMalloc() returns cudaErrorMemoryAllocation in case of failure." (https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html , checked 2026-08-30.) C does not warn on a discarded return value, so a program that throws every one of them away compiles clean.

The launch is the exception, and it is the one that matters. "Kernel launches using triple chevron notation do not return a cudaError_t." (https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html section 2.1.7, checked 2026-08-30.) There is no value to ignore, so there is nothing to wrap.

The runtime keeps one status per host thread, and each failure overwrites it. cudaGetLastError "Returns the last error that has been produced by any of the runtime calls in the same instance of the CUDA Runtime library in the host thread and resets it to cudaSuccess."

cudaPeekAtLastError reads the same status without resetting it: "This call does not reset the error to cudaSuccess like cudaGetLastError()." (Both https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__ERROR.html , checked 2026-08-30.)

Use cudaGetLastError when you handle a failure, so a later check does not report the old status. Use cudaPeekAtLastError when a logger must read the status without clearing it.

Queued work runs after the launch statement returns, so a fault inside a kernel cannot report itself at that point. NVIDIA instead warns on each runtime function that may report an earlier fault.

"Note that this function may also return error codes from previous, asynchronous launches" appeared 56 times on the memory management page when parsed on 2026-08-30. The note appears on cudaMalloc, cudaMemcpy, cudaFree, and most other functions on that page.

Each launch in this course therefore has two checks.

myKernel<<<blocks, kThreadsPerBlock>>>(d_in, d_out, kElems);
CUDA_CHECK(cudaGetLastError());        // a launch the driver refused
CUDA_CHECK(cudaDeviceSynchronize());   // what the kernel then did

cudaGetLastError catches a launch the driver rejected, such as a block size above 1024. cudaDeviceSynchronize waits for the kernel and reports errors that occurred while it ran.

You need both checks. This course keeps the sync visible because it changes what a timed region measures.

Three failures surface at different distances: a null copy at 0 calls, a refused launch at 1 check, and an illegal write at the first synchronizing call, 2 calls later here.

The habit that breaks here: one mistake, one message

In most CPU code, the failing call reports its own error. errno belongs to the call that failed, and a Python traceback points to the line that raised.

CUDA may instead point you to a valid copy that reported an earlier kernel fault.

Two NVIDIA passages describe why the message repeats.

From the programming guide: "When errors are returned by CUDA runtime API functions, the error state is not cleared. This means that error code from an asynchronous error, such as an invalid memory access by a kernel, will be returned by every CUDA runtime API until the error state has been cleared by calling cudaGetLastError." From the runtime reference, on the code an out-of-bounds write produces: "This leaves the process in an inconsistent state and any further CUDA work will return the same error. To continue using CUDA, the process must be terminated and relaunched." (https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html , checked 2026-08-30.)

The first quote describes the status variable, which cudaGetLastError resets. The second describes a damaged context, which the process cannot repair.

After such a fault, clearing the status does not help. The next CUDA call fails and sets the status again.

These are called sticky and non-sticky errors. Eleven of the 137 cudaError enumerators carried the "inconsistent state" sentence when parsed on 2026-08-30.

The other 126 enumerators do not carry that sentence. cudaErrorIllegalAddress (700) and cudaErrorLaunchFailure (719) are sticky. cudaErrorInvalidValue (1) and cudaErrorMemoryAllocation (2) are not, so they do not damage the context.

The count also matches the 137 errors that research/ERROR-PAGES.md section 2 verified on 2026-08-29.

Why the macro is shaped the way it is

The full program is in code/day06-error-checking/error_checking.cu. It makes two errors on purpose, in four parts, and prints which call reports each one. It does not measure time because most parts synchronize the device.

Day 5 gave you the macro and told you to copy it. Here is what each piece is doing.

#define CUDA_CHECK(call)                                                 \
    do {                                                                 \
        cudaError_t err_ = (call);                                       \
        if (err_ != cudaSuccess) {                                       \
            std::fprintf(stderr, "CUDA error %s:%d: %s: %s\n", __FILE__, \
                         __LINE__, #call, cudaGetErrorString(err_));     \
            std::exit(EXIT_FAILURE);                                     \
        }                                                                \
    } while (0)

do { } while (0) makes the macro one statement. With a bare block, if (x) CUDA_CHECK(...); else y(); does not compile because the semicolon leaves the else with no matching if. Compilers remove the while (0), so it has no run-time cost.

err_ carries a trailing underscore so the macro cannot collide with a variable called err at the call site. A macro pastes its body into your scope, and the underscore is what stops it shadowing something of yours.

#call is the call's own source text. The preprocessor turns the argument into a string literal, so the message names what failed and not only where. NVIDIA's own samples use the same trick: #define checkCudaErrors(val) check((val), #val, __FILE__, __LINE__), line 598 of https://github.com/NVIDIA/cuda-samples/blob/master/Common/helper_cuda.h (checked 2026-08-30).

It prints cudaGetErrorString and not the number, because the string is what a person pastes into a search box. cudaGetErrorName gives you the enum name instead, and the program prints both side by side so you can see which is which.

Part 1 of the program reads the same failure four times to show what the two readers do:

    // A null destination. This fails synchronously with
    // cudaErrorInvalidValue, the runtime rejects it before anything is
    // queued, and the context is untouched, which is what part 2 needs.
    //
    // The obvious choice, passing the wrong cudaMemcpyKind, does not work.
    // Measured on a Tesla T4 with CUDA 12.6: a host-to-device copy tagged
    // cudaMemcpyDeviceToHost returns cudaSuccess and copies correctly,
    // because unified addressing lets the runtime read the real direction
    // off the pointers and the kind argument is a hint it can overrule.
    // That is worth knowing on its own: the direction flag will not catch
    // your mistake for you.
    const cudaError_t badCopy =
        cudaMemcpy(nullptr, h_in.data(), bytes, cudaMemcpyHostToDevice);
    const cudaError_t peek1 = cudaPeekAtLastError();
    const cudaError_t peek2 = cudaPeekAtLastError();
    const cudaError_t taken = cudaGetLastError();
    const cudaError_t peek3 = cudaPeekAtLastError();

Part 3 is the one worth reading twice. The two lines from the code style are deliberately missing after this launch, so the fault has to find its own way out:

    writePastEnd<<<blocks, kThreadsPerBlock>>>(d_out, kElems, kOvershoot);
    const cudaError_t atLaunch = cudaPeekAtLastError();
    report("peek, straight after the launch", atLaunch);

    // Nothing is wrong with this copy. It reads a buffer that exists, into a
    // vector that exists, with the direction its pointers agree on. It is
    // just the next runtime call, which is the whole reason a CUDA error can
    // name a line that did nothing.
    const cudaError_t atNextCall =
        cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost);

Note. Two honest limits on this program. The reversed copy in part 1 is documented as undefined behaviour, "Calling cudaMemcpy() with dst and src pointers that do not match the direction of the copy results in an undefined behavior" (https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html , checked 2026-08-30), so the code it returns is not promised and the part is written to demonstrate the distance rather than the number. And the out-of-bounds write in part 3 overshoots by a whole gibibyte, not by a few elements, because cudaMalloc rounds an allocation up and a short overrun can land in memory the driver already owns and never fault. Both are stated in the source, and both parts return a non-zero exit code if the card does something else.

Results

Re-verified without behavioral drift on a Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88) on 2026-09-02. The original CUDA 12.6 transcript remains beside the new run in the page's evidence array.

GPU: Tesla T4 (compute capability 7.5)

part 1: one failed call, and who can still see it
  cudaMemcpy returned              cudaErrorInvalidValue          invalid argument
  cudaPeekAtLastError              cudaErrorInvalidValue          invalid argument
  cudaPeekAtLastError, again       cudaErrorInvalidValue          invalid argument
  cudaGetLastError                 cudaErrorInvalidValue          invalid argument
  cudaPeekAtLastError, after get   cudaSuccess                    no error

part 2: the same context after a non-sticky error
  all 611 elements match, so the context survived it

part 3: an error that arrives late
  peek, straight after the launch  cudaSuccess                    no error
  the next cudaMemcpy returned     cudaErrorIllegalAddress        an illegal memory access was encountered

part 4: what clearing a sticky error buys you
  cudaGetLastError returned        cudaErrorIllegalAddress        an illegal memory access was encountered
  peek, right after clearing it    cudaSuccess                    no error
  a fresh cudaMalloc returned      cudaErrorIllegalAddress        an illegal memory access was encountered
  cudaFree(d_in) returned          cudaErrorIllegalAddress        an illegal memory access was encountered
  cudaFree(d_out) returned         cudaErrorIllegalAddress        an illegal memory access was encountered

the context is gone. Only a new process gets CUDA back.

In part 1, cudaPeekAtLastError reports the same failure twice without clearing it. cudaGetLastError reports it once and clears it, so the next peek returns cudaSuccess.

In part 3, the launch check reports cudaSuccess because the asynchronous error has not happened yet. The next cudaMemcpy reports the fault even though the copy is valid.

Part 4 shows a sticky error. After the illegal access, the next cudaMalloc and both cudaFree calls return the same error even after the status is clear. Only a new process can use CUDA again, so cudaGetLastError diagnoses the fault but cannot recover from it.

Run it yourself

Compiler Explorer needs no local GPU, account, or install. Both programs use one file with no shared headers, so you can paste either one into its editor.

Target sm_75 or lower for the current runner. A higher target can compile but fail at run time with no kernel image is available for execution on the device. Both programs fit within its 20-second compile and run limits.

A Colab CUDA runtime or a local CUDA GPU works the same way. The build line is in the repo's README:

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

Exercise

Take code/day06-error-checking/starter/error_checking.cu, a vector add over 8,388,608 elements that compiles without a warning and prints a wrong answer. Add the checking, then find and fix both bugs.

Time: 25 to 40 minutes. Submit: your fixed error_checking.cu, plus two lines naming the call that reported each bug and how far it was from the line that caused it.

Check: build the starter, run it, and read which call reported each bug. It passes when the program prints all 8388608 elements match the CPU reference and exits 0. There is no separate harness on this day: the starter is a self-contained program with its own main(), and making it tell you what went wrong is the entire exercise.

Hint 1

You will need two runs, not one. The first failure stops the program, so the second bug stays invisible until the first is gone.

Both bugs are one token each, both are in main(), and neither is in the kernel.

Hint 2

For each call between the allocations and the comparison, ask whether the argument you passed agrees with the pointers you passed. Then, separately: how many bytes does each buffer hold, and how many does the kernel write into it?

Solution

Two edits. Here they are in source order. You meet the lower one first, because the allocation does not fail on the line that makes it.

-cudaMalloc(&d_out, kElems);
+CUDA_CHECK(cudaMalloc(&d_out, bytes));
...
-cudaMemcpy(d_b, h_b.data(), bytes, cudaMemcpyDeviceToHost);
+CUDA_CHECK(cudaMemcpy(d_b, h_b.data(), bytes, cudaMemcpyHostToDevice));

Run one: the copy is caught by the copy. The pointers say host to device and the argument says the opposite. Wrapped, the macro names that line and that call, and the distance is zero calls. Fix it and go again.

Run two: a later call reports the bad allocation size. cudaMalloc(&d_out, kElems) asks for 8,388,608 bytes, but the kernel writes 8,388,608 floats. The buffer is one quarter of the required size, so 24 MiB of writes go past its end.

The allocation and launch both succeed because each request is valid on its own. Your cudaDeviceSynchronize() reports an illegal memory access was encountered. Without that sync, the copy back reports it instead.

The if (i < n) check compares the index with the element count. It cannot know how much memory the caller allocated.

Use Compute Sanitizer to find this mismatch. compute-sanitizer --tool memcheck ./error_checking names the write, thread, and block even if the invalid address happens to be mapped.

A CUDA error names the call that reports it, not always the code that caused it. Check every call so the report stays close to the fault.

Pitfalls

Every call after the first failure returns the same code, so you fix four bugs that do not exist. An illegal access is sticky: the reference says the process "must be terminated and relaunched". Read the first message and ignore the rest, and see an illegal memory access was encountered.

A later cudaMemcpy reports invalid configuration argument. That code is raised synchronously at the launch, so if a copy is reporting it you have an unchecked launch earlier in the program and the copy is only the first call that looked. Put cudaGetLastError() under every launch and the report moves to where it belongs. Day 4 picks legal launch numbers; the error page ranks the other causes.

You check after the synchronise only, and a refused launch is blamed on the wrong line. A block size above 1024 never runs at all, and cudaDeviceSynchronize() will happily report it as though the kernel had failed. The two lines exist because they catch different things.

You get unspecified launch failure instead of an address. Treat 719 as 700 with less information: same causes, a card or driver that could not attribute the fault to an instruction. Both are on the list of five codes PyTorch tags with its "might be asynchronously reported" warning, alongside 710, 715 and 716. The error page has the split, and compute-sanitizer finds the write either way.

Your logger prints no error on every successful call. cudaGetErrorString(cudaSuccess) returns the string no error, not an empty one, so a line that formats the status unconditionally fills your output with it. Print only when the status is not cudaSuccess, which is what the macro's if is for.

You reach for CUDA_LAUNCH_BLOCKING=1 and leave it on. It is the right tool for finding which launch faulted, because it makes the reporting call the real one. NVIDIA is blunt about the price: "Disabling asynchronous execution results in slower execution but is useful for debugging" (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/environment-variables.html , checked 2026-08-30). A benchmark run with it set measures something else.

Go deeper

Next

Day 7 moves to two-dimensional grids, where the bounds check must cover both axes. Day 8 rewrites day 1's kernel so one grid can cover any n.

Day 61 uses Compute Sanitizer to find writes that do not fault. These lessons build on error checking in module 1.