cuda: sliding-window attention, sigmoid top-k routing and combine, dense rows for every family, checked against float64 by jschmied · Pull Request #798 · ashhart/TensorFold

tests/cuda/contract_attn: the reference's constructor frees what it took when a later allocation fails, with a host test that fails each allocation in turn

jschmied

committed

0ad4f17

cuda/block_attention: the replay takes every NaN and infinite value out of the mma and adds it back at its key's float64 weight, so an infinite value gives NaN where float64 does

jschmied

committed

2ae3b69

cuda/window_attention: the prompt replay takes every NaN and infinite value out of the mma and adds it back at its key's float64 weight (attn_odd), so an infinite value gives NaN where float64 does

jschmied

committed

462d952

cuda/block_attention: the replay's block tile keeps the infinities added back before it

jschmied

committed

3c82eb3

cuda/tree_attention: a NaN or infinite partial runs again with those values out of the mma, added back at float64's weights

jschmied

committed

0facb0f

cuda/window_attention: the decode kernels run a NaN or infinite tile again with those values out of the sums, added back at float64's weights

jschmied

committed

8cba69e

cuda/tree_attention, block_attention: the replay's comments on one line each

jschmied

committed

965c8cd

cuda/window_attention: the split decode and its merge move to window_split.cuh (the file was past 600 lines), comments on one line

jschmied

committed

74b749f

Navigation Menu

Repository navigation

cuda: sliding-window attention, sigmoid top-k routing and combine, dense rows for every family, checked against float64 - #798

cuda: sliding-window attention, sigmoid top-k routing and combine, dense rows for every family, checked against float64#798
jschmied wants to merge 23 commits into
ashhart:mainfrom
jschmied:cuda-moe-dense-window

Commits on Oct 16, 2026

cuda: tree attention and DFlash block attention for every family (tree_attention.cu, block_attention.cu), checked against core/attention_ref.zig

cuda/tree_attention: a group's 4 slices write their own partials, the merge combines them in order; a slice takes its own one-warp block where four-warp blocks would leave SMs idle (TreeAttention.s…

cuda/tree_attention: four-warp blocks combine a group's slices themselves again, only one-warp blocks leave them to the merge (a small fixed slice area), so the scratch stays at one partial a group

cuda/tree_attention: one-warp blocks fetch their tiles 2 ahead (3 cp.async stages), the same arithmetic in the same order

cuda/tree_attention: 512-key groups (4x the warps in flight at one row), partials sized to the query heads, one-warp blocks with 3 tile stages where they fit one wave of the SMs, else 2

cuda/tree_attention: the merge issues a chunk's partial loads (16 slots) before folding them, the same arithmetic in the same order

cuda: sliding-window attention, sigmoid top-k MoE routing and combine, and dense rows for every family, checked against float64

Commits on Oct 17, 2026

cuda/tree_attention: each row's own max flag, so a NaN head leaves other rows' keys

cuda/block_attention: keys a row cannot see add nothing, so an infinite value there stays out of its sum

cuda/tree_attention, block_attention: registered with the kernel contract

cuda/window_attention: the tensor-core decode keeps each head's own max flag, so a NaN head leaves the other heads' keys

cuda/window_attention: the split decode's slots outside the window hold zeros, not a copy of the window's first key, so an infinite value there reaches no sum

cuda/window_attention: prompt rows' keys a row cannot see add nothing, so a later row's infinite value stays out of earlier rows' sums

cuda/window_attention, moe_glue, dense_rows: registered with the kernel contract

cuda/window_attention: the tensor-core decode's per-head max needs no flag (a head with no key keeps m, exp2 of +0 is exactly 1), the old kernel's speed

Commits on Oct 18, 2026

tests/cuda/contract_attn: the reference's constructor frees what it took when a later allocation fails, with a host test that fails each allocation in turn

cuda/block_attention: the replay takes every NaN and infinite value out of the mma and adds it back at its key's float64 weight, so an infinite value gives NaN where float64 does

cuda/window_attention: the prompt replay takes every NaN and infinite value out of the mma and adds it back at its key's float64 weight (attn_odd), so an infinite value gives NaN where float64 does

cuda/block_attention: the replay's block tile keeps the infinities added back before it

cuda/tree_attention: a NaN or infinite partial runs again with those values out of the mma, added back at float64's weights

cuda/window_attention: the decode kernels run a NaN or infinite tile again with those values out of the sums, added back at float64's weights

cuda/tree_attention, block_attention: the replay's comments on one line each

cuda/window_attention: the split decode and its merge move to window_split.cuh (the file was past 600 lines), comments on one line