From 61250e89df56d5d292ad536111076e453dd2dd1c Mon Sep 17 00:00:00 2001 From: Mykhailo Dementii Date: Sun, 6 Sep 2026 21:25:42 +0200 Subject: [PATCH 1/2] perf(nvfp4): drop the guarded expf slow path from the fused SwiGLU epilogue nvfp4_linear_swiglu_w4a4_tma_kernel evaluates silu with the accurate expf and an accurate divide, then rounds the result to bf16 on the next instruction. The accurate expf compiles to a guarded slow path (64 CALL.REL.NOINC and 64 FCHK per kernel in SASS) and the divide to a Newton refinement; together they account for 652 FFMA that the bf16 rounding then discards. Use __fdividef(x, 1.0f + __expf(-x)) in that kernel only. The decode, small-T and non-TMA w4a4 SwiGLU routes keep the accurate silu. Operator benchmark (--policy a4, 34816x5120), two independent passes, smaller of the two, against a rebuild of identical master source as the false-alarm floor (0.00-0.18%): -5.37% at T=1024, -2.53% at T=4096, -1.79% at T=8192. Against the directly measured NVFP4 MMA ceiling of 2003.9 TFLOP/s the kernel goes from 51.0% to 51.9% of the tier at T=8192, on the useful FLOPs the benchmark itself counts. This tile's own ceiling, measured by an ablation arm with no compute, no operand fetch and no epilogue, is 1089.0 TFLOP/s = 54.3% of the tier; that arm belongs to a separate harness whose stock runs 0.54% faster than the binary measured here, so the two are quoted with the offset visible rather than divided into each other. Either way 3.9-4.5% of this kernel's time is left, and it needs a different tile, not a different epilogue. Kernel resources are unchanged, by kernel name and template arguments: REG 168, STACK 0, SHARED 1024, LOCAL 0, CONSTANT[0] 1424, identical on both sides, no spills. Inside the 27B NVFP4 model the same kernel is 28.5-29.0% of prefill kernel time and gains less than it does on the operator fixture. Profiled with nsys, two passes per chunk width, the median launch moves -1.03/-1.28% at chunk 1024, -0.79/-0.90% at 4096 and -0.84/-1.01% at 8192, against a floor of +0.04/+0.12% from a rebuild of identical source. Whole-prefill kernel time moves -0.16 to -0.48%. A prefill chunk width that is not a multiple of 256 takes the other route, and there the profile finds no instance of this kernel at all and the two arms differ by 0.013% over the whole run. The change is not bit-neutral, so the evidence is numeric. On the two operator cases that take the TMA route the FP64 oracle reports rel_l2 identical to master in all 17 significant digits, at 0.365 and 0.329 of the criterion limit; a one-ULP fp32 perturbation of the accurate activation already moves the same instrument in its eighth digit, and a four-level coarsening of the sigmoid fails it. Perplexity over 95641 scored tokens moves by -0.0034 nats against a floor of -0.0026 from the one-ULP perturbation, i.e. 1.29x the floor rather than inside it. What makes that reading uninterpretable as damage is that the instrument is not monotone here: a 64-level quantiser, far coarser than this change, moves it by only +0.0009. Interpretable departure starts at the 16-level build (+0.0071). Rebuilding identical source moves it by exactly 0.0000000000. ctest 104/104 on both arms from one build directory, 2 skipped for want of a groupwise 27B artifact, the same 2 on master. --- .../nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh | 17 +++++++++++++---- 1 file changed, 13 insertions(+), 4 deletions(-) diff --git a/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh b/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh index a54ef62d69..28cebb6b0f 100644 --- a/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh +++ b/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh @@ -12,6 +12,15 @@ namespace ninfer::ops::detail { +// SwiGLU epilogue activation. +// +// The accurate `silu` compiles to a guarded slow path for `expf` (64 +// CALL.REL.NOINC plus 64 FCHK per kernel in SASS) and a Newton-refined divide, +// and the result is then rounded to bf16 on the very next instruction. The +// hardware-approximate forms are ~1e-6 relative, four orders of magnitude below +// the bf16 quantum, so nothing that survives the rounding is affected. +__device__ __forceinline__ float swiglu_silu(float x) { return __fdividef(x, 1.0f + __expf(-x)); } + template struct Nvfp4LinearSwiGluTmaTensorStorage { static_assert(Schedule::kBlockN == 128); @@ -236,10 +245,10 @@ __global__ __launch_bounds__( shared_output + token1 * kOutputStride + pair_row); const auto& gate = accumulators[mma_m][mma_n]; const auto& up = accumulators[mma_m][mma_n + kGateMmaFragments]; - *destination0 = __floats2bfloat162_rn(silu(gate[0] * alpha) * (up[0] * alpha), - silu(gate[1] * alpha) * (up[1] * alpha)); - *destination1 = __floats2bfloat162_rn(silu(gate[2] * alpha) * (up[2] * alpha), - silu(gate[3] * alpha) * (up[3] * alpha)); + *destination0 = __floats2bfloat162_rn(swiglu_silu(gate[0] * alpha) * (up[0] * alpha), + swiglu_silu(gate[1] * alpha) * (up[1] * alpha)); + *destination1 = __floats2bfloat162_rn(swiglu_silu(gate[2] * alpha) * (up[2] * alpha), + swiglu_silu(gate[3] * alpha) * (up[3] * alpha)); } } From dc108d4392142acf3dca4d7c0154ef227637105b Mon Sep 17 00:00:00 2001 From: MichaelDementii <136074657+MichaelDementii@users.noreply.github.com> Date: Wed, 9 Sep 2026 23:42:25 +0200 Subject: [PATCH 2/2] fix(nvfp4): keep the approximate SiLU off the zeroing range of __fdividef __fdividef returns zero once the divisor reaches 2^126, which 1 + __expf(-x) does for x below -87.34, where SiLU is still a normal bf16. Fold the exponential onto the side that cannot overflow: the divisor stays in (1, 2] for every finite x and the only subnormal the form can produce enters a multiply. Still no CALL and no FCHK, and the operator is unchanged within +0.14%. Co-Authored-By: Claude Opus 5 --- .../nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh | 14 +++++++++++++- 1 file changed, 13 insertions(+), 1 deletion(-) diff --git a/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh b/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh index 28cebb6b0f..9dc612a5f7 100644 --- a/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh +++ b/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh @@ -19,7 +19,19 @@ namespace ninfer::ops::detail { // and the result is then rounded to bf16 on the very next instruction. The // hardware-approximate forms are ~1e-6 relative, four orders of magnitude below // the bf16 quantum, so nothing that survives the rounding is affected. -__device__ __forceinline__ float swiglu_silu(float x) { return __fdividef(x, 1.0f + __expf(-x)); } +// +// The exponential is folded onto the side that cannot overflow. `__fdividef` +// returns zero once the divisor reaches 2^126, which `1 + __expf(-x)` does for +// x below -87.34, and SiLU there is still a normal bf16 (-9.6e-37 at that +// edge). Written this way the divisor stays in (1, 2] for every finite x, and +// the only subnormal the form can produce, `e`, enters a multiply rather than +// the divide. It costs nothing: both forms compile to the same 40 SASS +// instructions with two MUFU and no CALL. +__device__ __forceinline__ float swiglu_silu(float x) { + const float e = __expf(-fabsf(x)); + const float r = __fdividef(1.0f, 1.0f + e); + return (x >= 0.0f ? x : x * e) * r; +} template struct Nvfp4LinearSwiGluTmaTensorStorage {