[CUDA] Add Triton Q4_K-to-FP8 prefill linear - #22192
Open
Gasoonjia wants to merge 4 commits into
Open
Conversation
🔗 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 FailuresAs of commit fc45c0d with merge base 5428092 ( 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. |
Gasoonjia
force-pushed
the
codex/stack-q4k-fp8-prefill
branch
from
August 26, 2026 21:56
0088e51 to
d11d0c6
Compare
Gasoonjia
changed the base branch from
codex/stack-online-splitk-tuning
to
codex/stack-d113-persistent-splitk
August 26, 2026 21:56
Gasoonjia
force-pushed
the
codex/stack-d113-persistent-splitk
branch
from
August 26, 2026 22:50
b8ef410 to
8c9f0cc
Compare
Gasoonjia
force-pushed
the
codex/stack-q4k-fp8-prefill
branch
from
August 26, 2026 22:50
d11d0c6 to
28a5119
Compare
Gasoonjia
force-pushed
the
codex/stack-d113-persistent-splitk
branch
from
August 26, 2026 23:03
8c9f0cc to
cfd96e1
Compare
Gasoonjia
force-pushed
the
codex/stack-q4k-fp8-prefill
branch
from
August 26, 2026 23:03
28a5119 to
37690cd
Compare
Gasoonjia
force-pushed
the
codex/stack-q4k-fp8-prefill
branch
from
August 26, 2026 23:39
37690cd to
5bc221a
Compare
Gasoonjia
force-pushed
the
codex/stack-d113-persistent-splitk
branch
from
August 26, 2026 23:39
cfd96e1 to
3ef87d1
Compare
digantdesai
reviewed
Aug 27, 2026
| tl.float32 | ||
| ) | ||
| value = (q - zc * zs) * (sc * ss) | ||
| tl.store(out + idx, value.to(tl.float8e4nv), mask=mask) |
Contributor
There was a problem hiding this comment.
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..
Contributor
Author
There was a problem hiding this comment.
I've add a accuracy test comparing with _dequant_matmul to guard the correctness. Let me do nll verfication on top of the stack
digantdesai
requested changes
Aug 27, 2026
Gasoonjia
force-pushed
the
codex/stack-d113-persistent-splitk
branch
from
August 27, 2026 19:41
281b666 to
65deb1a
Compare
Gasoonjia
force-pushed
the
codex/stack-q4k-fp8-prefill
branch
from
August 27, 2026 19:41
c6cd8e3 to
7722db7
Compare
Gasoonjia
force-pushed
the
codex/stack-d113-persistent-splitk
branch
from
September 1, 2026 07:42
65deb1a to
1044521
Compare
Gasoonjia
force-pushed
the
codex/stack-q4k-fp8-prefill
branch
from
September 1, 2026 07:42
7722db7 to
8f5d2a7
Compare
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
force-pushed
the
codex/stack-q4k-fp8-prefill
branch
from
September 1, 2026 18:25
c01c3df to
fc45c0d
Compare
Gasoonjia
changed the base branch from
codex/stack-d113-persistent-splitk
to
codex/stack-tma-prefill
September 1, 2026 18:26
Gasoonjia
force-pushed
the
codex/stack-tma-prefill
branch
2 times, most recently
from
September 1, 2026 18:27
51ff462 to
8da3f29
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
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.
TORCH_CUDA_ARCH_LISTtakes precedence, while native export falls back to the local device capability.The path is controlled by the
enable_q4k_fp8_prefillCUDA 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:
Operator-level performance
RTX 5090, representative Muse MLP projection (
N=39,936,K=6,656), Q4_K dequantization plus linear:NLL evaluation
notes_v5, 100 documents / 154,846 next-token labels. The supplied.ptstores one token chunk per document; labels arechunk[1:]for inputschunk[:-1](all chunks begin with BOS 200000).same/close/match/mismatch)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: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._dequant_matmulwith normalized max error below 5%.triton.q4k_fp8_linear, while default export retains the BF16 path.git diff --check.