Skip to content

[TRTLLM-13767][chore] upgrade CUTLASS DSL stack to 4.6.1 - #17274

Open
taianz-nv wants to merge 10 commits into
NVIDIA:mainfrom
taianz-nv:user/taianz/TRTLLM-13767-cutlass-dsl-4-6-1
Open

[TRTLLM-13767][chore] upgrade CUTLASS DSL stack to 4.6.1#17274
taianz-nv wants to merge 10 commits into
NVIDIA:mainfrom
taianz-nv:user/taianz/TRTLLM-13767-cutlass-dsl-4-6-1

Conversation

@taianz-nv

@taianz-nv taianz-nv commented Aug 5, 2026

Copy link
Copy Markdown
Collaborator

Summary

Upgrade the supported Blackwell CuTe DSL stack to nvidia-cutlass-dsl[cu13]==4.6.1. This version is required by the FP4 Conv3d integration in #17262, whose branch-merged shared-memory allocation uses the newer LaunchConfig API.

The upgrade is a coordinated dependency and source migration rather than a one-line package bump:

  • quack-kernels==0.6.2 is the release coupled exactly to CUTLASS DSL 4.6.1.
  • apache-tvm-ffi==0.1.13.post2 provides map_dataclass_to_tuple, which CUTLASS DSL 4.6.1 passes while compiling NVFP4 GEMM; the previous 0.1.6 pin fails at compile time. This pin also matches the version already selected by current main's security lock and remains inside every consumer's declared <0.2 compatibility range.
  • flash-attn-4==4.0.0b19 removes the old cute.core.ThrMma usage while remaining compatible with the released CUTLASS DSL 4.6.1 package. Its private forward API can append diagnostic values, so the VisualGen wrapper now consumes the output/LSE prefix only.
  • torch-c-dlpack-ext==0.1.3 remains pinned as part of the same FA4 runtime stack.

Scope

  • Align requirements.txt, constraints, the FA4 image installer, and the security-scanning Poetry project/lock file on the supported versions above.
  • Migrate all in-tree Python CuTeDSL consumers from the removed cute.make_fragment alias to cute.make_rmem_tensor, and from cute.core.ThrMma to cute.ThrMma.
  • Remove unused VSA PTX rounding helpers that depended on a CUTLASS 4.6-removed enum.
  • Keep optional VSA imports lazy only for unavailable modules/shared libraries; genuine CuTe API incompatibilities continue to fail loudly.
  • Register the existing FA4 key-padding-mask correctness test in the B200 L0 list.
  • Update the CuTe kernel-writing references to use the supported API.

This PR changes the shared CuTeDSL runtime stack and compatibility code; it does not add the FP4 VAE/Conv3d feature itself.

Validation

On an NSC B200 with CUTLASS DSL 4.6.1 and TVM-FFI 0.1.11 (the feature-bearing version used for compatibility validation), using real JIT compilation and kernel execution:

  • VSA correctness: 2/2 passed.
  • NVFP4 CuTeDSL GEMM through TVM-FFI: 6/6 passed.
  • VisualGen FA4 integration: 6/6 passed.
  • FA4 key-padding-mask correctness: 3/3 passed.
  • Representative non-VisualGen CuTeDSL coverage passed: FP8 block-scaled GEMM, FP4 paged MQA, FP8 paged MQA, GVR top-k decode, and argmax.
  • Follow-up NSC job 1636364 passed production dense GEMM native/TVM-FFI, filtered varlen top-k, and single-pass multi-CTA top-k.
  • Follow-up NSC job 1636394 passed the standalone BF16 GEMM runner's contiguous and strided branches.

The final TVM-FFI 0.1.13.post2 pin is newer than the NSC compatibility run above; it satisfies the declared CUTLASS DSL/Quack/FA4 ranges and is exercised by the final standard CI run.

Repository checks:

  • Full changed-file pre-commit suite passed.
  • Python syntax compilation passed for every changed Python file.
  • poetry check --lock passed for security_scanning/.
  • The branch was rebased onto main (33c6270c35) before the final validation run.

No new test implementation is needed: the relevant runtime/correctness tests already exist. This PR adds the missing B200 L0 registration for FA4 key-padding-mask behavior and expands the dependency-upgrade validation across the existing suites.

Compatibility notes

  • FA4 b19 may return optional diagnostic outputs after (output, lse); those are None on the VisualGen inference path used here and are intentionally ignored.
  • FA4 may emit an AuxData cannot be converted to JitArgument warning during first-shape JIT compilation. The tested kernels compile and execute successfully; this warning is non-fatal and does not occur in steady-state execution.

Review focus

Please review this as one supported CUTLASS/TVM-FFI/quack/FA4 version set. In particular, focus on the mechanical make_rmem_tensor API migration, VisualGen FA4 return compatibility, security-scanning lock consistency, and coverage of CuTeDSL consumers outside VisualGen.

Dev Engineer Review

  • Upgraded CUTLASS DSL to 4.6.1.
  • Pinned quack-kernels to 0.6.2.
  • Updated apache-tvm-ffi to 0.1.13.post2.
  • Updated FlashAttention 4 to 4.0.0b19.
  • Migrated removed CuTe APIs to cute.ThrMma and cute.make_rmem_tensor.
  • Improved optional VSA import handling for missing modules and shared libraries.
  • Removed obsolete PTX rounding helpers.
  • Removed the public VideoSparseAttentionForward export.
  • Updated FlashAttention result handling.
  • Updated CuTe DSL documentation for cute.make_rmem_tensor.
  • Updated standalone dense GEMM pointer and Constexpr handling.
  • Added the cu13 CUTLASS DSL extra to the devel image.
  • Dependency versions are consistent across the requirement, constraint, image, and security-scanning files.
  • Reported validation covered VSA, six NVFP4 CuTeDSL GEMM cases, FlashAttention 4, FP4 Conv3d consumers, and representative non-VisualGen CuTeDSL paths.
  • Changed-file pre-commit checks, Python syntax compilation, and the security-scanning lock check passed. Standard CI runs failed.

QA Engineer Review

  • Modified tests/integration/test_lists/test-db/l0_b200.yml.
  • Added unittest/_torch/visual_gen/test_fa4_key_padding_mask.py to the B200 pre-merge test list.
  • No test functions were added, modified, or removed.
  • Verdict: needs follow-up because CBTS coverage data is unavailable.

@taianz-nv

Copy link
Copy Markdown
Collaborator Author

/bot run

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #63913 [ run ] triggered by Bot. Commit: 95a8c32 Link to invocation

@coderabbitai

coderabbitai Bot commented Aug 5, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

Note

Reviews paused

It looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

  • @coderabbitai resume to resume automatic reviews.
  • @coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

  • ▶️ Resume reviews
  • 🔍 Trigger review

Walkthrough

The pull request updates CuTe DSL and FlashAttention dependencies, migrates kernel register allocations to cute.make_rmem_tensor, updates dense GEMM test pointers, improves Video Sparse Attention import diagnostics, and adjusts FlashAttention 4 handling and test coverage.

Changes

CuTe DSL and FlashAttention compatibility

Layer / File(s) Summary
Dependency and API alignment
requirements.txt, constraints.txt, security_scanning/pyproject.toml, docker/common/install_fa4.sh, docker/Dockerfile.multi, .claude/skills/kernel-cute-writing/references/*, tests/unittest/_torch/ray_orchestrator/multi_gpu/test_llm_update_weights_multi_gpu.py
Dependencies and Docker installation use updated CuTe DSL, FlashAttention 4, quack-kernels, and apache-tvm-ffi versions. The cute.make_rmem_tensor documentation now includes dtype.
Register tensor allocation migration
tensorrt_llm/_torch/cute_dsl_kernels/**/*.py
CuTe kernel temporaries, accumulators, caches, and output tensors now use cute.make_rmem_tensor instead of cute.make_fragment.
Video Sparse Attention compatibility
tensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/video_sparse_attention/*
Optional CuTe imports preserve ImportError and OSError causes. Internal kernel references and cute.ThrMma annotations are updated. Obsolete PTX helpers and package exports are removed.
FlashAttention 4 integration
tensorrt_llm/_torch/visual_gen/attention_backend/flash_attn4.py, tests/integration/test_lists/test-db/l0_b200.yml
FlashAttention 4 forward handling ignores extra returned diagnostics. The B200 key-padding-mask test is added to the pre-merge test list.
Dense GEMM pointer-based test updates
tests/scripts/cute_dsl_kernels/run_dense_bf16_gemm_persistent.py
Dense BF16 GEMM compilation, execution, and benchmark workspace generation now use aligned global-memory pointers for A and B. Strided invocation no longer passes max_active_clusters.

Estimated code review effort: 3 (Moderate) | ~25 minutes

Suggested reviewers: bowenfu, niukuo, cascade812

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 52.17% which is insufficient. The required threshold is 80.00%. Write docstrings for the functions missing them to satisfy the coverage threshold.
✅ Passed checks (4 passed)
Check name Status Explanation
Linked Issues check ✅ Passed Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check ✅ Passed Check skipped because no linked issues were found for this pull request.
Title check ✅ Passed The title clearly identifies the ticket, change type, and primary CUTLASS DSL stack upgrade to version 4.6.1.
Description check ✅ Passed The description clearly explains the scope, dependency changes, compatibility notes, validation results, and review focus; the template checklist is not explicitly reproduced but the required content is covered.
✨ Finishing Touches 💡 1
⚔️ Resolve merge conflicts 💡
  • Resolve merge conflict in branch user/taianz/TRTLLM-13767-cutlass-dsl-4-6-1
🧪 Generate unit tests (beta)
  • Create PR with unit tests

Comment @coderabbitai help to get the list of available commands.

@BowenFu

BowenFu commented Aug 5, 2026

Copy link
Copy Markdown

The AttributeError handling reads as defensive hygiene, but taken with the pins in the same file it may be describing a live break, and I couldn't settle that from the diff alone.

The main question. Your own comments say the skew is "flash_attn.cute / quack referencing a cute API removed in a newer cutlass-dsl". flash-attn-4 is a hard pin at 4.0.0b11 (requirements.txt:87) and this PR doesn't move it. If 4.0.0b11 really does touch an API that 4.6.1 removed, then after this merge the import fails in every environment, not just skewed ones — and FlashAttn4Attention.__init__ raises ImportError at construction (flash_attn4.py:107-110). That would take the FA4 backend and the parallel-combine path from working on 4.5.0 to hard-failing on 4.6.1, with the widened except making it look like a tolerated environment issue rather than a dependency that needs bumping alongside.

Can you state which it is? If FA4 needs a newer build to stay working on 4.6.1, that bump belongs in this PR; if 4.0.0b11 is fine and the catch is purely for out-of-tree installs, saying so in the comment would keep the next reader from having to re-derive it.

Same question for quack-kernels>=0.2.10 (requirements.txt:79, "required for MinimaxM3 MSA"). That's a floor rather than a pin, so a fresh resolve may pick up a 4.6.1-compatible build — but a pinned/cached 0.2.10 would not, and MSA has no equivalent graceful-degradation guard that I could find.

Scope of validation. nvidia-cutlass-dsl reaches well past VisualGen — the CuTe DSL MLA decode kernels, the NVFP4 GEMM autotuner tactic pruning (nvidia-matmul-heuristics right below it), and the DSv4 CuTe DSL paths all sit on it. "Exercised on NSC B200 by the dependent FP4 Conv3d test suite" covers the kernel that motivated the bump, not those. Your own PR body still says "Draft follow-up: Run the standard VisualGen import and attention-backend CI matrix before marking ready" — but the PR isn't a draft. Worth either running that matrix or moving the line out of the description.

Smaller ones:

  • The VSA __init__.py now re-exports VideoSparseAttentionForward from .interface, where it is None on import failure. A previously loud ImportError at import time becomes a None that surfaces later as a confusing AttributeError at the use site. block_sparse_attn_from_indices_cute guards on CUTE_AVAILABLE first, so the in-tree path is covered — it's direct importers of the symbol that lose the signal.
  • Heads-up on a merge conflict: [None][chore] Update flashinfer-python from 0.6.15 to 0.6.16 #17175 (approved) moves flashinfer-python to 0.6.16 in this same file, and this PR is based on 0.6.15. Whichever lands second will need a rebase.

@BowenFu

BowenFu commented Aug 5, 2026

Copy link
Copy Markdown

Following up on my own question — I pulled the wheel metadata, and it answers it. My FlashAttention-4 hypothesis was wrong; disregard that part.

flash-attn-4 4.0.0b11   Requires-Dist: nvidia-cutlass-dsl>=4.4.2
                        Requires-Dist: quack-kernels>=0.4.0
quack-kernels 0.2.10    Requires-Dist: nvidia-cutlass-dsl>=4.4.0
quack-kernels 0.4.0     Requires-Dist: nvidia-cutlass-dsl>=4.4.2
quack-kernels 0.6.2     Requires-Dist: nvidia-cutlass-dsl==4.6.1

4.0.0b11 declares an open lower bound, so 4.6.1 is inside its supported range. The widened except is defensive hygiene, not cover for a live break, and there's nothing for you to fix there.

But the same metadata shows a real side effect the PR doesn't mention. quack-kernels 0.6.2 pins nvidia-cutlass-dsl==4.6.1 exactly. So which quack gets installed is determined by your pin:

  • on main (==4.5.0), the 0.6.x line is excluded outright and the resolver settles on an older quack;
  • after this PR (==4.6.1), the 0.6.x line becomes the only one that satisfies both constraints.

So a one-line CUTLASS bump labelled [chore] also moves quack-kernels — the package requirements.txt:79 marks as "required for MinimaxM3 MSA" — across several minor versions, as a transitive consequence. That's worth saying out loud in the PR description, and it's a second reason the "FP4 Conv3d on B200" validation doesn't cover the blast radius: MSA is on a different kernel package version afterwards, and nothing in the diff exercises it.

Separately, requirements.txt:79 floors quack at >=0.2.10 while the flash-attn-4==4.0.0b11 you pin at :87 itself requires quack-kernels>=0.4.0. The repo's own floor is looser than what its pinned FA4 needs — worth tightening while you're in this file.

My remaining asks are unchanged and both cheap: run the VisualGen attention-backend CI matrix your description still calls for (the PR isn't a draft), and confirm the in-tree CuTe DSL consumers outside VisualGen — MLA decode, the NVFP4 GEMM autotuner — are exercised somewhere on 4.6.1.

@taianz-nv

Copy link
Copy Markdown
Collaborator Author

Thanks for tracing the dependency metadata and the broader CuTe impact. I addressed this as one supported dependency/API set rather than extending the fallback: quack-kernels is now pinned to 0.6.2 alongside CUTLASS DSL 4.6.1; the in-tree VSA cute.core.ThrMma uses are migrated to cute.ThrMma; the broad AttributeError catches were removed; and the nullable internal VSA class export was removed in favor of the guarded functional interface. I also updated the PR description to call out the dependency transition and cross-consumer validation scope. Commit: b5049f8.

@taianz-nv

Copy link
Copy Markdown
Collaborator Author

/bot run

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🧹 Nitpick comments (1)
requirements.txt (1)

77-79: 🗄️ Data Integrity & Integration | 🔵 Trivial | 💤 Low value

Document the quack-kernels coupling.

quack-kernels==0.6.2 declares nvidia-cutlass-dsl==4.6.1, so the exact pin is required. Update the inline comment to document this coupling; flash-attn-4 alone requires only quack-kernels>=0.4.0.

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@requirements.txt` around lines 77 - 79, Update the inline comment on the
quack-kernels==0.6.2 requirement to document that it declares and therefore
requires nvidia-cutlass-dsl==4.6.1; distinguish this coupling from flash-attn-4,
which only requires quack-kernels>=0.4.0.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Nitpick comments:
In `@requirements.txt`:
- Around line 77-79: Update the inline comment on the quack-kernels==0.6.2
requirement to document that it declares and therefore requires
nvidia-cutlass-dsl==4.6.1; distinguish this coupling from flash-attn-4, which
only requires quack-kernels>=0.4.0.

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 41364dc7-9c6e-4f65-bc76-b65fcbf8416d

📥 Commits

Reviewing files that changed from the base of the PR and between 95a8c32 and b5049f8.

📒 Files selected for processing (4)
  • requirements.txt
  • tensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/video_sparse_attention/__init__.py
  • tensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/video_sparse_attention/block_sparse_attn_dsl_fwd.py
  • tensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/video_sparse_attention/interface.py
🚧 Files skipped from review as they are similar to previous changes (1)
  • tensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/video_sparse_attention/interface.py

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #63929 [ run ] triggered by Bot. Commit: b5049f8 Link to invocation

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #63913 [ run ] completed with state ABORTED. Commit: 95a8c32
/LLM/main/L0_MergeRequest_PR pipeline #51852 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

  • Please check the failed tests and fix your PR
  • If you cannot view the failures, ask the CI triggerer to share details
  • Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

@taianz-nv

Copy link
Copy Markdown
Collaborator Author

Updated the inline dependency note to make the exact quack-kernels 0.6.2 / CUTLASS DSL 4.6.1 coupling explicit. Commit: 323fc25.

@taianz-nv

Copy link
Copy Markdown
Collaborator Author

/bot run

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #63934 [ run ] triggered by Bot. Commit: 323fc25 Link to invocation

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #63929 [ run ] completed with state ABORTED. Commit: b5049f8

Link to invocation

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #63934 [ run ] completed with state FAILURE. Commit: 323fc25
/LLM/main/L0_MergeRequest_PR pipeline #51869 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

  • Please check the failed tests and fix your PR
  • If you cannot view the failures, ask the CI triggerer to share details
  • Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

@BowenFu

BowenFu commented Aug 5, 2026

Copy link
Copy Markdown

Re-reviewed at 323fc2589. Both of my points are addressed, and the guard is now better than what it replaced.

The transitive coupling is declared: quack-kernels>=0.2.10==0.6.2 with the comment naming the CUTLASS DSL 4.6.1 pin. That was the substantive concern — on main the >=0.2.10 floor plus nvidia-cutlass-dsl==4.5.0 excluded the 0.6.x line entirely, so bumping CUTLASS silently made 0.6.x the only satisfiable resolution. Pinning it exactly makes the move visible in the diff instead of in a resolver, and it also closes the gap where the repo floored quack below what the pinned FA4 requires.

The import guard went the right direction rather than wider: except (ImportError, OSError) as error instead of adding AttributeError. That matters — an AttributeError from cutlass.cute is exactly the signal that the API moved under us, which is the thing you'd want to fail loudly on, and swallowing it would have turned a version-skew break into a silent fallback. Capturing _cute_import_error and re-raising with from also means the eventual RuntimeError carries the real cause instead of a generic "not importable".

Privatizing VideoSparseAttentionForward_VideoSparseAttentionForward and dropping it from the package __all__ is fine: it's under _torch/visual_gen/ (internal, not the tensorrt_llm/visual_gen/ public surface), and nothing outside that package imports it.

Still not approving as first reviewer — this moves two pinned dependencies and there are no approvals yet. The one ask left over from before is unchanged: the description still calls for the VisualGen attention-backend matrix, and it'd be good to confirm the non-VisualGen CuTe DSL consumers were exercised on 4.6.1, since nvidia-cutlass-dsl isn't only used here. Also still textually conflicts with #17175 in requirements.txt.

@zhenhuaw-me zhenhuaw-me left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM. Thanks!

@taianz-nv

Copy link
Copy Markdown
Collaborator Author

/bot run

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #64018 [ run ] triggered by Bot. Commit: 323fc25 Link to invocation

@taianz-nv

Copy link
Copy Markdown
Collaborator Author

/bot run

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #64296 [ run ] triggered by Bot. Commit: b1c0acb Link to invocation

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #64296 [ run ] completed with state FAILURE. Commit: b1c0acb
/LLM/main/L0_MergeRequest_PR pipeline #52196 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

  • Please check the failed tests and fix your PR
  • If you cannot view the failures, ask the CI triggerer to share details
  • Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

@brnguyen2 brnguyen2 left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Approving — the comments below are optional touch-ups, not blockers.

Clean, well-documented upgrade — the description matches the diff closely and the validation list is unusually thorough for a dependency bump. Two non-blocking asks:

  1. The make_fragmentmake_rmem_tensor rename is only checked at JIT-compile time, and three migrated kernels aren't in the validation list: dense_gemm_persistent.py, filtered_top_k_varlen_util.py, and single_pass_multi_cta_radix_topk.py. Can you confirm existing B200 CI exercises their compile paths, or run them once? The rename is mechanical so risk is low, but an uncovered kernel would fail loudly on first customer use rather than in CI.

  2. See the inline note on the missing [cu13] extra in constraints.txt — with 4.6.1 this now determines which GPU runtime libs land in the devel image.

For the record I verified no bare cute.make_fragment( calls remain in the tree (the make_fragment_A/B/C/_like hits are different, still-supported APIs), and that fmha_helpers.py's exp2 emulation uses the string-based rnd="rm" form rather than the removed nvvm.RoundingModeKind enum, so the ptx.py deletion leaves no broken sibling.

Comment thread constraints.txt
# Upgrade base image nvidia-cutlass-dsl 4.3.5 to 4.4.2
nvidia-cutlass-dsl>=4.4.2
# Upgrade the base-image nvidia-cutlass-dsl to the supported 4.6 API.
nvidia-cutlass-dsl==4.6.1

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

requirements.txt pins nvidia-cutlass-dsl[cu13]==4.6.1 but this pin drops the [cu13] extra. That mattered less before, but per the 4.6.1 metadata (visible in the poetry.lock in this PR) the base package now unconditionally depends on nvidia-cutlass-dsl-libs-cu12, with the cu13 libs only behind the extra — and Dockerfile.multi consumes this file with pip3 install -r, i.e. it actually installs these packages into the devel image after uninstalling the base image's copy. So the devel image gets cu12 runtime libs on a CUDA 13 base and no cu13 libs until requirements.txt is installed later. Suggest nvidia-cutlass-dsl[cu13]==4.6.1 here to keep the two files interchangeable.

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Good catch—the devel stage was installing the constraints as requirements, so the cu13 extra was lost. I did not add the extra to constraints.txt because that file is also consumed via -c, where pip rejects extras. In 1336c2e I request nvidia-cutlass-dsl[cu13]==4.6.1 at this Docker install site instead.

@taianz-nv

Copy link
Copy Markdown
Collaborator Author

@brnguyen2 Thanks for calling out the three kernels. I verified all three on an NSC B200 with CUTLASS DSL 4.6.1: the production dense GEMM native/TVM-FFI paths, filtered varlen top-k, and single-pass multi-CTA top-k all passed (job 1636364).

The standalone dense runner did expose stale Tensor/Pointer and runtime-Constexpr handling under 4.6.1. I fixed it in 1336c2e; its contiguous and strided numerical checks both pass (job 1636394). The same commit also fixes the devel image's missing cu13 extra without making constraints.txt invalid as a pip constraint file.

@taianz-nv

Copy link
Copy Markdown
Collaborator Author

/bot run

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 3

🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@docker/Dockerfile.multi`:
- Around line 55-60: Add the repository-standard NVIDIA copyright header at the
beginning of docker/Dockerfile.multi, using 2026 as the copyright year; leave
the existing Docker build commands unchanged.

In `@tests/scripts/cute_dsl_kernels/run_dense_bf16_gemm_persistent.py`:
- Around line 346-383: Add the return annotation -> cute.testing.JitArguments to
both generate_tensors functions, including the function shown and the second
generate_tensors definition nearby. Keep their existing implementations and
workspace_generator usage unchanged.
- Around line 221-233: Register run_dense_bf16_gemm_persistent.py in the
appropriate CI or QA test list. Configure the entry to exercise the run entry
point and both generate_tensors branches, reusing the test-list conventions for
similar CuTe DSL kernel scripts.
🪄 Autofix

Fix all unresolved CodeRabbit comments on this PR:

  • Push a commit to this branch (recommended)
  • Create a new PR with the fixes

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: b9e5ee51-3863-4b4b-9281-078cebc21d0b

📥 Commits

Reviewing files that changed from the base of the PR and between b1c0acb and 1336c2e.

📒 Files selected for processing (2)
  • docker/Dockerfile.multi
  • tests/scripts/cute_dsl_kernels/run_dense_bf16_gemm_persistent.py

Comment thread docker/Dockerfile.multi
Comment on lines +221 to +233
_, a_torch = cutlass_torch.cute_tensor_like(
a_ref, ab_dtype, is_dynamic_layout=True, assumed_align=16
)
b_tensor, b_torch = cutlass_torch.cute_tensor_like(
_, b_torch = cutlass_torch.cute_tensor_like(
b_ref, ab_dtype, is_dynamic_layout=True, assumed_align=16
)
c_tensor, c_torch = cutlass_torch.cute_tensor_like(
_, c_torch = cutlass_torch.cute_tensor_like(
c_ref, c_dtype, is_dynamic_layout=True, assumed_align=16
)

# Mark tensor to be byte aligned
a_tensor.mark_compact_shape_dynamic(
mode=1 if a_major == "k" else 0,
stride_order=(2, 0, 1) if a_major == "k" else (2, 1, 0),
divisibility=1,
)
b_tensor.mark_compact_shape_dynamic(
mode=1 if b_major == "k" else 0,
stride_order=(2, 0, 1) if b_major == "k" else (2, 1, 0),
divisibility=1,
)
c_tensor.mark_compact_shape_dynamic(
mode=1 if c_major == "n" else 0,
stride_order=(2, 0, 1) if c_major == "n" else (2, 1, 0),
divisibility=1,
a_ptr = make_ptr(ab_dtype, a_torch.data_ptr(), cute.AddressSpace.gmem, assumed_align=16)
b_ptr = make_ptr(ab_dtype, b_torch.data_ptr(), cute.AddressSpace.gmem, assumed_align=16)
c_tensor = from_dlpack(c_torch, assumed_align=16).mark_layout_dynamic(
leading_dim=1 if c_major == "n" else 0

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win

🧩 Analysis chain

🏁 Script executed:

#!/bin/bash
set -euo pipefail

fd -t f -e yml . tests/integration/test_lists 2>/dev/null | sort || true
rg -n -C 2 \
  'run_dense_bf16_gemm_persistent|dense_bf16_gemm_persistent|PersistentDenseGemmKernel' \
  tests/integration/test_lists/test-db \
  tests/integration/test_lists/qa || true

Repository: NVIDIA/TensorRT-LLM

Length of output: 4305


🏁 Script executed:

#!/bin/bash
set -euo pipefail

printf '%s\n' '--- changed files ---'
git diff --name-status

printf '%s\n' '--- relevant source structure ---'
ast-grep outline tests/scripts/cute_dsl_kernels/run_dense_bf16_gemm_persistent.py

printf '%s\n' '--- reviewed source ranges ---'
sed -n '180,255p' tests/scripts/cute_dsl_kernels/run_dense_bf16_gemm_persistent.py
sed -n '320,435p' tests/scripts/cute_dsl_kernels/run_dense_bf16_gemm_persistent.py

printf '%s\n' '--- repository-wide references ---'
rg -n -C 3 \
  'run_dense_bf16_gemm_persistent|dense_bf16_gemm_persistent|PersistentDenseGemmKernel' \
  . \
  -g '!build/**' -g '!dist/**' -g '!node_modules/**' || true

printf '%s\n' '--- test-list entry shapes ---'
for f in $(fd -t f . tests/integration/test_lists | sort | head -n 8); do
  printf '\n### %s\n' "$f"
  sed -n '1,80p' "$f"
done

Repository: NVIDIA/TensorRT-LLM

Length of output: 50377


🏁 Script executed:

#!/bin/bash
set -euo pipefail

printf '%s\n' '--- all test-list files ---'
git ls-files tests/integration/test_lists | sort

printf '%s\n' '--- direct script and directory references in test lists ---'
rg -n -i \
  'cute[_ -]?dsl|tests/scripts|run_dense|dense_gemm_persistent|gemm_persistent' \
  tests/integration/test_lists || true

printf '%s\n' '--- references outside test lists ---'
rg -n -i \
  'run_dense_bf16_gemm_persistent\.py|tests/scripts/cute_dsl_kernels|cute_dsl_kernels' \
  .github tests scripts tests/integration \
  -g '!tests/integration/test_lists/**' \
  -g '!**/__pycache__/**' || true

printf '%s\n' '--- function definitions in the reviewed script ---'
rg -n '^(async[[:space:]]+def|def)[[:space:]]+' \
  tests/scripts/cute_dsl_kernels/run_dense_bf16_gemm_persistent.py

Repository: NVIDIA/TensorRT-LLM

Length of output: 29471


Register the CuTe DSL test script.

No CI or QA test-list entry invokes tests/scripts/cute_dsl_kernels/run_dense_bf16_gemm_persistent.py. Add an appropriate entry that exercises run and both generate_tensors branches.

Test coverage summary: modified entry point run; modified workspace generators: both generate_tensors definitions; no test functions added or removed; coverage verdict: insufficient.

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@tests/scripts/cute_dsl_kernels/run_dense_bf16_gemm_persistent.py` around
lines 221 - 233, Register run_dense_bf16_gemm_persistent.py in the appropriate
CI or QA test list. Configure the entry to exercise the run entry point and both
generate_tensors branches, reusing the test-list conventions for similar CuTe
DSL kernel scripts.

Source: Path instructions

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't plan to register this standalone benchmark runner directly in a CI/QA test list. The lists contain pytest test nodes, while this is one of 12 manual CuTe DSL run_*.py utilities and none is registered directly. The production BF16 GEMM path is covered by TestDeepSeekV3Lite::test_cute_dsl_bf16_gemm[...] in tests/integration/test_lists/qa/llm_function_core.txt; for this change, both the contiguous and strided runner branches were also executed with numerical checks on an NSC B200 (job 1636394). Adding a new recurring hardware test for the manual runner would be disproportionate to this dependency-upgrade follow-up.

Comment thread tests/scripts/cute_dsl_kernels/run_dense_bf16_gemm_persistent.py Outdated
@taianz-nv

Copy link
Copy Markdown
Collaborator Author

/bot kill

@taianz-nv

Copy link
Copy Markdown
Collaborator Author

/bot run

@taianz-nv

Copy link
Copy Markdown
Collaborator Author

/bot kill

…tlass-dsl version skew

The visual_gen FlashAttn4 / parallel attention backends import flash_attn.cute
(which pulls quack); both already guard that optional import but only catch
(ImportError, OSError). A newer nvidia-cutlass-dsl (e.g. 4.6, required by the
NVFP4 conv kernel) removes cute.core.ThrMma, which flash_attn.cute / quack
reference at import time, raising AttributeError and hard-breaking
`import tensorrt_llm` for every code path — including the VAE, which never uses
these backends.

Add AttributeError to the two guards so the optional cute-dsl backends degrade
gracefully (re-imported lazily only when explicitly selected) instead of
breaking the whole package import under a cutlass-dsl version mismatch.

Signed-off-by: Taian Zhang <taianz@nvidia.com>
(cherry picked from commit 87f0ef71a43549e69f8109d621e0fe9ad665da8b)
Signed-off-by: Taian Zhang <taianz@nvidia.com>
Signed-off-by: Taian Zhang <taianz@nvidia.com>
Signed-off-by: Taian Zhang <taianz@nvidia.com>
Signed-off-by: Taian Zhang <taianz@nvidia.com>
Signed-off-by: Taian Zhang <taianz@nvidia.com>
Signed-off-by: Taian Zhang <taianz@nvidia.com>
Signed-off-by: Taian Zhang <taianz@nvidia.com>
Signed-off-by: Taian Zhang <taianz@nvidia.com>
Signed-off-by: Taian Zhang <taianz@nvidia.com>
Signed-off-by: Taian Zhang <taianz@nvidia.com>
@taianz-nv
taianz-nv force-pushed the user/taianz/TRTLLM-13767-cutlass-dsl-4-6-1 branch from f82cb2c to 8be2042 Compare August 7, 2026 05:51
@taianz-nv
taianz-nv requested review from a team as code owners August 7, 2026 05:51
@taianz-nv

Copy link
Copy Markdown
Collaborator Author

/bot run

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #64521 [ run ] triggered by Bot. Commit: 8be2042 Link to invocation

@tensorrt-cicd

Copy link
Copy Markdown
Collaborator

PR_Github #64521 [ run ] completed with state FAILURE. Commit: 8be2042
/LLM/main/L0_MergeRequest_PR pipeline #52389 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

  • Please check the failed tests and fix your PR
  • If you cannot view the failures, ask the CI triggerer to share details
  • Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

8 participants