code/day83-cudnn/conv_graph.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 83: a Conv2D forward built with the cuDNN graph API.
//
// The shape of every modern cuDNN program is here once, in order: three
// tensor descriptors carrying UIDs, a convolution descriptor, one operation
// that references them, an operation graph holding that operation, a
// heuristics query that returns engine configurations ranked best first, the
// first configuration that finalizes into an execution plan, the workspace
// that plan asks for, and a variant pack binding the UIDs to device pointers
// at execution time.
//
// This is the C backend API. NVIDIA's own guidance is the C++ frontend:
// "Use the cuDNN frontend to access the cuDNN Graph API unless you want to
// use legacy fixed-function routines or if you need a C-only graph API"
// (cuDNN Backend docs, Graph API page, checked 2026-09-01; the URL is in
// README.md). The frontend is a header-only wrapper over exactly the calls
// below, and it is not in the conda-forge cudnn package. Writing them out
// once is what stops the wrapper being magic.
//
// Nothing here is timed. Day 83's check is H, the harness, not M.
//
// The program grades itself against a float64 CPU reference with Kahan
// summation, then writes x, w and y to disk so check_against_pytorch.py can
// run the same convolution through torch and compare the two GPU answers.
//
// Build: nvcc -std=c++17 -O3 -arch=sm_75 -o conv_graph conv_graph.cu -lcudnn
// Run: ./conv_graph
//
// VERIFIED: Tesla T4, driver 580.173.02, CUDA 12.6, cuDNN 9.13.1,
// 2026-09-02. See evidence/run-2026-09-02.txt.
#include <algorithm>
#include <cmath>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <cuda_runtime.h>
#include <cudnn.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)
// cuDNN returns cudnnStatus_t, not cudaError_t, so CUDA_CHECK cannot wrap it,
// and CUDA-CODE-STYLE.md allows no second error macro. Day 44 met the same
// problem with cuBLAS and sent every status to a real branch; this file does
// the same through one function, so that a status is never discarded. Day 44
// shipped a bug by dropping a cuBLAS status inside a lambda, which is the
// whole reason this is a named function and not an expression.
// snippet: status-check
static bool cudnnOk(cudnnStatus_t status, const char* what) {
if (status == CUDNN_STATUS_SUCCESS) {
return true;
}
std::fprintf(stderr, "cuDNN error: %s: %s\n", what,
cudnnGetErrorString(status));
return false;
}
// end snippet
// One ResNet-shaped layer. 56x56 with 64 channels each side is the second
// stage of a ResNet-50 at batch 2, small enough that the float64 reference
// below finishes in seconds and large enough that cuDNN picks a real engine
// rather than a fallback.
constexpr int kBatch = 2;
constexpr int kInC = 64;
constexpr int kInH = 56;
constexpr int kInW = 56;
constexpr int kOutC = 64;
constexpr int kFiltR = 3;
constexpr int kFiltS = 3;
constexpr int kPad = 1;
constexpr int kStride = 1;
constexpr int kDilation = 1;
constexpr int kOutH =
(kInH + 2 * kPad - (kDilation * (kFiltR - 1) + 1)) / kStride + 1;
constexpr int kOutW =
(kInW + 2 * kPad - (kDilation * (kFiltS - 1) + 1)) / kStride + 1;
// Every output element is a sum of this many products. It sets the tolerance
// below and it is the reason the tolerance is not the table's flat 1e-5.
constexpr int kAccumTerms = kInC * kFiltR * kFiltS;
constexpr size_t kXElems = static_cast<size_t>(kBatch) * kInC * kInH * kInW;
constexpr size_t kWElems = static_cast<size_t>(kOutC) * kInC * kFiltR * kFiltS;
constexpr size_t kYElems = static_cast<size_t>(kBatch) * kOutC * kOutH * kOutW;
// UIDs. Any distinct int64_t values work; these are readable in a transcript.
// A UID is the only thing connecting a tensor descriptor to a device pointer,
// and the two are supplied in different calls minutes apart in wall time.
constexpr int64_t kUidX = 'x';
constexpr int64_t kUidW = 'w';
constexpr int64_t kUidY = 'y';
// The heuristics return a ranked list. Asking for eight and walking it is the
// documented workflow: "Look for the first engine config with functional
// support" (Graph API developer guide, checked 2026-09-01). Asking for one and
// finding it unsupported leaves you with nothing to fall back to.
constexpr int kMaxEngineConfigs = 8;
// Three tensors, one convolution, one operation, one graph, one heuristics
// query, eight engine configs, one engine, up to eight plans, one variant
// pack. 40 is comfortably above that and the bag refuses to overflow.
constexpr int kMaxDescriptors = 40;
// cudaMalloc returns memory aligned to at least 256 bytes, so 16 is a promise
// the allocator already keeps. cuDNN uses it to decide whether a vectorized
// engine is legal, so understating it silently costs performance.
constexpr int64_t kByteAlignment = 16;
// Every descriptor created goes in here and every one of them is destroyed on
// every exit path, success or failure. cuDNN descriptors are heap allocations
// behind an opaque pointer; leaking them leaks memory the sanitizer will not
// attribute to a line of this file.
struct DescriptorBag {
cudnnBackendDescriptor_t items[kMaxDescriptors];
int count;
};
static bool bagCreate(DescriptorBag& bag, cudnnBackendDescriptorType_t type,
cudnnBackendDescriptor_t* out) {
if (bag.count >= kMaxDescriptors) {
std::fprintf(stderr, "descriptor bag full at %d\n", kMaxDescriptors);
return false;
}
cudnnBackendDescriptor_t desc = nullptr;
if (!cudnnOk(cudnnBackendCreateDescriptor(type, &desc), "create")) {
return false;
}
bag.items[bag.count++] = desc;
*out = desc;
return true;
}
static void bagDestroy(DescriptorBag& bag) {
for (int i = bag.count - 1; i >= 0; --i) {
// Nothing to do about a failure here except say so: the process is on
// its way out and there is no second chance to free this.
(void)cudnnOk(cudnnBackendDestroyDescriptor(bag.items[i]), "destroy");
}
bag.count = 0;
}
// A tensor descriptor is dims, strides, dtype, alignment and a UID. There is
// no layout enum: NCHW versus NHWC is expressed entirely by the strides, which
// is why a transposed input is a stride change and not a different API.
// snippet: make-tensor
static bool makeTensor(DescriptorBag& bag, int64_t* dims, int64_t* strides,
int64_t uid, cudnnBackendDescriptor_t* out) {
cudnnDataType_t dtype = CUDNN_DATA_FLOAT;
int64_t alignment = kByteAlignment;
int64_t rank = 4;
if (!bagCreate(bag, CUDNN_BACKEND_TENSOR_DESCRIPTOR, out)) {
return false;
}
return cudnnOk(cudnnBackendSetAttribute(*out, CUDNN_ATTR_TENSOR_DATA_TYPE,
CUDNN_TYPE_DATA_TYPE, 1, &dtype),
"tensor data type") &&
cudnnOk(cudnnBackendSetAttribute(*out, CUDNN_ATTR_TENSOR_DIMENSIONS,
CUDNN_TYPE_INT64, rank, dims),
"tensor dimensions") &&
cudnnOk(cudnnBackendSetAttribute(*out, CUDNN_ATTR_TENSOR_STRIDES,
CUDNN_TYPE_INT64, rank, strides),
"tensor strides") &&
cudnnOk(cudnnBackendSetAttribute(*out, CUDNN_ATTR_TENSOR_UNIQUE_ID,
CUDNN_TYPE_INT64, 1, &uid),
"tensor unique id") &&
cudnnOk(
cudnnBackendSetAttribute(*out, CUDNN_ATTR_TENSOR_BYTE_ALIGNMENT,
CUDNN_TYPE_INT64, 1, &alignment),
"tensor byte alignment") &&
cudnnOk(cudnnBackendFinalize(*out), "finalize tensor");
}
// end snippet
// CUDNN_CROSS_CORRELATION, not CUDNN_CONVOLUTION. The two differ by a 180
// degree rotation of the filter, and every deep learning framework, PyTorch
// included, means cross-correlation when it says convolution. Picking the
// other one produces a plausible tensor full of wrong numbers.
static bool makeConvolution(DescriptorBag& bag, cudnnBackendDescriptor_t* out) {
cudnnDataType_t compType = CUDNN_DATA_FLOAT;
cudnnConvolutionMode_t mode = CUDNN_CROSS_CORRELATION;
int64_t spatialDims = 2;
int64_t pads[2] = {kPad, kPad};
int64_t filterStrides[2] = {kStride, kStride};
int64_t dilations[2] = {kDilation, kDilation};
if (!bagCreate(bag, CUDNN_BACKEND_CONVOLUTION_DESCRIPTOR, out)) {
return false;
}
return cudnnOk(
cudnnBackendSetAttribute(*out, CUDNN_ATTR_CONVOLUTION_COMP_TYPE,
CUDNN_TYPE_DATA_TYPE, 1, &compType),
"convolution compute type") &&
cudnnOk(
cudnnBackendSetAttribute(*out, CUDNN_ATTR_CONVOLUTION_CONV_MODE,
CUDNN_TYPE_CONVOLUTION_MODE, 1, &mode),
"convolution mode") &&
cudnnOk(cudnnBackendSetAttribute(*out,
CUDNN_ATTR_CONVOLUTION_SPATIAL_DIMS,
CUDNN_TYPE_INT64, 1, &spatialDims),
"convolution spatial dims") &&
cudnnOk(cudnnBackendSetAttribute(*out,
CUDNN_ATTR_CONVOLUTION_PRE_PADDINGS,
CUDNN_TYPE_INT64, 2, pads),
"convolution pre paddings") &&
cudnnOk(cudnnBackendSetAttribute(
*out, CUDNN_ATTR_CONVOLUTION_POST_PADDINGS,
CUDNN_TYPE_INT64, 2, pads),
"convolution post paddings") &&
cudnnOk(cudnnBackendSetAttribute(
*out, CUDNN_ATTR_CONVOLUTION_FILTER_STRIDES,
CUDNN_TYPE_INT64, 2, filterStrides),
"convolution filter strides") &&
cudnnOk(
cudnnBackendSetAttribute(*out, CUDNN_ATTR_CONVOLUTION_DILATIONS,
CUDNN_TYPE_INT64, 2, dilations),
"convolution dilations") &&
cudnnOk(cudnnBackendFinalize(*out), "finalize convolution");
}
// One operation, then a graph holding it. A graph with one node looks like
// pointless indirection until you want conv, bias and ReLU in one kernel,
// which is these same three calls with two more operations in the array.
// snippet: operation-graph
static bool makeOperationGraph(DescriptorBag& bag, cudnnHandle_t handle,
cudnnBackendDescriptor_t xDesc,
cudnnBackendDescriptor_t wDesc,
cudnnBackendDescriptor_t yDesc,
cudnnBackendDescriptor_t convDesc,
cudnnBackendDescriptor_t* out) {
// alpha and beta are the scaling factors in y = alpha * conv(x, w) +
// beta * y. beta = 0 means the output buffer is written, not accumulated
// into, so its previous contents never reach the answer.
float alpha = 1.0f;
float beta = 0.0f;
cudnnBackendDescriptor_t op = nullptr;
if (!bagCreate(bag, CUDNN_BACKEND_OPERATION_CONVOLUTION_FORWARD_DESCRIPTOR,
&op)) {
return false;
}
const bool built =
cudnnOk(cudnnBackendSetAttribute(
op, CUDNN_ATTR_OPERATION_CONVOLUTION_FORWARD_X,
CUDNN_TYPE_BACKEND_DESCRIPTOR, 1, &xDesc),
"operation x") &&
cudnnOk(cudnnBackendSetAttribute(
op, CUDNN_ATTR_OPERATION_CONVOLUTION_FORWARD_W,
CUDNN_TYPE_BACKEND_DESCRIPTOR, 1, &wDesc),
"operation w") &&
cudnnOk(cudnnBackendSetAttribute(
op, CUDNN_ATTR_OPERATION_CONVOLUTION_FORWARD_Y,
CUDNN_TYPE_BACKEND_DESCRIPTOR, 1, &yDesc),
"operation y") &&
cudnnOk(cudnnBackendSetAttribute(
op, CUDNN_ATTR_OPERATION_CONVOLUTION_FORWARD_CONV_DESC,
CUDNN_TYPE_BACKEND_DESCRIPTOR, 1, &convDesc),
"operation convolution descriptor") &&
cudnnOk(cudnnBackendSetAttribute(
op, CUDNN_ATTR_OPERATION_CONVOLUTION_FORWARD_ALPHA,
CUDNN_TYPE_FLOAT, 1, &alpha),
"operation alpha") &&
cudnnOk(cudnnBackendSetAttribute(
op, CUDNN_ATTR_OPERATION_CONVOLUTION_FORWARD_BETA,
CUDNN_TYPE_FLOAT, 1, &beta),
"operation beta") &&
cudnnOk(cudnnBackendFinalize(op), "finalize operation");
if (!built) {
return false;
}
if (!bagCreate(bag, CUDNN_BACKEND_OPERATIONGRAPH_DESCRIPTOR, out)) {
return false;
}
return cudnnOk(
cudnnBackendSetAttribute(*out, CUDNN_ATTR_OPERATIONGRAPH_OPS,
CUDNN_TYPE_BACKEND_DESCRIPTOR, 1, &op),
"operation graph ops") &&
cudnnOk(
cudnnBackendSetAttribute(*out, CUDNN_ATTR_OPERATIONGRAPH_HANDLE,
CUDNN_TYPE_HANDLE, 1, &handle),
"operation graph handle") &&
cudnnOk(cudnnBackendFinalize(*out), "finalize operation graph");
}
// end snippet
// The engine's global index is the number to quote when a cuDNN result is
// discussed, because it names the kernel family that produced it. It is two
// queries deep and the second one is not guaranteed on every build, so a
// failure here reports and continues rather than sinking the run.
static void printEngineIndex(DescriptorBag& bag,
cudnnBackendDescriptor_t engCfg) {
cudnnBackendDescriptor_t engine = nullptr;
if (!bagCreate(bag, CUDNN_BACKEND_ENGINE_DESCRIPTOR, &engine)) {
return;
}
int64_t count = 0;
if (cudnnBackendGetAttribute(engCfg, CUDNN_ATTR_ENGINECFG_ENGINE,
CUDNN_TYPE_BACKEND_DESCRIPTOR, 1, &count,
&engine) != CUDNN_STATUS_SUCCESS) {
std::printf("engine global index: not reported by this build\n");
return;
}
int64_t globalIndex = -1;
if (cudnnBackendGetAttribute(engine, CUDNN_ATTR_ENGINE_GLOBAL_INDEX,
CUDNN_TYPE_INT64, 1, &count,
&globalIndex) != CUDNN_STATUS_SUCCESS) {
std::printf("engine global index: not reported by this build\n");
return;
}
std::printf("engine global index: %lld\n",
static_cast<long long>(globalIndex));
}
// Heuristics, then the first engine configuration that finalizes into a plan.
// A configuration that fails to finalize is not an error: it is the engine
// saying it cannot run this graph, which is exactly what the walk is for.
// snippet: heuristics-and-plan
static bool makePlan(DescriptorBag& bag, cudnnHandle_t handle,
cudnnBackendDescriptor_t opGraph,
cudnnBackendDescriptor_t* out, int64_t* workspaceBytes) {
cudnnBackendHeurMode_t heurMode = CUDNN_HEUR_MODE_A;
cudnnBackendDescriptor_t heur = nullptr;
if (!bagCreate(bag, CUDNN_BACKEND_ENGINEHEUR_DESCRIPTOR, &heur)) {
return false;
}
if (!cudnnOk(cudnnBackendSetAttribute(
heur, CUDNN_ATTR_ENGINEHEUR_OPERATION_GRAPH,
CUDNN_TYPE_BACKEND_DESCRIPTOR, 1, &opGraph),
"heuristics operation graph") ||
!cudnnOk(cudnnBackendSetAttribute(heur, CUDNN_ATTR_ENGINEHEUR_MODE,
CUDNN_TYPE_HEUR_MODE, 1, &heurMode),
"heuristics mode") ||
!cudnnOk(cudnnBackendFinalize(heur), "finalize heuristics")) {
return false;
}
// GetAttribute fills descriptors the caller already created, so the
// configurations have to exist before the query that produces them.
cudnnBackendDescriptor_t engCfgs[kMaxEngineConfigs];
for (int i = 0; i < kMaxEngineConfigs; ++i) {
if (!bagCreate(bag, CUDNN_BACKEND_ENGINECFG_DESCRIPTOR, &engCfgs[i])) {
return false;
}
}
int64_t returned = 0;
if (!cudnnOk(
cudnnBackendGetAttribute(heur, CUDNN_ATTR_ENGINEHEUR_RESULTS,
CUDNN_TYPE_BACKEND_DESCRIPTOR,
kMaxEngineConfigs, &returned, engCfgs),
"heuristics results")) {
return false;
}
std::printf("heuristics mode A returned %lld engine configuration(s)\n",
static_cast<long long>(returned));
if (returned <= 0) {
std::fprintf(stderr, "no engine supports this graph\n");
return false;
}
for (int64_t i = 0; i < returned; ++i) {
cudnnBackendDescriptor_t plan = nullptr;
if (!bagCreate(bag, CUDNN_BACKEND_EXECUTION_PLAN_DESCRIPTOR, &plan)) {
return false;
}
if (!cudnnOk(
cudnnBackendSetAttribute(plan, CUDNN_ATTR_EXECUTION_PLAN_HANDLE,
CUDNN_TYPE_HANDLE, 1, &handle),
"plan handle") ||
!cudnnOk(cudnnBackendSetAttribute(
plan, CUDNN_ATTR_EXECUTION_PLAN_ENGINE_CONFIG,
CUDNN_TYPE_BACKEND_DESCRIPTOR, 1, &engCfgs[i]),
"plan engine config")) {
return false;
}
const cudnnStatus_t finalized = cudnnBackendFinalize(plan);
if (finalized != CUDNN_STATUS_SUCCESS) {
std::printf(" rank %lld cannot run this graph: %s\n",
static_cast<long long>(i),
cudnnGetErrorString(finalized));
continue;
}
int64_t count = 0;
if (!cudnnOk(cudnnBackendGetAttribute(
plan, CUDNN_ATTR_EXECUTION_PLAN_WORKSPACE_SIZE,
CUDNN_TYPE_INT64, 1, &count, workspaceBytes),
"plan workspace size")) {
return false;
}
std::printf("chose heuristic rank %lld, workspace %lld bytes\n",
static_cast<long long>(i),
static_cast<long long>(*workspaceBytes));
printEngineIndex(bag, engCfgs[i]);
*out = plan;
return true;
}
std::fprintf(stderr, "all %lld engine configurations refused this graph\n",
static_cast<long long>(returned));
return false;
}
// end snippet
// One output element, in double, with Kahan compensation because 576 terms is
// well past the 64 that EXERCISE-DESIGN.md sets as the threshold.
static double convOneOutput(const std::vector<float>& x,
const std::vector<float>& w, int n, int k, int p,
int q) {
double sum = 0.0;
double comp = 0.0;
for (int c = 0; c < kInC; ++c) {
for (int r = 0; r < kFiltR; ++r) {
const int h = p * kStride - kPad + r * kDilation;
if (h < 0 || h >= kInH) {
continue;
}
for (int s = 0; s < kFiltS; ++s) {
const int v = q * kStride - kPad + s * kDilation;
if (v < 0 || v >= kInW) {
continue;
}
const size_t xi =
((static_cast<size_t>(n) * kInC + c) * kInH + h) * kInW + v;
const size_t wi =
((static_cast<size_t>(k) * kInC + c) * kFiltR + r) *
kFiltS +
s;
const double term =
static_cast<double>(x[xi]) * static_cast<double>(w[wi]);
const double y = term - comp;
const double t = sum + y;
comp = (t - sum) - y;
sum = t;
}
}
}
return sum;
}
static void conv2dCpu(const std::vector<float>& x, const std::vector<float>& w,
std::vector<double>& y) {
for (int n = 0; n < kBatch; ++n) {
for (int k = 0; k < kOutC; ++k) {
for (int p = 0; p < kOutH; ++p) {
for (int q = 0; q < kOutW; ++q) {
const size_t yi =
((static_cast<size_t>(n) * kOutC + k) * kOutH + p) *
kOutW +
q;
y[yi] = convOneOutput(x, w, n, k, p, q);
}
}
}
}
}
// A fixed LCG, so the run is the same on every machine and a mismatch is
// never the input. Values land in [-1, 1).
static void fillDeterministic(std::vector<float>& v, uint32_t seed) {
uint32_t state = seed;
for (size_t i = 0; i < v.size(); ++i) {
state = state * 1664525u + 1013904223u;
v[i] = static_cast<float>(state >> 8) / 8388608.0f - 1.0f;
}
}
static bool writeRaw(const char* path, const void* data, size_t bytes) {
std::FILE* f = std::fopen(path, "wb");
if (f == nullptr) {
std::fprintf(stderr, "cannot open %s for writing\n", path);
return false;
}
const size_t written = std::fwrite(data, 1, bytes, f);
const int closed = std::fclose(f);
if (written != bytes || closed != 0) {
std::fprintf(stderr, "short write to %s: %zu of %zu bytes\n", path,
written, bytes);
return false;
}
return true;
}
// The Python half reads the shapes from here rather than repeating them, so
// the two programs cannot disagree about what they convolved.
static bool writeShapes(const char* path) {
std::FILE* f = std::fopen(path, "w");
if (f == nullptr) {
std::fprintf(stderr, "cannot open %s for writing\n", path);
return false;
}
std::fprintf(f, "n %d\nc %d\nh %d\nw %d\nk %d\nr %d\ns %d\n", kBatch, kInC,
kInH, kInW, kOutC, kFiltR, kFiltS);
std::fprintf(f, "pad %d\nstride %d\ndilation %d\np %d\nq %d\n", kPad,
kStride, kDilation, kOutH, kOutW);
return std::fclose(f) == 0;
}
int main() {
int device = 0;
CUDA_CHECK(cudaGetDevice(&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);
// Two versions, not one. CUDNN_VERSION is what the header said at compile
// time; cudnnGetVersion() is what the shared object says right now. A
// mismatch means the headers and the library came from different installs,
// which is the single most reported cuDNN problem there is.
std::printf("cuDNN header %d, runtime %zu\n", CUDNN_VERSION,
cudnnGetVersion());
std::printf("conv2d: x[%d,%d,%d,%d] * w[%d,%d,%d,%d] -> y[%d,%d,%d,%d]\n",
kBatch, kInC, kInH, kInW, kOutC, kInC, kFiltR, kFiltS, kBatch,
kOutC, kOutH, kOutW);
std::printf("pad %d, stride %d, dilation %d, cross-correlation, ", kPad,
kStride, kDilation);
std::printf("fp32 in, fp32 compute, fp32 out\n\n");
std::vector<float> h_x(kXElems);
std::vector<float> h_w(kWElems);
std::vector<float> h_y(kYElems);
std::vector<double> h_want(kYElems);
fillDeterministic(h_x, 20260901u);
fillDeterministic(h_w, 83u);
float* d_x = nullptr;
float* d_w = nullptr;
float* d_y = nullptr;
void* d_workspace = nullptr;
CUDA_CHECK(cudaMalloc(&d_x, kXElems * sizeof(float)));
CUDA_CHECK(cudaMalloc(&d_w, kWElems * sizeof(float)));
CUDA_CHECK(cudaMalloc(&d_y, kYElems * sizeof(float)));
CUDA_CHECK(cudaMemcpy(d_x, h_x.data(), kXElems * sizeof(float),
cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_w, h_w.data(), kWElems * sizeof(float),
cudaMemcpyHostToDevice));
int status = EXIT_SUCCESS;
DescriptorBag bag{};
cudnnHandle_t handle = nullptr;
if (!cudnnOk(cudnnCreate(&handle), "cudnnCreate")) {
status = EXIT_FAILURE;
}
cudnnBackendDescriptor_t plan = nullptr;
int64_t workspaceBytes = 0;
if (status == EXIT_SUCCESS) {
int64_t xDims[4] = {kBatch, kInC, kInH, kInW};
int64_t xStrides[4] = {static_cast<int64_t>(kInC) * kInH * kInW,
static_cast<int64_t>(kInH) * kInW, kInW, 1};
int64_t wDims[4] = {kOutC, kInC, kFiltR, kFiltS};
int64_t wStrides[4] = {static_cast<int64_t>(kInC) * kFiltR * kFiltS,
static_cast<int64_t>(kFiltR) * kFiltS, kFiltS,
1};
int64_t yDims[4] = {kBatch, kOutC, kOutH, kOutW};
int64_t yStrides[4] = {static_cast<int64_t>(kOutC) * kOutH * kOutW,
static_cast<int64_t>(kOutH) * kOutW, kOutW, 1};
cudnnBackendDescriptor_t xDesc = nullptr;
cudnnBackendDescriptor_t wDesc = nullptr;
cudnnBackendDescriptor_t yDesc = nullptr;
cudnnBackendDescriptor_t convDesc = nullptr;
cudnnBackendDescriptor_t opGraph = nullptr;
if (!makeTensor(bag, xDims, xStrides, kUidX, &xDesc) ||
!makeTensor(bag, wDims, wStrides, kUidW, &wDesc) ||
!makeTensor(bag, yDims, yStrides, kUidY, &yDesc) ||
!makeConvolution(bag, &convDesc) ||
!makeOperationGraph(bag, handle, xDesc, wDesc, yDesc, convDesc,
&opGraph) ||
!makePlan(bag, handle, opGraph, &plan, &workspaceBytes)) {
status = EXIT_FAILURE;
}
}
if (status == EXIT_SUCCESS && workspaceBytes > 0) {
CUDA_CHECK(
cudaMalloc(&d_workspace, static_cast<size_t>(workspaceBytes)));
}
if (status == EXIT_SUCCESS) {
// The variant pack is the late binding. Nothing before this point knew
// an address; the UIDs and the pointers are two arrays in the same
// order, and swapping two entries is a silent wrong answer.
// snippet: variant-pack
void* devPtrs[3] = {d_x, d_w, d_y};
int64_t uids[3] = {kUidX, kUidW, kUidY};
cudnnBackendDescriptor_t varPack = nullptr;
if (!bagCreate(bag, CUDNN_BACKEND_VARIANT_PACK_DESCRIPTOR, &varPack) ||
!cudnnOk(cudnnBackendSetAttribute(
varPack, CUDNN_ATTR_VARIANT_PACK_DATA_POINTERS,
CUDNN_TYPE_VOID_PTR, 3, devPtrs),
"variant pack data pointers") ||
!cudnnOk(cudnnBackendSetAttribute(
varPack, CUDNN_ATTR_VARIANT_PACK_UNIQUE_IDS,
CUDNN_TYPE_INT64, 3, uids),
"variant pack unique ids") ||
!cudnnOk(cudnnBackendSetAttribute(
varPack, CUDNN_ATTR_VARIANT_PACK_WORKSPACE,
CUDNN_TYPE_VOID_PTR, 1, &d_workspace),
"variant pack workspace") ||
!cudnnOk(cudnnBackendFinalize(varPack), "finalize variant pack") ||
!cudnnOk(cudnnBackendExecute(handle, plan, varPack), "execute")) {
status = EXIT_FAILURE;
}
// end snippet
}
if (status == EXIT_SUCCESS) {
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(h_y.data(), d_y, kYElems * sizeof(float),
cudaMemcpyDeviceToHost));
std::printf("\nreference: float64, Kahan, %d terms per output\n",
kAccumTerms);
conv2dCpu(h_x, h_w, h_want);
double maxAbsRef = 0.0;
for (size_t i = 0; i < kYElems; ++i) {
maxAbsRef = std::max(maxAbsRef, std::fabs(h_want[i]));
}
// EXERCISE-DESIGN.md: f32 rtol is 1e-5, and any output built from K
// accumulated terms gets max(rtol, 4 * eps * sqrt(K)) instead. cuDNN
// and this reference do not add the 576 products in the same order,
// and neither does PyTorch, so the bound has to cover a reordering
// rather than a bug. Day 68 is the lesson on why.
const double eps = 1.1920928955078125e-7; // 2^-23
const double rtol =
std::max(1e-5, 4.0 * eps * std::sqrt(1.0 * kAccumTerms));
const double atol = rtol * maxAbsRef;
std::printf("tolerance |got-exp| <= %.3e + %.3e*|exp|\n", atol, rtol);
std::printf(" rtol = max(1e-5, 4*2^-23*sqrt(%d))\n",
kAccumTerms);
size_t bad = 0;
double worst = 0.0;
size_t worstIndex = 0;
for (size_t i = 0; i < kYElems; ++i) {
const double diff = std::fabs(h_y[i] - h_want[i]);
if (diff > worst) {
worst = diff;
worstIndex = i;
}
if (diff > atol + rtol * std::fabs(h_want[i])) {
if (bad == 0) {
std::fprintf(stderr,
"first mismatch at %zu: cuDNN %.9g, "
"reference %.9g, diff %.3e\n",
i, static_cast<double>(h_y[i]), h_want[i],
diff);
}
++bad;
}
}
std::printf("largest |cuDNN - reference| = %.3e at index %zu\n", worst,
worstIndex);
if (bad > 0) {
std::fprintf(stderr, "%zu of %zu outputs outside tolerance\n", bad,
kYElems);
status = EXIT_FAILURE;
} else {
std::printf("all %zu outputs inside tolerance\n", kYElems);
}
}
if (status == EXIT_SUCCESS) {
if (!writeShapes("conv_shapes.txt") ||
!writeRaw("conv_x.f32", h_x.data(), kXElems * sizeof(float)) ||
!writeRaw("conv_w.f32", h_w.data(), kWElems * sizeof(float)) ||
!writeRaw("conv_y.f32", h_y.data(), kYElems * sizeof(float))) {
status = EXIT_FAILURE;
} else {
std::printf(
"wrote conv_shapes.txt, conv_x.f32, conv_w.f32, conv_y.f32\n"
"next: python3 check_against_pytorch.py\n");
}
}
bagDestroy(bag);
if (handle != nullptr && !cudnnOk(cudnnDestroy(handle), "cudnnDestroy")) {
status = EXIT_FAILURE;
}
if (d_workspace != nullptr) {
CUDA_CHECK(cudaFree(d_workspace));
}
CUDA_CHECK(cudaFree(d_x));
CUDA_CHECK(cudaFree(d_w));
CUDA_CHECK(cudaFree(d_y));
return status;
}