From 918d4d0a4363726c5700a76bcc64534de7c8102b Mon Sep 17 00:00:00 2001 From: Chris Lee Date: Sun, 19 Jul 2026 11:55:33 -0600 Subject: [PATCH] M1.7b.13: naive DeltaNet barrier reduction -- 599 tok/s (+1.3% over M1.7b.12) Cut per-token barriers in the naive recurrence from 12 to 3: - Removed the redundant pre-L2-norm barrier: each thread reads only its own ql[col]/kl[col] (written by itself) for the reduce_over_group input; reduce_over_group handles its own synchronization. - Replaced the 7-step rmsnorm sumsq tree reduction (7 barriers) with a single reduce_over_group + 1 SLM broadcast via pr (1 barrier). Reused pr for the sumsq so the lss local_accessor is gone. - Removed the redundant end-of-token barrier: the next iteration's first SLM op writes each thread's own ql/kl slot, and the first cross-lane read is guarded by the L2-norm broadcast barrier. Per-token barriers: 12 -> 3. SSM phase: 1.2 -> 1.1 ms/layer (-8%). Prefill (bench/session_4mb.txt): 8.95s -> 8.82s = 599 tok/s (was 591). Bit-exact correctness preserved (deltanet_naive_test 8/8 PASS, abs ~3e-6). All tests pass; output unchanged (' Paris. The'). --- src/qxmx_deltanet_naive.cpp | 32 ++++++++++++++++++-------------- 1 file changed, 18 insertions(+), 14 deletions(-) diff --git a/src/qxmx_deltanet_naive.cpp b/src/qxmx_deltanet_naive.cpp index b52d0bf..c90638a 100644 --- a/src/qxmx_deltanet_naive.cpp +++ b/src/qxmx_deltanet_naive.cpp @@ -51,11 +51,9 @@ void gpu_deltanet_batched(sycl::queue& q, q.submit([&](sycl::handler& h) { /* ql/kl: L2-normalized q,k for this v-head's key head, shared across * all HD columns via SLM (each thread reads its col). pr holds the - * (decay, beta_s) pair computed by lane 0. lss is the rmsnorm sumsq - * reduction buffer. */ + * (decay, beta_s) pair AND the rmsnorm sumsq (broadcast by col==0). */ sycl::local_accessor ql(sycl::range<1>(HD),h), kl(sycl::range<1>(HD),h); sycl::local_accessor pr(sycl::range<1>(2),h); - sycl::local_accessor lss(sycl::range<1>(LWS),h); h.parallel_for(sycl::nd_range<1>(NVH*LWS, LWS), [=](sycl::nd_item<1> it){ int hv=it.get_group(0), hk=hv%NKH, col=it.get_local_id(0); float* S = d_state + (size_t)hv*HD*HD; @@ -68,13 +66,18 @@ void gpu_deltanet_batched(sycl::queue& q, const float* qk_t = d_qk + (size_t)t * qk_s; const float* v_t = d_v + (size_t)t * v_s; /* Load raw q,k for this key head into SLM, L2-normalize via - * group reduce (identical to M1.5 / the decode oracle). */ + * group reduce (identical to M1.5 / the decode oracle). No + * barrier before the reduce: each thread reads only its own + * ql[col]/kl[col] (written by itself), and reduce_over_group + * handles its own synchronization. */ ql[col]=qk_t[hk*HD+col]; kl[col]=qk_t[NKH*HD+hk*HD+col]; - it.barrier(sycl::access::fence_space::local_space); double qq=(double)(ql[col]*ql[col]), kk=(double)(kl[col]*kl[col]); double qs=sycl::reduce_over_group(it.get_group(),qq,sycl::plus()); double ks=sycl::reduce_over_group(it.get_group(),kk,sycl::plus()); ql[col]*=1.0f/sycl::sqrt((float)(qs+eps)); kl[col]*=1.0f/sycl::sqrt((float)(ks+eps)); + /* Barrier: the ql/kl SLM writes above must be visible to all + * lanes before pass 1/2 read other lanes' values via kl[row]/ + * ql[row]. Then col==0 computes (decay, beta_s) into pr. */ it.barrier(sycl::access::fence_space::local_space); if(col==0){float apb=d_alpha[(size_t)t*NVH+hv]+d_dtb[hv]; float sp=apb>20?apb:(apb<-20?sycl::exp(apb):sycl::log1p(sycl::exp(apb))); @@ -100,17 +103,18 @@ void gpu_deltanet_batched(sycl::queue& q, acc+=s*ql[row]; } float out_val = scale_qk*acc; - /* Folded group_rmsnorm: sumsq reduction over this head's HD - * output values (one per column). Same reduction as M1.5. */ - lss[col] = (double)out_val * out_val; + /* Folded group_rmsnorm: hardware reduce_over_group for the + * sumsq (one op, no 7-step tree), then broadcast inv via pr. + * 2 barriers instead of 8. */ + double ssq = (double)out_val * out_val; + double sumsq = sycl::reduce_over_group(it.get_group(), ssq, sycl::plus()); + if(col==0) pr[0] = (float)sumsq; it.barrier(sycl::access::fence_space::local_space); - for (int s = LWS / 2; s > 0; s >>= 1) { - if (col < s) lss[col] += lss[col + s]; - it.barrier(sycl::access::fence_space::local_space); - } - float inv = 1.0f / sycl::sqrt((float)(lss[0] / HD) + QX_RMS_EPS); + float inv = 1.0f / sycl::sqrt((float)(pr[0] / HD) + QX_RMS_EPS); d_out[(size_t)t * out_s + hv*HD+col] = out_val * inv * d_norm_w[col]; - it.barrier(sycl::access::fence_space::local_space); + /* No end-of-token barrier: the next iteration's first SLM op + * writes ql[col]/kl[col] (each thread its own slot), and the + * first cross-lane read is after the L2-norm barrier above. */ } /* Write final S to global once. */ #pragma unroll -- 2.51.2