export default "// Row scatter for encoder row-packing: packed activations [T, N] → padded\n// [B·S, N], where T = Σ lens[b] and padded row (b, m < lens[b]) comes from\n// packed row starts[b] + m. Padding rows are left untouched — the destination\n// arena buffer is zero-initialized by WebGPU, so they read as zeros\n// downstream (decoder cross-attention masks j ≥ len and never reads them).\n//\n// One workgroup per padded row, vec4 element copies (N % 4 == 0 enforced by\n// the dispatch). Early returns are uniform: wid-derived plus read-only\n// storage loads at workgroup-uniform indices; there are no barriers.\n//\n// Template placeholders (buildShader in pipelines.js):\n//   ENABLE_F16  f16 enable directive when T is f16, else empty\n//   T           storage type of X/Y (f16|f32)\n//   WG          workgroup size\n{{ENABLE_IMMEDIATE}}\n{{ENABLE_F16}}\n\nstruct Params { B: u32, S: u32, N4: u32, _pad: u32 }\n\n{{PARAM_BINDING}}var<{{PARAM_ADDRESS}}> params: Params;\n@group(0) @binding(1) var<storage, read> starts: array<u32>; // packed row offsets [B]\n@group(0) @binding(2) var<storage, read> lens: array<u32>;   // sequence lengths [B]\n@group(0) @binding(3) var<storage, read> X: array<vec4<{{T}}>>;\n@group(0) @binding(4) var<storage, read_write> Y: array<vec4<{{T}}>>;\n\n@compute @workgroup_size({{WG}})\nfn main(@builtin(workgroup_id) wid: vec3<u32>, @builtin(local_invocation_id) lid: vec3<u32>) {\n  if (wid.x >= params.B * params.S) { return; }\n  let b = wid.x / params.S;\n  let m = wid.x % params.S;\n  if (m >= lens[b]) { return; }\n  let src = (starts[b] + m) * params.N4;\n  let dst = wid.x * params.N4;\n  for (var i = lid.x; i < params.N4; i = i + {{WG}}u) {\n    Y[dst + i] = X[src + i];\n  }\n}\n";