From b510857d37640209a4a3e01a13260ac26f2a93f3 Mon Sep 17 00:00:00 2001 From: trienebutdiene <59642603+kimseunghyun-kr@users.noreply.github.com> Date: Sun, 18 Oct 2026 15:14:06 +0800 Subject: [PATCH] [Bugfix][Build] Build fused GDN MTP decode for SM80-SM90 with CUDA < 13 (#60896) Signed-off-by: kimseunghyun Signed-off-by: kimseunghyun Co-authored-by: Claude Opus 5.5 Co-authored-by: Shengqi Chen Co-authored-by: kimseunghyun --- CMakeLists.txt | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/CMakeLists.txt b/CMakeLists.txt index e67221978c..c84e2b9eed 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1119,6 +1119,11 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP") "9.0a;10.0f;12.0f" "${CUDA_ARCHS}") cuda_archs_loose_intersection(FUSED_GDN_DECODE_ARCHS "8.0;8.6;8.9;9.0a;10.0f;11.0f;12.0f" "${CUDA_ARCHS}") + else() + # The fused GDN decode kernel needs only cp.async and BF16 math. Only the + # family-specific targets (10.0f, 11.0f, 12.0f) need CUDA >= 13.0. + cuda_archs_loose_intersection(FUSED_GDN_DECODE_ARCHS + "8.0;8.6;8.9;9.0a" "${CUDA_ARCHS}") endif() if(FUSED_KDA_DECODE_ARCHS) set(FUSED_KDA_DECODE_SRC