[kernel] sm_100a CUDA backward for VSA block-sparse attention (blk64) - #1819

Merged

SolitaryThinker merged 10 commits into

hao-ai-lab:mainfrom

lpc0220:pengchengl/vsa-sm100a-backward

Sep 7, 2026

Commits on Sep 5, 2026

[kernel] sm_100a CUDA backward for VSA block-sparse attention (64-token blocks)

lpc0220
and
claude
committed
749a9c5
[kernel] sm_100a backward: keep each softmax warp's bf16 overlay inside its own fp32 columns

lpc0220
and
claude
committed
88319b7
[kernel] sm_100a backward: drop the cached identity order array, nullptr remap = identity

lpc0220
and
claude
committed
57bb232
[kernel] sm_100a backward: drop tcgen05.wait::ld after the TMEM loads

lpc0220
and
claude
committed
d2c9d61

Commits on Sep 6, 2026

[fix]: restore tcgen05.wait::ld after the TMEM loads in the sm_100a backward

SolitaryThinker
and
claude
committed
60165c7
[misc]: pin the identity-order/chunked-launch invariant; fix set_extension docstring

SolitaryThinker
and
claude
committed
62f0bc1
[test]: sm_100a backward parity and run-to-run stability at 65536 tokens

SolitaryThinker
and
claude
committed
2d58c9b

Commits on Sep 7, 2026

[kernel] sm_100a backward: L2 transition as the per-head dQ accumulator bytes; order kernel whenever the launch is chunked

root
and
claude
committed
c3b05a8
[kernel] sm_100a backward: order gate back to nb >= ORDER_MIN_KV_BLOCKS; shorter comments

root
and
claude
committed
0ab2655
[test]: sm_100a backward at the smallest legal shape, num_blocks = 2

root
and
claude
committed
f947567