diff --git a/src/common/internal/file_loader.ts b/src/common/internal/file_loader.ts index b5e1b1a4460b..aae4b8799549 100644 --- a/src/common/internal/file_loader.ts +++ b/src/common/internal/file_loader.ts @@ -73,8 +73,9 @@ export abstract class TestFileLoader extends EventTarget { query: TestQuery, { subqueriesToExpand = [], + fullyExpandSubtrees = [], maxChunkTime = Infinity, - }: { subqueriesToExpand?: string[]; maxChunkTime?: number } = {} + }: { subqueriesToExpand?: string[]; fullyExpandSubtrees?: string[]; maxChunkTime?: number } = {} ): Promise { const tree = await loadTreeForQuery(this, query, { subqueriesToExpand: subqueriesToExpand.map(s => { @@ -82,6 +83,7 @@ export abstract class TestFileLoader extends EventTarget { assert(q.level >= 2, () => `subqueriesToExpand entries should not be multi-file:\n ${q}`); return q; }), + fullyExpandSubtrees: fullyExpandSubtrees.map(s => parseQuery(s)), maxChunkTime, }); this.dispatchEvent(new MessageEvent('finish')); diff --git a/src/common/internal/tree.ts b/src/common/internal/tree.ts index 594837059ca7..f2fad590373c 100644 --- a/src/common/internal/tree.ts +++ b/src/common/internal/tree.ts @@ -286,8 +286,9 @@ export async function loadTreeForQuery( queryToLoad: TestQuery, { subqueriesToExpand, + fullyExpandSubtrees = [], maxChunkTime = Infinity, - }: { subqueriesToExpand: TestQuery[]; maxChunkTime?: number } + }: { subqueriesToExpand: TestQuery[]; fullyExpandSubtrees?: TestQuery[]; maxChunkTime?: number } ): Promise { const suite = queryToLoad.suite; const specs = await loader.listing(suite); @@ -303,6 +304,10 @@ export async function loadTreeForQuery( // If toExpand == subquery, no expansion is needed (but it's still "seen"). if (ordering === Ordering.Equal) seenSubqueriesToExpand[i] = true; return ordering !== Ordering.StrictSubset; + }) && + fullyExpandSubtrees.every(toExpand => { + const ordering = compareQueries(toExpand, subquery); + return ordering === Ordering.Unordered; }); // L0 = suite-level, e.g. suite:* diff --git a/src/common/tools/gen_wpt_cts_html.ts b/src/common/tools/gen_wpt_cts_html.ts index e8161304e90f..90c7cd4ef4c5 100644 --- a/src/common/tools/gen_wpt_cts_html.ts +++ b/src/common/tools/gen_wpt_cts_html.ts @@ -71,6 +71,13 @@ interface ConfigJSON { /** The prefix to trim from every line of the expectations_file. */ prefix: string; }; + /** Expend all subtrees for provided queries */ + fullyExpandSubtrees?: { + file: string; + prefix: string; + }; + /*No long path assert */ + noLongPathAssert?: boolean; } interface Config { @@ -79,10 +86,15 @@ interface Config { template: string; maxChunkTimeMS: number; argumentsPrefixes: string[]; + noLongPathAssert: boolean; expectations?: { file: string; prefix: string; }; + fullyExpandSubtrees?: { + file: string; + prefix: string; + }; } let config: Config; @@ -101,6 +113,7 @@ let config: Config; template: path.resolve(jsonFileDir, configJSON.template), maxChunkTimeMS: configJSON.maxChunkTimeMS ?? Infinity, argumentsPrefixes: configJSON.argumentsPrefixes ?? ['?q='], + noLongPathAssert: configJSON.noLongPathAssert ?? false, }; if (configJSON.expectations) { config.expectations = { @@ -108,6 +121,12 @@ let config: Config; prefix: configJSON.expectations.prefix, }; } + if (configJSON.fullyExpandSubtrees) { + config.fullyExpandSubtrees = { + file: path.resolve(jsonFileDir, configJSON.fullyExpandSubtrees.file), + prefix: configJSON.fullyExpandSubtrees.prefix, + }; + } break; } case 4: @@ -130,6 +149,7 @@ let config: Config; suite, maxChunkTimeMS: Infinity, argumentsPrefixes: ['?q='], + noLongPathAssert: false, }; if (process.argv.length >= 7) { config.argumentsPrefixes = (await fs.readFile(argsPrefixesFile, 'utf8')) @@ -153,29 +173,16 @@ let config: Config; config.argumentsPrefixes.sort((a, b) => b.length - a.length); // Load expectations (if any) - let expectationLines = new Set(); - if (config.expectations) { - expectationLines = new Set( - (await fs.readFile(config.expectations.file, 'utf8')).split(/\r?\n/).filter(l => l.length) - ); - } + const expectations: Map = await loadQueryFile( + config.argumentsPrefixes, + config.expectations + ); - const expectations: Map = new Map(); - for (const prefix of config.argumentsPrefixes) { - expectations.set(prefix, []); - } - - expLoop: for (const exp of expectationLines) { - // Take each expectation for the longest prefix it matches. - for (const argsPrefix of config.argumentsPrefixes) { - const prefix = config.expectations!.prefix + argsPrefix; - if (exp.startsWith(prefix)) { - expectations.get(argsPrefix)!.push(exp.substring(prefix.length)); - continue expLoop; - } - } - console.log('note: ignored expectation: ' + exp); - } + // Load fullyExpandSubtrees queries (if any) + const fullyExpand: Map = await loadQueryFile( + config.argumentsPrefixes, + config.fullyExpandSubtrees + ); const loader = new DefaultTestFileLoader(); const lines = []; @@ -183,6 +190,7 @@ let config: Config; const rootQuery = new TestQueryMultiFile(config.suite, []); const tree = await loader.loadTree(rootQuery, { subqueriesToExpand: expectations.get(prefix), + fullyExpandSubtrees: fullyExpand.get(prefix), maxChunkTime: config.maxChunkTimeMS, }); @@ -199,19 +207,21 @@ let config: Config; alwaysExpandThroughLevel, })) { assert(query instanceof TestQueryMultiCase); - const queryString = query.toString(); - // Check for a safe-ish path length limit. Filename must be <= 255, and on Windows the whole - // path must be <= 259. Leave room for e.g.: - // 'c:\b\s\w\xxxxxxxx\layout-test-results\external\wpt\webgpu\cts_worker=0_q=...-actual.txt' - assert( - queryString.length < 185, - `Generated test variant would produce too-long -actual.txt filename. Possible solutions: + if (!config.noLongPathAssert) { + const queryString = query.toString(); + // Check for a safe-ish path length limit. Filename must be <= 255, and on Windows the whole + // path must be <= 259. Leave room for e.g.: + // 'c:\b\s\w\xxxxxxxx\layout-test-results\external\wpt\webgpu\cts_worker=0_q=...-actual.txt' + assert( + queryString.length < 185, + `Generated test variant would produce too-long -actual.txt filename. Possible solutions: - Reduce the length of the parts of the test query - Reduce the parameterization of the test - Make the test function faster and regenerate the listing_meta entry - Reduce the specificity of test expectations (if you're using them) ${queryString}` - ); + ); + } lines.push({ urlQueryString: prefix + query.toString(), // "?worker=0&q=..." @@ -232,6 +242,39 @@ ${queryString}` process.exit(1); }); +async function loadQueryFile( + argumentsPrefixes: string[], + queryFile?: { + file: string; + prefix: string; + } +): Promise> { + let lines = new Set(); + if (queryFile) { + lines = new Set( + (await fs.readFile(queryFile.file, 'utf8')).split(/\r?\n/).filter(l => l.length) + ); + } + + const result: Map = new Map(); + for (const prefix of argumentsPrefixes) { + result.set(prefix, []); + } + + expLoop: for (const exp of lines) { + // Take each expectation for the longest prefix it matches. + for (const argsPrefix of argumentsPrefixes) { + const prefix = queryFile!.prefix + argsPrefix; + if (exp.startsWith(prefix)) { + result.get(argsPrefix)!.push(exp.substring(prefix.length)); + continue expLoop; + } + } + console.log('note: ignored expectation: ' + exp); + } + return result; +} + async function generateFile( lines: Array<{ urlQueryString?: string; comment?: string } | undefined> ): Promise { diff --git a/src/webgpu/api/operation/storage_texture/read_write.spec.ts b/src/webgpu/api/operation/storage_texture/read_write.spec.ts new file mode 100644 index 000000000000..306031339183 --- /dev/null +++ b/src/webgpu/api/operation/storage_texture/read_write.spec.ts @@ -0,0 +1,329 @@ +export const description = ` +Tests for the behavior of read-write storage textures. + +TODO: +- Test resource usage transitions with read-write storage textures +- Test 1D and 3D textures +`; + +import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { assert, unreachable } from '../../../../common/util/util.js'; +import { kColorTextureFormats, kTextureFormatInfo } from '../../../format_info.js'; +import { GPUTest } from '../../../gpu_test.js'; +import { align } from '../../../util/math.js'; + +const kShaderStagesForReadWriteStorageTexture = ['fragment', 'compute'] as const; +type ShaderStageForReadWriteStorageTexture = + (typeof kShaderStagesForReadWriteStorageTexture)[number]; + +class F extends GPUTest { + GetInitialData(storageTexture: GPUTexture): ArrayBuffer { + const format = storageTexture.format; + const bytesPerBlock = kTextureFormatInfo[format].bytesPerBlock; + assert(bytesPerBlock !== undefined); + + const width = storageTexture.width; + const height = storageTexture.height; + const arrayLayers = storageTexture.depthOrArrayLayers; + const initialData = new ArrayBuffer(bytesPerBlock * width * height * arrayLayers); + const initialTypedData = this.GetTypedArrayBuffer(initialData, format); + for (let z = 0; z < arrayLayers; ++z) { + for (let y = 0; y < height; ++y) { + for (let x = 0; x < width; ++x) { + const index = z * width * height + y * width + x; + switch (format) { + case 'r32sint': + initialTypedData[index] = (index & 1 ? 1 : -1) * (2 * index + 1); + break; + case 'r32uint': + initialTypedData[index] = 2 * index + 1; + break; + case 'r32float': + initialTypedData[index] = (2 * index + 1) / 10.0; + break; + } + } + } + } + return initialData; + } + + GetTypedArrayBuffer(arrayBuffer: ArrayBuffer, format: GPUTextureFormat) { + switch (format) { + case 'r32sint': + return new Int32Array(arrayBuffer); + case 'r32uint': + return new Uint32Array(arrayBuffer); + case 'r32float': + return new Float32Array(arrayBuffer); + default: + unreachable(); + return new Uint8Array(arrayBuffer); + } + } + + GetExpectedData( + shaderStage: ShaderStageForReadWriteStorageTexture, + storageTexture: GPUTexture, + initialData: ArrayBuffer + ): ArrayBuffer { + const format = storageTexture.format; + const bytesPerBlock = kTextureFormatInfo[format].bytesPerBlock; + assert(bytesPerBlock !== undefined); + + const width = storageTexture.width; + const height = storageTexture.height; + const arrayLayers = storageTexture.depthOrArrayLayers; + const bytesPerRowAlignment = align(bytesPerBlock * width, 256); + const itemsPerRow = bytesPerRowAlignment / bytesPerBlock; + + const expectedData = new ArrayBuffer( + bytesPerRowAlignment * (height * arrayLayers - 1) + bytesPerBlock * width + ); + const expectedTypedData = this.GetTypedArrayBuffer(expectedData, format); + const initialTypedData = this.GetTypedArrayBuffer(initialData, format); + for (let z = 0; z < arrayLayers; ++z) { + for (let y = 0; y < height; ++y) { + for (let x = 0; x < width; ++x) { + const expectedIndex = z * itemsPerRow * height + y * itemsPerRow + x; + switch (shaderStage) { + case 'compute': { + // In the compute shader we flip the texture along the diagonal. + const initialIndex = + (arrayLayers - 1 - z) * width * height + (height - 1 - y) * width + (width - 1 - x); + expectedTypedData[expectedIndex] = initialTypedData[initialIndex]; + break; + } + case 'fragment': { + // In the fragment shader we double the original texel value of the read-write storage + // texture. + const initialIndex = z * width * height + y * width + x; + expectedTypedData[expectedIndex] = initialTypedData[initialIndex] * 2; + break; + } + } + } + } + } + return expectedData; + } + + RecordCommandsToTransform( + device: GPUDevice, + shaderStage: ShaderStageForReadWriteStorageTexture, + commandEncoder: GPUCommandEncoder, + rwTexture: GPUTexture + ) { + const isArray = rwTexture.depthOrArrayLayers > 1; + const textureDeclaration = ` + @group(0) @binding(0) var rwTexture: + texture_storage_2d${isArray ? '_array' : ''}<${rwTexture.format}, read_write>; + `; + + switch (shaderStage) { + case 'fragment': { + const vertexShader = ` + @vertex + fn main(@builtin(vertex_index) VertexIndex : u32) -> @builtin(position) vec4f { + var pos = array( + vec2f(-1.0, -1.0), + vec2f(-1.0, 1.0), + vec2f( 1.0, -1.0), + vec2f(-1.0, 1.0), + vec2f( 1.0, -1.0), + vec2f( 1.0, 1.0)); + return vec4f(pos[VertexIndex], 0.0, 1.0); + } + `; + const fragmentShader = ` + ${textureDeclaration} + @fragment + fn main(@builtin(position) fragCoord: vec4f) -> @location(0) vec4f { + let textureCoord = vec2u(fragCoord.xy); + + for (var z = 0; z < ${rwTexture.depthOrArrayLayers}; z++) { + let initialValue = textureLoad(rwTexture, textureCoord${isArray ? ', z' : ''}); + let outputValue = initialValue * 2; + textureStore(rwTexture, textureCoord, ${isArray ? 'z, ' : ''}outputValue); + } + + return vec4f(0.0, 1.0, 0.0, 1.0); + } + `; + const renderPipeline = device.createRenderPipeline({ + layout: 'auto', + vertex: { + module: device.createShaderModule({ + code: vertexShader, + }), + }, + fragment: { + module: device.createShaderModule({ + code: fragmentShader, + }), + targets: [ + { + format: 'rgba8unorm', + }, + ], + }, + primitive: { + topology: 'triangle-list', + }, + }); + + const bindGroup = device.createBindGroup({ + layout: renderPipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: rwTexture.createView(), + }, + ], + }); + + const dummyColorTexture = device.createTexture({ + size: [rwTexture.width, rwTexture.height, 1], + usage: GPUTextureUsage.RENDER_ATTACHMENT, + format: 'rgba8unorm', + }); + + const renderPassEncoder = commandEncoder.beginRenderPass({ + colorAttachments: [ + { + view: dummyColorTexture.createView(), + loadOp: 'clear', + clearValue: { r: 0, g: 0, b: 0, a: 0 }, + storeOp: 'store', + }, + ], + }); + renderPassEncoder.setPipeline(renderPipeline); + renderPassEncoder.setBindGroup(0, bindGroup); + renderPassEncoder.draw(6); + renderPassEncoder.end(); + break; + } + case 'compute': { + const computeShader = ` + ${textureDeclaration} + @compute + @workgroup_size(${rwTexture.width}, ${rwTexture.height}, ${rwTexture.depthOrArrayLayers}) + fn main(@builtin(local_invocation_id) invocationID: vec3u) { + let dimension = textureDimensions(rwTexture); + + let initialIndex = vec2u( + dimension.x - 1u - invocationID.x, dimension.y - 1u - invocationID.y); + let initialValue = textureLoad( + rwTexture, + initialIndex${isArray ? ', textureNumLayers(rwTexture) - 1u - invocationID.z' : ''}); + + textureBarrier(); + + textureStore(rwTexture, invocationID.xy, ${ + isArray ? 'invocationID.z, ' : '' + }initialValue); + }`; + + const computePipeline = device.createComputePipeline({ + compute: { + module: device.createShaderModule({ + code: computeShader, + }), + }, + layout: 'auto', + }); + const bindGroup = device.createBindGroup({ + layout: computePipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: rwTexture.createView(), + }, + ], + }); + const computePassEncoder = commandEncoder.beginComputePass(); + computePassEncoder.setPipeline(computePipeline); + computePassEncoder.setBindGroup(0, bindGroup); + computePassEncoder.dispatchWorkgroups(1); + computePassEncoder.end(); + break; + } + } + } +} + +export const g = makeTestGroup(F); + +g.test('basic') + .desc( + `The basic functionality tests for read-write storage textures. In the test we read data from + the read-write storage texture, do transforms and write the data back to the read-write storage + texture. textureBarrier() is also called in the tests using compute pipelines.` + ) + .params(u => + u + .combine('format', kColorTextureFormats) + .filter(p => kTextureFormatInfo[p.format].color?.readWriteStorage === true) + .combine('shaderStage', kShaderStagesForReadWriteStorageTexture) + .combine('arrayLayers', [1, 2] as const) + ) + .fn(t => { + const { format, shaderStage, arrayLayers } = t.params; + + const kWidth = 16; + const kHeight = 8; + const kTextureSize = [kWidth, kHeight, arrayLayers] as const; + const storageTexture = t.device.createTexture({ + format, + size: kTextureSize, + usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.COPY_DST | GPUTextureUsage.STORAGE_BINDING, + }); + + const bytesPerBlock = kTextureFormatInfo[format].bytesPerBlock; + const initialData = t.GetInitialData(storageTexture); + t.queue.writeTexture( + { texture: storageTexture }, + initialData, + { + bytesPerRow: bytesPerBlock * kWidth, + rowsPerImage: kHeight, + }, + kTextureSize + ); + + const commandEncoder = t.device.createCommandEncoder(); + + t.RecordCommandsToTransform(t.device, shaderStage, commandEncoder, storageTexture); + + const expectedData = t.GetExpectedData(shaderStage, storageTexture, initialData); + const readbackBuffer = t.device.createBuffer({ + size: expectedData.byteLength, + usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, + }); + const bytesPerRow = align(bytesPerBlock * kWidth, 256); + commandEncoder.copyTextureToBuffer( + { + texture: storageTexture, + }, + { + buffer: readbackBuffer, + bytesPerRow, + rowsPerImage: kHeight, + }, + kTextureSize + ); + t.queue.submit([commandEncoder.finish()]); + + switch (format) { + case 'r32sint': + t.expectGPUBufferValuesEqual(readbackBuffer, new Int32Array(expectedData)); + break; + case 'r32uint': + t.expectGPUBufferValuesEqual(readbackBuffer, new Uint32Array(expectedData)); + break; + case 'r32float': + t.expectGPUBufferValuesEqual(readbackBuffer, new Float32Array(expectedData)); + break; + } + }); diff --git a/src/webgpu/capability_info.ts b/src/webgpu/capability_info.ts index 1bd5d3b7c698..d26d93ca2e8d 100644 --- a/src/webgpu/capability_info.ts +++ b/src/webgpu/capability_info.ts @@ -821,6 +821,7 @@ export const kKnownWGSLLanguageFeatures = [ 'readonly_and_readwrite_storage_textures', 'packed_4x8_integer_dot_product', 'unrestricted_pointer_parameters', + 'pointer_composite_access', ] as const; export type WGSLLanguageFeature = (typeof kKnownWGSLLanguageFeatures)[number]; diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index e02364588779..c14c0912e793 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -198,6 +198,7 @@ "webgpu:api,operation,shader_module,compilation_info:getCompilationInfo_returns:*": { "subcaseMS": 0.284 }, "webgpu:api,operation,shader_module,compilation_info:line_number_and_position:*": { "subcaseMS": 1.867 }, "webgpu:api,operation,shader_module,compilation_info:offset_and_length:*": { "subcaseMS": 1.648 }, + "webgpu:api,operation,storage_texture,read_write:basic:*": { "subcaseMS": 5.000 }, "webgpu:api,operation,texture_view,format_reinterpretation:render_and_resolve_attachment:*": { "subcaseMS": 14.488 }, "webgpu:api,operation,texture_view,format_reinterpretation:texture_binding:*": { "subcaseMS": 17.225 }, "webgpu:api,operation,texture_view,read:aspect:*": { "subcaseMS": 0.601 }, diff --git a/src/webgpu/shader/execution/shader_io/compute_builtins.spec.ts b/src/webgpu/shader/execution/shader_io/compute_builtins.spec.ts index fcf3159c642c..a40b42633283 100644 --- a/src/webgpu/shader/execution/shader_io/compute_builtins.spec.ts +++ b/src/webgpu/shader/execution/shader_io/compute_builtins.spec.ts @@ -1,7 +1,6 @@ export const description = `Test compute shader builtin variables`; import { makeTestGroup } from '../../../../common/framework/test_group.js'; -import { iterRange } from '../../../../common/util/util.js'; import { GPUTest } from '../../../gpu_test.js'; export const g = makeTestGroup(GPUTest); @@ -98,17 +97,14 @@ g.test('inputs') // WGSL shader that stores every builtin value to a buffer, for every invocation in the grid. const wgsl = ` - struct S { - data : array + struct Outputs { + local_id: vec3u, + local_index: u32, + global_id: vec3u, + group_id: vec3u, + num_groups: vec3u, }; - struct V { - data : array> - }; - @group(0) @binding(0) var local_id_out : V; - @group(0) @binding(1) var local_index_out : S; - @group(0) @binding(2) var global_id_out : V; - @group(0) @binding(3) var group_id_out : V; - @group(0) @binding(4) var num_groups_out : V; + @group(0) @binding(0) var outputs : array; ${structures} @@ -122,11 +118,13 @@ g.test('inputs') ) { let group_index = ((${group_id}.z * ${num_groups}.y) + ${group_id}.y) * ${num_groups}.x + ${group_id}.x; let global_index = group_index * ${invocationsPerGroup}u + ${local_index}; - local_id_out.data[global_index] = ${local_id}; - local_index_out.data[global_index] = ${local_index}; - global_id_out.data[global_index] = ${global_id}; - group_id_out.data[global_index] = ${group_id}; - num_groups_out.data[global_index] = ${num_groups}; + var o: Outputs; + o.local_id = ${local_id}; + o.local_index = ${local_index}; + o.global_id = ${global_id}; + o.group_id = ${group_id}; + o.num_groups = ${num_groups}; + outputs[global_index] = o; } `; @@ -140,35 +138,24 @@ g.test('inputs') }, }); - // Helper to create a `size`-byte buffer with binding number `binding`. - function createBuffer(size: number, binding: number) { - const buffer = t.device.createBuffer({ - size, - usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, - }); - t.trackForCleanup(buffer); - - bindGroupEntries.push({ - binding, - resource: { - buffer, - }, - }); - - return buffer; - } + // Offsets are in u32 size units + const kLocalIdOffset = 0; + const kLocalIndexOffset = 3; + const kGlobalIdOffset = 4; + const kGroupIdOffset = 8; + const kNumGroupsOffset = 12; + const kOutputElementSize = 16; // Create the output buffers. - const bindGroupEntries: GPUBindGroupEntry[] = []; - const localIdBuffer = createBuffer(totalInvocations * 16, 0); - const localIndexBuffer = createBuffer(totalInvocations * 4, 1); - const globalIdBuffer = createBuffer(totalInvocations * 16, 2); - const groupIdBuffer = createBuffer(totalInvocations * 16, 3); - const numGroupsBuffer = createBuffer(totalInvocations * 16, 4); + const outputBuffer = t.device.createBuffer({ + size: totalInvocations * kOutputElementSize * 4, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + t.trackForCleanup(outputBuffer); const bindGroup = t.device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), - entries: bindGroupEntries, + entries: [{ binding: 0, resource: { buffer: outputBuffer } }], }); // Run the shader. @@ -204,11 +191,7 @@ g.test('inputs') // Helper to check that the vec3 value at each index of the provided `output` buffer // matches the expected value for that invocation, as generated by the `getBuiltinValue` // function. The `name` parameter is the builtin name, used for error messages. - const checkEachIndex = ( - output: Uint32Array, - name: string, - getBuiltinValue: (groupId: vec3, localId: vec3) => vec3 - ) => { + const checkEachIndex = (output: Uint32Array) => { // Loop over workgroups. for (let gz = 0; gz < t.params.numGroups.z; gz++) { for (let gy = 0; gy < t.params.numGroups.y; gy++) { @@ -220,30 +203,44 @@ g.test('inputs') const groupIndex = (gz * t.params.numGroups.y + gy) * t.params.numGroups.x + gx; const localIndex = (lz * t.params.groupSize.y + ly) * t.params.groupSize.x + lx; const globalIndex = groupIndex * invocationsPerGroup + localIndex; - const expected = getBuiltinValue( - { x: gx, y: gy, z: gz }, - { x: lx, y: ly, z: lz } - ); - if (output[globalIndex * 4 + 0] !== expected.x) { - return new Error( - `${name}.x failed at group(${gx},${gy},${gz}) local(${lx},${ly},${lz}))\n` + - ` expected: ${expected.x}\n` + - ` got: ${output[globalIndex * 4 + 0]}` - ); - } - if (output[globalIndex * 4 + 1] !== expected.y) { - return new Error( - `${name}.y failed at group(${gx},${gy},${gz}) local(${lx},${ly},${lz}))\n` + - ` expected: ${expected.y}\n` + - ` got: ${output[globalIndex * 4 + 1]}` + const globalOffset = globalIndex * kOutputElementSize; + + const expectEqual = (name: string, expected: number, actual: number) => { + if (actual !== expected) { + return new Error( + `${name} failed at group(${gx},${gy},${gz}) local(${lx},${ly},${lz}))\n` + + ` expected: ${expected}\n` + + ` got: ${actual}` + ); + } + return undefined; + }; + + const checkVec3Value = (name: string, fieldOffset: number, expected: vec3) => { + const offset = globalOffset + fieldOffset; + return ( + expectEqual(`${name}.x`, expected.x, output[offset + 0]) || + expectEqual(`${name}.y`, expected.y, output[offset + 1]) || + expectEqual(`${name}.z`, expected.z, output[offset + 2]) ); - } - if (output[globalIndex * 4 + 2] !== expected.z) { - return new Error( - `${name}.z failed at group(${gx},${gy},${gz}) local(${lx},${ly},${lz}))\n` + - ` expected: ${expected.z}\n` + - ` got: ${output[globalIndex * 4 + 2]}` + }; + + const error = + checkVec3Value('local_id', kLocalIdOffset, { x: lx, y: ly, z: lz }) || + checkVec3Value('global_id', kGlobalIdOffset, { + x: gx * t.params.groupSize.x + lx, + y: gy * t.params.groupSize.y + ly, + z: gz * t.params.groupSize.z + lz, + }) || + checkVec3Value('group_id', kGroupIdOffset, { x: gx, y: gy, z: gz }) || + checkVec3Value('num_groups', kNumGroupsOffset, t.params.numGroups) || + expectEqual( + 'local_index', + localIndex, + output[globalOffset + kLocalIndexOffset] ); + if (error) { + return error; } } } @@ -254,44 +251,8 @@ g.test('inputs') return undefined; }; - // Check @builtin(local_invocation_index) values. - t.expectGPUBufferValuesEqual( - localIndexBuffer, - new Uint32Array([...iterRange(totalInvocations, x => x % invocationsPerGroup)]) - ); - - // Check @builtin(local_invocation_id) values. - t.expectGPUBufferValuesPassCheck( - localIdBuffer, - outputData => checkEachIndex(outputData, 'local_invocation_id', (_, localId) => localId), - { type: Uint32Array, typedLength: totalInvocations * 4 } - ); - - // Check @builtin(global_invocation_id) values. - const getGlobalId = (groupId: vec3, localId: vec3) => { - return { - x: groupId.x * t.params.groupSize.x + localId.x, - y: groupId.y * t.params.groupSize.y + localId.y, - z: groupId.z * t.params.groupSize.z + localId.z, - }; - }; - t.expectGPUBufferValuesPassCheck( - globalIdBuffer, - outputData => checkEachIndex(outputData, 'global_invocation_id', getGlobalId), - { type: Uint32Array, typedLength: totalInvocations * 4 } - ); - - // Check @builtin(workgroup_id) values. - t.expectGPUBufferValuesPassCheck( - groupIdBuffer, - outputData => checkEachIndex(outputData, 'workgroup_id', (groupId, _) => groupId), - { type: Uint32Array, typedLength: totalInvocations * 4 } - ); - - // Check @builtin(num_workgroups) values. - t.expectGPUBufferValuesPassCheck( - numGroupsBuffer, - outputData => checkEachIndex(outputData, 'num_workgroups', () => t.params.numGroups), - { type: Uint32Array, typedLength: totalInvocations * 4 } - ); + t.expectGPUBufferValuesPassCheck(outputBuffer, outputData => checkEachIndex(outputData), { + type: Uint32Array, + typedLength: outputBuffer.size / 4, + }); }); diff --git a/src/webgpu/shader/validation/uniformity/uniformity.spec.ts b/src/webgpu/shader/validation/uniformity/uniformity.spec.ts index 41249e445d1c..aaba1f95aa54 100644 --- a/src/webgpu/shader/validation/uniformity/uniformity.spec.ts +++ b/src/webgpu/shader/validation/uniformity/uniformity.spec.ts @@ -367,7 +367,14 @@ function generatePointerCheck(check: string): string { } } -const kPointerCases = { +interface PointerCase { + code: string; + check: 'address' | 'contents'; + uniform: boolean | 'never'; + needs_deref_sugar?: boolean; +} + +const kPointerCases: Record = { address_uniform_literal: { code: `let ptr = &wg_array[0];`, check: `address`, @@ -585,6 +592,168 @@ const kPointerCases = { check: `contents`, uniform: false, }, + contents_lhs_ref_pointer_deref1: { + code: `*&func_scalar = uniform_value; + let test_val = func_scalar;`, + check: `contents`, + uniform: true, + }, + contents_lhs_ref_pointer_deref1a: { + code: `*&func_scalar = nonuniform_value; + let test_val = func_scalar;`, + check: `contents`, + uniform: false, + }, + contents_lhs_ref_pointer_deref2: { + code: `*&(func_array[nonuniform_value]) = uniform_value; + let test_val = func_array[0];`, + check: `contents`, + uniform: false, + }, + contents_lhs_ref_pointer_deref2a: { + code: `(func_array[nonuniform_value]) = uniform_value; + let test_val = func_array[0];`, + check: `contents`, + uniform: false, + }, + contents_lhs_ref_pointer_deref3: { + code: `*&(func_array[needs_uniform(uniform_value)]) = uniform_value; + let test_val = func_array[0];`, + check: `contents`, + uniform: true, + }, + contents_lhs_ref_pointer_deref3a: { + code: `*&(func_array[needs_uniform(nonuniform_value)]) = uniform_value; + let test_val = func_array[0];`, + check: `contents`, + uniform: 'never', + }, + contents_lhs_ref_pointer_deref4: { + code: `*&((*&(func_struct.x[uniform_value])).x[uniform_value].x[uniform_value]) = uniform_value; + let test_val = func_struct.x[0].x[0].x[0];`, + check: `contents`, + uniform: true, + }, + contents_lhs_ref_pointer_deref4a: { + code: `*&((*&(func_struct.x[uniform_value])).x[uniform_value].x[uniform_value]) = nonuniform_value; + let test_val = func_struct.x[0].x[0].x[0];`, + check: `contents`, + uniform: false, + }, + contents_lhs_ref_pointer_deref4b: { + code: `*&((*&(func_struct.x[uniform_value])).x[uniform_value].x[nonuniform_value]) = uniform_value; + let test_val = func_struct.x[0].x[0].x[0];`, + check: `contents`, + uniform: false, + }, + contents_lhs_ref_pointer_deref4c: { + code: `*&((*&(func_struct.x[uniform_value])).x[nonuniform_value]).x[uniform_value] = uniform_value; + let test_val = func_struct.x[0].x[0].x[0];`, + check: `contents`, + uniform: false, + }, + contents_lhs_ref_pointer_deref4d: { + code: `*&((*&(func_struct.x[nonuniform_value])).x[uniform_value].x)[uniform_value] = uniform_value; + let test_val = func_struct.x[0].x[0].x[0];`, + check: `contents`, + uniform: false, + }, + contents_lhs_ref_pointer_deref4e: { + code: `*&((*&(func_struct.x[uniform_value])).x[needs_uniform(nonuniform_value)].x[uniform_value]) = uniform_value; + let test_val = func_struct.x[0].x[0].x[0];`, + check: `contents`, + uniform: 'never', + }, + + // The following cases require the 'pointer_composite_access' language feature. + contents_lhs_pointer_deref2: { + code: `(&func_array)[uniform_value] = uniform_value; + let test_val = func_array[0];`, + check: `contents`, + uniform: true, + needs_deref_sugar: true, + }, + contents_lhs_pointer_deref2a: { + code: `(&func_array)[nonuniform_value] = uniform_value; + let test_val = func_array[0];`, + check: `contents`, + uniform: false, + needs_deref_sugar: true, + }, + contents_lhs_pointer_deref3: { + code: `(&func_array)[needs_uniform(uniform_value)] = uniform_value; + let test_val = func_array[0];`, + check: `contents`, + uniform: true, + needs_deref_sugar: true, + }, + contents_lhs_pointer_deref3a: { + code: `(&func_array)[needs_uniform(nonuniform_value)] = uniform_value; + let test_val = func_array[0];`, + check: `contents`, + uniform: 'never', + needs_deref_sugar: true, + }, + contents_lhs_pointer_deref4: { + code: `(&((&(func_struct.x[uniform_value])).x[uniform_value]).x)[uniform_value] = uniform_value; + let test_val = func_struct.x[0].x[0].x[0];`, + check: `contents`, + uniform: true, + needs_deref_sugar: true, + }, + contents_lhs_pointer_deref4a: { + code: `(&((&(func_struct.x[uniform_value])).x[uniform_value]).x)[uniform_value] = nonuniform_value; + let test_val = func_struct.x[0].x[0].x[0];`, + check: `contents`, + uniform: false, + needs_deref_sugar: true, + }, + contents_lhs_pointer_deref4b: { + code: `(&((&(func_struct.x[uniform_value])).x)[uniform_value]).x[nonuniform_value] = uniform_value; + let test_val = func_struct.x[0].x[0].x[0];`, + check: `contents`, + uniform: false, + needs_deref_sugar: true, + }, + contents_lhs_pointer_deref4c: { + code: `(&((&(func_struct.x[uniform_value])).x[nonuniform_value]).x)[uniform_value] = uniform_value; + let test_val = func_struct.x[0].x[0].x[0];`, + check: `contents`, + uniform: false, + needs_deref_sugar: true, + }, + contents_lhs_pointer_deref4d: { + code: `(&((&(func_struct.x[nonuniform_value])).x[uniform_value]).x)[uniform_value] = uniform_value; + let test_val = func_struct.x[0].x[0].x[0];`, + check: `contents`, + uniform: false, + needs_deref_sugar: true, + }, + contents_lhs_pointer_deref4e: { + code: `(&((&(func_struct.x[uniform_value])).x)[needs_uniform(nonuniform_value)].x[uniform_value]) = uniform_value; + let test_val = func_struct.x[0].x[0].x[0];`, + check: `contents`, + uniform: 'never', + needs_deref_sugar: true, + }, + contents_rhs_pointer_deref1: { + code: `let test_val = (&func_array)[uniform_value];`, + check: `contents`, + uniform: true, + needs_deref_sugar: true, + }, + contents_rhs_pointer_deref1a: { + code: `let test_val = (&func_array)[nonuniform_value];`, + check: `contents`, + uniform: false, + needs_deref_sugar: true, + }, + contents_rhs_pointer_deref2: { + code: `let test_val = (&func_array)[needs_uniform(nonuniform_value)];`, + check: `contents`, + uniform: `never`, + needs_deref_sugar: true, + }, }; g.test('pointers') @@ -612,6 +781,13 @@ var uniform_value : u32; @group(0) @binding(1) var nonuniform_value : u32; +fn needs_uniform(val : u32) -> u32{ + if val == 0 { + workgroupBarrier(); + } + return val; +} + @compute @workgroup_size(16, 1, 1) fn main(@builtin(local_invocation_id) lid : vec3, @builtin(global_invocation_id) gid : vec3) { @@ -627,11 +803,16 @@ fn main(@builtin(local_invocation_id) lid : vec3, ` ${generatePointerCheck(testcase.check)} }`; - if (!testcase.uniform) { + + if (testcase.needs_deref_sugar === true) { + t.skipIfLanguageFeatureNotSupported('pointer_composite_access'); + } + // Explicitly check false to distinguish from never. + if (testcase.uniform === false) { const without_check = code + `}\n`; t.expectCompileResult(true, without_check); } - t.expectCompileResult(testcase.uniform, with_check); + t.expectCompileResult(testcase.uniform === true, with_check); }); function expectedUniformity(uniform: string, init: string): boolean {