ONNX Runtime #32128 — Why Eight FP4 Values per Thread Were Nearly Twice as Fast as Thirty-Two

ONNX Runtime NVFP4, Part 2 of 3 — How warp-contiguous memory access and one uint4 store defined the effective vector width

Part 1 examined why NVFP4 prefill on H200 was dominated not by matrix multiplication, but by the kernel that expanded packed 4-bit weights into FP16 or BF16.

The fallback path was:

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

For the recorded Qwen3.8-27B workload, the original dequantization kernel accounted for 46.9% of total prefill GPU time.

ONNX Runtime PR #32128 vectorized that kernel.

Previously, each thread processed two FP4 values.

The new path processes eight FP4 values per thread.

Before

One thread
→ one packed byte
→ two FP4 values

After

One thread
→ one packed 32-bit word
→ eight FP4 values

It may seem natural to widen the vector further.

If eight values are faster,
would thirty-two values be even faster?

The experiment produced the opposite result.

With the same indexing logic, a copy-only kernel using eight values per thread reached approximately:

3.9 TB/s

Widening the per-thread chunk to thirty-two values reduced the measured throughput to approximately:

1.9 TB/s

Nearly half.

The difference was not determined by how much total work each thread eventually completed.

It was determined by a lower-level question:

When one memory instruction executes, which addresses do the thirty-two lanes of the warp access together?

Vectorization on a GPU is not merely the act of giving one thread more values.

It is the act of arranging the addresses produced by the entire warp into a pattern the memory system can serve efficiently.

Eight FP4 values occupy exactly four packed bytes

An FP4 payload uses four bits per value.

Eight values therefore occupy:

8 values
×
4 bits
=
32 bits
=
4 bytes

The new kernel assigns one such K-axis chunk to each thread.

Conceptually:

uint32_t packed_word =
    packed_weights[row][chunk];

The implementation reads the packed input through a 32-bit pointer:

const uint32_t word =
    reinterpret_cast<const uint32_t*>(
        b_packed + row_offset
    )[chunk];

A CUDA warp contains thirty-two threads.

If each thread reads the next adjacent four-byte chunk, the warp covers:

32 threads
×
4 bytes
=
128 bytes

The lane addresses can be arranged like this:

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

All lanes participating in the load point into one continuous 128-byte input region.

Lane 0 ─┐
Lane 1  │
Lane 2  │
...     ├→ one contiguous packed region
Lane 31 ┘

This is the first reason eight is a useful unit.

The thread-local chunk size aligns naturally with the warp-level input layout.

The eight values expand into exactly one 16-byte store

After dequantization, each FP4 value becomes an FP16 or BF16 value.

Both output formats use two bytes per element.

8 values
×
2 bytes
=
16 bytes

CUDA’s uint4 type contains four 32-bit words:

4
×
4 bytes
=
16 bytes

The reconstructed eight-element output therefore fits exactly into one uint4 store per thread.

Conceptually:

reinterpret_cast<uint4*>(
    output + row_offset
)[chunk] = output_values;

Across the full warp:

32 threads
×
16 bytes
=
512 bytes

The lane addresses become:

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

Every lane writes the region immediately following the previous lane’s region.

Lane 0 ─┐
Lane 1  │
Lane 2  │
...     ├→ one contiguous 512-byte output region
Lane 31 ┘

The important property is not that the hardware necessarily services all 512 bytes as one physical transaction.

The important property is that the address sequence contains no logical gaps between lanes.

The memory system can split the region into the transactions it requires without first resolving a sparse, strided pattern.

FP4 dequantization writes four times more payload than it reads

The kernel expands a packed four-bit representation into a sixteen-bit representation.

Per thread:

Input

Eight FP4 values
→ 4 bytes

Output

Eight FP16/BF16 values
→ 16 bytes

The output payload is four times wider than the input payload.

At warp scale:

Packed input:
128 bytes

Expanded output:
512 bytes

The kernel’s store pattern therefore matters at least as much as its packed-load pattern.

Even a perfectly organized FP4 input load would not be enough if the four-times-larger output were written inefficiently.

The eight-element design creates a useful equivalence:

Eight reconstructed values
=
one 16-byte vector
=
one uint4 store per thread

That lets the output expansion preserve a contiguous warp-level address pattern.

What changes when one thread handles thirty-two values?

Now suppose one thread processes thirty-two FP4 values.

The packed input size would be:

32 values
×
4 bits
=
128 bits
=
16 bytes

That appears attractive.

A thread could load a wider packed vector.

But the reconstructed output is:

32 values
×
2 bytes
=
64 bytes

One uint4 store writes only 16 bytes.

The thread therefore needs four back-to-back stores:

Store 0
→ output values 0–7

Store 1
→ output values 8–15

Store 2
→ output values 16–23

Store 3
→ output values 24–31

From the viewpoint of one thread, this remains a contiguous 64-byte output.

One thread
→ writes bytes 0–63 of its own chunk

But warp coalescing does not combine all four future instructions before deciding how to service the first one.

Each memory instruction has its own lane-address pattern.

The first `uint4` store becomes strided across lanes

Under the thirty-two-element design, each thread owns a 64-byte output chunk.

During the first uint4 store:

Lane 0
→ bytes 0–15 of its 64-byte chunk

Lane 1
→ bytes 64–79 of the next chunk

Lane 2
→ bytes 128–143

Lane 3
→ bytes 192–207

The start address advances by 64 bytes between lanes.

Lane 0  →   0–15
             48-byte gap

Lane 1  →  64–79
             48-byte gap

Lane 2  → 128–143
             48-byte gap

One instruction therefore produces:

write 16 bytes

skip 48 bytes

write 16 bytes

skip 48 bytes

across the warp.

The second store later fills another part of each thread’s 64-byte chunk.

The third and fourth stores fill the remaining parts.

After all four instructions complete, the overall output region may be fully populated.

But each individual instruction still spans a sparse set of addresses.

Filling the gaps later does not make the first instruction contiguous

The thirty-two-element version effectively produces four separate warp address patterns.

Instruction 1
→ first 16-byte segment from each 64-byte thread chunk

Instruction 2
→ second 16-byte segment from each 64-byte thread chunk

Instruction 3
→ third 16-byte segment from each 64-byte thread chunk

Instruction 4
→ fourth 16-byte segment from each 64-byte thread chunk

Each instruction sees the same 64-byte spacing between lanes.

By contrast, the eight-element version produces:

One instruction
→ one adjacent 16-byte segment per lane
→ contiguous warp region

The comparison is:

FP4 values per thread	Expanded output per thread	`uint4` stores per thread	Lane-to-lane start-address spacing within one store
8	16 bytes	1	16 bytes
32	64 bytes	4	64 bytes

The thirty-two-element version gives each thread more work.

It also gives each store instruction a less favorable warp-level address pattern.

Thread-local contiguity is not warp-level contiguity

This distinction is easy to miss.

Thread-local view

The thirty-two-element thread writes:

64 consecutive bytes

That appears fully contiguous.

Warp-level view

During one store instruction, the warp writes:

16 bytes from Lane 0

then a 48-byte hole

then 16 bytes from Lane 1

then another 48-byte hole

The thread-local regions are contiguous internally.

The addresses generated by the warp instruction are strided.

Those are not the same property.

A thread can own a contiguous output range while the warp still emits a sparse memory instruction.

GPU vectorization needs to account for both levels.

The copy-only experiment isolated the address-pattern cost

The full dequantization kernel contains several kinds of work:

Load packed FP4 values

Decode the FP4 bit patterns

Read scales

Multiply values by scales

Calculate addresses

Write FP16/BF16 output

If the complete kernel alone were benchmarked, it would be difficult to know how much of the difference came from:

FP4 conversion
scale handling
integer arithmetic
memory access

The change therefore records a copy-only comparison using the same index structure.

The conversion and scaling work were removed so that the memory layout itself could be observed.

The H200 measurements were approximately:

Eight elements per thread

One uint4 store per thread
→ 3.9 TB/s

Thirty-two elements per thread

Four uint4 stores per thread
→ 1.9 TB/s

The performance gap remained even without the FP4 arithmetic.

That is strong evidence that the address pattern was not a secondary detail.

It was a primary design boundary.

Tile tuning could not repair the lane-address geometry

A natural reaction to a slow GPU kernel is to adjust:

block dimensions
warp count
tile size
occupancy
grid shape

Those parameters matter.

But they do not necessarily change the relationship between the addresses emitted by adjacent lanes in the same store instruction.

In the thirty-two-element layout:

Lane stride within one store
=
64 bytes

Changing the number of blocks does not automatically convert that into:

Lane stride within one store
=
16 bytes

The issue is embedded in how thread ownership maps to memory.

The patch description therefore notes that the approximately two-times bandwidth difference was not something ordinary tile tuning could recover.

The data layout and instruction layout had to agree first.

This does not mean smaller vectors are always faster

The result should not be generalized into:

Small vectors always outperform large vectors.

Nor does it mean:

Eight values per thread is optimal for every GPU kernel.

The result belongs to a specific format conversion:

Input:
4-bit packed values

Output:
16-bit FP16 or BF16 values

Thread ownership:
contiguous K-axis chunk

Output vector:
16-byte uint4

Warp width:
32 lanes

Eight is special here because:

8 FP4 inputs
=
4 packed bytes

and:

8 FP16/BF16 outputs
=
16 bytes
=
one uint4

which produces:

32 lanes
×
one adjacent uint4 per lane
=
one contiguous warp-level output region

The value eight was not chosen because smaller is inherently better.

It was chosen because the input width, output width, vector-store width, and warp width align at that point.

Thread-level vector width and warp-level vector width are different

Vectorization is often described as:

Scalar

One thread
→ one value

versus:

Vectorized

One thread
→ multiple values

That is only the first level of GPU execution.

A second level exists:

Thirty-two threads execute the same instruction as one warp.

Two widths therefore matter.

Thread-local vector width

How many values does one thread own?

Warp-level address width

How do the addresses from all thirty-two lanes align
during one memory instruction?

In this case:

Thread-local width = 8
→ warp stores are contiguous

while:

Thread-local width = 32
→ each store instruction is strided across lanes

Increasing the first width damaged the second.

GPU vectorization must be designed for the warp, not only for the thread.

One `uint4` store became the performance boundary

The output assignment can look like a minor C++ implementation detail:

reinterpret_cast<uint4*>(out)[chunk] =
    *reinterpret_cast<const uint4*>(values);

But that line depends on several contracts:

Does each thread produce exactly 16 output bytes?

Is the destination aligned appropriately?

Does the next lane own the immediately following uint4?

Does the eight-element chunk remain inside one row?

Do all participating lanes use the same store width?

When those conditions hold:

Lane 0’s uint4

then Lane 1’s uint4

then Lane 2’s uint4

form an adjacent sequence.

If one thread instead owns several uint4 segments, each individual store instruction sees a wider lane stride.

In this kernel, uint4 is not merely a convenient type.

It is the bridge between:

one thread’s reconstructed FP4 chunk

and:

the warp’s contiguous output layout

The new path also replaces software-emulated FP4 conversion

Memory access was not the only improvement.

The new kernel also changes how packed FP4 codes become FP16 or BF16 values.

The scalar path used:

__nv_cvt_fp4x2_to_halfraw2()

On pre-Blackwell hardware, that conversion can be software-emulated.

The recorded H200 SASS contained branches and a subnormal-normalization loop.

The vectorized path reuses the Fp4Cvt logic already used by the decode GEMV path.

It processes a 32-bit packed word by separating relevant magnitude and sign bits.

A simplified representation is:

uint32_t magnitude =
    packed_word & magnitude_mask;

uint32_t sign =
    (packed_word >> shift) & sign_mask;

The conversion then reconstructs four FP4 pairs through branch-free bit manipulation and prmt-based logic.

One packed word
        ↓
Separate sign and magnitude fields
        ↓
Decode four pairs
        ↓
Produce eight values

This avoids the value-sensitive branch and normalization sequence associated with the previous Hopper conversion path.

A four-bit format has a small, enumerable input domain

A single FP4 code has only sixteen possible bit patterns.

A packed byte containing two FP4 values has:

16 × 16
=
256 combinations

That limited domain makes exhaustive primitive-level validation practical.

The change compares the new Fp4Cvt result with the old conversion for all 256 packed-byte values.

The mismatch count was:

The check included the negative-zero representation associated with code 0x8.

This matters because a branch-free implementation is useful only if it preserves the exact representation contract.

The optimization was not allowed to change the bit pattern merely because the numerical value looked approximately equal.

Row calculation moved from division into the grid

The original scalar path could derive the row and K position from one linear index.

Conceptually:

row
=
linear_index / packed_row_width

position
=
linear_index % packed_row_width

For large matrices, the row calculation could involve a 64-bit division.

The new kernel uses a two-dimensional launch structure.

blockIdx.x
→ selects a K-axis chunk

blockIdx.y
→ selects a weight row

Its structure resembles:

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

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

The row identity is encoded directly in the launch geometry.

Threads no longer need to repeatedly recover it from a global linear index.

Scale-block lookup became an incremental state machine

NVFP4 groups values under block scales.

A straightforward implementation can calculate the relevant scale as:

scale_block
=
element_index / block_size

Repeating that division for every value or packed pair is expensive.

The vectorized kernel instead calculates the initial scale block once.

It also tracks how many pairs remain before the next boundary.

At chunk start:

determine current scale block

determine pairs remaining in the block

Then:

decode one FP4 pair

remaining pairs -= 1

if remaining pairs == 0:
    advance scale block
    reset remaining-pair count

Repeated division becomes:

one initial calculation
+
small increments at actual boundaries

The scale relation is preserved with less arithmetic.

The tensor-level scale is hoisted out of the row loop

NVFP4 also uses a tensor-level scale, represented here by weight_scale_2.

The new kernel reads it before iterating over assigned rows:

const float global_scale =
    *weight_scale_2;

The value can remain in a register while the thread processes multiple rows.

Before

Repeated global loads of the same scale

After

One load
→ local reuse

The optimization appears small when viewed for one thread.

Across billions of weights, repeated loads and calculations become system-level costs.

The even-`block_size` guard preserves the packed-pair scale contract

Two FP4 values share one packed byte.

The vectorized path processes them as a pair and assumes that they share one block scale.

When block_size is even, the scale boundary aligns with pair boundaries.

For example:

block_size = 16

Values 0–15
→ scale 0

Values 16–31
→ scale 1

The packed pairs are:

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

No byte crosses the scale boundary.

With an odd block size:

block_size = 15

one byte can contain:

Value 14
→ scale 0

Value 15
→ scale 1

The two nibbles in that byte require different scales.

The pair-level fast-path assumption is no longer valid.

The dispatch therefore requires:

block_size % 2 == 0

An odd block size uses the scalar fallback.

This is a correctness condition, not merely a performance heuristic.

The `K % 8 == 0` guard preserves complete chunks

The vectorized kernel assigns exactly eight K-axis values to each thread.

If K is divisible by eight:

K = 128

128 / 8
=
16 complete chunks

Every thread performs:

one complete 32-bit packed load

eight valid conversions

one complete 16-byte output store

If K is not divisible by eight:

K = 12

Chunk 0
→ values 0–7

Chunk 1
→ values 8–11 are valid
→ four positions are outside the row

Executing the same unmasked load and store could cross the logical row boundary.

A separate masked-tail implementation could have been developed.

This patch instead preserves the scalar fallback.

K % 8 == 0
→ vectorized fast path

K % 8 != 0
→ scalar path

The common case is accelerated without broadening the fast path beyond what its memory contract proves.

The two guards define the proof boundary of the optimization

The full condition is:

K % 8 == 0
+
block_size % 2 == 0

The first statement guarantees:

Every thread owns one complete eight-value chunk.

The second guarantees:

Both FP4 values in a packed byte share the same scale.

Together, they justify:

one aligned packed 32-bit load

four pair decodes

one scale at a time per pair

one complete uint4 output store

The guard is therefore not an arbitrary tuning choice.

It is the boundary inside which the vectorized memory and scale assumptions are valid.

The complete dequantization kernel became about four times faster

The copy-only experiment compared eight values against thirty-two and isolated the warp-store geometry.

The full kernel comparison is different.

It compares the new eight-value vectorized implementation with the original two-value scalar dequantization path.

For:

M = 1024
Output dtype = BF16
block_size = 16

the H200 measurements were:

`N = 4096`, `K = 4096`

Original scalar kernel:
60.7 μs

Vectorized kernel:
15.5 μs

Speedup:
3.93×

`N = 6144`, `K = 2048`

Original scalar kernel:
46.1 μs

Vectorized kernel:
12.1 μs

Speedup:
3.81×

`N = 2048`, `K = 6144`

Original scalar kernel:
46.2 μs

Vectorized kernel:
11.9 μs

Speedup:
3.88×

These two comparisons must not be mixed together.

Eight versus thirty-two values
→ copy-only memory-pattern experiment
→ about 3.9 versus 1.9 TB/s

New eight-value kernel versus original scalar kernel
→ complete dequantization comparison
→ about 3.8–3.9× faster

The approximately four-times kernel speedup includes:

memory-layout improvement
FP4 conversion changes
indexing changes
scale reuse
reduced integer arithmetic

The optimization was a collection of aligned contracts

The patch cannot be reduced accurately to:

Use a wider load.

It changed several layers together.

Per-thread work

2 FP4 values
→ 8 FP4 values

Input access

one-byte work
→ aligned 32-bit chunk
→ contiguous 128-byte warp region

Output access

two narrow stores
→ one uint4 store
→ contiguous 512-byte warp region

FP4 conversion

software-emulated intrinsic path
→ branch-free Fp4Cvt

Row calculation

linear-index division
→ blockIdx.y

Scale-block calculation

repeated division
→ initial state plus incremental boundary tracking

Tensor-level scale

repeated global access
→ local reuse

The changes reinforced one another.

Optimizing only one layer would have left other per-thread costs intact.

The faster path remained bitwise equivalent

The patch changed:

the width of each thread’s work
the address pattern
the FP4 conversion implementation
scale-index calculations

A numerical tolerance alone would provide limited evidence.

The implementation was therefore checked at several levels.

Exhaustive packed-byte comparison

Every possible packed pair was tested:

256 packed-byte values

The old and new conversions produced:

0 mismatches

Full output SHA-256 comparison

The complete operator output was hashed across eight shapes covering:

FP16 and BF16
block sizes 16 and 32
several K alignments
dimensions ranging from 128 to 5120

The before-and-after hashes matched.

Qwen-generated-token comparison

The Qwen3.8-27B test used:

8K prompt

128 generated tokens

MTP N = 3

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

The optimization therefore preserved the existing value contract from the conversion primitive through the final model output.

Eight values were a balance point for this layout

The value eight satisfies three useful relationships.

Packed input

8 FP4 values
=
4 bytes
=
one uint32 load per thread

Expanded output

8 FP16/BF16 values
=
16 bytes
=
one uint4 store per thread

Warp layout

32 lanes

→ 128 contiguous packed input bytes

→ 512 contiguous expanded output bytes

A smaller per-thread chunk might still produce contiguous addresses, but it would change instruction width and work distribution.

A sixteen-element chunk would produce:

32 output bytes per thread
→ two uint4 stores

Each store instruction would have a 32-byte lane stride.

A thirty-two-element chunk creates:

64 output bytes per thread
→ four uint4 stores

Each store instruction has a 64-byte lane stride.

Eight is the point where:

one thread
→ one wide output store

and:

one warp
→ one adjacent output sequence

hold simultaneously.

Address regularity mattered more than maximum per-thread work

GPU optimization often favors:

larger tiles
more values per thread
fewer loop iterations
fewer thread-level instructions

Those are useful principles.

They are not sufficient.

The thirty-two-element design may reduce some loop overhead.

The copy-only measurement shows that the memory-layout loss was larger.

Possible reduction in thread-local overhead
<
Loss from strided warp stores

The central design question was not:

How much work can one thread perform?

It was:

How can all lanes perform that work
while preserving an adjacent memory pattern?

Input vector width alone would have selected the wrong design

Thirty-two FP4 values occupy sixteen packed bytes.

From the input side alone, that may appear ideal:

One thread
→ one 16-byte packed load

But dequantization expands the output by four times:

32 FP4 values
→ 32 FP16/BF16 values
→ 64 output bytes

The output requires four separate 16-byte stores.

Optimizing only for the packed input would overlook the wider output path.

The correct vector width must consider both:

Packed input representation

and:

Expanded output representation

In this kernel, output traffic is larger.

The best input vector is not automatically the best full-path vector.

What this patch established

The source and recorded measurements support the following conclusions:

Eight FP4 values occupy four packed bytes.
Eight reconstructed FP16 or BF16 values occupy 16 bytes.
The fast path maps each thread to one 32-bit input load and one uint4 output store.
Adjacent lanes therefore read adjacent four-byte chunks and write adjacent 16-byte chunks.
The warp covers a contiguous 128-byte packed region and a contiguous 512-byte output region.
A thirty-two-element thread requires four uint4 stores.
During each such store, adjacent lane start addresses are separated by 64 bytes.
The copy-only H200 measurement was approximately 3.9 TB/s for the eight-element layout and 1.9 TB/s for the thirty-two-element layout.
The full vectorized dequantization kernel was approximately 3.8–3.9× faster than the original scalar kernel.
The fast path is limited to K % 8 == 0 and even block_size.

What this patch does not establish

The evidence does not prove that:

eight values per thread is optimal for every GPU or kernel
smaller thread-local vectors are generally better
every part of the thirty-two-element slowdown comes from one factor alone
the same bandwidth figures apply to every H200 system
every NVFP4 layout enters the vectorized path
odd block sizes can safely use the pair-level scale assumption
ragged K tails can use the unmasked eight-value contract
the scratch tensor itself has been eliminated
Blackwell native FP4 matmul should use the same vector width

The result belongs to the specific packed-FP4-to-FP16/BF16 conversion path and the recorded H200 tests.

Wider per-thread work was less important than contiguous collective work

The important numbers are not eight and thirty-two by themselves.

The important difference is the address pattern they create.

Eight FP4 values per thread

Output per lane:
16 bytes

Lane spacing:
16 bytes

Result:
warp-contiguous store

Thirty-two FP4 values per thread

Output per lane:
64 bytes across four instructions

Lane spacing within each instruction:
64 bytes

Result:
warp-strided stores

From the viewpoint of one thread, thirty-two is a wider vector.

From the viewpoint of the warp memory instruction, eight is the more contiguous vector.

That second viewpoint determined the measured bandwidth.

GPU vectorization is not the act of placing more values inside one thread. It is the act of making all lanes participating in one instruction form a coherent memory layout.

In ONNX Runtime #32128, eight values per thread were not a small arbitrary tuning choice.

They were the point where:

one 4-byte packed load

eight FP4 decodes

one 16-byte expanded store

thirty-two adjacent warp lanes

fit together.

Widening the thread’s ownership to thirty-two values increased local work.

It also broke the warp’s contiguous store pattern.

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

Related material

Patch status: Merged into ONNX Runtime main
Selected vector width: Eight FP4 values per thread
Packed input: Four bytes per thread; one contiguous 128-byte region per warp
Expanded output: One 16-byte uint4 per thread; one contiguous 512-byte region per warp
Thirty-two-element comparison: Four uint4 stores per thread, with 64-byte lane spacing in each store instruction
Recorded copy-only bandwidth: Approximately 3.9 TB/s versus 1.9 TB/s on H200
Fast-path conditions: K % 8 == 0 and an even block_size

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

Part 1 examined why H200 expanded packed NVFP4 weights into an FP16 or BF16 scratch tensor before calling cuBLAS, and why that conversion consumed nearly half of the recorded prefill GPU time.

This article examined why eight FP4 values mapped naturally to one uint32 input load and one uint4 output store, while widening the per-thread chunk to thirty-two values created strided warp addresses during every output instruction.

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

#ONNXRuntime #NVIDIA #H200 #Hopper #NVFP4 #FP4 #CUDA #GPUProgramming #MemoryBandwidth #Vectorization #CodeAnalysis