Skip to content

Follow standard CMake semantics for CUDA architecture suffixes #29601

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
Sammy-Dabbas wants to merge 2 commits into
base: main
Choose a base branch
from
+35 −18
Open

Follow standard CMake semantics for CUDA architecture suffixes #29601

Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
3781c0b
Follow standard CMake semantics for CUDA architecture suffixes
Sammy-Dabbas Jul 7, 2026
cf624f2
Keep PTX forward compat and preserve family codes in CUDA arch setup
Sammy-Dabbas Jul 10, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
47 changes: 30 additions & 17 deletions cmake/external/cuda_configuration.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down Expand Up @@ -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)
Expand Down Expand Up @@ -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}")
Expand All @@ -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})

Expand Down
Loading