From f1e8d8a6cd21220b1e9e309f9da73480837c23f6 Mon Sep 17 00:00:00 2001 From: Jiawei Shao Date: Mon, 23 Dec 2024 14:06:48 +0800 Subject: [PATCH 1/6] Test build checks --- src/webgpu/api/validation/createPipelineLayout.spec.ts | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/webgpu/api/validation/createPipelineLayout.spec.ts b/src/webgpu/api/validation/createPipelineLayout.spec.ts index 366d06993ac..df80a2a0813 100644 --- a/src/webgpu/api/validation/createPipelineLayout.spec.ts +++ b/src/webgpu/api/validation/createPipelineLayout.spec.ts @@ -223,7 +223,7 @@ g.test('bind_group_layouts,create_pipeline_with_null_bind_group_layouts') the shaders. ` ) - .params(u => + .paramsSubcasesOnly(u => u .combine('pipelineType', ['Render', 'Compute'] as const) .combine('emptyBindGroupLayoutType', ['Null', 'Undefined'] as const) From 0458008677f9973fb6f62e435efa22bdec15086c Mon Sep 17 00:00:00 2001 From: Jiawei Shao Date: Mon, 23 Dec 2024 15:17:52 +0800 Subject: [PATCH 2/6] Add tests on render and compute pipeline --- ...reated_with_null_bind_group_layout.spec.ts | 301 ++++++++++++++++++ .../validation/createPipelineLayout.spec.ts | 2 +- 2 files changed, 302 insertions(+), 1 deletion(-) create mode 100644 src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts diff --git a/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts b/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts new file mode 100644 index 00000000000..e10ccf118ff --- /dev/null +++ b/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts @@ -0,0 +1,301 @@ +export const description = ` +Tests for the creation of pipeline layouts with null bind group layouts. +`; + +import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { GPUConst } from '../../../constants.js'; +import { GPUTest } from '../../../gpu_test.js'; + +export const g = makeTestGroup(GPUTest); + +g.test('pipeline_layout_with_null_bind_group_layout,rendering') + .desc( + ` +Tests that using a render pipeline created with a pipeline layout that has null bind group layout +works correctly. +` + ) + .params(u => + u + .combine('emptyBindGroupLayoutType', ['Null', 'Undefined'] as const) + .combine('emptyBindGroupLayoutIndex', [0, 1, 2, 3] as const) + ) + .fn(t => { + const { emptyBindGroupLayoutType, emptyBindGroupLayoutIndex } = t.params; + + const colors = [ + [0.2, 0, 0, 0.2], + [0, 0.2, 0, 0.2], + [0, 0, 0.2, 0.2], + [0.4, 0, 0, 0.2], + ] as const; + const outputColor = [0.0, 0.0, 0.0, 0.0]; + + let declarations = ''; + let statement = 'return vec4(0.0, 0.0, 0.0, 0.0)'; + const bindGroupLayouts: (GPUBindGroupLayout | null | undefined)[] = []; + const bindGroups: (GPUBindGroup | null)[] = []; + for (let bindGroupIndex = 0; bindGroupIndex < 4; ++bindGroupIndex) { + if (bindGroupIndex === emptyBindGroupLayoutIndex) { + switch (emptyBindGroupLayoutType) { + case 'Null': + bindGroupLayouts.push(null); + break; + case 'Undefined': + bindGroupLayouts.push(undefined); + break; + } + bindGroups.push(null); + continue; + } + + declarations += `@group(${bindGroupIndex}) @binding(0) var input${bindGroupIndex} : vec4f;\n`; + statement += ` + input${bindGroupIndex}`; + + const bindGroupLayout = t.device.createBindGroupLayout({ + entries: [ + { + binding: 0, + visibility: GPUConst.ShaderStage.FRAGMENT, + buffer: { + type: 'uniform', + minBindingSize: 16, + }, + }, + ], + }); + bindGroupLayouts.push(bindGroupLayout); + + const color = colors[bindGroupIndex]; + const buffer = t.createBufferTracked({ + usage: GPUBufferUsage.UNIFORM, + size: 16, + mappedAtCreation: true, + }); + const bufferData = new Float32Array(buffer.getMappedRange()); + for (let i = 0; i < color.length; ++i) { + bufferData[i] = color[i]; + + outputColor[i] += color[i]; + } + buffer.unmap(); + + const bindGroup = t.device.createBindGroup({ + layout: bindGroupLayout, + entries: [ + { + binding: 0, + resource: { + buffer, + }, + }, + ], + }); + bindGroups.push(bindGroup); + } + + const pipelineLayout = t.device.createPipelineLayout({ + bindGroupLayouts, + }); + + const format = 'rgba8unorm'; + const code = ` + ${declarations} + @vertex + fn vert_main() -> @builtin(position) vec4f { + return vec4f(0.0, 0.0, 0.0, 1.0); + } + @fragment + fn frag_main() -> @location(0) vec4f { + ${statement}; + } + `; + const shaderModule = t.device.createShaderModule({ + code, + }); + const renderPipeline = t.device.createRenderPipeline({ + layout: pipelineLayout, + vertex: { + module: shaderModule, + }, + fragment: { + module: shaderModule, + targets: [ + { + format, + }, + ], + }, + primitive: { + topology: 'point-list', + }, + }); + + const renderTarget = t.createTextureTracked({ + usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.COPY_SRC, + size: [1, 1, 1], + format, + }); + const commandEncoder = t.device.createCommandEncoder(); + const renderPassEncoder = commandEncoder.beginRenderPass({ + colorAttachments: [ + { + view: renderTarget.createView(), + loadOp: 'load', + storeOp: 'store', + }, + ], + }); + for (let i = 0; i < 4; ++i) { + renderPassEncoder.setBindGroup(i, bindGroups[i]); + } + renderPassEncoder.setPipeline(renderPipeline); + renderPassEncoder.draw(1); + renderPassEncoder.end(); + + t.queue.submit([commandEncoder.finish()]); + + t.expectSingleColor(renderTarget, format, { + size: [1, 1, 1], + exp: { R: outputColor[0], G: outputColor[1], B: outputColor[2], A: outputColor[3] }, + }); + }); + +g.test('pipeline_layout_with_null_bind_group_layout,compute') + .desc( + ` +Tests that using a compute pipeline created with a pipeline layout that has null bind group layout +works correctly. +` + ) + .params(u => + u + .combine('emptyBindGroupLayoutType', ['Null', 'Undefined'] as const) + .combine('emptyBindGroupLayoutIndex', [0, 1, 2, 3] as const) + ) + .fn(t => { + const { emptyBindGroupLayoutType, emptyBindGroupLayoutIndex } = t.params; + + let declarations = ''; + let statement = 'output = 0u '; + + const outputBuffer = t.createBufferTracked({ + size: 4, + usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.STORAGE, + }); + let expectedValue = 0; + + const bindGroupLayouts: (GPUBindGroupLayout | null | undefined)[] = []; + const bindGroups: (GPUBindGroup | null)[] = []; + let outputDeclared = false; + for (let bindGroupIndex = 0; bindGroupIndex < 4; ++bindGroupIndex) { + if (bindGroupIndex === emptyBindGroupLayoutIndex) { + switch (emptyBindGroupLayoutType) { + case 'Null': + bindGroupLayouts.push(null); + break; + case 'Undefined': + bindGroupLayouts.push(undefined); + break; + } + bindGroups.push(null); + continue; + } + + declarations += `@group(${bindGroupIndex}) @binding(0) var input${bindGroupIndex} : u32;\n`; + statement += ` + input${bindGroupIndex}`; + + const inputBuffer = t.createBufferTracked({ + usage: GPUBufferUsage.UNIFORM, + size: 4, + mappedAtCreation: true, + }); + const bufferData = new Uint32Array(inputBuffer.getMappedRange()); + bufferData[0] = bindGroupIndex + 1; + expectedValue += bindGroupIndex + 1; + inputBuffer.unmap(); + + const bindGroupLayoutEntries: GPUBindGroupLayoutEntry[] = []; + const bindGroupEntries: GPUBindGroupEntry[] = []; + bindGroupLayoutEntries.push({ + binding: 0, + visibility: GPUConst.ShaderStage.COMPUTE, + buffer: { + type: 'uniform', + minBindingSize: 4, + }, + }); + bindGroupEntries.push({ + binding: 0, + resource: { + buffer: inputBuffer, + }, + }); + + if (!outputDeclared) { + bindGroupLayoutEntries.push({ + binding: 1, + visibility: GPUConst.ShaderStage.COMPUTE, + buffer: { + type: 'storage', + minBindingSize: 4, + }, + }); + bindGroupEntries.push({ + binding: 1, + resource: { + buffer: outputBuffer, + }, + }); + declarations += `@group(${bindGroupIndex}) @binding(1) var output : u32;\n`; + outputDeclared = true; + } + + const bindGroupLayout = t.device.createBindGroupLayout({ + entries: bindGroupLayoutEntries, + }); + bindGroupLayouts.push(bindGroupLayout); + + const bindGroup = t.device.createBindGroup({ + layout: bindGroupLayout, + entries: bindGroupEntries, + }); + bindGroups.push(bindGroup); + } + + const pipelineLayout = t.device.createPipelineLayout({ + bindGroupLayouts, + }); + + const code = ` + ${declarations} + @compute @workgroup_size(1, 1) + fn main() { + ${statement}; + } + `; + const module = t.device.createShaderModule({ + code, + }); + const computePipeline = t.device.createComputePipeline({ + layout: pipelineLayout, + compute: { + module, + }, + }); + + const commandEncoder = t.device.createCommandEncoder(); + const computePassEncoder = commandEncoder.beginComputePass(); + for (let i = 0; i < bindGroups.length; ++i) { + computePassEncoder.setBindGroup(i, bindGroups[i]); + } + computePassEncoder.setPipeline(computePipeline); + computePassEncoder.dispatchWorkgroups(1); + computePassEncoder.end(); + + t.queue.submit([commandEncoder.finish()]); + + const expectedValues = new Uint32Array(1); + expectedValues[0] = expectedValue; + t.expectGPUBufferValuesEqual(outputBuffer, expectedValues); + }); diff --git a/src/webgpu/api/validation/createPipelineLayout.spec.ts b/src/webgpu/api/validation/createPipelineLayout.spec.ts index df80a2a0813..366d06993ac 100644 --- a/src/webgpu/api/validation/createPipelineLayout.spec.ts +++ b/src/webgpu/api/validation/createPipelineLayout.spec.ts @@ -223,7 +223,7 @@ g.test('bind_group_layouts,create_pipeline_with_null_bind_group_layouts') the shaders. ` ) - .paramsSubcasesOnly(u => + .params(u => u .combine('pipelineType', ['Render', 'Compute'] as const) .combine('emptyBindGroupLayoutType', ['Null', 'Undefined'] as const) From 1733951413f0068e1b1f30b38ba6c55b73f91782 Mon Sep 17 00:00:00 2001 From: Jiawei Shao Date: Tue, 24 Dec 2024 14:58:59 +0800 Subject: [PATCH 3/6] Small fix --- .../pipeline_layout_created_with_null_bind_group_layout.spec.ts | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts b/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts index e10ccf118ff..7bbf9a58897 100644 --- a/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts +++ b/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts @@ -269,7 +269,7 @@ works correctly. const code = ` ${declarations} - @compute @workgroup_size(1, 1) + @compute @workgroup_size(1) fn main() { ${statement}; } From 0b1e5c1a91ce719732c9f2a8710a0763483e4e24 Mon Sep 17 00:00:00 2001 From: Jiawei Shao Date: Tue, 24 Dec 2024 16:43:49 +0800 Subject: [PATCH 4/6] Also set bind group on null bind group layout --- ...reated_with_null_bind_group_layout.spec.ts | 99 ++++++++++++------- 1 file changed, 61 insertions(+), 38 deletions(-) diff --git a/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts b/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts index 7bbf9a58897..50552ea672e 100644 --- a/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts +++ b/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts @@ -34,24 +34,8 @@ works correctly. let declarations = ''; let statement = 'return vec4(0.0, 0.0, 0.0, 0.0)'; const bindGroupLayouts: (GPUBindGroupLayout | null | undefined)[] = []; - const bindGroups: (GPUBindGroup | null)[] = []; + const bindGroups: GPUBindGroup[] = []; for (let bindGroupIndex = 0; bindGroupIndex < 4; ++bindGroupIndex) { - if (bindGroupIndex === emptyBindGroupLayoutIndex) { - switch (emptyBindGroupLayoutType) { - case 'Null': - bindGroupLayouts.push(null); - break; - case 'Undefined': - bindGroupLayouts.push(undefined); - break; - } - bindGroups.push(null); - continue; - } - - declarations += `@group(${bindGroupIndex}) @binding(0) var input${bindGroupIndex} : vec4f;\n`; - statement += ` + input${bindGroupIndex}`; - const bindGroupLayout = t.device.createBindGroupLayout({ entries: [ { @@ -64,7 +48,6 @@ works correctly. }, ], }); - bindGroupLayouts.push(bindGroupLayout); const color = colors[bindGroupIndex]; const buffer = t.createBufferTracked({ @@ -75,11 +58,11 @@ works correctly. const bufferData = new Float32Array(buffer.getMappedRange()); for (let i = 0; i < color.length; ++i) { bufferData[i] = color[i]; - - outputColor[i] += color[i]; } buffer.unmap(); + // Still create and set the bind group when the corresponding bind group layout in the + // pipeline is null. The output color should not be affected by the buffer in this bind group const bindGroup = t.device.createBindGroup({ layout: bindGroupLayout, entries: [ @@ -92,6 +75,30 @@ works correctly. ], }); bindGroups.push(bindGroup); + + // Set `null` or `undefined` in `bindGroupLayouts` which is used in the creation of pipeline + // layout + if (bindGroupIndex === emptyBindGroupLayoutIndex) { + switch (emptyBindGroupLayoutType) { + case 'Null': + bindGroupLayouts.push(null); + break; + case 'Undefined': + bindGroupLayouts.push(undefined); + break; + } + continue; + } + + // Set the uniform buffers used in the shader + bindGroupLayouts.push(bindGroupLayout); + declarations += `@group(${bindGroupIndex}) @binding(0) var input${bindGroupIndex} : vec4f;\n`; + statement += ` + input${bindGroupIndex}`; + + // Compute the expected output color + for (let i = 0; i < color.length; ++i) { + outputColor[i] += color[i]; + } } const pipelineLayout = t.device.createPipelineLayout({ @@ -186,25 +193,9 @@ works correctly. let expectedValue = 0; const bindGroupLayouts: (GPUBindGroupLayout | null | undefined)[] = []; - const bindGroups: (GPUBindGroup | null)[] = []; + const bindGroups: GPUBindGroup[] = []; let outputDeclared = false; for (let bindGroupIndex = 0; bindGroupIndex < 4; ++bindGroupIndex) { - if (bindGroupIndex === emptyBindGroupLayoutIndex) { - switch (emptyBindGroupLayoutType) { - case 'Null': - bindGroupLayouts.push(null); - break; - case 'Undefined': - bindGroupLayouts.push(undefined); - break; - } - bindGroups.push(null); - continue; - } - - declarations += `@group(${bindGroupIndex}) @binding(0) var input${bindGroupIndex} : u32;\n`; - statement += ` + input${bindGroupIndex}`; - const inputBuffer = t.createBufferTracked({ usage: GPUBufferUsage.UNIFORM, size: 4, @@ -212,7 +203,6 @@ works correctly. }); const bufferData = new Uint32Array(inputBuffer.getMappedRange()); bufferData[0] = bindGroupIndex + 1; - expectedValue += bindGroupIndex + 1; inputBuffer.unmap(); const bindGroupLayoutEntries: GPUBindGroupLayoutEntry[] = []; @@ -232,6 +222,35 @@ works correctly. }, }); + // Set `null` or `undefined` in `bindGroupLayouts` which is used in the creation of pipeline + // layout + if (bindGroupIndex === emptyBindGroupLayoutIndex) { + switch (emptyBindGroupLayoutType) { + case 'Null': + bindGroupLayouts.push(null); + break; + case 'Undefined': + bindGroupLayouts.push(undefined); + break; + } + + // Still create and set the bind group when the corresponding bind group layout in the + // compute pipeline is null. The value in the output buffer should not be affected by the + // buffer in this bind group + const bindGroup = t.device.createBindGroup({ + layout: t.device.createBindGroupLayout({ + entries: bindGroupLayoutEntries, + }), + entries: bindGroupEntries, + }); + bindGroups.push(bindGroup); + continue; + } + + declarations += `@group(${bindGroupIndex}) @binding(0) var input${bindGroupIndex} : u32;\n`; + statement += ` + input${bindGroupIndex}`; + + // Set the output storage buffer if (!outputDeclared) { bindGroupLayoutEntries.push({ binding: 1, @@ -251,6 +270,7 @@ works correctly. outputDeclared = true; } + // Set the input uniform buffers const bindGroupLayout = t.device.createBindGroupLayout({ entries: bindGroupLayoutEntries, }); @@ -261,6 +281,9 @@ works correctly. entries: bindGroupEntries, }); bindGroups.push(bindGroup); + + // Compute the expected output value in the output storage buffer + expectedValue += bindGroupIndex + 1; } const pipelineLayout = t.device.createPipelineLayout({ From 108258197e44969c1b64db67aa8c809a3c86dd0c Mon Sep 17 00:00:00 2001 From: Jiawei Shao Date: Wed, 25 Dec 2024 16:04:06 +0800 Subject: [PATCH 5/6] Test empty bind group layouts --- ...reated_with_null_bind_group_layout.spec.ts | 26 ++++++++++++++----- 1 file changed, 20 insertions(+), 6 deletions(-) diff --git a/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts b/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts index 50552ea672e..bd4b848b2c0 100644 --- a/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts +++ b/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts @@ -17,7 +17,7 @@ works correctly. ) .params(u => u - .combine('emptyBindGroupLayoutType', ['Null', 'Undefined'] as const) + .combine('emptyBindGroupLayoutType', ['Null', 'Undefined', 'Empty'] as const) .combine('emptyBindGroupLayoutIndex', [0, 1, 2, 3] as const) ) .fn(t => { @@ -76,8 +76,8 @@ works correctly. }); bindGroups.push(bindGroup); - // Set `null` or `undefined` in `bindGroupLayouts` which is used in the creation of pipeline - // layout + // Set `null`, `undefined` or empty bind group layout in `bindGroupLayouts` which is used in + // the creation of pipeline layout if (bindGroupIndex === emptyBindGroupLayoutIndex) { switch (emptyBindGroupLayoutType) { case 'Null': @@ -86,6 +86,13 @@ works correctly. case 'Undefined': bindGroupLayouts.push(undefined); break; + case 'Empty': + bindGroupLayouts.push( + t.device.createBindGroupLayout({ + entries: [], + }) + ); + break; } continue; } @@ -177,7 +184,7 @@ works correctly. ) .params(u => u - .combine('emptyBindGroupLayoutType', ['Null', 'Undefined'] as const) + .combine('emptyBindGroupLayoutType', ['Null', 'Undefined', 'Empty'] as const) .combine('emptyBindGroupLayoutIndex', [0, 1, 2, 3] as const) ) .fn(t => { @@ -222,8 +229,8 @@ works correctly. }, }); - // Set `null` or `undefined` in `bindGroupLayouts` which is used in the creation of pipeline - // layout + // Set `null`, `undefined` or empty bind group layout in `bindGroupLayouts` which is used in + // the creation of pipeline layout if (bindGroupIndex === emptyBindGroupLayoutIndex) { switch (emptyBindGroupLayoutType) { case 'Null': @@ -232,6 +239,13 @@ works correctly. case 'Undefined': bindGroupLayouts.push(undefined); break; + case 'Empty': + bindGroupLayouts.push( + t.device.createBindGroupLayout({ + entries: [], + }) + ); + break; } // Still create and set the bind group when the corresponding bind group layout in the From da6b1248d5b4c1b29cb28f84516baf3fb53ee68c Mon Sep 17 00:00:00 2001 From: Jiawei Shao Date: Thu, 9 Jan 2025 15:04:45 +0800 Subject: [PATCH 6/6] Address reviewers' comments --- ...reated_with_null_bind_group_layout.spec.ts | 26 +++++-------------- 1 file changed, 6 insertions(+), 20 deletions(-) diff --git a/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts b/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts index bd4b848b2c0..10416e2b4d0 100644 --- a/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts +++ b/src/webgpu/api/operation/pipeline/pipeline_layout_created_with_null_bind_group_layout.spec.ts @@ -50,16 +50,7 @@ works correctly. }); const color = colors[bindGroupIndex]; - const buffer = t.createBufferTracked({ - usage: GPUBufferUsage.UNIFORM, - size: 16, - mappedAtCreation: true, - }); - const bufferData = new Float32Array(buffer.getMappedRange()); - for (let i = 0; i < color.length; ++i) { - bufferData[i] = color[i]; - } - buffer.unmap(); + const buffer = t.makeBufferWithContents(new Float32Array(color), GPUBufferUsage.UNIFORM); // Still create and set the bind group when the corresponding bind group layout in the // pipeline is null. The output color should not be affected by the buffer in this bind group @@ -203,14 +194,10 @@ works correctly. const bindGroups: GPUBindGroup[] = []; let outputDeclared = false; for (let bindGroupIndex = 0; bindGroupIndex < 4; ++bindGroupIndex) { - const inputBuffer = t.createBufferTracked({ - usage: GPUBufferUsage.UNIFORM, - size: 4, - mappedAtCreation: true, - }); - const bufferData = new Uint32Array(inputBuffer.getMappedRange()); - bufferData[0] = bindGroupIndex + 1; - inputBuffer.unmap(); + const inputBuffer = t.makeBufferWithContents( + new Uint32Array([bindGroupIndex + 1]), + GPUBufferUsage.UNIFORM + ); const bindGroupLayoutEntries: GPUBindGroupLayoutEntry[] = []; const bindGroupEntries: GPUBindGroupEntry[] = []; @@ -332,7 +319,6 @@ works correctly. t.queue.submit([commandEncoder.finish()]); - const expectedValues = new Uint32Array(1); - expectedValues[0] = expectedValue; + const expectedValues = new Uint32Array([expectedValue]); t.expectGPUBufferValuesEqual(outputBuffer, expectedValues); });