code/day02-gpu-vs-cpu/warp_slots.cuThis is the source used by the lesson and its recorded evidence. Compile commands and expected output live in the directory README.
// Day 2: what one SM does with your launch configuration.
//
// The two questions this day asks are questions about hardware state the
// source code never shows you: what 8 blocks of 16 threads costs against
// 4 blocks of 32, and whether one block can be split across SMs. This program
// answers the first from your own card.
//
// It prints the per-SM limits the driver reports, then, for a list of block
// sizes, how many warp slots each one takes and how many blocks of it a
// single SM can hold at once. The residency column comes from
// cudaOccupancyMaxActiveBlocksPerMultiprocessor, which answers about the
// compiled kernel, so it accounts for a register count the source does not
// show. The rest is arithmetic the page also does in words.
//
// What it does not do: launch a kernel, time anything, or say which
// configuration is faster. Day 10 sweeps block sizes and measures the curve;
// day 45 shows two configurations with different occupancy and equal speed.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o warp_slots warp_slots.cu
// Run: ./warp_slots
//
// 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 <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>
#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)
// 32 on every GPU this course targets. warpSize is a run-time built-in, so it
// cannot size an array or appear in a static_assert; the compile-time copy
// lives here and main() checks the two agree on the card you ran on.
constexpr int kWarpSize = 32;
// The block sizes the table walks. 16 and 32 are the two halves of the
// question this lesson opens with. 48 is the size that leaves 16 lanes of a
// warp idle for the block's whole lifetime. 1024 is the ceiling on every
// compute capability the course covers.
constexpr int kBlockSizes[] = {16, 32, 48, 64, 128, 256, 512, 1024};
constexpr int kNumBlockSizes =
static_cast<int>(sizeof(kBlockSizes) / sizeof(kBlockSizes[0]));
// The 128 threads of the opening question, launched two ways.
constexpr int kQuizThreads = 128;
constexpr int kSmallBlock = 16;
// out[i] = a * x[i] + y[i]. One thread owns one element.
//
// Memory: consecutive threads take consecutive elements, so one warp's 32
// addresses cover 128 contiguous bytes. Day 11 measures what it costs when
// they do not.
//
// Launch assumption: none, because this program never launches it. The
// occupancy API needs a real compiled function to answer about, and this one
// is deliberately small, so warp and block slots rather than registers are
// what limit residency on most cards. Day 45 goes the other way on purpose.
__global__ void scaleAdd(float a, const float* x, const float* y, 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 * x[i] + y[i];
}
}
// snippet: warps-per-block
// Warp slots a block of t threads occupies, which is ceil(t / 32). This is the
// guide's own formula, and it rounds up: a block of 48 threads takes two slots
// and the second one runs with 16 of its 32 lanes switched off.
// constexpr so the two claims below can be checked at compile time. They
// follow only from the constants in this file, so a run-time check would look
// like a measurement of your GPU and be nothing of the kind.
constexpr int warpsPerBlock(int t) {
return (t + kWarpSize - 1) / kWarpSize;
}
// end snippet
int main() {
const int device = 0;
CUDA_CHECK(cudaSetDevice(device));
cudaDeviceProp prop;
CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
// Counts checks that failed, so one bad row still reports every other.
int wrong = 0;
// cudaDeviceProp carries no maxWarpsPerSM field. It reports resident
// threads per SM, and the warp count is that over the warp size.
const int maxWarpsPerSM = prop.maxThreadsPerMultiProcessor / prop.warpSize;
std::printf("GPU: %s (compute capability %d.%d)\n\n", prop.name, prop.major,
prop.minor);
std::printf("What one SM holds\n");
std::printf(" SMs on the device %d\n",
prop.multiProcessorCount);
std::printf(" warp size %d threads\n", prop.warpSize);
std::printf(" max threads per block %d\n", prop.maxThreadsPerBlock);
std::printf(" max resident threads per SM %d\n",
prop.maxThreadsPerMultiProcessor);
std::printf(" max resident warps per SM %d\n", maxWarpsPerSM);
std::printf(" max resident blocks per SM %d\n",
prop.maxBlocksPerMultiProcessor);
std::printf(" 32-bit registers per SM %d\n",
prop.regsPerMultiprocessor);
std::printf(" shared memory per SM %zu KiB\n",
prop.sharedMemPerMultiprocessor / 1024);
std::printf(" shared memory per block %zu KiB\n",
prop.sharedMemPerBlock / 1024);
std::printf(" L2 cache %d KiB\n",
prop.l2CacheSize / 1024);
std::printf(
"\nOne SM, one block size at a time. The blocks/SM column is what the\n"
"driver reports for scaleAdd, not a formula this program made up.\n\n");
std::printf("%11s %10s %11s %11s %10s %10s\n", "threads/blk", "warps/blk",
"idle lanes", "blocks/SM", "warps/SM", "occupancy");
// snippet: residency
for (int c = 0; c < kNumBlockSizes; ++c) {
const int t = kBlockSizes[c];
const int warps = warpsPerBlock(t);
const int idle = warps * kWarpSize - t;
int resident = 0;
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&resident, scaleAdd, t, 0));
const int residentWarps = resident * warps;
// The two caps the page states in words, checked against the driver's
// own answer rather than asserted at the reader. A residency above
// either of them would mean the model on the page is wrong.
//
// These are real branches, not asserts. CI builds Release, which
// defines NDEBUG, and an assert under NDEBUG is deleted. A check that
// vanishes in the build that matters would let a wrong table reach the
// page with CI green, which is the one outcome this program exists to
// prevent.
if (resident > prop.maxBlocksPerMultiProcessor) {
std::fprintf(stderr,
"residency %d exceeds the per-SM block cap %d at %d "
"threads per block\n",
resident, prop.maxBlocksPerMultiProcessor, t);
++wrong;
}
if (residentWarps > maxWarpsPerSM) {
std::fprintf(stderr,
"resident warps %d exceeds the per-SM warp cap %d at "
"%d threads per block\n",
residentWarps, maxWarpsPerSM, t);
++wrong;
}
std::printf("%11d %10d %11d %11d %10d %9.0f%%\n", t, warps, idle,
resident, residentWarps,
100.0 * residentWarps / maxWarpsPerSM);
}
// end snippet
// The opening question, answered from the same arithmetic as the table.
// Both launches create 128 threads. They do not cost the same, because the
// hardware allocates warp slots and not threads.
const int smallBlocks = kQuizThreads / kSmallBlock;
const int largeBlocks = kQuizThreads / kWarpSize;
const int smallSlots = smallBlocks * warpsPerBlock(kSmallBlock);
const int largeSlots = largeBlocks * warpsPerBlock(kWarpSize);
std::printf(
"\n%d threads, launched two ways\n"
" %d blocks x 16 threads: %d warp slots, %d of %d lanes idle\n"
" %d blocks x 32 threads: %d warp slots, %d of %d lanes idle\n",
kQuizThreads, smallBlocks, smallSlots,
smallSlots * kWarpSize - kQuizThreads, smallSlots * kWarpSize,
largeBlocks, largeSlots, largeSlots * kWarpSize - kQuizThreads,
largeSlots * kWarpSize);
// Half the threads per block means twice the warp slots for the same work.
// That relation and the ceiling behaviour of warpsPerBlock follow only
// from the constants above, so they are compile-time claims. Asserting
// them at run time would look like a check against your GPU and be no
// such thing.
static_assert(warpsPerBlock(48) == 2,
"48 threads occupy 2 warp slots, 16 lanes idle");
static_assert(
(kQuizThreads / kSmallBlock) * warpsPerBlock(kSmallBlock) ==
2 * ((kQuizThreads / kWarpSize) * warpsPerBlock(kWarpSize)),
"16-thread blocks take twice the warp slots of 32-thread ones");
// This one does consult the device, so it is a runtime check.
if (prop.warpSize != kWarpSize) {
std::fprintf(stderr,
"this GPU reports warpSize %d; every number on the page "
"assumes %d\n",
prop.warpSize, kWarpSize);
++wrong;
}
std::printf(
"\nEvery row above describes one SM. Multiply blocks/SM by the SM\n"
"count to see how much of a grid is resident at once; the rest of the\n"
"grid is queued, not running.\n");
if (wrong > 0) {
std::fprintf(stderr,
"\n%d check(s) failed. The model on the lesson page does "
"not match this GPU.\n",
wrong);
return EXIT_FAILURE;
}
return EXIT_SUCCESS;
}