Skip to content

fix(nsa): tilelang sparse kernel hardcoded head offset + sm_120 smem-fit geometry - #70

Open
architehc wants to merge 1 commit into
kvcache-ai:mainfrom
architehc:sm120-tilelang-head-offset
Open

fix(nsa): tilelang sparse kernel hardcoded head offset + sm_120 smem-fit geometry#70
architehc wants to merge 1 commit into
kvcache-ai:mainfrom
architehc:sm120-tilelang-head-offset

Conversation

@architehc

Copy link
Copy Markdown
  • `sparse_attention_fwd_kernel_v1` hardcodes `(bx % REPLICATE_H) * 64` for the head offset, which breaks any head split where `H_per_block != 64`. Fixed to `* H_per_block`.
  • Adds an sm_120 geometry: `H_per_block=16`, `threads=128`, v1 with `num_stages=1` — fits the 99 KB shared-memory budget of consumer/workstation Blackwell (the default geometries need 154–226 KB and only fit Hopper/SM100).
  • `kv_dtype` parametrization so the kernel can load fp8 KV pools directly.
  • Drops the `wg_wait` kwarg removed in tilelang ≥ 0.1.9.

With this, tilelang NSA prefill and decode run on sm_120 (validated with tilelang 0.1.12 + apache-tvm-ffi 0.1.11 + nvidia-cutlass-dsl 4.5.0, serving GLM-5.2 at 1M context).

🤖 Generated with Claude Code

…plit

sparse_attention_fwd_kernel_v1 hardcodes a *64 head offset that breaks
any non-64 head split; fix to *H_per_block. Adds an sm_120 geometry
(H_per_block=16, threads=128, v1/num_stages=1) that fits the 99KB smem
budget, kv_dtype parametrization for fp8 KV pools, and drops the wg_wait
kwarg removed in tilelang >= 0.1.9.
@gemini-code-assist

Copy link
Copy Markdown

Caution

The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant