| {{ 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 %} |
| } |
| |
