code/day53-pinned/pinned.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 53: pinned memory and async copies.
//
// Three measurements, one binary:
//
// 1 Host-to-device bandwidth, pageable against pinned, five sizes.
// 2 How long cudaMemcpyAsync blocks the host before returning, pageable
// against pinned, which is where the word "Async" earns or loses its
// name.
// 3 A kernel reading device-resident memory against the same kernel
// reading mapped host memory (zero-copy), fourteen launches each (one
// correctness gate, three warm-ups, ten timed), so the per-access
// PCIe cost of zero-copy is visible.
//
// This is the second file in the course that reads a host clock (day 9 is
// the first, and it exists to discredit the practice). Part 2 has no other
// instrument: the question is how long the API call blocks the calling
// thread, which no CUDA event can observe. The clock is only ever read
// after the call has returned or after an explicit cudaStreamSynchronize,
// and the comments at each read say which. Kernels are timed with events.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o pinned pinned.cu
// Run: ./pinned
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <ctime>
#include <vector>
#include <cuda_runtime.h>
#include <nvtx3/nvToolsExt.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)
constexpr size_t kMiB = 1024ull * 1024ull;
// The sweep stops at 256 MiB because the pinned buffer has to be allocated
// at the largest size, and pinned memory is a scarce resource: these pages
// are locked away from the OS for the life of the program. Day 9 used
// 64 MiB buffers; the probe below reuses that size so the two lessons'
// numbers sit on the same scale.
constexpr size_t kSweepMiB[] = {1, 4, 16, 64, 256};
constexpr int kNumSweep = 5;
constexpr size_t kMaxSweepBytes = 256 * kMiB;
constexpr size_t kProbeBytes = 64 * kMiB;
// 611 is 13 x 47, so no block size divides it and the kernel's bounds check
// runs on every launch. 16 Mi floats is 64 MiB, the probe size again.
constexpr size_t kKernelElems = 16ull * 1024ull * 1024ull + 611ull;
constexpr int kThreadsPerBlock = 256; // 8 warps
constexpr int kWarmupRuns = 3;
constexpr int kTimedRuns = 10;
constexpr float kScale = 2.0f;
constexpr float kRelTolerance = 1e-5f;
static_assert(sizeof(kSweepMiB) / sizeof(kSweepMiB[0]) == kNumSweep,
"kNumSweep must count the sweep sizes");
static_assert(kThreadsPerBlock % 32 == 0,
"block size must be a whole number of warps");
static_assert(kProbeBytes <= kMaxSweepBytes,
"the probe reuses the sweep's device buffer");
static_assert(kKernelElems * sizeof(float) <= kMaxSweepBytes,
"part 3 stages the mapped buffer into the sweep's device "
"buffer");
static_assert(kKernelElems % kThreadsPerBlock != 0,
"n must not divide evenly, or the bounds check never runs");
// The course's host clock, copied from day 9 with the same defence:
// CLOCK_MONOTONIC does not jump when the system clock is adjusted. Part 2
// measures host-side blocking, which is the one thing a device-side event
// cannot see.
static double hostMs() {
struct timespec ts;
clock_gettime(CLOCK_MONOTONIC, &ts);
return static_cast<double>(ts.tv_sec) * 1.0e3 +
static_cast<double>(ts.tv_nsec) * 1.0e-6;
}
// out[i] = kScale * in[i]. One thread owns one element.
//
// Memory: consecutive threads read consecutive floats, so one warp's 32
// addresses cover 128 contiguous bytes. That coalescing is load-bearing in
// part 3: mapped host memory is not cached on the GPU, so every sector a
// warp touches through the mapped pointer is a separate trip across PCIe,
// and nothing amortises a sloppy pattern across launches.
//
// Launch assumption: gridDim.x * blockDim.x >= n.
// snippet: kernel
__global__ void scaleCopy(const float* __restrict__ in, float* __restrict__ out,
size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
out[i] = kScale * in[i];
}
}
// end snippet
// CPU reference. Written for obvious correctness, not speed: plain loop, no
// OpenMP, no intrinsics. It never allocates; the caller owns every buffer.
static void scaleCopyCpu(const float* in, float* out, size_t n) {
for (size_t i = 0; i < n; ++i) {
out[i] = kScale * in[i];
}
}
// Returns the first index where got and want differ by more than the
// relative tolerance, or n if they agree everywhere.
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;
}
// Times a launch with CUDA events and returns the mean milliseconds per run.
// The lambda may hold a kernel launch or a blocking cudaMemcpy: both are
// work on the default stream, so the two event records bracket them either
// way. Warm-up inside, because lazy module loading makes each kernel's
// first launch pay its own load.
template <typename LaunchFn>
static float timeKernel(LaunchFn launch) {
cudaEvent_t start, stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));
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;
}
static double gbPerS(size_t bytes, double ms) {
return static_cast<double>(bytes) / (ms * 1.0e-3) / 1.0e9;
}
int main() {
// cudaDeviceMapHost has to be set before the context exists, so this is
// the first CUDA call in the program. It asks the driver to make pinned
// allocations mappable into the device's address space, which part 3
// needs and parts 1 and 2 ignore.
CUDA_CHECK(cudaSetDeviceFlags(cudaDeviceMapHost));
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);
if (prop.canMapHostMemory == 0) {
std::fprintf(stderr,
"this device cannot map host memory; part 3 is "
"impossible here\n");
return EXIT_FAILURE;
}
const size_t maxElems = kMaxSweepBytes / sizeof(float);
const size_t kernelBytes = kKernelElems * sizeof(float);
// Pageable host memory: an ordinary allocation the OS may page out,
// which is exactly why the driver cannot DMA from it directly.
std::vector<float> h_pageable(maxElems);
std::vector<float> h_back(maxElems);
std::vector<float> h_out(kKernelElems);
std::vector<float> h_want(kKernelElems);
for (size_t i = 0; i < maxElems; ++i) {
h_pageable[i] = static_cast<float>(i % 977) * 0.25f;
}
// Pinned host memory for parts 1 and 2. Same size as the pageable
// buffer, same fill, so every copy comparison moves identical bytes.
float* h_pinned = nullptr;
CUDA_CHECK(cudaMallocHost(&h_pinned, kMaxSweepBytes));
for (size_t i = 0; i < maxElems; ++i) {
h_pinned[i] = h_pageable[i];
}
// Mapped pinned memory for part 3. cudaHostGetDevicePointer hands back
// the address a kernel dereferences to reach these host pages.
// snippet: zero-copy-setup
float* h_mapped = nullptr;
float* d_mapped = nullptr;
CUDA_CHECK(cudaHostAlloc(&h_mapped, kernelBytes, cudaHostAllocMapped));
CUDA_CHECK(cudaHostGetDevicePointer(&d_mapped, h_mapped, 0));
// end snippet
for (size_t i = 0; i < kKernelElems; ++i) {
h_mapped[i] = static_cast<float>(i % 353) - 176.0f;
}
float* d_buf = nullptr;
float* d_out = nullptr;
CUDA_CHECK(cudaMalloc(&d_buf, kMaxSweepBytes));
CUDA_CHECK(cudaMalloc(&d_out, kernelBytes));
cudaStream_t stream;
CUDA_CHECK(cudaStreamCreate(&stream));
// Every return path below this line frees everything above it.
auto freeAll = [&] {
CUDA_CHECK(cudaStreamDestroy(stream));
CUDA_CHECK(cudaFree(d_out));
CUDA_CHECK(cudaFree(d_buf));
CUDA_CHECK(cudaFreeHost(h_mapped));
CUDA_CHECK(cudaFreeHost(h_pinned));
};
// Gate: both copy paths actually copy, checked at the full 256 MiB
// before anything is timed. A benchmark of a copy that did not happen
// is day 11's oldest trap.
const char* names[2] = {"pageable", "pinned"};
const float* srcs[2] = {h_pageable.data(), h_pinned};
for (int s = 0; s < 2; ++s) {
CUDA_CHECK(
cudaMemcpy(d_buf, srcs[s], kMaxSweepBytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(h_back.data(), d_buf, kMaxSweepBytes,
cudaMemcpyDeviceToHost));
const size_t bad =
firstMismatch(h_back.data(), srcs[s], maxElems, 0.0f);
if (bad != maxElems) {
std::fprintf(stderr,
"%s H2D copy wrong at %zu: got %.9g, "
"want %.9g\n",
names[s], bad, h_back[bad], srcs[s][bad]);
freeAll();
return EXIT_FAILURE;
}
}
// Part 1. Blocking cudaMemcpy, both sources, five sizes. The events
// bracket ten copies on the default stream; timeKernel divides by ten.
std::printf("\nPart 1: host-to-device bandwidth, cudaMemcpy\n");
std::printf("mean of %d runs after %d warm-ups per cell\n", kTimedRuns,
kWarmupRuns);
std::printf("%9s %21s %21s %7s\n", "size", "pageable", "pinned",
"ratio");
for (int sz = 0; sz < kNumSweep; ++sz) {
const size_t bytes = kSweepMiB[sz] * kMiB;
double ms[2] = {0.0, 0.0};
for (int s = 0; s < 2; ++s) {
char label[64];
std::snprintf(label, sizeof(label), "h2d-%s-%zuMiB", names[s],
kSweepMiB[sz]);
nvtxRangePushA(label);
ms[s] = timeKernel([&] {
CUDA_CHECK(
cudaMemcpy(d_buf, srcs[s], bytes, cudaMemcpyHostToDevice));
});
nvtxRangePop();
}
std::printf(
"%6zu MiB %9.3f ms %5.1f GB/s %9.3f ms %5.1f GB/s"
" %6.2fx\n",
kSweepMiB[sz], ms[0], gbPerS(bytes, ms[0]), ms[1],
gbPerS(bytes, ms[1]), ms[0] / ms[1]);
}
// Part 2. The same 64 MiB copy issued with cudaMemcpyAsync on a
// non-default stream. Two durations per source: how long the call
// blocked the host, and how long the copy took to complete. For the
// pageable source the runtime stages through a pinned bounce buffer
// and the call blocks while it does; for the pinned source the call
// queues a DMA and returns.
std::printf("\nPart 2: cudaMemcpyAsync, 64.0 MiB, call vs completion\n");
std::printf("mean of %d runs after 1 warm-up per source\n", kTimedRuns);
// snippet: async-probe
for (int s = 0; s < 2; ++s) {
CUDA_CHECK(cudaMemcpyAsync(d_buf, srcs[s], kProbeBytes,
cudaMemcpyHostToDevice, stream));
CUDA_CHECK(cudaStreamSynchronize(stream)); // warm-up, not timed
char label[64];
std::snprintf(label, sizeof(label), "async-%s", names[s]);
nvtxRangePushA(label);
double callSum = 0.0;
double totalSum = 0.0;
for (int r = 0; r < kTimedRuns; ++r) {
const double t0 = hostMs();
CUDA_CHECK(cudaMemcpyAsync(d_buf, srcs[s], kProbeBytes,
cudaMemcpyHostToDevice, stream));
// The call has returned; how long did it hold the thread?
const double t1 = hostMs();
CUDA_CHECK(cudaStreamSynchronize(stream));
// The copy is done; only now may the clock judge the copy.
const double t2 = hostMs();
callSum += t1 - t0;
totalSum += t2 - t0;
}
nvtxRangePop();
const double callMs = callSum / kTimedRuns;
const double totalMs = totalSum / kTimedRuns;
std::printf(
"%9s call %8.3f ms completion %8.3f ms "
"call is %5.1f%%\n",
names[s], callMs, totalMs, 100.0 * callMs / totalMs);
}
// end snippet
// Part 3. One kernel, two input pointers. The resident variant reads
// device memory that one staging copy filled; the zero-copy variant
// dereferences the mapped host pages, so all ten timed launches and
// all three warm-ups cross PCIe again.
CUDA_CHECK(
cudaMemcpy(d_buf, h_mapped, kernelBytes, cudaMemcpyHostToDevice));
const int blocks = static_cast<int>((kKernelElems + kThreadsPerBlock - 1) /
kThreadsPerBlock);
scaleCopyCpu(h_mapped, h_want.data(), kKernelElems);
const float* ins[2] = {d_buf, d_mapped};
const char* variants[2] = {"device-resident", "zero-copy"};
float kernelMs[2] = {0.0f, 0.0f};
for (int v = 0; v < 2; ++v) {
scaleCopy<<<blocks, kThreadsPerBlock>>>(ins[v], d_out, kKernelElems);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, kernelBytes,
cudaMemcpyDeviceToHost));
const size_t bad = firstMismatch(h_out.data(), h_want.data(),
kKernelElems, kRelTolerance);
if (bad != kKernelElems) {
std::fprintf(stderr, "%s wrong at %zu: got %.9g, want %.9g\n",
variants[v], bad, h_out[bad], h_want[bad]);
freeAll();
return EXIT_FAILURE;
}
nvtxRangePushA(variants[v]);
kernelMs[v] = timeKernel([&] {
scaleCopy<<<blocks, kThreadsPerBlock>>>(ins[v], d_out,
kKernelElems);
});
nvtxRangePop();
}
std::printf("\nPart 3: scaleCopy over %zu floats (%.1f MiB read)\n",
kKernelElems,
static_cast<double>(kernelBytes) / (1024.0 * 1024.0));
std::printf("mean of %d launches after %d warm-ups per variant\n",
kTimedRuns, kWarmupRuns);
for (int v = 0; v < 2; ++v) {
std::printf("%16s %8.3f ms %6.1f GB/s effective\n", variants[v],
kernelMs[v],
gbPerS(2 * kKernelElems * sizeof(float), kernelMs[v]));
}
std::printf(
"zero-copy costs %.2fx the resident kernel; its read "
"crossed PCIe on every one of the %d launches (correctness "
"gate, warm-ups and timed runs), the resident read crossed "
"once, in the staging copy\n",
kernelMs[1] / kernelMs[0], kTimedRuns + kWarmupRuns + 1);
freeAll();
return EXIT_SUCCESS;
}