code/day93-unified-memory/managed_tuning.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 93: tuning a managed workload with cudaMemAdvise and
// cudaMemPrefetchAsync.
//
// Day 19 measured one launch over cold managed memory and one after a
// prefetch. This program measures a workload instead: four rounds of a
// gather-accumulate kernel, with a host pass over the lookup table between
// them, which is the shape that makes page faulting a repeating cost rather
// than a one-off.
//
// Five rows over the same arithmetic and the same inputs:
//
// explicit cudaMalloc plus cudaMemcpy, the ceiling
// managed naive first touch on the host, no advice, no prefetch
// managed prefetch cudaMemPrefetchAsync before every kernel
// managed advise cudaMemAdvise only, so nothing is moved up front
// managed both advice and prefetch together
//
// Three managed buffers, each with a different access pattern, because the
// three pieces of advice are answers to three different patterns:
//
// u_table read by the device every round and by the host every round
// u_keys read by the device every round, touched by the host only once
// u_out written by the device every round, read by the host at the end
//
// What it does not measure: the kernel's own access pattern. `keys[i]` walks
// the table with stride 3, so a warp's 32 lanes cover 96 consecutive floats
// and every row pays the same day 11 penalty. The stride is there to make
// every page of the table necessary, not to be fast, and it cancels between
// rows.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o managed_tuning managed_tuning.cu
// Run: ./managed_tuning
//
// VERIFIED: compiled and run on a Tesla T4 on 2026-09-02. The transcript is
// evidence/run-2026-09-02.txt; publish only the values captured there.
#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 table entries is 64 MiB, which is 16,384 host pages at 4 KiB and 32
// device pages at 2 MiB. Big enough that a fault storm has a shape, small
// enough that six buffers of this order fit a 15 GiB T4 and a free Colab
// instance. The output count carries day 5's 611 so no block size divides
// it and the kernel's bounds check runs on every launch.
constexpr size_t kTableElems = 16ull * 1024ull * 1024ull;
constexpr size_t kElems = 16ull * 1024ull * 1024ull + 611ull;
constexpr size_t kTableBytes = kTableElems * sizeof(float);
constexpr size_t kKeyBytes = kElems * sizeof(unsigned int);
constexpr size_t kOutBytes = kElems * sizeof(float);
constexpr int kThreadsPerBlock = 256; // 8 warps
constexpr int kRounds = 4;
constexpr int kWarmupRuns = 3;
constexpr int kConfigs = 5;
constexpr size_t kKeyStride = 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");
static_assert(kTableElems % kKeyStride != 0,
"a stride that divides the table would revisit a subset of it");
// 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 day 19's, unchanged, and everything below calls
// them rather than the runtime function.
#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
// Three timings per row, plus the per-round kernel times the advice claim is
// read from. A plain struct holding data, which modules 4 and up allow.
struct RowTimes {
float totalMs;
float kernelMs;
float moveMs;
float roundMs[kRounds];
};
// out[i] += scale * table[keys[i]]. One thread owns one output element and
// one table entry, and the accumulation is in place across rounds.
//
// Memory: consecutive threads take consecutive outputs, so the store and the
// key load are contiguous, while `table[keys[i]]` walks with stride 3, so a
// warp's 32 lanes touch 96 consecutive floats. Nothing here is a cache
// resident: the table is 64 MiB against a 4 MiB L2 on a T4. On a cold
// managed buffer every one of those addresses can fault, and the load cannot
// retire until its page has arrived.
//
// Launch assumption: gridDim.x * blockDim.x >= n, and every key is already
// below kTableElems, which the host fill guarantees. The guard is a plain
// `if` and not an early return.
__global__ void gatherAccumulate(const float* __restrict__ table,
const unsigned int* __restrict__ keys,
float* __restrict__ out, size_t n,
float scale) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
out[i] += scale * table[keys[i]];
}
}
// The weight each round applies. Four exact values, so the only rounding in
// the reference comparison comes from the table entries and the sum.
static float scaleFor(int r) {
return 0.5f * static_cast<float>(r + 1);
}
// CPU reference. Accumulates in double where the kernel accumulates in
// float, written for obvious correctness rather than speed, and it never
// allocates.
static void gatherAccumulateCpu(const float* table, const unsigned int* keys,
double* out, size_t n, double scale) {
for (size_t i = 0; i < n; ++i) {
out[i] += scale * static_cast<double>(table[keys[i]]);
}
}
// The host pass over the table, once per round. It is the reason the table
// is read-shared rather than device-only, and without SetReadMostly it is
// what drags the pages back off the card between kernels.
//
// The sum is returned and checked, not thrown away, for two reasons: an
// unused loop over 64 MiB is a loop the compiler may delete, and a
// read-mostly region that ever disagrees with the reference means a
// duplicate went stale, which is the one way this advice can hurt you.
static double tableChecksum(const float* table) {
double sum = 0.0;
for (size_t j = 0; j < kTableElems; ++j) {
sum += static_cast<double>(table[j]);
}
return sum;
}
// Milliseconds between two events, after waiting for the later one.
static float elapsedMs(cudaEvent_t start, cudaEvent_t stop) {
CUDA_CHECK(cudaEventSynchronize(stop));
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
return ms;
}
// Returns the first index where a float result leaves the tolerance band
// around the double reference, or n if it never does. Returning the index
// rather than a bool is the point: "wrong at 512" names the block.
static size_t firstMismatch(const float* got, const double* want, size_t n,
double rtol, double atol) {
for (size_t i = 0; i < n; ++i) {
const double diff = std::fabs(static_cast<double>(got[i]) - want[i]);
if (diff > atol + rtol * std::fabs(want[i])) {
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: every row runs the
// same kernel over the same inputs in the same order, so where the pages
// happened to live cannot change a single bit of the answer.
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;
}
// The three pieces of advice, one per access pattern. None of them moves a
// byte: advice changes where a fault resolves to and which copies may exist,
// and cudaMemPrefetchAsync is the only call here that transfers anything.
// snippet: advice
static void applyAdvice(float* u_table, unsigned int* u_keys, float* u_out,
int device) {
// Read by both processors, written by neither, so each may keep its own
// read-only copy and the host pass stops pulling pages off the card.
// The device argument is ignored for this advice.
CUDA_CHECK(
adviseFor(u_table, kTableBytes, cudaMemAdviseSetReadMostly, device));
// Read by the device on every round, touched by the host once at fill
// time. A preferred location does not migrate anything; it decides where
// a fault resolves, so these pages settle on the card and stay.
CUDA_CHECK(adviseFor(u_keys, kKeyBytes, cudaMemAdviseSetPreferredLocation,
device));
// Written by the device every round and read by the host once, at the
// end. Preferred location keeps it on the card; AccessedBy maps it for
// the host so that last read can cross the link instead of migrating.
CUDA_CHECK(
adviseFor(u_out, kOutBytes, cudaMemAdviseSetPreferredLocation, device));
CUDA_CHECK(adviseFor(u_out, kOutBytes, cudaMemAdviseSetAccessedBy,
cudaCpuDeviceId));
}
// end snippet
// Clears every hint a previous row set and puts all three buffers back on
// the host, so "cold" means the same thing at the top of every row instead
// of meaning whatever the row before it left behind.
static void resetManaged(float* u_table, unsigned int* u_keys, float* u_out,
int device) {
CUDA_CHECK(
adviseFor(u_table, kTableBytes, cudaMemAdviseUnsetReadMostly, device));
CUDA_CHECK(adviseFor(u_table, kTableBytes,
cudaMemAdviseUnsetPreferredLocation, device));
CUDA_CHECK(adviseFor(u_keys, kKeyBytes, cudaMemAdviseUnsetPreferredLocation,
device));
CUDA_CHECK(adviseFor(u_out, kOutBytes, cudaMemAdviseUnsetPreferredLocation,
device));
CUDA_CHECK(adviseFor(u_out, kOutBytes, cudaMemAdviseUnsetAccessedBy,
cudaCpuDeviceId));
CUDA_CHECK(prefetchTo(u_table, kTableBytes, cudaCpuDeviceId));
CUDA_CHECK(prefetchTo(u_keys, kKeyBytes, cudaCpuDeviceId));
CUDA_CHECK(prefetchTo(u_out, kOutBytes, cudaCpuDeviceId));
CUDA_CHECK(cudaDeviceSynchronize());
}
// One managed row. The timed region opens with every page on the host and
// closes after the host has read the whole output, so it contains the
// migration wherever the runtime decides to put it.
static RowTimes runManaged(cudaEvent_t start, cudaEvent_t stop,
cudaEvent_t phaseStart, cudaEvent_t phaseStop,
float* u_table, unsigned int* u_keys, float* u_out,
float* h_got, double want, int device, bool advise,
bool prefetch, int* badRounds) {
RowTimes row = {0.0f, 0.0f, 0.0f, {0.0f}};
resetManaged(u_table, u_keys, u_out, device);
if (advise) {
applyAdvice(u_table, u_keys, u_out, device);
}
// First touch of the output is on the host, which is what a program that
// clears a managed buffer from CPU code actually does.
for (size_t i = 0; i < kElems; ++i) {
u_out[i] = 0.0f;
}
CUDA_CHECK(cudaEventRecord(start));
// snippet: round-loop
for (int r = 0; r < kRounds; ++r) {
if (prefetch) {
CUDA_CHECK(cudaEventRecord(phaseStart));
CUDA_CHECK(prefetchTo(u_table, kTableBytes, device));
CUDA_CHECK(prefetchTo(u_keys, kKeyBytes, device));
CUDA_CHECK(prefetchTo(u_out, kOutBytes, device));
CUDA_CHECK(cudaEventRecord(phaseStop));
row.moveMs += elapsedMs(phaseStart, phaseStop);
}
CUDA_CHECK(cudaEventRecord(phaseStart));
gatherAccumulate<<<blocksFor(kElems), kThreadsPerBlock>>>(
u_table, u_keys, u_out, kElems, scaleFor(r));
CUDA_CHECK(cudaEventRecord(phaseStop));
CUDA_CHECK(cudaGetLastError());
// elapsedMs waits on phaseStop, so the launch has finished before
// the host touches the table below. That is the required
// synchronisation, not an accident of the timing.
row.roundMs[r] = elapsedMs(phaseStart, phaseStop);
row.kernelMs += row.roundMs[r];
if (tableChecksum(u_table) != want) {
*badRounds += 1;
}
}
// end snippet
// The final host read of the output, inside the timed region because it
// is part of the workload. On a naive row this faults 64 MiB back one
// page at a time.
for (size_t i = 0; i < kElems; ++i) {
h_got[i] = u_out[i];
}
CUDA_CHECK(cudaEventRecord(stop));
row.totalMs = elapsedMs(start, stop);
return row;
}
// The explicit row, and the ceiling. Two copies in, four kernels over
// resident device memory, one copy out, and a host pass that reads the
// host's own table and therefore migrates nothing.
static RowTimes runExplicit(cudaEvent_t start, cudaEvent_t stop,
cudaEvent_t phaseStart, cudaEvent_t phaseStop,
const float* h_table, const unsigned int* h_keys,
float* d_table, unsigned int* d_keys, float* d_out,
float* h_got, double want, int* badRounds) {
RowTimes row = {0.0f, 0.0f, 0.0f, {0.0f}};
CUDA_CHECK(cudaMemset(d_out, 0, kOutBytes));
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaEventRecord(start));
CUDA_CHECK(cudaEventRecord(phaseStart));
CUDA_CHECK(
cudaMemcpy(d_table, h_table, kTableBytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_keys, h_keys, kKeyBytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaEventRecord(phaseStop));
row.moveMs += elapsedMs(phaseStart, phaseStop);
for (int r = 0; r < kRounds; ++r) {
CUDA_CHECK(cudaEventRecord(phaseStart));
gatherAccumulate<<<blocksFor(kElems), kThreadsPerBlock>>>(
d_table, d_keys, d_out, kElems, scaleFor(r));
CUDA_CHECK(cudaEventRecord(phaseStop));
CUDA_CHECK(cudaGetLastError());
row.roundMs[r] = elapsedMs(phaseStart, phaseStop);
row.kernelMs += row.roundMs[r];
if (tableChecksum(h_table) != want) {
*badRounds += 1;
}
}
CUDA_CHECK(cudaEventRecord(phaseStart));
CUDA_CHECK(cudaMemcpy(h_got, d_out, kOutBytes, cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaEventRecord(phaseStop));
row.moveMs += elapsedMs(phaseStart, phaseStop);
CUDA_CHECK(cudaEventRecord(stop));
row.totalMs = elapsedMs(start, stop);
return row;
}
int main() {
const int device = 0;
CUDA_CHECK(cudaSetDevice(device));
cudaDeviceProp prop;
CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
// Ask the device, not the operating system. concurrentManagedAccess is
// the attribute that decides whether this page has anything to teach:
// where it is 0 the runtime migrates in bulk at launch and back at
// synchronize, so there is no fault storm for a hint to fix and both
// hint APIs refuse the device.
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));
std::printf("GPU: %s (compute capability %d.%d)\n", prop.name, prop.major,
prop.minor);
std::printf(
"table %zu floats (%.1f MiB), n = %zu, %d rounds, %d threads/block, "
"%d blocks\n",
kTableElems, static_cast<double>(kTableBytes) / (1024.0 * 1024.0),
kElems, kRounds, kThreadsPerBlock, blocksFor(kElems));
std::printf(" managedMemory %d\n", managedMemory);
std::printf(" concurrentManagedAccess %d\n", concurrent);
std::printf(" pageableMemoryAccess %d\n", pageable);
std::printf(" pageableMemoryAccessUsesHostPageTables %d\n\n",
hostPageTables);
// Exit before any allocation, so this path has nothing to free.
if (managedMemory == 0 || concurrent == 0) {
std::fprintf(stderr,
"this device reports managedMemory %d and "
"concurrentManagedAccess %d, so cudaMemAdvise and "
"cudaMemPrefetchAsync are unavailable and day 93 has "
"nothing to tune here. Use a Linux host with a discrete "
"NVIDIA GPU, for example a free Colab T4.\n",
managedMemory, concurrent);
return EXIT_FAILURE;
}
std::vector<float> h_table(kTableElems);
std::vector<unsigned int> h_keys(kElems);
std::vector<float> h_got(kElems);
std::vector<float> h_first(kElems);
std::vector<double> h_want(kElems, 0.0);
// Reciprocals of small integers, so almost no table entry is exactly
// representable in float and the tolerance below has something to do.
for (size_t j = 0; j < kTableElems; ++j) {
h_table[j] = 1.0f / static_cast<float>((j % 97) + 1);
}
for (size_t i = 0; i < kElems; ++i) {
h_keys[i] =
static_cast<unsigned int>((i * kKeyStride + 1) % kTableElems);
}
for (int r = 0; r < kRounds; ++r) {
gatherAccumulateCpu(h_table.data(), h_keys.data(), h_want.data(),
kElems, static_cast<double>(scaleFor(r)));
}
double maxAbsRef = 0.0;
for (size_t i = 0; i < kElems; ++i) {
maxAbsRef = std::fmax(maxAbsRef, std::fabs(h_want[i]));
}
// Each output is the sum of kRounds products, so the worst-case relative
// error grows like kRounds * eps with eps = 2^-23 for float. The factor
// of four is headroom for the ordering difference against the double
// reference. This is the deterministic bound rather than the sqrt(depth)
// statistical one, because a depth of four is too shallow for the
// statistical form to mean anything.
const double rtol = 4.0 * static_cast<double>(kRounds) * 0x1p-23;
const double atol = rtol * maxAbsRef;
const double wantChecksum = tableChecksum(h_table.data());
float* d_table = nullptr;
unsigned int* d_keys = nullptr;
float* d_out = nullptr;
CUDA_CHECK(cudaMalloc(&d_table, kTableBytes));
CUDA_CHECK(cudaMalloc(&d_keys, kKeyBytes));
CUDA_CHECK(cudaMalloc(&d_out, kOutBytes));
float* u_table = nullptr;
unsigned int* u_keys = nullptr;
float* u_out = nullptr;
CUDA_CHECK(cudaMallocManaged(&u_table, kTableBytes));
CUDA_CHECK(cudaMallocManaged(&u_keys, kKeyBytes));
CUDA_CHECK(cudaMallocManaged(&u_out, kOutBytes));
for (size_t j = 0; j < kTableElems; ++j) {
u_table[j] = h_table[j];
}
for (size_t i = 0; i < kElems; ++i) {
u_keys[i] = h_keys[i];
}
cudaEvent_t start;
cudaEvent_t stop;
cudaEvent_t phaseStart;
cudaEvent_t phaseStop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));
CUDA_CHECK(cudaEventCreate(&phaseStart));
CUDA_CHECK(cudaEventCreate(&phaseStop));
// Warm up the one kernel this program times, on the explicit buffers, so
// that no row below pays the module load lazy loading defers to the
// first launch. The output it dirties is zeroed inside runExplicit.
for (int i = 0; i < kWarmupRuns; ++i) {
gatherAccumulate<<<blocksFor(kElems), kThreadsPerBlock>>>(
d_table, d_keys, d_out, kElems, scaleFor(0));
}
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaGetLastError());
const char* names[kConfigs] = {"explicit cudaMalloc", "managed, naive",
"managed, prefetch", "managed, advise",
"managed, advise + prefetch"};
RowTimes rows[kConfigs];
int badRounds[kConfigs] = {0, 0, 0, 0, 0};
size_t badIndex[kConfigs];
size_t disagree[kConfigs] = {kElems, kElems, kElems, kElems, kElems};
// The offending values are read out here, inside the loop, because
// h_got is overwritten by the next row and the gates run after all five.
double gotBad[kConfigs] = {0.0, 0.0, 0.0, 0.0, 0.0};
double gotDiff[kConfigs] = {0.0, 0.0, 0.0, 0.0, 0.0};
for (int c = 0; c < kConfigs; ++c) {
if (c == 0) {
rows[c] =
runExplicit(start, stop, phaseStart, phaseStop, h_table.data(),
h_keys.data(), d_table, d_keys, d_out, h_got.data(),
wantChecksum, &badRounds[c]);
} else {
rows[c] =
runManaged(start, stop, phaseStart, phaseStop, u_table, u_keys,
u_out, h_got.data(), wantChecksum, device,
c == 3 || c == 4, c == 2 || c == 4, &badRounds[c]);
}
badIndex[c] =
firstMismatch(h_got.data(), h_want.data(), kElems, rtol, atol);
if (badIndex[c] != kElems) {
gotBad[c] = static_cast<double>(h_got[badIndex[c]]);
}
if (c == 0) {
for (size_t i = 0; i < kElems; ++i) {
h_first[i] = h_got[i];
}
} else {
disagree[c] = firstDifference(h_got.data(), h_first.data(), kElems);
if (disagree[c] != kElems) {
gotDiff[c] = static_cast<double>(h_got[disagree[c]]);
}
}
}
// One cleanup block, reached by every path that allocated anything. The
// failure branches are all below this line for that reason.
CUDA_CHECK(cudaEventDestroy(start));
CUDA_CHECK(cudaEventDestroy(stop));
CUDA_CHECK(cudaEventDestroy(phaseStart));
CUDA_CHECK(cudaEventDestroy(phaseStop));
CUDA_CHECK(cudaFree(d_table));
CUDA_CHECK(cudaFree(d_keys));
CUDA_CHECK(cudaFree(d_out));
CUDA_CHECK(cudaFree(u_table));
CUDA_CHECK(cudaFree(u_keys));
CUDA_CHECK(cudaFree(u_out));
std::printf("Where the time goes, one 4-round workload per row\n");
std::printf(" %-28s %10s %10s %10s %10s\n", "row", "total ms", "kernel ms",
"move ms", "host ms");
for (int c = 0; c < kConfigs; ++c) {
const double other = static_cast<double>(rows[c].totalMs) -
static_cast<double>(rows[c].kernelMs) -
static_cast<double>(rows[c].moveMs);
std::printf(" %-28s %10.3f %10.3f %10.3f %10.3f\n", names[c],
static_cast<double>(rows[c].totalMs),
static_cast<double>(rows[c].kernelMs),
static_cast<double>(rows[c].moveMs), other);
}
std::printf(
"\nmove ms is cudaMemcpy for the explicit row and\n"
"cudaMemPrefetchAsync for the rows that prefetch. host ms is what is\n"
"left: the four table passes, the final read of the output, and any\n"
"migration the runtime did outside a call this program timed.\n\n");
std::printf("Kernel milliseconds per round\n");
std::printf(" %-28s", "row");
for (int r = 0; r < kRounds; ++r) {
std::printf(" %9d", r + 1);
}
std::printf("\n");
for (int c = 0; c < kConfigs; ++c) {
std::printf(" %-28s", names[c]);
for (int r = 0; r < kRounds; ++r) {
std::printf(" %9.3f", static_cast<double>(rows[c].roundMs[r]));
}
std::printf("\n");
}
std::printf(
"\nEach row is one pass, not a mean, because a page migrates once "
"and\nthe second measurement would be a different one. Read a row "
"across,\nnot a column down.\n\n");
// The 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 them is a prediction about your hardware. Two are the answer,
// one is arithmetic, and one is the property read-mostly must not break.
for (int c = 0; c < kConfigs; ++c) {
if (rows[c].totalMs <= 0.0f || rows[c].kernelMs <= 0.0f) {
std::fprintf(stderr,
"%s: a timed region measured %.3f ms total and %.3f "
"ms of kernel, so the event pair never separated\n",
names[c], static_cast<double>(rows[c].totalMs),
static_cast<double>(rows[c].kernelMs));
return EXIT_FAILURE;
}
if (badRounds[c] != 0) {
std::fprintf(stderr,
"%s: the host read a different table on %d of %d "
"rounds, so a copy of a read-mostly region went "
"stale\n",
names[c], badRounds[c], kRounds);
return EXIT_FAILURE;
}
if (badIndex[c] != kElems) {
std::fprintf(stderr, "%s: wrong at %zu, got %.9g, want %.9g\n",
names[c], badIndex[c], gotBad[c], h_want[badIndex[c]]);
return EXIT_FAILURE;
}
if (disagree[c] != kElems) {
std::fprintf(stderr,
"%s: differs from the explicit row at %zu, %.9g "
"against %.9g\n",
names[c], disagree[c], gotDiff[c],
static_cast<double>(h_first[disagree[c]]));
return EXIT_FAILURE;
}
}
std::printf("all %zu outputs match the double reference in every row\n",
kElems);
std::printf("rtol %.3g, atol %.3g, from a reduction depth of %d\n", rtol,
atol, kRounds);
std::printf("every managed row agrees with the explicit row bit for bit\n");
return EXIT_SUCCESS;
}