Skip to content

Add support for CUDA architecture family codes #27278

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

Merged
tianleiwu merged 1 commit into from
Feb 8, 2026
+15 −7
Merged

Add support for CUDA architecture family codes #27278

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
22 changes: 15 additions & 7 deletions cmake/external/cuda_configuration.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -85,6 +85,11 @@ macro(setup_cuda_architectures)
# * Always use accelerated (`-a` suffix) target for supported real architectures.
# cmake-format: on

# Allow override via CUDAARCHS environment variable (standard CMake variable)
if(NOT CMAKE_CUDA_ARCHITECTURES AND DEFINED ENV{CUDAARCHS})
set(CMAKE_CUDA_ARCHITECTURES "$ENV{CUDAARCHS}")
endif()

if(CMAKE_CUDA_ARCHITECTURES STREQUAL "native")
# Detect highest available compute capability
set(OUTPUTFILE ${PROJECT_BINARY_DIR}/detect_cuda_arch)
Expand Down Expand Up @@ -142,12 +147,12 @@ macro(setup_cuda_architectures)
continue()
endif()

if(CUDA_ARCH MATCHES "^([1-9])([0-9])+a?-virtual$")
if(CUDA_ARCH MATCHES "^([1-9])([0-9])+[af]?-virtual$")
set(CMAKE_CUDA_ARCHITECTURES_LAST_VIRTUAL ${CUDA_ARCH})
elseif(CUDA_ARCH MATCHES "^(([1-9])([0-9])+)a?-real$")
list(APPEND CMAKE_CUDA_ARCHITECTURES_CLEAN ${CMAKE_MATCH_1})
elseif(CUDA_ARCH MATCHES "^(([1-9])([0-9])+)a?$")
elseif(CUDA_ARCH MATCHES "^(([1-9])([0-9])+)[af]?-real$")
list(APPEND CMAKE_CUDA_ARCHITECTURES_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})
else()
message(FATAL_ERROR "Unrecognized CUDA architecture: ${CUDA_ARCH}")
endif()
Expand All @@ -159,7 +164,7 @@ macro(setup_cuda_architectures)
set(CMAKE_CUDA_ARCHITECTURES_ORIG "${CMAKE_CUDA_ARCHITECTURES}")
message(STATUS "GPU architectures: ${CMAKE_CUDA_ARCHITECTURES_ORIG}")

set(ARCHITECTURES_WITH_KERNELS "80" "86" "89" "90" "100" "120")
set(ARCHITECTURES_WITH_KERNELS "80" "86" "89" "90" "100" "110" "120")
foreach(CUDA_ARCH IN LISTS ARCHITECTURES_WITH_KERNELS)
if(NOT "${CUDA_ARCH}" IN_LIST CMAKE_CUDA_ARCHITECTURES_ORIG)
add_definitions("-DEXCLUDE_SM_${CUDA_ARCH}")
Expand All @@ -168,10 +173,13 @@ macro(setup_cuda_architectures)
endforeach()

# Enable accelerated features (like WGMMA, TMA and setmaxnreg) for SM >= 90.
set(ARCHITECTURES_WITH_ACCEL "90" "100" "101" "120")
set(ARCHITECTURES_WITH_ACCEL "90" "100" "101" "110" "120")
unset(CMAKE_CUDA_ARCHITECTURES_NORMALIZED)
foreach(CUDA_ARCH IN LISTS CMAKE_CUDA_ARCHITECTURES)
if("${CUDA_ARCH}" IN_LIST ARCHITECTURES_WITH_ACCEL)
if(CUDA_ARCH MATCHES "^([0-9]+)f$")
# Family code, no -real 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-real")
else()
list(APPEND CMAKE_CUDA_ARCHITECTURES_NORMALIZED "${CUDA_ARCH}-real")
Expand Down
Loading