Function createCentroidReduceShader

e2e/webgpu-render.spec.ts:438–485

()

`;

      const createCentroidReduceShader = () => `
struct Params {
  inputCount: u32,
}

@group(0) @binding(0) var<storage, read> inputPartials: array<vec4<u32>>;
@group(0) @binding(1) var<storage, read_write> outputPartials: array<vec4<u32>>;
@group(0) @binding(2) var<uniform> params: Params;

var<workgroup> partialSums: array<vec4<u32>, ${centroidWorkgroupSize}>;

fn addPartial(a: vec4<u32>, b: vec4<u32>) -> vec4<u32> {
  return vec4<u32>(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);
}

@compute @workgroup_size(${centroidWorkgroupSize})
fn main(
  @builtin(global_invocation_id) globalId: vec3<u32>,
  @builtin(local_invocation_id) localId: vec3<u32>,
  @builtin(workgroup_id) workgroupId: vec3<u32>
) {
  let index = globalId.x;
  var partial = vec4<u32>(0u, 0u, 0u, 0u);
  if (index < params.inputCount) {
    partial = inputPartials[index];
  }

  partialSums[localId.x] = partial;
  workgroupBarrier();

  var stride = ${centroidWorkgroupSize / 2}u;
  loop {
    if (localId.x < stride) {
      partialSums[localId.x] = addPartial(partialSums[localId.x], partialSums[localId.x + stride]);
    }
    workgroupBarrier();

    if (stride == 1u) {
      break;
    }
    stride = stride / 2u;
  }

  if (localId.x == 0u) {
    outputPartials[workgroupId.x] = partialSums[0];
  }
}
`;

      const centroidComparePipeline = device.createComputePipeline({
        layout: 'auto',