From b1f4372038d4935bb6054e86283a9d103c00ee38 Mon Sep 17 00:00:00 2001 From: Mihai Chiorean Date: Thu, 2 Apr 2026 11:46:59 -0700 Subject: [PATCH 1/7] [#11932][fix] Filter CUTLASS MoE GEMM tile configs by device shared memory on SM121 SM121 (DGX Spark GB10) has 99 KiB shared memory per block vs 228 KiB on SM120 (B200). CUTLASS StageCountAutoCarveout computes pipeline stages assuming 228 KiB, causing TMA warp-specialized MoE grouped GEMM tactics to fail with opaque "Error Internal" on gemm.initialize(). This patch adds two fixes: 1. Runtime SMEM guard in moe_gemm_tma_ws_launcher.inl: checks kernel SharedStorage size against cudaDevAttrMaxSharedMemoryPerBlockOptin before launch, converting opaque CUTLASS errors into clear diagnostics that the autotuner can skip. 2. Heuristic filter in get_candidate_configs_sm120(): on devices with < 120 KiB SMEM, removes tile configs whose SharedStorage exceeds the device limit. Keeps only CtaShape128x128x64B which fits within 99 KiB including FINALIZE epilogue overhead (~80 KiB total). Signed-off-by: Mihai Chiorean --- .../cutlass_kernels/cutlass_heuristic.cpp | 27 ++++++++++++++++--- .../launchers/moe_gemm_tma_ws_launcher.inl | 14 +++++++++- 2 files changed, 37 insertions(+), 4 deletions(-) diff --git a/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp b/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp index 50794769b5ec..dc03bcd8b43a 100644 --- a/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp +++ b/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. All rights reserved. + * Copyright (c) 2020-2026, NVIDIA CORPORATION. All rights reserved. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -31,6 +31,7 @@ #pragma GCC diagnostic pop #endif // __GNUC +#include #include #include #include @@ -533,7 +534,6 @@ std::vector get_candidate_configs_sm120(CutlassGemmConfig::Ca // Mixed FP8 x FP4: restrict to 128x128x128B only candidate_configs.push_back(CutlassGemmConfig{CutlassTileConfigSM120::CtaShape128x128x128B, MainloopScheduleType::AUTO, EpilogueScheduleType::AUTO, ClusterShape::ClusterShape_1x1x1}); - return candidate_configs; } else if (config & CutlassGemmConfig::FP4_ONLY) { @@ -546,8 +546,29 @@ std::vector get_candidate_configs_sm120(CutlassGemmConfig::Ca MainloopScheduleType::AUTO, EpilogueScheduleType::AUTO, ClusterShape::ClusterShape_1x1x1}); candidate_configs.push_back(CutlassGemmConfig{CutlassTileConfigSM120::CtaShape256x128x64B, MainloopScheduleType::AUTO, EpilogueScheduleType::AUTO, ClusterShape::ClusterShape_1x1x1}); - return candidate_configs; } + // Filter configs by device shared memory. SM121 (GB10) has 99 KiB vs + // SM120 (B200) 228 KiB. On constrained devices, keep only CtaShape128x128x64B + // which fits within 99 KiB including FINALIZE epilogue (~80 KiB total). + // CtaShape128x256x64B/256x128x64B overflow with FINALIZE (~100 KiB). + // CtaShape128x128x128B also exceeds 99 KiB at typical stage counts. + { + constexpr int kMinSmemForFullTileSet = 120 * 1024; + int device = 0; + cudaGetDevice(&device); + int maxSmem = 0; + cudaDeviceGetAttribute(&maxSmem, cudaDevAttrMaxSharedMemoryPerBlockOptin, device); + + if (maxSmem < kMinSmemForFullTileSet) + { + auto it = std::remove_if(candidate_configs.begin(), candidate_configs.end(), + [](CutlassGemmConfig const& config) + { return config.tile_config_sm120 != CutlassTileConfigSM120::CtaShape128x128x64B; }); + candidate_configs.erase(it, candidate_configs.end()); + } + } + return candidate_configs; + TLLM_THROW("Not Implemented: SM120 group GEMM only supports mxfp8-mxfp4 mixed or nvfp4."); } else diff --git a/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl b/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl index 56552a484b5c..c48f941f340d 100644 --- a/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl +++ b/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2025, NVIDIA CORPORATION. All rights reserved. + * Copyright (c) 2020-2026, NVIDIA CORPORATION. All rights reserved. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -668,6 +668,18 @@ using namespace cutlass::epilogue; "Workspace is size %zu but only %zu were allocated", calculated_ws_size, \ tma_ws_input.gemm_workspace_size); \ \ + /* SM121 guard: check if kernel SMEM fits on this device before launch */ \ + { \ + using GemmKernel_ = typename GemmGrouped::GemmKernel; \ + int smem_size = static_cast(sizeof(typename GemmKernel_::SharedStorage)); \ + int device_ = 0; \ + cudaGetDevice(&device_); \ + int maxSmem_ = 0; \ + cudaDeviceGetAttribute(&maxSmem_, cudaDevAttrMaxSharedMemoryPerBlockOptin, device_); \ + TLLM_CHECK_WITH_INFO(smem_size <= maxSmem_, \ + "MoE grouped GEMM requires %d bytes shared memory but device supports %d", smem_size, maxSmem_); \ + } \ + \ auto can_implement = gemm.can_implement(args); \ TLLM_CHECK_WITH_INFO(can_implement == cutlass::Status::kSuccess, \ "Grouped GEMM kernel will fail for params. Error: " \ From fcd5096c9ddaf6c95c5469a7509bc1b2cbfd1f71 Mon Sep 17 00:00:00 2001 From: Mihai Chiorean Date: Thu, 11 Jun 2026 09:53:58 -0700 Subject: [PATCH 2/7] [#11932][doc] Correct SM family labels in the SMEM-filter comment MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The original comment said 'SM121 (GB10) has 99 KiB vs SM120 (B200) 228 KiB', but B200 is SM100, not SM120. SM120 is consumer Blackwell (RTX PRO 6000) which has the SAME 99 KiB SMEM as SM121 GB10 — confirmed independently in issue #11932 by @waynehacking8. The runtime code is unchanged and correct (it queries cudaDevAttrMaxSharedMemoryPerBlockOptin), only the prose is fixed to label arches accurately. Signed-off-by: Mihai Chiorean --- .../kernels/cutlass_kernels/cutlass_heuristic.cpp | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp b/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp index 26a80a7bba33..91a3d8660ab9 100644 --- a/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp +++ b/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp @@ -547,9 +547,10 @@ std::vector get_candidate_configs_sm120(CutlassGemmConfig::Ca candidate_configs.push_back(CutlassGemmConfig{CutlassTileConfigSM120::CtaShape256x128x64B, MainloopScheduleType::AUTO, EpilogueScheduleType::AUTO, ClusterShape::ClusterShape_1x1x1}); } - // Filter configs by device shared memory. SM121 (GB10) has 99 KiB vs - // SM120 (B200) 228 KiB. On constrained devices, keep only CtaShape128x128x64B - // which fits within 99 KiB including FINALIZE epilogue (~80 KiB total). + // Filter configs by device shared memory. SM100 (B200) has 228 KiB, but + // consumer Blackwell (SM120 RTX PRO 6000, SM121 GB10 / DGX Spark) has only + // 99 KiB. On these constrained devices, keep only CtaShape128x128x64B which + // fits within 99 KiB including FINALIZE epilogue (~80 KiB total). // CtaShape128x256x64B/256x128x64B overflow with FINALIZE (~100 KiB). // CtaShape128x128x128B also exceeds 99 KiB at typical stage counts. { From 3c39abe6ca80d7720fc9d2717e7d58d27597898d Mon Sep 17 00:00:00 2001 From: Mihai Chiorean Date: Mon, 29 Jun 2026 16:13:40 -0700 Subject: [PATCH 3/7] [#11932][fix] Preserve SM120 mixed MoE low-SMEM candidates Signed-off-by: Mihai Chiorean --- .../kernels/cutlass_kernels/cutlass_heuristic.cpp | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) diff --git a/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp b/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp index 91a3d8660ab9..1d31218b20b6 100644 --- a/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp +++ b/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp @@ -531,9 +531,12 @@ std::vector get_candidate_configs_sm120(CutlassGemmConfig::Ca std::vector candidate_configs; if (config & CutlassGemmConfig::FP8FP4_MIXED) { - // Mixed FP8 x FP4: restrict to 128x128x128B only + // Mixed FP8 x FP4: prefer 128x128x128B on devices with enough + // shared memory, but keep a low-SMEM fallback for SM120/SM121. candidate_configs.push_back(CutlassGemmConfig{CutlassTileConfigSM120::CtaShape128x128x128B, MainloopScheduleType::AUTO, EpilogueScheduleType::AUTO, ClusterShape::ClusterShape_1x1x1}); + candidate_configs.push_back(CutlassGemmConfig{CutlassTileConfigSM120::CtaShape128x128x64B, + MainloopScheduleType::AUTO, EpilogueScheduleType::AUTO, ClusterShape::ClusterShape_1x1x1}); } else if (config & CutlassGemmConfig::FP4_ONLY) { @@ -547,6 +550,10 @@ std::vector get_candidate_configs_sm120(CutlassGemmConfig::Ca candidate_configs.push_back(CutlassGemmConfig{CutlassTileConfigSM120::CtaShape256x128x64B, MainloopScheduleType::AUTO, EpilogueScheduleType::AUTO, ClusterShape::ClusterShape_1x1x1}); } + else + { + TLLM_THROW("Not Implemented: SM120 group GEMM only supports mxfp8-mxfp4 mixed or nvfp4."); + } // Filter configs by device shared memory. SM100 (B200) has 228 KiB, but // consumer Blackwell (SM120 RTX PRO 6000, SM121 GB10 / DGX Spark) has only // 99 KiB. On these constrained devices, keep only CtaShape128x128x64B which @@ -569,8 +576,6 @@ std::vector get_candidate_configs_sm120(CutlassGemmConfig::Ca } } return candidate_configs; - - TLLM_THROW("Not Implemented: SM120 group GEMM only supports mxfp8-mxfp4 mixed or nvfp4."); } else { From cc71dfb6ef5d33a0ec0d9932f1e80f5c86873bb7 Mon Sep 17 00:00:00 2001 From: Mihai Chiorean Date: Mon, 29 Jun 2026 17:04:29 -0700 Subject: [PATCH 4/7] [#11932][fix] Check current-device SMEM for MoE tactics Signed-off-by: Mihai Chiorean --- .../kernels/cutlass_kernels/cutlass_heuristic.cpp | 8 +++++--- .../launchers/moe_gemm_tma_ws_launcher.inl | 7 +++++-- .../moe_gemm_tma_ws_mixed_input_launcher.inl | 15 ++++++++++++++- 3 files changed, 24 insertions(+), 6 deletions(-) diff --git a/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp b/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp index 1d31218b20b6..76a3904e2ce7 100644 --- a/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp +++ b/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp @@ -17,6 +17,7 @@ #include "tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.h" #include "tensorrt_llm/common/config.h" #include "tensorrt_llm/common/cudaBf16Wrapper.h" +#include "tensorrt_llm/common/cudaUtils.h" #ifdef __GNUC__ // Check if the compiler is GCC or Clang #pragma GCC diagnostic push @@ -563,13 +564,14 @@ std::vector get_candidate_configs_sm120(CutlassGemmConfig::Ca { constexpr int kMinSmemForFullTileSet = 120 * 1024; int device = 0; - cudaGetDevice(&device); + tensorrt_llm::common::check_cuda_error(cudaGetDevice(&device)); int maxSmem = 0; - cudaDeviceGetAttribute(&maxSmem, cudaDevAttrMaxSharedMemoryPerBlockOptin, device); + tensorrt_llm::common::check_cuda_error( + cudaDeviceGetAttribute(&maxSmem, cudaDevAttrMaxSharedMemoryPerBlockOptin, device)); if (maxSmem < kMinSmemForFullTileSet) { - auto it = std::remove_if(candidate_configs.begin(), candidate_configs.end(), + auto const it = std::remove_if(candidate_configs.begin(), candidate_configs.end(), [](CutlassGemmConfig const& config) { return config.tile_config_sm120 != CutlassTileConfigSM120::CtaShape128x128x64B; }); candidate_configs.erase(it, candidate_configs.end()); diff --git a/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl b/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl index eed1fcd790ae..61f5093ef5a6 100644 --- a/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl +++ b/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl @@ -15,6 +15,8 @@ */ #pragma once +#include "tensorrt_llm/common/cudaUtils.h" + #include "cutlass/array.h" #include "cutlass/numeric_conversion.h" @@ -675,9 +677,10 @@ using namespace cutlass::epilogue; using GemmKernel_ = typename GemmGrouped::GemmKernel; \ int smem_size = static_cast(sizeof(typename GemmKernel_::SharedStorage)); \ int device_ = 0; \ - cudaGetDevice(&device_); \ + tensorrt_llm::common::check_cuda_error(cudaGetDevice(&device_)); \ int maxSmem_ = 0; \ - cudaDeviceGetAttribute(&maxSmem_, cudaDevAttrMaxSharedMemoryPerBlockOptin, device_); \ + tensorrt_llm::common::check_cuda_error( \ + cudaDeviceGetAttribute(&maxSmem_, cudaDevAttrMaxSharedMemoryPerBlockOptin, device_)); \ TLLM_CHECK_WITH_INFO(smem_size <= maxSmem_, \ "MoE grouped GEMM requires %d bytes shared memory but device supports %d", smem_size, maxSmem_); \ } \ diff --git a/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_mixed_input_launcher.inl b/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_mixed_input_launcher.inl index f37920dcf73c..cdf5ea8dc3c9 100644 --- a/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_mixed_input_launcher.inl +++ b/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_mixed_input_launcher.inl @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. All rights reserved. + * Copyright (c) 2020-2026, NVIDIA CORPORATION. All rights reserved. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,6 +19,8 @@ #pragma GCC diagnostic ignored "-Wstrict-aliasing" #endif // __GNUC__ +#include "tensorrt_llm/common/cudaUtils.h" + #include "cutlass/epilogue/collective/default_epilogue.hpp" #include "cutlass/epilogue/thread/linear_combination.h" #include "cutlass/gemm/collective/collective_builder.hpp" @@ -273,6 +275,17 @@ void sm90_generic_mixed_moe_gemm_kernelLauncher(GroupedGemmInput(sizeof(typename GemmKernel::SharedStorage)); + int device = 0; + tensorrt_llm::common::check_cuda_error(cudaGetDevice(&device)); + int maxSmem = 0; + tensorrt_llm::common::check_cuda_error( + cudaDeviceGetAttribute(&maxSmem, cudaDevAttrMaxSharedMemoryPerBlockOptin, device)); + TLLM_CHECK_WITH_INFO(smem_size <= maxSmem, + "Mixed dtype WS grouped GEMM requires %d bytes shared memory but device supports %d", smem_size, maxSmem); + } + auto can_implement = gemm.can_implement(arguments); if (can_implement != cutlass::Status::kSuccess) { From 2d1335012358f3bb10efcc3f9d7ebcda34203313 Mon Sep 17 00:00:00 2001 From: Mihai Chiorean Date: Tue, 14 Jul 2026 15:27:26 -0700 Subject: [PATCH 5/7] [#11932][chore] Clarify MoE SMEM guard comment Signed-off-by: Mihai Chiorean --- .../moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl b/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl index 6e7498270682..9c5ebbdaa19d 100644 --- a/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl +++ b/cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/launchers/moe_gemm_tma_ws_launcher.inl @@ -679,7 +679,7 @@ using namespace cutlass::epilogue; "Workspace is size %zu but only %zu were allocated", calculated_ws_size, \ tma_ws_input.gemm_workspace_size); \ \ - /* SM121 guard: check if kernel SMEM fits on this device before launch */ \ + /* Check if kernel SMEM fits on the active device before launch. */ \ { \ using GemmKernel_ = typename GemmGrouped::GemmKernel; \ int smem_size = static_cast(sizeof(typename GemmKernel_::SharedStorage)); \ From b7e605a59cd4da197dd9b97150eef630302cee32 Mon Sep 17 00:00:00 2001 From: Mihai Chiorean Date: Fri, 17 Jul 2026 12:41:53 -0700 Subject: [PATCH 6/7] [#11932][fix] Preserve SM120 mixed FP8-FP4 MoE tactic Signed-off-by: Mihai Chiorean --- .../kernels/cutlass_kernels/cutlass_heuristic.cpp | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp b/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp index c5a0c23ec69a..bcc0eb1165ac 100644 --- a/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp +++ b/cpp/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp @@ -575,12 +575,10 @@ std::vector get_candidate_configs_sm120(CutlassGemmConfig::Ca std::vector candidate_configs; if (config & CutlassGemmConfig::FP8FP4_MIXED) { - // Mixed FP8 x FP4: prefer 128x128x128B on devices with enough - // shared memory, but keep a low-SMEM fallback for SM120/SM121. + // Mixed FP8 x FP4 only supports the 128x128x128B tile. candidate_configs.push_back(CutlassGemmConfig{CutlassTileConfigSM120::CtaShape128x128x128B, MainloopScheduleType::AUTO, EpilogueScheduleType::AUTO, ClusterShape::ClusterShape_1x1x1}); - candidate_configs.push_back(CutlassGemmConfig{CutlassTileConfigSM120::CtaShape128x128x64B, - MainloopScheduleType::AUTO, EpilogueScheduleType::AUTO, ClusterShape::ClusterShape_1x1x1}); + return candidate_configs; } else if (config & CutlassGemmConfig::FP4_ONLY) { From e44f5c703b9e17201ca7fc3e138dfa78f3050bc7 Mon Sep 17 00:00:00 2001 From: Mihai Chiorean Date: Mon, 20 Jul 2026 09:40:52 -0700 Subject: [PATCH 7/7] Add 'siyidNV' to the list of contributors revert accidental change Signed-off-by: Mihai Chiorean --- .github/workflows/blossom-ci.yml | 1 + 1 file changed, 1 insertion(+) diff --git a/.github/workflows/blossom-ci.yml b/.github/workflows/blossom-ci.yml index e398b8591d3c..f5a60658d8dd 100644 --- a/.github/workflows/blossom-ci.yml +++ b/.github/workflows/blossom-ci.yml @@ -310,6 +310,7 @@ jobs: "shuyixiong", "shyeh25", "SimengLiu-nv", + "siyidNV", "sklevtsov-nvidia", "StanleySun639", "stnie",