perf(nvfp4): drop the guarded expf slow path from the fused SwiGLU epilogue

[MichaelDementii]

Base origin/master @ 487f8977. Rebased onto master 487f8977; the diff applies to it cleanly, with the same diff on its own base as the control. The measurements below were taken on ad0f3d38, which was master when the work was done, 83 commits back. I have not re-taken them here and I say so rather than implying they are fresh.

What

nvfp4_linear_swiglu_w4a4_tma_kernel computes its activation with the accurate
silu(x) = x / (1 + expf(-x)) and then rounds the result to bf16 on the very next
instruction. The accurate expf compiles to a guarded slow path and the accurate divide to
a Newton refinement, and both are paid for a value that is about to lose 16 mantissa bits.

This replaces the activation in that one kernel with the hardware-approximate forms,
__fdividef(x, 1.0f + __expf(-x)). No other SwiGLU route is touched: the decode, small-T
and non-TMA w4a4 kernels keep the accurate silu, and they serve as the control below.

SASS of nvfp4_linear_swiglu_w4a4_tma.cu.o, nvcc 13.1, -gencode arch=compute_120a,code=sm_120a:

instruction	before	after
CALL.REL.NOINC	64	0
FCHK	64	0
FFMA	652	0
FMUL	192	448
MUFU.EX2	64	64
MUFU.RCP	65	64
F2FP.BF16.F32.PACK_AB	32	32

Why it is worth doing

The kernel is feed-bound. Nine NINFER_ABL arms on the same base, same operator benchmark, two
passes each, from the campaign that motivated this change
(§3.2 of its write-up; its cells are two families of per-arm CSVs,
A_a*.csv and the matching B_a*.csv, which live outside this repository and are not files of
this package — supplied on request; that harness is a separate set of binaries from the ones
measured below, and its stock arm runs 0.54% faster than the e0 binary of this package):

• removing the epilogue alone (arm a7, the whole consumer inner loop still in place) takes
  5.73% off at T=8192: 2842.6 us -> 2679.8;
• removing the consumer inner loop as well — 24 ldmatrix, 20 scale reads and 64 OMMA per
  k-tile (arm a6) — gives 2681.9 us, 0.08% slower than a7. All of the compute and all of
  the shared-memory operand fetch are hidden behind the feed and cost nothing.

So the epilogue is the one non-zero item that is not the global-to-shared feed, and it is worth
162.8 us of the kernel's 2842.6. Of that, the accurate expf plus the accurate divide are
53.3 us; the remaining 109 us are the shared-memory staging and the store, which this change
does not touch and cannot.

That 162.8 us is a fixture figure and is data-dependent. On near-zero accumulators the same
epilogue costs 31.8 us instead — the slow expf path is not taken on zeros — so the ratio of
this change to the kernel is not a constant of the kernel. That is one of the reasons the model
gains less than the fixture, and the model numbers are in "End-to-end" below.

Measurements

ninfer_nvfp4_linear_swiglu_bench --policy a4 --t-sweep 1024,4096,8192 --warmup 3 --repeat 25,
RTX 5090 (sm_120a), CUDA 13.1, driver 580.159.03, 525 W cap.
Two independent passes with the arms interleaved; the table quotes the smaller of the two.

T	master	rebuild of identical master source	this branch	delta
1024	419.840 us	419.840 us (+0.00%)	397.312 us	-5.37%
4096	1458.176 us	1458.176 us (+0.00%)	1421.312 us	-2.53%
8192	2857.984 us	2863.104 us (+0.18%)	2806.784 us	-1.79%

The middle column is the instrument's own false-alarm floor: the same source rebuilt into a
different binary. It is 0.00-0.18%, i.e. below every delta reported.

Claim withdrawn 2026-09-05. An earlier revision said here that the same measurement had
been taken on a second RTX 5090 in another container of the same host and gave
-5.89 / -2.38 / -1.87%, reproducing the table to within 0.52 percentage points. Those three
numbers exist as one line of campaign notes and nowhere else: there is no raw file for them in
this package or in the campaign, and the stand they were taken on has been reassigned, so they
cannot be re-read. A number that cannot be read in a raw file is not quoted here, whichever
direction it points. The evidence for this table is the table itself and its own rebuild
floor, on one card. No second-card claim is made.

Roofline, on the fixture above (the model's own numbers are in the next section). Against
ceilings measured directly on this card (the NVFP4 MMA tier at
2003.9 TFLOP/s from the product register form of
mma.sync.aligned.kind::mxf4nvf4.block_scale..., two runs agreeing to 0.63%; L2 read at
7132-7461 GB/s). Useful FLOPs are 2 * 34816 * 5120 * T, which is what the benchmark's own
TFLOP/s column reports, and the two arms below are the two arms of the table above:

T	master, TFLOP/s	% of the MMA tier	this branch, TFLOP/s	% of the tier
1024	869.55	43.4%	918.86	45.9%
4096	1001.45	50.0%	1027.42	51.3%
8192	1021.90	51.0%	1040.54	51.9%

The 1027.42 in the 4096 row is a genuine coincidence with a number that appears below, and it is
how an error got into an earlier revision of this description: the ablation harness's stock arm
also reads 1027.4 TFLOP/s, at T=8192. They are different cells of different binaries.

The feed is unchanged (the epilogue reads no global memory), so the share of the L2 read ceiling
stays at 53-56% — the kernel moves 3988 GB/s against an L2 read plateau of 7132-7461 GB/s, both
from the same campaign, §3.3.

The kernel's own ceiling at this tile is 1089.0 TFLOP/s = 54.3% of the tier, measured by an
arm with no compute, no shared-memory operand fetch and no epilogue. That arm is not in this
package: it comes from the ablation harness of the campaign that motivated this change (arm a6
of that campaign, which is my own work outside this repository), and that harness is a
different binary — its own unmodified-master arm runs T=8192 in 2842.6 us against the
2857.984 us of the e0 binary measured here, a 0.54% offset, which is three times the
false-alarm floor of the speed instrument. So the two must be quoted with that offset visible:

• inside the ablation harness, where numerator and denominator come from the same binary, stock
  takes 94.3% of the kernel ceiling (1027.42 of 1089.01, i.e. 2681.856 us of 2842.624);
• this branch's 1040.54 TFLOP/s against the same ceiling is 95.5%, and master's 1021.90 is
  93.8% — but that pair crosses the 0.54% offset and is worth about that much less. Taking the
  offset back out, by scaling this branch's 2806.784 us with the 2842.624/2857.984 ratio of the
  two stock arms, puts it at 96.1% of the ceiling. That 96.1% is a numerical coincidence
  with a figure disproved below and arrived at a different way; they are not the same claim.

So what is left in this kernel is 3.9-4.5% of its time depending on which of the two
conventions is used, and it cannot be had without changing bytes per FLOP, i.e. without
changing the tile.

Provenance: those ceilings are direct measurements of the instructions in their product
register form on this card, two independent runs agreeing to within 0.2-0.63%. That is more
honest than quoting a constant compiled into a benchmark, but it is our own measurement, not
a vendor figure.

Resources. Identical on both arms, compared by kernel name and template arguments
(Nvfp4GemvGeometry<34816, 5120>, Nvfp4W4a4TmaSchedule<256, 3, 1>):
REG:168 STACK:0 SHARED:1024 LOCAL:0 CONSTANT[0]:1424. No spills either way.

Re-dumped from the submitted branch 2026-09-05 with
cuobjdump -res-usage, on the single translation unit
src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cu built with the project's own
flags (nvcc 13.1.115, -O3 -std=c++20 --generate-code=arch=compute_120a,code=[compute_120a,sm_120a],
taken verbatim from compile_commands.json). Raw:
raw/fmt_res/resusage_base.txt and raw/fmt_res/resusage_cand.txt. The TU instantiates exactly
one kernel; the Function line and the resource line that follows it are byte-identical on the
two sides (md5 d7d1389002533c25da7a5e9b19afba71 over both lines on either side). What does
change is the object: 199 568 bytes on master, 151 984 with this change, which is the removed
slow path.

The cand tree the dump was taken in holds blob
2be97137e7354c52398edf9447986fca470fd4ae for the changed file — the same blob as the submitted
commit, so this is the submitted code and not a rebuild of the same patch; the branch-location
note of this package records the chain.

End-to-end, on the model this kernel serves

The operator benchmark above is a fixture. The numbers below are the same kernel inside
Qwen3.6-27B NVFP4 driven through ninfer_bench (-p 16384, --kv-dtype fp8,
--max-ctx 20480, -r 7 --warmup 2), on the same card.

The gain is smaller in the model than on the fixture, and this is stated up front.
nsys cuda_gpu_kern_sum, two independent passes per chunk width, arms interleaved and the
order reversed in the second pass:

prefill chunk	share of prefill kernel time	median launch in pass A, master -> branch	delta, pass A / pass B	whole-prefill kernel time
1024	28.5% -> 28.2%	419518.0 -> 415198.0 ns	-1.030% / -1.278%	-0.218% / -0.483%
4096	29.0% -> 28.8%	1617556.0 -> 1604709.0 ns	-0.794% / -0.900%	-0.158% / -0.312%
8192	28.5-28.6% -> 28.3%	3210494.0 -> 3183503.5 ns	-0.841% / -1.010%	-0.240% / -0.346%

The false-alarm floor of this instrument is measured on the same target: a rebuild of
identical master source profiled the same way moves the median by +0.119% and +0.041%
and whole-prefill kernel time by +0.196% and +0.010%.

Two things about this table that a reviewer listing raw/e2e/prof/ would find, so they are said
here.

There is a third 4096 triple in that directory and it is not in the table. Files E0, E1,
ENUL without an A_/B_ prefix are an earlier pilot run, taken before the witness sampler was
started (raw/e2e/prof/card_nsys.samples covers only the A_/B_ profiles). It is excluded for
that reason and not for its numbers, so its numbers are given: median -1.387%, kernel sum
-0.830%, and whole-run kernel time +0.354% — the opposite sign to the two witnessed passes —
against a rebuild floor on the same triple of -0.193% / -0.229% / -0.233%, which is worse than the
witnessed floor in every column. That is what an unwitnessed window looks like on this host, and
it is the reason the table quotes the other two.

The 1024 and 8192 profiles ran under the exclusive card lock but without a per-arm witness
sample. raw/e2e/nsys_c1024.sh starts no sampler; only the 4096 script does. The lock is the
same one that refused five wall-clock arms outright, so the arms were not sharing the card, but
the per-window verdict that backs every other table in this description is absent for those four
profiles. Re-taking them under a sampler is a cheap fix and has not been done.

Against the fixture's -5.37 / -2.53 / -1.79% at the same three values of T, the model gives
-1.0…-1.3 / -0.8…-0.9 / -0.8…-1.0%. The fixture overstates this change by 1.9x to 4.4x,
most at small T. The epilogue's cost is data-dependent — on near-zero accumulators it is
31.8 us instead of 162.8 us — so a fixture number should not be read as a model number, and
the number to plan with is the one in the table above.

Control that the change cannot touch, proven by profile rather than asserted. A prefill
chunk width that is not a multiple of 256 leaves the TMA route (resolve_route requires
tokens >= 256 && tokens % 256 == 0). At chunk 3968 the profile contains no instance of
this kernel at all — the work goes through silu_and_mul_dim0_split_kernel — and the two
arms differ by +0.013% over the whole run (6186757431 vs 6187537720 ns). The decode
column of the same benchmark, which uses a different kernel, moves -0.030 / -0.029 /
-0.061% across three passes.

Wall clock, for completeness, and it does not resolve this change. Prefill throughput,
three independent passes, arm order permuted per pass:

chunk	pass 1	pass 2	pass 3	same three passes, rebuild of identical source
1024	+0.080%	+0.493%	+0.372%	-0.251% / +0.182% / -0.221%
4096	+0.265%	+0.427%	+0.427%	-0.208% / +0.160% / +0.141%
8192	+0.512%	+0.174%	+0.843%	+0.292% / -0.423% / +0.314%
3968 (control)	-0.085%	+0.245%	+0.370%	-0.179% / +0.034% / +0.134%

Every reading is in the right direction, but the false-alarm floor of the wall clock spans
-0.42% to +0.31%, which is the size of the effect. Wall-clock throughput on this benchmark
cannot grade a change of this size, and it is quoted here only to show it does not
contradict the profile. The claim rests on the kernel-time table, whose floor is ten times
smaller.

All 36 timed arms ran under an exclusive card lock with a per-arm witness sampling at 1 s;
all 36 windows are clean, and each window starts after the lock is taken. Five arms that the
lock refused outright because a foreign process held the card were re-taken rather than
patched over; both affected triples were re-taken whole, so every triple sits in one
contiguous window.

Correctness evidence

The change is not bit-neutral by construction, so a bit-exact gate is not offered as
evidence. Two numeric instruments are, each with its own control.

Numeric oracle (FP64), quoted as a margin, not as a verdict.
ninfer_linear_swiglu_nvfp4_test with NINFER_OP_REPORT_STATS=1 reports rel_l2 against
the criterion limit. Exactly two of its ten cases take the TMA route (NVFP4_A4 T=256 and
T=1024); on both, this branch reports the same rel_l2 as master to all 17 significant
digits (0.36488986484368158 and 0.32930753840228488). That is an aggregate, not a
byte-for-byte comparison, so the claim is exactly this: an instrument that resolves one fp32
ULP sees no difference.

The test fails on whichever of the two reported ratios crosses 1 first, and on this operator
that is gross_ratio, not rel_l2_ratio: the four-level build below fails with
gross_ratio 1.1736 while its rel_l2_ratio is still 0.9586. On the ratio that actually
decides, this branch reports exactly master's value — gross_ratio 0.48562588625839476 at
T=256 and 0.37263683208756154 at T=1024, digit for digit — so it consumes none of the
headroom master already has, which is 2.06x and 2.68x respectively.

That headroom is not a distance to failure, and it is quoted here only as headroom.
gross_ratio does not respond monotonically to this class of perturbation: coarsening the
sigmoid to 256 levels moves it down to 0.4627, to 16 levels down to 0.4110, and 64 levels
leaves it at master's value to all digits — and then 4 levels jumps it to 1.1736. So "2.06x"
does not mean the change could be 2.06 times larger and still pass. What the column does support
is the comparison actually being made, which is between builds at one ratio: master, the rebuild
and this branch are identical on it, and the only build that fails is the one four levels coarse.
rel_l2_ratio is the finer of the two readings — it is the one that resolves a single fp32 ULP —
but it is not the one that fails first here, which is why both are given.

The oracle is shown to be capable of resolving less than this change and of failing.
Builds that are never submitted, differing from master only in that same activation:

build	rel_l2 / limit, T=256	rel_l2 / limit, T=1024	test
master	0.36488986484368158	0.32930753840228488	pass
rebuild of identical source	0.36488986484368158	0.32930753840228488	pass
this branch	0.36488986484368158	0.32930753840228488	pass
master silu scaled by one fp32 ULP	0.36488986484368158	0.32930724077522605	pass
sigmoid coarsened to 256 levels	0.36177961657739593	0.32769740138392539	pass
sigmoid coarsened to 64 levels	0.37027962314467949	0.33414825872610487	pass
sigmoid coarsened to 16 levels	0.39752263648331654	0.36839657445992646	pass
sigmoid coarsened to 4 levels	0.95857446380370759	0.87328270048861767	FAIL

A one-ULP perturbation of the accurate result already moves the instrument in its eighth
digit; this branch does not move it at all. The four-level build fails the test
(gross_ratio 1.1736 against a limit of 1).

Control inside the same table. The eight cases that do not take the TMA route are
byte-identical across every build above, including the failing one. The instrument moves
where the change is and nowhere else.

End-to-end perplexity, with its own noise floor measured.
ninfer-perplexity <27B NVFP4 artifact> --text <405 KB of English prose> --context 4096 --stride 2048, 95 641 scored tokens per build.

build	mean_nll	delta vs master, nats
master	0.3484259516	-
rebuild of identical source	0.3484259516	+0.0000000000
sigmoid coarsened to 64 levels	0.3493674121	+0.0009415
master silu scaled by one fp32 ULP	0.3458178740	-0.0026081
this branch	0.3450712370	-0.0033547
sigmoid coarsened to 256 levels	0.3449070169	-0.0035189
sigmoid coarsened to 16 levels	0.3554921430	+0.0070662
sigmoid coarsened to 4 levels	0.5192349552	+0.1708090

Rebuilding the identical source into a different binary moves mean_nll by exactly
0.0000000000, so the instrument is deterministic across builds and the whole spread below is
numerical, not run-to-run.

The band is built from the two controls only, and this branch sits just outside it. The
controls are the rebuild (0.0000000000) and the one-ULP fp32 perturbation of the accurate
activation, which cannot matter and moves mean_nll by -0.0026081. That is the measured
floor: 0.0026 nats. This branch moves it -0.0033547, which is 1.29x the floor — inside
the same order of magnitude, in the same direction, and not inside the band.

What makes that reading uninterpretable as damage is that the instrument is not monotone in the
size of the perturbation over this range: a 64-level sigmoid — vastly coarser than anything
this change does — moves it by only +0.0009, about a third of what one ULP moves it, and in
the opposite direction. An instrument that ranks one ULP above a 64-level quantiser is not
resolving perturbation size at 0.003 nats. Monotone, interpretable departure begins at the
16-level build (+0.0071) and is unmistakable at 4 levels (+0.1708).

So the honest statement is: mean_nll does not separate this change from a one-ULP rounding
difference, and it rules out anything as coarse as a 16-level sigmoid. It does not, on its own,
certify the change; the FP64 oracle above does that, and this is the end-to-end corroboration
that nothing large is happening.

The bit-exact output gate is not cited at all, and neither is needle-in-a-haystack
retrieval; on a change of this kind neither instrument grades anything.

Throughput is not cited as correctness.

Tests

ctest from a single build directory, both arms, run serially with
NINFER_QWEN3_6_27B_NVFP4_WEIGHTS and NINFER_QWEN3_6_35B_A3B_WEIGHTS set:
104/104 on each (102 run and passed, 2 skipped for want of a groupwise 27B artifact).
The real-model tests need to be run one at a time: two of them co-scheduled do not fit in
32 GB of device memory, and that is true of master as well.

clang-format --dry-run on the changed file reports 2 diagnostics on this branch and the
same 2 on the base commit; the change adds none. Re-run 2026-09-05 in two separate worktrees
of the same repository, both sides real files on disk — reading a blob through a pipe does not
find the repository .clang-format and invents hundreds of diagnostics. Raw:
raw/fmt_res/cf_base.log and raw/fmt_res/cf_cand.log, clang-format 23.1.0, the same
.clang-format (md5 2b1cd100967456f321f6dd5ef76b8244) present in both trees.

Both diagnostics are the same two lines of alignment in an untouched block, and they simply
move by the 9 lines this change inserts above them:

side	line	text
base	195, 196	const int row_mod32 = scale_pair_row & 31; / const int row_quartile = scale_pair_row >> 5;
branch	204, 205	the same two lines

By the added-lines rule this is a zero: the lines this change adds are 15-23 and 248-251, and
neither diagnostic falls on any of them. Nothing here is reformatted.

🤖 Generated with Claude Code

[chatgpt-codex-connector]

Codex Review Summary

This comment shows the latest Codex review activity on this pull request.

Review	Status	Commit	Review trigger
📝 Code Review	✅ Completed 2026-09-09T21:50:54.180752Z	dc108d4	Manual request

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review" or "@codex security review".

Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings.

[MichaelDementii]

Re-measured on b88c0f6. The arm is this branch merged into today's master; ctest 114/114 on both
sides.

Operator, --policy a4, four passes with the arm order mirrored between them:

T	256	512	1024	2048	4096	8192
kernel time	−9.99 %	−12.13 %	−11.69 %	−8.97 %	−5.64 %	−2.74 %

24 of 24 cells faster, median −9.23 %. The −5.37 / −2.53 / −1.79 % in the description were taken
on ad0f3d38; please read these instead.

The mechanism claim also checks out on this base, exactly: CALL.REL.NOINC in the binary goes
76 → 12 and FCHK 8548 → 8484 — 64 of each, which is what the accurate silu was costing per
kernel.

@codex review

[chatgpt-codex-connector]

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 61250e89df

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

[chatgpt-codex-connector]

Avoid the large-divisor zeroing range of __fdividef

When the scaled gate is roughly between -87.34 and -88.72, __expf(-x) is finite but lies in [2^126, 2^128), where CUDA documents __fdividef as returning zero instead of the correctly rounded quotient. For example, x=-87.5 should produce SiLU near -8.7e-37, which is still representable in BF16, but this helper collapses it to zero; this also makes the TMA route numerically inconsistent with the other routes for valid finite inputs. Use an overflow-safe negative-side formulation such as x * __expf(x) / (1 + __expf(x)) while retaining approximate intrinsics.

AGENTS.md reference: AGENTS.md:L83-L88

Useful? React with 👍 / 👎.

[MichaelDementii]

The finding is right and is fixed in dc108d43. The suggested formula is not the fix, though — it
puts back exactly what this PR removes.

Probed on the card, x from −95 to −75 in steps of 0.005, against a double reference:

form	returns 0 where SiLU is a normal bf16
this branch before the fix	every x ≤ −87.34 (your window)
the accurate silu it replaces	every x ≤ −88.72 — its denominator overflows to inf
x * __expf(x) / (1 + __expf(x))	nowhere
committed form	nowhere

So the routes did disagree, on [−88.72, −87.34], and below that the base is wrong too. What the
suggestion costs, measured in this kernel rather than argued: / is a correctly-rounded divide, so
Nvfp4LinearSwiGluW4a4Tma... goes 1360 → 2232 instructions with 64 CALL and 64 FCHK — the same
64 of each the description reports removing.

Committed instead, keeping the divisor away from 2^126 rather than the numerator:

const float e = __expf(-fabsf(x));
const float r = __fdividef(1.0f, 1.0f + e);
return (x >= 0.0f ? x : x * e) * r;

__fdividef sees 1 over a divisor in (1, 2] for every finite x, and the only subnormal the form can
produce enters a multiply. In the kernel: 1232 → 1360 instructions, CALL and FCHK stay 0, MUFU
unchanged at 128, every other kernel in the binary byte-identical. Max relative error on [−20, 20]
is 9.2e-7 and 0 of 2001 points differ after rounding to bf16.

Cost: ninfer_nvfp4_linear_swiglu_bench --policy a4 --t-sweep 256,1024,4096,8192, two passes with
the arm order mirrored — median +0.03 % over the eight cells, range +0.00 … +0.14 %; the
kernel is feed-bound, so the extra 128 instructions do not show. ninfer_linear_swiglu_nvfp4_test
passes and its OP_ERROR_STATS are byte-identical on both arms.

@codex review

[chatgpt-codex-connector]

Codex Review: Didn't find any major issues. Breezy!

Reviewed commit: dc108d4392

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".