100 DAYS OF GPU Course preview
Day 83Module 9
in-technical-review

cuDNN with the graph API

Many cuDNN convolution examples still use this API:

cudnnFindConvolutionForwardAlgorithm(handle, xDesc, wDesc, convDesc,
                                     yDesc, 1, &count, &perf);
cudnnGetConvolutionForwardWorkspaceSize(handle, xDesc, wDesc, convDesc,
                                        yDesc, perf.algo, &bytes);
cudnnConvolutionForward(handle, &alpha, xDesc, x, wDesc, w, convDesc,
                        perf.algo, workspace, bytes, &beta, yDesc, y);

All three calls are on cuDNN 9's deprecated list. NVIDIA's API overview says "We recommend users to make API calls through the graph API" (cuDNN Backend API overview, https://docs.nvidia.com/deeplearning/cudnn/backend/v9.15.1/api/overview.html , checked 2026-09-01). The graph API uses a different model, not renamed functions.

This lesson builds a Conv2D forward graph and checks its output against PyTorch with a derived tolerance.

Six descriptors, then one call

cuDNN 9's general entry point is a graph builder. Everything is a cudnnBackendDescriptor_t, an opaque pointer whose type you name at creation, whose attributes you set one at a time, and which you then finalize. Finalize is where validation happens.

A tensor descriptor is dimensions, strides, a data type, a byte alignment and a unique id. There is no layout enum, so NCHW and NHWC are two stride arrays and nothing else. The unique id is an int64_t that connects the descriptor to a device address at execution.

The convolution descriptor carries padding, stride, dilation, the spatial dimension count and the compute type. An operation references three tensors, one convolution, and the scalars alpha and beta. An operation graph holds an array of operations and a handle.

The graph says what you want computed and nothing about how. You hand it to a heuristics descriptor, which returns a ranked array of engine configurations, best first. Each one is an offer, not a promise: you build an execution plan from it, and a plan that fails to finalize is the engine telling you it cannot run your graph, so you try the next.

The plan that does finalize reports a workspace size, scratch you allocate and hand back.

Only at execution does anything learn an address. A variant pack carries an array of UIDs and a parallel array of device pointers, plus the workspace, and cudnnBackendExecute runs the plan against it. Build once, execute many times with different data.

Diagram: where an address enters a cuDNN graph. Three vertical bands left to right, with one arrow between each pair, and a dashed horizontal line marking where device pointers first appear (only in band three). Band 1, describe: three tensor boxes labelled x, w, y each showing dims, strides and a UID, plus a convolution box, feeding one operation box, which feeds one operation graph box. Caption "Six descriptors, no addresses." Band 2, choose: the graph feeding a heuristics box that fans out to a ranked column of engine configurations, the first one that finalizes highlighted, arrow to a plan box carrying a workspace size. Caption "A ranked list; the first that finalizes wins." Band 3, bind: a variant pack box holding two parallel arrays, three UIDs and three device pointers, joined by three horizontal lines, arrow to cudnnBackendExecute. Caption "Three UIDs, three pointers, by position." Alt text: "A cuDNN graph is built, planned and only then bound to memory. Six descriptors carry no addresses; the heuristics rank engine configurations; the variant pack matches three UIDs to three device pointers by position."

Why the verbosity is the feature

Compared with one cuBLAS cublasSgemm call, the graph API has more setup. Build one graph by hand before wrapping the setup in a helper.

cuDNN can fuse convolution, bias, and activation into one kernel, or fuse a whole attention block. A graph with several operations uses the same three calls with a longer array, and the heuristics are then choosing between kernels that exist as no separate function and could not be named by an enum. The old API had one algorithm enum per operation, so it could not express a fusion at all.

That is day 48's argument about your own kernels, applied to somebody else's.

The other thing to unlearn is that a heuristic result is a guarantee. It is a ranked list, and the documented workflow is to walk it: query mode A or B, take the first configuration whose plan finalizes, fall back to CUDNN_HEUR_MODE_FALLBACK if none does. Code that asks for exactly one configuration fails on the first unusual shape, at plan-build time, with a message about support that reads like a bug in your descriptors.

One convolution, graded twice

Full program in code/day83-cudnn/conv_graph.cu, with the PyTorch half in check_against_pytorch.py. The shape is a ResNet-ish layer at batch 2: x[2,64,56,56] convolved with w[64,64,3,3], pad 1, stride 1, dilation 1.

The numeric contract is written in the program, not in this paragraph. fp32 in, fp32 compute, fp32 out, on both sides. CUDNN_CROSS_CORRELATION rather than CUDNN_CONVOLUTION, because that is what every framework means by convolution. On the Python side torch.backends.cudnn.allow_tf32 is set to False because TF32 is on by default for convolutions on Ampere and newer and would reduce the inputs to ten mantissa bits.

Turing GPUs have no TF32 path, but the setting keeps the comparison consistent on newer architectures.

Every cuDNN status goes to a real branch. There is no second error macro, because CUDA-CODE-STYLE.md forbids one and day 44 shipped a bug by discarding a cuBLAS status inside a lambda. One named function does the work:

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;
}

A tensor is five attributes and a finalize. Note that rank is 4 for both the image and the filter, and that the UID goes in here, hundreds of lines before any pointer exists:

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");
}

And the late binding, which is the whole model in nine lines. The UIDs and the pointers are two arrays in the same order, so swapping two entries is a wrong answer with no error anywhere:

        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;
        }

The tolerance is derived, not chosen. Each output is a sum of C*R*S = 576 products, and cuDNN, PyTorch and the float64 CPU reference are free to add those 576 in three different orders. The course's rule for a reduction of depth K is rtol = max(1e-5, 4 * 2^-23 * sqrt(K)), which at K = 576 beats the flat fp32 figure, so the depth sets the bar. Both programs print that arithmetic next to the diff, so you can see the bound was not tuned until the test went green.

Day 68 is why reordering a sum moves the last bits.

Nothing here is timed. Day 83's check is the harness, and a cuDNN-versus-anything number would need the engine, the layout and the math mode pinned on both sides before it meant anything.

Results

Verified on hardware. The graph ran on a Tesla T4 (driver 580.173.02, CUDA 12.6) against standalone cuDNN 9.13.1. The separate PyTorch check used torch 2.13.0+cu126 and its bundled cuDNN 9.10.2. Both programs exited 0.

What the run reports Where it comes from
CUDNN_VERSION and cudnnGetVersion() header and shared object
engine configurations returned by mode A heuristics query
heuristic rank chosen, and any ranks refused plan finalize walk
engine global index the chosen engine config
workspace bytes the finalized plan
largest |cuDNN minus reference| float64 Kahan CPU check
largest |cuDNN minus torch| check_against_pytorch.py
The four predictions all held:
  1. The two standalone version lines agreed. CUDNN_VERSION and cudnnGetVersion() both printed 91301.
  2. Mode A returned eight configurations and rank 0 finalized. The chosen plan reported engine global index 12.
  3. Both checks passed, and the GPU paths were closer to each other. The largest cuDNN-to-float64-reference difference was 1.268e-05. The largest standalone-cuDNN-to-PyTorch difference was 0.000e+00 across all 401408 outputs, despite the two processes loading different cuDNN versions.
  4. The workspace was nonzero: 409856 bytes.

Record the engine global index with the output differences. It names the kernel family that produced the answer and makes two cuDNN runs easier to compare.

The standalone C++ graph was also built and run natively with CUDA 13.0 (V13.0.88) and cuDNN 9.13.1. It reproduced every reported value above, including engine index 12, 409856 workspace bytes and the 1.268e-05 largest CPU-reference difference, then exited 0. A separate torch 2.13.0+cu126 environment, built against CUDA 12.6 and reporting cuDNN 91301, consumed the files written by that CUDA 13 C++ run and matched all 401408 outputs exactly.

That is a cross-framework correctness check across two toolkit environments. It is not evidence that the PyTorch process used CUDA 13.

Run it yourself

The main setup constraint is the library package, not a specific GPU. The exact CUDA 12 packages used by this run have moved to conda-forge's broken label, so they remain downloadable but no longer resolve from live solver repodata. Install the three archived builds by URL:

micromamba install -c conda-forge "cuda-version=12.6" \
  https://api.anaconda.org/download/conda-forge/cudnn/9.13.1.26/linux-64/cudnn-9.13.1.26-hbcb9cd8_0.conda \
  https://api.anaconda.org/download/conda-forge/libcudnn/9.13.1.26/linux-64/libcudnn-9.13.1.26-hf7e9902_0.conda \
  https://api.anaconda.org/download/conda-forge/libcudnn-dev/9.13.1.26/linux-64/libcudnn-dev-9.13.1.26-h58dd1b1_0.conda

Same version number, two builds, one per CUDA major. Confirm the exact package builds and that ldd conv_graph resolves libcudnn.so.9 from your chosen environment before trusting output. PyTorch goes in a separate environment, because its wheel bundles its own CUDA runtime and its own cuDNN.

Both recipes, and what the run must capture, are in the README.

Exercise

Change the input tensor from NCHW to NHWC by editing its stride array only, leave the dimension array alone, and rerun. Report the engine global index before and after, and say in one sentence what the heuristics did with the extra freedom.

Time: 25 to 40 minutes. Submit: the two engine indices, the two workspace sizes, and the sentence.

Check: the harness is the program you already have. The CPU reference indexes x by ((n*C + c)*H + h)*W + v, so changing the strides without changing how the input is laid out in memory prints a first-mismatch line and exits nonzero. Getting it right means writing the input in the new order too.

Hint 1 The dimension array says what the tensor is. The stride array says where each element lives. Which of the two does an NCHW to NHWC change actually alter, and what does that imply about the bytes you copied to the device?
Hint 2 For NHWC with dims still ordered `{N, C, H, W}`, the stride of `C` becomes 1 and the stride of `W` becomes `C`. Write the strides out for both layouts on paper before you touch the file, then look at what the engine index does rather than at the clock.
Solution The strides become `{H*W*C, 1, W*C, C}` and the host buffer has to be filled in that order, which is what the reference check enforces. The engine index usually moves, because NHWC puts the channel dimension contiguous and that is the layout the [tensor core](/glossary/tensor-core) paths want. On a T4 at fp32 the heuristics may pick the same family anyway, and an index that does not move is a real answer, not a failed exercise.

A layout is a stride array, and the choice of engine follows from it. That is why frameworks fight about memory format at all, and why the graph API has no layout enum to argue with.

Pitfalls

Your build links the wrong CUDA major and nothing says so until runtime. conda-forge ships cudnn 9.13.1.26 twice, build hbcb9cd8_0 against CUDA 12 and build h886f0b6_0 against CUDA 13. The tested CUDA 12 build has since moved out of live solver repodata, so use the README's three exact archived URLs and confirm the package builds and loaded libcudnn.so.9. Day 44 lost a session to the cuBLAS form of this.

cudnnGetVersion() and CUDNN_VERSION disagree. The macro is what the header said at compile time, the call is what the loaded shared object says now, and since cuDNN 9.0 both encode as MAJOR*10000 + MINOR*100 + PATCH. A mismatch means two installs are on the path. Printing both is the useful answer to "how do I verify my cuDNN installation" (https://stackoverflow.com/questions/31326015/how-to-verify-cudnn-installation , one of the most-viewed questions in this corner of Stack Overflow, checked 2026-09-01); reading a header tells you nothing about the library that will load.

Your output is a plausible tensor of wrong numbers. You set CUDNN_CONVOLUTION where every framework means CUDNN_CROSS_CORRELATION. The two differ by rotating the filter 180 degrees, so a symmetric filter hides the bug completely and an asymmetric one fails everywhere. Check against a framework, not against your own eye.

Plan finalize fails and you read it as a bug in your descriptors. An engine configuration that cannot run your graph reports that at plan-build time. It is the documented way to ask "do you support this", so treat a failure as a signal to try the next rank and only give up after the fallback heuristic. The program prints every refused rank for that reason.

Your kernel gets an answer, then leaks. Descriptors are heap allocations behind an opaque pointer, and cudnnBackendDestroyDescriptor is the only way back. Compute Sanitizer will not point at the line that made one, so the program registers every descriptor in one bag and empties it on every exit path. Day 61 covers finding the ones you miss.

Go deeper

Next

Day 84 goes one layer down to CUTLASS and CuTe, where the layout algebra this exercise poked at with two stride arrays becomes the whole programming model, and where you can read a state-of-the-art kernel instead of asking a heuristic for one. Day 87 comes back up: the same convolution, wrapped as a PyTorch operator, so your kernel and NVIDIA's sit behind the same call.