diff --git a/src/resources/cache/hashes.json b/src/resources/cache/hashes.json index ea368705c448..e0459422560f 100644 --- a/src/resources/cache/hashes.json +++ b/src/resources/cache/hashes.json @@ -97,7 +97,7 @@ "webgpu/shader/execution/unary/f16_arithmetic.bin": "4a20db6d", "webgpu/shader/execution/unary/f16_conversion.bin": "31f72f5a", "webgpu/shader/execution/unary/f32_arithmetic.bin": "f1c311cb", - "webgpu/shader/execution/unary/f32_conversion.bin": "57dc324c", + "webgpu/shader/execution/unary/f32_conversion.bin": "7539cdb3", "webgpu/shader/execution/unary/i32_arithmetic.bin": "de945eec", "webgpu/shader/execution/unary/i32_conversion.bin": "1728a03e", "webgpu/shader/execution/unary/u32_conversion.bin": "9e6ca0ce", diff --git a/src/resources/cache/webgpu/shader/execution/unary/f32_conversion.bin b/src/resources/cache/webgpu/shader/execution/unary/f32_conversion.bin index bdcc0c72988e..66b2bc73f889 100644 Binary files a/src/resources/cache/webgpu/shader/execution/unary/f32_conversion.bin and b/src/resources/cache/webgpu/shader/execution/unary/f32_conversion.bin differ diff --git a/src/webgpu/api/validation/buffer/mapping.spec.ts b/src/webgpu/api/validation/buffer/mapping.spec.ts index dfe5f4517f49..5bf1ba0a45c4 100644 --- a/src/webgpu/api/validation/buffer/mapping.spec.ts +++ b/src/webgpu/api/validation/buffer/mapping.spec.ts @@ -943,7 +943,7 @@ g.test('getMappedRange,disjointRanges') t.testGetMappedRangeCall(success, buffer, offset2, size2); }); -g.test('getMappedRange,disjoinRanges_many') +g.test('getMappedRange,disjointRanges_many') .desc('Test getting a lot of small ranges, and that the disjoint check checks them all.') .fn(async t => { const kStride = 256; diff --git a/src/webgpu/format_info.ts b/src/webgpu/format_info.ts index dcf7f6b77572..5bdd5fe4e569 100644 --- a/src/webgpu/format_info.ts +++ b/src/webgpu/format_info.ts @@ -1773,6 +1773,18 @@ export function isCompressedTextureFormat(format: GPUTextureFormat) { return format in kCompressedTextureFormatInfo; } +export function isDepthTextureFormat(format: GPUTextureFormat) { + return !!kTextureFormatInfo[format].depth; +} + +export function isStencilTextureFormat(format: GPUTextureFormat) { + return !!kTextureFormatInfo[format].stencil; +} + +export function isDepthOrStencilTextureFormat(format: GPUTextureFormat) { + return isDepthTextureFormat(format) || isStencilTextureFormat(format); +} + export const kCompatModeUnsupportedStorageTextureFormats: readonly GPUTextureFormat[] = [ 'rg32float', 'rg32sint', @@ -1796,6 +1808,13 @@ export function isRegularTextureFormat(format: GPUTextureFormat) { return format in kRegularTextureFormatInfo; } +/** + * Returns true of format is both compressed and a float format, for example 'bc6h-rgb-ufloat'. + */ +export function isCompressedFloatTextureFormat(format: GPUTextureFormat) { + return isCompressedTextureFormat(format) && format.includes('float'); +} + export const kFeaturesForFormats = getFeaturesForFormats(kAllTextureFormats); /** diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index 19451ceccf79..250ddfd17b5f 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -242,7 +242,7 @@ "webgpu:api,validation,buffer,destroy:while_mapped:*": { "subcaseMS": 1.150 }, "webgpu:api,validation,buffer,mapping:gc_behavior,mapAsync:*": { "subcaseMS": 32.200 }, "webgpu:api,validation,buffer,mapping:gc_behavior,mappedAtCreation:*": { "subcaseMS": 76.200 }, - "webgpu:api,validation,buffer,mapping:getMappedRange,disjoinRanges_many:*": { "subcaseMS": 73.700 }, + "webgpu:api,validation,buffer,mapping:getMappedRange,disjointRanges_many:*": { "subcaseMS": 73.700 }, "webgpu:api,validation,buffer,mapping:getMappedRange,disjointRanges:*": { "subcaseMS": 2.257 }, "webgpu:api,validation,buffer,mapping:getMappedRange,offsetAndSizeAlignment,mapped:*": { "subcaseMS": 3.119 }, "webgpu:api,validation,buffer,mapping:getMappedRange,offsetAndSizeAlignment,mappedAtCreation:*": { "subcaseMS": 5.611 }, @@ -1552,6 +1552,7 @@ "webgpu:shader,execution,expression,call,builtin,textureLoad:sampled_1d:*": { "subcaseMS": 83.312 }, "webgpu:shader,execution,expression,call,builtin,textureLoad:sampled_2d:*": { "subcaseMS": 96.737 }, "webgpu:shader,execution,expression,call,builtin,textureLoad:sampled_3d:*": { "subcaseMS": 158.534 }, + "webgpu:shader,execution,expression,call,builtin,textureLoad:storage_texel_formats:*": { "subcaseMS": 471.569 }, "webgpu:shader,execution,expression,call,builtin,textureNumLayers:arrayed:*": { "subcaseMS": 8.102 }, "webgpu:shader,execution,expression,call,builtin,textureNumLayers:sampled:*": { "subcaseMS": 2.101 }, "webgpu:shader,execution,expression,call,builtin,textureNumLayers:storage:*": { "subcaseMS": 8.000 }, @@ -1594,12 +1595,14 @@ "webgpu:shader,execution,expression,call,builtin,textureSampleLevel:sampled_3d_coords:*": { "subcaseMS": 118.901 }, "webgpu:shader,execution,expression,call,builtin,textureSampleLevel:sampled_array_2d_coords:*": { "subcaseMS": 822.400 }, "webgpu:shader,execution,expression,call,builtin,textureSampleLevel:sampled_array_3d_coords:*": { "subcaseMS": 817.200 }, + "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 }, "webgpu:shader,execution,expression,call,builtin,transpose:f32:*": { "subcaseMS": 75.887 }, @@ -1678,6 +1681,9 @@ "webgpu:shader,execution,expression,unary,f16_conversion:i32:*": { "subcaseMS": 24.557 }, "webgpu:shader,execution,expression,unary,f16_conversion:u32:*": { "subcaseMS": 84.500 }, "webgpu:shader,execution,expression,unary,f32_arithmetic:negation:*": { "subcaseMS": 16.400 }, + "webgpu:shader,execution,expression,unary,f32_conversion:abstract_float:*": { "subcaseMS": 688.718 }, + "webgpu:shader,execution,expression,unary,f32_conversion:abstract_float_mat:*": { "subcaseMS": 1409.951 }, + "webgpu:shader,execution,expression,unary,f32_conversion:abstract_int:*": { "subcaseMS": 506.131 }, "webgpu:shader,execution,expression,unary,f32_conversion:bool:*": { "subcaseMS": 7.182 }, "webgpu:shader,execution,expression,unary,f32_conversion:f16:*": { "subcaseMS": 107.463 }, "webgpu:shader,execution,expression,unary,f32_conversion:f16_mat:*": { "subcaseMS": 60.170 }, 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 30cc4fff5286..079e82b66cc0 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts @@ -14,12 +14,60 @@ If an out of bounds access occurs, the built-in function returns one of: * The data for some texel within bounds of the texture * A vector (0,0,0,0) or (0,0,0,1) of the appropriate type for non-depth textures * 0.0 for depth textures + +TODO: Test textureLoad with depth textures as texture_2d, etc... `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { unreachable, iterRange } from '../../../../../../common/util/util.js'; +import { + isCompressedFloatTextureFormat, + isDepthTextureFormat, + kCompressedTextureFormats, + kEncodableTextureFormats, +} from '../../../../../format_info.js'; import { GPUTest } from '../../../../../gpu_test.js'; +import { + kFloat32Format, + kFloat16Format, + numberToFloatBits, + pack4x8unorm, + pack4x8snorm, +} from '../../../../../util/conversion.js'; +import { TexelFormats } from '../../../../types.js'; + +import { + TextureCall, + checkCallResults, + chooseTextureSize, + createTextureWithRandomDataAndGetTexels, + doTextureCalls, + appendComponentTypeForFormatToTextureType, + vec2, +} from './texture_utils.js'; +import { + Boundary, + LevelSpec, + generateCoordBoundaries, + getCoordinateForBoundaries, + getMipLevelFromLevelSpec, + isBoundaryNegative, + isLevelSpecNegative, +} from './utils.js'; + +const kTestableColorFormats = [...kEncodableTextureFormats, ...kCompressedTextureFormats] as const; + +function filterOutDepthAndCompressedFloatTextureFormats({ format }: { format: GPUTextureFormat }) { + return !isDepthTextureFormat(format) && !isCompressedFloatTextureFormat(format); +} -import { generateCoordBoundaries } from './utils.js'; +function filterOutU32WithNegativeValues(t: { + C: 'i32' | 'u32'; + level: LevelSpec; + coordsBoundary: Boundary; +}) { + return t.C === 'i32' || (!isLevelSpecNegative(t.level) && !isBoundaryNegative(t.coordsBoundary)); +} export const g = makeTestGroup(GPUTest); @@ -50,8 +98,9 @@ g.test('sampled_2d') .desc( ` C is i32 or u32 +L is i32 or u32 -fn textureLoad(t: texture_2d, coords: vec2, level: C) -> vec4 +fn textureLoad(t: texture_2d, coords: vec2, level: L) -> vec4 Parameters: * t: The sampled texture to read from @@ -61,11 +110,58 @@ Parameters: ) .params(u => u + .combine('format', kTestableColorFormats) + .filter(filterOutDepthAndCompressedFloatTextureFormats) + .beginSubcases() .combine('C', ['i32', 'u32'] as const) - .combine('coords', generateCoordBoundaries(2)) - .combine('level', [-1, 0, `numlevels-1`, `numlevels`] as const) + .combine('L', ['i32', 'u32'] as const) + .combine('coordsBoundary', generateCoordBoundaries(2)) + .combine('level', [-1, 0, `numLevels-1`, `numLevels`] as const) + .filter(filterOutU32WithNegativeValues) ) - .unimplemented(); + .beforeAllSubcases(t => { + const { format } = t.params; + t.skipIfTextureFormatNotSupported(format); + t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); + }) + .fn(async t => { + const { format, C, L, coordsBoundary, level } = 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 }); + + const descriptor: GPUTextureDescriptor = { + format, + size: { width, height }, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + const mipLevel = getMipLevelFromLevelSpec(texture.mipLevelCount, level); + const coords = getCoordinateForBoundaries(texture, mipLevel, coordsBoundary); + + const calls: TextureCall[] = [ + { + builtin: 'textureLoad', + coordType: C === 'i32' ? 'i' : 'u', + levelType: L === 'i32' ? 'i' : 'u', + mipLevel, + coords, + }, + ]; + const textureType = appendComponentTypeForFormatToTextureType('texture_2d', texture.format); + const viewDescriptor = {}; + const sampler = undefined; + const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const res = await checkCallResults( + t, + { texels, descriptor, viewDescriptor }, + textureType, + sampler, + calls, + results + ); + t.expectOK(res); + }); g.test('sampled_3d') .specURL('https://www.w3.org/TR/WGSL/#textureload') @@ -85,7 +181,7 @@ Parameters: u .combine('C', ['i32', 'u32'] as const) .combine('coords', generateCoordBoundaries(3)) - .combine('level', [-1, 0, `numlevels-1`, `numlevels`] as const) + .combine('level', [-1, 0, `numLevels-1`, `numLevels`] as const) ) .unimplemented(); @@ -135,7 +231,7 @@ Parameters: u .combine('C', ['i32', 'u32'] as const) .combine('coords', generateCoordBoundaries(2)) - .combine('level', [-1, 0, `numlevels-1`, `numlevels`] as const) + .combine('level', [-1, 0, `numLevels-1`, `numLevels`] as const) ) .unimplemented(); @@ -180,6 +276,334 @@ Parameters: .combine('C', ['i32', 'u32'] as const) .combine('coords', generateCoordBoundaries(2)) .combine('array_index', [-1, 0, `numlayers-1`, `numlayers`] as const) - .combine('level', [-1, 0, `numlevels-1`, `numlevels`] as const) + .combine('level', [-1, 0, `numLevels-1`, `numLevels`] as const) ) .unimplemented(); + +// Returns texel values to use as inputs for textureLoad. +// Values are kept simple to avoid rounding issues. +function shaderValues(format: string, type: string) { + switch (type) { + case 'f32': { + switch (format) { + case 'rbga8snorm': + // prettier-ignore + return [ + { r: 0.0, g: 0.0, b: 0.0, a: 0.0, }, + { r: 0.2, g: 0.4, b: 0.6, a: 0.8, }, + { r: -0.2, g: -0.4, b: -0.6, a: -0.8, }, + { r: 0.2, g: -0.4, b: 0.6, a: -0.8, }, + { r: -0.2, g: 0.4, b: -0.6, a: 0.8, }, + { r: 0.2, g: 0.2, b: 0.2, a: 0.2, }, + { r: -0.2, g: -0.2, b: -0.2, a: -0.2, }, + { r: 0.4, g: 0.4, b: 0.4, a: 0.4, }, + { r: -0.4, g: -0.4, b: -0.4, a: -0.4, }, + { r: 0.6, g: 0.6, b: 0.6, a: 0.6, }, + { r: -0.6, g: -0.6, b: -0.6, a: -0.6, }, + { r: 0.8, g: 0.8, b: 0.8, a: 0.8, }, + { r: -0.8, g: -0.8, b: -0.8, a: -0.8, }, + ]; + case 'rgba8unorm': + case 'bgra8unorm': + // prettier-ignore + return [ + { r: 0.0, g: 0.0, b: 0.0, a: 0.0, }, + { r: 0.2, g: 0.4, b: 0.6, a: 0.8, }, + { r: 0.9, g: 0.4, b: 0.6, a: 0.8, }, + { r: 0.2, g: 0.9, b: 0.6, a: 0.8, }, + { r: 0.2, g: 0.4, b: 0.9, a: 0.8, }, + { r: 0.2, g: 0.4, b: 0.6, a: 0.9, }, + { r: 0.2, g: 0.2, b: 0.2, a: 0.2, }, + { r: 0.4, g: 0.4, b: 0.4, a: 0.4, }, + { r: 0.6, g: 0.6, b: 0.6, a: 0.6, }, + { r: 0.8, g: 0.8, b: 0.8, a: 0.8, }, + ]; + default: + // Stick within 16-bit ranges. + // prettier-ignore + return [ + { r: 100, g: 128, b: 100, a: 128, }, + { r: 64, g: 32, b: 32, a: 64, }, + { r: 8, g: 0, b: 8, a: 0, }, + { r: 0, g: 0, b: 0, a: 0, }, + { r: -100, g: 128, b: 100, a: 128, }, + { r: -64, g: 32, b: 32, a: 64, }, + { r: -8, g: 0, b: 8, a: 0, }, + { r: 100, g: -128, b: 100, a: 128, }, + { r: 64, g: -32, b: 32, a: 64, }, + { r: 8, g: 0, b: 8, a: 0, }, + { r: 100, g: 128, b: -100, a: 128, }, + { r: 64, g: 32, b: -32, a: 64, }, + { r: 8, g: 0, b: -8, a: 0, }, + { r: 100, g: 128, b: 100, a: -128, }, + { r: 64, g: 32, b: 32, a: -64, }, + { r: 8, g: 0, b: 8, a: 0, }, + ]; + } + break; + } + case 'u32': + // Keep all ranges within u8. + // prettier-ignore + return [ + { r: 0, g: 0, b: 0, a: 0, }, + { r: 0, g: 8, b: 16, a: 128, }, + { r: 8, g: 16, b: 32, a: 64, }, + { r: 16, g: 32, b: 64, a: 128, }, + { r: 255, g: 254, b: 253, a: 252, }, + { r: 255, g: 255, b: 255, a: 255, }, + { r: 128, g: 64, b: 32, a: 16, }, + { r: 64, g: 32, b: 16, a: 8, }, + { r: 32, g: 16, b: 8, a: 0, }, + ]; + case 'i32': + // Keep all ranges i8 + // prettier-ignore + return [ + { r: 0, g: 0, b: 0, a: 0, }, + { r: 0, g: -8, b: 16, a: 127, }, + { r: 8, g: 16, b: -32, a: 64, }, + { r: -16, g: 32, b: 64, a: -128, }, + { r: 127, g: 126, b: 125, a: 124, }, + { r: -128, g: -127, b: -126, a: -125, }, + { r: 127, g: 127, b: 127, a: 127, }, + { r: -128, g: -128, b: -128, a: -128, }, + ]; + default: + unreachable(`unhandled shader type ${type}`); + break; + } + return []; +} + +g.test('storage_texel_formats') + .desc('Test loading of texel formats') + .params(u => u.combineWithParams([...TexelFormats, { format: 'bgra8unorm', _shaderType: 'f32' }])) + .beforeAllSubcases(t => { + t.skipIf(!t.hasLanguageFeature('readonly_and_readwrite_storage_textures')); + if (t.params.format === 'bgra8unorm') { + t.selectDeviceOrSkipTestCase('bgra8unorm-storage'); + } else { + t.skipIfTextureFormatNotUsableAsStorageTexture(t.params.format as GPUTextureFormat); + } + }) + .fn(t => { + const { format, _shaderType } = t.params; + const values = shaderValues(format, _shaderType); + + // To avoid rounding issues, unorm and snorm values are repacked in the shader. + let useType = _shaderType; + let assignValue = `v`; + if (format === 'bgra8unorm' || format === 'rgba8unorm') { + useType = 'u32'; + assignValue = `vec4u(pack4x8unorm(v),0,0,0)`; + } else if (format === 'rgba8snorm') { + useType = 'u32'; + assignValue = `vec4u(pack4x8snorm(v),0,0,0)`; + } + const wgsl = ` +requires readonly_and_readwrite_storage_textures; + +@group(0) @binding(0) +var tex : texture_storage_1d<${format}, read>; + +@group(0) @binding(1) +var out : array>; + +@compute @workgroup_size(${values.length}) +fn main(@builtin(global_invocation_id) gid : vec3u) { + let v = textureLoad(tex, gid.x); + out[gid.x] = ${assignValue}; +}`; + + const bytesPerRow = 256; + let bytesPerTexel = 4; + switch (format) { + case 'rgba16uint': + case 'rgba16sint': + case 'rgba16float': + case 'rg32uint': + case 'rg32sint': + case 'rg32float': + bytesPerTexel = 8; + break; + case 'rgba32uint': + case 'rgba32sint': + case 'rgba32float': + bytesPerTexel = 16; + break; + default: + break; + } + + const textureSize: GPUExtent3D = { + width: bytesPerRow / bytesPerTexel, + height: 1, + depthOrArrayLayers: 1, + }; + const texture = t.createTextureTracked({ + format: format as GPUTextureFormat, + dimension: '1d', + size: textureSize, + mipLevelCount: 1, + usage: GPUTextureUsage.STORAGE_BINDING | GPUTextureUsage.COPY_DST, + }); + const outputBuffer = t.makeBufferWithContents( + new Uint32Array([...iterRange(values.length * 4, x => 0)]), + GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE + ); + t.trackForCleanup(outputBuffer); + + const transformed = values.flatMap(x => { + switch (format) { + case 'rgba8unorm': + return pack4x8unorm(x.r, x.g, x.b, x.a); + case 'bgra8unorm': + return pack4x8unorm(x.b, x.g, x.r, x.a); + case 'rgba8snorm': + return pack4x8snorm(x.r, x.g, x.b, x.a); + case 'r32uint': + case 'r32sint': + return x.r; + case 'rg32uint': + case 'rg32sint': + return [x.r, x.g]; + case 'rgba32uint': + case 'rgba32sint': + return [x.r, x.g, x.b, x.a]; + case 'rgba8uint': + case 'rgba8sint': + return (x.r & 0xff) | ((x.g & 0xff) << 8) | ((x.b & 0xff) << 16) | ((x.a & 0xff) << 24); + case 'rgba16uint': + case 'rgba16sint': + return [(x.r & 0xffff) | ((x.g & 0xffff) << 16), (x.b & 0xffff) | ((x.a & 0xffff) << 16)]; + case 'r32float': + return numberToFloatBits(x.r, kFloat32Format); + case 'rg32float': + return [numberToFloatBits(x.r, kFloat32Format), numberToFloatBits(x.g, kFloat32Format)]; + case 'rgba32float': + return [ + numberToFloatBits(x.r, kFloat32Format), + numberToFloatBits(x.g, kFloat32Format), + numberToFloatBits(x.b, kFloat32Format), + numberToFloatBits(x.a, kFloat32Format), + ]; + case 'rgba16float': + return [ + (numberToFloatBits(x.r, kFloat16Format) & 0xffff) | + ((numberToFloatBits(x.g, kFloat16Format) & 0xffff) << 16), + (numberToFloatBits(x.b, kFloat16Format) & 0xffff) | + ((numberToFloatBits(x.a, kFloat16Format) & 0xffff) << 16), + ]; + default: + unreachable(`unhandled format ${format}`); + break; + } + return 0; + }); + + const texelBuffer = t.makeBufferWithContents( + new Uint32Array([ + ...iterRange(bytesPerRow, x => { + if (x < transformed.length) { + return transformed[x]; + } else { + return 0; + } + }), + ]), + GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE + ); + t.trackForCleanup(texelBuffer); + + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ + code: wgsl, + }), + entryPoint: 'main', + }, + }); + const bg = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: texture.createView({ + format: format as GPUTextureFormat, + dimension: '1d', + }), + }, + { + binding: 1, + resource: { + buffer: outputBuffer, + }, + }, + ], + }); + + const encoder = t.device.createCommandEncoder(); + encoder.copyBufferToTexture( + { + buffer: texelBuffer, + offset: 0, + bytesPerRow, + rowsPerImage: 1, + }, + { texture }, + textureSize + ); + + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup(0, bg); + pass.dispatchWorkgroups(1, 1, 1); + pass.end(); + t.queue.submit([encoder.finish()]); + + const expected = new Uint32Array( + values.flatMap(x => { + switch (format) { + case 'r32uint': + case 'r32sint': + return [x.r, 0, 0, 1]; + case 'rg32uint': + case 'rg32sint': + return [x.r, x.g, 0, 1]; + case 'r32float': + return [ + numberToFloatBits(x.r, kFloat32Format), + 0, + 0, + numberToFloatBits(1, kFloat32Format), + ]; + case 'rg32float': + return [ + numberToFloatBits(x.r, kFloat32Format), + numberToFloatBits(x.g, kFloat32Format), + 0, + numberToFloatBits(1, kFloat32Format), + ]; + case 'rgba32float': + case 'rgba16float': + return [ + numberToFloatBits(x.r, kFloat32Format), + numberToFloatBits(x.g, kFloat32Format), + numberToFloatBits(x.b, kFloat32Format), + numberToFloatBits(x.a, kFloat32Format), + ]; + case 'rgba8unorm': + case 'bgra8unorm': + return [pack4x8unorm(x.r, x.g, x.b, x.a), 0, 0, 0]; + case 'rgba8snorm': + return [pack4x8snorm(x.r, x.g, x.b, x.a), 0, 0, 0]; + default: + break; + } + return [x.r, x.g, x.b, x.a]; + }) + ); + t.expectGPUBufferValuesEqual(outputBuffer, expected); + }); 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 b8453011614a..ca7ae3d0655c 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts @@ -5,9 +5,35 @@ Returns the number of layers (elements) of an array texture. `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; -import { GPUTest } from '../../../../../gpu_test.js'; +import { kTextureFormatInfo } from '../../../../../format_info.js'; +import { TexelFormats } from '../../../../types.js'; -export const g = makeTestGroup(GPUTest); +import { kSampleTypeInfo, WGSLTextureQueryTest } from './texture_utils.js'; + +const kNumLayers = 36; + +function getLayerSettingsAndExpected({ + view_type, + isCubeArray, +}: { + view_type: 'full' | 'partial'; + isCubeArray?: boolean; +}) { + const divisor = isCubeArray ? 6 : 1; + return view_type === 'partial' + ? { + baseArrayLayer: 11, + arrayLayerCount: 6, + expected: [6 / divisor], + } + : { + baseArrayLayer: 0, + arrayLayerCount: kNumLayers, + expected: [kNumLayers / divisor], + }; +} + +export const g = makeTestGroup(WGSLTextureQueryTest); g.test('sampled') .specURL('https://www.w3.org/TR/WGSL/#texturenumlayers') @@ -26,9 +52,49 @@ Parameters u .combine('texture_type', ['texture_2d_array', 'texture_cube_array'] as const) .beginSubcases() - .combine('sampled_type', ['f32-only', 'i32', 'u32'] as const) + .combine('sampled_type', ['f32', 'i32', 'u32'] as const) + .combine('view_type', ['full', 'partial'] as const) ) - .unimplemented(); + .beforeAllSubcases(t => { + t.skipIf( + t.isCompatibility && t.params.view === 'partial', + 'compatibility mode does not support partial layer views' + ); + t.skipIf( + t.isCompatibility && t.params.texture_type === 'texture_cube_array', + 'compatibility mode does not support cube arrays' + ); + }) + .fn(t => { + const { texture_type, sampled_type, view_type } = t.params; + const { format } = kSampleTypeInfo[sampled_type]; + + const texture = t.createTextureTracked({ + format, + usage: GPUTextureUsage.TEXTURE_BINDING, + size: [1, 1, kNumLayers], + }); + + 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); +} + `; + + const { baseArrayLayer, arrayLayerCount, expected } = getLayerSettingsAndExpected({ + view_type, + isCubeArray: texture_type === 'texture_cube_array', + }); + const view = texture.createView({ + dimension: texture_type === 'texture_2d_array' ? '2d-array' : 'cube-array', + baseArrayLayer, + arrayLayerCount, + }); + + t.executeAndExpectResult(code, view, expected); + }); g.test('arrayed') .specURL('https://www.w3.org/TR/WGSL/#texturenumlayers') @@ -42,9 +108,50 @@ Parameters ` ) .params(u => - u.combine('texture_type', ['texture_depth_2d_array', 'texture_depth_cube_array'] as const) + u + .combine('texture_type', ['texture_depth_2d_array', 'texture_depth_cube_array'] as const) + .beginSubcases() + .combine('view_type', ['full', 'partial'] as const) ) - .unimplemented(); + .beforeAllSubcases(t => { + t.skipIf( + t.isCompatibility && t.params.view === 'partial', + 'compatibility mode does not support partial layer views' + ); + t.skipIf( + t.isCompatibility && t.params.texture_type === 'texture_depth_cube_array', + 'compatibility mode does not support cube arrays' + ); + }) + .fn(t => { + const { texture_type, view_type } = t.params; + + const texture = t.createTextureTracked({ + format: 'depth32float', + usage: GPUTextureUsage.TEXTURE_BINDING, + size: [1, 1, kNumLayers], + }); + + 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); +} + `; + + const { baseArrayLayer, arrayLayerCount, expected } = getLayerSettingsAndExpected({ + view_type, + isCubeArray: texture_type === 'texture_depth_cube_array', + }); + const view = texture.createView({ + dimension: texture_type === 'texture_depth_2d_array' ? '2d-array' : 'cube-array', + baseArrayLayer, + arrayLayerCount, + }); + + t.executeAndExpectResult(code, view, expected); + }); g.test('storage') .specURL('https://www.w3.org/TR/WGSL/#texturenumlayers') @@ -76,25 +183,40 @@ Parameters ) .params(u => u + .combineWithParams(TexelFormats) .beginSubcases() - .combine('texel_format', [ - 'rgba8unorm', - 'rgba8snorm', - 'rgba8uint', - 'rgba8sint', - 'rgba16uint', - 'rgba16sint', - 'rgba16float', - 'r32uint', - 'r32sint', - 'r32float', - 'rg32uint', - 'rg32sint', - 'rg32float', - 'rgba32uint', - 'rgba32sint', - 'rgba32float', - ] as const) .combine('access_mode', ['read', 'write', 'read_write'] as const) + .filter( + t => t.access_mode !== 'read_write' || kTextureFormatInfo[t.format].color?.readWriteStorage + ) + .combine('view_type', ['full', 'partial'] as const) ) - .unimplemented(); + .beforeAllSubcases(t => t.skipIfTextureFormatNotUsableAsStorageTexture(t.params.format)) + .fn(t => { + const { format, access_mode, view_type } = t.params; + + const texture = t.createTextureTracked({ + format, + usage: GPUTextureUsage.STORAGE_BINDING, + size: [1, 1, kNumLayers], + }); + + 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); +} + `; + + const { baseArrayLayer, arrayLayerCount, expected } = getLayerSettingsAndExpected({ + view_type, + }); + const view = texture.createView({ + dimension: '2d-array', + baseArrayLayer, + arrayLayerCount, + }); + + t.executeAndExpectResult(code, view, 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 4204397b23b8..5610701601cb 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureNumLevels.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureNumLevels.spec.ts @@ -5,9 +5,38 @@ Returns the number of mip levels of a texture. `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; -import { GPUTest } from '../../../../../gpu_test.js'; +import { getTextureDimensionFromView } from '../../../../../util/texture/base.js'; -export const g = makeTestGroup(GPUTest); +import { kSampleTypeInfo, WGSLTextureQueryTest } from './texture_utils.js'; + +function getLevelSettingsAndExpected(viewType: 'full' | 'partial', mipLevelCount: number) { + return viewType === 'partial' + ? { + baseMipLevel: 1, + mipLevelCount: 2, + expected: [2], + } + : { + baseMipLevel: 0, + mipLevelCount, + expected: [mipLevelCount], + }; +} + +const kTextureTypeToViewDimension = { + texture_1d: '1d', + texture_2d: '2d', + texture_2d_array: '2d-array', + texture_3d: '3d', + texture_cube: 'cube', + texture_cube_array: 'cube-array', + texture_depth_2d: '2d', + texture_depth_2d_array: '2d-array', + texture_depth_cube: 'cube', + texture_depth_cube_array: 'cube-array', +} as const; + +export const g = makeTestGroup(WGSLTextureQueryTest); g.test('sampled') .specURL('https://www.w3.org/TR/WGSL/#texturenumlevels') @@ -34,12 +63,60 @@ Parameters 'texture_2d_array', 'texture_3d', 'texture_cube', - 'texture_cube_array`', + 'texture_cube_array', ] as const) .beginSubcases() - .combine('sampled_type', ['f32-only', 'i32', 'u32'] as const) + .combine('sampled_type', ['f32', 'i32', 'u32'] as const) + .combine('view_type', ['full', 'partial'] as const) + // 1d textures can't have mipLevelCount > 0 + .filter(t => t.texture_type !== 'texture_1d' || t.view_type !== 'partial') ) - .unimplemented(); + .beforeAllSubcases(t => { + t.skipIfTextureViewDimensionNotSupported(kTextureTypeToViewDimension[t.params.texture_type]); + }) + .fn(t => { + const { texture_type, sampled_type, view_type } = t.params; + const { format } = kSampleTypeInfo[sampled_type]; + + const viewDimension = kTextureTypeToViewDimension[texture_type]; + const dimension = getTextureDimensionFromView(viewDimension); + const isCube = texture_type.includes('cube'); + const width = 64; + const height = dimension === '1d' ? 1 : width; + const depthOrArrayLayers = isCube ? 6 : 1; + const mipCount = dimension === '1d' ? 1 : 4; + const texture = t.createTextureTracked({ + format, + dimension, + usage: GPUTextureUsage.TEXTURE_BINDING, + size: { + width, + height, + depthOrArrayLayers, + }, + mipLevelCount: mipCount, + }); + + 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); +} + `; + + const { baseMipLevel, mipLevelCount, expected } = getLevelSettingsAndExpected( + view_type, + mipCount + ); + const view = texture.createView({ + dimension: viewDimension, + baseMipLevel, + mipLevelCount, + }); + + t.executeAndExpectResult(code, view, expected); + }); g.test('depth') .specURL('https://www.w3.org/TR/WGSL/#texturenumlevels') @@ -55,11 +132,57 @@ Parameters ` ) .params(u => - u.combine('texture_type', [ - 'texture_depth_2d', - 'texture_depth_2d_array', - 'texture_depth_cube', - 'texture_depth_cube_array', - ] as const) + u + .combine('texture_type', [ + 'texture_depth_2d', + 'texture_depth_2d_array', + 'texture_depth_cube', + 'texture_depth_cube_array', + ] as const) + .combine('view_type', ['full', 'partial'] as const) ) - .unimplemented(); + .beforeAllSubcases(t => { + t.skipIfTextureViewDimensionNotSupported(kTextureTypeToViewDimension[t.params.texture_type]); + }) + .fn(t => { + const { texture_type, view_type } = t.params; + + const viewDimension = kTextureTypeToViewDimension[texture_type]; + const dimension = getTextureDimensionFromView(viewDimension); + const isCube = texture_type.includes('cube'); + const width = 64; + const height = dimension === '1d' ? 1 : width; + const depthOrArrayLayers = isCube ? 6 : 1; + const mipCount = dimension === '1d' ? 1 : 4; + const texture = t.createTextureTracked({ + format: 'depth32float', + dimension, + usage: GPUTextureUsage.TEXTURE_BINDING, + size: { + width, + height, + depthOrArrayLayers, + }, + mipLevelCount: mipCount, + }); + + 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); +} + `; + + const { baseMipLevel, mipLevelCount, expected } = getLevelSettingsAndExpected( + view_type, + mipCount + ); + const view = texture.createView({ + dimension: viewDimension, + baseMipLevel, + mipLevelCount, + }); + + t.executeAndExpectResult(code, view, 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 26bda6cd48e4..a6314198529b 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureNumSamples.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureNumSamples.spec.ts @@ -5,9 +5,10 @@ Returns the number samples per texel in a multisampled texture. `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; -import { GPUTest } from '../../../../../gpu_test.js'; -export const g = makeTestGroup(GPUTest); +import { kSampleTypeInfo, WGSLTextureQueryTest } from './texture_utils.js'; + +export const g = makeTestGroup(WGSLTextureQueryTest); g.test('sampled') .specURL('https://www.w3.org/TR/WGSL/#texturenumsamples') @@ -21,8 +22,32 @@ Parameters * t The multisampled texture. ` ) - .params(u => u.beginSubcases().combine('sampled_type', ['f32-only', 'i32', 'u32'] as const)) - .unimplemented(); + .params(u => u.beginSubcases().combine('sampled_type', ['f32', 'i32', 'u32'] as const)) + .fn(t => { + const { sampled_type } = t.params; + const { format } = kSampleTypeInfo[sampled_type]; + + const sampleCount = 4; + const texture = t.createTextureTracked({ + format, + usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.RENDER_ATTACHMENT, + size: [1, 1, 1], + sampleCount, + }); + + 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); +} + `; + + const expected = [sampleCount]; + const view = texture.createView({}); + + t.executeAndExpectResult(code, view, expected); + }); g.test('depth') .specURL('https://www.w3.org/TR/WGSL/#texturenumsamples') @@ -34,4 +59,25 @@ Parameters * t The multisampled texture. ` ) - .unimplemented(); + .fn(t => { + const sampleCount = 4; + const texture = t.createTextureTracked({ + format: 'depth32float', + usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.RENDER_ATTACHMENT, + size: [1, 1, 1], + sampleCount, + }); + + 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); +} + `; + + const expected = [sampleCount]; + const view = texture.createView({}); + + t.executeAndExpectResult(code, view, expected); + }); 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 eaef386314d3..09b48b13ce63 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureStore.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureStore.spec.ts @@ -12,8 +12,15 @@ If an out-of-bounds access occurs, the built-in function should not be executed. `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; -import { iterRange } from '../../../../../../common/util/util.js'; +import { unreachable, iterRange } from '../../../../../../common/util/util.js'; import { GPUTest, TextureTestMixin } from '../../../../../gpu_test.js'; +import { + kFloat32Format, + kFloat16Format, + numberToFloatBits, + pack4x8unorm, + pack4x8snorm, +} from '../../../../../util/conversion.js'; import { virtualMipSize } from '../../../../../util/texture/base.js'; import { TexelFormats } from '../../../../types.js'; @@ -121,6 +128,340 @@ Parameters: ) .unimplemented(); +// Returns shader input values for texel format tests. +// Values are intentionally simple to avoid rounding issues. +function inputArray(format: string): number[] { + switch (format) { + case 'rgba8snorm': + return [-1.1, 1.0, -0.6, -0.3, 0, 0.3, 0.6, 1.0, 1.1]; + case 'rgba8unorm': + case 'bgra8unorm': + return [-0.1, 0, 0.2, 0.4, 0.6, 0.8, 1.0, 1.1]; + case 'rgba8uint': + case 'rgba16uint': + case 'rgba32uint': + case 'r32uint': + case 'rg32uint': + // Stick within 8-bit ranges for simplicity. + return [0, 8, 16, 24, 32, 64, 100, 128, 200, 255]; + case 'rgba8sint': + case 'rgba16sint': + case 'rgba32sint': + 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 'rgba16float': + case 'rgba32float': + case 'r32float': + case 'rg32float': + // Stick with simple values. + return [-100, -50, -32, -16, -8, -1, 0, 1, 8, 16, 32, 50, 100]; + default: + unreachable(`unhandled format ${format}`); + break; + } + return []; +} + +g.test('texel_formats') + .desc(`Test storage of texel formats`) + .params(u => u.combineWithParams([...TexelFormats, { format: 'bgra8unorm', _shaderType: 'f32' }])) + .beforeAllSubcases(t => { + if (t.params.format === 'bgra8unorm') { + t.selectDeviceOrSkipTestCase('bgra8unorm-storage'); + } else { + t.skipIfTextureFormatNotUsableAsStorageTexture(t.params.format as GPUTextureFormat); + } + }) + .fn(t => { + const { format, _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 += ` +); + +@group(0) @binding(0) +var tex : texture_storage_1d<${format}, write>; + +@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); +} +`; + + const numTexels = values.length; + const textureSize: GPUExtent3D = { width: numTexels, height: 1, depthOrArrayLayers: 1 }; + const texture = t.createTextureTracked({ + format: format as GPUTextureFormat, + dimension: '1d', + size: textureSize, + mipLevelCount: 1, + usage: GPUTextureUsage.STORAGE_BINDING | GPUTextureUsage.COPY_SRC, + }); + + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ + code: wgsl, + }), + entryPoint: 'main', + }, + }); + const bg = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: texture.createView({ + format: format as GPUTextureFormat, + dimension: '1d', + }), + }, + ], + }); + + const encoder = t.device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup(0, bg); + pass.dispatchWorkgroups(1, 1, 1); + pass.end(); + t.queue.submit([encoder.finish()]); + + let bytesPerTexel = 4; + switch (format) { + case 'rgba16uint': + case 'rgba16sint': + case 'rgba16float': + case 'rg32uint': + case 'rg32sint': + case 'rg32float': + bytesPerTexel = 8; + break; + case 'rgba32uint': + case 'rgba32sint': + case 'rgba32float': + bytesPerTexel = 16; + break; + default: + break; + } + + let zeroChannel = 0; + const buffer = t.copyWholeTextureToNewBufferSimple(texture, 0); + const uintsPerTexel = bytesPerTexel / 4; + 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]; + switch (format) { + case 'rgba8unorm': { + const vals = [shaderVal, shaderVal, shaderVal, shaderVal]; + vals[zeroChannel++] = 0; + return pack4x8unorm(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]); + } + case 'rgba8snorm': { + const vals = [shaderVal, shaderVal, shaderVal, shaderVal]; + vals[zeroChannel++] = 0; + return pack4x8snorm(vals[0], vals[1], vals[2], vals[3]); + } + case 'r32uint': + case 'r32sint': + return shaderVal; + case 'rg32uint': + case 'rgba32uint': + case 'rg32sint': + case 'rgba32sint': { + const maskedVal = channel === zeroChannel++ ? 0 : shaderVal; + return maskedVal; + } + case 'rgba8uint': + case 'rgba8sint': { + const vals = [shaderVal, shaderVal, shaderVal, shaderVal]; + vals[zeroChannel++] = 0; + return ( + ((vals[3] & 0xff) << 24) | + ((vals[2] & 0xff) << 16) | + ((vals[1] & 0xff) << 8) | + (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++; + } + return ((vals[1] & 0xffff) << 16) | (vals[0] & 0xffff); + } + case 'r32float': { + return numberToFloatBits(shaderVal, kFloat32Format); + } + case 'rg32float': + case 'rgba32float': { + const maskedVal = channel === zeroChannel++ ? 0 : shaderVal; + return numberToFloatBits(maskedVal, 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++; + } + return ((vals[1] & 0xffff) << 16) | (vals[0] & 0xffff); + } + default: + unreachable(`unhandled format ${format}`); + break; + } + return 0; + }), + ]); + t.expectGPUBufferValuesEqual(buffer, expected); + }); + +g.test('bgra8unorm_swizzle') + .desc('Test bgra8unorm swizzling') + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('bgra8unorm-storage'); + }) + .fn(t => { + const values = [ + { r: -1.1, g: 0.6, b: 0.4, a: 1 }, + { r: 1.1, g: 0.6, b: 0.4, a: 1 }, + { r: 0.4, g: -1.1, b: 0.6, a: 1 }, + { r: 0.4, g: 1.1, b: 0.6, a: 1 }, + { r: 0.6, g: 0.4, b: -1.1, a: 1 }, + { r: 0.6, g: 0.4, b: 1.1, a: 1 }, + { r: 0.2, g: 0.4, b: 0.6, a: 1 }, + { r: -0.2, g: -0.4, b: -0.6, a: 1 }, + ]; + let wgsl = ` +@group(0) @binding(0) var tex : texture_storage_1d; + +const values = array(`; + for (const v of values) { + wgsl += `vec4(${v.r},${v.g},${v.b},${v.a}),\n`; + } + wgsl += `); + +@compute @workgroup_size(${values.length}) +fn main(@builtin(global_invocation_id) gid : vec3u) { + let value = values[gid.x]; + textureStore(tex, gid.x, value); +}`; + + const numTexels = values.length; + const textureSize: GPUExtent3D = { width: numTexels, height: 1, depthOrArrayLayers: 1 }; + const texture = t.createTextureTracked({ + format: 'bgra8unorm', + dimension: '1d', + size: textureSize, + mipLevelCount: 1, + usage: GPUTextureUsage.STORAGE_BINDING | GPUTextureUsage.COPY_SRC, + }); + + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ + code: wgsl, + }), + entryPoint: 'main', + }, + }); + const bg = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: texture.createView({ + format: 'bgra8unorm', + dimension: '1d', + }), + }, + ], + }); + + const encoder = t.device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup(0, bg); + pass.dispatchWorkgroups(1, 1, 1); + pass.end(); + t.queue.submit([encoder.finish()]); + + const buffer = t.copyWholeTextureToNewBufferSimple(texture, 0); + const expected = new Uint32Array([ + ...iterRange(numTexels, x => { + const { r, g, b, a } = values[x]; + return pack4x8unorm(b, g, r, a); + }), + ]); + t.expectGPUBufferValuesEqual(buffer, expected); + }); + // Texture width for dimensions >1D. // Sized such that mip level 2 will be at least 256 bytes/row. const kWidth = 256; 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 815761a709d5..831a05f4d1ba 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -2,7 +2,9 @@ import { keysOf } from '../../../../../../common/util/data_tables.js'; import { assert, range, unreachable } from '../../../../../../common/util/util.js'; import { EncodableTextureFormat, + isCompressedFloatTextureFormat, isCompressedTextureFormat, + isDepthOrStencilTextureFormat, kEncodableTextureFormats, kTextureFormatInfo, } from '../../../../../format_info.js'; @@ -20,17 +22,76 @@ import { import { effectiveViewDimensionForDimension, physicalMipSizeFromTexture, + reifyTextureDescriptor, virtualMipSize, } from '../../../../../util/texture/base.js'; import { kTexelRepresentationInfo, + NumericRange, + PerComponentNumericRange, PerTexelComponent, + TexelComponent, TexelRepresentationInfo, } from '../../../../../util/texture/texel_data.js'; import { TexelView } from '../../../../../util/texture/texel_view.js'; import { createTextureFromTexelViews } from '../../../../../util/texture.js'; import { reifyExtent3D } from '../../../../../util/unions.js'; +export type SampledType = 'f32' | 'i32' | 'u32'; + +export const kSampleTypeInfo = { + f32: { + format: 'rgba8unorm', + }, + i32: { + format: 'rgba8sint', + }, + u32: { + format: 'rgba8uint', + }, +} as const; + +/** + * Used for textureDimension, textureNumLevels, textureNumLayers + */ +export class WGSLTextureQueryTest extends GPUTest { + executeAndExpectResult(code: string, view: GPUTextureView, expected: number[]) { + const { device } = this; + const module = device.createShaderModule({ code }); + const pipeline = device.createComputePipeline({ + layout: 'auto', + compute: { + module, + }, + }); + + const resultBuffer = this.createBufferTracked({ + size: 16, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + + const bindGroup = device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: view }, + { binding: 1, resource: { buffer: resultBuffer } }, + ], + }); + + 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()]); + + const e = new Uint32Array(4); + e.set(expected); + this.expectGPUBufferValuesEqual(resultBuffer, e); + } +} + function getLimitValue(v: number) { switch (v) { case Number.POSITIVE_INFINITY: @@ -44,13 +105,15 @@ function getLimitValue(v: number) { function getValueBetweenMinAndMaxTexelValueInclusive( rep: TexelRepresentationInfo, + component: TexelComponent, normalized: number ) { - return lerp( - getLimitValue(rep.numericRange!.min), - getLimitValue(rep.numericRange!.max), - normalized - ); + assert(!!rep.numericRange); + const perComponentRanges = rep.numericRange as PerComponentNumericRange; + const perComponentRange = perComponentRanges[component]; + const range = rep.numericRange as NumericRange; + const { min, max } = perComponentRange ? perComponentRange : range; + return lerp(getLimitValue(min), getLimitValue(max), normalized); } /** @@ -63,6 +126,49 @@ export function getTexelViewFormatForTextureFormat(format: GPUTextureFormat) { return format.endsWith('-srgb') ? 'rgba8unorm-srgb' : 'rgba32float'; } +const kTextureTypeInfo = { + depth: { + componentType: 'f32', + resultType: 'vec4f', + resultFormat: 'rgba32float', + }, + float: { + componentType: 'f32', + resultType: 'vec4f', + resultFormat: 'rgba32float', + }, + 'unfilterable-float': { + componentType: 'f32', + resultType: 'vec4f', + resultFormat: 'rgba32float', + }, + sint: { + componentType: 'i32', + resultType: 'vec4i', + resultFormat: 'rgba32sint', + }, + uint: { + componentType: 'u32', + resultType: 'vec4u', + resultFormat: 'rgba32uint', + }, +} as const; + +function getTextureFormatTypeInfo(format: GPUTextureFormat) { + const info = kTextureFormatInfo[format]; + const type = info.color?.type ?? info.depth?.type ?? info.stencil?.type; + assert(!!type); + return kTextureTypeInfo[type]; +} + +/** + * given a texture type 'base', returns the base with the correct component for the given texture format. + * eg: `getTextureType('texture_2d', someUnsignedIntTextureFormat)` -> `texture_2d` + */ +export function appendComponentTypeForFormatToTextureType(base: string, format: GPUTextureFormat) { + return `${base}<${getTextureFormatTypeInfo(format).componentType}>`; +} + /** * Creates a TexelView filled with random values. */ @@ -76,7 +182,7 @@ export function createRandomTexelView(info: { for (const component of rep.componentOrder) { const rnd = hashU32(coords.x, coords.y, coords.z, component.charCodeAt(0)); const normalized = clamp(rnd / 0xffffffff, { min: 0, max: 1 }); - texel[component] = getValueBetweenMinAndMaxTexelValueInclusive(rep, normalized); + texel[component] = getValueBetweenMinAndMaxTexelValueInclusive(rep, component, normalized); } return quantize(texel, rep); }; @@ -122,6 +228,7 @@ export interface TextureCallArgs { coords?: T; mipLevel?: number; arrayIndex?: number; + sampleIndex?: number; ddx?: T; ddy?: T; offset?: T; @@ -129,7 +236,8 @@ export interface TextureCallArgs { export interface TextureCall extends TextureCallArgs { builtin: 'textureSample' | 'textureLoad'; - coordType: 'f'; + coordType: 'f' | 'i' | 'u'; + levelType?: 'i' | 'u'; } function toArray(coords: Dimensionality): number[] { @@ -189,6 +297,46 @@ export interface Texture { viewDescriptor: GPUTextureViewDescriptor; } +/** + * Converts the src texel representation to an RGBA representation. + */ +function convertPerTexelComponentToResultFormat( + src: PerTexelComponent, + format: EncodableTextureFormat +): PerTexelComponent { + const rep = kTexelRepresentationInfo[format]; + const out: PerTexelComponent = { R: 0, G: 0, B: 0, A: 1 }; + for (const component of rep.componentOrder) { + switch (component) { + case 'Stencil': + case 'Depth': + out.R = src[component]; + break; + default: + assert(out[component] !== undefined); // checks that component = R, G, B or A + out[component] = src[component]; + } + } + return out; +} + +/** + * Convert RGBA result format to texel view format of src texture. + * Effectively this converts something like { R: 0.1, G: 0, B: 0, A: 1 } + * to { Depth: 0.1 } + */ +function convertResultFormatToTexelViewFormat( + src: PerTexelComponent, + format: EncodableTextureFormat +): PerTexelComponent { + const rep = kTexelRepresentationInfo[format]; + const out: PerTexelComponent = {}; + for (const component of rep.componentOrder) { + out[component] = src[component] ?? src.R; + } + return out; +} + /** * Returns the expect value for a WGSL builtin texture function for a single * mip level @@ -196,19 +344,20 @@ export interface Texture { export function softwareTextureReadMipLevel( call: TextureCall, texture: Texture, - sampler: GPUSamplerDescriptor, + sampler: GPUSamplerDescriptor | undefined, mipLevel: number ): PerTexelComponent { - const rep = kTexelRepresentationInfo[texture.texels[mipLevel].format]; + const { format } = texture.texels[mipLevel]; + const rep = kTexelRepresentationInfo[format]; const textureSize = virtualMipSize( texture.descriptor.dimension || '2d', texture.descriptor.size, mipLevel ); const addressMode = [ - sampler.addressModeU ?? 'clamp-to-edge', - sampler.addressModeV ?? 'clamp-to-edge', - sampler.addressModeW ?? 'clamp-to-edge', + sampler?.addressModeU ?? 'clamp-to-edge', + sampler?.addressModeV ?? 'clamp-to-edge', + sampler?.addressModeW ?? 'clamp-to-edge', ]; const load = (at: number[]) => @@ -249,7 +398,7 @@ export function softwareTextureReadMipLevel( const samples: { at: number[]; weight: number }[] = []; - const filter = sampler.minFilter; + const filter = sampler?.minFilter ?? 'nearest'; switch (filter) { case 'linear': { // 'p0' is the lower texel for 'at' @@ -361,10 +510,11 @@ export function softwareTextureReadMipLevel( } } - return out; + return convertPerTexelComponentToResultFormat(out, format); } case 'textureLoad': { - return load(toArray(call.coords!)); + const c = applyAddressModesToCoords(addressMode, textureSize, call.coords!); + return convertPerTexelComponentToResultFormat(load(c), format); } } } @@ -440,6 +590,149 @@ export type TextureTestOptions = { offset?: readonly [number, number]; // a constant offset }; +/** + * out of bounds is defined as any of the following being true + * + * * coords is outside the range [0, textureDimensions(t, level)) + * * array_index is outside the range [0, textureNumLayers(t)) + * * level is outside the range [0, textureNumLevels(t)) + * * sample_index is outside the range [0, textureNumSamples(s)) + */ +function isOutOfBoundsCall(texture: Texture, call: TextureCall) { + assert(call.mipLevel !== undefined); + assert(call.coords !== undefined); + assert(call.offset === undefined); + + const desc = reifyTextureDescriptor(texture.descriptor); + + const { coords, mipLevel, arrayIndex, sampleIndex } = call; + + if (mipLevel < 0 || mipLevel >= desc.mipLevelCount) { + return true; + } + + const size = virtualMipSize( + texture.descriptor.dimension || '2d', + texture.descriptor.size, + mipLevel + ); + + for (let i = 0; i < coords.length; ++i) { + const v = coords[i]; + if (v < 0 || v >= size[i]) { + return true; + } + } + + if (arrayIndex !== undefined) { + const size = reifyExtent3D(desc.size); + if (arrayIndex < 0 || arrayIndex >= size.depthOrArrayLayers) { + return true; + } + } + + if (sampleIndex !== undefined) { + if (sampleIndex < 0 || sampleIndex >= desc.sampleCount) { + return true; + } + } + + return false; +} + +/** + * For a texture builtin with no sampler (eg textureLoad), + * any out of bounds access is allowed to return one of: + * + * * the value of any texel in the texture + * * 0,0,0,0 or 0,0,0,1 if not a depth texture + * * 0 if a depth texture + */ +function okBecauseOutOfBounds( + texture: Texture, + call: TextureCall, + gotRGBA: PerTexelComponent, + maxFractionalDiff: number +) { + if (!isOutOfBoundsCall(texture, call)) { + return false; + } + + if (texture.descriptor.format.includes('depth')) { + if (gotRGBA.R === 0) { + return true; + } + } else { + if ( + gotRGBA.R === 0 && + gotRGBA.B === 0 && + gotRGBA.G === 0 && + (gotRGBA.A === 0 || gotRGBA.A === 1) + ) { + return true; + } + } + + for (let mipLevel = 0; mipLevel < texture.texels.length; ++mipLevel) { + const mipTexels = texture.texels[mipLevel]; + const size = virtualMipSize( + texture.descriptor.dimension || '2d', + texture.descriptor.size, + mipLevel + ); + for (let z = 0; z < size[2]; ++z) { + for (let y = 0; y < size[1]; ++y) { + for (let x = 0; x < size[0]; ++x) { + const texel = mipTexels.color({ x, y, z }); + const rgba = convertPerTexelComponentToResultFormat(texel, mipTexels.format); + if (texelsApproximatelyEqual(gotRGBA, rgba, mipTexels.format, maxFractionalDiff)) { + return true; + } + } + } + } + } + + return false; +} + +const kRGBAComponents = [ + TexelComponent.R, + TexelComponent.G, + TexelComponent.B, + TexelComponent.A, +] as const; + +const kRComponent = [TexelComponent.R] as const; + +function texelsApproximatelyEqual( + gotRGBA: PerTexelComponent, + expectRGBA: PerTexelComponent, + format: EncodableTextureFormat, + maxFractionalDiff: number +) { + const rep = kTexelRepresentationInfo[format]; + const got = convertResultFormatToTexelViewFormat(gotRGBA, format); + const expect = convertResultFormatToTexelViewFormat(expectRGBA, format); + const gULP = rep.bitsToULPFromZero(rep.numberToBits(got)); + const eULP = rep.bitsToULPFromZero(rep.numberToBits(expect)); + + const rgbaComponentsToCheck = isDepthOrStencilTextureFormat(format) + ? kRComponent + : kRGBAComponents; + + for (const component of rgbaComponentsToCheck) { + const g = gotRGBA[component]!; + const e = expectRGBA[component]!; + const absDiff = Math.abs(g - e); + const ulpDiff = Math.abs(gULP[component]! - eULP[component]!); + if (ulpDiff > 3 && absDiff > maxFractionalDiff) { + return false; + } + } + return true; +} + /** * Checks the result of each call matches the expected result. */ @@ -447,18 +740,36 @@ export async function checkCallResults( t: GPUTest, texture: Texture, textureType: string, - sampler: GPUSamplerDescriptor, + sampler: GPUSamplerDescriptor | undefined, calls: TextureCall[], results: PerTexelComponent[] ) { const errs: string[] = []; const rep = kTexelRepresentationInfo[texture.texels[0].format]; - const maxFractionalDiff = getMaxFractionalDiffForTextureFormat(texture.descriptor.format); + const maxFractionalDiff = + sampler?.minFilter === 'linear' || + sampler?.magFilter === 'linear' || + sampler?.mipmapFilter === 'linear' + ? getMaxFractionalDiffForTextureFormat(texture.descriptor.format) + : 0; + for (let callIdx = 0; callIdx < calls.length; callIdx++) { const call = calls[callIdx]; - const got = results[callIdx]; - const expect = softwareTextureReadMipLevel(call, texture, sampler, 0); + const gotRGBA = results[callIdx]; + const expectRGBA = softwareTextureReadMipLevel(call, texture, sampler, 0); + + if ( + texelsApproximatelyEqual(gotRGBA, expectRGBA, texture.texels[0].format, maxFractionalDiff) + ) { + continue; + } + + if (!sampler && okBecauseOutOfBounds(texture, call, gotRGBA, maxFractionalDiff)) { + continue; + } + const got = convertResultFormatToTexelViewFormat(gotRGBA, texture.texels[0].format); + const expect = convertResultFormatToTexelViewFormat(expectRGBA, texture.texels[0].format); const gULP = rep.bitsToULPFromZero(rep.numberToBits(got)); const eULP = rep.bitsToULPFromZero(rep.numberToBits(expect)); for (const component of rep.componentOrder) { @@ -477,40 +788,42 @@ export async function checkCallResults( abs diff: ${absDiff.toFixed(4)} rel diff: ${(relDiff * 100).toFixed(2)}% ulp diff: ${ulpDiff} - sample points: `); - const expectedSamplePoints = [ - 'expected:', - ...(await identifySamplePoints(texture, (texels: TexelView) => { - return Promise.resolve( - softwareTextureReadMipLevel( - call, - { - texels: [texels], - descriptor: texture.descriptor, - viewDescriptor: texture.viewDescriptor, - }, - sampler, - 0 - ) - ); - })), - ]; - const gotSamplePoints = [ - 'got:', - ...(await identifySamplePoints(texture, async (texels: TexelView) => { - const gpuTexture = createTextureFromTexelViews(t, [texels], texture.descriptor); - const result = ( - await doTextureCalls(t, gpuTexture, texture.viewDescriptor, textureType, sampler, [ - call, - ]) - )[0]; - gpuTexture.destroy(); - return result; - })), - ]; - errs.push(layoutTwoColumns(expectedSamplePoints, gotSamplePoints).join('\n')); - errs.push('', ''); + if (sampler) { + const expectedSamplePoints = [ + 'expected:', + ...(await identifySamplePoints(texture, (texels: TexelView) => { + return Promise.resolve( + softwareTextureReadMipLevel( + call, + { + texels: [texels], + descriptor: texture.descriptor, + viewDescriptor: texture.viewDescriptor, + }, + sampler, + 0 + ) + ); + })), + ]; + const gotSamplePoints = [ + 'got:', + ...(await identifySamplePoints(texture, async (texels: TexelView) => { + const gpuTexture = createTextureFromTexelViews(t, [texels], texture.descriptor); + const result = ( + await doTextureCalls(t, gpuTexture, texture.viewDescriptor, textureType, sampler, [ + call, + ]) + )[0]; + gpuTexture.destroy(); + return result; + })), + ]; + errs.push(' sample points:'); + errs.push(layoutTwoColumns(expectedSamplePoints, gotSamplePoints).join('\n')); + errs.push('', ''); + } } } } @@ -739,7 +1052,8 @@ function getMaxFractionalDiffForTextureFormat(format: GPUTextureFormat) { } else if (format.endsWith('float')) { return 44; } else { - unreachable(); + // It's likely an integer format. In any case, zero tolerance is passable. + return 0; } } @@ -856,6 +1170,7 @@ function getBlockFiller(format: GPUTextureFormat) { * Fills a texture with random data. */ export function fillTextureWithRandomData(device: GPUDevice, texture: GPUTexture) { + assert(!isCompressedFloatTextureFormat(texture.format)); const info = kTextureFormatInfo[texture.format]; const hashBase = sumOfCharCodesOfString(texture.format) + @@ -1299,6 +1614,20 @@ function layoutTwoColumns(columnA: string[], columnB: string[]) { return out; } +function getDepthOrArrayLayersForViewDimension(viewDimension?: GPUTextureViewDimension) { + switch (viewDimension) { + case undefined: + case '2d': + return 1; + case '3d': + return 8; + case 'cube': + return 6; + default: + unreachable(); + } +} + /** * Choose a texture size based on the given parameters. * The size will be in a multiple of blocks. If it's a cube @@ -1320,9 +1649,10 @@ export function chooseTextureSize({ const height = align(Math.max(minSize, blockHeight * minBlocks), blockHeight); if (viewDimension === 'cube') { const size = lcm(width, height); - return [size, size]; + return [size, size, 6]; } - return [width, height]; + const depthOrArrayLayers = getDepthOrArrayLayersForViewDimension(viewDimension); + return [width, height, depthOrArrayLayers]; } export const kSamplePointMethods = ['texel-centre', 'spiral'] as const; @@ -1815,6 +2145,22 @@ function wgslExpr(data: number | vec1 | vec2 | vec3 | vec4): string { return data.toString(); } +function wgslExprFor(data: number | vec1 | vec2 | vec3 | vec4, type: 'f' | 'i' | 'u'): string { + if (Array.isArray(data)) { + switch (data.length) { + case 1: + return `${type}(${data[0].toString()})`; + case 2: + return `vec2${type}(${data.map(v => v.toString()).join(', ')})`; + case 3: + return `vec3${type}(${data.map(v => v.toString()).join(', ')})`; + default: + unreachable(); + } + } + return `${type}32(${data.toString()})`; +} + function binKey(call: TextureCall): string { const keys: string[] = []; for (const name of kTextureCallArgNames) { @@ -1848,8 +2194,9 @@ function buildBinnedCalls(calls: TextureCall[]) { if (name === 'offset') { args.push(`/* offset */ ${wgslExpr(value)}`); } else { + const type = name === 'mipLevel' ? prototype.levelType! : prototype.coordType; args.push(`args.${name}`); - fields.push(`@align(16) ${name} : ${wgslTypeFor(value, prototype.coordType)}`); + fields.push(`@align(16) ${name} : ${wgslTypeFor(value, type)}`); } } } @@ -1912,7 +2259,13 @@ export function describeTextureCall(call: TextureCall< for (const name of kTextureCallArgNames) { const value = call[name]; if (value !== undefined) { - args.push(`${name}: ${wgslExpr(value)}`); + if (name === 'coords') { + args.push(`${name}: ${wgslExprFor(value, call.coordType)}`); + } else if (name === 'mipLevel') { + args.push(`${name}: ${wgslExprFor(value, call.levelType!)}`); + } else { + args.push(`${name}: ${wgslExpr(value)}`); + } } } return `${call.builtin}(${args.join(', ')})`; @@ -1937,7 +2290,7 @@ export async function doTextureCalls( gpuTexture: GPUTexture, viewDescriptor: GPUTextureViewDescriptor, textureType: string, - sampler: GPUSamplerDescriptor, + sampler: GPUSamplerDescriptor | undefined, calls: TextureCall[] ) { let structs = ''; @@ -1972,9 +2325,11 @@ export async function doTextureCalls( }); t.device.queue.writeBuffer(dataBuffer, 0, new Uint32Array(data)); + const { resultType, resultFormat } = getTextureFormatTypeInfo(gpuTexture.format); + const rtWidth = 256; const renderTarget = t.createTextureTracked({ - format: 'rgba32float', + format: resultFormat, size: { width: rtWidth, height: Math.ceil(calls.length / rtWidth) }, usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.RENDER_ATTACHMENT, }); @@ -1996,13 +2351,13 @@ fn vs_main(@builtin(vertex_index) vertex_index : u32) -> @builtin(position) vec4 } @group(0) @binding(0) var T : ${textureType}; -@group(0) @binding(1) var S : sampler; +${sampler ? '@group(0) @binding(1) var S : sampler' : ''}; @group(0) @binding(2) var data : Data; @fragment -fn fs_main(@builtin(position) frag_pos : vec4f) -> @location(0) vec4f { +fn fs_main(@builtin(position) frag_pos : vec4f) -> @location(0) ${resultType} { let frag_idx = u32(frag_pos.x) + u32(frag_pos.y) * ${renderTarget.width}; - var result : vec4f; + var result : ${resultType}; ${body} return result; } @@ -2028,13 +2383,13 @@ ${body} pipelines.set(code, pipeline); } - const gpuSampler = t.device.createSampler(sampler); + const gpuSampler = sampler ? t.device.createSampler(sampler) : undefined; const bindGroup = t.device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [ { binding: 0, resource: gpuTexture.createView(viewDescriptor) }, - { binding: 1, resource: gpuSampler }, + ...(sampler ? [{ binding: 1, resource: gpuSampler! }] : []), { binding: 2, resource: { buffer: dataBuffer } }, ], }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/utils.ts b/src/webgpu/shader/execution/expression/call/builtin/utils.ts index 9cbee0093926..a13e22c0a81b 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/utils.ts @@ -1,11 +1,34 @@ +import { assert, unreachable } from '../../../../../../common/util/util.js'; +import { virtualMipSize } from '../../../../../util/texture/base.js'; + +/* Valid types of Boundaries */ +export type Boundary = + | 'in-bounds' + | 'x-min-wrap' + | 'x-min-boundary' + | 'x-max-wrap' + | 'x-max-boundary' + | 'y-min-wrap' + | 'y-min-boundary' + | 'y-max-wrap' + | 'y-max-boundary' + | 'z-min-wrap' + | 'z-min-boundary' + | 'z-max-wrap' + | 'z-max-boundary'; + +export function isBoundaryNegative(boundary: Boundary) { + return boundary.endsWith('min-wrap'); +} + /** * Generates the boundary entries for the given number of dimensions * * @param numDimensions: The number of dimensions to generate for * @returns an array of generated coord boundaries */ -export function generateCoordBoundaries(numDimensions: number) { - const ret = ['in-bounds']; +export function generateCoordBoundaries(numDimensions: number): Boundary[] { + const ret: Boundary[] = ['in-bounds']; if (numDimensions < 1 || numDimensions > 3) { throw new Error(`invalid numDimensions: ${numDimensions}`); @@ -15,7 +38,7 @@ export function generateCoordBoundaries(numDimensions: number) { for (let i = 0; i < numDimensions; ++i) { for (const j of ['min', 'max']) { for (const k of ['wrap', 'boundary']) { - ret.push(`${name[i]}-${j}-${k}`); + ret.push(`${name[i]}-${j}-${k}` as Boundary); } } } @@ -23,18 +46,91 @@ export function generateCoordBoundaries(numDimensions: number) { return ret; } +export type LevelSpec = -1 | 0 | 'numLevels-1' | 'numLevels'; + +export function getMipLevelFromLevelSpec(mipLevelCount: number, levelSpec: LevelSpec): number { + switch (levelSpec) { + case -1: + return -1; + case 0: + return 0; + case 'numLevels': + return mipLevelCount; + case 'numLevels-1': + return mipLevelCount - 1; + default: + unreachable(); + } +} + +export function isLevelSpecNegative(levelSpec: LevelSpec) { + return levelSpec === -1; +} + +function getCoordForSize(size: [number, number, number], boundary: Boundary) { + const coord = size.map(v => Math.floor(v / 2)); + switch (boundary) { + case 'in-bounds': + break; + default: { + const axis = boundary[0]; + const axisIndex = axis.charCodeAt(0) - 'x'.charCodeAt(0); + const axisSize = size[axisIndex]; + const location = boundary.substring(2); + let v = 0; + switch (location) { + case 'min-wrap': + v = -1; + break; + case 'min-boundary': + v = 0; + break; + case 'max-wrap': + v = axisSize; + break; + case 'max-boundary': + v = axisSize - 1; + break; + default: + unreachable(); + } + coord[axisIndex] = v; + } + } + return coord; +} + +function getNumDimensions(dimension: GPUTextureDimension) { + switch (dimension) { + case '1d': + return 1; + case '2d': + return 2; + case '3d': + return 3; + } +} + +export function getCoordinateForBoundaries( + texture: GPUTexture, + mipLevel: number, + boundary: Boundary +) { + const size = virtualMipSize(texture.dimension, texture, mipLevel); + const coord = getCoordForSize(size, boundary); + return coord.slice(0, getNumDimensions(texture.dimension)) as T; +} + /** - * Generates a set of offset values to attempt in the range [-9, 8]. + * Generates a set of offset values to attempt in the range [-8, 7]. * * @param numDimensions: The number of dimensions to generate for * @return an array of generated offset values */ export function generateOffsets(numDimensions: number) { - if (numDimensions < 2 || numDimensions > 3) { - throw new Error(`generateOffsets: invalid numDimensions: ${numDimensions}`); - } + assert(numDimensions >= 2 && numDimensions <= 3); const ret: Array> = [undefined]; - for (const val of [-9, -8, 0, 1, 7, 8]) { + for (const val of [-8, 0, 1, 7]) { const v = []; for (let i = 0; i < numDimensions; ++i) { v.push(val); diff --git a/src/webgpu/shader/execution/expression/expression.ts b/src/webgpu/shader/execution/expression/expression.ts index a609b4b73df4..f486723baf56 100644 --- a/src/webgpu/shader/execution/expression/expression.ts +++ b/src/webgpu/shader/execution/expression/expression.ts @@ -867,7 +867,7 @@ ${body} // Runtime eval ////////////////////////////////////////////////////////////////////////// let operation = ''; - if (inputSource === 'storage_rw') { + if (inputSource === 'storage_rw' && objectEquals(resultType, storageType(resultType))) { operation = ` outputs[i].value = ${storageType(resultType)}(inputs[i].lhs); outputs[i].value ${op} ${rhsType}(inputs[i].rhs);`; diff --git a/src/webgpu/shader/execution/expression/unary/f32_conversion.cache.ts b/src/webgpu/shader/execution/expression/unary/f32_conversion.cache.ts index f61435f07ce4..8ce39eccfd1b 100644 --- a/src/webgpu/shader/execution/expression/unary/f32_conversion.cache.ts +++ b/src/webgpu/shader/execution/expression/unary/f32_conversion.cache.ts @@ -1,7 +1,8 @@ -import { bool, f16, f32, i32, u32 } from '../../../../util/conversion.js'; -import { FP } from '../../../../util/floating_point.js'; +import { abstractInt, bool, f16, f32, i32, u32 } from '../../../../util/conversion.js'; +import { FP, FPInterval } from '../../../../util/floating_point.js'; import { fullI32Range, + fullI64Range, fullU32Range, scalarF16Range, scalarF32Range, @@ -10,6 +11,12 @@ import { } from '../../../../util/math.js'; import { makeCaseCache } from '../case_cache.js'; +const f32FiniteRangeInterval = new FPInterval( + 'f32', + FP.f32.constants().negative.min, + FP.f32.constants().positive.max +); + // Cases: f32_matCxR_[non_]const const f32_mat_cases = ([2, 3, 4] as const) .flatMap(cols => @@ -46,6 +53,23 @@ const f16_mat_cases = ([2, 3, 4] as const) ) .reduce((a, b) => ({ ...a, ...b }), {}); +// Cases: abstract_float_matCxR +// Note that abstract float values may be not exactly representable in f32 +// and/or out of range. +const abstract_float_mat_cases = ([2, 3, 4] as const) + .flatMap(cols => + ([2, 3, 4] as const).map(rows => ({ + [`abstract_float_mat${cols}x${rows}`]: () => { + return FP.abstract.generateMatrixToMatrixCases( + FP.abstract.sparseMatrixRange(cols, rows), + 'finite', + FP.f32.correctlyRoundedMatrix + ); + }, + })) + ) + .reduce((a, b) => ({ ...a, ...b }), {}); + export const d = makeCaseCache('unary/f32_conversion', { bool: () => { return [ @@ -63,6 +87,13 @@ export const d = makeCaseCache('unary/f32_conversion', { return { input: i32(i), expected: FP.f32.correctlyRoundedInterval(i) }; }); }, + abstract_int: () => { + return [...fullI64Range()] + .filter(v => f32FiniteRangeInterval.contains(Number(v))) + .map(i => { + return { input: abstractInt(i), expected: FP.f32.correctlyRoundedInterval(Number(i)) }; + }); + }, f32: () => { return scalarF32Range().map(f => { return { input: f32(f), expected: FP.f32.correctlyRoundedInterval(f) }; @@ -74,6 +105,15 @@ export const d = makeCaseCache('unary/f32_conversion', { return { input: f16(f), expected: FP.f32.correctlyRoundedInterval(f) }; }); }, + // Note that abstract float values may be not exactly representable in f32. + abstract_float: () => { + return FP.abstract.generateScalarToIntervalCases( + [...FP.abstract.scalarRange()], + 'finite', + FP.f32.correctlyRoundedInterval + ); + }, ...f32_mat_cases, ...f16_mat_cases, + ...abstract_float_mat_cases, }); diff --git a/src/webgpu/shader/execution/expression/unary/f32_conversion.spec.ts b/src/webgpu/shader/execution/expression/unary/f32_conversion.spec.ts index 464fdee44e79..7a42334cb860 100644 --- a/src/webgpu/shader/execution/expression/unary/f32_conversion.spec.ts +++ b/src/webgpu/shader/execution/expression/unary/f32_conversion.spec.ts @@ -5,7 +5,7 @@ Execution Tests for the f32 conversion operations import { makeTestGroup } from '../../../../../common/framework/test_group.js'; import { GPUTest } from '../../../../gpu_test.js'; import { Type } from '../../../../util/conversion.js'; -import { ShaderBuilder, allInputSources, run } from '../expression.js'; +import { ShaderBuilder, allInputSources, run, onlyConstInputSource } from '../expression.js'; import { d } from './f32_conversion.cache.js'; import { unary } from './unary.js'; @@ -73,6 +73,32 @@ Converted to f32 await run(t, vectorizeToExpression(t.params.vectorize), [Type.i32], Type.f32, t.params, cases); }); +g.test('abstract_int') + .specURL('https://www.w3.org/TR/WGSL/#value-constructor-builtin-function') + .desc( + ` +f32(e), where e is an AbstractInt + +Converted to f32, +/-Inf if out of range +` + ) + .params(u => + u + .combine('inputSource', onlyConstInputSource) + .combine('vectorize', [undefined, 2, 3, 4] as const) + ) + .fn(async t => { + const cases = await d.get('abstract_int'); + await run( + t, + vectorizeToExpression(t.params.vectorize), + [Type.abstractInt], + Type.f32, + t.params, + cases + ); + }); + g.test('f32') .specURL('https://www.w3.org/TR/WGSL/#value-constructor-builtin-function') .desc( @@ -166,3 +192,52 @@ g.test('f16_mat') cases ); }); + +g.test('abstract_float') + .specURL('https://www.w3.org/TR/WGSL/#value-constructor-builtin-function') + .desc( + ` +f32(e), where e is an AbstractFloat + +Correctly rounded to f32 +` + ) + .params(u => + u + .combine('inputSource', onlyConstInputSource) + .combine('vectorize', [undefined, 2, 3, 4] as const) + ) + .fn(async t => { + const cases = await d.get('abstract_float'); + await run( + t, + vectorizeToExpression(t.params.vectorize), + [Type.abstractFloat], + Type.f32, + t.params, + cases + ); + }); + +g.test('abstract_float_mat') + .specURL('https://www.w3.org/TR/WGSL/#matrix-builtin-functions') + .desc(`AbstractFloat matrix to f32 matrix tests`) + .params(u => + u + .combine('inputSource', onlyConstInputSource) + .combine('cols', [2, 3, 4] as const) + .combine('rows', [2, 3, 4] as const) + ) + .fn(async t => { + const cols = t.params.cols; + const rows = t.params.rows; + const cases = await d.get(`abstract_float_mat${cols}x${rows}`); + await run( + t, + matrixExperession(cols, rows), + [Type.mat(cols, rows, Type.abstractFloat)], + Type.mat(cols, rows, Type.f32), + t.params, + cases + ); + }); diff --git a/src/webgpu/shader/types.ts b/src/webgpu/shader/types.ts index 76b094310d79..a5a385deb289 100644 --- a/src/webgpu/shader/types.ts +++ b/src/webgpu/shader/types.ts @@ -195,7 +195,7 @@ export const TexelFormats = [ { format: 'rg32uint', _shaderType: 'u32' }, { format: 'rg32sint', _shaderType: 'i32' }, { format: 'rg32float', _shaderType: 'f32' }, - { format: 'rgba32uint', _shaderType: 'i32' }, + { format: 'rgba32uint', _shaderType: 'u32' }, { format: 'rgba32sint', _shaderType: 'i32' }, { format: 'rgba32float', _shaderType: 'f32' }, ] as const; diff --git a/src/webgpu/shader/validation/decl/context_dependent_resolution.spec.ts b/src/webgpu/shader/validation/decl/context_dependent_resolution.spec.ts index b6bdfc4fecab..2d51a7c6bfb6 100644 --- a/src/webgpu/shader/validation/decl/context_dependent_resolution.spec.ts +++ b/src/webgpu/shader/validation/decl/context_dependent_resolution.spec.ts @@ -28,7 +28,7 @@ const kAttributeCases = { // diagnostic is a keyword group: `@group(0) @binding(0) var s : sampler;`, id: `@id(1) override x : i32;`, - interpolate: `@fragment fn main(@location(0) @interpolate(flat) x : i32) { }`, + interpolate: `@fragment fn main(@location(0) @interpolate(flat, either) x : i32) { }`, invariant: `@fragment fn main(@builtin(position) @invariant pos : vec4f) { }`, location: `@fragment fn main(@location(0) x : f32) { }`, must_use: `@must_use fn foo() -> u32 { return 0; }`, @@ -102,6 +102,17 @@ g.test('builtin_value_names') .beginSubcases() .combine('decl', ['override', 'const', 'var'] as const) ) + .beforeAllSubcases(t => { + const wgsl = kBuiltinCases[t.params.case]; + t.skipIf( + t.isCompatibility && wgsl.includes('sample_mask'), + 'sample_mask is not supported in compatibility mode' + ); + t.skipIf( + t.isCompatibility && wgsl.includes('sample_index'), + 'sample_index is not supported in compatibility mode' + ); + }) .fn(t => { const code = ` ${t.params.decl} ${t.params.case} : u32 = 0; @@ -303,10 +314,20 @@ g.test('interpolation_type_names') .beginSubcases() .combine('decl', ['override', 'const', 'var'] as const) ) + .beforeAllSubcases(t => { + t.skipIf( + t.isCompatibility && t.params.case === 'linear', + 'compatibility mode does not support linear interpolation type' + ); + }) .fn(t => { + const attr = + t.isCompatibility && t.params.case === 'flat' + ? `@interpolate(flat, either)` + : `@interpolate(${t.params.case})`; const code = ` ${t.params.decl} ${t.params.case} : u32 = 0; - @fragment fn main(@location(0) @interpolate(${t.params.case}) x : f32) { } + @fragment fn main(@location(0) ${attr} x : f32) { } fn use_var() -> u32 { return ${t.params.case}; } @@ -325,6 +346,12 @@ g.test('interpolation_sampling_names') .beginSubcases() .combine('decl', ['override', 'const', 'var'] as const) ) + .beforeAllSubcases(t => { + t.skipIf( + t.isCompatibility && t.params.case === 'sample', + 'compatibility mode does not support sample sampling' + ); + }) .fn(t => { const code = ` ${t.params.decl} ${t.params.case} : u32 = 0; @@ -347,6 +374,12 @@ g.test('interpolation_flat_names') .beginSubcases() .combine('decl', ['override', 'const', 'var'] as const) ) + .beforeAllSubcases(t => { + t.skipIf( + t.isCompatibility && t.params.case === 'first', + 'compatibility mode does not support first sampling' + ); + }) .fn(t => { const code = ` ${t.params.decl} ${t.params.case} : u32 = 0; diff --git a/src/webgpu/shader/validation/extension/dual_source_blending.spec.ts b/src/webgpu/shader/validation/extension/dual_source_blending.spec.ts index a8122067462f..a7aacef21d80 100644 --- a/src/webgpu/shader/validation/extension/dual_source_blending.spec.ts +++ b/src/webgpu/shader/validation/extension/dual_source_blending.spec.ts @@ -560,3 +560,64 @@ ${ `; t.expectCompileResult(kUsageValidationTests[t.params.attr].pass, code); }); + +const kValidLocationTypes = [ + 'f16', + 'f32', + 'i32', + 'u32', + 'vec2h', + 'vec2f', + 'vec2i', + 'vec2u', + 'vec3h', + 'vec3f', + 'vec3i', + 'vec3u', + 'vec4h', + 'vec4f', + 'vec4i', + 'vec4u', +] as const; + +const kF16TypesSet = new Set(['f16', 'vec2h', 'vec3h', 'vec4h']); + +g.test('blend_src_same_type') + .desc(`Test that the struct member with @blend_src(0) and @blend_src(1) must have same type.`) + .params(u => + u.combine('blendSrc0Type', kValidLocationTypes).combine('blendSrc1Type', kValidLocationTypes) + ) + .beforeAllSubcases(t => { + const requiredFeatures: GPUFeatureName[] = ['dual-source-blending']; + const needF16Extension = + kF16TypesSet.has(t.params.blendSrc0Type) || kF16TypesSet.has(t.params.blendSrc1Type); + if (needF16Extension) { + requiredFeatures.push('shader-f16'); + } + t.selectDeviceOrSkipTestCase({ requiredFeatures }); + }) + .fn(t => { + const { blendSrc0Type, blendSrc1Type } = t.params; + + const needF16Extension = kF16TypesSet.has(blendSrc0Type) || kF16TypesSet.has(blendSrc1Type); + const code = ` +enable dual_source_blending; + +${needF16Extension ? 'enable f16;' : ''} + +struct BlendSrcOutput { + @location(0) @blend_src(0) color : ${blendSrc0Type}, + @location(0) @blend_src(1) blend : ${blendSrc1Type}, +} + +@fragment fn main() -> BlendSrcOutput { + var output : BlendSrcOutput; + output.color = ${blendSrc0Type}(); + output.blend = ${blendSrc1Type}(); + return output; +} +`; + + const success = blendSrc0Type === blendSrc1Type; + t.expectCompileResult(success, code); + }); diff --git a/src/webgpu/util/texture/texel_data.ts b/src/webgpu/util/texture/texel_data.ts index 0555ac5920d8..4c88d9c2182a 100644 --- a/src/webgpu/util/texture/texel_data.ts +++ b/src/webgpu/util/texture/texel_data.ts @@ -78,12 +78,15 @@ function makePerTexelComponent(components: TexelComponent[], value: T): PerTe * @returns {ComponentMapFn} The map function which clones the input component values, and applies * `fn` to each of component of `components`. */ -function applyEach(fn: (value: number) => number, components: TexelComponent[]): ComponentMapFn { +function applyEach( + fn: (value: number, component: TexelComponent) => number, + components: TexelComponent[] +): ComponentMapFn { return (values: PerTexelComponent) => { values = Object.assign({}, values); for (const c of components) { assert(values[c] !== undefined); - values[c] = fn(values[c]!); + values[c] = fn(values[c]!, c); } return values; }; @@ -122,7 +125,13 @@ const decodeSRGB: ComponentMapFn = components => { export function makeClampToRange(format: EncodableTextureFormat): ComponentMapFn { const repr = kTexelRepresentationInfo[format]; assert(repr.numericRange !== null, 'Format has unknown numericRange'); - return applyEach(x => clamp(x, repr.numericRange!), repr.componentOrder); + const perComponentRanges = repr.numericRange as PerComponentNumericRange; + const range = repr.numericRange as NumericRange; + + return applyEach((x, component) => { + const perComponentRange = perComponentRanges[component]; + return clamp(x, perComponentRange ? perComponentRange : range); + }, repr.componentOrder); } // MAINTENANCE_TODO: Look into exposing this map to the test fixture so that it can be GCed at the @@ -601,6 +610,23 @@ const kFloat11Format = { signed: 0, exponentBits: 5, mantissaBits: 6, bias: 15 } const kFloat10Format = { signed: 0, exponentBits: 5, mantissaBits: 5, bias: 15 } as const; export type PerComponentFiniteMax = Record; +export type NumericRange = { + min: number; + max: number; + finiteMin: number; + finiteMax: number | PerComponentFiniteMax; +}; +export type PerComponentNumericRange = Partial< + Record< + TexelComponent, + { + min: number; + max: number; + finiteMin: number; + finiteMax: number; + } + > +>; export type TexelRepresentationInfo = { /** Order of components in the packed representation. */ readonly componentOrder: TexelComponent[]; @@ -628,15 +654,11 @@ export type TexelRepresentationInfo = { /** Convert integer bit representations into ULPs-from-zero, e.g. unorm8 255 -> 255 ULPs */ readonly bitsToULPFromZero: ComponentMapFn; /** The valid range of numeric "color" values, e.g. [0, Infinity] for ufloat. */ - readonly numericRange: null | { - min: number; - max: number; - finiteMin: number; - finiteMax: number | PerComponentFiniteMax; - }; + readonly numericRange: null | NumericRange | PerComponentNumericRange; // Add fields as needed }; + export const kTexelRepresentationInfo: { readonly [k in UncompressedTextureFormat]: TexelRepresentationInfo; } = { @@ -726,7 +748,12 @@ export const kTexelRepresentationInfo: { return components; }, bitsToULPFromZero: components => components, - numericRange: null, + numericRange: { + R: { min: 0, max: 0x3ff, finiteMin: 0, finiteMax: 0x3ff }, + G: { min: 0, max: 0x3ff, finiteMin: 0, finiteMax: 0x3ff }, + B: { min: 0, max: 0x3ff, finiteMin: 0, finiteMax: 0x3ff }, + A: { min: 0, max: 0x3, finiteMin: 0, finiteMax: 0x3 }, + }, }, rgb10a2unorm: { componentOrder: kRGBA,