diff --git a/zig/build/hip.zig b/zig/build/hip.zig
index f165be7f23..22bd1cf9e7 100644
--- a/zig/build/hip.zig
+++ b/zig/build/hip.zig
@@ -312,15 +312,15 @@ const groups = [_][]const u8{
 
 /// What the group sources include, so an edit to one rebuilds them.
 const headers = [_][]const u8{
-    "attention/attention_fa.hip", "decode/pages.hpp",      "decode/plan.hpp",       "ops/attention.hpp",
-    "ops/common.hpp",             "ops/draw.hpp",          "ops/tp.hpp",            "ops/elementwise.hpp",
-    "ops/linear.hpp",             "ops/moe.hpp",           "ops/norms.hpp",         "ops/rope.hpp",
-    "common/arch.hpp",            "common/dot2.hpp",       "common/vec.hpp",        "common/wmma.hpp",
-    "quant/act.hpp",              "quant/mlx.hpp",         "quant/mlx_decoder.hpp", "quant/mlx_pieces.hpp",
-    "quant/mlx_tiles.hpp",        "tiles/dot2.hpp",        "tiles/epilogue.hpp",    "tiles/gemm.hpp",
-    "tiles/gemm_kp.hpp",          "tiles/matrix_gemm.hpp", "tiles/plan.hpp",        "tiles/stream.hpp",
-    "flashnext/streams.hpp",      "flashnext/ngram.hpp",   "flashnext/indexer.hpp", "flashnext/sparse.hpp",
-    "ops/nemotron_moe.hpp",
+    "attention/attention_fa.hip",   "decode/pages.hpp",      "decode/plan.hpp",       "ops/attention.hpp",
+    "ops/common.hpp",               "ops/draw.hpp",          "ops/tp.hpp",            "ops/elementwise.hpp",
+    "ops/linear.hpp",               "ops/moe.hpp",           "ops/norms.hpp",         "ops/rope.hpp",
+    "common/arch.hpp",              "common/dot2.hpp",       "common/vec.hpp",        "common/wmma.hpp",
+    "quant/act.hpp",                "quant/mlx.hpp",         "quant/mlx_decoder.hpp", "quant/mlx_pieces.hpp",
+    "quant/mlx_tiles.hpp",          "tiles/dot2.hpp",        "tiles/epilogue.hpp",    "tiles/gemm.hpp",
+    "tiles/gemm_kp.hpp",            "tiles/matrix_gemm.hpp", "tiles/plan.hpp",        "tiles/stream.hpp",
+    "flashnext/streams.hpp",        "flashnext/ngram.hpp",   "flashnext/indexer.hpp", "flashnext/sparse.hpp",
+    "flashnext/sparse_prefill.hpp", "ops/nemotron_moe.hpp",
 };
 
 /// The flags of the kernels' first build: no contraction, wave32 on RDNA, C++20.
diff --git a/zig/kernels/hip/flashnext/flashnext.hip b/zig/kernels/hip/flashnext/flashnext.hip
index 19260651c6..468a5a9f77 100644
--- a/zig/kernels/hip/flashnext/flashnext.hip
+++ b/zig/kernels/hip/flashnext/flashnext.hip
@@ -4,3 +4,4 @@
 #include "ngram.hpp"
 #include "indexer.hpp"
 #include "sparse.hpp"
+#include "sparse_prefill.hpp"
diff --git a/zig/kernels/hip/flashnext/sparse_prefill.hpp b/zig/kernels/hip/flashnext/sparse_prefill.hpp
new file mode 100644
index 0000000000..483c91ba20
--- /dev/null
+++ b/zig/kernels/hip/flashnext/sparse_prefill.hpp
@@ -0,0 +1,448 @@
+#pragma once
+
+#include <hip/hip_bfloat16.h>
+#include <hip/hip_fp16.h>
+#include <hip/hip_runtime.h>
+
+#include "../common/arch.hpp"
+#include "../ops/common.hpp"
+
+// Sparse prefill attention on gfx11 WMMA: a tile of query rows times a kv head's heads shares every gathered key.
+// q and P stay fp32: each is carried as three 16-bit parts and the products take one matrix op a part.
+
+namespace {
+
+constexpr int kFxRowsPer = 4;                     // query rows a block, times the kv head's query heads
+constexpr int kFxMostGroup = 12;                  // the whole model's query heads a kv head, as sparse.hpp's cap
+constexpr int kFxRows = kFxRowsPer * kFxMostGroup;   // 48, three 16-row WMMA tiles
+constexpr int kFxRowTiles = kFxRows / 16;
+constexpr int kFxKeys = 16;                       // keys a tile: one WMMA column tile, ratio 4 selection blocks
+constexpr int kFxRatio = 4;                       // the indexer's compress_ratio
+constexpr int kFxKeyBlocks = kFxKeys / kFxRatio;
+constexpr int kFxMostDims = 256;                  // the sparse attention's head dim, as sparse.hpp's
+constexpr int kFxKSteps = kFxMostDims / 16;
+constexpr int kFxDimTiles = kFxMostDims / 16;
+constexpr int kFxThreads = 512;
+constexpr int kFxWaves = kFxThreads / 32;
+constexpr int kFxHeadSplit = 4;                 // waves that split one S tile's head dimension
+constexpr int kFxKStepsPerWave = kFxKSteps / kFxHeadSplit;
+constexpr int kFxDimsPerWave = kFxDimTiles / kFxWaves;  // output dimension tiles a wave owns
+constexpr int kFxSoftLanes = 4;                  // threads a probability row, one selection block apiece
+constexpr int kFxKStride = kFxMostDims + 8;      // staged key rows, so 16-byte fragment reads miss no bank group
+constexpr int kFxVtStride = kFxKeys + 8;          // V transposed: a dimension's row of keys, padded for the reads
+constexpr int kFxPStride = kFxKeys + 8;
+constexpr int kFxScoreStride = kFxKeys + 4;      // a float4 of scores per softmax lane
+constexpr int kFxSplit = 3;                      // 16-bit parts a q or P value is carried in
+constexpr int kFxMostWords = 160;                // a part's block range: 5,120 blocks, 20,480 tokens at ratio 4
+
+typedef _Float16 f16x16 __attribute__((ext_vector_type(16)));
+typedef __bf16 bf16x16 __attribute__((ext_vector_type(16)));
+typedef unsigned u32x4 __attribute__((ext_vector_type(4)));  // vec.hpp's lives in tf::rocm, not this namespace
+typedef float f32x8 __attribute__((ext_vector_type(8)));
+
+// The sixteen contiguous k of one WMMA fragment: two 16-byte LDS reads joined by shufflevector.
+__device__ __forceinline__ f16x16 fxfrag16(const uint16_t* p) {
+    const u32x4 lo = *reinterpret_cast<const u32x4*>(p);
+    const u32x4 hi = *reinterpret_cast<const u32x4*>(p + 8);
+    return __builtin_bit_cast(f16x16, __builtin_shufflevector(lo, hi, 0, 1, 2, 3, 4, 5, 6, 7));
+}
+
+__device__ __forceinline__ bf16x16 fxfrag16bf(const uint16_t* p) {
+    const u32x4 lo = *reinterpret_cast<const u32x4*>(p);
+    const u32x4 hi = *reinterpret_cast<const u32x4*>(p + 8);
+    return __builtin_bit_cast(bf16x16, __builtin_shufflevector(lo, hi, 0, 1, 2, 3, 4, 5, 6, 7));
+}
+
+// The cache kind on the matrix cores: its fragment type, its narrowing into the cache type and its fp32 mma.
+template <int KIND>
+struct FxWmma;
+
+template <>
+struct FxWmma<1> {
+    using frag = f16x16;
+    static __device__ __forceinline__ uint16_t near(float v) { return __half_as_ushort(__float2half_rn(v)); }
+    static __device__ __forceinline__ float wide(uint16_t b) { return __half2float(__ushort_as_half(b)); }
+    static __device__ __forceinline__ bool finite(uint16_t b) { return (b & 0x7c00u) != 0x7c00u; }
+    static __device__ __forceinline__ f32x8 mma(frag a, frag b, f32x8 c) {
+#if TF_DEVICE_WMMA_GFX11
+        return __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a, b, c);
+#else
+        return c;  // the launcher takes this kernel only where the caps table says matrix = wmma11
+#endif
+    }
+    static __device__ __forceinline__ frag load(const uint16_t* p) { return fxfrag16(p); }
+};
+
+template <>
+struct FxWmma<2> {
+    using frag = bf16x16;
+    static __device__ __forceinline__ uint16_t near(float v) { return float_to_bf16(v); }
+    static __device__ __forceinline__ float wide(uint16_t b) { return __uint_as_float(static_cast<unsigned>(b) << 16); }
+    static __device__ __forceinline__ bool finite(uint16_t b) { return (b & 0x7f80u) != 0x7f80u; }
+    static __device__ __forceinline__ f32x8 mma(frag a, frag b, f32x8 c) {
+#if TF_DEVICE_WMMA_GFX11
+        return __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32(a, b, c);
+#else
+        return c;  // the launcher takes this kernel only where the caps table says matrix = wmma11
+#endif
+    }
+    static __device__ __forceinline__ frag load(const uint16_t* p) { return fxfrag16bf(p); }
+};
+
+// x as kFxSplit 16-bit parts, largest first: their fp32 sum is x, as three parts hold all 24 bits of a mantissa.
+template <int KIND>
+__device__ __forceinline__ void fxSplit(float x, uint16_t* out) {
+    float rest = x;
+#pragma unroll
+    for (int i = 0; i < kFxSplit; ++i) {
+        out[i] = FxWmma<KIND>::near(rest);
+        rest -= FxWmma<KIND>::wide(out[i]);
+    }
+}
+
+// WMMA column c is query row c % kFxRowsPer and head c / kFxRowsPer; on a tail tile live columns are scattered.
+__device__ __forceinline__ int fxCol(int i, int rows_in) { return (i / rows_in) * kFxRowsPer + (i % rows_in); }
+
+// Whether column c maps to a row and a head inside the tile; equals c < group * rows_in only on a full tile.
+__device__ __forceinline__ bool fxLive(int c, int rows_in, int group) {
+    return (c % kFxRowsPer) < rows_in && (c / kFxRowsPer) < group;
+}
+
+// Block (tile, g, part): query rows tile * kFxRowsPer .. of kv head g, this part's range of the cache's blocks.
+template <int KIND>
+__device__ void sparsePrefillWmma(const float* __restrict__ q, const uint16_t* __restrict__ k,
+                                   const uint16_t* __restrict__ v, const int* __restrict__ ids, long long id_stride,
+                                   const int* __restrict__ count, const int* __restrict__ sparse,
+                                   const int* __restrict__ ends, float* __restrict__ parts_o,
+                                   float* __restrict__ parts_ms, int rows, int heads, int kv_heads, int dims,
+                                   long long cap, float scale, int parts) {
+    using T = FxWmma<KIND>;
+    const int tile = blockIdx.x, g = blockIdx.y, part = blockIdx.z;
+    const int group = heads / kv_heads;
+    const int tiles = (group * kFxRowsPer + 15) / 16;  // 16-row WMMA tiles with a live column, of kFxRowTiles
+    const int t = threadIdx.x, lane = t & 31, wave = t >> 5, sub = lane & 15, half = lane >> 4;
+    const int q0 = tile * kFxRowsPer;
+    const int rows_in = min(kFxRowsPer, rows - q0);
+    const int cols = group * rows_in;  // how many of the 48 WMMA columns are live; see fxCol for WHICH ones
+    const long long head_base = static_cast<long long>(g) * cap * dims;
+
+    // The part's block range, word-aligned outward: every block of the cache falls in exactly one part's range.
+    const int nblocks = static_cast<int>((cap + kFxRatio - 1) / kFxRatio);
+    const int blo = static_cast<int>((static_cast<long long>(part) * nblocks) / parts);
+    const int bhi = min(nblocks, static_cast<int>((static_cast<long long>(part + 1) * nblocks) / parts));
+    const int wbase = blo >> 5;
+    const int words = bhi > blo ? (((bhi - 1) >> 5) - wbase + 1) : 0;
+
+    __shared__ unsigned maps[kFxRowsPer * kFxMostWords];  // the tile's rows' blocks, always kFxMostWords apart
+    __shared__ unsigned uni[kFxMostWords];                 // their union: the key walk
+    __shared__ alignas(16) uint16_t stage[kFxMostDims * kFxVtStride];  // K row-major, then V transposed: one buffer
+    __shared__ alignas(16) uint16_t p_lds[kFxSplit][kFxRows * kFxPStride];
+    __shared__ alignas(16) float s_lds[kFxHeadSplit][kFxRowTiles][16][kFxScoreStride];
+    __shared__ float row_max[kFxRows], row_sum[kFxRows], row_scale[kFxRows];
+    __shared__ unsigned sel_keys[kFxRows];                 // per column: the tile keys it selected, one bit each
+    __shared__ unsigned v_bad[kFxMostDims];                // per dim: the tile keys whose V there is not finite
+    __shared__ float new_max[kFxRows], new_sum[kFxRows];  // staged across the softmax's barrier, then folded in
+
+    // K occupies stage[0 .. kFxKeys * kFxKStride); V^T is staged into the same buffer after the scores are in LDS.
+    uint16_t* k_lds = stage;
+    uint16_t* v_lds = stage;
+
+    for (int i = t; i < kFxRowsPer * kFxMostWords; i += kFxThreads) maps[i] = 0u;
+    // p_lds' padded columns feed P V too: they hold probability zero, so a masked-out key contributes nothing
+    for (int i = t; i < tiles * 16 * kFxPStride; i += kFxThreads) {
+        for (int n = 0; n < kFxSplit; ++n) p_lds[n][i] = 0u;
+    }
+    for (int i = t; i < kFxRows; i += kFxThreads) {
+        row_max[i] = -INFINITY;
+        row_sum[i] = 0.f;
+        row_scale[i] = 1.f;
+        sel_keys[i] = 0u;
+    }
+    __syncthreads();
+
+    // The rows' bitmaps: the block of every listed token, and for a dense row every block with a token below ends.
+    for (int rl = 0; rl < kFxRowsPer; ++rl) {
+        const int r = q0 + rl;
+        if (r < rows) {
+            unsigned* map = maps + rl * kFxMostWords;
+            const bool listed = sparse != nullptr && ids != nullptr && sparse[r] != 0;
+            const int end = ends[r];
+            if (listed) {
+                const int n = count[r];
+                for (int j = t; j < n; j += kFxThreads) {
+                    const int b = ids[r * id_stride + j] >> 2;
+                    const int iw = (b >> 5) - wbase;
+                    if (iw >= 0 && iw < words) atomicOr(&map[iw], 1u << (b & 31));
+                }
+            } else {
+                const int last = min(bhi, (end + kFxRatio - 1) / kFxRatio);
+                for (int b = blo + t; b < last; b += kFxThreads) atomicOr(&map[(b >> 5) - wbase], 1u << (b & 31));
+            }
+        }
+        __syncthreads();  // every row's atomicOr of this pass is visible before the union folds them
+    }
+    for (int iw = t; iw < words; iw += kFxThreads) {
+        unsigned u = 0u;
+        for (int rl = 0; rl < kFxRowsPer; ++rl) u |= maps[rl * kFxMostWords + iw];
+        const int base = (wbase + iw) << 5;
+        if (base + 32 > bhi) u &= (1u << (bhi - base)) - 1u;  // past the part is another part's
+        uni[iw] = u;
+    }
+    __syncthreads();
+
+    // Q into fragments: wave w scores row tile w / 4 over head dimension quarter w % 4; its A row is sub.
+    const int rb = wave / kFxHeadSplit, dh = wave % kFxHeadSplit;
+    typename T::frag q_frag[kFxSplit][kFxKStepsPerWave];
+    {
+        const int c = rb * 16 + sub;
+        const bool live = (c % kFxRowsPer) < rows_in && c / kFxRowsPer < group;  // not c < cols: see fxCol
+        const float* qrow = q + (static_cast<long long>(live ? q0 + (c % kFxRowsPer) : 0) * heads +
+                                 g * group + (live ? c / kFxRowsPer : 0)) *
+                                    dims;
+#pragma unroll
+        for (int ks = 0; ks < kFxKStepsPerWave; ++ks) {
+            const int d0 = (dh * kFxKStepsPerWave + ks) * 16;
+            alignas(16) uint16_t tmp[kFxSplit][16];  // T::load reads it as two uint4: 16 bytes, so it must be aligned
+#pragma unroll
+            for (int i = 0; i < 16; ++i) {
+                const int d = d0 + i;
+                uint16_t part[kFxSplit];
+                fxSplit<KIND>((live && d < dims) ? qrow[d] * scale : 0.f, part);
+#pragma unroll
+                for (int n = 0; n < kFxSplit; ++n) tmp[n][i] = part[n];
+            }
+#pragma unroll
+            for (int n = 0; n < kFxSplit; ++n) q_frag[n][ks] = T::load(&tmp[n][0]);
+        }
+    }
+
+    f32x8 o_acc[kFxRowTiles][kFxDimsPerWave];
+#pragma unroll
+    for (int rt = 0; rt < kFxRowTiles; ++rt)
+#pragma unroll
+        for (int dt = 0; dt < kFxDimsPerWave; ++dt) o_acc[rt][dt] = f32x8{0, 0, 0, 0, 0, 0, 0, 0};
+
+    // The union walk, uniform across the block: set bits of uni in increasing block order, [blo, bhi) only.
+    int cw = 0;
+    unsigned bits = words > 0 ? (uni[0] & (~0u << (blo & 31))) : 0u;
+
+    while (true) {
+        int blk[kFxKeyBlocks];
+        int n = 0;
+        while (n < kFxKeyBlocks) {
+            while (bits == 0u) {
+                ++cw;
+                if (cw >= words) break;
+                bits = uni[cw];
+            }
+            if (bits == 0u) break;
+            blk[n++] = ((wbase + cw) << 5) + __builtin_ctz(bits);
+            bits &= bits - 1u;
+        }
+        if (n == 0) break;                        // the part's union is spent: every other part holds these blocks
+        while (n < kFxKeyBlocks) blk[n++] = -1;  // pad: its keys read key 0, score -inf, hold probability zero
+
+        // The tile's keys: key i is block blk[i / ratio] at offset i % ratio, from uniform blk[].
+        int key[kFxKeys];
+#pragma unroll
+        for (int i = 0; i < kFxKeys; ++i) {
+            const int b = blk[i / kFxRatio];
+            key[i] = b < 0 ? 0 : ((b << 2) + (i % kFxRatio));
+        }
+        for (int d = t; d < kFxMostDims; d += kFxThreads) v_bad[d] = 0u;
+        __syncthreads();
+
+        // stage K: one 16-byte load per lane, a wave over one key row, rows conflict-free at stride kFxKStride
+        constexpr int kKRegs = (kFxKeys * (kFxMostDims / 8)) / kFxThreads;
+#pragma unroll
+        for (int r = 0; r < kKRegs; ++r) {
+            const int idx = t + r * kFxThreads;
+            const int krow = idx / (kFxMostDims / 8), d8 = (idx % (kFxMostDims / 8)) * 8;
+            if (d8 + 8 <= dims) {
+                const uint4 packed =
+                    *reinterpret_cast<const uint4*>(k + head_base + static_cast<long long>(key[krow]) * dims + d8);
+                *reinterpret_cast<uint4*>(&k_lds[krow * kFxKStride + d8]) = packed;
+            }
+        }
+        __syncthreads();
+
+        // S = Q K^T: wave w scores row tile rb against the 16 keys over head quarter dh.
+        if (rb < tiles) {
+            f32x8 s_acc = f32x8{0, 0, 0, 0, 0, 0, 0, 0};
+#pragma unroll
+            for (int ks = 0; ks < kFxKStepsPerWave; ++ks) {
+                const int d0 = (dh * kFxKStepsPerWave + ks) * 16;
+                if (d0 >= dims) break;
+                const auto kf = T::load(&k_lds[sub * kFxKStride + d0]);
+#pragma unroll
+                for (int n = kFxSplit - 1; n >= 0; --n) s_acc = T::mma(q_frag[n][ks], kf, s_acc);  // smallest first
+            }
+#pragma unroll
+            for (int i = 0; i < 8; ++i) s_lds[dh][rb][2 * i + half][sub] = s_acc[i];
+        }
+        __syncthreads();
+
+        // The fp32 online softmax of every live column: kFxSoftLanes threads a column, one selection block each.
+        if (t < cols * kFxSoftLanes) {
+            const int c = fxCol(t / kFxSoftLanes, rows_in), seg = t % kFxSoftLanes;
+            const int rt = c >> 4, row = c & 15;
+            const int rl = c % kFxRowsPer, r = q0 + rl;
+            const int b = blk[seg];
+            const int iw = (b >> 5) - wbase;
+            const bool marked =
+                b >= 0 && iw >= 0 && iw < words && ((maps[rl * kFxMostWords + iw] >> (b & 31)) & 1u) != 0u;
+            float vals[kFxSoftLanes];
+            unsigned chosen = 0u;
+#pragma unroll
+            for (int m = 0; m < kFxSoftLanes; ++m) {
+                const int kpos = (b << 2) + m;
+                chosen |= (marked && kpos < ends[r]) ? 1u << (seg * kFxSoftLanes + m) : 0u;
+                vals[m] = (marked && kpos < ends[r])
+                            ? ((s_lds[0][rt][row][seg * kFxSoftLanes + m] + s_lds[1][rt][row][seg * kFxSoftLanes + m]) +
+                               s_lds[2][rt][row][seg * kFxSoftLanes + m]) +
+                                  s_lds[3][rt][row][seg * kFxSoftLanes + m]
+                            : -INFINITY;
+            }
+            float pmax = -INFINITY;
+#pragma unroll
+            for (int m = 0; m < kFxSoftLanes; ++m) pmax = fmaxf(pmax, vals[m]);
+#pragma unroll
+            for (int off = 1; off < kFxSoftLanes; off <<= 1) pmax = fmaxf(pmax, __shfl_xor(pmax, off, kFxSoftLanes));
+            const float prev = row_max[c];
+            const float top = fmaxf(prev, pmax);
+            const float f = prev == -INFINITY ? 0.f : expf(prev - top);
+            float psum = 0.f;
+#pragma unroll
+            for (int m = 0; m < kFxSoftLanes; ++m) {
+                const float e = vals[m] == -INFINITY ? 0.f : expf(vals[m] - top);
+                psum += e;
+                uint16_t part[kFxSplit];
+                fxSplit<KIND>(e, part);
+#pragma unroll
+                for (int n = 0; n < kFxSplit; ++n) p_lds[n][c * kFxPStride + seg * kFxSoftLanes + m] = part[n];
+            }
+#pragma unroll
+            for (int off = 1; off < kFxSoftLanes; off <<= 1) psum += __shfl_xor(psum, off, kFxSoftLanes);
+            // the running state is staged in new_* and folded after the barrier, since every lane reads row_max above
+            for (int off = 1; off < kFxSoftLanes; off <<= 1) chosen |= __shfl_xor(chosen, off, kFxSoftLanes);
+            if ((lane & (kFxSoftLanes - 1)) == 0) {
+                sel_keys[c] = chosen;
+                new_max[c] = top;
+                new_sum[c] = row_sum[c] * f + psum;
+                row_scale[c] = f;
+            }
+        }
+        __syncthreads();
+        // the same packed→column walk the softmax used: on a tail tile the live columns are not contiguous
+        for (int i = t; i < cols; i += kFxThreads) {
+            const int c = fxCol(i, rows_in);
+            row_max[c] = new_max[c];
+            row_sum[c] = new_sum[c];
+        }
+
+        // Stage V transposed into the buffer K vacated: a thread takes one key and 16 dims, storing v_lds[dim][key].
+        constexpr int kVRegs = (kFxKeys * kFxMostDims) / (kFxThreads * 8);
+        const int v_key = t % kFxKeys;
+        const int v_off = t / kFxKeys;  // 0 .. 15: which 16-dim span of the head this thread covers
+#pragma unroll
+        for (int r = 0; r < kVRegs; ++r) {
+            const int d0 = v_off * (kVRegs * 8) + r * 8;
+            if (d0 + 8 > dims) break;
+            const uint4 packed =
+                *reinterpret_cast<const uint4*>(v + head_base + static_cast<long long>(key[v_key]) * dims + d0);
+            const uint16_t* e = reinterpret_cast<const uint16_t*>(&packed);
+#pragma unroll
+            for (int m = 0; m < 8; ++m) {
+                const bool bad = !T::finite(e[m]);  // staged as zero: P's zero weight times a NaN would reach every row
+                if (bad) atomicOr(&v_bad[d0 + m], 1u << v_key);
+                v_lds[(d0 + m) * kFxVtStride + v_key] = bad ? static_cast<uint16_t>(0) : e[m];
+            }
+        }
+        __syncthreads();
+
+        // O = P V on the accumulator rescaled by the tile factors; wave w takes dim tile w.
+#pragma unroll
+        for (int rt = 0; rt < kFxRowTiles; ++rt) {
+#pragma unroll
+            for (int i = 0; i < 8; ++i) {
+                const float f = row_scale[rt * 16 + 2 * i + half];
+#pragma unroll
+                for (int dt = 0; dt < kFxDimsPerWave; ++dt) o_acc[rt][dt][i] *= f;
+            }
+        }
+#pragma unroll
+        for (int dt = 0; dt < kFxDimsPerWave; ++dt) {
+            const int dim_tile = wave + dt * kFxWaves;
+            if (dim_tile * 16 >= dims) continue;
+            const auto vf = T::load(&v_lds[(dim_tile * 16 + sub) * kFxVtStride]);
+#pragma unroll
+            for (int rt = 0; rt < kFxRowTiles; ++rt) {
+                if (rt >= tiles) continue;
+#pragma unroll
+                for (int n = kFxSplit - 1; n >= 0; --n) {
+                    const auto pf = T::load(&p_lds[n][(rt * 16 + sub) * kFxPStride]);
+                    o_acc[rt][dt] = T::mma(pf, vf, o_acc[rt][dt]);
+                }
+            }
+            const unsigned bad = v_bad[dim_tile * 16 + sub];
+            if (bad != 0u) {  // a row that selected a non-finite V gets NaN in that dim, as the scalar kernel's 0 * NaN
+#pragma unroll
+                for (int rt = 0; rt < kFxRowTiles; ++rt) {
+#pragma unroll
+                    for (int i = 0; i < 8; ++i) {
+                        if ((sel_keys[rt * 16 + 2 * i + half] & bad) != 0u) o_acc[rt][dt][i] = NAN;
+                    }
+                }
+            }
+        }
+        __syncthreads();
+    }
+
+    // The part's unnormalized output and its (max, sum) for tf_fn_sparse_merge; an empty part stores zeros, -inf, 0.
+#pragma unroll
+    for (int dt = 0; dt < kFxDimsPerWave; ++dt) {
+        const int dim_tile = wave + dt * kFxWaves;
+        if (dim_tile * 16 >= dims) continue;
+#pragma unroll
+        for (int rt = 0; rt < kFxRowTiles; ++rt) {
+            if (rt >= tiles) continue;
+#pragma unroll
+            for (int i = 0; i < 8; ++i) {
+                const int c = rt * 16 + 2 * i + half;
+                if (!fxLive(c, rows_in, group)) continue;  // padded columns belong to no row: store nothing
+                const int d = dim_tile * 16 + sub;
+                if (d >= dims) continue;
+                const long long out =
+                    (static_cast<long long>(q0 + (c % kFxRowsPer)) * heads + g * group + (c / kFxRowsPer)) * parts +
+                    part;
+                parts_o[out * dims + d] = reinterpret_cast<const float*>(&o_acc[rt][dt])[i];
+            }
+        }
+    }
+    if (t < cols) {
+        const int c = fxCol(t, rows_in);
+        const long long out =
+            (static_cast<long long>(q0 + (c % kFxRowsPer)) * heads + g * group + (c / kFxRowsPer)) * parts + part;
+        parts_ms[out * 2] = row_max[c];
+        parts_ms[out * 2 + 1] = row_sum[c];
+    }
+}
+
+// The two 16-bit cache kinds; cache_kind stays in the signature so the argument list matches tf_fn_sparse_parts.
+#define TF_FX_ENTRY(NAME, KIND)                                                                        \
+    extern "C" __global__ void __launch_bounds__(kFxThreads) NAME(                                     \
+        const float* __restrict__ q, const void* __restrict__ k, const void* __restrict__ v,          \
+        int cache_kind, const int* __restrict__ ids, long long id_stride, const int* __restrict__ count, \
+        const int* __restrict__ sparse, const int* __restrict__ ends, float* __restrict__ parts_o,     \
+        float* __restrict__ parts_ms, int rows, int heads, int kv_heads, int dims, long long cap,      \
+        float scale, int parts) {                                                                      \
+        (void)cache_kind;                                                                              \
+        sparsePrefillWmma<KIND>(q, static_cast<const uint16_t*>(k), static_cast<const uint16_t*>(v),   \
+                                ids, id_stride, count, sparse, ends, parts_o, parts_ms, rows, heads,   \
+                                kv_heads, dims, cap, scale, parts);                                    \
+    }
+
+TF_FX_ENTRY(tf_fn_sparse_prefill_wmma_bf16, 2)
+TF_FX_ENTRY(tf_fn_sparse_prefill_wmma_f16, 1)
+
+}  // namespace
diff --git a/zig/src/core/contract/hip_records.zig b/zig/src/core/contract/hip_records.zig
index 6d8ee06f85..d509834fd9 100644
--- a/zig/src/core/contract/hip_records.zig
+++ b/zig/src/core/contract/hip_records.zig
@@ -23,6 +23,21 @@ fn relu2Tile(comptime rows: u32, comptime cb: u32) limits.HipLaunch {
     };
 }
 
+/// Flash Next's WMMA sparse prefill attention for one cache kind; gfx1030 and gfx1201 compile it to a stub.
+fn sparsePrefill(comptime kind: []const u8) Record {
+    const symbol = "tf_fn_sparse_prefill_wmma_" ++ kind;
+    const stub: U = .{ .vgprs = 80, .sgprs = 58, .lds = 20736, .private = 0, .bound = 512, .wave = 32 };
+    return .{
+        .launch = .{ .name = "flashnext " ++ symbol, .symbol = symbol, .image = "flashnext", .threads = 512 },
+        .usage = .{
+            stub,
+            .{ .vgprs = 190, .sgprs = 70, .lds = 39936, .private = 0, .bound = 512, .wave = 32 },
+            .{ .vgprs = 189, .sgprs = 71, .lds = 39936, .private = 0, .bound = 512, .wave = 32 },
+            .{ .vgprs = 90, .sgprs = 59, .lds = 20736, .private = 0, .bound = 512, .wave = 32 },
+        },
+    };
+}
+
 /// HIP kernels as launched; tf-hip-kernel-contract limits DIR re-derives the usage records from the code objects.
 const own_records = [_]Record{
     .{
@@ -461,6 +476,8 @@ const own_records = [_]Record{
             .{ .vgprs = 36, .sgprs = 32, .lds = 50688, .private = 0, .bound = 1024, .wave = 32 },
         },
     },
+    sparsePrefill("f16"),
+    sparsePrefill("bf16"),
 } ++ @import("hip_records_pair.zig").records;
 
 pub const records = own_records ++ @import("hip_records_stream.zig").records;
diff --git a/zig/src/families/flashnext/hip/engine.zig b/zig/src/families/flashnext/hip/engine.zig
index ce3b74fd5e..d27adc4f82 100644
--- a/zig/src/families/flashnext/hip/engine.zig
+++ b/zig/src/families/flashnext/hip/engine.zig
@@ -116,6 +116,10 @@ pub const Engine = struct {
         _ = try steps.done(.{ 0, 0 }, .first);
         e.cfg = try config.Config.read(gpa, io, dir);
         errdefer e.cfg.deinit();
+        model.checkSparseRatio(o.policy.sparse == .wmma, e.cfg.index_ratio) catch |err| {
+            std.log.err("sparse=wmma needs indexer_compress_ratio {d}; this model has {d}", .{ model.sparse_wmma_ratio, e.cfg.index_ratio });
+            return err;
+        };
         e.ck = try core.checkpoint.Checkpoint.openModel(gpa, io, dir);
         errdefer e.ck.close();
         e.keep = .init(gpa);
diff --git a/zig/src/families/flashnext/hip/mixers.zig b/zig/src/families/flashnext/hip/mixers.zig
index 6c25d79412..29f01e1b89 100644
--- a/zig/src/families/flashnext/hip/mixers.zig
+++ b/zig/src/families/flashnext/hip/mixers.zig
@@ -404,6 +404,8 @@ fn attend(
             @intCast(key_stride),
             @ptrFromInt(count),
             if (sparse_flags) |f| @ptrFromInt(f) else null,
+            // a prompt's row ends select the WMMA tile; a lane round's rows keep the scalar kernel's bits
+            if (o.prefill) @ptrFromInt(dev_ends) else null,
             @ptrFromInt(po),
             @ptrFromInt(pm),
             @ptrFromInt(att + r0 * d.heads * hd * 4),
diff --git a/zig/src/families/flashnext/hip/model.zig b/zig/src/families/flashnext/hip/model.zig
index 56ecce8cc9..65a67e1981 100644
--- a/zig/src/families/flashnext/hip/model.zig
+++ b/zig/src/families/flashnext/hip/model.zig
@@ -148,6 +148,23 @@ pub const Model = struct {
     mtp: ?Head = null,
 };
 
+/// The WMMA sparse tile is written for the model's compress ratio; any other ratio must keep the scalar kernel.
+pub const sparse_wmma_ratio: usize = 4;
+
+/// Refuses `sparse=wmma` on a model whose indexer compress ratio is not the tile's.
+pub fn checkSparseRatio(wmma: bool, ratio: usize) error{SparseWmmaRatio}!void {
+    if (wmma and ratio != sparse_wmma_ratio) return error.SparseWmmaRatio;
+}
+
+test "sparse=wmma is refused on a compress ratio other than the tile's, and the scalar kernel takes any" {
+    var c = try config.parse(std.testing.allocator, @embedFile("../fixtures/config.json"));
+    defer c.deinit();
+    try checkSparseRatio(true, c.index_ratio);
+    try checkSparseRatio(false, 2);
+    try std.testing.expectError(error.SparseWmmaRatio, checkSparseRatio(true, 2));
+    try std.testing.expectError(error.SparseWmmaRatio, checkSparseRatio(true, 8));
+}
+
 test "a rank's dimensions: heads, value heads and the vocabulary split, the kv heads held two ranks a head" {
     var c = try config.parse(std.testing.allocator, @embedFile("../fixtures/config.json"));
     defer c.deinit();
diff --git a/zig/src/hip/kernel_tests.zig b/zig/src/hip/kernel_tests.zig
index 78dc1da06f..f5819f2758 100644
--- a/zig/src/hip/kernel_tests.zig
+++ b/zig/src/hip/kernel_tests.zig
@@ -31,6 +31,7 @@ test {
     _ = @import("kernel_tests/graph_split.zig");
     _ = @import("kernel_tests/flashnext.zig");
     _ = @import("kernel_tests/flashnext_attn.zig");
+    _ = @import("kernel_tests/flashnext_sparse_prefill.zig");
     _ = @import("kernel_tests/nemotron/moe.zig");
     _ = @import("kernel_tests/nemotron/pick.zig");
     _ = @import("kernel_tests/nemotron/attention.zig");
diff --git a/zig/src/hip/kernel_tests/flashnext_attn.zig b/zig/src/hip/kernel_tests/flashnext_attn.zig
index 9523dd1d6e..0f2b36cff0 100644
--- a/zig/src/hip/kernel_tests/flashnext_attn.zig
+++ b/zig/src/hip/kernel_tests/flashnext_attn.zig
@@ -341,6 +341,8 @@ fn sparseCase(t: *Rig, comptime heads: usize, comptime kv_heads: usize) !void {
         stride,
         ptr(?[*]const i32, dc),
         ptr(?[*]const i32, ds),
+        // ends: null, so this case stays on the scalar kernel -- its ids are not the indexer's block selections
+        null,
         ptr(?[*]f32, po),
         ptr(?[*]f32, pm),
         ptr(?[*]f32, out),
@@ -362,6 +364,7 @@ fn sparseCase(t: *Rig, comptime heads: usize, comptime kv_heads: usize) !void {
         stride,
         ptr(?[*]const i32, dc),
         ptr(?[*]const i32, ds),
+        null,
         ptr(?[*]f32, po),
         ptr(?[*]f32, pm),
         ptr(?[*]f32, one),
diff --git a/zig/src/hip/kernel_tests/flashnext_sparse_prefill.zig b/zig/src/hip/kernel_tests/flashnext_sparse_prefill.zig
new file mode 100644
index 0000000000..97f7dcd739
--- /dev/null
+++ b/zig/src/hip/kernel_tests/flashnext_sparse_prefill.zig
@@ -0,0 +1,215 @@
+//! Flash Next's sparse prefill attention, the scalar and the WMMA kernel, against fp64, on idx_select's lists.
+const std = @import("std");
+const rig = @import("rig.zig");
+const fx = @import("flashnext.zig");
+const Rig = rig.Rig;
+const gpa = rig.gpa;
+const ptr = fx.ptr;
+
+test "sparse prefill: the scalar and the WMMA kernels match fp64 on full and tail slices" {
+    const t = try Rig.open(1 << 22);
+    defer t.close();
+    // 12 query heads a kv head over a full slice, and 6 over a slice of rows that is not a multiple of 4
+    try case(t, 24, 2, 64, false);
+    try case(t, 6, 1, 17, false); // rows % 4 != 0: a tile of two rows, and group 6 pads its 16-row tile
+}
+
+test "sparse prefill: a deep cache with the indexer's full budget, the two kernels' time and their difference" {
+    const t = try Rig.open(1 << 27);
+    defer t.close();
+    try case(t, 24, 2, 64, true);
+}
+
+fn nowNs() u64 {
+    var ts: std.c.timespec = undefined;
+    _ = std.c.clock_gettime(.MONOTONIC, &ts);
+    return @as(u64, @intCast(ts.sec)) * 1_000_000_000 + @as(u64, @intCast(ts.nsec));
+}
+
+fn case(t: *Rig, comptime heads: usize, comptime kv_heads: usize, comptime rows: usize, comptime deep: bool) !void {
+    var rng: rig.Rng = .{ .state = 0x5EED + heads + rows };
+    const dims = 256;
+    const ratio = 4;
+    const parts = 16;
+    const cap = if (deep) 16384 else 1024;
+    const stride = if (deep) 2064 else 520; // id_stride: more than the deepest row's count
+    const scale: f32 = 1.0 / @sqrt(@as(f32, dims));
+    const group = heads / kv_heads;
+
+    const hk = try fx.acts(t, &rng, kv_heads * cap * dims, 1.0);
+    defer gpa.free(hk);
+    const hv = try fx.acts(t, &rng, kv_heads * cap * dims, 1.0);
+    defer gpa.free(hv);
+    const hq = try gpa.alloc(f32, rows * heads * dims);
+    defer gpa.free(hq);
+    for (hq) |*x| x.* = rng.unit() * 4.0;
+
+    // Ends march forward and neighbours share blocks; blocks are distinct, a repeat would count twice in one kernel.
+    var ids = try gpa.alloc(i32, rows * stride);
+    defer gpa.free(ids);
+    @memset(ids, 0);
+    const count = try gpa.alloc(i32, rows);
+    defer gpa.free(count);
+    const sparse = try gpa.alloc(i32, rows);
+    defer gpa.free(sparse);
+    const ends = try gpa.alloc(i32, rows);
+    defer gpa.free(ends);
+    for (0..rows) |r| {
+        // ends is not a multiple of ratio on three rows in four, so the incomplete tail is 1..3 tokens
+        ends[r] = @as(i32, @intCast((if (deep) cap - rows * 12 else 200) + r * 12 + (r % 4)));
+        const nb: usize = @as(usize, @intCast(ends[r])) / ratio; // the blocks the row may select
+        if (r == 0) { // dense: the first ends[0] keys, unlisted
+            sparse[r] = 0;
+            count[r] = ends[r];
+            continue;
+        }
+        sparse[r] = 1;
+        const want: usize = @min(if (deep) 512 else if (r < 8) @as(usize, 5) else 120, nb - 8);
+        const start: usize = if (nb > want) (r * 3) % (nb - want) else 0; // neighbours differ by three blocks
+        var w: usize = 0;
+        for (0..want) |i| {
+            const b = start + i; // distinct, increasing, all below nb
+            for (0..ratio) |o| ids[r * stride + w + o] = @intCast(b * ratio + o);
+            w += ratio;
+        }
+        var k = nb * ratio; // the incomplete tail past every complete block
+        while (k < @as(usize, @intCast(ends[r]))) : (k += 1) {
+            ids[r * stride + w] = @intCast(k);
+            w += 1;
+        }
+        count[r] = @intCast(w);
+    }
+
+    var dk = try t.upload(hk);
+    defer dk.free();
+    var dv = try t.upload(hv);
+    defer dv.free();
+    var dq = try t.upload(hq);
+    defer dq.free();
+    var di = try t.upload(ids);
+    defer di.free();
+    var dc = try t.upload(count);
+    defer dc.free();
+    var ds = try t.upload(sparse);
+    defer ds.free();
+    var de = try t.upload(ends);
+    defer de.free();
+    var po = try t.alloc(rows * heads * parts * dims * 4);
+    defer po.free();
+    var pm = try t.alloc(rows * heads * parts * 2 * 4);
+    defer pm.free();
+    var out = try t.alloc(rows * heads * dims * 4);
+    defer out.free();
+    const ck: c_int = if (t.bf16) 2 else 1;
+
+    var err: [2]f64 = undefined;
+    var ms: [2]f64 = undefined; // the kernel pair's time a launch
+    var dev: [2]f64 = undefined; // vs the scalar kernel's output, the largest |difference|
+    var scalar_out: []f32 = &.{};
+    defer gpa.free(scalar_out);
+    for (0..2) |which| {
+        t.on.sparse_wmma = which != 0;
+        try launch(t, dq, dk, dv, ck, di, stride, dc, ds, de, po, pm, out, rows, heads, kv_heads, dims, cap, scale, parts);
+        try t.stream.synchronize();
+        const t0 = nowNs();
+        for (0..10) |_| try launch(t, dq, dk, dv, ck, di, stride, dc, ds, de, po, pm, out, rows, heads, kv_heads, dims, cap, scale, parts);
+        try t.stream.synchronize();
+        ms[which] = @as(f64, @floatFromInt(nowNs() - t0)) / 10e6;
+        const got = try rig.download(f32, out);
+        defer if (which != 0) gpa.free(got);
+        if (which == 0) scalar_out = got;
+        dev[which] = 0;
+        for (got, scalar_out) |a, b| dev[which] = @max(dev[which], @abs(@as(f64, a) - b));
+        err[which] = 0;
+        if (deep) continue; // the fp64 reference over 2,048 keys a row is slow; the shallow cases hold the error
+        for (0..rows) |r| {
+            for (0..heads) |h| {
+                const kvh = h / group;
+                const listed = sparse[r] != 0;
+                const n: usize = @intCast(count[r]);
+                const end: usize = @intCast(ends[r]);
+                const w = try gpa.alloc(f64, n);
+                defer gpa.free(w);
+                var top: f64 = -std.math.inf(f64);
+                for (0..n) |j| {
+                    const key: usize = if (listed) @intCast(ids[r * stride + j]) else j;
+                    var sc: f64 = 0;
+                    for (0..dims) |d| sc += @as(f64, hq[(r * heads + h) * dims + d] * scale) *
+                        t.value(hk[(kvh * cap + key) * dims + d]);
+                    w[j] = if (key < end) sc else -std.math.inf(f64);
+                    top = @max(top, w[j]);
+                }
+                var total: f64 = 0;
+                for (w) |*x| {
+                    x.* = @exp(x.* - top);
+                    total += x.*;
+                }
+                for (0..dims) |d| {
+                    var acc: f64 = 0;
+                    for (0..n) |j| {
+                        const key: usize = if (listed) @intCast(ids[r * stride + j]) else j;
+                        acc += w[j] * t.value(hv[(kvh * cap + key) * dims + d]);
+                    }
+                    const e = @abs(@as(f64, got[(r * heads + h) * dims + d]) - acc / total);
+                    err[which] = @max(err[which], e);
+                }
+            }
+        }
+    }
+    std.debug.print("sparse prefill {d} keys a row, ms a launch: scalar {d:.3} wmma {d:.3}\n", .{ if (deep) 2048 else 480, ms[0], ms[1] });
+    if (deep) {
+        std.debug.print("sparse prefill deep: wmma differs {e} from scalar\n", .{dev[1]});
+        try rig.expect(dev[1] <= 1e-5, "the WMMA tile differs {e} from the scalar kernel", .{dev[1]});
+        return;
+    }
+    std.debug.print("sparse prefill heads={d} rows={d}: |err| fp64 scalar {e} wmma {e}; wmma vs scalar {e}\n", .{ heads, rows, err[0], err[1], dev[1] });
+    // both sum in fp32 over the cache's values; the tile carries q and p as three 16-bit parts, so it is held as tight
+    try rig.expect(err[0] <= 1e-5, "the scalar kernel drifts {e} from fp64", .{err[0]});
+    try rig.expect(err[1] <= 1e-5, "the WMMA tile drifts {e} from fp64", .{err[1]});
+}
+
+fn launch(
+    t: *Rig,
+    dq: anytype,
+    dk: anytype,
+    dv: anytype,
+    ck: c_int,
+    di: anytype,
+    stride: usize,
+    dc: anytype,
+    ds: anytype,
+    de: anytype,
+    po: anytype,
+    pm: anytype,
+    out: anytype,
+    rows: usize,
+    heads: usize,
+    kv_heads: usize,
+    dims: usize,
+    cap: usize,
+    scale: f32,
+    parts: usize,
+) !void {
+    try t.on.fn_sparse_attention(
+        ptr(?[*]const f32, dq),
+        ptr(?*const anyopaque, dk),
+        ptr(?*const anyopaque, dv),
+        ck,
+        ptr(?[*]const i32, di),
+        @intCast(stride),
+        ptr(?[*]const i32, dc),
+        ptr(?[*]const i32, ds),
+        ptr(?[*]const i32, de),
+        ptr(?[*]f32, po),
+        ptr(?[*]f32, pm),
+        ptr(?[*]f32, out),
+        @intCast(rows),
+        @intCast(heads),
+        @intCast(kv_heads),
+        @intCast(dims),
+        @intCast(cap),
+        scale,
+        @intCast(parts),
+        t.stream.handle,
+    );
+}
diff --git a/zig/src/hip/launches.zig b/zig/src/hip/launches.zig
index ace7c08e7c..6a6b338cce 100644
--- a/zig/src/hip/launches.zig
+++ b/zig/src/hip/launches.zig
@@ -80,6 +80,8 @@ pub const Launcher = struct {
     fuse: bool,
     /// The 64-row prefill attention tile is on.
     wide: bool,
+    /// Flash Next's sparse prefill attention takes the WMMA tile.
+    sparse_wmma: bool,
     /// The MLX affine products and the registry that picks their tiles.
     affine: affine.Kernels,
     /// Flash Next's kernels: its residual streams, n-gram rows, indexer and sparse attention.
@@ -146,8 +148,8 @@ pub const Launcher = struct {
         token_prob: Function,
     };
 
-    /// Which kernels a run takes where two give the same bits; every one defaults on.
-    pub const Choices = struct { fuse: bool = true, wide: bool = true, products: affine.Choices = .{} };
+    /// Which kernels a run takes: fuse and wide give the same bits and default on; sparse_wmma is opt-in.
+    pub const Choices = struct { fuse: bool = true, wide: bool = true, sparse_wmma: bool = false, products: affine.Choices = .{} };
 
     /// Loads the code objects on the current device and resolves every kernel the launchers use.
     pub fn load(r: *const runtime.Runtime, choices: Choices, images: [kernels.group_count][]align(8) const u8) Error!Launcher {
@@ -168,6 +170,8 @@ pub const Launcher = struct {
         const pl = l.mods[@backingInt(kernels.Group.plan)];
         l.fuse = choices.fuse;
         l.wide = choices.wide;
+        // the WMMA tile exists only for gfx11's matrix cores; any other chip keeps the scalar kernel
+        l.sparse_wmma = choices.sparse_wmma and (caps.Caps.of(kernels.arch) orelse return error.Invalid).matrix == .wmma11;
         l.dec = .{ .router = .{ try dec.function("tf_router_decode_f16"), try dec.function("tf_router_decode_bf16") }, .tail = try dec.function("tf_tail"), .conv_split = try dec.function("tf_conv_split"), .rms2 = try dec.function("tf_rms2"), .gnorm_out = try dec.function("tf_gnorm_out"), .select = try dec.function("tf_select_decode") };
         l.plan = .{
             .kv_write = try pl.function("tf_plan_kv_write"),
diff --git a/zig/src/hip/launches/flashnext.zig b/zig/src/hip/launches/flashnext.zig
index dd272698dc..5d7cc1586e 100644
--- a/zig/src/hip/launches/flashnext.zig
+++ b/zig/src/hip/launches/flashnext.zig
@@ -29,6 +29,8 @@ pub const Fns = struct {
     idx_scores: Function,
     idx_select: Function,
     sparse_parts: [2]Function, // groups of up to 6 and up to 12 query heads
+    /// The WMMA prefill tile over 4 rows and a kv head's query heads, q and P kept fp32; by cache kind (fp16, bf16).
+    sparse_prefill_wmma: [2]Function,
     sparse_merge: Function,
 
     pub fn load(m: Module) Error!Fns {
@@ -44,6 +46,10 @@ pub const Fns = struct {
             .idx_scores = try m.function("tf_fn_idx_scores"),
             .idx_select = try m.function("tf_fn_idx_select"),
             .sparse_parts = .{ try m.function("tf_fn_sparse_parts6"), try m.function("tf_fn_sparse_parts12") },
+            .sparse_prefill_wmma = .{
+                try m.function("tf_fn_sparse_prefill_wmma_f16"),
+                try m.function("tf_fn_sparse_prefill_wmma_bf16"),
+            },
             .sparse_merge = try m.function("tf_fn_sparse_merge"),
         };
     }
@@ -381,7 +387,7 @@ pub fn idxSelect(
     try l.go(l.fnx.idx_select, dim(rows, 1, 1), dim(1024, 1, 1), 0, s, &a);
 }
 
-/// Attention of each row's heads over its listed keys (or its first count[r] keys) in `parts` parts, then merged.
+/// Each row's heads over its listed keys (or first count[r]) in `parts` parts, merged; sparse=wmma tiles prefill rows.
 pub fn sparseAttention(
     l: *const Launcher,
     q: CF,
@@ -392,6 +398,7 @@ pub fn sparseAttention(
     id_stride: i64,
     count: CI,
     sparse: CI,
+    ends: CI,
     parts_o: F,
     parts_ms: F,
     out: F,
@@ -406,6 +413,48 @@ pub fn sparseAttention(
 ) Error!void {
     if (dims > 256 or @rem(heads, kv_heads) != 0 or @divExact(heads, kv_heads) > 12)
         return util.invalid("sparseAttention");
+    const wide: usize = if (@divExact(heads, kv_heads) <= 6) 0 else 1;
+    // A part's selection bitmap must fit the kernel's 160 words: its ratio-4 blocks, a word wider at both ends
+    const nblocks: i64 = @divTrunc(cap + 3, @as(i64, 4));
+    const span: i64 = @divTrunc(nblocks, @as(i64, parts)) + 1;
+    const words: i64 = @divTrunc(span + 31, 32) + 1;
+    if (l.sparse_wmma and rows > 1 and ends != null and (cache_kind == 1 or cache_kind == 2) and
+        words <= 160)
+    {
+        var a: Args = .{};
+        inline for (.{
+            ad(q),
+            ad(k),
+            ad(v),
+            cache_kind,
+            ad(ids),
+            id_stride,
+            ad(count),
+            ad(sparse),
+            ad(ends),
+            ad(parts_o),
+            ad(parts_ms),
+            rows,
+            heads,
+            kv_heads,
+            dims,
+            cap,
+            scale,
+            parts,
+        }) |x| try a.add(x);
+        try l.go(
+            l.fnx.sparse_prefill_wmma[@intFromBool(cache_kind == 2)],
+            dim(cdiv(rows, 4), kv_heads, parts),
+            dim(512, 1, 1),
+            0,
+            s,
+            &a,
+        );
+        var b: Args = .{};
+        inline for (.{ ad(parts_o), ad(parts_ms), ad(out), heads, dims, parts }) |x| try b.add(x);
+        try l.go(l.fnx.sparse_merge, dim(heads, rows, 1), dim(256, 1, 1), 0, s, &b);
+        return;
+    }
     var a: Args = .{};
     inline for (.{
         ad(q),
@@ -425,7 +474,6 @@ pub fn sparseAttention(
         scale,
         parts,
     }) |x| try a.add(x);
-    const wide: usize = if (@divExact(heads, kv_heads) <= 6) 0 else 1;
     try l.go(l.fnx.sparse_parts[wide], dim(kv_heads, rows, parts), dim(256, 1, 1), 0, s, &a);
     var b: Args = .{};
     inline for (.{ ad(parts_o), ad(parts_ms), ad(out), heads, dims, parts }) |x| try b.add(x);
diff --git a/zig/src/hip/policy.zig b/zig/src/hip/policy.zig
index 827e99720a..cef3607168 100644
--- a/zig/src/hip/policy.zig
+++ b/zig/src/hip/policy.zig
@@ -7,6 +7,8 @@ const affine = @import("launches/affine.zig");
 pub const Choice = enum { auto, on, off };
 /// The prefill attention tile: `f32` keeps the one-row kernel.
 pub const Attention = enum { auto, f32 };
+/// The sparse prefill attention of Flash Next: `wmma` takes the matrix tile, `scalar` keeps the per-row kernel.
+pub const Sparse = enum { scalar, wmma };
 /// One kernel family switched to its reference, apart from `kernels=reference` that switches them all.
 pub const Pick = enum { auto, reference };
 
@@ -32,6 +34,7 @@ pub const Path = struct {
 pub const Policy = struct {
     matrix: Choice = .auto,
     attention: Attention = .auto,
+    sparse: Sparse = .scalar,
     kernels: Pick = .auto,
     graphs: Choice = .auto,
     /// Reference kernels one at a time: decode stream tile, prefill GEMM tile, fused decode tails.
@@ -74,7 +77,7 @@ pub const Policy = struct {
             .on => .on,
             .off => .off,
         };
-        return .{ .fuse = p.fused(), .wide = p.wideAttention(), .products = .{ .matrix = matrix, .stream = p.streamOn(), .gemm = p.gemmOn() } };
+        return .{ .fuse = p.fused(), .wide = p.wideAttention(), .sparse_wmma = p.sparse == .wmma, .products = .{ .matrix = matrix, .stream = p.streamOn(), .gemm = p.gemmOn() } };
     }
 
     /// Whether the 64-row prefill attention tile is on.
@@ -107,7 +110,7 @@ pub const Policy = struct {
 
     fn set(p: *Policy, key: []const u8, value: []const u8) Error!void {
         const eql = std.mem.eql;
-        if (eql(u8, key, "matrix")) p.matrix = try parseEnum(Choice, value) else if (eql(u8, key, "attention")) p.attention = try parseEnum(Attention, value) else if (eql(u8, key, "kernels")) p.kernels = try parseEnum(Pick, value) else if (eql(u8, key, "graphs")) p.graphs = try parseEnum(Choice, value) else if (eql(u8, key, "stream")) p.stream = try parseEnum(Pick, value) else if (eql(u8, key, "gemm")) p.gemm = try parseEnum(Pick, value) else if (eql(u8, key, "fuse")) p.fuse = try parseEnum(Pick, value) else if (eql(u8, key, "prefill_step")) {
+        if (eql(u8, key, "matrix")) p.matrix = try parseEnum(Choice, value) else if (eql(u8, key, "attention")) p.attention = try parseEnum(Attention, value) else if (eql(u8, key, "sparse")) p.sparse = try parseEnum(Sparse, value) else if (eql(u8, key, "kernels")) p.kernels = try parseEnum(Pick, value) else if (eql(u8, key, "graphs")) p.graphs = try parseEnum(Choice, value) else if (eql(u8, key, "stream")) p.stream = try parseEnum(Pick, value) else if (eql(u8, key, "gemm")) p.gemm = try parseEnum(Pick, value) else if (eql(u8, key, "fuse")) p.fuse = try parseEnum(Pick, value) else if (eql(u8, key, "prefill_step")) {
             p.prefill_step = std.fmt.parseInt(u32, value, 10) catch return error.BadValue;
         } else if (eql(u8, key, "fill_rows")) {
             p.fill_rows = std.fmt.parseInt(u32, value, 10) catch return error.BadValue;
@@ -122,7 +125,7 @@ pub const Policy = struct {
 
     /// One line for the start-up log and server info; default and local fields are left out of the keys.
     pub fn format(p: Policy, w: *std.Io.Writer) std.Io.Writer.Error!void {
-        try w.print("matrix={t},attention={t},kernels={t},graphs={t},prefill_step={d},fill_rows={d}", .{ p.matrix, p.attention, p.kernels, p.graphs, p.prefill_step, p.fill_rows });
+        try w.print("matrix={t},attention={t},sparse={t},kernels={t},graphs={t},prefill_step={d},fill_rows={d}", .{ p.matrix, p.attention, p.sparse, p.kernels, p.graphs, p.prefill_step, p.fill_rows });
         if (p.stream != .auto) try w.print(",stream={t}", .{p.stream});
         if (p.gemm != .auto) try w.print(",gemm={t}", .{p.gemm});
         if (p.fuse != .auto) try w.print(",fuse={t}", .{p.fuse});
@@ -138,7 +141,7 @@ pub const Policy = struct {
         const conf: u64 = @bitCast(p.mtp.confidence);
         const pack = [_]u32{
             @backingInt(p.matrix), @backingInt(p.attention), @backingInt(p.kernels), @backingInt(p.graphs),
-            @backingInt(p.stream), @backingInt(p.gemm),      @backingInt(p.fuse),
+            @backingInt(p.stream), @backingInt(p.gemm),      @backingInt(p.fuse),    @backingInt(p.sparse),
         };
         var nibbles: [2]u32 = .{ 0, 0 };
         for (pack, 0..) |v, i| nibbles[i / 8] |= v << @intCast(4 * (i % 8));
@@ -148,7 +151,7 @@ pub const Policy = struct {
     /// `local` with every field rank 0 sent; a rank keeps its own graph fault and library path.
     pub fn fromWords(w: [word_count]u32, local: Policy) Policy {
         var p = local;
-        var vals: [7]u32 = undefined;
+        var vals: [8]u32 = undefined;
         for (&vals, 0..) |*v, i| v.* = (w[i / 8] >> @intCast(4 * (i % 8))) & 0xf;
         p.matrix = @fromBackingInt(@intCast(vals[0]));
         p.attention = @fromBackingInt(@intCast(vals[1]));
@@ -157,6 +160,7 @@ pub const Policy = struct {
         p.stream = @fromBackingInt(@intCast(vals[4]));
         p.gemm = @fromBackingInt(@intCast(vals[5]));
         p.fuse = @fromBackingInt(@intCast(vals[6]));
+        p.sparse = @fromBackingInt(@intCast(vals[7]));
         p.mtp = .{ .drafts = @truncate(w[1] >> 24), .confidence = @bitCast(@as(u64, w[3]) | @as(u64, w[4]) << 32) };
         p.prefill_step = w[2];
         p.slots = w[5];
@@ -225,7 +229,7 @@ test "apply sets fields and refuses what it does not know" {
 
 test "a policy survives its words" {
     var p: Policy = .{ .mtp = .{ .drafts = 2, .confidence = 0.25 }, .slots = 4, .graph_fail = 1 };
-    try p.apply("matrix=on,attention=f32,graphs=off,fuse=reference,fill_rows=128");
+    try p.apply("matrix=on,attention=f32,sparse=wmma,graphs=off,fuse=reference,fill_rows=128");
     const back = Policy.fromWords(p.words(), .{ .graph_fail = 1 });
     try std.testing.expectEqualDeep(p, back);
 }
diff --git a/zig/tests/hip/contract_main.zig b/zig/tests/hip/contract_main.zig
index 5b07bfefb3..eff82b88c7 100644
--- a/zig/tests/hip/contract_main.zig
+++ b/zig/tests/hip/contract_main.zig
@@ -16,6 +16,7 @@ const relu2_k = @import("contract_relu2.zig");
 const kv_k = @import("contract_kv_write.zig");
 const argmax_k = @import("contract_argmax.zig");
 const pair_k = @import("contract_relu2_pair.zig");
+const sparse_k = @import("contract_sparse_prefill.zig");
 
 const g = core.contract.gates;
 
@@ -136,6 +137,9 @@ fn run(a: std.mem.Allocator, io: std.Io) !usize {
     var pair = pair_k.Pair.init(a, &r, s, &l);
     defer pair.deinit();
     try kernels.append(a, pair.kernel());
+    var sparse = sparse_k.SparsePrefill.init(a, &r, s, &l);
+    defer sparse.deinit();
+    try kernels.append(a, sparse.kernel());
     var t: g.Tally = .{};
     const t0 = ms(io);
     try g.runAll(a, kernels.items, 0x6b63, &t);
diff --git a/zig/tests/hip/contract_sparse_prefill.zig b/zig/tests/hip/contract_sparse_prefill.zig
new file mode 100644
index 0000000000..198863883b
--- /dev/null
+++ b/zig/tests/hip/contract_sparse_prefill.zig
@@ -0,0 +1,273 @@
+//! Flash Next's sparse prefill attention on the WMMA tile under the contract on HIP: rows against an fp64 reference.
+const std = @import("std");
+const hip = @import("hip");
+const core = @import("core");
+const rig_ = @import("contract_rig.zig");
+const kit = @import("contract_kit.zig");
+
+const contract = core.contract;
+const c = contract.spec;
+const g = contract.gates;
+
+const dims = 256;
+const ratio = 4;
+const parts = 16;
+const stride = 512; // ids a row: at most 120 blocks and a tail of three tokens
+const scale: f32 = 1.0 / 16.0;
+
+/// Fp32 sums over a row's keys and the parts merge.
+fn depth(s: c.Shape) usize {
+    return s.k + 2 * parts;
+}
+
+const names = [2][]const u8{
+    "hip tf_fn_sparse_prefill_wmma_f16 and tf_fn_sparse_merge",
+    "hip tf_fn_sparse_prefill_wmma_bf16 and tf_fn_sparse_merge",
+};
+const rule = "a NaN a row did not select never reaches it; a selected value reaches its dim, a selected key the head";
+const why = "q and P stay fp32 as three 16-bit parts, but the matrix cores accumulate in their own order: the fp32 bound";
+
+pub fn spec(k: kit.Kind) c.Spec {
+    return .{
+        .name = names[@intFromBool(k == .bf16)],
+        .accum = .f32,
+        .depth = depth,
+        .ops = 6, // exp, the running rescale, the parts merge and the final divide
+        .out = .f32,
+        .nan = .{ .stated = rule },
+        .tiles = .{ .{ .size = 4, .min = 2 }, .{}, .{} },
+        // a full slice of 12 heads a kv head, a tail slice of a rank's 6, and a deep cache
+        .served = &.{
+            .{ .rows = 64, .n = 24, .k = 1024 },
+            .{ .rows = 17, .n = 6, .k = 1024 },
+            .{ .rows = 16, .n = 12, .k = 4096 },
+        },
+        .sums = .{ .f16 = 0, .bf16 = 0, .why = why, .exact = false },
+    };
+}
+
+/// A slice of rows over a kv head's cache of the kind, in the format tf_fn_idx_select writes: row 0 dense.
+pub const SparsePrefill = struct {
+    rig: rig_.Rig,
+    l: hip.Launcher,
+    k: kit.Kind,
+    arena: std.heap.ArenaAllocator,
+    shape: c.Shape = undefined,
+    fill: g.Fill = .random,
+    heads: usize = 0,
+    group: usize = 0,
+    q: []f32 = &.{},
+    kc: []f64 = &.{},
+    vc: []f64 = &.{},
+    ids: []i32 = &.{},
+    counts: []i32 = &.{},
+    sparse: []i32 = &.{},
+    ends: []i32 = &.{},
+    dev: [9]u64 = @splat(0),
+    out: u64 = 0,
+    msg: [200]u8 = undefined,
+
+    pub fn init(a: std.mem.Allocator, rt: *const hip.Runtime, s: hip.Stream, l: *const hip.Launcher) SparsePrefill {
+        var own = l.*;
+        own.sparse_wmma = true;
+        return .{ .rig = .{ .a = a, .r = rt, .s = s }, .l = own, .k = kit.act(), .arena = .init(a) };
+    }
+
+    pub fn deinit(self: *SparsePrefill) void {
+        self.arena.deinit();
+        self.rig.deinit();
+    }
+
+    pub fn kernel(self: *SparsePrefill) g.Kernel {
+        return .{ .spec = spec(self.k), .ctx = self, .vt = &vt };
+    }
+
+    const vt = g.hooks(SparsePrefill);
+
+    fn rows(self: *const SparsePrefill) usize {
+        return self.shape.rows;
+    }
+
+    /// Causal ends march forward, neighbours share most blocks; rows 1..7 select five blocks, the rest 120.
+    fn select(self: *SparsePrefill, a: std.mem.Allocator) !void {
+        const n = self.rows();
+        const cap = self.shape.k;
+        self.ids = try a.alloc(i32, n * stride);
+        @memset(self.ids, 0);
+        self.counts = try a.alloc(i32, n);
+        self.sparse = try a.alloc(i32, n);
+        self.ends = try a.alloc(i32, n);
+        for (0..n) |r| {
+            const end = cap - n * 12 + r * 12 + r % 4; // not a multiple of ratio on three rows in four
+            self.ends[r] = @intCast(end);
+            const nb = end / ratio;
+            if (r == 0) {
+                self.sparse[r] = 0;
+                self.counts[r] = @intCast(end);
+                continue;
+            }
+            self.sparse[r] = 1;
+            const want = @min(if (r < 8) @as(usize, 5) else 120, nb - 8);
+            const start = (r * 3) % (nb - want);
+            var w: usize = 0;
+            for (0..want * ratio) |i| {
+                self.ids[r * stride + w] = @intCast(start * ratio + i);
+                w += 1;
+            }
+            for (nb * ratio..end) |t| {
+                self.ids[r * stride + w] = @intCast(t);
+                w += 1;
+            }
+            self.counts[r] = @intCast(w);
+        }
+    }
+
+    /// Random: values in [-1, 1) exact in the kind. Stress: q = 0 and integer values, so every weight is 1.
+    pub fn prepare(self: *SparsePrefill, shape: c.Shape, fill: g.Fill, seed: u64) !bool {
+        const group: usize = if (shape.n % 12 == 0) 12 else 6;
+        if (shape.rows < 2 or shape.n % group != 0 or shape.k < 12 * shape.rows + 64 or shape.k > 20000) return false;
+        self.rig.reset();
+        _ = self.arena.reset(.retain_capacity);
+        const a = self.arena.allocator();
+        self.shape = shape;
+        self.fill = fill;
+        self.heads = shape.n;
+        self.group = group;
+        const kv = shape.n / group;
+        var prng = std.Random.DefaultPrng.init(seed);
+        const rng = prng.random();
+        const stress = fill == .stress;
+        self.q = try a.alloc(f32, shape.rows * shape.n * dims);
+        for (self.q) |*x| x.* = if (stress) 0 else @floatCast(rng.float(f64) * 2 - 1);
+        const cache = kv * shape.k * dims;
+        self.kc, const kb = try kit.values(a, rng, self.k, cache, -1, 1);
+        self.vc, const vb = try kit.values(a, rng, self.k, cache, -1, 1);
+        if (stress) for (self.vc, 0..) |*x, i| {
+            x.* = @floatFromInt(rng.intRangeAtMost(i32, -8, 8));
+            kit.put(self.k, vb, i, x.*);
+        };
+        try self.select(a);
+        if (fill == .nonfinite) try self.plant(kb, vb);
+        const ids = try self.rig.input(std.mem.sliceAsBytes(self.ids));
+        const cnt = try self.rig.input(std.mem.sliceAsBytes(self.counts));
+        const sparse = try self.rig.input(std.mem.sliceAsBytes(self.sparse));
+        const ends = try self.rig.input(std.mem.sliceAsBytes(self.ends));
+        const po = try self.rig.scratch(shape.rows * shape.n * parts * dims * 4);
+        const pm = try self.rig.scratch(shape.rows * shape.n * parts * 2 * 4);
+        const dq = try self.rig.input(std.mem.sliceAsBytes(self.q));
+        self.dev = .{ dq, try self.rig.input(kb), try self.rig.input(vb), ids, cnt, sparse, ends, po, pm };
+        self.out = try self.rig.output(shape.rows * shape.n * dims * 4);
+        return true;
+    }
+
+    /// NaN in kv head 0: a key in block 12, a value in block 16 and a value in row 1's tail, past row 0's end.
+    fn plant(self: *SparsePrefill, kb: []u8, vb: []u8) !void {
+        const nan = std.math.nan(f64);
+        const bad_key = 12 * ratio * dims;
+        const val = 16 * ratio * dims + 1;
+        kit.put(self.k, kb, bad_key, nan);
+        kit.put(self.k, vb, val, nan);
+        self.kc[bad_key] = nan;
+        self.vc[val] = nan;
+        const tail: usize = @intCast(@divTrunc(self.ends[1], ratio) * ratio);
+        kit.put(self.k, vb, tail * dims + 5, nan);
+        self.vc[tail * dims + 5] = nan;
+    }
+
+    pub fn run(self: *SparsePrefill) !void {
+        const s = self.shape;
+        const d = self.dev;
+        try self.l.fn_sparse_attention(
+            @ptrFromInt(d[0]),
+            @ptrFromInt(d[1]),
+            @ptrFromInt(d[2]),
+            @backingInt(self.k),
+            @ptrFromInt(d[3]),
+            stride,
+            @ptrFromInt(d[4]),
+            @ptrFromInt(d[5]),
+            @ptrFromInt(d[6]),
+            @ptrFromInt(d[7]),
+            @ptrFromInt(d[8]),
+            @ptrFromInt(self.out),
+            @intCast(s.rows),
+            @intCast(s.n),
+            @intCast(s.n / self.group),
+            dims,
+            @intCast(s.k),
+            scale,
+            parts,
+            self.rig.s.handle,
+        );
+        try self.rig.s.synchronize();
+    }
+
+    pub fn count(self: *SparsePrefill) usize {
+        return self.shape.rows * self.shape.n * dims;
+    }
+
+    /// The row's selected keys below its end: its listed tokens, or all of a dense row's.
+    fn key(self: *const SparsePrefill, r: usize, j: usize) ?usize {
+        const t: usize = if (self.sparse[r] != 0) @intCast(self.ids[r * stride + j]) else j;
+        return if (t < self.ends[r]) t else null;
+    }
+
+    /// fp64 attention over the selected keys only; a NaN score zeroes the head, as the merge does.
+    pub fn reference(self: *SparsePrefill, want: []f64, mag: []f64) !void {
+        const a = self.arena.allocator();
+        const w = try a.alloc(f64, @intCast(std.mem.max(i32, self.counts)));
+        defer a.free(w);
+        for (0..self.rows()) |r| for (0..self.heads) |h| {
+            const kvh = h / self.group;
+            const n: usize = @intCast(self.counts[r]);
+            var top = -std.math.inf(f64);
+            var nan = false;
+            for (0..n) |j| {
+                w[j] = std.math.nan(f64);
+                const t = self.key(r, j) orelse continue;
+                var s: f64 = 0;
+                for (0..dims) |d| {
+                    const kv: f64 = self.kc[(kvh * self.shape.k + t) * dims + d];
+                    const p = @as(f64, self.q[(r * self.heads + h) * dims + d]) * scale * kv;
+                    s += p;
+                }
+                w[j] = s;
+                nan = nan or std.math.isNan(s);
+                top = @max(top, s);
+            }
+            var total: f64 = 0;
+            for (0..n) |j| {
+                if (self.key(r, j) == null) {
+                    w[j] = 0;
+                    continue;
+                }
+                w[j] = if (nan) 0 else @exp(w[j] - top);
+                total += w[j];
+            }
+            for (0..dims) |d| {
+                var acc: f64 = 0;
+                var m: f64 = 0;
+                const i = (r * self.heads + h) * dims + d;
+                if (nan) {
+                    want[i] = 0;
+                    mag[i] = 0;
+                    continue;
+                }
+                for (0..n) |j| {
+                    const t = self.key(r, j) orelse continue;
+                    const v = self.vc[(kvh * self.shape.k + t) * dims + d];
+                    acc += w[j] * v;
+                    m += w[j] * @abs(v);
+                }
+                want[i] = acc / total;
+                mag[i] = m / total;
+            }
+        };
+    }
+
+    /// Every sum is exact when every weight is 1 and the values are integers, in any order and either accumulator.
+    pub fn simulate(self: *SparsePrefill, _: c.Accum, _: g.Order, out: []f64) !void {
+        const mag = try self.arena.allocator().alloc(f64, out.len);
+        try self.reference(out, mag);
+    }
+};