Realization.Nvidia.SM86.StreamingAttentionSM86

=========================================================================== The forward pass with each key tile's K and V^T staged in shared memory. The register kernel above has every warp fetch its own B fragments from global memory, then wait for them, before each tile's products: with two warps per scheduler little hides that latency. Here the block's 128 threads copy tile t + 1 (cp.async, 16 bytes a thread a step) into one shared buffer while the warps compute tile t from the other, and LDSM.x4 fills exactly the B-fragment registers the register kernel loads (same values, same lanes), so every product and every result word is unchanged. A row of 64 halves takes 144 bytes (128 + 16 of padding), which keeps LDSM's eight row reads and the copies' 16-byte stores on distinct banks. The loop body is emitted once per buffer (LDSM's address register is never rewritten); an odd tile count leaves after the first body by a forward branch.

575def sfRowBytes : Nat = 144