ONNX Runtime #32128 — Why Dequantizing 4-Bit Weights Became the Most Expensive Kernel in H200 Prefill

ONNX Runtime NVFP4, Part 1 of 3 — The dequantize → scratch → cuBLAS path that expands compressed weights before computation

FP4 is supposed to make AI computation lighter.

An FP16 weight requires 16 bits of storage.

An FP4 payload requires only 4 bits.

For the same model, that can substantially reduce weight storage and memory traffic.

On hardware that can consume FP4 values directly, lower precision can also increase arithmetic throughput.

FP4 is therefore often summarized like this:

Smaller weights
+
Less memory traffic
+
Higher arithmetic throughput
→ Faster AI

But the actual software execution path is not always that simple.

The fact that a GPU can store weights in a 4-bit representation does not mean the current workload can feed those 4-bit values directly into matrix multiplication.

ONNX Runtime PR #32128 exposed that distinction clearly.

When Qwen3.8-27B with NVFP4 weights was executed on an H200, the largest prefill bottleneck was not the cuBLAS matrix multiplication.

It was the kernel that expanded the packed FP4 weights back into FP16 or BF16.

Under the recorded workload, that dequantization kernel accounted for 46.9% of total prefill GPU time.

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

The result was structurally counterintuitive:

FP4 reduced weight storage.

But expanding those weights before computation
cost more than the matrix multiplications themselves.

The bottleneck in 4-bit inference was not 4-bit arithmetic.

It was the transition from the compressed FP4 representation into the wider format required by the fallback GEMM path.

H200 could not use the native SM120 FP4 matmul path

ONNX Runtime’s MatMulBlockQuantizedFp4Weight operator can follow several execution paths depending on matrix shape, hardware generation, and configuration.

A simplified dispatch structure looks like this:

Small decode workload
→ fused GEMV path

Blackwell native conditions
→ native FP4 matmul path

Other prefill workloads
→ dequantize the full FP4 weight
→ create an FP16/BF16 scratch tensor
→ call cuBLAS GEMM

The H200 belongs to the Hopper generation.

It does not satisfy the native SM120 Blackwell FP4 matmul path.

Prefill also processes multiple prompt tokens together, so it does not normally use the small-M decode GEMV route.

The H200 NVFP4 prefill path therefore became:

Packed NVFP4 weights
        ↓
Expand into an FP16 or BF16 [N, K] scratch tensor
        ↓
Run cuBLAS GEMM
        ↓
Produce the output

The weights remain compact while stored.

Two FP4 values can be packed into one byte.

But the cuBLAS fallback consumes FP16 or BF16 operands.

The entire weight matrix must therefore be reconstructed in a wider representation before matrix multiplication can begin.

Four-bit weights expand by at least four times before GEMM

Consider the payload width alone.

FP4 payload
→ 4 bits per value

FP16 or BF16 scratch
→ 16 bits per value

The dequantized payload is four times wider.

NVFP4 also includes block-scale and tensor-level scale metadata, but the central expansion remains:

Read compact FP4 weights
        ↓
Read scale metadata
        ↓
Apply the scales
        ↓
Convert values into FP16 or BF16
        ↓
Write a much larger scratch tensor

The matrix multiplication then reads that scratch tensor again.

The full memory path includes:

Read packed FP4 payload

Read scale metadata

Write FP16/BF16 scratch

Read FP16/BF16 scratch in cuBLAS

The model may occupy less storage, yet execution still creates a large temporary representation and an additional memory round trip.

Part of FP4’s storage advantage is spent again during format conversion.

Preparing the weights cost more than multiplying them

The workload recorded in PR #32128 was:

Model:
Qwen3.8-27B NVFP4

MatMulBlockQuantizedFp4Weight nodes:
168

Weights held by those nodes:
14.97 billion values

Prompt:
8K tokens

Generated tokens:
128

Multi-token prediction:
N = 3

Hardware:
NVIDIA H200

Before the patch, the dequantization kernel consumed:

Dequantization time:
1,750 ms

Total prefill GPU time:
3,732 ms

The ratio was:

1,750 / 3,732
≈ 46.9%

Almost half of prefill GPU time was not spent performing the final matrix multiplications.

It was spent preparing the weights so that those matrix multiplications could run.

The change description states that the dequantization kernel took more time than every cuBLAS GEMM in the model combined.

FP4 discussions often begin with tensor-core throughput.

In this workload, however, the real bottleneck was:

Native FP4 arithmetic throughput
→ not the dominant issue

Expanding packed FP4 into a wider operand
→ dominant issue

This exposes a broader rule of low-precision execution:

A compressed format being small does not mean the full path that consumes it is inexpensive.

The original kernel processed only two FP4 values per thread

The original DequantizeNvFp4Kernel assigned one packed byte to each thread.

One packed byte contains two FP4 codes.

One thread
        ↓
Load one packed byte
        ↓
Extract two FP4 values
        ↓
Apply the scale
        ↓
Store two FP16 or BF16 values

The structure was simple.

Across a very large weight matrix, however, several costs accumulated.

The patch description identifies the following characteristics:

One-byte load

Two separate two-byte stores

Integer division for row calculation

Integer division for scale-block calculation

A global tensor-scale load per thread

Software-emulated FP4 conversion on pre-Blackwell hardware

Small loads and separated stores limited memory efficiency

A thread loading one byte and performing two separate two-byte stores does not naturally create a wide and regular memory operation.

Across a warp, the runtime repeatedly performs:

Tiny individual loads

Multiple narrow stores

Per-thread address calculations

The data is contiguous at the logical tensor level.

The instructions used to move that data are still narrow and fragmented.

That distinction matters when the kernel processes billions of values.

Each thread repeated integer divisions

The scalar kernel used a linear index to determine:

which weight row the thread belonged to
which position inside the row it handled
which scale block applied to those values

Conceptually:

row
=
linear_index / row_width

and:

scale_block
=
element_position / block_size

Integer division is more expensive than simple addition, shifting, or masking.

One or two divisions are not necessarily important.

Repeating them across billions of weight elements is different.

A small per-thread cost can become a model-level bottleneck.

The same tensor-level scale could be loaded repeatedly

NVFP4 reconstruction uses both local block scales and a higher-level tensor scale.

The tensor-level value, called weight_scale_2 in this path, is shared across many values.

The scalar implementation could nevertheless load it repeatedly at thread granularity.

Many threads

Same tensor-level scale

Repeated global-memory loads

The information is shared.

The execution path did not fully exploit that reuse.

FP4 conversion was software-emulated on Hopper

The old kernel used a conversion based on:

__nv_cvt_fp4x2_to_halfraw2()

On Blackwell, FP4 receives more direct architectural support.

On pre-Blackwell hardware such as H200, this conversion can lower into a software-emulated instruction sequence.

According to the patch description, the resulting SASS included branches and a subnormal-normalization loop.

The Hopper path therefore combined:

Narrow memory operations
+
Repeated integer division
+
Repeated scale loading
+
Software-emulated FP4 conversion

The fact that the original values occupied only four bits was not enough to offset those costs.

The new kernel assigns eight FP4 values to each thread

PR #32128 introduces a vectorized fast path for common NVFP4 prefill layouts.

The fast path is selected when:

K % 8 == 0

and

block_size is even

Under those conditions, ONNX Runtime uses:

DequantizeNvFp4Vec8Kernel

Otherwise, the original scalar kernel remains as the fallback.

The per-thread work changes from:

Before

One thread
→ two FP4 values

to:

After

One thread
→ eight FP4 values

Eight FP4 values occupy exactly four packed bytes:

8 values
×
4 bits
=
32 bits
=
4 bytes

Each thread can therefore load one 32-bit packed word.

Across a 32-thread warp:

32 threads
×
4 bytes
=
128 contiguous bytes

The address pattern can become:

Lane 0  → bytes   0–3
Lane 1  → bytes   4–7
Lane 2  → bytes   8–11
...
Lane 31 → bytes 124–127

The warp reads one continuous packed region.

The expanded output becomes one 16-byte store per thread

After dequantization, eight FP4 values become eight FP16 or BF16 values.

8 values
×
2 bytes
=
16 bytes

Sixteen bytes fit exactly into one uint4 store.

The new kernel can therefore write one vector per thread.

Across the warp:

32 threads
×
16 bytes
=
512 contiguous bytes

The output pattern becomes:

Lane 0  → bytes    0–15
Lane 1  → bytes   16–31
Lane 2  → bytes   32–47
...
Lane 31 → bytes 496–511

The important property is not simply that each thread writes more data.

It is that all lanes participating in the same store instruction write adjacent regions.

Packed input:
128 contiguous bytes per warp

Expanded output:
512 contiguous bytes per warp

Thread-local vector width and warp-level memory layout now align.

The conversion path became branch-free

The vectorized kernel does not use the old software-emulated conversion path in the same way.

It reuses the Fp4Cvt mechanism already present in the decode GEMV implementation.

This path decodes packed FP4 codes through branch-free bit manipulation and prmt-based lookup logic.

Conceptually:

Packed 32-bit word
        ↓
Separate magnitude and sign bits
        ↓
Decode four FP4 pairs
        ↓
Produce eight FP16/BF16-compatible values

The conversion no longer depends on a value-sensitive branch and normalization loop for every packed pair.

The domain is also small.

Two packed FP4 values form one byte, so there are only:

256 possible packed-byte combinations

The implementation can exploit that limited representation space.

The row identity moved into `blockIdx.y`

The original scalar path derived a row from a linear index.

The new kernel uses a two-dimensional grid.

blockIdx.x
→ selects an eight-element K chunk

blockIdx.y
→ selects a weight row

The structure is conceptually:

chunk =
    blockIdx.x * blockDim.x
    + threadIdx.x;

for (row = blockIdx.y;
     row < N;
     row += gridDim.y) {
    ...
}

The row is no longer reconstructed through repeated division of a global linear index.

This removes one of the expensive index calculations from the common path.

Scale-block tracking became incremental

Each group of NVFP4 values shares a block scale.

A straightforward implementation can repeatedly calculate:

scale_block
=
element_index / block_size

The vectorized kernel instead computes the starting block once and keeps track of how many packed pairs remain before the next scale boundary.

Determine the initial scale block

Determine how many pairs remain in that block

Decode one pair

Decrease the remaining-pair counter

When it reaches zero:
→ advance to the next scale
→ reset the counter

Repeated division becomes a small state machine.

The calculation occurs only when crossing a scale-block boundary.

The tensor-level scale is loaded once and reused

The new kernel reads weight_scale_2 into a local value outside the row loop.

Conceptually:

const float global_scale =
    *weight_scale_2;

The compiler can keep that value in a register while the thread processes multiple rows.

Before

Repeated global-scale loads

After

One load
→ register-local reuse

Again, the individual optimization is small.

Applied to the complete weight matrix, it contributes to a much larger effect.

Why the block size must be even

Two FP4 values share one packed byte.

The vectorized path assumes that both values in the byte also share the same block scale.

When block_size is even, a scale boundary cannot divide 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)

Each pair remains inside one scale block.

With an odd block size, however, the boundary can split a packed byte.

block_size = 15

Value 14
→ scale block 0

Value 15
→ scale block 1

The two nibbles in one byte now require different scales.

The vectorized pair-level assumption no longer holds.

The fast path therefore requires:

block_size % 2 == 0

Odd block sizes continue through the scalar kernel.

The optimization does not silently broaden its correctness assumptions.

It states the conditions under which those assumptions are valid.

Why `K` must be divisible by eight

Each thread handles exactly eight K-axis values.

When K is divisible by eight, every thread owns a complete chunk.

K = 128

128 / 8
=
16 complete chunks

If K is not divisible by eight, the last chunk is incomplete.

K = 12

Thread 0
→ values 0–7

Thread 1
→ values 8–11 are valid
→ four positions would be outside the row

The vectorized kernel expects:

one complete 32-bit packed load

eight valid decoded values

one complete 16-byte output store

Using that contract on an incomplete tail could read or write beyond the logical row.

A masked tail path could have been added.

This patch instead keeps the existing scalar fallback:

K % 8 == 0
→ vectorized kernel

K % 8 != 0
→ scalar kernel

The fast path remains narrow and provable.

The kernel became approximately four times faster

The patch reports H200 measurements for:

M = 1024
Output dtype = BF16
block_size = 16

`N = 4096`, `K = 4096`

Scalar:
60.7 μs

Vectorized:
15.5 μs

Speedup:
3.93×

`N = 6144`, `K = 2048`

Scalar:
46.1 μs

Vectorized:
12.1 μs

Speedup:
3.81×

`N = 2048`, `K = 6144`

Scalar:
46.2 μs

Vectorized:
11.9 μs

Speedup:
3.88×

The gain did not come from one isolated technique.

It combined:

Wider per-thread work

Warp-contiguous input loads

Warp-contiguous output stores

Branch-free conversion

Fewer integer divisions

Incremental scale tracking

Register reuse of the tensor-level scale

The optimization reorganized the full data path.

Model-level time to first token fell by 32.5%

The original Qwen3.8-27B workload spent:

1,750 ms

inside the dequantization kernel.

After the vectorized path, that fell to:

409 ms

Total prefill GPU time changed from:

Before:
3,732 ms

After:
2,394 ms

End-to-end time to first token changed from:

Before:
3,876 ms

After:
2,618 ms

The recorded reduction was approximately:

32.5%

The kernel improvement therefore propagated through the system:

Faster dequantization
        ↓
Lower total prefill GPU time
        ↓
Lower user-visible TTFT

This is the concrete version of a broader FP4 question.

The important question is not only:

Does the stack support FP4?

It is:

How much overhead remains in creating, expanding, moving, and consuming FP4 data?

Decode throughput did not change

The patch does not accelerate every NVFP4 execution phase.

It targets the prefill fallback:

FP4 weight dequantization
→ FP16/BF16 scratch
→ cuBLAS GEMM

Decode uses a separate fused GEMV path for small M.

The reported results therefore show:

Prefill TTFT
→ improved

Decode throughput
→ unchanged

MTP acceptance
→ unchanged

It would be inaccurate to say that the patch accelerated all token generation.

The precise claim is:

It removed a major dequantization bottleneck from the Hopper NVFP4 prefill fallback.

Bitwise equivalence was checked at three levels

Changing the memory layout and the FP4 conversion implementation raises a critical question:

Does the faster kernel produce exactly the same values?

The patch verifies this at three levels.

All 256 packed-byte values

Two FP4 values fit into one byte.

Every possible packed input is therefore in the range:

0 through 255

The new Fp4Cvt output was compared with the previous conversion across all 256 combinations.

The number of mismatches was:

The comparison included the -0.0 representation associated with code 0x8.

Full operator-output hashes

The complete output tensor was hashed across eight configurations covering:

FP16 and BF16
block sizes 16 and 32
aligned and non-aligned K cases
N and K values from 128 to 5120

The SHA-256 outputs before and after the optimization were identical.

Generated-token hashes

The Qwen3.8-27B workload was also executed with:

8K prompt

128 generated tokens

MTP N = 3

The generated-token SHA-256 remained unchanged across three runs per comparison arm.

The evidence chain was:

Conversion primitive
→ all possible packed-byte inputs

Operator
→ full tensor hashes across multiple shapes

Model
→ final generated-token hash

The patch did not merely replace the old result with a numerically similar approximation.

It preserved the conversion contract while changing how that contract was executed.

Existing FP4 tests did not reach the prefill kernel

FP4 operator tests already existed before this change.

But the existing cases all used:

M <= 8

That shape selected the decode GEMV path.

M <= 8
→ fused decode GEMV

The following path was therefore not executed:

Full FP4 weight dequantization
→ [N, K] FP16/BF16 scratch
→ cuBLAS GEMM

The production workload spent nearly half of its prefill GPU time in a kernel that the existing FP4 test matrix did not enter.

The new tests use:

M > 8

so that the decode GEMV path is skipped and the intended prefill fallback is reached.

They also distinguish:

Vectorized conditions
→ new kernel

Odd block size
→ scalar fallback

K not divisible by eight
→ scalar fallback

That dispatch blind spot is the subject of Part 3.

The central lesson is:

Testing the relevant dtype is not the same as testing the relevant production dispatch path.

After optimization, the kernel reached the memory-bandwidth boundary

The patch estimates that one full dequantization pass over the Qwen3.8-27B weights moves approximately:

34.86 GiB

At roughly 3 TB/s, the theoretical transfer time is:

about 12.5 ms

The measured cost after optimization was:

about 12.4 ms per pass

The kernel had moved close to the bandwidth cost of the required data movement.

Before the patch, the bottleneck included:

Address calculation overhead

Conversion overhead

Narrow load/store behavior

Repeated scale operations

After the patch, the dominant cost became:

Moving the data that the fallback path inherently requires

That is an important performance boundary.

Before:
Instruction and access-pattern overhead dominated

After:
Mandatory memory traffic dominated

Further improvement while preserving the same fallback structure may now be limited.

The next questions become architectural:

Does the full [N, K] scratch need to be materialized?

Can dequantization and GEMM be fused more deeply?

Can the runtime move to a native FP4 path?

Can the scratch write-and-read round trip be eliminated?

The same FP4 model has different costs on different GPU generations

Blackwell can use a native FP4 execution route under the appropriate conditions.

Hopper may need to expand the same NVFP4 weight into a wider type first.

Blackwell native path

Packed FP4
→ native FP4 matmul

Hopper fallback path

Packed FP4
→ dequantize
→ FP16/BF16 scratch
→ cuBLAS GEMM

The model format can be identical.

The execution economics are not.

A software stack must manage different fallback contracts for each target architecture.

This is why FP4 cannot be understood as a dtype name alone.

FP4 performance depends not only on how many bits store the values, but on how far the target hardware can consume that representation natively.

Small weights and low latency are different properties

This patch separates three ideas that are often combined.

Storage efficiency

How small are the model file and resident weights?

NVFP4 can provide a strong advantage.

Compute efficiency

How quickly can the hardware perform FP4 arithmetic?

This matters most when a native FP4 path is available.

Conversion efficiency

When the backend cannot consume FP4 directly,
how cheaply can it transform the data into a usable operand?

On the H200 prefill fallback, the third category dominated system latency.

Weight storage:
4-bit payload

Actual GEMM operand:
16-bit

Conversion between the two:
almost half of prefill GPU time

The conclusion:

Smaller dtype
→ automatically faster model

does not hold.

The more useful questions are:

How long does the tensor remain in FP4 form?

Where is it expanded?

Is the expanded representation materialized?

Does another kernel have to read it again?

What this patch directly changed

PR #32128 directly changed:

NVFP4 weight dequantization in the pre-Blackwell prefill fallback
the common case where K % 8 == 0
the common case where block_size is even
per-thread processing from two FP4 values to eight
packed input access from byte-sized work to aligned 32-bit chunks
output writes to one 16-byte vector per thread
FP4 conversion to the branch-free Fp4Cvt path
row selection through a two-dimensional grid
scale-block tracking through incremental state
reuse of the tensor-level scale
test coverage for actual prefill shapes

What this patch does not prove

The patch does not establish that:

every GPU achieves a 3.8–3.9× dequantization speedup
every NVFP4 model receives a 32.5% TTFT reduction
Blackwell’s native FP4 path has the same bottleneck
FP4 is always faster than FP8 or BF16
decode throughput improves
the scratch tensor itself has been eliminated
every block size and every K shape uses the vectorized path
the H200 measurements generalize to all Hopper deployments

The reported results belong to the documented model, hardware, and execution configuration.

Additional workloads require additional evidence.

The real competition in the four-bit era appears in the conversion path

The full sequence behind the original bottleneck was:

Store the model in NVFP4
        ↓
Run prefill on H200
        ↓
Native SM120 FP4 matmul is unavailable
        ↓
Expand the entire weight into FP16/BF16
        ↓
Write the scratch tensor
        ↓
Read it again through cuBLAS
        ↓
Dequantization consumes 46.9% of prefill GPU time

The patch reorganized that path:

Tiny per-thread accesses
        ↓
Eight-value vectorized chunks

Repeated division
        ↓
2D grid and incremental scale tracking

Software-emulated conversion
        ↓
Branch-free lookup

Repeated scale load
        ↓
Register reuse

The recorded result was:

Dequantization kernel:
about 3.8–3.9× faster

Total prefill GPU time:
3,732 ms → 2,394 ms

Time to first token:
3,876 ms → 2,618 ms

FP4 performance is not completed by one FP4 instruction.

The way compressed values are read, scaled, reconstructed, stored, and passed into the next kernel can determine the performance of the entire system.

Related material

Patch status: Merged into ONNX Runtime main
Target path: NVFP4 prefill fallback when the native SM120 FP4 matmul path does not apply
Primary validation hardware: NVIDIA H200
Vectorized fast-path conditions: K % 8 == 0 and an even block_size
Scalar fallback: Odd block_size or K % 8 != 0
Recorded model-level result: Approximately 32.5% lower TTFT for Qwen3.8-27B with an 8K prompt
Paths not directly changed: Decode GEMV and the Blackwell native FP4 matmul path

This is Part 1 of a three-part series on ONNX Runtime’s NVFP4 prefill path.

Part 2 examines why assigning eight FP4 values to each thread produced warp-contiguous loads and stores, while widening the per-thread chunk to 32 values cut the measured copy bandwidth roughly in half.

Part 3 examines why the existing FP4 tests all used M <= 8, repeatedly selected the decode GEMV path, and never executed the prefill dequantization kernel that dominated the production workload.

#ONNXRuntime #NVIDIA #H200 #Hopper #NVFP4 #FP4 #CUDA #Quantization #Prefill #TTFT #CodeAnalysis