code/day66-testing/ci_gate.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, part 2: why a sanitizer step needs --error-exitcode.
//
// Two kernels. squareInBounds is correct. squarePastEnd also produces a
// correct visible answer, and on top of it writes 4 bytes one element past
// the end of the output allocation. The program checks only the bytes it
// allocated, so it prints PASS twice and exits 0: from CI's point of view
// this binary is green.
//
// compute-sanitizer sees the out-of-bounds write. But its documented
// default is to exit with the application's own code even when it found
// errors ("--error-exitcode", default 0), so a CI step that runs the tool
// without that flag stays green too. The three commands in the README run
// this binary bare, under the tool, and under the tool with
// --error-exitcode 1, and only the third one goes red.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o ci_gate ci_gate.cu
// Run: ./ci_gate
//
#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)
// 611 elements is 2,444 bytes, and cudaMalloc rounds an allocation up, so
// the one deliberate 4-byte overrun below lands in memory the driver
// already owns and is expected not to fault (day 6 measured exactly this
// class of overrun staying silent). If this card faults on it anyway, the
// bare run exits nonzero and the lesson's prediction dies; that would be a
// result to record, not to hide.
constexpr size_t kElems = 611;
constexpr int kThreadsPerBlock = 256; // 8 warps
static_assert(kThreadsPerBlock % 32 == 0,
"block size must be a whole number of warps");
// out[i] = in[i] * in[i]. One thread owns one element. Consecutive threads
// take consecutive elements, so a warp's loads cover 128 contiguous bytes.
// Launch assumption: gridDim.x * blockDim.x >= n.
__global__ void squareInBounds(const float* __restrict__ in,
float* __restrict__ out, size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
out[i] = in[i] * in[i];
}
}
// Same kernel, plus a deliberate bug (day 66): thread 0 also writes the
// element at out[n], 4 bytes past the allocation. Everything a correctness
// test can read is still right, which is the point: the only witness to
// this write is the sanitizer, and CI only hears the sanitizer through its
// exit code.
__global__ void squarePastEnd(const float* __restrict__ in,
float* __restrict__ out, size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
out[i] = in[i] * in[i];
}
if (i == 0) {
out[n] = in[0] * in[0]; // deliberate (day 66): 4 bytes out of bounds
}
}
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);
const size_t bytes = kElems * sizeof(float);
const int blocks =
static_cast<int>((kElems + kThreadsPerBlock - 1) / kThreadsPerBlock);
// Small whole numbers, so every square and every comparison is exact.
std::vector<float> h_in(kElems);
std::vector<float> h_out(kElems);
for (size_t i = 0; i < kElems; ++i) {
h_in[i] = static_cast<float>(i % 31);
}
float* d_in = nullptr;
float* d_out = nullptr;
CUDA_CHECK(cudaMalloc(&d_in, bytes));
CUDA_CHECK(cudaMalloc(&d_out, bytes));
CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), bytes, cudaMemcpyHostToDevice));
// Checks only the kElems elements the program allocated, which is all
// any host-side check can do: the overrun below is invisible from
// here. Returns the first bad index or kElems.
int wrongKernels = 0;
for (int pass = 0; pass < 2; ++pass) {
// The in-bounds kernel runs first, so if the overrun does fault,
// the clean kernel's verdict is already printed.
if (pass == 0) {
squareInBounds<<<blocks, kThreadsPerBlock>>>(d_in, d_out, kElems);
} else {
squarePastEnd<<<blocks, kThreadsPerBlock>>>(d_in, d_out, kElems);
}
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
size_t bad = kElems;
for (size_t i = 0; i < kElems; ++i) {
const float want = h_in[i] * h_in[i];
if (h_out[i] != want) {
bad = i;
break;
}
}
const char* name = (pass == 0) ? "squareInBounds" : "squarePastEnd";
if (bad != kElems) {
std::printf("%-16s FAIL at index %zu\n", name, bad);
++wrongKernels;
} else {
std::printf("%-16s PASS: all %zu visible elements correct\n", name,
kElems);
}
}
CUDA_CHECK(cudaFree(d_in));
CUDA_CHECK(cudaFree(d_out));
if (wrongKernels != 0) {
return EXIT_FAILURE;
}
std::printf(
"exit 0: this binary looks green to CI. Only the sanitizer "
"knows better,\nand only its exit code can say so.\n");
return EXIT_SUCCESS;
}