[WIP]: feat(cutedsl-sm100): BWD LoopK with warp-specialized pipeline + IndexSparse/BlockSparse support

[cennn]

Summary

在 CuTe-DSL SM100 BWD 内核中新增 InnerLoopK 路径（swap_bwd_qk_loop=True），支持外层遍历 Q blocks、内层流式迭代 K/V blocks。在此基础上支持 IndexSparse 和 BlockSparse 的 BWD LoopK，使得 IndexAttn 等稀疏注意力场景可以选择 LoopK 方向以获得更好的 dQ 存储效率（dQ TMEM 累积 + 单次 store，而非 LoopQ 的 atomicAdd dQ）。

核心改动

1. BWD LoopK warp-specialized pipeline

   • 新增 load_loop_k（producer）、mma_loop_k（consumer）、dKV_reduce_loop_k（reduce warps）三个 LoopK 专用方法
   • 独立 pipeline_K_load（与 pipeline_Q 分离，避免 mbarrier 冲突）
   • TMEM 布局重新编排：dQ persistent @0，S/P→dK @hdim，dP/dS @hdim+tile_m，dV @hdim+2*tile_m
   • dKV per-K-iter 通过 reduce warps 做 TMEM→SMEM→cpasync_reduce_bulk_add_f32 原子归约到 GMEM
   • dQ 在外层 tile 结束后一次性 reduce 到 GMEM

2. IndexSparse + BlockSparse BWD LoopK

   • 新增 index_attn_indices_to_block_sparse()：token-level indices (total_q, NHK, topk) → block-level BlockSparseTensors（GPU 矢量化，scatter_ 构建 presence matrix）
   • 新增 prepare_block_sparse_bwd_loopk()：forward-direction normalize（per M-block → N-block list）
   • LoopK 所有 warp（load/mma/compute/reduce）的 sparse 路径共享 get_curr_blocksparse_tensors + get_m_block_from_iter_bwd

3. 其他

   • LoopK 强制 dKV_postprocess=True（fp32 累积 → dtype 转换 + softmax_scale 缩放）
   • LoopK 禁用 2-CTA mode（disable_2cta 条件扩展）
   • K_stage=1 for D≥128（避免 SMEM 超限）；K_stage=2 for D<128

正确性验证

B300 SM103 环境，D=128, bf16:

Dense LoopK（vs LoopQ, bit-identical, max_diff=0.000000）:

配置	dQ	dK	dV
MHA full B=2 Sq=256	PASS	PASS	PASS
MHA full B=1 Sq=1024	PASS	PASS	PASS
MHA causal B=2 Sq=256	PASS	PASS	PASS
GQA full B=2 NHQ=8 NHK=2	PASS	PASS	PASS
GQA causal B=2 NHQ=8 NHK=2	PASS	PASS	PASS

BlockSparse LoopK（vs LoopQ reference）:

配置	dQ	dK	dV
MHA B=1 Sq=512 sparsity=0.5	PASS (0.0005)	PASS (0.001)	PASS (0.0005)
MHA B=2 Sq=512 sparsity=0.75	PASS (0.001)	PASS (0.001)	PASS (0.001)

IndexSparse LoopK（vs LoopQ reference）:

配置	dQ	dK	dV
MHA B=1 Sq=512 topk=256	PASS (0.002)	PASS (0.0005)	PASS (0.0005)
MHA B=2 Sq=512 topk=256	PASS (0.001)	PASS (0.001)	PASS (0.001)

限制

• LoopK 当前仅支持 1-CTA、non-varlen、hdim≤128
• LoopK 始终走 dKV_postprocess 路径（fp32 原子累加），尚未做 direct TMA store 优化