code/day09-timing/timing.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 9: one vector add, timed five ways.
//
// All five numbers are correct measurements. Four of them measure something
// other than the kernel, and those four are what gets quoted when someone
// reports that their GPU lost to their CPU:
//
// 1 a wall clock around the whole process
// 2 a wall clock around the launch, which is not the kernel
// 3 CUDA events around the first launch, which still includes a module load
// 4 CUDA events after a warm-up, which is the kernel
// 5 CUDA events around the copies and the kernel, which is the wait a user
// of your program actually experiences
//
// This is the one file in the course that reads a host clock on purpose.
// CUDA-CODE-STYLE.md bans host clocks and this lesson is the reason the ban
// exists, so the program has to produce the bad number before the page can
// take it apart. It calls clock_gettime rather than the C++ chrono clocks
// because the style rule names those specifically and a CI grep enforces it.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o timing timing.cu
// Run: ./timing
// time ./timing the number the forum posts quote
//
// Verified 2026-08-30 on a Tesla T4 (compute capability 7.5), driver
// 595.84, CUDA 12.6 (V12.6.85). Transcript: evidence/run-2026-08-30.txt
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <ctime>
#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 Mi elements plus 611, so the element count does not divide by any block
// size and the kernel's bounds check runs on every launch. 611 is day 5's
// number and it is here for the same reason.
//
// The size is large enough that the kernel is not pure launch overhead and
// small enough that three buffers fit on a 15 GiB T4 with room to spare.
constexpr size_t kElems = 16ull * 1024ull * 1024ull + 611ull;
constexpr int kThreadsPerBlock = 256; // 8 warps
constexpr int kWarmupRuns = 3;
constexpr int kTimedRuns = 10;
constexpr float kRelTolerance = 1e-5f;
// The sizes the crossover sweep walks. Each one is a slice of the buffers
// allocated for kElems, so no row reallocates and no row pays an allocation
// the other rows did not.
constexpr size_t kSweep[] = {1024ull, 64ull * 1024ull, 1024ull * 1024ull,
4ull * 1024ull * 1024ull};
constexpr int kNumSweep = 4;
// Blocks needed to cover n elements, by integer ceiling division.
//
// constexpr because the static_assert below calls it. A run-time check of
// something the compiler already knows is not a measurement, and an assert()
// would be deleted outright: CI builds Release, Release defines NDEBUG, and
// assert() under NDEBUG expands to nothing.
constexpr int blocksFor(size_t n) {
return static_cast<int>((n + kThreadsPerBlock - 1) / kThreadsPerBlock);
}
static_assert(kThreadsPerBlock % 32 == 0,
"block size must be a whole number of warps");
static_assert(kElems % kThreadsPerBlock != 0,
"n must not divide evenly, or the bounds check never runs");
static_assert(kSweep[kNumSweep - 1] <= kElems,
"the sweep reuses the buffers allocated for kElems");
static_assert(static_cast<size_t>(blocksFor(kElems)) * kThreadsPerBlock >=
kElems,
"the grid must cover every element");
// The only host clock in the course, and this lesson is about what it does
// not measure. CLOCK_MONOTONIC does not jump when the system clock is
// adjusted, which a stopwatch must not do.
//
// A C++ steady_clock would give exactly the same wrong answers below. The
// choice here is about the site's own grep, not about the clock.
static double hostMs() {
struct timespec ts;
clock_gettime(CLOCK_MONOTONIC, &ts);
return static_cast<double>(ts.tv_sec) * 1.0e3 +
static_cast<double>(ts.tv_nsec) * 1.0e-6;
}
// out[i] = a[i] + b[i]. One thread owns one element.
//
// Memory: consecutive threads take consecutive elements, so one warp's 32
// addresses cover 128 contiguous bytes. Three floats move per element, two
// read and one written, for a single add. That ratio is the reason no amount
// of tuning makes this kernel fast: it is waiting on memory, not on the
// adders. Day 11 measures what the access pattern is worth.
//
// Launch assumption: gridDim.x * blockDim.x >= n. The guard is a plain `if`
// and not an early return, because an early return hangs a kernel the moment
// a __syncthreads() appears below it, which happens on day 13.
__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];
}
}
// out[i] = a[i] * 2. What it computes does not matter.
//
// Memory: the same contiguous pattern as vectorAdd, two floats per element
// instead of three.
//
// It exists to be cold. It is a second piece of device code the driver has
// not loaded, so its first launch pays its own load even though the process
// is long past initialisation and vectorAdd is already warm. That is what
// makes "warm up the first kernel" the wrong rule and "warm up every kernel
// you time" the right one.
//
// Launch assumption: the same as vectorAdd.
__global__ void vectorScale(const float* a, 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] * 2.0f;
}
}
// CPU reference. Written for obvious correctness, not speed: plain loop, no
// OpenMP, no intrinsics, one thread. It never allocates; the caller owns
// every buffer.
//
// It is therefore a weak baseline, and the page says so. A real CPU version
// of this loop would use every core and every vector lane. The comparison
// below is already unflattering to the GPU with the CPU handicapped.
static void vectorAddCpu(const float* a, const float* b, float* out, size_t n) {
for (size_t i = 0; i < n; ++i) {
out[i] = a[i] + b[i];
}
}
// Returns the first index where got and want differ by more than the relative
// tolerance, or n if they agree everywhere. Returning the index rather than a
// bool is the point: "wrong at 512" names the block, "wrong" does not.
static size_t firstMismatch(const float* got, const float* want, size_t n,
float relTolerance) {
for (size_t i = 0; i < n; ++i) {
const float scale = (want[i] == 0.0f) ? 1.0f : std::fabs(want[i]);
if (std::fabs(got[i] - want[i]) > relTolerance * scale) {
return i;
}
}
return n;
}
// Times a launch with CUDA events and returns the mean milliseconds per run.
//
// This is the one template and the one lambda allowed in module 1 to 3 code.
// Copy it verbatim; the alternative is six copies of the event boilerplate,
// which is how a warm-up goes missing from one of them.
// snippet: time-kernel
template <typename LaunchFn>
static float timeKernel(LaunchFn launch) {
cudaEvent_t start, stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));
// Warm up this kernel, not just the first kernel in the program. Lazy
// module loading has been the default since CUDA 12.2 on Linux, so the
// first launch of each kernel pays its own load.
for (int i = 0; i < kWarmupRuns; ++i) {
launch();
}
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaEventRecord(start));
for (int i = 0; i < kTimedRuns; ++i) {
launch();
}
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
CUDA_CHECK(cudaGetLastError());
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
CUDA_CHECK(cudaEventDestroy(start));
CUDA_CHECK(cudaEventDestroy(stop));
return ms / kTimedRuns;
}
// end snippet
// Times one full round trip: both inputs up, one launch, the result back.
// Warms up first, then returns the mean of kTimedRuns.
//
// What it measures: the wait a user of your program experiences. What it does
// not measure: the kernel. Everywhere else in this course the copies stay
// outside the timed region, because there the question is how fast the kernel
// is. Here the question is how long the whole thing takes, and the answer
// includes the bus.
static float timeRoundTrip(const float* h_a, const float* h_b, float* h_out,
float* d_a, float* d_b, float* d_out, size_t n) {
const size_t bytes = n * sizeof(float);
const int blocks = blocksFor(n);
cudaEvent_t start;
cudaEvent_t stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));
for (int i = 0; i < kWarmupRuns + kTimedRuns; ++i) {
if (i == kWarmupRuns) {
CUDA_CHECK(cudaEventRecord(start));
}
CUDA_CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));
vectorAdd<<<blocks, kThreadsPerBlock>>>(d_a, d_b, d_out, n);
CUDA_CHECK(cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost));
}
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
CUDA_CHECK(cudaGetLastError());
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
CUDA_CHECK(cudaEventDestroy(start));
CUDA_CHECK(cudaEventDestroy(stop));
return ms / kTimedRuns;
}
int main() {
// Measurement 1 starts here. It cannot start earlier: process creation,
// the dynamic linker and libcudart's own load all happen before main, and
// `time ./timing` is the only way to see them. The README says to run it.
const double tProcessStart = hostMs();
// cudaSetDevice is the first CUDA call in this process, and as of CUDA
// 12.0 it "initialize[s] the runtime and the primary context associated
// with the specified device". Timing it separately is the only way to see
// a cost that otherwise hides inside whatever CUDA call you happen to
// make first.
const double tCtx0 = hostMs();
CUDA_CHECK(cudaSetDevice(0));
const double contextMs = hostMs() - tCtx0;
cudaDeviceProp prop;
CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
const size_t bytes = kElems * sizeof(float);
const int blocks = blocksFor(kElems);
std::vector<float> h_a(kElems);
std::vector<float> h_b(kElems);
std::vector<float> h_out(kElems);
std::vector<float> h_want(kElems);
for (size_t i = 0; i < kElems; ++i) {
h_a[i] = static_cast<float>(i % 97) * 0.5f;
h_b[i] = static_cast<float>(i % 13);
}
float* d_a = nullptr;
float* d_b = nullptr;
float* d_out = nullptr;
CUDA_CHECK(cudaMalloc(&d_a, bytes));
CUDA_CHECK(cudaMalloc(&d_b, bytes));
CUDA_CHECK(cudaMalloc(&d_out, bytes));
CUDA_CHECK(cudaMemcpy(d_a, h_a.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_b, h_b.data(), bytes, cudaMemcpyHostToDevice));
cudaEvent_t start;
cudaEvent_t stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));
// Measurement 3, and it has to come first. This is the first launch of
// vectorAdd in the process, so the driver loads its module here. Anything
// above this line that launched vectorAdd would make the number a lie.
//
// The events do not time the load itself, which happens on the host. They
// time the gap the load leaves in the device timeline: the start event
// completes, the GPU then sits idle while the host loads the module, and
// only then does the kernel run.
// snippet: cold-launch
CUDA_CHECK(cudaEventRecord(start));
vectorAdd<<<blocks, kThreadsPerBlock>>>(d_a, d_b, d_out, kElems);
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
CUDA_CHECK(cudaGetLastError());
float coldMs = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&coldMs, start, stop));
// end snippet
// Measurement 2. The module is loaded by now, so this isolates one thing:
// what a host clock around a launch statement actually covers. The launch
// returns to the host as soon as the work is queued, so the first number
// is the cost of queueing. The second adds cudaDeviceSynchronize() and so
// waits for the kernel, which is what the beginner meant to measure.
// snippet: launch-clock
const double tLaunch0 = hostMs();
vectorAdd<<<blocks, kThreadsPerBlock>>>(d_a, d_b, d_out, kElems);
const double tLaunch1 = hostMs();
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
const double launchOnlyMs = tLaunch1 - tLaunch0;
const double launchPlusSyncMs = hostMs() - tLaunch0;
// end snippet
// Measurement 4. The honest kernel time: warmed up, mean of ten, copies
// outside the timed region.
const float warmMs = timeKernel([&] {
vectorAdd<<<blocks, kThreadsPerBlock>>>(d_a, d_b, d_out, kElems);
});
// Measurement 5. The same kernel with the copies put back.
const float roundTripMs = timeRoundTrip(
h_a.data(), h_b.data(), h_out.data(), d_a, d_b, d_out, kElems);
// The CPU, on the same data, with the same host clock. This one is a fair
// use of a host clock: nothing here is asynchronous. The same pass fills
// h_want, so the check below costs nothing extra.
const double tCpu0 = hostMs();
vectorAddCpu(h_a.data(), h_b.data(), h_want.data(), kElems);
const double cpuMs = hostMs() - tCpu0;
// Checked here, while h_out still holds the whole round trip at kElems.
// The sweep below overwrites the front of both host buffers.
const size_t bad =
firstMismatch(h_out.data(), h_want.data(), kElems, kRelTolerance);
// The second kernel, cold and then warm. vectorAdd has run dozens of
// times by now and the process is well past initialisation, so if the
// load cost were per process this pair would be equal.
CUDA_CHECK(cudaEventRecord(start));
vectorScale<<<blocks, kThreadsPerBlock>>>(d_a, d_out, kElems);
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
CUDA_CHECK(cudaGetLastError());
float scaleColdMs = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&scaleColdMs, start, stop));
const float scaleWarmMs = timeKernel(
[&] { vectorScale<<<blocks, kThreadsPerBlock>>>(d_a, d_out, kElems); });
// The crossover sweep. Three numbers per size: the CPU, the kernel alone,
// and the round trip. Every row reuses the buffers above.
double sweepCpuMs[kNumSweep];
float sweepKernelMs[kNumSweep];
float sweepTripMs[kNumSweep];
for (int s = 0; s < kNumSweep; ++s) {
const size_t n = kSweep[s];
const int sweepBlocks = blocksFor(n);
const double t0 = hostMs();
vectorAddCpu(h_a.data(), h_b.data(), h_want.data(), n);
sweepCpuMs[s] = hostMs() - t0;
sweepKernelMs[s] = timeKernel([&] {
vectorAdd<<<sweepBlocks, kThreadsPerBlock>>>(d_a, d_b, d_out, n);
});
sweepTripMs[s] = timeRoundTrip(h_a.data(), h_b.data(), h_out.data(),
d_a, d_b, d_out, n);
}
CUDA_CHECK(cudaEventDestroy(start));
CUDA_CHECK(cudaEventDestroy(stop));
CUDA_CHECK(cudaFree(d_a));
CUDA_CHECK(cudaFree(d_b));
CUDA_CHECK(cudaFree(d_out));
if (bad != kElems) {
std::fprintf(stderr, "vectorAdd wrong at %zu: got %.9g, want %.9g\n",
bad, h_out[bad], h_want[bad]);
return EXIT_FAILURE;
}
// Measurement 1 ends here, so it covers everything above and none of the
// printing below.
const double wholeMs = hostMs() - tProcessStart;
std::printf("GPU: %s (compute capability %d.%d)\n", prop.name, prop.major,
prop.minor);
std::printf(
"n = %zu floats, %.1f MiB per buffer, %d threads/block, %d "
"blocks\n\n",
kElems, static_cast<double>(bytes) / (1024.0 * 1024.0),
kThreadsPerBlock, blocks);
std::printf(
"cudaSetDevice(0), the first CUDA call in this process: "
"%9.3f ms\n",
contextMs);
std::printf("No arithmetic of yours has run yet. That is the context.\n\n");
std::printf("Five ways to time one vector add\n");
std::printf(" 1 %-40s %10.3f ms\n", "wall clock, whole process", wholeMs);
std::printf(" 2 %-40s %10.3f ms\n", "wall clock, around the launch",
launchOnlyMs);
std::printf(" %-40s %10.3f ms\n", "the same clock, plus a sync",
launchPlusSyncMs);
std::printf(" 3 %-40s %10.3f ms\n", "events, kernel, first launch",
static_cast<double>(coldMs));
std::printf(" 4 %-40s %10.3f ms\n", "events, kernel, after warm-up",
static_cast<double>(warmMs));
std::printf(" 5 %-40s %10.3f ms\n", "events, copies + kernel + copy back",
static_cast<double>(roundTripMs));
std::printf(" %-40s %10.3f ms\n\n", "CPU, same n, one thread, no SIMD",
cpuMs);
std::printf("Rows 4 and 5 are means of %d runs after %d warm-ups.\n",
kTimedRuns, kWarmupRuns);
std::printf(
"Rows 1, 2 and 3 are single measurements, because each one is "
"about\nsomething that happens once.\n\n");
std::printf("A second kernel is cold too\n");
std::printf(" %-40s %10.3f ms\n", "vectorScale, first launch",
static_cast<double>(scaleColdMs));
std::printf(" %-40s %10.3f ms\n\n", "vectorScale, after warm-up",
static_cast<double>(scaleWarmMs));
std::printf("Where the CPU wins\n");
std::printf("%12s %12s %12s %14s\n", "n", "CPU (ms)", "kernel (ms)",
"round trip (ms)");
for (int s = 0; s < kNumSweep; ++s) {
std::printf("%12zu %12.3f %12.3f %14.3f\n", kSweep[s], sweepCpuMs[s],
static_cast<double>(sweepKernelMs[s]),
static_cast<double>(sweepTripMs[s]));
}
std::printf(
"\nRow 4 is the kernel. Row 5 is what a user waits for. Row 1 "
"is what\n`time ./timing` prints, and it is never the answer "
"to how fast a\nkernel is.\n");
// Two consistency checks, and they are real branches returning
// EXIT_FAILURE rather than assert()s. CI builds Release, Release defines
// NDEBUG, and an assert() under NDEBUG is deleted, so a check written
// that way would vanish in exactly the build that matters.
//
// Neither of these is a prediction about the hardware. Both are
// arithmetic: measurement 5 contains measurement 4, and measurement 1
// contains both. If either fires, the timing is broken and every number
// on the page above it is worthless.
if (roundTripMs < warmMs) {
std::fprintf(stderr,
"timing broken: copies + kernel (%.3f ms) came out "
"faster than the kernel alone (%.3f ms)\n",
static_cast<double>(roundTripMs),
static_cast<double>(warmMs));
return EXIT_FAILURE;
}
if (wholeMs < static_cast<double>(roundTripMs)) {
std::fprintf(stderr,
"timing broken: the whole process (%.3f ms) came out "
"shorter than one round trip (%.3f ms)\n",
wholeMs, static_cast<double>(roundTripMs));
return EXIT_FAILURE;
}
return EXIT_SUCCESS;
}