code/day45-occupancy/occupancy.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 45: occupancy is not the goal.
//
// One kernel body, ten launch configurations. Every configuration reads the
// same two arrays, writes the same one, and moves the same bytes. Only two
// things move: how many independent loads one thread has outstanding, and how
// many warps the SM is allowed to hold.
//
// The three knobs, in the order the lesson uses them:
//
// items per thread more independent loads per warp, so fewer warps are
// needed to keep the same number of requests in flight
// dynamic shared mem reserved and never read, purely to cap residency. The
// kernel's instructions do not change, so this is the
// one knob that moves occupancy and nothing else
// __launch_bounds__ a register budget, which forces residency up and
// makes the compiler pay for it somewhere else
//
// What it does not measure. It never varies the block size: that is day 10,
// and two knobs at once means no row can be attributed to either. It says
// nothing about a compute-bound kernel, where warps wait on the math pipes
// rather than on memory and the trade looks different. And it does not
// compute the register ceiling itself, because the register allocation
// granularity is not published per compute capability; the driver's answer
// from cudaOccupancyMaxActiveBlocksPerMultiprocessor is used instead.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o occupancy occupancy.cu
// Run: ./occupancy
//
#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. Three buffers of 64 MiB, so 192 MiB moves per
// pass, which is far above the half-microsecond event resolution and far
// above launch overhead. The 611 is there so that no items-per-thread setting
// divides the element count evenly and the bounds guard runs on every launch
// instead of never.
constexpr size_t kElems = 16ull * 1024ull * 1024ull + 611ull;
// Fixed for the whole sweep. Day 10 varied this and held the rest still; this
// page holds it still and varies the rest. 256 threads is 8 warps, which
// divides the 1024-thread maximum on every compute capability in the course.
constexpr int kThreadsPerBlock = 256;
// 32 on every GPU this course targets. warpSize is a run-time built-in, so it
// cannot 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 second argument of __launch_bounds__: the number of blocks of
// kThreadsPerBlock threads the compiler must leave room for. Four blocks of
// 256 is 1024 resident threads, which fills a Turing SM, so this is the
// setting that asks for 100 percent occupancy and lets the compiler decide
// what to give up for it.
constexpr int kMinBlocksPerSm = 4;
constexpr int kWarmupRuns = 3;
constexpr int kTimedRuns = 10;
constexpr float kAlpha = 2.5f;
constexpr float kRelTolerance = 1e-5f;
// Both hold at every block size the exercise invites, including 32, which is
// the smallest a whole warp can fill.
static_assert(kThreadsPerBlock % kWarpSize == 0,
"block size must be a whole number of warps");
static_assert(kThreadsPerBlock <= 1024,
"1024 threads per block is the ceiling on every compute "
"capability this course covers");
// out[i] = a * x[i] + y[i], for kItems elements per thread, spaced one whole
// grid apart.
//
// Memory: for a fixed j, consecutive threads take consecutive elements, so
// each of the 2 * kItems loads covers 128 contiguous bytes per warp and costs
// four 32-byte sectors, exactly as day 11 measured. Spacing by a whole grid
// rather than giving each thread a contiguous run is what keeps that true;
// day 11 measured the contiguous-run version losing seven-eighths of its
// bandwidth.
//
// The two loops are separate on purpose, and that separation is the whole
// experiment. All 2 * kItems loads issue before the first multiply, so one
// warp has that many requests outstanding at once. A single loop doing
// load, multiply, store per item would let the store of item j wait on the
// load of item j, and the warp would hold one request in flight whatever
// kItems is.
//
// Launch assumption: gridDim.x * blockDim.x * kItems >= n. Any grid at or
// above that covers [0, n) exactly once, because i = base + j * step over
// base in [0, gridDim.x * blockDim.x) and j in [0, kItems) is a bijection
// onto [0, gridDim.x * blockDim.x * kItems). So the grid is free to change
// without changing the answer.
// snippet: ilp-body
template <int kItems>
__device__ void scaleAddBody(float a, const float* __restrict__ x,
const float* __restrict__ y,
float* __restrict__ out, size_t n) {
const size_t step = gridDim.x * static_cast<size_t>(blockDim.x);
const size_t base =
blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
float xv[kItems];
float yv[kItems];
#pragma unroll
for (int j = 0; j < kItems; ++j) {
const size_t i = base + static_cast<size_t>(j) * step;
xv[j] = (i < n) ? x[i] : 0.0f;
yv[j] = (i < n) ? y[i] : 0.0f;
}
#pragma unroll
for (int j = 0; j < kItems; ++j) {
const size_t i = base + static_cast<size_t>(j) * step;
if (i < n) {
out[i] = a * xv[j] + yv[j];
}
}
}
// end snippet
// The plain kernel. Whatever registers ptxas wants, it gets.
template <int kItems>
__global__ void scaleAdd(float a, const float* __restrict__ x,
const float* __restrict__ y, float* __restrict__ out,
size_t n) {
scaleAddBody<kItems>(a, x, y, out, n);
}
// The same body under a register budget. __launch_bounds__(T, B) is not a
// description of the launch; it is a promise the compiler has to make room
// for, so ptxas divides the register file by T times B and fits the kernel
// into the result, spilling to local memory if it has to. Day 17 measured
// that trade going the wrong way on a different kernel.
// snippet: bounded-kernel
template <int kItems>
__global__ __launch_bounds__(kThreadsPerBlock, kMinBlocksPerSm) void
scaleAddBounded(float a, const float* __restrict__ x,
const float* __restrict__ y, float* __restrict__ out,
size_t n) {
scaleAddBody<kItems>(a, x, y, out, n);
}
// 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 scaleAddCpu(float a, const float* x, const float* y, float* out,
size_t n) {
for (size_t i = 0; i < n; ++i) {
out[i] = a * x[i] + y[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 whole point: "wrong at 65536" 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;
}
// 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;
}
// Every configuration reads two arrays and writes one, so the figure comes
// from one constant rather than from a per-row count that could drift.
static double bandwidthGBs(float ms) {
const double bytes = 3.0 * static_cast<double>(kElems) * sizeof(float);
return bytes / (static_cast<double>(ms) * 1.0e-3) / 1.0e9;
}
// One row of the sweep.
//
// `smemBlocks` is the residency the shared-memory reservation aims at: the
// program asks the device how much shared memory that many blocks can each
// have, reserves exactly that, and never reads a byte of it. Zero means no
// reservation. The reservation is a launch parameter, not a code change, so
// two rows that differ only in this field run identical instructions.
struct Config {
const char* label;
int items;
int smemBlocks;
bool bounded;
};
// tlp = one item a thread, so latency is hidden by having many warps.
// ilp = several independent items a thread, so latency is hidden inside one
// warp. -bN caps residency at N blocks an SM with the shared-memory dial.
// -lb is the __launch_bounds__ row, which pushes residency the other way.
constexpr Config kConfigs[] = {
{"tlp1", 1, 0, false}, {"tlp1-b3", 1, 3, false},
{"tlp1-b2", 1, 2, false}, {"tlp1-b1", 1, 1, false},
{"ilp2", 2, 0, false}, {"ilp4", 4, 0, false},
{"ilp8", 8, 0, false}, {"ilp4-b1", 4, 1, false},
{"ilp8-b1", 8, 1, false}, {"ilp8-lb", 8, 0, true},
};
constexpr int kNumConfigs =
static_cast<int>(sizeof(kConfigs) / sizeof(kConfigs[0]));
// The switch in main() dispatches on items per thread, so every value in the
// table above needs a case and an instantiation. Checking that here rather
// than at run time keeps the default branch unreachable.
constexpr bool everyItemsValueInstantiated() {
for (int c = 0; c < kNumConfigs; ++c) {
const int v = kConfigs[c].items;
if (v != 1 && v != 2 && v != 4 && v != 8) {
return false;
}
}
return true;
}
static_assert(everyItemsValueInstantiated(),
"every items-per-thread value in kConfigs needs a case in the "
"switch in main() and an instantiation of scaleAdd");
// The two rows the lesson pairs off, as indices into the table above so that
// an edit to one is an edit to the other. Same loads in flight by
// construction: 32 warps holding 2 each against 8 warps holding 8 each on a
// card whose SM takes 32 warps.
constexpr int kPairManyWarps = 0; // tlp1
constexpr int kPairFewWarps = 7; // ilp4-b1
static_assert(kConfigs[kPairManyWarps].items == 1 &&
kConfigs[kPairManyWarps].smemBlocks == 0,
"kPairManyWarps must be the one-item row with no reservation");
static_assert(kConfigs[kPairFewWarps].items == 4 &&
kConfigs[kPairFewWarps].smemBlocks == 1,
"kPairFewWarps must be the four-item row capped at one block");
// Bytes of dynamic shared memory to reserve so that `target` blocks holding
// that much can share one SM's shared memory.
//
// Rounded down to 256 bytes because the hardware rounds an allocation up to a
// granularity the compute-capability tables do not publish, and rounding down
// keeps the request under the ceiling whatever that granularity turns out to
// be. Clamped to the per-block cap, which on a Turing card is 48 KiB of the
// SM's 64 KiB unless the kernel opts in with cudaFuncSetAttribute.
// snippet: smem-dial
static size_t sharedForBlocks(const cudaDeviceProp& prop, int target) {
if (target < 1) {
return 0;
}
size_t want = prop.sharedMemPerMultiprocessor / static_cast<size_t>(target);
want &= ~static_cast<size_t>(255);
if (want > prop.sharedMemPerBlock) {
want = prop.sharedMemPerBlock;
}
return want;
}
// end snippet
// One measured configuration. Collected first and printed afterwards, because
// the "fraction of the best row" column cannot be filled until every row has
// been measured.
struct Row {
int items;
int regs;
size_t localBytes;
size_t smemBytes;
int byWarps;
int byBlocks;
int byShared; // -1 when nothing is reserved, so no shared ceiling applies
int blocksPerSm;
int residentWarps;
int loadsInFlight;
float ms;
};
// Measures one configuration: asks the driver for the residency, checks the
// answer against a CPU reference at that configuration, then times it.
//
// The correctness check is repeated at every configuration rather than done
// once, because the claim the whole table rests on is that the kernel covers
// [0, n) for any grid and any items-per-thread setting. One check at one
// configuration would not test that claim. d_out is zeroed first, so a
// configuration that covered only part of the array is caught here instead of
// passing on the previous row's answer.
template <int kItems>
static Row runConfig(const Config& cfg, const cudaDeviceProp& prop,
int maxWarpsPerSm, const float* d_x, const float* d_y,
float* d_out, float* h_out, const float* h_want,
int* wrong) {
const size_t smem = sharedForBlocks(prop, cfg.smemBlocks);
const int warpsPerBlock = (kThreadsPerBlock + kWarpSize - 1) / kWarpSize;
// Two questions for the compiled kernel, not for the source: how many
// registers ptxas gave each thread, and how many blocks of this shape the
// driver will place on one SM given that register count and this shared
// memory reservation. Both answers change when the kernel body changes,
// which is why they are read rather than computed on the page.
cudaFuncAttributes attrs;
int blocksPerSm = 0;
if (cfg.bounded) {
CUDA_CHECK(cudaFuncGetAttributes(&attrs, scaleAddBounded<kItems>));
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&blocksPerSm, scaleAddBounded<kItems>, kThreadsPerBlock, smem));
} else {
// snippet: occupancy-query
CUDA_CHECK(cudaFuncGetAttributes(&attrs, scaleAdd<kItems>));
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&blocksPerSm, scaleAdd<kItems>, kThreadsPerBlock, smem));
// end snippet
}
// Enough blocks that every element is covered exactly once. This grows as
// items per thread falls, which is the point: the work is fixed and the
// shape of the launch that covers it is not.
const size_t perBlock =
static_cast<size_t>(kThreadsPerBlock) * static_cast<size_t>(kItems);
const int gridBlocks = static_cast<int>((kElems + perBlock - 1) / perBlock);
// Real branches, not asserts. 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 and let a wrong table
// reach the page with CI green.
if (attrs.maxThreadsPerBlock < kThreadsPerBlock) {
std::fprintf(stderr,
"%s: the kernel accepts at most %d threads per block but "
"the sweep launches %d; the launch bound is too tight\n",
cfg.label, attrs.maxThreadsPerBlock, kThreadsPerBlock);
++*wrong;
}
if (blocksPerSm < 1) {
std::fprintf(stderr,
"%s: no block of %d threads fits on an SM with %zu bytes "
"of shared memory reserved, so this row is not a "
"measurement\n",
cfg.label, kThreadsPerBlock, smem);
++*wrong;
}
if (blocksPerSm > prop.maxBlocksPerMultiProcessor) {
std::fprintf(stderr,
"%s: residency %d exceeds the per-SM block cap %d\n",
cfg.label, blocksPerSm, prop.maxBlocksPerMultiProcessor);
++*wrong;
}
if (blocksPerSm * warpsPerBlock > maxWarpsPerSm) {
std::fprintf(stderr,
"%s: resident warps %d exceeds the per-SM warp cap %d\n",
cfg.label, blocksPerSm * warpsPerBlock, maxWarpsPerSm);
++*wrong;
}
const size_t bytes = kElems * sizeof(float);
CUDA_CHECK(cudaMemset(d_out, 0, bytes));
if (cfg.bounded) {
scaleAddBounded<kItems><<<gridBlocks, kThreadsPerBlock, smem>>>(
kAlpha, d_x, d_y, d_out, kElems);
} else {
scaleAdd<kItems><<<gridBlocks, kThreadsPerBlock, smem>>>(
kAlpha, d_x, d_y, d_out, kElems);
}
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost));
const size_t bad = firstMismatch(h_out, h_want, kElems, kRelTolerance);
if (bad != kElems) {
std::fprintf(stderr, "%s: wrong at %zu: got %.9g, want %.9g\n",
cfg.label, bad, h_out[bad], h_want[bad]);
++*wrong;
}
float ms = 0.0f;
if (cfg.bounded) {
ms = timeKernel([&] {
scaleAddBounded<kItems><<<gridBlocks, kThreadsPerBlock, smem>>>(
kAlpha, d_x, d_y, d_out, kElems);
});
} else {
ms = timeKernel([&] {
scaleAdd<kItems><<<gridBlocks, kThreadsPerBlock, smem>>>(
kAlpha, d_x, d_y, d_out, kElems);
});
}
Row row;
row.items = kItems;
row.regs = attrs.numRegs;
row.localBytes = attrs.localSizeBytes;
row.smemBytes = smem;
row.byWarps = maxWarpsPerSm / warpsPerBlock;
row.byBlocks = prop.maxBlocksPerMultiProcessor;
row.byShared =
(smem == 0) ? -1
: static_cast<int>(prop.sharedMemPerMultiprocessor / smem);
row.blocksPerSm = blocksPerSm;
row.residentWarps = blocksPerSm * warpsPerBlock;
// Little's law, applied by counting rather than by timing: each thread
// has 2 * kItems loads outstanding before its first multiply, and there
// are this many warps on the SM to issue them. Two rows with the same
// product should keep the memory system equally busy however their
// occupancy differs, and that is the claim the table tests.
row.loadsInFlight = row.residentWarps * kItems * 2;
row.ms = ms;
return row;
}
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", prop.name, prop.major,
prop.minor);
std::printf(" SMs %d, resident warps/SM %d, resident blocks/SM %d\n",
prop.multiProcessorCount, maxWarpsPerSm,
prop.maxBlocksPerMultiProcessor);
std::printf(
" 32-bit registers/SM %d, shared/SM %zu B, shared/block %zu B "
"(%zu B opt-in)\n",
prop.regsPerMultiprocessor, prop.sharedMemPerMultiprocessor,
prop.sharedMemPerBlock, prop.sharedMemPerBlockOptin);
std::printf("n = %zu elements, %d threads per block (%d warps)\n", kElems,
kThreadsPerBlock, kThreadsPerBlock / kWarpSize);
std::printf(
"%.1f MiB moved per pass, the same for every row below\n",
3.0 * static_cast<double>(kElems) * sizeof(float) / (1024.0 * 1024.0));
std::printf("Timing: mean of %d runs after %d warm-ups, CUDA events\n\n",
kTimedRuns, kWarmupRuns);
const size_t bytes = kElems * sizeof(float);
std::vector<float> h_x(kElems);
std::vector<float> h_y(kElems);
std::vector<float> h_out(kElems);
std::vector<float> h_want(kElems);
// Small whole-ish numbers with three different periods, so a kernel that
// read the wrong array or skipped a stripe would be caught rather than
// averaging out.
for (size_t i = 0; i < kElems; ++i) {
h_x[i] = static_cast<float>(i % 97) * 0.5f;
h_y[i] = static_cast<float>(i % 13);
}
scaleAddCpu(kAlpha, h_x.data(), h_y.data(), h_want.data(), kElems);
float* d_x = nullptr;
float* d_y = nullptr;
float* d_out = nullptr;
CUDA_CHECK(cudaMalloc(&d_x, bytes));
CUDA_CHECK(cudaMalloc(&d_y, bytes));
CUDA_CHECK(cudaMalloc(&d_out, bytes));
CUDA_CHECK(cudaMemcpy(d_x, h_x.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_y, h_y.data(), bytes, cudaMemcpyHostToDevice));
Row rows[kNumConfigs];
for (int c = 0; c < kNumConfigs; ++c) {
const Config& cfg = kConfigs[c];
switch (cfg.items) {
case 1:
rows[c] =
runConfig<1>(cfg, prop, maxWarpsPerSm, d_x, d_y, d_out,
h_out.data(), h_want.data(), &wrong);
break;
case 2:
rows[c] =
runConfig<2>(cfg, prop, maxWarpsPerSm, d_x, d_y, d_out,
h_out.data(), h_want.data(), &wrong);
break;
case 4:
rows[c] =
runConfig<4>(cfg, prop, maxWarpsPerSm, d_x, d_y, d_out,
h_out.data(), h_want.data(), &wrong);
break;
case 8:
rows[c] =
runConfig<8>(cfg, prop, maxWarpsPerSm, d_x, d_y, d_out,
h_out.data(), h_want.data(), &wrong);
break;
default:
// Unreachable: everyItemsValueInstantiated() is a
// static_assert over the same table. The sentinel time is
// large so that a row which never ran cannot win the
// comparison below, and the run fails either way.
std::fprintf(stderr,
"%s: no kernel instantiated for %d items per "
"thread; add a case here and above\n",
cfg.label, cfg.items);
++wrong;
rows[c] = Row{cfg.items, 0, 0, 0, 0, 0, -1, 0, 0, 0, 1.0e30f};
break;
}
}
std::printf(
"Ceilings on residency, before the driver is asked. \"by shared\"\n"
"ignores the allocation granularity, which the compute-capability\n"
"tables do not publish; the driver's answer is the one that counts.\n"
"\n");
std::printf("%-9s %6s %6s %8s %11s %9s %10s %10s %7s\n", "label", "items",
"regs", "local B", "smem B/blk", "by warps", "by blocks",
"by shared", "driver");
for (int c = 0; c < kNumConfigs; ++c) {
const Row& row = rows[c];
char shared[16];
if (row.byShared < 0) {
std::snprintf(shared, sizeof(shared), "%s", "-");
} else {
std::snprintf(shared, sizeof(shared), "%d", row.byShared);
}
std::printf("%-9s %6d %6d %8zu %11zu %9d %10d %10s %7d\n",
kConfigs[c].label, row.items, row.regs, row.localBytes,
row.smemBytes, row.byWarps, row.byBlocks, shared,
row.blocksPerSm);
}
int best = 0;
for (int c = 1; c < kNumConfigs; ++c) {
if (rows[c].ms < rows[best].ms) {
best = c;
}
}
std::printf(
"\nMeasured. \"in flight\" is resident warps times the loads one\n"
"thread issues before its first multiply, which is 2 per item.\n\n");
std::printf("%-9s %7s %9s %6s %10s %9s %8s %8s\n", "label", "blk/SM",
"warps/SM", "occ", "in flight", "ms", "GB/s", "of best");
for (int c = 0; c < kNumConfigs; ++c) {
const Row& row = rows[c];
std::printf("%-9s %7d %9d %5.0f%% %10d %9.3f %8.1f %8.2f\n",
kConfigs[c].label, row.blocksPerSm, row.residentWarps,
100.0 * row.residentWarps / maxWarpsPerSm,
row.loadsInFlight, row.ms, bandwidthGBs(row.ms),
rows[best].ms / row.ms);
}
// The API that answers the launch question for you. It returns the block
// size reaching the highest occupancy this kernel can reach, and the
// smallest grid that fills the device at that size. What it maximises is
// occupancy; what the table above measures is time, and the two agree
// only when the kernel is short of warps.
int suggestedGrid = 0;
int suggestedBlock = 0;
CUDA_CHECK(cudaOccupancyMaxPotentialBlockSize(
&suggestedGrid, &suggestedBlock, scaleAdd<1>));
std::printf(
"\ncudaOccupancyMaxPotentialBlockSize suggests %d threads per block "
"and a grid of %d blocks for scaleAdd<1>.\n",
suggestedBlock, suggestedGrid);
std::printf("Fastest measured: %s at %.3f ms.\n", kConfigs[best].label,
rows[best].ms);
// The pair the lesson is built on: one item a thread at whatever
// occupancy the card allows, against four items a thread with residency
// dialled down to one block an SM. Different occupancy, same loads in
// flight. If their times differ by much, the claim that loads in flight
// is the quantity that matters is wrong on this card and the page has to
// say so.
const Row& many = rows[kPairManyWarps];
const Row& few = rows[kPairFewWarps];
std::printf(
"\nThe pair: %s at %.0f%% occupancy and %d loads in flight took "
"%.3f ms;\n%s at %.0f%% occupancy and %d loads in flight took %.3f "
"ms, a ratio of %.2f.\n",
kConfigs[kPairManyWarps].label,
100.0 * many.residentWarps / maxWarpsPerSm, many.loadsInFlight, many.ms,
kConfigs[kPairFewWarps].label,
100.0 * few.residentWarps / maxWarpsPerSm, few.loadsInFlight, few.ms,
few.ms / many.ms);
// This one consults the device, so it is a runtime check rather than a
// static_assert.
if (prop.warpSize != kWarpSize) {
std::fprintf(stderr,
"this GPU reports warpSize %d; every number on the page "
"assumes %d\n",
prop.warpSize, kWarpSize);
++wrong;
}
CUDA_CHECK(cudaFree(d_x));
CUDA_CHECK(cudaFree(d_y));
CUDA_CHECK(cudaFree(d_out));
if (wrong > 0) {
std::fprintf(stderr,
"\n%d check(s) failed. The tables above are not a "
"measurement of what the lesson claims.\n",
wrong);
return EXIT_FAILURE;
}
return EXIT_SUCCESS;
}