Skip to content

[CUDA] Add Triton Q4_K-to-FP8 prefill linear - #22192

Open
Gasoonjia wants to merge 4 commits into
codex/stack-tma-prefillfrom
codex/stack-q4k-fp8-prefill
Open

[CUDA] Add Triton Q4_K-to-FP8 prefill linear#22192
Gasoonjia wants to merge 4 commits into
codex/stack-tma-prefillfrom
codex/stack-q4k-fp8-prefill

Conversation

@Gasoonjia

@Gasoonjia Gasoonjia commented Aug 26, 2026

Copy link
Copy Markdown
Contributor

Stack dependency

Depends on #22414.

Summary

Add an SM90+ pure-Triton Q4_K-to-FP8 prefill linear path on top of the TMA attention PR.

  • Dequantize Q4_K weights directly to FP8 and feed an FP8 matrix multiplication with BF16 output.
  • Register the custom-op schema unconditionally so exported programs remain loadable on every host.
  • Select the optimized path from the requested CUDA target: explicit TORCH_CUDA_ARCH_LIST takes precedence, while native export falls back to the local device capability.
  • Avoid CUDA context creation at module import and explicitly exclude ROCm.
  • Retain the existing low-token decode kernel and BF16 prefill fallback for pre-SM90 targets, non-BF16 activations, non-contiguous activations, and unsupported shapes.
  • Apply the same input validation to the real and fake implementations.
  • Add float8 dtype support to the slim AOT runtime shims.

The path is controlled by the enable_q4k_fp8_prefill CUDA compile spec and is disabled by default. Users must explicitly opt into the lossy performance path. Unsupported targets warn and retain the BF16 fallback. This change remains Triton-only and does not introduce CUDA or CuteDSL kernels.

This is the top of the reordered stack:

  1. [CUDA] Elide dense causal masks in long-context prefill #22189: D111 causal-mask elision.
  2. [CUDA] Tune hardware-aware split-K decode scheduling #22190: hardware-aware online-softmax split-K scheduling.
  3. [CUDA] Use Triton TMA for long-context causal prefill #22414: opt-in TMA long-context causal prefill.
  4. [CUDA] Add Triton Q4_K-to-FP8 prefill linear #22192: opt-in Q4_K-to-FP8 prefill linear.

Operator-level performance

RTX 5090, representative Muse MLP projection (N=39,936, K=6,656), Q4_K dequantization plus linear:

Prefill rows M BF16 dequant + linear Q4_K-to-FP8 path Speedup
128 5.033 ms 0.507 ms 9.93x
512 5.818 ms 0.945 ms 6.16x
2,048 9.182 ms 2.688 ms 3.42x

NLL evaluation

notes_v5, 100 documents / 154,846 next-token labels. The supplied .pt stores one token chunk per document; labels are chunk[1:] for inputs chunk[:-1] (all chunks begin with BOS 200000).

Comparison Mean token abs NLL gap Max token gap Mean document gap Max document gap Document buckets (same/close/match/mismatch)
A100 BF16 reference vs #22190 0.023180 1.462394 0.001020 0.003171 56 / 44 / 0 / 0
A100 BF16 reference vs final TMA+FP8 top 0.159700 - 0.022701 0.067222 1 / 12 / 87 / 0

Buckets use document mean-NLL gap: <1e-3, [1e-3,1e-2), [1e-2,1e-1), and >=1e-1. The FP8 path materially changes NLL, so it is intentionally opt-in rather than the default.

End-to-end top-stack performance

Both optional prefill paths enabled. Single otherwise-idle RTX 5090, single-GPU llama-cli, P2409802035 protocol, ET and llama.cpp interleaved per repetition, runs 3-7 mean:

Prompt length ET prefill tok/s llama.cpp prefill tok/s ET vs llama.cpp prefill ET decode tok/s llama.cpp decode tok/s ET vs llama.cpp decode
512 3216.25 2877.14 +11.79% 77.69 78.34 -0.83%
2,048 5378.22 3615.14 +48.77% 77.21 77.12 +0.11%
8,192 5740.71 4060.96 +41.36% 74.77 76.62 -2.42%
32,768 5378.54 4012.94 +34.03% 74.70 74.94 -0.32%
130,048 4167.30 3455.06 +20.61% 70.63 69.54 +1.56%

Across the five lengths, ET prefill is 30.61% faster geometric mean than single-GPU llama-cli; decode is effectively tied overall at -0.39% geometric mean.

For comparison, the lower TMA-only PR is 11.75% slower geometric mean than llama.cpp in prefill. Adding this FP8 linear path above TMA supplies the large model-wide prefill uplift; TMA supplies the additional long-context attention gain.

The isolated 8K decode miss is primarily launch overlap rather than SDPA kernel latency. In a matched Nsight trace, ET's stable decode graph is serialized (12.421 ms summed kernel time, 12.742 ms wall time), while llama.cpp overlaps up to seven kernels (16.114 ms summed kernel time, 12.460 ms wall time). llama.cpp enables Programmatic Dependent Launch by default: an idle-GPU 7-run A/B measured 76.58 tok/s with PDL and 75.24 tok/s with GGML_CUDA_PDL=0 (+1.78%). PDL remains intentionally outside this stack.

Test plan

  • lintrunner init && lintrunner -a: clean on both reordered branches.
  • Focused final FP8/TMA routing tests: passed.
  • Full int4 dispatch suite: 26 passed before the reorder; the reordered focused dispatch suite also passed.
  • Accuracy is checked against _dequant_matmul with normalized max error below 5%.
  • Covered compile-spec default/off/on behavior, invalid values, unsupported targets, explicit target-architecture selection, ROCm exclusion, non-BF16 fallback, and non-contiguous fallback.
  • Verified that an opted-in SM90+ export captures triton.q4k_fp8_linear, while default export retains the BF16 path.
  • Final opted-in PTE profile contains 208 FP8 matmul and 208 Q4_K-to-FP8 dequant kernel calls.
  • git diff --check.

@pytorch-bot

pytorch-bot Bot commented Aug 26, 2026

Copy link
Copy Markdown

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/pytorch/executorch/22192

Note: Links to docs will display an error until the docs builds have been completed.

⏳ 13 Pending, 2 Unrelated Failures

As of commit fc45c0d with merge base 5428092 (image):

FLAKY - The following job failed but was likely due to flakiness present on trunk:

BROKEN TRUNK - The following job failed but was present on the merge base:

👉 Rebase onto the `viable/strict` branch to avoid these failures

This comment was automatically generated by Dr. CI and updates every 15 minutes.

@meta-cla meta-cla Bot added the CLA Signed This label is managed by the Facebook bot. Authors need to sign the CLA before a PR can be reviewed. label Aug 26, 2026
@Gasoonjia
Gasoonjia force-pushed the codex/stack-q4k-fp8-prefill branch from 0088e51 to d11d0c6 Compare August 26, 2026 21:56
@Gasoonjia
Gasoonjia changed the base branch from codex/stack-online-splitk-tuning to codex/stack-d113-persistent-splitk August 26, 2026 21:56
@Gasoonjia
Gasoonjia force-pushed the codex/stack-d113-persistent-splitk branch from b8ef410 to 8c9f0cc Compare August 26, 2026 22:50
@Gasoonjia
Gasoonjia force-pushed the codex/stack-q4k-fp8-prefill branch from d11d0c6 to 28a5119 Compare August 26, 2026 22:50
@Gasoonjia
Gasoonjia force-pushed the codex/stack-d113-persistent-splitk branch from 8c9f0cc to cfd96e1 Compare August 26, 2026 23:03
@Gasoonjia
Gasoonjia force-pushed the codex/stack-q4k-fp8-prefill branch from 28a5119 to 37690cd Compare August 26, 2026 23:03
@Gasoonjia
Gasoonjia force-pushed the codex/stack-q4k-fp8-prefill branch from 37690cd to 5bc221a Compare August 26, 2026 23:39
@Gasoonjia
Gasoonjia force-pushed the codex/stack-d113-persistent-splitk branch from cfd96e1 to 3ef87d1 Compare August 26, 2026 23:39
tl.float32
)
value = (q - zc * zs) * (sc * ss)
tl.store(out + idx, value.to(tl.float8e4nv), mask=mask)

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.

why does this feels like this can be lossy or clamp, let's do some evals? NLL might be a good one just comparing these two..

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.

I've add a accuracy test comparing with _dequant_matmul to guard the correctness. Let me do nll verfication on top of the stack

@Gasoonjia
Gasoonjia force-pushed the codex/stack-d113-persistent-splitk branch from 281b666 to 65deb1a Compare August 27, 2026 19:41
@Gasoonjia
Gasoonjia force-pushed the codex/stack-q4k-fp8-prefill branch from c6cd8e3 to 7722db7 Compare August 27, 2026 19:41
@Gasoonjia
Gasoonjia force-pushed the codex/stack-d113-persistent-splitk branch from 65deb1a to 1044521 Compare September 1, 2026 07:42
@Gasoonjia
Gasoonjia force-pushed the codex/stack-q4k-fp8-prefill branch from 7722db7 to 8f5d2a7 Compare September 1, 2026 07:42
Add an SM90+ Triton path that dequantizes Q4_K weights directly to FP8 and feeds FP8 matrix multiplication during prefill. Extend the AOT slim dtype shims for float8 and make export-time dispatch select the CUDA target while tracing CPU example tensors.
@Gasoonjia
Gasoonjia force-pushed the codex/stack-q4k-fp8-prefill branch from c01c3df to fc45c0d Compare September 1, 2026 18:25
@Gasoonjia
Gasoonjia changed the base branch from codex/stack-d113-persistent-splitk to codex/stack-tma-prefill September 1, 2026 18:26
@Gasoonjia
Gasoonjia force-pushed the codex/stack-tma-prefill branch 2 times, most recently from 51ff462 to 8da3f29 Compare September 1, 2026 18:27
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

CLA Signed This label is managed by the Facebook bot. Authors need to sign the CLA before a PR can be reviewed.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants