diff --git a/package-lock.json b/package-lock.json index cc4e97a2ce64..cd59eb95b21c 100644 --- a/package-lock.json +++ b/package-lock.json @@ -24,7 +24,7 @@ "@types/w3c-image-capture": "^1.0.10", "@typescript-eslint/eslint-plugin": "^6.9.1", "@typescript-eslint/parser": "^6.9.1", - "@webgpu/types": "^0.1.49", + "@webgpu/types": "^0.1.51", "ansi-colors": "4.1.3", "babel-plugin-add-header-comment": "^1.0.3", "babel-plugin-const-enum": "^1.2.0", @@ -1539,10 +1539,17 @@ "dev": true }, "node_modules/@webgpu/types": { - "version": "0.1.49", - "resolved": "https://registry.npmjs.org/@webgpu/types/-/types-0.1.49.tgz", - "integrity": "sha512-NMmS8/DofhH/IFeW+876XrHVWel+J/vdcFCHLDqeJgkH9x0DeiwjVd8LcBdaxdG/T7Rf8VUAYsA8X1efMzLjRQ==", - "dev": true +<<<<<<< HEAD + "version": "0.1.51", + "resolved": "https://registry.npmjs.org/@webgpu/types/-/types-0.1.51.tgz", + "integrity": "sha512-ktR3u64NPjwIViNCck+z9QeyN0iPkQCUOQ07ZCV1RzlkfP+olLTeEZ95O1QHS+v4w9vJeY9xj/uJuSphsHy5rQ==", +======= + "version": "0.1.50", + "resolved": "https://registry.npmjs.org/@webgpu/types/-/types-0.1.50.tgz", + "integrity": "sha512-GjG3CQV7SyWk/lEXqFPuKchRPHIBbD317Gj8NUqqB+UOnQlOYtjGLCTRIWzO9Ta698LVzlBCSE9XKqBSWpIDmg==", +>>>>>>> 8a80203467ffcb3f87bb6a7f74655a6e1043d051 + "dev": true, + "license": "BSD-3-Clause" }, "node_modules/abbrev": { "version": "1.1.1", @@ -10076,9 +10083,15 @@ "dev": true }, "@webgpu/types": { - "version": "0.1.49", - "resolved": "https://registry.npmjs.org/@webgpu/types/-/types-0.1.49.tgz", - "integrity": "sha512-NMmS8/DofhH/IFeW+876XrHVWel+J/vdcFCHLDqeJgkH9x0DeiwjVd8LcBdaxdG/T7Rf8VUAYsA8X1efMzLjRQ==", +<<<<<<< HEAD + "version": "0.1.51", + "resolved": "https://registry.npmjs.org/@webgpu/types/-/types-0.1.51.tgz", + "integrity": "sha512-ktR3u64NPjwIViNCck+z9QeyN0iPkQCUOQ07ZCV1RzlkfP+olLTeEZ95O1QHS+v4w9vJeY9xj/uJuSphsHy5rQ==", +======= + "version": "0.1.50", + "resolved": "https://registry.npmjs.org/@webgpu/types/-/types-0.1.50.tgz", + "integrity": "sha512-GjG3CQV7SyWk/lEXqFPuKchRPHIBbD317Gj8NUqqB+UOnQlOYtjGLCTRIWzO9Ta698LVzlBCSE9XKqBSWpIDmg==", +>>>>>>> 8a80203467ffcb3f87bb6a7f74655a6e1043d051 "dev": true }, "abbrev": { diff --git a/package.json b/package.json index 3ef62315db25..cef3de27ed1d 100644 --- a/package.json +++ b/package.json @@ -50,7 +50,7 @@ "@types/w3c-image-capture": "^1.0.10", "@typescript-eslint/eslint-plugin": "^6.9.1", "@typescript-eslint/parser": "^6.9.1", - "@webgpu/types": "^0.1.49", + "@webgpu/types": "^0.1.51", "ansi-colors": "4.1.3", "babel-plugin-add-header-comment": "^1.0.3", "babel-plugin-const-enum": "^1.2.0", diff --git a/src/webgpu/api/operation/adapter/info.spec.ts b/src/webgpu/api/operation/adapter/info.spec.ts index ee5c4e86c57f..4fef9a56ddf7 100644 --- a/src/webgpu/api/operation/adapter/info.spec.ts +++ b/src/webgpu/api/operation/adapter/info.spec.ts @@ -1,11 +1,12 @@ export const description = ` -Tests GPUAdapter.info members formatting. +Tests for GPUAdapterInfo. `; import { Fixture } from '../../../../common/framework/fixture.js'; import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { keysOf } from '../../../../common/util/data_tables.js'; import { getGPU } from '../../../../common/util/navigator_gpu.js'; -import { assert } from '../../../../common/util/util.js'; +import { assert, objectEquals } from '../../../../common/util/util.js'; export const g = makeTestGroup(Fixture); @@ -39,3 +40,99 @@ g.test('adapter_info') `adapterInfo.device should be a normalized identifier. But it's '${adapterInfo.device}'` ); }); + +g.test('same_object') + .desc( + ` +GPUAdapter.info and GPUDevice.adapterInfo provide the same object each time they're accessed, +but different objects from one another.` + ) + .fn(async t => { + const gpu = getGPU(t.rec); + const adapter = await gpu.requestAdapter(); + assert(adapter !== null); + assert(adapter.info instanceof GPUAdapterInfo); + + const adapterInfo1 = adapter.info; + const adapterInfo2 = adapter.info; + t.expect(adapterInfo1 === adapterInfo2, 'adapter.info should obey [SameObject]'); + + const device = await t.requestDeviceTracked(adapter); + assert(device !== null); + assert(device.adapterInfo instanceof GPUAdapterInfo); + + const deviceAdapterInfo1 = device.adapterInfo; + const deviceAdapterInfo2 = device.adapterInfo; + t.expect( + deviceAdapterInfo1 === deviceAdapterInfo2, + 'device.adapterInfo should obey [SameObject]' + ); + + t.expect( + adapter.info !== device.adapterInfo, + 'adapter.info and device.adapterInfo should NOT return the same object' + ); + }); + +g.test('device_matches_adapter') + .desc( + ` +Test that GPUDevice.adapterInfo matches GPUAdapter.info. Cases access the members in +different orders to make sure that they are consistent regardless of the access order.` + ) + .paramsSubcasesOnly(u => + u.combine('testDeviceFirst', [true, false]).combine('testMembersFirst', [true, false]) + ) + .fn(async t => { + const { testDeviceFirst, testMembersFirst } = t.params; + + const gpu = getGPU(t.rec); + const adapter = await gpu.requestAdapter(); + assert(adapter !== null); + + const device = await t.requestDeviceTracked(adapter); + assert(device !== null); + + const deviceInfo: unknown[] = []; + const adapterInfo: unknown[] = []; + + const kGPUAdapterInfoKeys = keysOf(GPUAdapterInfo.prototype); + if (testMembersFirst) { + if (testDeviceFirst) { + assert(device.adapterInfo instanceof GPUAdapterInfo); + for (const k of kGPUAdapterInfoKeys) { + deviceInfo.push(device.adapterInfo[k]); + } + assert(adapter.info instanceof GPUAdapterInfo); + for (const k of kGPUAdapterInfoKeys) { + adapterInfo.push(adapter.info[k]); + } + } else { + assert(adapter.info instanceof GPUAdapterInfo); + for (const k of kGPUAdapterInfoKeys) { + adapterInfo.push(adapter.info[k]); + } + assert(device.adapterInfo instanceof GPUAdapterInfo); + for (const k of kGPUAdapterInfoKeys) { + deviceInfo.push(device.adapterInfo[k]); + } + } + } else { + if (testDeviceFirst) { + assert(device.adapterInfo instanceof GPUAdapterInfo); + assert(adapter.info instanceof GPUAdapterInfo); + for (const k of kGPUAdapterInfoKeys) { + deviceInfo.push(device.adapterInfo[k]); + adapterInfo.push(adapter.info[k]); + } + } else { + assert(adapter.info instanceof GPUAdapterInfo); + assert(device.adapterInfo instanceof GPUAdapterInfo); + for (const k of kGPUAdapterInfoKeys) { + adapterInfo.push(adapter.info[k]); + deviceInfo.push(device.adapterInfo[k]); + } + } + t.expect(objectEquals(deviceInfo, adapterInfo)); + } + }); diff --git a/src/webgpu/api/operation/command_buffer/image_copy.spec.ts b/src/webgpu/api/operation/command_buffer/image_copy.spec.ts index b27ecbbb91d0..73f1ff751b4c 100644 --- a/src/webgpu/api/operation/command_buffer/image_copy.spec.ts +++ b/src/webgpu/api/operation/command_buffer/image_copy.spec.ts @@ -126,7 +126,7 @@ const altDataGenerator = new DataArrayGenerator(); class ImageCopyTest extends TextureTestMixin(GPUTest) { /** - * This is used for testing passing undefined members of `GPUImageDataLayout` instead of actual + * This is used for testing passing undefined members of `GPUTexelCopyBufferLayout` instead of actual * values where possible. Passing arguments as values and not as objects so that they are passed * by copy and not by reference. */ @@ -135,7 +135,7 @@ class ImageCopyTest extends TextureTestMixin(GPUTest) { rowsPerImage: number | undefined, bytesPerRow: number | undefined, changeBeforePass: ChangeBeforePass - ): GPUImageDataLayout { + ): GPUTexelCopyBufferLayout { if (changeBeforePass === 'undefined') { if (offset === 0) { offset = undefined; @@ -151,7 +151,7 @@ class ImageCopyTest extends TextureTestMixin(GPUTest) { } /** - * This is used for testing passing undefined members of `GPUImageCopyTexture` instead of actual + * This is used for testing passing undefined members of `GPUTexelCopyTextureInfo` instead of actual * values where possible and also for testing passing the origin as `[number, number, number]`. * Passing arguments as values and not as objects so that they are passed by copy and not by * reference. @@ -163,7 +163,7 @@ class ImageCopyTest extends TextureTestMixin(GPUTest) { origin_z: number | undefined, mipLevel: number | undefined, changeBeforePass: ChangeBeforePass - ): GPUImageCopyTexture { + ): GPUTexelCopyTextureInfo { let origin: GPUOrigin3D | undefined = { x: origin_x, y: origin_y, z: origin_z }; if (changeBeforePass === 'undefined') { @@ -225,7 +225,7 @@ class ImageCopyTest extends TextureTestMixin(GPUTest) { buffer: GPUBuffer, format: ColorTextureFormat, size: Required, - dataLayout: Required + dataLayout: Required ) { if (isCompressedTextureFormat(format)) { this.expectGPUBufferValuesEqual(buffer, expected); @@ -323,7 +323,7 @@ class ImageCopyTest extends TextureTestMixin(GPUTest) { /** Run a CopyT2B command with appropriate arguments corresponding to `ChangeBeforePass` */ copyTextureToBufferWithAppliedArguments( buffer: GPUBuffer, - { offset, rowsPerImage, bytesPerRow }: Required, + { offset, rowsPerImage, bytesPerRow }: Required, { width, height, depthOrArrayLayers }: Required, { texture, mipLevel, origin }: TextureCopyViewWithRequiredOrigin, changeBeforePass: ChangeBeforePass @@ -363,7 +363,7 @@ class ImageCopyTest extends TextureTestMixin(GPUTest) { /** Put data into a part of the texture with an appropriate method. */ uploadLinearTextureDataToTextureSubBox( textureCopyView: TextureCopyViewWithRequiredOrigin, - textureDataLayout: GPUImageDataLayout & { bytesPerRow: number }, + textureDataLayout: GPUTexelCopyBufferLayout & { bytesPerRow: number }, copySize: Required, partialData: Uint8Array, method: InitMethod, @@ -432,7 +432,7 @@ class ImageCopyTest extends TextureTestMixin(GPUTest) { copySize: Required, format: ColorTextureFormat, expected: Uint8Array, - expectedDataLayout: Required + expectedDataLayout: Required ): void { const size = [ actualTexture.width, @@ -507,7 +507,7 @@ class ImageCopyTest extends TextureTestMixin(GPUTest) { checkSize: Required, format: ColorTextureFormat, expected: Uint8Array, - expectedDataLayout: Required, + expectedDataLayout: Required, changeBeforePass: ChangeBeforePass = 'none' ): void { // The alignment is necessary because we need to copy and map data from this buffer. @@ -573,7 +573,7 @@ class ImageCopyTest extends TextureTestMixin(GPUTest) { copyWholeTextureToBufferAndCheckContentsWithUpdatedData( { texture, mipLevel, origin }: TextureCopyViewWithRequiredOrigin, fullTextureCopyLayout: TextureCopyLayout, - texturePartialDataLayout: Required, + texturePartialDataLayout: Required, copySize: Required, format: ColorTextureFormat, fullData: GPUBuffer, @@ -632,7 +632,7 @@ class ImageCopyTest extends TextureTestMixin(GPUTest) { checkMethod, changeBeforePass = 'none', }: { - textureDataLayout: Required; + textureDataLayout: Required; copySize: Required; dataSize: number; mipLevel?: number; diff --git a/src/webgpu/api/operation/vertex_state/correctness.spec.ts b/src/webgpu/api/operation/vertex_state/correctness.spec.ts index 7af06e150616..1e023b09fb62 100644 --- a/src/webgpu/api/operation/vertex_state/correctness.spec.ts +++ b/src/webgpu/api/operation/vertex_state/correctness.spec.ts @@ -427,25 +427,51 @@ struct VSOutputs { case 'unorm': { if (formatInfo.bytesPerComponent === 'packed') { - assert(format === 'unorm10-10-10-2'); // This is the only packed format for now. assert(bitSize === 0); - /* prettier-ignore */ - const data = [ - [ 0, 0, 0, 0], - [1023, 1023, 1023, 3], - [ 243, 567, 765, 2], - ]; - const vertexData = new Uint32Array(data.map(makeRgb10a2)).buffer; - const expectedData = new Float32Array(data.flat().map(normalizeRgb10a2)).buffer; - - return { - shaderBaseType: 'f32', - testComponentCount: data.flat().length, - expectedData, - vertexData, - floatTolerance: 0.1 / 1023, - }; + switch (format) { + case 'unorm10-10-10-2': { + /* prettier-ignore */ + const data = [ + [ 0, 0, 0, 0], + [1023, 1023, 1023, 3], + [ 243, 567, 765, 2], + ]; + const vertexData = new Uint32Array(data.map(makeRgb10a2)).buffer; + const expectedData = new Float32Array(data.flat().map(normalizeRgb10a2)).buffer; + + return { + shaderBaseType: 'f32', + testComponentCount: data.flat().length, + expectedData, + vertexData, + floatTolerance: 0.1 / 1023, + }; + } + + case 'unorm8x4-bgra': { + const data = [42, 0, 1, 2, 3, 4, 128, 255]; + const vertexData = new Uint8Array(data).buffer; + const expectedData = new Float32Array( + data.map(v => normalizedIntegerAsFloat(v, 8, false)) + ); + + for (let i = 0; i + 2 < expectedData.length; i += 4) { + const r = expectedData[i + 0]; + const b = expectedData[i + 2]; + expectedData[i + 0] = b; + expectedData[i + 2] = r; + } + + return { + shaderBaseType: 'f32', + testComponentCount: data.length, + expectedData: expectedData.buffer, + vertexData, + floatTolerance: 0.1 / 255, + }; + } + } } /* prettier-ignore */ diff --git a/src/webgpu/api/validation/capability_checks/limits/maxTextureDimension2D.spec.ts b/src/webgpu/api/validation/capability_checks/limits/maxTextureDimension2D.spec.ts index d67b165a0fd0..cfe9fe18a4a9 100644 --- a/src/webgpu/api/validation/capability_checks/limits/maxTextureDimension2D.spec.ts +++ b/src/webgpu/api/validation/capability_checks/limits/maxTextureDimension2D.spec.ts @@ -59,7 +59,7 @@ g.test('configure,at_over') // This should not fail, even if the size is too large but it might fail // if we're in a worker and HTMLCanvasElement does not exist. - const canvas = createCanvas(t, canvasType, size[0], size[1])!; + const canvas = createCanvas(t, canvasType, size[0], size[1]); if (canvas) { const context = canvas.getContext('webgpu') as GPUCanvasContext; t.expect(!!context, 'should not fail to create context even if size is too large'); @@ -96,7 +96,7 @@ g.test('getCurrentTexture,at_over') // Start with a small size so configure will succeed. // This should not fail, even if the size is too large but it might fail // if we're in a worker and HTMLCanvasElement does not exist. - const canvas = createCanvas(t, canvasType, 1, 1)!; + const canvas = createCanvas(t, canvasType, 1, 1); if (canvas) { const context = canvas.getContext('webgpu') as GPUCanvasContext; t.expect(!!context, 'should not fail to create context even if size is too large'); diff --git a/src/webgpu/api/validation/compute_pipeline.spec.ts b/src/webgpu/api/validation/compute_pipeline.spec.ts index 790f25897a94..704f09fc2475 100644 --- a/src/webgpu/api/validation/compute_pipeline.spec.ts +++ b/src/webgpu/api/validation/compute_pipeline.spec.ts @@ -718,6 +718,7 @@ g.test('resource_compatibility') !t.hasLanguageFeature('readonly_and_readwrite_storage_textures'), 'Storage textures require language feature' ); + t.skipIfTextureViewDimensionNotSupported(wgslResource.texture?.viewDimension); const layout = t.device.createPipelineLayout({ bindGroupLayouts: [ diff --git a/src/webgpu/api/validation/encoding/cmds/copyTextureToTexture.spec.ts b/src/webgpu/api/validation/encoding/cmds/copyTextureToTexture.spec.ts index c49261c0a055..f29ca3d9057c 100644 --- a/src/webgpu/api/validation/encoding/cmds/copyTextureToTexture.spec.ts +++ b/src/webgpu/api/validation/encoding/cmds/copyTextureToTexture.spec.ts @@ -19,8 +19,8 @@ import { ValidationTest } from '../../validation_test.js'; class F extends ValidationTest { TestCopyTextureToTexture( - source: GPUImageCopyTexture, - destination: GPUImageCopyTexture, + source: GPUTexelCopyTextureInfo, + destination: GPUTexelCopyTextureInfo, copySize: GPUExtent3D, expectation: 'Success' | 'FinishError' | 'SubmitError' ): void { @@ -691,7 +691,7 @@ TODO: Extend to 1D and 3D textures.` g.test('copy_aspects') .desc( ` -Test the validations on the member 'aspect' of GPUImageCopyTexture in CopyTextureToTexture(). +Test the validations on the member 'aspect' of GPUTexelCopyTextureInfo in CopyTextureToTexture(). - for all the color and depth-stencil formats: the texture copy aspects must be both 'all'. - for all the depth-only formats: the texture copy aspects must be either 'all' or 'depth-only'. - for all the stencil-only formats: the texture copy aspects must be either 'all' or 'stencil-only'. diff --git a/src/webgpu/api/validation/image_copy/buffer_texture_copies.spec.ts b/src/webgpu/api/validation/image_copy/buffer_texture_copies.spec.ts index 937861cea03e..d417f23137fa 100644 --- a/src/webgpu/api/validation/image_copy/buffer_texture_copies.spec.ts +++ b/src/webgpu/api/validation/image_copy/buffer_texture_copies.spec.ts @@ -18,8 +18,8 @@ import { ValidationTest } from '../validation_test.js'; class ImageCopyTest extends ValidationTest { testCopyBufferToTexture( - source: GPUImageCopyBuffer, - destination: GPUImageCopyTexture, + source: GPUTexelCopyBufferInfo, + destination: GPUTexelCopyTextureInfo, copySize: GPUExtent3DStrict, isSuccess: boolean ): void { @@ -29,8 +29,8 @@ class ImageCopyTest extends ValidationTest { } testCopyTextureToBuffer( - source: GPUImageCopyTexture, - destination: GPUImageCopyBuffer, + source: GPUTexelCopyTextureInfo, + destination: GPUTexelCopyBufferInfo, copySize: GPUExtent3DStrict, isSuccess: boolean ): void { @@ -40,9 +40,9 @@ class ImageCopyTest extends ValidationTest { } testWriteTexture( - destination: GPUImageCopyTexture, + destination: GPUTexelCopyTextureInfo, uploadData: Uint8Array, - dataLayout: GPUImageDataLayout, + dataLayout: GPUTexelCopyBufferLayout, copySize: GPUExtent3DStrict, isSuccess: boolean ): void { diff --git a/src/webgpu/api/validation/image_copy/image_copy.ts b/src/webgpu/api/validation/image_copy/image_copy.ts index 1a86fac68794..42c40c42c536 100644 --- a/src/webgpu/api/validation/image_copy/image_copy.ts +++ b/src/webgpu/api/validation/image_copy/image_copy.ts @@ -11,8 +11,8 @@ import { ValidationTest } from '../validation_test.js'; export class ImageCopyTest extends ValidationTest { testRun( - textureCopyView: GPUImageCopyTexture, - textureDataLayout: GPUImageDataLayout, + textureCopyView: GPUTexelCopyTextureInfo, + textureDataLayout: GPUTexelCopyBufferLayout, size: GPUExtent3D, { method, @@ -122,7 +122,7 @@ export class ImageCopyTest extends ValidationTest { testBuffer( buffer: GPUBuffer, texture: GPUTexture, - textureDataLayout: GPUImageDataLayout, + textureDataLayout: GPUTexelCopyBufferLayout, size: GPUExtent3D, { method, diff --git a/src/webgpu/api/validation/non_filterable_texture.spec.ts b/src/webgpu/api/validation/non_filterable_texture.spec.ts new file mode 100644 index 000000000000..11057c397704 --- /dev/null +++ b/src/webgpu/api/validation/non_filterable_texture.spec.ts @@ -0,0 +1,126 @@ +export const description = ` +Tests that non-filterable textures used with filtering samplers generate a validation error. +`; + +import { makeTestGroup } from '../../../common/framework/test_group.js'; +import { keysOf } from '../../../common/util/data_tables.js'; + +import { ValidationTest } from './validation_test.js'; + +const kNonFilterableCaseInfo: Record = { + sint: { type: 'i32', component: '0,' }, + uint: { type: 'u32', component: '0,' }, + float: { type: 'f32', component: '0,' }, // no error for f32 + 'unfilterable-float': { type: 'f32', component: '0,' }, // no error for f32 + depth: { type: 'depth', component: '' }, +}; +const kNonFilterableCases = keysOf(kNonFilterableCaseInfo); + +export const g = makeTestGroup(ValidationTest); + +g.test('non_filterable_texture_with_filtering_sampler') + .desc( + 'test that createXXXPipeline generates a validation error if a depth/u32/i32 texture binding is used with a filtering sampler binding' + ) + .params(u => + u + .combine('pipeline', ['compute', 'render']) + .combine('async', [true, false] as const) + .combine('sampleType', kNonFilterableCases) + .combine('viewDimension', ['2d', '2d-array', 'cube', 'cube-array'] as const) + .combine('sameGroup', [true, false] as const) + ) + .beforeAllSubcases(t => t.skipIfTextureViewDimensionNotSupported(t.params.viewDimension)) + .fn(t => { + const { device } = t; + const { pipeline, async, sampleType, viewDimension, sameGroup } = t.params; + const { type, component } = kNonFilterableCaseInfo[sampleType]; + + const coord = viewDimension.startsWith('2d') ? 'vec2f(0)' : 'vec3f(0)'; + const dimensionSuffix = viewDimension.replace('-', '_'); + const textureType = + type === 'depth' ? `texture_depth_${dimensionSuffix}` : `texture_${dimensionSuffix}<${type}>`; + const layer = viewDimension.endsWith('-array') ? ', 0' : ''; + + const groupNdx = sameGroup ? 0 : 1; + + const module = device.createShaderModule({ + code: ` + @group(0) @binding(0) var t: ${textureType}; + @group(${groupNdx}) @binding(1) var s: sampler; + + fn test() { + _ = textureGather(${component} t, s, ${coord}${layer}); + } + + @compute @workgroup_size(1) fn cs() { + test(); + } + + @vertex fn vs() -> @builtin(position) vec4f { + return vec4f(0); + } + + @fragment fn fs() -> @location(0) vec4f { + test(); + return vec4f(0); + } + `, + }); + + const bindGroup0LayoutEntries: GPUBindGroupLayoutEntry[] = [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE | GPUShaderStage.FRAGMENT, + texture: { + sampleType, + viewDimension, + multisampled: false, + }, + }, + ]; + + const samplerBGLEntry: GPUBindGroupLayoutEntry = { + binding: 1, + visibility: GPUShaderStage.COMPUTE | GPUShaderStage.FRAGMENT, + sampler: { + type: 'filtering', + }, + }; + + if (sameGroup) { + bindGroup0LayoutEntries.push(samplerBGLEntry); + } + + const bindGroupLayout0 = device.createBindGroupLayout({ + entries: bindGroup0LayoutEntries, + }); + + const pipelineLayoutDesc = { + bindGroupLayouts: [bindGroupLayout0], + }; + + if (!sameGroup) { + const bindGroupLayout1 = device.createBindGroupLayout({ + entries: [samplerBGLEntry], + }); + pipelineLayoutDesc.bindGroupLayouts.push(bindGroupLayout1); + } + + const layout = device.createPipelineLayout(pipelineLayoutDesc); + + const success = sampleType === 'float'; + + if (pipeline === 'compute') { + t.doCreateComputePipelineTest(async, success, { + layout, + compute: { module }, + }); + } else { + t.doCreateRenderPipelineTest(async, success, { + layout, + vertex: { module }, + fragment: { module, targets: [{ format: 'rgba8unorm' }] }, + }); + } + }); diff --git a/src/webgpu/api/validation/queue/copyToTexture/CopyExternalImageToTexture.spec.ts b/src/webgpu/api/validation/queue/copyToTexture/CopyExternalImageToTexture.spec.ts index 5677a81cc03c..622133721bed 100644 --- a/src/webgpu/api/validation/queue/copyToTexture/CopyExternalImageToTexture.spec.ts +++ b/src/webgpu/api/validation/queue/copyToTexture/CopyExternalImageToTexture.spec.ts @@ -178,8 +178,8 @@ class CopyExternalImageToTextureTest extends ValidationTest { } runTest( - imageBitmapCopyView: GPUImageCopyExternalImage, - textureCopyView: GPUImageCopyTextureTagged, + imageBitmapCopyView: GPUCopyExternalImageSourceInfo, + textureCopyView: GPUCopyExternalImageDestInfo, copySize: GPUExtent3D, validationScopeSuccess: boolean, exceptionName?: string diff --git a/src/webgpu/api/validation/render_pipeline/depth_stencil_state.spec.ts b/src/webgpu/api/validation/render_pipeline/depth_stencil_state.spec.ts index 165e6f9c405b..73467657eedc 100644 --- a/src/webgpu/api/validation/render_pipeline/depth_stencil_state.spec.ts +++ b/src/webgpu/api/validation/render_pipeline/depth_stencil_state.spec.ts @@ -232,6 +232,10 @@ g.test('depth_bias') .fn(t => { const { isAsync, topology, depthBias, depthBiasSlopeScale, depthBiasClamp } = t.params; + if (t.isCompatibility && !!depthBiasClamp) { + t.skip('depthBiasClamp must be 0 on compatibility mode'); + } + const isTriangleTopology = topology === 'triangle-list' || topology === 'triangle-strip'; const hasDepthBias = !!depthBias || !!depthBiasSlopeScale || !!depthBiasClamp; const shouldSucceed = !hasDepthBias || isTriangleTopology; diff --git a/src/webgpu/api/validation/render_pipeline/resource_compatibility.spec.ts b/src/webgpu/api/validation/render_pipeline/resource_compatibility.spec.ts index 1d2c582dc036..8c516ebb3b85 100644 --- a/src/webgpu/api/validation/render_pipeline/resource_compatibility.spec.ts +++ b/src/webgpu/api/validation/render_pipeline/resource_compatibility.spec.ts @@ -55,6 +55,7 @@ g.test('resource_compatibility') wgslResource.storageTexture.access !== 'read-only')), 'Storage buffers and textures cannot be used in vertex shaders' ); + t.skipIfTextureViewDimensionNotSupported(wgslResource.texture?.viewDimension); const emptyVS = ` @vertex fn main() -> @builtin(position) vec4f { diff --git a/src/webgpu/capability_info.ts b/src/webgpu/capability_info.ts index 5af80c45fa4c..7fcab69d1b97 100644 --- a/src/webgpu/capability_info.ts +++ b/src/webgpu/capability_info.ts @@ -248,7 +248,7 @@ export type VertexFormatInfo = { /** Number of components. */ readonly componentCount: 1 | 2 | 3 | 4; /** Size in bytes. */ - readonly byteSize: 2 | 4 | 8 | 12 | 16; + readonly byteSize: 1 | 2 | 4 | 8 | 12 | 16; /** The completely matching WGSL type for vertex format */ readonly wgslType: | 'f32' @@ -273,23 +273,32 @@ export const kVertexFormatInfo: { ['bytesPerComponent', 'type', 'componentCount', 'byteSize', 'wgslType'] as const, [ , , , , ] as const, { // 8 bit components + 'uint8': [ 1, 'uint', 1, 1, 'u32'], 'uint8x2': [ 1, 'uint', 2, 2, 'vec2'], 'uint8x4': [ 1, 'uint', 4, 4, 'vec4'], + 'sint8': [ 1, 'sint', 1, 1, 'i32'], 'sint8x2': [ 1, 'sint', 2, 2, 'vec2'], 'sint8x4': [ 1, 'sint', 4, 4, 'vec4'], + 'unorm8': [ 1, 'unorm', 1, 1, 'f32'], 'unorm8x2': [ 1, 'unorm', 2, 2, 'vec2'], 'unorm8x4': [ 1, 'unorm', 4, 4, 'vec4'], + 'snorm8': [ 1, 'snorm', 1, 1, 'f32'], 'snorm8x2': [ 1, 'snorm', 2, 2, 'vec2'], 'snorm8x4': [ 1, 'snorm', 4, 4, 'vec4'], // 16 bit components + 'uint16': [ 2, 'uint', 1, 2, 'u32'], 'uint16x2': [ 2, 'uint', 2, 4, 'vec2'], 'uint16x4': [ 2, 'uint', 4, 8, 'vec4'], + 'sint16': [ 2, 'sint', 1, 2, 'i32'], 'sint16x2': [ 2, 'sint', 2, 4, 'vec2'], 'sint16x4': [ 2, 'sint', 4, 8, 'vec4'], + 'unorm16': [ 2, 'unorm', 1, 2, 'f32'], 'unorm16x2': [ 2, 'unorm', 2, 4, 'vec2'], 'unorm16x4': [ 2, 'unorm', 4, 8, 'vec4'], + 'snorm16': [ 2, 'snorm', 1, 2, 'f32'], 'snorm16x2': [ 2, 'snorm', 2, 4, 'vec2'], 'snorm16x4': [ 2, 'snorm', 4, 8, 'vec4'], + 'float16': [ 2, 'float', 1, 2, 'f32'], 'float16x2': [ 2, 'float', 2, 4, 'vec2'], 'float16x4': [ 2, 'float', 4, 8, 'vec4'], // 32 bit components @@ -306,7 +315,8 @@ export const kVertexFormatInfo: { 'sint32x3': [ 4, 'sint', 3, 12, 'vec3'], 'sint32x4': [ 4, 'sint', 4, 16, 'vec4'], // 32 bit packed - 'unorm10-10-10-2': [ 'packed', 'unorm', 4, 4, 'vec4'] + 'unorm10-10-10-2': [ 'packed', 'unorm', 4, 4, 'vec4'], + 'unorm8x4-bgra': [ 'packed', 'unorm', 4, 4, 'vec4'], } as const); /** List of all GPUVertexFormat values. */ export const kVertexFormats = keysOf(kVertexFormatInfo); diff --git a/src/webgpu/compat/api/validation/pipeline_creation.spec.ts b/src/webgpu/compat/api/validation/pipeline_creation.spec.ts new file mode 100644 index 000000000000..7e70a036f536 --- /dev/null +++ b/src/webgpu/compat/api/validation/pipeline_creation.spec.ts @@ -0,0 +1,146 @@ +export const description = ` +Tests that createComputePipeline(async), and createRenderPipeline(async) +reject pipelines that are invalid in compat mode + +- test that depth textures can not be used with non-comparison samplers + +TODO: +- test that a shader that has more than min(maxSamplersPerShaderStage, maxSampledTexturesPerShaderStage) + texture+sampler combinations generates a validation error. +`; + +import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { kShaderStages } from '../../../shader/validation/decl/util.js'; +import { CompatibilityTest } from '../../compatibility_test.js'; + +export const g = makeTestGroup(CompatibilityTest); + +g.test('depth_textures') + .desc('Tests that depth textures can not be used with non-comparison samplers in compat mode.') + .params(u => + u // + .combineWithParams([ + { + sampleWGSL: 'textureSample(t, s, vec2f(0))', // should pass + textureType: 'texture_2d', + }, + { + sampleWGSL: 'textureSample(t, s, vec2f(0))', + textureType: 'texture_depth_2d', + }, + { + sampleWGSL: 'textureSample(t, s, vec3f(0))', + textureType: 'texture_depth_cube', + }, + { + sampleWGSL: 'textureSample(t, s, vec2f(0), 0)', + textureType: 'texture_depth_2d_array', + }, + { + sampleWGSL: 'textureSample(t, s, vec2f(0), vec2i(0, 0))', + textureType: 'texture_depth_2d', + }, + { + sampleWGSL: 'textureSample(t, s, vec2f(0), 0, vec2i(0, 0))', + textureType: 'texture_depth_2d_array', + }, + { + sampleWGSL: 'textureSampleLevel(t, s, vec2f(0), 0)', + textureType: 'texture_depth_2d', + }, + { + sampleWGSL: 'textureSampleLevel(t, s, vec3f(0), 0)', + textureType: 'texture_depth_cube', + }, + { + sampleWGSL: 'textureSampleLevel(t, s, vec2f(0), 0, 0)', + textureType: 'texture_depth_2d_array', + }, + { + sampleWGSL: 'textureSampleLevel(t, s, vec2f(0), 0, vec2i(0, 0))', + textureType: 'texture_depth_2d', + }, + { + sampleWGSL: 'textureSampleLevel(t, s, vec2f(0), 0, 0, vec2i(0, 0))', + textureType: 'texture_depth_2d_array', + }, + { + sampleWGSL: 'textureGather(t, s, vec2f(0))', + textureType: 'texture_depth_2d', + }, + { + sampleWGSL: 'textureGather(t, s, vec3f(0))', + textureType: 'texture_depth_cube', + }, + { + sampleWGSL: 'textureGather(t, s, vec2f(0), 0)', + textureType: 'texture_depth_2d_array', + }, + { + sampleWGSL: 'textureGather(t, s, vec2f(0), vec2i(0, 0))', + textureType: 'texture_depth_2d', + }, + { + sampleWGSL: 'textureGather(t, s, vec2f(0), 0, vec2i(0, 0))', + textureType: 'texture_depth_2d_array', + }, + ]) + .combine('stage', kShaderStages) + .filter(t => t.sampleWGSL.startsWith('textureGather') || t.stage === 'fragment') + .combine('async', [false, true] as const) + ) + .fn(t => { + const { sampleWGSL, textureType, stage, async } = t.params; + + const usageWGSL = `_ = ${sampleWGSL};`; + const module = t.device.createShaderModule({ + code: ` + @group(0) @binding(0) var t: ${textureType}; + @group(1) @binding(0) var s: sampler; + + // make sure it's fine such a combination exists but it's not used. + fn unused() { + ${usageWGSL}; + } + + @vertex fn vs() -> @builtin(position) vec4f { + ${stage === 'vertex' ? usageWGSL : ''} + return vec4f(0); + } + + @fragment fn fs() -> @location(0) vec4f { + ${stage === 'fragment' ? usageWGSL : ''} + return vec4f(0); + } + + @compute @workgroup_size(1) fn cs() { + ${stage === 'compute' ? usageWGSL : ''}; + } + `, + }); + + const success = !t.isCompatibility || textureType === 'texture_2d'; + switch (stage) { + case 'compute': + t.doCreateComputePipelineTest(async, success, { + layout: 'auto', + compute: { + module, + }, + }); + break; + case 'fragment': + case 'vertex': + t.doCreateRenderPipelineTest(async, success, { + layout: 'auto', + vertex: { + module, + }, + fragment: { + module, + targets: [{ format: 'rgba8unorm' }], + }, + }); + break; + } + }); diff --git a/src/webgpu/format_info.ts b/src/webgpu/format_info.ts index cbebf8873428..e65838fd2a10 100644 --- a/src/webgpu/format_info.ts +++ b/src/webgpu/format_info.ts @@ -1793,6 +1793,23 @@ export function canUseAsRenderTarget(format: GPUTextureFormat) { return kTextureFormatInfo[format].colorRender || isDepthOrStencilTextureFormat(format); } +export function is32Float(format: GPUTextureFormat) { + return format === 'r32float' || format === 'rg32float' || format === 'rgba32float'; +} + +/** + * Returns true if texture is filterable as `texture_xxx` + * + * examples: + * * 'rgba8unorm' -> true + * * 'depth16unorm' -> false + * * 'rgba32float' -> true (you need to enable feature 'float32-filterable') + */ +export function isFilterableAsTextureF32(format: GPUTextureFormat) { + const info = kTextureFormatInfo[format]; + return info.color?.type === 'float' || is32Float(format); +} + export const kCompatModeUnsupportedStorageTextureFormats: readonly GPUTextureFormat[] = [ 'rg32float', 'rg32sint', @@ -1817,12 +1834,21 @@ export function isRegularTextureFormat(format: GPUTextureFormat) { } /** - * Returns true of format is both compressed and a float format, for example 'bc6h-rgb-ufloat'. + * Returns true if format is both compressed and a float format, for example 'bc6h-rgb-ufloat'. */ export function isCompressedFloatTextureFormat(format: GPUTextureFormat) { return isCompressedTextureFormat(format) && format.includes('float'); } +/** + * Returns true if format is sint or uint + */ +export function isSintOrUintFormat(format: GPUTextureFormat) { + const info = kTextureFormatInfo[format]; + const type = info.color?.type ?? info.depth?.type ?? info.stencil?.type; + return type === 'sint' || type === 'uint'; +} + /** * Returns true of format can be multisampled. */ diff --git a/src/webgpu/gpu_test.ts b/src/webgpu/gpu_test.ts index ef210712feed..ed8c170d93d2 100644 --- a/src/webgpu/gpu_test.ts +++ b/src/webgpu/gpu_test.ts @@ -300,6 +300,14 @@ export class GPUTestSubcaseBatchState extends SubcaseBatchState { } } + /** Skips this test case if a depth texture can not be used with a non-comparison sampler. */ + skipIfDepthTextureCanNotBeUsedWithNonComparisonSampler() { + this.skipIf( + this.isCompatibility, + 'depth textures are not usable with non-comparison samplers in compatibility mode' + ); + } + /** Skips this test case if the `langFeature` is *not* supported. */ skipIfLanguageFeatureNotSupported(langFeature: WGSLLanguageFeature) { if (!this.hasLanguageFeature(langFeature)) { @@ -1296,7 +1304,7 @@ export interface TextureTestMixinType { * to the expected TexelView passes without error. */ expectTexelViewComparisonIsOkInTexture( - src: GPUImageCopyTexture, + src: GPUTexelCopyTextureInfo, exp: TexelView, size: GPUExtent3D, comparisonOptions?: TexelCompareOptions @@ -1307,7 +1315,7 @@ export interface TextureTestMixinType { * their expected colors without error. */ expectSinglePixelComparisonsAreOkInTexture( - src: GPUImageCopyTexture, + src: GPUTexelCopyTextureInfo, exp: PerPixelComparison[], comparisonOptions?: TexelCompareOptions ): void; @@ -1383,7 +1391,7 @@ export interface TextureTestMixinType { * Gets a byte offset to a texel */ getTexelOffsetInBytes( - textureDataLayout: Required, + textureDataLayout: Required, format: ColorTextureFormat, texel: Required, origin?: Required @@ -1481,7 +1489,7 @@ function getPipelineToRenderTextureToRGB8UnormTexture( } type LinearCopyParameters = { - dataLayout: Required; + dataLayout: Required; origin: Required; data: Uint8Array; }; @@ -1511,7 +1519,7 @@ export function TextureTestMixin>( } expectTexelViewComparisonIsOkInTexture( - src: GPUImageCopyTexture, + src: GPUTexelCopyTextureInfo, exp: TexelView, size: GPUExtent3D, comparisonOptions = { @@ -1526,7 +1534,7 @@ export function TextureTestMixin>( } expectSinglePixelComparisonsAreOkInTexture( - src: GPUImageCopyTexture, + src: GPUTexelCopyTextureInfo, exp: PerPixelComparison[], comparisonOptions = { maxIntDiff: 0, @@ -1836,7 +1844,7 @@ export function TextureTestMixin>( /** Offset for a particular texel in the linear texture data */ getTexelOffsetInBytes( - textureDataLayout: Required, + textureDataLayout: Required, format: ColorTextureFormat, texel: Required, origin: Required = { x: 0, y: 0, z: 0 } diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index c455224da996..a45fa21e1ef8 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -274,6 +274,7 @@ "webgpu:api,validation,buffer,mapping:unmap,state,mappingPending:*": { "subcaseMS": 22.951 }, "webgpu:api,validation,buffer,mapping:unmap,state,unmapped:*": { "subcaseMS": 74.200 }, "webgpu:api,validation,capability_checks,features,clip_distances:createRenderPipeline,at_over:*": { "subcaseMS": 13.700 }, + "webgpu:api,validation,capability_checks,features,clip_distances:createRenderPipeline,max_vertex_output_location:*": { "subcaseMS": 267.295 }, "webgpu:api,validation,capability_checks,features,query_types:createQuerySet:*": { "subcaseMS": 10.451 }, "webgpu:api,validation,capability_checks,features,query_types:timestamp:*": { "subcaseMS": 1.200 }, "webgpu:api,validation,capability_checks,features,texture_formats:canvas_configuration:*": { "subcaseMS": 4.339 }, @@ -434,6 +435,7 @@ "webgpu:api,validation,createView:format:*": { "subcaseMS": 0.742 }, "webgpu:api,validation,createView:mip_levels:*": { "subcaseMS": 0.436 }, "webgpu:api,validation,createView:texture_state:*": { "subcaseMS": 0.400 }, + "webgpu:api,validation,createView:texture_view_usage:*": { "subcaseMS": 3106.634 }, "webgpu:api,validation,debugMarker:push_pop_call_count_unbalance,command_encoder:*": { "subcaseMS": 1.522 }, "webgpu:api,validation,debugMarker:push_pop_call_count_unbalance,render_compute_pass:*": { "subcaseMS": 0.601 }, "webgpu:api,validation,encoding,beginComputePass:timestampWrites,invalid_query_set:*": { "subcaseMS": 0.201 }, @@ -652,6 +654,9 @@ "webgpu:api,validation,queue,destroyed,texture:setBindGroup:*": { "subcaseMS": 5.783 }, "webgpu:api,validation,queue,destroyed,texture:writeTexture:*": { "subcaseMS": 16.601 }, "webgpu:api,validation,queue,submit:command_buffer,device_mismatch:*": { "subcaseMS": 0.467 }, + "webgpu:api,validation,queue,submit:command_buffer,duplicate_buffers:*": { "subcaseMS": 0.981 }, + "webgpu:api,validation,queue,submit:command_buffer,invalid_submit_invalidates:*": { "subcaseMS": 0.820 }, + "webgpu:api,validation,queue,submit:command_buffer,submit_invalidates:*": { "subcaseMS": 1.120 }, "webgpu:api,validation,queue,writeBuffer:buffer,device_mismatch:*": { "subcaseMS": 16.000 }, "webgpu:api,validation,queue,writeBuffer:buffer_state:*": { "subcaseMS": 6.201 }, "webgpu:api,validation,queue,writeBuffer:ranges:*": { "subcaseMS": 17.600 }, @@ -706,13 +711,17 @@ "webgpu:api,validation,render_pass,resolve:resolve_attachment:*": { "subcaseMS": 6.205 }, "webgpu:api,validation,render_pipeline,depth_stencil_state:depthCompare_optional:*": { "subcaseMS": 21.401 }, "webgpu:api,validation,render_pipeline,depth_stencil_state:depthWriteEnabled_optional:*": { "subcaseMS": 16.950 }, + "webgpu:api,validation,render_pipeline,depth_stencil_state:depth_bias:*": { "subcaseMS": 45.563 }, "webgpu:api,validation,render_pipeline,depth_stencil_state:depth_test:*": { "subcaseMS": 3.407 }, "webgpu:api,validation,render_pipeline,depth_stencil_state:depth_write,frag_depth:*": { "subcaseMS": 6.465 }, "webgpu:api,validation,render_pipeline,depth_stencil_state:depth_write:*": { "subcaseMS": 4.113 }, "webgpu:api,validation,render_pipeline,depth_stencil_state:format:*": { "subcaseMS": 3.521 }, "webgpu:api,validation,render_pipeline,depth_stencil_state:stencil_test:*": { "subcaseMS": 3.124 }, "webgpu:api,validation,render_pipeline,depth_stencil_state:stencil_write:*": { "subcaseMS": 3.183 }, + "webgpu:api,validation,render_pipeline,float32_blendable:create_render_pipeline:*": { "subcaseMS": 78.452 }, "webgpu:api,validation,render_pipeline,fragment_state:color_target_exists:*": { "subcaseMS": 29.150 }, + "webgpu:api,validation,render_pipeline,fragment_state:dual_source_blending,color_target_count:*": { "subcaseMS": 38.712 }, + "webgpu:api,validation,render_pipeline,fragment_state:dual_source_blending,use_blend_src:*": { "subcaseMS": 316.418 }, "webgpu:api,validation,render_pipeline,fragment_state:limits,maxColorAttachmentBytesPerSample,aligned:*": { "subcaseMS": 0.991 }, "webgpu:api,validation,render_pipeline,fragment_state:limits,maxColorAttachmentBytesPerSample,unaligned:*": { "subcaseMS": 14.750 }, "webgpu:api,validation,render_pipeline,fragment_state:limits,maxColorAttachments:*": { "subcaseMS": 9.351 }, @@ -734,6 +743,7 @@ "webgpu:api,validation,render_pipeline,inter_stage:type:*": { "subcaseMS": 6.170 }, "webgpu:api,validation,render_pipeline,misc:basic:*": { "subcaseMS": 0.901 }, "webgpu:api,validation,render_pipeline,misc:external_texture:*": { "subcaseMS": 35.189 }, + "webgpu:api,validation,render_pipeline,misc:no_attachment:*": { "subcaseMS": 2.264 }, "webgpu:api,validation,render_pipeline,misc:pipeline_layout,device_mismatch:*": { "subcaseMS": 8.700 }, "webgpu:api,validation,render_pipeline,misc:vertex_state_only:*": { "subcaseMS": 1.125 }, "webgpu:api,validation,render_pipeline,multisample_state:alpha_to_coverage,count:*": { "subcaseMS": 3.200 }, @@ -796,6 +806,7 @@ "webgpu:api,validation,resource_usages,texture,in_render_misc:subresources,set_bind_group_on_same_index_depth_stencil_texture:*": { "subcaseMS": 0.925 }, "webgpu:api,validation,resource_usages,texture,in_render_misc:subresources,set_unused_bind_group:*": { "subcaseMS": 6.200 }, "webgpu:api,validation,resource_usages,texture,in_render_misc:subresources,texture_usages_in_copy_and_render_pass:*": { "subcaseMS": 4.763 }, + "webgpu:api,validation,resource_usages,texture,in_render_misc:subresources,texture_view_usages:*": { "subcaseMS": 24.999 }, "webgpu:api,validation,shader_module,entry_point:compute:*": { "subcaseMS": 4.439 }, "webgpu:api,validation,shader_module,entry_point:compute_undefined_entry_point_and_extra_stage:*": { "subcaseMS": 17.075 }, "webgpu:api,validation,shader_module,entry_point:fragment:*": { "subcaseMS": 5.865 }, @@ -1556,6 +1567,7 @@ "webgpu:shader,execution,expression,call,builtin,subgroupBitwise:compute,split:*": { "subcaseMS": 1743.045 }, "webgpu:shader,execution,expression,call,builtin,subgroupBitwise:data_types:*": { "subcaseMS": 5081.792 }, "webgpu:shader,execution,expression,call,builtin,subgroupBitwise:fragment,all_active:*": { "subcaseMS": 9079.446 }, + "webgpu:shader,execution,expression,call,builtin,subgroupBitwise:fragment,split:*": { "subcaseMS": 0.347 }, "webgpu:shader,execution,expression,call,builtin,subgroupBroadcast:data_types:*": { "subcaseMS": 252.374 }, "webgpu:shader,execution,expression,call,builtin,subgroupBroadcast:fragment:*": { "subcaseMS": 0.108 }, "webgpu:shader,execution,expression,call,builtin,subgroupBroadcast:workgroup_uniform_load:*": { "subcaseMS": 109.832 }, @@ -1639,10 +1651,6 @@ "webgpu:shader,execution,expression,call,builtin,textureStore:bgra8unorm_swizzle:*": { "subcaseMS": 30.325 }, "webgpu:shader,execution,expression,call,builtin,textureStore:out_of_bounds:*": { "subcaseMS": 942.418 }, "webgpu:shader,execution,expression,call,builtin,textureStore:out_of_bounds_array:*": { "subcaseMS": 609.565 }, - "webgpu:shader,execution,expression,call,builtin,textureStore:store_1d_coords:*": { "subcaseMS": 19.907 }, - "webgpu:shader,execution,expression,call,builtin,textureStore:store_2d_coords:*": { "subcaseMS": 28.809 }, - "webgpu:shader,execution,expression,call,builtin,textureStore:store_3d_coords:*": { "subcaseMS": 37.206 }, - "webgpu:shader,execution,expression,call,builtin,textureStore:store_array_2d_coords:*": { "subcaseMS": 98.804 }, "webgpu:shader,execution,expression,call,builtin,textureStore:texel_formats:*": { "subcaseMS": 86.179 }, "webgpu:shader,execution,expression,call,builtin,transpose:abstract_float:*": { "subcaseMS": 64537.678 }, "webgpu:shader,execution,expression,call,builtin,transpose:f16:*": { "subcaseMS": 33.311 }, diff --git a/src/webgpu/shader/execution/expression/call/builtin/quadBroadcast.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/quadBroadcast.spec.ts index 29a3ec47c860..691bc5e8a30e 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/quadBroadcast.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/quadBroadcast.spec.ts @@ -447,6 +447,9 @@ predication filters are skipped. const wgsl = ` enable subgroups; +diagnostic(off, subgroup_uniformity); +diagnostic(off, subgroup_branching); + @group(0) @binding(0) var inputs : u32; // unused diff --git a/src/webgpu/shader/execution/expression/call/builtin/quadSwap.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/quadSwap.spec.ts index e6b6863a8e94..49489df0085c 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/quadSwap.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/quadSwap.spec.ts @@ -466,6 +466,9 @@ predication filters are skipped. const wgsl = ` enable subgroups; +diagnostic(off, subgroup_uniformity); +diagnostic(off, subgroup_branching); + @group(0) @binding(0) var inputs : u32; // unused diff --git a/src/webgpu/shader/execution/expression/call/builtin/subgroupAdd.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/subgroupAdd.spec.ts index a35981a31d5a..86510fe87738 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/subgroupAdd.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/subgroupAdd.spec.ts @@ -311,6 +311,9 @@ g.test('compute,split') const wgsl = ` enable subgroups; +diagnostic(off, subgroup_uniformity); +diagnostic(off, subgroup_branching); + @group(0) @binding(0) var input : array; diff --git a/src/webgpu/shader/execution/expression/call/builtin/subgroupAll.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/subgroupAll.spec.ts index 0aa461c4a578..2b762053028e 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/subgroupAll.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/subgroupAll.spec.ts @@ -208,6 +208,9 @@ g.test('compute,split') const wgsl = ` enable subgroups; +diagnostic(off, subgroup_uniformity); +diagnostic(off, subgroup_branching); + @group(0) @binding(0) var inputs : array; diff --git a/src/webgpu/shader/execution/expression/call/builtin/subgroupAny.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/subgroupAny.spec.ts index 5d5b9de11420..5254ade1739a 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/subgroupAny.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/subgroupAny.spec.ts @@ -208,6 +208,9 @@ g.test('compute,split') const wgsl = ` enable subgroups; +diagnostic(off, subgroup_uniformity); +diagnostic(off, subgroup_branching); + @group(0) @binding(0) var inputs : array; diff --git a/src/webgpu/shader/execution/expression/call/builtin/subgroupBallot.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/subgroupBallot.spec.ts index 9b55405f88ba..4bd79dbdc15a 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/subgroupBallot.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/subgroupBallot.spec.ts @@ -188,6 +188,9 @@ g.test('compute,split') const wgsl = ` enable subgroups; +diagnostic(off, subgroup_uniformity); +diagnostic(off, subgroup_branching); + @group(0) @binding(0) var size : u32; @@ -224,6 +227,8 @@ g.test('predicate') const wgsl = ` enable subgroups; +diagnostic(off, subgroup_branching); + @group(0) @binding(0) var size : u32; @@ -313,6 +318,9 @@ g.test('predicate_and_control_flow') const wgsl = ` enable subgroups; +diagnostic(off, subgroup_branching); +diagnostic(off, subgroup_uniformity); + @group(0) @binding(0) var size : u32; diff --git a/src/webgpu/shader/execution/expression/call/builtin/subgroupBitwise.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/subgroupBitwise.spec.ts index c50fd08a1570..e95c486c3e59 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/subgroupBitwise.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/subgroupBitwise.spec.ts @@ -378,6 +378,9 @@ g.test('compute,split') const wgsl = ` enable subgroups; +diagnostic(off, subgroup_uniformity); +diagnostic(off, subgroup_branching); + @group(0) @binding(0) var inputs : array; diff --git a/src/webgpu/shader/execution/expression/call/builtin/subgroupBroadcast.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/subgroupBroadcast.spec.ts index 75fe27e8cb5d..d62abb830328 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/subgroupBroadcast.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/subgroupBroadcast.spec.ts @@ -242,6 +242,8 @@ g.test('workgroup_uniform_load') const wgsl = ` enable subgroups; +diagnostic(off, subgroup_branching); + var wgmem : u32; @group(0) @binding(0) diff --git a/src/webgpu/shader/execution/expression/call/builtin/subgroupMul.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/subgroupMul.spec.ts index 28734afeb7d6..dfcaf481fcab 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/subgroupMul.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/subgroupMul.spec.ts @@ -334,6 +334,9 @@ g.test('compute,split') const wgsl = ` enable subgroups; +diagnostic(off, subgroup_uniformity); +diagnostic(off, subgroup_branching); + @group(0) @binding(0) var input : array; diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureDimensions.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureDimensions.spec.ts index b807ca97708b..b80c8334a55b 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureDimensions.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureDimensions.spec.ts @@ -15,10 +15,12 @@ import { sampleTypeForFormatAndAspect, textureDimensionAndFormatCompatible, } from '../../../../../format_info.js'; -import { GPUTest } from '../../../../../gpu_test.js'; import { align } from '../../../../../util/math.js'; +import { kShaderStages, ShaderStage } from '../../../../validation/decl/util.js'; -export const g = makeTestGroup(GPUTest); +import { WGSLTextureQueryTest } from './texture_utils.js'; + +export const g = makeTestGroup(WGSLTextureQueryTest); /// The maximum number of texture mipmap levels to test. /// Keep this small to reduce memory and test permutations. @@ -218,8 +220,10 @@ function testValues(params: { * `values.expected`. */ function run( - t: GPUTest, - view: GPUTextureView, + t: WGSLTextureQueryTest, + stage: ShaderStage, + texture: GPUTexture | GPUExternalTexture, + viewDescriptor: GPUTextureViewDescriptor | undefined, textureType: string, levelArg: number | undefined, values: TestValues @@ -227,44 +231,16 @@ function run( const outputType = values.expected.length > 1 ? `vec${values.expected.length}u` : 'u32'; const wgsl = ` @group(0) @binding(0) var texture : ${textureType}; -@group(0) @binding(1) var output : ${outputType}; -@compute @workgroup_size(1) -fn main() { -output = ${ +fn getValue() -> ${outputType} { + return ${ levelArg !== undefined ? `textureDimensions(texture, ${levelArg})` : 'textureDimensions(texture)' }; } `; - const module = t.device.createShaderModule({ - code: wgsl, - }); - const pipeline = t.device.createComputePipeline({ - compute: { module }, - layout: 'auto', - }); - const outputBuffer = t.createBufferTracked({ - size: 32, - usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.STORAGE, - }); - const bindgroup = t.device.createBindGroup({ - layout: pipeline.getBindGroupLayout(0), - entries: [ - { binding: 0, resource: view }, - { binding: 1, resource: { buffer: outputBuffer } }, - ], - }); - const encoder = t.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(pipeline); - pass.setBindGroup(0, bindgroup); - pass.dispatchWorkgroups(1); - pass.end(); - t.device.queue.submit([encoder.finish()]); - - t.expectGPUBufferValuesEqual(outputBuffer, new Uint32Array(values.expected)); + t.executeAndExpectResult(stage, wgsl, texture, viewDescriptor, values.expected); } /** @returns true if the GPUTextureViewDimension is valid for a storage texture */ @@ -314,6 +290,7 @@ Parameters: .expand('aspect', u => aspectsForFormat(u.format)) .expand('samples', u => samplesForFormat(u.format)) .beginSubcases() + .combine('stage', kShaderStages) .expand('dimensions', viewDimensions) .expand('textureMipCount', textureMipCount) .expand('baseMipLevel', baseMipLevel) @@ -339,11 +316,11 @@ Parameters: sampleCount: t.params.samples, mipLevelCount: t.params.textureMipCount, }); - const textureView = texture.createView({ + const viewDescriptor: GPUTextureViewDescriptor = { dimension: t.params.dimensions, aspect: t.params.aspect, baseMipLevel: t.params.baseMipLevel, - }); + }; function wgslSampledTextureType(): string { const base = t.params.samples !== 1 ? 'texture_multisampled' : 'texture'; @@ -362,7 +339,15 @@ Parameters: } } - run(t, textureView, wgslSampledTextureType(), t.params.textureDimensionsLevel, values); + run( + t, + t.params.stage, + texture, + viewDescriptor, + wgslSampledTextureType(), + t.params.textureDimensionsLevel, + values + ); }); g.test('depth') @@ -394,6 +379,7 @@ Parameters: .unless(u => u.aspect === 'stencil-only') .expand('samples', u => samplesForFormat(u.format)) .beginSubcases() + .combine('stage', kShaderStages) .expand('dimensions', viewDimensions) .expand('textureMipCount', textureMipCount) .expand('baseMipLevel', baseMipLevel) @@ -419,11 +405,11 @@ Parameters: sampleCount: t.params.samples, mipLevelCount: t.params.textureMipCount, }); - const textureView = texture.createView({ + const viewDescriptor: GPUTextureViewDescriptor = { dimension: t.params.dimensions, aspect: t.params.aspect, baseMipLevel: t.params.baseMipLevel, - }); + }; function wgslDepthTextureType(): string { const base = t.params.samples !== 1 ? 'texture_depth_multisampled' : 'texture_depth'; @@ -431,7 +417,15 @@ Parameters: return `${base}_${dimensions}`; } - run(t, textureView, wgslDepthTextureType(), t.params.textureDimensionsLevel, values); + run( + t, + t.params.stage, + texture, + viewDescriptor, + wgslDepthTextureType(), + t.params.textureDimensionsLevel, + values + ); }); g.test('storage') @@ -471,6 +465,15 @@ Parameters: .filter(p => kTextureFormatInfo[p.format].color?.storage === true) .expand('aspect', u => aspectsForFormat(u.format)) .beginSubcases() + .combine('stage', kShaderStages) + .combine('access', ['read', 'write', 'read_write'] as const) + // vertex stage can not use writable storage. + .unless(t => t.stage === 'vertex' && t.access !== 'read') + // Only some formats support write + .unless( + t => + kTextureFormatInfo[t.format].color.readWriteStorage === false && t.access === 'read_write' + ) .expand('dimensions', u => viewDimensions(u).filter(dimensionsValidForStorage)) .expand('textureMipCount', textureMipCount) .expand('baseMipLevel', baseMipLevel) @@ -490,19 +493,19 @@ Parameters: format: t.params.format, mipLevelCount: t.params.textureMipCount, }); - const textureView = texture.createView({ + const viewDescriptor: GPUTextureViewDescriptor = { dimension: t.params.dimensions, aspect: t.params.aspect, mipLevelCount: 1, baseMipLevel: t.params.baseMipLevel, - }); + }; function wgslStorageTextureType(): string { const dimensions = t.params.dimensions.replace('-', '_'); - return `texture_storage_${dimensions}<${t.params.format}, write>`; + return `texture_storage_${dimensions}<${t.params.format}, ${t.params.access}>`; } - run(t, textureView, wgslStorageTextureType(), undefined, values); + run(t, t.params.stage, texture, viewDescriptor, wgslStorageTextureType(), undefined, values); }); g.test('external') @@ -515,4 +518,23 @@ Parameters: * t: the external texture ` ) - .unimplemented(); + .params(u => + u + .beginSubcases() + .combine('stage', kShaderStages) + .combine('width', [8, 16, 24] as const) + .combine('height', [8, 16, 24] as const) + ) + .fn(t => { + const { stage, width, height } = t.params; + const canvas = new OffscreenCanvas(width, height); + // We have to make a context for VideoFrame to accept the canvas. + canvas.getContext('2d'); + const videoFrame = new VideoFrame(canvas, { timestamp: 0 }); + const texture = t.device.importExternalTexture({ source: videoFrame }); + + run(t, stage, texture, undefined, 'texture_external', undefined, { + size: [width, height], + expected: [width, height], + }); + }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts index d2ba15adb969..5aae5bafbe57 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts @@ -27,10 +27,9 @@ A texture gather operation reads from a 2D, 2D array, cube, or cube array textur import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; import { isDepthTextureFormat, - isEncodableTextureFormat, - kCompressedTextureFormats, + isFilterableAsTextureF32, kDepthStencilFormats, - kEncodableTextureFormats, + kAllTextureFormats, } from '../../../../../format_info.js'; import { @@ -54,8 +53,6 @@ import { WGSLTextureSampleTest, } from './texture_utils.js'; -const kTestableColorFormats = [...kEncodableTextureFormats, ...kCompressedTextureFormats] as const; - export const g = makeTestGroup(WGSLTextureSampleTest); g.test('sampled_2d_coords') @@ -87,9 +84,10 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('offset', [false, true] as const) @@ -184,9 +182,10 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('mode', kShortAddressModes) .beginSubcases() .combine('C', ['i32', 'u32'] as const) @@ -292,9 +291,10 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('offset', [false, true] as const) @@ -398,9 +398,10 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('mode', kShortAddressModes) .beginSubcases() .combine('samplePoints', kCubeSamplePointMethods) @@ -505,17 +506,18 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) - .combine('filt', ['nearest', 'linear'] as const) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('offset', [false, true] as const) .beginSubcases() .combine('samplePoints', kSamplePointMethods) ) + .beforeAllSubcases(t => { + t.skipIfDepthTextureCanNotBeUsedWithNonComparisonSampler(); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); + }) .fn(async t => { - const { format, stage, samplePoints, modeU, modeV, filt: minFilter, offset } = t.params; + const { format, stage, samplePoints, modeU, modeV, offset } = t.params; // We want at least 4 blocks or something wide enough for 3 mip levels. const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); @@ -529,9 +531,6 @@ Parameters: const sampler: GPUSamplerDescriptor = { addressModeU: kShortAddressModeToAddressMode[modeU], addressModeV: kShortAddressModeToAddressMode[modeV], - minFilter, - magFilter: minFilter, - mipmapFilter: minFilter, }; const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { @@ -540,7 +539,7 @@ Parameters: sampler, descriptor, offset, - hashInputs: [stage, format, samplePoints, modeU, modeV, minFilter, offset], + hashInputs: [stage, format, samplePoints, modeU, modeV, offset], }).map(({ coords, offset }) => { return { builtin: 'textureGather', @@ -591,15 +590,16 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) - .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .beginSubcases() .combine('samplePoints', kCubeSamplePointMethods) ) + .beforeAllSubcases(t => { + t.skipIfDepthTextureCanNotBeUsedWithNonComparisonSampler(); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); + }) .fn(async t => { - const { format, stage, samplePoints, mode, filt: minFilter } = t.params; + const { format, stage, samplePoints, mode } = t.params; const viewDimension: GPUTextureViewDimension = 'cube'; const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); @@ -617,9 +617,6 @@ Parameters: addressModeU: kShortAddressModeToAddressMode[mode], addressModeV: kShortAddressModeToAddressMode[mode], addressModeW: kShortAddressModeToAddressMode[mode], - minFilter, - magFilter: minFilter, - mipmapFilter: minFilter, }; const calls: TextureCall[] = generateSamplePointsCube(50, { @@ -627,7 +624,7 @@ Parameters: sampler, descriptor, textureBuiltin: 'textureGather', - hashInputs: [stage, format, samplePoints, mode, minFilter], + hashInputs: [stage, format, samplePoints, mode], }).map(({ coords, component }) => { return { builtin: 'textureGather', @@ -689,9 +686,6 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) - .combine('filt', ['nearest', 'linear'] as const) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('offset', [false, true] as const) @@ -700,11 +694,12 @@ Parameters: .combine('A', ['i32', 'u32'] as const) ) .beforeAllSubcases(t => { + t.skipIfDepthTextureCanNotBeUsedWithNonComparisonSampler(); t.skipIfTextureFormatNotSupported(t.params.format); - skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.filt, t.params.format); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { - const { format, stage, samplePoints, A, modeU, modeV, filt: minFilter, offset } = t.params; + const { format, stage, samplePoints, A, modeU, modeV, offset } = t.params; // We want at least 4 blocks or something wide enough for 3 mip levels. const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); @@ -720,9 +715,6 @@ Parameters: const sampler: GPUSamplerDescriptor = { addressModeU: kShortAddressModeToAddressMode[modeU], addressModeV: kShortAddressModeToAddressMode[modeV], - minFilter, - magFilter: minFilter, - mipmapFilter: minFilter, }; const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { @@ -732,7 +724,7 @@ Parameters: descriptor, arrayIndex: { num: texture.depthOrArrayLayers, type: A }, offset, - hashInputs: [stage, format, samplePoints, A, modeU, modeV, minFilter, offset], + hashInputs: [stage, format, samplePoints, A, modeU, modeV, offset], }).map(({ coords, arrayIndex, offset }) => { return { builtin: 'textureGather', @@ -788,19 +780,18 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) - .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .beginSubcases() .combine('samplePoints', kCubeSamplePointMethods) .combine('A', ['i32', 'u32'] as const) ) .beforeAllSubcases(t => { + t.skipIfDepthTextureCanNotBeUsedWithNonComparisonSampler(); t.skipIfTextureViewDimensionNotSupported('cube-array'); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { - const { format, A, stage, samplePoints, mode, filt: minFilter } = t.params; + const { format, A, stage, samplePoints, mode } = t.params; const viewDimension: GPUTextureViewDimension = 'cube-array'; const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); @@ -817,9 +808,6 @@ Parameters: addressModeU: kShortAddressModeToAddressMode[mode], addressModeV: kShortAddressModeToAddressMode[mode], addressModeW: kShortAddressModeToAddressMode[mode], - minFilter, - magFilter: minFilter, - mipmapFilter: minFilter, }; const calls: TextureCall[] = generateSamplePointsCube(50, { @@ -828,7 +816,7 @@ Parameters: descriptor, textureBuiltin: 'textureGather', arrayIndex: { num: texture.depthOrArrayLayers / 6, type: A }, - hashInputs: [stage, format, samplePoints, mode, minFilter], + hashInputs: [stage, format, samplePoints, mode], }).map(({ coords, arrayIndex }) => { return { builtin: 'textureGather', diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts index f86a152c19bc..e8a2be30bad4 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts @@ -18,11 +18,7 @@ A texture gather compare operation performs a depth comparison on four texels in import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; import { kCompareFunctions } from '../../../../../capability_info.js'; -import { - isDepthTextureFormat, - isEncodableTextureFormat, - kDepthStencilFormats, -} from '../../../../../format_info.js'; +import { isDepthTextureFormat, kDepthStencilFormats } from '../../../../../format_info.js'; import { checkCallResults, @@ -74,8 +70,6 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combine('filt', ['nearest', 'linear'] as const) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) @@ -87,6 +81,7 @@ Parameters: ) .beforeAllSubcases(t => { t.skipIfTextureFormatNotSupported(t.params.format); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { const { @@ -187,8 +182,6 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .beginSubcases() @@ -198,6 +191,7 @@ Parameters: ) .beforeAllSubcases(t => { t.skipIfTextureViewDimensionNotSupported('cube-array'); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { const { format, A, stage, samplePoints, mode, filt: minFilter, compare } = t.params; @@ -294,8 +288,6 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .combine('offset', [false, true] as const) @@ -304,6 +296,7 @@ Parameters: .combine('samplePoints', kSamplePointMethods) .combine('compare', kCompareFunctions) ) + .beforeAllSubcases(t => t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format)) .fn(async t => { const { format, C, stage, samplePoints, mode, compare, filt: minFilter, offset } = t.params; @@ -385,14 +378,13 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .beginSubcases() .combine('samplePoints', kCubeSamplePointMethods) .combine('compare', kCompareFunctions) ) + .beforeAllSubcases(t => t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format)) .fn(async t => { const { format, stage, samplePoints, mode, filt: minFilter, compare } = t.params; diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts index 689df4feb084..62378111c0a5 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts @@ -26,12 +26,10 @@ import { canUseAsRenderTarget, isCompressedFloatTextureFormat, isDepthTextureFormat, - isEncodableTextureFormat, isMultisampledTextureFormat, isStencilTextureFormat, - kCompressedTextureFormats, kDepthStencilFormats, - kEncodableTextureFormats, + kAllTextureFormats, kTextureFormatInfo, textureDimensionAndFormatCompatible, } from '../../../../../format_info.js'; @@ -58,8 +56,6 @@ import { createVideoFrameWithRandomDataAndGetTexels, } from './texture_utils.js'; -const kTestableColorFormats = [...kEncodableTextureFormats, ...kCompressedTextureFormats] as const; - export function normalizedCoordToTexelLoadTestCoord( descriptor: GPUTextureDescriptor, mipLevel: number, @@ -92,7 +88,7 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => textureDimensionAndFormatCompatible('1d', t.format)) // 1d textures can't have a height !== 1 .filter(t => kTextureFormatInfo[t.format].blockHeight === 1) @@ -179,7 +175,7 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => !isCompressedFloatTextureFormat(t.format)) .beginSubcases() .combine('samplePoints', kSamplePointMethods) @@ -261,7 +257,7 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => textureDimensionAndFormatCompatible('3d', t.format)) .beginSubcases() .combine('samplePoints', kSamplePointMethods) @@ -351,7 +347,7 @@ Parameters: 'texture_multisampled_2d', 'texture_depth_multisampled_2d', ] as const) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isMultisampledTextureFormat(t.format)) .filter(t => !isStencilTextureFormat(t.format)) // Filter out texture_depth_multisampled_2d with non-depth formats @@ -445,8 +441,6 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .beginSubcases() .combine('samplePoints', kSamplePointMethods) .combine('C', ['i32', 'u32'] as const) @@ -454,6 +448,7 @@ Parameters: ) .beforeAllSubcases(t => { t.skipIfTextureLoadNotSupportedForTextureType('texture_depth_2d'); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { const { format, stage, samplePoints, C, L } = t.params; @@ -604,7 +599,7 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) // MAINTENANCE_TODO: Update createTextureFromTexelViews to support stencil8 and remove this filter. .filter(t => t.format !== 'stencil8' && !isCompressedFloatTextureFormat(t.format)) .combine('texture_type', ['texture_2d_array', 'texture_depth_2d_array'] as const) diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts index 500376321444..0cb94b798a5d 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts @@ -7,6 +7,7 @@ Returns the number of layers (elements) of an array texture. import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; import { kTextureFormatInfo } from '../../../../../format_info.js'; import { TexelFormats } from '../../../../types.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; import { kSampleTypeInfo, WGSLTextureQueryTest } from './texture_utils.js'; @@ -54,6 +55,7 @@ Parameters .combine('view_type', ['full', 'partial'] as const) .beginSubcases() .combine('sampled_type', ['f32', 'i32', 'u32'] as const) + .combine('stage', kShaderStages) ) .beforeAllSubcases(t => { t.skipIf( @@ -66,7 +68,7 @@ Parameters ); }) .fn(t => { - const { texture_type, sampled_type, view_type } = t.params; + const { stage, texture_type, sampled_type, view_type } = t.params; const { format } = kSampleTypeInfo[sampled_type]; const texture = t.createTextureTracked({ @@ -77,9 +79,8 @@ Parameters const code = ` @group(0) @binding(0) var t: ${texture_type}<${sampled_type}>; -@group(0) @binding(1) var result: u32; -@compute @workgroup_size(1) fn cs() { - result = textureNumLayers(t); +fn getValue() -> u32 { + return textureNumLayers(t); } `; @@ -87,13 +88,13 @@ Parameters view_type, isCubeArray: texture_type === 'texture_cube_array', }); - const view = texture.createView({ + const viewDescription: GPUTextureViewDescriptor = { dimension: texture_type === 'texture_2d_array' ? '2d-array' : 'cube-array', baseArrayLayer, arrayLayerCount, - }); + }; - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, code, texture, viewDescription, expected); }); g.test('arrayed') @@ -111,6 +112,8 @@ Parameters u .combine('texture_type', ['texture_depth_2d_array', 'texture_depth_cube_array'] as const) .combine('view_type', ['full', 'partial'] as const) + .beginSubcases() + .combine('stage', kShaderStages) ) .beforeAllSubcases(t => { t.skipIf( @@ -123,7 +126,7 @@ Parameters ); }) .fn(t => { - const { texture_type, view_type } = t.params; + const { stage, texture_type, view_type } = t.params; const texture = t.createTextureTracked({ format: 'depth32float', @@ -134,8 +137,8 @@ Parameters const code = ` @group(0) @binding(0) var t: ${texture_type}; @group(0) @binding(1) var result: u32; -@compute @workgroup_size(1) fn cs() { - result = textureNumLayers(t); +fn getValue() -> u32 { + return textureNumLayers(t); } `; @@ -143,13 +146,13 @@ Parameters view_type, isCubeArray: texture_type === 'texture_depth_cube_array', }); - const view = texture.createView({ + const viewDescription: GPUTextureViewDescriptor = { dimension: texture_type === 'texture_depth_2d_array' ? '2d-array' : 'cube-array', baseArrayLayer, arrayLayerCount, - }); + }; - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, code, texture, viewDescription, expected); }); g.test('storage') @@ -185,10 +188,13 @@ Parameters .combineWithParams(TexelFormats) .combine('view_type', ['full', 'partial'] as const) .beginSubcases() + .combine('stage', kShaderStages) .combine('access_mode', ['read', 'write', 'read_write'] as const) .filter( t => t.access_mode !== 'read_write' || kTextureFormatInfo[t.format].color?.readWriteStorage ) + // Vertex stage can not use writable storage textures. + .unless(t => t.stage === 'vertex' && t.access_mode !== 'read') ) .beforeAllSubcases(t => { t.skipIf( @@ -198,7 +204,7 @@ Parameters t.skipIfTextureFormatNotUsableAsStorageTexture(t.params.format); }) .fn(t => { - const { format, access_mode, view_type } = t.params; + const { stage, format, access_mode, view_type } = t.params; const texture = t.createTextureTracked({ format, @@ -209,19 +215,19 @@ Parameters const code = ` @group(0) @binding(0) var t: texture_storage_2d_array<${format}, ${access_mode}>; @group(0) @binding(1) var result: u32; -@compute @workgroup_size(1) fn cs() { - result = textureNumLayers(t); +fn getValue() -> u32 { + return textureNumLayers(t); } `; const { baseArrayLayer, arrayLayerCount, expected } = getLayerSettingsAndExpected({ view_type, }); - const view = texture.createView({ + const viewDescription: GPUTextureViewDescriptor = { dimension: '2d-array', baseArrayLayer, arrayLayerCount, - }); + }; - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, code, texture, viewDescription, expected); }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureNumLevels.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureNumLevels.spec.ts index 471a462504d4..cc509233399c 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureNumLevels.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureNumLevels.spec.ts @@ -6,6 +6,7 @@ Returns the number of mip levels of a texture. import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; import { getTextureDimensionFromView } from '../../../../../util/texture/base.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; import { kSampleTypeInfo, WGSLTextureQueryTest } from './texture_utils.js'; @@ -66,6 +67,7 @@ Parameters 'texture_cube_array', ] as const) .beginSubcases() + .combine('stage', kShaderStages) .combine('sampled_type', ['f32', 'i32', 'u32'] as const) .combine('view_type', ['full', 'partial'] as const) // 1d textures can't have mipLevelCount > 0 @@ -75,7 +77,7 @@ Parameters t.skipIfTextureViewDimensionNotSupported(kTextureTypeToViewDimension[t.params.texture_type]); }) .fn(t => { - const { texture_type, sampled_type, view_type } = t.params; + const { stage, texture_type, sampled_type, view_type } = t.params; const { format } = kSampleTypeInfo[sampled_type]; const viewDimension = kTextureTypeToViewDimension[texture_type]; @@ -101,8 +103,8 @@ Parameters const code = ` @group(0) @binding(0) var t: ${texture_type}<${sampled_type}>; @group(0) @binding(1) var result: u32; -@compute @workgroup_size(1) fn cs() { - result = textureNumLevels(t); +fn getValue() -> u32 { + return textureNumLevels(t); } `; @@ -110,13 +112,13 @@ Parameters view_type, mipCount ); - const view = texture.createView({ + const viewDescription = { dimension: viewDimension, baseMipLevel, mipLevelCount, - }); + }; - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, code, texture, viewDescription, expected); }); g.test('depth') @@ -141,12 +143,14 @@ Parameters 'texture_depth_cube_array', ] as const) .combine('view_type', ['full', 'partial'] as const) + .beginSubcases() + .combine('stage', kShaderStages) ) .beforeAllSubcases(t => { t.skipIfTextureViewDimensionNotSupported(kTextureTypeToViewDimension[t.params.texture_type]); }) .fn(t => { - const { texture_type, view_type } = t.params; + const { stage, texture_type, view_type } = t.params; const viewDimension = kTextureTypeToViewDimension[texture_type]; const dimension = getTextureDimensionFromView(viewDimension); @@ -171,8 +175,8 @@ Parameters const code = ` @group(0) @binding(0) var t: ${texture_type}; @group(0) @binding(1) var result: u32; -@compute @workgroup_size(1) fn cs() { - result = textureNumLevels(t); +fn getValue() -> u32 { + return textureNumLevels(t); } `; @@ -180,11 +184,11 @@ Parameters view_type, mipCount ); - const view = texture.createView({ + const viewDescription = { dimension: viewDimension, baseMipLevel, mipLevelCount, - }); + }; - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, code, texture, viewDescription, expected); }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureNumSamples.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureNumSamples.spec.ts index a6314198529b..24f7f9b6997f 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureNumSamples.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureNumSamples.spec.ts @@ -5,6 +5,7 @@ Returns the number samples per texel in a multisampled texture. `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; import { kSampleTypeInfo, WGSLTextureQueryTest } from './texture_utils.js'; @@ -22,9 +23,14 @@ Parameters * t The multisampled texture. ` ) - .params(u => u.beginSubcases().combine('sampled_type', ['f32', 'i32', 'u32'] as const)) + .params(u => + u + .beginSubcases() + .combine('stage', kShaderStages) + .combine('sampled_type', ['f32', 'i32', 'u32'] as const) + ) .fn(t => { - const { sampled_type } = t.params; + const { stage, sampled_type } = t.params; const { format } = kSampleTypeInfo[sampled_type]; const sampleCount = 4; @@ -38,15 +44,13 @@ Parameters const code = ` @group(0) @binding(0) var t: texture_multisampled_2d<${sampled_type}>; @group(0) @binding(1) var result: u32; -@compute @workgroup_size(1) fn cs() { - result = textureNumSamples(t); +fn getValue() -> u32 { + return textureNumSamples(t); } `; const expected = [sampleCount]; - const view = texture.createView({}); - - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, code, texture, {}, expected); }); g.test('depth') @@ -59,7 +63,9 @@ Parameters * t The multisampled texture. ` ) + .params(u => u.beginSubcases().combine('stage', kShaderStages)) .fn(t => { + const { stage } = t.params; const sampleCount = 4; const texture = t.createTextureTracked({ format: 'depth32float', @@ -71,13 +77,11 @@ Parameters const code = ` @group(0) @binding(0) var t: texture_depth_multisampled_2d; @group(0) @binding(1) var result: u32; -@compute @workgroup_size(1) fn cs() { - result = textureNumSamples(t); +fn getValue() -> u32 { + return textureNumSamples(t); } `; const expected = [sampleCount]; - const view = texture.createView({}); - - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, code, texture, {}, expected); }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSample.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSample.spec.ts index 36492276a4b4..386101462eb8 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureSample.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureSample.spec.ts @@ -10,11 +10,10 @@ note: uniformity validation is covered in src/webgpu/shader/validation/uniformit import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; import { isDepthTextureFormat, - isEncodableTextureFormat, - kCompressedTextureFormats, kDepthStencilFormats, - kEncodableTextureFormats, + kAllTextureFormats, textureDimensionAndFormatCompatible, + isFilterableAsTextureF32, } from '../../../../../format_info.js'; import { TextureTestMixin } from '../../../../../gpu_test.js'; @@ -44,8 +43,6 @@ import { skipIfNeedsFilteringAndIsUnfilterable, } from './texture_utils.js'; -const kTestableColorFormats = [...kEncodableTextureFormats, ...kCompressedTextureFormats] as const; - export const g = makeTestGroup(TextureTestMixin(WGSLTextureSampleTest)); g.test('sampled_1d_coords') @@ -62,7 +59,7 @@ Parameters: ) .params(u => u - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => textureDimensionAndFormatCompatible('1d', t.format)) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) @@ -151,9 +148,10 @@ Parameters: ) .params(u => u - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('offset', [false, true] as const) @@ -249,11 +247,12 @@ Parameters: ) .params(u => u - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('dim', ['3d', 'cube'] as const) .filter(t => isSupportedViewFormatCombo(t.format, t.dim)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('modeW', kShortAddressModes) @@ -337,7 +336,7 @@ Parameters: const viewDescriptor = { dimension: viewDimension, }; - const textureType = getTextureTypeForTextureViewDimension(viewDimension)!; + const textureType = getTextureTypeForTextureViewDimension(viewDimension); const results = await doTextureCalls( t, texture, @@ -384,17 +383,18 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) - .combine('filt', ['nearest', 'linear'] as const) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('offset', [false, true] as const) .beginSubcases() .combine('samplePoints', kSamplePointMethods) ) + .beforeAllSubcases(t => { + t.skipIfDepthTextureCanNotBeUsedWithNonComparisonSampler(); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); + }) .fn(async t => { - const { format, samplePoints, modeU, modeV, filt: minFilter, offset } = t.params; + const { format, samplePoints, modeU, modeV, offset } = t.params; // We want at least 4 blocks or something wide enough for 3 mip levels. const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); @@ -408,9 +408,6 @@ Parameters: const sampler: GPUSamplerDescriptor = { addressModeU: kShortAddressModeToAddressMode[modeU], addressModeV: kShortAddressModeToAddressMode[modeV], - minFilter, - magFilter: minFilter, - mipmapFilter: minFilter, }; const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { @@ -419,7 +416,7 @@ Parameters: descriptor, derivatives: true, offset, - hashInputs: [format, samplePoints, modeU, modeV, minFilter, offset], + hashInputs: [format, samplePoints, modeU, modeV, offset], }).map(({ coords, derivativeMult, offset }) => { return { builtin: 'textureSample', @@ -478,9 +475,10 @@ Parameters: ) .params(u => u - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('offset', [false, true] as const) @@ -574,9 +572,10 @@ Parameters: ) .params(u => u - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('mode', kShortAddressModes) .beginSubcases() .combine('samplePoints', kCubeSamplePointMethods) @@ -674,23 +673,22 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combineWithParams([ { viewDimension: 'cube' }, { viewDimension: 'cube-array', A: 'i32' }, { viewDimension: 'cube-array', A: 'u32' }, ] as const) - .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .beginSubcases() .combine('samplePoints', kCubeSamplePointMethods) ) .beforeAllSubcases(t => { + t.skipIfDepthTextureCanNotBeUsedWithNonComparisonSampler(); t.skipIfTextureViewDimensionNotSupported(t.params.viewDimension); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { - const { format, viewDimension, samplePoints, A, mode, filt: minFilter } = t.params; + const { format, viewDimension, samplePoints, A, mode } = t.params; const size = chooseTextureSize({ minSize: 32, @@ -711,9 +709,6 @@ Parameters: addressModeU: kShortAddressModeToAddressMode[mode], addressModeV: kShortAddressModeToAddressMode[mode], addressModeW: kShortAddressModeToAddressMode[mode], - minFilter, - magFilter: minFilter, - mipmapFilter: minFilter, }; const calls: TextureCall[] = generateSamplePointsCube(50, { @@ -722,7 +717,7 @@ Parameters: descriptor, derivatives: true, arrayIndex: A ? { num: texture.depthOrArrayLayers / 6, type: A } : undefined, - hashInputs: [format, viewDimension, samplePoints, mode, minFilter], + hashInputs: [format, viewDimension, samplePoints, mode], }).map(({ coords, derivativeMult, arrayIndex }) => { return { builtin: 'textureSample', @@ -788,9 +783,6 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) - .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .combine('offset', [false, true] as const) .beginSubcases() @@ -798,8 +790,12 @@ Parameters: .combine('A', ['i32', 'u32'] as const) .combine('L', ['i32', 'u32'] as const) ) + .beforeAllSubcases(t => { + t.skipIfDepthTextureCanNotBeUsedWithNonComparisonSampler(); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); + }) .fn(async t => { - const { format, samplePoints, mode, filt: minFilter, A, L, offset } = t.params; + const { format, samplePoints, mode, A, L, offset } = t.params; // We want at least 4 blocks or something wide enough for 3 mip levels. const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); @@ -814,9 +810,6 @@ Parameters: const sampler: GPUSamplerDescriptor = { addressModeU: kShortAddressModeToAddressMode[mode], addressModeV: kShortAddressModeToAddressMode[mode], - minFilter, - magFilter: minFilter, - mipmapFilter: minFilter, }; const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { @@ -826,7 +819,7 @@ Parameters: derivatives: true, arrayIndex: { num: texture.depthOrArrayLayers, type: A }, offset, - hashInputs: [format, samplePoints, mode, minFilter, L, A, offset], + hashInputs: [format, samplePoints, mode, L, A, offset], }).map(({ coords, derivativeMult, arrayIndex, offset }) => { return { builtin: 'textureSample', @@ -882,19 +875,18 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) - .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .beginSubcases() .combine('samplePoints', kCubeSamplePointMethods) .combine('A', ['i32', 'u32'] as const) ) .beforeAllSubcases(t => { + t.skipIfDepthTextureCanNotBeUsedWithNonComparisonSampler(); t.skipIfTextureViewDimensionNotSupported('cube-array'); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { - const { format, samplePoints, A, mode, filt: minFilter } = t.params; + const { format, samplePoints, A, mode } = t.params; const viewDimension: GPUTextureViewDimension = 'cube-array'; const size = chooseTextureSize({ @@ -916,9 +908,6 @@ Parameters: addressModeU: kShortAddressModeToAddressMode[mode], addressModeV: kShortAddressModeToAddressMode[mode], addressModeW: kShortAddressModeToAddressMode[mode], - minFilter, - magFilter: minFilter, - mipmapFilter: minFilter, }; const calls: TextureCall[] = generateSamplePointsCube(50, { @@ -927,7 +916,7 @@ Parameters: descriptor, derivatives: true, arrayIndex: A ? { num: texture.depthOrArrayLayers / 6, type: A } : undefined, - hashInputs: [format, viewDimension, samplePoints, mode, minFilter], + hashInputs: [format, viewDimension, samplePoints, mode], }).map(({ coords, derivativeMult, arrayIndex }) => { return { builtin: 'textureSample', diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSampleBias.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSampleBias.spec.ts index f49322f878d6..8d36c13e928d 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleBias.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleBias.spec.ts @@ -8,7 +8,7 @@ Samples a texture with a bias to the mip level. `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; -import { kCompressedTextureFormats, kEncodableTextureFormats } from '../../../../../format_info.js'; +import { isFilterableAsTextureF32, kAllTextureFormats } from '../../../../../format_info.js'; import { TextureTestMixin } from '../../../../../gpu_test.js'; import { @@ -35,10 +35,12 @@ import { skipIfNeedsFilteringAndIsUnfilterable, } from './texture_utils.js'; -const kTestableColorFormats = [...kEncodableTextureFormats, ...kCompressedTextureFormats] as const; - export const g = makeTestGroup(TextureTestMixin(WGSLTextureSampleTest)); +// See comment "Issues with textureSampleBias" in texture_utils.ts +// 3 was chosen because it shows errors on M1 Mac +const kMinBlocksForTextureSampleBias = 3; + g.test('sampled_2d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturesamplebias') .desc( @@ -61,9 +63,10 @@ Parameters: ) .params(u => u - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('offset', [false, true] as const) @@ -77,8 +80,12 @@ Parameters: const { format, samplePoints, modeU, modeV, filt: minFilter, offset } = t.params; skipIfNeedsFilteringAndIsUnfilterable(t, minFilter, format); - // We want at least 4 blocks or something wide enough for 3 mip levels. - const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); + // We want at least something wide enough for 3 mip levels with more than 1 pixel at the smallest level + const [width, height] = chooseTextureSize({ + minSize: 8, + minBlocks: kMinBlocksForTextureSampleBias, + format, + }); const descriptor: GPUTextureDescriptor = { format, @@ -159,11 +166,12 @@ Parameters: ) .params(u => u - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('dim', ['3d', 'cube'] as const) .filter(t => isSupportedViewFormatCombo(t.format, t.dim)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('modeW', kShortAddressModes) @@ -248,7 +256,7 @@ Parameters: const viewDescriptor = { dimension: viewDimension, }; - const textureType = getTextureTypeForTextureViewDimension(viewDimension)!; + const textureType = getTextureTypeForTextureViewDimension(viewDimension); const results = await doTextureCalls( t, texture, @@ -296,9 +304,10 @@ Parameters: ) .params(u => u - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('offset', [false, true] as const) @@ -313,8 +322,12 @@ Parameters: const { format, samplePoints, A, modeU, modeV, filt: minFilter, offset } = t.params; skipIfNeedsFilteringAndIsUnfilterable(t, minFilter, format); - // We want at least 4 blocks or something wide enough for 3 mip levels. - const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); + // We want at least something wide enough for 3 mip levels with more than 1 pixel at the smallest level + const [width, height] = chooseTextureSize({ + minSize: 8, + minBlocks: kMinBlocksForTextureSampleBias, + format, + }); const depthOrArrayLayers = 4; const descriptor: GPUTextureDescriptor = { @@ -400,9 +413,10 @@ Parameters: ) .params(u => u - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('mode', kShortAddressModes) .beginSubcases() .combine('samplePoints', kCubeSamplePointMethods) diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompare.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompare.spec.ts index 27e55a8b189b..e34e66383f5c 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompare.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompare.spec.ts @@ -7,11 +7,7 @@ Samples a depth texture and compares the sampled depth values against a referenc import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; import { kCompareFunctions } from '../../../../../capability_info.js'; -import { - isDepthTextureFormat, - isEncodableTextureFormat, - kDepthStencilFormats, -} from '../../../../../format_info.js'; +import { isDepthTextureFormat, kDepthStencilFormats } from '../../../../../format_info.js'; import { checkCallResults, @@ -58,8 +54,6 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combine('filt', ['nearest', 'linear'] as const) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) @@ -68,6 +62,7 @@ Parameters: .combine('samplePoints', kSamplePointMethods) .combine('compare', kCompareFunctions) ) + .beforeAllSubcases(t => t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format)) .fn(async t => { const { format, samplePoints, modeU, modeV, filt: minFilter, compare, offset } = t.params; @@ -152,14 +147,13 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .beginSubcases() .combine('samplePoints', kCubeSamplePointMethods) .combine('compare', kCompareFunctions) ) + .beforeAllSubcases(t => t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format)) .fn(async t => { const { format, samplePoints, mode, filt: minFilter, compare } = t.params; @@ -258,8 +252,6 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combine('filt', ['nearest', 'linear'] as const) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) @@ -271,6 +263,7 @@ Parameters: ) .beforeAllSubcases(t => { t.skipIfTextureFormatNotSupported(t.params.format); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { const { format, samplePoints, A, modeU, modeV, filt: minFilter, compare, offset } = t.params; @@ -363,8 +356,6 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .beginSubcases() @@ -374,6 +365,7 @@ Parameters: ) .beforeAllSubcases(t => { t.skipIfTextureViewDimensionNotSupported('cube-array'); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { const { format, A, samplePoints, mode, filt: minFilter, compare } = t.params; diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompareLevel.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompareLevel.spec.ts index 61d093a638cd..81855908fddc 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompareLevel.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompareLevel.spec.ts @@ -13,11 +13,7 @@ The textureSampleCompareLevel function is the same as textureSampleCompare, exce import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; import { kCompareFunctions } from '../../../../../capability_info.js'; -import { - isDepthTextureFormat, - isEncodableTextureFormat, - kDepthStencilFormats, -} from '../../../../../format_info.js'; +import { isDepthTextureFormat, kDepthStencilFormats } from '../../../../../format_info.js'; import { checkCallResults, @@ -66,8 +62,6 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combine('filt', ['nearest', 'linear'] as const) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) @@ -76,6 +70,7 @@ Parameters: .combine('samplePoints', kSamplePointMethods) .combine('compare', kCompareFunctions) ) + .beforeAllSubcases(t => t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format)) .fn(async t => { const { format, @@ -170,14 +165,13 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .beginSubcases() .combine('samplePoints', kCubeSamplePointMethods) .combine('compare', kCompareFunctions) ) + .beforeAllSubcases(t => t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format)) .fn(async t => { const { format, stage, samplePoints, mode, filt: minFilter, compare } = t.params; @@ -277,8 +271,6 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combine('filt', ['nearest', 'linear'] as const) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) @@ -290,6 +282,7 @@ Parameters: ) .beforeAllSubcases(t => { t.skipIfTextureFormatNotSupported(t.params.format); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { const { @@ -393,8 +386,6 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .beginSubcases() @@ -404,6 +395,7 @@ Parameters: ) .beforeAllSubcases(t => { t.skipIfTextureViewDimensionNotSupported('cube-array'); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { const { format, A, stage, samplePoints, mode, filt: minFilter, compare } = t.params; diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSampleGrad.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSampleGrad.spec.ts index 8da6ffdfe9c1..2019cd3600b2 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleGrad.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleGrad.spec.ts @@ -6,7 +6,7 @@ Samples a texture using explicit gradients. `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; -import { kCompressedTextureFormats, kEncodableTextureFormats } from '../../../../../format_info.js'; +import { isFilterableAsTextureF32, kAllTextureFormats } from '../../../../../format_info.js'; import { appendComponentTypeForFormatToTextureType, @@ -34,8 +34,6 @@ import { WGSLTextureSampleTest, } from './texture_utils.js'; -const kTestableColorFormats = [...kEncodableTextureFormats, ...kCompressedTextureFormats] as const; - export const g = makeTestGroup(WGSLTextureSampleTest); g.test('sampled_2d_coords') @@ -62,9 +60,10 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('offset', [false, true] as const) @@ -161,11 +160,12 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('dim', ['3d', 'cube'] as const) .filter(t => isSupportedViewFormatCombo(t.format, t.dim)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('modeW', kShortAddressModes) @@ -251,7 +251,7 @@ Parameters: const viewDescriptor = { dimension: viewDimension, }; - const textureType = getTextureTypeForTextureViewDimension(viewDimension)!; + const textureType = getTextureTypeForTextureViewDimension(viewDimension); const results = await doTextureCalls( t, texture, @@ -301,9 +301,10 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('offset', [false, true] as const) @@ -407,9 +408,10 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('mode', kShortAddressModes) .beginSubcases() .combine('samplePoints', kCubeSamplePointMethods) diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts index 840bafcab223..78a3c6361379 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts @@ -7,10 +7,9 @@ Samples a texture. import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; import { isDepthTextureFormat, - isEncodableTextureFormat, - kCompressedTextureFormats, + isFilterableAsTextureF32, + kAllTextureFormats, kDepthStencilFormats, - kEncodableTextureFormats, } from '../../../../../format_info.js'; import { @@ -40,8 +39,6 @@ import { WGSLTextureSampleTest, } from './texture_utils.js'; -const kTestableColorFormats = [...kEncodableTextureFormats, ...kCompressedTextureFormats] as const; - export const g = makeTestGroup(WGSLTextureSampleTest); g.test('sampled_2d_coords') @@ -71,9 +68,10 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('offset', [false, true] as const) @@ -175,9 +173,10 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('modeU', kShortAddressModes) .combine('modeV', kShortAddressModes) .combine('offset', [false, true] as const) @@ -283,11 +282,12 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('dim', ['3d', 'cube'] as const) .filter(t => isSupportedViewFormatCombo(t.format, t.dim)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('mode', kShortAddressModes) .combine('offset', [false, true] as const) .filter(t => t.dim !== 'cube' || t.offset !== true) @@ -416,9 +416,10 @@ Parameters: .params(u => u .combine('stage', kShortShaderStages) - .combine('format', kTestableColorFormats) + .combine('format', kAllTextureFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('filt', ['nearest', 'linear'] as const) + .filter(t => t.filt === 'nearest' || isFilterableAsTextureF32(t.format)) .combine('mode', kShortAddressModes) .beginSubcases() .combine('samplePoints', kCubeSamplePointMethods) @@ -531,20 +532,18 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) - .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .combine('offset', [false, true] as const) .beginSubcases() .combine('samplePoints', kSamplePointMethods) .combine('L', ['i32', 'u32'] as const) ) - .beforeAllSubcases(t => - skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) - ) + .beforeAllSubcases(t => { + t.skipIfDepthTextureCanNotBeUsedWithNonComparisonSampler(); + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format); + }) .fn(async t => { - const { format, stage, samplePoints, mode, filt: minFilter, L, offset } = t.params; + const { format, stage, samplePoints, mode, L, offset } = t.params; // We want at least 4 blocks or something wide enough for 3 mip levels. const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); @@ -558,9 +557,6 @@ Parameters: const sampler: GPUSamplerDescriptor = { addressModeU: kShortAddressModeToAddressMode[mode], addressModeV: kShortAddressModeToAddressMode[mode], - minFilter, - magFilter: minFilter, - mipmapFilter: minFilter, }; const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { @@ -569,7 +565,7 @@ Parameters: descriptor, mipLevel: { num: texture.mipLevelCount, type: L }, offset, - hashInputs: [stage, format, samplePoints, mode, minFilter, L, offset], + hashInputs: [stage, format, samplePoints, mode, L, offset], }).map(({ coords, mipLevel, offset }) => { return { builtin: 'textureSampleLevel', @@ -638,9 +634,6 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) - .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .combine('offset', [false, true] as const) .beginSubcases() @@ -648,11 +641,12 @@ Parameters: .combine('A', ['i32', 'u32'] as const) .combine('L', ['i32', 'u32'] as const) ) - .beforeAllSubcases(t => - skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) - ) + .beforeAllSubcases(t => { + t.skipIfDepthTextureCanNotBeUsedWithNonComparisonSampler(); + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format); + }) .fn(async t => { - const { format, stage, samplePoints, mode, filt: minFilter, A, L, offset } = t.params; + const { format, stage, samplePoints, mode, A, L, offset } = t.params; // We want at least 4 blocks or something wide enough for 3 mip levels. const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); @@ -667,9 +661,6 @@ Parameters: const sampler: GPUSamplerDescriptor = { addressModeU: kShortAddressModeToAddressMode[mode], addressModeV: kShortAddressModeToAddressMode[mode], - minFilter, - magFilter: minFilter, - mipmapFilter: minFilter, }; const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { @@ -679,7 +670,7 @@ Parameters: arrayIndex: { num: texture.depthOrArrayLayers, type: A }, mipLevel: { num: texture.mipLevelCount, type: L }, offset, - hashInputs: [stage, format, samplePoints, mode, minFilter, L, A, offset], + hashInputs: [stage, format, samplePoints, mode, L, A, offset], }).map(({ coords, mipLevel, arrayIndex, offset }) => { return { builtin: 'textureSampleLevel', @@ -749,25 +740,23 @@ Parameters: .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) - // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. - .filter(t => isEncodableTextureFormat(t.format)) .combineWithParams([ { viewDimension: 'cube' }, { viewDimension: 'cube-array', A: 'i32' }, { viewDimension: 'cube-array', A: 'u32' }, ] as const) - .combine('filt', ['nearest', 'linear'] as const) .combine('mode', kShortAddressModes) .beginSubcases() .combine('samplePoints', kCubeSamplePointMethods) .combine('L', ['i32', 'u32'] as const) ) .beforeAllSubcases(t => { + t.skipIfDepthTextureCanNotBeUsedWithNonComparisonSampler(); skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format); t.skipIfTextureViewDimensionNotSupported(t.params.viewDimension); }) .fn(async t => { - const { format, stage, viewDimension, samplePoints, A, L, mode, filt: minFilter } = t.params; + const { format, stage, viewDimension, samplePoints, A, L, mode } = t.params; const size = chooseTextureSize({ minSize: 32, @@ -787,9 +776,6 @@ Parameters: addressModeU: kShortAddressModeToAddressMode[mode], addressModeV: kShortAddressModeToAddressMode[mode], addressModeW: kShortAddressModeToAddressMode[mode], - minFilter, - magFilter: minFilter, - mipmapFilter: minFilter, }; const calls: TextureCall[] = generateSamplePointsCube(50, { @@ -798,7 +784,7 @@ Parameters: descriptor, mipLevel: { num: texture.mipLevelCount - 1, type: L }, arrayIndex: A ? { num: texture.depthOrArrayLayers / 6, type: A } : undefined, - hashInputs: [stage, format, viewDimension, samplePoints, mode, minFilter], + hashInputs: [stage, format, viewDimension, samplePoints, mode], }).map(({ coords, mipLevel, arrayIndex }) => { return { builtin: 'textureSampleLevel', diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureStore.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureStore.spec.ts index e955b82ed603..1dc7f8139d79 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureStore.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureStore.spec.ts @@ -12,7 +12,8 @@ If an out-of-bounds access occurs, the built-in function should not be executed. `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; -import { unreachable, iterRange } from '../../../../../../common/util/util.js'; +import { unreachable, iterRange, range } from '../../../../../../common/util/util.js'; +import { kTextureFormatInfo } from '../../../../../format_info.js'; import { GPUTest, TextureTestMixin } from '../../../../../gpu_test.js'; import { kFloat32Format, @@ -21,115 +22,17 @@ import { pack4x8unorm, pack4x8snorm, } from '../../../../../util/conversion.js'; -import { virtualMipSize } from '../../../../../util/texture/base.js'; +import { align, clamp } from '../../../../../util/math.js'; +import { getTextureDimensionFromView, virtualMipSize } from '../../../../../util/texture/base.js'; import { TexelFormats } from '../../../../types.js'; -import { generateCoordBoundaries } from './utils.js'; +const kDims = ['1d', '2d', '3d'] as const; +const kViewDimensions = ['1d', '2d', '2d-array', '3d'] as const; export const g = makeTestGroup(TextureTestMixin(GPUTest)); -g.test('store_1d_coords') - .specURL('https://www.w3.org/TR/WGSL/#texturestore') - .desc( - ` -C is i32 or u32 - -fn textureStore(t: texture_storage_1d, coords: C, value: vec4) - -Parameters: - * t The sampled, depth, or external texture to sample. - * s The sampler type. - * coords The texture coordinates used for sampling. - * value The new texel value -` - ) - .params(u => - u - .combineWithParams(TexelFormats) - .beginSubcases() - .combine('coords', generateCoordBoundaries(1)) - .combine('C', ['i32', 'u32'] as const) - ) - .unimplemented(); - -g.test('store_2d_coords') - .specURL('https://www.w3.org/TR/WGSL/#texturestore') - .desc( - ` -C is i32 or u32 - -fn textureStore(t: texture_storage_2d, coords: vec2, value: vec4) - -Parameters: - * t The sampled, depth, or external texture to sample. - * s The sampler type. - * coords The texture coordinates used for sampling. - * value The new texel value -` - ) - .params(u => - u - .combineWithParams(TexelFormats) - .beginSubcases() - .combine('coords', generateCoordBoundaries(2)) - .combine('C', ['i32', 'u32'] as const) - ) - .unimplemented(); - -g.test('store_array_2d_coords') - .specURL('https://www.w3.org/TR/WGSL/#texturestore') - .desc( - ` -C is i32 or u32 - -fn textureStore(t: texture_storage_2d_array, coords: vec2, array_index: C, value: vec4) - -Parameters: - * t The sampled, depth, or external texture to sample. - * s The sampler type. - * array_index The 0-based texture array index - * coords The texture coordinates used for sampling. - * value The new texel value -` - ) - .params( - u => - u - .combineWithParams(TexelFormats) - .beginSubcases() - .combine('coords', generateCoordBoundaries(2)) - .combine('C', ['i32', 'u32'] as const) - .combine('C_value', [-1, 0, 1, 2, 3, 4] as const) - /* array_index not param'd as out-of-bounds is implementation specific */ - ) - .unimplemented(); - -g.test('store_3d_coords') - .specURL('https://www.w3.org/TR/WGSL/#texturestore') - .desc( - ` -C is i32 or u32 - -fn textureStore(t: texture_storage_3d, coords: vec3, value: vec4) - -Parameters: - * t The sampled, depth, or external texture to sample. - * s The sampler type. - * coords The texture coordinates used for sampling. - * value The new texel value -` - ) - .params(u => - u - .combineWithParams(TexelFormats) - .beginSubcases() - .combine('coords', generateCoordBoundaries(3)) - .combine('C', ['i32', 'u32'] as const) - ) - .unimplemented(); - -// Returns shader input values for texel format tests. -// Values are intentionally simple to avoid rounding issues. +// We require a few values that are out of range for a given type +// so we can check clamping behavior. function inputArray(format: string): number[] { switch (format) { case 'rgba8snorm': @@ -138,24 +41,26 @@ function inputArray(format: string): number[] { case 'bgra8unorm': return [-0.1, 0, 0.2, 0.4, 0.6, 0.8, 1.0, 1.1]; case 'rgba8uint': + return [0, 8, 16, 24, 32, 64, 100, 128, 200, 255, 256, 512]; case 'rgba16uint': + return [0, 8, 16, 24, 32, 64, 100, 128, 200, 255, 0xffff, 0x1ffff]; case 'rgba32uint': case 'r32uint': case 'rg32uint': - // Stick within 8-bit ranges for simplicity. - return [0, 8, 16, 24, 32, 64, 100, 128, 200, 255]; + return [0, 8, 16, 24, 32, 64, 100, 128, 200, 255, 256, 512, 0xffffffff]; case 'rgba8sint': + return [-128, -100, -64, -32, -16, -8, 0, 8, 16, 32, 64, 100, 127]; case 'rgba16sint': - case 'rgba32sint': + return [-32768, -32769, -100, -64, -32, -16, -8, 0, 8, 16, 32, 64, 100, 127, 0x7fff, 0x8000]; case 'r32sint': case 'rg32sint': - // Stick within 8-bit ranges for simplicity. - return [-128, -100, -64, -32, -16, -8, 0, 8, 16, 32, 64, 100, 127]; + case 'rgba32sint': + return [-0x8000000, -32769, -100, -64, -32, -16, -8, 0, 8, 16, 32, 64, 100, 127, 0x7ffffff]; case 'rgba16float': case 'rgba32float': case 'r32float': case 'rg32float': - // Stick with simple values. + // Stick with simple values to avoid rounding issues. return [-100, -50, -32, -16, -8, -1, 0, 1, 8, 16, 32, 50, 100]; default: unreachable(`unhandled format ${format}`); @@ -165,8 +70,29 @@ function inputArray(format: string): number[] { } g.test('texel_formats') - .desc(`Test storage of texel formats`) - .params(u => u.combineWithParams([...TexelFormats, { format: 'bgra8unorm', _shaderType: 'f32' }])) + .desc( + ` + Test storage of texel formats + + - test values make it through. + - test out of range values get clamped. + - test 1d, 2d, 2d-array, 3d. + - test all storage formats. + ` + ) + .params(u => + u + .combineWithParams([...TexelFormats, { format: 'bgra8unorm', _shaderType: 'f32' }]) + .combine('viewDimension', kViewDimensions) + // Note: We can't use writable storage textures in a vertex stage. + .combine('stage', ['compute', 'fragment'] as const) + .combine('access', ['write', 'read_write'] as const) + .unless( + t => + t.access === 'read_write' && + !kTextureFormatInfo[t.format as GPUTextureFormat].color?.readWriteStorage + ) + ) .beforeAllSubcases(t => { if (t.params.format === 'bgra8unorm') { t.selectDeviceOrSkipTestCase('bgra8unorm-storage'); @@ -175,72 +101,84 @@ g.test('texel_formats') } }) .fn(t => { - const { format, _shaderType } = t.params; + const { format, stage, access, viewDimension, _shaderType } = t.params; const values = inputArray(format); - let numChannels = 4; - switch (format) { - case 'r32uint': - case 'r32sint': - case 'r32float': - numChannels = 1; - break; - case 'rg32uint': - case 'rg32sint': - case 'rg32float': - numChannels = 2; - break; - default: - break; - } - - let zeroVal = ``; - if (numChannels > 1) { - zeroVal = `val[idx % ${numChannels}] = 0;`; - } - - let wgsl = ` -const range = array(`; - for (const v of values) { - wgsl += `${v},\n`; - } - - wgsl += ` -); + const suffix = format.endsWith('sint') ? 'i' : format.endsWith('uint') ? 'u' : 'f'; + const swizzleWGSL = viewDimension === '1d' ? 'x' : viewDimension === '3d' ? 'xyz' : 'xy'; + const layerWGSL = viewDimension === '2d-array' ? ', gid.z' : ''; + const wgsl = ` +const range = array(${values.map(v => `${v}${suffix}`).join(',')}); @group(0) @binding(0) -var tex : texture_storage_1d<${format}, write>; +var tex : texture_storage_${viewDimension.replace('-', '_')}<${format}, ${access}>; + +fn setValue(gid: vec3u) { + let ndx = gid.x + gid.y + gid.z; + let vecVal = vec4( + range[(ndx + 0) % ${values.length}], + range[(ndx + 1) % ${values.length}], + range[(ndx + 2) % ${values.length}], + range[(ndx + 3) % ${values.length}], + ); + var val = vec4<${_shaderType}>(vecVal); + let coord = gid.${swizzleWGSL}; + textureStore(tex, coord${layerWGSL}, val); +} @compute @workgroup_size(${values.length}) -fn main(@builtin(global_invocation_id) gid : vec3u) { - let idx = gid.x; - let scalarVal = range[idx]; - let vecVal = vec4(scalarVal); - var val = vec4<${_shaderType}>(vecVal); - ${zeroVal} - textureStore(tex, gid.x, val); +fn cs(@builtin(global_invocation_id) gid : vec3u) { + setValue(gid); +} + +struct VOut { + @builtin(position) pos: vec4f, + @location(0) @interpolate(flat, either) z: u32, +} +@vertex fn vs( + @builtin(vertex_index) vNdx: u32, + @builtin(instance_index) iNdx: u32, +) -> VOut { + let pos = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1)); + return VOut(vec4f(pos[vNdx], 0, 1), iNdx); +} + +@fragment fn fs(v: VOut) -> @location(0) vec4f { + setValue(vec3u(u32(v.pos.x), u32(v.pos.y), v.z)); + return vec4f(0); } `; - const numTexels = values.length; - const textureSize: GPUExtent3D = { width: numTexels, height: 1, depthOrArrayLayers: 1 }; + const textureSize = [ + values.length, + viewDimension === '1d' ? 1 : values.length, + viewDimension === '2d-array' || viewDimension === '3d' ? values.length : 1, + ] as const; + const dimension = getTextureDimensionFromView(viewDimension); const texture = t.createTextureTracked({ format: format as GPUTextureFormat, - dimension: '1d', size: textureSize, mipLevelCount: 1, + dimension, usage: GPUTextureUsage.STORAGE_BINDING | GPUTextureUsage.COPY_SRC, }); - const pipeline = t.device.createComputePipeline({ - layout: 'auto', - compute: { - module: t.device.createShaderModule({ - code: wgsl, - }), - entryPoint: 'main', - }, + const module = t.device.createShaderModule({ + code: wgsl, }); + + const pipeline = + stage === 'compute' + ? t.device.createComputePipeline({ + layout: 'auto', + compute: { module }, + }) + : t.device.createRenderPipeline({ + layout: 'auto', + vertex: { module }, + fragment: { module, targets: [{ format: 'rgba8unorm' }] }, + }); + const bg = t.device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [ @@ -248,18 +186,44 @@ fn main(@builtin(global_invocation_id) gid : vec3u) { binding: 0, resource: texture.createView({ format: format as GPUTextureFormat, - dimension: '1d', + dimension: viewDimension, }), }, ], }); const encoder = t.device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(pipeline); - pass.setBindGroup(0, bg); - pass.dispatchWorkgroups(1, 1, 1); - pass.end(); + switch (stage) { + case 'compute': { + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline as GPUComputePipeline); + pass.setBindGroup(0, bg); + pass.dispatchWorkgroups(...textureSize); + pass.end(); + break; + } + case 'fragment': { + const renderTarget = t.createTextureTracked({ + size: textureSize.slice(0, 2), + format: 'rgba8unorm', + usage: GPUTextureUsage.RENDER_ATTACHMENT, + }); + const pass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: renderTarget.createView(), + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + pass.setPipeline(pipeline as GPURenderPipeline); + pass.setBindGroup(0, bg); + pass.draw(3, textureSize[2]); + pass.end(); + break; + } + } t.queue.submit([encoder.finish()]); let bytesPerTexel = 4; @@ -281,45 +245,67 @@ fn main(@builtin(global_invocation_id) gid : vec3u) { break; } - let zeroChannel = 0; const buffer = t.copyWholeTextureToNewBufferSimple(texture, 0); - const uintsPerTexel = bytesPerTexel / 4; + const u32sPerTexel = bytesPerTexel / 4; + const bytesPerRow = align(textureSize[0] * bytesPerTexel, 256); + const texelsPerRow = bytesPerRow / bytesPerTexel; + const texelsPerSlice = texelsPerRow * textureSize[1]; + const getValue = (i: number) => values[i % values.length]; + const clampedPack4x8unorm = (...v: number[]) => { + const c = v.map(v => clamp(v, { min: 0, max: 1 })); + return pack4x8unorm(c[0], c[1], c[2], c[3]); + }; + const clampedPack4x8snorm = (...v: number[]) => { + const c = v.map(v => clamp(v, { min: -1, max: 1 })); + return pack4x8snorm(c[0], c[1], c[2], c[3]); + }; const expected = new Uint32Array([ - ...iterRange(numTexels * uintsPerTexel, x => { - const idx = Math.floor(x / uintsPerTexel); - const channel = idx % numChannels; - zeroChannel = zeroChannel % numChannels; - const shaderVal = values[idx]; + // iterate over each u32 + ...iterRange(buffer.size / 4, i => { + const texelId = (i / u32sPerTexel) | 0; + const z = (texelId / texelsPerSlice) | 0; + const y = ((texelId / texelsPerRow) | 0) % textureSize[1]; + const x = texelId % texelsPerRow; + // buffer is padded to 256 per row so when x is out of range just return 0 + if (x >= textureSize[0]) { + return 0; + } + const id = x + y + z; + const unit = i % u32sPerTexel; switch (format) { case 'rgba8unorm': { - const vals = [shaderVal, shaderVal, shaderVal, shaderVal]; - vals[zeroChannel++] = 0; - return pack4x8unorm(vals[0], vals[1], vals[2], vals[3]); + const vals = range(4, i => getValue(id + i)); + return clampedPack4x8unorm(vals[0], vals[1], vals[2], vals[3]); } case 'bgra8unorm': { - const vals = [shaderVal, shaderVal, shaderVal, shaderVal]; - vals[zeroChannel++] = 0; - return pack4x8unorm(vals[2], vals[1], vals[0], vals[3]); + const vals = range(4, i => getValue(id + i)); + return clampedPack4x8unorm(vals[2], vals[1], vals[0], vals[3]); } case 'rgba8snorm': { - const vals = [shaderVal, shaderVal, shaderVal, shaderVal]; - vals[zeroChannel++] = 0; - return pack4x8snorm(vals[0], vals[1], vals[2], vals[3]); + const vals = range(4, i => getValue(id + i)); + return clampedPack4x8snorm(vals[0], vals[1], vals[2], vals[3]); } case 'r32uint': + return clamp(getValue(id), { min: 0, max: 0xffffffff }); case 'r32sint': - return shaderVal; + return clamp(getValue(id), { min: -0x80000000, max: 0x7fffffff }); case 'rg32uint': case 'rgba32uint': + return clamp(getValue(id + unit), { min: 0, max: 0xffffffff }); case 'rg32sint': - case 'rgba32sint': { - const maskedVal = channel === zeroChannel++ ? 0 : shaderVal; - return maskedVal; + case 'rgba32sint': + return clamp(getValue(id + unit), { min: -0x80000000, max: 0x7fffffff }); + case 'rgba8uint': { + const vals = range(4, i => clamp(getValue(id + i), { min: 0, max: 255 })); + return ( + ((vals[3] & 0xff) << 24) | + ((vals[2] & 0xff) << 16) | + ((vals[1] & 0xff) << 8) | + (vals[0] & 0xff) + ); } - case 'rgba8uint': case 'rgba8sint': { - const vals = [shaderVal, shaderVal, shaderVal, shaderVal]; - vals[zeroChannel++] = 0; + const vals = range(4, i => clamp(getValue(id + i), { min: -0x80, max: 0x7f })); return ( ((vals[3] & 0xff) << 24) | ((vals[2] & 0xff) << 16) | @@ -327,55 +313,31 @@ fn main(@builtin(global_invocation_id) gid : vec3u) { (vals[0] & 0xff) ); } - case 'rgba16uint': - case 'rgba16sint': { - // 4 channels split over 2 uint32s. - // Determine if this pair has the zero channel. - const vals = [shaderVal, shaderVal]; - const lowChannels = (x & 0x1) === 0; - if (lowChannels) { - if (zeroChannel < 2) { - vals[zeroChannel] = 0; - } - } else { - if (zeroChannel >= 2) { - vals[zeroChannel - 2] = 0; - } - zeroChannel++; - } + case 'rgba16uint': { + const vals = range(2, i => clamp(getValue(id + unit * 2 + i), { min: 0, max: 0xffff })); return ((vals[1] & 0xffff) << 16) | (vals[0] & 0xffff); } - case 'r32float': { - return numberToFloatBits(shaderVal, kFloat32Format); + case 'rgba16sint': { + const vals = range(2, i => + clamp(getValue(id + unit * 2 + i), { min: -0x8000, max: 0x7fff }) + ); + return ((vals[1] & 0xffff) << 16) | (vals[0] & 0xffff); } + case 'r32float': case 'rg32float': case 'rgba32float': { - const maskedVal = channel === zeroChannel++ ? 0 : shaderVal; - return numberToFloatBits(maskedVal, kFloat32Format); + return numberToFloatBits(getValue(id + unit), kFloat32Format); } case 'rgba16float': { - // 4 channels split over 2 uint32s. - // Determine if this pair has the zero channel. - const bits = numberToFloatBits(shaderVal, kFloat16Format); - const vals = [bits, bits]; - const lowChannels = (x & 0x1) === 0; - if (lowChannels) { - if (zeroChannel < 2) { - vals[zeroChannel] = 0; - } - } else { - if (zeroChannel >= 2) { - vals[zeroChannel - 2] = 0; - } - zeroChannel++; - } + const vals = range(2, i => + numberToFloatBits(getValue(id + unit * 2 + i), kFloat16Format) + ); return ((vals[1] & 0xffff) << 16) | (vals[0] & 0xffff); } default: unreachable(`unhandled format ${format}`); break; } - return 0; }), ]); t.expectGPUBufferValuesEqual(buffer, expected); @@ -599,8 +561,6 @@ function getMipTexels(numTexels: number, dim: GPUTextureDimension, mip: number): return texels; } -const kDims = ['1d', '2d', '3d'] as const; - g.test('out_of_bounds') .desc('Test that textureStore on out-of-bounds coordinates have no effect') .params(u => diff --git a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.spec.ts new file mode 100644 index 000000000000..483b8f36a456 --- /dev/null +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.spec.ts @@ -0,0 +1,291 @@ +export const description = ` +Tests for texture_utils.ts +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { assert } from '../../../../../../common/util/util.js'; +import { isMultisampledTextureFormat, kDepthStencilFormats } from '../../../../../format_info.js'; +import { GPUTest } from '../../../../../gpu_test.js'; +import { getTextureDimensionFromView, virtualMipSize } from '../../../../../util/texture/base.js'; +import { + kTexelRepresentationInfo, + PerTexelComponent, + TexelRepresentationInfo, +} from '../../../../../util/texture/texel_data.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; + +import { + chooseTextureSize, + createTextureWithRandomDataAndGetTexels, + graphWeights, + isSupportedViewFormatCombo, + makeRandomDepthComparisonTexelGenerator, + queryMipLevelMixWeightsForDevice, + readTextureToTexelViews, + texelsApproximatelyEqual, +} from './texture_utils.js'; + +export const g = makeTestGroup(GPUTest); + +function texelFormat(texel: Readonly>, rep: TexelRepresentationInfo) { + return rep.componentOrder.map(component => `${component}: ${texel[component]}`).join(', '); +} + +g.test('createTextureWithRandomDataAndGetTexels_with_generator') + .desc( + ` + Test createTextureWithRandomDataAndGetTexels with a generator. Generators + are only used with textureXXXCompare builtins as we need specific random + values to test these builtins with a depth reference value. + ` + ) + .params(u => + u + .combine('format', kDepthStencilFormats) + .combine('viewDimension', ['2d', '2d-array', 'cube', 'cube-array'] as const) + .filter(t => isSupportedViewFormatCombo(t.format, t.viewDimension)) + ) + .beforeAllSubcases(t => { + t.skipIfTextureViewDimensionNotSupported(t.params.viewDimension); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); + }) + .fn(async t => { + const { format, viewDimension } = t.params; + const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format, viewDimension }); + const descriptor: GPUTextureDescriptor = { + format, + dimension: getTextureDimensionFromView(viewDimension), + size, + mipLevelCount: 3, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + ...(t.isCompatibility && { textureBindingViewDimension: viewDimension }), + }; + await createTextureWithRandomDataAndGetTexels(t, descriptor, { + generator: makeRandomDepthComparisonTexelGenerator(descriptor, 'equal'), + }); + // We don't expect any particular results. We just expect no validation errors. + }); + +g.test('readTextureToTexelViews') + .desc('test readTextureToTexelViews for various formats and dimensions') + .params(u => + u + .combineWithParams([ + { srcFormat: 'r8unorm', texelViewFormat: 'rgba32float' }, + { srcFormat: 'r8sint', texelViewFormat: 'rgba32sint' }, + { srcFormat: 'r8uint', texelViewFormat: 'rgba32uint' }, + { srcFormat: 'rgba32float', texelViewFormat: 'rgba32float' }, + { srcFormat: 'rgba32uint', texelViewFormat: 'rgba32uint' }, + { srcFormat: 'rgba32sint', texelViewFormat: 'rgba32sint' }, + { srcFormat: 'depth24plus', texelViewFormat: 'rgba32float' }, + { srcFormat: 'depth24plus', texelViewFormat: 'r32float' }, + { srcFormat: 'depth24plus-stencil8', texelViewFormat: 'r32float' }, + { srcFormat: 'stencil8', texelViewFormat: 'rgba32sint' }, + ] as const) + .combine('viewDimension', ['1d', '2d', '2d-array', '3d', 'cube', 'cube-array'] as const) + .filter(t => isSupportedViewFormatCombo(t.srcFormat, t.viewDimension)) + .combine('sampleCount', [1, 4] as const) + .unless( + t => + t.sampleCount > 1 && + (!isMultisampledTextureFormat(t.srcFormat) || t.viewDimension !== '2d') + ) + ) + .beforeAllSubcases(t => { + t.skipIfTextureViewDimensionNotSupported(t.params.viewDimension); + }) + .fn(async t => { + const { srcFormat, texelViewFormat, viewDimension, sampleCount } = t.params; + const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format: srcFormat, viewDimension }); + const descriptor: GPUTextureDescriptor = { + format: srcFormat, + dimension: getTextureDimensionFromView(viewDimension), + size, + mipLevelCount: viewDimension === '1d' || sampleCount > 1 ? 1 : 3, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + sampleCount, + ...(t.isCompatibility && { textureBindingViewDimension: viewDimension }), + }; + const { texels: expectedTexelViews, texture } = await createTextureWithRandomDataAndGetTexels( + t, + descriptor + ); + const actualTexelViews = await readTextureToTexelViews(t, texture, descriptor, texelViewFormat); + + assert(actualTexelViews.length === expectedTexelViews.length, 'num mip levels match'); + + const errors = []; + for (let mipLevel = 0; mipLevel < actualTexelViews.length; ++mipLevel) { + const actualMipLevelTexelView = actualTexelViews[mipLevel]; + const expectedMipLevelTexelView = expectedTexelViews[mipLevel]; + const mipLevelSize = virtualMipSize(texture.dimension, size, mipLevel); + + const actualRep = kTexelRepresentationInfo[actualMipLevelTexelView.format]; + const expectedRep = kTexelRepresentationInfo[expectedMipLevelTexelView.format]; + + for (let z = 0; z < mipLevelSize[2]; ++z) { + for (let y = 0; y < mipLevelSize[1]; ++y) { + for (let x = 0; x < mipLevelSize[0]; ++x) { + const actual = actualMipLevelTexelView.color({ x, y, z }); + const expected = expectedMipLevelTexelView.color({ x, y, z }); + // This currently expects the exact same values in actual vs expected. + // It's possible this needs to be relaxed slightly but only for non-integer formats. + // For now, if the tests pass everywhere, we'll keep it at 0 tolerance. + const maxFractionalDiff = 0; + if ( + !texelsApproximatelyEqual( + actual, + actualMipLevelTexelView.format, + expected, + expectedMipLevelTexelView.format, + maxFractionalDiff + ) + ) { + const actualStr = texelFormat(actual, actualRep); + const expectedStr = texelFormat(expected, expectedRep); + errors.push( + `texel at ${x}, ${y}, ${z}, expected: ${expectedStr}, actual: ${actualStr}` + ); + } + } + } + } + + assert(errors.length === 0, errors.join('\n')); + } + }); + +function validateWeights(t: GPUTest, stage: string, builtin: string, weights: number[]) { + const kNumMixSteps = weights.length - 1; + const showWeights = () => ` +${weights.map((v, i) => `${i.toString().padStart(2)}: ${v}`).join('\n')} + +e = expected +A = actual +${graphWeights(32, weights)} +`; + + t.expect( + weights[0] === 0, + `stage: ${stage}, ${builtin}, weight 0 expected 0 but was ${weights[0]}\n${showWeights()}` + ); + t.expect( + weights[kNumMixSteps] === 1, + `stage: ${stage}, ${builtin}, top weight expected 1 but was ${ + weights[kNumMixSteps] + }\n${showWeights()}` + ); + + const dx = 1 / kNumMixSteps; + for (let i = 0; i < kNumMixSteps; ++i) { + const dy = weights[i + 1] - weights[i]; + // dy / dx because dy might be 0 + const slope = dy / dx; + + // Validate the slope is not going down. + assert( + slope >= 0, + `stage: ${stage}, ${builtin}, weight[${i}] was not <= weight[${i + 1}]\n${showWeights()}` + ); + + // Validate the slope is not going up too steeply. + // The correct slope is 1 / kNumMixSteps but Mac AMD and Mac Intel + // have the wrong mix weights. 2 is enough to pass Mac AMD which we + // decided is ok but will fail on Mac Intel in compute stage which we + // decides is not ok. + assert( + slope <= 2, + `stage: ${stage}, ${builtin}, slope from weight[${i}] to weight[${ + i + 1 + }] is > 2.\n${showWeights()}` + ); + } + + // Test that we don't have a mostly flat set of weights. + // Note: Ideally every value is unique but 66% is enough to pass AMD Mac + // which we decided was ok but high enough to fail Intel Mac in a compute stage + // which we decided is not ok. + const kMinPercentUniqueWeights = 66; + t.expect( + new Set(weights).size >= ((weights.length * kMinPercentUniqueWeights * 0.01) | 0), + `stage: ${stage}, ${builtin}, expected at least ~${kMinPercentUniqueWeights}% unique weights\n${showWeights()}` + ); +} + +g.test('weights') + .desc( + ` +Test the mip level weights are linear. + +Given 2 mip levels, textureSampleLevel(....., mipLevel) should return +mix(colorFromLevel0, colorFromLevel1, mipLevel). + +Similarly, textureSampleGrad(...., ddx, ...) where ddx is +vec2(mix(1.0, 2.0, mipLevel) / textureWidth, 0) should so return +mix(colorFromLevel0, colorFromLevel1, mipLevel). + +If we put 0,0,0,0 in level 0 and 1,1,1,1 in level 1 then we should arguably +be able to assert + + for (mipLevel = 0; mipLevel <= 1, mipLevel += 0.01) { + assert(textureSampleLevel(t, s, vec2f(0.5), mipLevel) === mipLevel) + ddx = vec2(mix(1.0, 2.0, mipLevel) / textureWidth, 0) + assert(textureSampleGrad(t, s, vec2f(0.5), ddx, vec2f(0)) === mipLevel) + } + +Unfortunately, the GPUs do not do this. In particular: + +AMD Mac goes like this: Not great but we allow it + + +----------------+ + | ***| + | ** | + | * | + | ** | + | ** | + | * | + | ** | + |*** | + +----------------+ + + Intel Mac goes like this in a compute stage + + +----------------+ + | *******| + | * | + | * | + | * | + | * | + | * | + | * | + |******* | + +----------------+ + +Where as they should go like this + + +----------------+ + | **| + | ** | + | ** | + | ** | + | ** | + | ** | + | ** | + |** | + +----------------+ + +To make the texture builtin tests pass, they use the mix weights we query from the GPU +even if they are arguably bad. This test is to surface the failure of the GPU +to use mix weights the approximate a linear interpolation. + +We allow the AMD case as but disallow extreme Intel case. WebGPU implementations +are supposed to work around this issue by poly-filling on devices that fail this test. +` + ) + .params(u => u.combine('stage', kShaderStages)) + .fn(async t => { + const { stage } = t.params; + const weights = await queryMipLevelMixWeightsForDevice(t, t.params.stage); + validateWeights(t, stage, 'textureSampleLevel', weights.sampleLevelWeights); + validateWeights(t, stage, 'textureSampleGrad', weights.softwareMixToGPUMixGradWeights); + }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts index b01f3a5e758f..29259f5d32a3 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -3,11 +3,13 @@ import { assert, range, unreachable } from '../../../../../../common/util/util.j import { Float16Array } from '../../../../../../external/petamoriken/float16/float16.js'; import { EncodableTextureFormat, + is32Float, isCompressedFloatTextureFormat, isCompressedTextureFormat, isDepthOrStencilTextureFormat, isDepthTextureFormat, isEncodableTextureFormat, + isSintOrUintFormat, isStencilTextureFormat, kEncodableTextureFormats, kTextureFormatInfo, @@ -79,8 +81,8 @@ export function isSupportedViewFormatCombo( viewDimension: GPUTextureViewDimension ) { return !( - (isCompressedTextureFormat(format) || isDepthTextureFormat(format)) && - viewDimension === '3d' + (isCompressedTextureFormat(format) || isDepthOrStencilTextureFormat(format)) && + (viewDimension === '3d' || viewDimension === '1d') ); } @@ -106,8 +108,10 @@ export function getTextureTypeForTextureViewDimension(viewDimension: GPUTextureV } } -const is32Float = (format: GPUTextureFormat) => - format === 'r32float' || format === 'rg32float' || format === 'rgba32float'; +const isUnencodableDepthFormat = (format: GPUTextureFormat) => + format === 'depth24plus' || + format === 'depth24plus-stencil8' || + format === 'depth32float-stencil8'; /** * Skips a subcase if the filter === 'linear' and the format is type @@ -186,6 +190,12 @@ export function skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable( } } +const builtinNeedsMipLevelWeights = (builtin: TextureBuiltin) => + builtin !== 'textureLoad' && + builtin !== 'textureGather' && + builtin !== 'textureGatherCompare' && + builtin !== 'textureSampleBaseClampToEdge'; + /** * Splits in array into multiple arrays where every Nth value goes to a different array */ @@ -236,7 +246,10 @@ function* linear0to1OverN(n: number) { } } -function graphWeights(height: number, weights: number[]) { +/** + * Generates an ascii graph of weights + */ +export function graphWeights(height: number, weights: number[]) { const graph = makeGraph(weights.length, height); graph.plotValues(linear0to1OverN(weights.length - 1), 1); graph.plotValues(weights, 2); @@ -261,96 +274,20 @@ ${graphWeights(32, weights)} `stage: ${stage}, weight 0 expected 0 but was ${weights[0]}\n${showWeights()}` ); assert( - weights[kMipGradientSteps] === 1, - `stage: ${stage}, top weight expected 1 but was ${weights[kMipGradientSteps]}\n${showWeights()}` + weights[kMipLevelWeightSteps] === 1, + `stage: ${stage}, top weight expected 1 but was ${ + weights[kMipLevelWeightSteps] + }\n${showWeights()}` ); - // Note: for 16 steps, these are the AMD weights - // - // standard - // step mipLevel gpu AMD - // ---- -------- -------- ---------- - // 0: 0 0 0 - // 1: 0.0625 0.0625 0 - // 2: 0.125 0.125 0.03125 - // 3: 0.1875 0.1875 0.109375 - // 4: 0.25 0.25 0.1875 - // 5: 0.3125 0.3125 0.265625 - // 6: 0.375 0.375 0.34375 - // 7: 0.4375 0.4375 0.421875 - // 8: 0.5 0.5 0.5 - // 9: 0.5625 0.5625 0.578125 - // 10: 0.625 0.625 0.65625 - // 11: 0.6875 0.6875 0.734375 - // 12: 0.75 0.75 0.8125 - // 13: 0.8125 0.8125 0.890625 - // 14: 0.875 0.875 0.96875 - // 15: 0.9375 0.9375 1 - // 16: 1 1 1 - // - // notice step 1 is 0 and step 15 is 1. - // so we only check the 1 through 14. - // - // Note: these 2 changes are effectively here to catch Intel Mac - // issues and require implementations to work around them. - // - // Ideally the weights should form a straight line - // - // +----------------+ - // | **| - // | ** | - // | ** | - // | ** | - // | ** | - // | ** | - // | ** | - // |** | - // +----------------+ - // - // AMD Mac goes like this: Not great but we allow it - // - // +----------------+ - // | ***| - // | ** | - // | * | - // | ** | - // | ** | - // | * | - // | ** | - // |*** | - // +----------------+ - // - // Intel Mac goes like this: Unacceptable - // - // +----------------+ - // | *******| - // | * | - // | * | - // | * | - // | * | - // | * | - // | * | - // |******* | - // +----------------+ - // - const dx = 1 / kMipGradientSteps; - for (let i = 0; i < kMipGradientSteps; ++i) { - const dy = weights[i + 1] - weights[i]; - // dy / dx because dy might be 0 - const slope = dy / dx; - assert( - slope >= 0, - `stage: ${stage}, weight[${i}] was not <= weight[${i + 1}]\n${showWeights()}` - ); - assert( - slope <= 2, - `stage: ${stage}, slope from weight[${i}] to weight[${i + 1}] is > 2.\n${showWeights()}` - ); - } - + // Test that we don't have a mostly flat set of weights. + // This is also some small guarantee that we actually read something. + // Note: Ideally every value is unique but 25% is about how many an Intel Mac + // returns in a compute stage. + const kMinPercentUniqueWeights = 25; assert( - new Set(weights).size >= ((weights.length * 0.66) | 0), - `stage: ${stage}, expected more unique weights\n${showWeights()}` + new Set(weights).size >= ((weights.length * kMinPercentUniqueWeights * 0.01) | 0), + `stage: ${stage}, expected at least ~${kMinPercentUniqueWeights}% unique weights\n${showWeights()}` ); } @@ -455,7 +392,7 @@ ${graphWeights(32, weights)} * +--------+--------+--------+--------+ */ -async function queryMipGradientValuesForDevice(t: GPUTest, stage: ShaderStage) { +export async function queryMipLevelMixWeightsForDevice(t: GPUTest, stage: ShaderStage) { const { device } = t; const kNumWeightTypes = 2; const module = device.createShaderModule({ @@ -471,7 +408,7 @@ async function queryMipGradientValuesForDevice(t: GPUTest, stage: ShaderStage) { }; fn getMixLevels(wNdx: u32) -> vec4f { - let mipLevel = f32(wNdx) / ${kMipGradientSteps}; + let mipLevel = f32(wNdx) / ${kMipLevelWeightSteps}; let size = textureDimensions(tex); let g = mix(1.0, 2.0, mipLevel) / f32(size.x); let ddx = vec2f(g, 0); @@ -550,7 +487,7 @@ async function queryMipGradientValuesForDevice(t: GPUTest, stage: ShaderStage) { }); const storageBuffer = t.createBufferTracked({ - size: 4 * (kMipGradientSteps + 1) * kNumWeightTypes, + size: 4 * (kMipLevelWeightSteps + 1) * kNumWeightTypes, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, }); @@ -579,7 +516,7 @@ async function queryMipGradientValuesForDevice(t: GPUTest, stage: ShaderStage) { const pass = encoder.beginComputePass(); pass.setPipeline(pipeline); pass.setBindGroup(0, createBindGroup(pipeline)); - pass.dispatchWorkgroups(kMipGradientSteps + 1); + pass.dispatchWorkgroups(kMipLevelWeightSteps + 1); pass.end(); break; } @@ -600,7 +537,7 @@ async function queryMipGradientValuesForDevice(t: GPUTest, stage: ShaderStage) { }); pass.setPipeline(pipeline); pass.setBindGroup(0, createBindGroup(pipeline)); - pass.draw(3, kMipGradientSteps + 1); + pass.draw(3, kMipLevelWeightSteps + 1); pass.end(); break; } @@ -621,7 +558,7 @@ async function queryMipGradientValuesForDevice(t: GPUTest, stage: ShaderStage) { }); pass.setPipeline(pipeline); pass.setBindGroup(0, createBindGroup(pipeline)); - pass.draw(3, kMipGradientSteps + 1); + pass.draw(3, kMipLevelWeightSteps + 1); pass.end(); break; } @@ -758,20 +695,20 @@ const euclideanModulo = (n: number, m: number) => ((n % m) + m) % m; * for subcase 1's "query the weights" step. Otherwise, all subcases would do the * "get the weights" step separately. */ -const kMipGradientSteps = 64; -const s_deviceToMipGradientValuesPromise = new WeakMap< +const kMipLevelWeightSteps = 64; +const s_deviceToMipLevelWeightsPromise = new WeakMap< GPUDevice, Record> >(); -const s_deviceToMipGradientValues = new WeakMap>(); +const s_deviceToMipLevelWeights = new WeakMap>(); -async function initMipGradientValuesForDevice(t: GPUTest, stage: ShaderStage) { +async function initMipLevelWeightsForDevice(t: GPUTest, stage: ShaderStage) { const { device } = t; // Get the per stage promises (or make them) const stageWeightsP = - s_deviceToMipGradientValuesPromise.get(device) ?? + s_deviceToMipLevelWeightsPromise.get(device) ?? ({} as Record>); - s_deviceToMipGradientValuesPromise.set(device, stageWeightsP); + s_deviceToMipLevelWeightsPromise.set(device, stageWeightsP); let weightsP = stageWeightsP[stage]; if (!weightsP) { @@ -779,12 +716,12 @@ async function initMipGradientValuesForDevice(t: GPUTest, stage: ShaderStage) { // and add a then clause so the first thing that will happen // when the promise resolves is that we'll record the weights for // that stage. - weightsP = queryMipGradientValuesForDevice(t, stage); + weightsP = queryMipLevelMixWeightsForDevice(t, stage); weightsP .then(weights => { const stageWeights = - s_deviceToMipGradientValues.get(device) ?? ({} as Record); - s_deviceToMipGradientValues.set(device, stageWeights); + s_deviceToMipLevelWeights.get(device) ?? ({} as Record); + s_deviceToMipLevelWeights.set(device, stageWeights); stageWeights[stage] = weights; }) .catch(e => { @@ -805,7 +742,7 @@ function getMixWeightByTypeForMipLevel( return euclideanModulo(mipLevel, 1); } // linear interpolate between weights - const weights = s_deviceToMipGradientValues.get(t.device)![stage][weightType]; + const weights = s_deviceToMipLevelWeights.get(t.device)![stage][weightType]; assert( !!weights, 'you must use WGSLTextureSampleTest or call initializeDeviceMipWeights before calling this function' @@ -832,39 +769,240 @@ function getWeightForMipLevel( } /** - * Used for textureDimension, textureNumLevels, textureNumLayers + * Used for textureNumSamples, textureNumLevels, textureNumLayers, textureDimension */ export class WGSLTextureQueryTest extends GPUTest { - executeAndExpectResult(code: string, view: GPUTextureView, expected: number[]) { + executeAndExpectResult( + stage: ShaderStage, + code: string, + texture: GPUTexture | GPUExternalTexture, + viewDescriptor: GPUTextureViewDescriptor | undefined, + expected: number[] + ) { const { device } = this; - const module = device.createShaderModule({ code }); - const pipeline = device.createComputePipeline({ - layout: 'auto', - compute: { - module, - }, - }); + const returnType = `vec4`; + const castWGSL = `${returnType}(getValue()${range(4 - expected.length, () => ', 0').join('')})`; + const stageWGSL = + stage === 'vertex' + ? ` +// --------------------------- vertex stage shaders -------------------------------- +@vertex fn vsVertex( + @builtin(vertex_index) vertex_index : u32, + @builtin(instance_index) instance_index : u32) -> VOut { + let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1)); + return VOut(vec4f(positions[vertex_index], 0, 1), + instance_index, + ${castWGSL}); +} - const resultBuffer = this.createBufferTracked({ - size: 16, - usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, +@fragment fn fsVertex(v: VOut) -> @location(0) vec4u { + return bitcast(v.result); +} +` + : stage === 'fragment' + ? ` +// --------------------------- fragment stage shaders -------------------------------- +@vertex fn vsFragment( + @builtin(vertex_index) vertex_index : u32, + @builtin(instance_index) instance_index : u32) -> VOut { + let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1)); + return VOut(vec4f(positions[vertex_index], 0, 1), instance_index, ${returnType}(0)); +} + +@fragment fn fsFragment(v: VOut) -> @location(0) vec4u { + return bitcast(${castWGSL}); +} +` + : ` +// --------------------------- compute stage shaders -------------------------------- +@group(1) @binding(0) var results: array<${returnType}>; + +@compute @workgroup_size(1) fn csCompute(@builtin(global_invocation_id) id: vec3u) { + results[id.x] = ${castWGSL}; +} +`; + const wgsl = ` + ${code} + +struct VOut { + @builtin(position) pos: vec4f, + @location(0) @interpolate(flat, either) ndx: u32, + @location(1) @interpolate(flat, either) result: ${returnType}, +}; + + ${stageWGSL} + `; + const module = device.createShaderModule({ code: wgsl }); + + const visibility = + stage === 'compute' + ? GPUShaderStage.COMPUTE + : stage === 'fragment' + ? GPUShaderStage.FRAGMENT + : GPUShaderStage.VERTEX; + + const entries: GPUBindGroupLayoutEntry[] = []; + if (texture instanceof GPUExternalTexture) { + entries.push({ + binding: 0, + visibility, + externalTexture: {}, + }); + } else if (code.includes('texture_storage')) { + entries.push({ + binding: 0, + visibility, + storageTexture: { + access: code.includes(', read>') + ? 'read-only' + : code.includes(', write>') + ? 'write-only' + : 'read-write', + viewDimension: viewDescriptor?.dimension ?? '2d', + format: texture.format, + }, + }); + } else { + const sampleType = + viewDescriptor?.aspect === 'stencil-only' + ? 'uint' + : code.includes('texture_depth') + ? 'depth' + : isDepthTextureFormat(texture.format) + ? 'unfilterable-float' + : isStencilTextureFormat(texture.format) + ? 'uint' + : texture.sampleCount > 1 && kTextureFormatInfo[texture.format].color?.type === 'float' + ? 'unfilterable-float' + : kTextureFormatInfo[texture.format].color?.type ?? 'unfilterable-float'; + entries.push({ + binding: 0, + visibility, + texture: { + sampleType, + viewDimension: viewDescriptor?.dimension ?? '2d', + multisampled: texture.sampleCount > 1, + }, + }); + } + + const bindGroupLayouts: GPUBindGroupLayout[] = [device.createBindGroupLayout({ entries })]; + + if (stage === 'compute') { + bindGroupLayouts.push( + device.createBindGroupLayout({ + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'storage', + hasDynamicOffset: false, + minBindingSize: 16, + }, + }, + ], + }) + ); + } + + const layout = device.createPipelineLayout({ + bindGroupLayouts, }); - const bindGroup = device.createBindGroup({ + let pipeline: GPUComputePipeline | GPURenderPipeline; + + switch (stage) { + case 'compute': + pipeline = device.createComputePipeline({ + layout, + compute: { module }, + }); + break; + case 'fragment': + case 'vertex': + pipeline = device.createRenderPipeline({ + layout, + vertex: { module }, + fragment: { + module, + targets: [{ format: 'rgba32uint' }], + }, + }); + break; + } + + const bindGroup0 = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [ - { binding: 0, resource: view }, - { binding: 1, resource: { buffer: resultBuffer } }, + { + binding: 0, + resource: + texture instanceof GPUExternalTexture ? texture : texture.createView(viewDescriptor), + }, ], }); + const renderTarget = this.createTextureTracked({ + format: 'rgba32uint', + size: [expected.length, 1], + usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.RENDER_ATTACHMENT, + }); + + const resultBuffer = this.createBufferTracked({ + size: align(expected.length * 4, 256), + usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC, + }); + + let storageBuffer: GPUBuffer | undefined; const encoder = device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(1); - pass.end(); - device.queue.submit([encoder.finish()]); + + if (stage === 'compute') { + storageBuffer = this.createBufferTracked({ + size: resultBuffer.size, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + + const bindGroup1 = device.createBindGroup({ + layout: pipeline!.getBindGroupLayout(1), + entries: [{ binding: 0, resource: { buffer: storageBuffer } }], + }); + + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline! as GPUComputePipeline); + pass.setBindGroup(0, bindGroup0); + pass.setBindGroup(1, bindGroup1); + pass.dispatchWorkgroups(expected.length); + pass.end(); + encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, storageBuffer.size); + } else { + const pass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: renderTarget.createView(), + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + + pass.setPipeline(pipeline! as GPURenderPipeline); + pass.setBindGroup(0, bindGroup0); + for (let i = 0; i < expected.length; ++i) { + pass.setViewport(i, 0, 1, 1, 0, 1); + pass.draw(3, 1, 0, i); + } + pass.end(); + encoder.copyTextureToBuffer( + { texture: renderTarget }, + { + buffer: resultBuffer, + bytesPerRow: resultBuffer.size, + }, + [renderTarget.width, 1] + ); + } + this.device.queue.submit([encoder.finish()]); const e = new Uint32Array(4); e.set(expected); @@ -920,6 +1058,11 @@ function getMinAndMaxTexelValueForComponent( * or something similar to TexelView. */ export function getTexelViewFormatForTextureFormat(format: GPUTextureFormat) { + if (format.endsWith('sint')) { + return 'rgba32sint'; + } else if (format.endsWith('uint')) { + return 'rgba32uint'; + } return format.endsWith('-srgb') ? 'rgba8unorm-srgb' : 'rgba32float'; } @@ -982,7 +1125,8 @@ export function makeRandomDepthComparisonTexelGenerator( }, comparison: GPUCompareFunction ) { - const rep = kTexelRepresentationInfo[info.format as EncodableTextureFormat]; + const format = isUnencodableDepthFormat(info.format) ? 'depth32float' : info.format; + const rep = kTexelRepresentationInfo[format as EncodableTextureFormat]; const size = reifyExtent3D(info.size); const comparisonIsEqualOrNotEqual = comparison === 'equal' || comparison === 'not-equal'; @@ -994,7 +1138,7 @@ export function makeRandomDepthComparisonTexelGenerator( // The problem with comparing equal is other than 0.0 and 1.0, no other // values are guaranteed to be equal. const fixedValues = [0, 0.6, 1, 1]; - const format = comparisonIsEqualOrNotEqual + const encode = comparisonIsEqualOrNotEqual ? (norm: number) => fixedValues[(norm * (fixedValues.length - 1)) | 0] : (norm: number) => norm; @@ -1012,7 +1156,7 @@ export function makeRandomDepthComparisonTexelGenerator( size.depthOrArrayLayers ); const normalized = clamp(rnd / 0xffffffff, { min: 0, max: 1 }); - texel[component] = format(normalized); + texel[component] = encode(normalized); } return quantize(texel, rep); }; @@ -1243,6 +1387,9 @@ const builtinNeedsDerivatives = (builtin: TextureBuiltin) => const isCubeViewDimension = (viewDescriptor?: GPUTextureViewDescriptor) => viewDescriptor?.dimension === 'cube' || viewDescriptor?.dimension === 'cube-array'; +const isViewDimensionCubeOrCubeArray = (viewDimension: GPUTextureViewDimension) => + viewDimension === 'cube' || viewDimension === 'cube-array'; + const s_u32 = new Uint32Array(1); const s_f32 = new Float32Array(s_u32.buffer); const s_i32 = new Int32Array(s_u32.buffer); @@ -1982,7 +2129,15 @@ function isValidOutOfBoundsValue( for (let sampleIndex = 0; sampleIndex < sampleCount; ++sampleIndex) { const texel = mipTexels.color({ x, y, z, sampleIndex }); const rgba = convertPerTexelComponentToResultFormat(texel, mipTexels.format); - if (texelsApproximatelyEqual(gotRGBA, rgba, mipTexels.format, maxFractionalDiff)) { + if ( + texelsApproximatelyEqual( + gotRGBA, + texture.descriptor.format, + rgba, + mipTexels.format, + maxFractionalDiff + ) + ) { return true; } } @@ -2024,25 +2179,29 @@ const kRGBAComponents = [ const kRComponent = [TexelComponent.R] as const; -function texelsApproximatelyEqual( +/** + * Compares two Texels + */ +export function texelsApproximatelyEqual( gotRGBA: PerTexelComponent, + gotFormat: GPUTextureFormat, expectRGBA: PerTexelComponent, - format: EncodableTextureFormat, + expectedFormat: EncodableTextureFormat, maxFractionalDiff: number ) { - const rep = kTexelRepresentationInfo[format]; - const got = convertResultFormatToTexelViewFormat(gotRGBA, format); - const expect = convertResultFormatToTexelViewFormat(expectRGBA, format); + const rep = kTexelRepresentationInfo[expectedFormat]; + const got = convertResultFormatToTexelViewFormat(gotRGBA, expectedFormat); + const expect = convertResultFormatToTexelViewFormat(expectRGBA, expectedFormat); const gULP = convertPerTexelComponentToResultFormat( rep.bitsToULPFromZero(rep.numberToBits(got)), - format + expectedFormat ); const eULP = convertPerTexelComponentToResultFormat( rep.bitsToULPFromZero(rep.numberToBits(expect)), - format + expectedFormat ); - const rgbaComponentsToCheck = isDepthOrStencilTextureFormat(format) + const rgbaComponentsToCheck = isDepthOrStencilTextureFormat(gotFormat) ? kRComponent : kRGBAComponents; @@ -2116,7 +2275,9 @@ export async function checkCallResults( gpuTexture?: GPUTexture ) { const stage = kShortShaderStageToShaderStage[shortShaderStage]; - await initMipGradientValuesForDevice(t, stage); + if (builtinNeedsMipLevelWeights(calls[0].builtin)) { + await initMipLevelWeightsForDevice(t, stage); + } let haveComparisonCheckInfo = false; let checkInfo = { @@ -2141,6 +2302,93 @@ export async function checkCallResults( const call = calls[callIdx]; const gotRGBA = results.results[callIdx]; const expectRGBA = softwareTextureRead(t, stage, call, texture, sampler); + // Issues with textureSampleBias + // + // textureSampleBias tests start to get unexpected results when bias >= ~12 + // where the mip level selected by the GPU is off by +/- 0.41. + // + // The issue is probably an internal precision issue. In order to test a bias of 12 + // we choose a target mip level between 0 and mipLevelCount - 1. For example 0.4. + // We then compute what mip level we need the derivatives to select such that when + // we add in the bias it will result in a mip level of 0.4. For a bias of 12 + // that's means we need the derivatives to select mip level -11.4. That means + // the derivatives are `pow(2, -11.4) / textureSize` so for a texture that's 16 + // pixels wide that's `0.00002312799936691891`. I'm just guessing some of that + // gets rounded off leading. For example, if we round it ourselves. + // + // | derivative | mip level | + // +------------------------+-----------+ + // | 0.00002312799936691891 | -11.4 | + // | 0.000022 | -11.47 | + // | 0.000023 | -11.408 | + // | 0.000024 | -11.34 | + // +------------------------+-----------+ + // + // Note: As an example of a bad case: set `callSpecificMaxFractionalDiff = maxFractionalDiff` below + // then run `webgpu:shader,execution,expression,call,builtin,textureSampleBias:sampled_2d_coords:format="astc-6x6-unorm";filt="linear";modeU="m";modeV="m";offset=false` + // on an M1 Mac. + // + // ``` + // EXPECTATION FAILED: subcase: samplePoints="spiral" + // result was not as expected: + // size: [18, 18, 1] + // mipCount: 3 + // call: textureSampleBias(texture: T, sampler: S, coords: vec2f(0.1527777777777778, 1.4166666666666667) + derivativeBase * derivativeMult(vec2f(0.00002249990733551491, 0)), bias: f32(15.739721414633095)) // #32 + // : as texel coord @ mip level[0]: (2.750, 25.500) + // : as texel coord @ mip level[1]: (1.375, 12.750) + // : as texel coord @ mip level[2]: (0.611, 5.667) + // implicit derivative based mip level: -15.439721414633095 (without bias) + // clamped bias: 15.739721414633095 + // mip level with bias: 0.3000000000000007 + // got: 0.555311381816864, 0.7921856045722961, 0.8004884123802185, 0.38046398758888245 + // expected: 0.6069580801937625, 0.7999182825318225, 0.8152446179041957, 0.335314491045024 + // max diff: 0.027450980392156862 + // abs diffs: 0.0516466983768985, 0.007732677959526368, 0.014756205523977162, 0.04514949654385847 + // rel diffs: 8.51%, 0.97%, 1.81%, 11.87% + // ulp diffs: 866488, 129733, 247568, 1514966 + // + // sample points: + // expected: | got: + // ... + // a: mip(0) at: [ 2, 10, 0], weight: 0.52740 | a: mip(0) at: [ 2, 10, 0], weight: 0.60931 + // b: mip(0) at: [ 3, 10, 0], weight: 0.17580 | b: mip(0) at: [ 3, 10, 0], weight: 0.20319 + // a: value: R: 0.46642, G: 0.77875, B: 0.77509, A: 0.45788 | a: value: R: 0.46642, G: 0.77875, B: 0.77509, A: 0.45788 + // b: value: R: 0.46642, G: 0.77875, B: 0.77509, A: 0.45788 | b: value: R: 0.46642, G: 0.77875, B: 0.77509, A: 0.45788 + // mip level (0) weight: 0.70320 | mip level (0) weight: 0.81250 + // ``` + // + // Notice above the "expected" level weight (0.7) matches the "mip level with bias (0.3)" which is + // the mip level we expected the GPU to select. Selecting mip level 0.3 will do `mix(level0, level1, 0.3)` + // which is 0.7 of level 0 and 0.3 of level 1. Notice the "got" level weight is 0.81 which is pretty far off. + // + // Just looking at the failures, the largest formula below makes most of the tests pass + // + // MAINTENANCE_TODO: Consider different solutions for this issue + // + // 1. Try to figure out what the exact rounding issue is the take it into account + // + // 2. The code currently samples the texture once via the GPU and once via softwareTextureRead. These values are + // "got:" and "expected:" above. The test only fails if they are too different. We could rather get the bilinear + // sample from every mip level and then check the "got" value is between 2 of the levels (or equal if nearest). + // In other words. + // + // if (bias >= 12) + // colorForEachMipLevel = range(mipLevelCount, mipLevel => softwareTextureReadLevel(..., mipLevel)) + // if nearest + // pass = got === one of colorForEachMipLevel + // else // linear + // pass = false; + // for (i = 0; !pass && i < mipLevelCount - 1; i) + // pass = got is between colorForEachMipLevel[i] and colorForEachMipLevel[i + 1] + // + // This would check "something" but effectively it would no longer be checking "bias" for values > 12. Only that + // textureSampleBias returns some possible answer vs some completely wrong answer. + // + // 3. It's possible this check is just not possible given the precision required. We could just check bias -16 to 12 + // and ignore values > 12. We won't be able to test clamping but maybe that's irrelevant. + // + const callSpecificMaxFractionalDiff = + call.bias! >= 12 ? maxFractionalDiff * (2 + call.bias! - 12) : maxFractionalDiff; // The spec says depth and stencil have implementation defined values for G, B, and A // so if this is `textureGather` and component > 0 then there's nothing to check. @@ -2152,11 +2400,19 @@ export async function checkCallResults( continue; } - if (texelsApproximatelyEqual(gotRGBA, expectRGBA, format, maxFractionalDiff)) { + if ( + texelsApproximatelyEqual( + gotRGBA, + texture.descriptor.format, + expectRGBA, + format, + callSpecificMaxFractionalDiff + ) + ) { continue; } - if (!sampler && okBecauseOutOfBounds(texture, call, gotRGBA, maxFractionalDiff)) { + if (!sampler && okBecauseOutOfBounds(texture, call, gotRGBA, callSpecificMaxFractionalDiff)) { continue; } @@ -2179,7 +2435,7 @@ export async function checkCallResults( assert(!Number.isNaN(ulpDiff)); const maxAbs = Math.max(Math.abs(g), Math.abs(e)); const relDiff = maxAbs > 0 ? absDiff / maxAbs : 0; - if (ulpDiff > 3 && absDiff > maxFractionalDiff) { + if (ulpDiff > 3 && absDiff > callSpecificMaxFractionalDiff) { bad = true; } return { absDiff, relDiff, ulpDiff }; @@ -2245,7 +2501,7 @@ export async function checkCallResults( errs.push(`\ got: ${fix5v(rgbaToArray(gotRGBA))} expected: ${fix5v(rgbaToArray(expectRGBA))} - max diff: ${maxFractionalDiff} + max diff: ${callSpecificMaxFractionalDiff} abs diffs: ${fix5v(diffs.map(({ absDiff }) => absDiff))} rel diffs: ${diffs.map(({ relDiff }) => `${(relDiff * 100).toFixed(2)}%`).join(', ')} ulp diffs: ${diffs.map(({ ulpDiff }) => ulpDiff).join(', ')} @@ -2326,6 +2582,21 @@ export async function checkCallResults( const callForSamplePoints = checkInfo.calls[callIdx]; + // We're going to create textures with black and white texels + // but if it's a compressed texture we use an encodable texture. + // It's not perfect but we already know it failed. We're just hoping + // to get sample points. + const useTexelFormatForGPUTexture = isCompressedTextureFormat(texture.descriptor.format); + + if (useTexelFormatForGPUTexture) { + errs.push(` +### WARNING: sample points are derived from un-compressed textures and may not match the +actual GPU results of sampling a compressed texture. The test itself failed at this point +(see expected: and got: above). We're only trying to determine what the GPU sampled, but +we can not do that easily with compressed textures. ### +`); + } + const expectedSamplePoints = [ 'expected:', ...(await identifySamplePoints( @@ -2360,7 +2631,11 @@ export async function checkCallResults( call, gpuTexels, async (texels: TexelView[]) => { - const gpuTexture = createTextureFromTexelViewsLocal(t, texels, texture.descriptor); + const descriptor = { ...texture.descriptor }; + if (useTexelFormatForGPUTexture) { + descriptor.format = texels[0].format; + } + const gpuTexture = createTextureFromTexelViewsLocal(t, texels, descriptor); const result = (await checkInfo.runner.run(gpuTexture))[callIdx]; gpuTexture.destroy(); return result; @@ -2575,6 +2850,13 @@ function getEffectiveViewDimension( ); } +/** + * Reads a texture to an array of TexelViews, one per mip level. + * format is the format of the TexelView you want. Often this is + * same as the texture.format but if the texture.format is not + * "Encodable" then you need to choose a different format. + * Example: depth24plus -> r32float, bc1-rgba-unorm to rgba32float + */ export async function readTextureToTexelViews( t: GPUTest, texture: GPUTexture, @@ -2587,78 +2869,95 @@ export async function readTextureToTexelViews( new Map(); s_readTextureToRGBA32DeviceToPipeline.set(device, viewDimensionToPipelineMap); + const { componentType, resultType } = getTextureFormatTypeInfo(texture.format); const viewDimension = getEffectiveViewDimension(t, descriptor); - const id = `${viewDimension}:${texture.sampleCount}`; + const id = `${texture.format}:${viewDimension}:${texture.sampleCount}`; let pipeline = viewDimensionToPipelineMap.get(id); if (!pipeline) { let textureWGSL; let loadWGSL; - let dimensionWGSL = 'textureDimensions(tex, uni.mipLevel)'; + let dimensionWGSL = 'textureDimensions(tex, 0)'; switch (viewDimension) { case '2d': if (texture.sampleCount > 1) { - textureWGSL = 'texture_multisampled_2d'; + textureWGSL = `texture_multisampled_2d<${componentType}>`; loadWGSL = 'textureLoad(tex, coord.xy, sampleIndex)'; dimensionWGSL = 'textureDimensions(tex)'; } else { - textureWGSL = 'texture_2d'; - loadWGSL = 'textureLoad(tex, coord.xy, mipLevel)'; + textureWGSL = `texture_2d<${componentType}>`; + loadWGSL = 'textureLoad(tex, coord.xy, 0)'; } break; case 'cube-array': // cube-array doesn't exist in compat so we can just use 2d_array for this case '2d-array': - textureWGSL = 'texture_2d_array'; + textureWGSL = `texture_2d_array<${componentType}>`; loadWGSL = ` textureLoad( tex, coord.xy, coord.z, - mipLevel)`; + 0)`; break; case '3d': - textureWGSL = 'texture_3d'; - loadWGSL = 'textureLoad(tex, coord.xyz, mipLevel)'; + textureWGSL = `texture_3d<${componentType}>`; + loadWGSL = 'textureLoad(tex, coord.xyz, 0)'; break; case 'cube': - textureWGSL = 'texture_cube'; + textureWGSL = `texture_cube<${componentType}>`; loadWGSL = ` - textureLoadCubeAs2DArray(tex, coord.xy, coord.z, mipLevel); + textureLoadCubeAs2DArray(tex, coord.xy, coord.z); `; break; + case '1d': + textureWGSL = `texture_1d<${componentType}>`; + loadWGSL = `textureLoad(tex, coord.x, 0)`; + dimensionWGSL = `vec2u(textureDimensions(tex), 1)`; + break; default: unreachable(`unsupported view: ${viewDimension}`); } + + const textureLoadCubeWGSL = ` + const faceMat = array( + mat3x3f( 0, 0, -2, 0, -2, 0, 1, 1, 1), // pos-x + mat3x3f( 0, 0, 2, 0, -2, 0, -1, 1, -1), // neg-x + mat3x3f( 2, 0, 0, 0, 0, 2, -1, 1, -1), // pos-y + mat3x3f( 2, 0, 0, 0, 0, -2, -1, -1, 1), // neg-y + mat3x3f( 2, 0, 0, 0, -2, 0, -1, 1, 1), // pos-z + mat3x3f(-2, 0, 0, 0, -2, 0, 1, 1, -1)); // neg-z + + // needed for compat mode. + fn textureLoadCubeAs2DArray(tex: texture_cube<${componentType}>, coord: vec2u, layer: u32) -> ${resultType} { + // convert texel coord normalized coord + let size = textureDimensions(tex, 0); + let uv = (vec2f(coord) + 0.5) / vec2f(size.xy); + + // convert uv + layer into cube coord + let cubeCoord = faceMat[layer] * vec3f(uv, 1.0); + + // We have to use textureGather as it's the only texture builtin that works on cubemaps + // with integer texture formats. + let r = textureGather(0, tex, smp, cubeCoord); + let g = textureGather(1, tex, smp, cubeCoord); + let b = textureGather(2, tex, smp, cubeCoord); + let a = textureGather(3, tex, smp, cubeCoord); + + // element 3 is the texel corresponding to cubeCoord + return ${resultType}(r[3], g[3], b[3], a[3]); + } + `; + const module = device.createShaderModule({ code: ` - const faceMat = array( - mat3x3f( 0, 0, -2, 0, -2, 0, 1, 1, 1), // pos-x - mat3x3f( 0, 0, 2, 0, -2, 0, -1, 1, -1), // neg-x - mat3x3f( 2, 0, 0, 0, 0, 2, -1, 1, -1), // pos-y - mat3x3f( 2, 0, 0, 0, 0, -2, -1, -1, 1), // neg-y - mat3x3f( 2, 0, 0, 0, -2, 0, -1, 1, 1), // pos-z - mat3x3f(-2, 0, 0, 0, -2, 0, 1, 1, -1)); // neg-z - - // needed for compat mode. - fn textureLoadCubeAs2DArray(tex: texture_cube, coord: vec2u, layer: u32, mipLevel: u32) -> vec4f { - // convert texel coord normalized coord - let size = textureDimensions(tex, mipLevel); - let uv = (vec2f(coord) + 0.5) / vec2f(size.xy); - - // convert uv + layer into cube coord - let cubeCoord = faceMat[layer] * vec3f(uv, 1.0); - - return textureSampleLevel(tex, smp, cubeCoord, f32(mipLevel)); - } - + ${isViewDimensionCubeOrCubeArray(viewDimension) ? textureLoadCubeWGSL : ''} struct Uniforms { - mipLevel: u32, sampleCount: u32, }; @group(0) @binding(0) var uni: Uniforms; @group(0) @binding(1) var tex: ${textureWGSL}; @group(0) @binding(2) var smp: sampler; - @group(0) @binding(3) var data: array; + @group(0) @binding(3) var data: array<${resultType}>; @compute @workgroup_size(1) fn cs( @builtin(global_invocation_id) global_invocation_id : vec3) { @@ -2669,12 +2968,56 @@ export async function readTextureToTexelViews( global_invocation_id.x; let coord = vec3u(global_invocation_id.x / uni.sampleCount, global_invocation_id.yz); let sampleIndex = global_invocation_id.x % uni.sampleCount; - let mipLevel = uni.mipLevel; data[ndx] = ${loadWGSL}; } `, }); - pipeline = device.createComputePipeline({ layout: 'auto', compute: { module } }); + const info = kTextureFormatInfo[texture.format]; + const sampleType = info.depth + ? 'unfilterable-float' // depth only supports unfilterable-float if not a comparison. + : info.stencil + ? 'uint' + : info.color.type === 'float' + ? 'unfilterable-float' + : info.color.type; + const bindGroupLayout = device.createBindGroupLayout({ + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'uniform', + }, + }, + { + binding: 1, + visibility: GPUShaderStage.COMPUTE, + texture: { + sampleType, + viewDimension, + multisampled: texture.sampleCount > 1, + }, + }, + { + binding: 2, + visibility: GPUShaderStage.COMPUTE, + sampler: { + type: 'non-filtering', + }, + }, + { + binding: 3, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'storage', + }, + }, + ], + }); + const layout = device.createPipelineLayout({ + bindGroupLayouts: [bindGroupLayout], + }); + pipeline = device.createComputePipeline({ layout, compute: { module } }); viewDimensionToPipelineMap.set(id, pipeline); } @@ -2684,7 +3027,7 @@ export async function readTextureToTexelViews( for (let mipLevel = 0; mipLevel < texture.mipLevelCount; ++mipLevel) { const size = virtualMipSize(texture.dimension, texture, mipLevel); - const uniformValues = new Uint32Array([mipLevel, texture.sampleCount, 0, 0]); // min size is 16 bytes + const uniformValues = new Uint32Array([texture.sampleCount, 0, 0, 0]); // min size is 16 bytes const uniformBuffer = t.createBufferTracked({ size: uniformValues.byteLength, usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, @@ -2704,11 +3047,20 @@ export async function readTextureToTexelViews( const sampler = device.createSampler(); + const aspect = getAspectForTexture(texture); const bindGroup = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [ { binding: 0, resource: { buffer: uniformBuffer } }, - { binding: 1, resource: texture.createView({ dimension: viewDimension }) }, + { + binding: 1, + resource: texture.createView({ + dimension: viewDimension, + aspect, + baseMipLevel: mipLevel, + mipLevelCount: 1, + }), + }, { binding: 2, resource: sampler }, { binding: 3, resource: { buffer: storageBuffer } }, ], @@ -2730,7 +3082,9 @@ export async function readTextureToTexelViews( await readBuffer.mapAsync(GPUMapMode.READ); // need a copy of the data since unmapping will nullify the typedarray view. - const data = new Float32Array(readBuffer.getMappedRange()).slice(); + const Ctor = + componentType === 'i32' ? Int32Array : componentType === 'u32' ? Uint32Array : Float32Array; + const data = new Ctor(readBuffer.getMappedRange()).slice(); readBuffer.unmap(); const { sampleCount } = texture; @@ -2756,11 +3110,11 @@ export async function readTextureToTexelViews( function createTextureFromTexelViewsLocal( t: GPUTest, texelViews: TexelView[], - desc: Omit + desc: GPUTextureDescriptor ): GPUTexture { const modifiedDescriptor = { ...desc }; // If it's a depth or stencil texture we need to render to it to fill it with data. - if (isDepthOrStencilTextureFormat(texelViews[0].format)) { + if (isDepthOrStencilTextureFormat(desc.format) || desc.sampleCount! > 1) { modifiedDescriptor.usage = desc.usage | GPUTextureUsage.RENDER_ATTACHMENT; } return createTextureFromTexelViews(t, texelViews, modifiedDescriptor); @@ -2794,6 +3148,25 @@ export async function createTextureWithRandomDataAndGetTexels( getTexelViewFormatForTextureFormat(texture.format) ); return { texture, texels }; + } else if (isUnencodableDepthFormat(descriptor.format)) { + // This is round about. We can't directly write to depth24plus, depth24plus-stencil8, depth32float-stencil8 + // and they are not encodable. So: (1) we make random data using `depth32float`. We create a texture with + // that data (createTextureFromTexelViewsLocal will render the data into the texture rather than copy). + // We then need to read it back out but as rgba32float since that is encodable but, since it round tripped + // through the GPU it's now been quantized. + const d32Descriptor = { + ...descriptor, + format: 'depth32float' as GPUTextureFormat, + }; + const tempTexels = createRandomTexelViewMipmap(d32Descriptor, options); + const texture = createTextureFromTexelViewsLocal(t, tempTexels, descriptor); + const texels = await readTextureToTexelViews( + t, + texture, + descriptor, + getTexelViewFormatForTextureFormat(texture.format) + ); + return { texture, texels }; } else { const texels = createRandomTexelViewMipmap(descriptor, options); const texture = createTextureFromTexelViewsLocal(t, texels, descriptor); @@ -2960,6 +3333,8 @@ async function identifySamplePoints( const format = ( kEncodableTextureFormats.includes(info.format as EncodableTextureFormat) ? info.format + : isDepthTextureFormat(info.format) + ? 'depth16unorm' : 'rgba8unorm' ) as EncodableTextureFormat; const rep = kTexelRepresentationInfo[format]; @@ -3037,19 +3412,29 @@ async function identifySamplePoints( layerEntries.set(xyId, weight); } - // +---+---+---+---+ - // | a | | | | - // +---+---+---+---+ - // | | | | | - // +---+---+---+---+ - // | | | | | - // +---+---+---+---+ - // | | | | b | - // +---+---+---+---+ + // example when blockWidth = 2, blockHeight = 2 + // + // 0 1 2 3 + // +===+===+===+===+ + // 0 # a | # | # + // +---+---+---+---+ + // 1 # | # | # + // +===+===+===+===+ + // 2 # | # | # + // +---+---+---+---+ + // 3 # | # | b # + // +===+===+===+===+ + const lines: string[] = []; const letter = (idx: number) => String.fromCodePoint(idx < 30 ? 97 + idx : idx + 9600 - 30); // 97: 'a' let idCount = 0; + const { blockWidth, blockHeight } = kTextureFormatInfo[texture.descriptor.format]; + const [blockHChar, blockVChar] = Math.max(blockWidth, blockHeight) > 1 ? ['=', '#'] : ['-', '|']; + const blockHCell = '+'.padStart(4, blockHChar); // generates ---+ or ===+ + // range + concatenate results. + const rangeCat = (num: number, fn: (i: number) => T) => range(num, fn).join(''); + for (let mipLevel = 0; mipLevel < mipLevelCount; ++mipLevel) { const level = levels[mipLevel]; if (!level) { @@ -3067,66 +3452,54 @@ async function identifySamplePoints( const unSampled = layerEntries ? '' : 'un-sampled'; if (isCube) { const face = kFaceNames[layer % 6]; - lines.push(`layer: ${layer}, cube-layer: ${(layer / 6) | 0} (${face}) ${unSampled}`); + lines.push( + `layer: ${layer} mip(${mipLevel}), cube-layer: ${(layer / 6) | 0} (${face}) ${unSampled}` + ); } else { - lines.push(`layer: ${layer} ${unSampled}`); + lines.push(`layer: ${layer} mip(${mipLevel}) ${unSampled}`); } if (!layerEntries) { continue; } - { - let line = ' '; - for (let x = 0; x < width; x++) { - line += ` ${x.toString().padEnd(2)}`; - } - lines.push(line); - } - { - let line = ' +'; - for (let x = 0; x < width; x++) { - line += x === width - 1 ? '---+' : '---+'; - } - lines.push(line); - } + lines.push(` ${rangeCat(width, x => ` ${x.toString().padEnd(2)}`)}`); + lines.push(` +${rangeCat(width, () => blockHCell)}`); for (let y = 0; y < height; y++) { { - let line = `${y.toString().padEnd(2)}|`; + let line = `${y.toString().padStart(2)} ${blockVChar}`; for (let x = 0; x < width; x++) { + const colChar = (x + 1) % blockWidth === 0 ? blockVChar : '|'; const texelIdx = x + y * texelsPerRow; const weight = layerEntries.get(texelIdx); if (weight !== undefined) { - line += ` ${letter(idCount + orderedTexelIndices.length)} |`; + line += ` ${letter(idCount + orderedTexelIndices.length)} ${colChar}`; orderedTexelIndices.push(texelIdx); } else { - line += ' |'; + line += ` ${colChar}`; } } lines.push(line); } if (y < height - 1) { - let line = ' +'; - for (let x = 0; x < width; x++) { - line += x === width - 1 ? '---+' : '---+'; - } - lines.push(line); - } - } - { - let line = ' +'; - for (let x = 0; x < width; x++) { - line += x === width - 1 ? '---+' : '---+'; + lines.push( + ` +${rangeCat(width, () => ((y + 1) % blockHeight === 0 ? blockHCell : '---+'))}` + ); } - lines.push(line); } + lines.push(` +${range(width, () => blockHCell).join('')}`); const pad2 = (n: number) => n.toString().padStart(2); - const fix5 = (n: number) => n.toFixed(5); + const pad3 = (n: number) => n.toString().padStart(3); + const fix5 = (n: number) => { + const s = n.toFixed(5); + return s === '0.00000' && n !== 0 ? n.toString() : s; + }; + const formatValue = isSintOrUintFormat(format) ? pad3 : fix5; const formatTexel = (texel: PerTexelComponent | undefined) => texel ? Object.entries(texel) - .map(([k, v]) => `${k}: ${fix5(v)}`) + .map(([k, v]) => `${k}: ${formatValue(v)}`) .join(', ') : '*texel values unavailable*'; @@ -3168,7 +3541,7 @@ async function identifySamplePoints( lines.push(...colorLines); lines.push(...compareLines); if (!isNaN(levelWeight)) { - lines.push(`level weight: ${fix5(levelWeight)}`); + lines.push(`mip level (${mipLevel}) weight: ${fix5(levelWeight)}`); } idCount += orderedTexelIndices.length; } @@ -4263,6 +4636,15 @@ function describeTextureCall(call: TextureCall): st return `${call.builtin}(${args.join(', ')})`; } +const getAspectForTexture = (texture: GPUTexture | GPUExternalTexture): GPUTextureAspect => + texture instanceof GPUExternalTexture + ? 'all' + : isDepthTextureFormat(texture.format) + ? 'depth-only' + : isStencilTextureFormat(texture.format) + ? 'stencil-only' + : 'all'; + const s_deviceToPipelines = new WeakMap< GPUDevice, Map @@ -4400,7 +4782,7 @@ function createTextureCallsRunner( const samplerType = isCompare ? 'sampler_comparison' : 'sampler'; const renderTarget = t.createTextureTracked({ - format: resultFormat, + format: 'rgba32uint', size: [calls.length, 1], usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.RENDER_ATTACHMENT, }); @@ -4440,8 +4822,8 @@ function createTextureCallsRunner( getResult(instance_index, ${derivativeType}(0))); } -@fragment fn fsVertex(v: VOut) -> @location(0) ${returnType} { - return v.result; +@fragment fn fsVertex(v: VOut) -> @location(0) vec4u { + return bitcast(v.result); } ` : stage === 'fragment' @@ -4454,9 +4836,9 @@ function createTextureCallsRunner( return VOut(vec4f(positions[vertex_index], 0, 1), instance_index, ${returnType}(0)); } -@fragment fn fsFragment(v: VOut) -> @location(0) ${returnType} { +@fragment fn fsFragment(v: VOut) -> @location(0) vec4u { ${derivativeBaseWGSL} - return getResult(v.ndx, derivativeBase); + return bitcast(getResult(v.ndx, derivativeBase)); } ` : ` @@ -4577,12 +4959,11 @@ ${stageWGSL} } if (sampler) { + const type = isCompare ? 'comparison' : isFiltering ? 'filtering' : 'non-filtering'; entries.push({ binding: 1, visibility, - sampler: { - type: isCompare ? 'comparison' : isFiltering ? 'filtering' : 'non-filtering', - }, + sampler: { type }, }); } @@ -4626,7 +5007,7 @@ ${stageWGSL} vertex: { module }, fragment: { module, - targets: [{ format: renderTarget.format }], + targets: [{ format: 'rgba32uint' }], }, }); break; @@ -4642,6 +5023,12 @@ ${stageWGSL} usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, }); + const aspect = getAspectForTexture(gpuTexture); + const runViewDescriptor = { + ...viewDescriptor, + aspect, + }; + const bindGroup0 = t.device.createBindGroup({ layout: pipeline!.getBindGroupLayout(0), entries: [ @@ -4650,7 +5037,7 @@ ${stageWGSL} resource: gpuTexture instanceof GPUExternalTexture ? gpuTexture - : gpuTexture.createView(viewDescriptor), + : gpuTexture.createView(runViewDescriptor), }, ...(sampler ? [{ binding: 1, resource: gpuSampler! }] : []), { binding: 2, resource: { buffer: dataBuffer } }, diff --git a/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts b/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts index 7a6aa8901e28..bb7f3b113e26 100644 --- a/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts +++ b/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts @@ -345,7 +345,7 @@ function generateFragmentInputs({ const cw = isTriangleClockwise(windowPoints2D); const frontFacing = frontFace === 'cw' ? cw : !cw; - const fragmentOffsets = getMultisampleFragmentOffsets(sampleCount)!; + const fragmentOffsets = getMultisampleFragmentOffsets(sampleCount); for (let y = 0; y < height; ++y) { for (let x = 0; x < width; ++x) { @@ -1453,76 +1453,6 @@ function popcount(input: number): number { return (((n + (n >> 4)) & 0xf0f0f0f) * 0x1010101) >> 24; } -/** - * Checks subgroup_size builtin value consistency. - * - * The builtin subgroup_size is not assumed to be uniform in fragment shaders. - * Therefore, this function checks the value is a power of two within the device - * limits and that the ballot size is less than the stated size. - * @param data An array of vec4u that contains (per texel): - * * builtin value - * * ballot size - * * comparison to other invocations - * * 0 - * @param format The texture format for data - * @param min The minimum subgroup size from the device - * @param max The maximum subgroup size from the device - * @param width The width of the framebuffer - * @param height The height of the framebuffer - */ -function checkSubgroupSizeConsistency( - data: Uint32Array, - format: GPUTextureFormat, - min: number, - max: number, - width: number, - height: number -): Error | undefined { - const { blockWidth, blockHeight, bytesPerBlock } = kTextureFormatInfo[format]; - const blocksPerRow = width / blockWidth; - // Image copies require bytesPerRow to be a multiple of 256. - const bytesPerRow = align(blocksPerRow * (bytesPerBlock ?? 1), 256); - const uintsPerRow = bytesPerRow / 4; - const uintsPerTexel = (bytesPerBlock ?? 1) / blockWidth / blockHeight / 4; - - for (let row = 0; row < height; row++) { - for (let col = 0; col < width; col++) { - const offset = uintsPerRow * row + col * uintsPerTexel; - const builtinSize = data[offset]; - const ballotSize = data[offset + 1]; - const comparison = data[offset + 2]; - if (builtinSize === 0) { - continue; - } - - if (popcount(builtinSize) !== 1) { - return new Error(`Subgroup size '${builtinSize}' is not a power of two`); - } - - if (builtinSize < min) { - return new Error(`Subgroup size '${builtinSize}' is less than minimum '${min}'`); - } - if (max < builtinSize) { - return new Error(`Subgroup size '${builtinSize}' is greater than maximum '${max}'`); - } - - if (builtinSize < ballotSize) { - return new Error(`Inconsistent subgroup ballot size -- icoord: (${row}, ${col}) -- expected: ${builtinSize} -- got: ${ballotSize}`); - } - - if (comparison !== 1) { - return new Error(`Not all invocations in subgroup have same view of the size -- icoord: (${row}, ${col})`); - } - } - } - - return undefined; -} - /** * Runs a subgroup builtin test for fragment shaders * @@ -1578,23 +1508,6 @@ fn vsMain(@builtin(vertex_index) index : u32) -> @builtin(position) vec4f { const byteLength = bytesPerRow * blocksPerColumn; const uintLength = byteLength / 4; - const buffer = t.makeBufferWithContents( - new Uint32Array([1]), - GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST - ); - - const bg = t.device.createBindGroup({ - layout: pipeline.getBindGroupLayout(0), - entries: [ - { - binding: 0, - resource: { - buffer, - }, - }, - ], - }); - for (let i = 0; i < 2; i++) { const framebuffer = t.createTextureTracked({ size: [width, height], @@ -1617,8 +1530,8 @@ fn vsMain(@builtin(vertex_index) index : u32) -> @builtin(position) vec4f { ], }); pass.setPipeline(pipeline); - pass.setBindGroup(0, bg); - pass.draw(3, 1, i); + // Draw the uperr-left triangle (vertices 0-2) or the lower-right triangle (vertices 3-5) + pass.draw(3, 1, i * 3); pass.end(); t.queue.submit([encoder.finish()]); @@ -1635,6 +1548,101 @@ fn vsMain(@builtin(vertex_index) index : u32) -> @builtin(position) vec4f { } } +const kMaximiumSubgroupSize = 128; +// A non-zero magic number indicating no expectation error, in order to prevent the false no-error +// result from zero-initialization. +const kSubgroupShaderNoError = 17; + +/** + * Checks subgroup_size builtin value consistency. + * + * The builtin subgroup_size is not assumed to be uniform in fragment shaders. + * Therefore, this function checks the value is a power of two within the device + * limits and that the ballot size is less than the stated size. + * @param data An array of vec4u that contains (per texel): + * * subgroup_size builtin value + * * balloted active invocations number + * * balloted subgroup size all active invocations agreed on, otherwise 0 + * * error flag, should be equal to kSubgroupShaderNoError or shader found + * expection failed otherwise. + * @param format The texture format for data + * @param min The minimum subgroup size from the device + * @param max The maximum subgroup size from the device + * @param width The width of the framebuffer + * @param height The height of the framebuffer + */ +function checkSubgroupSizeConsistency( + data: Uint32Array, + format: GPUTextureFormat, + min: number, + max: number, + width: number, + height: number +): Error | undefined { + const { blockWidth, blockHeight, bytesPerBlock } = kTextureFormatInfo[format]; + const blocksPerRow = width / blockWidth; + // Image copies require bytesPerRow to be a multiple of 256. + const bytesPerRow = align(blocksPerRow * (bytesPerBlock ?? 1), 256); + const uintsPerRow = bytesPerRow / 4; + const uintsPerTexel = (bytesPerBlock ?? 1) / blockWidth / blockHeight / 4; + + for (let row = 0; row < height; row++) { + for (let col = 0; col < width; col++) { + const offset = uintsPerRow * row + col * uintsPerTexel; + const subgroupSize = data[offset]; + const countActive = data[offset + 1]; + const ballotedSubgroupSize = data[offset + 2]; + const error = data[offset + 3]; + + if (error === 0) { + // Inactive fragment get error `0` instead of noError. Check all output being zero. + if (subgroupSize !== 0 || countActive !== 0 || ballotedSubgroupSize !== 0) { + return new Error( + `Unexpected zero error with non-zero outputs for (${row}, ${col}): got output [${subgroupSize}, ${countActive}, ${ballotedSubgroupSize}, ${error}]` + ); + } + continue; + } + + if (popcount(subgroupSize) !== 1) { + return new Error(`Subgroup size '${subgroupSize}' is not a power of two`); + } + + if (subgroupSize < min) { + return new Error(`Subgroup size '${subgroupSize}' is less than minimum '${min}'`); + } + if (max < subgroupSize) { + return new Error(`Subgroup size '${subgroupSize}' is greater than maximum '${max}'`); + } + + if (subgroupSize < countActive) { + return new Error(`Unexpected active invocations number larger than subgroup size +- icoord: (${row}, ${col}) +- subgroupSize: ${subgroupSize} +- countActive: ${countActive}`); + } + + if (subgroupSize !== ballotedSubgroupSize) { + return new Error(`Inconsistent subgroup size +- icoord: (${row}, ${col}) +- subgroupSize: ${subgroupSize} +- balloted subgroup size: ${ballotedSubgroupSize}`); + } + + if (error !== kSubgroupShaderNoError) { + return new Error( + `Unexpected error value +- icoord: (${row}, ${col}) +- expected: noError (${kSubgroupShaderNoError}) +- got: ${error}` + ); + } + } + } + + return undefined; +} + g.test('subgroup_size') .desc('Tests subgroup_size values') .params(u => @@ -1656,28 +1664,46 @@ g.test('subgroup_size') const fsShader = ` enable subgroups; +const maxSubgroupSize = ${kMaximiumSubgroupSize}u; +const noError = ${kSubgroupShaderNoError}u; + const width = ${t.params.size[0]}; const height = ${t.params.size[1]}; -@group(0) @binding(0) var for_layout : u32; - @fragment fn fsMain( @builtin(position) pos : vec4f, @builtin(subgroup_size) sg_size : u32, ) -> @location(0) vec4u { - _ = for_layout; + var error: u32 = noError; - let ballot = countOneBits(subgroupBallot(true)); - let ballotSize = ballot.x + ballot.y + ballot.z + ballot.w; + let ballotActive = countOneBits(subgroupBallot(true)); + let countActive = ballotActive.x + ballotActive.y + ballotActive.z + ballotActive.w; + // Validate that balloted active invocations number no larger than subgroup size + if (countActive > sg_size) { + error++; + } - // Do all invocations in the subgroup see the same subgroup size? - let firstSize = subgroupBroadcast(sg_size, 0); - let compareBallot = countOneBits(subgroupBallot(firstSize == sg_size)); - let compareSize = compareBallot.x + compareBallot.y + compareBallot.z + compareBallot.w; - let sameSize = select(0u, 1u, compareSize == ballotSize); + var subgroupSizeBallotedInvocations: u32 = 0u; + var ballotedSubgroupSize: u32 = 0u; + for (var i: u32 = 0; i <= maxSubgroupSize; i++) { + let ballotSubgroupSizeEqualI = countOneBits(subgroupBallot(sg_size == i)); + let countSubgroupSizeEqualI = ballotSubgroupSizeEqualI.x + ballotSubgroupSizeEqualI.y + ballotSubgroupSizeEqualI.z + ballotSubgroupSizeEqualI.w; + subgroupSizeBallotedInvocations += countSubgroupSizeEqualI; + // Validate that all active invocations see the same subgroup size, i.e. ballotedSubgroupSize + ballotedSubgroupSize = select(ballotedSubgroupSize, i, countSubgroupSizeEqualI == countActive); + error = select(error, error + 1, countSubgroupSizeEqualI != countActive && countSubgroupSizeEqualI != 0); + } + // Validate that all active invocations balloted in previous loop + if (subgroupSizeBallotedInvocations != countActive) { + error++; + } + // Validate that ballotedSubgroupSize is identical to subgroup_size + if (ballotedSubgroupSize != sg_size) { + error++; + } - return vec4u(sg_size, ballotSize, sameSize, 0); + return vec4u(sg_size, countActive, ballotedSubgroupSize, error); }`; await runSubgroupTest( @@ -1704,12 +1730,14 @@ fn fsMain( * * Very little uniformity is expected for subgroup_invocation_id. * This function checks that all ids are less than the subgroup size - * and no id is repeated. + * (not the ballot size, since the subgroup id can be allocated to + * inactivate invocations between active ones) and no id is repeated. * @param data An array of vec4u that contains (per texel): * * subgroup_invocation_id - * * ballot size - * * non-zero ID unique to each subgroup - * * 0 + * * subgroup size + * * ballot active invocation number + * * error flag, should be equal to kSubgroupShaderNoError or shader found + * expection failed otherwise. * @param format The texture format of data * @param width The width of the framebuffer * @param height The height of the framebuffer @@ -1726,31 +1754,44 @@ function checkSubgroupInvocationIdConsistency( const uintsPerRow = bytesPerRow / 4; const uintsPerTexel = (bytesPerBlock ?? 1) / blockWidth / blockHeight / 4; - const mappings = new Map(); for (let row = 0; row < height; row++) { for (let col = 0; col < width; col++) { const offset = uintsPerRow * row + col * uintsPerTexel; const id = data[offset]; - const size = data[offset + 1]; - const repId = data[offset + 2]; - - if (repId === 0) { + const sgSize = data[offset + 1]; + const ballotSize = data[offset + 2]; + const error = data[offset + 3]; + + if (error === 0) { + // Inactive fragment get error `0` instead of noError. Check all output being zero. + if (id !== 0 || sgSize !== 0 || ballotSize !== 0) { + return new Error( + `Unexpected zero error with non-zero outputs for (${row}, ${col}): got output [${id}, ${sgSize}, ${ballotSize}, ${error}]` + ); + } continue; } - if (size < id) { + if (sgSize < id) { + return new Error( + `Invocation id '${id}' is greater than subgroup size '${sgSize}' for (${row}, ${col})` + ); + } + + if (sgSize < ballotSize) { return new Error( - `Invocation id '${id}' is greater than subgroup size '${size}' for (${row}, ${col})` + `Ballot size '${ballotSize}' is greater than subgroup size '${sgSize}' for (${row}, ${col})` ); } - let v = mappings.get(repId) ?? 0n; - const mask = 1n << BigInt(id); - if ((mask & v) !== 0n) { - return new Error(`Multiple invocations with id '${id}' in subgroup '${repId}'`); + if (error !== kSubgroupShaderNoError) { + return new Error( + `Unexpected error value +- icoord: (${row}, ${col}) +- expected: noError (${kSubgroupShaderNoError}) +- got: ${error}` + ); } - v |= mask; - mappings.set(repId, v); } } @@ -1775,7 +1816,10 @@ enable subgroups; const width = ${t.params.size[0]}; const height = ${t.params.size[1]}; -@group(0) @binding(0) var counter : atomic; +const maxSubgroupSize = ${kMaximiumSubgroupSize}u; +// A non-zero magic number indicating no expectation error, in order to prevent the +// false no-error result from zero-initialization. +const noError = ${kSubgroupShaderNoError}u; @fragment fn fsMain( @@ -1783,14 +1827,40 @@ fn fsMain( @builtin(subgroup_invocation_id) id : u32, @builtin(subgroup_size) sg_size : u32, ) -> @location(0) vec4u { - let ballot = countOneBits(subgroupBallot(true)); - let ballotSize = ballot.x + ballot.y + ballot.z + ballot.w; - // Generate representative id for this subgroup. - var repId = atomicAdd(&counter, 1); - repId = subgroupBroadcast(repId, 0); + var error: u32 = noError; + + // Validate that reported subgroup size is no larger than maxSubgroupSize + if (sg_size > maxSubgroupSize) { + error++; + } + + // Validate that reported subgroup invocation id is smaller than subgroup size + if (id >= sg_size) { + error++; + } + + // Validate that each subgroup id is assigned to at most one active invocation + // in the subgroup + var countAssignedId: u32 = 0u; + for (var i: u32 = 0; i < maxSubgroupSize; i++) { + let ballotIdEqualsI = countOneBits(subgroupBallot(id == i)); + let countInvocationIdEqualsI = ballotIdEqualsI.x + ballotIdEqualsI.y + ballotIdEqualsI.z + ballotIdEqualsI.w; + // Validate an id assigned at most once + error += select(1u, 0u, countInvocationIdEqualsI <= 1); + // Validate id larger than subgroup size will not get balloted + error += select(1u, 0u, (id < sg_size) || (countInvocationIdEqualsI == 0)); + // Sum up the assigned invocation number of each id + countAssignedId += countInvocationIdEqualsI; + } + // Validate that all active invocation get counted during the above loop + let ballotActive = countOneBits(subgroupBallot(true)); + let activeInvocations = ballotActive.x + ballotActive.y + ballotActive.z + ballotActive.w; + if (activeInvocations != countAssignedId) { + error++; + } - return vec4u(id, ballotSize, repId, 0); + return vec4u(id, sg_size, activeInvocations, error); }`; await runSubgroupTest( diff --git a/src/webgpu/shader/validation/expression/call/builtin/const_override_validation.ts b/src/webgpu/shader/validation/expression/call/builtin/const_override_validation.ts index 6ea2caf3c694..83cc1c76f8a0 100644 --- a/src/webgpu/shader/validation/expression/call/builtin/const_override_validation.ts +++ b/src/webgpu/shader/validation/expression/call/builtin/const_override_validation.ts @@ -137,7 +137,7 @@ export type ExecutionStage = 'constant' | 'override' | 'runtime'; * @returns true if evaluation stage `stage` supports expressions of type @p. */ export function stageSupportsType(stage: ConstantOrOverrideStage, type: Type) { - if (stage === 'override' && isAbstractType(elementTypeOf(type)!)) { + if (stage === 'override' && isAbstractType(elementTypeOf(type))) { // Abstract numerics are concretized before being used in an override expression. return false; } @@ -162,7 +162,7 @@ export function validateConstOrOverrideBuiltinEval( stage: ConstantOrOverrideStage, returnType?: Type ) { - const elTys = args.map(arg => elementTypeOf(arg.type)!); + const elTys = args.map(arg => elementTypeOf(arg.type)); const enables = elTys.some(ty => ty === Type.f16) ? 'enable f16;' : ''; const optionalVarType = returnType ? `: ${returnType.toString()}` : ''; @@ -225,7 +225,7 @@ export function validateConstOrOverrideBinaryOpEval( right: Value ) { const allArgs = [left, right]; - const elTys = allArgs.map(arg => elementTypeOf(arg.type)!); + const elTys = allArgs.map(arg => elementTypeOf(arg.type)); const enables = elTys.some(ty => ty === Type.f16) ? 'enable f16;' : ''; const codeLines = [enables]; diff --git a/src/webgpu/shader/validation/shader_io/align.spec.ts b/src/webgpu/shader/validation/shader_io/align.spec.ts index e65dfd4e3a24..b61a1c0f574c 100644 --- a/src/webgpu/shader/validation/shader_io/align.spec.ts +++ b/src/webgpu/shader/validation/shader_io/align.spec.ts @@ -13,7 +13,10 @@ const kTests = { }, one: { src: '@align(1)', - pass: true, + pass: false, + // EXCEPTION: Error: Unexpected validation error occurred: + // Error while parsing WGSL: :6:10 error: alignment must be a + // multiple of '4' bytes for the 'uniform' address space @align(1) a: i32, }, four_a: { src: '@align(4)', @@ -45,6 +48,14 @@ const kTests = { }, const_expr: { src: '@align(i_val + 4 - 6)', + pass: false, + // EXCEPTION: Error: Unexpected validation error occurred: + // Error while parsing WGSL: :6:10 error: alignment must be a + // multiple of '4' bytes for the 'uniform' address space + // @align(i_val + 4 - 6) a: i32 + }, + const_expr_2: { + src: '@align(i_val + 8 - 4)', pass: true, }, large: { @@ -185,7 +196,7 @@ g.test('required_alignment') { name: 'mat3x4', storage: 8, uniform: 8 }, { name: 'mat4x4', storage: 8, uniform: 8 }, { name: 'array, 2>', storage: 8, uniform: 16 }, - { name: 'array, 2>', storage: 8, uniform: 16 }, + { name: 'array, 2>', storage: 16, uniform: 16 }, { name: 'S', storage: 8, uniform: 16 }, ]) .beginSubcases() @@ -218,15 +229,12 @@ g.test('required_alignment') `; } - let align = t.params.align; - if (t.params.align === 'alignment') { - // Alignment value listed in the spec - if (t.params.address_space === 'storage') { - align = `${t.params.type.storage}`; - } else { - align = `${t.params.type.uniform}`; - } - } + // Alignment value listed in the spec + const min_align = + t.params.address_space === 'storage' + ? `${t.params.type.storage}` + : `${t.params.type.uniform}`; + const align = t.params.align === 'alignment' ? min_align : t.params.align; let address_space = 'uniform'; if (t.params.address_space === 'storage') { @@ -252,7 +260,8 @@ g.test('required_alignment') // requires that inner vec2 to have an align 16 which can only be done by specifying `vec4` // instead. const fails = - t.params.address_space === 'uniform' && t.params.type.name.startsWith('array, texelCompareOptions: TexelCompareOptions diff --git a/src/webgpu/util/texture.ts b/src/webgpu/util/texture.ts index 20e99fdfad4d..0bde14ff0549 100644 --- a/src/webgpu/util/texture.ts +++ b/src/webgpu/util/texture.ts @@ -1,5 +1,10 @@ import { assert } from '../../common/util/util.js'; -import { isDepthOrStencilTextureFormat, kTextureFormatInfo } from '../format_info.js'; +import { + isDepthOrStencilTextureFormat, + isDepthTextureFormat, + isStencilTextureFormat, + kTextureFormatInfo, +} from '../format_info.js'; import { GPUTest } from '../gpu_test.js'; import { getTextureCopyLayout } from './texture/layout.js'; @@ -16,8 +21,6 @@ const kLoadValueFromStorageInfo: Partial<{ storageType: string; texelType: string; unpackWGSL: string; - useFragDepth?: boolean; - discardWithStencil?: boolean; }; }> = { r8unorm: { @@ -223,7 +226,6 @@ const kLoadValueFromStorageInfo: Partial<{ let v = unpack2x16unorm(src[byteOffset / 4])[byteOffset % 4 / 2]; return vec4f(v, 0.123, 0.123, 0.123) `, - useFragDepth: true, }, depth32float: { storageType: 'f32', @@ -232,7 +234,6 @@ const kLoadValueFromStorageInfo: Partial<{ let v = src[byteOffset / 4]; return vec4f(v, 0.123, 0.123, 0.123) `, - useFragDepth: true, }, stencil8: { storageType: 'u32', @@ -240,14 +241,25 @@ const kLoadValueFromStorageInfo: Partial<{ unpackWGSL: ` return vec4u(unpack4xU8(src[byteOffset / 4])[byteOffset % 4], 123, 123, 123) `, - discardWithStencil: true, }, }; -function getCopyBufferToTextureViaRenderCode(format: GPUTextureFormat) { - const info = kLoadValueFromStorageInfo[format]; +function getDepthStencilOptionsForFormat(format: GPUTextureFormat) { + // Note: For now we prefer depth over stencil. To fix this would require passing GPUTextureAspect all the way down. + return { + useFragDepth: isDepthTextureFormat(format), + discardWithStencil: isStencilTextureFormat(format) && !isDepthTextureFormat(format), + }; +} + +function getCopyBufferToTextureViaRenderCode( + srcFormat: GPUTextureFormat, + dstFormat: GPUTextureFormat +) { + const info = kLoadValueFromStorageInfo[srcFormat]; assert(!!info); - const { storageType, texelType, unpackWGSL, useFragDepth, discardWithStencil } = info; + const { storageType, texelType, unpackWGSL } = info; + const { useFragDepth, discardWithStencil } = getDepthStencilOptionsForFormat(dstFormat); const [depthDecl, depthCode] = useFragDepth ? ['@builtin(frag_depth) d: f32,', 'fs.d = fs.v[0];'] @@ -318,25 +330,24 @@ const s_copyBufferToTextureViaRenderPipelines = new WeakMap< function copyBufferToTextureViaRender( t: GPUTest, encoder: GPUCommandEncoder, - source: GPUImageCopyBuffer, - dest: GPUImageCopyTexture, + source: GPUTexelCopyBufferInfo, + sourceFormat: GPUTextureFormat, + dest: GPUTexelCopyTextureInfo, size: GPUExtent3D ) { - const { format, sampleCount } = dest.texture; + const { format: textureFormat, sampleCount } = dest.texture; const origin = reifyOrigin3D(dest.origin ?? [0]); const copySize = reifyExtent3D(size); - - const msInfo = kLoadValueFromStorageInfo[format]; - assert(!!msInfo); - const { useFragDepth, discardWithStencil } = msInfo; + const { useFragDepth, discardWithStencil } = getDepthStencilOptionsForFormat(dest.texture.format); const { device } = t; const numBlits = discardWithStencil ? 8 : 1; for (let blitCount = 0; blitCount < numBlits; ++blitCount) { - const code = getCopyBufferToTextureViaRenderCode(format); + const code = getCopyBufferToTextureViaRenderCode(sourceFormat, dest.texture.format); const stencilWriteMask = 1 << blitCount; const id = JSON.stringify({ - format, + textureFormat, + sourceFormat, useFragDepth, stencilWriteMask, discardWithStencil, @@ -350,7 +361,7 @@ function copyBufferToTextureViaRender( if (!pipeline) { const module = device.createShaderModule({ code }); pipeline = device.createRenderPipeline({ - label: `blitCopyFor-${format}`, + label: `blitCopyFor-${textureFormat}`, layout: 'auto', vertex: { module }, ...(discardWithStencil @@ -362,7 +373,7 @@ function copyBufferToTextureViaRender( depthStencil: { depthWriteEnabled: false, depthCompare: 'always', - format, + format: textureFormat, stencilWriteMask, stencilFront: { passOp: 'replace', @@ -378,13 +389,13 @@ function copyBufferToTextureViaRender( depthStencil: { depthWriteEnabled: true, depthCompare: 'always', - format, + format: textureFormat, }, } : { fragment: { module, - targets: [{ format }], + targets: [{ format: textureFormat }], }, }), primitive: { @@ -395,7 +406,7 @@ function copyBufferToTextureViaRender( pipelines.set(id, pipeline); } - const info = kTextureFormatInfo[format]; + const info = kTextureFormatInfo[sourceFormat]; const uniforms = new Uint32Array([ copySize.height, // numTexelRows: u32, source.bytesPerRow!, // bytesPerRow: u32, @@ -445,6 +456,7 @@ function copyBufferToTextureViaRender( depthClearValue: 0, depthLoadOp: 'clear', depthStoreOp: 'store', + stencilReadOnly: true, }, } : { @@ -490,17 +502,19 @@ function copyBufferToTextureViaRender( export function createTextureFromTexelViews( t: GPUTest, texelViews: TexelView[], - desc: Omit + desc: Omit & { format?: GPUTextureFormat } ): GPUTexture { // All texel views must be the same format for mipmaps. assert(texelViews.length > 0 && texelViews.every(e => e.format === texelViews[0].format)); - const format = texelViews[0].format; + const viewsFormat = texelViews[0].format; + const textureFormat = desc.format ?? viewsFormat; + const isTextureFormatDifferentThanTexelViewFormat = textureFormat !== viewsFormat; const { width, height, depthOrArrayLayers } = reifyExtent3D(desc.size); // Create the texture and then initialize each mipmap level separately. const texture = t.createTextureTracked({ ...desc, - format, + format: textureFormat, usage: desc.usage | GPUTextureUsage.COPY_DST, mipLevelCount: texelViews.length, }); @@ -513,9 +527,14 @@ export function createTextureFromTexelViews( bytesPerRow, rowsPerImage, mipSize: [mipWidth, mipHeight, mipDepthOrArray], - } = getTextureCopyLayout(format, desc.dimension ?? '2d', [width, height, depthOrArrayLayers], { - mipLevel, - }); + } = getTextureCopyLayout( + viewsFormat, + desc.dimension ?? '2d', + [width, height, depthOrArrayLayers], + { + mipLevel, + } + ); // Create a staging buffer to upload the texture mip level contents. const stagingBuffer = t.createBufferTracked({ @@ -535,11 +554,16 @@ export function createTextureFromTexelViews( }); stagingBuffer.unmap(); - if (texture.sampleCount > 1 || isDepthOrStencilTextureFormat(format)) { + if ( + isTextureFormatDifferentThanTexelViewFormat || + texture.sampleCount > 1 || + isDepthOrStencilTextureFormat(textureFormat) + ) { copyBufferToTextureViaRender( t, commandEncoder, { buffer: stagingBuffer, bytesPerRow, rowsPerImage }, + viewsFormat, { texture, mipLevel }, [mipWidth, mipHeight, mipDepthOrArray] ); diff --git a/src/webgpu/util/texture/layout.ts b/src/webgpu/util/texture/layout.ts index 24a0cdf0040e..8c6c101ae014 100644 --- a/src/webgpu/util/texture/layout.ts +++ b/src/webgpu/util/texture/layout.ts @@ -291,7 +291,7 @@ function validateRowsPerImage({ } interface DataBytesForCopyArgs { - layout: GPUImageDataLayout; + layout: GPUTexelCopyBufferLayout; format: SizedTextureFormat; copySize: Readonly | readonly number[]; method: ImageCopyType; diff --git a/src/webgpu/util/texture/texel_data.ts b/src/webgpu/util/texture/texel_data.ts index 4c88d9c2182a..734cfc00eef7 100644 --- a/src/webgpu/util/texture/texel_data.ts +++ b/src/webgpu/util/texture/texel_data.ts @@ -1001,7 +1001,7 @@ export function getSingleDataType(format: UncompressedTextureFormat): ComponentD assert(cur !== undefined); assert(acc === undefined || acc === cur.dataType); return cur.dataType; - }, infos[0]!.dataType); + }, infos[0].dataType); } /** diff --git a/src/webgpu/util/texture/texture_ok.ts b/src/webgpu/util/texture/texture_ok.ts index abfec1f1d7c2..3eb94c2deea0 100644 --- a/src/webgpu/util/texture/texture_ok.ts +++ b/src/webgpu/util/texture/texture_ok.ts @@ -167,7 +167,7 @@ function comparePerComponent( /** Create a new mappable GPUBuffer, and copy a subrectangle of GPUTexture data into it. */ function createTextureCopyForMapRead( t: GPUTest, - source: GPUImageCopyTexture, + source: GPUTexelCopyTextureInfo, copySize: GPUExtent3D, { format }: { format: EncodableTextureFormat } ): { buffer: GPUBuffer; bytesPerRow: number; rowsPerImage: number } { @@ -298,7 +298,7 @@ ${generatePrettyTable(opts, [ */ export async function textureContentIsOKByT2B( t: GPUTest, - source: GPUImageCopyTexture, + source: GPUTexelCopyTextureInfo, copySize_: GPUExtent3D, { expTexelView }: { expTexelView: TexelView }, texelCompareOptions: TexelCompareOptions, diff --git a/src/webgpu/web_platform/canvas/configure.spec.ts b/src/webgpu/web_platform/canvas/configure.spec.ts index 06e590751b33..3e9905dd5bd8 100644 --- a/src/webgpu/web_platform/canvas/configure.spec.ts +++ b/src/webgpu/web_platform/canvas/configure.spec.ts @@ -49,7 +49,7 @@ g.test('defaults') t.expect(configuration.usage === GPUTextureUsage.RENDER_ATTACHMENT); t.expect(configuration.viewFormats.length === 0); t.expect(configuration.colorSpace === 'srgb'); - t.expect(configuration.toneMapping.mode === 'standard'); + t.expect(configuration.toneMapping?.mode === 'standard'); t.expect(configuration.alphaMode === 'opaque'); const currentTexture = ctx.getCurrentTexture(); @@ -109,7 +109,7 @@ g.test('device') t.expect(configuration.usage === GPUTextureUsage.RENDER_ATTACHMENT); t.expect(configuration.viewFormats.length === 0); t.expect(configuration.colorSpace === 'srgb'); - t.expect(configuration.toneMapping.mode === 'standard'); + t.expect(configuration.toneMapping?.mode === 'standard'); t.expect(configuration.alphaMode === 'opaque'); // getCurrentTexture will succeed with a valid device. @@ -140,7 +140,7 @@ g.test('device') t.expect(newConfiguration.usage === GPUTextureUsage.RENDER_ATTACHMENT); t.expect(newConfiguration.viewFormats.length === 0); t.expect(newConfiguration.colorSpace === 'srgb'); - t.expect(newConfiguration.toneMapping.mode === 'standard'); + t.expect(newConfiguration.toneMapping?.mode === 'standard'); t.expect(newConfiguration.alphaMode === 'premultiplied'); }); diff --git a/src/webgpu/web_platform/copyToTexture/ImageBitmap.spec.ts b/src/webgpu/web_platform/copyToTexture/ImageBitmap.spec.ts index b7638cd08b02..5036c8ef20bc 100644 --- a/src/webgpu/web_platform/copyToTexture/ImageBitmap.spec.ts +++ b/src/webgpu/web_platform/copyToTexture/ImageBitmap.spec.ts @@ -31,17 +31,17 @@ g.test('from_ImageData') Then call copyExternalImageToTexture() to do a full copy to the 0 mipLevel of dst texture, and read the contents out to compare with the ImageBitmap contents. - Do premultiply alpha during copy if 'premultipliedAlpha' in 'GPUImageCopyTextureTagged' + Do premultiply alpha during copy if 'premultipliedAlpha' in 'GPUCopyExternalImageDestInfo' is set to 'true' and do unpremultiply alpha if it is set to 'false'. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped. The tests covers: - Valid dstFormat of copyExternalImageToTexture() - Valid source image alphaMode - Valid dest alphaMode - - Valid 'flipY' config in 'GPUImageCopyExternalImage' (named 'srcFlipYInCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo' (named 'srcFlipYInCopy' in cases) And the expected results are all passed. ` @@ -151,10 +151,10 @@ g.test('from_canvas') Then call copyExternalImageToTexture() to do a full copy to the 0 mipLevel of dst texture, and read the contents out to compare with the ImageBitmap contents. - Do premultiply alpha during copy if 'premultipliedAlpha' in 'GPUImageCopyTextureTagged' + Do premultiply alpha during copy if 'premultipliedAlpha' in 'GPUCopyExternalImageDestInfo' is set to 'true' and do unpremultiply alpha if it is set to 'false'. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped. The tests covers: @@ -162,7 +162,7 @@ g.test('from_canvas') - Valid dstFormat of copyExternalImageToTexture() - Valid source image alphaMode - Valid dest alphaMode - - Valid 'flipY' config in 'GPUImageCopyExternalImage' (named 'srcFlipYInCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo' (named 'srcFlipYInCopy' in cases) And the expected results are all passed. ` @@ -299,10 +299,10 @@ g.test('copy_subrect_from_ImageData') rect info list, to the 0 mipLevel of dst texture, and read the contents out to compare with the ImageBitmap contents. - Do premultiply alpha during copy if 'premultipliedAlpha' in 'GPUImageCopyTextureTagged' + Do premultiply alpha during copy if 'premultipliedAlpha' in 'GPUCopyExternalImageDestInfo' is set to 'true' and do unpremultiply alpha if it is set to 'false'. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped, and origin is top-left consistantly. The tests covers: @@ -310,7 +310,7 @@ g.test('copy_subrect_from_ImageData') - Valid dstFormat of copyExternalImageToTexture() - Valid source image alphaMode - Valid dest alphaMode - - Valid 'flipY' config in 'GPUImageCopyExternalImage' (named 'srcFlipYInCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo' (named 'srcFlipYInCopy' in cases) - Valid subrect copies. And the expected results are all passed. @@ -416,10 +416,10 @@ g.test('copy_subrect_from_2D_Canvas') rect info list, to the 0 mipLevel of dst texture, and read the contents out to compare with the ImageBitmap contents. - Do premultiply alpha during copy if 'premultipliedAlpha' in 'GPUImageCopyTextureTagged' + Do premultiply alpha during copy if 'premultipliedAlpha' in 'GPUCopyExternalImageDestInfo' is set to 'true' and do unpremultiply alpha if it is set to 'false'. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped, and origin is top-left consistantly. The tests covers: @@ -427,7 +427,7 @@ g.test('copy_subrect_from_2D_Canvas') - Valid dstFormat of copyExternalImageToTexture() - Valid source image alphaMode - Valid dest alphaMode - - Valid 'flipY' config in 'GPUImageCopyExternalImage' (named 'srcFlipYInCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo' (named 'srcFlipYInCopy' in cases) - Valid subrect copies. And the expected results are all passed. diff --git a/src/webgpu/web_platform/copyToTexture/ImageData.spec.ts b/src/webgpu/web_platform/copyToTexture/ImageData.spec.ts index 03e8f9a893b3..a9a0a441653a 100644 --- a/src/webgpu/web_platform/copyToTexture/ImageData.spec.ts +++ b/src/webgpu/web_platform/copyToTexture/ImageData.spec.ts @@ -24,15 +24,15 @@ g.test('from_ImageData') of dst texture, and read the contents out to compare with the ImageData contents. Expect alpha to get premultiplied in the copy if, and only if, 'premultipliedAlpha' - in 'GPUImageCopyTextureTagged' is set to 'true'. + in 'GPUCopyExternalImageDestInfo' is set to 'true'. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped. The tests covers: - Valid dstColorFormat of copyExternalImageToTexture() - Valid dest alphaMode - - Valid 'flipY' config in 'GPUImageCopyExternalImage' (named 'srcDoFlipYDuringCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo' (named 'srcDoFlipYDuringCopy' in cases) And the expected results are all passed. ` @@ -130,16 +130,16 @@ g.test('copy_subrect_from_ImageData') with the ImageBitmap contents. Expect alpha to get premultiplied in the copy if, and only if, 'premultipliedAlpha' - in 'GPUImageCopyTextureTagged' is set to 'true'. + in 'GPUCopyExternalImageDestInfo' is set to 'true'. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped, and origin is top-left consistantly. The tests covers: - Source WebGPU Canvas lives in the same GPUDevice or different GPUDevice as test - Valid dstColorFormat of copyExternalImageToTexture() - Valid dest alphaMode - - Valid 'flipY' config in 'GPUImageCopyExternalImage' (named 'srcDoFlipYDuringCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo' (named 'srcDoFlipYDuringCopy' in cases) - Valid subrect copies. And the expected results are all passed. diff --git a/src/webgpu/web_platform/copyToTexture/canvas.spec.ts b/src/webgpu/web_platform/copyToTexture/canvas.spec.ts index 210a63abaff9..84334df603e3 100644 --- a/src/webgpu/web_platform/copyToTexture/canvas.spec.ts +++ b/src/webgpu/web_platform/copyToTexture/canvas.spec.ts @@ -466,10 +466,10 @@ g.test('copy_contents_from_2d_context_canvas') Then call copyExternalImageToTexture() to do a full copy to the 0 mipLevel of dst texture, and read the contents out to compare with the canvas contents. - Provide premultiplied input if 'premultipliedAlpha' in 'GPUImageCopyTextureTagged' + Provide premultiplied input if 'premultipliedAlpha' in 'GPUCopyExternalImageDestInfo' is set to 'true' and unpremultiplied input if it is set to 'false'. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped. The tests covers: @@ -477,7 +477,7 @@ g.test('copy_contents_from_2d_context_canvas') - Valid 2d context type - Valid dstColorFormat of copyExternalImageToTexture() - Valid dest alphaMode - - Valid 'flipY' config in 'GPUImageCopyExternalImage' (named 'srcDoFlipYDuringCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo' (named 'srcDoFlipYDuringCopy' in cases) - TODO(#913): color space tests need to be added And the expected results are all passed. @@ -527,10 +527,10 @@ g.test('copy_contents_from_gl_context_canvas') Then call copyExternalImageToTexture() to do a full copy to the 0 mipLevel of dst texture, and read the contents out to compare with the canvas contents. - Provide premultiplied input if 'premultipliedAlpha' in 'GPUImageCopyTextureTagged' + Provide premultiplied input if 'premultipliedAlpha' in 'GPUCopyExternalImageDestInfo' is set to 'true' and unpremultiplied input if it is set to 'false'. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped. The tests covers: @@ -539,7 +539,7 @@ g.test('copy_contents_from_gl_context_canvas') - Valid dstColorFormat of copyExternalImageToTexture() - Valid source image alphaMode - Valid dest alphaMode - - Valid 'flipY' config in 'GPUImageCopyExternalImage'(named 'srcDoFlipYDuringCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo'(named 'srcDoFlipYDuringCopy' in cases) - TODO: color space tests need to be added And the expected results are all passed. @@ -595,10 +595,10 @@ g.test('copy_contents_from_gpu_context_canvas') Then call copyExternalImageToTexture() to do a full copy to the 0 mipLevel of dst texture, and read the contents out to compare with the canvas contents. - Provide premultiplied input if 'premultipliedAlpha' in 'GPUImageCopyTextureTagged' + Provide premultiplied input if 'premultipliedAlpha' in 'GPUCopyExternalImageDestInfo' is set to 'true' and unpremultiplied input if it is set to 'false'. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped. The tests covers: @@ -607,7 +607,7 @@ g.test('copy_contents_from_gpu_context_canvas') - Valid dstColorFormat of copyExternalImageToTexture() - TODO: test more source image alphaMode - Valid dest alphaMode - - Valid 'flipY' config in 'GPUImageCopyExternalImage'(named 'srcDoFlipYDuringCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo'(named 'srcDoFlipYDuringCopy' in cases) - TODO: color space tests need to be added And the expected results are all passed. @@ -665,10 +665,10 @@ g.test('copy_contents_from_bitmaprenderer_context_canvas') Then call copyExternalImageToTexture() to do a full copy to the 0 mipLevel of dst texture, and read the contents out to compare with the canvas contents. - Provide premultiplied input if 'premultipliedAlpha' in 'GPUImageCopyTextureTagged' + Provide premultiplied input if 'premultipliedAlpha' in 'GPUCopyExternalImageDestInfo' is set to 'true' and unpremultiplied input if it is set to 'false'. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped. The tests covers: @@ -676,7 +676,7 @@ g.test('copy_contents_from_bitmaprenderer_context_canvas') - Valid ImageBitmapRendering context type - Valid dstColorFormat of copyExternalImageToTexture() - Valid dest alphaMode - - Valid 'flipY' config in 'GPUImageCopyExternalImage' (named 'srcDoFlipYDuringCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo' (named 'srcDoFlipYDuringCopy' in cases) - TODO(#913): color space tests need to be added And the expected results are all passed. @@ -736,10 +736,10 @@ g.test('color_space_conversion') Then call copyExternalImageToTexture() to do a full copy to the 0 mipLevel of dst texture, and read the contents out to compare with the canvas contents. - Provide premultiplied input if 'premultipliedAlpha' in 'GPUImageCopyTextureTagged' + Provide premultiplied input if 'premultipliedAlpha' in 'GPUCopyExternalImageDestInfo' is set to 'true' and unpremultiplied input if it is set to 'false'. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped. If color space from source input and user defined dstTexture color space are different, the @@ -748,7 +748,7 @@ g.test('color_space_conversion') The tests covers: - Valid dstColorFormat of copyExternalImageToTexture() - Valid dest alphaMode - - Valid 'flipY' config in 'GPUImageCopyExternalImage' (named 'srcDoFlipYDuringCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo' (named 'srcDoFlipYDuringCopy' in cases) - Valid 'colorSpace' config in 'dstColorSpace' And the expected results are all passed. diff --git a/src/webgpu/web_platform/copyToTexture/image.spec.ts b/src/webgpu/web_platform/copyToTexture/image.spec.ts index 2bb6eab8f176..7bb577ae5d07 100644 --- a/src/webgpu/web_platform/copyToTexture/image.spec.ts +++ b/src/webgpu/web_platform/copyToTexture/image.spec.ts @@ -35,10 +35,10 @@ g.test('from_image') Then call copyExternalImageToTexture() to do a full copy to the 0 mipLevel of dst texture, and read the contents out to compare with the HTMLImageElement contents. - Do premultiply alpha during copy if 'premultipliedAlpha' in 'GPUImageCopyTextureTagged' + Do premultiply alpha during copy if 'premultipliedAlpha' in 'GPUCopyExternalImageDestInfo' is set to 'true' and do unpremultiply alpha if it is set to 'false'. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped. The tests covers: @@ -46,7 +46,7 @@ g.test('from_image') - Valid dstColorFormat of copyExternalImageToTexture() - Valid source image alphaMode - Valid dest alphaMode - - Valid 'flipY' config in 'GPUImageCopyExternalImage' (named 'srcDoFlipYDuringCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo' (named 'srcDoFlipYDuringCopy' in cases) And the expected results are all passed. ` @@ -229,10 +229,10 @@ g.test('copy_subrect_from_2D_Canvas') rect info list, to the 0 mipLevel of dst texture, and read the contents out to compare with the HTMLImageElement contents. - Do premultiply alpha during copy if 'premultipliedAlpha' in 'GPUImageCopyTextureTagged' + Do premultiply alpha during copy if 'premultipliedAlpha' in 'GPUCopyExternalImageDestInfo' is set to 'true' and do unpremultiply alpha if it is set to 'false'. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped, and origin is top-left consistantly. The tests covers: @@ -240,7 +240,7 @@ g.test('copy_subrect_from_2D_Canvas') - Valid dstColorFormat of copyExternalImageToTexture() - Valid source image alphaMode - Valid dest alphaMode - - Valid 'flipY' config in 'GPUImageCopyExternalImage' (named 'srcDoFlipYDuringCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo' (named 'srcDoFlipYDuringCopy' in cases) - Valid subrect copies. And the expected results are all passed. diff --git a/src/webgpu/web_platform/copyToTexture/video.spec.ts b/src/webgpu/web_platform/copyToTexture/video.spec.ts index e6e641d982e3..c5e2695f79a9 100644 --- a/src/webgpu/web_platform/copyToTexture/video.spec.ts +++ b/src/webgpu/web_platform/copyToTexture/video.spec.ts @@ -34,12 +34,12 @@ It creates HTMLVideoElement with videos under Resource folder. Then call copyExternalImageToTexture() to do a full copy to the 0 mipLevel of dst texture, and read the contents out to compare with the ImageBitmap contents. - If 'flipY' in 'GPUImageCopyExternalImage' is set to 'true', copy will ensure the result + If 'flipY' in 'GPUCopyExternalImageSourceInfo' is set to 'true', copy will ensure the result is flipped. The tests covers: - Video comes from different color spaces. - - Valid 'flipY' config in 'GPUImageCopyExternalImage' (named 'srcDoFlipYDuringCopy' in cases) + - Valid 'flipY' config in 'GPUCopyExternalImageSourceInfo' (named 'srcDoFlipYDuringCopy' in cases) - TODO: partial copy tests should be added - TODO: all valid dstColorFormat tests should be added. `