code/day50-checklist/mystery.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 50: three mystery kernels for the performance checklist.
//
// The names carry no information on purpose. Every other file in this course
// names a kernel for what it does to memory. This one does not, because the
// exercise is to reach a diagnosis from a profiler report rather than from
// the source, and a name like copyStrided would hand over the answer. The
// source is right here and reading it still will not tell you which of the
// five checklist steps catches which kernel. That is the lesson.
//
// Three kernels, four launches, one per failure the checklist separates:
// mysteryA scales a kDim by kDim matrix with consecutive lanes kDim
// floats apart, so every lane owns a 32-byte sector of its own.
// mysteryB runs four independent chains of fused multiply-adds behind a
// branch on the lane index, so every warp splits in half.
// mysteryC runs the same four chains with no branch, launched twice: once
// asking for 48000 bytes of dynamic shared memory it never
// reads, and once asking for none. Those two launches differ in
// occupancy and in nothing else at all.
//
// mysteryB and mysteryC run exactly kChains * kIters fused multiply-adds per
// thread and move exactly one float in and one float out per thread, so the
// only variables between them are the branch and the shared memory request.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o mystery mystery.cu
// Run: ./mystery
// Profile: sudo ncu --set full --kernel-name regex:'mystery' \
// --launch-count 4 --clock-control base \
// --export profile/day50-mystery ./mystery
//
// All four correctness launches happen before any timing, so the first four
// launches ncu matches are one of each configuration and --launch-count 4
// captures exactly the set the lesson ships.
//
// Plain ncu returns ERR_NVGPUCTRPERM on a stock driver install, because the
// driver parameter RmProfilingAdminOnly defaults to 1. sudo works where you
// have root. A hosted tier gives you neither, which is why the report lives
// in profile/ in this repository.
//
// Not yet run. When it runs, the only output that may be published as this
// program's output is the transcript in evidence/. See
// research/REVIEW-PROCESS.md.
#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)
// The matrix mysteryA walks. A power of two, so the row and the column come
// out of a mask and a shift rather than a division. Day 11's strided kernel
// paid for a 64-bit modulo that its coalesced sibling did not, and that one
// difference broke the claim that the address pattern was the only variable.
constexpr size_t kDim = 2048;
constexpr int kDimLog2 = 11;
constexpr size_t kElems = kDim * kDim;
constexpr int kThreadsPerBlock = 256; // 8 warps
constexpr int kBlocks = static_cast<int>(kElems / kThreadsPerBlock);
constexpr int kWarmupRuns = 3;
constexpr int kTimedRuns = 10;
constexpr float kRelTolerance = 1e-5f;
// Four independent chains per thread, so neither mysteryB nor mysteryC needs
// resident warps to keep the arithmetic pipe busy. That is the whole reason
// mysteryC's two launches are a fair test of what occupancy is worth here.
constexpr int kChains = 4;
constexpr int kIters = 256;
constexpr float kChainOffset = 0.25f;
// mysteryA's arithmetic: one fused multiply-add per element, so the kernel
// has a non-zero arithmetic intensity and still moves eight bytes to get it.
constexpr float kScale = 2.5f;
constexpr float kBias = 1.25f;
// Both multipliers are one plus a power of two and both addends are a power
// of two, so every step is exact in binary floating point and the host
// reference agrees with the device to the bit rather than to a tolerance.
constexpr float kMulA = 1.000244140625f; // 1 + 2^-12
constexpr float kAddA = 0.000244140625f; // 2^-12
constexpr float kMulB = 1.00048828125f; // 1 + 2^-11
constexpr float kAddB = 0.00048828125f; // 2^-11
// The input repeats with this period. Seven is odd, so it is coprime with
// the lane parity mysteryB branches on and every combination of the two
// actually occurs somewhere in the buffer.
constexpr int kInputPeriod = 7;
// 48000 bytes is under the 49152 a block gets by default on every card in
// this course's hardware matrix, and over half the 65536 a Turing SM can
// hand out, so exactly one block of mysteryC fits per SM under it. Eight
// warps out of the 32 an SM holds is 25 percent occupancy.
constexpr size_t kScratchBytes = 48000;
constexpr int kNumRows = 4;
static_assert(kThreadsPerBlock % 32 == 0,
"a block must be a whole number of warps, or the tail warp is "
"half idle for the block's whole life");
static_assert((kDim & (kDim - 1)) == 0,
"mysteryA takes the row from the low bits of the thread index, "
"which is only a permutation when kDim is a power of two");
static_assert((static_cast<size_t>(1) << kDimLog2) == kDim,
"kDimLog2 must be the base-two logarithm of kDim");
static_assert(kElems % kThreadsPerBlock == 0,
"the grid covers the buffer exactly, so no launch relies on "
"the tail guard doing anything");
static_assert(kScratchBytes <= 49152,
"48 KiB per block is the default cap; more needs an explicit "
"cudaFuncSetAttribute opt-in, which this program does not do");
static_assert(kInputPeriod % 2 == 1,
"an even period would pair one input value with one lane "
"parity forever and half the reference table would be dead");
// The input pattern, written once so the fill, the reference table and the
// page cannot disagree about it. Every value is a multiple of a half, so it
// is exact in binary floating point.
static float inputValue(size_t i) {
return static_cast<float>(i % kInputPeriod) * 0.5f;
}
// One thread scales one element of a kDim by kDim matrix.
//
// One warp: lane L takes row (base + L) and column (base >> kDimLog2), so
// the warp's 32 addresses are kDim floats apart. Each lane lands in a
// 32-byte sector of its own, and in a 128-byte line of its own.
//
// Launch assumption: none. The grid covers n exactly, so the guard never
// fires here, and it is written anyway because this kernel is where the
// exercise starts editing.
//
// The mapping is a permutation by construction: t runs over [0, kDim * kDim),
// row takes its low kDimLog2 bits and col takes the rest, so every (row, col)
// pair is produced exactly once. main() poisons the output buffer before the
// launch, so an element this kernel never writes fails the comparison rather
// than passing quietly. Day 11 shipped the version of that mistake which did
// not fail, and it reported half the traffic the hardware moved.
// snippet: mystery-a
__global__ void mysteryA(const float* __restrict__ in, float* __restrict__ out,
size_t n) {
const size_t t = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (t < n) {
const size_t row = t & (kDim - 1);
const size_t col = t >> kDimLog2;
const size_t j = row * kDim + col;
out[j] = fmaf(in[j], kScale, kBias);
}
}
// end snippet
// Four chains of fused multiply-adds, iters steps each.
//
// The chains are independent, so chain q never waits on chain q - 1 and one
// thread on its own keeps four multiply-adds in flight. That is
// instruction-level parallelism, and it is why the kernels below do not need
// resident warps to hide their own arithmetic latency.
//
// iters arrives as an argument rather than as kIters so that nvcc cannot
// unroll all 1024 multiply-adds into straight-line code and turn this into a
// measurement of the instruction cache. Day 30's FP32 ceiling kernel takes
// its iteration count the same way and for the same reason.
__device__ inline float runChains(float x, float mul, float add, int iters) {
float acc[kChains];
#pragma unroll
for (int q = 0; q < kChains; ++q) {
acc[q] = x + static_cast<float>(q) * kChainOffset;
}
for (int k = 0; k < iters; ++k) {
#pragma unroll
for (int q = 0; q < kChains; ++q) {
acc[q] = fmaf(acc[q], mul, add);
}
}
float sum = 0.0f;
#pragma unroll
for (int q = 0; q < kChains; ++q) {
sum += acc[q];
}
return sum;
}
// One thread reads one float, runs kChains * iters fused multiply-adds down
// one of two paths, and writes one float.
//
// One warp: lane L reads and writes element (base + L), so both ends are 128
// contiguous bytes and four sectors. Memory is a constant here, not the
// variable.
//
// Launch assumption: a whole number of warps per block, which the
// static_assert above covers. The predicate reads threadIdx.x directly, so
// the even and the odd lanes of every warp take different paths.
// snippet: mystery-b
__global__ void mysteryB(const float* __restrict__ in, float* __restrict__ out,
size_t n, int iters) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
const float x = (i < n) ? in[i] : 0.0f;
float y = 0.0f;
if ((threadIdx.x & 1u) == 0u) {
y = runChains(x, kMulA, kAddA, iters);
} else {
y = runChains(x, kMulB, kAddB, iters);
}
if (i < n) {
out[i] = y;
}
}
// end snippet
// The same four chains as mysteryB with no branch anywhere in them.
//
// One warp: the same two coalesced accesses, 128 bytes and four sectors at
// each end.
//
// Launch assumption: none. This kernel declares no shared memory and reads
// none. The dynamic shared memory main() launches it with is a launch
// parameter and nothing else, which is exactly what makes the two launches
// of it comparable: same instructions, same bytes, different occupancy.
// snippet: mystery-c
__global__ void mysteryC(const float* __restrict__ in, float* __restrict__ out,
size_t n, int iters) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
const float x = (i < n) ? in[i] : 0.0f;
const float y = runChains(x, kMulA, kAddA, iters);
if (i < n) {
out[i] = y;
}
}
// end snippet
// CPU reference for mysteryA. The kernel's mapping is a permutation of
// [0, n), so scaling every element in index order produces the same buffer.
// Written for obvious correctness rather than speed: a plain loop, no
// blocking, no intrinsics.
static void scaleCpu(const float* in, float* want, size_t n) {
for (size_t i = 0; i < n; ++i) {
want[i] = std::fmaf(in[i], kScale, kBias);
}
}
// CPU reference for the chains. It mirrors runChains rather than sharing it,
// because the device copy carries #pragma unroll that a host compiler has no
// use for and would warn about on every build. Both run the same fused
// multiply-adds in the same order, so they agree to the bit and
// kRelTolerance is slack rather than necessity.
static float chainsCpu(float x, float mul, float add, int iters) {
float acc[kChains];
for (int q = 0; q < kChains; ++q) {
acc[q] = x + static_cast<float>(q) * kChainOffset;
}
for (int k = 0; k < iters; ++k) {
for (int q = 0; q < kChains; ++q) {
acc[q] = std::fmaf(acc[q], mul, add);
}
}
float sum = 0.0f;
for (int q = 0; q < kChains; ++q) {
sum += acc[q];
}
return sum;
}
// Fills want[] for mysteryB from fourteen distinct answers.
//
// A chain's result depends only on its input value and its path. The input
// repeats with period kInputPeriod and the path depends only on the parity
// of the thread index inside its block, so there are kInputPeriod * 2
// distinct results and the host computes each one once. Running 4.3 billion
// fused multiply-adds on one CPU thread to check a kernel that takes about a
// millisecond is not a reference, it is a wait.
static void mysteryBCpu(float* want, size_t n, int iters) {
std::vector<float> table(static_cast<size_t>(kInputPeriod) * 2);
for (int v = 0; v < kInputPeriod; ++v) {
const float x = inputValue(static_cast<size_t>(v));
table[static_cast<size_t>(v) * 2] = chainsCpu(x, kMulA, kAddA, iters);
table[static_cast<size_t>(v) * 2 + 1] =
chainsCpu(x, kMulB, kAddB, iters);
}
for (size_t i = 0; i < n; ++i) {
const size_t tid = i % static_cast<size_t>(kThreadsPerBlock);
const size_t slot = (i % kInputPeriod) * 2 + (tid & 1u);
want[i] = table[slot];
}
}
// The same for mysteryC, which has no branch, so seven answers cover it.
static void mysteryCCpu(float* want, size_t n, int iters) {
std::vector<float> table(static_cast<size_t>(kInputPeriod));
for (int v = 0; v < kInputPeriod; ++v) {
table[static_cast<size_t>(v)] =
chainsCpu(inputValue(static_cast<size_t>(v)), kMulA, kAddA, iters);
}
for (size_t i = 0; i < n; ++i) {
want[i] = table[i % kInputPeriod];
}
}
// 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 33" names a lane, "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.
// A host clock around a launch measures the launch, because launches are
// asynchronous; day 9 takes that apart. Copy this helper verbatim.
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 row of the launch table. Everything except the name comes from the
// driver's own view of the compiled kernel and of this device, so the page
// cannot quote a register count or an occupancy the binary does not have.
// This is the whole of checklist step 5 without a profiler.
static void printLaunchRow(const char* name, const cudaFuncAttributes& attr,
size_t dynamicBytes, int blocksPerSm,
int maxThreadsPerSm) {
const int warpsPerSm = blocksPerSm * kThreadsPerBlock / 32;
const double occupancy = 100.0 * blocksPerSm * kThreadsPerBlock /
static_cast<double>(maxThreadsPerSm);
std::printf("%-26s %5d %10d %11zu %7d %8d %8.1f\n", name, attr.numRegs,
static_cast<int>(attr.sharedSizeBytes), dynamicBytes,
blocksPerSm, warpsPerSm, occupancy);
}
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("%d SMs, warp size %d, %d resident threads/SM\n",
prop.multiProcessorCount, prop.warpSize,
prop.maxThreadsPerMultiProcessor);
std::printf("shared memory: %zu B per block by default, %zu B per SM\n",
prop.sharedMemPerBlock, prop.sharedMemPerMultiprocessor);
std::printf("%zu elements, %zu MiB per buffer, %d blocks x %d threads\n",
kElems, kElems * sizeof(float) / (1024 * 1024), kBlocks,
kThreadsPerBlock);
std::printf(
"mysteryB and mysteryC each run %d fused multiply-adds "
"per thread\n\n",
kChains * kIters);
// Step 5 of the checklist, answered before anything runs. Occupancy is
// decided by the compiled kernel and the launch configuration, and both
// are knowable from the runtime API on any card, with no counters and no
// root.
cudaFuncAttributes attrA;
cudaFuncAttributes attrB;
cudaFuncAttributes attrC;
CUDA_CHECK(cudaFuncGetAttributes(&attrA, mysteryA));
CUDA_CHECK(cudaFuncGetAttributes(&attrB, mysteryB));
CUDA_CHECK(cudaFuncGetAttributes(&attrC, mysteryC));
int blocksA = 0;
int blocksB = 0;
int blocksCScratch = 0;
int blocksCPlain = 0;
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&blocksA, mysteryA, kThreadsPerBlock, 0));
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&blocksB, mysteryB, kThreadsPerBlock, 0));
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&blocksCScratch, mysteryC, kThreadsPerBlock, kScratchBytes));
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&blocksCPlain, mysteryC, kThreadsPerBlock, 0));
std::printf("Part 1: what the launch configuration alone decides\n");
std::printf("%-26s %5s %10s %11s %7s %8s %8s\n", "kernel", "regs",
"static B", "dynamic B", "blk/SM", "warps/SM", "occ %");
std::printf("%-26s %5s %10s %11s %7s %8s %8s\n",
"--------------------------", "-----", "----------",
"-----------", "-------", "--------", "--------");
printLaunchRow("mysteryA", attrA, 0, blocksA,
prop.maxThreadsPerMultiProcessor);
printLaunchRow("mysteryB", attrB, 0, blocksB,
prop.maxThreadsPerMultiProcessor);
printLaunchRow("mysteryC, 48000 B dynamic", attrC, kScratchBytes,
blocksCScratch, prop.maxThreadsPerMultiProcessor);
printLaunchRow("mysteryC, 0 B dynamic", attrC, 0, blocksCPlain,
prop.maxThreadsPerMultiProcessor);
// Every failure below records itself and falls through to the one
// cleanup block at the bottom, so no path can return with device memory
// still allocated.
int failures = 0;
int rows = 0;
const size_t bytes = kElems * sizeof(float);
float* d_in = nullptr;
float* d_out = nullptr;
CUDA_CHECK(cudaMalloc(&d_in, bytes));
CUDA_CHECK(cudaMalloc(&d_out, bytes));
std::vector<float> h_in(kElems);
std::vector<float> h_out(kElems);
std::vector<float> h_want(kElems);
for (size_t i = 0; i < kElems; ++i) {
h_in[i] = inputValue(i);
}
CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), bytes, cudaMemcpyHostToDevice));
// The correctness pass runs every configuration once, in order, before
// anything is timed. Two reasons. A number that came from a wrong kernel
// never reaches the table, and the first four launches ncu sees are one
// of each, so `--launch-count 4` captures exactly the set the lesson
// ships.
//
// 0x7F in every byte is 3.4e38, not a NaN, and that is deliberate: a NaN
// fails every comparison written as `fabs(got - want) > tol`, because a
// comparison against a NaN is false. A huge finite number fails it
// loudly. Any element a kernel forgets to write is caught here.
const char* names[kNumRows] = {"mysteryA", "mysteryB",
"mysteryC, 48000 B dynamic",
"mysteryC, 0 B dynamic"};
scaleCpu(h_in.data(), h_want.data(), kElems);
CUDA_CHECK(cudaMemset(d_out, 0x7F, bytes));
mysteryA<<<kBlocks, kThreadsPerBlock>>>(d_in, d_out, kElems);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
size_t bad =
firstMismatch(h_out.data(), h_want.data(), kElems, kRelTolerance);
if (bad != kElems) {
std::fprintf(stderr, "%s wrong at %zu: got %.9g, want %.9g\n", names[0],
bad, h_out[bad], h_want[bad]);
++failures;
}
mysteryBCpu(h_want.data(), kElems, kIters);
CUDA_CHECK(cudaMemset(d_out, 0x7F, bytes));
mysteryB<<<kBlocks, kThreadsPerBlock>>>(d_in, d_out, kElems, kIters);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
bad = firstMismatch(h_out.data(), h_want.data(), kElems, kRelTolerance);
if (bad != kElems) {
std::fprintf(stderr, "%s wrong at %zu: got %.9g, want %.9g\n", names[1],
bad, h_out[bad], h_want[bad]);
++failures;
}
// The third chevron argument is dynamic shared memory per block. It is
// the experiment on this page, so it is written out rather than left at
// its default, and the same kernel is launched again below with 0.
mysteryCCpu(h_want.data(), kElems, kIters);
CUDA_CHECK(cudaMemset(d_out, 0x7F, bytes));
mysteryC<<<kBlocks, kThreadsPerBlock, kScratchBytes>>>(d_in, d_out, kElems,
kIters);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
bad = firstMismatch(h_out.data(), h_want.data(), kElems, kRelTolerance);
if (bad != kElems) {
std::fprintf(stderr, "%s wrong at %zu: got %.9g, want %.9g\n", names[2],
bad, h_out[bad], h_want[bad]);
++failures;
}
CUDA_CHECK(cudaMemset(d_out, 0x7F, bytes));
mysteryC<<<kBlocks, kThreadsPerBlock, 0>>>(d_in, d_out, kElems, kIters);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
bad = firstMismatch(h_out.data(), h_want.data(), kElems, kRelTolerance);
if (bad != kElems) {
std::fprintf(stderr, "%s wrong at %zu: got %.9g, want %.9g\n", names[3],
bad, h_out[bad], h_want[bad]);
++failures;
}
// Only launches go inside the timed region. The allocations and the
// copies are above it, so every number below is kernel time and nothing
// else. Copies are not included and the page says so.
float ms[kNumRows] = {0.0f, 0.0f, 0.0f, 0.0f};
ms[0] = timeKernel(
[&] { mysteryA<<<kBlocks, kThreadsPerBlock>>>(d_in, d_out, kElems); });
ms[1] = timeKernel([&] {
mysteryB<<<kBlocks, kThreadsPerBlock>>>(d_in, d_out, kElems, kIters);
});
ms[2] = timeKernel([&] {
mysteryC<<<kBlocks, kThreadsPerBlock, kScratchBytes>>>(d_in, d_out,
kElems, kIters);
});
ms[3] = timeKernel([&] {
mysteryC<<<kBlocks, kThreadsPerBlock, 0>>>(d_in, d_out, kElems, kIters);
});
rows = kNumRows;
// The counts are worked from the constants at the top of this file and
// from nowhere else, so a wrong count is visible on the page rather than
// folded silently into a ratio. Bytes is the traffic the algorithm asks
// global memory for, one float in and one float out per thread, which is
// not the traffic DRAM sees once the caches have served part of it. Day
// 30 is where that gap is taken apart.
const double n = static_cast<double>(kElems);
const double chainFlops = n * kChains * kIters * 2.0;
const double flops[kNumRows] = {2.0 * n, chainFlops, chainFlops,
chainFlops};
const double moved = 2.0 * n * sizeof(float);
std::printf(
"\nPart 2: mean of %d runs after %d warm-ups, copies not "
"included\n",
kTimedRuns, kWarmupRuns);
std::printf("%-26s %10s %10s %9s %10s\n", "kernel", "FLOP/byte", "time ms",
"GB/s", "GFLOP/s");
std::printf("%-26s %10s %10s %9s %10s\n", "--------------------------",
"----------", "----------", "---------", "----------");
for (int r = 0; r < kNumRows; ++r) {
const double seconds = static_cast<double>(ms[r]) * 1.0e-3;
std::printf("%-26s %10.3f %10.3f %9.1f %10.1f\n", names[r],
flops[r] / moved, static_cast<double>(ms[r]),
moved / seconds / 1.0e9, flops[r] / seconds / 1.0e9);
}
std::printf(
"\nEvery row moved the same %.0f bytes. The two mysteryC "
"rows also ran\nthe same instructions; only the dynamic "
"shared memory request changed.\n",
moved);
// Every check below is a real branch ending in EXIT_FAILURE, never an
// assert. CI builds Release, Release defines NDEBUG, and NDEBUG deletes
// assert(), so a gate written that way would be missing from exactly the
// build that matters.
for (int r = 0; r < kNumRows; ++r) {
if (!(ms[r] > 0.0f)) {
std::fprintf(stderr,
"%s timed at %.6f ms, which is about to be "
"a denominator\n",
names[r], static_cast<double>(ms[r]));
++failures;
}
}
// The two mysteryC launches are the same kernel with the same arguments,
// so a difference in their occupancy that does not come from the shared
// memory request would mean the runtime and this program disagree about
// what a launch parameter does.
if (blocksCPlain <= blocksCScratch) {
std::fprintf(stderr,
"mysteryC fits %d block(s)/SM with %zu dynamic bytes and "
"%d with none; the request bought no occupancy back\n",
blocksCScratch, kScratchBytes, blocksCPlain);
++failures;
}
// The row count is checked so the lesson's tables and this program
// cannot drift apart on how many rows there are.
if (rows != kNumRows) {
std::fprintf(stderr,
"printed %d rows, expected %d; the lesson's tables and "
"this program disagree\n",
rows, kNumRows);
++failures;
}
CUDA_CHECK(cudaFree(d_in));
CUDA_CHECK(cudaFree(d_out));
if (failures != 0) {
std::fprintf(stderr, "%d check(s) failed\n", failures);
return EXIT_FAILURE;
}
return EXIT_SUCCESS;
}