Your 3090 Isn't Missing Anything for FP8/NVFP4 Weight-Only Inference

After I published a quantized model, a user with two RTX 5090s asked, in short: “If I use an INT4 quant, don’t I lose the NVFP4/FP8 advantage of Blackwell?”

My answer is: within weight-only (A16) quantization, you lose nothing. This post shows the evidence in the vLLM source, in the SASS of the shipped binary, and in measurements on an RTX 3090. It also takes apart a claim that comes up often, “the 3090 has no FP8/FP4 tensor cores, so FP8/NVFP4 models are slow on it”, and states which parts of it are correct and which are not.

The conclusion first. The comparison is between a real RTX 3090 and a hypothetical 3090 with FP8/FP4 tensor cores added. Everything else (memory bandwidth, SM count, clocks) is the same.

The only row where the 3090 loses is the last one. That row is not weight-only inference. It is a different quantization scheme that also quantizes the activations.

The hypothetical 3090 does not exist, so it cannot be measured directly. Instead, I measured these three points:

• All 480 A16 Marlin kernels built for the 5090 (sm_120) issue only BF16/FP16 HMMA. They contain no FP8/FP4 MMA instructions and no instructions that convert from FP8/FP4 (Section 2.1).
• In decode (M=1), Marlin uses 91–94% of DRAM bandwidth and only 13–33% of the instruction issue slots. FP8 and INT8, and NVFP4 and INT4, have different dequantization instructions but run at the same effective bandwidth (Section 4.3).
• In prefill (M=8192), Marlin runs at 95–97% of the speed of BF16 cuBLAS at the same clock. The total cost of weight-only, dequantization included, is 3–5% (Section 4.4).

1. “FP8 model” and “NVFP4 model” mean two different things

The phrase “an FP8/NVFP4 model” does not distinguish between these:

Most checkpoints on Hugging Face named “FP8” or “NVFP4” are built as the last two rows (W8A8 / W4A4). Blackwell’s FP8/FP4 tensor cores make a difference to speed only for those two rows.

When a 3090 loads a W8A8 / W4A4 checkpoint, vLLM does not quantize the activations and runs the model as weight-only (Section 3). An “FP8 model” running on a 3090 is always W8A16.

2. Tensor cores need both inputs in the same family of types

The input type combinations accepted by the PTX mma instruction are roughly:

There is no bf16 × e4m3 or bf16 × e2m1 combination. If the activations stay in BF16, the only option is to convert the weights to BF16 and issue a bf16 × bf16 mma. This holds on every GPU generation. As long as the activations are 16-bit, the FP8/FP4 tensor cores are not used, even when the hardware has them.

vLLM’s Marlin kernel is implemented exactly this way. For BF16 activations it issues this instruction (csrc/libtorch_stable/quantization/marlin/marlin_mma.h:71, vLLM b22afe45a):

"mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 "

The same file also contains an f32.e4m3.e4m3.f32 mma (lines 79 and 97). It is used only when the activation type (a_type) is FP8, which means W4A8-FP8. Below sm_89 it is removed at compile time (marlin_template.h:283).

2.1 SASS of the shipped binary

I also checked the binary that is actually shipped, not only the source. The target is _C_stable_libtorch.abi3.so in the vLLM 0.29.1rc1.dev47+gdc36fcce9 image. I disassembled every Marlin kernel in it with cuobjdump -sass and counted the types of MMA instructions. A 3090 (sm_86) executes the sm_80 binary.

• The A16 kernels for the 5090 issue only the same HMMA (BF16/FP16 tensor cores) as the ones for the 3090.
• The FP8 tensor core instruction QMMA appears only in kernels whose activations are FP8.
• The only conversion instruction in the A16 kernels is F2FP.BF16.PACK_AB (it packs the FP32 accumulators into the BF16 output). There is no instruction that converts FP8/FP4 weights.

Comparing one kernel with the same configuration (256 threads, W8A16 FP8), the instruction types are the same. The differences in counts come from compiler scheduling.

NVIDIA’s own implementation does the same. FlashInfer (0.6.18.post1), which vLLM uses, has BF16 × NVFP4 kernels for Blackwell. The SM100 version (dense_gemm_bf16_fp4_sm100.py) describes itself as follows:

The kernel decodes each 16-value E2M1 weight block and its E4M3 scale to BF16, then executes BF16 tcgen05 MMA with FP32 accumulation.

The GeForce Blackwell (SM12x) version (dense_gemm_bf16_fp4_sm12x.py) uses cute.nvgpu.warp.MmaF16BF16Op (an F16/BF16 MMA). Even W4A16 kernels written specifically for Blackwell do not use the FP4 tensor cores.

3. With A16, vLLM picks the same kernel on a 5090 as on a 3090

The code below is from vLLM 0.29.1rc1.dev47+gdc36fcce9, the version used for the measurements.

NVFP4 weight-only (NVFP4A16). SM100/SM103 (B200 etc.) use the FlashInfer CuTe-DSL kernel, and every other GPU uses Marlin (vllm/model_executor/kernels/linear/__init__.py:1081-1092). The 5090 (SM120) uses Marlin.

elif linear_backend == "auto" and use_a16:
    ...
    # Weight-only: prefer FlashInfer CuTe-DSL W4A16 on SM100/103,
    # otherwise Marlin.
    ...
    if compute_capability in (100, 103) and cutedsl_ok:
        force_kernel = FlashInferCuteDslNvFp4W4A16LinearKernel
    else:
        force_kernel = MarlinNvFp4LinearKernel

Section 2.1 showed that the FlashInfer kernel uses a BF16 MMA.

FP8 weight-only (W8A16). There are only two candidates, HummingFP8ScaledMMLinearKernel and MarlinFP8ScaledMMLinearKernel (__init__.py:475-479). Both run on SM75 and later. I did not read the internals of Humming for this post. But because of the instruction set constraint in Section 2, an A16 kernel has no choice other than a BF16/FP16 mma.

MoE. The fused MoE Marlin kernel (csrc/libtorch_stable/moe/marlin_moe_wna16/marlin_template.h) includes the same dequant.h and marlin_mma.h.

So for an A16 checkpoint, a 3090 and a 5090 execute the same kernel code. The speed difference between them comes from general GPU performance: memory bandwidth, SM count, clocks, cache size. It does not come from FP8/FP4 tensor cores.

When a 3090 loads a W8A8 / W4A4 checkpoint (compressed-tensors format), this is what happens:

• W8A8 FP8: the decision is made before kernel selection. The W8A8 scheme requires compute capability 8.9 or higher (compressed_tensors_w8a8_fp8.py:83-85). So on a 3090 the W8A16 scheme is used instead (compressed_tensors.py:840-857), and the layer runs on Humming or Marlin, the same as W8A16.
• W4A4 NVFP4: when I ran the support checks of the FP4 kernel candidates (__init__.py:544-) in order on a 3090, every candidate that uses FP4 tensor cores reported “not supported”, and the first one that passed was Marlin.

In both cases the activations are not quantized and stay in BF16.

4. Measuring the cost of dequantization

An A16 GEMM always includes a step that converts the weights to BF16/FP16 (dequantization). This step is the only place where the hypothetical 3090 could differ, so I measured its cost.

4.1 The Marlin implementation

Marlin does not expand the weights to BF16 in VRAM. For every GEMM, it loads a tile of weights from shared memory into registers, converts it there, and passes it to mma. The only one-time work at startup is a repack (a layout permutation).

The FP8 (E4M3) to FP16 conversion is (dequant.h:321-336, vLLM b22afe45a):

template <>
__device__ inline void dequant<half2, vllm::kFE4M3fn.id(), true>(
    int q, half2* frag_b) {
  constexpr int FP8_EXPONENT = 4, FP16_EXPONENT = 5;
  constexpr int RIGHT_SHIFT = FP16_EXPONENT - FP8_EXPONENT;
  constexpr int MASK = 0x7F007F00;

  int Out1 = (q & 0x80008000) | ((q & MASK) >> RIGHT_SHIFT);
  q <<= 8;
  int Out2 = (q & 0x80008000) | ((q & MASK) >> RIGHT_SHIFT);

  frag_b[1] = *reinterpret_cast<const half2*>(&Out1);
  frag_b[0] = *reinterpret_cast<const half2*>(&Out2);
}

• Four elements packed in one 32-bit integer become FP16 bit patterns using only and/shift/or.
• The exponent bias difference (2^8 for FP16, 2^120 for BF16) is multiplied into the scales at load time (marlin_utils_fp8.py:35). No multiply is needed per conversion.
• Every E4M3 and E2M1 value is exactly representable in FP16/BF16. The conversion does not round.
• The only architecture branch in dequant.h is __CUDA_ARCH__ >= 750 (line 70). The hardware FP8 → FP16 conversion instructions available on sm_89 and later are not used. This matches the SASS in Section 2.1.

4.2 Measurement setup

The NVFP4 weights are random values made with a vLLM test helper. This post covers speed only, not accuracy. For FP8 W8A16, vLLM’s first choice is Humming (Section 3). I measured only Marlin.

4.3 Decode (small M)

Time for one GEMM (μs) as M varies:

• FP and INT of the same bit width (FP8 and INT8, NVFP4 and INT4) have the same effective bandwidth at M=1, within 1.5%. Their dequantization instructions are completely different (next table), and that difference does not show up in decode speed.
• Marlin reads the weights at a higher effective bandwidth than BF16 cuBLAS.

Nsight Compute results for each kernel at M=1:

Instruction counts are per thread, divided by the number of weight elements (K × N).

• DRAM bandwidth determines decode time. Marlin uses 91–94% of the bandwidth. Meanwhile, even NVFP4, which has the most instructions, uses only 32.5% of the issue slots.
• Dequantization is 1.8–2.9 instructions per element. FP8 is 1.81, which matches the count from the source in 4.1 (7 instructions per 4 elements).
• Even if the hypothetical 3090 reduced the dequantization instructions with hardware conversion instructions, that would only leave more issue slots idle. The DRAM bandwidth, which determines the time, would not change.

Effective bandwidth at M=1 (GB/s) with the SM clock locked to reduce compute capacity:

The clocks in the “not locked” column differ by format because the GPU reached its 280 W power limit.

• At 1,200 MHz and above, the effective bandwidth of all four Marlin formats stays within 1%.
• At 900 MHz, the two 4-bit formats drop by 15–16%. But BF16 cuBLAS, which does no dequantization, drops by 16% at 1,200 MHz and by 36% at 900 MHz. Whether a kernel does dequantization does not determine its sensitivity to clock.
• Even at 900 MHz, the difference between NVFP4 and INT4 is 1.6%, and between FP8 and INT8 it is 1.3%.

4.4 Prefill (large M)

Nsight Compute results at M=8192. ncu locks the clock, so every format ran at 1.69 GHz.

• At the same clock, Marlin runs at 95–97% of the speed of cuBLAS. These 3–5% include dequantization and every other difference, such as the tile configurations of Marlin and cuBLAS. The cost of dequantization alone is smaller.
• Marlin’s tensor pipe utilization is at about the same level as cuBLAS.

Results with the 280 W power limit and no clock lock (M=8192, two runs):

• cuBLAS’s 41.3–41.9 TFLOPS/GHz is 99% of the 3090’s published figure (71 TFLOPS BF16 with FP32 accumulate at 1,695 MHz, which is 41.9 TFLOPS/GHz).
• Under the power limit, Marlin’s clock (especially for 8-bit) is about 10% lower than cuBLAS’s. As a result, the raw TFLOPS gap becomes 6–15%. This is because dequantization and memory reads consume power, and on a power-limited GPU it is a real cost. But as long as the activations are 16-bit, this power does not depend on whether FP8/FP4 tensor cores exist.

5. What the 3090 actually lacks

What it lacks is the option to choose W8A8 / W4A4.

• Prefill and large batches: where the GEMM is compute-bound, FP8 tensor cores generally have 2x the compute ceiling of BF16 on the same GPU, and FP4 has 4x. As shown in 4.4, A16 Marlin already runs close to the BF16 tensor core ceiling. The only way to go faster is to quantize the activations as well.
• Decode (small batch): it is limited by memory bandwidth, so A8/A4 does not make it faster. The extra step of quantizing the activations can even make it slower. vLLM’s Marlin allows W4A8-FP8 only on SM89 and SM12x, and its error message gives the reason: “It is slower than Marlin W4A16 on other devices” (marlin.cu:419, vLLM b22afe45a).
• Quality: W8A8 / W4A4 adds activation rounding error on top of the weight error. I did not measure the KLD of W4A4 for this post.

So “NVFP4 is fast on Blackwell” is not a comparison of the same model on two GPUs. It is the result of choosing a different quantization scheme, which trades activation error for prefill speed. As long as you use a weight-only checkpoint, Blackwell does not use this option either.

6. Expected objections and answers

6.1 “vLLM itself warns that performance degrades on a 3090”

When a 3090 loads an FP8/NVFP4 checkpoint and Marlin is selected, vLLM prints this warning (kernels/linear/nvfp4/marlin.py:34-39; for FP8, marlin_utils_fp8.py:110-115):

Your GPU does not have native support for FP4 computation but FP4 quantization is being used. Weight-only FP4 compression will be used leveraging the Marlin kernel. This may degrade performance for compute-heavy workloads.

The warning is correct. But what it compares against is W4A4 executed on FP4 tensor cores, and it is limited to “compute-heavy workloads” (prefill). It says the same thing as Section 5. It is not evidence that weight-only is slow on a 3090.

Also, this warning is printed unconditionally inside MarlinNvFp4LinearKernel.process_weights_after_loading. As shown in Section 3, NVFP4A16 uses Marlin on every GPU except SM100/SM103. According to the code, a 5090 loading an NVFP4A16 checkpoint also prints “Your GPU does not have native support for FP4 computation”. The warning is attached to the selected kernel, not to the capability of the GPU.

6.2 “Blackwell has FP8/FP4 → BF16 conversion instructions, so dequantization is faster”

As shown in 4.1 and 2.1, Marlin does not use those instructions. Even if it did, they could remove only part of the 1.8–2.9 dequantization instructions per element. In decode, more than 67% of the issue slots are idle, and DRAM bandwidth determines the time (4.3). In prefill, the entire cost of weight-only, dequantization included, is 3–5% (4.4).

6.3 “NVFP4 is 4x faster than BF16 on Blackwell”

That is the compute ceiling of a W4A4 GEMM. A weight-only checkpoint does not take that path on any GPU (Section 2). The FlashInfer W4A16 kernels that NVIDIA wrote specifically for Blackwell also use a BF16 MMA (2.1).

6.4 “In practice an NVFP4 model on a 5090 is far faster than on a 3090”

That comparison mixes two factors:

1. General GPU performance. Memory bandwidth alone is 1,792 GB/s on the 5090 and 936 GB/s on the 3090 (published figures).
2. The quantization scheme. The 5090 runs W4A4, and the 3090 runs W4A16.

To isolate the effect of the FP8/FP4 tensor cores, compare W4A16 and W4A4 on the same GPU. To compare across GPUs, run the same A16 checkpoint on both.

6.5 “FP8 on a 3090 is emulation, so it loses accuracy”

As shown in 4.1, the conversion from FP8/FP4 to FP16/BF16 is exact and does not round. The Marlin kernel on a 3090 is the same kernel a 5090 runs for the same A16 checkpoint (the sm_120 build in 2.1). In terms of accuracy, the scheme at a disadvantage is W8A8 / W4A4, which also rounds the activations.

6.6 “W4A4 loses very little quality, so Blackwell users should use A4”

This is outside the claim of this post. Choosing W4A4 is a trade between prefill speed and activation error, and it should be decided by measuring KLD with the same method. I did not make that comparison here. The only claim of this post is that once you choose weight-only, Blackwell’s FP8/FP4 tensor cores are not used on any GPU.

6.7 “4-bit decode gets slower when the clock is lowered. That proves dequantization matters”

As shown in 4.3, 4-bit Marlin gets 15–16% slower at 900 MHz. But BF16 cuBLAS, which does no dequantization, gets 36% slower under the same condition. At 1,200 MHz, Marlin does not change and cuBLAS is 16% slower. Dequantization does not explain how much a kernel slows down at low clocks. Also, comparing FP and INT (which have different dequantization instructions) at the same clock, the difference is within 1.6% at every clock.

6.8 “You only measured a 3090. It may be different on a 5090”

I did not measure on a 5090. What I can confirm:

• The sm_120 binary that a 5090 executes consists of the same types of instructions (HMMA and integer operations) as the 3090’s sm_80 binary (2.1).
• From the 5090’s published figures (170 SMs, 2.41 GHz boost clock, 1,792 GB/s), its instruction issue ceiling is about 2.9x the 3090’s, and its memory bandwidth is about 1.9x. With the NVFP4 instruction count from 4.3 (3.66 per element), issue slot usage at full bandwidth would be about 22%. Putting the 3090’s measured bandwidth utilization and clock into the same formula gives 32%, which matches the ncu measurement of 32.5%. The 5090 has more compute headroom, not less.

Summary

• FP8/FP4 tensor core instructions require both inputs to be FP8/FP4-family types. There is no BF16 × FP8 instruction. All 480 A16 Marlin kernels built for the 5090 issue only HMMA (BF16/FP16). The FlashInfer W4A16 kernels NVIDIA wrote for Blackwell also use a BF16 MMA.
• For A16, vLLM selects the same kernel on a 3090 and a 5090 (Marlin for NVFP4A16; Humming or Marlin for FP8).
• Measured on a 3090, Marlin in decode uses 91–94% of DRAM bandwidth and only 13–33% of the issue slots. FP and INT formats with different dequantization instructions run at the same effective bandwidth. In prefill, Marlin runs at 95–97% of the speed of BF16 cuBLAS at the same clock.
• What the 3090 lacks is the option to choose W8A8 / W4A4: the prefill compute ceiling (2x for FP8, 4x for FP4). That is a different quantization scheme that trades activation error for speed. It is not weight-only inference.
• vLLM’s “This may degrade performance” warning is about prefill compared with W4A4. According to the code, a 5090 prints the same warning for NVFP4A16.