summaryrefslogtreecommitdiffstats
path: root/dom/webgpu/tests/cts/checkout/src/stress/compute
diff options
context:
space:
mode:
authorDaniel Baumann <daniel.baumann@progress-linux.org>2024-04-07 19:33:14 +0000
committerDaniel Baumann <daniel.baumann@progress-linux.org>2024-04-07 19:33:14 +0000
commit36d22d82aa202bb199967e9512281e9a53db42c9 (patch)
tree105e8c98ddea1c1e4784a60a5a6410fa416be2de /dom/webgpu/tests/cts/checkout/src/stress/compute
parentInitial commit. (diff)
downloadfirefox-esr-36d22d82aa202bb199967e9512281e9a53db42c9.tar.xz
firefox-esr-36d22d82aa202bb199967e9512281e9a53db42c9.zip
Adding upstream version 115.7.0esr.upstream/115.7.0esrupstream
Signed-off-by: Daniel Baumann <daniel.baumann@progress-linux.org>
Diffstat (limited to 'dom/webgpu/tests/cts/checkout/src/stress/compute')
-rw-r--r--dom/webgpu/tests/cts/checkout/src/stress/compute/README.txt1
-rw-r--r--dom/webgpu/tests/cts/checkout/src/stress/compute/compute_pass.spec.ts243
2 files changed, 244 insertions, 0 deletions
diff --git a/dom/webgpu/tests/cts/checkout/src/stress/compute/README.txt b/dom/webgpu/tests/cts/checkout/src/stress/compute/README.txt
new file mode 100644
index 0000000000..b41aabc66b
--- /dev/null
+++ b/dom/webgpu/tests/cts/checkout/src/stress/compute/README.txt
@@ -0,0 +1 @@
+Stress tests covering operations specific to GPUComputePipeline and GPUComputePass.
diff --git a/dom/webgpu/tests/cts/checkout/src/stress/compute/compute_pass.spec.ts b/dom/webgpu/tests/cts/checkout/src/stress/compute/compute_pass.spec.ts
new file mode 100644
index 0000000000..76979f9fbb
--- /dev/null
+++ b/dom/webgpu/tests/cts/checkout/src/stress/compute/compute_pass.spec.ts
@@ -0,0 +1,243 @@
+export const description = `
+Stress tests covering GPUComputePassEncoder usage.
+`;
+
+import { makeTestGroup } from '../../common/framework/test_group.js';
+import { assert, iterRange } from '../../common/util/util.js';
+import { GPUTest } from '../../webgpu/gpu_test.js';
+
+export const g = makeTestGroup(GPUTest);
+
+g.test('many')
+ .desc(
+ `Tests execution of a huge number of compute passes using the same
+GPUComputePipeline.`
+ )
+ .fn(async t => {
+ const kNumElements = 64;
+ const data = new Uint32Array([...iterRange(kNumElements, x => x)]);
+ const buffer = t.makeBufferWithContents(data, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
+ const pipeline = t.device.createComputePipeline({
+ layout: 'auto',
+ compute: {
+ module: t.device.createShaderModule({
+ code: `
+ struct Buffer { data: array<u32>, };
+ @group(0) @binding(0) var<storage, read_write> buffer: Buffer;
+ @compute @workgroup_size(1) fn main(
+ @builtin(global_invocation_id) id: vec3<u32>) {
+ buffer.data[id.x] = buffer.data[id.x] + 1u;
+ }
+ `,
+ }),
+ entryPoint: 'main',
+ },
+ });
+ const bindGroup = t.device.createBindGroup({
+ layout: pipeline.getBindGroupLayout(0),
+ entries: [{ binding: 0, resource: { buffer } }],
+ });
+ const kNumIterations = 250_000;
+ for (let i = 0; i < kNumIterations; ++i) {
+ const encoder = t.device.createCommandEncoder();
+ const pass = encoder.beginComputePass();
+ pass.setPipeline(pipeline);
+ pass.setBindGroup(0, bindGroup);
+ pass.dispatchWorkgroups(kNumElements);
+ pass.end();
+ t.device.queue.submit([encoder.finish()]);
+ }
+ t.expectGPUBufferValuesEqual(
+ buffer,
+ new Uint32Array([...iterRange(kNumElements, x => x + kNumIterations)])
+ );
+ });
+
+g.test('pipeline_churn')
+ .desc(
+ `Tests execution of a huge number of compute passes which each use a different
+GPUComputePipeline.`
+ )
+ .fn(async t => {
+ const buffer = t.makeBufferWithContents(
+ new Uint32Array([0]),
+ GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
+ );
+ const kNumIterations = 10_000;
+ const stages = iterRange(kNumIterations, i => ({
+ module: t.device.createShaderModule({
+ code: `
+ struct Buffer { data: u32, };
+ @group(0) @binding(0) var<storage, read_write> buffer: Buffer;
+ @compute @workgroup_size(1) fn main${i}() {
+ buffer.data = buffer.data + 1u;
+ }
+ `,
+ }),
+ entryPoint: `main${i}`,
+ }));
+ for (const compute of stages) {
+ const encoder = t.device.createCommandEncoder();
+ const pipeline = t.device.createComputePipeline({ layout: 'auto', compute });
+ const bindGroup = t.device.createBindGroup({
+ layout: pipeline.getBindGroupLayout(0),
+ entries: [{ binding: 0, resource: { buffer } }],
+ });
+ const pass = encoder.beginComputePass();
+ pass.setPipeline(pipeline);
+ pass.setBindGroup(0, bindGroup);
+ pass.dispatchWorkgroups(1);
+ pass.end();
+ t.device.queue.submit([encoder.finish()]);
+ }
+ t.expectGPUBufferValuesEqual(buffer, new Uint32Array([kNumIterations]));
+ });
+
+g.test('bind_group_churn')
+ .desc(
+ `Tests execution of compute passes which switch between a huge number of bind
+groups.`
+ )
+ .fn(async t => {
+ const kNumElements = 64;
+ const data = new Uint32Array([...iterRange(kNumElements, x => x)]);
+ const buffer1 = t.makeBufferWithContents(
+ data,
+ GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
+ );
+ const buffer2 = t.makeBufferWithContents(
+ data,
+ GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
+ );
+ const module = t.device.createShaderModule({
+ code: `
+ struct Buffer { data: array<u32>, };
+ @group(0) @binding(0) var<storage, read_write> buffer1: Buffer;
+ @group(0) @binding(1) var<storage, read_write> buffer2: Buffer;
+ @compute @workgroup_size(1) fn main(
+ @builtin(global_invocation_id) id: vec3<u32>) {
+ buffer1.data[id.x] = buffer1.data[id.x] + 1u;
+ buffer2.data[id.x] = buffer2.data[id.x] + 2u;
+ }
+ `,
+ });
+ const kNumIterations = 250_000;
+ const pipeline = t.device.createComputePipeline({
+ layout: 'auto',
+ compute: { module, entryPoint: 'main' },
+ });
+ const encoder = t.device.createCommandEncoder();
+ const pass = encoder.beginComputePass();
+ pass.setPipeline(pipeline);
+ for (let i = 0; i < kNumIterations; ++i) {
+ const buffer1Binding = i % 2;
+ const buffer2Binding = buffer1Binding ^ 1;
+ const bindGroup = t.device.createBindGroup({
+ layout: pipeline.getBindGroupLayout(0),
+ entries: [
+ { binding: buffer1Binding, resource: { buffer: buffer1 } },
+ { binding: buffer2Binding, resource: { buffer: buffer2 } },
+ ],
+ });
+ pass.setBindGroup(0, bindGroup);
+ pass.dispatchWorkgroups(kNumElements);
+ }
+ pass.end();
+ t.device.queue.submit([encoder.finish()]);
+ const kTotalAddition = (kNumIterations / 2) * 3;
+ t.expectGPUBufferValuesEqual(
+ buffer1,
+ new Uint32Array([...iterRange(kNumElements, x => x + kTotalAddition)])
+ );
+ t.expectGPUBufferValuesEqual(
+ buffer2,
+ new Uint32Array([...iterRange(kNumElements, x => x + kTotalAddition)])
+ );
+ });
+
+g.test('many_dispatches')
+ .desc(`Tests execution of compute passes with a huge number of dispatch calls`)
+ .fn(async t => {
+ const kNumElements = 64;
+ const data = new Uint32Array([...iterRange(kNumElements, x => x)]);
+ const buffer = t.makeBufferWithContents(data, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
+ const module = t.device.createShaderModule({
+ code: `
+ struct Buffer { data: array<u32>, };
+ @group(0) @binding(0) var<storage, read_write> buffer: Buffer;
+ @compute @workgroup_size(1) fn main(
+ @builtin(global_invocation_id) id: vec3<u32>) {
+ buffer.data[id.x] = buffer.data[id.x] + 1u;
+ }
+ `,
+ });
+ const kNumIterations = 1_000_000;
+ const pipeline = t.device.createComputePipeline({
+ layout: 'auto',
+ compute: { module, entryPoint: 'main' },
+ });
+ const encoder = t.device.createCommandEncoder();
+ const pass = encoder.beginComputePass();
+ pass.setPipeline(pipeline);
+ const bindGroup = t.device.createBindGroup({
+ layout: pipeline.getBindGroupLayout(0),
+ entries: [{ binding: 0, resource: { buffer } }],
+ });
+ pass.setBindGroup(0, bindGroup);
+ for (let i = 0; i < kNumIterations; ++i) {
+ pass.dispatchWorkgroups(kNumElements);
+ }
+ pass.end();
+ t.device.queue.submit([encoder.finish()]);
+ t.expectGPUBufferValuesEqual(
+ buffer,
+ new Uint32Array([...iterRange(kNumElements, x => x + kNumIterations)])
+ );
+ });
+
+g.test('huge_dispatches')
+ .desc(`Tests execution of compute passes with huge dispatch calls`)
+ .fn(async t => {
+ const kDimensions = [512, 512, 128];
+ kDimensions.forEach(x => {
+ assert(x <= t.device.limits.maxComputeWorkgroupsPerDimension);
+ });
+
+ const kNumElements = kDimensions[0] * kDimensions[1] * kDimensions[2];
+ const data = new Uint32Array([...iterRange(kNumElements, x => x)]);
+ const buffer = t.makeBufferWithContents(data, GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC);
+ const module = t.device.createShaderModule({
+ code: `
+ struct Buffer { data: array<u32>, };
+ @group(0) @binding(0) var<storage, read_write> buffer: Buffer;
+ @compute @workgroup_size(1) fn main(
+ @builtin(global_invocation_id) id: vec3<u32>) {
+ let index = (id.z * 512u + id.y) * 512u + id.x;
+ buffer.data[index] = buffer.data[index] + 1u;
+ }
+ `,
+ });
+ const kNumIterations = 16;
+ const pipeline = t.device.createComputePipeline({
+ layout: 'auto',
+ compute: { module, entryPoint: 'main' },
+ });
+ const bindGroup = t.device.createBindGroup({
+ layout: pipeline.getBindGroupLayout(0),
+ entries: [{ binding: 0, resource: { buffer } }],
+ });
+ for (let i = 0; i < kNumIterations; ++i) {
+ const encoder = t.device.createCommandEncoder();
+ const pass = encoder.beginComputePass();
+ pass.setBindGroup(0, bindGroup);
+ pass.setPipeline(pipeline);
+ pass.dispatchWorkgroups(kDimensions[0], kDimensions[1], kDimensions[2]);
+ pass.end();
+ t.device.queue.submit([encoder.finish()]);
+ await t.device.queue.onSubmittedWorkDone();
+ }
+ t.expectGPUBufferValuesEqual(
+ buffer,
+ new Uint32Array([...iterRange(kNumElements, x => x + kNumIterations)])
+ );
+ });