Skip to content

feat: support FP8 KV NoPE MLA on SM90 - #4373

Merged
leejnau merged 9 commits into
flashinfer-ai:mainfrom
JustinTong0323:feat/native-nope-mla-head-dim-kpe-0
Aug 14, 2026
Merged

leejnau merged 9 commits into
flashinfer-ai:mainfrom
JustinTong0323:feat/native-nope-mla-head-dim-kpe-0

Conversation

@JustinTong0323

@JustinTong0323 JustinTong0323 commented Aug 6, 2026

Copy link
Copy Markdown
Contributor

Description

Allow BatchMLAPagedAttentionWrapper to run with head_dim_kpe=0, including FP8 KV caches on SM90.

When the positional-key-embedding dimension is zero, the KPE MMA loop is empty. The FA2 and FA3 paths initialize the QK accumulator from compressed KV instead. For FP8 KV, both native backends repack to BF16 and can consume an optional contiguous FP32 scale tensor with one scale per physical KV token and 128 CKV channels.

This change:

  • supports zero-width KPE in FA2 and FA3;
  • supports FP8 E4M3 KV with BF16 queries on SM90 for head_dim_ckv=512 and head_dim_kpe in {0, 64};
  • adds optional ckv_scale_arr with shape ckv_cache.shape[:-1] + (4,) for per-token, per-group CKV dequantization;
  • preserves existing per-tensor ckv_scale / kpe_scale behavior;
  • validates backend, device, dtype, dimensions, scale shape, dtype, device, and contiguity;
  • adds paged-MLA correctness coverage with a non-identity page table for FA2 and FA3.

Related Issues

None.

Pull Request Checklist

  • Pre-commit hooks pass on the changed files.
  • Tests have been added or updated as needed.
  • GPU tests pass.

Tests

NVIDIA H200, CUDA 12.8, PyTorch 2.11.0+cu128:

  • test_batch_mla_without_kpe: FA2 and FA3 passed.
  • test_batch_mla_fp8_nope_group_scales_matches_bf16_reference: FA2 and FA3 passed.
  • test_fp8_kv_scales_are_keyword_only and test_fp8_kv_requires_scales: passed.
  • Representative existing KPE=64 per-tensor FP8 and row-aliasing regressions: passed.

Reviewer Notes

  • ckv_scale_arr follows the physical paged-cache layout used by the downstream SGLang integration: [num_pages, page_size, head_dim_ckv / 128].
  • Existing BF16/FP16 and FP8 per-tensor scale behavior is unchanged.
  • The latest external attention CI failures before this update were not stable across reruns: HCA failures were classified as pre-existing, while the new failures alternated between GB200 multi-node NCCL setup and RTX Pro 6000 cudaErrorUnknown failures.

Summary by CodeRabbit

  • New Features

    • Added FP8 batched paged attention with optional per-token/per-channel CKV scaling.
    • Added support for FA2 and FA3 FP8 execution.
    • Expanded supported configurations to include zero- or 64-dimensional key positional encoding.
  • Bug Fixes

    • Fixed attention computation without key positional encoding dimensions.
    • Added validation for incompatible scale inputs, shapes, layouts, and configurations.
  • Tests

    • Added coverage for grouped CKV scaling and trace inputs.
    • Added BF16 reference comparisons across supported implementations.

@coderabbitai

coderabbitai Bot commented Aug 6, 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

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info
⚙️ Run configuration

Configuration used: defaults

Review profile: CHILL

Plan: Pro Plus

Run ID: 96a9f21e-a089-4cc1-867a-4487e4ff43d2

📥 Commits

Reviewing files that changed from the base of the PR and between 9219186 and b586789.

📒 Files selected for processing (5)
  • flashinfer/trace/templates/attention.py
  • tests/trace/fi_trace_out/mla_paged_decode_h16_ckv512_kpe64_ps1.json
  • tests/trace/fi_trace_out/mla_paged_decode_h16_ckv512_kpe64_ps64.json
  • tests/trace/test_fi_trace.py
  • tests/trace/test_mla_paged_decode_reference_correctness.py
🚧 Files skipped from review as they are similar to previous changes (1)
  • tests/trace/test_fi_trace.py

📝 Walkthrough

Walkthrough

MLA now supports optional per-token or per-group FP8 CKV scales and zero-dimensional KPE layouts. The wrapper validates and forwards scale arrays. CUDA bindings and kernels apply the scales during FP8 repacking. Trace schemas and tests cover FA2 and FA3 paths.

Changes

FP8 CKV scaling and zero-KPE MLA

Layer / File(s) Summary
CKV scale-array API wiring
flashinfer/mla/_core.py, csrc/batch_mla_*, include/flashinfer/attention/mla_params.cuh
The wrapper accepts, validates, and forwards ckv_scale_arr. CUDA bindings store its pointer in MLAParams. CUTLASS rejects unsupported scale arguments.
FP8 KV repacking and computation
include/flashinfer/attention/mla.cuh, include/flashinfer/attention/mla_hopper.cuh
FP8 KV data is repacked to BF16. CKV dequantization uses scalar or per-token/group scale values. QK and PV paths use the repacked storage.
Zero-KPE execution
flashinfer/mla/_core.py, include/flashinfer/attention/mla_hopper.cuh
Planning and Hopper trait validation accept head_dim_kpe=0. The kernels allocate valid staging storage and initialize CKV scores without KPE MMA operations.
Trace schema and regression coverage
flashinfer/trace/templates/attention.py, tests/trace/*, tests/attention/test_deepseek_mla.py
Trace definitions expose ckv_scale_groups and ckv_scale_arr. Tests cover FA2 and FA3 execution, grouped FP8 scales, BF16 references, keyword-only arguments, and invalid shapes.

Estimated code review effort: 4 (Complex) | ~60 minutes

Mergeability Score: ⚪ Minimal · up to b5867

The change adds FP8 KV and zero-width positional-key support with validation and targeted GPU coverage; no actionable merge-blocking risk remains beyond normal checks and review.

Sequence Diagram(s)

sequenceDiagram
  participant MLAWrapper
  participant CUDABinding
  participant MLAParams
  participant FP8MLAKernel
  MLAWrapper->>CUDABinding: pass validated ckv_scale_arr
  CUDABinding->>MLAParams: store scale-array pointer
  MLAParams->>FP8MLAKernel: provide CKV scale metadata
  FP8MLAKernel->>FP8MLAKernel: repack FP8 CKV to BF16
Loading

Suggested labels: run-ci

Suggested reviewers: saltyminty, qsang-nv, yyihuang

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 33.33% 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
Title check ✅ Passed The title clearly summarizes the main change: FP8 KV NoPE MLA support on SM90.
Description check ✅ Passed The description covers the change, related issues, checks, tests, and reviewer notes, with clear validation and CI details.
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.
✨ Finishing Touches 💡 1
🛠️ Fix failing CI checks 💡
  • Create stacked PR
  • Commit on current branch
🧪 Generate unit tests (beta)
  • Create PR with unit tests

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

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

@JustinTong0323
JustinTong0323 marked this pull request as ready for review August 6, 2026 15:30
@flashinfer-bot

Copy link
Copy Markdown
Collaborator

[FAILED] Pipeline #62257276 — 10/18 executed test jobs passed

Compared with nightly #62109159.

Unit Tests

GPU CUDA 12.9 CUDA 13.0 Notes
5090 ✅ Pass ✅ Pass
B300 🟡 Old 🟡 Old Old: tests.attention.test_cute_dsl_hca_dsv4 (10 failures; CUDA 12.9, CUDA 13.0)
GB200 🟡 Old 🟡 Old Old: tests.attention.test_cute_dsl_hca_dsv4 (10 failures; CUDA 12.9, CUDA 13.0)
GB300 🟡 Old 🟡 Old Old: tests.attention.test_cute_dsl_hca_dsv4 (10 failures; CUDA 12.9, CUDA 13.0)
H100 ✅ Pass ✅ Pass
RTX Pro 6000 Blackwell ❌ New ❌ New New: tests.attention.test_attention_sink (765 failures; CUDA 12.9, CUDA 13.0)

✅ Pass · 🟡 Old failure · ❌ New failure · ⏱ Test timeout · ⚠️ Infrastructure · ❔ Unknown or unclassified · — Not run

Multi-GPU and Multi-Node Tests — 6/6 passed

GPU CUDA 12.9 CUDA 13.0 Notes
B300 (multi-GPU) ✅ Pass ✅ Pass
GB200 (multi-node) ✅ Pass ✅ Pass
GB300 (multi-node) ✅ Pass ✅ Pass
Failure details

New relative to nightly (attribution uncertain)

Pre-existing failures

  • tests.attention.test_cute_dsl_hca_dsv4 — 30 failures on B300 / CUDA 12.9, B300 / CUDA 13.0, GB200 / CUDA 12.9, GB200 / CUDA 13.0, GB300 / CUDA 12.9, GB300 / CUDA 13.0
    • AttributeError: module 'cutlass.cute.nvgpu.cpasync' has no attribute 'CopyBulkTensor2DGather4G2SOp'. Did you mean: 'CopyBulkTensorTileG2SOp'?

@JustinTong0323 JustinTong0323 changed the title feat: support zero KPE dimension in native MLA feat: support FP8 KV NoPE MLA on SM90 Aug 12, 2026
@JustinTong0323

Copy link
Copy Markdown
Contributor Author

I added the requested SM90 FP8 KV NoPE support, including optional per-token/per-128-channel CKV scales for both FA2 and FA3, and rebased the branch onto the latest main. On an H200 with CUDA 12.8 / PyTorch 2.11.0+cu128, the new non-identity-page-table group-scale test passes on both FA2 and FA3; the existing KPE=64 per-tensor FP8 and row-aliasing regressions also pass. The current required check is red only because CI was skipped pending authorization for the new commit (b72a3c14); could you please run @flashinfer-bot run (and /bot run tests/attention if needed)? The prior external failures were not stable across reruns: HCA was classified as old, while the new failures alternated between GB200 NCCL setup and RTX Pro 6000 cudaErrorUnknown, so they do not appear attributable to this MLA change.

@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: 2

🧹 Nitpick comments (3)
flashinfer/mla/_core.py (1)

2523-2523: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

Use tuple unpacking instead of tuple concatenation.

Ruff reports RUF005 on this line.

♻️ Proposed change
-            expected_scale_shape = ckv_cache.shape[:-1] + (self._head_dim_ckv // 128,)
+            expected_scale_shape = (*ckv_cache.shape[:-1], self._head_dim_ckv // 128)
🤖 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 `@flashinfer/mla/_core.py` at line 2523, Update the expected_scale_shape
assignment to construct the shape tuple using tuple unpacking instead of
concatenating ckv_cache.shape[:-1] with a one-element tuple, while preserving
the same resulting dimensions.

Source: Linters/SAST tools

include/flashinfer/attention/mla_hopper.cuh (1)

392-409: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

Merge the duplicated scale-array branch.

Lines 393 and 399 test ckv_scale_arr twice in a row. Fold the lookup and the scaled conversion into one branch. Add a short comment that states why the scale-array path uses scalar FP32 multiplication instead of __hmul2.

🤖 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 `@include/flashinfer/attention/mla_hopper.cuh` around lines 392 - 409, The
scale-array handling in the surrounding attention conversion code duplicates the
ckv_scale_arr branch. Merge the scale lookup and conversion loop into a single
ckv_scale_arr branch, retaining scalar FP32 multiplication there and the
existing packed __hmul2 path in the else branch. Add a brief comment explaining
why the scale-array path requires scalar FP32 multiplication.
include/flashinfer/attention/mla.cuh (1)

296-348: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Duplicated CKV scale-lookup branch in repack_fp8_kv_to_bf16. Both kernels copy the same block: a first if (ckv_scale_arr != nullptr) that resolves the physical row and reads the group scale, then a second identical test that selects scalar FP32 multiplication or packed __hmul2. The shared root cause is the copied scale-lookup logic. Extract one helper that returns the CKV scale for a (row, col) pair and applies it, then call it from both kernels.

  • include/flashinfer/attention/mla.cuh#L296-L348: merge the two ckv_scale_arr != nullptr tests at Lines 323 and 331 into one branch, and add the hot-path rationale comment for the FP8-to-BF16 staging and the scalar-scale path.
  • include/flashinfer/attention/mla_hopper.cuh#L392-L409: merge the two ckv_scale_arr tests at Lines 393 and 399 into the same shared helper.
🤖 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 `@include/flashinfer/attention/mla.cuh` around lines 296 - 348, Extract the
duplicated CKV scale lookup and application logic into one shared helper for
`(row, col)` pairs, preserving paged `physical_row` resolution and both scalar
and packed multiplication paths; invoke it from `repack_fp8_kv_to_bf16` in
include/flashinfer/attention/mla.cuh:296-348, merging the two `ckv_scale_arr`
branches and adding the requested hot-path rationale comment. Apply the same
helper and branch consolidation in
include/flashinfer/attention/mla_hopper.cuh:392-409.
🤖 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 `@flashinfer/mla/_core.py`:
- Around line 2518-2532: Add ckv_scale_arr as an optional tensor in
mla_paged_decode_trace, using the existing shape, dtype, device, and contiguity
validation behavior. Update the trace template and regenerate both generated
fixtures so they include this optional tensor and preserve the Invalid shape of
ckv_scale_arr validation prefix.

In `@include/flashinfer/attention/mla.cuh`:
- Around line 350-373: The FP8 KPE repacking path in
load_kv/repack_fp8_kv_to_bf16 uses the colliding standard KPE swizzle for FA2
head_dim_kpe=64. Mirror SWIZZLE_MODE_KPE_RAW, LANES_PER_ROW_KPE, and
INNER_LOADS_KPE from mla_hopper.cuh, use the raw swizzle when reading KPE, and
add FA2 coverage for head_dim_kpe=64.

---

Nitpick comments:
In `@flashinfer/mla/_core.py`:
- Line 2523: Update the expected_scale_shape assignment to construct the shape
tuple using tuple unpacking instead of concatenating ckv_cache.shape[:-1] with a
one-element tuple, while preserving the same resulting dimensions.

In `@include/flashinfer/attention/mla_hopper.cuh`:
- Around line 392-409: The scale-array handling in the surrounding attention
conversion code duplicates the ckv_scale_arr branch. Merge the scale lookup and
conversion loop into a single ckv_scale_arr branch, retaining scalar FP32
multiplication there and the existing packed __hmul2 path in the else branch.
Add a brief comment explaining why the scale-array path requires scalar FP32
multiplication.

In `@include/flashinfer/attention/mla.cuh`:
- Around line 296-348: Extract the duplicated CKV scale lookup and application
logic into one shared helper for `(row, col)` pairs, preserving paged
`physical_row` resolution and both scalar and packed multiplication paths;
invoke it from `repack_fp8_kv_to_bf16` in
include/flashinfer/attention/mla.cuh:296-348, merging the two `ckv_scale_arr`
branches and adding the requested hot-path rationale comment. Apply the same
helper and branch consolidation in
include/flashinfer/attention/mla_hopper.cuh:392-409.
🪄 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: defaults

Review profile: CHILL

Plan: Pro Plus

Run ID: 20e6738e-cbe3-4e92-982d-c77a728082f5

📥 Commits

Reviewing files that changed from the base of the PR and between 3422059 and b72a3c1.

📒 Files selected for processing (9)
  • csrc/batch_mla_binding.cu
  • csrc/batch_mla_run.cu
  • csrc/batch_mla_sm90_binding.cu
  • csrc/batch_mla_sm90_run.cu
  • flashinfer/mla/_core.py
  • include/flashinfer/attention/mla.cuh
  • include/flashinfer/attention/mla_hopper.cuh
  • include/flashinfer/attention/mla_params.cuh
  • tests/attention/test_deepseek_mla.py

Comment thread flashinfer/mla/_core.py
Comment thread include/flashinfer/attention/mla.cuh
@leejnau

leejnau commented Aug 12, 2026

Copy link
Copy Markdown
Contributor

/bot run tests/attention

@flashinfer-bot

Copy link
Copy Markdown
Collaborator

GitLab MR !1155 has been updated with latest changes, and the CI pipeline #62355852 is currently running. I'll report back once the pipeline job completes.

@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.

Caution

Some comments are outside the diff and can’t be posted inline due to platform limitations.

⚠️ Outside diff range comments (1)
flashinfer/trace/templates/attention.py (1)

1428-1437: 🗄️ Data Integrity & Integration | 🟡 Minor | ⚡ Quick win

Add FP8 scale support to mla_paged_decode init and reference functions.

The trace schema declares optional FP8 inputs (ckv_scale_arr, ckv_scale, kpe_scale), but _mla_paged_decode_init discards ckv_scale_groups without constructing scale tensors, and _mla_paged_decode_reference does not accept these scale parameters. If trace replay forwards these optional inputs, the reference invocation fails. Either implement FP8 dequantization in the reference and scale tensor construction in the init, or move FP8 variants to a separate trace template and mark this template for non-quantized KV only.

🤖 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 `@flashinfer/trace/templates/attention.py` around lines 1428 - 1437, The MLA
paged decode trace template must handle the schema’s optional FP8 inputs
consistently. Update _mla_paged_decode_init to construct the required scale
tensors from ckv_scale_groups, and extend _mla_paged_decode_reference and its
invocation to accept and apply ckv_scale_arr, ckv_scale, and kpe_scale;
alternatively, restrict this template to non-quantized KV and route FP8 variants
through a separate template.
🧹 Nitpick comments (1)
tests/trace/test_fi_trace.py (1)

683-703: 🗄️ Data Integrity & Integration | 🔵 Trivial | ⚡ Quick win

Assert the complete ckv_scale_arr contract.

The test checks only the symbolic shape. It does not check that the emitted input remains optional and uses float32. Add assertions for defn["inputs"]["ckv_scale_arr"]["optional"] is True and defn["inputs"]["ckv_scale_arr"]["dtype"] == "float32".

🤖 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/trace/test_fi_trace.py` around lines 683 - 703, Update the assertions
for ckv_scale_arr in the BatchMLAPagedAttentionWrapper trace test to also
require optional == True and dtype == "float32", while preserving the existing
symbolic shape assertion.
🤖 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.

Outside diff comments:
In `@flashinfer/trace/templates/attention.py`:
- Around line 1428-1437: The MLA paged decode trace template must handle the
schema’s optional FP8 inputs consistently. Update _mla_paged_decode_init to
construct the required scale tensors from ckv_scale_groups, and extend
_mla_paged_decode_reference and its invocation to accept and apply
ckv_scale_arr, ckv_scale, and kpe_scale; alternatively, restrict this template
to non-quantized KV and route FP8 variants through a separate template.

---

Nitpick comments:
In `@tests/trace/test_fi_trace.py`:
- Around line 683-703: Update the assertions for ckv_scale_arr in the
BatchMLAPagedAttentionWrapper trace test to also require optional == True and
dtype == "float32", while preserving the existing symbolic shape assertion.

ℹ️ Review info
⚙️ Run configuration

Configuration used: defaults

Review profile: CHILL

Plan: Pro Plus

Run ID: 3ca1ce68-c82a-48d7-b479-312419b21510

📥 Commits

Reviewing files that changed from the base of the PR and between b72a3c1 and 9219186.

📒 Files selected for processing (8)
  • flashinfer/mla/_core.py
  • flashinfer/trace/templates/attention.py
  • include/flashinfer/attention/mla.cuh
  • include/flashinfer/attention/mla_hopper.cuh
  • tests/attention/test_deepseek_mla.py
  • tests/trace/fi_trace_out/mla_paged_decode_h16_ckv512_kpe64_ps1.json
  • tests/trace/fi_trace_out/mla_paged_decode_h16_ckv512_kpe64_ps64.json
  • tests/trace/test_fi_trace.py
🚧 Files skipped from review as they are similar to previous changes (4)
  • flashinfer/mla/_core.py
  • tests/attention/test_deepseek_mla.py
  • include/flashinfer/attention/mla_hopper.cuh
  • include/flashinfer/attention/mla.cuh

@leejnau

leejnau commented Aug 12, 2026

Copy link
Copy Markdown
Contributor

/bot run tests/attention

@flashinfer-bot

Copy link
Copy Markdown
Collaborator

GitLab MR !1155 has been updated with latest changes, and the CI pipeline #62369148 is currently running. I'll report back once the pipeline job completes.

@flashinfer-bot

Copy link
Copy Markdown
Collaborator

[FAILED] Pipeline #62369148 — 12/18 executed test jobs passed

Compared with nightly #62109159.

Unit Tests

GPU CUDA 12.9 CUDA 13.0 Notes
5090 ✅ Pass ✅ Pass
B300 🟡 Old 🟡 Old Old: tests.attention.test_cute_dsl_hca_dsv4 (10 failures; CUDA 12.9, CUDA 13.0)
GB200 🟡 Old 🟡 Old Old: tests.attention.test_cute_dsl_hca_dsv4 (10 failures; CUDA 12.9, CUDA 13.0)
GB300 🟡 Old 🟡 Old Old: tests.attention.test_cute_dsl_hca_dsv4 (10 failures; CUDA 12.9, CUDA 13.0)
H100 ✅ Pass ✅ Pass
RTX Pro 6000 Blackwell ✅ Pass ✅ Pass

✅ Pass · 🟡 Old failure · ❌ New failure · ⏱ Test timeout · ⚠️ Infrastructure · ❔ Unknown or unclassified · — Not run

Multi-GPU and Multi-Node Tests — 6/6 passed

GPU CUDA 12.9 CUDA 13.0 Notes
B300 (multi-GPU) ✅ Pass ✅ Pass
GB200 (multi-node) ✅ Pass ✅ Pass
GB300 (multi-node) ✅ Pass ✅ Pass
Failure details

Pre-existing failures

  • tests.attention.test_cute_dsl_hca_dsv4 — 30 failures on B300 / CUDA 12.9, B300 / CUDA 13.0, GB200 / CUDA 12.9, GB200 / CUDA 13.0, GB300 / CUDA 12.9, GB300 / CUDA 13.0
    • AttributeError: module 'cutlass.cute.nvgpu.cpasync' has no attribute 'CopyBulkTensor2DGather4G2SOp'. Did you mean: 'CopyBulkTensorTileG2SOp'?

@leejnau

leejnau commented Aug 13, 2026

Copy link
Copy Markdown
Contributor

/bot run tests/attention

@flashinfer-bot

Copy link
Copy Markdown
Collaborator

GitLab MR !1155 has been updated with latest changes, and the CI pipeline #62545365 is currently running. I'll report back once the pipeline job completes.

@saltyminty saltyminty 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.

Left one non-blocking comment.

Comment thread flashinfer/mla/_core.py
*,
ckv_scale: Optional[float] = None,
kpe_scale: Optional[float] = None,
ckv_scale_arr: Optional[torch.Tensor] = None,

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.

Non-blocking but just want to hear thoughts: the intention is that we will accept only one of either ckv_scale or ckv_scale_arr, correct?

  • Since these arguments are already keyword only, we are free to reorder them (i.e. putting ckv_scale_arr directly after ckv_scale)
  • would it be cleaner to only keep one argument as ckv_scale: Union[float, torch.Tensor]?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Yes, they are intended to be mutually exclusive. I reordered ckv_scale_arr directly after ckv_scale and now require exactly one of them for FP8 KV.

I kept them as separate arguments because the trace schema represents scalar and tensor inputs with different descriptors. A single Union[float, torch.Tensor] parameter cannot faithfully describe both forms in the generated benchmark definition. Internally, the grouped-scale path still passes 1.0 through the unused scalar launcher slot, so the C++ ABI remains unchanged.

@flashinfer-bot

Copy link
Copy Markdown
Collaborator

[SUCCESS] Pipeline #62545365: 18/18 executed test jobs passed

@github-actions

Copy link
Copy Markdown
Contributor

🚨 POTENTIAL BREAKING PUBLIC API CHANGE DETECTED 🚨

Caution

THIS PR APPEARS TO BREAK THE PUBLIC API. AUTHORS AND REVIEWERS: DO NOT MISS THIS.

This is an advisory warning and does not gate merging. Confirm compatibility and provide a deprecation or migration path, or track the fix in a follow-up PR.

1 public API finding(s):

  • flashinfer/mla/_core.py:2308 — Public API flashinfer.mla._core.BatchMLAPagedAttentionWrapper.run signature changed in a potentially breaking way. Before: `def run(self, q_nope: torch.Tensor, q_pe: torch.Tensor, ckv_cache: torch.Tensor, kpe_cache: torch.Tensor, out: Optional[torch.Tensor]=None, lse: Optional[torch.Tensor]=None, return_lse: bool=False, profiler_buffer: Optional[torch.Tensor]=None, kv_len: Optional[torch.Tensor]=None, page_table: Optional[torch.Tensor]=None, return_lse_base_on_e: bool=False, o_scale: Optional[

View the full check run

@leejnau

leejnau commented Aug 14, 2026

Copy link
Copy Markdown
Contributor

/bot run tests/attention

@flashinfer-bot

Copy link
Copy Markdown
Collaborator

GitLab MR !1155 has been updated with latest changes, and the CI pipeline #62737821 is currently running. I'll report back once the pipeline job completes.

@JustinTong0323

Copy link
Copy Markdown
Contributor Author

Regarding the public API advisory above: this change is backward-compatible. ckv_scale_arr is an optional keyword-only parameter defaulting to None. Existing per-tensor callers continue to use ckv_scale=... and kpe_scale=... unchanged, while grouped-scale callers explicitly opt in with ckv_scale_arr=.... No existing parameter is removed, renamed, or narrowed, so no deprecation or migration path is needed.

@leejnau

leejnau commented Aug 14, 2026

Copy link
Copy Markdown
Contributor

@flashinfer-bot run

@saltyminty saltyminty 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.

Supposed API breaking change is in a keyword only section and is not an issue.

@leejnau
leejnau merged commit 38ebb15 into flashinfer-ai:main Aug 14, 2026
30 of 31 checks passed
jefby pushed a commit to jefby/flashinfer that referenced this pull request Aug 19, 2026
## Description

Allow `BatchMLAPagedAttentionWrapper` to run with `head_dim_kpe=0`,
including FP8 KV caches on SM90.

When the positional-key-embedding dimension is zero, the KPE MMA loop is
empty. The FA2 and FA3 paths initialize the QK accumulator from
compressed KV instead. For FP8 KV, both native backends repack to BF16
and can consume an optional contiguous FP32 scale tensor with one scale
per physical KV token and 128 CKV channels.

This change:

- supports zero-width KPE in FA2 and FA3;
- supports FP8 E4M3 KV with BF16 queries on SM90 for `head_dim_ckv=512`
and `head_dim_kpe` in `{0, 64}`;
- adds optional `ckv_scale_arr` with shape `ckv_cache.shape[:-1] + (4,)`
for per-token, per-group CKV dequantization;
- preserves existing per-tensor `ckv_scale` / `kpe_scale` behavior;
- validates backend, device, dtype, dimensions, scale shape, dtype,
device, and contiguity;
- adds paged-MLA correctness coverage with a non-identity page table for
FA2 and FA3.

## Related Issues

None.

## Pull Request Checklist

- [x] Pre-commit hooks pass on the changed files.
- [x] Tests have been added or updated as needed.
- [x] GPU tests pass.

## Tests

NVIDIA H200, CUDA 12.8, PyTorch 2.11.0+cu128:

- `test_batch_mla_without_kpe`: FA2 and FA3 passed.
- `test_batch_mla_fp8_nope_group_scales_matches_bf16_reference`: FA2 and
FA3 passed.
- `test_fp8_kv_scales_are_keyword_only` and
`test_fp8_kv_requires_scales`: passed.
- Representative existing KPE=64 per-tensor FP8 and row-aliasing
regressions: passed.

## Reviewer Notes

- `ckv_scale_arr` follows the physical paged-cache layout used by the
downstream SGLang integration: `[num_pages, page_size, head_dim_ckv /
128]`.
- Existing BF16/FP16 and FP8 per-tensor scale behavior is unchanged.
- The latest external attention CI failures before this update were not
stable across reruns: HCA failures were classified as pre-existing,
while the new failures alternated between GB200 multi-node NCCL setup
and RTX Pro 6000 `cudaErrorUnknown` failures.


<!-- This is an auto-generated comment: release notes by coderabbit.ai
-->
## Summary by CodeRabbit

* **New Features**
* Added FP8 batched paged attention with optional per-token/per-channel
CKV scaling.
  * Added support for FA2 and FA3 FP8 execution.
* Expanded supported configurations to include zero- or 64-dimensional
key positional encoding.

* **Bug Fixes**
* Fixed attention computation without key positional encoding
dimensions.
* Added validation for incompatible scale inputs, shapes, layouts, and
configurations.

* **Tests**
  * Added coverage for grouped CKV scaling and trace inputs.
  * Added BF16 reference comparisons across supported implementations.
<!-- end of auto-generated comment: release notes by coderabbit.ai -->

---------

Co-authored-by: Lee Nau <lnau@nvidia.com>
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.

4 participants