From bc32dbc5d4685a43b018bed153d328469b270572 Mon Sep 17 00:00:00 2001 From: James Price Date: Mon, 29 Jan 2024 15:48:04 -0500 Subject: [PATCH 1/5] shader/validation: Test function addrspace at module-scope (#3347) Issue #1570 --- src/webgpu/listing_meta.json | 1 + src/webgpu/shader/validation/decl/var.spec.ts | 10 ++++++++++ 2 files changed, 11 insertions(+) diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index 614e023f2035..879634ebb294 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -1730,6 +1730,7 @@ "webgpu:shader,validation,decl,ptr_spelling:ptr_not_instantiable:*": { "subcaseMS": 1.310 }, "webgpu:shader,validation,decl,var:module_scope_types:*": { "subcaseMS": 1.000 }, "webgpu:shader,validation,decl,var:function_scope_types:*": { "subcaseMS": 1.000 }, + "webgpu:shader,validation,decl,var:function_addrspace_at_module_scope:*": { "subcaseMS": 1.000 }, "webgpu:shader,validation,decl,var_access_mode:explicit_access_mode:*": { "subcaseMS": 1.373 }, "webgpu:shader,validation,decl,var_access_mode:implicit_access_mode:*": { "subcaseMS": 1.000 }, "webgpu:shader,validation,decl,var_access_mode:read_access:*": { "subcaseMS": 1.177 }, diff --git a/src/webgpu/shader/validation/decl/var.spec.ts b/src/webgpu/shader/validation/decl/var.spec.ts index bf2772efa1c8..8e3f69c7b4df 100644 --- a/src/webgpu/shader/validation/decl/var.spec.ts +++ b/src/webgpu/shader/validation/decl/var.spec.ts @@ -328,3 +328,13 @@ g.test('function_scope_types') t.expectCompileResult(shouldPass, wgsl); }); + +g.test('function_addrspace_at_module_scope') + .desc('Test that the function address space is not allowed at module scope.') + .params(u => u.combine('addrspace', ['private', 'function'])) + .fn(t => { + t.expectCompileResult( + t.params.addrspace === 'private', + `var<${t.params.addrspace}> foo : i32;` + ); + }); From eeb30fbf741ab978fd338feb2c171ae6b981014e Mon Sep 17 00:00:00 2001 From: David Neto Date: Mon, 29 Jan 2024 15:52:24 -0500 Subject: [PATCH 2/5] wgsl: Test mixing stage attributes on the same function (#3341) Add to the existing tests for parsing stage attributes @compute, @fragment, @vertex --- src/webgpu/listing_meta.json | 6 +-- .../validation/parse/pipeline_stage.spec.ts | 42 ++++++++++++------- 2 files changed, 30 insertions(+), 18 deletions(-) diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index 879634ebb294..bac7c63ca75e 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -1924,9 +1924,9 @@ "webgpu:shader,validation,parse,must_use:declaration:*": { "subcaseMS": 1.523 }, "webgpu:shader,validation,parse,must_use:ignore_result_of_non_must_use_that_returns_call_of_must_use:*": { "subcaseMS": 0.0 }, "webgpu:shader,validation,parse,pipeline_stage:compute_parsing:*": { "subcaseMS": 1.000 }, - "webgpu:shader,validation,parse,pipeline_stage:duplicate_compute_on_function:*": { "subcaseMS": 2.651 }, - "webgpu:shader,validation,parse,pipeline_stage:duplicate_fragment_on_function:*": { "subcaseMS": 1.001 }, - "webgpu:shader,validation,parse,pipeline_stage:duplicate_vertex_on_function:*": { "subcaseMS": 1.000 }, + "webgpu:shader,validation,parse,pipeline_stage:extra_on_compute_function:*": { "subcaseMS": 2.651 }, + "webgpu:shader,validation,parse,pipeline_stage:extra_on_fragment_function:*": { "subcaseMS": 1.001 }, + "webgpu:shader,validation,parse,pipeline_stage:extra_on_vertex_function:*": { "subcaseMS": 1.000 }, "webgpu:shader,validation,parse,pipeline_stage:fragment_parsing:*": { "subcaseMS": 2.600 }, "webgpu:shader,validation,parse,pipeline_stage:multiple_entry_points:*": { "subcaseMS": 1.100 }, "webgpu:shader,validation,parse,pipeline_stage:placement:*": { "subcaseMS": 1.388 }, diff --git a/src/webgpu/shader/validation/parse/pipeline_stage.spec.ts b/src/webgpu/shader/validation/parse/pipeline_stage.spec.ts index 78dcb9578233..f492121f25e4 100644 --- a/src/webgpu/shader/validation/parse/pipeline_stage.spec.ts +++ b/src/webgpu/shader/validation/parse/pipeline_stage.spec.ts @@ -73,34 +73,46 @@ g.test('multiple_entry_points') t.expectCompileResult(true, code); }); -g.test('duplicate_compute_on_function') - .desc(`Test that duplcate @compute attributes are not allowed.`) - .params(u => u.combine('dupe', ['', '@compute'])) +g.test('extra_on_compute_function') + .desc(`Test that an extra stage attribute on @compute functions are not allowed.`) + .params(u => + u.combine('extra', ['', '@compute', '@fragment', '@vertex']).combine('before', [false, true]) + ) .fn(t => { + const before = t.params.before ? t.params.extra : ''; + const after = t.params.before ? '' : t.params.extra; const code = ` -@compute ${t.params.dupe} @workgroup_size(1) fn compute_1() {} +${before} @compute ${after} @workgroup_size(1) fn main() {} `; - t.expectCompileResult(t.params.dupe === '', code); + t.expectCompileResult(t.params.extra === '', code); }); -g.test('duplicate_fragment_on_function') - .desc(`Test that duplcate @fragment attributes are not allowed.`) - .params(u => u.combine('dupe', ['', '@fragment'])) +g.test('extra_on_fragment_function') + .desc(`Test that an extra stage attribute on @fragment functions are not allowed.`) + .params(u => + u.combine('extra', ['', '@compute', '@fragment', '@vertex']).combine('before', [false, true]) + ) .fn(t => { + const before = t.params.before ? t.params.extra : ''; + const after = t.params.before ? '' : t.params.extra; const code = ` -@fragment ${t.params.dupe} fn vtx() -> @location(0) vec4f { return vec4f(1); } +${before} @fragment ${after} fn main() -> @location(0) vec4f { return vec4f(1); } `; - t.expectCompileResult(t.params.dupe === '', code); + t.expectCompileResult(t.params.extra === '', code); }); -g.test('duplicate_vertex_on_function') - .desc(`Test that duplcate @vertex attributes are not allowed.`) - .params(u => u.combine('dupe', ['', '@vertex'])) +g.test('extra_on_vertex_function') + .desc(`Test that an extra stage attribute on @vertex functions are not allowed.`) + .params(u => + u.combine('extra', ['', '@compute', '@fragment', '@vertex']).combine('before', [false, true]) + ) .fn(t => { + const before = t.params.before ? t.params.extra : ''; + const after = t.params.before ? '' : t.params.extra; const code = ` -@vertex ${t.params.dupe} fn vtx() -> @builtin(position) vec4f { return vec4f(1); } +${before} @vertex ${after} fn main() -> @builtin(position) vec4f { return vec4f(1); } `; - t.expectCompileResult(t.params.dupe === '', code); + t.expectCompileResult(t.params.extra === '', code); }); g.test('placement') From 1969f842b956d09fef2bf9ab60f88b781e782a61 Mon Sep 17 00:00:00 2001 From: David Neto Date: Mon, 29 Jan 2024 15:58:51 -0500 Subject: [PATCH 3/5] wgsl: Add trivial execution test for each shader stage (#3343) --- src/webgpu/listing_meta.json | 2 + src/webgpu/shader/execution/stage.spec.ts | 133 ++++++++++++++++++++++ 2 files changed, 135 insertions(+) create mode 100644 src/webgpu/shader/execution/stage.spec.ts diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index bac7c63ca75e..dcebf8fdadc7 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -1692,6 +1692,8 @@ "webgpu:shader,execution,shadow:loop:*": { "subcaseMS": 4.901 }, "webgpu:shader,execution,shadow:switch:*": { "subcaseMS": 4.601 }, "webgpu:shader,execution,shadow:while:*": { "subcaseMS": 7.400 }, + "webgpu:shader,execution,stage:basic_compute:*": { "subcaseMS": 1.000 }, + "webgpu:shader,execution,stage:basic_render:*": { "subcaseMS": 1.000 }, "webgpu:shader,execution,statement,increment_decrement:frexp_exp_increment:*": { "subcaseMS": 4.700 }, "webgpu:shader,execution,statement,increment_decrement:scalar_i32_decrement:*": { "subcaseMS": 20.301 }, "webgpu:shader,execution,statement,increment_decrement:scalar_i32_decrement_underflow:*": { "subcaseMS": 4.900 }, diff --git a/src/webgpu/shader/execution/stage.spec.ts b/src/webgpu/shader/execution/stage.spec.ts new file mode 100644 index 000000000000..6e06e67e3776 --- /dev/null +++ b/src/webgpu/shader/execution/stage.spec.ts @@ -0,0 +1,133 @@ +export const description = `Test trivial shaders for each shader stage kind`; + +// There are many many more shaders executed in other tests. + +import { makeTestGroup } from '../../../common/framework/test_group.js'; +import { GPUTest } from '../../gpu_test.js'; +import { checkElementsEqual } from '../../util/check_contents.js'; + +export const g = makeTestGroup(GPUTest); + +g.test('basic_compute') + .desc(`Test a trivial compute shader`) + .fn(async t => { + const code = ` + +@group(0) @binding(0) +var v : vec4u; + +@compute @workgroup_size(1) +fn main() { + v = vec4u(1,2,3,42); +}`; + + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ + code, + }), + entryPoint: 'main', + }, + }); + + const buffer = t.makeBufferWithContents( + new Uint32Array([0, 0, 0, 0]), + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST + ); + t.trackForCleanup(buffer); + + const bg = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer, + }, + }, + ], + }); + + const encoder = t.device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup(0, bg); + pass.dispatchWorkgroups(1, 1, 1); + pass.end(); + t.queue.submit([encoder.finish()]); + + const bufferReadback = await t.readGPUBufferRangeTyped(buffer, { + srcByteOffset: 0, + type: Uint32Array, + typedLength: 4, + method: 'copy', + }); + const got: Uint32Array = bufferReadback.data; + const expected = new Uint32Array([1, 2, 3, 42]); + + t.expectOK(checkElementsEqual(got, expected)); + }); + +g.test('basic_render') + .desc(`Test trivial vertex and fragment shaders`) + .fn(t => { + const code = ` +@vertex +fn vert_main(@builtin(vertex_index) idx: u32) -> @builtin(position) vec4f { + // A right triangle covering the whole framebuffer. + const pos = array( + vec2f(-1,-3), + vec2f(3,1), + vec2f(-1,1)); + return vec4f(pos[idx], 0, 1); +} + +@fragment +fn frag_main() -> @location(0) vec4f { + return vec4(0, 1, 0, 1); // green +} +`; + const module = t.device.createShaderModule({ code }); + + const [width, height] = [8, 8] as const; + const format = 'rgba8unorm' as const; + const texture = t.device.createTexture({ + size: { width, height }, + usage: + GPUTextureUsage.RENDER_ATTACHMENT | + GPUTextureUsage.TEXTURE_BINDING | + GPUTextureUsage.COPY_SRC, + format, + }); + + // We'll copy one pixel only. + const dst = t.device.createBuffer({ + size: 4, + usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, + }); + + const pipeline = t.device.createRenderPipeline({ + layout: 'auto', + vertex: { module, entryPoint: 'vert_main' }, + fragment: { module, entryPoint: 'frag_main', targets: [{ format }] }, + }); + + const encoder = t.device.createCommandEncoder(); + const pass = encoder.beginRenderPass({ + colorAttachments: [{ view: texture.createView(), loadOp: 'clear', storeOp: 'store' }], + }); + pass.setPipeline(pipeline); + pass.draw(3); + pass.end(); + + encoder.copyTextureToBuffer( + { texture, mipLevel: 0, origin: { x: 0, y: 0, z: 0 } }, + { buffer: dst, bytesPerRow: 256 }, + { width: 1, height: 1, depthOrArrayLayers: 1 } + ); + t.queue.submit([encoder.finish()]); + + // Expect one green pixel. + t.expectGPUBufferValuesEqual(dst, new Uint8Array([0x00, 0xff, 0x00, 0xff])); + }); From ecf7770c1f127b52e972e6db6a65bd77418aeebf Mon Sep 17 00:00:00 2001 From: David Neto Date: Mon, 29 Jan 2024 16:05:56 -0500 Subject: [PATCH 4/5] wgsl: Add compound statement tests (#3337) Validation: - Check parsing of compound statements - Update break-if statement to add a case for a plain compound statement surrounding a break-if. - For a declaration in a compound statement, check potential uses and potentially conflicting declarations. Execution: Add execution tests for compound statement. Observe a value at various points relative to its declaration inside a compound statement. Add timings for new tests Fixed: #1634 --- src/webgpu/listing_meta.json | 4 + .../execution/statement/compound.spec.ts | 137 ++++++++++++++++++ .../decl/compound_statement.spec.ts | 98 +++++++++++++ .../shader/validation/parse/break_if.spec.ts | 4 + .../shader/validation/parse/compound.spec.ts | 52 +++++++ 5 files changed, 295 insertions(+) create mode 100644 src/webgpu/shader/execution/statement/compound.spec.ts create mode 100644 src/webgpu/shader/validation/decl/compound_statement.spec.ts create mode 100644 src/webgpu/shader/validation/parse/compound.spec.ts diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index dcebf8fdadc7..2dc721e85989 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -1694,6 +1694,7 @@ "webgpu:shader,execution,shadow:while:*": { "subcaseMS": 7.400 }, "webgpu:shader,execution,stage:basic_compute:*": { "subcaseMS": 1.000 }, "webgpu:shader,execution,stage:basic_render:*": { "subcaseMS": 1.000 }, + "webgpu:shader,execution,statement,compound:decl:*": { "subcaseMS": 29.767 }, "webgpu:shader,execution,statement,increment_decrement:frexp_exp_increment:*": { "subcaseMS": 4.700 }, "webgpu:shader,execution,statement,increment_decrement:scalar_i32_decrement:*": { "subcaseMS": 20.301 }, "webgpu:shader,execution,statement,increment_decrement:scalar_i32_decrement_underflow:*": { "subcaseMS": 4.900 }, @@ -1717,6 +1718,8 @@ "webgpu:shader,validation,const_assert,const_assert:constant_expression_logical_or_no_assert:*": { "subcaseMS": 1.373 }, "webgpu:shader,validation,const_assert,const_assert:constant_expression_no_assert:*": { "subcaseMS": 1.655 }, "webgpu:shader,validation,const_assert,const_assert:evaluation_stage:*": { "subcaseMS": 3.367 }, + "webgpu:shader,validation,decl,compound_statement:decl_conflict:*": { "subcaseMS": 5.225 }, + "webgpu:shader,validation,decl,compound_statement:decl_use:*": { "subcaseMS": 0.625 }, "webgpu:shader,validation,decl,const:no_direct_recursion:*": { "subcaseMS": 0.951 }, "webgpu:shader,validation,decl,const:no_indirect_recursion:*": { "subcaseMS": 0.950 }, "webgpu:shader,validation,decl,const:no_indirect_recursion_via_array_size:*": { "subcaseMS": 2.601 }, @@ -1889,6 +1892,7 @@ "webgpu:shader,validation,parse,comments:line_comment_eof:*": { "subcaseMS": 4.500 }, "webgpu:shader,validation,parse,comments:line_comment_terminators:*": { "subcaseMS": 1.021 }, "webgpu:shader,validation,parse,comments:unterminated_block_comment:*": { "subcaseMS": 8.950 }, + "webgpu:shader,validation,parse,compound:parse:*": { "subcaseMS": 4.315 }, "webgpu:shader,validation,parse,const:placement:*": { "subcaseMS": 1.167 }, "webgpu:shader,validation,parse,const_assert:parse:*": { "subcaseMS": 1.400 }, "webgpu:shader,validation,parse,diagnostic:after_other_directives:*": { "subcaseMS": 1.000 }, diff --git a/src/webgpu/shader/execution/statement/compound.spec.ts b/src/webgpu/shader/execution/statement/compound.spec.ts new file mode 100644 index 000000000000..aed0cc224557 --- /dev/null +++ b/src/webgpu/shader/execution/statement/compound.spec.ts @@ -0,0 +1,137 @@ +export const description = ` +Compound statement execution. +`; + +import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { keysOf } from '../../../../common/util/data_tables.js'; +import { TypedArrayBufferView } from '../../../../common/util/util.js'; +import { GPUTest } from '../../../gpu_test.js'; + +export const g = makeTestGroup(GPUTest); + +/** + * Builds, runs then checks the output of a statement shader test. + * + * @param t The test object + * @param ty The WGSL scalar type to be written + * @param values The expected output values of type ty + * @param wgsl_main The body of the WGSL entry point. + */ +export function runStatementTest( + t: GPUTest, + ty: string, + values: TypedArrayBufferView, + wgsl_main: string +) { + const wgsl = ` +struct Outputs { + data : array<${ty}>, +}; +var count: u32 = 0; + +@group(0) @binding(1) var outputs : Outputs; + +fn put(value : ${ty}) { + outputs.data[count] = value; + count += 1; +} + +@compute @workgroup_size(1) +fn main() { + _ = &outputs; + ${wgsl_main} +} +`; + + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ code: wgsl }), + entryPoint: 'main', + }, + }); + + const maxOutputValues = 1000; + const outputBuffer = t.device.createBuffer({ + size: 4 * (1 + maxOutputValues), + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + + const bindGroup = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [{ binding: 1, resource: { buffer: outputBuffer } }], + }); + + // Run the shader. + const encoder = t.device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup(0, bindGroup); + pass.dispatchWorkgroups(1); + pass.end(); + t.queue.submit([encoder.finish()]); + + t.expectGPUBufferValuesEqual(outputBuffer, values); +} + +// Consider a declaration X of identifier 'x' inside a compound statement. +// Check the value of 'x' at various places relative to X: +// a { b; X=c; d; { e; } } f; + +const kTests = { + uses: { + // Observe values without conflicting declarations. + src: `let x = 1; + put(x); + { + put(x); + let x = x+1; // The declaration in question + put(x); + { + put(x); + } + put(x); + } + put(x);`, + values: [1, 1, 2, 2, 2, 1], + }, + shadowed: { + // Observe values when shadowed + src: `let x = 1; + put(x); + { + put(x); + let x = x+1; // The declaration in question + put(x); + { + let x = x+1; // A shadow + put(x); + } + put(x); + } + put(x);`, + values: [1, 1, 2, 3, 2, 1], + }, + gone: { + // The declaration goes out of scope. + src: `{ + let x = 2; // The declaration in question + put(x); + } + let x = 1; + put(x);`, + values: [2, 1], + }, +} as const; + +g.test('decl') + .desc('Tests the value of a declared value in a compound statment.') + .params(u => u.combine('case', keysOf(kTests))) + .fn(t => { + runStatementTest( + t, + 'i32', + new Int32Array(kTests[t.params.case].values), + kTests[t.params.case].src + ); + }); diff --git a/src/webgpu/shader/validation/decl/compound_statement.spec.ts b/src/webgpu/shader/validation/decl/compound_statement.spec.ts new file mode 100644 index 000000000000..8ad89f48a8aa --- /dev/null +++ b/src/webgpu/shader/validation/decl/compound_statement.spec.ts @@ -0,0 +1,98 @@ +export const description = ` +Validation tests for declarations in compound statements. +`; + +import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { keysOf } from '../../../../common/util/data_tables.js'; +import { ShaderValidationTest } from '../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +// 9.1 Compound Statements +// When a declaration is one of those statements, its identifier is in scope from +// the start of the next statement until the end of the compound statement. +// +// Enumerating cases: Consider a declaration X inside a compound statement. +// The X declaration should be tested with potential uses and potentially +// conflicting declarations in positions [a, b, c, d, e], in the following: +// a { b; X; c; { d; } } e; + +const kConflictTests = { + a: { + src: 'let x = 1; { let x = 1; }', + pass: true, + }, + bc: { + src: '{let x = 1; let x = 1; }', + pass: false, + }, + d: { + src: '{let x = 1; { let x = 1; }}', + pass: true, + }, + e: { + src: '{let x = 1; } let x = 1;', + pass: true, + }, +}; + +g.test('decl_conflict') + .desc( + 'Test a potentially conflicting declaration relative to a declaration in a compound statement' + ) + .params(u => u.combine('case', keysOf(kConflictTests))) + .fn(t => { + const wgsl = ` +@vertex fn vtx() -> @builtin(position) vec4f { + ${kConflictTests[t.params.case].src} + return vec4f(1); +}`; + t.expectCompileResult(kConflictTests[t.params.case].pass, wgsl); + }); + +const kUseTests = { + a: { + src: 'let y = x; { let x = 1; }', + pass: false, // not visible + }, + b: { + src: '{ let y = x; let x = 1; }', + pass: false, // not visible + }, + self: { + src: '{ let x = (x);}', + pass: false, // not visible + }, + c_yes: { + src: '{ const x = 1; const_assert x == 1; }', + pass: true, + }, + c_no: { + src: '{ const x = 1; const_assert x == 2; }', + pass: false, + }, + d_yes: { + src: '{ const x = 1; { const_assert x == 1; }}', + pass: true, + }, + d_no: { + src: '{ const x = 1; { const_assert x == 2; }}', + pass: false, + }, + e: { + src: '{ const x = 1; } let y = x;', + pass: false, // not visible + }, +}; + +g.test('decl_use') + .desc('Test a use of a declaration in a compound statement') + .params(u => u.combine('case', keysOf(kUseTests))) + .fn(t => { + const wgsl = ` +@vertex fn vtx() -> @builtin(position) vec4f { + ${kUseTests[t.params.case].src} + return vec4f(1); +}`; + t.expectCompileResult(kUseTests[t.params.case].pass, wgsl); + }); diff --git a/src/webgpu/shader/validation/parse/break_if.spec.ts b/src/webgpu/shader/validation/parse/break_if.spec.ts index 41d448437e45..97a625f625cc 100644 --- a/src/webgpu/shader/validation/parse/break_if.spec.ts +++ b/src/webgpu/shader/validation/parse/break_if.spec.ts @@ -7,6 +7,10 @@ import { ShaderValidationTest } from '../shader_validation_test.js'; export const g = makeTestGroup(ShaderValidationTest); const kTests = { + compound_break: { + src: '{ break if true; }', + pass: false, + }, loop_break: { src: 'loop { break if true; }', pass: false, diff --git a/src/webgpu/shader/validation/parse/compound.spec.ts b/src/webgpu/shader/validation/parse/compound.spec.ts new file mode 100644 index 000000000000..b3627c2e5b9e --- /dev/null +++ b/src/webgpu/shader/validation/parse/compound.spec.ts @@ -0,0 +1,52 @@ +export const description = `Validation tests for compound statements`; + +import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { keysOf } from '../../../../common/util/data_tables.js'; +import { ShaderValidationTest } from '../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +const kTests = { + missing_start: { + src: '}', + pass: false, + }, + missing_end: { + src: '{', + pass: false, + }, + empty: { + src: '{}', + pass: true, + }, + semicolon: { + src: '{;}', + pass: true, + }, + semicolons: { + src: '{;;}', + pass: true, + }, + decl: { + src: '{const c = 1;}', + pass: true, + }, + nested: { + src: '{ {} }', + pass: true, + }, +}; + +g.test('parse') + .desc('Test that compound statments parse') + .params(u => u.combine('stmt', keysOf(kTests))) + .fn(t => { + const code = ` +@vertex +fn vtx() -> @builtin(position) vec4f { + ${kTests[t.params.stmt].src} + return vec4f(1); +} + `; + t.expectCompileResult(kTests[t.params.stmt].pass, code); + }); From df460fa1ad87214fc74886cf223bebdd8993290c Mon Sep 17 00:00:00 2001 From: ShrekShao <5031596+shrekshao@users.noreply.github.com> Date: Mon, 29 Jan 2024 14:30:07 -0800 Subject: [PATCH 5/5] webgpu,api,operation,texture_view,write:format:* (#3333) --- .../api/operation/texture_view/write.spec.ts | 347 +++++++++++++++++- src/webgpu/util/shader.ts | 21 ++ 2 files changed, 365 insertions(+), 3 deletions(-) diff --git a/src/webgpu/api/operation/texture_view/write.spec.ts b/src/webgpu/api/operation/texture_view/write.spec.ts index 034012133479..b4ce6f4cec64 100644 --- a/src/webgpu/api/operation/texture_view/write.spec.ts +++ b/src/webgpu/api/operation/texture_view/write.spec.ts @@ -1,6 +1,9 @@ export const description = ` Test the result of writing textures through texture views with various options. +Reads value from a shader array, writes the value via various write methods. +Check the texture result with the expected texel view. + All x= every possible view write method: { - storage write {fragment, compute} - render pass store @@ -13,20 +16,358 @@ TODO: Write helper for this if not already available (see resource_init, buffer_ `; import { makeTestGroup } from '../../../../common/framework/test_group.js'; -import { GPUTest } from '../../../gpu_test.js'; +import { unreachable } from '../../../../common/util/util.js'; +import { + kRegularTextureFormats, + kTextureFormatInfo, + RegularTextureFormat, +} from '../../../format_info.js'; +import { GPUTest, TextureTestMixin } from '../../../gpu_test.js'; +import { kFullscreenQuadVertexShaderCode } from '../../../util/shader.js'; +import { TexelView } from '../../../util/texture/texel_view.js'; + +export const g = makeTestGroup(TextureTestMixin(GPUTest)); + +const kTextureViewWriteMethods = [ + 'storage-write-fragment', + 'storage-write-compute', + 'render-pass-store', + 'render-pass-resolve', +] as const; +type TextureViewWriteMethod = (typeof kTextureViewWriteMethods)[number]; + +// Src color values to read from a shader array. +const kColorsFloat = [ + { R: 1.0, G: 0.0, B: 0.0, A: 0.8 }, + { R: 0.0, G: 1.0, B: 0.0, A: 0.7 }, + { R: 0.0, G: 0.0, B: 0.0, A: 0.6 }, + { R: 0.0, G: 0.0, B: 0.0, A: 0.5 }, + { R: 1.0, G: 1.0, B: 1.0, A: 0.4 }, + { R: 0.7, G: 0.0, B: 0.0, A: 0.3 }, + { R: 0.0, G: 0.8, B: 0.0, A: 0.2 }, + { R: 0.0, G: 0.0, B: 0.9, A: 0.1 }, + { R: 0.1, G: 0.2, B: 0.0, A: 0.3 }, + { R: 0.4, G: 0.3, B: 0.6, A: 0.8 }, +]; + +function FloatToIntColor(c: number) { + return Math.floor(c * 100); +} + +const kColorsInt = kColorsFloat.map(c => { + return { + R: FloatToIntColor(c.R), + G: FloatToIntColor(c.G), + B: FloatToIntColor(c.B), + A: FloatToIntColor(c.A), + }; +}); -export const g = makeTestGroup(GPUTest); +const kTextureSize = 16; + +function writeTextureAndGetExpectedTexelView( + t: GPUTest, + method: TextureViewWriteMethod, + view: GPUTextureView, + format: RegularTextureFormat, + sampleCount: number +) { + const info = kTextureFormatInfo[format]; + const isFloatType = info.color.type === 'float' || info.color.type === 'unfilterable-float'; + const kColors = isFloatType ? kColorsFloat : kColorsInt; + const expectedTexelView = TexelView.fromTexelsAsColors( + format, + coords => { + const pixelPos = coords.y * kTextureSize + coords.x; + return kColors[pixelPos % kColors.length]; + }, + { clampToFormatRange: true } + ); + const vecType = isFloatType ? 'vec4f' : info.color.type === 'sint' ? 'vec4i' : 'vec4u'; + const kColorArrayShaderString = `array<${vecType}, ${kColors.length}>( + ${kColors.map(t => `${vecType}(${t.R}, ${t.G}, ${t.B}, ${t.A}) `).join(',')} + )`; + + switch (method) { + case 'storage-write-compute': + { + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ + code: ` + @group(0) @binding(0) var dst: texture_storage_2d<${format}, write>; + @compute @workgroup_size(1, 1) fn main( + @builtin(global_invocation_id) global_id: vec3, + ) { + const src = ${kColorArrayShaderString}; + let coord = vec2u(global_id.xy); + let idx = coord.x + coord.y * ${kTextureSize}; + textureStore(dst, coord, src[idx % ${kColors.length}]); + }`, + }), + entryPoint: 'main', + }, + }); + const commandEncoder = t.device.createCommandEncoder(); + const pass = commandEncoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup( + 0, + t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: view, + }, + ], + }) + ); + pass.dispatchWorkgroups(kTextureSize, kTextureSize); + pass.end(); + t.device.queue.submit([commandEncoder.finish()]); + } + break; + + case 'storage-write-fragment': + { + // Create a placeholder color attachment texture, + // The size of which equals that of format texture we are testing, + // so that we have the same number of fragments and texels. + const kPlaceholderTextureFormat = 'rgba8unorm'; + const placeholderTexture = t.trackForCleanup( + t.device.createTexture({ + format: kPlaceholderTextureFormat, + size: [kTextureSize, kTextureSize], + usage: GPUTextureUsage.RENDER_ATTACHMENT, + }) + ); + + const pipeline = t.device.createRenderPipeline({ + layout: 'auto', + vertex: { + module: t.device.createShaderModule({ + code: kFullscreenQuadVertexShaderCode, + }), + }, + fragment: { + module: t.device.createShaderModule({ + code: ` + @group(0) @binding(0) var dst: texture_storage_2d<${format}, write>; + @fragment fn main( + @builtin(position) fragCoord: vec4, + ) { + const src = ${kColorArrayShaderString}; + let coord = vec2u(fragCoord.xy); + let idx = coord.x + coord.y * ${kTextureSize}; + textureStore(dst, coord, src[idx % ${kColors.length}]); + }`, + }), + // Set writeMask to 0 as the fragment shader has no output. + targets: [ + { + format: kPlaceholderTextureFormat, + writeMask: 0, + }, + ], + }, + }); + const commandEncoder = t.device.createCommandEncoder(); + const pass = commandEncoder.beginRenderPass({ + colorAttachments: [ + { + view: placeholderTexture.createView(), + loadOp: 'clear', + storeOp: 'discard', + }, + ], + }); + pass.setPipeline(pipeline); + pass.setBindGroup( + 0, + t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: view, + }, + ], + }) + ); + pass.draw(6); + pass.end(); + t.device.queue.submit([commandEncoder.finish()]); + } + break; + + case 'render-pass-store': + case 'render-pass-resolve': + { + // Create a placeholder color attachment texture for the store target when tesing texture is used as resolve target. + const targetView = + method === 'render-pass-store' + ? view + : t + .trackForCleanup( + t.device.createTexture({ + format, + size: [kTextureSize, kTextureSize], + usage: GPUTextureUsage.RENDER_ATTACHMENT, + sampleCount: 4, + }) + ) + .createView(); + const resolveView = method === 'render-pass-store' ? undefined : view; + const multisampleCount = method === 'render-pass-store' ? sampleCount : 4; + + const pipeline = t.device.createRenderPipeline({ + layout: 'auto', + vertex: { + module: t.device.createShaderModule({ + code: kFullscreenQuadVertexShaderCode, + }), + }, + fragment: { + module: t.device.createShaderModule({ + code: ` + @fragment fn main( + @builtin(position) fragCoord: vec4, + ) -> @location(0) ${vecType} { + const src = ${kColorArrayShaderString}; + let coord = vec2u(fragCoord.xy); + let idx = coord.x + coord.y * ${kTextureSize}; + return src[idx % ${kColors.length}]; + }`, + }), + targets: [ + { + format, + }, + ], + }, + multisample: { + count: multisampleCount, + }, + }); + const commandEncoder = t.device.createCommandEncoder(); + const pass = commandEncoder.beginRenderPass({ + colorAttachments: [ + { + view: targetView, + resolveTarget: resolveView, + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + pass.setPipeline(pipeline); + pass.draw(6); + pass.end(); + t.device.queue.submit([commandEncoder.finish()]); + } + break; + default: + unreachable(); + } + + return expectedTexelView; +} g.test('format') .desc( `Views of every allowed format. +Read values from color array in the shader, and write it to the texture view via different write methods. + - x= every texture format - x= sampleCount {1, 4} if valid - x= every possible view write method (see above) + +TODO: Test sampleCount > 1 for 'render-pass-store' after extending copySinglePixelTextureToBufferUsingComputePass + to read multiple pixels from multisampled textures. [1] +TODO: Test rgb10a2uint when TexelRepresentation.numericRange is made per-component. [2] ` ) - .unimplemented(); + .params(u => + u // + .combine('method', kTextureViewWriteMethods) + .combine('format', kRegularTextureFormats) + .combine('sampleCount', [1, 4]) + .filter(({ format, method, sampleCount }) => { + const info = kTextureFormatInfo[format]; + + if (sampleCount > 1 && !info.multisample) { + return false; + } + + // [2] + if (format === 'rgb10a2uint') { + return false; + } + + switch (method) { + case 'storage-write-compute': + case 'storage-write-fragment': + return info.color?.storage && sampleCount === 1; + case 'render-pass-store': + // [1] + if (sampleCount > 1) { + return false; + } + return !!info.colorRender; + case 'render-pass-resolve': + return !!info.colorRender?.resolve && sampleCount === 1; + } + return true; + }) + ) + .beforeAllSubcases(t => { + const { format, method } = t.params; + t.skipIfTextureFormatNotSupported(format); + + switch (method) { + case 'storage-write-compute': + case 'storage-write-fragment': + // Still need to filter again for compat mode. + t.skipIfTextureFormatNotUsableAsStorageTexture(format); + break; + } + }) + .fn(t => { + const { format, method, sampleCount } = t.params; + + const usage = + GPUTextureUsage.COPY_SRC | + (method.includes('storage') + ? GPUTextureUsage.STORAGE_BINDING + : GPUTextureUsage.RENDER_ATTACHMENT); + + const texture = t.trackForCleanup( + t.device.createTexture({ + format, + usage, + size: [kTextureSize, kTextureSize], + sampleCount, + }) + ); + + const view = texture.createView(); + const expectedTexelView = writeTextureAndGetExpectedTexelView( + t, + method, + view, + format, + sampleCount + ); + + // [1] Use copySinglePixelTextureToBufferUsingComputePass to check multisampled texture. + t.expectTexelViewComparisonIsOkInTexture({ texture }, expectedTexelView, [ + kTextureSize, + kTextureSize, + ]); + }); g.test('dimension') .desc( diff --git a/src/webgpu/util/shader.ts b/src/webgpu/util/shader.ts index 2a09061527c5..6f291ec4d1c3 100644 --- a/src/webgpu/util/shader.ts +++ b/src/webgpu/util/shader.ts @@ -11,6 +11,27 @@ export const kDefaultFragmentShaderCode = ` return vec4(1.0, 1.0, 1.0, 1.0); }`; +// MAINTENANCE_TODO(#3344): deduplicate fullscreen quad shader code. +export const kFullscreenQuadVertexShaderCode = ` + struct VertexOutput { + @builtin(position) Position : vec4 + }; + + @vertex fn main(@builtin(vertex_index) VertexIndex : u32) -> VertexOutput { + var pos = array, 6>( + vec2( 1.0, 1.0), + vec2( 1.0, -1.0), + vec2(-1.0, -1.0), + vec2( 1.0, 1.0), + vec2(-1.0, -1.0), + vec2(-1.0, 1.0)); + + var output : VertexOutput; + output.Position = vec4(pos[VertexIndex], 0.0, 1.0); + return output; + } +`; + const kPlainTypeInfo = { i32: { suffix: '',