[CUDA] Fix CUDA architecture gating for SM12x, family/arch-specific targets and -virtual entries by ConnorBaker · Pull Request #33373 · microsoft/onnxruntime

Connor Baker (ConnorBaker) · 2026-10-14T19:22:25Z

Description

cmake/external/cuda_configuration.cmake and the CUDA provider cmake files decide which CUDA kernels to build by
string-matching CMAKE_CUDA_ARCHITECTURES entries. Several of those tests mishandle architecture spellings that CMake
and nvcc accept: a/f suffixes, -real/-virtual, and compute capabilities other than the exact one tested. The
results range from kernels silently left out, to kernels built that can never run, to dispatch throwing at run time.

This adds two helpers to cuda_configuration.cmake:

onnxruntime_parse_cuda_arch() splits an entry <cc>[a|f][-real|-virtual] into compute capability, suffix and
whether it is -virtual.
onnxruntime_cuda_archs_have_sass() / onnxruntime_cuda_archs_have_code() report whether any entry produces SASS
(or SASS or PTX) for a range of compute capabilities.

It then uses them consistently. One commit per change:

SM120 kernels for 121a/120f/121f/120a (cmake: match CUDA architectures with a/f/-real suffixes).
- Kernels left out: the SM120 TMA grouped GEMMs, matmul_block_scaled_fp4_sm120.cu,
  COMPILE_BLACKWELL_SM120_TMA_GROUPED_GEMMS and ORT_ENABLE_BLOCKQUANT_SM120 were gated on
  "120" IN_LIST ... OR "121" IN_LIST CMAKE_CUDA_ARCHITECTURES_ORIG. Bare entries keep their suffix there, so
  121a (DGX Spark/GB10), 120f and 121f silently built none of it.
- Kernels emptied: conversely, the EXCLUDE_SM_* loop matched ^120[af]?$, so any SM121-only build defined
  EXCLUDE_SM_120. That empties the SM120 MoE GEMM instantiations, which the dispatcher still references for
  121/121-real.
- Fix: ORT_HAS_SM12X (12.0–12.1, any spelling) replaces the three IN_LIST tests, and EXCLUDE_SM_120
  follows it. SM121 then uses the SM120 kernels compiled for the requested 12.x targets; [CUDA] Fix QMoE SM120 configuration dispatch on SM121 devices #33225 lets the runtime
  dispatch accept SM121 for SM120 configurations.
SM90 TMA kernels only when SM90 is targeted (cmake: only build SM90 TMA kernels when SM90 is targeted).
- The problem: the SM90 TMA OBJECT library (TMA warp-specialized MoE GEMMs, fpA_intB SM90 launchers, DeepGEMM)
  is compiled only for 90a-real, but it was built for any architecture ≥ 90. Code built for sm_90a runs only on
  compute capability 9.0 (CUDA C++ Programming Guide, compute capabilities appendix; Blackwell compatibility guide),
  so for a build without SM90 these 341 translation units produce code that can never run. [CUDA] Fix QMoE MXFP4 W4A16 on SM120/SM121 and speed up FP4 decode #33067 removed the last
  runtime path that used them on SM120.
- Fix: ORT_HAS_SM90 gates the library, COMPILE_HOPPER_TMA_GEMMS/COMPILE_HOPPER_TMA_GROUPED_GEMMS, and
  DeepGEMM/USE_DEEP_GEMM. That's the configuration MSVC builds already use. EXCLUDE_SM_90 is defined exactly
  when ORT_HAS_SM90 is not.
SM80-generation kernels on newer GPUs (cmake: keep SM80-generation kernels whenever an SM80+ GPU is targeted).
- The problem: EXCLUDE_SM_80 was defined unless 80 itself was listed. So 121a, 120, 90a, 100a, and
  even 86;89 compiled the fused_moe_gemm_sm80_* launchers to nothing. Dispatch still selects them, and they hit
  ORT_THROW("Fused MoE SM80 kernels are not available ...") (from QMoE CUDA EP — FP4/FP8/WFP4AFP8 Quantized Mixture-of-Experts + MoE GEMM Refactor #28467). On SM100 and SM12x that path is the only
  one for unquantized FP16/BF16 MoE with a gated activation.
- Why the kernels are fine on newer GPUs: they are CuTe SM80 atoms (mma.sync, cp.async), which CUTLASS
  enables for any __CUDA_ARCH__ >= 800. They are compiled in the llm library for every requested architecture
  ≥ 7.5. TensorRT-LLM, where they come from, has no such throwing branch.
- Fix: 80/86/89 are kept when any requested architecture, SASS or PTX, has at least that compute capability.
  90/100/110/120 still require native code for that family. The gates now read the normalized architecture list,
  which also keeps them correct with Follow standard CMake semantics for CUDA architecture suffixes #29601.
SM100 family (cmake: clear EXCLUDE_SM_100 for every SM100-family architecture).
- The problem: 103a (B300) and 107a builds defined EXCLUDE_SM_100, although compute_100f covers 10.0,
  10.3 and 10.7 (Programming Guide, "Family-Specific Compatibility").
- Fix: use the 10.0–10.7 range. No C++ reads EXCLUDE_SM_100/EXCLUDE_SM_110 today; this makes them mean the
  same thing as EXCLUDE_SM_120.
Keep a/f on -real entries (cmake: keep explicit a/f suffixes of -real CUDA architectures).
- The problem: setup_cuda_architectures() dropped the suffix of every -real entry (Add support for CUDA architecture family codes #27278 widened the regex
  to accept f but still discarded it). 120f-real became 120a-real, which doesn't run on 12.1, and 100f-real
  became 100a-real, which doesn't run on 10.3/10.7. 107a-real became 107-real.
- Fix: <cc>f-real is kept verbatim, an explicit a is kept, and the normalized list is de-duplicated. The
  regex and the de-duplication line match Follow standard CMake semantics for CUDA architecture suffixes #29601's text, to keep the two changes compatible. The commit message
  describes how they combine.
XQA on SM121: dropped when this branch was rebased onto main on 2026-11-03. [CUDA] Add sm_121 to XQA arch-specific branches #33762 made the same
|| __CUDA_ARCH__ == 1210 change to utils.cuh and mha_impl.cuh, so the commit became empty.
-virtual in HAS_SM*_OR_LATER (cmake: count the -virtual architecture in HAS_SM*_OR_LATER).
- The problem: these were computed from CMAKE_CUDA_ARCHITECTURES_ORIG, which leaves out the -virtual entry.
  For 75;121-virtual, HAS_SM80_OR_LATER was undefined, so XQA's Launch() returned "XQA is only supported on
  Ampere (SM80) or newer GPUs" on the 12.1 GPU the PTX targets.
- Fix: include _LAST_VIRTUAL. This becomes a no-op on top of Follow standard CMake semantics for CUDA architecture suffixes #29601.
Native SASS fallback in the MSVC plugin LLM library (cmake: add native SASS for every SM12x entry in the MSVC plugin LLM library, from review).
- The problem: under MSVC the plugin's LLM library turns native SM120+ entries into <cc>-virtual and adds
  -gencode=arch=compute_120,code=sm_120 so the ordinary LLM kernels don't depend on JIT. That fallback only
  checked for the literal 120/120a. With change 1 the filter also converts 120f/120f-real, which became
  PTX-only, and no SM121 spelling ever had the fallback.
- Fix: onnxruntime_filter_cuda_archs() reports the compute capabilities it converted (REPLACED_SM120_CCS),
  and the plugin adds -gencode=arch=compute_<cc>,code=sm_<cc> for each.

Builds that spell their architectures the way the gates already expected (e.g. 80;86;89;90;100;120) are unchanged,
apart from the SM80 change (3), the -virtual change (7) and the MSVC plugin fallback (8), which only add kernels
that were previously missing. The
literal __CUDA_ARCH__ and sm == checks elsewhere (fpA_intB == 890, XQA == 900, runtime sm == 90) select
SM89- or sm_90a-only code and are left as they are.

Motivation and Context

Building ONNX Runtime for DGX Spark / GB10 (compute capability 12.1) with CMAKE_CUDA_ARCHITECTURES=121a, as Nixpkgs
does, produced a CUDA provider that:

contained no native SM120/SM121 MoE or FP4 block-quant kernels. There is no onnxruntime_providers_cuda_sm120_tma
target in the build;
contained 291 sm_90a cubins next to 286 sm_121a ones in libonnxruntime_providers_cuda.so (cuobjdump --list-elf,
v1.30.0). The sm_90a ones can't run on the target GPU and are among the most expensive translation units to
compile;
threw at run time for unquantized FP16/BF16 fused MoE, because of EXCLUDE_SM_80.

Testing

A cmake -P harness runs setup_cuda_architectures() and the gates from all four changed cmake files on 30
architecture lists, with MSVC on and off. It checks each result against an independent model of which GPUs can run
the generated code: CUDA binary/PTX compatibility plus the family table.
- Lists tested include 121a, 121a-real, 121, 120, 120f, 121f, 120a-real, 121-virtual, 90a;121a,
  90, 90a-real, 90-virtual, 100a, 103a, 107a-real, 100f;120f, 86, 86;89, 75;121-virtual and
  80;86;89;90;100;120.
- Before: 59 violations on Linux and 51 with MSVC; 2 of 30 lists pass.
- After: 0 violations; 30 of 30 pass on both.
For change 8, a second cmake -P harness runs setup_cuda_architectures() and the plugin's MSVC LLM block (read
from onnxruntime_providers_cuda_plugin.cmake) on 22 lists, and checks that every GPU covered by a requested native
SM12x entry gets plain (non-a/f) native SASS. Before: 12 of 22 fail (every 120f and 121 spelling,
90a;121a, 100f;120f, ...). After: 0. This is configuration-level only; there is no Windows build.

Same changes backported to v1.30.0, with [CUDA] Fix QMoE MXFP4 W4A16 on SM120/SM121 and speed up FP4 decode #33067, [CUDA] Fix QMoE SM120 configuration dispatch on SM121 devices #33225 and [CUDA] Fix SM120 QMoE MoE GEMM TMA build failure with GCC #32584 (ConnorBaker/onnxruntime:v1.30.0-sm12x-fixes),
built with CMAKE_CUDA_ARCHITECTURES=121a on a GB10 (aarch64, CUDA 13.4, GCC 16) and compared with unpatched
v1.30.0 built the same way:

	v1.30.0	with these changes
`onnxruntime_providers_cuda_sm120_tma` / `_sm90_tma` built	no / yes	yes / no
cubins in `libonnxruntime_providers_cuda.so` (`cuobjdump --list-elf`)	577: 291 `sm_90a`, 286 `sm_121a`	287: all `sm_121a`
SM120 kernel symbols (`cuobjdump --dump-elf-symbols`)	0	15
`fused_moe` (SM80 fused MoE) kernel symbols	0	60
`libonnxruntime_providers_cuda.so` size	163 MB	114 MB
build wall time, 8 cores	about 57 min (other builds ran concurrently)	27 min

libonnxruntime_providers_cuda.so has no undefined MoeGemmRunner symbols.
A MatMul+Add+Relu session on the CUDA EP and on the TensorRT EP matches a NumPy reference (max abs error
5.5e-3, the same as unpatched v1.30.0).
Without [CUDA] Fix SM120 QMoE MoE GEMM TMA build failure with GCC #32584, the newly enabled SM120 MoE GEMM files failed to compile with GCC
(moe_gemm_tma_ws_launcher.inl:82: error: expression list treated as compound expression in functional cast).
That fix is already on main.
Not tested: running the newly built MoE kernels (QMoE / fused MoE) on the GPU.

This branch itself (on main) has not been compiled.

Related: #29601 (CUDA architecture suffix semantics; these changes are written to stack with it), #33067 and #33225
(SM120/SM121 MoE dispatch), #28467 (where the EXCLUDE_SM_80 throw came from), #27278 (family codes).

AI disclosure: this change was developed with AI assistance (Claude Code): tracing the gating logic, writing the
changes and the cmake -P harness, and the backport build. I reviewed the changes and the results.

azure-pipelines · 2026-10-14T19:22:35Z

Azure Pipelines: There may be pipelines that require an authorized user to comment /azp run to run.

Connor Baker (ConnorBaker) · 2026-10-14T20:02:56Z

@microsoft-github-policy-service agree

Copilot

🔵 Needs a closer look

The current branch remains uncompiled, and newly enabled MoE kernels lack GPU execution validation.

0 open findings

What changed in this PR

Corrects CUDA architecture handling so kernel selection matches requested targets, including SM121, architecture/family suffixes, and PTX-only entries.

Changes:

Adds shared architecture parsing and capability checks.
Aligns standard and plugin CUDA provider gates for SM80+, SM90, and SM12x kernels.
Explicitly applies SM120 XQA constants to SM121.

File	Description
onnxruntime/contrib_ops/cuda/bert/xqa/utils.cuh	Applies the shared-memory limit to SM121.
onnxruntime/contrib_ops/cuda/bert/xqa/mha_impl.cuh	Applies SM120 tile constants to SM121.
cmake/onnxruntime_providers_cuda.cmake	Corrects standard provider kernel gates.
cmake/onnxruntime_providers_cuda_plugin.cmake	Corrects plugin provider kernel gates.
cmake/onnxruntime_cuda_source_filters.cmake	Updates source selection and architecture filtering.
cmake/external/cuda_configuration.cmake	Adds parsing helpers and corrects normalization and capability gates.

🧠 Review effort: Balanced

💡 Configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

Connor Baker (ConnorBaker) · 2026-10-24T22:20:51Z

Failures should be unrelated. Are they flaky?

Connor Baker (ConnorBaker) · 2026-10-30T06:34:19Z

Bump

## Description Ports the architecture-table part of NVIDIA/TensorRT-LLM#12170 ("Compile XQA with sm_120f") to ORT's XQA decode kernels: `__CUDA_ARCH__ == 1210` (sm_121 — GB10 / DGX Spark and other Blackwell Arm devices) is now listed next to sm_120 in the two compile-time arch tables. sm_121 has the same 99 KB opt-in shared-memory limit as sm_120 and uses the same K/V tile configuration. This is explicit-parity hardening, not a user-visible bug fix: sm_121 already built and ran through the "unknown architecture" `#else` default, which has the same values as sm_120. See Testing for the evidence that the generated code is identical before and after. ## Summary of Changes ### XQA arch tables | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/bert/xqa/mha_impl.cuh` | Add `__CUDA_ARCH__ == 1210` to the `860 / 890 / 1200` branch that selects `preferedKHeadPartBytes = 64` and `cacheVTileSeqLen = 32` (previously reached only via the `#else` default). Mention sm_121 in the related comment and in the shared-memory-limit hint text. | | `onnxruntime/contrib_ops/cuda/bert/xqa/utils.cuh` | Add `__CUDA_ARCH__ == 1210` to the 99 KB `kMAX_SMEM_SIZE` branch. Before, 1210 fell through to the 48 KB default. `kMAX_SMEM_SIZE` is currently only used by a `#if 0` `static_assert`, so this has no runtime effect today; it just keeps the table correct. | ### Not ported from TRT-LLM #12170 (not applicable) The rest of #12170 changes TensorRT-LLM's NVRTC JIT wrapper (`-arch=sm_120f`, dropping the compute_89 PTX to sm_120 two-stage compile through nvPTXCompiler, and the `nvptxcompiler_static` link). ORT compiles XQA ahead of time with nvcc, so there is nothing to port. In TRT-LLM the default arch branch was an `#error` and `kMAX_SMEM_SIZE` was undefined for 1210, which is why sm_121 failed to compile there; ORT's `#else` fallbacks mean it never did. ## Do the shipped libraries include sm_121? Yes, so no packaging or CMake change is needed: - Linux aarch64 CUDA 13 wheel: `87-real;89-real;90-real;100-real;103-real;110-real;120-real;121-real` (`py-cuda13-packaging-pipeline.yml`). - Unified Linux aarch64 wheel (`py-cuda13-unified-packaging-pipeline.yml`): passes no `-a`, so it takes the `build_linux_python_package.sh` default for CUDA 13.0 on aarch64, which is the same list. - CUDA plugin (`plugin-cuda13-pipeline.yml` / `plugin-cuda-packaging-stage.yml`): Linux aarch64 uses the same list; Windows ARM64 uses `120-real;121-real`. - x86_64 lists only carry `120-real` (sm_121 parts are Arm). - CMake normalizes `121-real` to `121a-real`, and XQA sources are compiled for every selected arch >= 80, so XQA gets sm_121a SASS (no PTX). The CUDA plugin's XQA target (`onnxruntime_providers_cuda_plugin_xqa`) uses the same `MIN_SM 80` filter, with no virtual replacement (that option is only used for the plugin's separate LLM target), so it gets the same SASS on Linux and Windows. ## Testing Built the XQA object library of the CUDA provider (`onnxruntime_providers_cuda_xqa`, CUDA 13.0) with `CMAKE_CUDA_ARCHITECTURES='120-real;121-real'`. The configure log shows `CMAKE_CUDA_ARCHITECTURES: 120a-real;121a-real`. - All 44 XQA objects contain both `sm_120a` and `sm_121a` cubins (`cuobjdump -lelf`), with 158 `kernel_mha` kernels each. Per-kernel register / shared / local / constant usage is identical between sm_120a and sm_121a. - A/B against `main` (without this change): SASS is identical for all 44 objects on sm_120a and sm_121a (600,232 instructions compared; all 88 cubins byte-identical), and also identical on sm_90a for 8 representative translation units (loader, paged, spec-dec, fp8/int8/int4, head sizes 64-512). So existing archs are unaffected and sm_121 behaves exactly as before. - The XQA kernels are present in the final `libonnxruntime_providers_cuda.so` (checked on an sm_90a build: 158 `kernel_mha` kernels). - `lintrunner` and `clang-format` are clean. - Not run on real sm_121 hardware (none available), so a smoke test on a DGX Spark would be welcome. To reproduce the cubin check: ``` python tools/ci_build/build.py --build_dir build/xqa_sm121 --config Release --use_cuda \ --cuda_home <cuda13> --cudnn_home <cudnn> --build_shared_lib --skip_tests --update \ --cmake_extra_defines 'CMAKE_CUDA_ARCHITECTURES=120-real;121-real' ninja -C build/xqa_sm121/Release onnxruntime_providers_cuda_xqa cuobjdump -lelf <obj dir>/xqa_loader_fp16_128.cu.o # lists sm_120a and sm_121a cubins ``` ## Motivation and Context Keeps ORT's XQA arch tables in sync with TensorRT-LLM (NVIDIA/TensorRT-LLM#12170) so sm_121 is handled explicitly rather than by the unknown-arch fallback, and so the `kMAX_SMEM_SIZE` table is right if that check is re-enabled. Independent of the XQA race fixes in #33703. Overlap with #33373: its commit "xqa: treat SM121 like SM120 in the per-architecture constants" makes the same two `|| __CUDA_ARCH__ == 1210` edits with identical spelling. This PR carries the XQA part on its own so it does not have to wait for that larger CMake change. The two PRs merge cleanly in either order (checked with `git merge-tree` against the head of #33373), so whichever lands second just has a no-op hunk. The extra edits here (the sm_121 mention in the comment and the shared-memory-limit hint text) are in different hunks. ## Checklist - [ ] Tests added/updated (N/A: arch-table-only change with identical generated code; verified by cubin comparison above) - [x] No breaking changes - [x] Documentation updated (N/A: only an error-message hint and comments changed) Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Tianlei Wu (tianleiwu) · 2026-11-03T07:49:14Z

Connor Baker (@ConnorBaker), please rebase upon latest main branch to trigger CI.

The SM120 kernels were enabled with literal list tests: "120" IN_LIST CMAKE_CUDA_ARCHITECTURES_ORIG OR "121" IN_LIST ... setup_cuda_architectures() strips -real (and with it any a/f suffix), but keeps the a/f suffix on bare entries, so CMAKE_CUDA_ARCHITECTURES_ORIG holds "121a", "120f", "121f", ... for those spellings. With CMAKE_CUDA_ARCHITECTURES=121a (DGX Spark / GB10, as Nixpkgs passes it) the tests fail: COMPILE_BLACKWELL_SM120_TMA_GROUPED_GEMMS and ORT_ENABLE_BLOCKQUANT_SM120 are not defined, and neither the SM120 TMA grouped GEMMs nor matmul_block_scaled_fp4_sm120.cu are built, even though the 121a-real code is otherwise compiled. The same happens for 120a, 120f and 121f. The EXCLUDE_SM_<arch> loop has the opposite problem for SM121: it matches "^120[af]?$", so a build targeting only SM121 (in any spelling, including 121-real) defines EXCLUDE_SM_120. That empties the SM120 MoE GEMM instantiations (they are wrapped in #ifndef EXCLUDE_SM_120), while for "121"/"121-real" COMPILE_BLACKWELL_SM120_TMA_GROUPED_GEMMS is still defined and the dispatcher references them. Add two helpers to cuda_configuration.cmake: - onnxruntime_parse_cuda_arch() splits an entry of the form <cc>[a|f][-real|-virtual] into compute capability, suffix and whether it is -virtual. - onnxruntime_cuda_archs_have_sass() reports whether any entry generates SASS for a compute capability in a given range. and use them for: - ORT_HAS_SM12X (SM120 or SM121 in any spelling), which replaces the three IN_LIST tests in onnxruntime_providers_cuda.cmake, onnxruntime_providers_cuda_plugin.cmake and onnxruntime_cuda_source_filters.cmake; - EXCLUDE_SM_<arch>, where SM121 now keeps the SM120 kernels. They are compiled for the requested 12.x architectures (121a-real for a 121a build) by onnxruntime_filter_cuda_archs(... MIN_SM 120), and the runtime dispatch already accepts SM121 for SM120 configurations (microsoft#33225); - onnxruntime_filter_cuda_archs(), whose MSVC-only REPLACE_SM120_REAL_WITH_VIRTUAL / EXCLUDE_SM120_REAL options only matched "-real$" and so missed family codes such as "120f", which are kept without -real and also generate SASS. -virtual entries never count: they only embed PTX, and these gates guard native kernels that need arch- or family-specific SASS (sm_120a/sm_121a block-scaled MMA, TMA). This matches the existing behaviour, where -virtual entries are already left out of CMAKE_CUDA_ARCHITECTURES_ORIG. Spellings that already worked ("120", "121", "120a-real", ...) and builds without SM12x are unaffected. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The SM90 TMA OBJECT library (the TMA warp-specialized MoE GEMMs, the fpA_intB SM90 launchers and DeepGEMM) is compiled only for 90a-real, but it was built, and COMPILE_HOPPER_TMA_GEMMS, COMPILE_HOPPER_TMA_GROUPED_GEMMS and USE_DEEP_GEMM were defined, whenever any architecture >= 90 was requested (ORT_HAS_SM90_OR_LATER, or any arch >= 90 in onnxruntime_filter_cuda_archs, which includes -virtual entries). Code built for an arch-specific target such as sm_90a runs only on that exact compute capability: it is neither forward nor backward compatible (CUDA C Programming Guide, "Feature Availability"; Blackwell compatibility guide). So for a build targeting only SM100/SM110/SM120/SM121 (e.g. CMAKE_CUDA_ARCHITECTURES=121a for DGX Spark) these 341 heavy CUTLASS and DeepGEMM translation units (324 TMA MoE GEMMs, 16 fpA_intB launchers, deep_gemm_sm90.cu) produce code that can never run; microsoft#33067 removed the last runtime path that tried to use them on SM120. Add ORT_HAS_SM90, set via onnxruntime_cuda_archs_have_sass() when SM90 itself is requested in any spelling (90, 90a, 90a-real, ...; not 90-virtual), and gate on it: - the onnxruntime_providers_cuda[_plugin]_sm90_tma OBJECT libraries; - COMPILE_HOPPER_TMA_GEMMS / COMPILE_HOPPER_TMA_GROUPED_GEMMS; - include(deep_gemm) and USE_DEEP_GEMM. Nothing else references the skipped code: the TMA MoE dispatch throws when COMPILE_HOPPER_TMA_GROUPED_GEMMS is undefined, DeepGEMM callers are behind USE_DEEP_GEMM, and the fpA_intB SM90 launchers are behind EXCLUDE_SM_90, which is defined exactly when ORT_HAS_SM90 is not (both use the same helper). MSVC builds already run without COMPILE_HOPPER_* and DeepGEMM, and builds without any arch >= 90 (e.g. 86-real) already skip the library. Builds that include SM90 are unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

EXCLUDE_SM_80 was defined unless compute capability 8.0 itself was in CMAKE_CUDA_ARCHITECTURES, so every build without "80" (121a, 120, 90a, 100a, and also 86;89) compiled the SM80 fused MoE GEMM launchers (fused_moe_gemm_sm80_*.generated.cu, "#ifndef EXCLUDE_SM_80") down to nothing. Dispatch still selects them: MoeGemmRunner::dispatchToArch (moe_gemm_template_dispatch.h) uses the Sm80 grouped GEMM on SM80-SM89 and, on SM90 and later, as the fallback for every non-TMA config, and genericMoeGemmKernelLauncher routes FP16/BF16 SiLU/SwiGLU/GeLU/GeGLU epilogues to sm80_generic_fused_moe_gemm_kernelLauncher, which under EXCLUDE_SM_80 is ORT_THROW("Fused MoE SM80 kernels are not available ...") (added in microsoft#28467, 548ab6e). On SM100 (COMPILE_BLACKWELL_TMA_* is never defined) and SM12x (isValidSM120MOESpecialisation only accepts FP4xFP4/FP8xFP4) that fallback is the only path, so unquantized FP16/BF16 MoE with a gated activation always threw on a 12.x- or 10.x-only build. The macro conflated "the kernel is tagged Sm80" with "the kernel runs only on CC 8.0". These kernels are CuTe SM80 atoms (mma.sync SM80_16x8x16_*, cp.async; fused_moe_kernel_traits.cuh), which CUTLASS enables for any __CUDA_ARCH__ >= 800 (CUTE_ARCH_MMA_SM80_ENABLED and CUTE_ARCH_CP_ASYNC_SM80_ENABLED in cute/arch/{mma,copy}_sm80.hpp), and they live in the llm object library, which is compiled for every requested architecture >= 7.5. Compiled for sm_121a they are ordinary sm_121a kernels. There is no separate "SM80 build" that dispatch needs. Define EXCLUDE_SM_<sm> by what dispatch can reach: * 80, 86, 89 (CUTLASS 2.x / CuTe SM80-style, no arch-specific features): keep them if any requested architecture, SASS or PTX, has at least that compute capability (new helper onnxruntime_cuda_archs_have_code). PTX counts because compute_<cc> PTX is JIT-compiled for every GPU >= <cc>. * 90 (sm_90a only), 100, 110, 120: unchanged, native code for that compute capability (or 12.0/12.1 for 120). The gates now read the normalized CMAKE_CUDA_ARCHITECTURES instead of CMAKE_CUDA_ARCHITECTURES_ORIG. The normalized list states what nvcc generates for every entry ("-real", "-virtual", bare "<cc>f") and also holds the -virtual entry, which _ORIG drops. That also keeps ORT_HAS_SM12X/ORT_HAS_SM90 correct if _ORIG ever includes virtual architectures without their suffix, as microsoft#29601 proposes. The gates moved below the normalization for that reason; their logic is otherwise unchanged. TensorRT-LLM, where these kernels come from, also defines EXCLUDE_SM_<sm> only by exact match against the requested architectures (cpp/cmake/modules/cuda_configuration.cmake), but its dispatcher calls the SM80 fused MoE launcher unconditionally, with no EXCLUDE_SM_80 branch. So the throwing branch exists only in onnxruntime. Checked with a cmake -P harness that runs setup_cuda_architectures() on 30 architecture lists and compares the result against an independent model of which GPUs can run the generated code. Before: 28 lists (MSVC off) had "SM80 fused MoE reachable but EXCLUDE_SM_80 throws", including 121a, 120f, 90a, 100a and 86;89. After: none. A 75-only build still excludes SM80, and "75;121-virtual" now keeps it. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

EXCLUDE_SM_100 was cleared only by native code for compute capability 10.0 itself, so a 103a (B300) or 107a build defined it, although both belong to the SM100 family. The CUDA C++ Programming Guide's "Family-Specific Compatibility" table (Compute Capabilities appendix) lists compute_100f as compatible with 10.0, 10.3 and 10.7, compute_110f with 11.0, and compute_120f with 12.0 and 12.1. EXCLUDE_SM_120 already uses the family (ORT_HAS_SM12X, 12.0-12.1). Use the same CC-range helper for 100 (10.0-10.7). 110 stays 11.0 only, which is its whole family. No C++ tests EXCLUDE_SM_100 or EXCLUDE_SM_110 today. This makes the macros mean the same thing as EXCLUDE_SM_120 before anything starts relying on them; upstream's SM100 kernels in TensorRT-LLM are generated for "any(has_arch(sm) for sm in range(100, 110))" (cpp/tensorrt_llm/kernels/cutlass_kernels/ python/generate_kernels.py). cmake -P harness: "103a", "103a-real" and "107a-real" defined EXCLUDE_SM_100 before and do not after. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

setup_cuda_architectures() dropped the suffix of every "-real" entry ("^(([1-9])([0-9])+)[af]?-real$" keeps only the number) and then re-added "a" for the numbers in ARCHITECTURES_WITH_ACCEL. That has two effects: * "120f-real" became "120a-real", and "100f-real" became "100a-real". Per the CUDA C++ Programming Guide (Compute Capabilities appendix), compute_120a/sm_120a code runs only on compute capability 12.0, while the family target compute_120f covers 12.0 and 12.1, and compute_100f covers 10.0, 10.3 and 10.7. So a build asking for portable 12.x SASS got code that does not run on 12.1 (GB10/DGX Spark), and a 10.x one got code that does not run on 10.3 (B300). Bare "120f" was already kept verbatim (family codes were added in microsoft#27278, 6625856, which extended the -real regex to [af]? but still discarded the match). * "107a-real" (any "a" entry whose number is not in ARCHITECTURES_WITH_ACCEL) silently became "107-real". Keep the suffix: "<cc>f-real" is kept verbatim (the "-real" tells it apart from "<cc>f", which also embeds PTX), and "<cc>a-real" keeps its "a". The normalization loop passes "<cc>f-real" through like "<cc>f". _ORIG now holds "120f-real" for such an entry. Its consumers (HAS_SM*_OR_LATER, which matches the leading number, and the onnxruntime_parse_cuda_arch-based gates) handle that. Since "121" and "121a-real" now normalize to the same entry, the normalized list is de-duplicated. Coordination with the open microsoft#29601 ("Follow standard CMake semantics for CUDA architecture suffixes"): that PR splits the parsing into real and virtual lists, keeps every -virtual entry, and gives suffix-less entries PTX as well. It uses the same "-real" regex as the second branch here, "^(([1-9])([0-9])+[af]?)-real$" (suffix captured), and adds the same list(REMOVE_DUPLICATES CMAKE_CUDA_ARCHITECTURES_NORMALIZED), so the two changes overlap textually without contradicting each other. One difference: microsoft#29601 still emits family entries without a kind suffix ("120f"), so "120f-real" gains PTX and "120f-virtual" gains SASS. On top of microsoft#29601 the equivalent fix is to emit "<cc>f-real" for real-list and "<cc>f-virtual" for virtual-list family entries. cmake -P harness (independent model of which GPUs each entry generates SASS for, plus "an explicit a/f suffix survives"): before, 7 violations on 5 of 30 lists ("120f-real", "121f-real", "100f-real", "107a-real", "120f-real;121a-real"). After, none. Only lists that contain a "-real" entry normalize differently. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

ORT_HAS_SM{80,90,100,120}_OR_LATER (and the HAS_SM*_OR_LATER definitions) were computed from CMAKE_CUDA_ARCHITECTURES_ORIG, which leaves out the -virtual entry that setup_cuda_architectures() keeps in CMAKE_CUDA_ARCHITECTURES_LAST_VIRTUAL and appends to the normalized list. For "75;121-virtual" or "90-virtual" the CUDA sources are still compiled for compute_121/compute_90 PTX, but HAS_SM80_OR_LATER was not defined. xqa_impl_gen.cuh and xqa_paged_impl_gen.cuh rely on it in the host pass ("If any SM80+ arch is enabled, the host stub must be emitted"), so their Launch() returned "XQA is only supported on Ampere (SM80) or newer GPUs" on the 12.1 GPU that the embedded PTX targets. Iterate over _ORIG plus _LAST_VIRTUAL. The "^([0-9]+)" match already handles a "-virtual" entry. microsoft#29601 would get the same result by making _ORIG the union of real and virtual architectures. On top of it this line is a no-op (it removes _LAST_VIRTUAL). This was the only other gate that a sweep of cmake/ and the CUDA sources found to mishandle an architecture spelling. The literal __CUDA_ARCH__ comparisons outside XQA (fpA_intB_gemm.h "== 890", XQA's "== 900" occupancy) and the runtime "sm == 90" checks select sm_90a-only or SM89-only code paths and are correct as written. cmake -P harness: "121-virtual", "90-virtual" and "75;121-virtual" no longer violate "an architecture >= N is built but HAS_SMN_OR_LATER is not defined" (10 violations before, 0 after); the rest of the harness output is identical. In those builds ORT_HAS_SM90_OR_LATER now also adds -DCUTLASS_ENABLE_GDC_FOR_SM90=1 and moves the SM90 TMA sources out of the main list. The SM90 library itself still requires ORT_HAS_SM90 (sm_90a SASS). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Tianlei Wu (tianleiwu)

The shared parsing/gating direction looks good. I exercised the checked-out CMake functions: nine normalization/native-gating cases passed, but the Windows plugin probes expose a missing native fallback for family spellings, detailed inline.

This is configuration-level validation, not a full CUDA build/runtime test. The only current-head reported check is CLA. No approval submitted.

…ibrary Under MSVC, onnxruntime_filter_cuda_archs(... REPLACE_SM120_REAL_WITH_VIRTUAL) turns every native SM120+ entry of the plugin's LLM OBJECT library into <cc>-virtual, and the plugin then adds non-"a" native SASS (-gencode=arch=compute_120,code=sm_120) so the ordinary LLM kernels do not depend on the driver JIT-compiling the toolkit's PTX. That fallback only checked for the literal entries "120" and "120a" in CMAKE_CUDA_ARCHITECTURES_ORIG. Since the filter now also converts family codes, "120f" and "120f-real" became PTX-only, and every SM121 spelling ("121", "121a", "121f", "121a-real", ...) never had the fallback. Have the filter report the compute capabilities it converted (REPLACED_SM120_CCS) and add -gencode=arch=compute_<cc>,code=sm_<cc> for each of them, so the fallback always matches the conversion. sm_120 SASS runs on 12.0 and 12.1, which covers what compute_120f covers. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Tianlei Wu (tianleiwu)

Re-reviewed 3afda6615b01172a502c798d1ff717893601b896. The previous Windows plugin native-SASS concern is addressed; no additional actionable code findings.

Fresh local configuration validation executed the real architecture normalization/filter functions and the unchanged MSVC plugin branch: 22 architecture-list cases and 136 assertions passed. Coverage includes SM120/SM121 native spellings, a/f and -real variants, virtual-only entries, mixed lists, Hopper gating, and duplicate fallback removal.

Validation limit: these are CMake configuration checks, not CUDA compilation or a Windows build. The current head reports only license/cla; please obtain Linux and Windows CUDA/plugin CI before merging. No approval submitted.

Tianlei Wu (tianleiwu) · 2026-11-10T23:33:01Z

Connor Baker (@ConnorBaker), Thanks for the contribution.
This PR falls behind main so it cannot pass CI. I've cherry-picked your commits to #34089 and merge to main.

Tianlei Wu (tianleiwu) requested a balanced review from Copilot October 15, 2026 02:49

Copilot started reviewing on behalf of Tianlei Wu (tianleiwu) October 15, 2026 02:50

View session

Copilot AI reviewed Oct 15, 2026

View reviewed changes

Connor Baker (ConnorBaker) marked this pull request as ready for review October 15, 2026 03:04

Gaétan Lepage (GaetanLepage) mentioned this pull request Oct 28, 2026

onnxruntime: 1.27.1 -> 1.31.1 NixOS/nixpkgs#547208

Open

1 of 13 tasks

Tianlei Wu (tianleiwu) mentioned this pull request Oct 30, 2026

[CUDA] Add sm_121 to XQA arch-specific branches #33762

Merged

2 of 3 tasks

Connor Baker (ConnorBaker) and others added 6 commits November 3, 2026 14:56

Connor Baker (ConnorBaker) force-pushed the main-sm12x-fixes branch from d44dea2 to 7082d93 Compare November 3, 2026 22:58

Tianlei Wu (tianleiwu) requested changes Nov 5, 2026

View reviewed changes

Comment thread cmake/onnxruntime_cuda_source_filters.cmake Resolved

Connor Baker (ConnorBaker) requested a review from Tianlei Wu (tianleiwu) November 5, 2026 16:20

Tianlei Wu (tianleiwu) reviewed Nov 9, 2026

View reviewed changes

Tianlei Wu (tianleiwu) approved these changes Nov 10, 2026

View reviewed changes

Tianlei Wu (tianleiwu) mentioned this pull request Nov 10, 2026

[CUDA] Build the aarch64 wheel with family targets and fix CUDA arch gating #34089

Merged

3 tasks done

Tianlei Wu (tianleiwu) closed this Nov 10, 2026

Navigation Menu

Repository navigation

[CUDA] Fix CUDA architecture gating for SM12x, family/arch-specific targets and -virtual entries - #33373

[CUDA] Fix CUDA architecture gating for SM12x, family/arch-specific targets and -virtual entries#33373
Connor Baker (ConnorBaker) wants to merge 7 commits into
microsoft:mainfrom
ConnorBaker:main-sm12x-fixes

Connor Baker (ConnorBaker) commented Oct 14, 2026 •

edited

Loading

Uh oh!

azure-pipelines Bot commented Oct 14, 2026

Uh oh!

Connor Baker (ConnorBaker) commented Oct 14, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

Connor Baker (ConnorBaker) commented Oct 24, 2026

Uh oh!

Connor Baker (ConnorBaker) commented Oct 30, 2026

Uh oh!

Tianlei Wu (tianleiwu) commented Nov 3, 2026

Uh oh!

Tianlei Wu (tianleiwu) left a comment

Uh oh!

Uh oh!

Tianlei Wu (tianleiwu) left a comment

Uh oh!

Tianlei Wu (tianleiwu) commented Nov 10, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

Connor Baker (ConnorBaker) commented Oct 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Motivation and Context

Testing

Uh oh!

azure-pipelines Bot commented Oct 14, 2026

Uh oh!

Connor Baker (ConnorBaker) commented Oct 14, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🔵 Needs a closer look

Uh oh!

Connor Baker (ConnorBaker) commented Oct 24, 2026

Uh oh!

Connor Baker (ConnorBaker) commented Oct 30, 2026

Uh oh!

Tianlei Wu (tianleiwu) commented Nov 3, 2026

Uh oh!

Tianlei Wu (tianleiwu) left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Tianlei Wu (tianleiwu) left a comment

Choose a reason for hiding this comment

Uh oh!

Tianlei Wu (tianleiwu) commented Nov 10, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Connor Baker (ConnorBaker) commented Oct 14, 2026 •

edited

Loading