// Row scatter for encoder row-packing: packed activations [T, N] → padded
// [B·S, N], where T = Σ lens[b] and padded row (b, m < lens[b]) comes from
// packed row starts[b] + m. Padding rows are left untouched — the destination
// arena buffer is zero-initialized by WebGPU, so they read as zeros
// downstream (decoder cross-attention masks j ≥ len and never reads them).
//
// One workgroup per padded row, vec4 element copies (N % 4 == 0 enforced by
// the dispatch). Early returns are uniform: wid-derived plus read-only
// storage loads at workgroup-uniform indices; there are no barriers.
//
// Template placeholders (buildShader in pipelines.js):
//   ENABLE_F16  f16 enable directive when T is f16, else empty
//   T           storage type of X/Y (f16|f32)
//   WG          workgroup size
{{ENABLE_IMMEDIATE}}
{{ENABLE_F16}}

struct Params { B: u32, S: u32, N4: u32, _pad: u32 }

{{PARAM_BINDING}}var<{{PARAM_ADDRESS}}> params: Params;
@group(0) @binding(1) var<storage, read> starts: array<u32>; // packed row offsets [B]
@group(0) @binding(2) var<storage, read> lens: array<u32>;   // sequence lengths [B]
@group(0) @binding(3) var<storage, read> X: array<vec4<{{T}}>>;
@group(0) @binding(4) var<storage, read_write> Y: array<vec4<{{T}}>>;

@compute @workgroup_size({{WG}})
fn main(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {
  if (wid.x >= params.B * params.S) { return; }
  let b = wid.x / params.S;
  let m = wid.x % params.S;
  if (m >= lens[b]) { return; }
  let src = (starts[b] + m) * params.N4;
  let dst = wid.x * params.N4;
  for (var i = lid.x; i < params.N4; i = i + {{WG}}u) {
    Y[dst + i] = X[src + i];
  }
}