         atomicLoad(&histogram[ndx][0]),
          atomicLoad(&histogram[ndx][1]),
          atomicLoad(&histogram[ndx][2]),
          atomicLoad(&histogram[ndx][3]),
        );
      }
    `,
  });

  const chunkSumModule = device.createShaderModule({
    label: 'chunk sum shader',
    code: `
      @group(0) @binding(0) var<storage, read> histogramChunks: array<array<vec4u, 256>>;
      @group(0) @binding(1) var<storage, read_write> histogram: array<vec4u, 256>;

      @compute @workgroup_size(256, 1, 1) 