code/day04-indexing/indexing.cuThis is the source used by the lesson and its recorded evidence. Compile commands and expected output live in the directory README.
// Day 4: which thread touches element 1337?
//
// Every thread writes down who it is into the element it owns, and nothing
// else. Nothing is computed, so no arithmetic can hide an indexing bug behind
// a right-looking answer. The mapping is the output.
//
// The same kernel runs three ways: with the truncating grid size that leaves
// the tail of the array with no owner at all, with the ceiling grid size that
// covers it, and with 1024-thread blocks, which hands element 1337 to a
// different thread without one line of the kernel changing.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o indexing indexing.cu
// Run: ./indexing
//
// 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 <vector>
#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)
// 2000 is deliberately not a multiple of 256. 2000 = 7 * 256 + 208, so a grid
// sized by integer division is one block short and 208 elements end up with
// no owner, while the ceiling grid over-provides by 48 threads.
constexpr size_t kElems = 2000;
constexpr int kThreadsPerBlock = 256; // 8 warps
// The element the exercise asks about. It is inside kElems, so a correct
// launch always has exactly one thread that owns it.
constexpr size_t kQuery = 1337;
// The hardware ceiling on threads per block for every compute capability this
// course targets, from Table 30 of the compute capability appendix. Day 4 is
// not the block-size lesson (day 10 is); this exists only so the page can show
// the same element changing hands when the launch changes.
constexpr int kMaxThreadsPerBlock = 1024;
// 32 on every GPU this course targets. warpSize is a run-time built-in and
// cannot size anything at compile time, so the constant is spelled out here.
constexpr int kWarpSize = 32;
// No thread can write -1, so an element still holding it after a launch is an
// element nobody owned. cudaMemset writes bytes, and 0xff in all four bytes of
// an int is -1, which is why the fill below and this value agree.
constexpr int kUntouched = -1;
// snippet: index-kernel
// One thread writes its own two coordinates into the element it owns.
//
// Memory: consecutive threads in a warp have consecutive threadIdx.x, so one
// warp's 32 addresses cover 128 contiguous bytes of each output array. That is
// the coalesced pattern day 11 measures.
//
// Launch assumption: a 1D grid of 1D blocks. The guard is what makes an
// over-sized grid safe, so it is not optional.
__global__ void writeOwner(int* blockOf, int* threadOf, size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
blockOf[i] = static_cast<int>(blockIdx.x);
threadOf[i] = static_cast<int>(threadIdx.x);
}
}
// end snippet
// Refills both device arrays with kUntouched, runs one launch, and copies the
// result back. Every configuration goes through here, so no configuration can
// read another one's leftovers and look correct.
static void runConfig(int blocks, int threads, int* d_blockOf, int* d_threadOf,
std::vector<int>& h_blockOf,
std::vector<int>& h_threadOf) {
const size_t bytes = kElems * sizeof(int);
CUDA_CHECK(cudaMemset(d_blockOf, 0xff, bytes));
CUDA_CHECK(cudaMemset(d_threadOf, 0xff, bytes));
writeOwner<<<blocks, threads>>>(d_blockOf, d_threadOf, kElems);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(h_blockOf.data(), d_blockOf, bytes, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(h_threadOf.data(), d_threadOf, bytes,
cudaMemcpyDeviceToHost));
}
// How many elements no thread wrote. Zero is the only answer a correct launch
// gives; anything else is the truncating-grid bug, counted rather than argued.
static size_t countUntouched(const std::vector<int>& h_blockOf) {
size_t untouched = 0;
for (size_t i = 0; i < kElems; ++i) {
if (h_blockOf[i] == kUntouched) {
++untouched;
}
}
return untouched;
}
// Checks every written element against blockIdx * blockDim + threadIdx == i,
// which is the formula the lesson prints. It reports the first element that
// disagrees and returns false, because a page that states a formula the
// hardware did not follow is the failure this program exists to catch.
static bool ownersMatchFormula(const std::vector<int>& h_blockOf,
const std::vector<int>& h_threadOf,
int threads) {
for (size_t i = 0; i < kElems; ++i) {
if (h_blockOf[i] == kUntouched) {
continue;
}
const size_t owner =
static_cast<size_t>(h_blockOf[i]) * static_cast<size_t>(threads) +
static_cast<size_t>(h_threadOf[i]);
if (owner != i) {
std::fprintf(stderr,
"element %zu records block %d thread %d, which is "
"%zu, not %zu\n",
i, h_blockOf[i], h_threadOf[i], owner, i);
return false;
}
}
return true;
}
// Prints one configuration: its grid, how many elements it missed, how many
// threads it over-provided, and which thread ended up owning kQuery. The warp
// and lane come along because the warp is what the hardware schedules, which
// is day 21.
// Returns the number of elements no thread wrote, so main can gate on it.
// The README states these counts as facts, and a fact the program prints but
// never checks is a fact that drifts the first time somebody edits a constant.
static size_t report(int blocks, int threads, const std::vector<int>& h_blockOf,
const std::vector<int>& h_threadOf) {
const size_t launched =
static_cast<size_t>(blocks) * static_cast<size_t>(threads);
const size_t untouched = countUntouched(h_blockOf);
const size_t spare = (launched > kElems) ? launched - kElems : 0;
std::printf("%5d x %5d = %7zu threads unwritten %4zu idle %4zu ",
blocks, threads, launched, untouched, spare);
if (h_blockOf[kQuery] == kUntouched) {
std::printf("element %zu has no owner\n", kQuery);
return untouched;
}
const int b = h_blockOf[kQuery];
const int t = h_threadOf[kQuery];
std::printf("element %zu <- block %d, thread %d (warp %d, lane %d)\n",
kQuery, b, t, t / kWarpSize, t % kWarpSize);
return untouched;
}
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("%zu elements, asking who owns element %zu\n\n", kElems,
kQuery);
const size_t bytes = kElems * sizeof(int);
int* d_blockOf = nullptr;
int* d_threadOf = nullptr;
CUDA_CHECK(cudaMalloc(&d_blockOf, bytes));
CUDA_CHECK(cudaMalloc(&d_threadOf, bytes));
std::vector<int> h_blockOf(kElems);
std::vector<int> h_threadOf(kElems);
// snippet: grid-size
// Integer division truncates, so this grid is one block short whenever
// kElems is not a multiple of the block size. Nothing warns you. The
// launch succeeds, the kernel returns, and the tail of the array keeps
// whatever was in it before.
const int shortBlocks = static_cast<int>(kElems / kThreadsPerBlock);
// Ceiling division. One more block than you strictly need, and the guard
// inside the kernel switches off the threads that block over-provides.
const int blocks =
static_cast<int>((kElems + kThreadsPerBlock - 1) / kThreadsPerBlock);
// end snippet
const int bigBlocks = static_cast<int>((kElems + kMaxThreadsPerBlock - 1) /
kMaxThreadsPerBlock);
// The truncating grid is expected to lose exactly the elements its short
// launch cannot reach. Anything else means the lesson's arithmetic moved.
const size_t expectedLost =
kElems - static_cast<size_t>(shortBlocks) *
static_cast<size_t>(kThreadsPerBlock);
runConfig(shortBlocks, kThreadsPerBlock, d_blockOf, d_threadOf, h_blockOf,
h_threadOf);
const size_t lostShort =
report(shortBlocks, kThreadsPerBlock, h_blockOf, h_threadOf);
runConfig(blocks, kThreadsPerBlock, d_blockOf, d_threadOf, h_blockOf,
h_threadOf);
const size_t lostCeil =
report(blocks, kThreadsPerBlock, h_blockOf, h_threadOf);
const bool ok = ownersMatchFormula(h_blockOf, h_threadOf, kThreadsPerBlock);
runConfig(bigBlocks, kMaxThreadsPerBlock, d_blockOf, d_threadOf, h_blockOf,
h_threadOf);
const size_t lostBig =
report(bigBlocks, kMaxThreadsPerBlock, h_blockOf, h_threadOf);
CUDA_CHECK(cudaFree(d_blockOf));
CUDA_CHECK(cudaFree(d_threadOf));
int wrong = 0;
if (lostShort != expectedLost) {
std::fprintf(stderr,
"truncating grid lost %zu elements, expected %zu\n",
lostShort, expectedLost);
++wrong;
}
if (lostCeil != 0 || lostBig != 0) {
std::fprintf(stderr,
"a rounded-up grid lost elements (%zu and %zu); the "
"bounds check or the grid arithmetic is wrong\n",
lostCeil, lostBig);
++wrong;
}
if (!ok || wrong > 0) {
return EXIT_FAILURE;
}
std::printf(
"\nEvery written element satisfied blockIdx.x * blockDim.x +\n"
"threadIdx.x == its own index. The only row that lost elements is the\n"
"one whose grid was sized with a truncating division.\n");
return EXIT_SUCCESS;
}