Why can the same NVFP4 GEMM report 1,200 TFLOPS on one input and 1,700 on another?

We thought we had a kernel-efficiency problem. On our RTX 5090, top-dispatcher—a shape-based selection of CUTLASS configurations—reported about 1,146–1,338 TFLOPS across eight matrix shapes. Our instruction microbenchmark was near 2,000 TFLOPS. We looked at the gap and thought: where did a third of the hardware go?

That was an uncomfortable conclusion to revisit. We had compared two measurements without checking whether the GPU was running at the same clock.

Was the instruction itself being held back?

We started with a register-only microbenchmark, normalized by clock64. The NVFP4 block-scaled MMA (mma.sync … kind::mxf4nvf4 … m16n8k64) delivered 4,057 FLOP per SM per cycle—about 99% of the 4,096 we expected—at roughly 2.9 GHz. We did not observe an extra throughput limit on this instruction path.

The same test showed the familiar GeForce pattern on older paths: FP16 and FP8 with FP32 accumulation run at half the rate of their FP16-accumulate versions.

A separate instruction probe also included structured-sparse paths. We keep it separate here: its dense NVFP4 observation was 4,093.1 FLOP per SM per cycle, rather than the earlier 4,057. The sparse rows count dense-equivalent arithmetic, so their higher rates do not mean the same dense model automatically runs twice as fast.

We also inspected the listed tensors from a Qwen 27B NVFP4 weight set. The file reports that only 2.77% of the inspected groups already contain at least two zeros. Its reported extra-zeroing requirement is a structural check, not evidence that pruning would preserve model quality.

What clock was the GEMM actually running at?

This is where our original comparison fell apart. During repeated large GEMMs, the clock fell far below 2.9 GHz. The same 16,384³ kernel reported 1,203 TFLOPS at about 1.83 GHz with random inputs, and 1,734 TFLOPS at about 2.67 GHz with all-zero inputs. As far as the run records show, nothing else changed; how operands and scales were generated, and the starting state, are still to verify.

We left the power limit at 530 W and did not lock clocks. These runs show a large data-dependent power and frequency effect. They do not establish that every NVFP4 GEMM is power-limited: the Gaussian and zero conditions drew 481 W and 391 W.

Zero-filled and constant-filled matrices are diagnostic inputs, not model workloads or evidence for a performance claim. We include them because they expose something a single throughput number leaves out. Without stating the input values, a benchmark can report almost any number in this range.

Run length gave us another reason to be careful. A 4,096³ ping-pong kernel reported 1,264 TFLOPS over 40 calls and 1,144 over 4,000. Even 4,000 calls lasted under half a second. Neither measurement was a thermal steady state.

What did we get wrong about “57–66% of peak”?

We treated the gap as if it measured time lost to operand loads, synchronization or the epilogue. It did not. Dividing application throughput by a peak measured at a different frequency cannot tell us that.

Profiling helps us locate where warps are waiting, but those samples have their own denominator. The following two captures use different shapes and schedules; their sample shares cannot be read as fractions of elapsed kernel time.

For frequency-normalized efficiency, we need effective FLOPs divided by SM count × per-cycle peak × the clock integrated over the same measurement window. That requires clock telemetry from the same run; a separate profiler replay will not supply the missing denominator.

We also have to apply that scrutiny to our own dispatcher. Its 1.03–1.16× over torch._scaled_mm came from a benchmark that always timed our kernel first. We need a balanced-order retest before calling it a speedup. That is the next note.

How much can we conclude from these runs?

Each condition is one historical sequence on one card (n = 1), with low-rate nvidia-smi sampling and no locked clock. The results were enough to make us withdraw our original explanation. Strengthening the causal claim still needs repeats, time-based windows, logged throttle reasons, and a fixed-clock run at a frequency every input can hold.