Bump CuTe DSL to 4.6.0 (#31714)
This commit is contained in:
@@ -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
|
"easydict", # Required by remote model code (e.g. DeepSeek-OCR) loaded via trust_remote_code; validated by transformers 5.4+ check_imports
|
||||||
"einops",
|
"einops",
|
||||||
"fastapi",
|
"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
|
"flashinfer_python[cu13]==0.6.14", # keep it aligned with jit-cache version in Dockerfile
|
||||||
"gguf",
|
"gguf",
|
||||||
"helion==0.2.6",
|
"helion==0.2.6",
|
||||||
@@ -45,7 +45,7 @@ dependencies = [
|
|||||||
"ninja",
|
"ninja",
|
||||||
"numba==0.65.1",
|
"numba==0.65.1",
|
||||||
"numpy",
|
"numpy",
|
||||||
"nvidia-cutlass-dsl[cu13]==4.5.2",
|
"nvidia-cutlass-dsl[cu13]==4.6.0",
|
||||||
"nvidia-mathdx==25.6.0",
|
"nvidia-mathdx==25.6.0",
|
||||||
"nvidia-ml-py",
|
"nvidia-ml-py",
|
||||||
"openai==2.6.1",
|
"openai==2.6.1",
|
||||||
@@ -62,7 +62,7 @@ dependencies = [
|
|||||||
"pydantic",
|
"pydantic",
|
||||||
"python-multipart",
|
"python-multipart",
|
||||||
"pyzmq>=25.1.2",
|
"pyzmq>=25.1.2",
|
||||||
"quack-kernels>=0.4.1",
|
"quack-kernels>=0.6.1",
|
||||||
"requests",
|
"requests",
|
||||||
"scipy",
|
"scipy",
|
||||||
"sentencepiece",
|
"sentencepiece",
|
||||||
|
|||||||
@@ -11,8 +11,8 @@ from cutlass.cutlass_dsl import dsl_user_op
|
|||||||
|
|
||||||
NVVM_CTA_GROUP_MAP = [
|
NVVM_CTA_GROUP_MAP = [
|
||||||
None,
|
None,
|
||||||
nvvm.Tcgen05GroupKind.CTA_1,
|
nvvm.CTAGroupKind.CTA_1,
|
||||||
nvvm.Tcgen05GroupKind.CTA_2,
|
nvvm.CTAGroupKind.CTA_2,
|
||||||
]
|
]
|
||||||
LDST_MAP = {
|
LDST_MAP = {
|
||||||
"32x32b": (nvvm.Tcgen05LdStShape.SHAPE_32X32B, 1),
|
"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)
|
mbar_llvm = mbar.to_llvm_ptr(loc=loc, ip=ip)
|
||||||
group = NVVM_CTA_GROUP_MAP[cta_group]
|
group = NVVM_CTA_GROUP_MAP[cta_group]
|
||||||
if cutlass.const_expr(cta_mask is not None):
|
if cutlass.const_expr(cta_mask is not None):
|
||||||
nvvm.tcgen05_commit_arrive(
|
nvvm.tcgen05_commit(
|
||||||
mbar_llvm,
|
mbar_llvm,
|
||||||
multicast_mask=cta_mask.ir_value(loc=loc, ip=ip),
|
multicast_mask=cta_mask.ir_value(loc=loc, ip=ip),
|
||||||
group=group,
|
group=group,
|
||||||
@@ -142,7 +142,7 @@ def commit(mbar, cta_mask=None, cta_group: int = 1, *, loc=None, ip=None):
|
|||||||
ip=ip,
|
ip=ip,
|
||||||
)
|
)
|
||||||
else:
|
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
|
@dsl_user_op
|
||||||
|
|||||||
@@ -34,7 +34,7 @@ if TYPE_CHECKING:
|
|||||||
|
|
||||||
logger = init_logger(__name__)
|
logger = init_logger(__name__)
|
||||||
|
|
||||||
SGL_TEST_FILES_CI_DATA_REVISION = "d51ca9623e0bb27087da243a44c942fdda5aafe5"
|
SGL_TEST_FILES_CI_DATA_REVISION = "320949ecc2587474a2f535229ffc8f47ed16ee51"
|
||||||
|
|
||||||
if current_platform.is_npu():
|
if current_platform.is_npu():
|
||||||
SGL_TEST_FILES_CI_DATA_REVISION = "6b62f4b6825c76a25fd2ba28248df68f2b400e65"
|
SGL_TEST_FILES_CI_DATA_REVISION = "6b62f4b6825c76a25fd2ba28248df68f2b400e65"
|
||||||
|
|||||||
@@ -1820,7 +1820,7 @@ def suppress_noisy_warnings():
|
|||||||
cutlass_dsl_noisy = {
|
cutlass_dsl_noisy = {
|
||||||
(
|
(
|
||||||
DeprecationWarning,
|
DeprecationWarning,
|
||||||
"Use explicit `struct.scalar.ptr` for pointer instead.",
|
"Using `struct.scalar` as pointer is deprecated.",
|
||||||
),
|
),
|
||||||
(
|
(
|
||||||
UserWarning,
|
UserWarning,
|
||||||
|
|||||||
@@ -417,34 +417,6 @@ download_flashinfer_cache() {
|
|||||||
mark_step_done "${FUNCNAME[0]}"
|
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() {
|
stabilize_flashinfer_jit_paths() {
|
||||||
# In venv mode, FlashInfer JIT writes build.ninja with hardcoded -isystem
|
# 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
|
# paths. Per-job venvs get unique paths, but the JIT cache is shared on the
|
||||||
@@ -624,7 +596,6 @@ main() {
|
|||||||
install_sglang_router
|
install_sglang_router
|
||||||
install_flashinfer_cubin
|
install_flashinfer_cubin
|
||||||
download_flashinfer_cache
|
download_flashinfer_cache
|
||||||
force_reinstall_cutlass_dsl_libs_cu13
|
|
||||||
stabilize_flashinfer_jit_paths
|
stabilize_flashinfer_jit_paths
|
||||||
install_extra_deps
|
install_extra_deps
|
||||||
install_test_tools
|
install_test_tools
|
||||||
|
|||||||
Reference in New Issue
Block a user