From 3781c0b547f5a9246a88e44b5b328ebfe9247964 Mon Sep 17 00:00:00 2001 From: Sammy Dabbas Date: Tue, 7 Jul 2026 15:47:55 -0400 Subject: [PATCH 1/2] Follow standard CMake semantics for CUDA architecture suffixes 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. --- cmake/external/cuda_configuration.cmake | 47 ++++++++++++++++--------- 1 file changed, 30 insertions(+), 17 deletions(-) diff --git a/cmake/external/cuda_configuration.cmake b/cmake/external/cuda_configuration.cmake index 9d7bfb0e5ad61..35b36843d24e0 100644 --- a/cmake/external/cuda_configuration.cmake +++ b/cmake/external/cuda_configuration.cmake @@ -79,10 +79,11 @@ macro(setup_cuda_architectures) # Special values: # (1) `native` is resolved to HIGHEST available architecture. Fallback to `all` if detection failed. # (2) `all` / `all-major` / unset is resolved to a default set of architectures we optimized and compiler supports. - # Numerical architectures: - # * For `-virtual` architectures, the last one is kept as it is, and the others are ignored. - # * `-real` suffix is automatically added for other cases. - # * Always use accelerated (`-a` suffix) target for supported real architectures. + # Numerical architectures follow standard CMake semantics: + # * A suffix-less architecture (e.g. `86`) enables BOTH `86-real` (SASS) and `86-virtual` (PTX). + # * A `-real` architecture is kept real only, and a `-virtual` architecture is kept virtual only. + # Every `-virtual` entry embeds its own PTX (the whole list is preserved, not just the last one). + # * Always use accelerated (`-a` suffix) target for supported real and virtual architectures (SM >= 90). # cmake-format: on # Allow override via CUDAARCHS environment variable (standard CMake variable) @@ -137,28 +138,33 @@ macro(setup_cuda_architectures) endif() endif() - unset(CMAKE_CUDA_ARCHITECTURES_CLEAN) - unset(CMAKE_CUDA_ARCHITECTURES_LAST_VIRTUAL) + unset(CMAKE_CUDA_ARCHITECTURES_REAL_CLEAN) + unset(CMAKE_CUDA_ARCHITECTURES_VIRTUAL_CLEAN) foreach(CUDA_ARCH IN LISTS CMAKE_CUDA_ARCHITECTURES) if(CUDA_ARCH STREQUAL "") continue() endif() - if(CUDA_ARCH MATCHES "^([1-9])([0-9])+[af]?-virtual$") - 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$") - list(APPEND CMAKE_CUDA_ARCHITECTURES_CLEAN ${CMAKE_MATCH_1}) + list(APPEND CMAKE_CUDA_ARCHITECTURES_REAL_CLEAN ${CMAKE_MATCH_1}) elseif(CUDA_ARCH MATCHES "^(([1-9])([0-9])+)([af]?)$") - list(APPEND CMAKE_CUDA_ARCHITECTURES_CLEAN ${CMAKE_MATCH_1}${CMAKE_MATCH_4}) + # Suffix-less architecture: standard CMake enables both real (SASS) and virtual (PTX). + list(APPEND CMAKE_CUDA_ARCHITECTURES_REAL_CLEAN ${CMAKE_MATCH_1}${CMAKE_MATCH_4}) + list(APPEND CMAKE_CUDA_ARCHITECTURES_VIRTUAL_CLEAN ${CMAKE_MATCH_1}${CMAKE_MATCH_4}) else() message(FATAL_ERROR "Unrecognized CUDA architecture: ${CUDA_ARCH}") endif() endforeach() - list(REMOVE_DUPLICATES CMAKE_CUDA_ARCHITECTURES_CLEAN) - set(CMAKE_CUDA_ARCHITECTURES ${CMAKE_CUDA_ARCHITECTURES_CLEAN}) + list(REMOVE_DUPLICATES CMAKE_CUDA_ARCHITECTURES_REAL_CLEAN) + list(REMOVE_DUPLICATES CMAKE_CUDA_ARCHITECTURES_VIRTUAL_CLEAN) + set(CMAKE_CUDA_ARCHITECTURES ${CMAKE_CUDA_ARCHITECTURES_REAL_CLEAN}) # CMAKE_CUDA_ARCHITECTURES_ORIG contains all architectures enabled, without automatically added -real or -a suffix. - set(CMAKE_CUDA_ARCHITECTURES_ORIG "${CMAKE_CUDA_ARCHITECTURES}") + set(CMAKE_CUDA_ARCHITECTURES_ORIG "${CMAKE_CUDA_ARCHITECTURES_REAL_CLEAN}") + list(APPEND CMAKE_CUDA_ARCHITECTURES_ORIG ${CMAKE_CUDA_ARCHITECTURES_VIRTUAL_CLEAN}) + list(REMOVE_DUPLICATES CMAKE_CUDA_ARCHITECTURES_ORIG) message(STATUS "GPU architectures: ${CMAKE_CUDA_ARCHITECTURES_ORIG}") unset(ORT_HAS_SM80_OR_LATER) @@ -205,7 +211,7 @@ macro(setup_cuda_architectures) # Enable accelerated features (like WGMMA, TMA and setmaxnreg) for SM >= 90. set(ARCHITECTURES_WITH_ACCEL "90" "100" "101" "110" "120") unset(CMAKE_CUDA_ARCHITECTURES_NORMALIZED) - foreach(CUDA_ARCH IN LISTS CMAKE_CUDA_ARCHITECTURES) + foreach(CUDA_ARCH IN LISTS CMAKE_CUDA_ARCHITECTURES_REAL_CLEAN) if(CUDA_ARCH MATCHES "^([0-9]+)f$") # Family code, no -real suffix list(APPEND CMAKE_CUDA_ARCHITECTURES_NORMALIZED "${CUDA_ARCH}") @@ -216,9 +222,16 @@ macro(setup_cuda_architectures) endif() endforeach() - if(DEFINED CMAKE_CUDA_ARCHITECTURES_LAST_VIRTUAL) - list(APPEND CMAKE_CUDA_ARCHITECTURES_NORMALIZED "${CMAKE_CUDA_ARCHITECTURES_LAST_VIRTUAL}") - endif() + foreach(CUDA_ARCH IN LISTS CMAKE_CUDA_ARCHITECTURES_VIRTUAL_CLEAN) + if(CUDA_ARCH MATCHES "^([0-9]+)f$") + # Family code, no -virtual suffix + list(APPEND CMAKE_CUDA_ARCHITECTURES_NORMALIZED "${CUDA_ARCH}") + elseif("${CUDA_ARCH}" IN_LIST ARCHITECTURES_WITH_ACCEL) + list(APPEND CMAKE_CUDA_ARCHITECTURES_NORMALIZED "${CUDA_ARCH}a-virtual") + else() + list(APPEND CMAKE_CUDA_ARCHITECTURES_NORMALIZED "${CUDA_ARCH}-virtual") + endif() + endforeach() set(CMAKE_CUDA_ARCHITECTURES ${CMAKE_CUDA_ARCHITECTURES_NORMALIZED}) From cf624f22d4b7b28a3dcbd367ced5c5cb6eef544a Mon Sep 17 00:00:00 2001 From: Sammy Dabbas Date: Fri, 10 Jul 2026 14:03:16 -0400 Subject: [PATCH 2/2] Keep PTX forward compat and preserve family codes in CUDA arch setup 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. --- cmake/external/cuda_configuration.cmake | 12 ++++++++---- 1 file changed, 8 insertions(+), 4 deletions(-) diff --git a/cmake/external/cuda_configuration.cmake b/cmake/external/cuda_configuration.cmake index 35b36843d24e0..84a7522324128 100644 --- a/cmake/external/cuda_configuration.cmake +++ b/cmake/external/cuda_configuration.cmake @@ -83,7 +83,9 @@ macro(setup_cuda_architectures) # * A suffix-less architecture (e.g. `86`) enables BOTH `86-real` (SASS) and `86-virtual` (PTX). # * A `-real` architecture is kept real only, and a `-virtual` architecture is kept virtual only. # Every `-virtual` entry embeds its own PTX (the whole list is preserved, not just the last one). - # * Always use accelerated (`-a` suffix) target for supported real and virtual architectures (SM >= 90). + # * Supported real (SASS) architectures (SM >= 90) use the accelerated (`-a` suffix) target for + # features like WGMMA and TMA. Virtual (PTX) architectures keep the user-specified variant so the + # embedded PTX stays forward compatible on newer GPUs. # cmake-format: on # Allow override via CUDAARCHS environment variable (standard CMake variable) @@ -147,7 +149,7 @@ macro(setup_cuda_architectures) 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$") + elseif(CUDA_ARCH MATCHES "^(([1-9])([0-9])+[af]?)-real$") list(APPEND CMAKE_CUDA_ARCHITECTURES_REAL_CLEAN ${CMAKE_MATCH_1}) elseif(CUDA_ARCH MATCHES "^(([1-9])([0-9])+)([af]?)$") # Suffix-less architecture: standard CMake enables both real (SASS) and virtual (PTX). @@ -226,13 +228,15 @@ macro(setup_cuda_architectures) if(CUDA_ARCH MATCHES "^([0-9]+)f$") # Family code, no -virtual suffix list(APPEND CMAKE_CUDA_ARCHITECTURES_NORMALIZED "${CUDA_ARCH}") - elseif("${CUDA_ARCH}" IN_LIST ARCHITECTURES_WITH_ACCEL) - list(APPEND CMAKE_CUDA_ARCHITECTURES_NORMALIZED "${CUDA_ARCH}a-virtual") else() + # Keep the user-specified virtual arch and do not force the accelerated `a` suffix: + # `Na-virtual` PTX is not forward compatible, while `N-virtual` PTX JITs onto newer GPUs. list(APPEND CMAKE_CUDA_ARCHITECTURES_NORMALIZED "${CUDA_ARCH}-virtual") endif() endforeach() + list(REMOVE_DUPLICATES CMAKE_CUDA_ARCHITECTURES_NORMALIZED) + set(CMAKE_CUDA_ARCHITECTURES ${CMAKE_CUDA_ARCHITECTURES_NORMALIZED}) message(STATUS "CMAKE_CUDA_ARCHITECTURES: ${CMAKE_CUDA_ARCHITECTURES}")