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<float>()) * 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