Skip to content

[TRTLLM-6286] [feat] Update CUTLASS to 4.2 and enable SM103 group gemm #7832

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
merged 4 commits into from
Sep 19, 2025
Merged

[TRTLLM-6286] [feat] Update CUTLASS to 4.2 and enable SM103 group gemm #7832

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
61973a9
[TRTLLM-6286] [feat] Update CUTLASS to 4.2 and enable SM103 group gemm
VALLIS-NERIA Sep 16, 2025
ed4ef2c
Merge remote-tracking branch 'origin/main' into user/xiweny/cutlass4.2
VALLIS-NERIA Sep 18, 2025
cad9621
AI suggestion
VALLIS-NERIA Sep 16, 2025
bd62825
Merge branch 'main' into user/xiweny/cutlass4.2
VALLIS-NERIA Sep 18, 2025
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
2 changes: 1 addition & 1 deletion 3rdparty/cutlass
Open in desktop
Submodule cutlass updated 175 files
13 changes: 7 additions & 6 deletions cpp/tensorrt_llm/kernels/cutlass_kernels/moe_gemm/moe_gemm_template_dispatch_tma_ws.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -138,22 +138,23 @@ void dispatchMoeGemmFinalDispatchTmaWarpSpecialized(TmaWarpSpecializedGroupedGem
TLLM_THROW("Please recompile with support for hopper by passing 90-real as an arch to build_wheel.py.");
}
#endif
// #ifndef COMPILE_BLACKWELL_SM103_TMA_GROUPED_GEMMS
#ifndef COMPILE_BLACKWELL_SM103_TMA_GROUPED_GEMMS
else if constexpr (Arch::kMinComputeCapability == 103)
{
static std::once_flag flag;
std::call_once(flag,
[]()
{
TLLM_LOG_WARNING(
"Falling back to sm100f version due to a bug in cutlass." /*"For best performance please recompile with support for blackwell by "
"passing 103-real as an arch to build_wheel.py."*/);
"For best performance please recompile with support for blackwell by "
"passing 103-real as an arch to build_wheel.py.");
});
return dispatchMoeGemmFinalDispatchTmaWarpSpecialized<cutlass::arch::Sm100, T, WeightType, OutputType,
EpilogueTag, FUSION, TileShape, ClusterShape>(
dispatchMoeGemmFinalDispatchTmaWarpSpecialized<cutlass::arch::Sm100, T, WeightType, OutputType, EpilogueTag,
FUSION, TileShape, ClusterShape>(
hopper_input, num_experts, gemm_config, multi_processor_count, stream, occupancy, workspace_size);
return;
}
// #endif
#endif
#ifndef COMPILE_BLACKWELL_TMA_GROUPED_GEMMS
else if constexpr (Arch::kMinComputeCapability >= 100 && Arch::kMinComputeCapability < 120)
{
Expand Down
13 changes: 0 additions & 13 deletions cpp/tensorrt_llm/kernels/cutlass_kernels/python/generate_kernels.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,19 +3,6 @@
import os
from itertools import chain, product

file_to_patch = os.path.abspath(
os.path.join(
os.path.dirname(__file__),
"../../../../../3rdparty/cutlass/python/cutlass_library/heuristics_provider.py"
))
# replace "from library import" to "from cutlass_library.library import"
with open(file_to_patch, "r") as f:
file_contents = f.read()
with open(file_to_patch, "w") as f:
f.write(
file_contents.replace("from library import",
"from cutlass_library.library import"))

from cutlass_library import *


Expand Down
Loading