code/day28-cooperative-groups/cooperative_reduce.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 28: cooperative groups and grid-wide sync.
//
// Sums one array two ways that do identical arithmetic on an identical grid
// and differ only in how the second stage is reached: a second kernel
// launch, or cooperative_groups::grid_group::sync() inside one launch.
//
// It also asks the driver the three questions the lesson is built on. How
// many blocks of this kernel fit on the device at once. What
// cg::this_grid().is_valid() reports under each kind of launch. And what the
// runtime does when a cooperative launch asks for one block more than fits.
//
// What this program will not do: run a grid barrier that might not complete.
// The oversized launch uses a probe kernel with no barrier in it. A grid
// barrier in a grid the device cannot hold has no clean failure, and a wedged
// GPU is not a teachable result. Day 14 made the same call about a
// non-uniform __syncthreads().
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o cooperative_reduce \
// cooperative_reduce.cu
// Run: ./cooperative_reduce
//
// 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 <cooperative_groups.h>
#include <cooperative_groups/reduce.h>
#include <cuda_runtime.h>
// One short alias for one long namespace, which is the spelling NVIDIA's own
// documentation and samples use. It is not `using namespace`: every group
// type below still says where it came from.
namespace cg = cooperative_groups;
// 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.
#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)
// kElems is 2^22 plus 611. The 611 makes the last step of every grid-stride
// loop partial, so the loop condition is exercised on every launch instead
// of never.
//
// The input pattern is (i % 8) * 0.25f, so every value and every partial sum
// is a multiple of 0.25. The static_assert below checks that the exact total
// times four stays at or below 2^24, the largest integer a float holds
// exactly. That makes every reduction order bit-exact, so kRelTolerance is
// slack rather than necessity and a mismatch is an indexing bug.
constexpr size_t kElems = 4194304ull + 611ull;
constexpr int kThreadsPerBlock = 256; // 8 warps
constexpr int kWarpSize = 32; // 32 on every GPU this course targets
constexpr int kWarpsPerBlock = kThreadsPerBlock / kWarpSize;
constexpr int kWarmupRuns = 3;
constexpr int kTimedRuns = 10;
constexpr float kRelTolerance = 1e-6f;
constexpr int kReportThreads = 64; // exactly two warps
// The exact sum of the input pattern, times four, as an integer. Every input
// is a multiple of 0.25, so this is the count of quarter units in the total.
// constexpr because the static_assert below calls it, and a static_assert may
// only call a constexpr function.
constexpr long long inputSumTimesFour(size_t n) {
return static_cast<long long>(n / 8) * 28 +
(n % 8 == 0 ? 0LL
: static_cast<long long>(n % 8) *
(static_cast<long long>(n % 8) - 1) / 2);
}
static_assert(kThreadsPerBlock % kWarpSize == 0,
"the block must be a whole number of 32-lane tiles");
static_assert(kThreadsPerBlock <= 1024,
"1024 threads per block is the ceiling on every compute "
"capability this course covers");
static_assert(kWarpsPerBlock <= kWarpSize,
"tile 0 reduces one value per tile in a single pass, so the "
"block may hold no more than 32 tiles");
static_assert(kElems % kThreadsPerBlock != 0,
"pick a size no block size divides, so the tail of the "
"grid-stride loop runs");
static_assert(inputSumTimesFour(kElems) <= (1LL << 24),
"the exact total must stay inside the range where a float "
"represents every quarter unit, or the reduction is no longer "
"order-independent and the tolerance starts doing real work");
static_assert(kReportThreads % kWarpSize == 0,
"the coalesced-group report reads one row per warp");
// Sums one value per thread across the whole block.
//
// What one thread does: hands its value to a 32-lane tile reduction, then one
// lane per tile parks the tile total in shared memory and tile 0 reduces
// those.
//
// What one warp's 32 addresses look like: the only global traffic is in the
// callers. Here a tile writes one float per 32 lanes to shared memory and
// tile 0 reads kWarpsPerBlock contiguous floats back, which is one bank per
// lane and no conflict.
//
// Launch assumption: exactly kThreadsPerBlock threads per block, and every
// thread in the block calls this. cg::reduce is a collective: leave one
// thread out and the result is undefined, the same contract __syncthreads()
// has.
//
// The total is returned in the threads of tile 0 and zero elsewhere. Callers
// store it from thread 0, which is in tile 0.
//
// The group handles arrive by const reference because the guide says to:
// "It is recommended that you pass group handles by reference to functions."
__device__ float blockSum(const cg::thread_block& block,
const cg::thread_block_tile<32>& tile, float v) {
// One shared array per kernel, not per call. Two calls inside one kernel
// are safe only because a barrier separates them, which is true of the
// cooperative kernel below and is why this comment exists.
__shared__ float tileSums[kWarpsPerBlock];
v = cg::reduce(tile, v, cg::plus<float>());
if (tile.thread_rank() == 0) {
tileSums[tile.meta_group_rank()] = v;
}
block.sync();
float total = 0.0f;
if (tile.meta_group_rank() == 0) {
// Uniform across tile 0, so all 32 lanes reach the reduce below.
// thread_rank() returns unsigned long long, so the cap is cast
// rather than compared across signedness.
const unsigned long long tiles =
static_cast<unsigned long long>(kWarpsPerBlock);
const float partial =
(tile.thread_rank() < tiles) ? tileSums[tile.thread_rank()] : 0.0f;
total = cg::reduce(tile, partial, cg::plus<float>());
}
return total;
}
// Sums this block's share of `in` into out[blockIdx.x]. The two-launch path
// calls it twice: once over the input, once over the partials it produced.
//
// What one thread does: walks the array with a grid stride, adding into a
// register, then joins the block reduction.
//
// What one warp's 32 addresses look like: consecutive lanes hold consecutive
// elements at every step of the loop, so each load is 128 contiguous bytes
// and costs four 32-byte sectors. Day 11 measures what that is worth.
//
// Launch assumption: kThreadsPerBlock threads per block, and out has at least
// gridDim.x elements. No early return anywhere above block.sync().
__global__ void reduceBlocks(const float* __restrict__ in,
float* __restrict__ out, size_t n) {
cg::thread_block block = cg::this_thread_block();
cg::thread_block_tile<32> tile = cg::tiled_partition<32>(block);
const size_t step = gridDim.x * static_cast<size_t>(blockDim.x);
float v = 0.0f;
for (size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
i < n; i += step) {
v += in[i];
}
const float total = blockSum(block, tile, v);
if (block.thread_rank() == 0) {
out[blockIdx.x] = total;
}
}
// The same sum in one launch. Every block writes its partial, the whole grid
// meets at grid.sync(), then block 0 reduces the partials.
//
// What one thread does: the same grid-stride walk as reduceBlocks, then it
// waits for every other block in the grid, then it either helps block 0 with
// the partials or it is finished.
//
// What one warp's 32 addresses look like: identical to reduceBlocks, because
// grid.thread_rank() and grid.num_threads() are the cooperative-groups
// spellings of the two lines above. The names change, the addresses do not.
//
// Launch assumption: a cooperative launch. grid.sync() is only a barrier when
// the runtime has promised every block is resident, and the only way to get
// that promise is cudaLaunchKernelEx with cudaLaunchAttributeCooperative set,
// or cudaLaunchCooperativeKernel. Under <<<>>> this kernel compiles, links
// and is undefined.
// snippet: reduce-cooperative
__global__ void reduceCooperativeGrid(const float* __restrict__ in,
float* __restrict__ partials,
float* __restrict__ out, size_t n) {
cg::grid_group grid = cg::this_grid();
cg::thread_block block = cg::this_thread_block();
cg::thread_block_tile<32> tile = cg::tiled_partition<32>(block);
const size_t step = grid.num_threads();
float v = 0.0f;
for (size_t i = grid.thread_rank(); i < n; i += step) {
v += in[i];
}
float total = blockSum(block, tile, v);
if (block.thread_rank() == 0) {
partials[blockIdx.x] = total;
}
// Every write above is visible to every thread below. That is the second
// half of the sync() contract and it is what replaces the kernel
// boundary the two-launch version pays for.
grid.sync();
if (blockIdx.x == 0) {
const size_t m = grid.num_blocks();
float w = 0.0f;
for (size_t i = block.thread_rank(); i < m; i += block.num_threads()) {
w += partials[i];
}
total = blockSum(block, tile, w);
if (block.thread_rank() == 0) {
out[0] = total;
}
}
}
// end snippet
// Reports what cg::this_grid() knows about the launch that produced it, and
// synchronizes nothing.
//
// What one thread does: nothing, unless it is thread 0 of the grid, which
// writes two words.
//
// What one warp's 32 addresses look like: one lane writes eight bytes and the
// other 31 write nothing. This kernel is a question, not a workload.
//
// Launch assumption: none, deliberately. It carries no barrier, so it is safe
// to launch with a grid too large to be co-resident, which is the whole
// reason it exists.
__global__ void probeCooperative(unsigned int* __restrict__ report) {
cg::grid_group grid = cg::this_grid();
if (grid.thread_rank() == 0) {
report[0] = grid.is_valid() ? 1u : 0u;
report[1] = static_cast<unsigned int>(grid.num_blocks());
}
}
// Records what cg::coalesced_threads() returns inside a data-dependent
// branch, so the page can print the answer instead of asserting it.
//
// What one thread does: reads its flag, and if the flag is set, names the
// group of lanes that took the same branch and writes that group's size and
// its own rank in it.
//
// What one warp's 32 addresses look like: 32 consecutive ints in, two arrays
// of 32 consecutive words out. Coalesced, and irrelevant: nothing here is
// timed.
//
// Launch assumption: one block of exactly n threads, n a multiple of 32.
__global__ void reportCoalesced(const int* __restrict__ take,
unsigned int* __restrict__ groupSize,
unsigned int* __restrict__ groupRank,
size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
if (take[i] != 0) {
cg::coalesced_group active = cg::coalesced_threads();
groupSize[i] = static_cast<unsigned int>(active.num_threads());
groupRank[i] = static_cast<unsigned int>(active.thread_rank());
} else {
groupSize[i] = 0u;
groupRank[i] = 0u;
}
}
}
// CPU reference. Written for obvious correctness, not speed: plain loop, no
// OpenMP, no intrinsics. It never allocates; the caller owns every buffer.
// It accumulates in double even though the kernels accumulate in float,
// because the reference's job is to be right.
static double sumCpu(const float* x, size_t n) {
double total = 0.0;
for (size_t i = 0; i < n; ++i) {
total += static_cast<double>(x[i]);
}
return total;
}
// 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;
}
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), %d SMs\n", prop.name,
prop.major, prop.minor, prop.multiProcessorCount);
// The support question comes before the sizing question. This attribute
// is 0 on a Windows configuration that cannot do it and on any device
// below compute capability 6.0, and the failure it produces otherwise is
// not obviously about cooperative launches at all.
int supportsCoop = 0;
CUDA_CHECK(cudaDeviceGetAttribute(&supportsCoop,
cudaDevAttrCooperativeLaunch, device));
std::printf("cudaDevAttrCooperativeLaunch: %d\n", supportsCoop);
if (supportsCoop == 0) {
std::fprintf(stderr,
"this device or platform does not support cooperative "
"launches, so there is nothing for this program to "
"measure\n");
return EXIT_FAILURE;
}
// How many blocks of each kernel the device holds at once. The cap is a
// property of the kernel, not of the card, because it is decided by the
// registers and the shared memory that kernel needs, so both kernels are
// asked separately.
int coopBlocksPerSM = 0;
int probeBlocksPerSM = 0;
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&coopBlocksPerSM, reduceCooperativeGrid, kThreadsPerBlock, 0));
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&probeBlocksPerSM, probeCooperative, kThreadsPerBlock, 0));
const int maxCoopBlocks = coopBlocksPerSM * prop.multiProcessorCount;
const int maxProbeBlocks = probeBlocksPerSM * prop.multiProcessorCount;
if (maxCoopBlocks < 1 || maxProbeBlocks < 1) {
std::fprintf(stderr, "occupancy query returned no resident blocks\n");
return EXIT_FAILURE;
}
const size_t bytes = kElems * sizeof(float);
const int blocksNeeded =
static_cast<int>((kElems + kThreadsPerBlock - 1) / kThreadsPerBlock);
const int blocks =
(blocksNeeded < maxCoopBlocks) ? blocksNeeded : maxCoopBlocks;
std::printf("n = %zu floats, %.1f MiB, %d threads/block\n\n", kElems,
static_cast<double>(bytes) / (1024.0 * 1024.0),
kThreadsPerBlock);
std::printf("How big a cooperative grid fits\n");
std::printf(" reduceCooperativeGrid %d blocks/SM x %d SMs = %d\n",
coopBlocksPerSM, prop.multiProcessorCount, maxCoopBlocks);
std::printf(" probeCooperative %d blocks/SM x %d SMs = %d\n",
probeBlocksPerSM, prop.multiProcessorCount, maxProbeBlocks);
std::printf(" a <<<>>> launch sized to n would use %d blocks\n",
blocksNeeded);
std::printf(" both sums below run %d blocks, the same grid\n\n", blocks);
std::vector<float> h_in(kElems);
for (size_t i = 0; i < kElems; ++i) {
h_in[i] = static_cast<float>(i % 8) * 0.25f;
}
std::vector<int> h_take(kReportThreads, 0);
// Warp 0 activates lanes 2, 4 and 8, which is the case the programming
// guide's own coalesced_threads example describes. Warp 1 activates its
// lower half, so the two warps cannot report the same group size unless
// something is very wrong.
h_take[2] = 1;
h_take[4] = 1;
h_take[8] = 1;
for (int t = kWarpSize; t < kWarpSize + 16; ++t) {
h_take[t] = 1;
}
float* d_in = nullptr;
float* d_partials = nullptr;
float* d_out = nullptr;
unsigned int* d_report = nullptr;
int* d_take = nullptr;
unsigned int* d_size = nullptr;
unsigned int* d_rank = nullptr;
CUDA_CHECK(cudaMalloc(&d_in, bytes));
CUDA_CHECK(
cudaMalloc(&d_partials, static_cast<size_t>(blocks) * sizeof(float)));
CUDA_CHECK(cudaMalloc(&d_out, sizeof(float)));
CUDA_CHECK(cudaMalloc(&d_report, 2 * sizeof(unsigned int)));
CUDA_CHECK(cudaMalloc(&d_take, kReportThreads * sizeof(int)));
CUDA_CHECK(cudaMalloc(&d_size, kReportThreads * sizeof(unsigned int)));
CUDA_CHECK(cudaMalloc(&d_rank, kReportThreads * sizeof(unsigned int)));
CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_take, h_take.data(), kReportThreads * sizeof(int),
cudaMemcpyHostToDevice));
// Every failure below increments this instead of returning, so one run
// reports every problem it found and every allocation is still freed by
// the single cleanup block at the end. These are real branches, not
// asserts: CI builds Release, Release defines NDEBUG, and NDEBUG deletes
// assert(), so a gate written that way would not exist in the build that
// matters.
int wrong = 0;
// The launch configuration this whole lesson is about. The grid, block,
// shared-memory and stream fields are what <<<>>> already gives you. The
// attribute array is the part it cannot express.
// snippet: launch-config
cudaLaunchAttribute coopAttr[1] = {};
coopAttr[0].id = cudaLaunchAttributeCooperative;
coopAttr[0].val.cooperative = 1;
cudaLaunchConfig_t config = {};
config.gridDim = dim3(static_cast<unsigned int>(blocks));
config.blockDim = dim3(static_cast<unsigned int>(kThreadsPerBlock));
config.dynamicSmemBytes = 0;
config.stream = nullptr;
config.attrs = coopAttr;
config.numAttrs = 1;
// end snippet
std::printf("What cg::this_grid() reports\n");
std::vector<unsigned int> h_report(2, 0u);
probeCooperative<<<blocks, kThreadsPerBlock>>>(d_report);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_report.data(), d_report, 2 * sizeof(unsigned int),
cudaMemcpyDeviceToHost));
std::printf(" <<<%d, %d>>> is_valid=%u blocks=%u\n",
blocks, kThreadsPerBlock, h_report[0], h_report[1]);
cudaLaunchConfig_t probeConfig = config;
probeConfig.gridDim = dim3(static_cast<unsigned int>(maxProbeBlocks));
CUDA_CHECK(cudaLaunchKernelEx(&probeConfig, probeCooperative, d_report));
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_report.data(), d_report, 2 * sizeof(unsigned int),
cudaMemcpyDeviceToHost));
std::printf(" cooperative, %d blocks is_valid=%u blocks=%u\n",
maxProbeBlocks, h_report[0], h_report[1]);
// One block more than the device can hold at once. The runtime knows the
// number, so this is the case where it can refuse rather than misbehave.
probeConfig.gridDim = dim3(static_cast<unsigned int>(maxProbeBlocks + 1));
const cudaError_t tooBig =
cudaLaunchKernelEx(&probeConfig, probeCooperative, d_report);
std::printf(" cooperative, %d blocks %s: %s\n", maxProbeBlocks + 1,
cudaGetErrorName(tooBig), cudaGetErrorString(tooBig));
if (tooBig == cudaSuccess) {
std::fprintf(stderr,
"a cooperative launch of %d blocks was accepted, but "
"the documented cap for this kernel is %d\n",
maxProbeBlocks + 1, maxProbeBlocks);
++wrong;
CUDA_CHECK(cudaDeviceSynchronize());
} else {
// The error is checked by being printed on the line above. Discard
// the sticky copy so the checks below report their own failures.
(void)cudaGetLastError();
}
std::printf("\n");
// Both sums, on the same grid, against a reference in double.
const double want = sumCpu(h_in.data(), kElems);
float h_twoPass = 0.0f;
reduceBlocks<<<blocks, kThreadsPerBlock>>>(d_in, d_partials, kElems);
CUDA_CHECK(cudaGetLastError());
reduceBlocks<<<1, kThreadsPerBlock>>>(d_partials, d_out,
static_cast<size_t>(blocks));
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(&h_twoPass, d_out, sizeof(float), cudaMemcpyDeviceToHost));
float h_coop = 0.0f;
CUDA_CHECK(cudaLaunchKernelEx(&config, reduceCooperativeGrid, d_in,
d_partials, d_out, kElems));
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(
cudaMemcpy(&h_coop, d_out, sizeof(float), cudaMemcpyDeviceToHost));
std::printf("Sum of %zu floats\n", kElems);
std::printf(" CPU reference, double %.6f\n", want);
std::printf(" two launches %.6f\n",
static_cast<double>(h_twoPass));
std::printf(" one cooperative launch %.6f\n\n",
static_cast<double>(h_coop));
const double tol = kRelTolerance * std::fabs(want);
if (std::fabs(static_cast<double>(h_twoPass) - want) > tol) {
std::fprintf(stderr, "two-launch sum wrong: got %.6f, want %.6f\n",
static_cast<double>(h_twoPass), want);
++wrong;
}
if (std::fabs(static_cast<double>(h_coop) - want) > tol) {
std::fprintf(stderr, "cooperative sum wrong: got %.6f, want %.6f\n",
static_cast<double>(h_coop), want);
++wrong;
}
// What cg::coalesced_threads() returns under divergence.
std::vector<unsigned int> h_size(kReportThreads, 0u);
std::vector<unsigned int> h_rank(kReportThreads, 0u);
reportCoalesced<<<1, kReportThreads>>>(d_take, d_size, d_rank,
kReportThreads);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_size.data(), d_size,
kReportThreads * sizeof(unsigned int),
cudaMemcpyDeviceToHost));
CUDA_CHECK(cudaMemcpy(h_rank.data(), d_rank,
kReportThreads * sizeof(unsigned int),
cudaMemcpyDeviceToHost));
std::printf("What cg::coalesced_threads() returns\n");
std::printf(" warp lane group size rank in group\n");
for (int t = 0; t < kReportThreads; ++t) {
if (h_take[t] != 0) {
std::printf(" %4d %4d %10u %13u\n", t / kWarpSize,
t % kWarpSize, h_size[t], h_rank[t]);
}
}
std::printf("\n");
// A coalesced group never crosses a warp and never contains a thread
// that did not take the branch, so its size is at least one and at most
// the number of active lanes in that warp, and every member's rank is
// inside it. Which lanes end up together is not checked, because the
// documentation promises nothing about that and the run is here to
// report it.
for (int w = 0; w * kWarpSize < kReportThreads; ++w) {
int activeInWarp = 0;
for (int lane = 0; lane < kWarpSize; ++lane) {
activeInWarp += (h_take[w * kWarpSize + lane] != 0) ? 1 : 0;
}
for (int lane = 0; lane < kWarpSize; ++lane) {
const int t = w * kWarpSize + lane;
if (h_take[t] == 0) {
continue;
}
if (h_size[t] < 1u ||
h_size[t] > static_cast<unsigned int>(activeInWarp)) {
std::fprintf(stderr,
"warp %d lane %d: group size %u, but only %d "
"lanes of that warp took the branch\n",
w, lane, h_size[t], activeInWarp);
++wrong;
}
if (h_rank[t] >= h_size[t]) {
std::fprintf(stderr,
"warp %d lane %d: rank %u in a group of %u\n", w,
lane, h_rank[t], h_size[t]);
++wrong;
}
}
}
// Only launches sit inside the timed regions. Allocation and every copy
// are above them, so these numbers are kernel time and nothing else.
const float twoPassMs = timeKernel([&] {
reduceBlocks<<<blocks, kThreadsPerBlock>>>(d_in, d_partials, kElems);
reduceBlocks<<<1, kThreadsPerBlock>>>(d_partials, d_out,
static_cast<size_t>(blocks));
});
const float coopMs = timeKernel([&] {
CUDA_CHECK(cudaLaunchKernelEx(&config, reduceCooperativeGrid, d_in,
d_partials, d_out, kElems));
});
std::printf("Time, mean of %d runs after %d warm-ups, copies excluded\n",
kTimedRuns, kWarmupRuns);
std::printf(" two launches %.4f ms\n", twoPassMs);
std::printf(" one cooperative launch %.4f ms\n", coopMs);
// One cleanup block that every path falls through to, so a failure frees
// as much as a success does.
CUDA_CHECK(cudaFree(d_in));
CUDA_CHECK(cudaFree(d_partials));
CUDA_CHECK(cudaFree(d_out));
CUDA_CHECK(cudaFree(d_report));
CUDA_CHECK(cudaFree(d_take));
CUDA_CHECK(cudaFree(d_size));
CUDA_CHECK(cudaFree(d_rank));
if (wrong != 0) {
std::fprintf(stderr, "%d check(s) failed\n", wrong);
return EXIT_FAILURE;
}
std::printf("\nall checks passed\n");
return EXIT_SUCCESS;
}