Skip to content

[moe training] Support cutlass-dsl 4.6 in the CuTe DSL utils - #11

Closed
wolfcomos wants to merge 1 commit into
mainfrom
4over6/ao5-cutedsl-compat
Closed

[moe training] Support cutlass-dsl 4.6 in the CuTe DSL utils#11
wolfcomos wants to merge 1 commit into
mainfrom
4over6/ao5-cutedsl-compat

Conversation

@wolfcomos

@wolfcomos wolfcomos commented Aug 30, 2026

Copy link
Copy Markdown
Owner

Stack (emulated ghstack, oldest at bottom):

Summary

cutlass-dsl 4.6.0 dropped the leading result-type argument from the nvvm wrapper calls (cvt_packfloat, cvt_packfloat_f32, fmax) that the MXFP8 CuTe DSL utils emit, so kernels built against the <= 4.5.x call form fail to trace on 4.6 images. This PR branches on the installed cutlass-dsl version and emits the matching call form; the version probe (importlib.metadata.version("nvidia-cutlass-dsl")) treats PackageNotFoundError as the current API rather than failing the whole package import.

Design notes

  • Container image churn ships cutlass under changed API and distribution names: the runtime availability probe checks module importability, not distribution names, so cutlass can be importable while the nvidia-cutlass-dsl distribution name is absent (re-vendored under another name). The version probe must not crash unrelated moe_training imports in that case; it assumes the current (4.6+) API and lets tracing surface any residual mismatch.
  • The version branch is evaluated once at module scope, next to the existing runtime-availability gate, so per-op call sites stay unchanged.

Test plan

CPU (no GPU; import + collection inside the CI-equivalent container, this branch first on PYTHONPATH):

python3 -c "import torchao.prototype.moe_training.kernels.mxfp8.cute_utils"  # OK
python3 -m pytest test/prototype/moe_training/ --collect-only -q
# 427 tests collected

Lint: ruff check and ruff format --check (ruff 0.11.6, the version pinned by .github/workflows/ruff_linter.yml) pass on the changed file.

GPU results come from the assembled stack tip on GB200 (full moe_training GPU suite: 207 passed / 0 failed), since ao CI has no SM100 runner. Runtime GPU evidence comes from cutlass-dsl 4.6.x environments (the new call form); the pre-4.6 branch is validated by CPU collection/inspection only on this branch.

Draft on the fork to stage the upstream submission; supersedes the exploratory stack (#7/#8/#9)

Review pass (2026-08-30)

  • Version compare is now major/minor only: prerelease/local version strings ("4.6.0.dev0", "+git...") of 4.6+ wheels were misclassified as pre-4.6 by the full-triple compare.
  • The probe also catches ValueError from unparsable version strings (re-vendored builds), not just PackageNotFoundError.

cutlass-dsl 4.6.0 dropped the leading result-type argument from the
nvvm cvt_packfloat / cvt_packfloat_f32 / fmax wrappers, so the MXFP8
CuTe DSL utils built against <=4.5.x fail to trace on 4.6 images.
Branch on the installed version and emit the matching call form.

The version probe reads importlib.metadata for nvidia-cutlass-dsl,
but the runtime availability probe checks importability rather than
distribution names, so cutlass can be importable while that exact
distribution name is absent (container images re-vendor it under
changed names). Treat PackageNotFoundError as the current API instead
of crashing every unrelated moe_training import.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
@wolfcomos
wolfcomos force-pushed the 4over6/ao5-cutedsl-compat branch from ab985fb to 20c78fe Compare August 31, 2026 02:00
@wolfcomos

Copy link
Copy Markdown
Owner Author

Closing — decoupled from the four-over-six series (nothing in it is needed by pytorch#4851pytorch#4853); branch preserved.

@wolfcomos wolfcomos closed this Aug 31, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant