CUDA's default failure mode is silence — notes from a tester's first week

In my quest to learn how AI infrastructure is built, I started reading about GPUs and learning by doing. I'd measured GPUs from the outside before — utilization, throughput, the numbers you read off the card while it runs. This time I wanted to see inside the kernel. This article is about my experience as a first-time learner, and the thing that struck me during the exercise wasn't how the GPU worked — it was how quietly it broke.

A small program

I wrote about the simplest thing you can write on a GPU: add two lists of numbers. Take 1000 numbers in list A, 1000 in list B, and produce a third list C where each element is the sum of the matching pair. On a CPU you'd write a loop. On a GPU you write what one thread does to one element, and then launch a thousand copies of it.

The launch looks like this:

add<<<4, 256>>>(d_a, d_b, d_c, n);

Those two numbers are the whole configuration. 4 is the number of blocks, 256 is the number of threads per block. Four blocks of 256 threads is 1024 threads — one for each of my 1000 elements, with 24 spare threads left over doing nothing. Every element gets an owner.

It ran. c[999] came back as 2997, which is correct (999 + 2×999). So far, so good.

One character

Then I changed one character. 4 became 1:

add<<<1, 256>>>(d_a, d_b, d_c, n);

One block of 256 threads is 256 threads. Everything else in the program was identical — same kernel, same data, same compiler, same GPU.

256 threads can only cover 256 elements. Threads 0 through 255 computed their positions, found their index was inside the array, and correctly wrote c[0] through c[255]. Elements 256 through 999 — 744 of them — had no thread assigned. Nobody wrote to that memory.

Here is the part that stayed with me. Those 744 slots weren't empty, and they weren't flagged. They held whatever bytes the GPU memory allocator left behind, and that happened to be zeros. So c[500] read back as 0.0 — a clean, ordinary-looking number. Not garbage, not NaN, nothing that shouts "wrong." And c[0] was 0.0 too, which is actually correct (0 + 0 = 0) — so the first element I'd instinctively spot-check looked completely fine.

A wrong answer that looks exactly like a right answer. If I'd eyeballed the output instead of checking all 1000, I'd have shipped it.

I added error checking, and it caught nothing

My first instinct was that I must be missing some error handling. Every CUDA tutorial teaches this: the calls return error codes, so you wrap them and check. I built the usual macro and put it on everything:

CHECK(cudaMalloc(&d_a, bytes));
CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));

And after the launch, the two checks the tutorials specifically tell you to add:

add<<<1, 256>>>(d_a, d_b, d_c, n);
CHECK(cudaGetLastError());        // did the launch fail?
CHECK(cudaDeviceSynchronize());   // did it fault while running?

I ran it. Every check passed. Same 744 wrong answers.

That was the moment the real lesson landed. Error checking detects failures — something goes wrong and reports it. But nothing went wrong here:

• cudaMalloc succeeded. Memory really was allocated.
• cudaMemcpy succeeded. Bytes really were copied.
• cudaGetLastError was clean. The launch was valid — <<<1, 256>>> is a perfectly legal configuration.
• cudaDeviceSynchronize was clean. No thread went out of bounds or hit an illegal instruction.
• Every one of the 256 threads that ran did correct arithmetic and stayed in bounds.

There was no failure to catch. The bug wasn't something going wrong — it was something not happening. 744 elements never got a thread, and nothing in CUDA is watching for work you never asked it to do. There is no "you under-launched" error, because the GPU has no idea how many elements I meant to process. I told it to run 256 threads. It ran 256 threads. Flawlessly.

What actually caught it

The only thing that found the bug was a loop that ignored the GPU entirely and compared the output against what I already knew the answer should be:

int errors = 0;
for (int i = 0; i < n; i++) {
    if (h_c[i] != h_a[i] + h_b[i]) errors++;
}
printf("mismatches: %d of %d\n", errors, n);

mismatches: 744 of 1000.

This doesn't ask the GPU whether anything went wrong. It asks a different question — not "did it error?" but "is the answer right?" And that gap, between "no errors reported" and "correct answer," is the whole story.

The takeaway

After this experience, I feel that verifying output against what you expect isn't a nice-to-have — it's a requirement. And it's not a new idea: comparing actual results against a known-correct answer is ordinary testing practice, the kind of thing testers do every day. What's different on a GPU is the temptation to skip it. GPUs are fast, checking a million results back on the CPU can feel like it defeats the point, so the check quietly gets dropped — which is exactly when a silent gap slips through.

I came to GPUs to learn how AI infrastructure is built. I left my first week with a testing lesson instead: on hardware this fast, the answer looking plausible is not the same as the answer being right, and the only thing that tells them apart is checking.