code/day66-testing/testing.cuThis is the source used by the lesson and its recorded evidence. Compile commands and expected output live in the directory README.
// SPDX-License-Identifier: MIT
//
// Day 66: testing GPU kernels.
//
// One reduction, two variants, two test suites. The drop-tail variant
// carries a deliberate one-expression bug in its grid arithmetic. The
// supplied suite (four tests, the kind that accumulate in real projects)
// passes it anyway: a golden snapshot captured from the kernel itself, a
// determinism check, an oracle behind a tolerance a hundred times too
// loose, and one honest property test that only runs at a size where the
// bug cannot show. The hardened suite (double-precision oracle on a size
// ladder, a permutation swap, tolerance scaled to the reduction length)
// exposes it. The program gates on exactly that split and exits nonzero
// if any side of it fails to happen.
//
// Nothing here is timed. The subject is evidence, not speed.
//
// starter/testing.cu mirrors this file's buggy variant plus the four
// supplied tests; the verification run diffs the two so they cannot
// drift apart silently.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o testing testing.cu
// Run: ./testing
//
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <cuda_runtime.h>
// The one error macro. This file is standalone, the way a Compiler Explorer
// embed is, so it carries its own verbatim copy. `err_` carries a trailing
// underscore so it cannot collide with a variable at the call site, and the
// do/while makes the macro one statement so it survives a braceless `if`.
#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)
// 16,483 is 169 whole cycles of the 97-value input pattern plus 90, and it
// is not a multiple of the block size, so the tail exists at the largest
// size. 611 is day 5's canonical odd size (13 x 47). 1,024 is the size that
// hides the bug: four whole blocks, no remainder, which is why the supplied
// suite's one honest test runs there and sees nothing.
constexpr size_t kMaxElems = 16483;
constexpr size_t kOddElems = 611;
constexpr size_t kExactElems = 1024;
constexpr int kThreadsPerBlock = 256; // 8 warps
constexpr int kMaxBlocks =
static_cast<int>((kMaxElems + kThreadsPerBlock - 1) / kThreadsPerBlock);
constexpr int kSwapCount = 99; // elements swapped between head and tail
constexpr float kBaseRtol = 1e-5f; // the f32 default, about 84 ulps
constexpr float kLooseRtol = 1e-2f; // the supplied suite's "room for floats"
static_assert(kThreadsPerBlock % 32 == 0,
"block size must be a whole number of warps");
static_assert((kThreadsPerBlock & (kThreadsPerBlock - 1)) == 0,
"sumBlocks halves the stride, so the block size must be a "
"power of two");
// Sums one block's slice of `in` into out[blockIdx.x]. Day 24's tree, the
// day 5 harness shapes around it, nothing new: this day is about the tests.
//
// One thread loads one element, then the block folds kThreadsPerBlock
// values in shared memory. A warp's 32 loads cover 128 contiguous bytes.
//
// Launch requirement: exactly kThreadsPerBlock threads per block. The guard
// covers the load, not the barrier, and there is no return above a barrier.
__global__ void sumBlocks(const float* __restrict__ in, float* __restrict__ out,
size_t n) {
__shared__ float tile[kThreadsPerBlock];
const unsigned int tid = threadIdx.x;
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + tid;
tile[tid] = (i < n) ? in[i] : 0.0f;
__syncthreads();
for (unsigned int half = kThreadsPerBlock / 2; half > 0; half /= 2) {
if (tid < half) {
tile[tid] += tile[tid + half];
}
__syncthreads();
}
if (tid == 0) {
out[blockIdx.x] = tile[0];
}
}
// Both variants launch the same kernel and fold its partials the same way.
// The only difference between them is the grid arithmetic, which is the
// point: the kernel under test is really the kernel plus its launch, and a
// suite that only ever exercises the kernel body cannot see this bug.
static float foldPartials(const float* d_in, size_t n, int blocks,
float* d_partials, float* h_partials) {
sumBlocks<<<blocks, kThreadsPerBlock>>>(d_in, d_partials, n);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_partials, d_partials,
static_cast<size_t>(blocks) * sizeof(float),
cudaMemcpyDeviceToHost));
float total = 0.0f; // float on purpose: the whole pipeline stays f32
for (int b = 0; b < blocks; ++b) {
total += h_partials[b];
}
return total;
}
// snippet: drop-tail
// Deliberate bug (day 66): integer division rounds down, so the last
// n % kThreadsPerBlock elements are never covered by any block. At 1,024
// the remainder is zero and the variant is exactly right; at 611 it
// silently ignores 99 elements. sumGuarded below differs by one
// expression. Below 256 elements this grid is zero blocks and the launch
// itself fails, which is why the size ladder here stops at 611.
static float sumDropTail(const float* d_in, size_t n, float* d_partials,
float* h_partials) {
const int blocks = static_cast<int>(n / kThreadsPerBlock);
return foldPartials(d_in, n, blocks, d_partials, h_partials);
}
// end snippet
// The fixed variant: day 5's ceiling division, so the last partial block
// exists and the kernel's own bounds guard does the rest.
static float sumGuarded(const float* d_in, size_t n, float* d_partials,
float* h_partials) {
const int blocks =
static_cast<int>((n + kThreadsPerBlock - 1) / kThreadsPerBlock);
return foldPartials(d_in, n, blocks, d_partials, h_partials);
}
using SumFn = float (*)(const float*, size_t, float*, float*);
// The oracle. Double accumulation with Kahan compensation: the slowest
// correct sum we can write, deliberately nothing like the kernel. An
// oracle that shares the kernel's order or its precision shares its bugs.
static double sumCpu(const float* x, size_t n) {
double total = 0.0;
double comp = 0.0;
for (size_t i = 0; i < n; ++i) {
const double y = static_cast<double>(x[i]) - comp;
const double t = total + y;
comp = (t - total) - y;
total = t;
}
return total;
}
// snippet: scaled-tolerance
// A sum of K float terms accumulates rounding error. For a tree or any
// shuffled order the per-step errors are uncorrelated, so the bound grows
// like sqrt(K), not K; the factor 4 is slack for a different but valid
// summation order. One fixed rtol is wrong at both ends: 1e-5 fails
// correct code at K in the millions, and a value loosened until every
// size passes stops being a test.
static float scaledRtol(size_t k) {
const float eps = 1.1920929e-7f; // 2^-23, float spacing at 1.0
const float grown = 4.0f * eps * std::sqrt(static_cast<float>(k));
return grown > kBaseRtol ? grown : kBaseRtol;
}
// end snippet
// One comparison, one printed line, one bool. Every test funnels through
// here so a failure always shows got, want, the relative error and the
// tolerance that judged it, with the arithmetic visible.
static bool checkRel(const char* name, double got, double want, double rtol) {
const double scale = (want == 0.0) ? 1.0 : std::fabs(want);
const double rel = std::fabs(got - want) / scale;
const bool ok = rel <= rtol;
std::printf(" %-34s %s got %.6f want %.6f rel %.3e rtol %.3e\n", name,
ok ? "PASS" : "FAIL", got, want, rel, rtol);
return ok;
}
// ---- the supplied suite: what the exercise starts from ----
// snippet: golden
// Tautological (test 1 of the supplied suite): this constant was captured
// from the kernel under test, so the test can only fail when the kernel
// changes. It does not certify the sum of 611 elements; it certifies
// whatever the drop-tail variant did the day someone ran it. 11,815.5 is
// the sum of elements 0..511 of this input, the two whole blocks the
// buggy grid covers. The true 611-element sum is 14,171.
constexpr float kGolden611 = 11815.5f;
static bool testGoldenSnapshot(SumFn sum, const float* d_in, float* d_partials,
float* h_partials) {
const float got = sum(d_in, kOddElems, d_partials, h_partials);
return checkRel("golden snapshot, n=611", got, kGolden611, 0.0);
}
// end snippet
// Tautological (test 2): the kernel against itself. A deterministic wrong
// answer passes this forever; day 27's racy kernel would usually pass it
// too. Determinism is worth one test, but it is not correctness.
static bool testDeterminism(SumFn sum, const float* d_in, float* d_partials,
float* h_partials) {
const float a = sum(d_in, kMaxElems, d_partials, h_partials);
const float b = sum(d_in, kMaxElems, d_partials, h_partials);
return checkRel("determinism, n=16483", a, b, 0.0);
}
// Under-toleranced (test 3): a real double-precision oracle, ruined by
// rtol = 1e-2 "to be safe with floats". The dropped tail at this size is
// about 6.1e-3 of the total, which sails under a one percent gate. The
// scaled rule for this K is 6.1e-5.
static bool testLooseOracle(SumFn sum, const float* d_in, const float* h_in,
float* d_partials, float* h_partials) {
const float got = sum(d_in, kMaxElems, d_partials, h_partials);
const double want = sumCpu(h_in, kMaxElems);
return checkRel("loose oracle, n=16483", got, want, kLooseRtol);
}
// Honest but blind (test 4): linearity, sum(2x) == 2 * sum(x), is a real
// property test. It runs only at n = 1024, where four whole blocks leave
// no tail, and dropping a tail is linear anyway, so it cannot see this
// bug at any size. A good test is still only a filter for its own class.
static bool testLinearityExact(SumFn sum, const float* d_in, const float* d_in2,
float* d_partials, float* h_partials) {
const float one = sum(d_in, kExactElems, d_partials, h_partials);
const float two = sum(d_in2, kExactElems, d_partials, h_partials);
return checkRel("linearity, n=1024", two, 2.0 * one, kBaseRtol);
}
// ---- the hardened suite: what the exercise builds ----
// Oracle on a size ladder with the scaled tolerance. 611 has a 99-element
// tail, 1,024 has none, 16,483 has a 99-element tail again at a size
// where it is small enough to test the tolerance rather than the reader.
static int runOracleLadder(SumFn sum, const float* d_in, const float* h_in,
float* d_partials, float* h_partials,
bool* fail611) {
const size_t sizes[3] = {kOddElems, kExactElems, kMaxElems};
const char* names[3] = {"oracle, n=611", "oracle, n=1024",
"oracle, n=16483"};
int failures = 0;
for (int s = 0; s < 3; ++s) {
const float got = sum(d_in, sizes[s], d_partials, h_partials);
const double want = sumCpu(h_in, sizes[s]);
const bool ok = checkRel(names[s], got, want, scaledRtol(sizes[s]));
if (!ok) {
++failures;
if (s == 0) {
*fail611 = true;
}
}
}
return failures;
}
// snippet: swap-ends
// Permutation invariance: a sum must not care where its elements sit.
// The permutation is a deliberate one, not a random shuffle: the first
// and last 99 elements trade places, so whatever a kernel does to the
// tail of the buffer, it now does it to different values. A random
// shuffle would also work but would make the failure a different number
// every seed, and a test you cannot reproduce exactly is a test you
// cannot debug.
static bool testSwapEnds(SumFn sum, const float* d_in, const float* d_swap,
float* d_partials, float* h_partials) {
const float plain = sum(d_in, kMaxElems, d_partials, h_partials);
const float moved = sum(d_swap, kMaxElems, d_partials, h_partials);
return checkRel("swap-ends permutation, n=16483", moved, plain,
scaledRtol(kMaxElems));
}
// end snippet
int main() {
const int device = 0;
CUDA_CHECK(cudaSetDevice(device));
cudaDeviceProp prop;
CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
std::printf("GPU: %s (compute capability %d.%d)\n", prop.name, prop.major,
prop.minor);
std::printf(
"n up to %zu, %d threads per block, partials folded on the "
"host in float\n\n",
kMaxElems, kThreadsPerBlock);
// Input: h_a[i] = (i % 97) * 0.5f. Every value is a multiple of 0.5
// and every sum stays far below 2^24, so all float arithmetic in this
// program is exact and a mismatch can only be an indexing or launch
// bug. That is the day 5 harness's int-exact doctrine: tolerance
// questions are real, but never on the case that hunts index bugs.
std::vector<float> h_a(kMaxElems);
std::vector<float> h_a2(kMaxElems);
std::vector<float> h_swap(kMaxElems);
for (size_t i = 0; i < kMaxElems; ++i) {
h_a[i] = static_cast<float>(i % 97) * 0.5f;
h_a2[i] = 2.0f * h_a[i];
}
for (size_t i = 0; i < kMaxElems; ++i) {
h_swap[i] = h_a[i];
}
for (size_t i = 0; i < kSwapCount; ++i) {
h_swap[i] = h_a[kMaxElems - kSwapCount + i];
h_swap[kMaxElems - kSwapCount + i] = h_a[i];
}
const size_t bytes = kMaxElems * sizeof(float);
std::vector<float> h_partials(kMaxBlocks);
float* d_a = nullptr;
float* d_a2 = nullptr;
float* d_swap = nullptr;
float* d_partials = nullptr;
CUDA_CHECK(cudaMalloc(&d_a, bytes));
CUDA_CHECK(cudaMalloc(&d_a2, bytes));
CUDA_CHECK(cudaMalloc(&d_swap, bytes));
CUDA_CHECK(cudaMalloc(&d_partials, kMaxBlocks * sizeof(float)));
CUDA_CHECK(cudaMemcpy(d_a, h_a.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_a2, h_a2.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(
cudaMemcpy(d_swap, h_swap.data(), bytes, cudaMemcpyHostToDevice));
// Phase 1: the supplied suite against the buggy variant. The gate is
// inverted on purpose: all four must PASS, because the lesson's claim
// is that this suite certifies a wrong kernel.
std::printf("supplied suite vs sumDropTail (expected: all four pass)\n");
int weakPassed = 0;
weakPassed +=
testGoldenSnapshot(sumDropTail, d_a, d_partials, h_partials.data());
weakPassed +=
testDeterminism(sumDropTail, d_a, d_partials, h_partials.data());
weakPassed += testLooseOracle(sumDropTail, d_a, h_a.data(), d_partials,
h_partials.data());
weakPassed += testLinearityExact(sumDropTail, d_a, d_a2, d_partials,
h_partials.data());
// Phase 2: the hardened suite against the buggy variant. Linearity is
// included and expected to pass both variants: the property is blind
// to a dropped tail, and printing that is part of the lesson.
std::printf("\nhardened suite vs sumDropTail (expected: bug exposed)\n");
bool buggyFail611 = false;
int buggyFailures =
runOracleLadder(sumDropTail, d_a, h_a.data(), d_partials,
h_partials.data(), &buggyFail611);
buggyFailures +=
testSwapEnds(sumDropTail, d_a, d_swap, d_partials, h_partials.data())
? 0
: 1;
buggyFailures += testLinearityExact(sumDropTail, d_a, d_a2, d_partials,
h_partials.data())
? 0
: 1;
// Phase 3: the hardened suite against the fixed variant. All pass.
std::printf("\nhardened suite vs sumGuarded (expected: all pass)\n");
bool unusedFail611 = false;
int fixedFailures = runOracleLadder(sumGuarded, d_a, h_a.data(), d_partials,
h_partials.data(), &unusedFail611);
fixedFailures +=
testSwapEnds(sumGuarded, d_a, d_swap, d_partials, h_partials.data())
? 0
: 1;
fixedFailures +=
testLinearityExact(sumGuarded, d_a, d_a2, d_partials, h_partials.data())
? 0
: 1;
// Phase 4: the golden snapshot against the fixed variant. It fails,
// which is the tautology running backwards: a test captured from a
// bug defends the bug against the fix.
std::printf("\ngolden snapshot vs sumGuarded (expected: fails)\n");
const bool goldenPassesFixed =
testGoldenSnapshot(sumGuarded, d_a, d_partials, h_partials.data());
CUDA_CHECK(cudaFree(d_a));
CUDA_CHECK(cudaFree(d_a2));
CUDA_CHECK(cudaFree(d_swap));
CUDA_CHECK(cudaFree(d_partials));
// The program's own gate, as real branches. Each names what the day
// claims; if the hardware disagrees with any of them, that is a
// finding for the lesson, not something to paper over.
int bad = 0;
if (weakPassed != 4) {
std::fprintf(stderr,
"GATE: supplied suite passed %d of 4 against the buggy "
"variant; the lesson claims all 4\n",
weakPassed);
++bad;
}
if (!buggyFail611) {
std::fprintf(stderr,
"GATE: hardened oracle at n=611 did not expose the "
"dropped tail\n");
++bad;
}
if (buggyFailures == 0) {
std::fprintf(stderr, "GATE: hardened suite passed the buggy variant\n");
++bad;
}
if (fixedFailures != 0) {
std::fprintf(stderr,
"GATE: hardened suite failed the fixed variant %d "
"time(s)\n",
fixedFailures);
++bad;
}
if (goldenPassesFixed) {
std::fprintf(stderr,
"GATE: the golden snapshot passed the fixed variant, "
"so it was not tautological after all\n");
++bad;
}
if (bad != 0) {
return EXIT_FAILURE;
}
std::printf(
"\nall gates hold: four green tests certified a wrong "
"kernel, and the hardened suite told them apart\n");
return EXIT_SUCCESS;
}