{{ env.wgsl.resourceDeclarations }}

// Same-shape vec4 bitwise binary (and/or/xor): 4 contiguous elements per lane
// (128-bit loads/stores). uint8 storage uses one u32 slot per element, so
// the result is masked to the low byte per lane. Same semantics as the scalar
// broadcast kernel when A, B, C share a shape.
{% set vec4PerThread = vec4PerThread %}
{% if vec4PerThread > 1 %}
const ITEMS: u32 = {{ vec4PerThread }}u;
{% endif %}

@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
  // per-axis dispatch fold width (the dispatch caps x and spills the rest into y).
{% if vec4PerThread > 1 %}
  // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
  // access consecutive words on every step, while each lane can keep several
  // independent loads in flight.
  let tid = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
  let span = (params.count + ITEMS - 1u) / ITEMS;
  for (var j = 0u; j < ITEMS; j = j + 1u) {
    let i = tid + j * span;
    if (i >= params.count) {
      break;
    }
{% else %}
  let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
  if (i >= params.count) {
    return;
  }
{% endif %}

  var value = a[i] & b[i];
{% if cDtype == "uint8" %}
  value = value & vec4<u32>(0xffu);
{% endif %}
  c[i] = value;
{% if vec4PerThread > 1 %}
  }
{% endif %}
}