From e7a2706513ba1b9fbcfb27bab42e7eec9fa4153a Mon Sep 17 00:00:00 2001 From: Chris Lee Date: Tue, 21 Jul 2026 03:31:03 -0600 Subject: [PATCH] attn: native::exp in decode softmax (6.7% attn at 33k ctx) sycl::exp -> sycl::native::exp for the two exp() calls per token in the FlashDecoding scan's online-softmax recurrence. Inputs are bounded softmax diffs (score - m_new, max_s - m_new) -- the same numerical class the M1.7c fused FA kernel (qxmx_fa_fused.cpp:284) already uses native::exp for. Measured on B50 at 33k context (QXMX_PROFILE=2): attn 122.8 -> 114.6 ms (-6.7%), total decode step 186.7 -> 178.7 ms. ~12.7M exp calls/step at that depth, so the cheaper intrinsic matters on the scan's critical path. Correctness: qxmx_diff mean|diff|=0.1064 (was 0.1111 on this B50; B70 reference 0.108) -- slightly better, gate green, greedy argmax unchanged. --- src/qxmx_gpu.cpp | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/src/qxmx_gpu.cpp b/src/qxmx_gpu.cpp index 8cc4019..3ca825f 100644 --- a/src/qxmx_gpu.cpp +++ b/src/qxmx_gpu.cpp @@ -581,8 +581,15 @@ void gpu_attn_softmax(sycl::queue& q, const float* q_vec, float dot = ql[lid] * k_val; float score = sycl::reduce_over_group(it.get_group(), dot, sycl::plus()) * scale; float m_new = sycl::fmax(max_s, score); - float f = sycl::exp(max_s - m_new); - float e = sycl::exp(score - m_new); + /* native::exp: the M1.7c fused FA kernel (qxmx_fa_fused.cpp:284) + * uses the fast hardware exp for the softmax here; this decode + * path is the same recurrence and the inputs (score - m_new, + * max_s - m_new) are bounded softmax diffs, so native is safe. + * Was sycl::exp -- 2 of these per token, ~12.7M calls/step at + * 33k ctx, so the cheaper intrinsic matters on the scan's + * critical path. */ + float f = sycl::native::exp(max_s - m_new); + float e = sycl::native::exp(score - m_new); sum_s = f * sum_s + e; /* Phase 1 V dequant — T13 deferred WHT: fold vscale and * vnorm_corr into the pre-WHT dequant; skip the butterfly -- 2.51.2