Adding a more complex compute pass memory sync test
diff --git a/src/webgpu/api/operation/memory_sync/buffer/single_buffer.spec.ts b/src/webgpu/api/operation/memory_sync/buffer/single_buffer.spec.ts
index 817d646..0936ba9 100644
--- a/src/webgpu/api/operation/memory_sync/buffer/single_buffer.spec.ts
+++ b/src/webgpu/api/operation/memory_sync/buffer/single_buffer.spec.ts
@@ -14,6 +14,7 @@
`;
import { makeTestGroup } from '../../../../../common/framework/test_group.js';
+import { checkElementsEqual } from '../../../../util/check_contents.js';
import {
kOperationBoundaries,
kBoundaryInfo,
@@ -255,3 +256,72 @@
t.device.queue.submit([encoder.finish()]);
t.verifyData(buffer, 2);
});
+
+g.test('multiple_dispatches_in_the_same_compute_pass')
+ .desc(
+ `Test multiple write-after-write operations with same compute pass. The first write will write
+ index * 2 into a storage buffer. The second write will write value + 2 into the same buffer in
+ the same pass. The third will write value / 2 into the same buffer in the same pass. The last
+ will write value - 1 into the same buffer in the same pass. Expected data in buffer is index.`
+ )
+ .fn(async t => {
+ const kBufferElementCount = 512;
+ const kWorkgroupSize = 64;
+
+ const buffer = t.trackForCleanup(
+ t.device.createBuffer({
+ size: Uint32Array.BYTES_PER_ELEMENT * kBufferElementCount,
+ usage:
+ GPUBufferUsage.COPY_SRC |
+ GPUBufferUsage.COPY_DST |
+ GPUBufferUsage.STORAGE,
+ })
+ );
+
+ const createBufferOperationComputePipeline = (operation: string) => {
+ const module = t.device.createShaderModule({
+ code: `
+ @group(0) @binding(0) var<storage, read_write> data : array<u32>;
+ @compute @workgroup_size(${kWorkgroupSize})
+ fn main(@builtin(global_invocation_id) globalIndex: vec3u) {
+ let i = globalIndex.x;
+ data[i] = ${operation};
+ }
+ `
+ });
+ return t.device.createComputePipeline({
+ layout: 'auto',
+ compute: {
+ module,
+ entryPoint: 'main',
+ },
+ });
+ }
+
+
+ const pipelines = [];
+ pipelines.push(createBufferOperationComputePipeline('i * 2'));
+ pipelines.push(createBufferOperationComputePipeline('data[i] + 2'));
+ pipelines.push(createBufferOperationComputePipeline('data[i] / 2'));
+ pipelines.push(createBufferOperationComputePipeline('data[i] - 1'));
+
+ const encoder = t.device.createCommandEncoder();
+ const pass = encoder.beginComputePass();
+
+ for (const pipeline of pipelines) {
+ const bindGroup = t.createBindGroup(pipeline, buffer);
+ pass.setPipeline(pipeline);
+ pass.setBindGroup(0, bindGroup);
+ pass.dispatchWorkgroups(kBufferElementCount / kWorkgroupSize);
+ }
+
+ pass.end();
+ t.device.queue.submit([encoder.finish()]);
+
+ const expected = new Uint32Array(kBufferElementCount);
+ for (let i = 0; i < kBufferElementCount; ++i) {
+ expected[i] = i;
+ }
+
+ t.expectGPUBufferValuesEqual(buffer, expected);
+ });
