1 files changed, 938 insertions, 0 deletions
diff --git a/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/memory_sync/buffer/buffer_sync_test.ts b/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/memory_sync/buffer/buffer_sync_test.ts
new file mode 100644
index 0000000000..e18dc59abf
--- /dev/null
+++ b/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/memory_sync/buffer/buffer_sync_test.ts
@@ -0,0 +1,938 @@
+import { assert, unreachable } from '../../../../../common/util/util.js';
+import { GPUTest } from '../../../../gpu_test.js';
+import { checkElementsEqualEither } from '../../../../util/check_contents.js';
+import { OperationContext, OperationContextHelper } from '../operation_context_helper.js';
+
+export const kAllWriteOps = ['storage', 'b2b-copy', 't2b-copy', 'write-buffer'] as const;
+
+export const kAllReadOps = [
+  'input-vertex',
+  'input-index',
+  'input-indirect',
+  'input-indirect-index',
+  'input-indirect-dispatch',
+
+  'constant-uniform',
+
+  'storage-read',
+
+  'b2b-copy',
+  'b2t-copy',
+] as const;
+
+export type ReadOp = typeof kAllReadOps[number];
+export type WriteOp = typeof kAllWriteOps[number];
+
+export type Op = ReadOp | WriteOp;
+
+interface OpInfo {
+  readonly contexts: OperationContext[];
+}
+
+const kOpInfo: {
+  readonly [k in Op]: OpInfo;
+} = /* prettier-ignore */ {
+  'write-buffer': {
+    contexts: [ 'queue' ],
+  },
+  'b2t-copy': {
+    contexts: [ 'command-encoder' ],
+  },
+  'b2b-copy': {
+    contexts: [ 'command-encoder' ],
+  },
+  't2b-copy': {
+    contexts: [ 'command-encoder' ],
+  },
+  'storage': {
+    contexts: [ 'compute-pass-encoder', 'render-pass-encoder', 'render-bundle-encoder' ],
+  },
+  'storage-read': {
+    contexts: [ 'compute-pass-encoder', 'render-pass-encoder', 'render-bundle-encoder' ],
+  },
+  'input-vertex': {
+    contexts: [ 'render-pass-encoder', 'render-bundle-encoder' ],
+  },
+  'input-index': {
+    contexts: [ 'render-pass-encoder', 'render-bundle-encoder' ],
+  },
+  'input-indirect': {
+    contexts: [ 'render-pass-encoder', 'render-bundle-encoder' ],
+  },
+  'input-indirect-index': {
+    contexts: [ 'render-pass-encoder', 'render-bundle-encoder' ],
+  },
+  'input-indirect-dispatch': {
+    contexts: [ 'compute-pass-encoder' ],
+  },
+  'constant-uniform': {
+    contexts: [ 'render-pass-encoder', 'render-bundle-encoder' ],
+  },
+};
+
+export function checkOpsValidForContext(
+  ops: [Op, Op],
+  context: [OperationContext, OperationContext]
+) {
+  const valid =
+    kOpInfo[ops[0]].contexts.includes(context[0]) && kOpInfo[ops[1]].contexts.includes(context[1]);
+  if (!valid) return false;
+
+  if (
+    context[0] === 'render-bundle-encoder' ||
+    context[0] === 'render-pass-encoder' ||
+    context[1] === 'render-bundle-encoder' ||
+    context[1] === 'render-pass-encoder'
+  ) {
+    // In a render pass, it is invalid to use a resource as both writable and another usage.
+    // Also, for storage+storage usage, the application is opting into racy behavior.
+    // The storage+storage case is also skipped as the results cannot be reliably tested.
+    const checkImpl = (op1: Op, op2: Op) => {
+      switch (op1) {
+        case 'storage':
+          switch (op2) {
+            case 'storage':
+            case 'storage-read':
+            case 'input-vertex':
+            case 'input-index':
+            case 'input-indirect':
+            case 'input-indirect-index':
+            case 'constant-uniform':
+              // Write+other, or racy.
+              return false;
+            case 'b2t-copy':
+            case 't2b-copy':
+            case 'b2b-copy':
+            case 'write-buffer':
+              // These don't occur in a render pass.
+              return true;
+          }
+          break;
+        case 'input-vertex':
+        case 'input-index':
+        case 'input-indirect':
+        case 'input-indirect-index':
+        case 'constant-uniform':
+        case 'b2t-copy':
+        case 't2b-copy':
+        case 'b2b-copy':
+        case 'write-buffer':
+          // These are not write usages, or don't occur in a render pass.
+          break;
+      }
+      return true;
+    };
+    return checkImpl(ops[0], ops[1]) && checkImpl(ops[1], ops[0]);
+  }
+  return true;
+}
+
+const kDummyVertexShader = `
+@vertex fn vert_main() -> @builtin(position) vec4<f32> {
+  return vec4<f32>(0.5, 0.5, 0.0, 1.0);
+}
+`;
+
+// Note: If it would be useful to have any of these helpers be separate from the fixture,
+// they can be refactored into standalone functions.
+export class BufferSyncTest extends GPUTest {
+  // Vertex and index buffers used in read render pass
+  vertexBuffer?: GPUBuffer;
+  indexBuffer?: GPUBuffer;
+
+  // Temp buffer and texture with values for buffer/texture copy write op
+  // There can be at most 2 write op
+  tmpValueBuffers: (GPUBuffer | undefined)[] = [undefined, undefined];
+  tmpValueTextures: (GPUTexture | undefined)[] = [undefined, undefined];
+
+  // These intermediate buffers/textures are created before any read/write op
+  // to avoid extra memory synchronization between ops introduced by await on buffer/texture creations.
+  // Create extra buffers/textures needed by write operation
+  async createIntermediateBuffersAndTexturesForWriteOp(
+    writeOp: WriteOp,
+    slot: number,
+    value: number
+  ) {
+    switch (writeOp) {
+      case 'b2b-copy':
+        this.tmpValueBuffers[slot] = await this.createBufferWithValue(value);
+        break;
+      case 't2b-copy':
+        this.tmpValueTextures[slot] = await this.createTextureWithValue(value);
+        break;
+      default:
+        break;
+    }
+  }
+
+  // Create extra buffers/textures needed by read operation
+  async createBuffersForReadOp(readOp: ReadOp, srcValue: number, opValue: number) {
+    // This helps create values that will be written into dst buffer by the readop
+    switch (readOp) {
+      case 'input-index':
+        // The index buffer will be the src buffer of the read op.
+        // The src value for readOp will be 0
+        // If the index buffer value is 0, the src value is written into the dst buffer.
+        // If the index buffer value is 1, the op value is written into the dst buffer.
+        this.vertexBuffer = await this.createBufferWithValues([srcValue, opValue]);
+        break;
+      case 'input-indirect':
+        // The indirect buffer for the draw cmd will be the src buffer of the read op.
+        // If the first value in the indirect buffer is 1, then the op value in vertex buffer will be written into dst buffer.
+        // If the first value in indirect buffer is 0, then nothing will be write into dst buffer.
+        this.vertexBuffer = await this.createBufferWithValues([opValue]);
+        break;
+      case 'input-indirect-index':
+        // The indirect buffer for draw indexed cmd will be the src buffer of the read op.
+        // If the first value in the indirect buffer is 1, then the opValue in vertex buffer will be written into dst buffer.
+        // If the first value in indirect buffer is 0, then nothing will be write into dst buffer.
+        this.vertexBuffer = await this.createBufferWithValues([opValue]);
+        this.indexBuffer = await this.createBufferWithValues([0]);
+        break;
+      default:
+        break;
+    }
+
+    let srcBuffer: GPUBuffer;
+    switch (readOp) {
+      case 'input-indirect':
+        // vertexCount = {0, 1}
+        // instanceCount = 1
+        // firstVertex = 0
+        // firstInstance = 0
+        srcBuffer = await this.createBufferWithValues([srcValue, 1, 0, 0]);
+        break;
+      case 'input-indirect-index':
+        // indexCount = {0, 1}
+        // instanceCount = 1
+        // firstIndex = 0
+        // baseVertex = 0
+        // firstInstance = 0
+        srcBuffer = await this.createBufferWithValues([srcValue, 1, 0, 0, 0]);
+        break;
+      case 'input-indirect-dispatch':
+        // workgroupCountX = {0, 1}
+        // workgroupCountY = 1
+        // workgroupCountZ = 1
+        srcBuffer = await this.createBufferWithValues([srcValue, 1, 1]);
+        break;
+      default:
+        srcBuffer = await this.createBufferWithValue(srcValue);
+        break;
+    }
+
+    const dstBuffer = this.trackForCleanup(
+      this.device.createBuffer({
+        size: Uint32Array.BYTES_PER_ELEMENT,
+        usage:
+          GPUBufferUsage.COPY_SRC |
+          GPUBufferUsage.COPY_DST |
+          GPUBufferUsage.STORAGE |
+          GPUBufferUsage.VERTEX |
+          GPUBufferUsage.INDEX |
+          GPUBufferUsage.INDIRECT |
+          GPUBufferUsage.UNIFORM,
+      })
+    );
+
+    return { srcBuffer, dstBuffer };
+  }
+
+  // Create a buffer with 1 uint32 element, and initialize it to a specified value.
+  async createBufferWithValue(initValue: number): Promise<GPUBuffer> {
+    const buffer = this.trackForCleanup(
+      this.device.createBuffer({
+        mappedAtCreation: true,
+        size: Uint32Array.BYTES_PER_ELEMENT,
+        usage:
+          GPUBufferUsage.COPY_SRC |
+          GPUBufferUsage.COPY_DST |
+          GPUBufferUsage.STORAGE |
+          GPUBufferUsage.VERTEX |
+          GPUBufferUsage.INDEX |
+          GPUBufferUsage.INDIRECT |
+          GPUBufferUsage.UNIFORM,
+      })
+    );
+    new Uint32Array(buffer.getMappedRange()).fill(initValue);
+    buffer.unmap();
+    await this.queue.onSubmittedWorkDone();
+    return buffer;
+  }
+
+  // Create a buffer, and initialize it to the specified values.
+  async createBufferWithValues(initValues: number[]): Promise<GPUBuffer> {
+    const buffer = this.trackForCleanup(
+      this.device.createBuffer({
+        mappedAtCreation: true,
+        size: Uint32Array.BYTES_PER_ELEMENT * initValues.length,
+        usage:
+          GPUBufferUsage.COPY_SRC |
+          GPUBufferUsage.COPY_DST |
+          GPUBufferUsage.STORAGE |
+          GPUBufferUsage.VERTEX |
+          GPUBufferUsage.INDEX |
+          GPUBufferUsage.INDIRECT |
+          GPUBufferUsage.UNIFORM,
+      })
+    );
+    const bufferView = new Uint32Array(buffer.getMappedRange());
+    bufferView.set(initValues);
+    buffer.unmap();
+    await this.queue.onSubmittedWorkDone();
+    return buffer;
+  }
+
+  // Create a 1x1 texture, and initialize it to a specified value for all elements.
+  async createTextureWithValue(initValue: number): Promise<GPUTexture> {
+    const data = new Uint32Array(1).fill(initValue);
+    const texture = this.trackForCleanup(
+      this.device.createTexture({
+        size: { width: 1, height: 1, depthOrArrayLayers: 1 },
+        format: 'r32uint',
+        usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.COPY_DST,
+      })
+    );
+    this.device.queue.writeTexture(
+      { texture, mipLevel: 0, origin: { x: 0, y: 0, z: 0 } },
+      data,
+      { offset: 0, bytesPerRow: 256, rowsPerImage: 1 },
+      { width: 1, height: 1, depthOrArrayLayers: 1 }
+    );
+    await this.queue.onSubmittedWorkDone();
+    return texture;
+  }
+
+  createBindGroup(
+    pipeline: GPURenderPipeline | GPUComputePipeline,
+    buffer: GPUBuffer
+  ): GPUBindGroup {
+    return this.device.createBindGroup({
+      layout: pipeline.getBindGroupLayout(0),
+      entries: [{ binding: 0, resource: { buffer } }],
+    });
+  }
+
+  // Create a compute pipeline and write given data into storage buffer.
+  createStorageWriteComputePipeline(value: number): GPUComputePipeline {
+    const wgslCompute = `
+      struct Data {
+        a : u32
+      };
+
+      @group(0) @binding(0) var<storage, read_write> data : Data;
+      @compute @workgroup_size(1) fn main() {
+        data.a = ${value}u;
+      }
+    `;
+
+    return this.device.createComputePipeline({
+      layout: 'auto',
+      compute: {
+        module: this.device.createShaderModule({
+          code: wgslCompute,
+        }),
+        entryPoint: 'main',
+      },
+    });
+  }
+
+  createTrivialRenderPipeline(wgslShaders: { vertex: string; fragment: string }) {
+    return this.device.createRenderPipeline({
+      layout: 'auto',
+      vertex: {
+        module: this.device.createShaderModule({
+          code: wgslShaders.vertex,
+        }),
+        entryPoint: 'vert_main',
+      },
+      fragment: {
+        module: this.device.createShaderModule({
+          code: wgslShaders.fragment,
+        }),
+        entryPoint: 'frag_main',
+        targets: [{ format: 'rgba8unorm' }],
+      },
+      primitive: { topology: 'point-list' },
+    });
+  }
+
+  // Create a render pipeline and write given data into storage buffer at fragment stage.
+  createStorageWriteRenderPipeline(value: number): GPURenderPipeline {
+    const wgslShaders = {
+      vertex: kDummyVertexShader,
+      fragment: `
+      struct Data {
+        a : u32
+      };
+
+      @group(0) @binding(0) var<storage, read_write> data : Data;
+      @fragment fn frag_main() -> @location(0) vec4<f32> {
+        data.a = ${value}u;
+        return vec4<f32>();  // result does't matter
+      }
+    `,
+    };
+
+    return this.createTrivialRenderPipeline(wgslShaders);
+  }
+
+  beginSimpleRenderPass(encoder: GPUCommandEncoder): GPURenderPassEncoder {
+    const view = this.trackForCleanup(
+      this.device.createTexture({
+        size: { width: 1, height: 1, depthOrArrayLayers: 1 },
+        format: 'rgba8unorm',
+        usage: GPUTextureUsage.RENDER_ATTACHMENT,
+      })
+    ).createView();
+    return encoder.beginRenderPass({
+      colorAttachments: [
+        {
+          view,
+          clearValue: { r: 0.0, g: 1.0, b: 0.0, a: 1.0 },
+          loadOp: 'clear',
+          storeOp: 'store',
+        },
+      ],
+    });
+  }
+
+  // Write buffer via draw call in render pass. Use bundle if needed.
+  encodeWriteAsStorageBufferInRenderPass(
+    renderer: GPURenderPassEncoder | GPURenderBundleEncoder,
+    buffer: GPUBuffer,
+    value: number
+  ) {
+    const pipeline = this.createStorageWriteRenderPipeline(value);
+    const bindGroup = this.createBindGroup(pipeline, buffer);
+
+    renderer.setBindGroup(0, bindGroup);
+    renderer.setPipeline(pipeline);
+    renderer.draw(1, 1, 0, 0);
+  }
+
+  // Write buffer via dispatch call in compute pass.
+  encodeWriteAsStorageBufferInComputePass(
+    pass: GPUComputePassEncoder,
+    buffer: GPUBuffer,
+    value: number
+  ) {
+    const pipeline = this.createStorageWriteComputePipeline(value);
+    const bindGroup = this.createBindGroup(pipeline, buffer);
+    pass.setPipeline(pipeline);
+    pass.setBindGroup(0, bindGroup);
+    pass.dispatchWorkgroups(1);
+  }
+
+  // Write buffer via BufferToBuffer copy.
+  encodeWriteByB2BCopy(encoder: GPUCommandEncoder, buffer: GPUBuffer, slot: number) {
+    const tmpBuffer = this.tmpValueBuffers[slot];
+    assert(tmpBuffer !== undefined);
+    // The write operation via b2b copy is just encoded into command encoder, it doesn't write immediately.
+    encoder.copyBufferToBuffer(tmpBuffer, 0, buffer, 0, Uint32Array.BYTES_PER_ELEMENT);
+  }
+
+  // Write buffer via TextureToBuffer copy.
+  encodeWriteByT2BCopy(encoder: GPUCommandEncoder, buffer: GPUBuffer, slot: number) {
+    const tmpTexture = this.tmpValueTextures[slot];
+    assert(tmpTexture !== undefined);
+    // The write operation via t2b copy is just encoded into command encoder, it doesn't write immediately.
+    encoder.copyTextureToBuffer(
+      { texture: tmpTexture, mipLevel: 0, origin: { x: 0, y: 0, z: 0 } },
+      { buffer, bytesPerRow: 256 },
+      { width: 1, height: 1, depthOrArrayLayers: 1 }
+    );
+  }
+
+  // Write buffer via writeBuffer API on queue
+  writeByWriteBuffer(buffer: GPUBuffer, value: number) {
+    const data = new Uint32Array(1).fill(value);
+    this.device.queue.writeBuffer(buffer, 0, data);
+  }
+
+  // Issue write operation via render pass, compute pass, copy, etc.
+  encodeWriteOp(
+    helper: OperationContextHelper,
+    operation: WriteOp,
+    context: OperationContext,
+    buffer: GPUBuffer,
+    writeOpSlot: number,
+    value: number
+  ) {
+    helper.ensureContext(context);
+
+    switch (operation) {
+      case 'write-buffer':
+        this.writeByWriteBuffer(buffer, value);
+        break;
+      case 'storage':
+        switch (context) {
+          case 'render-pass-encoder':
+            assert(helper.renderPassEncoder !== undefined);
+            this.encodeWriteAsStorageBufferInRenderPass(helper.renderPassEncoder, buffer, value);
+            break;
+          case 'render-bundle-encoder':
+            assert(helper.renderBundleEncoder !== undefined);
+            this.encodeWriteAsStorageBufferInRenderPass(helper.renderBundleEncoder, buffer, value);
+            break;
+          case 'compute-pass-encoder':
+            assert(helper.computePassEncoder !== undefined);
+            this.encodeWriteAsStorageBufferInComputePass(helper.computePassEncoder, buffer, value);
+            break;
+          default:
+            unreachable();
+        }
+        break;
+      case 'b2b-copy':
+        assert(helper.commandEncoder !== undefined);
+        this.encodeWriteByB2BCopy(helper.commandEncoder, buffer, writeOpSlot);
+        break;
+      case 't2b-copy':
+        assert(helper.commandEncoder !== undefined);
+        this.encodeWriteByT2BCopy(helper.commandEncoder, buffer, writeOpSlot);
+        break;
+      default:
+        unreachable();
+    }
+  }
+
+  // Create a compute pipeline: read from src buffer and write it into the storage buffer.
+  createStorageReadComputePipeline(): GPUComputePipeline {
+    const wgslCompute = `
+      struct Data {
+        a : u32
+      };
+
+      @group(0) @binding(0) var<storage, read> srcData : Data;
+      @group(0) @binding(1) var<storage, read_write> dstData : Data;
+
+      @compute @workgroup_size(1) fn main() {
+        dstData.a = srcData.a;
+      }
+    `;
+
+    return this.device.createComputePipeline({
+      layout: 'auto',
+      compute: {
+        module: this.device.createShaderModule({
+          code: wgslCompute,
+        }),
+        entryPoint: 'main',
+      },
+    });
+  }
+
+  createBindGroupSrcDstBuffer(
+    pipeline: GPURenderPipeline | GPUComputePipeline,
+    srcBuffer: GPUBuffer,
+    dstBuffer: GPUBuffer
+  ): GPUBindGroup {
+    return this.device.createBindGroup({
+      layout: pipeline.getBindGroupLayout(0),
+      entries: [
+        { binding: 0, resource: { buffer: srcBuffer } },
+        { binding: 1, resource: { buffer: dstBuffer } },
+      ],
+    });
+  }
+
+  // Create a render pipeline: read from vertex/index buffer and write it into the storage dst buffer at fragment stage.
+  createVertexReadRenderPipeline(): GPURenderPipeline {
+    const wgslShaders = {
+      vertex: `
+      struct VertexOutput {
+        @builtin(position) position : vec4<f32>,
+        @location(0) @interpolate(flat) data : u32,
+      };
+
+      @vertex fn vert_main(@location(0) input: u32) -> VertexOutput {
+        var output : VertexOutput;
+        output.position = vec4<f32>(0.5, 0.5, 0.0, 1.0);
+        output.data = input;
+        return output;
+      }
+      `,
+      fragment: `
+      struct Data {
+        a : u32
+      };
+
+      @group(0) @binding(0) var<storage, read_write> data : Data;
+
+      @fragment fn frag_main(@location(0) @interpolate(flat) input : u32) -> @location(0) vec4<f32> {
+        data.a = input;
+        return vec4<f32>();  // result does't matter
+      }
+      `,
+    };
+
+    return this.device.createRenderPipeline({
+      layout: 'auto',
+      vertex: {
+        module: this.device.createShaderModule({
+          code: wgslShaders.vertex,
+        }),
+        entryPoint: 'vert_main',
+        buffers: [
+          {
+            arrayStride: Uint32Array.BYTES_PER_ELEMENT,
+            attributes: [
+              {
+                shaderLocation: 0,
+                offset: 0,
+                format: 'uint32',
+              },
+            ],
+          },
+        ],
+      },
+      fragment: {
+        module: this.device.createShaderModule({
+          code: wgslShaders.fragment,
+        }),
+        entryPoint: 'frag_main',
+        targets: [{ format: 'rgba8unorm' }],
+      },
+      primitive: { topology: 'point-list' },
+    });
+  }
+
+  // Create a render pipeline: read from uniform buffer and write it into the storage dst buffer at fragment stage.
+  createUniformReadRenderPipeline(): GPURenderPipeline {
+    const wgslShaders = {
+      vertex: kDummyVertexShader,
+      fragment: `
+      struct Data {
+        a : u32
+      };
+
+      @group(0) @binding(0) var<uniform> constant: Data;
+      @group(0) @binding(1) var<storage, read_write> data : Data;
+
+      @fragment fn frag_main() -> @location(0) vec4<f32> {
+        data.a = constant.a;
+        return vec4<f32>();  // result does't matter
+      }
+      `,
+    };
+
+    return this.createTrivialRenderPipeline(wgslShaders);
+  }
+
+  // Create a render pipeline: read from storage src buffer and write it into the storage dst buffer at fragment stage.
+  createStorageReadRenderPipeline(): GPURenderPipeline {
+    const wgslShaders = {
+      vertex: kDummyVertexShader,
+      fragment: `
+        struct Data {
+          a : u32
+        };
+
+        @group(0) @binding(0) var<storage, read> srcData : Data;
+        @group(0) @binding(1) var<storage, read_write> dstData : Data;
+
+        @fragment fn frag_main() -> @location(0) vec4<f32> {
+          dstData.a = srcData.a;
+          return vec4<f32>();  // result does't matter
+        }
+      `,
+    };
+
+    return this.device.createRenderPipeline({
+      layout: 'auto',
+      vertex: {
+        module: this.device.createShaderModule({
+          code: wgslShaders.vertex,
+        }),
+        entryPoint: 'vert_main',
+      },
+      fragment: {
+        module: this.device.createShaderModule({
+          code: wgslShaders.fragment,
+        }),
+        entryPoint: 'frag_main',
+        targets: [{ format: 'rgba8unorm' }],
+      },
+      primitive: { topology: 'point-list' },
+    });
+  }
+
+  // Write buffer via dispatch call in compute pass.
+  encodeReadAsStorageBufferInComputePass(
+    pass: GPUComputePassEncoder,
+    srcBuffer: GPUBuffer,
+    dstBuffer: GPUBuffer
+  ) {
+    const pipeline = this.createStorageReadComputePipeline();
+    const bindGroup = this.createBindGroupSrcDstBuffer(pipeline, srcBuffer, dstBuffer);
+    pass.setPipeline(pipeline);
+    pass.setBindGroup(0, bindGroup);
+    pass.dispatchWorkgroups(1);
+  }
+
+  // Write buffer via dispatchWorkgroupsIndirect call in compute pass.
+  encodeReadAsIndirectBufferInComputePass(
+    pass: GPUComputePassEncoder,
+    srcBuffer: GPUBuffer,
+    dstBuffer: GPUBuffer,
+    value: number
+  ) {
+    const pipeline = this.createStorageWriteComputePipeline(value);
+    const bindGroup = this.createBindGroup(pipeline, dstBuffer);
+    pass.setPipeline(pipeline);
+    pass.setBindGroup(0, bindGroup);
+    pass.dispatchWorkgroupsIndirect(srcBuffer, 0);
+  }
+
+  // Read as vertex input and write buffer via draw call in render pass. Use bundle if needed.
+  encodeReadAsVertexBufferInRenderPass(
+    renderer: GPURenderPassEncoder | GPURenderBundleEncoder,
+    srcBuffer: GPUBuffer,
+    dstBuffer: GPUBuffer
+  ) {
+    const pipeline = this.createVertexReadRenderPipeline();
+    const bindGroup = this.device.createBindGroup({
+      layout: pipeline.getBindGroupLayout(0),
+      entries: [{ binding: 0, resource: { buffer: dstBuffer } }],
+    });
+
+    renderer.setBindGroup(0, bindGroup);
+    renderer.setPipeline(pipeline);
+    renderer.setVertexBuffer(0, srcBuffer);
+    renderer.draw(1);
+  }
+
+  // Read as index input and write buffer via draw call in render pass. Use bundle if needed.
+  encodeReadAsIndexBufferInRenderPass(
+    renderer: GPURenderPassEncoder | GPURenderBundleEncoder,
+    srcBuffer: GPUBuffer,
+    dstBuffer: GPUBuffer,
+    vertexBuffer: GPUBuffer
+  ) {
+    const pipeline = this.createVertexReadRenderPipeline();
+    const bindGroup = this.device.createBindGroup({
+      layout: pipeline.getBindGroupLayout(0),
+      entries: [{ binding: 0, resource: { buffer: dstBuffer } }],
+    });
+
+    renderer.setBindGroup(0, bindGroup);
+    renderer.setPipeline(pipeline);
+    renderer.setVertexBuffer(0, vertexBuffer);
+    renderer.setIndexBuffer(srcBuffer, 'uint32');
+    renderer.drawIndexed(1);
+  }
+
+  // Read as indirect input and write buffer via draw call in render pass. Use bundle if needed.
+  encodeReadAsIndirectBufferInRenderPass(
+    renderer: GPURenderPassEncoder | GPURenderBundleEncoder,
+    srcBuffer: GPUBuffer,
+    dstBuffer: GPUBuffer,
+    vertexBuffer: GPUBuffer
+  ) {
+    const pipeline = this.createVertexReadRenderPipeline();
+    const bindGroup = this.device.createBindGroup({
+      layout: pipeline.getBindGroupLayout(0),
+      entries: [{ binding: 0, resource: { buffer: dstBuffer } }],
+    });
+
+    renderer.setBindGroup(0, bindGroup);
+    renderer.setPipeline(pipeline);
+    renderer.setVertexBuffer(0, vertexBuffer);
+    renderer.drawIndirect(srcBuffer, 0);
+  }
+
+  // Read as indexed indirect input and write buffer via draw call in render pass. Use bundle if needed.
+  encodeReadAsIndexedIndirectBufferInRenderPass(
+    renderer: GPURenderPassEncoder | GPURenderBundleEncoder,
+    srcBuffer: GPUBuffer,
+    dstBuffer: GPUBuffer,
+    vertexBuffer: GPUBuffer,
+    indexBuffer: GPUBuffer
+  ) {
+    const pipeline = this.createVertexReadRenderPipeline();
+    const bindGroup = this.device.createBindGroup({
+      layout: pipeline.getBindGroupLayout(0),
+      entries: [{ binding: 0, resource: { buffer: dstBuffer } }],
+    });
+
+    renderer.setBindGroup(0, bindGroup);
+    renderer.setPipeline(pipeline);
+    renderer.setVertexBuffer(0, vertexBuffer);
+    renderer.setIndexBuffer(indexBuffer, 'uint32');
+    renderer.drawIndexedIndirect(srcBuffer, 0);
+  }
+
+  // Read as uniform buffer and write buffer via draw call in render pass. Use bundle if needed.
+  encodeReadAsUniformBufferInRenderPass(
+    renderer: GPURenderPassEncoder | GPURenderBundleEncoder,
+    srcBuffer: GPUBuffer,
+    dstBuffer: GPUBuffer
+  ) {
+    const pipeline = this.createUniformReadRenderPipeline();
+    const bindGroup = this.device.createBindGroup({
+      layout: pipeline.getBindGroupLayout(0),
+      entries: [
+        { binding: 0, resource: { buffer: srcBuffer } },
+        { binding: 1, resource: { buffer: dstBuffer } },
+      ],
+    });
+
+    renderer.setBindGroup(0, bindGroup);
+    renderer.setPipeline(pipeline);
+    renderer.draw(1);
+  }
+
+  // Read as storage buffer and write buffer via draw call in render pass. Use bundle if needed.
+  encodeReadAsStorageBufferInRenderPass(
+    renderer: GPURenderPassEncoder | GPURenderBundleEncoder,
+    srcBuffer: GPUBuffer,
+    dstBuffer: GPUBuffer
+  ) {
+    const pipeline = this.createStorageReadRenderPipeline();
+    const bindGroup = this.createBindGroupSrcDstBuffer(pipeline, srcBuffer, dstBuffer);
+
+    renderer.setBindGroup(0, bindGroup);
+    renderer.setPipeline(pipeline);
+    renderer.draw(1, 1, 0, 0);
+  }
+
+  // Read and write via BufferToBuffer copy.
+  encodeReadByB2BCopy(encoder: GPUCommandEncoder, srcBuffer: GPUBuffer, dstBuffer: GPUBuffer) {
+    // The b2b copy is just encoded into command encoder, it doesn't write immediately.
+    encoder.copyBufferToBuffer(srcBuffer, 0, dstBuffer, 0, Uint32Array.BYTES_PER_ELEMENT);
+  }
+
+  // Read and Write texture via BufferToTexture copy.
+  encodeReadByB2TCopy(encoder: GPUCommandEncoder, srcBuffer: GPUBuffer, dstBuffer: GPUBuffer) {
+    const tmpTexture = this.trackForCleanup(
+      this.device.createTexture({
+        size: { width: 1, height: 1, depthOrArrayLayers: 1 },
+        format: 'r32uint',
+        usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.COPY_DST,
+      })
+    );
+
+    // The b2t copy is just encoded into command encoder, it doesn't write immediately.
+    encoder.copyBufferToTexture(
+      { buffer: srcBuffer, bytesPerRow: 256 },
+      { texture: tmpTexture, mipLevel: 0, origin: { x: 0, y: 0, z: 0 } },
+      { width: 1, height: 1, depthOrArrayLayers: 1 }
+    );
+    // The t2b copy is just encoded into command encoder, it doesn't write immediately.
+    encoder.copyTextureToBuffer(
+      { texture: tmpTexture, mipLevel: 0, origin: { x: 0, y: 0, z: 0 } },
+      { buffer: dstBuffer, bytesPerRow: 256 },
+      { width: 1, height: 1, depthOrArrayLayers: 1 }
+    );
+  }
+
+  encodeReadOp(
+    helper: OperationContextHelper,
+    operation: ReadOp,
+    context: OperationContext,
+    srcBuffer: GPUBuffer,
+    dstBuffer: GPUBuffer
+  ) {
+    helper.ensureContext(context);
+
+    const renderer =
+      context === 'render-bundle-encoder' ? helper.renderBundleEncoder : helper.renderPassEncoder;
+    const computePass = context === 'compute-pass-encoder' ? helper.computePassEncoder : undefined;
+
+    switch (operation) {
+      case 'input-vertex':
+        // The srcBuffer is used as vertexBuffer.
+        // draw writes the same value in srcBuffer[0] to dstBuffer[0].
+        assert(renderer !== undefined);
+        this.encodeReadAsVertexBufferInRenderPass(renderer, srcBuffer, dstBuffer);
+        break;
+      case 'input-index':
+        // The srcBuffer is used as indexBuffer.
+        // With this vertexBuffer, drawIndexed writes the same value in srcBuffer[0] to dstBuffer[0].
+        assert(renderer !== undefined);
+        assert(this.vertexBuffer !== undefined);
+        this.encodeReadAsIndexBufferInRenderPass(renderer, srcBuffer, dstBuffer, this.vertexBuffer);
+        break;
+      case 'input-indirect':
+        // The srcBuffer is used as indirectBuffer for drawIndirect.
+        // srcBuffer[0] = 0 or 1 (vertexCount), which will decide the value written into dstBuffer to be either 0 or 1.
+        assert(renderer !== undefined);
+        assert(this.vertexBuffer !== undefined);
+        this.encodeReadAsIndirectBufferInRenderPass(
+          renderer,
+          srcBuffer,
+          dstBuffer,
+          this.vertexBuffer
+        );
+        break;
+      case 'input-indirect-index':
+        // The srcBuffer is used as indirectBuffer for drawIndexedIndirect.
+        // srcBuffer[0] = 0 or 1 (indexCount), which will decide the value written into dstBuffer to be either 0 or 1.
+        assert(renderer !== undefined);
+        assert(this.vertexBuffer !== undefined);
+        assert(this.indexBuffer !== undefined);
+        this.encodeReadAsIndexedIndirectBufferInRenderPass(
+          renderer,
+          srcBuffer,
+          dstBuffer,
+          this.vertexBuffer,
+          this.indexBuffer
+        );
+        break;
+      case 'input-indirect-dispatch':
+        // The srcBuffer is used as indirectBuffer for dispatch.
+        // srcBuffer[0] = 0 or 1 (workgroupCountX), which will decide the value written into dstBuffer to be either 0 or 1.
+        assert(computePass !== undefined);
+        this.encodeReadAsIndirectBufferInComputePass(computePass, srcBuffer, dstBuffer, 1);
+        break;
+      case 'constant-uniform':
+        // The srcBuffer is used as uniform buffer.
+        assert(renderer !== undefined);
+        this.encodeReadAsUniformBufferInRenderPass(renderer, srcBuffer, dstBuffer);
+        break;
+      case 'storage-read':
+        switch (context) {
+          case 'render-pass-encoder':
+          case 'render-bundle-encoder':
+            assert(renderer !== undefined);
+            this.encodeReadAsStorageBufferInRenderPass(renderer, srcBuffer, dstBuffer);
+            break;
+          case 'compute-pass-encoder':
+            assert(computePass !== undefined);
+            this.encodeReadAsStorageBufferInComputePass(computePass, srcBuffer, dstBuffer);
+            break;
+          default:
+            unreachable();
+        }
+        break;
+      case 'b2b-copy':
+        assert(helper.commandEncoder !== undefined);
+        this.encodeReadByB2BCopy(helper.commandEncoder, srcBuffer, dstBuffer);
+        break;
+      case 'b2t-copy':
+        assert(helper.commandEncoder !== undefined);
+        this.encodeReadByB2TCopy(helper.commandEncoder, srcBuffer, dstBuffer);
+        break;
+      default:
+        unreachable();
+    }
+  }
+
+  verifyData(buffer: GPUBuffer, expectedValue: number) {
+    const bufferData = new Uint32Array(1);
+    bufferData[0] = expectedValue;
+    this.expectGPUBufferValuesEqual(buffer, bufferData);
+  }
+
+  verifyDataTwoValidValues(buffer: GPUBuffer, expectedValue1: number, expectedValue2: number) {
+    const bufferData1 = new Uint32Array(1);
+    bufferData1[0] = expectedValue1;
+    const bufferData2 = new Uint32Array(1);
+    bufferData2[0] = expectedValue2;
+    this.expectGPUBufferValuesPassCheck(
+      buffer,
+      a => checkElementsEqualEither(a, [bufferData1, bufferData2]),
+      { type: Uint32Array, typedLength: 1 }
+    );
+  }
+}