sycl: isolate opt-in mixed-IQ3 decode optimizations

Keep production-covered expert, projection, HC, planning and attention changes with their targeted checks. Exclude local profiler instrumentation and prefill experiments.

Co-authored-by: Copilot <[email protected]>

docs/INTEL.md

@@ -510,3 +510,35 @@ a type pun the SYCL compiler does not honour. A prompt that borrowed cache slots
 Level Zero v2). An A750 prompt of 2,744 tokens hung (direct copy out of pageable memory; regression test
 `sycl/tools/test_prompt_deadlock.sh`). The first request after a start returned token 0 on `i915` when the spin bound was too short.
 An A770 JIT build misread memory past 4 GiB into an allocation. The xe driver aliased pages in a 22 GiB arena (NaN at layer 4).
+
+### Opt-in mixed-IQ3 decode experiments
+
+These controls select separate decode paths. They do not enable INT8 prefill,
+change the expert cache size or change other devices' defaults. Set runtime
+controls before starting the engine; restart after changing them.
+
+Build with `-DSTRATA_SYCL_IQ_DECODE_FAST=ON` to select the packed Q2/IQ4 and
+sign helpers. The build option defaults to `OFF`. Test the same model and
+request sequence against a build without that option before using it.
+
+| Runtime control | Selected work |
+|---|---|
+| `STRATA_IQ3S_GU_SLM=1` | Stage native IQ3_S gate/up activations |
+| `STRATA_IQ4NL_DOWN_SLM=1` | Stage native IQ4_NL down activations |
+| `STRATA_IQ2S_GU_ROW_REUSE=1` | Reuse inputs across two IQ2_S gate/up rows |
+| `STRATA_IQ3XXS_GU_ROW_REUSE=1` | Reuse inputs across two IQ3_XXS gate/up rows |
+| `STRATA_EXPERT_ROW_E3=1` | Process eligible three-entry tails in the row-reuse path |
+| `STRATA_IQ3S_GU_ROW_REUSE=1` | Reuse IQ3_S gate/up row inputs |
+| `STRATA_RESIDENT_PLAN_COMPACT=1` | Size the resident routing work-group for the entry count |
+| `STRATA_GR_CACHED_NORM=1` | Reuse HC normalized input on supported layouts |
+| `STRATA_GR_LO_REUSE=1` | Reuse HC low-rank SiLU output on supported layouts |
+| `STRATA_GR_REDUCE_ILP=1` | Keep partial loads in flight in the cached-lo HC reduce |
+| `STRATA_MMVQ_SHAPE_RPW=1` | Select row reuse for bounded Q4_K/Q5_K projection shapes |
+| `STRATA_Q6_SHAPE_GRF=1` | Select large GRF for bounded Q6_K projection shapes |
+| `STRATA_GDN_QK_REG=1` | Register-broadcast QK for supported BMG GDN geometry |
+| `STRATA_ATTN_MERGE_ACTIVE=1` | Merge only active chunks in bounded batched QSA decode |
+
+Unsupported layouts retain their existing paths. Explicit projection overrides
+take precedence over shape selection. These options need workload-specific
+validation; enabling a control does not establish that a matching kernel ran
+or that the full request became faster.

sycl/CMakeLists.txt

@@ -29,6 +29,11 @@ endif()
 add_compile_definitions(STRATA_VERSION="${PROJECT_VERSION}")
 
 set(STRATA_SYCL_AOT "" CACHE STRING "AOT device for the GPU code, e.g. bmg-g31 (empty: SPIR-V, JIT at run time)")
+option(STRATA_SYCL_IQ_DECODE_FAST "Opt in to the packed IQ/Q2 decode helpers" OFF)
+if(STRATA_SYCL_IQ_DECODE_FAST)
+  add_compile_definitions(STRATA_Q2_WORD_FAST=1 STRATA_IQ4NL_LUT_SWAR=1 STRATA_IQ4XS_LUT_SWAR=1
+                         STRATA_IQ4_LUT_SWAR=1 STRATA_IQ3S_SIGN_FAST=1 STRATA_IQ3S_SIGN_APPLY_FAST=1)
+endif()
 # The device spins on host flags are bounded (include/strata/sycl_doorbell.hpp). The bound is chosen per device at run
 # time (20,000 reads on a B-series card, 2,000,000 elsewhere; STRATA_SPIN_MAX overrides). STRATA_SYCL_SPIN_MAX fixes one
 # bound for every device in this build.
@@ -355,6 +360,11 @@ if(STRATA_SYCL_PARITY)
   add_executable(native_grouped_parity ${_src})
   target_link_libraries(native_grouped_parity PRIVATE strata_kernels)
   add_test(NAME native_grouped_parity COMMAND native_grouped_parity)
+  if(STRATA_NATIVE_EXPERTS)
+    strata_resolve(_src src/kernels/native_expert_bench.cpp)
+    add_executable(native_expert_bench ${_src})
+    target_link_libraries(native_expert_bench PRIVATE strata_kernels strata_kernels_cpu)
+  endif()
   # 0.1.40 (#783): the batched verify-window kernels against the ones they replace, bitwise (synthetic, no model)
   foreach(p IN ITEMS verify native_multi kv_steps resident_plan gr_multi)
     strata_resolve(_src src/kernels/${p}_parity.cpp)
@@ -396,6 +406,8 @@ if(STRATA_SYCL_PARITY)
   target_link_libraries(q6k_align_bench PRIVATE strata_kernels)
   add_executable(gr_bench ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/gr_bench.cpp)   # fused GR read timing
   target_link_libraries(gr_bench PRIVATE strata_kernels)
+  add_executable(fused_gr_bench ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/fused_gr_bench.cpp)
+  target_link_libraries(fused_gr_bench PRIVATE strata_kernels)
   add_executable(sycl_router_fast ${CMAKE_CURRENT_SOURCE_DIR}/probe/router_fast.cpp)
   target_link_libraries(sycl_router_fast PRIVATE strata_kernels)
   add_test(NAME sycl_router_fast COMMAND sycl_router_fast)

sycl/include/strata/kernels/iq4_lut.hpp

@@ -0,0 +1,28 @@
+#pragma once
+
+#include <sycl/sycl.hpp>
+#include <cstdint>
+
+namespace strata::kernels {
+inline uint32_t iq4_lut4_swar(uint32_t q4) {
+    constexpr uint32_t table[] = {
+        0x81818181u, 0x98989898u, 0xADADADADu, 0xBFBFBFBFu,
+        0xCFCFCFCFu, 0xDDDDDDDDu, 0xEAEAEAEAu, 0xF6F6F6F6u,
+        0x01010101u, 0x0D0D0D0Du, 0x19191919u, 0x26262626u,
+        0x35353535u, 0x45454545u, 0x59595959u, 0x71717171u,
+    };
+    uint32_t mask[4], a[8], b[4], c[2];
+#pragma unroll
+    for (int i = 0; i < 4; ++i) {
+        const uint32_t bit = (q4 >> i) & 0x01010101u;
+        mask[i] = (bit << 8) - bit;
+    }
+#pragma unroll
+    for (int i = 0; i < 8; ++i) a[i] = sycl::bitselect(table[2 * i], table[2 * i + 1], mask[0]);
+#pragma unroll
+    for (int i = 0; i < 4; ++i) b[i] = sycl::bitselect(a[2 * i], a[2 * i + 1], mask[1]);
+#pragma unroll
+    for (int i = 0; i < 2; ++i) c[i] = sycl::bitselect(b[2 * i], b[2 * i + 1], mask[2]);
+    return sycl::bitselect(c[0], c[1], mask[3]);
+}
+}

sycl/src/kernels/cuda/fused_gr.dp.cpp

@@ -252,8 +252,43 @@ struct GrMulti {
     float* part;   ///< [DOWN_SPLIT][kFusedGrMaxT][LR] partial row sums of the down kernel
     int nsplit = DOWN_SPLIT;   ///< SYCL port: splits the up kernel sums (1 after the sliced down kernel's reduce)
     float* part2 = nullptr;    ///< SYCL port: [GRS_SLICES][kFusedGrMaxT][LR + HC] slice partials of the sliced kernel
+    bool cached_norm = false;
+    bool cached_lo = false;
 };
 
+bool gr_overlap(const void* a, size_t as, const void* b, size_t bs) {
+    if (!a || !b || !as || !bs) return false;
+    const uintptr_t av = reinterpret_cast<uintptr_t>(a), bv = reinterpret_cast<uintptr_t>(b);
+    return av <= bv ? bv - av < as : av - bv < bs;
+}
+
+bool gr_reuse_safe(const FusedGrArgs* a, int T, const float* xn) {
+    for (int k = 0; k < T; ++k) {
+        const std::pair<const void*, size_t> spans[] = {
+            {a[k].R, D * sizeof(float)}, {a[k].R_out, D * sizeof(float)},
+            {a[k].bo_prev, N * sizeof(float)}, {a[k].inj_prev, HC * sizeof(float)},
+            {a[k].w_norm, D * sizeof(float)}, {a[k].w_down, size_t(LR) * D * sizeof(uint16_t)},
+            {a[k].w_up, size_t(D) * LR * sizeof(uint16_t)},
+            {a[k].w_inject, size_t(HC) * D * sizeof(uint16_t)},
+            {a[k].rs, HC * sizeof(float)}, {a[k].inject_out, HC * sizeof(float)},
+            {a[k].mixed, N * sizeof(float)},
+        };
+        if (gr_overlap(xn, size_t(T) * D * sizeof(float), a[k].lo, LR * sizeof(float)))
+            return false;
+        for (const auto& span : spans) {
+            if (gr_overlap(xn, size_t(T) * D * sizeof(float), span.first, span.second))
+                return false;
+            for (int j = 0; j < T; ++j)
+                if (gr_overlap(a[j].lo, LR * sizeof(float), span.first, span.second))
+                    return false;
+        }
+        for (int j = 0; j < k; ++j)
+            if (gr_overlap(a[j].lo, LR * sizeof(float), a[k].lo, LR * sizeof(float)))
+                return false;
+    }
+    return true;
+}
+
 // SYCL port (Arc): a RUNTIME index into the by-value kernel argument (m.a[k]) makes IGC copy the whole GrMulti
 // (~1.2 KB) into every work-item's private scratch - unitrace: ~40 KB private per SIMD32 thread on the A770 for
 // gr_up_multi / gr_down_reduce / gr_norm_split. GR_PICK selects ONE field with compile-time indices only (a chain
@@ -596,7 +631,7 @@ __dpct_inline__ void gr_down_sliced_kernel(GrMulti m, float* tile) {
     }
 }
 static_assert(GRS_SLICES % 8 == 0, "the 8-at-a-time reduce reads whole groups of 8 slices");
-template <bool MORE = false>
+template <bool MORE = false, bool REUSE_LO = false>
 __dpct_inline__ void gr_down_reduce_kernel(GrMulti m) {
     auto item = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
     const int i = (int) item.get_global_id(2);
@@ -619,8 +654,15 @@ __dpct_inline__ void gr_down_reduce_kernel(GrMulti m) {
     } else {
         for (int sl = 0; sl < GRS_SLICES; ++sl) x += m.part2[((size_t) sl * kFusedGrMaxT + k) * GRS_ROWS + r];
     }
-    if (r < LR) m.part[(size_t) k * LR + r] = x;   // split 0
-    else GR_PICK(m, k, inject_out)[r - LR] = x;
+    if (r < LR) {
+        m.part[(size_t) k * LR + r] = x;   // split 0
+        if constexpr (REUSE_LO) {
+            float v = 0.0f;
+            v += x;
+            v /= (float) HC;
+            GR_PICK(m, k, lo)[r] = v / (1.0f + sycl::native::exp(-v));
+        }
+    } else GR_PICK(m, k, inject_out)[r - LR] = x;
 }
 // Large-GRF mode (256 registers per thread instead of 128) for the sliced down kernel and the multi-token up kernel.
 // On the A770 the up kernel spills ~3.4 KB per thread at 128 (unitrace). Large GRF halves the threads an Xe core keeps
@@ -681,7 +723,7 @@ Consult with your hardware vendor to find the total register size available and
 adjust the code, or use smaller sub-group size to avoid high register pressure.
 */
 // COLS: columns per block (UPM_COLS, or UPM_COLS_MORE with LO_ONCE under gr_more_groups)
-template <int COLS = UPM_COLS, bool ILP = false, bool LO_ONCE = false>   // columns per block: 16 (the original), 8 or 4 - each row's dot, butterfly and the per-column
+template <int COLS = UPM_COLS, bool ILP = false, bool LO_ONCE = false, bool REUSE = false>
 __dpct_inline__ void gr_up_multi_kernel(GrMulti m) {   // stream sum unchanged, so the same bits; more blocks, fewer serial rows per warp
     auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
 auto &lo = *sycl::ext::oneapi::group_local_memory_for_overwrite<
@@ -696,6 +738,12 @@ auto &lo = *sycl::ext::oneapi::group_local_memory_for_overwrite<
 #pragma unroll
     for (int i = t; i < T * LR; i += THREADS) {   // the down kernel's partial sums, in split order, then silu
         const int k = i / LR, r = i % LR;
+        if constexpr (REUSE) {
+            if (m.cached_lo) {
+                lo[k][r] = GR_PICK(m, k, lo)[r];
+                continue;
+            }
+        }
         float x = 0.0f;
 #pragma unroll
         for (int sp = 0; sp < DOWN_SPLIT; ++sp) if (sp < m.nsplit) x += m.part[((size_t) sp * kFusedGrMaxT + k) * LR + r];
@@ -740,9 +788,11 @@ auto &lo = *sycl::ext::oneapi::group_local_memory_for_overwrite<
         float rv = 0.0f, wn = 0.0f, rsc = 0.0f, bo = 0.0f, ip = 0.0f;
         bool apply = false;
         if (lane < T) {
-            rv = pR[i];
-            wn = pWn[i];
-            rsc = pRs[c];
+            if (!REUSE || !m.cached_norm || pApply) rv = pR[i];
+            if (!REUSE || !m.cached_norm) {
+                wn = pWn[i];
+                rsc = pRs[c];
+            }
             apply = pApply;
             if (apply) { bo = pBo[d0 + dd]; ip = pIp[c]; }
         }
@@ -781,7 +831,7 @@ auto &lo = *sycl::ext::oneapi::group_local_memory_for_overwrite<
                 rv = sycl::fma(bo, 2.0f * sigmoidf_(ip / (float)HC), rv);
                 pRout[i] = rv;
             }
-            const float x = rv * wn * rsc;
+            const float x = REUSE && m.cached_norm ? m.xn[(size_t) lane * D + i] : rv * wn * rsc;
             g[lane][c][dd] = x * sigmoidf_(mine);
         }
     }
@@ -850,6 +900,15 @@ template <int C>
 void launch_gr_up_cols(dpct::queue_ptr st, const GrMulti& m, bool grf) {
     auto exp_props = sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync};
     const sycl::nd_range<3> r(sycl::range(1, 1, N / C) * sycl::range(1, 1, THREADS), sycl::range(1, 1, THREADS));
+    if constexpr (C == UPM_COLS) {
+        if (m.cached_norm || m.cached_lo) {
+            st->parallel_for<dpct_kernel_name<class gr_up_multi_reuse>>(r, exp_props,
+                [=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(32)]] {
+                    gr_up_multi_kernel<C, false, false, true>(m);
+                });
+            return;
+        }
+    }
     if (grf) {   // Arc A-series: large-GRF mode (256 regs/thread); at 128 this kernel spills ~3.4 KB/thread on the A770
         st->parallel_for<dpct_kernel_name<class gr_up_multi_cols_grf256, dpct_kernel_scalar<C>>>(r,
             sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync,
@@ -1918,6 +1977,27 @@ bool fused_gr_read_multi(const FusedGrArgs* a, int n_tok, float* xn_scratch, voi
         return false;   // SYCL port: never writes the q8_1 images (STRATA_QFUSE is not supported here)
     }
     m.part = down_partials(st);
+    static const bool reuse_norm = [] {
+        const char* v = std::getenv("STRATA_GR_CACHED_NORM");
+        return v && v[0] == '1';
+    }();
+    static const bool reuse_lo = [] {
+        const char* v = std::getenv("STRATA_GR_LO_REUSE");
+        return v && v[0] == '1';
+    }();
+    if ((reuse_norm || reuse_lo) && gr_norm_split() && gr_down_sliced() &&
+        !gr_more_groups(st) && !gr_large_grf(st) && !gr_up_ilp() && gr_up_cols() == UPM_COLS &&
+        strata::intel_gpu_gen(st->get_device()) == strata::IntelGpuGen::Battlemage &&
+        gr_reuse_safe(a, n_tok, xn_scratch)) {
+        m.cached_norm = reuse_norm;
+        m.cached_lo = reuse_lo;
+        static const bool note = [] {
+            std::fprintf(stderr, "strata: GR reuse enabled: normalized activation %d, low-rank SiLU %d\n",
+                         (int) reuse_norm, (int) reuse_lo);
+            return true;
+        }();
+        (void) note;
+    }
     {
         auto exp_props = sycl::ext::oneapi::experimental::properties{
             sycl::ext::oneapi::experimental::use_root_sync};
@@ -1978,9 +2058,27 @@ bool fused_gr_read_multi(const FusedGrArgs* a, int n_tok, float* xn_scratch, voi
                     });
         });
         const unsigned nred = unsigned((n_tok * GRS_ROWS + 127) / 128);
-        st->parallel_for<dpct_kernel_name<class gr_down_reduce_k>>(
-            sycl::nd_range<3>(sycl::range(1, 1, nred) * sycl::range(1, 1, 128), sycl::range(1, 1, 128)),
-            [=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(32)]] { gr_down_reduce_kernel(m); });
+        const sycl::nd_range<3> nr(sycl::range(1, 1, nred) * sycl::range(1, 1, 128),
+                                    sycl::range(1, 1, 128));
+        if (m.cached_lo) {
+            static const bool reduce_ilp = [] {
+                const char* v = std::getenv("STRATA_GR_REDUCE_ILP");
+                return v && v[0] == '1';
+            }();
+            if (reduce_ilp) {
+                st->parallel_for<dpct_kernel_name<class gr_down_reduce_lo_ilp>>(nr,
+                    [=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(32)]] {
+                        gr_down_reduce_kernel<true, true>(m);
+                    });
+            } else {
+                st->parallel_for<dpct_kernel_name<class gr_down_reduce_lo>>(nr,
+                    [=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(32)]] {
+                        gr_down_reduce_kernel<false, true>(m);
+                    });
+            }
+        } else
+            st->parallel_for<dpct_kernel_name<class gr_down_reduce_k>>(nr,
+                [=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(32)]] { gr_down_reduce_kernel(m); });
     } else if (gr_down_direct()) {
         st->parallel_for<dpct_kernel_name<class gr_down_multi_direct>>(
             sycl::nd_range<3>(sycl::range(1, 1, DOWN_BLOCKS * DOWN_SPLIT + 1) * sycl::range(1, 1, THREADS),

sycl/src/kernels/cuda/iq_kernels.dp.cpp

@@ -13,6 +13,7 @@
 #include "strata/kernels/iq_dequant_expert.hpp"
 #include "strata/kernels/dp4a.hpp"
 #include "strata/kernels/q8_1_finite.hpp"
+#include "strata/kernels/iq4_lut.hpp"
 #include "s26_tsum.dp.hpp"
 
 #define GGML_COMMON_DECL_SYCL
@@ -100,6 +101,26 @@ __dpct_inline__ sycl::int2 get_int_from_table_16(const int &q4,
 #define ggml_cuda_dp4a(a, b, c) strata::dp4a((a), (b), (c))   // SYCL port: the one dp4a (strata/sycl_math.hpp)
 
 // ---------------------------------------------------------------- the dot products (vecdotq.cuh)
+#ifndef STRATA_Q2_WORD_FAST
+#define STRATA_Q2_WORD_FAST 0
+#endif
+__dpct_inline__ sycl::int2 q2_decode8(uint16_t q) {
+#if STRATA_Q2_WORD_FAST
+    const auto decode4 = [](uint32_t c) {
+        c &= 0xFFu;
+        c = (c | (c << 12)) & 0x000F000Fu;
+        c = (c | (c << 6)) & 0x03030303u;
+        return (int) ((c + 0x7F7F7F7Fu) ^ 0x80808080u);
+    };
+    return sycl::int2(decode4(q), decode4(q >> 8));
+#else
+    const int qe = dpct::byte_level_permute(0x020100FF, 0x020100FF, q);
+    const int qo = dpct::byte_level_permute(0x020100FF, 0x020100FF, q >> 2);
+    return sycl::int2(dpct::byte_level_permute(qe, qo, 0x5140),
+                      dpct::byte_level_permute(qe, qo, 0x7362));
+#endif
+}
+
 __dpct_inline__ float vec_dot_q2_0_q8_1(const void *__restrict__ vbq,
                                         const block_q8_1 *__restrict__ bq8_1,
                                         const int &kbx, const int &iqs) {
@@ -113,12 +134,9 @@ __dpct_inline__ float vec_dot_q2_0_q8_1(const void *__restrict__ vbq,
         const int q = qs[j];
         const int u = get_int_b4(bq8_1_chunk->qs, j * 2 + 0);
         const int v = get_int_b4(bq8_1_chunk->qs, j * 2 + 1);
-        const int qe = dpct::byte_level_permute(0x020100FF, 0x020100FF, q >> 0);
-        const int qo = dpct::byte_level_permute(0x020100FF, 0x020100FF, q >> 2);
-        const int qx = dpct::byte_level_permute(qe, qo, 0x5140);
-        const int qy = dpct::byte_level_permute(qe, qo, 0x7362);
-        sumi = ggml_cuda_dp4a(u, qx, sumi);
-        sumi = ggml_cuda_dp4a(v, qy, sumi);
+        const sycl::int2 w = q2_decode8((uint16_t) q);
+        sumi = ggml_cuda_dp4a(u, w.x(), sumi);
+        sumi = ggml_cuda_dp4a(v, w.y(), sumi);
     }
     const float d8 = bq8_1_chunk->ds[0];
     return d2 * d8 * sumi;
@@ -147,6 +165,29 @@ __dpct_inline__ int swar_ne4(unsigned x) { return dpct::vectorized_binary<sycl::
 __dpct_inline__ int swar_sub4(unsigned a, unsigned s) { return dpct::vectorized_binary<sycl::uchar4>(a, s, std::minus<>()); }
 #endif
 
+#ifndef STRATA_IQ3S_SIGN_FAST
+#define STRATA_IQ3S_SIGN_FAST 0
+#endif
+#ifndef STRATA_IQ3S_SIGN_APPLY_FAST
+#define STRATA_IQ3S_SIGN_APPLY_FAST 0
+#endif
+__dpct_inline__ int iq3s_sign_mask(unsigned nibble) {
+#if STRATA_IQ3S_SIGN_FAST
+    const unsigned bits = (nibble & 0x0Fu) * 0x00204081u & 0x01010101u;
+    return (int) ((bits << 8) - bits);
+#else
+    return swar_ne4((nibble & 1u) | ((nibble & 2u) << 8) |
+                    ((nibble & 4u) << 16) | ((nibble & 8u) << 24));
+#endif
+}
+__dpct_inline__ int iq3s_apply_sign(unsigned grid, unsigned sign) {
+#if STRATA_IQ3S_SIGN_APPLY_FAST
+    return (int) ((grid ^ sign) + (sign & 0x01010101u));
+#else
+    return swar_sub4(grid ^ sign, sign);
+#endif
+}
+
 __dpct_inline__ float vec_dot_iq2_xxs_q8_1(const void *__restrict__ vbq,
                                            const block_q8_1 *__restrict__ bq8_1,
                                            const int &kbx, const int &iqs) {
@@ -194,11 +235,11 @@ __dpct_inline__ float vec_dot_iq2_xs_q8_1(const void *__restrict__ vbq,
         const sycl::uint2 grid_pos =
             ((const sycl::uint2 *)iq2xs_grid)[q2[l0 / 2] & 0x1FF];
         const uint32_t signs = unpack_ksigns(q2[l0 / 2] >> 9);
-        const int signs0 = swar_ne4(signs & 0x08040201);
-        const int grid_l = swar_sub4(grid_pos.x() ^ signs0, signs0);
+        const int signs0 = iq3s_sign_mask(signs);
+        const int grid_l = iq3s_apply_sign(grid_pos.x(), signs0);
         const int u0 = get_int_b4(bq8_1[iqs / 2].qs, l0 + 0);
-        const int signs1 = swar_ne4(signs & 0x80402010);
-        const int grid_h = swar_sub4(grid_pos.y() ^ signs1, signs1);
+        const int signs1 = iq3s_sign_mask(signs >> 4);
+        const int grid_h = iq3s_apply_sign(grid_pos.y(), signs1);
         const int u1 = get_int_b4(bq8_1[iqs / 2].qs, l0 + 1);
         if (l0 < 4) {
             sumi0 = ggml_cuda_dp4a(grid_l, u0, sumi0);
@@ -230,12 +271,10 @@ __dpct_inline__ float vec_dot_iq2_s_q8_1(const void *__restrict__ vbq,
 #pragma unroll
     for (int l0 = 0; l0 < 8; l0 += 2) {
         const int* grid_pos = (const int*) (iq2s_grid + (qs[l0 / 2] | ((qh << (8 - l0)) & 0x300)));
-        const int signs0 = swar_ne4(((signs_packed_8[l0 / 2] & 0x03) << 7) |
-                ((signs_packed_8[l0 / 2] & 0x0C) << 21));
-        const int signs1 = swar_ne4(((signs_packed_8[l0 / 2] & 0x30) << 3) |
-                ((signs_packed_8[l0 / 2] & 0xC0) << 17));
-        const int grid_l = swar_sub4(grid_pos[0] ^ signs0, signs0);
-        const int grid_h = swar_sub4(grid_pos[1] ^ signs1, signs1);
+        const int signs0 = iq3s_sign_mask(signs_packed_8[l0 / 2]);
+        const int signs1 = iq3s_sign_mask(signs_packed_8[l0 / 2] >> 4);
+        const int grid_l = iq3s_apply_sign(grid_pos[0], signs0);
+        const int grid_h = iq3s_apply_sign(grid_pos[1], signs1);
         const int u0 = get_int_b4(bq8_1[iqs / 2].qs, l0 + 0);
         const int u1 = get_int_b4(bq8_1[iqs / 2].qs, l0 + 1);
         if (l0 < 4) {
@@ -358,7 +397,13 @@ __dpct_inline__ float vec_dot_iq1_m_q8_1(const void *__restrict__ vbq,
 // the codebook from global memory per lookup and each weight int as two 16-bit loads (block_iq4_nl is 18 bytes, its
 // qs 2-byte aligned); the expert down projection (IQ4_NL in 39 of 48 layers) was the largest expert kernel. Same ints
 // out, so the dots are bitwise unchanged.
+#ifndef STRATA_IQ4NL_LUT_SWAR
+#define STRATA_IQ4NL_LUT_SWAR 0
+#endif
 __dpct_inline__ uint32_t iq4nl_lut4(uint32_t q4) {   // four nibbles -> four kvalues_iq4nl bytes
+#if STRATA_IQ4NL_LUT_SWAR
+    return iq4_lut4_swar(q4);
+#else
     constexpr uint32_t t0 = 0xBFAD9881u, t1 = 0xF6EADDCFu, t2 = 0x26190D01u, t3 = 0x71594535u;
     uint32_t r = 0;
 #pragma unroll
@@ -368,6 +413,7 @@ __dpct_inline__ uint32_t iq4nl_lut4(uint32_t q4) {   // four nibbles -> four kva
         r |= ((t >> (8 * (n & 3))) & 0xFF) << (8 * i);
     }
     return r;
+#endif
 }
 __dpct_inline__ sycl::int2 iq4nl_pair(int aux) {   // = get_int_from_table_16(aux)
     return sycl::int2((int) iq4nl_lut4((uint32_t) aux & 0x0f0f0f0fu), (int) iq4nl_lut4(((uint32_t) aux >> 4) & 0x0f0f0f0fu));
@@ -423,6 +469,9 @@ __dpct_inline__ float vec_dot_iq4_nl_q8_1(const void *__restrict__ vbq,
     return d * sumi;
 }
 
+#ifndef STRATA_IQ4XS_LUT_SWAR
+#define STRATA_IQ4XS_LUT_SWAR 0
+#endif
 // IQ4_XS: 256 values as 8 sub-blocks of 32 (6-bit scale each); one call covers one sub-block (iqs = 4 * sub-block),
 // and `bq8_1` is the super-block's first q8_1 block, so the call's activation is bq8_1[iqs / 4].  The GSQ-RCO IQ3_S
 // file keeps one layer's routed gate/up experts in this format.
@@ -434,7 +483,11 @@ __dpct_inline__ float vec_dot_iq4_xs_q8_1(const void *__restrict__ vbq,
 #pragma unroll
     for (int j = 0; j < 4; ++j) {
         const int aux_q4 = get_int_b4(bq4->qs, iqs + j);
+#if STRATA_IQ4XS_LUT_SWAR
+        const sycl::int2 v = iq4nl_pair(aux_q4);
+#else
         const sycl::int2 v = get_int_from_table_16(aux_q4, kvalues_iq4nl);
+#endif
         const int u0 = get_int_b4(bq8_1[iqs / 4].qs, j + 0);
         const int u1 = get_int_b4(bq8_1[iqs / 4].qs, j + 4);
         sumi = ggml_cuda_dp4a(v.x(), u0, sumi);
@@ -1032,14 +1085,10 @@ template<> struct Split<17> : SplitLs2 {   // IQ2_XS
         for (int l0 = 0; l0 < 8; l0 += 2) {
             const sycl::uint2 grid_pos = grid_lut[q2[l0 / 2] & 0x1FF];
             const uint32_t signs = unpack_ksigns(q2[l0 / 2] >> 9);
-            const int signs0 = dpct::vectorized_binary<sycl::uchar4>(
-                signs & 0x08040201, 0, std::not_equal_to<>());
-            r.g[l0 + 0] = dpct::vectorized_binary<sycl::uchar4>(
-                grid_pos.x() ^ signs0, signs0, std::minus<>());
-            const int signs1 = dpct::vectorized_binary<sycl::uchar4>(
-                signs & 0x80402010, 0, std::not_equal_to<>());
-            r.g[l0 + 1] = dpct::vectorized_binary<sycl::uchar4>(
-                grid_pos.y() ^ signs1, signs1, std::minus<>());
+            const int signs0 = iq3s_sign_mask(signs);
+            r.g[l0 + 0] = iq3s_apply_sign(grid_pos.x(), signs0);
+            const int signs1 = iq3s_sign_mask(signs >> 4);
+            r.g[l0 + 1] = iq3s_apply_sign(grid_pos.y(), signs1);
         }
         r.dw = sycl::vec<sycl::half, 1>(bq2->d)
                    .convert<float, sycl::rounding_mode::automatic>()[0];
@@ -1065,18 +1114,10 @@ template<> struct Split<22> : SplitLs2 {   // IQ2_S
 #pragma unroll
         for (int l0 = 0; l0 < 8; l0 += 2) {
             const int* grid_pos = (const int*) (grid_lut + (qs[l0 / 2] | ((qh << (8 - l0)) & 0x300)));
-            const int signs0 = dpct::vectorized_binary<sycl::uchar4>(
-                ((signs_packed_8[l0 / 2] & 0x03) << 7) |
-                    ((signs_packed_8[l0 / 2] & 0x0C) << 21),
-                0x00000000, std::not_equal_to<>());
-            const int signs1 = dpct::vectorized_binary<sycl::uchar4>(
-                ((signs_packed_8[l0 / 2] & 0x30) << 3) |
-                    ((signs_packed_8[l0 / 2] & 0xC0) << 17),
-                0x00000000, std::not_equal_to<>());
-            r.g[l0 + 0] = dpct::vectorized_binary<sycl::uchar4>(
-                grid_pos[0] ^ signs0, signs0, std::minus<>());
-            r.g[l0 + 1] = dpct::vectorized_binary<sycl::uchar4>(
-                grid_pos[1] ^ signs1, signs1, std::minus<>());
+            const int signs0 = iq3s_sign_mask(signs_packed_8[l0 / 2]);
+            const int signs1 = iq3s_sign_mask(signs_packed_8[l0 / 2] >> 4);
+            r.g[l0 + 0] = iq3s_apply_sign(grid_pos[0], signs0);
+            r.g[l0 + 1] = iq3s_apply_sign(grid_pos[1], signs1);
         }
         r.dw = sycl::vec<sycl::half, 1>(bq2->d)
                    .convert<float, sycl::rounding_mode::automatic>()[0];
@@ -1280,12 +1321,9 @@ template<> struct Split<42> {   // Q2_0
 #pragma unroll
         for (int j = 0; j < 4; ++j) {
             const int q = qs[j];
-            const int qe =
-                dpct::byte_level_permute(0x020100FF, 0x020100FF, q >> 0);
-            const int qo =
-                dpct::byte_level_permute(0x020100FF, 0x020100FF, q >> 2);
-            r.qx[j] = dpct::byte_level_permute(qe, qo, 0x5140);
-            r.qy[j] = dpct::byte_level_permute(qe, qo, 0x7362);
+            const sycl::int2 w = q2_decode8((uint16_t) q);
+            r.qx[j] = w.x();
+            r.qy[j] = w.y();
         }
         return r;
     }
@@ -1712,10 +1750,10 @@ template <> struct Multi<18> {   // iq3_xxs
             const sycl::int2 grid_pos = G ? sycl::int2(G[q3[l0 + 0]], G[q3[l0 + 1]])
                                           : sycl::int2(iq3xxs_grid[q3[l0 + 0]], iq3xxs_grid[q3[l0 + 1]]);
             const uint32_t signs = unpack_ksigns(aux32 >> (7 * l0 / 2));
-            const int signs0 = swar_ne4(signs & 0x08040201);
-            const int signs1 = swar_ne4(signs & 0x80402010);
-            m.w[l0 + 0] = swar_sub4(grid_pos.x() ^ signs0, signs0);
-            m.w[l0 + 1] = swar_sub4(grid_pos.y() ^ signs1, signs1);
+            const int signs0 = iq3s_sign_mask(signs);
+            const int signs1 = iq3s_sign_mask(signs >> 4);
+            m.w[l0 + 0] = iq3s_apply_sign(grid_pos.x(), signs0);
+            m.w[l0 + 1] = iq3s_apply_sign(grid_pos.y(), signs1);
         }
         m.a = (int) (aux32 >> 28);
         m.d = sycl::vec<sycl::half, 1>(bq3->d).convert<float, sycl::rounding_mode::automatic>()[0];
@@ -1750,10 +1788,10 @@ template <> struct Multi<22> {   // iq2_s
         for (int l0 = 0; l0 < 8; l0 += 2) {
             const int gi = qs[l0 / 2] | ((qh << (8 - l0)) & 0x300);
             const int* grid_pos = G ? (const int*) (G + gi) : (const int*) (iq2s_grid + gi);
-            const int signs0 = swar_ne4(((sp[l0 / 2] & 0x03) << 7) | ((sp[l0 / 2] & 0x0C) << 21));
-            const int signs1 = swar_ne4(((sp[l0 / 2] & 0x30) << 3) | ((sp[l0 / 2] & 0xC0) << 17));
-            m.w[l0 + 0] = swar_sub4(grid_pos[0] ^ signs0, signs0);
-            m.w[l0 + 1] = swar_sub4(grid_pos[1] ^ signs1, signs1);
+            const int signs0 = iq3s_sign_mask(sp[l0 / 2]);
+            const int signs1 = iq3s_sign_mask(sp[l0 / 2] >> 4);
+            m.w[l0 + 0] = iq3s_apply_sign(grid_pos[0], signs0);
+            m.w[l0 + 1] = iq3s_apply_sign(grid_pos[1], signs1);
         }
         m.a = bq2->scales[iqs / 2] & 0x0F;
         m.b = bq2->scales[iqs / 2] >> 4;
@@ -1788,10 +1826,20 @@ template <> struct Multi<21> {   // iq3_s
         for (int l0 = 0; l0 < 8; l0 += 2) {
             const int i0 = qs[l0 + 0] | ((qh << (8 - l0)) & 0x100), i1 = qs[l0 + 1] | ((qh << (7 - l0)) & 0x100);
             const sycl::int2 grid_pos = G ? sycl::int2(G[i0], G[i1]) : sycl::int2(iq3s_grid[i0], iq3s_grid[i1]);
+#if STRATA_IQ3S_SIGN_FAST
+            const int signs0 = iq3s_sign_mask(sp[l0 / 2]);
+            const int signs1 = iq3s_sign_mask(sp[l0 / 2] >> 4);
+#else
             const int signs0 = swar_ne4(((sp[l0 / 2] & 0x03) << 7) | ((sp[l0 / 2] & 0x0C) << 21));
             const int signs1 = swar_ne4(((sp[l0 / 2] & 0x30) << 3) | ((sp[l0 / 2] & 0xC0) << 17));
+#endif
+#if STRATA_IQ3S_SIGN_APPLY_FAST
+            m.w[l0 + 0] = iq3s_apply_sign(grid_pos.x(), signs0);
+            m.w[l0 + 1] = iq3s_apply_sign(grid_pos.y(), signs1);
+#else
             m.w[l0 + 0] = swar_sub4(grid_pos.x() ^ signs0, signs0);
             m.w[l0 + 1] = swar_sub4(grid_pos.y() ^ signs1, signs1);
+#endif
         }
         m.a = 1 + 2 * ((bq3->scales[iqs / 4] >> ((iqs << 1) & 0x04)) & 0x0F);
         m.d = sycl::vec<sycl::half, 1>(bq3->d).convert<float, sycl::rounding_mode::automatic>()[0];
@@ -1843,10 +1891,9 @@ template <> struct Multi<42> {   // q2_0 (down)
 #pragma unroll
         for (int j = 0; j < 4; ++j) {
             const int q = qs[j];
-            const int qe = dpct::byte_level_permute(0x020100FF, 0x020100FF, q >> 0);
-            const int qo = dpct::byte_level_permute(0x020100FF, 0x020100FF, q >> 2);
-            m.w[2 * j + 0] = dpct::byte_level_permute(qe, qo, 0x5140);
-            m.w[2 * j + 1] = dpct::byte_level_permute(qe, qo, 0x7362);
+            const sycl::int2 w = q2_decode8((uint16_t) q);
+            m.w[2 * j + 0] = w.x();
+            m.w[2 * j + 1] = w.y();
         }
         m.d = bq2_0->d;
     }
@@ -1877,6 +1924,9 @@ template <int TY> struct KqSplitTq {
 template <> struct KqSplit<144> : KqSplitTq<144> {};
 template <> struct KqSplit<145> : KqSplitTq<145> {};
 template <> struct KqSplit<146> : KqSplitTq<146> {};
+#if STRATA_IQ4XS_LUT_SWAR
+template <> struct KqSplit<23> : Split<23> { static constexpr bool has = true; };
+#endif
 template <> struct KqSplit<12> {   // Q4_K (vec_dot_q4_K_q8_1)
     static constexpr bool has = true;
     struct W { int v[2]; uint16_t aux[2]; sycl::half2 dm; };
@@ -2163,6 +2213,71 @@ __dpct_inline__ void native_gu_kernel(
 // has halves of <= 4), so 4 takes such windows in one pass; 8 would take ~64-80 registers against ~48 (ptxas -v).
 constexpr int GRP_NC = 4;
 
+__dpct_inline__ void native_gu_iq3s_slm_kernel(
+    const unsigned long long *__restrict__ grp_ptr,
+    const int32_t *__restrict__ grp_start, const int32_t *__restrict__ n_groups,
+    const int32_t *__restrict__ ent_tok, const block_q8_1 *__restrict__ xq,
+    NativeExpertLayout L, float *__restrict__ gate, float *__restrict__ up) {
+    auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
+    auto &s_x_words = *sycl::ext::oneapi::group_local_memory_for_overwrite<
+        uint32_t[GRP_NC * 80 * 9]>(item_ct1.get_group());
+    const int ng = *n_groups;
+    const int rib = item_ct1.get_local_id(2) / 8;
+    const int sub = item_ct1.get_local_id(2) % 8;
+    const int row = item_ct1.get_group(2) * 32 + rib;
+    const bool is_up = row >= L.n_ff;
+    const int r = is_up ? row - (int) L.n_ff : row;
+    const size_t off = (is_up ? L.up_off : 0) + (size_t) r * L.gu_row;
+    const int nb = (int) (L.n_embd / Fmt<21>::qk);
+    const int xb = (int) (L.n_embd / 32);
+    auto *s_x = reinterpret_cast<const block_q8_1 *>(s_x_words);
+    for (int g = (int) item_ct1.get_group(1); g < ng; g += (int) item_ct1.get_group_range(1)) {
+        const uint8_t *wr = (const uint8_t *) grp_ptr[g] + off;
+        const int e0 = grp_start[g], e1 = grp_start[g + 1];
+        for (int e = e0; e < e1; e += GRP_NC) {
+            const int n = sycl::min(GRP_NC, e1 - e);
+            item_ct1.barrier(sycl::access::fence_space::local_space);
+            for (int j = 0; j < n; ++j) {
+                const auto *src = reinterpret_cast<const uint32_t *>(xq + (size_t) ent_tok[e + j] * xb);
+                for (int i = (int) item_ct1.get_local_id(2); i < xb * 9; i += 256)
+                    s_x_words[j * xb * 9 + i] = src[i];
+            }
+            item_ct1.barrier(sycl::access::fence_space::local_space);
+            float *dst = (is_up ? up : gate) + r;
+            if (n == 1) {
+                const block_q8_1 *xs[1] = {s_x};
+                float v[1];
+                row_dot_multi<21, 8, 1>(wr, xs, nb, sub, v);
+                if (sub == 0) dst[(size_t) e * L.n_ff] = v[0];
+            } else if (n == 2) {
+                const block_q8_1 *xs[2] = {s_x, s_x + xb};
+                float v[2];
+                row_dot_multi<21, 8, 2>(wr, xs, nb, sub, v);
+                if (sub == 0) {
+                    dst[(size_t) e * L.n_ff] = v[0];
+                    dst[(size_t) (e + 1) * L.n_ff] = v[1];
+                }
+            } else if (n == 3) {
+                const block_q8_1 *xs[3] = {s_x, s_x + xb, s_x + 2 * xb};
+                float v[3];
+                row_dot_multi<21, 8, 3>(wr, xs, nb, sub, v);
+                if (sub == 0) {
+                    dst[(size_t) e * L.n_ff] = v[0];
+                    dst[(size_t) (e + 1) * L.n_ff] = v[1];
+                    dst[(size_t) (e + 2) * L.n_ff] = v[2];
+                }
+            } else {
+                const block_q8_1 *xs[4] = {s_x, s_x + xb, s_x + 2 * xb, s_x + 3 * xb};
+                float v[4];
+                row_dot_multi<21, 8, 4>(wr, xs, nb, sub, v);
+                if (sub == 0)
+#pragma unroll
+                    for (int i = 0; i < 4; ++i) dst[(size_t) (e + i) * L.n_ff] = v[i];
+            }
+        }
+    }
+}
+
 template <int TG, bool STAGE_GRID = kStageIqGrid<TG>, bool SUB16 = true>
 /*
 DPCT1110: The total declared local variable size in device function
@@ -2369,6 +2484,66 @@ __dpct_inline__ void native_down_kernel(
     }
 }
 
+__dpct_inline__ void native_down_iq4nl_slm_kernel(
+    const unsigned long long *__restrict__ grp_ptr,
+    const int32_t *__restrict__ grp_start, const int32_t *__restrict__ n_groups,
+    const int32_t *__restrict__ ent_dst, const block_q8_1 *__restrict__ hq,
+    NativeExpertLayout L, float *__restrict__ out) {
+    auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
+    auto &s_hq_words = *sycl::ext::oneapi::group_local_memory_for_overwrite<
+        uint32_t[GRP_NC * 20 * 9]>(item_ct1.get_group());
+    const int ng = *n_groups;
+    const int rib = item_ct1.get_local_id(2) / 8;
+    const int sub = item_ct1.get_local_id(2) % 8;
+    const int r = item_ct1.get_group(2) * 32 + rib;
+    const size_t off = L.down_off + (size_t) r * L.d_row;
+    const int nb = (int) (L.n_ff / Fmt<20>::qk);
+    const int hb = (int) (L.n_ff / 32);
+    auto *s_hq = reinterpret_cast<const block_q8_1 *>(s_hq_words);
+    for (int g = (int) item_ct1.get_group(1); g < ng; g += (int) item_ct1.get_group_range(1)) {
+        const uint8_t *wr = (const uint8_t *) grp_ptr[g] + off;
+        const int e0 = grp_start[g], e1 = grp_start[g + 1];
+        for (int e = e0; e < e1; e += GRP_NC) {
+            const int n = sycl::min(GRP_NC, e1 - e);
+            const auto *src = reinterpret_cast<const uint32_t *>(hq + (size_t) e * hb);
+            const int words = n * hb * 9;
+            item_ct1.barrier(sycl::access::fence_space::local_space);
+            for (int i = (int) item_ct1.get_local_id(2); i < words; i += 256) s_hq_words[i] = src[i];
+            item_ct1.barrier(sycl::access::fence_space::local_space);
+            if (n == 1) {
+                const block_q8_1 *xs[1] = {s_hq};
+                float v[1];
+                row_dot_multi<20, 8, 1>(wr, xs, nb, sub, v);
+                if (sub == 0) out[(size_t) ent_dst[e] * L.n_embd + r] = v[0];
+            } else if (n == 2) {
+                const block_q8_1 *xs[2] = {s_hq, s_hq + hb};
+                float v[2];
+                row_dot_multi<20, 8, 2>(wr, xs, nb, sub, v);
+                if (sub == 0) {
+                    out[(size_t) ent_dst[e] * L.n_embd + r] = v[0];
+                    out[(size_t) ent_dst[e + 1] * L.n_embd + r] = v[1];
+                }
+            } else if (n == 3) {
+                const block_q8_1 *xs[3] = {s_hq, s_hq + hb, s_hq + 2 * hb};
+                float v[3];
+                row_dot_multi<20, 8, 3>(wr, xs, nb, sub, v);
+                if (sub == 0) {
+                    out[(size_t) ent_dst[e] * L.n_embd + r] = v[0];
+                    out[(size_t) ent_dst[e + 1] * L.n_embd + r] = v[1];
+                    out[(size_t) ent_dst[e + 2] * L.n_embd + r] = v[2];
+                }
+            } else {
+                const block_q8_1 *xs[4] = {s_hq, s_hq + hb, s_hq + 2 * hb, s_hq + 3 * hb};
+                float v[4];
+                row_dot_multi<20, 8, 4>(wr, xs, nb, sub, v);
+                if (sub == 0)
+#pragma unroll
+                    for (int i = 0; i < 4; ++i) out[(size_t) ent_dst[e + i] * L.n_embd + r] = v[i];
+            }
+        }
+    }
+}
+
 // native_down_kernel with the entries taken GRP_NC at a time
 template <int TD, bool SUB_DOWN = true>
 /*
@@ -3208,6 +3383,8 @@ bool g_old_kernels = env_on("STRATA_OLD_IQ_MMVQ");
 // every format; STRATA_EXPERT_SPLIT=1 takes upstream's multi kernels, launched with their own grid (GU_ROWS rows).
 bool g_split_multi = env_on("STRATA_EXPERT_SPLIT");
 bool g_no_sub16_gu = env_on("STRATA_NO_SUB16_GU");
+bool g_iq3s_gu_slm = env_on("STRATA_IQ3S_GU_SLM");
+bool g_iq4nl_down_slm = env_on("STRATA_IQ4NL_DOWN_SLM");
 // STRATA_IQ_STAGE_GRID=0 disables staging 64-bit i-quant codebook tables into shared memory
 bool g_stage_grid = [] {
     const char* v = std::getenv("STRATA_IQ_STAGE_GRID");
@@ -4483,6 +4660,141 @@ void s27_launch_ts(bool ts, const NativeExpertLayout &L, int64_t cap_groups,
 
 #endif
 
+template<int TY, int E>
+__dpct_inline__ void rows_dot_reuse(const uintptr_t (&rows)[2],
+                                   const block_q8_1* const* xs, int nb, int sub,
+                                   float (&out)[2][E]) {
+    static_assert(TY == 22 || TY == 18 || TY == 21);
+    using F = Fmt<TY>;
+    using M = Multi<TY>;
+    static_assert(F::qk % 32 == 0 && 8 % F::ipb == 0 && M::NW == 8);
+    constexpr int BS = TY == 22 ? sizeof(block_iq2_s) :
+                       TY == 21 ? sizeof(block_iq3_s) : sizeof(block_iq3_xxs);
+    constexpr int KSTEP = 8 / F::ipb, XB = F::qk / 32;
+    const int iqs = F::step * (sub % F::ipb), kb0 = sub / F::ipb;
+    uintptr_t wp[2] = {rows[0] + (size_t) kb0 * BS, rows[1] + (size_t) kb0 * BS};
+    size_t xoff = (size_t) kb0 * XB * sizeof(block_q8_1);
+    float sum[2][E];
+#pragma unroll
+    for (int r = 0; r < 2; ++r)
+#pragma unroll
+        for (int e = 0; e < E; ++e) sum[r][e] = 0.f;
+    for (int kb = kb0; kb < nb; kb += KSTEP) {
+        int u[E][M::NW];
+        float ds[E];
+#pragma unroll
+        for (int e = 0; e < E; ++e)
+            M::acts((const block_q8_1*) (uintptr_t(xs[e]) + xoff), iqs, u[e], ds[e]);
+#pragma unroll
+        for (int r = 0; r < 2; ++r) {
+            MultiW m;
+            M::prep((const void*) wp[r], 0, iqs, m);
+#pragma unroll
+            for (int e = 0; e < E; ++e) {
+                int s0 = 0, s1 = 0;
+#pragma unroll
+                for (int j = 0; j < M::NW / 2; ++j) s0 = ggml_cuda_dp4a(m.w[j], u[e][j], s0);
+#pragma unroll
+                for (int j = M::NW / 2; j < M::NW; ++j) s1 = ggml_cuda_dp4a(m.w[j], u[e][j], s1);
+                sum[r][e] += M::finish(s0, s1, m, ds[e]);
+            }
+            wp[r] += (size_t) KSTEP * BS;
+        }
+        xoff += (size_t) KSTEP * XB * sizeof(block_q8_1);
+    }
+#pragma unroll
+    for (int r = 0; r < 2; ++r)
+#pragma unroll
+        for (int e = 0; e < E; ++e) out[r][e] = lanes_sum<8>(sum[r][e]);
+}
+
+template<int TG, bool E3 = false>
+__dpct_inline__ void native_gu_row_reuse_kernel(
+    const unsigned long long* __restrict__ grp_ptr, const int32_t* __restrict__ grp_start,
+    const int32_t* __restrict__ n_groups, const int32_t* __restrict__ ent_tok,
+    const block_q8_1* __restrict__ xq, NativeExpertLayout L,
+    float* __restrict__ gate, float* __restrict__ up) {
+    auto it = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
+    const int ng = *n_groups;
+    if ((int) it.get_group(1) >= ng) return;
+    const int rib = (int) it.get_local_id(2) / 8, sub = (int) it.get_local_id(2) % 8;
+    const int row0 = (int) it.get_group(2) * 64 + rib, nrows = (int) (2 * L.n_ff);
+    if (row0 >= nrows) return;
+    const int nb = (int) (L.n_embd / Fmt<TG>::qk), xb = (int) (L.n_embd / 32);
+    size_t off[2];
+    float* dst[2];
+    bool live[2];
+#pragma unroll
+    for (int r = 0; r < 2; ++r) {
+        const int v = row0 + r * 32;
+        live[r] = v < nrows;
+        const int vv = live[r] ? v : row0;
+        const bool is_up = vv >= L.n_ff;
+        const int wr = is_up ? vv - (int) L.n_ff : vv;
+        off[r] = (is_up ? L.up_off : 0) + (size_t) wr * L.gu_row;
+        dst[r] = (is_up ? up : gate) + wr;
+    }
+    for (int g = (int) it.get_group(1); g < ng; g += (int) it.get_group_range(1)) {
+        const uintptr_t base = uintptr_t(grp_ptr[g]);
+        const uintptr_t rows[2] = {base + off[0], base + off[1]};
+        const int e1 = grp_start[g + 1];
+        int e = grp_start[g];
+        auto pass = [&](auto ec) {
+            constexpr int E = decltype(ec)::value;
+            const block_q8_1* xs[E];
+#pragma unroll
+            for (int j = 0; j < E; ++j) xs[j] = xq + (size_t) ent_tok[e + j] * xb;
+            float o[2][E];
+            rows_dot_reuse<TG, E>(rows, xs, nb, sub, o);
+            if (sub == 0)
+#pragma unroll
+                for (int r = 0; r < 2; ++r) if (live[r])
+#pragma unroll
+                    for (int j = 0; j < E; ++j) dst[r][(size_t) (e + j) * L.n_ff] = o[r][j];
+            e += E;
+        };
+        for (; e + 4 <= e1;) pass(std::integral_constant<int, 4>{});
+        if constexpr (E3) {
+            if (e + 3 == e1) pass(std::integral_constant<int, 3>{});
+        }
+        if (e + 2 <= e1) pass(std::integral_constant<int, 2>{});
+        if (e < e1) pass(std::integral_constant<int, 1>{});
+    }
+}
+
+template<int TG>
+void launch_gu_row_reuse(unsigned groups, dpct::queue_ptr s,
+                         const unsigned long long* ptr, const int32_t* start,
+                         const int32_t* ng, const int32_t* tok,
+                         const block_q8_1* xq, const NativeExpertLayout& L,
+                         float* gate, float* up) {
+    const unsigned gx = (unsigned) ((2 * L.n_ff + 63) / 64);
+    auto props = sycl::ext::oneapi::experimental::properties{
+        sycl::ext::oneapi::experimental::use_root_sync};
+    static const bool e3 = TG == 21 || env_on("STRATA_EXPERT_ROW_E3");
+    if (e3) {
+        static const bool logged = [] {
+            std::fprintf(stderr, "strata: %s exact three-entry row reuse enabled\n",
+                         TG == 22 ? "IQ2_S" : TG == 21 ? "IQ3_S" : "IQ3_XXS");
+            return true;
+        }();
+        (void) logged;
+        s->parallel_for<dpct_kernel_name<class native_gu_row_reuse_e3, dpct_kernel_scalar<TG>>>(
+            sycl::nd_range<3>(sycl::range<3>(1, groups, (size_t) gx * 256),
+                              sycl::range<3>(1, 1, 256)), props,
+            [=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(16)]] {
+                native_gu_row_reuse_kernel<TG, true>(ptr, start, ng, tok, xq, L, gate, up);
+            });
+        return;
+    }
+    s->parallel_for<dpct_kernel_name<class native_gu_row_reuse, dpct_kernel_scalar<TG>>>(
+        sycl::nd_range<3>(sycl::range<3>(1, groups, (size_t) gx * 256),
+                          sycl::range<3>(1, 1, 256)), props,
+        [=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(16)]] {
+            native_gu_row_reuse_kernel<TG>(ptr, start, ng, tok, xq, L, gate, up);
+        });
+}
+
 // SYCL port: lanes per row of the port's grouped expert kernels at run time (STRATA_GU_LANES / STRATA_DOWN_LANES:
 // 4, 8, 16 or 32; default STRATA_EXPERT_LANES), 256 / lanes rows per work-group. Gate/up rows are 80 dot calls
 // (2560 wide), the IQ4_NL down rows 40 (640 wide), so the best split can differ by role.
@@ -4501,6 +4813,22 @@ template <int TG, int LN, bool KM = true>
 void launch_gu_port(unsigned groups, dpct::queue_ptr s, const unsigned long long *grp_ptr, const int32_t *grp_start,
                     const int32_t *n_groups, const int32_t *ent_tok, const block_q8_1 *X,
                     const NativeExpertLayout &L, float *gate, float *up) {
+#if !STRATA_GRID_SLM
+    if constexpr ((TG == 22 || TG == 18) && LN == 8) {
+        static const bool reuse = env_on(TG == 22 ? "STRATA_IQ2S_GU_ROW_REUSE" : "STRATA_IQ3XXS_GU_ROW_REUSE");
+        if (reuse && !g_old_kernels && !g_split_multi &&
+            strata::intel_gpu_gen(s->get_device()) == strata::IntelGpuGen::Battlemage) {
+            static const bool note = [] {
+                std::fprintf(stderr, "strata: %s raw two-row GU reuse enabled (8 lanes, SIMD16)\n",
+                             TG == 22 ? "IQ2_S" : "IQ3_XXS");
+                return true;
+            }();
+            (void) note;
+            launch_gu_row_reuse<TG>(groups, s, grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up);
+            return;
+        }
+    }
+#endif
     constexpr int ROWS = 256 / LN;
     const unsigned gx = (unsigned) ((2 * L.n_ff + ROWS - 1) / ROWS);
     auto exp_props = sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync};
@@ -4519,6 +4847,13 @@ void gu_port_km(unsigned groups, dpct::queue_ptr s, const unsigned long long *gr
                 float *gate, float *up) {
     if constexpr (KqSplit<TG>::has) {
         if (!g_kq_split) { launch_gu_port<TG, LN, false>(groups, s, grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); return; }
+        if constexpr (TG == 23) {
+            static const bool logged = [] {
+                std::fprintf(stderr, "strata: IQ4_XS exact multi-entry decode reuse enabled\n");
+                return true;
+            }();
+            (void) logged;
+        }
     }
     launch_gu_port<TG, LN, true>(groups, s, grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up);
 }
@@ -4526,7 +4861,8 @@ template <int TG>
 void launch_gu_lanes(dpct::dim3 grid, dpct::queue_ptr s, const unsigned long long *grp_ptr, const int32_t *grp_start,
                      const int32_t *n_groups, const int32_t *ent_tok, const block_q8_1 *X,
                      const NativeExpertLayout &L, float *gate, float *up) {
-    switch (KqSplit<TG>::has ? kq_gu_lanes() : gu_lanes()) {
+    // IQ4_XS retains its original lane selection and reduction tree.
+    switch ((KqSplit<TG>::has && TG != 23) ? kq_gu_lanes() : gu_lanes()) {
         case 4: gu_port_km<TG, 4>(grid.y, s, grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
         case 16: gu_port_km<TG, 16>(grid.y, s, grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
         case 32: gu_port_km<TG, 32>(grid.y, s, grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
@@ -4590,6 +4926,33 @@ void launch_gu(dpct::dim3 grid, dpct::queue_ptr s,
                const int32_t *n_groups, const int32_t *ent_tok,
                const block_q8_1 *X, const NativeExpertLayout &L, float *gate,
                float *up) {
+    if constexpr (TG == 21) {
+#if !STRATA_GRID_SLM
+        static const bool reuse = env_on("STRATA_IQ3S_GU_ROW_REUSE");
+        if (reuse && !g_old_kernels && !g_split_multi && gu_lanes() == 8 &&
+            L.n_embd == 2560 && L.n_ff == 640 &&
+            strata::intel_gpu_gen(s->get_device()) == strata::IntelGpuGen::Battlemage) {
+            static const bool logged = [] {
+                std::fprintf(stderr, "strata: IQ3_S two-row GU reuse enabled (8 lanes, SIMD16, direct activation reads)\n");
+                return true;
+            }();
+            (void) logged;
+            launch_gu_row_reuse<21>(grid.y, s, grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up);
+            return;
+        }
+#endif
+        if (!g_old_kernels && g_iq3s_gu_slm && L.n_embd == 2560 && L.n_ff == 640) {
+            auto exp_props =
+                sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync};
+            const size_t gx = (size_t) ((2 * L.n_ff + 31) / 32);
+            s->parallel_for<dpct_kernel_name<class native_gu_iq3s_slm>>(
+                sycl::nd_range<3>(sycl::range<3>(1, grid.y, gx * 256), sycl::range<3>(1, 1, 256)), exp_props,
+                [=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(32)]] {
+                    native_gu_iq3s_slm_kernel(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up);
+                });
+            return;
+        }
+    }
     if constexpr (!kSplit<TG>) {
         launch_gu_lanes<TG>(grid, s, grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up);
     } else if (g_old_kernels || !g_split_multi) {
@@ -4618,12 +4981,30 @@ void launch_down_multi(dpct::dim3 grid, dpct::queue_ptr s, const unsigned long l
         });
 }
 
+void launch_down_iq4nl_slm(dpct::dim3 grid, dpct::queue_ptr s, const unsigned long long *grp_ptr,
+                            const int32_t *grp_start, const int32_t *n_groups, const int32_t *ent_dst,
+                            const block_q8_1 *hq, const NativeExpertLayout &L, float *out) {
+    auto exp_props = sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync};
+    const size_t gx = (size_t) ((L.n_embd + 31) / 32);
+    s->parallel_for<dpct_kernel_name<class native_down_iq4nl_slm>>(
+        sycl::nd_range<3>(sycl::range<3>(1, grid.y, gx * 256), sycl::range<3>(1, 1, 256)), exp_props,
+        [=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(32)]] {
+            native_down_iq4nl_slm_kernel(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out);
+        });
+}
+
 template <int TD>
 void launch_down(dpct::dim3 grid, dpct::queue_ptr s,
                  const unsigned long long *grp_ptr, const int32_t *grp_start,
                  const int32_t *n_groups, const int32_t *ent_dst,
                  const block_q8_1 *hq, const NativeExpertLayout &L,
                  float *out) {
+    if constexpr (TD == 20) {
+        if (!g_old_kernels && g_iq4nl_down_slm && L.n_embd == 2560 && L.n_ff == 640) {
+            launch_down_iq4nl_slm(grid, s, grp_ptr, grp_start, n_groups, ent_dst, hq, L, out);
+            return;
+        }
+    }
     if constexpr (!kSplit<TD>) {
         launch_down_lanes<TD>(grid, s, grp_ptr, grp_start, n_groups, ent_dst, hq, L, out);
     } else if (g_old_kernels || !g_split_multi) {

sycl/src/kernels/cuda/native_mmvq.dp.cpp

@@ -32,6 +32,7 @@
 #include "strata/kernels/native_mmvq.hpp"
 #include "strata/kernels/dp4a.hpp"
 #include "strata/kernels/q8_1_finite.hpp"
+#include "strata/kernels/iq4_lut.hpp"
 #include "strata/kernels/iq_kernels.hpp"
 #include "s26_tsum.dp.hpp"
 
@@ -1752,7 +1753,25 @@ void launch_q6k_wide(const void* weights, const void* x_q8_1, float* y, int n_in
     if (rpw == 1 && q6k_a2()) {
         const unsigned blocks = unsigned((std::size_t(n_out) + WARPS - 1) / WARPS);
         if constexpr (NCOLS > 1) {
-            if (mmvq_large_grf(s)) {
+            static const bool shape_grf = [] {
+                const char* v = std::getenv("STRATA_Q6_SHAPE_GRF");
+                return v && v[0] == '1';
+            }();
+            const bool eligible = (n_in == 6144 && n_out == 2560) ||
+                (n_in == 2560 && (n_out == 512 || n_out == 640)) ||
+                (NCOLS >= 3 && n_in == 2560 && n_out == 6144) ||
+                (NCOLS >= 4 && n_in == 2560 && (n_out == 10240 || n_out == 12288));
+            const bool bounded_grf = shape_grf && eligible &&
+                std::getenv("STRATA_MMVQ_LARGE_GRF") == nullptr &&
+                strata::intel_gpu_gen(s->get_device()) == strata::IntelGpuGen::Battlemage;
+            if (bounded_grf) {
+                static const bool note = [] {
+                    std::fprintf(stderr, "strata: shape-guarded Q6_K large-GRF dispatch enabled\n");
+                    return true;
+                }();
+                (void) note;
+            }
+            if (mmvq_large_grf(s) || bounded_grf) {
                 s->parallel_for<dpct_kernel_name<class native_mmvq_q6k_wide_a2_grf256, dpct_kernel_scalar<NCOLS>>>(
                     sycl::nd_range<3>(sycl::range(1, 1, blocks) * sycl::range(1, WARPS, WARP), sycl::range(1, WARPS, WARP)),
                     sycl::ext::oneapi::experimental::properties{sycl::ext::intel::experimental::grf_size<256>},
@@ -1874,7 +1893,13 @@ struct WideQ5K {
     static A act(const Q81Block* xb, int g) { return WideQ4K::act(xb, g); }
     static float dot(const W& r, const A& v) { return WideQ4K::dot(r, v); }
 };
+#ifndef STRATA_IQ4_LUT_SWAR
+#define STRATA_IQ4_LUT_SWAR 0
+#endif
 __dpct_inline__ uint32_t iq4_lut4(uint32_t q4) {   // four nibbles -> four codebook bytes, codebook in registers
+#if STRATA_IQ4_LUT_SWAR
+    return iq4_lut4_swar(q4);
+#else
     constexpr uint32_t t0 = 0xBFAD9881u, t1 = 0xF6EADDCFu, t2 = 0x26190D01u, t3 = 0x71594535u;
     uint32_t r = 0;
 #pragma unroll
@@ -1884,6 +1909,7 @@ __dpct_inline__ uint32_t iq4_lut4(uint32_t q4) {   // four nibbles -> four codeb
         r |= ((t >> (8 * (n & 3))) & 0xFF) << (8 * i);
     }
     return r;
+#endif
 }
 struct WideIQ4XS {
     using Block = IQ4XSBlock;
@@ -1985,7 +2011,29 @@ template <typename F, int NCOLS>
 void launch_wide(const void* weights, const void* x_q8_1, float* y, int n_in, int n_out, dpct::queue_ptr s) {
     const auto* w = static_cast<const typename F::Block*>(weights);
     const auto* x = static_cast<const Q81Block*>(x_q8_1);
-#define STRATA_WKR(RP) if (wide_k_rpw() == RP) { \
+    int rpw = wide_k_rpw();
+    static const bool shape_reuse = [] {
+        const char* v = std::getenv("STRATA_MMVQ_SHAPE_RPW");
+        return v && v[0] == '1';
+    }();
+    if (shape_reuse && std::getenv("STRATA_MMVQ_WIDE_K_RPW") == nullptr &&
+        strata::intel_gpu_gen(s->get_device()) == strata::IntelGpuGen::Battlemage) {
+        bool eligible = false;
+        if constexpr (std::is_same_v<F, WideQ4K>)
+            eligible = (n_in == 2560 && (n_out == 6144 || n_out == 10240 || n_out == 12288)) ||
+                       (n_in == 6144 && n_out == 2560);
+        if constexpr (std::is_same_v<F, WideQ5K>)
+            eligible = (n_in == 2560 && n_out == 12288) || (n_in == 6144 && n_out == 2560);
+        if (eligible && NCOLS <= 4) {
+            rpw = 2;
+            static const bool note = [] {
+                std::fprintf(stderr, "strata: shape-guarded Q4_K/Q5_K dense row reuse enabled\n");
+                return true;
+            }();
+            (void) note;
+        }
+    }
+#define STRATA_WKR(RP) if (rpw == RP) { \
         const unsigned b = unsigned((std::size_t(n_out) + WARPS * RP - 1) / (WARPS * RP)); \
         s->parallel_for<dpct_kernel_name<class native_mmvq_wide_rpw, F, dpct_kernel_scalar<NCOLS>, dpct_kernel_scalar<RP>>>( \
             sycl::nd_range<3>(sycl::range(1, 1, b) * sycl::range(1, WARPS, WARP), sycl::range(1, WARPS, WARP)), \

sycl/src/kernels/cuda/qsa_decode_attn.dp.cpp

@@ -615,6 +615,73 @@ __dpct_inline__ void attn_merge_kernel(const float *__restrict__ part_acc,
     attn[(size_t) h * HD + d] = L > 0.0f ? acc / L : 0.0f;
 }
 
+// Chunks beyond the device-side width contain only -FLT_MAX maxima. Keep the
+// original fold order while issuing eight independent partial reads together.
+__dpct_inline__ void attn_merge_active_kernel(
+    const float *__restrict__ part_acc, const float *__restrict__ part_m,
+    const float *__restrict__ part_l, int n_chunks, float *__restrict__ attn,
+    long long scratch_stride, const int32_t *__restrict__ steps) {
+    auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
+    const int z = item_ct1.get_group(1);
+    part_acc += (size_t) z * (size_t) scratch_stride;
+    part_m += (size_t) z * (size_t) scratch_stride;
+    part_l += (size_t) z * (size_t) scratch_stride;
+    attn += (size_t) z * (size_t) item_ct1.get_group_range(2) * HD;
+    const int width = steps[(size_t) z * kStepCount + kStepWidth];
+    const int reach = width <= 0 ? 0 : 1 + (width - 1) / CHUNK;
+    const int nc = sycl::min(reach, n_chunks);
+    const int h = item_ct1.get_group(2), kvh = h / G, hl = h % G;
+    const int d = item_ct1.get_local_id(2);
+    const float* pm_h = part_m + (size_t) kvh * n_chunks * G + hl;
+    const float* pl_h = part_l + (size_t) kvh * n_chunks * G + hl;
+    const float* pa_h = part_acc + ((size_t) kvh * n_chunks * G + hl) * HD + d;
+    constexpr int MU = 8;
+    float M = -FLT_MAX;
+    int c = 0;
+    for (; c + MU <= nc; c += MU) {
+        float pm[MU];
+#pragma unroll
+        for (int j = 0; j < MU; ++j) pm[j] = pm_h[(size_t) (c + j) * G];
+#pragma unroll
+        for (int j = 0; j < MU; ++j) M = sycl::fmax(M, pm[j]);
+    }
+    for (; c < nc; ++c) M = sycl::fmax(M, pm_h[(size_t) c * G]);
+    float L = 0.0f, acc = 0.0f;
+    c = 0;
+    for (; c + MU <= nc; c += MU) {
+        float pm[MU], pl[MU], pa[MU];
+#pragma unroll
+        for (int j = 0; j < MU; ++j) {
+            pm[j] = pm_h[(size_t) (c + j) * G];
+            pl[j] = pl_h[(size_t) (c + j) * G];
+            pa[j] = pa_h[(size_t) (c + j) * G * HD];
+        }
+#pragma unroll
+        for (int j = 0; j < MU; ++j) {
+            if (pm[j] == -FLT_MAX) continue;
+            const float w = sycl::native::exp(pm[j] - M);
+            L = sycl::fma(pl[j], w, L);
+            acc = sycl::fma(pa[j], w, acc);
+        }
+    }
+    for (; c < nc; ++c) {
+        const float m = pm_h[(size_t) c * G];
+        if (m == -FLT_MAX) continue;
+        const float w = sycl::native::exp(m - M);
+        L = sycl::fma(pl_h[(size_t) c * G], w, L);
+        acc = sycl::fma(pa_h[(size_t) c * G * HD], w, acc);
+    }
+    attn[(size_t) h * HD + d] = L > 0.0f ? acc / L : 0.0f;
+}
+
+bool merge_active_on() {
+    static const bool on = [] {
+        const char* v = std::getenv("STRATA_ATTN_MERGE_ACTIVE");
+        return v && v[0] == '1';
+    }();
+    return on;
+}
+
 #if defined(STRATA_EXPERIMENTAL_SM60)
 // the current device is below sm_75 (per device: a layer split can mix cards); STRATA_ATTN_PRE75=0 turns PR #540's
 // kernel off (A/B)
@@ -832,15 +899,27 @@ void qsa_decode_attn_batch(const float* q, const QsaAttnPools& pools, const int3
         auto exp_props = sycl::ext::oneapi::experimental::properties{
             sycl::ext::oneapi::experimental::use_root_sync};
 
-        st->parallel_for<dpct_kernel_name<class attn_merge_kernel_338175>>(
-            sycl::nd_range<3>(
-                sycl::range(1, (unsigned)n_q, (unsigned)s.n_head) *
-                    sycl::range(1, 1, HD),
-                sycl::range(1, 1, HD)),
-            exp_props, [=](sycl::nd_item<3> item_ct1) {
-                attn_merge_kernel(part_acc, part_m, part_l, n_chunks, attn,
-                                  stride);
-            });
+        if (merge_active_on() && n_q <= 8) {
+            st->parallel_for<dpct_kernel_name<class attn_merge_active>>(
+                sycl::nd_range<3>(
+                    sycl::range(1, (unsigned)n_q, (unsigned)s.n_head) *
+                        sycl::range(1, 1, HD),
+                    sycl::range(1, 1, HD)),
+                exp_props, [=](sycl::nd_item<3> item_ct1) {
+                    attn_merge_active_kernel(part_acc, part_m, part_l, n_chunks,
+                                             attn, stride, steps);
+                });
+        } else {
+            st->parallel_for<dpct_kernel_name<class attn_merge_kernel_338175>>(
+                sycl::nd_range<3>(
+                    sycl::range(1, (unsigned)n_q, (unsigned)s.n_head) *
+                        sycl::range(1, 1, HD),
+                    sycl::range(1, 1, HD)),
+                exp_props, [=](sycl::nd_item<3> item_ct1) {
+                    attn_merge_kernel(part_acc, part_m, part_l, n_chunks, attn,
+                                      stride);
+                });
+        }
     }
     /*
     DPCT1010: SYCL uses exceptions to report errors and does not use the

sycl/src/kernels/cuda/verify_kernels.dp.cpp

@@ -4,6 +4,7 @@
 // elementwise.cu) with the same operation order, so a verify window reproduces plain decode bit for bit.
 #define DPCT_PROFILING_ENABLED
 #include <sycl/sycl.hpp>
+#include <sycl/ext/intel/experimental/grf_size_properties.hpp>
 #include "strata/core/geometry.hpp"
 #include <dpct/dpct.hpp>
 #include "strata/sycl_queue.hpp"
@@ -336,7 +337,7 @@ __dpct_inline__ void gdn_q8_1_store(GdnQ81 *__restrict__ xq, size_t idx,
     if (idx % 32 == 0) xq[idx / 32].ds = sycl::half2(d, sum);
 }
 
-template <bool ALL_OUT, bool Q>
+template <bool ALL_OUT, bool Q, bool QK_REG = false>
 /*
 DPCT1110: The total declared local variable size in device function
 gdn_step_norm_multi_kernel exceeds 128 bytes and may cause high register
@@ -373,7 +374,7 @@ auto &sk = *sycl::ext::oneapi::group_local_memory_for_overwrite<float[2][S]>(
     const int qk = S * h_k;             // q at [0, qk), k at [qk, 2qk), v at [2qk, ...)
     const int value_dim = S * h_v;
     const int n = ALL_OUT ? T : (n_keep ? *n_keep : T);
-    if (n > 0) {
+    if (!QK_REG && n > 0) {
         if (rg == 0) sk[0][col] = hbuf[qk + qh * S + col];
         else if (rg == 1 && (ALL_OUT || 0 >= t_out_begin)) sq[0][col] = hbuf[qh * S + col];
     }
@@ -403,7 +404,14 @@ auto &sk = *sycl::ext::oneapi::group_local_memory_for_overwrite<float[2][S]>(
         better performance if there is no access to global memory.
         */
         item_ct1.barrier();
-        if (t + 1 < n) {
+        float k_own = 0.0f, q_own = 0.0f;
+        auto sg = item_ct1.get_sub_group();
+        if constexpr (QK_REG) {
+            const int i = rg * RPG + sg.get_local_linear_id();
+            k_own = ht[qk + qh * S + i];
+            if (need_out) q_own = ht[qh * S + i];
+        }
+        if (!QK_REG && t + 1 < n) {
             const float* ht_next = ht + C;
             const bool need_next = ALL_OUT || (t + 1 >= t_out_begin);
             if (rg == 0) sk[nxt][col] = ht_next[qk + qh * S + col];
@@ -413,7 +421,7 @@ auto &sk = *sycl::ext::oneapi::group_local_memory_for_overwrite<float[2][S]>(
         float kv = 0.0f;
 #pragma unroll
         for (int r = 0; r < RPG; ++r)
-            kv = sycl::fma(s[r], sk[cur][rg * RPG + r], kv);
+            kv = sycl::fma(s[r], QK_REG ? sycl::select_from_group(sg, k_own, r) : sk[cur][rg * RPG + r], kv);
         red_kv[rg][col] = kv;
         /*
         DPCT1118: SYCL group functions and algorithms must be encountered in
@@ -430,14 +438,16 @@ auto &sk = *sycl::ext::oneapi::group_local_memory_for_overwrite<float[2][S]>(
         if (!need_out) {
 #pragma unroll
             for (int r = 0; r < RPG; ++r)
-                s[r] = sycl::fma((float)g, s[r], sk[cur][rg * RPG + r] * delta);
+                s[r] = sycl::fma((float)g, s[r],
+                                (QK_REG ? sycl::select_from_group(sg, k_own, r) : sk[cur][rg * RPG + r]) * delta);
             continue;
         }
         float o = 0.0f;
 #pragma unroll
         for (int r = 0; r < RPG; ++r) {
-            s[r] = sycl::fma((float)g, s[r], sk[cur][rg * RPG + r] * delta);
-            o = sycl::fma(s[r], sq[cur][rg * RPG + r], o);
+            s[r] = sycl::fma((float)g, s[r],
+                            (QK_REG ? sycl::select_from_group(sg, k_own, r) : sk[cur][rg * RPG + r]) * delta);
+            o = sycl::fma(s[r], QK_REG ? sycl::select_from_group(sg, q_own, r) : sq[cur][rg * RPG + r], o);
         }
         red_o[rg][col] = o;
         /*
@@ -1839,6 +1849,31 @@ void gdn_step_norm_multi(float* state, const float* h, int conv_channels, const
         std::fprintf(stderr, "gdn_step_norm_multi: invalid arguments\n");
         std::exit(1);
     }
+    static const bool qk_reg = [] {
+        const char* v = std::getenv("STRATA_GDN_QK_REG");
+        return v && v[0] == '1';
+    }();
+    if (qk_reg && !xq && !n_keep && t_out_begin == 0 &&
+        h_k == 16 && h_v == 48 && conv_channels == 10240 &&
+        strata::intel_gpu_gen(strata::q_of(stream)->get_device()) == strata::IntelGpuGen::Battlemage) {
+        static const bool note = [] {
+            std::fprintf(stderr, "strata: GDN QK register broadcast enabled (GRF256)\n");
+            return true;
+        }();
+        (void) note;
+        auto props = sycl::ext::oneapi::experimental::properties{
+            sycl::ext::oneapi::experimental::use_root_sync,
+            sycl::ext::intel::experimental::grf_size<256>};
+        strata::q_of(stream)->parallel_for<dpct_kernel_name<class gdn_qk_register>>(
+            sycl::nd_range<3>(sycl::range(1, RG, (unsigned) h_v * S), sycl::range(1, RG, S)),
+            props, [=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(32)]] {
+                gdn_step_norm_multi_kernel<true, false, true>(
+                    state, h, conv_channels, gate, beta, z, gamma, eps, y,
+                    h_k, h_v, n_tok, nullptr, 0, nullptr);
+            });
+        check("gdn QK register");
+        return;
+    }
     static const bool split = [] { const char* v = std::getenv("STRATA_GDN_SPLIT"); return v && v[0] == '1'; }();
     // S26: the split runs where it is faster (gfx1151 harness, h_v 48: the commit 8-19% per call; a window of 2+ rows
     // 3-13%; one row: no gain, the old kernel stays) - bitwise equal in every case
@@ -2155,6 +2190,7 @@ auto &s_ids = *sycl::ext::oneapi::group_local_memory_for_overwrite<
     const int my_pack = (my_cnt << 16) | (is_first ? 1 : 0);
     const int lane = tid & 31;
     const int warp = tid >> 5;
+    const int n_warps = item_ct1.get_local_range(2) / 32;
     int pref = my_pack;
 #pragma unroll
     for (int d = 1; d < 32; d <<= 1) {
@@ -2190,7 +2226,7 @@ auto &s_ids = *sycl::ext::oneapi::group_local_memory_for_overwrite<
     if (is_first) {
         int tot = s_excl[tid];
 #pragma unroll
-        for (int w = 0; w < 4; ++w) {
+        for (int w = 0; w < n_warps; ++w) {
             if (w < warp) tot += s_wsum[w];
         }
         const int grp_idx = tot & 0xffff;
@@ -2203,7 +2239,7 @@ auto &s_ids = *sycl::ext::oneapi::group_local_memory_for_overwrite<
         const int fj_warp = first_j >> 5;
         int fj_tot = s_excl[first_j];
 #pragma unroll
-        for (int w = 0; w < 4; ++w) {
+        for (int w = 0; w < n_warps; ++w) {
             if (w < fj_warp) fj_tot += s_wsum[w];
         }
         const int out_idx = (fj_tot >> 16) + rank_in_group;
@@ -2211,7 +2247,9 @@ auto &s_ids = *sycl::ext::oneapi::group_local_memory_for_overwrite<
         tok[out_idx] = tid / k;
     }
     if (tid == 0) {
-        const int groups = (s_wsum[0] + s_wsum[1] + s_wsum[2] + s_wsum[3]) & 0xffff;
+        int groups = 0;
+        for (int w = 0; w < n_warps; ++w) groups += s_wsum[w];
+        groups &= 0xffff;
         start[groups] = n;
         start2[0] = n;
         counts[0] = groups;
@@ -2386,14 +2424,19 @@ void resident_plan(const int32_t* ids, int n_entries, int k, const int32_t* res_
     int32_t* use = nullptr;                    // the layer's slice of the device routing counts, if on
     if (g_mirror_table != nullptr && in_mirror_table(res_layer)) mir = g_mirror_table + (res_layer - g_mirror_res);
     if (g_usage != nullptr && in_mirror_table(res_layer)) use = g_usage + (res_layer - g_mirror_res);
+    static const bool compact = [] {
+        const char* e = std::getenv("STRATA_RESIDENT_PLAN_COMPACT");
+        return e != nullptr && std::atoi(e) != 0;
+    }();
+    const int threads = compact ? std::max(32, ((n_entries + 31) / 32) * 32) : kResidentPlanMax;
     {
         auto exp_props = sycl::ext::oneapi::experimental::properties{
             sycl::ext::oneapi::experimental::use_root_sync};
 
         strata::q_of(stream)
             ->parallel_for<dpct_kernel_name<class resident_plan_kernel_db6b7a>>(
-                sycl::nd_range<3>(sycl::range(1, 1, kResidentPlanMax),
-                                  sycl::range(1, 1, kResidentPlanMax)),
+                sycl::nd_range<3>(sycl::range(1, 1, threads),
+                                  sycl::range(1, 1, threads)),
                 exp_props,
                 [=](sycl::nd_item<3> item_ct1)
                     [[sycl::reqd_sub_group_size(32)]] {

sycl/src/kernels/fused_gr_bench.cpp

@@ -2,6 +2,7 @@
 // ones, bit for bit on every output (lo, rs, inject, mixed, the in-place R), and each one's time.
 //
 //     build/fused_gr_bench [iters] [T min] [T max]
+// On SYCL the fast setter is a no-op; compare hashes across builds or environment settings.
 #define DPCT_PROFILING_ENABLED
 #include <sycl/sycl.hpp>
 #include <dpct/dpct.hpp>
@@ -184,8 +185,15 @@ int main(int argc, char** argv) {
                              1000000.0f;
                         us[f] = std::fmin(us[f], 1e3 * ms / iters);
                     }
-                std::printf("T %d apply %d inject %d | old %6.1f us  fast %6.1f us | %s\n", T, apply, inject, us[0], us[1],
-                            same ? "bitwise equal" : "DIFFERS");
+                uint64_t hash = 1469598103934665603ull;
+                for (float value : out[1]) {
+                    uint32_t bits;
+                    std::memcpy(&bits, &value, sizeof(bits));
+                    hash = (hash ^ bits) * 1099511628211ull;
+                }
+                std::printf("T %d apply %d inject %d | old %6.1f us  fast %6.1f us | %s | hash %016llx\n",
+                            T, apply, inject, us[0], us[1], same ? "bitwise equal" : "DIFFERS",
+                            (unsigned long long) hash);
                 if (!same) ++failures;
             }
     std::printf("fused_gr_bench: %d differing\n", failures);

sycl/src/kernels/gdn_parity.cpp

@@ -918,6 +918,45 @@ int main(int argc, char** argv) {
                                 std::memcmp(st.data(), (c.n_keep < 0 ? st0 : sr).data(), st.size() * 4) != 0 ? ", state too" : "");
                     ++sn_bad;
                 }
+                if (std::getenv("STRATA_GDN_VERIFY_BENCH")) {
+                    auto hash = [](const std::vector<float>& values) {
+                        uint64_t h = 1469598103934665603ull;
+                        for (float v : values) {
+                            uint32_t u;
+                            std::memcpy(&u, &v, sizeof(u));
+                            h = (h ^ u) * 1099511628211ull;
+                        }
+                        return h;
+                    };
+                    std::printf("GDNSIGNATURE T=%d case=%s y=%016llx state=%016llx\n", T, c.name,
+                                (unsigned long long) hash(y), (unsigned long long) hash(st));
+                    if (c.n_keep < 0 && c.tb == 0) {
+                        auto* q = &dpct::get_in_order_queue();
+                        constexpr int calls = 32;
+                        dpct::experimental::command_graph_ptr graph;
+                        dpct::experimental::begin_recording(q);
+                        for (int i = 0; i < calls; ++i)
+                            strata::kernels::gdn_step_norm_multi(
+                                d_st, d_h, C, d_gate, d_beta, d_z, d_gm, eps,
+                                d_y, h_k, h_v, T, nullptr, q, 0);
+                        dpct::experimental::end_recording(q, &graph);
+                        auto executable = graph->finalize();
+                        delete graph;
+                        std::vector<double> times;
+                        for (int i = 0; i < 20; ++i) {
+                            auto begin = q->ext_oneapi_submit_barrier();
+                            q->ext_oneapi_graph(executable);
+                            auto end = q->ext_oneapi_submit_barrier();
+                            end.wait_and_throw();
+                            if (i >= 5)
+                                times.push_back(double(
+                                    end.get_profiling_info<sycl::info::event_profiling::command_end>() -
+                                    begin.get_profiling_info<sycl::info::event_profiling::command_start>()) / 1000 / calls);
+                        }
+                        std::sort(times.begin(), times.end());
+                        std::printf("GDNVERIFYBENCH T=%d us=%.3f\n", T, times[times.size() / 2]);
+                    }
+                }
             }
             sycl::free(d_h, dpct::get_in_order_queue());
                 sycl::free(d_gate, dpct::get_in_order_queue());

sycl/src/kernels/native_expert_bench.cpp

@@ -16,6 +16,7 @@
 #include "ggml.h"
 
 #include <cmath>
+#include <cstdint>
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
@@ -27,6 +28,16 @@
 namespace cpu = strata::kernels::cpu;
 namespace K = strata::kernels;
 
+uint64_t fnv1a(const void* data, size_t bytes) {
+    const auto* p = static_cast<const uint8_t*>(data);
+    uint64_t h = 1469598103934665603ull;
+    for (size_t i = 0; i < bytes; ++i) {
+        h ^= p[i];
+        h *= 1099511628211ull;
+    }
+    return h;
+}
+
 int main(int argc, char** argv) {
     setvbuf(stdout, nullptr, _IONBF, 0);
     if (argc < 7) {
@@ -48,9 +59,6 @@ int main(int argc, char** argv) {
     int failures = 0;
     dpct::queue_ptr s;
     s = dpct::get_current_device().create_queue(true);
-    dpct::event_ptr e0, e1;
-    e0 = new sycl::event();
-    e1 = new sycl::event();
     for (int l : layers) {
         const strata::TensorInfo* t[3] = {};
         const char* roles[3] = {"gate", "up", "down"};
@@ -91,16 +99,16 @@ int main(int argc, char** argv) {
         float* dout;
         unsigned long long* dptr;
         int32_t *dstart, *dn, *ddst, *dtok;
-        cudaMalloc((void**) &dblob, blobs.size());
-        cudaMalloc(&dx, x.size() * 4);
-        cudaMalloc(&dxq, (size_t) NTOK * H / 32 * 36);
-        cudaMalloc(&dscr, K::native_expert_scratch_bytes(NE, FF));
-        cudaMalloc((void**) &dout, (size_t) NE * H * 4);
-        cudaMalloc((void**) &dptr, (size_t) G * 8);
-        cudaMalloc((void**) &dstart, (size_t) (G + 1) * 4);
-        cudaMalloc((void**) &dn, 4);
-        cudaMalloc((void**) &ddst, (size_t) NE * 4);
-        cudaMalloc((void**) &dtok, (size_t) NE * 4);
+        dblob = sycl::malloc_device<uint8_t>(blobs.size(), *s);
+        dx = sycl::malloc_device(x.size() * sizeof(float), *s);
+        dxq = sycl::malloc_device((size_t) NTOK * H / 32 * 36, *s);
+        dscr = sycl::malloc_device(K::native_expert_scratch_bytes(NE, FF), *s);
+        dout = sycl::malloc_device<float>((size_t) NE * H, *s);
+        dptr = sycl::malloc_device<unsigned long long>(G, *s);
+        dstart = sycl::malloc_device<int32_t>(G + 1, *s);
+        dn = sycl::malloc_device<int32_t>(1, *s);
+        ddst = sycl::malloc_device<int32_t>(NE, *s);
+        dtok = sycl::malloc_device<int32_t>(NE, *s);
         for (int g = 0; g < G; ++g) {
             ptr[g] = (unsigned long long) (dblob + (size_t) g * f.bytes);
             start[g] = g * T;
@@ -110,26 +118,27 @@ int main(int argc, char** argv) {
             }
         }
         start[G] = NE;
-        cudaMemcpy(dblob, blobs.data(), blobs.size(), cudaMemcpyHostToDevice);
-        cudaMemcpy(dx, x.data(), x.size() * 4, cudaMemcpyHostToDevice);
-        cudaMemcpy(dptr, ptr.data(), (size_t) G * 8, cudaMemcpyHostToDevice);
-        cudaMemcpy(dstart, start.data(), (size_t) (G + 1) * 4, cudaMemcpyHostToDevice);
-        cudaMemcpy(dn, &G, 4, cudaMemcpyHostToDevice);
-        cudaMemcpy(ddst, dst.data(), (size_t) NE * 4, cudaMemcpyHostToDevice);
-        cudaMemcpy(dtok, tok.data(), (size_t) NE * 4, cudaMemcpyHostToDevice);
+        s->memcpy(dblob, blobs.data(), blobs.size());
+        s->memcpy(dx, x.data(), x.size() * sizeof(float));
+        s->memcpy(dptr, ptr.data(), (size_t) G * sizeof(*dptr));
+        s->memcpy(dstart, start.data(), (size_t) (G + 1) * sizeof(*dstart));
+        s->memcpy(dn, &G, sizeof(*dn));
+        s->memcpy(ddst, dst.data(), (size_t) NE * sizeof(*ddst));
+        s->memcpy(dtok, tok.data(), (size_t) NE * sizeof(*dtok));
         K::quantize_q8_1_rows((const float*) dx, NTOK, H, dxq, s);
+        s->wait_and_throw();
         auto run = [&](int m, int phase) {
             K::native_expert_set_mode(m, phase);
             K::native_expert_grouped(L, dptr, dstart, dn, ddst, dtok, G, NE, dxq, dscr, dout, s);
         };
         std::vector<float> ref((size_t) NE * H), got((size_t) NE * H);
-        cudaMemsetAsync(dout, 0xff, got.size() * 4, s);
+        s->memset(dout, 0xff, got.size() * sizeof(float));
         run(mref, 0);
-        cudaMemcpyAsync(ref.data(), dout, ref.size() * 4, cudaMemcpyDeviceToHost, s);
-        cudaMemsetAsync(dout, 0xff, got.size() * 4, s);
+        s->memcpy(ref.data(), dout, ref.size() * sizeof(float));
+        s->memset(dout, 0xff, got.size() * sizeof(float));
         run(mode, 0);
-        cudaMemcpyAsync(got.data(), dout, got.size() * 4, cudaMemcpyDeviceToHost, s);
-        if (cudaStreamSynchronize(s) != cudaSuccess) { std::printf("layer %d: CUDA error\n", l); return 1; }
+        s->memcpy(got.data(), dout, got.size() * sizeof(float));
+        s->wait_and_throw();
         size_t ndiff = 0;
         double worst = 0;
         for (size_t i = 0; i < ref.size(); ++i)
@@ -139,13 +148,13 @@ int main(int argc, char** argv) {
             }
         auto time = [&](int m, int phase) {
             for (int i = 0; i < 10; ++i) run(m, phase);
-            cudaEventRecord(e0, s);
+            sycl::event e0 = s->ext_oneapi_submit_barrier();
             for (int i = 0; i < iters; ++i) run(m, phase);
-            cudaEventRecord(e1, s);
-            cudaEventSynchronize(e1);
-            float ms = 0;
-            cudaEventElapsedTime(&ms, e0, e1);
-            return 1e3 * ms / iters;
+            sycl::event e1 = s->ext_oneapi_submit_barrier();
+            e1.wait_and_throw();
+            const auto begin = e0.get_profiling_info<sycl::info::event_profiling::command_start>();
+            const auto end = e1.get_profiling_info<sycl::info::event_profiling::command_end>();
+            return (double) (end - begin) / 1e3 / iters;
         };
         const double gu_b = (double) G * 2 * f.up_off, d_b = (double) G * dsz;
         double us[2][3] = {{1e30, 1e30, 1e30}, {1e30, 1e30, 1e30}};
@@ -153,18 +162,21 @@ int main(int argc, char** argv) {
             for (int k = 0; k < 2; ++k)
                 for (int p = 0; p < 3; ++p) us[k][p] = std::fmin(us[k][p], time(k ? mode : mref, p));
         std::printf("layer %2d %-7s/%-6s G %d T %d | mode %d: all %7.1f us  gu %7.1f us (%3.0f GB/s)  down %6.1f us (%3.0f GB/s)"
-                    " | mode %d: all %7.1f us  gu %7.1f us (%3.0f GB/s)  down %6.1f us (%3.0f GB/s) | %s\n",
+                    " | mode %d: all %7.1f us  gu %7.1f us (%3.0f GB/s)  down %6.1f us (%3.0f GB/s)"
+                    " | hash %016llx | %s\n",
                     l, ggml_type_name((ggml_type) f.gu_type), ggml_type_name((ggml_type) f.d_type), G, T,
                     mref, us[0][0], us[0][1], gu_b / us[0][1] / 1e3, us[0][2], d_b / us[0][2] / 1e3,
                     mode, us[1][0], us[1][1], gu_b / us[1][1] / 1e3, us[1][2], d_b / us[1][2] / 1e3,
+                    (unsigned long long) fnv1a(got.data(), got.size() * sizeof(float)),
                     ndiff ? "DIFFERS" : "bitwise equal");
         if (ndiff) {
             std::printf("          %zu of %zu outputs differ, worst rel %.2e\n", ndiff, ref.size(), worst);
             ++failures;
         }
         K::native_expert_set_mode(-1, 0);
-        cudaFree(dblob); cudaFree(dx); cudaFree(dxq); cudaFree(dscr); cudaFree(dout); cudaFree(dptr);
-        cudaFree(dstart); cudaFree(dn); cudaFree(ddst); cudaFree(dtok);
+        sycl::free(dblob, *s); sycl::free(dx, *s); sycl::free(dxq, *s); sycl::free(dscr, *s);
+        sycl::free(dout, *s); sycl::free(dptr, *s); sycl::free(dstart, *s); sycl::free(dn, *s);
+        sycl::free(ddst, *s); sycl::free(dtok, *s);
     }
     std::printf("native_expert_bench: %d differing\n", failures);
     return failures ? 1 : 0;

sycl/src/kernels/qsa_prompt_attn_parity.cpp

@@ -117,6 +117,11 @@ int run(int fmt, int64_t ctx, int64_t nq, int reps) {   // fmt 1 int8, 0 fp16, 2
     std::vector<int32_t> table(pages);
     for (int64_t i = 0; i < pages; ++i) table[i] = (int32_t) i;
     std::shuffle(table.begin(), table.end(), rng);
+    if (const char* mask = std::getenv("QA_MASK_PAGES")) {
+        const bool all = std::strcmp(mask, "all") == 0;
+        for (size_t i = 0; i < table.size(); ++i)
+            if (all || i % 7 == 0) table[i] = -1;
+    }
     // queries at positions ctx - nq .. ctx - 1
     const int64_t cap = k::qsa_selection_width(ctx, s);
     std::vector<int32_t> ids((size_t) (nq * cap), 0), steps((size_t) (nq * k::kStepCount), 0);
@@ -261,6 +266,11 @@ int run(int fmt, int64_t ctx, int64_t nq, int reps) {   // fmt 1 int8, 0 fp16, 2
                             .memcpy(nw.data(), d_new, nw.size() * 4)
                             .wait()),
        "down");
+    {
+        uint64_t hh = 1469598103934665603ull;
+        for (float x : o) { uint32_t u; std::memcpy(&u, &x, 4); hh = (hh ^ u) * 1099511628211ull; }
+        std::printf("FP32HASH %016llx\n", (unsigned long long) hh);
+    }
     {   // a hash of the new kernel's output bits (bitwise A/B of two builds / switches)
         uint64_t hh = 1469598103934665603ull;
         for (float x : nw) { uint32_t u; std::memcpy(&u, &x, 4); hh = (hh ^ u) * 1099511628211ull; }
@@ -277,6 +287,10 @@ int run(int fmt, int64_t ctx, int64_t nq, int reps) {   // fmt 1 int8, 0 fp16, 2
             double mx = -1e300;
             for (int64_t c = 0; c < w; ++c) {
                 const int64_t cell = sel[c], row = ((int64_t) table[cell / PS] * NKV + kvh) * PS + cell % PS;
+                if (table[cell / PS] < 0) {
+                    sco[c] = -INFINITY;
+                    continue;
+                }
                 double a = 0;
                 for (int64_t d = 0; d < HD; ++d) {
                     const double kv = fmt == 2 ? q4v(k4, row, d)
@@ -292,11 +306,12 @@ int run(int fmt, int64_t ctx, int64_t nq, int reps) {   // fmt 1 int8, 0 fp16, 2
                 double a = 0;
                 for (int64_t c = 0; c < w; ++c) {
                     const int64_t cell = sel[c], row = ((int64_t) table[cell / PS] * NKV + kvh) * PS + cell % PS;
+                    if (table[cell / PS] < 0) continue;
                     const double vv = fmt == 2 ? q4v(v4, row, d)
                                       : fmt == 1 ? (double) vq[row * HD + d] * h2f(vs[row * 4 + d / 64]) : h2f(vh[row * HD + d]);
                     a += sco[c] * vv;
                 }
-                const double r = a / l;
+                const double r = l > 0 ? a / l : 0.0;
                 const size_t at = (size_t) ((i * NH + h) * HD + d);
                 ref_scale = std::max(ref_scale, std::fabs(r));
                 err_old = std::max(err_old, std::fabs(o[at] - r));
@@ -331,6 +346,30 @@ int run(int fmt, int64_t ctx, int64_t nq, int reps) {   // fmt 1 int8, 0 fp16, 2
         (e1->get_profiling_info<sycl::info::event_profiling::command_end>() -
          e0->get_profiling_info<sycl::info::event_profiling::command_start>()) /
         1000000.0f;
+    if (std::getenv("QA_FP32_GRAPH")) {
+        auto* queue = &dpct::get_in_order_queue();
+        constexpr int calls = 32;
+        dpct::experimental::command_graph_ptr graph;
+        dpct::experimental::begin_recording(queue);
+        for (int i = 0; i < calls; ++i) old_run();
+        dpct::experimental::end_recording(queue, &graph);
+        auto executable = graph->finalize();
+        delete graph;
+        std::vector<double> times;
+        for (int i = 0; i < 20; ++i) {
+            auto begin = queue->ext_oneapi_submit_barrier();
+            queue->ext_oneapi_graph(executable);
+            auto end = queue->ext_oneapi_submit_barrier();
+            end.wait_and_throw();
+            if (i >= 5)
+                times.push_back(double(
+                    end.get_profiling_info<sycl::info::event_profiling::command_end>() -
+                    begin.get_profiling_info<sycl::info::event_profiling::command_start>()) / 1000 / calls);
+        }
+        std::sort(times.begin(), times.end());
+        std::printf("DECODEGRAPH fmt=%d ctx=%lld nq=%lld us=%.3f\n",
+                    fmt, (long long) ctx, (long long) nq, times[times.size() / 2]);
+    }
     const bool ok1 = err_new <= std::max(4.0 * err_old, 1e-6 * ref_scale);
     const bool ok2 = diff <= 1e-4 * scale;
     std::printf("%s %s ctx %lld, %lld queries: vs FP64 old %.3g new %.3g (output scale %.3g); new vs old %.3g (%.2g of "
@@ -378,12 +417,19 @@ int main(int argc, char** argv) {
     const int64_t ctx = argc > 1 ? std::atoll(argv[1]) : 32768;
     const int64_t nq = argc > 2 ? std::atoll(argv[2]) : 2048;
     const int reps = argc > 3 ? std::atoi(argv[3]) : 5;
+    const bool int8_only = argc > 4 && std::strcmp(argv[4], "--int8-only") == 0;
+    if (argc > 4 && !int8_only) {
+        std::fprintf(stderr, "usage: qsa_prompt_attn_parity [context] [queries] [reps] [--int8-only]\n");
+        return 2;
+    }
     int fails = 0;
     fails += run(1, ctx, nq, reps);
 #if !defined(__HIP_PLATFORM_AMD__)
-    fails += run(0, ctx, nq, reps);   // FP16 KV: the RDNA4 kernel takes int8 KV only
-    fails += run(2, ctx, nq, reps);   // Q4_0 KV (mode 4)
-    fails += run(2, 1500, std::min<int64_t>(nq, 1500), reps);
+    if (!int8_only) {
+        fails += run(0, ctx, nq, reps);   // FP16 KV: the RDNA4 kernel takes int8 KV only
+        fails += run(2, ctx, nq, reps);   // Q4_0 KV (mode 4)
+        fails += run(2, 1500, std::min<int64_t>(nq, 1500), reps);
+    }
 #endif
     fails += run(1, 1500, std::min<int64_t>(nq, 1500), reps);   // short context: the selection is every cell
     fails += run(1, 2100, std::min<int64_t>(nq, 256), reps);    // the identity-to-sparse edge