[Bugfix] Fix broken NVVM caused by CuteDSL 4.6.0 (#49108)
Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg>
Thien Tran committed
Jul 20, 2026 at 10:45 UTC
2730b657c4aa92a1a342a243514bbbd68804866f
2 files changed
+11
-4
.buildkite/test_areas/kernels.yaml
+7
@@ -237,6 +237,11 @@ steps:
237
- vllm/model_executor/kernels/linear/cute_dsl/ll_bf16.py
238
- vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_dotprod.py
239
- vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py
240
+ - vllm/cute_utils/
241
+ - vllm/model_executor/layers/mamba/ops/gdn_chunk_cutedsl/
242
+ - vllm/model_executor/layers/fused_moe/router/bf16x3_router_gemm_cutedsl.py
243
+ - tests/kernels/mamba/test_gdn_prefill_cutedsl.py
244
+ - tests/kernels/test_bf16x3_router_gemm_cutedsl.py
245
- tests/kernels/test_ll_bf16_gemm.py
246
- tests/kernels/test_top_k_per_row.py
247
commands:
@@ -266,6 +271,8 @@ steps:
271
- pytest -v -s tests/kernels/moe/test_flashinfer_moe.py
272
- pytest -v -s tests/kernels/moe/test_trtllm_nvfp4_moe.py
273
- pytest -v -s tests/kernels/moe/test_cutedsl_moe.py
274
+ - pytest -v -s tests/kernels/mamba/test_gdn_prefill_cutedsl.py
275
+ - pytest -v -s tests/kernels/test_bf16x3_router_gemm_cutedsl.py
276
- pytest -v -s tests/kernels/test_ll_bf16_gemm.py
277
# e2e
278
- pytest -v -s tests/models/quantization/test_nvfp4.py
vllm/cute_utils/_tcgen05.py
+4
-4
@@ -10,8 +10,8 @@ from cutlass.cutlass_dsl import dsl_user_op
10
11
NVVM_CTA_GROUP_MAP = [
12
None,
13
- nvvm.Tcgen05GroupKind.CTA_1,
14
- nvvm.Tcgen05GroupKind.CTA_2,
13
+ nvvm.CTAGroupKind.CTA_1,
14
+ nvvm.CTAGroupKind.CTA_2,
15
]
16
LDST_MAP = {
17
"32x32b": (nvvm.Tcgen05LdStShape.SHAPE_32X32B, 1),
@@ -136,7 +136,7 @@ def commit(mbar, cta_mask=None, cta_group: int = 1, *, loc=None, ip=None):
136
group = NVVM_CTA_GROUP_MAP[cta_group]
137
if cutlass.const_expr(cta_mask is not None):
138
with cute.arch.elect_one():
139
- nvvm.tcgen05_commit_arrive(
139
+ nvvm.tcgen05_commit(
140
mbar_llvm,
141
multicast_mask=cta_mask.ir_value(loc=loc, ip=ip),
142
group=group,
@@ -145,7 +145,7 @@ def commit(mbar, cta_mask=None, cta_group: int = 1, *, loc=None, ip=None):
145
)
146
else:
147
with cute.arch.elect_one():
148
- nvvm.tcgen05_commit_arrive(mbar_llvm, group=group, loc=loc, ip=ip)
148
+ nvvm.tcgen05_commit(mbar_llvm, group=group, loc=loc, ip=ip)
149
150
151
@dsl_user_op