Day 67in-technical-review

CUDA with CMake: separate compilation and which SMs to ship

"How to let CMake find CUDA" has 205,686 views on Stack Overflow (https://stackoverflow.com/questions/19980412/how-to-let-cmake-find-cuda , checked 2026-08-29), and the answers that era produced teach find_package(CUDA), a module CMake itself has marked "Deprecated since version 3.10. Superseded by first-class CUDA language support" (https://cmake.org/cmake/help/v3.28/module/FindCUDA.html , checked 2026-09-01).

Copy one today and you use an obsolete build method. This page builds day 43 and day 44 the current way, wraps day 43's tile in a static library that cannot build without relocatable device code, and then opens every binary with cuobjdump to see what shipped.

The architecture list defines what the binary can run

CMake has treated CUDA as a language, not a package, since 3.8: project(... LANGUAGES CXX CUDA) and every .cu file compiles like any other source. The one decision it will not make well for you is the CUDA_ARCHITECTURES target property, added in CMake 3.18, which since policy CMP0104 "must be set to a non-empty value on targets that compile CUDA sources" (https://cmake.org/cmake/help/v3.28/prop_tgt/CUDA_ARCHITECTURES.html , checked 2026-09-01).

Each entry in that list becomes one gencode pair on the nvcc command line. 75 becomes -gencode arch=compute_75,code=[compute_75,sm_75]: ptxas builds an sm_75 cubin, the real machine code day 46 taught you to read as SASS, and the compute_75 PTX is packed in beside it so a card newer than Turing can JIT its own cubin at first launch.

The suffixes narrow the output: "An architecture can be suffixed by either -real or -virtual to specify the kind of architecture to generate code for. If no suffix is given then code is generated for both" (same property page).

So 75-real ships machine code for exactly one compute capability and nothing a future card can use. 75-virtual ships only the PTX, so every card compiles it on first launch.

Three special values arrived later: all and all-major in CMake 3.23, native in 3.24. Under this project's 12.6 toolkit, all-major means what nvcc says it means: "a compiled code image for all supported major versions (sm_*0), plus the earliest supported, and adds a PTX program for the highest major virtual architecture" (https://docs.nvidia.com/cuda/archive/12.6.3/cuda-compiler-driver-nvcc/index.html , checked 2026-09-01), which for 12.6 is five cubins, sm_50 through sm_90. That is a good default for a release binary, though device code size grows with the list.

native asks the build machine which GPU it has. That can suit a local build, but not a build server with no GPU.

If nobody sets the list, CMake falls back to the architecture it detected the compiler with, which tracks nvcc's own default: sm_52 on 12.6 ("sm_52 is used as the default value", same 12.6 manual page), sm_75 on 13.x. An unset architecture list on a 12.6 toolchain can therefore ship Maxwell code. The compute_52 PTX beside it still JIT-compiles on newer cards, which can hide the bad default.

Diagram: three architecture lists and what lands in one fatbin. Three horizontal bands, each an arch setting on the left and the resulting fatbin drawn as labelled sections on the right; cubin sections solid, PTX sections dashed. Band 1, 75: one solid sm_75 section plus one dashed compute_75 section. Caption "1 cubin + 1 PTX: runs on Turing now, JITs anywhere later." Band 2, 80-real: one solid sm_80 section, no dashed section. Caption "1 cubin, 0 PTX: on a T4 the launch fails, nothing can be JIT-built." Band 3, all-major under CUDA 12.6: five solid sections, sm_50 to sm_90, one dashed compute_90. Caption "5 cubins + 1 PTX: the release build, at five times the device code." Alt text: "Plain 75 packs one cubin and one PTX. 80-real packs one cubin and no PTX, so a T4 has nothing to run or build. all-major on CUDA 12.6 packs five cubins plus compute_90 PTX."

Separate compilation has a device-code cost

In host C++, splitting code across translation units is routine. That can make CUDA_SEPARABLE_COMPILATION ON look like a harmless default. The property has existed since CMake added CUDA support in 3.8.

Device code has a different cost. nvcc's default is whole-program compilation per translation unit: ptxas sees every function body at once, inlines the small ones, and allocates registers across the result.

The moment a kernel calls a __device__ function that lives in another .cu, that stops being possible. The code must be compiled relocatable (-rdc=true, which is exactly what the CMake property passes) and stitched by a device link step, and "Relocatable device code must be linked before it can be executed" (nvcc 12.6 manual, same page as above).

The call the compiler could not see through stays a real call in the SASS, with ABI register costs, unless you enable device link-time optimisation. That is why this course's build keeps separate compilation off globally and turns it on only for this lesson's targets. Enable it when a library boundary needs it.

One project with a checked failure case

The full project is code/day67-cmake/, one CMakeLists.txt and three .cu files. A few rules keep its transcript easy to compare with the predictions below.

Days 43 and 44 are built in place, not copied. The targets name ../day43-matmul-1/matmul_steps.cu and ../day44-matmul-2/matmul_registers.cu where those lessons left them, so this build inherits their verified sources byte for byte, plus one extra target that rebuilds day 43 with relocatable device code forced on, so the run can diff build time and cuobjdump -res-usage between the two.

The library's need for -rdc=true is real, and the bug switch removes it. The mm67 library keeps day 43's tiled kernel in two variants that differ only in where a one-line FMA helper lives: same translation unit, or across in mm_step.cu. -DDAY67_RDC=OFF is a labelled, deliberate misconfiguration that turns separable compilation back off so you can watch ptxas fail on the cross-unit call.

A harness gates the library. day67-check diffs both kernels against a CPU reference at n = 500, chosen so the tile's edge guards fire, and returns EXIT_FAILURE at the first mismatch. Every experiment below ends by asking it.

Set the architecture list above project():

cmake_minimum_required(VERSION 3.24)

# CMAKE_CUDA_ARCHITECTURES must hold a value before the CUDA language is
# enabled, or CMake falls back to the architecture it detected the
# compiler with, which tracks nvcc's own default: sm_52 on the CUDA 12.6
# toolkit this project verifies with, hardware this course dropped on
# day 3. Setting a default here, above project(), is the difference
# between shipping Turing code and shipping Maxwell code nobody asked
# for. -DCMAKE_CUDA_ARCHITECTURES=... on the command line still wins.
if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
  set(CMAKE_CUDA_ARCHITECTURES 75)
endif()

project(day67-matmul-lib LANGUAGES CXX CUDA)

The reuse targets, which is all it takes to give two finished lessons a build system (CUDA::cublas comes from find_package(CUDAToolkit), the modern module, CMake 3.17):

add_executable(day43-matmul-steps
  ${CMAKE_CURRENT_SOURCE_DIR}/../day43-matmul-1/matmul_steps.cu)

add_executable(day44-matmul-registers
  ${CMAKE_CURRENT_SOURCE_DIR}/../day44-matmul-2/matmul_registers.cu)
target_link_libraries(day44-matmul-registers PRIVATE CUDA::cublas)

The library, the harness and the bug switch:

# Deliberate bug switch, day 67: configure with -DDAY67_RDC=OFF and the
# build stops inside ptxas on mm_lib.cu, which meets a call to fmaStep()
# and has no body for it. The transcript of that failure is one of this
# day's evidence artifacts. ON is the fix, kept in the same file so the
# two configurations differ by one cache variable.
option(DAY67_RDC "Compile mm67 with relocatable device code" ON)
if(DAY67_RDC)
  # The variable initialises CUDA_SEPARABLE_COMPILATION on every target
  # created after this line, library and harness alike.
  set(CMAKE_CUDA_SEPARABLE_COMPILATION ON)
endif()

add_library(mm67 STATIC mm_lib.cu mm_step.cu)
target_include_directories(mm67 PUBLIC ${CMAKE_CURRENT_SOURCE_DIR})

add_executable(day67-check check_mm.cu)
target_link_libraries(day67-check PRIVATE mm67)

These two lines of C++ require separate compilation. First, a declaration in mm_lib.cu:

extern __device__ float fmaStep(float acc, float a, float b);

with its body in mm_step.cu:

// One multiply-add. The body is nothing; the address is everything. Because
// it is defined here and called from mm_lib.cu, the call in the kernel is a
// real SASS CAL until link-time optimisation is asked for, and day 46's
// cuobjdump can show it.
__device__ float fmaStep(float acc, float a, float b) {
    return acc + a * b;
}

An FMA is costly to put behind a translation-unit boundary. It makes the cost of a device call that the compiler cannot inline easy to measure.

Results

Run on the project's Tesla T4 (driver 595.84, CUDA 12.6 V12.6.85, CMake 4.4.3, Ninja) on 2026-09-01. Six transcripts in code/day67-cmake/evidence/. Two predictions held exactly, one had a different error string, one had a different register count, and two are unsettled because the command that would have settled them was wrong.

The CUDA 13.0 matrix completed with the expected success and failure legs. The 75, all-major and native builds passed both 250,000-element checks; the 80-real build succeeded and then refused the T4 at run time; and the RDC-off build again failed on unresolved fmaStep.

The corrected cuobjdump -all -lptx command resolved the default build's open question by listing three sm_75 PTX files alongside its one sm_75 cubin. CUDA 13 no longer accepts architecture 52, so that configure leg now fails early with Unsupported gpu architecture 'compute_52', an expected compatibility boundary rather than a matrix failure.

Build	`cuobjdump -lelf`	binary size	`day67-check`
`75` (the default)	one `sm_75` cubin	1,035,600 bytes	both kernels pass, exit 0
`52` (what a bare 12.6 nvcc picks)	one `sm_52` cubin	not compared	not run
`80-real`	one `sm_80` cubin	not compared	fails, exit 1
`all-major`	`sm_50`, `sm_60`, `sm_70`, `sm_80`, `sm_90`	1,097,072 bytes	both kernels pass, exit 0
`native`	one `sm_75` cubin	not compared	both kernels pass, exit 0
`-DDAY67_RDC=OFF`	never built	never built	never built

The corrected command settled the default build. The CUDA 12.6 cuobjdump -lelf capture lists exactly one ELF, day67-check.1.sm_75.cubin. cuobjdump -lptx printed no PTX at all, and printed its own explanation: cuobjdump info : No PTX file found to extract from '.../build/t75/day67-check'. You may try with -all option. The same message came back from the 80-real, all-major and native builds, so an empty PTX listing here is a property of the command, not of the architecture list.

day67-check links a separately compiled library, which is the case -all covers. The CUDA 13 rerun used cuobjdump -all -lptx and listed three files: day67-check.1.sm_75.ptx, .2.sm_75.ptx, and .3.sm_75.ptx. The default 75 build therefore carries both one sm_75 cubin and three sm_75 PTX files.
The error string differed. The 80-real build did compile clean, did link clean, and did fail at run time with exit 1 after the allocations and copies, all as predicted. The string is not the no-kernel-image error:
```
CUDA error /tmp/100dc/code/day67-cmake/mm_lib.cu:140: cudaGetLastError(): named symbol not found
```
The check that fired is the cudaGetLastError() after the first launch in mm_lib.cu, so the shape of the failure was right and the name was wrong.

Day 69 made the same mistake in the same week on the same card, an sm_90-only build run on a T4, and got no kernel image is available for execution on the device. Two binaries, one wrong architecture list each, two different runtime complaints. This binary uses relocatable device code, while day 69's does not.

"The wrong-architecture error message" is not a thing you can grep for; the fix is the same either way, and it is in the arch list.
The cubin list and size direction held, but the size change was small. all-major shipped five ELFs, sm_50, sm_60, sm_70, sm_80, sm_90, in that order and nothing else. The ls -l comparison: 1,097,072 bytes against 1,035,600, so five times the device code costs 61,472 bytes, 5.9 percent of the file.

The binary also carries the host program and the static runtime, so device code is a small part of its size. This result supports shipping the full list.

The PTX half of this prediction is unsettled for the reason in point 1.
native selected the installed GPU, but its PTX result was unsettled. One sm_75 ELF, the same listing plain 75 produced, and the harness passed. Whether a virtual architecture rode along is the question -lptx -all will answer.
Held exactly, down to the mangled name. -DDAY67_RDC=OFF stopped the build at [1/12], inside ptxas, on mm_lib.cu:
```
FAILED: [code=255] CMakeFiles/mm67.dir/mm_lib.cu.o
ptxas fatal   : Unresolved extern function '_Z7fmaStepfff'
```
mm_step.cu compiled without complaint on the very next line, because it defines fmaStep and calls nothing across a unit boundary. Ninja exited 255.
The register count differed in the opposite direction. From cuobjdump -res-usage on the passing 75 build: matmulTiledExternFma takes 40 registers and matmulTiledLocalFma takes 45, both with 2,048 bytes of shared memory and no local memory. The prediction said the extern variant would be at least as high.

It is five registers lower. In an ABI call, the arithmetic that was inlined into the local variant's register budget now lives in the callee's frame, and fmaStep reports its own 24. Fewer caller registers do not show an improvement because the callee now holds part of the work.

The CAL half of the prediction is unsettled: the capture counted CAL over the whole binary, 16 of them, with no per-kernel attribution, so it cannot say the local variant has none. The build-time comparison and the plain-against-rdc -res-usage diff the README asks for were not captured at all, so "the rdc rebuild builds slower" has no evidence on this page and remains unverified.

Run cuobjdump -lelf -lptx -all to inspect what a build contains. Two of the six predictions remain open because the corrected PTX command was captured for the default build, but not for the all-major and native builds.

Run it yourself

You can configure, build, test -DDAY67_RDC=OFF, and run cuobjdump on any machine with a CUDA 12.x toolkit and CMake 3.24 or newer. The harness needs a GPU that meets min_cc 7.5.

The day67-check target runs that harness, and the day 44 target also needs cuBLAS. Allow a few extra minutes for all-major, which compiles every kernel five times.

Exercise

Break the harness without touching a single source file, then fix it two different ways and prove, with cuobjdump, that your two fixes ship different things.

Time: 30 to 40 minutes. Submit: the failing run's error line, the two cuobjdump -lelf -lptx listings from your fixes, and two sentences on which fix you would ship and why.

Check: day67-check is the harness. Broken, it must exit nonzero with a CUDA error line pointing into mm_lib.cu; after each fix it must print pass for both kernels and mm67 library check: PASS. If your "broken" build passes, look at your dump: you left a PTX section in, and the runtime built a kernel from it.

Hint 1

The harness takes no arguments and the sources are off limits, so the only thing you can change is what the build puts in the binary. What has to be missing for a launch to have nothing to run, and nothing to build one from?

Hint 2

80-real is one answer from the README. For the two fixes: one adds machine code for your card, one adds no machine code at all and still passes. What does the second one cost, and at which launch is it paid?

Solution

Break: -DCMAKE_CUDA_ARCHITECTURES=80-real (any list with no cubin for your card and no PTX). On a T4 every launch fails with the no-kernel-image error. Fix one: 75-real, and the dump shows one sm_75 ELF, no PTX; runs only on Turing class cards.

Fix two: 75-virtual, and the dump shows no ELF and one compute_75 PTX; it passes on the T4 and on newer cards, paying a JIT compilation on the first launch of each kernel.

Ship the fatter list (75 or a real list plus one virtual) for anything that leaves your machine: cubins for the cards you know, PTX for the ones that do not exist yet. The general rule: the arch list is not a compiler tuning knob, it is the population of machines your binary has promised to run on.

Pitfalls

Your build works and your kernels are mysteriously slow on a new card. Nobody set the architecture list, so CMake took nvcc's default, sm_52 on a 12.6 toolkit, and every launch runs JIT-built code from compute_52 PTX tuned for no card you own. Set CMAKE_CUDA_ARCHITECTURES explicitly, above project(). Day 3 covers the same trap on the raw nvcc command line.

You copied a find_package(CUDA) answer and it half works. That is the deprecated FindCUDA module (CMake 3.10), with its own variables and flags that fight the language support. Delete it: LANGUAGES CUDA in project() for the compiler, find_package(CUDAToolkit) for the libraries.

You asked an old toolkit for an architecture it never met, or a new one for an architecture it dropped. Both die at configure or compile with nvcc fatal : Unsupported gpu architecture 'sm_60', three spaces before the colon, sm_60 being real 13.x behaviour since Pascal support was removed (error page). The fix is matching your list to the toolkit's supported range, not downgrading CMake.

You turned separable compilation on globally "to be safe". Every kernel call the optimiser could have inlined is now a candidate ABI call, register counts drift up, and every target pays a device-link step. Turn it on for the target whose translation units actually share device code, the way this project does, and nowhere else.

Your static library builds but the executable fails at device link with undefined device symbols. The library compiled relocatable, but nothing performed the device link. In CMake that happens when the final link is done by a target with no CUDA in it; give the executable a .cu file or set CUDA_RESOLVE_DEVICE_SYMBOLS on the library. This day's harness links from a .cu, so the problem never appears here.

You proved nothing shipped sm_90 and closed the bug. A stale build directory happily links yesterday's objects. Fresh build directory per architecture experiment, as the README insists.

Go deeper

CMake, CUDA_ARCHITECTURES target property, including all, all-major, native and the -real/-virtual suffixes: https://cmake.org/cmake/help/v3.28/prop_tgt/CUDA_ARCHITECTURES.html (checked 2026-09-01)
CUDA Compiler Driver NVCC (12.6), "GPU Compilation", for gencode, the fatbin and the default architecture: https://docs.nvidia.com/cuda/archive/12.6.3/cuda-compiler-driver-nvcc/index.html (checked 2026-09-01)
CUDA Binary Utilities, for every cuobjdump flag this page leans on: https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html (checked 2026-08-30)
NVIDIA, "Building Cross-Platform CUDA Applications with CMake": https://developer.nvidia.com/blog/building-cuda-applications-cmake/ (checked 2026-09-01)

Day 68 stays inside the binary you just learned to ship and asks a harder question: whether the same build produces the same bits, run after run, card after card. And day 69 takes this page's arch-list experiments to their conclusion, one binary that runs on a T4 and an H100 and takes different paths on each.