code/day69-portability/portability.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 69: portability across GPU generations. One binary that runs on a
// Tesla T4 (sm_75) and an H100 (sm_90): sm_75 SASS for an exact match,
// compute_90 PTX for the driver to JIT on Hopper. The kernel is the same
// shuffle reduction everywhere; what changes per architecture is the
// launch bounds baked in at compile time and the grid the host sizes at
// run time, from the runtime's own occupancy probe rather than from a
// version table.
//
// The file keeps one deliberate bug, archSeenByHostBuggy(), so the page
// can put the broken host-side __CUDA_ARCH__ check next to the fixed
// runtime query and the harness can diff their behavior.
//
// Build (fat binary):
// nvcc -std=c++17 -O3 -gencode arch=compute_75,code=sm_75 \
// -gencode arch=compute_90,code=compute_90 -o portability \
// portability.cu
// Run:
// ./portability dispatch from the device's real CC
// ./portability --force-cc 90 dispatch decision forced to CC 9.0, same
// kernel, so the Hopper host path runs on
// the T4 too
//
// The README builds two more binaries from this same file (sm_90 SASS
// only, compute_75 PTX only) to capture the failure string and the JIT
// path on the T4.
//
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <vector>
#include <cuda_runtime.h>
// The one error macro. This file is standalone, so it carries its own
// verbatim copy. `err_` has 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)
// 4 Mi elements, every one 1.0f. Every partial sum is a whole number and
// the total, 4194304, sits far below 2^24, so float addition is exact in
// any order and the gate can demand equality rather than tolerance.
constexpr size_t kElems = 4 * 1024 * 1024;
constexpr int kThreadsPerBlock = 256; // 8 warps; the exercise varies this
static_assert(kThreadsPerBlock % 32 == 0,
"block size must be a whole number of warps");
static_assert(kThreadsPerBlock >= 32 && kThreadsPerBlock <= 1024,
"one warp minimum, hardware block ceiling maximum");
// __CUDA_ARCH__ exists only while nvcc compiles device code, once per
// -gencode target, so the block below is evaluated twice for the fat
// binary: once as 750, once as 900. CC 7.5 keeps 1024 resident threads
// and 16 resident blocks per SM; CC 9.0 keeps 2048 and 32. The second
// argument of __launch_bounds__ promises ptxas this many resident blocks
// per SM, so each architecture's SASS is register-budgeted for its own
// ceiling. Host code also compiles this block, sees no __CUDA_ARCH__, and
// lands in the #else branch. Nothing in host code is allowed to read
// these constants; the host asks the runtime instead (see main).
// snippet: launch-bounds
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 900
constexpr int kResidentThreadsPerSm = 2048;
constexpr int kResidentBlocksPerSm = 32;
#else
constexpr int kResidentThreadsPerSm = 1024;
constexpr int kResidentBlocksPerSm = 16;
#endif
constexpr int kWantedBlocksPerSm =
(kResidentThreadsPerSm / kThreadsPerBlock < kResidentBlocksPerSm)
? kResidentThreadsPerSm / kThreadsPerBlock
: kResidentBlocksPerSm;
// end snippet: launch-bounds
// Grid-stride shuffle reduction: butterfly-free __shfl_down_sync inside
// each warp, one shared slot per warp, warp 0 folds the slots, lane 0 adds
// the block's sum to *out. Day 23 built the shuffle, day 24 the shape.
// snippet: kernel
__global__ void __launch_bounds__(kThreadsPerBlock, kWantedBlocksPerSm)
reduceShuffle(const float* __restrict__ in, float* __restrict__ out, size_t n) {
__shared__ float warpSums[kThreadsPerBlock / 32];
const int lane = threadIdx.x & 31;
const int warp = threadIdx.x >> 5;
float sum = 0.0f;
const size_t stride = static_cast<size_t>(gridDim.x) * blockDim.x;
for (size_t elem =
blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
elem < n; elem += stride) {
sum += in[elem];
}
for (int offset = 16; offset > 0; offset >>= 1) {
sum += __shfl_down_sync(0xffffffffu, sum, offset);
}
if (lane == 0) {
warpSums[warp] = sum;
}
__syncthreads();
if (warp == 0) {
const int numWarps = static_cast<int>(blockDim.x) >> 5;
sum = (lane < numWarps) ? warpSums[lane] : 0.0f;
for (int offset = 16; offset > 0; offset >>= 1) {
sum += __shfl_down_sync(0xffffffffu, sum, offset);
}
if (lane == 0) {
atomicAdd(out, sum);
}
}
}
// end snippet
// snippet: host-arch-trap
// DELIBERATE BUG (day 69), kept so the harness can diff it against the
// fixed variant. __CUDA_ARCH__ is defined only while nvcc compiles
// device code. Here the preprocessor sees an undefined identifier,
// which #if evaluates as 0, so the first branch is dead on every
// machine that will ever run this program. No warning, no error.
static int archSeenByHostBuggy() {
#if __CUDA_ARCH__ >= 750
return __CUDA_ARCH__;
#else
return 0;
#endif
}
// Fixed: the host cannot know the device architecture at compile time,
// because the same host binary runs unchanged next to any card. It asks
// the runtime, in __CUDA_ARCH__'s units.
static int archSeenByHostFixed(const cudaDeviceProp& prop) {
return prop.major * 100 + prop.minor * 10;
}
// end snippet: host-arch-trap
int main(int argc, char** argv) {
// --force-cc 90 overrides the dispatch decision only. The kernel, the
// probes and the gates are identical, which is what makes the Hopper
// host path testable on a T4.
int forcedCc = 0;
if (argc == 3 && std::strcmp(argv[1], "--force-cc") == 0) {
forcedCc = std::atoi(argv[2]);
}
if ((argc != 1 && forcedCc == 0) || (forcedCc != 0 && forcedCc < 75)) {
std::fprintf(stderr, "usage: %s [--force-cc 75|90]\n", argv[0]);
return EXIT_FAILURE;
}
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);
// snippet: probes
// Feature probes, not version checks: the attribute answers "can
// this device do it", which survives new architectures the way a
// `major >= 9` test does not.
int coop = 0;
int cluster = 0;
CUDA_CHECK(
cudaDeviceGetAttribute(&coop, cudaDevAttrCooperativeLaunch, device));
CUDA_CHECK(
cudaDeviceGetAttribute(&cluster, cudaDevAttrClusterLaunch, device));
// The occupancy probe folds every per-architecture ceiling (threads,
// blocks, registers, shared memory) into one number for the code the
// driver actually loaded, so the host never carries the table that
// kWantedBlocksPerSm needed a #if for.
int blocksPerSm = 0;
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&blocksPerSm, reduceShuffle, kThreadsPerBlock, 0));
const int blocks = blocksPerSm * prop.multiProcessorCount;
// end snippet: probes
const int dispatchCc =
(forcedCc != 0) ? forcedCc : prop.major * 10 + prop.minor;
const char* path = (dispatchCc >= 90)
? "Hopper: compute_90 PTX, JIT at first launch"
: "Turing: sm_75 SASS, exact match, no JIT";
const int tableBlocksPerSm =
(dispatchCc >= 90)
? (2048 / kThreadsPerBlock < 32 ? 2048 / kThreadsPerBlock : 32)
: (1024 / kThreadsPerBlock < 16 ? 1024 / kThreadsPerBlock : 16);
std::printf("dispatch CC: %d.%d (%s)\n", dispatchCc / 10, dispatchCc % 10,
(forcedCc != 0) ? "forced" : "from device");
std::printf("path: %s\n", path);
std::printf("probe: cooperative launch %d, cluster launch %d\n", coop,
cluster);
std::printf("probe: %d blocks/SM x %d SMs -> grid %d\n", blocksPerSm,
prop.multiProcessorCount, blocks);
std::printf("version table said %d blocks/SM: %s\n", tableBlocksPerSm,
(tableBlocksPerSm == blocksPerSm)
? "probe agrees"
: "probe disagrees, probe wins (reported, not gated)");
const size_t bytes = kElems * sizeof(float);
std::vector<float> h_in(kElems, 1.0f);
const float want = static_cast<float>(kElems);
float* d_in = nullptr;
float* d_out = nullptr;
CUDA_CHECK(cudaMalloc(&d_in, bytes));
CUDA_CHECK(cudaMalloc(&d_out, sizeof(float)));
CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemset(d_out, 0, sizeof(float)));
reduceShuffle<<<blocks, kThreadsPerBlock>>>(d_in, d_out, kElems);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
float got = 0.0f;
CUDA_CHECK(cudaMemcpy(&got, d_out, sizeof(float), cudaMemcpyDeviceToHost));
// Free before gating, so the failure paths free too. Status
// fall-through, one return at the end, no early exits past this line.
CUDA_CHECK(cudaFree(d_in));
CUDA_CHECK(cudaFree(d_out));
// Three gates, all real branches. Gate 1 pins the deliberate bug to
// its known wrong answer: if it ever stops returning 0, the file was
// edited and the page's diff no longer matches the code.
int failures = 0;
const int buggy = archSeenByHostBuggy();
const int fixed = archSeenByHostFixed(prop);
std::printf("host arch, buggy #if __CUDA_ARCH__: %d\n", buggy);
std::printf("host arch, runtime query: %d\n", fixed);
if (buggy != 0) {
std::fprintf(stderr,
"gate 1 FAIL: buggy probe returned %d, expected the "
"dead-branch 0\n",
buggy);
failures += 1;
}
if (fixed < 750) {
std::fprintf(stderr,
"gate 2 FAIL: runtime probe returned %d, below the "
"sm_75 floor this binary ships\n",
fixed);
failures += 1;
}
if (got != want) {
std::fprintf(stderr, "gate 3 FAIL: gpu sum %.1f, reference %.1f\n",
static_cast<double>(got), static_cast<double>(want));
failures += 1;
}
if (failures != 0) {
return EXIT_FAILURE;
}
std::printf(
"all gates pass: sum %.0f == %.0f, buggy probe 0, runtime "
"probe %d\n",
static_cast<double>(got), static_cast<double>(want), fixed);
return EXIT_SUCCESS;
}