CUDA vector addition, end to end with error checking

A first CUDA program can fail to allocate, copy, or run a kernel, then print zeros and exit 0. The program did not check for errors.

NVIDIA's guide states: "A value of cudaSuccess when checking the error state immediately after a kernel launch does not mean the kernel has executed successfully or even started execution." (https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html , checked 2026-08-29.) The launch does not return an error code. A later call, often the copy back, reports a fault that occurred inside the kernel.

This page gives you the error macro and test harness before the kernel. You will break the program three times and record where each error appears.

Two memories, and a pointer that means nothing to your CPU

The host is your CPU and its RAM. The device is the GPU and its own DRAM, which the course calls global memory. Separate chips, separate memory controllers, and nothing crosses between them unless you copy it.

cudaMalloc(&d_a, bytes) allocates on the device. It takes the address of your pointer because it writes the new device address through it.

The GPU can use that address. Your CPU process cannot.

float* h_a = static_cast<float*>(std::malloc(bytes));  // your process reads it
std::free(h_a);

float* d_a = nullptr;
cudaMalloc(&d_a, bytes);   // the GPU reads it, and only the GPU
cudaFree(d_a);

The calls look alike, but they allocate different memory. std::printf("%f\n", d_a[0]) loads from an address your CPU process does not own, so the process dies.

No CUDA call runs, so CUDA error checks cannot catch this bug. Prefix host pointers with h_ and device pointers with d_ so the argument list shows the mistake.

Data moves with cudaMemcpy, in the order destination, source, bytes, direction. The direction has to agree with the pointers you passed, and the runtime API is explicit that this is your job: "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-29.) The copy is also synchronous: "The cudaMemcpy API is synchronous. That is, it does not return until the copy has completed." (https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html , checked 2026-08-29.)

Pair every cudaMalloc with cudaFree. This program allocates, copies data in, launches a kernel, copies data out, then frees the device memory.

Why the launch cannot tell you anything

In C, a call that fails hands you something back. fopen returns null, malloc returns null, read returns minus one. You learned to check the return value, and that habit is right everywhere except the one line that matters most here.

A kernel launch is not a function call. The guide says it in one sentence: "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 , checked 2026-08-29.) So the short version of this program, the one you copy on your first evening, has four calls that can fail and reports none of them.

cudaMalloc(&d_a, bytes);                                  // returns a code
cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);      // returns a code
vectorAdd<<<blocks, 256>>>(d_a, d_b, d_out, n);           // returns void
cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost);  // returns a code

Three of those four calls returned a cudaError_t, but the program ignored it. The launch returned nothing.

If a bad index makes the kernel fail, the copy on the last line reports the error. Without checks, h_out keeps its old data and the program exits 0.

The first failure can also affect later CUDA calls. From the same page: "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".

A run that reports four errors may show one error four times. Read the first report first. Day 6 explains sticky and non-sticky CUDA errors.

Checking every call, and one guard that has to run

The full program is in code/day05-vector-add/vector_add.cu. It follows three rules.

Put every call that returns cudaError_t inside CUDA_CHECK. Copy the macro without changing it. It prints the file, line, call, and string from cudaGetErrorString, then exits.

#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)

Follow every launch with two checks, in this order. cudaGetLastError() reports an invalid launch, such as a block size above 1024. cudaDeviceSynchronize() waits and reports errors that occurred while the kernel ran, such as an illegal address.

You need both checks. This course writes both lines instead of hiding them in CUDA_CHECK_KERNEL, because a sync changes a timed region and day 9 measures that cost.

n is 611, so some threads fail the guard. 611 is 13 times 47. At 256 threads per block, the launch uses 3 blocks and 768 threads.

The last 157 threads have no element to add, so if (i < n) stops them from writing past the buffer.

If you test only 1024 elements, the guard never rejects a thread. That test can miss a deleted bounds check.

__global__ void vectorAdd(const float* a, const float* b, float* out,
                          size_t n) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (i < n) {
        out[i] = a[i] + b[i];
    }
}

That index line is day 4's global thread index, with one addition: blockIdx.x and blockDim.x are both unsigned int, so without the cast the multiply happens in 32 bits and wraps at 2^32 threads. It is silent when it happens and the answer is wrong in the middle of the array.

The program checks 611 elements but cannot check the bytes after them. Use compute-sanitizer to prove that the 157 extra threads did not write past the buffer.

Note. Do not time this program yet, and in particular do not time it with time ./vector_add. A wall clock around a first CUDA run measures process start, context creation, two copies over PCIe, one kernel and one copy back, and the kernel is the smallest term in that sum by a long way. You will conclude that the GPU is slower than your CPU. For that measurement you will be right, and it will have taught you nothing. It is the archetypal beginner post on NVIDIA's own forums (https://forums.developer.nvidia.com/t/cuda-slower-than-cpu/263339 , checked 2026-08-29), usually from someone running the vector add in https://developer.nvidia.com/blog/even-easier-introduction-cuda/ (both checked 2026-08-29). Day 9 measures one kernel five ways and takes each number apart. Timing arrives there, and not before.

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 capture and the CUDA 13.2 addendum remain in the page's evidence array.

GPU: Tesla T4 (compute capability 7.5)
n = 611, 256 threads per block, 3 blocks, 768 threads
157 threads have no element to add
all 611 elements match the CPU reference

The grid rounds 611 elements up to 3 blocks of 256 threads, so 768 threads launch. The last 157 threads have no element to add. They fail the bounds check and exit without touching memory.

Most input sizes do not divide evenly by the block size. The if (i < n) guard prevents those extra threads from writing past the buffer.

This is why the lesson uses 611 rather than a round number. With 1024 elements and 256 threads the bounds check never fires and you could delete it and never notice, until the day your input is 611.

The first three lines come from source constants: 611 elements, 256 threads per block, and the rounded-up block count. The last line needs the device because it compares every GPU result with the CPU reference.

Run it yourself

Compiler Explorer needs no local GPU, account, or install. Target sm_75 or lower for its current runner; a higher target can compile but fail at run time with no kernel image is available for execution on the device.

This one-file program has no shared headers and fits within the runner's 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 vector_add vector_add.cu

Exercise

Write solve() for vector addition, then take the standalone program, break it the three ways listed in its README, and write down which call reported each failure.

Time: 30 to 40 minutes. Submit: your vector_add.cu, plus three lines naming the call that caught each break.

Check: the harness owns the buffers, stream, and main(). It runs your function on nine sizes, from 0 and 1 through 611, 1024, and 4,194,304, and reports the smallest failing case.

A failure shows the launch configuration, tolerance, arithmetic, and first ten wrong indices with their block and thread. If 1024 passes but 611 fails, the bounds check is missing. Full contract at /reference/harness.

const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
    out[i] = a[i] + b[i];
}

Pitfalls

The program runs, exits 0, and prints zeros. Every runtime call returned a code, but the program ignored each one. Wrap them all, then check the launch with the two lines above.

You check after the synchronize only, and a bad launch slips through. A block size above 1024 is refused at launch and reported as invalid configuration argument by cudaGetLastError(), synchronously. Wait for the synchronize and you see it attributed to the wrong line.

You pass the pointer to cudaMalloc instead of its address. cudaMalloc(d_a, bytes) does not compile, which is the good outcome. The bad one is reaching for the C-style cast older tutorials print, cudaMalloc((void**)d_a, bytes), which compiles, hands the runtime an uninitialised pointer value and comes back as invalid argument. One reason this course writes static_cast and never a C cast.

You delete the bounds check and nothing appears to go wrong. cudaMalloc rounds an allocation up, and nearby pages may already be mapped. A 157-element overrun past a 2,444-byte buffer may not fault even though it is invalid.

compute-sanitizer --tool memcheck ./vector_add reports the write. Compute Sanitizer appears on day 61.

Every call after the first failure returns the same error. An illegal memory access is sticky: the code "will be returned by every CUDA runtime API until the error state has been cleared by calling cudaGetLastError". Read the first message and ignore the ones after it.

A size calculation overflows before it widens. int n = 40000; size_t bytes = n * n * sizeof(float); multiplies in int, wraps, and asks for the wrong number, which comes back as out of memory for an allocation that would have fit.

Go deeper

• CUDA Programming Guide 2.1, "Introduction to CUDA C++", sections 2.1.3 on memory management and 2.1.7 on error handling: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html (checked 2026-08-29)
• CUDA Runtime API, Memory Management, for what cudaMalloc and cudaMemcpy return and when: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html (checked 2026-08-29)
• cuda-samples, cpp/0_Introduction/vectorAdd, NVIDIA's own version of this program: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Introduction/vectorAdd (checked 2026-08-29)
• Programming Massively Parallel Processors, 4th edition, chapter 2, on the host and device split and this exact kernel: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0

Next

Day 6 goes back over the macro you just used and says what it hides: which errors are sticky, why cudaPeekAtLastError exists, and how to find two bugs in a program that reports one. Day 8 rewrites this kernel so one grid covers any n, which is where 611 stops being a special case. Both sit in module 1, and the GPU kernel engineer hub maps where the rest of it goes.