Follow standard CMake semantics for CUDA architecture suffixes by Sammy-Dabbas · Pull Request #29601 · microsoft/onnxruntime

Sammy Dabbas (Sammy-Dabbas) · 2026-07-07T19:49:08Z

Description

setup_cuda_architectures in cmake/external/cuda_configuration.cmake was non-standard in two ways (#29598): only the last -virtual entry was kept, so 86-virtual;120-virtual dropped PTX for 8.6; and a suffix-less arch like 86 became 86-real only, dropping its virtual arch.

This tracks real and virtual architectures in separate lists. A suffix-less arch enables both (standard CMake: 86 gives 86-real and 86-virtual), -real/-virtual entries are kept as given, and all -virtual entries are preserved. The accelerated a and family-code f handling apply to both, and CMAKE_CUDA_ARCHITECTURES_ORIG is the deduplicated union.

Motivation and Context

Fixes #29598. Thanks Corentin Schreiber (@cschreib-ibex) for the reference patch; this builds on it, fixing that its -virtual regex only matched even digit counts (so 100-virtual/120-virtual fatal-errored) and predated the f family branch. Verified the list logic with a cmake -P harness; nvcc compilation should be confirmed by CI.

setup_cuda_architectures kept only the last -virtual entry and dropped PTX for suffix-less architectures, so 86-virtual;120-virtual embedded PTX for 12.0 only and bare 86 became 86-real without a virtual arch. Track real and virtual architectures in separate lists: a suffix-less arch enables both real and virtual (standard CMake), every -virtual entry is preserved, and the accelerated (a) and family-code (f) handling apply to both. CMAKE_CUDA_ARCHITECTURES_ORIG becomes the deduplicated union so PTX-enabled archs also get their SM-specific kernels.

Copilot

Pull request overview

This PR fixes ONNX Runtime’s CUDA architecture normalization in setup_cuda_architectures to follow standard CMake semantics, ensuring both SASS (-real) and PTX (-virtual) handling is correct and that multiple -virtual entries are preserved (fixing #29598).

Changes:

Split normalization into separate “real” and “virtual” cleaned lists so multiple -virtual entries are retained.
Make suffix-less architectures (e.g., 86) enable both 86-real and 86-virtual as standard CMake does.
Normalize accelerated (a) and family (f) handling across both real and virtual variants and expose the deduplicated union via CMAKE_CUDA_ARCHITECTURES_ORIG.

-      set(CMAKE_CUDA_ARCHITECTURES_LAST_VIRTUAL ${CUDA_ARCH})
+    if(CUDA_ARCH MATCHES "^(([1-9])([0-9])+[af]?)-virtual$")
+      list(APPEND CMAKE_CUDA_ARCHITECTURES_VIRTUAL_CLEAN ${CMAKE_MATCH_1})
    elseif(CUDA_ARCH MATCHES "^(([1-9])([0-9])+)[af]?-real$")


+    endif()
+  endforeach()

  set(CMAKE_CUDA_ARCHITECTURES ${CMAKE_CUDA_ARCHITECTURES_NORMALIZED})


Corentin Schreiber (cschreib-ibex) · 2026-07-10T08:23:12Z

Thanks for writing the PR.

FYI I also discovered today that the logic of forcing all "accelerated" archs (90, 100, 101, 110, 120) to the *a-virtual/*a-real architecture family, which is also non-standard, can be problematic because these architectures are not forward compatible. See CUDA docs:

Application binaries that include PTX version of kernels with architecture conditional features using sm_100a or compute_100a in order to take full advantage of Blackwell GPU architecture, are not forward or backward compatible. For example, PTX compiled for compute_90a (Hopper) are not supported on the Blackwell architecture.

So for example, if you select 120-virtual in CMake, this will be turned into 120a-virtual for onnxruntime, and the embedded PTX will not be forward compatible with newer GPUs. If this is the only virtual architecture in the list, effectively this means the resulting binaries will not work on newer GPUs.

At present this can be circumvented by adding at least one "non-accelerated" virtual architecture (e.g., 89-virtual) to the list.

The setup_cuda_architectures macro forced the accelerated `a` suffix onto virtual (PTX) architectures. Architecture-conditional PTX built as compute_XXa is not forward compatible, so a selected 120-virtual became 120a-virtual and its embedded PTX would not JIT onto newer GPUs. Drop the accel branch in the virtual normalization loop so virtual archs keep the user-specified variant. Real (SASS) architectures still use the `a` suffix because the ORT kernels need the accelerated features such as WGMMA and TMA, and SASS never JITs forward regardless, so this costs no forward compatibility. Fix the -real regex so the optional family or accel suffix sits inside the capture group. The previous pattern left [af] outside the group, so 100f-real captured only 100 and was rewritten to 100a-real, silently dropping the family code. It now mirrors the -virtual branch and yields 100f. Deduplicate the normalized architecture list. A suffix-less family code such as 100f is pushed onto both the real and virtual clean lists, and both loops emit a bare 100f, which produced a duplicate in the final output.

Sammy Dabbas (Sammy-Dabbas) · 2026-07-10T18:12:22Z

Thanks for catching this. The a-forcing on virtual architectures was introduced by this PR through the new virtual normalization loop, I have removed it, so a plain 120-virtual stays a 120-virtual and its compute_120 PTX remains forward compatible via JIT on newer GPUs. The a-forcing on real architectures is kept because the ORT kernels need the accelerated SASS features such as WGMMA and TMA, and SASS is architecture specific anyway, so this does not cost any forward compatibility that PTX would otherwise provide. The fix only stops auto adding the a suffix. An explicit Xa-virtual is preserved unchanged, so the accelerated PTX variant is still available when it's requested directly.

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>

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>

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>

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>

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>

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>

…gating (#34089) > Stacked on #34088 (base branch `tlwu/20261110/flash_attention_reuse_sm80`). It will be retargeted to `main` after #34088 merges. ## Description The Linux aarch64 CUDA 13 wheel is 389 MB because every architecture in its list costs a whole cubin set (34-80 MB). This PR builds it with one family-specific cubin per Blackwell family (`100f`, `120f`) instead of the `100a` + `103a` and `120a` + `121a` pairs, which saves ~82 MB (to ~310 MB) without dropping any GPU. That required fixing how the CUDA architecture list is normalized and how the SM-specific kernels are gated: `100f-real` was rewritten to `100a-real`, and several gates recognized only literal spellings such as `"120"`, which compiled out or broke kernels for other valid spellings. ## Summary of Changes ### Linux aarch64 wheel | File | Change | |------|--------| | `py-cuda13-packaging-pipeline.yml`, `plugin-cuda13-pipeline.yml`, `stages/plugin-cuda-packaging-stage.yml`, `linux/build_linux_python_package.sh` | `87-real;89-real;90-real;100-real;103-real;110-real;120-real;121-real` → `87-real;89-real;90-real;100f-real;110-real;120f-real` | ### CUDA architecture handling (`cuda_configuration.cmake`, `onnxruntime_cuda_source_filters.cmake`, provider and plugin cmake) | Commit | Fix | |---|---| | match CUDA architectures with a/f/-real suffixes | Add `onnxruntime_parse_cuda_arch()` and `onnxruntime_cuda_archs_have_sass()`. `ORT_HAS_SM12X` replaces the literal `"120"`/`"121" IN_LIST` tests, so `121a`, `120f` and `121f` builds get the SM120 TMA grouped GEMMs and block-scaled FP4 kernels, and SM121-only builds no longer define `EXCLUDE_SM_120`. | | only build SM90 TMA kernels when SM90 is targeted | `ORT_HAS_SM90` gates the SM90 TMA library, `COMPILE_HOPPER_TMA_*` and DeepGEMM. sm_90a code cannot run on SM100/SM12x, so builds without SM90 no longer compile those 341 units. | | keep SM80-generation kernels whenever an SM80+ GPU is targeted | `EXCLUDE_SM_80/86/89` only when no requested architecture (SASS or PTX) reaches that compute capability. Builds without `80` (the aarch64 wheel, `121a`, `120f`, `90a`, ...) threw `Fused MoE SM80 kernels are not available` for unquantized FP16/BF16 MoE. | | clear EXCLUDE_SM_100 for every SM100-family architecture | `103a` and `107a` builds excluded the SM100-family kernels. | | keep explicit a/f suffixes of -real CUDA architectures | `100f-real`/`120f-real` became `100a-real`/`120a-real`, which do not run on 10.3/12.1; `107a-real` lost its `a`. | | count the -virtual architecture in HAS_SM*_OR_LATER | `75;121-virtual` and `90-virtual` builds missed `HAS_SM80_OR_LATER`, so the XQA host stubs reported "XQA is only supported on Ampere (SM80) or newer GPUs". | | add native SASS for every SM12x entry / for one SM12x compute capability per major in the MSVC plugin LLM library | The MSVC plugin's native non-"a" SASS fallback covers every SM12x spelling and adds one SASS per major version (`120;121` → `sm_120`, as before). | | report the CMake version that family-specific CUDA architectures need | CMake accepts the `f` suffix in `CMAKE_CUDA_ARCHITECTURES` only since 3.31.8 (4.0.2 in the 4.0 series). Fail early with a clear message instead of a generic error in `enable_language(CUDA)`. | ### Tests and docs - `tools/ci_build/test_cuda_package_architectures.py`: aarch64 normalization, gates and definitions; the four aarch64 lists stay in sync; on the targets nvcc builds, the list covers Orin, Ada, GH200, B200/GB200, B300/GB300, Thor, RTX 50/RTX PRO and DGX Spark with no redundant entry; the CMake version check; `onnxruntime_cuda_lowest_cc_per_major()`. - `tools/ci_build/test_cuda_plugin_architectures.py`: checks the exact set of native SASS targets, including the Windows arm64 list. - `docs/Validate_CUDA_Package.md`: aarch64 coverage table and the CMake requirement. ## Testing - `python tools/ci_build/test_cuda_package_architectures.py` and `python tools/ci_build/test_cuda_plugin_architectures.py`, which the Tooling Tests workflow runs with CMake 3.31.8. Every commit of the stack passes them. - Before/after on the production lists: Linux x64 (CUDA 13 and 12.8) and Windows x64 normalize and gate as before (Windows x64 only skips the dead SM90 TMA library). Windows arm64 (MSVC) no longer defines `EXCLUDE_SM_80/86/89` and still adds only `sm_120` native LLM SASS. The aarch64 list normalizes to `87-real;89-real;90a-real;100f-real;110a-real;120f-real` with no `EXCLUDE_SM_*`. - `a` vs `f` targets, nvcc 13.0.48 and CUTLASS 4.7.0: the SM120 FP4 TMA MoE launcher and `matmul_block_scaled_fp4_sm120.cu` (`sm_120a` vs `sm_120f`), and `moe_kernels.cu` (`sm_100a` vs `sm_100f`), contain the same kernels with identical counts of every feature instruction (`OMMA.SF` block-scaled MMA, `USETMAXREG`, `UTMALDG/UTMASTG/UTMACCTL`, `SYNCS`, `F2FP...E2M1`) and no arch-conditional abort path. CUTLASS 4.7 pairs every `*A_ENABLED` feature gate with an `*F_ENABLED` one, and ORT's own device code has no guard that only an `a` target satisfies. - The re-enabled SM80 fused MoE kernels add ~0.6 MB (compressed) per architecture. - **To do before merging:** run the Linux aarch64 wheel and CUDA plugin packaging pipelines (their docker image needs CMake 3.31.8 or later), then the CUDA tests on GH200, B200 or B300, DGX Spark and an RTX 50 GPU. ## Motivation and Context - Family-specific targets need CUDA 12.9 or later (the wheel uses CUDA 13) and CMake 3.31.8 or later (4.0.2 or later in the 4.0 series). - The architecture parsing overlaps with the open #29601 ("Follow standard CMake semantics for CUDA architecture suffixes"); the commit messages describe how the two fit together. ## Checklist - [x] Tests added/updated - [x] No breaking changes (or documented in description) - [x] Documentation updated (if applicable) --------- Co-authored-by: Connor Baker <ConnorBaker01@gmail.com> Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Tianlei Wu (tianleiwu) requested a review from Copilot July 7, 2026 21:39

Copilot started reviewing on behalf of Tianlei Wu (tianleiwu) July 7, 2026 21:39

View session

Copilot AI reviewed Jul 7, 2026

View reviewed changes

Connor Baker (ConnorBaker) mentioned this pull request Oct 14, 2026

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

Closed

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

Navigation Menu

Repository navigation

Follow standard CMake semantics for CUDA architecture suffixes - #29601

Follow standard CMake semantics for CUDA architecture suffixes#29601
Sammy Dabbas (Sammy-Dabbas) wants to merge 2 commits into
microsoft:mainfrom
Sammy-Dabbas:fix/cuda-arch-suffix-semantics

Sammy Dabbas (Sammy-Dabbas) commented Jul 7, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

Corentin Schreiber (cschreib-ibex) commented Jul 10, 2026

Uh oh!

Sammy Dabbas (Sammy-Dabbas) commented Jul 10, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

Sammy Dabbas (Sammy-Dabbas) commented Jul 7, 2026

Description

Motivation and Context

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

Corentin Schreiber (cschreib-ibex) commented Jul 10, 2026

Uh oh!

Sammy Dabbas (Sammy-Dabbas) commented Jul 10, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants