{% macro sparse_schedule() %}
// How much history precedes this call's tokens. Prompt mode is determined from
// scalar total_sequence_length rather than per-batch values and forces the past
// length to zero; a padded prompt row must not be read as if it had history.
fn past_sequence_length(batch: u32) -> u32 {
  if (u32(total_sequence_length[0]) == params.seqLen) {
    return 0u;
  }
  let total = u32(key_total_sequence_lengths[batch]);
  return select(0u, total - params.seqLen, total >= params.seqLen);
}
{%- endmacro %}
{% macro sparse_rotary(interleaved) %}
// Which cos/sin entry a component uses, and which member of its rotation pair it is.
// The two layouts differ only here: the NeoX split pairs d with d + ROTARY_HALF, and the
// interleaved layout pairs the even lane with the odd one beside it.
fn rotary_slot(d: u32) -> u32 {
{% if interleaved %}
  return d / 2u;
{% else %}
  return select(d - ROTARY_HALF, d, d < ROTARY_HALF);
{% endif %}
}

fn rotary_partner(d: u32) -> u32 {
{% if interleaved %}
  return select(d - 1u, d + 1u, (d % 2u) == 0u);
{% else %}
  return select(d - ROTARY_HALF, d + ROTARY_HALF, d < ROTARY_HALF);
{% endif %}
}

fn rotary_is_first(d: u32) -> bool {
{% if interleaved %}
  return (d % 2u) == 0u;
{% else %}
  return d < ROTARY_HALF;
{% endif %}
}

// One component of the rotation, written from the component's own point of view so a
// single expression covers both members of the pair: the leading one subtracts its
// partner's sine term, the trailing one adds it.
fn rotary_value(own: f32, partner: f32, cs: f32, sn: f32, first: bool) -> f32 {
  return select(own * cs + partner * sn, own * cs - partner * sn, first);
}
{%- endmacro %}

{{ env.wgsl.resourceDeclarations }}

// com.microsoft.SparseAttention, KV append pass.
// past_key and present_key are the same allocation, so this writes only the new rows:
// one invocation per (batch, kv head, new token, component) of the BNSH cache, landing
// at absolute position past + t. Nothing rewrites the retained history.
const KV_HEADS: u32 = {{ kvNumHeads }}u;
const HEAD_DIM: u32 = {{ headSize }}u;
const MAX_CACHE_SEQ: u32 = {{ maxCacheSeq }}u;
{% if packedQkv %}
// Packed layout: query carries [Q | K | V] on one row, so K starts after the query
// heads and V after the key heads.
const Q_STRIDE: u32 = {{ packedStride }}u;
const PACKED_K_OFFSET: u32 = {{ numHeads * headSize }}u;
const PACKED_V_OFFSET: u32 = {{ (numHeads + kvNumHeads) * headSize }}u;
{% else %}
const KV_HIDDEN: u32 = {{ kvNumHeads * headSize }}u;
{% endif %}
const WG: u32 = {{ appendWorkgroupSize }}u;
{% if usesRotary %}
const ROTARY_HALF: u32 = {{ rotaryHalf }}u;
const ROTARY_DIM: u32 = {{ rotaryDim }}u;
{% endif %}

{{ sparse_schedule() }}
{% if usesRotary %}

{{ sparse_rotary(rotaryInterleaved) }}
{% endif %}

@compute @workgroup_size(WG, 1, 1)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  // 2D-folded flat index: gid.y carries the high bits past the
  // per-axis dispatch fold width. Reduces to gid.x when the dispatch does not fold.
  let index = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * WG;
  let count = params.batchSize * KV_HEADS * params.seqLen * HEAD_DIM;
  if (index >= count) {
    return;
  }
  let d = index % HEAD_DIM;
  let t = (index / HEAD_DIM) % params.seqLen;
  let kv_head = (index / (HEAD_DIM * params.seqLen)) % KV_HEADS;
  let batch = index / (HEAD_DIM * params.seqLen * KV_HEADS);

  let position = past_sequence_length(batch) + t;
  let dst = ((batch * KV_HEADS + kv_head) * MAX_CACHE_SEQ + position) * HEAD_DIM + d;
{% if packedQkv %}
  let row = (batch * params.seqLen + t) * Q_STRIDE;
  let k_base = row + PACKED_K_OFFSET + kv_head * HEAD_DIM;
  let v_base = row + PACKED_V_OFFSET + kv_head * HEAD_DIM;
{% else %}
  let k_base = ((batch * params.seqLen + t) * KV_HIDDEN) + kv_head * HEAD_DIM;
  let v_base = k_base;
{% endif %}

{% if usesRotary %}
  // The key is rotated at its own absolute position before it enters the cache, so the
  // cache holds post-rotary keys and the attention pass never rotates again.
  if (d < ROTARY_DIM) {
    let slot = rotary_slot(d);
    let cs = f32(cos_cache[position * ROTARY_HALF + slot]);
    let sn = f32(sin_cache[position * ROTARY_HALF + slot]);
    let own = f32({{ kvSource }}[k_base + d]);
    let partner = f32({{ kvSource }}[k_base + rotary_partner(d)]);
    present_key[dst] = {{ scalar }}(rotary_value(own, partner, cs, sn, rotary_is_first(d)));
  } else {
    present_key[dst] = {{ scalar }}({{ kvSource }}[k_base + d]);
  }
{% else %}
  present_key[dst] = {{ scalar }}({{ kvSource }}[k_base + d]);
{% endif %}
  present_value[dst] = {{ scalar }}({{ vSource }}[v_base + d]);
}