100 DAYS OF GPU Course preview
Day 69Module 7
in-technical-review

One CUDA binary for several GPU generations

PyTorch prints this to thousands of people who just bought a new card: "NVIDIA GeForce RTX 5090 with CUDA capability sm_120 is not compatible with the current PyTorch installation. The current PyTorch install supports CUDA capabilities sm_75 sm_80 sm_86 sm_90." (The template is incompatible_device_warn in https://github.com/pytorch/pytorch/blob/main/torch/cuda/__init__.py , checked 2026-08-29.) The library and card can both work, yet the first kernel still fails because the binary lacks code for that architecture. The build's architecture list causes the failure.

This lesson builds one binary with a native path for compute capability 7.5 and a PTX path for compute capability 9.0. The program prints which path it uses and why.

What nvcc packs into the file

nvcc compiles your device code once per -gencode pair, and each pair names two architectures. The virtual one, compute_75, fixes which instructions the compiler may use and produces PTX, the portable assembly from day 46.

The real one, sm_75, is a physical chip, and asking for it produces SASS, machine code that runs on exactly that generation. A gencode pair like arch=compute_75,code=sm_75 embeds only the SASS; a pair like arch=compute_90,code=compute_90 embeds only the PTX. All of it lands in one executable, the fat binary (https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#gpu-compilation , checked 2026-09-01).

At load time, the driver searches the binary for code the device can run. SASS whose architecture matches the device's compute capability runs without compilation.

If no matching SASS exists, the driver takes the best-matching PTX at or below the device's capability and compiles it on the spot, which is JIT compilation: a real compile, paid on first launch, cached on disk afterwards (https://developer.nvidia.com/blog/cuda-pro-tip-understand-fat-binaries-jit-caching/ , checked 2026-09-01). If neither exists, the launch fails with the string in the Pitfalls section.

The build line in this lesson, sm_75 SASS plus compute_90 PTX, therefore uses native code on a compute capability 7.5 device and JIT compilation on a compute capability 9.0 device.

One exception matters: architecture-specific targets with an a suffix, like sm_90a, do not JIT forward to anything (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions-wgmma-mma , checked 2026-08-29). Day 78 covers those targets; this lesson uses the portable path.

Diagram: one file, two payloads, two shopping trips. A fat binary drawn as a box with two slots on the left, sm_75 SASS and compute_90 PTX; a T4 and an H100 on the right, one arrow each back into the box. Band 1, the box alone. Caption "one executable, 2 device payloads". Band 2, the T4's arrow lands on the SASS slot. Caption "CC 7.5: exact SASS match, 0 compiles at load". Band 3, the H100's arrow skips the SASS, lands on the PTX slot, and a gear marks the driver compiling it. Caption "CC 9.0: no sm_90 SASS, so the driver JITs the PTX once and caches it". Alt text: "A fat binary holds sm_75 machine code and compute_90 PTX. The T4 loads the machine code directly. The H100 finds no matching machine code, so the driver compiles the PTX at first launch."

The macro that is never defined where you test it

__CUDA_ARCH__ looks like a platform macro from C, a _WIN32 you can check anywhere and branch on. It is not.

It exists only while nvcc compiles device code, once per -gencode target, and host code is compiled with the macro absent. The preprocessor evaluates an undefined identifier in #if as 0, so a host-side #if __CUDA_ARCH__ >= 900 is not an error. The branch is dead on every host.

The program below keeps this bug so you can see its output.

The subtler mistake is deciding what the device can do from prop.major >= 9. A version check answers "which generation is this", and every new generation forces you to ship an update that extends the table.

A feature probe asks the question you actually have. Day 28 measured cudaDevAttrCooperativeLaunch returning 1 on a compute capability 7.5 GPU, then measured the real constraint the attribute guards: a cooperative grid of 160 blocks launches and 161 is refused. The same probe style answers whether a device can launch thread block clusters, cudaDevAttrClusterLaunch, which the compute capability tables say arrives at CC 9.0 (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html , checked 2026-08-29).

The probe keeps being right on cards that did not exist when you compiled.

Where each number is allowed to come from, and one bug left in

Full program in code/day69-portability/portability.cu. It runs the warp shuffle reduction shape from day 23 and day 24 over 4 Mi ones, so the answer, 4194304, is exact in float and the gate demands equality. Everything else follows from where each number is allowed to come from.

Per-architecture numbers in device code come from __CUDA_ARCH__, once per target. CC 7.5 keeps 1024 resident threads and 16 resident blocks per SM; CC 9.0 keeps 2048 and 32. The launch bounds bake each ceiling into that target's SASS:

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

Per-architecture numbers in host code come from the runtime, never from the macro. The host asks the attributes what the device can do and asks the occupancy query how many blocks per SM the loaded code sustains, then sizes the grid from the answer:

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

The broken version stays in the file, labelled, next to its fix. The harness pins the bug to its known wrong answer, so the diff between the two functions is behavior you can run, not a comment you must trust:

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

A --force-cc 90 flag overrides only the dispatch decision, so the compute capability 9.0 host path, label, and version-table expectation run on an older GPU.

This tests host logic, not compute capability 9.0 hardware. The compute_90 PTX path has compiled but has not run on matching hardware.

Results

Measured on a Tesla T4 with driver 595.84 and CUDA 12.6 (V12.6.85) on 2026-09-01. All supported-image runs passed their gates.

CUDA 13.0 reproduced the artifact listings and every run-time verdict: the fat binary carried two sm_75 ELF images and one compute_90 PTX image, native and forced dispatch passed, the PTX-only binary JITed successfully with its cache disabled, and the sm_90-only image was rejected on the T4 with exit 1. No behavioral drift was observed, and no Hopper-silicon execution is claimed.

Step Binary Measured
cuobjdump -lelf / -lptx fat 2 sm_75 ELF images, 1 sm_90 PTX image
default run fat Turing path, cooperative 1, cluster 0, 4 blocks/SM x 40 = 160
--force-cc 90 run fat table 8, probe 4, probe wins, all gates pass
CUDA_CACHE_DISABLE=1 run compute_75 PTX only all gates pass at 4 blocks/SM x 40 = 160
run on T4 sm_90 SASS only occupancy query returns no kernel image is available, exit 1

Three predictions held: the default probe reproduced day 28's 160-block ceiling, the forced host path reported table 8 against probe 4 and obeyed the probe, and the PTX-only binary passed the same gates.

Two details differed from the predictions. cuobjdump listed two sm_75 cubins rather than one, both from the same requested target, plus the predicted sm_90 PTX. The sm_90-only binary failed before launch at cudaOccupancyMaxActiveBlocksPerMultiprocessor, and the exact error was no kernel image is available for execution on the device, not invalid device function.

An occupancy query can force the driver to load the kernel image, so unsupported-image handling belongs around probes as well as launches.

Portability depends on where each number came from. The macro answers once per compile target, while the probe answers once per device. Code that confuses the two may fail on a later generation.

Run it yourself

The experiment needs a shell, a supported GPU, and three builds. The repo README scripts the three nvcc lines and four runs:

nvcc -std=c++17 -O3 -gencode arch=compute_75,code=sm_75 \
    -gencode arch=compute_90,code=compute_90 \
    -o portability portability.cu

Nothing here needs an API newer than CUDA 12.x, and cuobjdump ships with the toolkit (a conda toolchain needs the cuda-cuobjdump package). Compiler Explorer cannot run the full exercise. It needs three builds plus an environment variable, so use a shell.

Exercise

Change kThreadsPerBlock from 256 to 32, rebuild and run, then to 1024, rebuild and run. Before each run, predict the occupancy probe from the two CC 7.5 ceilings the launch-bounds block names, then check the probe: line.

Time: 25 to 40 minutes. Submit: the three probe values and two sentences naming which hardware ceiling set each one.

Check: the program's three gates are real branches returning EXIT_FAILURE: the buggy host probe must still return 0, the runtime probe must be at least 750, and the sum must equal 4194304 exactly. A failure names the gate and both values.

All three must keep passing at every block size, because the grid is sized from the probe rather than from a constant that just changed. Harness contract at /reference/harness.

Hint 1

The probe never exceeds either of two quotients: resident threads per SM over your block size, and the resident block slots per SM. Write both down for each block size before running anything.

Hint 2

At 32 threads per block the thread ceiling allows 1024 / 32 = 32 blocks. The probe will not say 32. What other number in the launch-bounds block is smaller?

Solution

At 256: 1024 / 256 = 4, under the 16-slot limit, so 4 (thread ceiling). At 32: the thread ceiling allows 32, but a CC 7.5 SM has 16 resident block slots, so 16 (slot ceiling). At 1024: 1024 / 1024 = 1 (thread ceiling again).

On a compute capability 9.0 GPU, the same three answers become 8, 32, and 2 without a host-code change. The probe folds every per-architecture cap into one number, so the table lives in exactly one place, the device pass that compiled the kernel.

Pitfalls

Your program compiles cleanly and fails at the first kernel. You built above your device, so the binary holds no image the driver can use. For example, -arch=sm_80 is above a compute capability 7.5 device. Day 3 reproduced the exact failure: the launch fails with error 209, no kernel image is available for execution on the device.

Rebuild with a -gencode set whose floor is at or below your card. More on this error.

The same build error may first appear in a probe. If the code calls cudaFuncGetAttributes or an occupancy query before launching, the missing image surfaces there as invalid device function instead, and people hunt for a typo in the kernel name. Check what the binary contains, cuobjdump -lelf, before checking the code. More on this error.

nvcc refuses the architecture at build time. An older toolkit does not know newer targets; the shape of the message, from a verified run of a CUDA 13.3 nvcc handed -arch=sm_60 (checked 2026-08-29), is nvcc fatal : Unsupported gpu architecture 'sm_60', with three spaces before the colon, and your unsupported target quoted where sm_60 stands. New cards need new toolkits; day 3 covers the pairing. More on this error.

A host-side #if __CUDA_ARCH__ silently takes the else branch. The macro does not exist in host compilation and #if treats it as 0, so the check compiles, runs and always answers the same thing. Guard device code with it freely; in host code, read cudaDeviceProp.

You shipped sm_90a and expected later architectures to run it. Architecture-specific a targets never JIT forward, so a binary built only for sm_90a is Hopper-only in a way plain compute_90 PTX is not. Day 78 explains when to use the a targets.

Your dispatch table was right about every card you tested. PyTorch's table was also current until sm_120 shipped. A version check encodes the cards you knew about; an attribute probe encodes the question. Prefer the probe wherever one exists.

Go deeper

Next

Day 70 closes the module with broken programs: ten bugs from days 61 through 69 in one codebase. Use this module's tools to find them. The catalog of their error strings ships as its own page, because each string is a search query long before it is a lesson.