code/day23-shuffles/shuffles.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 23: warp shuffles and votes, and what the mask argument does.
//
// One warp reduces one row. A row holds kCols = 20 live values inside a
// 32-column stride, so the group is smaller than the warp and is not a power
// of two. That is the case where the mask stops being a formality.
//
// Three kernels over the same rows:
//
// reduceRowsPadded all 32 lanes participate, the lanes with no data
// carry the identity, the mask is 0xffffffff.
// Correct. Gated.
// reduceRowsShared the same sum and max through a shared tile and
// __syncwarp, so "no shared memory" on the page is a
// comparison and not a boast. Correct. Gated.
// reduceRowsBallotDown the mask computed the way NVIDIA's own warp
// primitives post computes it, with a
// __shfl_down_sync ladder under it. Every mask rule
// holds and the kernel is still wrong, because the
// ladder reads lanes the mask does not name.
// Reported, never gated.
//
// Nothing here is timed. The subject is correctness and which lane ends up
// holding the answer, not speed. Days 24 and 25 time the reduction ladder.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o shuffles shuffles.cu
// Run: ./shuffles
//
// Verified: run on a Tesla T4 (compute capability 7.5), driver 595.84,
// CUDA 12.6 (V12.6.85), on 2026-08-30. The only output that may be
// published as this program's output is the transcript in
// evidence/run-2026-08-30.txt. 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)
// 32 on every GPU this course targets. The built-in `warpSize` is a run-time
// value, so it cannot size an array or appear in a static_assert; this
// constant can do both. The program prints the driver's own answer next to
// it so the two can be compared rather than assumed equal.
constexpr int kWarpSize = 32;
// Every lane of the warp. It is the right mask here only because these
// kernels keep all 32 lanes inside every collective, which is the whole
// design choice this file exists to show.
constexpr unsigned int kFullMask = 0xffffffffu;
// 20 live values per row inside a 32-float stride. 20 is deliberately neither
// 32 nor a power of two: it is the case where a mask built from the live
// lanes is legal and a __shfl_down_sync ladder under it is still wrong.
//
// The stride is the warp width, so each row starts on a 128-byte boundary and
// one warp's 32 addresses cover exactly one row. Day 11 is why.
constexpr int kCols = 20;
constexpr int kStride = kWarpSize;
// 8205 is not a multiple of the 8 warps in a block, so the last block carries
// warps that own no row and the guard on `row` runs on every launch.
constexpr size_t kRows = 8205;
constexpr int kThreadsPerBlock = 256; // 8 warps
constexpr int kWarpsPerBlock = kThreadsPerBlock / kWarpSize;
// Below every value this program produces and still finite, so no lane can
// carry a NaN or an infinity into the max butterfly.
constexpr float kMaxIdentity = -3.0e38f;
// The threshold __any_sync votes on. The row maxima this input produces run
// from 57 to 60, so 58 splits the rows and the vote answer differs between
// them. A threshold outside that band would make the flag a constant, and a
// constant is not a check.
constexpr float kHighWater = 58.0f;
// The pad columns hold a sentinel rather than zero. No kernel reads them, and
// if a later edit widens a mask so that one does, the wrong sum is a multiple
// of 999 instead of a plausible number.
constexpr float kPadValue = 999.0f;
// Element (row, col) holds (7 * row + 3 * col) mod 61. For col in [0, 20) the
// twenty residues are distinct, because 61 is prime and 3 * 19 < 61, so no
// row is all zeros and every row sum is at least 190. Every value is a whole
// number below 61 and a row holds 20 of them, so every partial sum is under
// 1300 and exact in float. The sum gate below therefore compares bit for bit
// with no tolerance. The normalised output is a division, so it gets the
// house relative tolerance instead.
constexpr int kModulus = 61;
constexpr float kRelTolerance = 1e-5f;
// Steps a halving ladder takes over w lanes, which is log2(w). constexpr so
// the claim below is checked by the compiler. It follows only from the
// constants in this file, so a run-time check would look like a measurement
// of your GPU and be nothing of the kind.
constexpr int ladderSteps(int w) {
return (w <= 1) ? 0 : 1 + ladderSteps(w / 2);
}
static_assert(kThreadsPerBlock % kWarpSize == 0,
"block size must be a whole number of warps");
static_assert((kWarpSize & (kWarpSize - 1)) == 0,
"the halving ladder needs a power-of-two warp width");
static_assert(ladderSteps(kWarpSize) == 5,
"a 32-lane warp reduces itself in five shuffle steps");
static_assert(kCols > 0 && kCols < kWarpSize,
"the live group must be smaller than the warp, or the lesson "
"has no partial warp to talk about");
static_assert((kCols & (kCols - 1)) != 0,
"the live group must not be a power of two, or a shortened "
"ladder would accidentally be correct");
static_assert(kStride == kWarpSize,
"one warp owns one row, so the row stride is the warp width");
// One warp reduces one row of kCols live values: the sum, the max, how many
// lanes were live, and two vote answers. Every live lane then writes its own
// value divided by the row sum.
//
// What one thread does: loads one element of one row, takes part in five
// shuffle steps and two votes, and writes one normalised element.
//
// What one warp's 32 addresses look like: lanes 0 to 31 read
// in[row * 32 + lane], which is 128 contiguous bytes on a 128-byte boundary.
// One row, one warp, four 32-byte sectors.
//
// Launch assumption: blockDim.x is a whole number of warps. Nothing else.
//
// There is no __shared__ array and no __syncthreads() below. A shuffle reads
// another lane's register, so the only storage it uses is the register the
// value already lives in, and the only ordering it needs is the one the
// intrinsic performs itself.
//
// Every lane stays inside every collective and the lanes with no data carry
// the identity of the operation they are in: 0 for the sum, kMaxIdentity for
// the max, false for __any_sync, true for __all_sync. That is what makes
// kFullMask correct here. It is day 14's barrier rule in a new place: guard
// the loads and the stores, never the collective.
__global__ void reduceRowsPadded(const float* __restrict__ in,
float* __restrict__ norm,
float* __restrict__ sums,
float* __restrict__ maxes,
int* __restrict__ counts,
int* __restrict__ flags, size_t rows) {
const size_t t = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
// The grid is one-dimensional and one warp owns one row, so `row` comes
// from x here rather than from y.
const size_t row = t / kWarpSize;
const unsigned int lane = threadIdx.x % kWarpSize;
const bool live = (row < rows) && (lane < static_cast<unsigned int>(kCols));
const float value = live ? in[row * kStride + lane] : 0.0f;
// snippet: warp-sum
// Five steps, each halving the number of lanes that still hold a partial
// sum. Lane L reads lane L + offset, so after the last step lane 0 holds
// the total and no other lane does.
float sum = value;
for (unsigned int offset = kWarpSize / 2; offset > 0; offset /= 2) {
sum += __shfl_down_sync(kFullMask, sum, offset);
}
// Every lane needs the total to normalise its own value, and only lane 0
// has it, so one broadcast copies lane 0's register into all 32.
const float rowSum = __shfl_sync(kFullMask, sum, 0);
// end snippet
// snippet: warp-max
// The same five steps as a butterfly. Lane L swaps with lane L xor
// laneMask, so every lane ends up holding the max and no broadcast is
// needed. Same step count as the ladder above, different lane owns the
// answer at the end, and that is the whole difference between the two.
float best = live ? value : kMaxIdentity;
for (int laneMask = kWarpSize / 2; laneMask > 0; laneMask /= 2) {
const float other = __shfl_xor_sync(kFullMask, best, laneMask);
best = (other > best) ? other : best;
}
// end snippet
// snippet: votes
// Three questions about the whole warp, one instruction each. The lanes
// with no data vote the identity of the question they are in, which is
// how all 32 can stay named in the mask without changing an answer.
const unsigned int liveMask = __ballot_sync(kFullMask, live);
const int liveCount = __popc(liveMask);
const int anyHigh = __any_sync(kFullMask, live && value > kHighWater);
const int allPositive = __all_sync(kFullMask, !live || value > 0.0f);
// end snippet
if (live) {
norm[row * kStride + lane] = value / rowSum;
}
if (row < rows && lane == 0) {
sums[row] = rowSum;
maxes[row] = best;
counts[row] = liveCount;
flags[row] = (anyHigh != 0 ? 1 : 0) + (allPositive != 0 ? 2 : 0);
}
}
// The same row sum and row max through shared memory. This kernel exists to
// answer one question: does the shuffle version produce the same numbers with
// no shared memory and no barrier? It computes nothing else.
//
// What one thread does: writes its value into two shared slots, then reads a
// neighbour's slot once per step for five steps.
//
// What one warp's 32 addresses look like: the same 128 contiguous bytes as
// the shuffle kernel, because only the reduction changed.
//
// Launch assumption: blockDim.x is exactly kThreadsPerBlock, since both tiles
// are sized from that constant.
//
// The two __syncwarp() calls inside the loop are not decoration. Each step is
// a shared read followed by a shared write, and one barrier between steps
// leaves a race, because nothing then orders one lane's read against another
// lane's write inside the same step. That is listing 7 in NVIDIA's warp
// primitives post and listing 8 is this fix.
__global__ void reduceRowsShared(const float* __restrict__ in,
float* __restrict__ sums,
float* __restrict__ maxes, size_t rows) {
__shared__ float tileSum[kThreadsPerBlock];
__shared__ float tileMax[kThreadsPerBlock];
const size_t t = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
const size_t row = t / kWarpSize;
const unsigned int tid = threadIdx.x;
const unsigned int lane = tid % kWarpSize;
const bool live = (row < rows) && (lane < static_cast<unsigned int>(kCols));
const float value = live ? in[row * kStride + lane] : 0.0f;
tileSum[tid] = value;
tileMax[tid] = live ? value : kMaxIdentity;
__syncwarp();
for (unsigned int offset = kWarpSize / 2; offset > 0; offset /= 2) {
// A lane in the top half of the step has no partner inside its own
// warp, exactly as __shfl_down_sync leaves the upper `offset` lanes
// unchanged. It reads the identity instead of a neighbouring warp's
// slot, which would be both wrong and out of bounds at tid 240.
const bool reads = lane + offset < static_cast<unsigned int>(kWarpSize);
const float s = reads ? tileSum[tid + offset] : 0.0f;
const float m = reads ? tileMax[tid + offset] : kMaxIdentity;
__syncwarp();
tileSum[tid] += s;
tileMax[tid] = (m > tileMax[tid]) ? m : tileMax[tid];
__syncwarp();
}
if (row < rows && lane == 0) {
sums[row] = tileSum[tid];
maxes[row] = tileMax[tid];
}
}
// The trap, and the reason this page exists. The mask is computed exactly the
// way NVIDIA's warp primitives post computes it, with __ballot_sync over the
// live predicate, and then a __shfl_down_sync ladder runs under it.
//
// Every mask rule holds. Each calling lane has its own bit set, each
// non-calling lane has its bit clear, and all twenty callers pass the same
// value. The kernel is still wrong, because __shfl_down_sync reads lane
// L + offset: at offset 16, lane 4 asks for lane 20, which the mask does not
// name and which is not executing the intrinsic. The programming guide's rule
// is that a thread "may only read data from another thread that is actively
// participating in the intrinsics" and that the value is otherwise undefined.
//
// Why it is safe to ship. The rule this breaks is the source-lane rule, which
// yields an undefined value. It is not the mask rule, which is the one the
// guide documents as able to hang. Nothing here reads outside an allocation.
//
// Its answers are printed and never gated. A test that demanded a wrong
// number would be asserting that undefined behaviour is dependable, and if
// this kernel comes out right on the verification card, that is a fact about
// one driver and one compiler rather than about the code.
//
// What one thread does: loads one element and takes part in a five step
// ladder whose middle lanes read from lanes that are not running.
// What one warp's 32 addresses look like: the same 128 contiguous bytes.
// Launch assumption: blockDim.x is a whole number of warps.
__global__ void reduceRowsBallotDown(const float* __restrict__ in,
float* __restrict__ sums, size_t rows) {
const size_t t = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
const size_t row = t / kWarpSize;
const unsigned int lane = threadIdx.x % kWarpSize;
const bool live = (row < rows) && (lane < static_cast<unsigned int>(kCols));
// snippet: trap
// Every lane of the warp reaches the ballot, so kFullMask is right here.
// liveMask comes back as 0x000fffff, lanes 0 to 19, which is a legal mask
// and is what the canonical example tells you to pass.
const unsigned int liveMask = __ballot_sync(kFullMask, live);
if (live) {
float sum = in[row * kStride + lane];
for (unsigned int offset = kWarpSize / 2; offset > 0; offset /= 2) {
// Legal mask, undefined read. At offset 16 lane 4 asks for lane
// 20, which liveMask does not name and which is not running this
// line, so what comes back is whatever that lane's register
// happens to hold.
sum += __shfl_down_sync(liveMask, sum, offset);
}
if (lane == 0) {
sums[row] = sum;
}
}
// end snippet
}
// CPU reference. Written for obvious correctness, not speed: plain loops, no
// OpenMP, no intrinsics. It never allocates; the caller owns every buffer.
//
// It accumulates in double where the kernel accumulates in float, which is
// the house rule. That costs nothing here: a row sum is a whole number under
// 1300, so the conversion back to float is exact and the gate can still
// compare bit for bit.
static void reduceRowsCpu(const float* in, float* norm, float* sums,
float* maxes, int* counts, int* flags, size_t rows) {
for (size_t row = 0; row < rows; ++row) {
double total = 0.0;
float best = kMaxIdentity;
int liveCount = 0;
bool anyHigh = false;
bool allPositive = true;
for (int col = 0; col < kCols; ++col) {
const float v = in[row * kStride + col];
total += static_cast<double>(v);
if (v > best) {
best = v;
}
++liveCount;
if (v > kHighWater) {
anyHigh = true;
}
if (!(v > 0.0f)) {
allPositive = false;
}
}
const float rowSum = static_cast<float>(total);
sums[row] = rowSum;
maxes[row] = best;
counts[row] = liveCount;
flags[row] = (anyHigh ? 1 : 0) + (allPositive ? 2 : 0);
for (int col = 0; col < kStride; ++col) {
norm[row * kStride + col] =
(col < kCols) ? in[row * kStride + col] / rowSum : 0.0f;
}
}
}
// Returns how many entries of got and want differ by more than the relative
// tolerance, and writes the smallest such index into *firstBad. Pass a
// tolerance of zero for an exact comparison, which is what the sum and max
// gates use, because every value on both sides is a whole number under 1300
// and float holds those exactly.
static size_t countMismatches(const float* got, const float* want, size_t n,
float relTolerance, size_t* firstBad) {
size_t bad = 0;
*firstBad = n;
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) {
if (bad == 0) {
*firstBad = i;
}
++bad;
}
}
return bad;
}
// The same, for the two integer outputs. Kept separate rather than templated:
// modules 1 to 3 of this course use one function template in total and it is
// the timing helper, which this program does not need.
static size_t countIntMismatches(const int* got, const int* want, size_t n,
size_t* firstBad) {
size_t bad = 0;
*firstBad = n;
for (size_t i = 0; i < n; ++i) {
if (got[i] != want[i]) {
if (bad == 0) {
*firstBad = i;
}
++bad;
}
}
return bad;
}
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);
// The guide calls warpSize "a run-time value defined as the number of
// threads in a warp, commonly 32". Print the driver's answer next to the
// constant this file compiled against rather than assuming they agree.
std::printf("warpSize from the driver: %d, compiled against: %d\n",
prop.warpSize, kWarpSize);
const size_t elems = kRows * static_cast<size_t>(kStride);
const size_t bytes = elems * sizeof(float);
const size_t rowFloats = kRows * sizeof(float);
const size_t rowInts = kRows * sizeof(int);
const size_t threads = kRows * static_cast<size_t>(kWarpSize);
const int blocks =
static_cast<int>((threads + kThreadsPerBlock - 1) / kThreadsPerBlock);
const size_t warpsLaunched =
static_cast<size_t>(blocks) * static_cast<size_t>(kWarpsPerBlock);
std::printf(
"rows = %zu, %d live columns in a %d column stride, %d threads per "
"block, %d blocks\n",
kRows, kCols, kStride, kThreadsPerBlock, blocks);
std::printf("%zu warps launched, %zu of them own no row\n", warpsLaunched,
warpsLaunched - kRows);
std::vector<float> h_in(elems);
std::vector<float> h_norm(elems);
std::vector<float> h_wantNorm(elems);
std::vector<float> h_sums(kRows);
std::vector<float> h_wantSums(kRows);
std::vector<float> h_maxes(kRows);
std::vector<float> h_wantMaxes(kRows);
std::vector<int> h_counts(kRows);
std::vector<int> h_wantCounts(kRows);
std::vector<int> h_flags(kRows);
std::vector<int> h_wantFlags(kRows);
for (size_t row = 0; row < kRows; ++row) {
for (int col = 0; col < kStride; ++col) {
const size_t at = row * kStride + static_cast<size_t>(col);
if (col < kCols) {
const size_t v =
(7 * row + 3 * static_cast<size_t>(col)) % kModulus;
h_in[at] = static_cast<float>(v);
} else {
h_in[at] = kPadValue;
}
}
}
reduceRowsCpu(h_in.data(), h_wantNorm.data(), h_wantSums.data(),
h_wantMaxes.data(), h_wantCounts.data(), h_wantFlags.data(),
kRows);
float* d_in = nullptr;
float* d_norm = nullptr;
float* d_sums = nullptr;
float* d_maxes = nullptr;
int* d_counts = nullptr;
int* d_flags = nullptr;
CUDA_CHECK(cudaMalloc(&d_in, bytes));
CUDA_CHECK(cudaMalloc(&d_norm, bytes));
CUDA_CHECK(cudaMalloc(&d_sums, rowFloats));
CUDA_CHECK(cudaMalloc(&d_maxes, rowFloats));
CUDA_CHECK(cudaMalloc(&d_counts, rowInts));
CUDA_CHECK(cudaMalloc(&d_flags, rowInts));
CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), bytes, cudaMemcpyHostToDevice));
// The claim this page makes about the shuffle kernels is that they use no
// shared memory. The driver answers that, not the source, so the answer
// is printed. It is deliberately not a gate: encoding the claim the run
// exists to test would make the test unfalsifiable.
cudaFuncAttributes attrPadded;
cudaFuncAttributes attrShared;
cudaFuncAttributes attrTrap;
CUDA_CHECK(cudaFuncGetAttributes(&attrPadded, reduceRowsPadded));
CUDA_CHECK(cudaFuncGetAttributes(&attrShared, reduceRowsShared));
CUDA_CHECK(cudaFuncGetAttributes(&attrTrap, reduceRowsBallotDown));
std::printf("\nKernel attributes, from cudaFuncGetAttributes\n");
std::printf("%-22s %14s %10s %14s\n", "kernel", "shared/block", "registers",
"local/thread");
std::printf("%-22s %14zu %10d %14zu\n", "reduceRowsPadded",
attrPadded.sharedSizeBytes, attrPadded.numRegs,
attrPadded.localSizeBytes);
std::printf("%-22s %14zu %10d %14zu\n", "reduceRowsShared",
attrShared.sharedSizeBytes, attrShared.numRegs,
attrShared.localSizeBytes);
std::printf("%-22s %14zu %10d %14zu\n", "reduceRowsBallotDown",
attrTrap.sharedSizeBytes, attrTrap.numRegs,
attrTrap.localSizeBytes);
int status = EXIT_SUCCESS;
std::printf("\nCorrectness against the CPU reference, %zu rows\n", kRows);
std::printf("%-22s %11s %12s\n", "kernel", "sums wrong", "maxes wrong");
// Kernel 1: every lane in, the identity for the lanes with no data.
CUDA_CHECK(cudaMemset(d_norm, 0, bytes));
reduceRowsPadded<<<blocks, kThreadsPerBlock>>>(
d_in, d_norm, d_sums, d_maxes, d_counts, d_flags, kRows);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(h_norm.data(), d_norm, bytes, cudaMemcpyDeviceToHost));
CUDA_CHECK(
cudaMemcpy(h_sums.data(), d_sums, rowFloats, cudaMemcpyDeviceToHost));
CUDA_CHECK(
cudaMemcpy(h_maxes.data(), d_maxes, rowFloats, cudaMemcpyDeviceToHost));
CUDA_CHECK(
cudaMemcpy(h_counts.data(), d_counts, rowInts, cudaMemcpyDeviceToHost));
CUDA_CHECK(
cudaMemcpy(h_flags.data(), d_flags, rowInts, cudaMemcpyDeviceToHost));
size_t firstBad = kRows;
size_t firstMaxBad = kRows;
const size_t paddedSumsBad = countMismatches(
h_sums.data(), h_wantSums.data(), kRows, 0.0f, &firstBad);
const size_t paddedMaxesBad = countMismatches(
h_maxes.data(), h_wantMaxes.data(), kRows, 0.0f, &firstMaxBad);
if (firstMaxBad < firstBad) {
firstBad = firstMaxBad;
}
std::printf("%-22s %11zu %12zu\n", "reduceRowsPadded", paddedSumsBad,
paddedMaxesBad);
// The gate. These are real branches rather than asserts, because CI
// builds Release, Release defines NDEBUG, and NDEBUG deletes assert() out
// of exactly the build that matters.
if (paddedSumsBad != 0 || paddedMaxesBad != 0) {
std::fprintf(stderr,
"reduceRowsPadded wrong at row %zu: sum %.9g want %.9g, "
"max %.9g want %.9g\n",
firstBad, h_sums[firstBad], h_wantSums[firstBad],
h_maxes[firstBad], h_wantMaxes[firstBad]);
status = EXIT_FAILURE;
}
// Kernel 2: the same two answers through shared memory and __syncwarp.
CUDA_CHECK(cudaMemset(d_sums, 0, rowFloats));
CUDA_CHECK(cudaMemset(d_maxes, 0, rowFloats));
std::vector<float> h_sharedSums(kRows);
std::vector<float> h_sharedMaxes(kRows);
reduceRowsShared<<<blocks, kThreadsPerBlock>>>(d_in, d_sums, d_maxes,
kRows);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_sharedSums.data(), d_sums, rowFloats,
cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(h_sharedMaxes.data(), d_maxes, rowFloats,
cudaMemcpyDeviceToHost));
size_t sharedFirstBad = kRows;
size_t sharedFirstMaxBad = kRows;
const size_t sharedSumsBad = countMismatches(
h_sharedSums.data(), h_wantSums.data(), kRows, 0.0f, &sharedFirstBad);
const size_t sharedMaxesBad =
countMismatches(h_sharedMaxes.data(), h_wantMaxes.data(), kRows, 0.0f,
&sharedFirstMaxBad);
if (sharedFirstMaxBad < sharedFirstBad) {
sharedFirstBad = sharedFirstMaxBad;
}
std::printf("%-22s %11zu %12zu\n", "reduceRowsShared", sharedSumsBad,
sharedMaxesBad);
if (sharedSumsBad != 0 || sharedMaxesBad != 0) {
std::fprintf(stderr,
"reduceRowsShared wrong at row %zu: sum %.9g want %.9g, "
"max %.9g want %.9g\n",
sharedFirstBad, h_sharedSums[sharedFirstBad],
h_wantSums[sharedFirstBad], h_sharedMaxes[sharedFirstBad],
h_wantMaxes[sharedFirstBad]);
status = EXIT_FAILURE;
}
// Kernel 3: the trap. Reported, never gated, for the reason in its
// comment. It produces no max, so that column reads as a dash.
CUDA_CHECK(cudaMemset(d_sums, 0, rowFloats));
std::vector<float> h_trapSums(kRows);
reduceRowsBallotDown<<<blocks, kThreadsPerBlock>>>(d_in, d_sums, kRows);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_trapSums.data(), d_sums, rowFloats,
cudaMemcpyDeviceToHost));
size_t trapFirstBad = kRows;
const size_t trapSumsBad = countMismatches(
h_trapSums.data(), h_wantSums.data(), kRows, 0.0f, &trapFirstBad);
std::printf("%-22s %11zu %12s\n", "reduceRowsBallotDown", trapSumsBad, "-");
if (trapFirstBad != kRows) {
std::printf(
" first at row %zu: the ladder produced %.9g against "
"%.9g\n",
trapFirstBad, h_trapSums[trapFirstBad], h_wantSums[trapFirstBad]);
} else {
std::printf(
" every row came out right, which is a fact about this "
"driver and not about the code\n");
}
// The remaining outputs of kernel 1, each with its own gate.
size_t normFirstBad = elems;
const size_t normBad = countMismatches(h_norm.data(), h_wantNorm.data(),
elems, kRelTolerance, &normFirstBad);
size_t countFirstBad = kRows;
const size_t countBad = countIntMismatches(
h_counts.data(), h_wantCounts.data(), kRows, &countFirstBad);
size_t flagFirstBad = kRows;
const size_t flagBad = countIntMismatches(
h_flags.data(), h_wantFlags.data(), kRows, &flagFirstBad);
size_t crossFirstBad = kRows;
const size_t crossBad = countMismatches(h_sums.data(), h_sharedSums.data(),
kRows, 0.0f, &crossFirstBad);
std::printf(
"\nreduceRowsPadded against reduceRowsShared: %zu of %zu rows "
"differ bit for bit\n",
crossBad, kRows);
std::printf(
"normalised elements outside a %g relative tolerance: %zu of "
"%zu\n",
static_cast<double>(kRelTolerance), normBad, elems);
std::printf("__popc(__ballot_sync(...)) disagrees with %d on %zu rows\n",
kCols, countBad);
std::printf(
"__any_sync and __all_sync disagree with the reference on %zu "
"rows\n",
flagBad);
if (normBad != 0) {
std::fprintf(
stderr, "normalised output wrong at %zu: got %.9g, want %.9g\n",
normFirstBad, h_norm[normFirstBad], h_wantNorm[normFirstBad]);
status = EXIT_FAILURE;
}
if (countBad != 0) {
std::fprintf(stderr, "live count wrong at row %zu: got %d, want %d\n",
countFirstBad, h_counts[countFirstBad],
h_wantCounts[countFirstBad]);
status = EXIT_FAILURE;
}
if (flagBad != 0) {
std::fprintf(stderr, "vote flags wrong at row %zu: got %d, want %d\n",
flagFirstBad, h_flags[flagFirstBad],
h_wantFlags[flagFirstBad]);
status = EXIT_FAILURE;
}
if (crossBad != 0) {
std::fprintf(stderr,
"the shuffle and shared sums differ at row %zu: %.9g "
"against %.9g\n",
crossFirstBad, h_sums[crossFirstBad],
h_sharedSums[crossFirstBad]);
status = EXIT_FAILURE;
}
// One cleanup block that every path falls through to, so a wrong answer
// above still frees every allocation.
CUDA_CHECK(cudaFree(d_in));
CUDA_CHECK(cudaFree(d_norm));
CUDA_CHECK(cudaFree(d_sums));
CUDA_CHECK(cudaFree(d_maxes));
CUDA_CHECK(cudaFree(d_counts));
CUDA_CHECK(cudaFree(d_flags));
if (status != EXIT_SUCCESS) {
std::fprintf(stderr, "a gated kernel produced a wrong answer\n");
return status;
}
std::printf("\nevery gated kernel matched the CPU reference\n");
return EXIT_SUCCESS;
}