[None][feat] VisualGen: Add CuTe DSL JIT kernels for FMHA#15831
Conversation
|
/bot run |
📝 WalkthroughWalkthroughReplaces precompiled cubin-based Blackwell FMHA kernels with runtime CuTe DSL JIT compilation and a compile cache, adds block-scaled (MXFP8/NVFP4) Q/K quantization support with new ChangesCuTe DSL Runtime Compilation Migration
Estimated code review effort: 4 (Complex) | ~75 minutes Sequence Diagram(s)sequenceDiagram
participant Attn as CuTeDSLAttention
participant Fwd as cute_dsl_fmha_fwd
participant Cache as _COMPILE_CACHE
participant Kernel as BlackwellFusedMultiHeadAttentionForward
Attn->>Fwd: call with q, k, v, qk_sf_vec, q_sf, k_sf
Fwd->>Fwd: validate NHD inputs, map mask/window args
Fwd->>Cache: lookup compiled kernel by cache key
alt cache miss
Fwd->>Kernel: JIT-compile kernel (dense or block-scaled)
Kernel-->>Cache: store compiled kernel
end
Fwd->>Kernel: launch on current CUDA stream
Kernel-->>Fwd: output tensor, optional lse
Fwd-->>Attn: return output
Suggested reviewers: 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
There was a problem hiding this comment.
Actionable comments posted: 10
🧹 Nitpick comments (2)
tests/unittest/_torch/visual_gen/test_visual_gen_args.py (1)
110-155: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winAdd negative CUTEDSL mismatch coverage.
Coverage in
tests/unittest/_torch/visual_gen/test_visual_gen_args.pyis insufficient for the new validation matrix. Please add CUTEDSL rejection cases for mismatchedqk_dtype/qk_sf_vec, e.g. FP4 withqk_sf_vec=0or32, FP8 withqk_sf_vec=16, and an unsupported vector size like17.As per path instructions,
tests/**reviews should assess coverage and suggest concrete file names; coverage here needs this follow-up in the same test file.🤖 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/unittest/_torch/visual_gen/test_visual_gen_args.py` around lines 110 - 155, Add negative coverage in AttentionConfig tests for CUTEDSL quant validation in the same test module, since the current tests only cover valid fp8/fp4 combinations. Extend the existing quant config cases around test_supported_quant_config_cute_mxfp8, test_supported_quant_config_cute_nvfp4, and test_fp4_qk_dtype_rejected_on_trtllm to assert CUTEDSL rejects mismatched qk_dtype/qk_sf_vec inputs such as fp4 with qk_sf_vec=0 or 32, fp8 with qk_sf_vec=16, and an unsupported qk_sf_vec like 17, using ValidationError with the existing unsupported quant_attention_config-style validation.Source: Path instructions
tensorrt_llm/_torch/visual_gen/attention_backend/cute_dsl.py (1)
543-571: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winAdd return annotations on the new methods.
__init__()and_smooth_qk()are newly added but still unannotated, which goes against the repo’s Python typing rule and makes this backend harder to type-check.Suggested fix
def __init__( self, layer_idx: int = 0, num_heads: int = 8, @@ quant_attention_config: QuantAttentionConfig | None = None, skip_softmax_threshold_scale: float | None = None, **kwargs, - ): + ) -> None: @@ def _smooth_qk( self, q: torch.Tensor, k: torch.Tensor, alpha: float = 0.7, - ): + ) -> tuple[torch.Tensor, torch.Tensor]:As per coding guidelines, “Always annotate functions with return types (use
Noneif no return).”🤖 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 `@tensorrt_llm/_torch/visual_gen/attention_backend/cute_dsl.py` around lines 543 - 571, The new methods in Cute DSL backend are missing return type annotations, which violates the typing guideline. Add an explicit return annotation of None to the __init__ method in cute_dsl.py, and add the appropriate return annotation to _smooth_qk based on its actual return value so the class remains type-checkable. Use the method names __init__ and _smooth_qk to locate the changes.Source: Coding guidelines
🤖 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 `@tensorrt_llm/_torch/visual_gen/attention_backend/cute_dsl.py`:
- Around line 578-585: The smooth scale computation in the attention backend can
produce NaNs when a Q/K channel is entirely zero because `q_max.pow(alpha) /
k_max.pow(1 - alpha)` hits `0/0` and `clamp()` does not recover it. Update the
scale logic in `cute_dsl.py` to guard the `s` calculation in the same block that
uses `q_max` and `k_max`, so zero-valued channels produce a finite fallback
scale instead of NaN. Keep the existing `q` and `k` rescaling path intact, but
ensure `s` is sanitized before `q.reciprocal()` and `k * s` are applied.
- Around line 92-105: Add an explicit divisibility check in cute_dsl_fmha_fwd
before any head-group reshape logic, since the later num_heads_q // num_heads_kv
and q.view(..., num_heads_kv, num_head_groups, ...) assume exact divisibility.
If num_heads_q is not a multiple of num_heads_kv, raise a clear ValueError
alongside the existing tensor validation near the Q/K/V/O shape checks so the
failure is caught early and reported with a contract-specific message.
In `@tensorrt_llm/_torch/visual_gen/cute_dsl_kernels/helpers/fmha_helpers.py`:
- Around line 1019-1028: The backward-inference mask offset in fmha_helpers is
using the wrong sign, so the WINDOW_MASK_BWD_INFERENCE path shifts the attention
window in the opposite direction. Update the offset logic in the helper that
reads `cute.arch.thread_idx()` and computes `offset` so that
`WINDOW_MASK_BWD_INFERENCE` follows the same direction as the other BWD helpers,
using the `seqlen_q` minus `seqlen_k` form while keeping the inference mask
branch behavior otherwise unchanged.
- Around line 1143-1148: In the FP8 conversion branch that sets cvt_instruction,
replace the assert False fallback with an explicit TypeError so unsupported
types cannot pass silently if Python assertions are disabled. Keep the existing
cutlass.const_expr(fp8_type == cutlass.Float8E4M3FN) check, and in the else path
raise a clear error from this helper indicating the FP8 element type is
unsupported.
- Around line 827-831: The trailing-tile check in get_trip_count for
RESIDUAL_MASK_BWD is using the K dimension, which can produce the wrong mask
count. Update the RESIDUAL_MASK_BWD branch in fmha_helpers.py to base the
trailing-tile decision on the Q sequence length and its corresponding tile size
in the get_trip_count logic, so residual backward masking matches Q tiles
instead of seqlen_k.
In
`@tensorrt_llm/_torch/visual_gen/cute_dsl_kernels/helpers/static_persistent_tile_scheduler.py`:
- Around line 541-566: The work-item validity check in
StaticPersistentTileScheduler should also reject CTAs that land in the padded
tail of the final partial cluster, since current_work_linear_idx only validates
against problem_layout_ncluster_mnl. Update the logic in the work-tile
coordinate path that builds cur_cluster_coord and cur_tile_coord so it
additionally verifies the computed tile coordinate is within
problem_shape_ntile_mnl before returning WorkTileInfo. Use the existing
scheduler fields and helpers in this method, such as
problem_layout_ncluster_mnl, cluster_shape_mn, and cur_tile_coord, so tail CTAs
are marked invalid when the problem size is not divisible by the cluster shape.
- Around line 798-809: The runtime tile coordinate in
static_persistent_tile_scheduler.py is missing the computed cta_id_in_cluster,
causing CTAs within the same cluster to map to the same tile. Update the
cur_tile_coord construction in the scheduler logic to include cta_id_in_cluster
alongside cluster_tile_coord_mn, and verify the placement logic in the create()
and runtime coordinate path stays consistent for cluster_shape_mn values other
than (1, 1).
In `@tensorrt_llm/visual_gen/args.py`:
- Around line 61-62: Fix the field description string in args.py so the two
adjacent literals are separated by a space; update the description used for the
Q/K quantization dtype option in the relevant Args/pydantic field so generated
schema/docs render “unquantized. fp4” instead of “unquantized.fp4”.
In `@tests/unittest/_torch/visual_gen/test_attention_cute_dsl.py`:
- Around line 210-215: Add a block-scaled attention test case with a sequence
length that is not a multiple of 128 so the _quantize_blockscaled_one()
pad/unpad path is exercised end to end. Update the parametrized cases in
test_attention_cute_dsl.py (the attention test using batch_size, seq_len_q,
seq_len_kv, is_causal) to include at least one non-128-multiple
seq_len_q/seq_len_kv value, and mirror the same coverage in the related block
around the other referenced lines so the quantized path is actually validated.
- Around line 41-60: The smoke test only verifies the first `cute_dsl_fmha_fwd`
launch populates `_COMPILE_CACHE` once, but it does not confirm identical calls
reuse the same cached kernel. Update `test_cute_dsl_jit_compile_smoke` to invoke
`cute_dsl_fmha_fwd(q, k, v, out, is_causal=False, sm_scale=sm_scale)` a second
time after synchronization and assert `_COMPILE_CACHE` remains at size 1, so the
test covers cache-hit behavior and catches accidental recompilation under a new
key.
---
Nitpick comments:
In `@tensorrt_llm/_torch/visual_gen/attention_backend/cute_dsl.py`:
- Around line 543-571: The new methods in Cute DSL backend are missing return
type annotations, which violates the typing guideline. Add an explicit return
annotation of None to the __init__ method in cute_dsl.py, and add the
appropriate return annotation to _smooth_qk based on its actual return value so
the class remains type-checkable. Use the method names __init__ and _smooth_qk
to locate the changes.
In `@tests/unittest/_torch/visual_gen/test_visual_gen_args.py`:
- Around line 110-155: Add negative coverage in AttentionConfig tests for
CUTEDSL quant validation in the same test module, since the current tests only
cover valid fp8/fp4 combinations. Extend the existing quant config cases around
test_supported_quant_config_cute_mxfp8, test_supported_quant_config_cute_nvfp4,
and test_fp4_qk_dtype_rejected_on_trtllm to assert CUTEDSL rejects mismatched
qk_dtype/qk_sf_vec inputs such as fp4 with qk_sf_vec=0 or 32, fp8 with
qk_sf_vec=16, and an unsupported qk_sf_vec like 17, using ValidationError with
the existing unsupported quant_attention_config-style validation.
🪄 Autofix (Beta)
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: c3ad8c1a-872a-45d0-8076-600b278c749d
⛔ Files ignored due to path filters (64)
tensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_100a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/aarch64/sm_103a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_100a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_causal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_e4m3_bf16_h128_nocausal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_h128_causal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_lse_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_lse_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_skipsm_tvmffi.sois excluded by!**/*.sotensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/cubins/x86_64/sm_103a/cute_dsl_fmha_bf16_h128_nocausal_nonpersistent_varlen_tvmffi.sois excluded by!**/*.so
📒 Files selected for processing (16)
.gitattributes.gitignoresetup.pytensorrt_llm/_torch/visual_gen/attention_backend/cute_dsl.pytensorrt_llm/_torch/visual_gen/attention_backend/utils.pytensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/__init__.pytensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/fmha.pytensorrt_llm/_torch/visual_gen/cute_dsl_kernels/blackwell/attention/fmha_blockscaled.pytensorrt_llm/_torch/visual_gen/cute_dsl_kernels/helpers/__init__.pytensorrt_llm/_torch/visual_gen/cute_dsl_kernels/helpers/fmha_helpers.pytensorrt_llm/_torch/visual_gen/cute_dsl_kernels/helpers/static_persistent_tile_scheduler.pytensorrt_llm/visual_gen/args.pytests/unittest/_torch/visual_gen/test_attention_cute_dsl.pytests/unittest/_torch/visual_gen/test_attention_integration.pytests/unittest/_torch/visual_gen/test_attention_perf.pytests/unittest/_torch/visual_gen/test_visual_gen_args.py
💤 Files with no reviewable changes (3)
- .gitattributes
- .gitignore
- setup.py
|
PR_Github #56918 [ run ] triggered by Bot. Commit: |
9eb7623 to
a68cec4
Compare
|
PR_Github #56918 [ run ] completed with state
|
6e52ac7 to
ad46d3e
Compare
|
/bot run |
|
PR_Github #57200 [ run ] triggered by Bot. Commit: |
|
PR_Github #57200 [ run ] completed with state
|
|
/bot run |
|
PR_Github #57462 [ run ] triggered by Bot. Commit: |
|
PR_Github #57462 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #57479 [ run ] triggered by Bot. Commit: |
|
PR_Github #57479 [ run ] completed with state
|
|
/bot run |
|
PR_Github #57509 [ run ] triggered by Bot. Commit: |
|
/bot help |
GitHub Bot Help
Provide a user friendly way for developers to interact with a Jenkins server. Run See details below for each supported subcommand. Details
Launch build/test pipelines. All previously running jobs will be killed.
kill
Kill all running builds associated with pull request. skip
Skip testing for latest commit on pull request. reuse-pipeline
Reuse a previous pipeline to validate current commit. This action will also kill all currently running builds associated with the pull request. IMPORTANT NOTE: This is dangerous since lack of user care and validation can cause top of tree to break. |
5a11ff5 to
8fc990e
Compare
|
Thanks so much, Bowen ! Right now I have no causal xattn use case on my immediate list, but will check alignment whenever a use case comes up. Some future possibilities (currently out of scope):
Thanks for pointing out. Currently VisualGen's attention backends are tested through SDPA-reference + single-block runs + offline video quality checks. IIRC there are some ongoing plans to add e2e video-gen to the tests but still iterating on the time/coverage trade-off. |
Thanks for the comment, @pengbowang-nv ! I tried to rerun To make it 100% safe that the previous failure is only caused by absence of rebase, I'll re-trigger the test to see if |
|
/bot run |
|
PR_Github #59656 [ run ] triggered by Bot. Commit: |
|
PR_Github #59656 [ run ] completed with state
|
|
Hey. I've found the reason that caused skipSm's behavior mismatch against trtllmGen kernels. We need to fix a default value related to a CuTeDSL-specific optimization that allows skipping This fix is not kernel-side but unfortunately the value is set internally by host side of the kernel module. I'll raise a MR to our internal variant as well, but for TRTLLM here, please allow for one more CI loop as I'm about to push new changes. Thanks |
|
/bot run --disable-fail-fast |
|
PR_Github #59719 [ run ] triggered by Bot. Commit: |
|
PR_Github #59719 [ run ] completed with state
|
Replace the packaged Blackwell CuTe DSL FMHA binaries with runtime JIT compilation and port the forward kernels and scheduling helpers into VisualGen using public CuTe DSL interfaces. Add dense QK16PV8 and block-scaled MXFP8/NVFP4 attention paths, including quantized Q/K scale handling and QK smoothing. Extend the VisualGen attention configuration with validated backend-specific recipes while keeping internal kernel controls private. Update kernel integration, packaging metadata, and unit and performance coverage for the JIT and quantized execution paths. Signed-off-by: Ruqing Xu <7891482+xrq-phys@users.noreply.github.com>
Co-authored-by: Zhenhua Wang <4936589+zhenhuaw-me@users.noreply.github.com> Signed-off-by: RuQing Xu <7891482+xrq-phys@users.noreply.github.com>
Config: Use qk_dtype=mxfp8 instead of qk_dtype=fp8 + reusing sage block sizes for DSL MXFP8 Signed-off-by: Ruqing Xu <7891482+xrq-phys@users.noreply.github.com>
Signed-off-by: RuQing Xu <7891482+xrq-phys@users.noreply.github.com>
There is a default-to-on optimization specific to CuTeDSL: usually skip-correction only happens when new_row_max < old_row_max , but CuTeDSL fmha.py has an optimization to allow skip updating row_max whenever new_row_max < 8.0 * old_row_max. Usually this improves the performance as it reduces sm-correction frequency, but the same threshold is interacting with skipSoftmax threshold in a way that it sticks row_max to a lower value, making local_max/row_max < skip_sm_threshold harder to meet. This commit resets this 8.0 threshold to 1.0 whenever skip-softmax is enabled. Signed-off-by: Ruqing Xu <7891482+xrq-phys@users.noreply.github.com>
f6421b8 to
534cdfe
Compare
|
/bot run --disable-fail-fast |
|
PR_Github #60246 [ run ] triggered by Bot. Commit: |
|
PR_Github #60246 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #60311 [ run ] triggered by Bot. Commit: |
|
PR_Github #60311 [ run ] completed with state
|
|
/bot skip --comment "Error not related to VisualGen changes. Multi-GPU tests have all previously succeeded." |
|
PR_Github #60358 [ skip ] triggered by Bot. Commit: |
|
PR_Github #60358 [ skip ] completed with state |
|
@zhenhuaw-me I'm feeling confident about merging this now. |
Summary by CodeRabbit
New Features
Bug Fixes
Description
Replace the packaged Blackwell CuTe DSL FMHA binaries with runtime JIT compilation and port the forward kernels and scheduling helpers into VisualGen using public CuTe DSL interfaces.
Add dense QK16PV8 and block-scaled MXFP8/NVFP4 attention paths, including quantized Q/K scale handling and QK smoothing. Extend the VisualGen attention configuration with validated backend-specific recipes while keeping internal kernel controls private.
Update kernel integration, packaging metadata, and unit and performance coverage for the JIT and quantized execution paths.
Performance (B300, TFLOPS, higher is better)
Above results use CuTe DSL 4.5.0, as pinned by
requirements.txt.Measured from CUPTI device-kernel timestamps collected by PyTorch Profiler/Kineto. Compilation, quantization, allocation, and model execution are excluded.
if CuTe DSL version was upgraded from 4.5.0 to 4.6.0
Test Coverage
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-compatibleorapi-breaking. Forapi-breaking, includeBREAKINGin 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.