| {{ env.wgsl.resourceDeclarations }} |
| |
| {% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %} |
| {% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %} |
| // 2D-folded flat index: gid.y carries the high bits past the dispatch's |
| // per-axis workgroup fold width. |
| {% if bound == "" %} |
| let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};{% else %} |
| let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }}; |
| if ({{ name }} >= {{ bound }}) { |
| return; |
| }{% endif %}{% endmacro %} |
| {% if vec4PerThread > 1 %} |
| const ITEMS: u32 = {{ vec4PerThread }}u; |
| |
| {% endif %} |
| @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }}) |
| fn main(@builtin(global_invocation_id) gid: vec3<u32>) { |
| {% 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. |
| {{ flat_index_2d(tunables.WORKGROUP_SIZE, "tid", "") }} |
| let span = (params.count + ITEMS - 1u) / ITEMS; |
| // Lanes at or past span exist only because the dispatch rounds up to whole |
| // workgroups. Lane span + k would start on lane k's second group and rewrite |
| // up to ITEMS - 1 groups another lane already stored. |
| if (tid >= span) { |
| return; |
| } |
| for (var j = 0u; j < ITEMS; j = j + 1u) { |
| let i = tid + j * span; |
| if (i >= params.count) { |
| break; |
| } |
| {% else %} |
| {{ flat_index_2d(tunables.WORKGROUP_SIZE) }} |
| {% endif %} |
| {% set scalarOperand = scalarOperand if scalarOperand is defined else "" %} |
| {% if scalarOperand == "a" %} |
| // One-element operand: read once and splat across the vector. |
| let av = {{ vectorScalar }}(a[0]); |
| {% else %} |
| let av = a[i]; |
| {% endif %} |
| {% if scalarOperand == "b" %} |
| let bv = {{ vectorScalar }}(b[0]); |
| {% else %} |
| let bv = b[i]; |
| {% endif %} |
| {% if scalar == "i32" or scalar == "u32" %} |
| let r = av + bv; |
| // Narrow integer operations wrap modulo the logical dtype width; int8/uint8 |
| // use i32/u32 storage. |
| {% if cDtype == "int8" %} |
| c[i] = (r << vec4<u32>(24u)) >> vec4<u32>(24u); |
| {% elif cDtype == "uint8" %} |
| c[i] = r & vec4<u32>(0xFFu); |
| {% else %} |
| c[i] = r; |
| {% endif %} |
| {% elif scalar == "f16" %} |
| c[i] = vec4<f16>(vec4<f32>(av) + vec4<f32>(bv)); |
| {% else %} |
| c[i] = av + bv; |
| {% endif %} |
| {% if vec4PerThread > 1 %} |
| } |
| {% endif %} |
| } |
| |
