code/day21-warps/warps.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 21: what a warp is, read off the hardware instead of asserted.
//
// One kernel, launched once per block size as a single block. Every thread
// samples clock64() and __activemask() at the top of the kernel and writes
// both, plus the run-time warpSize, into its own slot. The host reads the
// schedule back afterwards. Day 64 is why none of this is printed from the
// device: printf order is not execution order.
//
// Nothing here is timed. clock64() returns "the value of a per-multiprocessor
// counter that is incremented every clock cycle", which makes it an ordering
// probe and not a stopwatch, so this file contains no cudaEvent and the output
// carries no milliseconds. Two consequences the program is built around. A
// reading is only comparable with another reading from the same SM, which is
// why every launch is a single block. And the absolute value means nothing, so
// every cycle count printed below is relative to the smallest sample in its
// own launch.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o warps warps.cu
// Run: ./warps
//
// Verified: run on a Tesla T4 (compute capability 7.5), driver 595.84,
// CUDA 12.6 (V12.6.85), on 2026-08-30. The only output that may be
// published as this program's output is the transcript in
// evidence/run-2026-08-30.txt. See research/REVIEW-PROCESS.md.
#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)
// snippet: warp-size
// 32 on every GPU this course targets, and the number the whole page is built
// on. The built-in `warpSize` is documented as "A run-time value defined as
// the number of threads in a warp, commonly 32", so it cannot size an array or
// appear in a static_assert. This constant can do both, and the program checks
// the two against each other at run time rather than assuming they agree.
constexpr int kWarpSize = 32;
// end snippet: warp-size
// The arrays are sized from this, not from the largest row of the sweep, so
// the exercise can add a 1023-thread block without touching an allocation.
// 1024 is the per-block ceiling at every compute capability in this course.
constexpr int kMaxBlockThreads = 1024;
// 32 is one full warp and is the baseline row. 33 and 48 are the point of the
// sweep: neither is a multiple of the warp size, so each ends in a warp that
// is partly empty and whose active mask says so. 64 is the same thread count
// as 33 rounded up to whole warps, and 256 is the course default.
constexpr int kBlockSizes[] = {32, 33, 48, 64, 256};
constexpr int kBlockSizeCount =
static_cast<int>(sizeof(kBlockSizes) / sizeof(kBlockSizes[0]));
// The one launch whose per-thread rows are printed in full, because a table
// that already grouped the threads by warp cannot be used to check the
// grouping. 48 threads is small enough to read and is the partial-warp case.
constexpr int kDumpBlockSize = 48;
// Warps a block of t threads occupies. This is the guide's own formula,
// ceil(T / Wsize), and it rounds up: a block of 48 threads occupies two warps
// and the second one carries 16 lanes that were never launched.
constexpr int warpsInBlock(int t) {
return (t + kWarpSize - 1) / kWarpSize;
}
// Threads that actually exist in warp `w` of a block of `t` threads. Every
// warp but the last holds 32; the last holds the remainder.
constexpr int lanesInWarp(int t, int w) {
return (t - w * kWarpSize < kWarpSize) ? (t - w * kWarpSize) : kWarpSize;
}
// constexpr so the static_asserts below can call them. Each follows only from
// the constants in this file and never consults the device, which is what
// keeps them compile-time claims rather than measurements wearing a disguise.
constexpr bool everyBlockSizeFits() {
for (int b = 0; b < kBlockSizeCount; ++b) {
if (kBlockSizes[b] < 1 || kBlockSizes[b] > kMaxBlockThreads) {
return false;
}
}
return true;
}
constexpr bool someBlockSizeIsPartial() {
for (int b = 0; b < kBlockSizeCount; ++b) {
if (kBlockSizes[b] % kWarpSize != 0) {
return true;
}
}
return false;
}
constexpr bool dumpSizeIsInTheSweep() {
for (int b = 0; b < kBlockSizeCount; ++b) {
if (kBlockSizes[b] == kDumpBlockSize) {
return true;
}
}
return false;
}
static_assert(kMaxBlockThreads <= 1024,
"1024 threads per block is the hardware ceiling at every "
"compute capability this course targets");
static_assert(everyBlockSizeFits(),
"every block size must be launchable and must fit the buffers, "
"which are sized from kMaxBlockThreads");
static_assert(someBlockSizeIsPartial(),
"at least one block size must not be a whole number of warps, "
"or the partial-warp rows this lesson turns on never run");
static_assert(dumpSizeIsInTheSweep(),
"the per-thread dump has to name a launch that happens");
// Samples the SM's cycle counter and the warp's active mask, then writes both
// plus the run-time warpSize into this thread's own slot.
//
// Memory: one warp's 32 stores to each array land in 32 consecutive slots, so
// every store is coalesced. That is not what the kernel is for, but a
// scattered write would put store latency inside the thing being sampled.
//
// Launch assumption: one block of at most kMaxBlockThreads threads, and three
// arrays of kMaxBlockThreads entries. There is no bounds guard, and that is
// deliberate rather than sloppy: a branch above the sample is part of the
// measurement, because __activemask() "only provides an instantaneous snapshot
// of the active threads within a warp". The check that would have been a guard
// is a static_assert instead, which costs nothing at run time and cannot end
// up between kernel entry and the sample.
// snippet: probe
__global__ void recordWarpSchedule(long long* clocks, unsigned int* masks,
int* warpSizes) {
const unsigned int tid = threadIdx.x;
const long long cycle = clock64();
const unsigned int active = __activemask();
clocks[tid] = cycle;
masks[tid] = active;
warpSizes[tid] = warpSize;
}
// end snippet: probe
// How many different values appear in the first n entries. Quadratic, over at
// most 1024 entries, on the host, once per launch. Sorting would be faster and
// would need a comparator the reader has not met yet.
static int countDistinct(const long long* values, int n) {
int distinct = 0;
for (int i = 0; i < n; ++i) {
bool seen = false;
for (int j = 0; j < i; ++j) {
if (values[j] == values[i]) {
seen = true;
break;
}
}
if (!seen) {
++distinct;
}
}
return distinct;
}
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("warpSize: %d from the driver, %d compiled into this file\n",
prop.warpSize, kWarpSize);
std::printf("maxThreadsPerBlock: %d\n", prop.maxThreadsPerBlock);
std::printf(
"one block per launch, so every clock64() reading below comes from "
"one SM\n");
const size_t slots = static_cast<size_t>(kMaxBlockThreads);
const size_t clockBytes = slots * sizeof(long long);
const size_t maskBytes = slots * sizeof(unsigned int);
const size_t warpSizeBytes = slots * sizeof(int);
std::printf("%zu slots per array, %zu bytes of device memory in total\n",
slots, clockBytes + maskBytes + warpSizeBytes);
std::vector<long long> h_clocks(slots);
std::vector<unsigned int> h_masks(slots);
std::vector<int> h_warpSizes(slots);
long long* d_clocks = nullptr;
unsigned int* d_masks = nullptr;
int* d_warpSizes = nullptr;
CUDA_CHECK(cudaMalloc(&d_clocks, clockBytes));
CUDA_CHECK(cudaMalloc(&d_masks, maskBytes));
CUDA_CHECK(cudaMalloc(&d_warpSizes, warpSizeBytes));
// One status variable and one exit. Every path below falls through to the
// frees at the bottom, including the failing ones, because a lesson that
// returns early past its own cudaFree teaches the opposite of day 5.
int status = EXIT_SUCCESS;
for (int b = 0; b < kBlockSizeCount; ++b) {
const int blockSize = kBlockSizes[b];
const int warps = warpsInBlock(blockSize);
// Zero first, so a slot the launch never covered is distinguishable
// from one it did. An active mask of 0 cannot come from a running
// thread: the caller is active by definition, so its own bit is set.
CUDA_CHECK(cudaMemset(d_clocks, 0, clockBytes));
CUDA_CHECK(cudaMemset(d_masks, 0, maskBytes));
CUDA_CHECK(cudaMemset(d_warpSizes, 0, warpSizeBytes));
recordWarpSchedule<<<1, blockSize>>>(d_clocks, d_masks, d_warpSizes);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_clocks.data(), d_clocks, clockBytes,
cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(h_masks.data(), d_masks, maskBytes,
cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(h_warpSizes.data(), d_warpSizes, warpSizeBytes,
cudaMemcpyDeviceToHost));
// The three gates. All of them are real branches setting `status`,
// never asserts: CI builds Release, Release defines NDEBUG, and NDEBUG
// deletes assert() out of the build that matters, so a gate written
// that way would let a wrong table reach the page with CI green.
//
// Gate 1: the warp size the kernel read equals the one the driver
// reports and the one this file compiled in. warpSize is documented as
// a run-time value, so this is a reading rather than a promise, and if
// a card ever answers anything else then every count below is wrong.
//
// Gate 2: every thread's own lane bit is set in the mask it recorded.
// The guide defines the Nth bit as set "if the Nth lane in the warp is
// active", and the calling thread is active. A zero here means the
// readback is not what the kernel wrote.
for (int t = 0; t < blockSize; ++t) {
const unsigned int lane = static_cast<unsigned int>(t % kWarpSize);
if (h_warpSizes[t] != kWarpSize ||
h_warpSizes[t] != prop.warpSize) {
std::fprintf(stderr,
"thread %d of the %d-thread block read warpSize "
"%d; the driver says %d and this file compiled "
"in %d\n",
t, blockSize, h_warpSizes[t], prop.warpSize,
kWarpSize);
status = EXIT_FAILURE;
break;
}
if (((h_masks[t] >> lane) & 1u) == 0u) {
std::fprintf(stderr,
"thread %d of the %d-thread block recorded mask "
"0x%08x, which does not include its own lane %u\n",
t, blockSize, h_masks[t], lane);
status = EXIT_FAILURE;
break;
}
}
// Gate 3: nothing outside the launch was written. This is the one that
// catches an indexing edit, which is the failure the exercise invites.
for (int t = blockSize; t < kMaxBlockThreads; ++t) {
if (h_masks[t] != 0u || h_clocks[t] != 0 || h_warpSizes[t] != 0) {
std::fprintf(stderr,
"the %d-thread launch wrote slot %d, which is "
"outside the block\n",
blockSize, t);
status = EXIT_FAILURE;
break;
}
}
long long blockMin = h_clocks[0];
for (int t = 1; t < blockSize; ++t) {
if (h_clocks[t] < blockMin) {
blockMin = h_clocks[t];
}
}
const int distinct = countDistinct(h_clocks.data(), blockSize);
std::printf(
"\n%d threads, 1 block: %d warp(s), %d distinct clock64 "
"value(s)\n",
blockSize, warps, distinct);
std::printf("%6s %6s %12s %9s %9s %10s\n", "warp", "lanes",
"active mask", "cycles", "one mask", "one clock");
for (int w = 0; w < warps; ++w) {
const int first = w * kWarpSize;
const int lanes = lanesInWarp(blockSize, w);
const unsigned int mask = h_masks[first];
long long warpMin = h_clocks[first];
bool oneMask = true;
bool oneClock = true;
for (int t = first; t < first + lanes; ++t) {
if (h_masks[t] != mask) {
oneMask = false;
}
if (h_clocks[t] != h_clocks[first]) {
oneClock = false;
}
if (h_clocks[t] < warpMin) {
warpMin = h_clocks[t];
}
}
std::printf("%6d %6d 0x%08x %9lld %9s %10s\n", w, lanes, mask,
warpMin - blockMin, oneMask ? "yes" : "no",
oneClock ? "yes" : "no");
}
// The raw rows, once. The table above already grouped the threads by
// warp, so it cannot be used to check that the grouping is real.
if (blockSize == kDumpBlockSize) {
std::printf("\n every thread of the %d-thread block, ungrouped\n",
kDumpBlockSize);
std::printf(" %6s %6s %6s %12s %9s\n", "thread", "warp", "lane",
"active mask", "cycles");
for (int t = 0; t < blockSize; ++t) {
std::printf(" %6d %6d %6d 0x%08x %9lld\n", t, t / kWarpSize,
t % kWarpSize, h_masks[t], h_clocks[t] - blockMin);
}
}
}
CUDA_CHECK(cudaFree(d_clocks));
CUDA_CHECK(cudaFree(d_masks));
CUDA_CHECK(cudaFree(d_warpSizes));
if (status != EXIT_SUCCESS) {
std::fprintf(stderr,
"\na gate failed; the tables above are not "
"trustworthy\n");
return status;
}
std::printf(
"\nevery gate passed: the warp size agrees three ways, every thread "
"is in its own mask, and no launch wrote outside its block\n");
return EXIT_SUCCESS;
}