code/day19-unified-memory/unified_memory.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 19: day 5's vector add on managed memory, timed by where the pages are.
//
// One kernel, one element count, one set of inputs. The only thing that
// changes from row to row is where the pages live when the launch happens:
//
// cold the CPU touched the data last, so the GPU faults it in
// advised the same, after cudaMemAdviseSetPreferredLocation
// prefetched cudaMemPrefetchAsync moved the pages before the launch
// resident the pages are already on the card from the launch before
//
// plus day 5's explicit cudaMalloc path as the anchor, and the migration
// timed on its own so the cold row's time has somewhere to go.
//
// The kernel is day 5's, unchanged, and carries no __restrict__ even though
// day 11 introduced it. Adding it would put a second variable into a lesson
// whose whole claim is that only the pointer changed.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o unified_memory unified_memory.cu
// Run: ./unified_memory
//
// 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 <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. That is day 9's size, so the resident row here can
// be read straight against day 9's warm kernel row on the same card, and 611
// is day 5's number: it stops the count dividing by any block size, so the
// kernel's bounds check runs on every launch.
//
// Six device-side buffers of 64 MiB each, three managed and three explicit,
// plus four host vectors. Roughly 640 MiB in total, which leaves room on a
// 15 GiB T4 and on a free Colab instance.
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;
constexpr int kManagedBuffers = 3;
// 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(static_cast<size_t>(blocksFor(kElems)) * kThreadsPerBlock >=
kElems,
"the grid must cover every element");
// CUDA 13.0 changed the shape of both hint APIs: the `int device` parameter
// became a `cudaMemLocation` struct, so a call written against 12.x does not
// compile on 13.x and the reverse is also true. NVIDIA made the same edit to
// its own samples (https://github.com/NVIDIA/cuda-samples , CHANGELOG.md:
// "changing the parameter int device to cudaMemLocation location").
//
// These two wrappers are the entire difference, and everything below calls
// them rather than the runtime function, so the version split lives in one
// place instead of at seven call sites.
#if CUDART_VERSION >= 13000
static cudaMemLocation locationOf(int device) {
cudaMemLocation loc;
loc.type = (device == cudaCpuDeviceId) ? cudaMemLocationTypeHost
: cudaMemLocationTypeDevice;
loc.id = (device == cudaCpuDeviceId) ? 0 : device;
return loc;
}
static cudaError_t prefetchTo(const void* p, size_t bytes, int device) {
return cudaMemPrefetchAsync(p, bytes, locationOf(device), 0, 0);
}
static cudaError_t adviseFor(const void* p, size_t bytes,
cudaMemoryAdvise advice, int device) {
return cudaMemAdvise(p, bytes, advice, locationOf(device));
}
#else
static cudaError_t prefetchTo(const void* p, size_t bytes, int device) {
return cudaMemPrefetchAsync(p, bytes, device, 0);
}
static cudaError_t adviseFor(const void* p, size_t bytes,
cudaMemoryAdvise advice, int device) {
return cudaMemAdvise(p, bytes, advice, device);
}
#endif
// out[i] = a[i] + b[i]. One thread owns one element. This is day 5's kernel
// and the pointers may be managed or explicit; the kernel cannot tell, which
// is the whole ergonomic case for managed memory.
//
// 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. On a cold managed buffer those same addresses fault:
// the page has to arrive before the load can retire, and the launch waits.
//
// 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];
}
}
// CPU reference. Written for obvious correctness, not speed: plain loop, no
// OpenMP, no intrinsics. It never allocates; the caller owns every buffer.
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;
}
// Returns the first index where two float arrays differ by a single bit, or n
// if they are identical. Not a tolerance comparison: the managed run and the
// explicit run execute the same kernel over the same inputs, one element per
// thread with no reduction, so the two results have to match exactly.
static size_t firstDifference(const float* left, const float* right, size_t n) {
for (size_t i = 0; i < n; ++i) {
if (left[i] != right[i]) {
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.
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;
}
// One launch, timed, with no warm-up inside it.
//
// The missing warm-up is deliberate and it is not the hole day 9 warns about.
// The kernel is already warm here: main() runs timeKernel over the explicit
// buffers before any of this, so the module is loaded and the code is hot.
// What is cold is the memory, and a page migrates once. Time it twice and the
// second answer is a different measurement.
static float timeOneLaunch(cudaEvent_t start, cudaEvent_t stop, const float* a,
const float* b, float* out, size_t n) {
CUDA_CHECK(cudaEventRecord(start));
vectorAdd<<<blocksFor(n), kThreadsPerBlock>>>(a, b, out, n);
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
CUDA_CHECK(cudaGetLastError());
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
return ms;
}
// Migrates `count` managed buffers to `device`, which may be cudaCpuDeviceId,
// and returns the milliseconds it took.
//
// Called with the CPU as the destination this is the reset between rows: it
// puts every page back on the host so the next launch has to fault them in
// again. Called with the GPU it is the optimisation the lesson is measuring.
// Same call either way, which is the useful thing about it.
//
// The prefetch is stream ordered, so events on the same stream bracket it:
// "The migration does not begin until all prior operations in the stream have
// completed, and completes before any subsequent operation in the stream."
// snippet: move-pages
static float movePages(cudaEvent_t start, cudaEvent_t stop,
float* const* buffers, int count, size_t bytes,
int device) {
CUDA_CHECK(cudaEventRecord(start));
for (int k = 0; k < count; ++k) {
CUDA_CHECK(prefetchTo(buffers[k], bytes, device));
}
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
return ms;
}
// end snippet
// Applies one piece of advice to every buffer. Separate from movePages
// because advice moves nothing: it changes where a fault resolves to, and the
// runtime documentation is explicit that "setting the preferred location does
// not cause data to migrate to that location immediately".
static void adviseAll(float* const* buffers, int count, size_t bytes,
cudaMemoryAdvise advice, int device) {
for (int k = 0; k < count; ++k) {
CUDA_CHECK(adviseFor(buffers[k], bytes, advice, device));
}
}
// Names which of the four unified memory paradigms this device is on, from
// the three attributes the programming guide's own decision tree reads
// (section 2.6.2.1, "Unified Memory Paradigms").
static void printParadigm(int concurrent, int pageable, int hostPageTables) {
if (concurrent == 0) {
std::printf(" paradigm limited: Windows, WSL 2 or Tegra\n");
return;
}
if (pageable == 0) {
std::printf(" paradigm full, for cudaMallocManaged allocations\n");
return;
}
std::printf(" paradigm full, for every allocation, %s coherence\n",
hostPageTables ? "hardware" : "software");
}
int main() {
const int device = 0;
CUDA_CHECK(cudaSetDevice(device));
cudaDeviceProp prop;
CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
// Ask the device rather than the operating system. This is the check the
// WSL user guide asks for by name: "CUDA queries will say whether it is
// supported or not and applications are expected to check this."
// snippet: paradigm-query
int managedMemory = 0;
int concurrent = 0;
int pageable = 0;
int hostPageTables = 0;
CUDA_CHECK(cudaDeviceGetAttribute(&managedMemory, cudaDevAttrManagedMemory,
device));
CUDA_CHECK(cudaDeviceGetAttribute(
&concurrent, cudaDevAttrConcurrentManagedAccess, device));
CUDA_CHECK(cudaDeviceGetAttribute(&pageable,
cudaDevAttrPageableMemoryAccess, device));
CUDA_CHECK(cudaDeviceGetAttribute(
&hostPageTables, cudaDevAttrPageableMemoryAccessUsesHostPageTables,
device));
// end snippet
const size_t bytes = kElems * sizeof(float);
const int blocks = blocksFor(kElems);
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("What this device says about unified memory\n");
std::printf(" managedMemory %d\n", managedMemory);
std::printf(" concurrentManagedAccess %d\n", concurrent);
std::printf(" pageableMemoryAccess %d\n", pageable);
std::printf(" pageableMemoryAccessUsesHostPageTables %d\n",
hostPageTables);
printParadigm(concurrent, pageable, hostPageTables);
std::printf("\n");
// Nothing below can run without managed memory, and there is no fallback
// worth writing, because the lesson is about the allocator. Exit before
// any allocation, so this path has nothing to free.
if (managedMemory == 0) {
std::fprintf(stderr,
"this device cannot allocate managed memory, so day 19 "
"has nothing to measure here\n");
return EXIT_FAILURE;
}
// Both hint APIs refuse a device whose concurrentManagedAccess is 0, and
// so does a prefetch to the CPU, because the stream belongs to that same
// device. That is Windows, WSL 2 and some Tegra parts. The run still
// checks the answer, says why the table is missing, and exits 0, because
// a learner on WSL 2 has not broken anything.
const bool canMigrate = (concurrent != 0);
if (!canMigrate) {
std::printf(
"concurrentManagedAccess is 0, so cudaMemPrefetchAsync and\n"
"cudaMemAdvise are unavailable here and the migration table is\n"
"skipped. Your build is fine and this is expected on Windows and\n"
"on WSL 2. NVIDIA: \"Full Managed Memory Support is not available\n"
"on Windows native and therefore WSL 2 will not support it for\n"
"the foreseeable future.\" The correctness check still runs.\n\n");
}
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));
// cudaMallocManaged takes the address of the pointer for the same reason
// cudaMalloc does, and hands back one address that both the CPU and the
// GPU may dereference. It does not decide where the memory lives: on a
// device with concurrentManagedAccess set, "managed memory may not be
// populated when this API returns and instead may be populated on
// access". The `u_` prefix is the course's third pointer prefix, after
// `h_` and `d_`, and it means exactly this allocator.
// snippet: managed-alloc
float* u_a = nullptr;
float* u_b = nullptr;
float* u_out = nullptr;
CUDA_CHECK(cudaMallocManaged(&u_a, bytes));
CUDA_CHECK(cudaMallocManaged(&u_b, bytes));
CUDA_CHECK(cudaMallocManaged(&u_out, bytes));
// First touch is on the host, through an ordinary pointer. No cudaMemcpy,
// no direction argument, no second pointer to keep in step.
for (size_t i = 0; i < kElems; ++i) {
u_a[i] = h_a[i];
u_b[i] = h_b[i];
}
// end snippet
float* managed[kManagedBuffers] = {u_a, u_b, u_out};
float* justOut[1] = {u_out};
cudaEvent_t start;
cudaEvent_t stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));
// The anchor, and it has to come first. timeKernel warms vectorAdd three
// times before it times anything, so after this line the module is loaded
// and the code is hot. Every managed row below is therefore cold in its
// memory and warm in its code, which is the only way to attribute a
// difference to migration.
const float explicitMs = timeKernel([&] {
vectorAdd<<<blocks, kThreadsPerBlock>>>(d_a, d_b, d_out, kElems);
});
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
float coldMs = 0.0f;
float residentOnceMs = 0.0f;
float residentMeanMs = 0.0f;
float advisedMs = 0.0f;
float prefetchedMs = 0.0f;
float migrateInMs = 0.0f;
float migrateOutMs = 0.0f;
if (canMigrate) {
// Cold: every page sits on the host, so the launch faults them in.
movePages(start, stop, managed, kManagedBuffers, bytes,
cudaCpuDeviceId);
coldMs = timeOneLaunch(start, stop, u_a, u_b, u_out, kElems);
// Resident: the launch above left all three buffers on the card and
// nothing has touched them from the host since.
residentOnceMs = timeOneLaunch(start, stop, u_a, u_b, u_out, kElems);
residentMeanMs = timeKernel([&] {
vectorAdd<<<blocks, kThreadsPerBlock>>>(u_a, u_b, u_out, kElems);
});
// Advice only. A preferred location is a policy for resolving a
// fault, not a migration, so this row is there to be compared against
// the cold row rather than against the prefetched one. It is unset
// afterwards so it cannot leak into the rows below.
movePages(start, stop, managed, kManagedBuffers, bytes,
cudaCpuDeviceId);
adviseAll(managed, kManagedBuffers, bytes,
cudaMemAdviseSetPreferredLocation, device);
advisedMs = timeOneLaunch(start, stop, u_a, u_b, u_out, kElems);
adviseAll(managed, kManagedBuffers, bytes,
cudaMemAdviseUnsetPreferredLocation, device);
// Prefetched: the same cold start, with the migration lifted out of
// the launch and measured on a line of its own.
movePages(start, stop, managed, kManagedBuffers, bytes,
cudaCpuDeviceId);
migrateInMs =
movePages(start, stop, managed, kManagedBuffers, bytes, device);
prefetchedMs = timeOneLaunch(start, stop, u_a, u_b, u_out, kElems);
// And the trip back. The host is about to read u_out, and where
// directManagedMemAccessFromHost is 0 that read faults every page
// across to the CPU. Doing it in one call is both faster and
// measurable.
migrateOutMs =
movePages(start, stop, justOut, 1, bytes, cudaCpuDeviceId);
} else {
vectorAdd<<<blocks, kThreadsPerBlock>>>(u_a, u_b, u_out, kElems);
CUDA_CHECK(cudaGetLastError());
}
// Required before the host reads u_out, and not optional. The launch is
// asynchronous, so without this the CPU races the kernel for the same
// pages: undefined behaviour where concurrentManagedAccess is 0, and a
// plain read of stale data where it is 1.
CUDA_CHECK(cudaDeviceSynchronize());
vectorAddCpu(h_a.data(), h_b.data(), h_want.data(), kElems);
const size_t badManaged =
firstMismatch(u_out, h_want.data(), kElems, kRelTolerance);
const size_t badExplicit =
firstMismatch(h_out.data(), h_want.data(), kElems, kRelTolerance);
const size_t disagree = firstDifference(u_out, h_out.data(), kElems);
// Read the offending values out before the frees below, because u_out is
// about to stop existing and the failure branches print from it.
const double gotManaged =
(badManaged == kElems) ? 0.0 : static_cast<double>(u_out[badManaged]);
const double gotDisagree =
(disagree == kElems) ? 0.0 : static_cast<double>(u_out[disagree]);
// One cleanup block, reached by every path that allocated anything. Every
// cudaMalloc and every cudaMallocManaged has its cudaFree here, and the
// failure branches are all below this line for that reason.
CUDA_CHECK(cudaEventDestroy(start));
CUDA_CHECK(cudaEventDestroy(stop));
CUDA_CHECK(cudaFree(d_a));
CUDA_CHECK(cudaFree(d_b));
CUDA_CHECK(cudaFree(d_out));
CUDA_CHECK(cudaFree(u_a));
CUDA_CHECK(cudaFree(u_b));
CUDA_CHECK(cudaFree(u_out));
if (canMigrate) {
std::printf("Where the pages are when the kernel starts\n");
std::printf(" %-40s %10.3f ms\n", "explicit cudaMalloc, resident",
static_cast<double>(explicitMs));
std::printf(" %-40s %10.3f ms\n", "managed, resident",
static_cast<double>(residentMeanMs));
std::printf(" %-40s %10.3f ms\n", "managed, resident, one launch",
static_cast<double>(residentOnceMs));
std::printf(" %-40s %10.3f ms\n", "managed, cold, one launch",
static_cast<double>(coldMs));
std::printf(" %-40s %10.3f ms\n",
"managed, cold, preferred location set",
static_cast<double>(advisedMs));
std::printf(" %-40s %10.3f ms\n\n", "managed, prefetched, one launch",
static_cast<double>(prefetchedMs));
std::printf("The migration on its own\n");
std::printf(" %-40s %10.3f ms\n", "3 buffers, host to device",
static_cast<double>(migrateInMs));
std::printf(" %-40s %10.3f ms\n\n", "1 buffer, device to host",
static_cast<double>(migrateOutMs));
std::printf(
"The first two rows are means of %d runs after %d warm-ups. The\n"
"other four are single launches, because a page migrates once and\n"
"you cannot get it back, so they carry more noise. Read those\n"
"four against each other, not against the means.\n\n",
kTimedRuns, kWarmupRuns);
}
// Three gates, all 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.
//
// None of the three is a prediction about your hardware. The first two
// are the answer, and the third is arithmetic: the same kernel over the
// same inputs, one element per thread and no reduction, has to produce
// the same bits whichever allocator held the memory.
if (badManaged != kElems) {
std::fprintf(stderr, "managed wrong at %zu: got %.9g, want %.9g\n",
badManaged, gotManaged,
static_cast<double>(h_want[badManaged]));
return EXIT_FAILURE;
}
if (badExplicit != kElems) {
std::fprintf(stderr, "explicit wrong at %zu: got %.9g, want %.9g\n",
badExplicit, static_cast<double>(h_out[badExplicit]),
static_cast<double>(h_want[badExplicit]));
return EXIT_FAILURE;
}
if (disagree != kElems) {
std::fprintf(
stderr, "managed and explicit disagree at %zu: %.9g against %.9g\n",
disagree, gotDisagree, static_cast<double>(h_out[disagree]));
return EXIT_FAILURE;
}
std::printf("all %zu elements match the CPU reference\n", kElems);
std::printf("managed and explicit agree bit for bit\n");
return EXIT_SUCCESS;
}