From e37889f489d2d357aba2367931b8ef02bdd10d19 Mon Sep 17 00:00:00 2001 From: Jiaxuan Bai Date: Wed, 8 Jul 2026 10:16:39 +0800 Subject: [PATCH] Fix grouped GEMM with CUTLASS DSL 4.7 (#3369) The Hopper grouped GEMM wrapper selected the current NVVM TMA bindings only when the CUTLASS version prefix was exactly 4.6. A 4.7 wheel therefore fell back to the removed CpAsyncBulkTensorLoadMode/loadMode interface. Detect the TMALoadMode capability directly, preserving the legacy fallback for older bindings. For example, 4.7 now uses the TMALoadMode/mode path and descriptor address-space cast. Verification: Python syntax compilation and diff checks. --- .../CuTeDSL/cute/hopper/kernel/grouped_gemm/grouped_gemm.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/examples/python/CuTeDSL/cute/hopper/kernel/grouped_gemm/grouped_gemm.py b/examples/python/CuTeDSL/cute/hopper/kernel/grouped_gemm/grouped_gemm.py index 9ef2a66b1..26a7e6e51 100644 --- a/examples/python/CuTeDSL/cute/hopper/kernel/grouped_gemm/grouped_gemm.py +++ b/examples/python/CuTeDSL/cute/hopper/kernel/grouped_gemm/grouped_gemm.py @@ -121,9 +121,9 @@ def CpAsyncBulkTensorGlobalToSharedClusterOpWrapper( loc = None, ip = None, ) -> None: - if cutlass.__version__[0:3] == "4.6": + if hasattr(_nvvm_d, "TMALoadMode"): from cutlass._mlir.dialects import llvm as _llvm_d - # Handle new nvvm API in 4.6+ + # Handle the current NVVM API when available. tmaDescriptor = _llvm_d.addrspacecast( _llvm_d.PointerType.get(_CuteAddressSpace.generic), tmaDescriptor, loc=loc, ip=ip )