diff --git a/tools/zig/cuda_sass_identity.py b/tools/zig/cuda_sass_identity.py
new file mode 100644
index 0000000000..1602578223
--- /dev/null
+++ b/tools/zig/cuda_sass_identity.py
@@ -0,0 +1,198 @@
+"""cuda_sass_identity.py BASE HEAD: cuobjdump -sass per function, base against head, of every image a change reaches."""
+import argparse
+import concurrent.futures as cf
+import hashlib
+import os
+import re
+import subprocess
+import sys
+import tempfile
+from pathlib import Path
+
+SMS = (80, 86, 89, 120, 121)
+NVCC = os.environ.get("NVCC", "/usr/local/cuda/bin/nvcc")
+CUOBJDUMP = os.environ.get("CUOBJDUMP", str(Path(NVCC).with_name("cuobjdump")))
+GENERATED = ("qwen4_gdn", "qwen4_ple")  # the build compiles a generated text; the checked-in unit stands in for it
+CONTROLS = (
+    ("fp4pair constant", "common/qmm_frag.cuh", r'"r"\(0x7E807E80u\)', '"r"(0x7E817E80u)'),
+    ("warp_sum order", "common/reduce.cuh", r"int o = 16; o > 0; o >>= 1\)", "int o = 1; o < 32; o <<= 1)"),
+)
+USAGE = """BASE and HEAD are git revisions of this repository (a merge commit works). Every catalog image
+(zig/build/cuda_kernels.zig) whose unit includes, directly or through other headers, a file under zig/kernels changed
+between them is built from both trees with that tree's catalog flags, torch's nvcc flags and -I zig/kernels/cuda, for
+sm_80/86/89/120/121 as the build passes them (min_sm, the "a" targets of arch-specific units), and compared function
+by function. Two planted controls then prove the comparison can go red: a changed fp4pair constant
+(common/qmm_frag.cuh) and a reversed warp_sum butterfly (common/reduce.cuh), each built against the head. Exit status
+0 only when every function matches and both controls change functions."""
+INCLUDE = re.compile(r'^\s*#\s*include\s*"([^"]+)"', re.M)
+
+
+def git(repo, *args):
+    """Run git in the repository and return its stdout."""
+    return subprocess.run(["git", "-C", repo, *args], capture_output=True, check=True).stdout
+
+
+def export(repo, rev, dest):
+    """The revision's zig/kernels and build files under dest."""
+    dest.mkdir(parents=True)
+    tar = git(repo, "archive", rev, "zig/kernels", "zig/build")
+    subprocess.run(["tar", "-x", "-C", str(dest)], input=tar, check=True)
+
+
+def catalog(tree):
+    """Image name -> (unit, flags, arch_specific, min_sm): the catalog, then images other build files name (h3, ...)."""
+    text = (tree / "zig/build/cuda_kernels.zig").read_text()
+    consts = {m[1]: re.findall(r'"([^"]+)"', m[2]) for m in re.finditer(r"const (\w+) = &\[_\]\[\]const u8\{([^}]*)\}", text)}
+    torch = re.findall(r'"([^"]+)"', re.search(r"torch_flags = \[_\]\[\]const u8\{([^}]*)\}", text)[1])
+    images = {}
+    others = sorted(p for p in (tree / "zig/build").glob("*.zig") if p.name != "cuda_kernels.zig")
+    for f in [tree / "zig/build/cuda_kernels.zig", *others]:
+        for m in re.finditer(r'\.\{ \.name = "(\w+)",((?:[^{}]|\{[^{}]*\})*)\}', f.read_text()):
+            name, rest = m[1], m[2]
+            fl = re.search(r"\.flags = (&\.\{[^}]*\}|\w+)", rest)
+            src = re.search(r'\.src = "([^"]+)"', rest)
+            unit = src[1] if src else name
+            if not fl or name in images or not (tree / "zig/kernels/cuda" / f"{unit}.cu").is_file():
+                continue
+            flags = consts[fl[1]] if fl[1] in consts else re.findall(r'"([^"]+)"', fl[1])
+            ms = re.search(r"\.min_sm = (\d+)", rest)
+            images[name] = (unit, flags, ".arch_specific = true" in rest, int(ms[1]) if ms else 0)
+    return torch, images
+
+
+def includes(tree, unit):
+    """Every file under zig/kernels the unit reaches through quoted includes (conditionals ignored: a superset)."""
+    root = tree / "zig/kernels/cuda"
+    seen, todo = set(), [root / f"{unit}.cu"]
+    while todo:
+        f = todo.pop().resolve()
+        if f in seen or not f.is_file():
+            continue
+        seen.add(f)
+        for inc in INCLUDE.findall(f.read_text(errors="replace")):
+            todo += [f.parent / inc, root / inc]
+    return {p.relative_to(tree.resolve()).as_posix() for p in seen}
+
+
+def build(tree, torch, img, out):
+    """nvcc -fatbin as zig/build/cuda_fatbin.zig runs it, then cuobjdump -sass; returns the dump or an error line."""
+    unit, flags, arch, min_sm = img
+    gens = [f"-gencode=arch=compute_{s}{'a' if arch and s >= 120 else ''},code=sm_{s}{'a' if arch and s >= 120 else ''}"
+            for s in SMS if s >= min_sm]
+    if not gens:
+        return ""
+    out.parent.mkdir(parents=True, exist_ok=True)
+    cmd = [NVCC, "-fatbin", *torch, *flags, f"-I{tree}/zig/kernels/cuda", *gens, "-o", str(out), f"{tree}/zig/kernels/cuda/{unit}.cu"]
+    r = subprocess.run(cmd, capture_output=True, text=True)
+    if r.returncode:
+        return "BUILD FAIL " + r.stderr[-600:]
+    dump = subprocess.run([CUOBJDUMP, "-sass", str(out)], capture_output=True, text=True, check=True).stdout
+    out.unlink()
+    return dump
+
+
+def functions(dump):
+    """(arch, function) -> each instruction's line (any address width) and continuation (second word), collapsed."""
+    res, cur, arch = {}, None, None
+    for line in dump.splitlines():
+        m = re.match(r"\s*arch = (sm_\w+)", line)
+        if m:
+            arch = m[1]
+            continue
+        m = re.match(r"\s*Function : (\S+)", line)
+        if m:
+            cur = (arch, m[1])
+            res[cur] = []
+            continue
+        if cur and re.match(r"\s*(/\*[0-9a-f]{4,}\*/|/\* 0x)", line):
+            res[cur].append(re.sub(r"\s+", " ", line.strip()))
+    return res
+
+
+def digest(lines):
+    """sha256 of a function's instruction lines."""
+    return hashlib.sha256("\n".join(lines).encode()).hexdigest()
+
+
+def compare(fb, fh):
+    """Changed, added and removed (arch, function) keys."""
+    return [k for k in fb if k in fh and fb[k] != fh[k]], [k for k in fh if k not in fb], [k for k in fb if k not in fh]
+
+
+def main():
+    ap = argparse.ArgumentParser(description=__doc__, epilog=USAGE)
+    ap.add_argument("base")
+    ap.add_argument("head")
+    ap.add_argument("--receipt", help="write image, arch, function and both SASS digests here")
+    ap.add_argument("--no-controls", action="store_true")
+    ap.add_argument("-j", type=int, default=4, help="parallel nvcc runs")
+    ap.add_argument("--repo", default=".")
+    a = ap.parse_args()
+    repo = git(a.repo, "rev-parse", "--show-toplevel").decode().strip()
+    revs = [git(repo, "rev-parse", r + "^{commit}").decode().strip() for r in (a.base, a.head)]
+    names = git(repo, "diff", "--name-only", "--no-renames", revs[0], revs[1], "--", "zig/kernels").decode().split()
+    changed = set(names)
+    nvcc = subprocess.run([NVCC, "--version"], capture_output=True, text=True).stdout.strip().splitlines()[-1]
+    print(f"base {revs[0][:12]} head {revs[1][:12]}; {len(changed)} files changed under zig/kernels; {nvcc}")
+    with tempfile.TemporaryDirectory(prefix="sassid-") as tmp:
+        tmp = Path(tmp)
+        trees = {"base": tmp / "base", "head": tmp / "head"}
+        for tag, rev in zip(trees, revs):
+            export(repo, rev, trees[tag])
+        cats = {tag: catalog(t) for tag, t in trees.items()}
+        both = sorted(set(cats["base"][1]) & set(cats["head"][1]))
+        for gone in sorted(set(cats["base"][1]) ^ set(cats["head"][1])):
+            print(f"{gone}: in one catalog only")
+        reach = {tag: {n: includes(trees[tag], cats[tag][1][n][0]) for n in both} for tag in trees}
+        imgs = [n for n in both if reach["base"][n] & changed or reach["head"][n] & changed or cats["base"][1][n] != cats["head"][1][n]]
+        jobs = [(tag, n) for n in imgs for tag in trees]
+        with cf.ThreadPoolExecutor(a.j) as ex:
+            dumps = dict(zip(jobs, ex.map(lambda j: build(trees[j[0]], cats[j[0]][0], cats[j[0]][1][j[1]], tmp / "o" / j[0] / f"{j[1]}.fatbin"), jobs)))
+        bad, total, heads, rows = 0, 0, {}, []
+        for n in imgs:
+            b, h = dumps[("base", n)], dumps[("head", n)]
+            if b.startswith("BUILD FAIL") or h.startswith("BUILD FAIL"):
+                bad += 1
+                print(f"{n}: {b if b.startswith('BUILD FAIL') else h}")
+                continue
+            fb, fh = functions(b), functions(h)
+            heads[n] = fh
+            ch, add, rem = compare(fb, fh)
+            total += len(fh)
+            bad += len(ch) + len(add) + len(rem)
+            note = " (checked-in unit; the build generates this one)" if n in GENERATED else ""
+            print(f"{n}: {len(fh)} functions over {','.join(sorted({k[0] for k in fh}))}; changed {len(ch)} added {len(add)} removed {len(rem)}{note}")
+            for k in (ch + add + rem)[:5]:
+                print("   ", *k)
+            rows += [(n, *k, digest(fb[k]) if k in fb else "-", digest(fh[k]) if k in fh else "-") for k in sorted(set(fb) | set(fh))]
+        print(f"images {len(imgs)} functions {total} differing {bad}")
+        if a.receipt:
+            with open(a.receipt, "w") as f:
+                f.write(f"# cuda_sass_identity base {revs[0]} head {revs[1]}; {nvcc}\n# image\tarch\tfunction\tbase sha256\thead sha256\n")
+                f.writelines("\t".join(r) + "\n" for r in rows)
+        red = True
+        if not a.no_controls:
+            for label, rel, pattern, repl in CONTROLS:
+                ctl = tmp / "ctl"
+                subprocess.run(["rm", "-rf", str(ctl)], check=True)
+                export(repo, revs[1], ctl)
+                f = ctl / "zig/kernels/cuda" / rel
+                text, k = re.subn(pattern, repl, f.read_text(), count=1)
+                hit = [n for n in imgs if f"zig/kernels/cuda/{rel}" in reach["head"][n] and n in heads]
+                if not k or not hit:
+                    print(f"control {label}: could not plant (pattern found {k}, images {len(hit)})")
+                    red = False
+                    continue
+                f.write_text(text)
+                with cf.ThreadPoolExecutor(a.j) as ex:
+                    cd = dict(zip(hit, ex.map(lambda n: build(ctl, cats["head"][0], cats["head"][1][n], tmp / "o" / "ctl" / f"{n}.fatbin"), hit)))
+                moved = {n: len(compare(heads[n], functions(cd[n]))[0]) for n in hit if not cd[n].startswith("BUILD FAIL")}
+                fails = [n for n in hit if cd[n].startswith("BUILD FAIL")]
+                print(f"control {label}: {sum(moved.values())} functions changed in {sum(1 for v in moved.values() if v)} of {len(hit)} images"
+                      + (f", build failed: {' '.join(fails)}" if fails else ""))
+                red = red and sum(moved.values()) > 0 and not fails
+    return 0 if bad == 0 and red else 1
+
+
+if __name__ == "__main__":
+    sys.exit(main())
diff --git a/tools/zig/cuda_sass_identity_test.py b/tools/zig/cuda_sass_identity_test.py
new file mode 100644
index 0000000000..b8900b47ac
--- /dev/null
+++ b/tools/zig/cuda_sass_identity_test.py
@@ -0,0 +1,69 @@
+"""cuda_sass_identity.functions must see every instruction: five-digit addresses and the second encoding word."""
+import unittest
+
+import cuda_sass_identity as sid
+
+WORD2 = " " * 84
+
+
+def dump(late_op, late_word2):
+    """A canned cuobjdump -sass dump: function `same` and function `late`, whose 0x10000 instruction varies."""
+    return f"""
+Fatbin elf code:
+================
+arch = sm_121
+code version = [1,8]
+host = linux
+compile_size = 64bit
+
+	code for sm_121
+		Function : same
+	.headerflags	@"EF_CUDA_SM121 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM121)"
+        /*0000*/                   LDC R1, c[0x0][0x37c] ;                           /* 0x0000df00ff017b82 */
+{WORD2}/* 0x000fe20000000800 */
+        /*10000*/                  EXIT ;                                            /* 0x000000000000794d */
+{WORD2}/* 0x000fea0003800000 */
+		..........
+
+
+		Function : late
+	.headerflags	@"EF_CUDA_SM121 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM121)"
+        /*0000*/                   LDC R1, c[0x0][0x37c] ;                           /* 0x0000df00ff017b82 */
+{WORD2}/* 0x000fe20000000800 */
+        /*fff0*/                   NOP ;                                             /* 0x0000000000007918 */
+{WORD2}/* 0x000fc00000000000 */
+        /*10000*/                  {late_op} ;                                       /* 0x0000000000007918 */
+{WORD2}/* {late_word2} */
+		..........
+"""
+
+
+class FunctionsTest(unittest.TestCase):
+    def changed(self, base, head):
+        return sorted(k[1] for k in sid.compare(sid.functions(base), sid.functions(head))[0])
+
+    def test_five_digit_address(self):
+        """Two functions differing only at 0x10000: the late one changes, the other does not."""
+        base = dump("FADD R2, R2, R3", "0x000fc00000000000")
+        head = dump("FMUL R2, R2, R3", "0x000fc00000000000")
+        self.assertEqual(self.changed(base, head), ["late"])
+        self.assertEqual(self.changed(base, base), [])
+
+    def test_second_encoding_word(self):
+        """Only the continuation line's word differs (a scheduling control change): still a change."""
+        base = dump("NOP", "0x000fc00000000000")
+        head = dump("NOP", "0x000fe20000000000")
+        self.assertEqual(self.changed(base, head), ["late"])
+
+    def test_lines_kept(self):
+        """Each instruction keeps its line and its continuation, in order."""
+        fs = sid.functions(dump("NOP", "0x000fc00000000000"))
+        self.assertEqual(sorted(fs), [("sm_121", "late"), ("sm_121", "same")])
+        self.assertEqual(len(fs[("sm_121", "same")]), 4)
+        self.assertEqual(len(fs[("sm_121", "late")]), 6)
+        self.assertTrue(fs[("sm_121", "same")][2].startswith("/*10000*/ EXIT ;"))
+        self.assertEqual(fs[("sm_121", "same")][3], "/* 0x000fea0003800000 */")
+
+
+if __name__ == "__main__":
+    unittest.main()
diff --git a/zig/build/qwen4_ple_cuda_host.zig b/zig/build/qwen4_ple_cuda_host.zig
index 97299a3f09..7db805af2a 100644
--- a/zig/build/qwen4_ple_cuda_host.zig
+++ b/zig/build/qwen4_ple_cuda_host.zig
@@ -15,11 +15,12 @@ pub fn generate(a: std.mem.Allocator) ![]const u8 {
     var text = try ctl.cuda(a, production);
     text = try edit.once(a, text, reduce, "");
     for ([_][]const u8{ "kss", "qss", "dot", "gss" }) |v| text = try std.mem.replaceOwned(u8, a, text, try std.fmt.allocPrint(a, "block_sum({s}, scratch)", .{v}), try std.fmt.allocPrint(a, "(co_await cpu_block_sum({s}))", .{v}));
-    var math = @embedFile("../kernels/cuda/glue_math.cuh").*;
+    var math = @embedFile("../kernels/cuda/common/glue_math.cuh").*;
     var helpers: []const u8 = &math;
-    helpers = try edit.once(a, helpers, try edit.function(helpers, "float warp_sum("), "");
+    helpers = try edit.once(a, helpers, "#include \"common/reduce.cuh\"", "");
+    helpers = try edit.once(a, helpers, "using tf_common::warp_sum;", "");
     const replacements = [_][2][]const u8{
-        .{ "#pragma once", "" },                       .{ "#include <cuda_bf16.h>", "" },                                                                .{ "#include <cuda_fp16.h>", "" },                                                                .{ "#include \"glue_math.cuh\"", "" },
+        .{ "#pragma once", "" },                       .{ "#include <cuda_bf16.h>", "" },                                                                .{ "#include <cuda_fp16.h>", "" },                                                                .{ "#include \"common/glue_math.cuh\"", "" },
         .{ "__device__ __forceinline__ ", "inline " }, .{ "extern \"C\" __global__ void __launch_bounds__(512) qwen4_ple_gate", "Task qwen4_ple_gate" }, .{ "extern \"C\" __global__ void __launch_bounds__(512) qwen4_ple_conv", "void qwen4_ple_conv" }, .{ "__shared__ float", "static float" },
         .{ "blockIdx.", "cpu_block." },                .{ "threadIdx.x", "cpu_lane" },                                                                   .{ "blockDim.x", "512" },
     };
diff --git a/zig/kernels/cuda/act_l2.cuh b/zig/kernels/cuda/act_l2.cuh
index 5b347c12f8..50d70f02fd 100644
--- a/zig/kernels/cuda/act_l2.cuh
+++ b/zig/kernels/cuda/act_l2.cuh
@@ -4,6 +4,7 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/vec.cuh"
 
 namespace tf_act_l2 {
 
@@ -21,23 +22,8 @@ __device__ __forceinline__ float point(float v) {
     else return v;
 }
 
-__device__ __forceinline__ void unpack8(const uint4& u, float* f) {
-    const __nv_bfloat162* h = reinterpret_cast<const __nv_bfloat162*>(&u);
-#pragma unroll
-    for (int i = 0; i < 4; ++i) {
-        const float2 t = __bfloat1622float2(h[i]);
-        f[2 * i] = t.x;
-        f[2 * i + 1] = t.y;
-    }
-}
-
-__device__ __forceinline__ uint4 pack8(const float* f) {
-    uint4 u;
-    __nv_bfloat162* h = reinterpret_cast<__nv_bfloat162*>(&u);
-#pragma unroll
-    for (int i = 0; i < 4; ++i) h[i] = __floats2bfloat162_rn(f[2 * i], f[2 * i + 1]);
-    return u;
-}
+using tf_common::unpack8;
+using tf_common::pack8;
 
 // torch's fp32 sigmoid and silu: 1 / (1 + exp(-x)), x / (1 + exp(-x)), IEEE division.
 __device__ __forceinline__ float sigmoid(float x) { return __fdiv_rn(1.f, __fadd_rn(1.f, expf(-x))); }
diff --git a/zig/kernels/cuda/act_round.cuh b/zig/kernels/cuda/act_round.cuh
index 0e00dd2ebb..9137c9e01f 100644
--- a/zig/kernels/cuda/act_round.cuh
+++ b/zig/kernels/cuda/act_round.cuh
@@ -4,6 +4,7 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 
 namespace tf_act_round {
 
@@ -51,36 +52,9 @@ __device__ __forceinline__ float softcap(float x, float cap, int keep) {
     return __fmul_rn(point(tanhf(t), keep, CAP_TANH), cap);
 }
 
-__device__ __forceinline__ void unpack8(const uint4& u, float* f) {
-    const __nv_bfloat162* h = reinterpret_cast<const __nv_bfloat162*>(&u);
-#pragma unroll
-    for (int i = 0; i < 4; ++i) {
-        const float2 t = __bfloat1622float2(h[i]);
-        f[2 * i] = t.x;
-        f[2 * i + 1] = t.y;
-    }
-}
-
-__device__ __forceinline__ uint4 pack8(const float* f) {
-    uint4 u;
-    __nv_bfloat162* h = reinterpret_cast<__nv_bfloat162*>(&u);
-#pragma unroll
-    for (int i = 0; i < 4; ++i) h[i] = __floats2bfloat162_rn(f[2 * i], f[2 * i + 1]);
-    return u;
-}
-
-// act_rows.cu's group_sum8: a thread's 8 stored values in order, then the group's 8 lanes in a fixed tree.
-__device__ __forceinline__ float group_sum8(const uint4& packed, bool live) {
-    float f[8];
-    unpack8(packed, f);
-    float s = 0.f;
-    if (live)
-#pragma unroll
-        for (int i = 0; i < 8; ++i) s = __fadd_rn(s, f[i]);
-    s = __fadd_rn(s, __shfl_xor_sync(0xFFFFFFFFu, s, 4));
-    s = __fadd_rn(s, __shfl_xor_sync(0xFFFFFFFFu, s, 2));
-    return __fadd_rn(s, __shfl_xor_sync(0xFFFFFFFFu, s, 1));
-}
+using tf_common::unpack8;
+using tf_common::pack8;
+using tf_common::group_sum8;
 
 // out[r, i] = bf16(act(gate[r * g_ld + i], up[r * u_ld + i])) for i < n (n % 64 == 0), GEGLU or not; xs its sums.
 template <bool GEGLU>
diff --git a/zig/kernels/cuda/act_rows.cu b/zig/kernels/cuda/act_rows.cu
index 230791ec34..2bb8ce861d 100644
--- a/zig/kernels/cuda/act_rows.cu
+++ b/zig/kernels/cuda/act_rows.cu
@@ -4,45 +4,19 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 
 namespace tf_act {
 
 constexpr int NT = 256;
 
-__device__ __forceinline__ void unpack8(const uint4& u, float* f) {
-    const __nv_bfloat162* h = reinterpret_cast<const __nv_bfloat162*>(&u);
-#pragma unroll
-    for (int i = 0; i < 4; ++i) {
-        const float2 t = __bfloat1622float2(h[i]);
-        f[2 * i] = t.x;
-        f[2 * i + 1] = t.y;
-    }
-}
-
-__device__ __forceinline__ uint4 pack8(const float* f) {
-    uint4 u;
-    __nv_bfloat162* h = reinterpret_cast<__nv_bfloat162*>(&u);
-#pragma unroll
-    for (int i = 0; i < 4; ++i) h[i] = __floats2bfloat162_rn(f[2 * i], f[2 * i + 1]);
-    return u;
-}
+using tf_common::unpack8;
+using tf_common::pack8;
+using tf_common::group_sum8;
 
 // Exactly rounded division on purpose: the same bits on every GPU, where div.full and ex2.approx are not.
 __device__ __forceinline__ float sigmoid(float x) { return __fdiv_rn(1.f, __fadd_rn(1.f, expf(-x))); }
 
-// A thread's 8 stored values in order, then the group's 8 lanes in a fixed tree (every lane of the warp calls).
-__device__ __forceinline__ float group_sum8(const uint4& packed, bool live) {
-    float f[8];
-    unpack8(packed, f);
-    float s = 0.f;
-    if (live)
-#pragma unroll
-        for (int i = 0; i < 8; ++i) s = __fadd_rn(s, f[i]);
-    s = __fadd_rn(s, __shfl_xor_sync(0xFFFFFFFFu, s, 4));
-    s = __fadd_rn(s, __shfl_xor_sync(0xFFFFFFFFu, s, 2));
-    return __fadd_rn(s, __shfl_xor_sync(0xFFFFFFFFu, s, 1));
-}
-
 }  // namespace tf_act
 
 using namespace tf_act;
diff --git a/zig/kernels/cuda/affine8.cu b/zig/kernels/cuda/affine8.cu
index ecd9aac2a8..5ec2c58ade 100644
--- a/zig/kernels/cuda/affine8.cu
+++ b/zig/kernels/cuda/affine8.cu
@@ -5,7 +5,7 @@
 #include <stdint.h>
 
 #include "affine_codes.cuh"
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace tf_affine8 {
 
diff --git a/zig/kernels/cuda/affine_prefill.cu b/zig/kernels/cuda/affine_prefill.cu
index bfa76037cc..2776b72035 100644
--- a/zig/kernels/cuda/affine_prefill.cu
+++ b/zig/kernels/cuda/affine_prefill.cu
@@ -5,7 +5,7 @@
 #include <stdint.h>
 
 #include "affine_codes.cuh"
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace tf_affine8 {
 
diff --git a/zig/kernels/cuda/attn_train.cu b/zig/kernels/cuda/attn_train.cu
index 44a9a62493..15e937541c 100644
--- a/zig/kernels/cuda/attn_train.cu
+++ b/zig/kernels/cuda/attn_train.cu
@@ -2,44 +2,17 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 
 typedef __nv_bfloat16 bf16;
 
 __device__ __forceinline__ float f32(bf16 v) { return __bfloat162float(v); }
 __device__ __forceinline__ bf16 b16(float v) { return __float2bfloat16(v); }
 
-__device__ __forceinline__ float warp_sum(float v) {
-#pragma unroll
-    for (int o = 16; o > 0; o >>= 1) v += __shfl_xor_sync(0xffffffffu, v, o);
-    return v;
-}
-
-__device__ __forceinline__ float warp_max(float v) {
-#pragma unroll
-    for (int o = 16; o > 0; o >>= 1) v = fmaxf(v, __shfl_xor_sync(0xffffffffu, v, o));
-    return v;
-}
-
-// The sum over a block of `groups` warps, the same in every thread.
-__device__ __forceinline__ float block_sum(float v, float* scratch, int groups) {
-    const int lane = threadIdx.x & 31, w = threadIdx.x >> 5;
-    v = warp_sum(v);
-    if (lane == 0) scratch[w] = v;
-    __syncthreads();
-    v = warp_sum(lane < groups ? scratch[lane] : 0.0f);
-    __syncthreads();
-    return v;
-}
-
-__device__ __forceinline__ float block_max(float v, float* scratch, int groups) {
-    const int lane = threadIdx.x & 31, w = threadIdx.x >> 5;
-    v = warp_max(v);
-    if (lane == 0) scratch[w] = v;
-    __syncthreads();
-    v = warp_max(lane < groups ? scratch[lane] : -INFINITY);
-    __syncthreads();
-    return v;
-}
+using tf_common::warp_sum;
+using tf_common::warp_max;
+using tf_common::block_sum_tree;
+using tf_common::block_max_tree;
 
 __device__ __forceinline__ float sigm(float v) { return 1.0f / (1.0f + expf(-v)); }
 
@@ -74,7 +47,7 @@ extern "C" __global__ void __launch_bounds__(256) tf_attn_prep(const bf16* pq, c
     const bool q = hd < heads;
     const float x = q ? f32(pq[(size_t)r * 2 * qw + hd * 2 * dim + d]) : f32(pk[(size_t)r * kw + (hd - heads) * dim + d]);
     const float* w = q ? qscale : kscale;
-    const float rinv = 1.0f / sqrtf(block_sum(x * x, scratch, dim / 32) / (float)dim + eps);
+    const float rinv = 1.0f / sqrtf(block_sum_tree(x * x, scratch, dim / 32) / (float)dim + eps);
     xn[d] = f32(b16(x * rinv * w[d]));
     __syncthreads();
     float out = xn[d];
@@ -117,20 +90,20 @@ extern "C" __global__ void __launch_bounds__(256) tf_attn_q(const bf16* qkv, con
     __syncthreads();
     float top = -INFINITY;
     for (int j = t; j <= at; j += 256) top = fmaxf(top, sc[j]);
-    top = block_max(top, scratch, 8);
+    top = block_max_tree(top, scratch, 8);
     float total = 0;
     for (int j = t; j <= at; j += 256) {
         const float e = expf(sc[j] - top);
         sc[j] = e;
         total += e;
     }
-    total = block_sum(total, scratch, 8);
+    total = block_sum_tree(total, scratch, 8);
     float mix = 0;
     for (int j = t; j <= at; j += 256) {
         sc[j] /= total;
         mix += sc[j] * dp[j];
     }
-    mix = block_sum(mix, scratch, 8);
+    mix = block_sum_tree(mix, scratch, 8);
     const size_t row = ((size_t)h * a.rows + i) * keys;
     for (int j = t; j < keys; j += 256) {
         const bool live = j <= at;
@@ -181,7 +154,7 @@ extern "C" __global__ void __launch_bounds__(256) tf_attn_prep_back(const bf16*
         gn = d < half ? go[d] * c + go[d + half] * s : go[d] * c - go[d - half] * s;
     }
     const float w = (q ? qscale : kscale)[d];
-    const float sq = block_sum(x * x, scratch, dim / 32), dot = block_sum(w * gn * x, scratch, dim / 32);
+    const float sq = block_sum_tree(x * x, scratch, dim / 32), dot = block_sum_tree(w * gn * x, scratch, dim / 32);
     const float rinv = 1.0f / sqrtf(sq / (float)dim + eps);
     const float g = rinv * w * gn - rinv * rinv * rinv * x * dot / (float)dim;
     if (q)
diff --git a/zig/kernels/cuda/block_attention.cu b/zig/kernels/cuda/block_attention.cu
index b8675f80d0..a3a0f70a46 100644
--- a/zig/kernels/cuda/block_attention.cu
+++ b/zig/kernels/cuda/block_attention.cu
@@ -8,13 +8,14 @@
 #include <stdint.h>
 
 #include "attn_odd.cuh"
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace tf_block {
 
 using qmm_frag::commit;
 using qmm_frag::cp16z;
 using qmm_frag::ldmatrix4;
+using qmm_frag::ldmatrix4t;
 using qmm_frag::mma;
 using qmm_frag::mma0;
 
@@ -23,11 +24,6 @@ constexpr float LOG2E = 1.4426950408889634f;
 
 __device__ __forceinline__ void wait_all() { asm volatile("cp.async.wait_group 0;\n" ::); }
 
-__device__ __forceinline__ void ldmatrix4t(uint32_t (&r)[4], const void* p) {
-    asm volatile("ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%0, %1, %2, %3}, [%4];\n"
-                 : "=r"(r[0]), "=r"(r[1]), "=r"(r[2]), "=r"(r[3]) : "r"(qmm_frag::smem(p)));
-}
-
 __device__ __forceinline__ float texp(float x) {
     float y;
     asm("ex2.approx.f32 %0, %1;\n" : "=f"(y) : "f"(__fmul_rn(x, LOG2E)));
diff --git a/zig/kernels/cuda/ck_experts.cu b/zig/kernels/cuda/ck_experts.cu
index e89a5dec2f..363230793d 100644
--- a/zig/kernels/cuda/ck_experts.cu
+++ b/zig/kernels/cuda/ck_experts.cu
@@ -6,7 +6,7 @@
 
 #include "mma4.cuh"
 #include "nvfp4q.cuh"
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace tf_ck4 {
 
diff --git a/zig/kernels/cuda/ck_lane.cu b/zig/kernels/cuda/ck_lane.cu
index ac1a732887..259a100ffd 100644
--- a/zig/kernels/cuda/ck_lane.cu
+++ b/zig/kernels/cuda/ck_lane.cu
@@ -6,7 +6,7 @@
 #include <stdint.h>
 
 #include "mma4.cuh"
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 using namespace qmm_frag;
 using namespace mma4;
diff --git a/zig/kernels/cuda/glue_math.cuh b/zig/kernels/cuda/common/glue_math.cuh
similarity index 91%
rename from zig/kernels/cuda/glue_math.cuh
rename to zig/kernels/cuda/common/glue_math.cuh
index e2fc6e59ed..7f94bf003c 100644
--- a/zig/kernels/cuda/glue_math.cuh
+++ b/zig/kernels/cuda/common/glue_math.cuh
@@ -4,6 +4,8 @@
 #include <cuda_bf16.h>
 #include <stdint.h>
 
+#include "common/reduce.cuh"
+
 namespace tf_glue {
 
 __device__ __forceinline__ float bf(__nv_bfloat16 v) { return __bfloat162float(v); }
@@ -61,11 +63,6 @@ __device__ __forceinline__ float tf_softplus(float v) {
     return __fadd_rn(fmaxf(v, 0.0f), tf_log1p_unit(tf_exp(-fabsf(v))));
 }
 
-// A warp's 32 values summed by xor butterfly (16, 8, 4, 2, 1): every lane ends with the same sum.
-__device__ __forceinline__ float warp_sum(float v) {
-#pragma unroll
-    for (int off = 16; off > 0; off >>= 1) v = __fadd_rn(v, __shfl_xor_sync(0xffffffffu, v, off));
-    return v;
-}
+using tf_common::warp_sum;
 
 }  // namespace tf_glue
diff --git a/zig/kernels/cuda/qmm_frag.cuh b/zig/kernels/cuda/common/qmm_frag.cuh
similarity index 82%
rename from zig/kernels/cuda/qmm_frag.cuh
rename to zig/kernels/cuda/common/qmm_frag.cuh
index be441212df..b02b6df37e 100644
--- a/zig/kernels/cuda/qmm_frag.cuh
+++ b/zig/kernels/cuda/common/qmm_frag.cuh
@@ -1,4 +1,4 @@
-// Device code of src/tensorfold/cuda/kernels/qmm_frag.cuh (lines 1-118, comments dropped), checked by zig/tests/cuda/copies.py.
+// Shared device primitives: cp.async, ldmatrix (x2, x4, x4 .trans), bf16 mma, int4 and e2m1 pair decoding, swizzle.
 #pragma once
 
 #include <cuda_bf16.h>
@@ -43,6 +43,14 @@ __device__ __forceinline__ void cp4(void* dst, const void* src) {
     asm volatile("cp.async.ca.shared.global [%0], [%1], 4;\n" ::"r"(smem(dst)), "l"(src));
 }
 
+__device__ __forceinline__ uint4 ld_nc(const uint4* p) {
+    uint4 r;
+    asm volatile("ld.global.nc.L1::no_allocate.v4.u32 {%0, %1, %2, %3}, [%4];\n"
+                 : "=r"(r.x), "=r"(r.y), "=r"(r.z), "=r"(r.w)
+                 : "l"(p));
+    return r;
+}
+
 __device__ __forceinline__ void commit() { asm volatile("cp.async.commit_group;\n" ::); }
 
 template <int N>
@@ -58,6 +66,11 @@ __device__ __forceinline__ void ldmatrix2(uint32_t (&r)[2], const void* p) {
     asm volatile("ldmatrix.sync.aligned.m8n8.x2.shared.b16 {%0, %1}, [%2];\n" : "=r"(r[0]), "=r"(r[1]) : "r"(smem(p)));
 }
 
+__device__ __forceinline__ void ldmatrix4t(uint32_t (&r)[4], const void* p) {
+    asm volatile("ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%0, %1, %2, %3}, [%4];\n"
+                 : "=r"(r[0]), "=r"(r[1]), "=r"(r[2]), "=r"(r[3]) : "r"(smem(p)));
+}
+
 __device__ __forceinline__ void mma(float (&d)[4], const uint32_t (&a)[4], uint32_t b0, uint32_t b1) {
     asm(
         "mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 {%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9}, "
@@ -85,6 +98,14 @@ __device__ __forceinline__ uint32_t pair(uint32_t w, int s) {
     return r;
 }
 
+__device__ __forceinline__ uint32_t fp4pair(uint32_t w, int s) {
+    const uint32_t v = w >> s;
+    const uint32_t t = ((v & 0x00070007u) << 6) | ((v & 0x00080008u) << 12);
+    uint32_t r;
+    asm("fma.rn.bf16x2 %0, %1, %2, %3;\n" : "=r"(r) : "r"(t), "r"(0x7E807E80u), "r"(0x80008000u));
+    return r;
+}
+
 __device__ __forceinline__ void grid_wait() {
 #if __CUDA_ARCH__ >= 900
     asm volatile("griddepcontrol.wait;\n" ::: "memory");
diff --git a/zig/kernels/cuda/common/reduce.cuh b/zig/kernels/cuda/common/reduce.cuh
new file mode 100644
index 0000000000..3d52f10d4b
--- /dev/null
+++ b/zig/kernels/cuda/common/reduce.cuh
@@ -0,0 +1,58 @@
+// Warp and block reductions, one fixed order each; every lane that calls gets the same total.
+#pragma once
+
+#include <math.h>
+#include <stdint.h>
+
+#include "common/vec.cuh"
+
+namespace tf_common {
+
+// xor butterfly, offsets 16 down to 1
+__device__ __forceinline__ float warp_sum(float v) {
+#pragma unroll
+    for (int o = 16; o > 0; o >>= 1) v += __shfl_xor_sync(0xffffffffu, v, o);
+    return v;
+}
+
+__device__ __forceinline__ float warp_max(float v) {
+#pragma unroll
+    for (int o = 16; o > 0; o >>= 1) v = fmaxf(v, __shfl_xor_sync(0xffffffffu, v, o));
+    return v;
+}
+
+// Tree order: each warp's butterfly, then one more butterfly over the `groups` warp sums (zeros past them).
+__device__ __forceinline__ float block_sum_tree(float v, float* scratch, int groups) {
+    const int lane = threadIdx.x & 31, w = threadIdx.x >> 5;
+    v = warp_sum(v);
+    if (lane == 0) scratch[w] = v;
+    __syncthreads();
+    v = warp_sum(lane < groups ? scratch[lane] : 0.0f);
+    __syncthreads();
+    return v;
+}
+
+__device__ __forceinline__ float block_max_tree(float v, float* scratch, int groups) {
+    const int lane = threadIdx.x & 31, w = threadIdx.x >> 5;
+    v = warp_max(v);
+    if (lane == 0) scratch[w] = v;
+    __syncthreads();
+    v = warp_max(lane < groups ? scratch[lane] : -INFINITY);
+    __syncthreads();
+    return v;
+}
+
+// A thread's 8 stored values in order, then the group's 8 lanes in a fixed tree (every lane of the warp calls).
+__device__ __forceinline__ float group_sum8(const uint4& packed, bool live) {
+    float f[8];
+    unpack8(packed, f);
+    float s = 0.f;
+    if (live)
+#pragma unroll
+        for (int i = 0; i < 8; ++i) s = __fadd_rn(s, f[i]);
+    s = __fadd_rn(s, __shfl_xor_sync(0xFFFFFFFFu, s, 4));
+    s = __fadd_rn(s, __shfl_xor_sync(0xFFFFFFFFu, s, 2));
+    return __fadd_rn(s, __shfl_xor_sync(0xFFFFFFFFu, s, 1));
+}
+
+}  // namespace tf_common
diff --git a/zig/kernels/cuda/common/vec.cuh b/zig/kernels/cuda/common/vec.cuh
new file mode 100644
index 0000000000..dc089ea3ac
--- /dev/null
+++ b/zig/kernels/cuda/common/vec.cuh
@@ -0,0 +1,27 @@
+// Eight bf16 values in one 16-byte word, widened to or rounded from float in index order.
+#pragma once
+
+#include <cuda_bf16.h>
+#include <stdint.h>
+
+namespace tf_common {
+
+__device__ __forceinline__ void unpack8(const uint4& u, float* f) {
+    const __nv_bfloat162* h = reinterpret_cast<const __nv_bfloat162*>(&u);
+#pragma unroll
+    for (int i = 0; i < 4; ++i) {
+        const float2 t = __bfloat1622float2(h[i]);
+        f[2 * i] = t.x;
+        f[2 * i + 1] = t.y;
+    }
+}
+
+__device__ __forceinline__ uint4 pack8(const float* f) {
+    uint4 u;
+    __nv_bfloat162* h = reinterpret_cast<__nv_bfloat162*>(&u);
+#pragma unroll
+    for (int i = 0; i < 4; ++i) h[i] = __floats2bfloat162_rn(f[2 * i], f[2 * i + 1]);
+    return u;
+}
+
+}  // namespace tf_common
diff --git a/zig/kernels/cuda/contract_controls.cu b/zig/kernels/cuda/contract_controls.cu
index ac513c1190..1cc1fd9c84 100644
--- a/zig/kernels/cuda/contract_controls.cu
+++ b/zig/kernels/cuda/contract_controls.cu
@@ -265,7 +265,7 @@ extern "C" __global__ void tf_ctl_q35_dense_half(const __nv_bfloat16 *X, const _
 // The pre-template GLM BF16 operators, kept apart from the candidate math for byte equality.
 #include "contract_glm_bf16_baseline.cuh"
 // Gate 4: flashnext_hc's fn_hc_dense_mix8 with each 32-column group summed in bf16, while its spec declares fp32 sums.
-#include "glue_math.cuh"
+#include "common/glue_math.cuh"
 
 extern "C" __global__ void __launch_bounds__(256) tf_ctl_hc_mix8_bf16(const __nv_bfloat16* UP, const __nv_bfloat16* NORMED,
                                                                       __nv_bfloat16* MIXED, float* XSM, int D, int SS) {
diff --git a/zig/kernels/cuda/drama3_codec.cu b/zig/kernels/cuda/drama3_codec.cu
index b311bbcd63..5c57b8a6af 100644
--- a/zig/kernels/cuda/drama3_codec.cu
+++ b/zig/kernels/cuda/drama3_codec.cu
@@ -2,15 +2,12 @@
 #include <stdint.h>
 #include <math.h>
 #include <cuda_bf16.h>
+#include "common/reduce.cuh"
 
 namespace {
 constexpr unsigned FULL = 0xFFFFFFFFu;
 
-__device__ __forceinline__ float warp_sum(float v) {
-#pragma unroll
-  for (int o = 16; o > 0; o >>= 1) v += __shfl_xor_sync(FULL, v, o);
-  return v;
-}
+using tf_common::warp_sum;
 
 __device__ float block_sum(float v) {
   __shared__ float part[32];
diff --git a/zig/kernels/cuda/drama3_common.cuh b/zig/kernels/cuda/drama3_common.cuh
index 4987d13ebe..7a468a48ab 100644
--- a/zig/kernels/cuda/drama3_common.cuh
+++ b/zig/kernels/cuda/drama3_common.cuh
@@ -4,6 +4,7 @@
 #include <cuda_fp8.h>
 #include <stdint.h>
 #include <math.h>
+#include "common/reduce.cuh"
 
 typedef __nv_bfloat16 bf16;
 
@@ -15,17 +16,8 @@ __device__ __forceinline__ float b2f(bf16 x) { return __bfloat162float(x); }
 __device__ __forceinline__ bf16 f2b(float x) { return __float2bfloat16(x); }
 __device__ __forceinline__ float rb(float x) { return __bfloat162float(__float2bfloat16(x)); }  // round to bf16
 
-__device__ __forceinline__ float warp_sum(float v) {
-#pragma unroll
-  for (int o = 16; o > 0; o >>= 1) v += __shfl_xor_sync(FULL, v, o);
-  return v;
-}
-
-__device__ __forceinline__ float warp_max(float v) {
-#pragma unroll
-  for (int o = 16; o > 0; o >>= 1) v = fmaxf(v, __shfl_xor_sync(FULL, v, o));
-  return v;
-}
+using tf_common::warp_sum;
+using tf_common::warp_max;
 
 // Sum over the block (blockDim.x a multiple of 32, at most 1024); every thread gets the total.
 __device__ float block_sum(float v) {
diff --git a/zig/kernels/cuda/drama3_fp8.cu b/zig/kernels/cuda/drama3_fp8.cu
index 993daf3524..b6a9300d28 100644
--- a/zig/kernels/cuda/drama3_fp8.cu
+++ b/zig/kernels/cuda/drama3_fp8.cu
@@ -4,6 +4,7 @@
 #include <stdint.h>
 
 #include "experts.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace tf_d3_fp8 {
 
@@ -11,13 +12,7 @@ constexpr int LANE4 = 2;
 constexpr int BLOCK4 = 32 * LANE4;
 constexpr int PER_SCALE = 4;
 
-__device__ __forceinline__ uint4 ld_nc(const uint4* p) {
-  uint4 r;
-  asm volatile("ld.global.nc.L1::no_allocate.v4.u32 {%0, %1, %2, %3}, [%4];\n"
-               : "=r"(r.x), "=r"(r.y), "=r"(r.z), "=r"(r.w)
-               : "l"(p));
-  return r;
-}
+using qmm_frag::ld_nc;
 
 __device__ __forceinline__ uint32_t fp8pair(uint32_t v) {
   const uint32_t t = ((v & 0x7Fu) << 4) | ((v & 0x80u) << 8) | ((v & 0x7F00u) << 12) | ((v & 0x8000u) << 16);
diff --git a/zig/kernels/cuda/flashnext_attn.cu b/zig/kernels/cuda/flashnext_attn.cu
index 09aa10cacc..ff4f7aabb7 100644
--- a/zig/kernels/cuda/flashnext_attn.cu
+++ b/zig/kernels/cuda/flashnext_attn.cu
@@ -6,8 +6,8 @@
 #include <math.h>
 #include <stdint.h>
 
-#include "glue_math.cuh"
-#include "qmm_frag.cuh"
+#include "common/glue_math.cuh"
+#include "common/qmm_frag.cuh"
 
 using namespace tf_glue;
 
@@ -42,12 +42,7 @@ __device__ __forceinline__ float rope(float xn, float xpn, int d, int half, int
     return d < half ? __fsub_rn(__fmul_rn(xn, c), __fmul_rn(xpn, s)) : __fadd_rn(__fmul_rn(xpn, s), __fmul_rn(xn, c));
 }
 
-// ldmatrix .x4 .trans: four 8x8 bf16 tiles of shared memory, each thread's pair down a column (the B layout of rows).
-__device__ __forceinline__ void ldmatrix4t(uint32_t (&r)[4], const void* p) {
-    asm volatile("ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%0, %1, %2, %3}, [%4];\n"
-                 : "=r"(r[0]), "=r"(r[1]), "=r"(r[2]), "=r"(r[3])
-                 : "r"(qmm_frag::smem(p)));
-}
+using qmm_frag::ldmatrix4t;
 
 __device__ __forceinline__ int partner(int d, int half) {
     return d < half ? d + half : (d < 2 * half ? d - half : d);
diff --git a/zig/kernels/cuda/flashnext_fact.cu b/zig/kernels/cuda/flashnext_fact.cu
index 50db0f35a2..efe02c1e64 100644
--- a/zig/kernels/cuda/flashnext_fact.cu
+++ b/zig/kernels/cuda/flashnext_fact.cu
@@ -1,5 +1,5 @@
 // Prompt GEMM on #952's factored arithmetic (qmm_kernel's bits): 2 groups a stage, mma per group then fp32 fmas.
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace fact {
 using namespace qmm_frag;
diff --git a/zig/kernels/cuda/flashnext_hc.cu b/zig/kernels/cuda/flashnext_hc.cu
index 9d3de8a145..925506cb7a 100644
--- a/zig/kernels/cuda/flashnext_hc.cu
+++ b/zig/kernels/cuda/flashnext_hc.cu
@@ -5,8 +5,8 @@
 #include <cuda_bf16.h>
 #include <stdint.h>
 
-#include "glue_math.cuh"
-#include "qmm_frag.cuh"
+#include "common/glue_math.cuh"
+#include "common/qmm_frag.cuh"
 
 using namespace tf_glue;
 
diff --git a/zig/kernels/cuda/flashnext_hcf.cu b/zig/kernels/cuda/flashnext_hcf.cu
index 3657cc63f6..cf73138756 100644
--- a/zig/kernels/cuda/flashnext_hcf.cu
+++ b/zig/kernels/cuda/flashnext_hcf.cu
@@ -1,6 +1,6 @@
 // #952's hyper-connection read-out (fn_hc_down / fn_hc_upmix bits) on prompt tiles: factored groups, norm in smem.
-#include "qmm_frag.cuh"
-#include "glue_math.cuh"
+#include "common/qmm_frag.cuh"
+#include "common/glue_math.cuh"
 
 namespace hcf {
 using namespace qmm_frag;
diff --git a/zig/kernels/cuda/flashnext_train.cu b/zig/kernels/cuda/flashnext_train.cu
index 7a5b57f717..1932beed37 100644
--- a/zig/kernels/cuda/flashnext_train.cu
+++ b/zig/kernels/cuda/flashnext_train.cu
@@ -4,28 +4,15 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 
 typedef __nv_bfloat16 bf16;
 
 __device__ __forceinline__ float f32(bf16 v) { return __bfloat162float(v); }
 __device__ __forceinline__ bf16 b16(float v) { return __float2bfloat16(v); }
 
-__device__ __forceinline__ float warp_sum(float v) {
-#pragma unroll
-    for (int o = 16; o > 0; o >>= 1) v += __shfl_xor_sync(0xffffffffu, v, o);
-    return v;
-}
-
-// The sum over a block of `groups` warps, the same in every thread.
-__device__ __forceinline__ float block_sum(float v, float* scratch, int groups) {
-    const int lane = threadIdx.x & 31, w = threadIdx.x >> 5;
-    v = warp_sum(v);
-    if (lane == 0) scratch[w] = v;
-    __syncthreads();
-    v = warp_sum(lane < groups ? scratch[lane] : 0.0f);
-    __syncthreads();
-    return v;
-}
+using tf_common::warp_sum;
+using tf_common::block_sum_tree;
 
 __device__ __forceinline__ float sigm(float v) { return 1.0f / (1.0f + expf(-v)); }
 __device__ __forceinline__ float silu(float v) { return v / (1.0f + expf(-v)); }
@@ -255,8 +242,8 @@ extern "C" __global__ void __launch_bounds__(256) tf_fn_pocket_grad(const bf16*
         a += gate[(size_t)j * d + i] * v;
         b += up[(size_t)j * d + i] * v;
     }
-    a = block_sum(a, scratch, 8);
-    b = block_sum(b, scratch, 8);
+    a = block_sum_tree(a, scratch, 8);
+    b = block_sum_tree(b, scratch, 8);
     const float k = w * wts[(size_t)r * slots + at] * silu(a) * b;
     for (int i = t; i < d; i += 256) grad[(size_t)j * d + i] += k * push[(size_t)r * d + i];
 }
@@ -282,7 +269,7 @@ extern "C" __global__ void __launch_bounds__(256) tf_fn_hc_branch_back(const flo
         dbranch[(size_t)r * d + i] = v;
     }
     for (int s = 0; s < streams; s++) {
-        const float v = block_sum(part[s], scratch, 8);
+        const float v = block_sum_tree(part[s], scratch, 8);
         if (t == 0) dinject[r * streams + s] = v;
     }
 }
@@ -332,8 +319,8 @@ extern "C" __global__ void __launch_bounds__(256) tf_fn_hc_norm_back(const bf16*
         sq += v * v;
         dot += w[i] * dnormed[base + i] * v;
     }
-    sq = block_sum(sq, scratch, 8);
-    dot = block_sum(dot, scratch, 8);
+    sq = block_sum_tree(sq, scratch, 8);
+    dot = block_sum_tree(dot, scratch, 8);
     const float inv = rsqrtf(sq / (float)d + eps), k = inv * inv * inv * dot / (float)d;
     for (int i = t; i < d; i += 256) g[base + i] += inv * w[i] * dnormed[base + i] - k * f32(h[base + i]);
 }
diff --git a/zig/kernels/cuda/gdn.cu b/zig/kernels/cuda/gdn.cu
index 2c96e1a3ad..f8c8a031ae 100644
--- a/zig/kernels/cuda/gdn.cu
+++ b/zig/kernels/cuda/gdn.cu
@@ -4,16 +4,13 @@
 #include <cuda_runtime.h>
 
 #include <cstdint>
+#include "common/reduce.cuh"
 
 namespace tf_gdn {
 
 constexpr int DK = 128;
 
-__device__ __forceinline__ float warp_sum(float x) {
-#pragma unroll
-    for (int m = 16; m; m >>= 1) x += __shfl_xor_sync(0xffffffffu, x, m);
-    return x;
-}
+using tf_common::warp_sum;
 
 template <int N>
 __device__ __forceinline__ float spread_sum(float (&v)[N], int lane) {
diff --git a/zig/kernels/cuda/gdn_train.cu b/zig/kernels/cuda/gdn_train.cu
index 9b908fc850..2119697d7e 100644
--- a/zig/kernels/cuda/gdn_train.cu
+++ b/zig/kernels/cuda/gdn_train.cu
@@ -2,6 +2,7 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 
 typedef __nv_bfloat16 bf16;
 
@@ -21,21 +22,8 @@ __device__ __forceinline__ float f32(bf16 v) { return __bfloat162float(v); }
 __device__ __forceinline__ bf16 b16(float v) { return __float2bfloat16_rn(v); }
 __device__ __forceinline__ float bf(float x) { return f32(b16(x)); }
 
-__device__ __forceinline__ float warp_sum(float x) {
-    for (int o = 16; o; o >>= 1) x += __shfl_xor_sync(0xffffffffu, x, o);
-    return x;
-}
-
-// The sum over a block of `groups` warps, the same in every thread.
-__device__ __forceinline__ float block_sum(float v, float* scratch, int groups) {
-    const int lane = threadIdx.x & 31, w = threadIdx.x >> 5;
-    v = warp_sum(v);
-    if (lane == 0) scratch[w] = v;
-    __syncthreads();
-    v = warp_sum(lane < groups ? scratch[lane] : 0.0f);
-    __syncthreads();
-    return v;
-}
+using tf_common::warp_sum;
+using tf_common::block_sum_tree;
 
 __device__ __forceinline__ float sigmoidf_(float x) { return 1.0f / (1.0f + expf(-x)); }
 __device__ __forceinline__ float softplusf_(float x) { return x > 20.0f ? x : log1pf(expf(x)); }
@@ -200,7 +188,7 @@ extern "C" __global__ void __launch_bounds__(128) tf_gdn_norm_back(const bf16* _
     const float sg = sigmoidf_(f32(P[zc])), d = dout[row + t];
     dproj[zc] = d * yn * sg * (1.0f - sg);
     const float dn = d * sg * w;
-    const float dot = block_sum(yv * dn, scratch, 4);
+    const float dot = block_sum_tree(yv * dn, scratch, 4);
     dy[row + t] = ri * dn - ri * ri * ri * yv * dot / (float)DV;
 }
 
@@ -371,8 +359,8 @@ extern "C" __global__ void __launch_bounds__(128) tf_gdn_qk_back(const bf16* __r
             for (int p = 0; p < SPLIT; ++p) d += part[(((size_t)r * s.nv + h) * SPLIT + p) * DK + t];
         const int c = which * s.nk * DK + hk * DK + t;
         const float x = f32(act[(size_t)r * s.conv + c]);
-        const float ss = block_sum(x * x, scratch, 4);
-        const float dot = block_sum(x * d, scratch, 4);
+        const float ss = block_sum_tree(x * x, scratch, 4);
+        const float dot = block_sum_tree(x * d, scratch, 4);
         const float inv = 1.0f / sqrtf(ss + 1e-6f), sc = which == 0 ? 1.0f / sqrtf((float)DK) : 1.0f;
         dact[(size_t)r * s.conv + c] = sc * inv * (d - inv * inv * x * dot);
     }
diff --git a/zig/kernels/cuda/glm_kda.cu b/zig/kernels/cuda/glm_kda.cu
index 4d4eb8349a..542870dca0 100644
--- a/zig/kernels/cuda/glm_kda.cu
+++ b/zig/kernels/cuda/glm_kda.cu
@@ -3,16 +3,14 @@
 #include <cuda_bf16.h>
 #include <cuda_runtime.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 namespace {
 
 constexpr int DK = 128, DV = 128, TAPS = 4;
 
 __device__ __forceinline__ float bf(float x) { return __bfloat162float(__float2bfloat16_rn(x)); }
 
-__device__ __forceinline__ float warp_sum(float x) {
-    for (int o = 16; o; o >>= 1) x += __shfl_xor_sync(0xffffffffu, x, o);
-    return x;
-}
+using tf_common::warp_sum;
 
 __device__ __forceinline__ float sigmoidf_(float x) { return 1.0f / (1.0f + expf(-x)); }
 
diff --git a/zig/kernels/cuda/h3.cu b/zig/kernels/cuda/h3.cu
index 1fbfa862f2..c64301f60d 100644
--- a/zig/kernels/cuda/h3.cu
+++ b/zig/kernels/cuda/h3.cu
@@ -1,5 +1,6 @@
 // MiniMax H3 / FastH3's transformer blocks on CUDA in the checkpoint's bf16: the row, routing and attention kernels.
 #include "h3_attention.cuh"
+#include "common/reduce.cuh"
 
 namespace tf_h3 {
 
@@ -17,16 +18,9 @@ __device__ __forceinline__ float block_sum(float v, float* scratch) {
     __syncthreads();
     return t;
 }
-__device__ __forceinline__ float warp_max(float v) {
-#pragma unroll
-    for (int o = 16; o > 0; o >>= 1) v = fmaxf(v, __shfl_xor_sync(0xffffffffu, v, o));
-    return v;
-}
-__device__ __forceinline__ float warp_sum(float v) {
-#pragma unroll
-    for (int o = 16; o > 0; o >>= 1) v += __shfl_xor_sync(0xffffffffu, v, o);
-    return v;
-}
+
+using tf_common::warp_max;
+using tf_common::warp_sum;
 
 // The modulated norm of one value, rounded to bf16 where torch's bf16 ops round: norm, 1 + scale, product, shift.
 __device__ __forceinline__ float modulated(float x, float inv, float w, float scale, float shift) {
diff --git a/zig/kernels/cuda/lane_gemv.cu b/zig/kernels/cuda/lane_gemv.cu
index 92db256443..802dc048f3 100644
--- a/zig/kernels/cuda/lane_gemv.cu
+++ b/zig/kernels/cuda/lane_gemv.cu
@@ -4,7 +4,7 @@
 #include <cuda_runtime.h>
 #include <stdint.h>
 
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace tf_lane_gemv {
 
diff --git a/zig/kernels/cuda/lfm2.cu b/zig/kernels/cuda/lfm2.cu
index f1d193c659..f7d418bc20 100644
--- a/zig/kernels/cuda/lfm2.cu
+++ b/zig/kernels/cuda/lfm2.cu
@@ -4,7 +4,7 @@
 #include <math.h>
 #include <stdint.h>
 
-#include "glue_math.cuh"
+#include "common/glue_math.cuh"
 
 using namespace tf_glue;
 
diff --git a/zig/kernels/cuda/nemotron_attention.cu b/zig/kernels/cuda/nemotron_attention.cu
index 24a3d5898c..313b096f28 100644
--- a/zig/kernels/cuda/nemotron_attention.cu
+++ b/zig/kernels/cuda/nemotron_attention.cu
@@ -4,7 +4,7 @@
 #include <math.h>
 #include <stdint.h>
 
-#include "glue_math.cuh"
+#include "common/glue_math.cuh"
 
 using namespace tf_glue;
 
diff --git a/zig/kernels/cuda/nemotron_latent.cu b/zig/kernels/cuda/nemotron_latent.cu
index c8e72884a6..594ca8fe20 100644
--- a/zig/kernels/cuda/nemotron_latent.cu
+++ b/zig/kernels/cuda/nemotron_latent.cu
@@ -4,7 +4,7 @@
 #include <math.h>
 #include <stdint.h>
 
-#include "glue_math.cuh"
+#include "common/glue_math.cuh"
 
 using namespace tf_glue;
 
diff --git a/zig/kernels/cuda/nemotron_mamba.cu b/zig/kernels/cuda/nemotron_mamba.cu
index 164b393d3e..a746d094a4 100644
--- a/zig/kernels/cuda/nemotron_mamba.cu
+++ b/zig/kernels/cuda/nemotron_mamba.cu
@@ -3,7 +3,7 @@
 #include <cuda_bf16.h>
 #include <stdint.h>
 
-#include "glue_math.cuh"
+#include "common/glue_math.cuh"
 
 using namespace tf_glue;
 
diff --git a/zig/kernels/cuda/nemotron_norms.cu b/zig/kernels/cuda/nemotron_norms.cu
index b3ed7e2e19..9188a0b504 100644
--- a/zig/kernels/cuda/nemotron_norms.cu
+++ b/zig/kernels/cuda/nemotron_norms.cu
@@ -3,7 +3,7 @@
 #include <cuda_bf16.h>
 #include <stdint.h>
 
-#include "glue_math.cuh"
+#include "common/glue_math.cuh"
 
 namespace tf_glue {
 
diff --git a/zig/kernels/cuda/nemotron_route.cu b/zig/kernels/cuda/nemotron_route.cu
index 03be60a19f..9d0e45d097 100644
--- a/zig/kernels/cuda/nemotron_route.cu
+++ b/zig/kernels/cuda/nemotron_route.cu
@@ -4,7 +4,7 @@
 #include <math.h>
 #include <stdint.h>
 
-#include "glue_math.cuh"
+#include "common/glue_math.cuh"
 
 using namespace tf_glue;
 
diff --git a/zig/kernels/cuda/norms_l2.cuh b/zig/kernels/cuda/norms_l2.cuh
index 8296101195..f05cde9fe1 100644
--- a/zig/kernels/cuda/norms_l2.cuh
+++ b/zig/kernels/cuda/norms_l2.cuh
@@ -4,6 +4,7 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/vec.cuh"
 
 namespace tf_l2 {
 
@@ -22,23 +23,8 @@ __device__ __forceinline__ float point(float v) {
     else return v;
 }
 
-__device__ __forceinline__ void unpack8(const uint4& u, float* f) {
-    const __nv_bfloat162* h = reinterpret_cast<const __nv_bfloat162*>(&u);
-#pragma unroll
-    for (int i = 0; i < 4; ++i) {
-        const float2 t = __bfloat1622float2(h[i]);
-        f[2 * i] = t.x;
-        f[2 * i + 1] = t.y;
-    }
-}
-
-__device__ __forceinline__ uint4 pack8(const float* f) {
-    uint4 u;
-    __nv_bfloat162* h = reinterpret_cast<__nv_bfloat162*>(&u);
-#pragma unroll
-    for (int i = 0; i < 4; ++i) h[i] = __floats2bfloat162_rn(f[2 * i], f[2 * i + 1]);
-    return u;
-}
+using tf_common::unpack8;
+using tf_common::pack8;
 
 // 1 / sqrt(sum + eps) of the row's squares: a thread's fma chain, the warp's xor tree, the 8 warps' tree (IEEE).
 __device__ __forceinline__ float inv_l2(const float (&v)[MOST][8], float eps) {
diff --git a/zig/kernels/cuda/norms_tail.cuh b/zig/kernels/cuda/norms_tail.cuh
index e27c642d5e..707f19d2fb 100644
--- a/zig/kernels/cuda/norms_tail.cuh
+++ b/zig/kernels/cuda/norms_tail.cuh
@@ -4,6 +4,7 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 
 namespace tf_rms_tail {
 
@@ -33,23 +34,9 @@ __device__ __forceinline__ float normed(float v, float inv, float w) {
     else return __fmul_rn(__fmul_rn(v, inv), w);
 }
 
-__device__ __forceinline__ void unpack8(const uint4& u, float* f) {
-    const __nv_bfloat162* h = reinterpret_cast<const __nv_bfloat162*>(&u);
-#pragma unroll
-    for (int i = 0; i < 4; ++i) {
-        const float2 t = __bfloat1622float2(h[i]);
-        f[2 * i] = t.x;
-        f[2 * i + 1] = t.y;
-    }
-}
-
-__device__ __forceinline__ uint4 pack8(const float* f) {
-    uint4 u;
-    __nv_bfloat162* h = reinterpret_cast<__nv_bfloat162*>(&u);
-#pragma unroll
-    for (int i = 0; i < 4; ++i) h[i] = __floats2bfloat162_rn(f[2 * i], f[2 * i + 1]);
-    return u;
-}
+using tf_common::unpack8;
+using tf_common::pack8;
+using tf_common::group_sum8;
 
 // 1 / sqrt(sum / width + eps), IEEE division and square root: norms.cu's inv_rms_v tree over the 8 warps.
 template <bool ONCE_ONLY>
@@ -66,19 +53,6 @@ __device__ __forceinline__ float inv_rms(float part, int width, float eps) {
     return r;
 }
 
-// A thread's 8 stored outputs summed in order, then the group's 8 lanes in a fixed tree (every lane of the warp calls).
-__device__ __forceinline__ float group_sum8(const uint4& packed, bool live) {
-    float f[8];
-    unpack8(packed, f);
-    float s = 0.f;
-    if (live)
-#pragma unroll
-        for (int i = 0; i < 8; ++i) s = __fadd_rn(s, f[i]);
-    s = __fadd_rn(s, __shfl_xor_sync(0xFFFFFFFFu, s, 4));
-    s = __fadd_rn(s, __shfl_xor_sync(0xFFFFFFFFu, s, 2));
-    return __fadd_rn(s, __shfl_xor_sync(0xFFFFFFFFu, s, 1));
-}
-
 // b = wb ? bf16(rms(x) wb) : x, h = bf16(sc ? (r + b) sc : r + b) (if r, wb or sc), y = bf16(rms(h) w), xs its sums.
 template <int R>
 __device__ __forceinline__ void tail(const __nv_bfloat16* __restrict__ x, const __nv_bfloat16* __restrict__ r,
diff --git a/zig/kernels/cuda/prefill_attention.cu b/zig/kernels/cuda/prefill_attention.cu
index a1853aa326..582e6dc041 100644
--- a/zig/kernels/cuda/prefill_attention.cu
+++ b/zig/kernels/cuda/prefill_attention.cu
@@ -5,7 +5,7 @@
 #include <cuda_runtime.h>
 #include <stdint.h>
 
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace tf_prefill_attention {
 
diff --git a/zig/kernels/cuda/prompt16.cu b/zig/kernels/cuda/prompt16.cu
index 0847158716..2c7f0d4ce0 100644
--- a/zig/kernels/cuda/prompt16.cu
+++ b/zig/kernels/cuda/prompt16.cu
@@ -4,7 +4,7 @@
 #include <cuda_bf16.h>
 #include <cuda_runtime.h>
 
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace tf_prompt16 {
 
diff --git a/zig/kernels/cuda/qmm.cu b/zig/kernels/cuda/qmm.cu
index c41c99b154..64066f79c0 100644
--- a/zig/kernels/cuda/qmm.cu
+++ b/zig/kernels/cuda/qmm.cu
@@ -5,7 +5,7 @@
 #include <cuda_bf16.h>
 #include <cuda_runtime.h>
 
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace tf_qmm {
 
diff --git a/zig/kernels/cuda/qmm_group.cu b/zig/kernels/cuda/qmm_group.cu
index 51a14e5d0d..79f432b50c 100644
--- a/zig/kernels/cuda/qmm_group.cu
+++ b/zig/kernels/cuda/qmm_group.cu
@@ -7,7 +7,7 @@
 #include <type_traits>
 #include <vector>
 
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace tf_qmm_group {
 
diff --git a/zig/kernels/cuda/qmm_prefill.cu b/zig/kernels/cuda/qmm_prefill.cu
index 9af0ad2e41..779ec58a37 100644
--- a/zig/kernels/cuda/qmm_prefill.cu
+++ b/zig/kernels/cuda/qmm_prefill.cu
@@ -4,7 +4,7 @@
 #include <cuda_bf16.h>
 #include <cuda_runtime.h>
 
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace tf_qmm_prefill {
 
diff --git a/zig/kernels/cuda/qmmf.cu b/zig/kernels/cuda/qmmf.cu
index 46291eda63..c18084c111 100644
--- a/zig/kernels/cuda/qmmf.cu
+++ b/zig/kernels/cuda/qmmf.cu
@@ -6,7 +6,7 @@
 #include <cuda_fp8.h>
 #include <cuda_runtime.h>
 
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace tf_qmmf {
 
@@ -16,14 +16,6 @@ enum Mode : int { FP4 = 0, FP8 = 1, MXFP8 = 2, FP8G = 3, BF16 = 4 };
 
 constexpr int GS = 64;
 
-__device__ __forceinline__ uint32_t fp4pair(uint32_t w, int s) {
-    const uint32_t v = w >> s;
-    const uint32_t t = ((v & 0x00070007u) << 6) | ((v & 0x00080008u) << 12);
-    uint32_t r;
-    asm("fma.rn.bf16x2 %0, %1, %2, %3;\n" : "=r"(r) : "r"(t), "r"(0x7E807E80u), "r"(0x80008000u));
-    return r;
-}
-
 __device__ __forceinline__ uint32_t fp8pair(uint32_t w) {
     const uint32_t x = (w & 0xFFu) | ((w & 0xFF00u) << 8);
     const uint32_t t = ((x & 0x007F007Fu) << 4) | ((x & 0x00800080u) << 8);
diff --git a/zig/kernels/cuda/qwen27_train.cu b/zig/kernels/cuda/qwen27_train.cu
index f81b441bb7..556416d433 100644
--- a/zig/kernels/cuda/qwen27_train.cu
+++ b/zig/kernels/cuda/qwen27_train.cu
@@ -2,6 +2,7 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 
 typedef __nv_bfloat16 bf16;
 
@@ -10,22 +11,8 @@ __device__ __forceinline__ bf16 b16(float v) { return __float2bfloat16(v); }
 __device__ __forceinline__ float wide(float v) { return v; }
 __device__ __forceinline__ float wide(bf16 v) { return __bfloat162float(v); }
 
-__device__ __forceinline__ float warp_sum(float v) {
-#pragma unroll
-    for (int o = 16; o > 0; o >>= 1) v += __shfl_xor_sync(0xffffffffu, v, o);
-    return v;
-}
-
-// The sum over a block of `groups` warps, the same in every thread.
-__device__ __forceinline__ float block_sum(float v, float* scratch, int groups) {
-    const int lane = threadIdx.x & 31, w = threadIdx.x >> 5;
-    v = warp_sum(v);
-    if (lane == 0) scratch[w] = v;
-    __syncthreads();
-    v = warp_sum(lane < groups ? scratch[lane] : 0.0f);
-    __syncthreads();
-    return v;
-}
+using tf_common::warp_sum;
+using tf_common::block_sum_tree;
 
 __device__ __forceinline__ float sigm(float v) { return 1.0f / (1.0f + expf(-v)); }
 __device__ __forceinline__ float silu(float v) { return v / (1.0f + expf(-v)); }
@@ -390,8 +377,8 @@ extern "C" __global__ void __launch_bounds__(256) tf_q27_pocket_grad(const bf16*
         a += gate[(size_t)j * d + i] * v;
         b += up[(size_t)j * d + i] * v;
     }
-    a = block_sum(a, scratch, 8);
-    b = block_sum(b, scratch, 8);
+    a = block_sum_tree(a, scratch, 8);
+    b = block_sum_tree(b, scratch, 8);
     const float k = w * silu(a) * b;
     for (int i = t; i < d; i += 256) grad[(size_t)j * d + i] += k * push[(size_t)r * d + i];
 }
diff --git a/zig/kernels/cuda/qwen35.cuh b/zig/kernels/cuda/qwen35.cuh
index 927c6129a9..6d4656f656 100644
--- a/zig/kernels/cuda/qwen35.cuh
+++ b/zig/kernels/cuda/qwen35.cuh
@@ -3,6 +3,7 @@
 #pragma once
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 
 namespace tf_q35 {
 
@@ -84,11 +85,7 @@ __device__ __forceinline__ int back(const int* __restrict__ rowpar, int w, int s
 __device__ __forceinline__ float bf(__nv_bfloat16 x) { return __bfloat162float(x); }
 __device__ __forceinline__ float rbf(float x) { return __bfloat162float(__float2bfloat16_rn(x)); }
 
-__device__ __forceinline__ float warp_sum(float x) {
-#pragma unroll
-    for (int m = 16; m; m >>= 1) x += __shfl_xor_sync(0xffffffffu, x, m);
-    return x;
-}
+using tf_common::warp_sum;
 
 // P (device counts): the row past the forward's rows, as qwen35.cu's other counts read it.
 __device__ __forceinline__ bool past_rows(const int* P, int m) { return P != nullptr && m >= P[2]; }
diff --git a/zig/kernels/cuda/qwen35_attn_tc.cuh b/zig/kernels/cuda/qwen35_attn_tc.cuh
index 4bd665ce76..7a0ea01a29 100644
--- a/zig/kernels/cuda/qwen35_attn_tc.cuh
+++ b/zig/kernels/cuda/qwen35_attn_tc.cuh
@@ -3,7 +3,7 @@
 #pragma once
 #include <cuda_fp8.h>
 
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 #include "qwen35.cuh"
 #include "qwen35_page_run.h"
 
diff --git a/zig/kernels/cuda/qwen35_draft.cu b/zig/kernels/cuda/qwen35_draft.cu
index dbc1436bd7..f1344734e3 100644
--- a/zig/kernels/cuda/qwen35_draft.cu
+++ b/zig/kernels/cuda/qwen35_draft.cu
@@ -2,6 +2,7 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 
 namespace tf_q35d {
 
@@ -10,11 +11,7 @@ constexpr int HD = 128;  // the drafter's head width (q, k, v)
 __device__ __forceinline__ float bf(__nv_bfloat16 x) { return __bfloat162float(x); }
 __device__ __forceinline__ float rbf(float x) { return __bfloat162float(__float2bfloat16_rn(x)); }
 
-__device__ __forceinline__ float warp_sum(float x) {
-#pragma unroll
-    for (int m = 16; m; m >>= 1) x += __shfl_xor_sync(0xffffffffu, x, m);
-    return x;
-}
+using tf_common::warp_sum;
 
 }  // namespace tf_q35d
 
diff --git a/zig/kernels/cuda/qwen36_mtp.cu b/zig/kernels/cuda/qwen36_mtp.cu
index 5118a8aee7..196af0d088 100644
--- a/zig/kernels/cuda/qwen36_mtp.cu
+++ b/zig/kernels/cuda/qwen36_mtp.cu
@@ -2,6 +2,7 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 
 namespace tf_q36m {
 
@@ -11,11 +12,7 @@ constexpr int ROT = 64;  // rotary dims: pairs (i, i + 32), the target's partial
 __device__ __forceinline__ float bf(__nv_bfloat16 x) { return __bfloat162float(x); }
 __device__ __forceinline__ float rbf(float x) { return __bfloat162float(__float2bfloat16_rn(x)); }
 
-__device__ __forceinline__ float warp_sum(float x) {
-#pragma unroll
-    for (int m = 16; m; m >>= 1) x += __shfl_xor_sync(0xffffffffu, x, m);
-    return x;
-}
+using tf_common::warp_sum;
 
 }  // namespace tf_q36m
 
diff --git a/zig/kernels/cuda/qwen4_hc.cu b/zig/kernels/cuda/qwen4_hc.cu
index 4d8a594e9e..2738d820a6 100644
--- a/zig/kernels/cuda/qwen4_hc.cu
+++ b/zig/kernels/cuda/qwen4_hc.cu
@@ -1,5 +1,5 @@
 // Shared hyper-connection activation only; the original Flash Next image remains unchanged.
-#include "glue_math.cuh"
+#include "common/glue_math.cuh"
 #include <cuda_fp16.h>
 #define TF_INLINE __device__ __forceinline__
 #define TF_DATA
diff --git a/zig/kernels/cuda/qwen4_ple.cu b/zig/kernels/cuda/qwen4_ple.cu
index 0c0372dcb5..712ccfad51 100644
--- a/zig/kernels/cuda/qwen4_ple.cu
+++ b/zig/kernels/cuda/qwen4_ple.cu
@@ -1,5 +1,5 @@
 #include <cuda_bf16.h>
-#include "glue_math.cuh"
+#include "common/glue_math.cuh"
 using namespace tf_glue;
 __device__ __forceinline__ float bsig(float x) { const float e = tf_exp(-fabsf(x)); return rbf(__fdiv_rn(x < 0 ? e : 1.0f, __fadd_rn(1.0f, e))); }
 __device__ __forceinline__ float bsilu(float x) { const float e = tf_exp(-fabsf(x)); return rbf(__fdiv_rn(x < 0 ? __fmul_rn(x, e) : x, __fadd_rn(1.0f, e))); }
diff --git a/zig/kernels/cuda/route_rows.cu b/zig/kernels/cuda/route_rows.cu
index 5489c57868..94779d3985 100644
--- a/zig/kernels/cuda/route_rows.cu
+++ b/zig/kernels/cuda/route_rows.cu
@@ -6,8 +6,8 @@
 #include <math.h>
 #include <stdint.h>
 
-#include "glue_math.cuh"
-#include "qmm_frag.cuh"
+#include "common/glue_math.cuh"
+#include "common/qmm_frag.cuh"
 
 using namespace tf_glue;
 
diff --git a/zig/kernels/cuda/sample.cu b/zig/kernels/cuda/sample.cu
index d27061e065..64192ad2f1 100644
--- a/zig/kernels/cuda/sample.cu
+++ b/zig/kernels/cuda/sample.cu
@@ -5,6 +5,7 @@
 #include <stdint.h>
 
 #include "logprob_rows.cuh"
+#include "common/reduce.cuh"
 
 namespace tf_sample {
 
@@ -51,18 +52,8 @@ __device__ __forceinline__ float race_uniform(unsigned long long seed, unsigned
     return fminf(uniform(seed, pos, id), __uint_as_float(0x3F7FFFFFu));
 }
 
-// Metal's simd_sum as the host reference has it: a 32-lane butterfly, every lane ending with the sum.
-__device__ __forceinline__ float warp_sum(float x) {
-#pragma unroll
-    for (int off = 16; off > 0; off >>= 1) x += __shfl_xor_sync(FULL, x, off);
-    return x;
-}
-
-__device__ __forceinline__ float warp_max(float x) {
-#pragma unroll
-    for (int off = 16; off > 0; off >>= 1) x = fmaxf(x, __shfl_xor_sync(FULL, x, off));
-    return x;
-}
+using tf_common::warp_sum;
+using tf_common::warp_max;
 
 __device__ __forceinline__ unsigned warp_sum_u(unsigned x) {
 #pragma unroll
diff --git a/zig/kernels/cuda/train.cu b/zig/kernels/cuda/train.cu
index 90cff26ede..1be9a10e6d 100644
--- a/zig/kernels/cuda/train.cu
+++ b/zig/kernels/cuda/train.cu
@@ -2,44 +2,17 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 
 typedef __nv_bfloat16 bf16;
 
 __device__ __forceinline__ float f32(bf16 v) { return __bfloat162float(v); }
 __device__ __forceinline__ bf16 b16(float v) { return __float2bfloat16(v); }
 
-__device__ __forceinline__ float warp_sum(float v) {
-#pragma unroll
-    for (int o = 16; o > 0; o >>= 1) v += __shfl_xor_sync(0xffffffffu, v, o);
-    return v;
-}
-
-__device__ __forceinline__ float warp_max(float v) {
-#pragma unroll
-    for (int o = 16; o > 0; o >>= 1) v = fmaxf(v, __shfl_xor_sync(0xffffffffu, v, o));
-    return v;
-}
-
-// The sum over a block of `groups` warps, the same in every thread.
-__device__ __forceinline__ float block_sum(float v, float* scratch, int groups) {
-    const int lane = threadIdx.x & 31, w = threadIdx.x >> 5;
-    v = warp_sum(v);
-    if (lane == 0) scratch[w] = v;
-    __syncthreads();
-    v = warp_sum(lane < groups ? scratch[lane] : 0.0f);
-    __syncthreads();
-    return v;
-}
-
-__device__ __forceinline__ float block_max(float v, float* scratch, int groups) {
-    const int lane = threadIdx.x & 31, w = threadIdx.x >> 5;
-    v = warp_max(v);
-    if (lane == 0) scratch[w] = v;
-    __syncthreads();
-    v = warp_max(lane < groups ? scratch[lane] : -INFINITY);
-    __syncthreads();
-    return v;
-}
+using tf_common::warp_sum;
+using tf_common::warp_max;
+using tf_common::block_sum_tree;
+using tf_common::block_max_tree;
 
 __device__ __forceinline__ float silu(float v) { return v / (1.0f + expf(-v)); }
 
@@ -60,10 +33,10 @@ extern "C" __global__ void __launch_bounds__(1024) tf_train_softmax(bf16* logits
     bf16* l = logits + (size_t)r * vocab;
     float m = -INFINITY;
     for (int v = t; v < vocab; v += 1024) m = fmaxf(m, f32(l[v]));
-    m = block_max(m, scratch, 32);
+    m = block_max_tree(m, scratch, 32);
     float s = 0;
     for (int v = t; v < vocab; v += 1024) s += expf(f32(l[v]) - m);
-    s = block_sum(s, scratch, 32);
+    s = block_sum_tree(s, scratch, 32);
     const unsigned target = targets[r];
     const float lt = f32(l[target]);
     if (t == 0) stats[r] = make_float2(logf(s) + m - lt, expf(lt - m) / s);
@@ -85,8 +58,8 @@ extern "C" __global__ void __launch_bounds__(256) tf_train_rms_back(const bf16*
         sq += v * v;
         dot += f32(w[j]) * dr[j] * v;
     }
-    sq = block_sum(sq, scratch, 8);
-    dot = block_sum(dot, scratch, 8);
+    sq = block_sum_tree(sq, scratch, 8);
+    dot = block_sum_tree(dot, scratch, 8);
     const float s = rsqrtf(sq / (float)dim + eps), k = s * s * s * dot / (float)dim;
     for (int j = t; j < dim; j += 256) g[(size_t)r * dim + j] += s * f32(w[j]) * dr[j] - k * f32(hr[j]);
 }
@@ -478,7 +451,7 @@ extern "C" __global__ void __launch_bounds__(256) tf_train_route_back(const floa
     for (int j = 0; j < top_k; j++) {
         float s = 0;
         for (int i = t; i < dim; i += 256) s += g[(size_t)r * dim + i] * f32(ys[((size_t)r * slots + j) * dim + i]);
-        a[j] = block_sum(s, scratch, 8);
+        a[j] = block_sum_tree(s, scratch, 8);
     }
     if (t == 0) {
         float p[16];
diff --git a/zig/kernels/cuda/train_attn.cu b/zig/kernels/cuda/train_attn.cu
index 29ab6ac54b..31e27a0772 100644
--- a/zig/kernels/cuda/train_attn.cu
+++ b/zig/kernels/cuda/train_attn.cu
@@ -4,44 +4,17 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 
 typedef __nv_bfloat16 bf16;
 
 __device__ __forceinline__ float f32(bf16 v) { return __bfloat162float(v); }
 __device__ __forceinline__ bf16 b16(float v) { return __float2bfloat16(v); }
 
-__device__ __forceinline__ float warp_sum(float v) {
-#pragma unroll
-    for (int o = 16; o > 0; o >>= 1) v += __shfl_xor_sync(0xffffffffu, v, o);
-    return v;
-}
-
-__device__ __forceinline__ float warp_max(float v) {
-#pragma unroll
-    for (int o = 16; o > 0; o >>= 1) v = fmaxf(v, __shfl_xor_sync(0xffffffffu, v, o));
-    return v;
-}
-
-// The sum over a block of `groups` warps, the same in every thread.
-__device__ __forceinline__ float block_sum(float v, float* scratch, int groups) {
-    const int lane = threadIdx.x & 31, w = threadIdx.x >> 5;
-    v = warp_sum(v);
-    if (lane == 0) scratch[w] = v;
-    __syncthreads();
-    v = warp_sum(lane < groups ? scratch[lane] : 0.0f);
-    __syncthreads();
-    return v;
-}
-
-__device__ __forceinline__ float block_max(float v, float* scratch, int groups) {
-    const int lane = threadIdx.x & 31, w = threadIdx.x >> 5;
-    v = warp_max(v);
-    if (lane == 0) scratch[w] = v;
-    __syncthreads();
-    v = warp_max(lane < groups ? scratch[lane] : -INFINITY);
-    __syncthreads();
-    return v;
-}
+using tf_common::warp_sum;
+using tf_common::warp_max;
+using tf_common::block_sum_tree;
+using tf_common::block_max_tree;
 
 __device__ __forceinline__ float sigm(float v) { return 1.0f / (1.0f + expf(-v)); }
 
@@ -74,7 +47,7 @@ extern "C" __global__ void __launch_bounds__(256) tf_gattn_prep(const bf16* pq,
     const bool q = hd < heads;
     const float x = q ? f32(pq[(size_t)r * 2 * qw + hd * 2 * dim + d]) : f32(pk[(size_t)r * kw + (hd - heads) * dim + d]);
     const float* w = q ? qscale : kscale;
-    const float rinv = 1.0f / sqrtf(block_sum(x * x, scratch, dim / 32) / (float)dim + eps);
+    const float rinv = 1.0f / sqrtf(block_sum_tree(x * x, scratch, dim / 32) / (float)dim + eps);
     xn[d] = f32(b16(x * rinv * w[d]));
     __syncthreads();
     float out = xn[d];
@@ -117,20 +90,20 @@ extern "C" __global__ void __launch_bounds__(256) tf_gattn_q(const bf16* qkv, co
     __syncthreads();
     float top = -INFINITY;
     for (int j = t; j <= at; j += 256) top = fmaxf(top, sc[j]);
-    top = block_max(top, scratch, 8);
+    top = block_max_tree(top, scratch, 8);
     float total = 0;
     for (int j = t; j <= at; j += 256) {
         const float e = expf(sc[j] - top);
         sc[j] = e;
         total += e;
     }
-    total = block_sum(total, scratch, 8);
+    total = block_sum_tree(total, scratch, 8);
     float mix = 0;
     for (int j = t; j <= at; j += 256) {
         sc[j] /= total;
         mix += sc[j] * dp[j];
     }
-    mix = block_sum(mix, scratch, 8);
+    mix = block_sum_tree(mix, scratch, 8);
     const size_t row = ((size_t)h * a.rows + i) * keys;
     for (int j = t; j < keys; j += 256) {
         const bool live = j <= at;
@@ -181,7 +154,7 @@ extern "C" __global__ void __launch_bounds__(256) tf_gattn_prep_back(const bf16*
         gn = d < half ? go[d] * c + go[d + half] * s : go[d] * c - go[d - half] * s;
     }
     const float w = (q ? qscale : kscale)[d];
-    const float sq = block_sum(x * x, scratch, dim / 32), dot = block_sum(w * gn * x, scratch, dim / 32);
+    const float sq = block_sum_tree(x * x, scratch, dim / 32), dot = block_sum_tree(w * gn * x, scratch, dim / 32);
     const float rinv = 1.0f / sqrtf(sq / (float)dim + eps);
     const float g = rinv * w * gn - rinv * rinv * rinv * x * dot / (float)dim;
     if (q)
diff --git a/zig/kernels/cuda/train_mixers.cuh b/zig/kernels/cuda/train_mixers.cuh
index c37bf06df9..c81b894312 100644
--- a/zig/kernels/cuda/train_mixers.cuh
+++ b/zig/kernels/cuda/train_mixers.cuh
@@ -126,7 +126,7 @@ extern "C" __global__ void __launch_bounds__(1024) tf_train_ssm_back(const bf16*
             column_sums(v, scratch, lane, w, t, s.n, db_part + ((size_t)r * s.heads + h) * s.n);
             dx = row_sum(dx);
             if (lane % 16 == 0) dact[(size_t)r * s.conv + h * s.dh + p] = d * dx + d_skip[h] * gy;
-            dd = block_sum(dd, scratch, 32);
+            dd = block_sum_tree(dd, scratch, 32);
             if (t == 0) ddt[r * s.heads + h] = dd;
             for (int i = 0; i < 8; i++) ds[i] *= a;
         }
@@ -157,8 +157,8 @@ extern "C" __global__ void __launch_bounds__(256) tf_train_gate_back(const float
         sq += gated * gated;
         dot += f32(w[j]) * dn[(size_t)r * s.inner + j] * gated;
     }
-    sq = block_sum(sq, scratch, 8);
-    dot = block_sum(dot, scratch, 8);
+    sq = block_sum_tree(sq, scratch, 8);
+    dot = block_sum_tree(dot, scratch, 8);
     const float inv = rsqrtf(sq / (float)width + eps), k = inv * inv * inv * dot / (float)width;
     for (int j = j0 + t; j < j0 + width; j += 256) {
         const float z = f32(proj[(size_t)r * s.proj + j]), yv = y[(size_t)r * s.inner + j];
@@ -228,11 +228,11 @@ extern "C" __global__ void __launch_bounds__(256) tf_train_attn_q(const bf16* qk
         score = sc * a.scale;
         dp = dv;
     }
-    const float top = block_max(score, scratch, 8);
+    const float top = block_max_tree(score, scratch, 8);
     const float e = live ? expf(score - top) : 0.0f;
-    const float total = block_sum(e, scratch, 8);
+    const float total = block_sum_tree(e, scratch, 8);
     const float pr = e / total;
-    const float mix = block_sum(pr * dp, scratch, 8);
+    const float mix = block_sum_tree(pr * dp, scratch, 8);
     const float g = pr * (dp - mix);
     ds[t] = g;
     if (t < keys) {
diff --git a/zig/kernels/cuda/train_q4.cu b/zig/kernels/cuda/train_q4.cu
index fda33541ee..3774d31353 100644
--- a/zig/kernels/cuda/train_q4.cu
+++ b/zig/kernels/cuda/train_q4.cu
@@ -4,6 +4,7 @@
 
 #include <cuda_bf16.h>
 #include <stdint.h>
+#include "common/reduce.cuh"
 
 typedef __nv_bfloat16 bf16;
 
@@ -12,22 +13,8 @@ __device__ __forceinline__ bf16 b16(float v) { return __float2bfloat16(v); }
 __device__ __forceinline__ float wide(float v) { return v; }
 __device__ __forceinline__ float wide(bf16 v) { return __bfloat162float(v); }
 
-__device__ __forceinline__ float warp_sum(float v) {
-#pragma unroll
-    for (int o = 16; o > 0; o >>= 1) v += __shfl_xor_sync(0xffffffffu, v, o);
-    return v;
-}
-
-// The sum over a block of `groups` warps, the same in every thread.
-__device__ __forceinline__ float block_sum(float v, float* scratch, int groups) {
-    const int lane = threadIdx.x & 31, w = threadIdx.x >> 5;
-    v = warp_sum(v);
-    if (lane == 0) scratch[w] = v;
-    __syncthreads();
-    v = warp_sum(lane < groups ? scratch[lane] : 0.0f);
-    __syncthreads();
-    return v;
-}
+using tf_common::warp_sum;
+using tf_common::block_sum_tree;
 
 __device__ __forceinline__ float sigm(float v) { return 1.0f / (1.0f + expf(-v)); }
 __device__ __forceinline__ float silu(float v) { return v / (1.0f + expf(-v)); }
@@ -409,7 +396,7 @@ extern "C" __global__ void __launch_bounds__(128) tf_q4_swiglu_back(const float*
         out[i] = wt * d * uv * silu_grad(gv);
         out[width + i] = wt * d * sg;
     }
-    dot = block_sum(dot, scratch, 4);
+    dot = block_sum_tree(dot, scratch, 4);
     if (t == 0) a[p] = dot;
 }
 
diff --git a/zig/kernels/cuda/tree_attention.cu b/zig/kernels/cuda/tree_attention.cu
index 5ab4bdf8f8..3f1164345e 100644
--- a/zig/kernels/cuda/tree_attention.cu
+++ b/zig/kernels/cuda/tree_attention.cu
@@ -6,13 +6,14 @@
 #include <stdint.h>
 
 #include "attn_odd.cuh"
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 
 namespace tf_tree {
 
 using qmm_frag::commit;
 using qmm_frag::cp16z;
 using qmm_frag::ldmatrix4;
+using qmm_frag::ldmatrix4t;
 using qmm_frag::mma;
 using qmm_frag::mma0;
 
@@ -23,11 +24,6 @@ constexpr float LOG2E = 1.4426950408889634f;
 template <int N>
 __device__ __forceinline__ void wait_group() { asm volatile("cp.async.wait_group %0;\n" ::"n"(N)); }
 
-__device__ __forceinline__ void ldmatrix4t(uint32_t (&r)[4], const void* p) {
-    asm volatile("ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%0, %1, %2, %3}, [%4];\n"
-                 : "=r"(r[0]), "=r"(r[1]), "=r"(r[2]), "=r"(r[3]) : "r"(qmm_frag::smem(p)));
-}
-
 __device__ __forceinline__ float texp(float x) {
     float y;
     asm("ex2.approx.f32 %0, %1;\n" : "=f"(y) : "f"(__fmul_rn(x, LOG2E)));
diff --git a/zig/kernels/cuda/vit.cu b/zig/kernels/cuda/vit.cu
index e9b8becc09..47faadecc8 100644
--- a/zig/kernels/cuda/vit.cu
+++ b/zig/kernels/cuda/vit.cu
@@ -4,7 +4,7 @@
 #include <math.h>
 #include <stdint.h>
 
-#include "glue_math.cuh"
+#include "common/glue_math.cuh"
 #include "vit_norm.cuh"
 #include "vit_gemma.cuh"
 
diff --git a/zig/kernels/cuda/vit_gemma.cuh b/zig/kernels/cuda/vit_gemma.cuh
index b610fb0ead..b6b40102a3 100644
--- a/zig/kernels/cuda/vit_gemma.cuh
+++ b/zig/kernels/cuda/vit_gemma.cuh
@@ -5,7 +5,7 @@
 #include <math.h>
 #include <stdint.h>
 
-#include "glue_math.cuh"
+#include "common/glue_math.cuh"
 
 namespace tf_vit_g4 {
 
diff --git a/zig/kernels/cuda/vit_norm.cuh b/zig/kernels/cuda/vit_norm.cuh
index f9c589fb69..4e331ac252 100644
--- a/zig/kernels/cuda/vit_norm.cuh
+++ b/zig/kernels/cuda/vit_norm.cuh
@@ -4,7 +4,7 @@
 #include <cuda_bf16.h>
 #include <stdint.h>
 
-#include "glue_math.cuh"
+#include "common/glue_math.cuh"
 
 namespace tf_vit_norm {
 
diff --git a/zig/kernels/cuda/w4a8_prefill.cu b/zig/kernels/cuda/w4a8_prefill.cu
index b60a7e7d64..76a7a7a2a8 100644
--- a/zig/kernels/cuda/w4a8_prefill.cu
+++ b/zig/kernels/cuda/w4a8_prefill.cu
@@ -2,7 +2,7 @@
 #include <cuda_bf16.h>
 #include <stdint.h>
 
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 #include "w4a8_quantize.cuh"
 #include "w4a8_store.cuh"
 
diff --git a/zig/kernels/cuda/window_attention.cu b/zig/kernels/cuda/window_attention.cu
index b4c749c105..5d34eb8779 100644
--- a/zig/kernels/cuda/window_attention.cu
+++ b/zig/kernels/cuda/window_attention.cu
@@ -6,7 +6,7 @@
 #include <stdint.h>
 
 #include "attn_odd.cuh"
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 #include "window_split.cuh"
 
 // Prompt rows: prefill_attention.cu's tensor-core kernel with a window (keys (pos - win, pos]) and a ring of keys.
@@ -15,6 +15,7 @@ namespace tf_window_prompt {
 using qmm_frag::commit;
 using qmm_frag::cp16z;
 using qmm_frag::ldmatrix4;
+using qmm_frag::ldmatrix4t;
 using qmm_frag::mma;
 using qmm_frag::mma0;
 
@@ -24,11 +25,6 @@ constexpr float LOG2E = 1.4426950408889634f;
 template <int N>
 __device__ __forceinline__ void wait_group() { asm volatile("cp.async.wait_group %0;\n" ::"n"(N)); }
 
-__device__ __forceinline__ void ldmatrix4t(uint32_t (&r)[4], const void* p) {
-    asm volatile("ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%0, %1, %2, %3}, [%4];\n"
-                 : "=r"(r[0]), "=r"(r[1]), "=r"(r[2]), "=r"(r[3]) : "r"(qmm_frag::smem(p)));
-}
-
 __device__ __forceinline__ float texp(float x) {
     float y;
     asm("ex2.approx.f32 %0, %1;\n" : "=f"(y) : "f"(__fmul_rn(x, LOG2E)));
diff --git a/zig/kernels/cuda/window_rows.cuh b/zig/kernels/cuda/window_rows.cuh
index 37402546d5..deb4efdb15 100644
--- a/zig/kernels/cuda/window_rows.cuh
+++ b/zig/kernels/cuda/window_rows.cuh
@@ -6,7 +6,7 @@
 #include <math.h>
 #include <stdint.h>
 
-#include "qmm_frag.cuh"
+#include "common/qmm_frag.cuh"
 #include "window_fallback.h"
 
 namespace tf_wrows {
@@ -39,6 +39,7 @@ struct Pg {
 using qmm_frag::commit;
 using qmm_frag::cp16z;
 using qmm_frag::ldmatrix4;
+using qmm_frag::ldmatrix4t;
 using qmm_frag::mma;
 using qmm_frag::mma0;
 
@@ -300,12 +301,6 @@ __device__ __forceinline__ void merge(const float* __restrict__ po, const float*
 
 namespace prompt {
 
-__device__ __forceinline__ void ldmatrix4t(uint32_t (&r)[4], const void* p) {
-    asm volatile("ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 {%0, %1, %2, %3}, [%4];\n"
-                 : "=r"(r[0]), "=r"(r[1]), "=r"(r[2]), "=r"(r[3])
-                 : "r"(qmm_frag::smem(p)));
-}
-
 template <int N>
 __device__ __forceinline__ void wait_group() { asm volatile("cp.async.wait_group %0;\n" ::"n"(N)); }
 
diff --git a/zig/src/families/flashnext/hybrid_ple_source.zig b/zig/src/families/flashnext/hybrid_ple_source.zig
index ec26a6c4f6..6ad8a94e75 100644
--- a/zig/src/families/flashnext/hybrid_ple_source.zig
+++ b/zig/src/families/flashnext/hybrid_ple_source.zig
@@ -16,5 +16,5 @@ pub fn cuda(a: std.mem.Allocator) ![]u8 {
     conv = try edit.once(mem, conv, "bf(CW[static_cast<long long>(e) * TAPS + t])", "CW[static_cast<long long>(e) * TAPS + t]");
     const sigmoid = try edit.once(mem, try edit.function(original, "__device__ __forceinline__ float bsig("), "return rbf(tf_sigmoid(x));", "const float e = tf_exp(-fabsf(x)); return rbf(__fdiv_rn(x < 0 ? e : 1.0f, __fadd_rn(1.0f, e)));");
     const silu = try edit.once(mem, try edit.function(original, "__device__ __forceinline__ float bsilu("), "return rbf(tf_silu(x));", "const float e = tf_exp(-fabsf(x)); return rbf(__fdiv_rn(x < 0 ? __fmul_rn(x, e) : x, __fadd_rn(1.0f, e)));");
-    return std.mem.concat(a, u8, &.{ "#include <cuda_bf16.h>\n#include \"glue_math.cuh\"\nusing namespace tf_glue;\n", sigmoid, "\n", silu, "\n", try edit.function(original, "__device__ __forceinline__ float block_sum("), "\n", gate, "\n", conv, "\n" });
+    return std.mem.concat(a, u8, &.{ "#include <cuda_bf16.h>\n#include \"common/glue_math.cuh\"\nusing namespace tf_glue;\n", sigmoid, "\n", silu, "\n", try edit.function(original, "__device__ __forceinline__ float block_sum("), "\n", gate, "\n", conv, "\n" });
 }
diff --git a/zig/tests/fixtures/gradient-serving.tsv b/zig/tests/fixtures/gradient-serving.tsv
index 6e59e0fe1a..fa1daab5b1 100644
--- a/zig/tests/fixtures/gradient-serving.tsv
+++ b/zig/tests/fixtures/gradient-serving.tsv
@@ -1,6 +1,6 @@
 zig/kernels/cuda/affine4_pack.cu	cbc03ae0622574b57ae266091ea9f26888b21b2b1c020a031940eacddcf8ff14	raw
-zig/kernels/cuda/ck_experts.cu	9f670bf9f98164cf8aeed2a7b21c22d68540620a8679f2fbf4aa2509e524aacf	raw	1.1.0 checkpoint NVFP4 serving kernels; Flash Next learning reads separate affine4 weights and adjoints.
-zig/kernels/cuda/ck_lane.cu	913a0df7bef408ed561c657de74356b3fcb67e328552d6b2b18d4d77df1252d3	raw	commits c0a58fd30; HELD: must be re-qualified by the LW device gates on the 1.1.1 candidate
+zig/kernels/cuda/ck_experts.cu	8f14bb5377efe3b7b7b7ab677807b2e224d0117ca73b468690305b6e0cf13fb9	raw	1.1.0 checkpoint NVFP4 serving kernels; Flash Next learning reads separate affine4 weights and adjoints.
+zig/kernels/cuda/ck_lane.cu	f2fcca72d0fe9009728c44fc09dd91a681ff5b978141a0f64f5ba52792ae13ec	raw	commits c0a58fd30; HELD: must be re-qualified by the LW device gates on the 1.1.1 candidate
 zig/kernels/cuda/contract_controls.cu	f8ac481cd87c630d4562a07dbebe926a2210d4b2a1245955a3ebb284d296bdd5	raw
 zig/kernels/cuda/embed_rows.cu	09c3eea45d826e9bc2cbf8a6719f03a0667134f7dd9e71a1b59d2db444b52768	raw
 zig/kernels/cuda/exlr8.cu	0eb19c8965aeb8188efd2b15a5844da860d9bf75364b12b62dd3a2589a89178c	raw
@@ -16,7 +16,7 @@ zig/kernels/cuda/flashnext_gdn.cu	979106b8f7424402d22373850763c46f3803ce03f6a259
 zig/kernels/cuda/flashnext_hc.cu	493a03dc898a0020ab024fea46b95e031a9b52c9905b46559394485d966e3caf	raw
 zig/kernels/cuda/flashnext_mtp.cu	051e151bb5b7cb8ec1059b41cc24a024ae00c8f8477d3b3afb95c10f47c89d4e	raw
 zig/kernels/cuda/fp8_experts.cu	c350a6e060bc35fab9c3f69ad0850ed6dca6ace370dc7431f7dc354ab90a0a65	raw
-zig/kernels/cuda/gdn.cu	a6bf09198ffd685f7392d4d433af9739c91d67f7039309d3fd5fe6785f810ab9	raw
+zig/kernels/cuda/gdn.cu	a9caeec5b270dcab933a5eae0f500d7aabcd31cde156238225ae064ac92ba2e8	raw
 zig/kernels/cuda/glm_glue.cu	8c71cb95fafdcc6a1540cb260edc40f5ef280b2d549e5b39a66ecdc60fbdaa9b	raw
 zig/kernels/cuda/glm_glue_attn.cuh	d8fea79905f2ee6773983f2b6244fb60599e608a47068f72c4b5f3b098966ffe	raw
 zig/kernels/cuda/glm_glue_bmm.cuh	103ad235e03c75482de99ad578d91fdd3e7467e1d522b895376e7d841afdc5d2	raw
@@ -27,32 +27,32 @@ zig/kernels/cuda/glm_glue_moe.cuh	e39f5de255684c81f9a005ea40048b474ce8910ee76e84
 zig/kernels/cuda/glm_glue_norm.cuh	c243a44d064dd464c57090703f0c98ff64b05ee005ea93f86bc529d9c32c16bc	raw
 zig/kernels/cuda/glm_glue_router.cuh	0ba25f8cd6842c30d328b5c2db3e2a84453ab52f97f236be2643c9fb7ac6285e	raw
 zig/kernels/cuda/glm_glue_rows.cuh	e7c81f4f5190f83644b2d14155db768fc1d5f862a76734c731626e3757d49a2c	raw
-zig/kernels/cuda/glm_kda.cu	d88d7b61ee168667352b8bfb4dec12e77c218d6087e5ec06a43157ad544e7572	raw
+zig/kernels/cuda/glm_kda.cu	04f88d8e221126e6b471fe272c808808532fd37a65784faa01d81e60cfa3a384	raw
 zig/kernels/cuda/glm_ops.cu	b96cf786a8ab0915a7c2c5f3df2d132fd2e0eb03b15ad03696e6e4b71cc61093	raw
 zig/kernels/cuda/glm_quant.cu	76c4575639d41d7097f76c3ade91150b6ed14359bf59a61222c84cf99f38ba81	raw
 zig/kernels/cuda/glm_triton.cuh	8eee12f24bbf283a08a29f9396e3db5588b48fce85abe74de7f85adc12fa94f4	raw
-zig/kernels/cuda/glue_math.cuh	e836bbca47980e3120584d5ad67519f7a6936fb7dc707103a10e63e769fa1938	raw
+zig/kernels/cuda/common/glue_math.cuh	b7234684fc786b3700473358632bf10f0b2d2240e5957eefddb13254dae5d746	raw
 zig/kernels/cuda/lane_gemv.cu	9551469ae616283ee893d7f117d143c43de08bc868650f539cf38570a47e7a13	raw
-zig/kernels/cuda/lfm2.cu	3f7771b095ba6bbbdc56c6595de9aaa3d926434e8cea85e23441ddc0c2d03cff	raw
+zig/kernels/cuda/lfm2.cu	d51f85b5b34cd5d172b8a0f989752706e2042d8c2178b942c1b9ae55a8b600ab	raw
 zig/kernels/cuda/logprob_rows.cuh	ee5711c4cb8656b7aa887ab4aa2b73ec4e8e65fdb5dc5aca8ab21f5ef499a5d2	raw	series patches 02-cuda-rounds.patch; HELD: must be re-qualified by the LW device gates on the 1.1.1 candidate
 zig/kernels/cuda/mma4.cuh	52152491c8b53572952030e059106b186db1b9b535d55f0a7543fbe9c807a7a9	raw
 zig/kernels/cuda/modelopt_pack.cu	2ceb2ab450f5583c2007f216569d2196d800607144240591b87ccb0d43ff65d6	raw
-zig/kernels/cuda/nemotron_attention.cu	37c96d91e69cac2be417106b28d154e278527ab632e264b3620efcb42676f6d5	raw
+zig/kernels/cuda/nemotron_attention.cu	c04ad68bea1202e5c6f3d289af3745b824b89a8aa2fe697d5c1c8cebefe07151	raw
 zig/kernels/cuda/nemotron_keyed.cu	ed8bdd7012970d9896ec703bd7c3d5ec83e02f41dba4ce1f896fc39c1ee349c9	raw
-zig/kernels/cuda/nemotron_latent.cu	0f2852d1940c7a25ccbf7159c5f8812e061e7239d51cac053e5386d171dbed0f	raw	commits e40bd6eae, 8a8ff3d64; HELD: must be re-qualified by the LW device gates on the 1.1.1 candidate
-zig/kernels/cuda/nemotron_mamba.cu	f5c6c4bfe99e8a01f01807173becd8121a486409187626c1f2f901586cf58981	raw	commits fa2af585f, f0ef17e69; HELD: must be re-qualified by the LW device gates on the 1.1.1 candidate
-zig/kernels/cuda/nemotron_norms.cu	8d1de19123d86cd17a819d1f3b4f1733386b047aa7ccc6632ecfe737cd837005	raw
+zig/kernels/cuda/nemotron_latent.cu	04fba97043d0948f82a3b9477f448dd76b21f6d11d1a445e020bd41cacb85c2b	raw	commits e40bd6eae, 8a8ff3d64; HELD: must be re-qualified by the LW device gates on the 1.1.1 candidate
+zig/kernels/cuda/nemotron_mamba.cu	dcaba12253d3931ae48c7ebd4e13dfb956f540ae19f4ab83130c3f21745be5ec	raw	commits fa2af585f, f0ef17e69; HELD: must be re-qualified by the LW device gates on the 1.1.1 candidate
+zig/kernels/cuda/nemotron_norms.cu	b3ac0d2dadc41c94e9bb77be5b9687e71886a55072e24d6d9b7ce1bc1eb79242	raw
 zig/kernels/cuda/nemotron_ops.cu	b2b1a4441202229cc6989b2ae67259ea680c328b2d8cb2284985cece3c27d38b	raw
-zig/kernels/cuda/nemotron_route.cu	89f782c9762e1038fc9b6560f5261adf3716d44a64cfb4d2e9fff9948efb056b	raw	commits 739c924de, 8cfdaa739, b77a33b08, 7f6f940eb, 05d0617af, bc76d90d6; HELD: must be re-qualified by the LW device gates on the 1.1.1 candidate
+zig/kernels/cuda/nemotron_route.cu	74cda9ea64addcc8b1288e1f0f0db296a5378dccfcbc7962c432c5cda32817fa	raw	commits 739c924de, 8cfdaa739, b77a33b08, 7f6f940eb, 05d0617af, bc76d90d6; HELD: must be re-qualified by the LW device gates on the 1.1.1 candidate
 zig/kernels/cuda/norms.cu	ed5a2dcb4dfc81ed048a422a707a4161a217c5b1aa06af82ec8b8e2d66bd6ae3	raw
 zig/kernels/cuda/nvfp4_experts.cu	ac07cf0a5799757b878b7add5e8bbe2673f8bc84af6c3c7ab6f52510e1c81b98	raw
 zig/kernels/cuda/nvfp4q.cuh	5fe4fa58568f2293bb72ccfa58b362ba2783b21a781e56946a5a2c328303d111	raw
 zig/kernels/cuda/prefill_attention.cu	71b9aba5ba586b12454a4c2cb8168e59227aed4db8bc78190ecc9f262ba3d62f	raw
 zig/kernels/cuda/prefill_attention_rows.cuh	ab6fc82651f1cababa480726f6bdf3689a6b999eb83d9cf37f370501b778a2de	raw
 zig/kernels/cuda/probe.cu	c5a6ca4513311d25282866de2185a056ec9746443a13094db80de62f7474255a	raw
-zig/kernels/cuda/prompt16.cu	13c24ebf15ad912d62f8830d59cbfa7c9fce0c76112dac3c84cffce7fc671b26	raw
+zig/kernels/cuda/prompt16.cu	cf373a963cf525d8f658820b31d4be8b52970a80adcc4b19bbd59f0539abed81	raw
 zig/kernels/cuda/qmm.cu	91741c5aecc2c97bcb4c00b83d22220be5658112b2f370c146ae8393353d668b	raw
-zig/kernels/cuda/qmm_frag.cuh	e8030c07c2926ecdc15d71a6e13591a43dac524cd3aa85e388cfafc746bf4c07	raw
+zig/kernels/cuda/common/qmm_frag.cuh	7f08de88be73341122d5a13474b29f87a973330cfd142add5e82435722f8e9ef	raw
 zig/kernels/cuda/qmm_group.cu	3b27b8e964f6e7453e962cd7a259253ae7e1940297439b4b518f05da04dd1041	raw
 zig/kernels/cuda/qmm_prefill.cu	6dbafecbf5c28b48da1bcd222f85494b471e728a482eb17fc5fffa8c522dee95	raw
 zig/kernels/cuda/qmmf.cu	40a2a342cfceca25a31f53903f9e3eb08aa38114433af094dcc1f60566614c1b	raw
@@ -77,7 +77,7 @@ zig/kernels/cuda/torch_ops/norm_rotary.cu	759fd44f9d1c32f586fab33e8f796a310cfc62
 zig/kernels/cuda/torch_ops/nucleus.cu	c3328dc3ad1fb28c820c343f22c10dd1b4f9a9d8c87c3d38c9f7fc0730d1a754	raw
 zig/kernels/cuda/torch_ops/pointwise.cu	3d3118159265538679968e56a8c4b5d61b01c07487530e95e4f93ba8662f98c8	raw
 zig/kernels/cuda/torch_ops/topk.cu	0b5ddb5b02af2ebd9a3992cb11bb80eefd7547c32d6dfdafdf216ab484dd1627	raw
-zig/kernels/cuda/train_mixers.cuh	8370208f737c65bce6b76531e50bf1f44a208eab1cfced3dd87b8924b039763d	raw
+zig/kernels/cuda/train_mixers.cuh	cec88446cede27c941da254d94f81afa1c92c6880e63bfd66a0efcdc7f95aa13	raw
 zig/kernels/cuda/vit.cu	9cc1bdb9cc3db2a6cbe07549f90d806004f286a6cea86b1e9f7c5bfdb0aa8ef7	raw
 zig/src/cuda/abi.zig	8ccdbb31eaac67a5e9c3fc2ca8bc3e4eb33c83e5c0c2ffd98e1aa4a966d94e2f	raw	commits ef23b5934; series patches 14-runtime-foundation.patch; HELD: must be re-qualified by the LW device gates on the 1.1.1 candidate
 zig/src/cuda/aot.zig	08324c7c982b1b4f95e7316662d015e809b28677ce3816b0c1a3c61d0993ae99	raw	commits 9349e38bb; series patches 14-runtime-foundation.patch, 45-runtime-one-fallback.patch, 46-preserve-nospec-and-capture.patch; HELD: must be re-qualified by the LW device gates on the 1.1.1 candidate
diff --git a/zig/tests/fixtures/pocket-serving-approved.tsv b/zig/tests/fixtures/pocket-serving-approved.tsv
index ecb3f88ddd..cd205fd910 100644
--- a/zig/tests/fixtures/pocket-serving-approved.tsv
+++ b/zig/tests/fixtures/pocket-serving-approved.tsv
@@ -41,7 +41,7 @@ zig/kernels/cuda/glm_glue_attn.cuh	daabfe78a2e747575aadc07bcee7324ab7e5fa43d2871
 zig/kernels/cuda/glm_glue_latent.cuh	85682ea59a41a5d4f6e07ee209bc51084b67a7aca20bd4b2be8fd3bab42c20a8	0	1.1.2 increments 118; approved 20 Oct 2026
 zig/kernels/cuda/glm_quant.cu	e27379a4192be52f390ad28f678157d592989af45b69d31f830671b2c54ca028	0	1.1.2 increments 126; approved 20 Oct 2026
 zig/kernels/cuda/norms.cu	0af0d9452803066422640857acc8509872dffe5ee5c029c6084ea7b22ddf8da8	0	1.1.2 increments 215 (#1056); approved 20 Oct 2026
-zig/kernels/cuda/route_rows.cu	64c1ac1064b2de6fcd70c68272295945c6ae4610deff84eb85d57d411b44ce5f	0	1.1.2 increments 234 (#1046), 241 (#1046), 243 (#1046); approved 20 Oct 2026
+zig/kernels/cuda/route_rows.cu	6a9c4eac0526465d8f6aae8f8bb9100bc85012b827e69e9c252896815383091a	0	1.1.2 increments 234 (#1046), 241 (#1046), 243 (#1046); approved 20 Oct 2026; #1440 moves its shared helpers to common/ (SASS unchanged)
 zig/src/cuda/embed_rows.zig	09fb2e612f43297c19d2cb77d45139dd98212a20075da0a15069a70544f72f6b	0	1.1.2 increments 220 (#1070); approved 20 Oct 2026
 zig/src/cuda/norms.zig	c87aeffa84b77119f714faa1aa764d72633f2809cbb8271d2f80be2455e34b37	0	1.1.2 increments 215 (#1056); approved 20 Oct 2026
 zig/src/cuda/route_rows.zig	938fa2311aa530a845c1218867734ea0a93f762bebb574d31b9aab83d2f3aa3b	0	1.1.2 increments 234 (#1046), 241 (#1046), 245 (#1046); approved 20 Oct 2026