[https://nvbugs/6595612][fix] Use CuTe DSL string rounding modes in VSA

[JiaxuanBai]

Dev Engineer Review

• Updated packed FP32 subtraction and addition helpers to pass "rn" and "rm" string literals to cute.arch.
• Preserved existing rounding behavior.
• The change supports CuTe DSL 4.4.0 and newer.
• py_compile and git diff --check passed.
• No configuration, test-list, or public API changes were made.

QA Engineer Review

No test changes.

Description

The Video Sparse Attention CuTe DSL kernel passes nvvm.RoundingModeKind values to public cute.arch wrappers. CuTe DSL 4.4.0 and later require string literals for these arguments, and newer generated NVVM bindings no longer export RoundingModeKind. This causes TensorRT-LLM import and test collection to fail with recent nvidia-cutlass-dsl releases.

Replace RoundingModeKind.RN and RoundingModeKind.RM with the supported "rn" and "rm" literals. This preserves the existing rounding behavior. It does not change the public API, numerical behavior, performance, dependencies, ownership, documentation, or architecture.

Test Coverage

• python3 -m py_compile tensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/video_sparse_attention/ptx.py — passed.
• git diff --check — passed.
• Existing coverage: tests/unittest/_torch/visual_gen/test_attention_cute_dsl_vsa.py::test_cute_kernel_matches_dense_at_full_topk. This was not run locally because the host does not have pytest or a compatible CUDA GPU.

PR Checklist

Please review the following before submitting your PR:

• PR description clearly explains what and why. If using CodeRabbit's summary, please make sure it makes sense.

• PR Follows TRT-LLM CODING GUIDELINES to the best of your knowledge.

• Test cases are provided for new code paths (see test instructions)

• If PR introduces API changes, an appropriate PR label is added - either api-compatible or api-breaking. For api-breaking, include BREAKING in the PR title.

• Any new dependencies have been scanned for license and vulnerabilities

• CODEOWNERS updated if ownership changes

• Documentation updated as needed

• Update tava architecture diagram if there is a significant design change in PR.

• The reviewers assigned automatically/manually are appropriate for the PR.

• Please check this after reviewing the above items as appropriate for this PR.

GitHub Bot Help

To see a list of available CI bot commands, please comment /bot help.

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 02983716-c1c7-4bd8-9e71-014c143ac4d7

📥 Commits

Reviewing files that changed from the base of the PR and between f274b6c and 24d5054.

📒 Files selected for processing (1)

• tensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/video_sparse_attention/ptx.py

---

Walkthrough

The change replaces NVVM rounding-mode enum constants with string values in two packed FP32 operations: round-nearest subtraction and round-down addition.

Changes

Rounding mode argument updates

Layer / File(s)	Summary
Packed FP32 rounding modes tensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/video_sparse_attention/ptx.py	sub_packed_f32x2 now uses "rn". exp2_emulation_2 now uses "rm" for packed addition.

Estimated code review effort: 1 (Trivial) | ~2 minutes

Mergeability Score: ⚪ Minimal · up to 24d50

This localized change preserves existing rounding behavior while restoring compatibility with newer CuTe DSL releases; no actionable merge-blocking risk remains beyond normal checks and review.

Suggested reviewers: bowenfu

🚥 Pre-merge checks | ✅ 5

✅ Passed checks (5 passed)

Check name	Status	Explanation
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.
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 follows the required format and clearly identifies the fix to use CuTe DSL string rounding modes in VSA.
Description check	✅ Passed	The description explains the issue, solution, test coverage, limitations, and checklist items relevant to this change.

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

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