code/day58-pdl/pdl_pipeline.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 58: programmatic dependent launch (PDL).
//
// One two-kernel pipeline, one stream, launched two ways:
//
// baseline produce then combine, serialized by stream order. combine's
// first instruction waits for produce's last block to retire.
// pdl combine is launched with programmaticStreamSerializationAllowed
// set, so its blocks may start their prologue while produce is
// still draining. combine calls cudaGridDependencySynchronize()
// before it reads anything produce wrote.
//
// The two modes launch the same two kernels from the same launch helper, and
// the only difference is one attribute value, so nothing else can explain a
// gap between the two times. The program requires the two output vectors to
// be bit-identical before it times anything.
//
// This is the one program in the course that does not run on the Tesla T4
// verification node. PDL needs compute capability 9.0 (Hopper, or Blackwell
// including RTX 50 cards); the program checks and says so rather than dying
// inside the driver.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_90 -lineinfo -o pdl_pipeline \
// pdl_pipeline.cu
// Run: ./pdl_pipeline
//
// -arch=sm_90, not sm_90a: PDL needs nothing architecture-specific, and
// sm_90 embeds compute_90 PTX that JIT-compiles forward onto CC 10.x and
// 12.x cards. An sm_90a binary loads on nothing but Hopper.
//
// UNVERIFIED: not yet compiled or run on real hardware. Do not publish any
// output as this program's output until it has run. See
// research/REVIEW-PROCESS.md.
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <vector>
#include <cuda_runtime.h>
#include <nvtx3/nvToolsExt.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)
// kElems is deliberately not a multiple of the block size, so every kernel's
// bounds check runs on every launch. Five float buffers at this size is
// about 20 MiB, small enough for any CC 9.0 card and large enough that a
// launch is not the whole kernel.
constexpr size_t kElems = 1024ull * 1024ull + 611ull;
constexpr int kThreadsPerBlock = 256; // 8 warps
constexpr int kWarmupRuns = 3;
constexpr int kTimedRuns = 10;
constexpr int kPairs = 64; // produce+combine pairs per timed run
constexpr int kInnerIters = 256; // FMA chain length per phase
constexpr float kRelTolerance = 1e-5f;
static_assert(kThreadsPerBlock % 32 == 0,
"block size must be a whole number of warps");
// Each phase iterates v = fma(v, mul, add) so the kernels have a real body
// whose cost is knowable: kInnerIters dependent FMAs per thread per phase.
// The multipliers sit just under 1 so the values stay bounded.
constexpr float kMulProduce = 0.999f;
constexpr float kAddProduce = 0.25f;
constexpr float kMulEpilogue = 0.997f;
constexpr float kAddEpilogue = 0.125f;
constexpr float kMulPrologue = 0.998f;
constexpr float kAddPrologue = 0.5f;
// A chain of dependent FMAs, computed identically on host (std::fmaf) and
// device (fmaf). Both are correctly rounded, so the reference match is
// exact and the tolerance below only covers habits, not need.
__host__ __device__ inline float fmaChain(float v, float mul, float add) {
for (int k = 0; k < kInnerIters; ++k) {
v = fmaf(v, mul, add);
}
return v;
}
// snippet: primary
// The primary kernel. One thread owns one element. Consecutive threads read
// and write consecutive elements, so every access coalesces; nothing here
// is memory-clever on purpose.
//
// Phase 1 computes mid[i], which the secondary kernel will read. The
// trigger then tells the driver the dependent launch may begin. Phase 2
// writes aux[i], work the secondary never reads, so it can overlap the
// secondary's prologue. The trigger promises nothing about memory: the
// secondary's cudaGridDependencySynchronize() is what makes mid visible.
__global__ void produceStage(const float* __restrict__ x,
float* __restrict__ mid, float* __restrict__ aux,
size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
mid[i] = fmaChain(x[i], kMulProduce, kAddProduce);
}
cudaTriggerProgrammaticLaunchCompletion();
if (i < n) {
aux[i] = fmaChain(x[i], kMulEpilogue, kAddEpilogue);
}
}
// end snippet
// snippet: secondary
// The secondary kernel. Same launch shape and access pattern as the
// primary. The prologue depends only on y, so it is legal before the
// synchronize; the read of mid[i] is not, and sits after it. Launched
// without the PDL attribute this kernel is still correct: the wait finds
// the primary already finished and costs nothing.
__global__ void combineStage(const float* __restrict__ mid,
const float* __restrict__ y,
float* __restrict__ out, size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
float p = 0.0f;
if (i < n) {
p = fmaChain(y[i], kMulPrologue, kAddPrologue);
}
cudaGridDependencySynchronize();
if (i < n) {
out[i] = mid[i] + p;
}
}
// end snippet
// CPU reference. Written for obvious correctness: the same three FMA chains
// in the same order, one loop, no tricks. It never allocates; the caller
// owns every buffer.
static void pipelineCpu(const float* x, const float* y, float* mid, float* aux,
float* out, size_t n) {
for (size_t i = 0; i < n; ++i) {
mid[i] = fmaChain(x[i], kMulProduce, kAddProduce);
aux[i] = fmaChain(x[i], kMulEpilogue, kAddEpilogue);
out[i] = mid[i] + fmaChain(y[i], kMulPrologue, kAddPrologue);
}
}
// Returns the first index where got and want differ by more than the
// relative tolerance, or n if they agree everywhere.
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;
}
// Launches kPairs produce+combine pairs into one stream. `allowPdl` is the
// single bit the two timed modes differ by. The secondary always goes
// through cudaLaunchKernelEx so the launch path is identical too; with the
// attribute value 0 the launch is an ordinary stream-ordered launch.
static void launchChain(cudaStream_t stream, const float* d_x, float* d_mid,
float* d_aux, const float* d_y, float* d_out,
int blocks, int allowPdl) {
// snippet: attribute
cudaLaunchAttribute attrs[1];
attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
attrs[0].val.programmaticStreamSerializationAllowed = allowPdl;
cudaLaunchConfig_t cfg = {};
cfg.gridDim = dim3(static_cast<unsigned int>(blocks));
cfg.blockDim = dim3(kThreadsPerBlock);
cfg.dynamicSmemBytes = 0;
cfg.stream = stream;
cfg.attrs = attrs;
cfg.numAttrs = 1;
// end snippet
for (int pair = 0; pair < kPairs; ++pair) {
produceStage<<<blocks, kThreadsPerBlock, 0, stream>>>(d_x, d_mid, d_aux,
kElems);
CUDA_CHECK(
cudaLaunchKernelEx(&cfg, combineStage, d_mid, d_y, d_out, kElems));
}
}
// Times a launch with CUDA events and returns the mean milliseconds per
// run. Copied verbatim from the course timing helper: warm-up inside, so
// every kernel it times pays its lazy-loading cost before the clock starts.
template <typename LaunchFn>
static float timeKernel(LaunchFn launch) {
cudaEvent_t start, stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));
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)\n", prop.name, prop.major,
prop.minor);
// PDL exists from CC 9.0. Failing here, before any allocation, beats a
// no-kernel-image failure ten lines further down on the course's own T4.
if (prop.major < 9) {
std::fprintf(stderr,
"programmatic dependent launch needs compute capability "
"9.0 or newer; %s is %d.%d\n",
prop.name, prop.major, prop.minor);
return EXIT_FAILURE;
}
const size_t bytes = kElems * sizeof(float);
const int blocks =
static_cast<int>((kElems + kThreadsPerBlock - 1) / kThreadsPerBlock);
std::printf("n = %zu, %d blocks of %d threads, %d pairs per run\n", kElems,
blocks, kThreadsPerBlock, kPairs);
std::vector<float> h_x(kElems);
std::vector<float> h_y(kElems);
std::vector<float> h_mid(kElems);
std::vector<float> h_aux(kElems);
std::vector<float> h_want(kElems);
std::vector<float> h_out(kElems);
std::vector<float> h_outPdl(kElems);
for (size_t i = 0; i < kElems; ++i) {
h_x[i] = static_cast<float>(i % 97) * 0.01f;
h_y[i] = static_cast<float>(i % 13) * 0.05f;
}
float* d_x = nullptr;
float* d_y = nullptr;
float* d_mid = nullptr;
float* d_aux = nullptr;
float* d_out = nullptr;
CUDA_CHECK(cudaMalloc(&d_x, bytes));
CUDA_CHECK(cudaMalloc(&d_y, bytes));
CUDA_CHECK(cudaMalloc(&d_mid, bytes));
CUDA_CHECK(cudaMalloc(&d_aux, 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));
cudaStream_t stream; // one stream; PDL is about order within it
CUDA_CHECK(cudaStreamCreate(&stream));
bool ok = true;
// Correctness first, both modes, before anything is timed. The baseline
// run checks against the CPU reference; the PDL run must then match the
// baseline bit for bit, because the attribute changes scheduling and
// must change nothing else.
nvtxRangePushA("verify");
launchChain(stream, d_x, d_mid, d_aux, d_y, d_out, blocks, 0);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaStreamSynchronize(stream));
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));
pipelineCpu(h_x.data(), h_y.data(), h_mid.data(), h_aux.data(),
h_want.data(), kElems);
const size_t bad =
firstMismatch(h_out.data(), h_want.data(), kElems, kRelTolerance);
if (bad != kElems) {
std::fprintf(stderr, "baseline wrong at %zu: got %.9g, want %.9g\n",
bad, h_out[bad], h_want[bad]);
ok = false;
}
launchChain(stream, d_x, d_mid, d_aux, d_y, d_out, blocks, 1);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaStreamSynchronize(stream));
CUDA_CHECK(
cudaMemcpy(h_outPdl.data(), d_out, bytes, cudaMemcpyDeviceToHost));
if (ok && std::memcmp(h_out.data(), h_outPdl.data(), bytes) != 0) {
std::fprintf(stderr,
"pdl output differs from baseline; the attribute must "
"not change results\n");
ok = false;
}
nvtxRangePop();
if (ok) {
// Time both modes. Same helper, same stream, same kernels; the one
// bit differs. The NVTX ranges are what day 41's nsys workflow reads.
nvtxRangePushA("baseline");
const float msBase = timeKernel([&] {
launchChain(stream, d_x, d_mid, d_aux, d_y, d_out, blocks, 0);
});
nvtxRangePop();
nvtxRangePushA("pdl");
const float msPdl = timeKernel([&] {
launchChain(stream, d_x, d_mid, d_aux, d_y, d_out, blocks, 1);
});
nvtxRangePop();
std::printf("%d pairs, mean of %d runs, copies not included\n", kPairs,
kTimedRuns);
std::printf("baseline: %.3f ms (%.2f us per pair)\n", msBase,
msBase * 1000.0f / kPairs);
std::printf("pdl: %.3f ms (%.2f us per pair)\n", msPdl,
msPdl * 1000.0f / kPairs);
std::printf("pdl/baseline: %.3f\n", msPdl / msBase);
}
CUDA_CHECK(cudaStreamDestroy(stream));
CUDA_CHECK(cudaFree(d_x));
CUDA_CHECK(cudaFree(d_y));
CUDA_CHECK(cudaFree(d_mid));
CUDA_CHECK(cudaFree(d_aux));
CUDA_CHECK(cudaFree(d_out));
return ok ? EXIT_SUCCESS : EXIT_FAILURE;
}