From 3d82dacd580a2313064dd4efce482b9c51ee8698 Mon Sep 17 00:00:00 2001 From: Mohammad Miadh Angkad <176301910+mmangkad@users.noreply.github.com> Date: Mon, 20 Jul 2026 17:11:59 +0800 Subject: [PATCH] Bump CuTe DSL to 4.6.0 (#31714) --- python/pyproject.toml | 6 ++-- .../ops/attention/cute_utils/_tcgen05.py | 8 ++--- .../sglang/multimodal_gen/test/test_utils.py | 2 +- python/sglang/srt/utils/common.py | 2 +- scripts/ci/cuda/ci_install_dependency.sh | 29 ------------------- 5 files changed, 9 insertions(+), 38 deletions(-) diff --git a/python/pyproject.toml b/python/pyproject.toml index 3fae86339..f0146ceba 100755 --- a/python/pyproject.toml +++ b/python/pyproject.toml @@ -30,7 +30,7 @@ dependencies = [ "easydict", # Required by remote model code (e.g. DeepSeek-OCR) loaded via trust_remote_code; validated by transformers 5.4+ check_imports "einops", "fastapi", - "flash-attn-4==4.0.0b15", + "flash-attn-4>=4.0.0b18", "flashinfer_python[cu13]==0.6.14", # keep it aligned with jit-cache version in Dockerfile "gguf", "helion==0.2.6", @@ -45,7 +45,7 @@ dependencies = [ "ninja", "numba==0.65.1", "numpy", - "nvidia-cutlass-dsl[cu13]==4.5.2", + "nvidia-cutlass-dsl[cu13]==4.6.0", "nvidia-mathdx==25.6.0", "nvidia-ml-py", "openai==2.6.1", @@ -62,7 +62,7 @@ dependencies = [ "pydantic", "python-multipart", "pyzmq>=25.1.2", - "quack-kernels>=0.4.1", + "quack-kernels>=0.6.1", "requests", "scipy", "sentencepiece", diff --git a/python/sglang/kernels/ops/attention/cute_utils/_tcgen05.py b/python/sglang/kernels/ops/attention/cute_utils/_tcgen05.py index 8c0bdd3bc..5e1f2a327 100644 --- a/python/sglang/kernels/ops/attention/cute_utils/_tcgen05.py +++ b/python/sglang/kernels/ops/attention/cute_utils/_tcgen05.py @@ -11,8 +11,8 @@ from cutlass.cutlass_dsl import dsl_user_op NVVM_CTA_GROUP_MAP = [ None, - nvvm.Tcgen05GroupKind.CTA_1, - nvvm.Tcgen05GroupKind.CTA_2, + nvvm.CTAGroupKind.CTA_1, + nvvm.CTAGroupKind.CTA_2, ] LDST_MAP = { "32x32b": (nvvm.Tcgen05LdStShape.SHAPE_32X32B, 1), @@ -134,7 +134,7 @@ def commit(mbar, cta_mask=None, cta_group: int = 1, *, loc=None, ip=None): mbar_llvm = mbar.to_llvm_ptr(loc=loc, ip=ip) group = NVVM_CTA_GROUP_MAP[cta_group] if cutlass.const_expr(cta_mask is not None): - nvvm.tcgen05_commit_arrive( + nvvm.tcgen05_commit( mbar_llvm, multicast_mask=cta_mask.ir_value(loc=loc, ip=ip), group=group, @@ -142,7 +142,7 @@ def commit(mbar, cta_mask=None, cta_group: int = 1, *, loc=None, ip=None): ip=ip, ) else: - nvvm.tcgen05_commit_arrive(mbar_llvm, group=group, loc=loc, ip=ip) + nvvm.tcgen05_commit(mbar_llvm, group=group, loc=loc, ip=ip) @dsl_user_op diff --git a/python/sglang/multimodal_gen/test/test_utils.py b/python/sglang/multimodal_gen/test/test_utils.py index 844220f1b..4e83f2716 100644 --- a/python/sglang/multimodal_gen/test/test_utils.py +++ b/python/sglang/multimodal_gen/test/test_utils.py @@ -34,7 +34,7 @@ if TYPE_CHECKING: logger = init_logger(__name__) -SGL_TEST_FILES_CI_DATA_REVISION = "d51ca9623e0bb27087da243a44c942fdda5aafe5" +SGL_TEST_FILES_CI_DATA_REVISION = "320949ecc2587474a2f535229ffc8f47ed16ee51" if current_platform.is_npu(): SGL_TEST_FILES_CI_DATA_REVISION = "6b62f4b6825c76a25fd2ba28248df68f2b400e65" diff --git a/python/sglang/srt/utils/common.py b/python/sglang/srt/utils/common.py index e697342f4..103a89c98 100644 --- a/python/sglang/srt/utils/common.py +++ b/python/sglang/srt/utils/common.py @@ -1820,7 +1820,7 @@ def suppress_noisy_warnings(): cutlass_dsl_noisy = { ( DeprecationWarning, - "Use explicit `struct.scalar.ptr` for pointer instead.", + "Using `struct.scalar` as pointer is deprecated.", ), ( UserWarning, diff --git a/scripts/ci/cuda/ci_install_dependency.sh b/scripts/ci/cuda/ci_install_dependency.sh index fa9645fc5..c819a8781 100755 --- a/scripts/ci/cuda/ci_install_dependency.sh +++ b/scripts/ci/cuda/ci_install_dependency.sh @@ -417,34 +417,6 @@ download_flashinfer_cache() { mark_step_done "${FUNCNAME[0]}" } -force_reinstall_cutlass_dsl_libs_cu13() { - # nvidia-cutlass-dsl[cu13] has additive PyPI extras: installing it pulls in - # both -libs-base and -libs-cu13. The two wheels ship intentionally-different - # content for the same paths (cutlass/_mlir/dialects/_gpu_ops_gen.py and - # cutlass/_mlir/_mlir_libs/_cutlass_ir.cpython-*.so) -- each Python wrapper - # is paired with a matching pybind11 .so. If install order leaves the .py - # from one wheel and the .so from the other, GPUModuleOp.__init__ raises - # TypeError: incompatible function arguments at kernel-compile time. - # - # Force-reinstall -libs-cu13 LAST so both files come from the same wheel - # (BOTH-cu13 state), eliminating the mismatch. The version is parsed from - # pyproject.toml so this stays in sync with whatever nvidia-cutlass-dsl - # version the project pins. - if [ "$CU_MAJOR" != "13" ]; then - return - fi - - CUTLASS_DSL_VERSION=$(grep -Po -m1 'nvidia-cutlass-dsl(\[[^]]+\])?==\K[0-9A-Za-z\.\-]+' "${REPO_ROOT}/python/pyproject.toml" || echo "") - if [ -z "$CUTLASS_DSL_VERSION" ]; then - echo "WARNING: could not detect nvidia-cutlass-dsl version from pyproject.toml; skipping libs-cu13 force-reinstall" - return - fi - - $PIP_CMD install --force-reinstall --no-deps "nvidia-cutlass-dsl-libs-cu13==${CUTLASS_DSL_VERSION}" $PIP_INSTALL_SUFFIX - - mark_step_done "${FUNCNAME[0]}" -} - stabilize_flashinfer_jit_paths() { # In venv mode, FlashInfer JIT writes build.ninja with hardcoded -isystem # paths. Per-job venvs get unique paths, but the JIT cache is shared on the @@ -624,7 +596,6 @@ main() { install_sglang_router install_flashinfer_cubin download_flashinfer_cache - force_reinstall_cutlass_dsl_libs_cu13 stabilize_flashinfer_jit_paths install_extra_deps install_test_tools