ONNX Runtime #32128 — Why the Existing FP4 Tests Never Reached the Prefill Dequantization Kernel

ONNX Runtime NVFP4, Part 3 of 3 — How M <= 8 kept selecting decode GEMV and left the production bottleneck outside regression coverage

The first two articles in this series examined a major performance bottleneck in ONNX Runtime’s NVFP4 prefill path.

On H200, the native Blackwell SM120 FP4 matmul path was not available.

For prefill workloads processing multiple prompt tokens, ONNX Runtime therefore expanded the packed NVFP4 weights into an FP16 or BF16 scratch tensor before calling cuBLAS.

Packed NVFP4 weights
        ↓
Dequantize into an FP16/BF16 [N, K] scratch tensor
        ↓
Run cuBLAS GEMM

This expansion was not a minor preprocessing step.

In the recorded Qwen3.8-27B workload with an 8K prompt, the original DequantizeNvFp4Kernel accounted for 46.9% of total prefill GPU time.

It consumed more time than all of the model’s cuBLAS GEMMs combined.

PR #32128 introduced DequantizeNvFp4Vec8Kernel, which processes eight FP4 values per thread for common layouts.

The new path reduced the dequantization-kernel time by approximately 3.8–3.9× in the reported H200 measurements.

The recorded end-to-end time to first token also fell from 3,876ms to 2,618ms.

But the patch exposed another important problem.

FP4 operator tests already existed.

They were passing.

Yet the prefill dequantization kernel responsible for nearly half of the recorded prefill GPU time had never executed in those tests.

The reason was not the dtype.

It was not the operator name.

It was the matrix shape used by every existing case.

Existing FP4 tests

M <= 8
        ↓
Decode GEMV selected
        ↓
Prefill dequantization never executed

The tests exercised the FP4 operator.

They did not exercise the FP4 execution path that contained the production bottleneck.

One operator name can hide several different kernels

At the graph level, the operation may appear as one operator:

MatMulBlockQuantizedFp4Weight

Internally, the runtime can select very different implementations according to:

matrix shape
GPU architecture
data type
scale layout
runtime configuration

A simplified dispatch structure is:

MatMulBlockQuantizedFp4Weight
        ↓
Evaluate dispatch conditions
        ↓
┌────────────────────────────────────┐
│ Small M                            │
│ → fused decode GEMV                │
├────────────────────────────────────┤
│ Native Blackwell conditions        │
│ → native FP4 matmul                │
├────────────────────────────────────┤
│ Other prefill workloads            │
│ → weight dequantization            │
│ → FP16/BF16 scratch                │
│ → cuBLAS GEMM                      │
└────────────────────────────────────┘

Calling the same operator does not prove that the same kernel executed.

Even with:

the same FP4 weight format

the same output dtype

the same operator name

a change in M can select a different execution contract.

In the existing tests, small values of M selected the decode-oriented GEMV path.

The new tests use M > 8 specifically to prevent that route from being selected and to enter the prefill dequantization path instead.

Decode and prefill are not the same matrix-multiplication workload

LLM inference is commonly divided into two stages.

Prefill

Prefill processes the prompt tokens together.

Long prompt
        ↓
Many token rows processed together
        ↓
Relatively large M

The workload provides enough parallel work for a conventional GEMM-oriented path.

Decode

Decode generates the next token using the existing KV cache.

Existing KV cache
+
one or a few new token rows
        ↓
small M

For small M, a specialized fused GEMV can avoid some of the overhead associated with a larger GEMM setup.

ONNX Runtime reflected that distinction in its dispatch logic:

Small M
→ decode GEMV

Larger M
→ prefill dequantization and GEMM

A test suite that always uses small M can therefore execute the FP4 operator repeatedly while observing only the decode path.

The prefill implementation may remain present in the source but unreachable from the test matrix.

Why the existing tests passed

The existing FP4 tests were not accepting incorrect results.

They correctly validated the path they actually executed.

Prepare FP4 input
        ↓
Call MatMulBlockQuantizedFp4Weight
        ↓
Decode GEMV executes
        ↓
Compare against the reference
        ↓
PASS

That PASS established:

Under the tested small-M conditions, the selected decode GEMV produced the expected output.

It did not establish:

Under larger-M conditions, the prefill dequantization path also works correctly and performs reasonably on production-like shapes.

The test result was not false.

The interpretation of its scope could become false.

Operator-level PASS
≠
Every dispatch branch passed

Dtype coverage and execution-path coverage are different

A test matrix containing FP4 cases can establish dtype coverage.

It may answer questions such as:

Can the operator accept an FP4 weight?

Can it produce FP16 output?

Can it produce BF16 output?

A multi-path runtime requires another set of questions:

Did the decode GEMV execute?

Did the vectorized prefill dequantizer execute?

Did the scalar prefill fallback execute?

Did the native Blackwell path execute?

These forms of coverage are related but not interchangeable.

Coverage type	Question
Dtype coverage	Can FP4 values and FP16/BF16 results be represented correctly?
Shape coverage	Are small-M, large-M, aligned-K, and ragged-K cases included?
Dispatch coverage	Which internal implementation was selected?
Fast-path coverage	Did the optimized kernel run when its conditions were satisfied?
Fallback coverage	Did unsupported fast-path inputs return to the safe implementation?
Model coverage	Did the full LLM preserve output and improve end-to-end behavior?

The existing tests covered FP4 values and the decode route.

They did not cover the prefill dispatch branch.

The most expensive production path was invisible to the tests

The deeper problem was the mismatch between test frequency and production cost.

Path repeatedly executed by existing tests

Decode GEMV
→ covered

Path consuming the most time in the recorded 8K prefill

DequantizeNvFp4Kernel
→ not reached

The Qwen3.8-27B workload contained:

168 MatMulBlockQuantizedFp4Weight nodes

14.97 billion weight values

During prefill, those weights were expanded through the scalar dequantization kernel.

The kernel consumed 1,750ms in the recorded run.

Yet every existing FP4 test used M <= 8.

The test suite could contain many cases and still repeat the same internal branch.

Many FP4 tests

all using small M

result:
the same decode path is exercised repeatedly

prefill path:
untested

A high test count does not necessarily imply broad dispatch coverage.

The first purpose of `M > 8` was to change the dispatch

PR #32128 added four prefill dequantization test cases.

Every one of them uses:

M > 8

The purpose was not merely to create a numerically larger test.

It was to avoid the decode GEMV route.

Before

M <= 8
→ decode GEMV

New tests

M > 8
→ decode GEMV skipped
→ prefill dequantization entered

The shape was therefore not only data.

It was a control input selecting the implementation under test.

In low-level runtimes, a shape can determine:

dispatch identity

kernel identity

memory layout

vector width

scratch allocation

synchronization behavior

A test shape can select a contract before it produces a numerical result.

The four new tests protect different boundaries

The patch did not add only one successful vectorized example.

It added positive fast-path cases and cases that must remain on the scalar fallback.

Test	Intended condition	Intended kernel
`PrefillDequantVectorizedFp16`	FP16 output, vectorized conditions satisfied	`DequantizeNvFp4Vec8Kernel<__half>`
`PrefillDequantVectorizedBiasBf16`	BF16 output with bias, vectorized conditions satisfied	`DequantizeNvFp4Vec8Kernel<__nv_bfloat16>`
`PrefillDequantOddBlockSizeFp16`	Odd `block_size`	Scalar FP16 dequantization
`PrefillDequantKNotMultipleOf8Bf16`	`K % 8 != 0`	Scalar BF16 dequantization

All four cases use M > 8.

They first enter the prefill branch.

Within that branch, the inputs then determine whether the vectorized fast path is valid.

Vectorization contract satisfied
→ new fast path

Vectorization contract not satisfied
→ scalar fallback

The change description says each case was confirmed under nsys to reach its intended kernel, and all 17 related MatMulBlockQuantizedFp4WeightOpTest cases passed.

The first positive case protects FP16 vectorization

PrefillDequantVectorizedFp16 is the most direct positive test for the new kernel.

Its contract is:

M > 8
→ avoid decode GEMV

K % 8 == 0
→ every thread receives one complete eight-element chunk

block_size is even
→ both FP4 codes in one packed byte share a scale

output dtype is FP16
→ DequantizeNvFp4Vec8Kernel<__half>

The test needs to protect more than the final FP16 values.

The complete contract includes:

correct dispatch

correct packed load

correct FP4 decode

correct scale selection

correct FP16 vector store

correct final matmul output

A correct output produced by another implementation would not provide the same evidence about the new fast path.

The second positive case protects BF16 and bias handling

PrefillDequantVectorizedBiasBf16 exercises another specialization.

M > 8

K % 8 == 0

even block_size

BF16 output

bias present
        ↓
DequantizeNvFp4Vec8Kernel<__nv_bfloat16>

A passing FP16 test does not automatically protect the BF16 specialization.

Changing the output type can affect:

conversion helpers
vector storage type
output bit patterns
reference comparison
the operand passed into the subsequent GEMM

The presence of bias can also affect argument wiring and the complete operator route.

The second positive case is therefore not merely redundant coverage.

It protects a separate fast-path specialization.

Positive tests alone cannot prove that the guard is narrow enough

The vectorized fast path requires:

K % 8 == 0

and

block_size % 2 == 0

Suppose a future change accidentally broadens the dispatch condition:

all K values

all block sizes
→ vectorized kernel

The ordinary positive cases would still pass.

They already satisfy the vectorization contract.

The failure would appear only when the fast path is selected for an input that violates one of its assumptions.

A complete regression matrix therefore needs:

Inputs that must select the fast path
+
Inputs that must not select the fast path

The fallback tests define the negative space of the optimization.

The odd-`block_size` test protects a scale boundary

PrefillDequantOddBlockSizeFp16 uses an odd block size.

This matters because two FP4 values share one packed byte.

Packed byte

first nibble
→ FP4 value A

second nibble
→ FP4 value B

The vectorized kernel processes the pair under the assumption that both values share the same block scale.

With an even block size, the scale boundary cannot split a two-value pair.

For example:

block_size = 16

Values 0–15
→ scale block 0

Values 16–31
→ scale block 1

The packed pairs are:

(0, 1)
(2, 3)
...
(14, 15)

Every pair stays within one scale block.

With an odd block size:

block_size = 15

a packed byte can contain:

Value 14
→ scale block 0

Value 15
→ scale block 1

One byte now contains two values requiring different scales.

one packed byte
+
two scale identities

The pair-level fast-path assumption no longer holds.

The correct dispatch is:

odd block_size
→ reject the vectorized path
→ use the scalar kernel

This test protects the boundary at which the scale contract changes.

It does not merely check another output shape.

The non-divisible-K test protects the row tail

PrefillDequantKNotMultipleOf8Bf16 uses a K dimension that is not divisible by eight.

The fast path assigns exactly eight K-axis values to each thread.

One thread
→ eight values

If K is divisible by eight, every chunk is complete.

K = 128

128 / 8
=
16 complete chunks

If K is not divisible by eight, the final chunk is partial.

K = 12

Chunk 0
→ values 0–7

Chunk 1
→ values 8–11 are valid
→ four positions fall beyond the logical row

The fast path assumes:

one full 32-bit packed load

eight valid decoded values

one full 16-byte output store

Using that unmasked contract on a partial tail could cause:

a read beyond the packed row

an output write into the next row

an incorrect scale lookup

PR #32128 did not add a masked vector tail.

It preserved the scalar fallback:

K % 8 != 0
→ scalar dequantization

The fallback test protects the rule:

If a complete eight-element chunk cannot be proven, the vectorized kernel must not be selected.

Negative tests define the legal scope of an optimization

Performance tests often focus on the successful fast path.

The correctness boundary frequently lives in the cases where that fast path must refuse to run.

The matrix here is:

Positive fast path

K % 8 == 0
+
even block_size
→ vectorized kernel

Negative boundary 1

odd block_size
→ a packed pair may cross the scale boundary
→ scalar kernel

Negative boundary 2

K % 8 != 0
→ the final eight-element chunk is incomplete
→ scalar kernel

The positive test asks:

Does the fast path work when its assumptions are true?

The fallback tests ask:

Does the runtime avoid the fast path when those assumptions are false?

Those are separate contracts.

Correct output does not prove that the intended kernel executed

The new tests use shape and parameter values designed to select particular kernels.

The author also used nsys to confirm the actual kernel reached by each test:

Vectorized FP16
→ DequantizeNvFp4Vec8Kernel<__half>

Vectorized BF16
→ DequantizeNvFp4Vec8Kernel<__nv_bfloat16>

Odd block size
→ DequantizeNvFp4Kernel<__half>

Non-divisible K
→ DequantizeNvFp4Kernel<__nv_bfloat16>

That confirmation matters because the following false sense of coverage is possible:

Test input is created

Dispatch behavior changes unexpectedly

Another correct fallback implementation runs

Output matches the reference

Test passes

The numerical result would be correct.

The intended fast path might no longer have any regression coverage.

The public change description establishes that nsys was used during validation to confirm the intended kernels.

It does not, by itself, establish that every CI execution asserts the selected kernel symbol directly.

The retained automated evidence appears to rely primarily on the dispatch-producing inputs and output comparison.

A stable runtime-path assertion could strengthen future coverage if the dispatch architecture changes substantially, but that would be a future design choice rather than a feature demonstrated by this patch.

Three evidence layers are needed

The test design can be divided into three layers.

1. Dispatch conditions

M > 8

K alignment

block-size parity

output dtype

bias presence

These conditions describe which implementation should be selected.

2. Runtime-path confirmation

The author used nsys to verify which kernel actually executed.

3. Numerical result

The operator output was compared against a reference.

Each layer answers a different question.

Dispatch conditions
→ Which path should this input choose?

Runtime-path confirmation
→ Which path did the runtime actually choose?

Output comparison
→ Did that selected path produce the correct result?

Output comparison alone cannot answer all three.

Model-level hashing checks a different boundary

PR #32128 also compares the generated-token SHA-256 for the actual Qwen3.8-27B workload.

An operator test covers:

Known input tensors
        ↓
dequantization and matmul
        ↓
reference output comparison

A model-level run covers a much wider chain:

168 quantized matmul nodes

many accumulated layers

attention and residual paths

token selection

multi-token prediction
        ↓
final generated-token sequence

A small difference in one conversion operation could propagate across layers and eventually change token generation.

The recorded generated-token hash remained unchanged across three runs per comparison arm.

The complete evidence chain became:

Primitive-level packed-byte comparison

Operator-level tensor hashing

Model-level generated-token hashing

A model hash cannot replace targeted fallback tests

The reverse is also true.

A full model producing the same tokens does not prove that every dispatch edge is protected.

Real models tend to use common layouts.

They may naturally satisfy:

even block_size

K % 8 == 0

The actual Qwen workload can repeatedly exercise the vectorized fast path while never presenting:

odd block_size

non-divisible K

The model test therefore protects realistic system behavior.

The targeted operator tests protect artificial boundary conditions.

Model test
→ realistic end-to-end workload

Targeted operator test
→ deliberately constructed dispatch edge

Both are needed.

Shape is a control input, not only a tensor dimension

In a simple unit test, shape can appear to be only the size of the data.

In a dispatching runtime, shape can control the code path.

M
→ decode or prefill

K % 8
→ vectorized or scalar

block_size parity
→ pair-scale assumption valid or invalid

output dtype
→ FP16 or BF16 specialization

The input fixture therefore has two roles:

Values to compute
+
Control information selecting the implementation

Choosing a numerically representative shape is not enough.

The test designer must also know what the shape causes the runtime to execute.

Small shapes can hide entire implementations

In CUTLASS #3017, kElementsPerAccess = 1 hid an address-calculation asymmetry because division by one did not change the result.

In this ONNX Runtime case, M <= 8 hid the prefill implementation because the runtime selected decode GEMV instead.

The common structure is:

Small or default condition
        ↓
One branch is selected repeatedly
        ↓
Another implementation remains unobserved
        ↓
The suite passes

Small shapes are not invalid.

Decode is a real and important workload.

The problem was that no larger-M counterpart existed.

Small-M test
→ proves the decode contract

Large-M test
→ proves the prefill contract

Both are required.

Count dispatch partitions, not only test cases

Imagine a test suite containing:

20 FP16 cases

20 BF16 cases

10 bias cases

30 combinations of N and K

The suite appears broad.

But if every case uses:

M <= 8

then all 80 tests may select the same decode GEMV branch.

80 tests

one dispatch partition

A smaller set of four tests can protect more execution contracts if it deliberately covers:

vectorized FP16

vectorized BF16 with bias

odd-block-size scalar fallback

non-divisible-K scalar fallback

The meaningful unit is not only the number of cases.

It is the number of distinct dispatch regions represented by those cases.

The dispatch matrix after the patch

A simplified coverage matrix now looks like this:

M condition	K condition	`block_size`	Selected path	Coverage
`M <= 8`	multiple	multiple	Decode GEMV	Existing tests
`M > 8`	`K % 8 == 0`	even	Vectorized prefill dequantization	Added
`M > 8`	`K % 8 == 0`	odd	Scalar prefill dequantization	Added
`M > 8`	`K % 8 != 0`	multiple	Scalar prefill dequantization	Added
Native SM120 conditions	separate	separate	Blackwell native FP4	Not the direct target of this patch

The new cases separate at least four contracts:

Decode path

Prefill fast path

Prefill scale-boundary fallback

Prefill tail fallback

Numerical tests alone may not catch a performance regression

The original scalar dequantization kernel was numerically correct.

Its problem was performance.

Correct output
+
46.9% of total prefill GPU time

A numerical operator test can execute the prefill path and still pass even if that path becomes dramatically slower.

PR #32128 therefore includes several forms of evidence:

Correctness test
→ Are the values unchanged?

Dispatch validation
→ Did the intended kernel execute?

Kernel measurement
→ Did the optimized kernel become faster?

Model benchmark
→ Did the user-visible TTFT improve?

These are separate claims.

Hardware performance thresholds can be difficult to enforce in CI because timing varies across machines and system conditions.

But the first requirement is unavoidable:

A kernel that never executes in the test suite cannot have either its correctness or its performance protected there.

Unreached code is effectively unprotected code

Suppose a future change introduces a bug in DequantizeNvFp4Vec8Kernel.

If every FP4 test still entered decode GEMV:

Vectorized prefill kernel becomes incorrect

Test suite runs

Only decode GEMV executes

All tests pass

The kernel exists in the repository.

It does not exist within the observed test execution.

Present in source
≠
Protected by regression tests

Protection requires at least:

Reach the branch

Exercise its memory and scale contract

Validate the resulting output

Fallback paths can survive longer without observation

The optimized path targets common production layouts.

It may naturally receive attention from benchmarks and real-model tests.

Fallback paths can be much rarer.

Odd block size

Ragged K

Uncommon layout

Those inputs may not appear in a major model workload for a long time.

This makes targeted fallback cases especially important.

The fast path can be exercised by performance benchmarks.

Fallback behavior often needs to be manufactured deliberately.

Test names can document the dispatch contract

The added test names describe more than the expected output.

PrefillDequantVectorizedFp16

encodes:

Prefill

Dequantization

Vectorized path

FP16 specialization

PrefillDequantOddBlockSizeFp16

encodes:

Prefill

Odd block size

Scalar fallback boundary

FP16 output

A useful test name can identify the protected execution contract.

When a failure occurs, the developer sees:

odd-block-size scalar fallback failed

rather than only:

FP4 output mismatch

That distinction narrows the investigation immediately.

What the patch added to the test contract

PR #32128 directly added:

prefill shapes using M > 8
an FP16 vectorized case
a BF16 vectorized case with bias
an odd-block_size scalar-fallback case
a K % 8 != 0 scalar-fallback case
author validation under nsys that each case reached the intended kernel
passing results for all 17 related operator-test cases

What the public evidence does not establish

The available material does not establish that:

every CI run asserts the exact selected kernel symbol
the recorded performance numbers are enforced as regression thresholds
every prefill shape or block size is covered
the native Blackwell SM120 path is protected by these four tests
future dispatch changes will preserve the same kernel selection for these inputs
H200 bandwidth and timing figures reproduce identically on every system
every possible scale layout and K tail is represented

The supported conclusion is narrower:

The new inputs were designed to enter the intended prefill branches, the author confirmed their kernel selection under nsys, and the corresponding numerical tests passed.

Test the execution path, not only the dtype

The entire event can be summarized in two sentences:

FP4 tests existed.

FP4 prefill tests did not.

More precisely:

MatMulBlockQuantizedFp4Weight tests existed.

Tests that reached DequantizeNvFp4Kernel
under actual prefill conditions did not.

Classifying coverage only by operator name and dtype can hide this difference.

A low-level runtime test matrix should ideally record:

Operator

Dtype

Shape

Hardware

Dispatch predicate

Selected kernel

Fallback reason

Reproducing the same operator is not enough

Suppose a slow production kernel is identified.

A regression test that calls the same high-level operator may still execute another implementation.

Production

large M
→ prefill path

Test

small M
→ decode path

The test would not reproduce the relevant contract.

The dispatch predicates that selected the production path need to be preserved in the fixture.

Production dispatch condition
        ↓
Explicitly represented in the test

For this patch, one such condition was:

M > 8

Regression design begins at the dispatch entrance

The defect was located inside the dequantization kernel.

That does not mean the test had to call the kernel directly.

Testing through the public operator can be stronger because it includes:

MatMulBlockQuantizedFp4Weight
        ↓
real dispatch
        ↓
vectorized or scalar dequantization
        ↓
scratch allocation
        ↓
cuBLAS
        ↓
final output

This route verifies:

dispatch predicates
argument wiring
scratch creation
dequantization
the subsequent GEMM
final numerical behavior

But that strength exists only when the input actually selects the intended branch.

The complete three-part structure

The ONNX Runtime #32128 series now forms one chain.

Part 1 — The bottleneck was not four-bit arithmetic

NVFP4 weights
        ↓
No native H200 prefill FP4 matmul
        ↓
Expand into FP16/BF16 scratch
        ↓
Dequantization consumes 46.9% of prefill GPU time

Small storage did not guarantee low execution cost.

Part 2 — More work per thread was not always better

Eight values per thread
→ one uint4 store
→ warp-contiguous output
→ about 3.9 TB/s

Thirty-two values per thread
→ four uint4 stores
→ strided addresses within each instruction
→ about 1.9 TB/s

The effective vectorization unit was the warp address pattern, not only the thread-local vector.

Part 3 — The tests exercised FP4 but not the production bottleneck

Existing tests:
M <= 8
→ decode GEMV

Production bottleneck:
M > 8
→ prefill dequantization

Testing the relevant dtype was not the same as testing the relevant kernel.

The patch restored observability as well as performance

The code change made the dequantization kernel faster.

The test change made the path observable.

Before

The most expensive production kernel
→ not reached by the FP4 test matrix

After

Production bottleneck path
→ reached by targeted prefill tests

Fast-path contract
→ positive cases

Fallback contract
→ negative boundary cases

Model behavior
→ generated-token hash

Performance
→ kernel timing and TTFT measurement

The optimization became a regression-testable contract rather than only a faster implementation.

Coverage should be judged by reachability

A relevant test file exists.

The operator name appears in it.

The dtype is present.

Every test passes.

Those facts may still be insufficient.

The decisive question is:

Did the test reach the code we intended to protect?

For this event, the answer changed from:

Before:
No

to:

After:
Yes

The old FP4 tests were not incorrect.

They were looking at decode.

The production bottleneck was in prefill.

A test does not protect every implementation that exists behind an operator. It protects only the execution paths it actually reaches.

ONNX Runtime #32128 — Why Dequantizing 4-Bit Weights Became the Most Expensive Kernel in H200 Prefill
ONNX Runtime #32128 — Why Eight FP4 Values per Thread Were Nearly Twice as Fast as Thirty-Two

Related material

Patch status: Merged into ONNX Runtime main
Previous test boundary: Every existing FP4 case used M <= 8 and entered decode GEMV
New prefill condition: All four newly added cases use M > 8
Fast-path coverage: FP16 and BF16-with-bias vectorized dequantization
Fallback coverage: Odd block_size and K % 8 != 0
Runtime-path verification: The author confirmed the intended kernels under nsys
Operator tests: All 17 related cases passed
Core regression contract: Test the actual prefill dequantization dispatch, not merely the FP4 operator

This is Part 3 and the final article in the ONNX Runtime NVFP4 series.

Part 1 examined why H200 expanded the complete packed NVFP4 weight into an FP16 or BF16 scratch tensor and why that conversion dominated the recorded model latency.

Part 2 examined why eight values per thread aligned one uint32 input load with one uint4 output store, while thirty-two values produced strided warp addresses in each store instruction.

This final article examined how every existing FP4 test used a small M, repeatedly selected decode GEMV, and left the production prefill dequantization path outside regression coverage.

#ONNXRuntime #NVIDIA #H200 #Hopper #NVFP4 #FP4 #CUDA #GPUProgramming #RegressionTesting #KernelDispatch #CodeAnalysis