diff --git a/.gitignore b/.gitignore index 09eb42eabcd7..1a359f82b041 100644 --- a/.gitignore +++ b/.gitignore @@ -11,6 +11,9 @@ *.tmp.txt /docs/tsdoc/ +# vim swap files +*.swp* + # `pnpm` lockfile pnpm-lock.yaml diff --git a/src/common/util/util.ts b/src/common/util/util.ts index 540a29f77ec1..4e19f9e32067 100644 --- a/src/common/util/util.ts +++ b/src/common/util/util.ts @@ -373,6 +373,7 @@ interface TypedArrayMap { Uint8ClampedArray: Uint8ClampedArray; Int32Array: Int32Array; Uint32Array: Uint32Array; + Float16Array: Float16Array; Float32Array: Float32Array; Float64Array: Float64Array; BigInt64Array: BigInt64Array; diff --git a/src/resources/cache/hashes.json b/src/resources/cache/hashes.json index 8c678b1a15ef..e62489909fad 100644 --- a/src/resources/cache/hashes.json +++ b/src/resources/cache/hashes.json @@ -1,112 +1,112 @@ { - "webgpu/shader/execution/binary/af_addition.bin": "6e41f81e", - "webgpu/shader/execution/binary/af_logical.bin": "254d66ec", - "webgpu/shader/execution/binary/af_division.bin": "e15d545b", - "webgpu/shader/execution/binary/af_matrix_addition.bin": "4bfc497f", - "webgpu/shader/execution/binary/af_matrix_subtraction.bin": "2bbdd5e6", - "webgpu/shader/execution/binary/af_multiplication.bin": "23c7c8ad", - "webgpu/shader/execution/binary/af_remainder.bin": "9b84ecb1", - "webgpu/shader/execution/binary/af_subtraction.bin": "bd4b2ec9", - "webgpu/shader/execution/binary/f16_addition.bin": "5fe1fb4e", - "webgpu/shader/execution/binary/f16_logical.bin": "225d12ee", - "webgpu/shader/execution/binary/f16_division.bin": "dffabf5f", - "webgpu/shader/execution/binary/f16_matrix_addition.bin": "acc779ee", - "webgpu/shader/execution/binary/f16_matrix_matrix_multiplication.bin": "dcfbfff9", - "webgpu/shader/execution/binary/f16_matrix_scalar_multiplication.bin": "4b73a544", - "webgpu/shader/execution/binary/f16_matrix_subtraction.bin": "bfe9a25d", - "webgpu/shader/execution/binary/f16_matrix_vector_multiplication.bin": "a2222aec", - "webgpu/shader/execution/binary/f16_multiplication.bin": "5d5265c9", - "webgpu/shader/execution/binary/f16_remainder.bin": "5284a25", - "webgpu/shader/execution/binary/f16_subtraction.bin": "b98fc99f", - "webgpu/shader/execution/binary/f32_addition.bin": "13d7f835", - "webgpu/shader/execution/binary/f32_logical.bin": "c1f1b636", - "webgpu/shader/execution/binary/f32_division.bin": "80fee721", - "webgpu/shader/execution/binary/f32_matrix_addition.bin": "a0d20c36", - "webgpu/shader/execution/binary/f32_matrix_matrix_multiplication.bin": "3cabc829", - "webgpu/shader/execution/binary/f32_matrix_scalar_multiplication.bin": "aed8659", - "webgpu/shader/execution/binary/f32_matrix_subtraction.bin": "6063545d", - "webgpu/shader/execution/binary/f32_matrix_vector_multiplication.bin": "e36324e4", - "webgpu/shader/execution/binary/f32_multiplication.bin": "40527d93", - "webgpu/shader/execution/binary/f32_remainder.bin": "d8738967", - "webgpu/shader/execution/binary/f32_subtraction.bin": "f5453a8", - "webgpu/shader/execution/binary/i32_arithmetic.bin": "37440f39", - "webgpu/shader/execution/binary/i32_comparison.bin": "9679ab05", - "webgpu/shader/execution/binary/u32_arithmetic.bin": "c9fe0078", - "webgpu/shader/execution/binary/u32_comparison.bin": "7d0ba49", - "webgpu/shader/execution/abs.bin": "320b5a02", - "webgpu/shader/execution/acos.bin": "bb357762", - "webgpu/shader/execution/acosh.bin": "1975ce7f", - "webgpu/shader/execution/asin.bin": "74096576", - "webgpu/shader/execution/asinh.bin": "cdb5dae3", - "webgpu/shader/execution/atan.bin": "4d8cd242", - "webgpu/shader/execution/atan2.bin": "213c0112", - "webgpu/shader/execution/atanh.bin": "ca3cda5b", - "webgpu/shader/execution/bitcast.bin": "2b639d22", - "webgpu/shader/execution/ceil.bin": "9d91caf6", - "webgpu/shader/execution/clamp.bin": "d7c7512f", - "webgpu/shader/execution/cos.bin": "832af30d", - "webgpu/shader/execution/cosh.bin": "1e9d346f", - "webgpu/shader/execution/cross.bin": "8db02194", - "webgpu/shader/execution/degrees.bin": "4e6926ca", - "webgpu/shader/execution/determinant.bin": "46c05eb0", - "webgpu/shader/execution/distance.bin": "d3d25ee0", - "webgpu/shader/execution/dot.bin": "aac0064d", - "webgpu/shader/execution/exp.bin": "4edc9995", - "webgpu/shader/execution/exp2.bin": "71f9b9f1", - "webgpu/shader/execution/faceForward.bin": "bbfa6f58", - "webgpu/shader/execution/floor.bin": "46a8f6f6", - "webgpu/shader/execution/fma.bin": "d8de76da", - "webgpu/shader/execution/fract.bin": "ce59af66", - "webgpu/shader/execution/frexp.bin": "2d856679", - "webgpu/shader/execution/inverseSqrt.bin": "c2e19ec5", - "webgpu/shader/execution/ldexp.bin": "85a3d6da", - "webgpu/shader/execution/length.bin": "6462de73", - "webgpu/shader/execution/log.bin": "9d3e5eac", - "webgpu/shader/execution/log2.bin": "f0e6911", - "webgpu/shader/execution/max.bin": "10d3b0f0", - "webgpu/shader/execution/min.bin": "8150d732", - "webgpu/shader/execution/mix.bin": "e80a07f6", - "webgpu/shader/execution/modf.bin": "4ec56dfa", - "webgpu/shader/execution/normalize.bin": "10f90ae0", - "webgpu/shader/execution/pack2x16float.bin": "f238e86f", - "webgpu/shader/execution/pow.bin": "342665d2", - "webgpu/shader/execution/quantizeToF16.bin": "e2736816", - "webgpu/shader/execution/radians.bin": "6520dcc1", - "webgpu/shader/execution/reflect.bin": "4183175", - "webgpu/shader/execution/refract.bin": "4915dbfe", - "webgpu/shader/execution/round.bin": "99b75791", - "webgpu/shader/execution/saturate.bin": "a14f35f0", - "webgpu/shader/execution/sign.bin": "36761bf9", - "webgpu/shader/execution/sin.bin": "74a2057b", - "webgpu/shader/execution/sinh.bin": "8353f6ea", - "webgpu/shader/execution/smoothstep.bin": "8122e6fb", - "webgpu/shader/execution/sqrt.bin": "4b04e554", - "webgpu/shader/execution/step.bin": "11ec30ed", - "webgpu/shader/execution/tan.bin": "67da16db", - "webgpu/shader/execution/tanh.bin": "2f62fb36", - "webgpu/shader/execution/transpose.bin": "29a099da", - "webgpu/shader/execution/trunc.bin": "ccf08bdc", - "webgpu/shader/execution/unpack2x16float.bin": "7bfbbdf6", - "webgpu/shader/execution/unpack2x16snorm.bin": "bed6a1e7", - "webgpu/shader/execution/unpack2x16unorm.bin": "a26a06f0", - "webgpu/shader/execution/unpack4x8snorm.bin": "756c46a0", - "webgpu/shader/execution/unpack4x8unorm.bin": "575c7feb", - "webgpu/shader/execution/unary/af_arithmetic.bin": "1176909d", - "webgpu/shader/execution/unary/af_assignment.bin": "73927f22", - "webgpu/shader/execution/unary/bool_conversion.bin": "e96f1762", - "webgpu/shader/execution/unary/f16_arithmetic.bin": "e77cc4e8", - "webgpu/shader/execution/unary/f16_conversion.bin": "4555bbc2", - "webgpu/shader/execution/unary/f32_arithmetic.bin": "8e28dc17", - "webgpu/shader/execution/unary/f32_conversion.bin": "1d219", - "webgpu/shader/execution/unary/i32_arithmetic.bin": "ca6a70d", - "webgpu/shader/execution/unary/i32_conversion.bin": "cf99c846", - "webgpu/shader/execution/unary/u32_conversion.bin": "7f09c0ce", - "webgpu/shader/execution/unary/ai_assignment.bin": "98de2fcc", - "webgpu/shader/execution/binary/ai_arithmetic.bin": "6b88d962", - "webgpu/shader/execution/unary/ai_arithmetic.bin": "e584bcb7", - "webgpu/shader/execution/binary/af_matrix_matrix_multiplication.bin": "c82b9644", - "webgpu/shader/execution/binary/af_matrix_scalar_multiplication.bin": "7ea675b6", - "webgpu/shader/execution/binary/af_matrix_vector_multiplication.bin": "28f11a02", - "webgpu/shader/execution/derivatives.bin": "93c7bf53", - "webgpu/shader/execution/fwidth.bin": "e9668101" + "webgpu/shader/execution/binary/af_addition.bin": "aa14ec26", + "webgpu/shader/execution/binary/af_logical.bin": "8890f3f1", + "webgpu/shader/execution/binary/af_division.bin": "ad842076", + "webgpu/shader/execution/binary/af_matrix_addition.bin": "c6d8db63", + "webgpu/shader/execution/binary/af_matrix_subtraction.bin": "c09b18cc", + "webgpu/shader/execution/binary/af_multiplication.bin": "f0d2c49c", + "webgpu/shader/execution/binary/af_remainder.bin": "781a9ff6", + "webgpu/shader/execution/binary/af_subtraction.bin": "2cf4b54a", + "webgpu/shader/execution/binary/f16_addition.bin": "3efff4fc", + "webgpu/shader/execution/binary/f16_logical.bin": "de9424e7", + "webgpu/shader/execution/binary/f16_division.bin": "7cd144a1", + "webgpu/shader/execution/binary/f16_matrix_addition.bin": "3b986de1", + "webgpu/shader/execution/binary/f16_matrix_matrix_multiplication.bin": "4e135781", + "webgpu/shader/execution/binary/f16_matrix_scalar_multiplication.bin": "4c21775c", + "webgpu/shader/execution/binary/f16_matrix_subtraction.bin": "47e76662", + "webgpu/shader/execution/binary/f16_matrix_vector_multiplication.bin": "2b557015", + "webgpu/shader/execution/binary/f16_multiplication.bin": "6a3e21d7", + "webgpu/shader/execution/binary/f16_remainder.bin": "70c9358c", + "webgpu/shader/execution/binary/f16_subtraction.bin": "4563dd62", + "webgpu/shader/execution/binary/f32_addition.bin": "1f35a4dd", + "webgpu/shader/execution/binary/f32_logical.bin": "723abe00", + "webgpu/shader/execution/binary/f32_division.bin": "c95c68a", + "webgpu/shader/execution/binary/f32_matrix_addition.bin": "75957d7a", + "webgpu/shader/execution/binary/f32_matrix_matrix_multiplication.bin": "baf57d3c", + "webgpu/shader/execution/binary/f32_matrix_scalar_multiplication.bin": "cdbe1ae4", + "webgpu/shader/execution/binary/f32_matrix_subtraction.bin": "d9e7057", + "webgpu/shader/execution/binary/f32_matrix_vector_multiplication.bin": "634f59e6", + "webgpu/shader/execution/binary/f32_multiplication.bin": "a1b88039", + "webgpu/shader/execution/binary/f32_remainder.bin": "51048d79", + "webgpu/shader/execution/binary/f32_subtraction.bin": "aa74d81b", + "webgpu/shader/execution/binary/i32_arithmetic.bin": "e268d3de", + "webgpu/shader/execution/binary/i32_comparison.bin": "5c84159e", + "webgpu/shader/execution/binary/u32_arithmetic.bin": "99c3cdcf", + "webgpu/shader/execution/binary/u32_comparison.bin": "ee2050d9", + "webgpu/shader/execution/abs.bin": "d3a0c673", + "webgpu/shader/execution/acos.bin": "c6f18083", + "webgpu/shader/execution/acosh.bin": "2499b5b7", + "webgpu/shader/execution/asin.bin": "6b98a030", + "webgpu/shader/execution/asinh.bin": "129f4a25", + "webgpu/shader/execution/atan.bin": "6deed949", + "webgpu/shader/execution/atan2.bin": "395d1e72", + "webgpu/shader/execution/atanh.bin": "1a0f331a", + "webgpu/shader/execution/bitcast.bin": "ca478e97", + "webgpu/shader/execution/ceil.bin": "a002b286", + "webgpu/shader/execution/clamp.bin": "664e4e71", + "webgpu/shader/execution/cos.bin": "995ba625", + "webgpu/shader/execution/cosh.bin": "860e0f00", + "webgpu/shader/execution/cross.bin": "706065a0", + "webgpu/shader/execution/degrees.bin": "bacef53d", + "webgpu/shader/execution/determinant.bin": "a5fa820f", + "webgpu/shader/execution/distance.bin": "ae309025", + "webgpu/shader/execution/dot.bin": "6fe2eaad", + "webgpu/shader/execution/exp.bin": "16c78e4a", + "webgpu/shader/execution/exp2.bin": "aa3bd37a", + "webgpu/shader/execution/faceForward.bin": "6ec092ba", + "webgpu/shader/execution/floor.bin": "6c8c8506", + "webgpu/shader/execution/fma.bin": "d0240050", + "webgpu/shader/execution/fract.bin": "eedd46ad", + "webgpu/shader/execution/frexp.bin": "4f86736f", + "webgpu/shader/execution/inverseSqrt.bin": "c472a6d", + "webgpu/shader/execution/ldexp.bin": "f79e5c4c", + "webgpu/shader/execution/length.bin": "61c438dc", + "webgpu/shader/execution/log.bin": "602e67ab", + "webgpu/shader/execution/log2.bin": "59531d51", + "webgpu/shader/execution/max.bin": "315ad00", + "webgpu/shader/execution/min.bin": "5e67b47a", + "webgpu/shader/execution/mix.bin": "7b9c30fb", + "webgpu/shader/execution/modf.bin": "a926d2ab", + "webgpu/shader/execution/normalize.bin": "7a0f4894", + "webgpu/shader/execution/pack2x16float.bin": "5bda6ee2", + "webgpu/shader/execution/pow.bin": "d860633f", + "webgpu/shader/execution/quantizeToF16.bin": "e2cc5f7c", + "webgpu/shader/execution/radians.bin": "a564cae8", + "webgpu/shader/execution/reflect.bin": "ee6dda3f", + "webgpu/shader/execution/refract.bin": "769439d5", + "webgpu/shader/execution/round.bin": "e5c69d73", + "webgpu/shader/execution/saturate.bin": "a2baca59", + "webgpu/shader/execution/sign.bin": "c3031854", + "webgpu/shader/execution/sin.bin": "32cdc685", + "webgpu/shader/execution/sinh.bin": "807d1b08", + "webgpu/shader/execution/smoothstep.bin": "643c83a", + "webgpu/shader/execution/sqrt.bin": "ac7e5381", + "webgpu/shader/execution/step.bin": "ac99a50e", + "webgpu/shader/execution/tan.bin": "5efc9f6f", + "webgpu/shader/execution/tanh.bin": "8385aa65", + "webgpu/shader/execution/transpose.bin": "f3685a2e", + "webgpu/shader/execution/trunc.bin": "7eb127a7", + "webgpu/shader/execution/unpack2x16float.bin": "4eef663c", + "webgpu/shader/execution/unpack2x16snorm.bin": "c48522c4", + "webgpu/shader/execution/unpack2x16unorm.bin": "2af3269", + "webgpu/shader/execution/unpack4x8snorm.bin": "d4dbf395", + "webgpu/shader/execution/unpack4x8unorm.bin": "4c534802", + "webgpu/shader/execution/unary/af_arithmetic.bin": "1fde55fa", + "webgpu/shader/execution/unary/af_assignment.bin": "38273056", + "webgpu/shader/execution/unary/bool_conversion.bin": "3a8b1d94", + "webgpu/shader/execution/unary/f16_arithmetic.bin": "12ceb8ec", + "webgpu/shader/execution/unary/f16_conversion.bin": "585809f2", + "webgpu/shader/execution/unary/f32_arithmetic.bin": "8ebcea6b", + "webgpu/shader/execution/unary/f32_conversion.bin": "d6681b46", + "webgpu/shader/execution/unary/i32_arithmetic.bin": "b2cc8056", + "webgpu/shader/execution/unary/i32_conversion.bin": "bd484400", + "webgpu/shader/execution/unary/u32_conversion.bin": "82e99fb5", + "webgpu/shader/execution/unary/ai_assignment.bin": "c2c1d4e1", + "webgpu/shader/execution/binary/ai_arithmetic.bin": "b01bab0a", + "webgpu/shader/execution/unary/ai_arithmetic.bin": "dd916b48", + "webgpu/shader/execution/binary/af_matrix_matrix_multiplication.bin": "ad511107", + "webgpu/shader/execution/binary/af_matrix_scalar_multiplication.bin": "b54158a8", + "webgpu/shader/execution/binary/af_matrix_vector_multiplication.bin": "32c18616", + "webgpu/shader/execution/derivatives.bin": "a7a6547c", + "webgpu/shader/execution/fwidth.bin": "a4db2b2c" } \ No newline at end of file diff --git a/src/webgpu/api/validation/dispatch.spec.ts b/src/webgpu/api/validation/dispatch.spec.ts index 3b06a2259afc..01061ddb241d 100644 --- a/src/webgpu/api/validation/dispatch.spec.ts +++ b/src/webgpu/api/validation/dispatch.spec.ts @@ -4,6 +4,8 @@ Compute dispatch validation tests. import { AllFeaturesMaxLimitsGPUTest } from '../.././gpu_test.js'; import { makeTestGroup } from '../../../common/framework/test_group.js'; +import { keysOf } from '../../../common/util/data_tables.js'; +import { WGSLLanguageFeature } from '../../capability_info.js'; export const g = makeTestGroup(AllFeaturesMaxLimitsGPUTest); @@ -141,3 +143,338 @@ fn main(@builtin(${t.params.builtin}) input : u32, const expected = t.params.size === 'max' ? 0 : kMagic; t.expectGPUBufferValuesEqual(outputBuffer, new Uint32Array([expected])); }); + +interface RequiredSizeCase { + code: string; + size: number; + binding_type: GPUBufferBindingType; + requires?: WGSLLanguageFeature[]; +} + +const kRequiredSizeCases: Record = { + ro_storage_32bytes: { + code: ` +@group(0) @binding(0) var v : array; +@compute @workgroup_size(1) +fn main() { + _ = v; +}`, + size: 32, + binding_type: 'read-only-storage', + }, + storage_32bytes: { + code: ` +@group(0) @binding(0) var v : array; +@compute @workgroup_size(1) +fn main() { + _ = v; +}`, + size: 32, + binding_type: 'storage', + }, + uniform_32bytes: { + code: ` +@group(0) @binding(0) var v : array; +@compute @workgroup_size(1) +fn main() { + _ = v; +}`, + size: 32, + binding_type: 'uniform', + }, + ro_storage_sized_buffer: { + code: ` +@group(0) @binding(0) var v : buffer<32>; +@compute @workgroup_size(1) +fn main() { + _ = &v; +}`, + size: 32, + binding_type: 'read-only-storage', + requires: ['buffer_view'], + }, + storage_sized_buffer: { + code: ` +@group(0) @binding(0) var v : buffer<32>; +@compute @workgroup_size(1) +fn main() { + _ = &v; +}`, + size: 32, + binding_type: 'storage', + requires: ['buffer_view'], + }, + uniform_sized_buffer: { + code: ` +@group(0) @binding(0) var v : buffer<32>; +@compute @workgroup_size(1) +fn main() { + _ = &v; +}`, + size: 32, + binding_type: 'uniform', + requires: ['buffer_view'], + }, + ro_storage_unsized_buffer_bufferView1: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let p = bufferView>(&v, 0); +}`, + size: 16 * 4, + binding_type: 'read-only-storage', + requires: ['buffer_view'], + }, + ro_storage_unsized_buffer_bufferView2: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let p = bufferView>(&v, 56); +}`, + size: 16 * 4 + 56, + binding_type: 'read-only-storage', + requires: ['buffer_view'], + }, + storage_unsized_buffer_bufferView3: { + code: ` +struct S { + a: vec4u, + b: vec2u, +} +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let x = 16u; + let p = bufferView(&v, x); +}`, + size: 32, + binding_type: 'storage', + requires: ['buffer_view'], + }, + storage_unsized_buffer_bufferView4: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let p = bufferView>(&v, 0); +}`, + size: 8, + binding_type: 'storage', + requires: ['buffer_view'], + }, + storage_unsized_buffer_bufferView5: { + code: ` +struct S { + a: array +} + +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let p = bufferView(&v, 0); + let q = bufferView(&v, 0); +}`, + size: 16, + binding_type: 'storage', + requires: ['buffer_view'], + }, + storage_unsized_buffer_bufferView6: { + code: ` +struct S { + a: array +} + +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let p = bufferView(&v, 16); + let q = bufferView(&v, 0); +}`, + size: 20, + binding_type: 'storage', + requires: ['buffer_view'], + }, + storage_unsized_buffer_bufferView7: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let p = bufferView(&v, 16); +} +@fragment +fn main2() { + let q = bufferView(&v, 32); +}`, + size: 20, + binding_type: 'storage', + requires: ['buffer_view'], + }, + storage_unsized_buffer_bufferView8: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let p = bufferView(&v, 16); +} +fn foo() { + let q = bufferView(&v, 32); +}`, + size: 20, + binding_type: 'storage', + requires: ['buffer_view'], + }, + storage_unsized_buffer_bufferView9: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let p = bufferView(&v, 16); + foo(&v); +} +fn foo(p : ptr) { + let q = bufferView(&v, 32); +}`, + size: 16 + 32, + binding_type: 'storage', + requires: ['buffer_view', 'unrestricted_pointer_parameters'], + }, + ro_storage_unsized_buffer_bufferArrayView1: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let p = bufferArrayView>(&v, 0, 32); +}`, + size: 32, + binding_type: 'read-only-storage', + requires: ['buffer_view'], + }, + ro_storage_unsized_buffer_bufferArrayView2: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let p = bufferArrayView>(&v, 16, 32); +}`, + size: 16 + 32, + binding_type: 'read-only-storage', + requires: ['buffer_view'], + }, + ro_storage_unsized_buffer_bufferArrayView3: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let o = 16; + let s = 32; + let p = bufferArrayView>(&v, o, s); +}`, + size: 8, + binding_type: 'read-only-storage', + requires: ['buffer_view'], + }, + ro_storage_unsized_buffer_bufferArrayView4: { + code: ` +struct S { + a: vec4u, + b: u32, +} +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let o = 16; + let s = 32; + let p = bufferArrayView>(&v, o, s); +}`, + size: 32, + binding_type: 'read-only-storage', + requires: ['buffer_view'], + }, +}; + +g.test('shader_required_buffer_size') + .desc('Test that dispatch time validation occurs about the required buffer size') + .params(u => + u + .combine('case', keysOf(kRequiredSizeCases)) + .beginSubcases() + .combine('valid', [false, true] as const) + .combine('layout', ['auto', 'explicit'] as const) + ) + .fn(t => { + const testcase = kRequiredSizeCases[t.params.case]; + const features = testcase.requires ?? []; + features.forEach(f => { + t.skipIfLanguageFeatureNotSupported(f); + }); + + const buffer = t.createBufferTracked({ + size: t.params.valid ? testcase.size : testcase.size - 4, + usage: testcase.binding_type === 'uniform' ? GPUBufferUsage.UNIFORM : GPUBufferUsage.STORAGE, + }); + + const bgLayout = t.device.createBindGroupLayout({ + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: testcase.binding_type, + minBindingSize: 0, + }, + }, + ], + }); + const layout = t.device.createPipelineLayout({ bindGroupLayouts: [bgLayout] }); + + const pipeline = t.device.createComputePipeline({ + layout: t.params.layout === 'auto' ? 'auto' : layout, + compute: { + module: t.device.createShaderModule({ code: testcase.code }), + }, + }); + + if (t.params.layout === 'auto' && !t.params.valid) { + // 'auto' layout get minBindingSize from the shader. + t.expectValidationError(() => { + t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer, + }, + }, + ], + }), + true; + }); + } else { + // Expect dispatch time validation. + const bg = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer, + }, + }, + ], + }); + + const commandEncoder = t.device.createCommandEncoder(); + const computePassEncoder = commandEncoder.beginComputePass(); + computePassEncoder.setPipeline(pipeline); + computePassEncoder.setBindGroup(0, bg); + computePassEncoder.dispatchWorkgroups(1); + computePassEncoder.end(); + + t.expectValidationError(() => { + commandEncoder.finish(); + }, !t.params.valid); + } + }); diff --git a/src/webgpu/capability_info.ts b/src/webgpu/capability_info.ts index 3d7160bb894e..a65c6e880068 100644 --- a/src/webgpu/capability_info.ts +++ b/src/webgpu/capability_info.ts @@ -987,6 +987,7 @@ export const kKnownWGSLLanguageFeatures = [ 'swizzle_assignment', 'linear_indexing', 'texture_formats_tier1', + 'buffer_view', ] as const; export type WGSLLanguageFeature = (typeof kKnownWGSLLanguageFeatures)[number]; diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index 76b1c3974bb1..6e21fdd49c1d 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -511,6 +511,7 @@ "webgpu:api,validation,debugMarker:push_pop_call_count_unbalance,render_compute_pass:*": { "subcaseMS": 0.601 }, "webgpu:api,validation,dispatch:dispatch,linear_indexing_range:*": { "subcaseMS": 359.656 }, "webgpu:api,validation,dispatch:dispatchIndirect,linear_indexing_range:*": { "subcaseMS": 320.426 }, + "webgpu:api,validation,dispatch:shader_required_buffer_size:*": { "subcaseMS": 133.401 }, "webgpu:api,validation,encoding,beginComputePass:timestampWrites,invalid_query_set:*": { "subcaseMS": 0.201 }, "webgpu:api,validation,encoding,beginComputePass:timestampWrites,query_index:*": { "subcaseMS": 0.201 }, "webgpu:api,validation,encoding,beginComputePass:timestampWrites,query_set_type:*": { "subcaseMS": 0.401 }, @@ -1359,6 +1360,21 @@ "webgpu:shader,execution,expression,call,builtin,bitcast:vec4h_to_vec2f:*": { "subcaseMS": 22.812 }, "webgpu:shader,execution,expression,call,builtin,bitcast:vec4h_to_vec2i:*": { "subcaseMS": 20.915 }, "webgpu:shader,execution,expression,call,builtin,bitcast:vec4h_to_vec2u:*": { "subcaseMS": 29.514 }, + "webgpu:shader,execution,expression,call,builtin,bufferArrayView:array_length,functions:*": { "subcaseMS": 2076.378 }, + "webgpu:shader,execution,expression,call,builtin,bufferArrayView:array_length:*": { "subcaseMS": 315103.476 }, + "webgpu:shader,execution,expression,call,builtin,bufferArrayView:read:*": { "subcaseMS": 27926.111 }, + "webgpu:shader,execution,expression,call,builtin,bufferArrayView:read_layout:*": { "subcaseMS": 1143.557 }, + "webgpu:shader,execution,expression,call,builtin,bufferArrayView:write:*": { "subcaseMS": 1861.783 }, + "webgpu:shader,execution,expression,call,builtin,bufferArrayView:write_layout:*": { "subcaseMS": 1284.188 }, + "webgpu:shader,execution,expression,call,builtin,bufferLength:max_size_buffer:*": { "subcaseMS": 3.422 }, + "webgpu:shader,execution,expression,call,builtin,bufferLength:sized_buffer:*": { "subcaseMS": 549.908 }, + "webgpu:shader,execution,expression,call,builtin,bufferLength:unsized_buffer:*": { "subcaseMS": 390.207 }, + "webgpu:shader,execution,expression,call,builtin,bufferView:array_length,functions:*": { "subcaseMS": 24721.845 }, + "webgpu:shader,execution,expression,call,builtin,bufferView:array_length:*": { "subcaseMS": 19.803 }, + "webgpu:shader,execution,expression,call,builtin,bufferView:read:*": { "subcaseMS": 1.096 }, + "webgpu:shader,execution,expression,call,builtin,bufferView:read_layout:*": { "subcaseMS": 423.901 }, + "webgpu:shader,execution,expression,call,builtin,bufferView:write:*": { "subcaseMS": 54.157 }, + "webgpu:shader,execution,expression,call,builtin,bufferView:write_layout:*": { "subcaseMS": 316.101 }, "webgpu:shader,execution,expression,call,builtin,ceil:abstract_float:*": { "subcaseMS": 15217.441 }, "webgpu:shader,execution,expression,call,builtin,ceil:f16:*": { "subcaseMS": 29.209 }, "webgpu:shader,execution,expression,call,builtin,ceil:f32:*": { "subcaseMS": 11.132 }, @@ -2231,6 +2247,22 @@ "webgpu:shader,validation,expression,call,builtin,bitcast:must_use:*": { "subcaseMS": 1.715 }, "webgpu:shader,validation,expression,call,builtin,bitcast:valid_vec2h:*": { "subcaseMS": 3.405 }, "webgpu:shader,validation,expression,call,builtin,bitcast:valid_vec4h:*": { "subcaseMS": 5.610 }, + "webgpu:shader,validation,expression,call,builtin,bufferArrayView:buffer_type:*": { "subcaseMS": 43.489 }, + "webgpu:shader,validation,expression,call,builtin,bufferArrayView:early_eval_errors:*": { "subcaseMS": 297.765 }, + "webgpu:shader,validation,expression,call,builtin,bufferArrayView:lhs_call:*": { "subcaseMS": 10.967 }, + "webgpu:shader,validation,expression,call,builtin,bufferArrayView:must_use:*": { "subcaseMS": 45.264 }, + "webgpu:shader,validation,expression,call,builtin,bufferArrayView:offset_type:*": { "subcaseMS": 79.468 }, + "webgpu:shader,validation,expression,call,builtin,bufferArrayView:return_type:*": { "subcaseMS": 1483.486 }, + "webgpu:shader,validation,expression,call,builtin,bufferArrayView:size_type:*": { "subcaseMS": 78.211 }, + "webgpu:shader,validation,expression,call,builtin,bufferLength:data_type:*": { "subcaseMS": 31.283 }, + "webgpu:shader,validation,expression,call,builtin,bufferLength:must_use:*": { "subcaseMS": 197.831 }, + "webgpu:shader,validation,expression,call,builtin,bufferLength:return_type:*": { "subcaseMS": 320.197 }, + "webgpu:shader,validation,expression,call,builtin,bufferView:buffer_type:*": { "subcaseMS": 36.775 }, + "webgpu:shader,validation,expression,call,builtin,bufferView:early_eval_errors:*": { "subcaseMS": 18.020 }, + "webgpu:shader,validation,expression,call,builtin,bufferView:lhs_call:*": { "subcaseMS": 1.547 }, + "webgpu:shader,validation,expression,call,builtin,bufferView:must_use:*": { "subcaseMS": 196.136 }, + "webgpu:shader,validation,expression,call,builtin,bufferView:offset_type:*": { "subcaseMS": 72.505 }, + "webgpu:shader,validation,expression,call,builtin,bufferView:return_type:*": { "subcaseMS": 2079.184 }, "webgpu:shader,validation,expression,call,builtin,ceil:arguments:*": { "subcaseMS": 70.393 }, "webgpu:shader,validation,expression,call,builtin,ceil:integer_argument:*": { "subcaseMS": 1.456 }, "webgpu:shader,validation,expression,call,builtin,ceil:must_use:*": { "subcaseMS": 3.370 }, @@ -3016,6 +3048,8 @@ "webgpu:shader,validation,types,atomics:parse:*": { "subcaseMS": 7.692 }, "webgpu:shader,validation,types,atomics:trailing_comma:*": { "subcaseMS": 2.053 }, "webgpu:shader,validation,types,atomics:type:*": { "subcaseMS": 1.050 }, + "webgpu:shader,validation,types,buffer:address_space:*": { "subcaseMS": 45.804 }, + "webgpu:shader,validation,types,buffer:parse:*": { "subcaseMS": 233.842 }, "webgpu:shader,validation,types,enumerant:decl_value:*": { "subcaseMS": 45.257 }, "webgpu:shader,validation,types,enumerant:type_declaration:*": { "subcaseMS": 310.813 }, "webgpu:shader,validation,types,enumerant:value_type:*": { "subcaseMS": 211.873 }, diff --git a/src/webgpu/shader/execution/expression/call/builtin/bufferArrayView.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/bufferArrayView.spec.ts new file mode 100644 index 000000000000..178b14c45b29 --- /dev/null +++ b/src/webgpu/shader/execution/expression/call/builtin/bufferArrayView.spec.ts @@ -0,0 +1,612 @@ +export const description = ` +Execution tests for bufferArrayView +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf } from '../../../../../../common/util/data_tables.js'; +import { assert } from '../../../../../../common/util/util.js'; +import { AllFeaturesMaxLimitsGPUTest } from '../../../../../gpu_test.js'; +import { Type } from '../../../../../util/conversion.js'; + +import { + kBufferSizes, + kArrayLengthTypes, + kOffsets, + kSizes, + kCalls, + kStructDecls, + isValidArrayLengthCase, + calculateArrayLength, + runLengthTest, + kLayoutCases, + runReadLayoutTest, + runWriteLayoutTest, + runReadWriteTest, +} from './buffer_view_utils.js'; + +export const g = makeTestGroup(AllFeaturesMaxLimitsGPUTest); + +g.test('array_length') + .desc('Tests arrayLength from a bufferView') + .params(u => + u + .combine('type', keysOf(kArrayLengthTypes)) + .combine('aspace', ['workgroup', 'storage', 'uniform'] as const) + .beginSubcases() + .combine('sized', [false, true] as const) + .combine('override', [false, true] as const) + .combine('bufferSize', kBufferSizes) + .combine('dynamic_offset', [0, 256] as const) + .filter(t => { + if (t.aspace === 'workgroup' && t.dynamic_offset !== 0) { + return false; + } + if (t.aspace !== 'workgroup' && t.override === true) { + return false; + } + return t.sized || t.aspace === 'storage'; + }) + ) + .fn(t => { + const testcase = kArrayLengthTypes[t.params.type]; + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (t.params.aspace === 'uniform' && testcase.uniformStdLayout === true) { + t.skipIfLanguageFeatureNotSupported('uniform_buffer_standard_layout'); + } + if (testcase.f16) { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + + const bSize = t.params.override ? 'bufferSize' : `${t.params.bufferSize}`; + const buffer_ty = `buffer${t.params.sized ? `<${bSize}>` : ``}`; + let decl = ''; + switch (t.params.aspace) { + case 'workgroup': + decl = `var v : ${buffer_ty};\n@group(0) @binding(0) var dummy : array;`; + break; + case 'uniform': + decl = `@group(0) @binding(0) var v : ${buffer_ty};`; + break; + case 'storage': + decl = `@group(0) @binding(0) var v : ${buffer_ty};`; + break; + } + + const values: number[] = []; + const offsets: number[] = []; + const sizes: number[] = []; + // Limit the number of assignments. + for (let o = 0; o < kOffsets.length; o += 2) { + for (const size of kSizes) { + const offset = kOffsets[o]; + + // Skip any case that results in an invalid memory reference. + // Split the writes up one per thread. + if (isValidArrayLengthCase(testcase, offset, size, t.params.bufferSize)) { + values.push(calculateArrayLength(testcase, offset, size, t.params.bufferSize)); + // Setup the offset and size buffer values so they can be indexed by invocation id. + offsets.push(offset); + sizes.push(size); + } + } + } + + assert(values.length > 0, 'no tests run'); + + const access = testcase.access ?? ''; + const addrOf = access === '' ? '' : '&'; + + const enables = testcase.f16 === true ? 'enable f16;' : ''; + + const wgsl = ` +${enables} + +${kStructDecls} + +override bufferSize : u32 = 0u; +${decl} +@group(0) @binding(1) var offsets : array; +@group(0) @binding(2) var sizes : array; // unused +@group(0) @binding(3) var out : array; + +override wgx : u32; +@compute @workgroup_size(wgx) +fn main(@builtin(global_invocation_id) gid : vec3u, + @builtin(local_invocation_index) lid : u32) { + ${t.params.aspace === 'workgroup' ? '_ = dummy[0];' : ''} + + if gid.x >= ${values.length} { + return; + } + + out[gid.x] = arrayLength(${addrOf}bufferArrayView<${ + testcase.type + }>(&v, offsets[gid.x], sizes[gid.x])${access}); +}`; + + runLengthTest( + t, + wgsl, + t.params.aspace === 'uniform' ? GPUBufferUsage.UNIFORM : GPUBufferUsage.STORAGE, + t.params.bufferSize, + t.params.dynamic_offset, + offsets, + sizes, + values + ); + }); + +g.test('array_length,functions') + .desc('Tests arrayLength from a bufferView') + .params(u => + u + .combine('type', keysOf(kArrayLengthTypes)) + .combine('aspace', ['workgroup', 'storage', 'uniform'] as const) + .beginSubcases() + .combine('call', kCalls) + .combine('bufferSize', kBufferSizes) + .combine('dynamic_offset', [0, 256] as const) + .filter(t => { + if (t.bufferSize === 128) { + return false; + } + return t.aspace !== 'workgroup' || t.dynamic_offset === 0; + }) + ) + .fn(t => { + const testcase = kArrayLengthTypes[t.params.type]; + t.skipIfLanguageFeatureNotSupported('buffer_view'); + t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters'); + if (t.params.aspace === 'uniform' && testcase.uniformStdLayout === true) { + t.skipIfLanguageFeatureNotSupported('uniform_buffer_standard_layout'); + } + if (testcase.f16) { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + + const buffer_ty = `buffer<${t.params.bufferSize}>`; + let decl = ''; + switch (t.params.aspace) { + case 'workgroup': + decl = `var v : ${buffer_ty};\n@group(0) @binding(0) var dummy : array;`; + break; + case 'uniform': + decl = `@group(0) @binding(0) var v : ${buffer_ty};`; + break; + case 'storage': + decl = `@group(0) @binding(0) var v : ${buffer_ty};`; + break; + } + + const values: number[] = []; + const offsets: number[] = []; + const sizes: number[] = []; + // Limit the number of assignments. + for (let o = 0; o < kOffsets.length; o += 2) { + for (const size of kSizes) { + const offset = kOffsets[o]; + const usedSize = + t.params.call === 'unsized' ? t.params.bufferSize : t.params.bufferSize / 2; + + // Skip any case that results in an invalid memory reference. + // Split the writes up one per thread. + if (isValidArrayLengthCase(testcase, offset, size, usedSize)) { + values.push(calculateArrayLength(testcase, offset, size, usedSize)); + // Setup the offset and size buffer values so they can be indexed by invocation id. + offsets.push(offset); + sizes.push(size); + } + } + } + + assert(values.length > 0, 'no tests run'); + + const access = testcase.access ?? ''; + const addrOf = access === '' ? '' : '&'; + + const enables = testcase.f16 === true ? 'enable f16;' : ''; + + const wgsl = ` +${enables} + +${kStructDecls} + +fn unsizedLength(p : ptr<${t.params.aspace}, buffer>, gidx : u32) -> u32 { + return arrayLength(${addrOf}bufferArrayView<${ + testcase.type + }>(p, offsets[gidx], sizes[gidx])${access}); +} + +fn sizedLength(p : ptr<${t.params.aspace}, buffer<${t.params.bufferSize / 2}>>, gidx : u32) -> u32 { + return arrayLength(${addrOf}bufferArrayView<${ + testcase.type + }>(p, offsets[gidx], sizes[gidx])${access}); +} + +fn sized_indirectLength(p : ptr<${t.params.aspace}, buffer<${ + t.params.bufferSize / 2 + }>>, gidx : u32) -> u32 { + return unsizedLength(p, gidx); +} + +override bufferSize : u32 = 0u; +${decl} +@group(0) @binding(1) var offsets : array; +@group(0) @binding(2) var sizes : array; // unused +@group(0) @binding(3) var out : array; + +override wgx : u32; +@compute @workgroup_size(wgx) +fn main(@builtin(global_invocation_id) gid : vec3u, + @builtin(local_invocation_index) lid : u32) { + ${t.params.aspace === 'workgroup' ? '_ = dummy[0];' : ''} + + if gid.x >= ${values.length} { + return; + } + + out[gid.x] = ${t.params.call}Length(&v, gid.x); +}`; + + runLengthTest( + t, + wgsl, + t.params.aspace === 'uniform' ? GPUBufferUsage.UNIFORM : GPUBufferUsage.STORAGE, + t.params.bufferSize, + t.params.dynamic_offset, + offsets, + sizes, + values + ); + }); + +g.test('read_layout') + .desc('Test reading memory layout from a bufferArrayView') + .params(u => + u + .combine('case', keysOf(kLayoutCases)) + .beginSubcases() + .combine('aspace', ['storage', 'ro_storage', 'uniform', 'workgroup'] as const) + .combine('offset', kOffsets) + .filter(t => { + const testcase = kLayoutCases[t.case]; + if ((t.offset & (testcase.align - 1)) !== 0) { + return false; + } + return testcase.offset + t.offset < 252; + }) + ) + .fn(t => { + const testcase = kLayoutCases[t.params.case]; + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (t.params.aspace === 'uniform') { + if (testcase.uniformStdLayoutArrayView === true || testcase.uniformStdLayoutView === true) { + t.skipIfLanguageFeatureNotSupported('uniform_buffer_standard_layout'); + } + } + if (testcase.f16) { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + + const enables = testcase.f16 === true ? 'enable f16;' : ''; + const v = t.params.aspace === 'workgroup' ? 'wg_var' : 'in'; + + const wgsl = ` +${enables} + +${testcase.decl ?? ''} + +@group(0) @binding(0) +var<${t.params.aspace === 'uniform' ? 'uniform' : 'storage'}${ + t.params.aspace === 'storage' ? ', read_write' : '' + }> in : buffer<256 * 4>; + +@group(0) @binding(1) var out : u32; + +var wg_var : buffer<256 * 4>; + +@compute @workgroup_size(1) +fn main() { + let in_ptr = bufferView>(&in, 0); + let wg_ptr = bufferView>(&wg_var, 0); + for (var i = 0; i < 256 / 4; i++) { + (*wg_ptr)[i] = (*in_ptr)[i]; + } + + workgroupBarrier(); + + let p = bufferArrayView>(&${v}, ${ + t.params.offset + }, bufferLength(&${v}) - ${t.params.offset}); + out = u32((*p)[0]${testcase.access}); +} +`; + + runReadLayoutTest(t, testcase, wgsl, t.params.aspace, t.params.offset); + }); + +g.test('write_layout') + .desc('Test writing memory layout via a bufferView') + .params(u => + u + .combine('case', keysOf(kLayoutCases)) + .beginSubcases() + .combine('assign', ['let', 'call'] as const) + .combine('aspace', ['storage', 'workgroup'] as const) + .combine('offset', kOffsets) + .filter(t => { + const testcase = kLayoutCases[t.case]; + if ((t.offset & (testcase.align - 1)) !== 0) { + return false; + } + return testcase.offset + t.offset < 252; + }) + ) + .fn(t => { + const testcase = kLayoutCases[t.params.case]; + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (testcase.f16) { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + + const enables = testcase.f16 === true ? 'enable f16;' : ''; + const v = t.params.aspace === 'workgroup' ? 'wg_var' : 'out'; + const call_expr = `bufferArrayView>(&${v}, ${t.params.offset}, bufferLength(&${v}) - ${t.params.offset})`; + + let post_assign = ''; + if (t.params.aspace === 'workgroup') { + post_assign = ` + workgroupBarrier(); + + let out_ptr = bufferView>(&out, 0); + let wg_ptr = bufferView>(&wg_var, 0); + for (var i = 0; i < 256; i++) { + (*out_ptr)[i] = (*wg_ptr)[i]; + } +`; + } + + const wgsl = ` +${enables} + +${testcase.decl ?? ''} + +@group(0) @binding(0) var in : u32; +@group(0) @binding(1) var out : buffer<256 * 4>; + + +var wg_var : buffer<256 * 4>; + +@compute @workgroup_size(1) +fn main() { + let val = ${testcase.f16 ? 'f16' : testcase.f32 ? 'f32' : 'u32'}(in); + let ptr = ${call_expr}; + (*${t.params.assign === 'call' ? call_expr : 'ptr'})[0]${testcase.access} = val; + + ${post_assign} +} +`; + + runWriteLayoutTest(t, testcase, wgsl, t.params.offset); + }); + +g.test('read') + .desc('Test reading various types from bufferArrayView') + .params(u => + u + .combine('base_type', ['u32', 'i32', 'f32', 'f16'] as const) + .beginSubcases() + .combine('wrap', ['none', 'vector', 'array', 'matrix'] as const) + .combine('width', [1, 2, 3, 4] as const) + .combine('aspace', ['workgroup', 'storage', 'ro_storage', 'uniform'] as const) + .combine('offset', [0, 4, 8, 12, 16, 32, 48, 64] as const) + .filter(t => { + if (t.wrap !== 'none' && t.width === 1) { + return false; + } + if (t.wrap === 'none' && t.width !== 1) { + return false; + } + if (t.wrap === 'matrix' && t.base_type !== 'f32' && t.base_type !== 'f16') { + return false; + } + if (t.aspace === 'uniform' && t.wrap === 'array') { + return false; + } + const ty = Type[t.base_type]; + let align = ty.alignment; + switch (t.wrap) { + case 'vector': + align = Type.vec(t.width, ty).alignment; + break; + case 'array': + align = Type.array(2, ty).alignment; + break; + case 'matrix': + align = Type.mat(t.width, 2, ty).alignment; + break; + case 'none': + break; + } + if (t.offset % align !== 0) { + return false; + } + if (t.aspace === 'uniform' && align % 16 !== 0) { + return false; + } + return true; + }) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (t.params.base_type === 'f16') { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + + const ele_ty = Type[t.params.base_type]; + let ty: Type = ele_ty; + switch (t.params.wrap) { + case 'vector': + ty = Type.vec(t.params.width, ele_ty); + break; + case 'array': + ty = Type.array(t.params.width, ele_ty); + break; + case 'matrix': + ty = Type.mat(t.params.width, 2, ele_ty); + break; + case 'none': + break; + } + + const bufferSize = 128; + const enables = t.params.base_type === 'f16' ? 'enable f16;' : ''; + const v = t.params.aspace === 'workgroup' ? 'wg_var' : 'in'; + const wgsl = ` +${enables} + + +@group(0) @binding(0) +var<${t.params.aspace === 'uniform' ? 'uniform' : 'storage'}${ + t.params.aspace === 'storage' ? ', read_write' : '' + }> in : buffer<${bufferSize}>; + +@group(0) @binding(1) var out : ${ty.toString()}; + +var wg_var : buffer<${bufferSize}>; + +@compute @workgroup_size(1) +fn main() { + let in_ptr = bufferView>(&in, 0); + let wg_ptr = bufferView>(&wg_var, 0); + for (var i = 0; i < 256 / 4; i++) { + (*wg_ptr)[i] = (*in_ptr)[i]; + } + + workgroupBarrier(); + + let p = bufferArrayView>(&${v}, ${ + t.params.offset + }, bufferLength(&${v}) - ${t.params.offset}); + out = (*p)[0]; +} +`; + + runReadWriteTest(true, t, wgsl, ele_ty, ty, t.params.aspace, t.params.offset, bufferSize); + }); + +g.test('write') + .desc('Test writing various types via bufferArrayView') + .params(u => + u + .combine('base_type', ['u32', 'i32', 'f32', 'f16'] as const) + .beginSubcases() + .combine('wrap', ['none', 'vector', 'array', 'matrix'] as const) + .combine('width', [1, 2, 3, 4] as const) + .combine('aspace', ['workgroup', 'storage'] as const) + .combine('offset', [0, 4, 8, 12, 16, 32, 48, 64] as const) + .combine('assign', ['let', 'call'] as const) + .combine('swizzle', [false, true] as const) + .filter(t => { + if (t.wrap !== 'none' && t.width === 1) { + return false; + } + if (t.wrap === 'none' && t.width !== 1) { + return false; + } + if (t.wrap === 'matrix' && t.base_type !== 'f32' && t.base_type !== 'f16') { + return false; + } + if (t.wrap !== 'vector' && t.swizzle) { + return false; + } + const ty = Type[t.base_type]; + let align = ty.alignment; + switch (t.wrap) { + case 'vector': + align = Type.vec(t.width, ty).alignment; + break; + case 'array': + align = Type.array(2, ty).alignment; + break; + case 'matrix': + align = Type.mat(t.width, 2, ty).alignment; + break; + case 'none': + break; + } + if (t.offset % align !== 0) { + return false; + } + return true; + }) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (t.params.base_type === 'f16') { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + if (t.params.swizzle) { + t.skipIfLanguageFeatureNotSupported('swizzle_assignment'); + } + + let swizzle = ''; + const ele_ty = Type[t.params.base_type]; + let ty: Type = ele_ty; + switch (t.params.wrap) { + case 'vector': + ty = Type.vec(t.params.width, ele_ty); + if (t.params.swizzle) { + swizzle = '.xyzw'.substring(0, 1 + t.params.width); + } + break; + case 'array': + ty = Type.array(t.params.width, ele_ty); + break; + case 'matrix': + ty = Type.mat(t.params.width, 2, ele_ty); + break; + case 'none': + break; + } + + const bufferSize = 128; + const enables = t.params.base_type === 'f16' ? 'enable f16;' : ''; + const v = t.params.aspace === 'workgroup' ? 'wg_var' : 'out'; + + let post_assign = ''; + if (t.params.aspace === 'workgroup') { + post_assign = ` + workgroupBarrier(); + + let out_ptr = bufferView>(&out, 0); + let wg_ptr = bufferView>(&wg_var, 0); + for (var i = 0; i < 256; i++) { + (*out_ptr)[i] = (*wg_ptr)[i]; + } +`; + } + + const call_expr = `bufferArrayView>(&${v}, ${ + t.params.offset + }, bufferLength(&${v}) - ${t.params.offset})`; + const wgsl = ` +${enables} + +@group(0) @binding(0) var in : ${ty.toString()}; +@group(0) @binding(1) var out : buffer<${bufferSize}>; + + +var wg_var : buffer<${bufferSize}>; + +@compute @workgroup_size(1) +fn main() { + let ptr = ${call_expr}; + (*${t.params.assign === 'call' ? call_expr : 'ptr'})[0]${swizzle} = in; + + ${post_assign} +} +`; + + runReadWriteTest(false, t, wgsl, ele_ty, ty, t.params.aspace, t.params.offset, bufferSize); + }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/bufferLength.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/bufferLength.spec.ts new file mode 100644 index 000000000000..57754cdb55bc --- /dev/null +++ b/src/webgpu/shader/execution/expression/call/builtin/bufferLength.spec.ts @@ -0,0 +1,485 @@ +export const description = ` +Execution tests for bufferLength + +Tests sized and unsized buffers across all address spaces (where applicable). +Tested against buffers with static and dynamic offsets. +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { GPUTest } from '../../../../../gpu_test.js'; + +export const g = makeTestGroup(GPUTest); + +g.test('sized_buffer') + .desc('Tests bufferLength directly on sized buffer variables') + .params(u => + u + .combine('size', [256, 512, 600, 1024, 2048] as const) + .beginSubcases() + .combine('param', ['none', 'unsized', 'sized'] as const) + .combine('padding', [false, true] as const) + .combine('dynamic_offset', [false, true] as const) + .combine('offset', [0, 256] as const) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (t.params.param !== 'none') { + t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters'); + } + + let fn_decls = ''; + const assigns = ['', '', '', '']; + switch (t.params.param) { + case 'none': + assigns[0] = 'bufferLength(&uniform_buffer)'; + assigns[1] = 'bufferLength(&ro_storage_buffer)'; + assigns[2] = 'bufferLength(&storage_buffer)'; + assigns[3] = 'bufferLength(&workgroup_buffer)'; + break; + case 'unsized': + fn_decls = ` +fn uniformUnsized(p : ptr) -> u32 { + return bufferLength(p); +} +fn roStorageUnsized(p : ptr) -> u32 { + return bufferLength(p); +} +fn storageUnsized(p : ptr) -> u32 { + return bufferLength(p); +} +fn workgroupUnsized(p : ptr) -> u32 { + return bufferLength(p); +}`; + assigns[0] = 'uniformUnsized(&uniform_buffer)'; + assigns[1] = 'roStorageUnsized(&ro_storage_buffer)'; + assigns[2] = 'storageUnsized(&storage_buffer)'; + assigns[3] = 'workgroupUnsized(&workgroup_buffer)'; + break; + case 'sized': + fn_decls = ` +fn uniformSized(p : ptr>) -> u32 { + return bufferLength(p); +} +fn roStorageSized(p : ptr>) -> u32 { + return bufferLength(p); +} +fn storageSized(p : ptr, read_write>) -> u32 { + return bufferLength(p); +} +fn workgroupSized(p : ptr>) -> u32 { + return bufferLength(p); +}`; + assigns[0] = 'uniformSized(&uniform_buffer)'; + assigns[1] = 'roStorageSized(&ro_storage_buffer)'; + assigns[2] = 'storageSized(&storage_buffer)'; + assigns[3] = 'workgroupSized(&workgroup_buffer)'; + break; + } + + const wgsl = ` +@group(0) @binding(0) var uniform_buffer : buffer<${t.params.size}>; +@group(0) @binding(1) var ro_storage_buffer : buffer<${t.params.size}>; +@group(0) @binding(2) var storage_buffer : buffer<${t.params.size}>; +var workgroup_buffer : buffer<${t.params.size}>; + +@group(0) @binding(3) var out : array; + +${fn_decls} + +@compute @workgroup_size(4) +fn main(@builtin(local_invocation_index) lid : u32) { + out[0] = ${assigns[0]}; + out[1] = ${assigns[1]}; + out[2] = ${assigns[2]}; + out[3] = ${assigns[3]}; +}`; + + const padding = t.params.padding ? 256 : 0; + const offset = t.params.offset; + const dynOffset = t.params.dynamic_offset ? 256 : 0; + const bufferSize = t.params.size + padding + offset + dynOffset; + const noOffsetSize = t.params.size + padding; + + const uniformBuffer = t.createBufferTracked({ + size: bufferSize, + usage: GPUBufferUsage.UNIFORM, + }); + const roStorageBuffer = t.createBufferTracked({ + size: bufferSize, + usage: GPUBufferUsage.STORAGE, + }); + const storageBuffer = t.createBufferTracked({ + size: bufferSize, + usage: GPUBufferUsage.STORAGE, + }); + const outputBuffer = t.createBufferTracked({ + size: 4 * 4, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + + const bgLayout = t.device.createBindGroupLayout({ + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'uniform', + hasDynamicOffset: t.params.dynamic_offset, + minBindingSize: t.params.size, + }, + }, + { + binding: 1, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'read-only-storage', + hasDynamicOffset: t.params.dynamic_offset, + minBindingSize: t.params.size, + }, + }, + { + binding: 2, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'storage', + hasDynamicOffset: t.params.dynamic_offset, + minBindingSize: t.params.size, + }, + }, + { + binding: 3, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'storage', + minBindingSize: 4 * 4, + }, + }, + ], + }); + + const pipelineLayout = t.device.createPipelineLayout({ + bindGroupLayouts: [bgLayout], + }); + const pipeline = t.device.createComputePipeline({ + layout: pipelineLayout, + compute: { + module: t.device.createShaderModule({ code: wgsl }), + }, + }); + + const bg = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer: uniformBuffer, + offset, + size: noOffsetSize, + }, + }, + { + binding: 1, + resource: { + buffer: roStorageBuffer, + offset, + size: noOffsetSize, + }, + }, + { + binding: 2, + resource: { + buffer: storageBuffer, + offset, + size: noOffsetSize, + }, + }, + { + binding: 3, + resource: { + buffer: outputBuffer, + size: 4 * 4, + }, + }, + ], + }); + + const encoder = t.device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + if (t.params.dynamic_offset) { + pass.setBindGroup(0, bg, [dynOffset, dynOffset, dynOffset]); + } else { + pass.setBindGroup(0, bg); + } + pass.dispatchWorkgroups(1, 1, 1); + pass.end(); + t.queue.submit([encoder.finish()]); + + const expectedSize = t.params.param === 'sized' ? t.params.size / 2 : t.params.size; + const expected = new Uint32Array([expectedSize, expectedSize, expectedSize, expectedSize]); + t.expectGPUBufferValuesEqual(outputBuffer, expected); + }); + +g.test('unsized_buffer') + .desc('Tests bufferLength directly on unsized buffer variables') + .params(u => + u + .combine('size', [256, 512, 600, 1024, 2048] as const) + .beginSubcases() + .combine('param', ['none', 'unsized'] as const) + .combine('padding', [false, true] as const) + .combine('dynamic_offset', [false, true] as const) + .combine('offset', [0, 256] as const) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (t.params.param !== 'none') { + t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters'); + } + + let fn_decls = ''; + const assigns = ['', '', '', '']; + switch (t.params.param) { + case 'none': + assigns[0] = 'bufferLength(&ro_storage_buffer)'; + assigns[1] = 'bufferLength(&storage_buffer)'; + break; + case 'unsized': + fn_decls = ` +fn roStorageUnsized(p : ptr) -> u32 { + return bufferLength(p); +} +fn storageUnsized(p : ptr) -> u32 { + return bufferLength(p); +}`; + assigns[0] = 'roStorageUnsized(&ro_storage_buffer)'; + assigns[1] = 'storageUnsized(&storage_buffer)'; + break; + } + + const wgsl = ` +@group(0) @binding(0) var ro_storage_buffer : buffer; +@group(0) @binding(1) var storage_buffer : buffer; + +@group(0) @binding(2) var out : array; + +${fn_decls} + +@compute @workgroup_size(4) +fn main(@builtin(local_invocation_index) lid : u32) { + out[0] = ${assigns[0]}; + out[1] = ${assigns[1]}; +}`; + + const padding = t.params.padding ? 256 : 0; + const offset = t.params.offset; + const dynOffset = t.params.dynamic_offset ? 256 : 0; + const bufferSize = t.params.size + padding + offset + dynOffset; + const noOffsetSize = t.params.size + padding; + + const roStorageBuffer = t.createBufferTracked({ + size: bufferSize, + usage: GPUBufferUsage.STORAGE, + }); + const storageBuffer = t.createBufferTracked({ + size: bufferSize, + usage: GPUBufferUsage.STORAGE, + }); + const outputBuffer = t.createBufferTracked({ + size: 4 * 4, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + + const bgLayout = t.device.createBindGroupLayout({ + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'read-only-storage', + hasDynamicOffset: t.params.dynamic_offset, + minBindingSize: t.params.size, + }, + }, + { + binding: 1, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'storage', + hasDynamicOffset: t.params.dynamic_offset, + minBindingSize: t.params.size, + }, + }, + { + binding: 2, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'storage', + minBindingSize: 4 * 4, + }, + }, + ], + }); + + const pipelineLayout = t.device.createPipelineLayout({ + bindGroupLayouts: [bgLayout], + }); + const pipeline = t.device.createComputePipeline({ + layout: pipelineLayout, + compute: { + module: t.device.createShaderModule({ code: wgsl }), + }, + }); + + const bg = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer: roStorageBuffer, + offset, + size: noOffsetSize, + }, + }, + { + binding: 1, + resource: { + buffer: storageBuffer, + offset, + size: noOffsetSize, + }, + }, + { + binding: 2, + resource: { + buffer: outputBuffer, + size: 4 * 4, + }, + }, + ], + }); + + const encoder = t.device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + if (t.params.dynamic_offset) { + pass.setBindGroup(0, bg, [dynOffset, dynOffset]); + } else { + pass.setBindGroup(0, bg); + } + pass.dispatchWorkgroups(1, 1, 1); + pass.end(); + t.queue.submit([encoder.finish()]); + + const expected = new Uint32Array([noOffsetSize, noOffsetSize]); + t.expectGPUBufferValuesEqual(outputBuffer, expected); + }); + +g.test('max_size_buffer') + .desc('Test with a maximum sized buffer') + .params(u => + u + .combine('aspace', ['workgroup', 'uniform', 'storage', 'ro_storage'] as const) + .combine('sized', [false, true] as const) + .filter(t => { + return t.sized === true || t.aspace === 'storage' || t.aspace === 'ro_storage'; + }) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + + const { + maxUniformBufferBindingSize, + maxStorageBufferBindingSize, + maxComputeWorkgroupStorageSize, + } = t.device.limits; + + let size = 16; + switch (t.params.aspace) { + case 'workgroup': + size = maxComputeWorkgroupStorageSize; + break; + case 'uniform': + size = maxUniformBufferBindingSize; + break; + case 'storage': + case 'ro_storage': + size = maxStorageBufferBindingSize; + break; + } + + const type = `buffer${t.params.sized ? `<${size}u>` : ``}`; + let decl = ''; + switch (t.params.aspace) { + case 'workgroup': + decl = `var v : ${type};\n@group(0) @binding(1) var dummy : u32;`; + break; + case 'uniform': + decl = `@group(0) @binding(1) var v: ${type};`; + break; + case 'storage': + decl = `@group(0) @binding(1) var v : ${type};`; + break; + case 'ro_storage': + decl = `@group(0) @binding(1) var v : ${type};`; + break; + } + + const wgsl = ` +@group(0) @binding(0) var out : u32; +${decl} + +@compute @workgroup_size(1) +fn main() { + out = bufferLength(&v); + ${t.params.aspace === 'workgroup' ? '_ = dummy;' : ''} +}`; + + const outputBuffer = t.createBufferTracked({ + size: 4, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + // For the workgroup case we create a dummy storage buffer to simplify bindings. + const inputBuffer = t.createBufferTracked({ + size, + usage: t.params.aspace === 'uniform' ? GPUBufferUsage.UNIFORM : GPUBufferUsage.STORAGE, + }); + + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ code: wgsl }), + }, + }); + + const bg = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer: outputBuffer, + }, + }, + { + binding: 1, + resource: { + buffer: inputBuffer, + }, + }, + ], + }); + + 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 expected = new Uint32Array([size]); + t.expectGPUBufferValuesEqual(outputBuffer, expected); + }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/bufferView.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/bufferView.spec.ts new file mode 100644 index 000000000000..b36fbc1c4533 --- /dev/null +++ b/src/webgpu/shader/execution/expression/call/builtin/bufferView.spec.ts @@ -0,0 +1,578 @@ +export const description = ` +Execution tests for bufferView +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf } from '../../../../../../common/util/data_tables.js'; +import { assert } from '../../../../../../common/util/util.js'; +import { AllFeaturesMaxLimitsGPUTest } from '../../../../../gpu_test.js'; +import { Type } from '../../../../../util/conversion.js'; + +import { + kBufferSizes, + kArrayLengthTypes, + kOffsets, + kCalls, + kStructDecls, + isValidArrayLengthCase, + calculateArrayLength, + runLengthTest, + kLayoutCases, + runReadLayoutTest, + runWriteLayoutTest, + runReadWriteTest, +} from './buffer_view_utils.js'; + +export const g = makeTestGroup(AllFeaturesMaxLimitsGPUTest); + +g.test('array_length') + .desc('Tests arrayLength from a bufferView') + .params(u => + u + .combine('type', keysOf(kArrayLengthTypes)) + .combine('aspace', ['workgroup', 'storage', 'uniform'] as const) + .beginSubcases() + .combine('sized', [false, true] as const) + .combine('override', [false, true] as const) + .combine('bufferSize', kBufferSizes) + .combine('dynamic_offset', [0, 256] as const) + .filter(t => { + if (t.aspace === 'workgroup' && t.dynamic_offset !== 0) { + return false; + } + if (t.aspace !== 'workgroup' && t.override === true) { + return false; + } + return t.sized || t.aspace === 'storage'; + }) + ) + .fn(t => { + const testcase = kArrayLengthTypes[t.params.type]; + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (t.params.aspace === 'uniform' && testcase.uniformStdLayout === true) { + t.skipIfLanguageFeatureNotSupported('uniform_buffer_standard_layout'); + } + if (testcase.f16) { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + + const bSize = t.params.override ? 'bufferSize' : `${t.params.bufferSize}`; + const buffer_ty = `buffer${t.params.sized ? `<${bSize}>` : ``}`; + let decl = ''; + switch (t.params.aspace) { + case 'workgroup': + decl = `var v : ${buffer_ty};\n@group(0) @binding(0) var dummy : array;`; + break; + case 'uniform': + decl = `@group(0) @binding(0) var v : ${buffer_ty};`; + break; + case 'storage': + decl = `@group(0) @binding(0) var v : ${buffer_ty};`; + break; + } + + const values: number[] = []; + const offsets: number[] = []; + const sizes: number[] = [0]; + for (const offset of kOffsets) { + // Skip any case that results in an invalid memory reference. + // Split the writes up one per thread. + if (isValidArrayLengthCase(testcase, offset, 0, t.params.bufferSize)) { + values.push(calculateArrayLength(testcase, offset, 0, t.params.bufferSize)); + // Setup the offset and size buffer values so they can be indexed by invocation id. + offsets.push(offset); + } + } + assert(values.length > 0, 'no tests run'); + + const access = testcase.access ?? ''; + const addrOf = access === '' ? '' : '&'; + + const enables = testcase.f16 === true ? 'enable f16;' : ''; + + const wgsl = ` +${enables} + +${kStructDecls} + +override bufferSize : u32 = 0u; +${decl} +@group(0) @binding(1) var offsets : array; +@group(0) @binding(2) var sizes : array; // unused +@group(0) @binding(3) var out : array; + +override wgx : u32; +@compute @workgroup_size(wgx) +fn main(@builtin(global_invocation_id) gid : vec3u) { + ${t.params.aspace === 'workgroup' ? '_ = dummy[0];' : ''} + + if gid.x >= ${values.length} { + return; + } + + out[gid.x] = arrayLength(${addrOf}bufferView<${testcase.type}>(&v, offsets[gid.x])${access}); +}`; + + runLengthTest( + t, + wgsl, + t.params.aspace === 'uniform' ? GPUBufferUsage.UNIFORM : GPUBufferUsage.STORAGE, + t.params.bufferSize, + t.params.dynamic_offset, + offsets, + sizes, + values + ); + }); + +g.test('array_length,functions') + .desc('Tests arrayLength from a bufferView through various function calls') + .params(u => + u + .combine('type', keysOf(kArrayLengthTypes)) + .combine('aspace', ['workgroup', 'storage', 'uniform'] as const) + .beginSubcases() + .combine('call', kCalls) + .combine('bufferSize', kBufferSizes) + .combine('dynamic_offset', [0, 256] as const) + .filter(t => { + if (t.type === 'array_S_2' || t.type === 'mat4x4f' || t.type === 'mat4x3f') { + return t.bufferSize > 128; + } + return t.aspace !== 'workgroup' || t.dynamic_offset === 0; + }) + ) + .fn(t => { + const testcase = kArrayLengthTypes[t.params.type]; + t.skipIfLanguageFeatureNotSupported('buffer_view'); + t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters'); + if (t.params.aspace === 'uniform' && testcase.uniformStdLayout === true) { + t.skipIfLanguageFeatureNotSupported('uniform_buffer_standard_layout'); + } + if (testcase.f16) { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + + const buffer_ty = `buffer<${t.params.bufferSize}>`; + let decl = ''; + switch (t.params.aspace) { + case 'workgroup': + decl = `var v : ${buffer_ty};\n@group(0) @binding(0) var dummy : array;`; + break; + case 'uniform': + decl = `@group(0) @binding(0) var v : ${buffer_ty};`; + break; + case 'storage': + decl = `@group(0) @binding(0) var v : ${buffer_ty};`; + break; + } + + const values: number[] = []; + const offsets: number[] = []; + const sizes: number[] = [0]; + for (const offset of kOffsets) { + const usedSize = t.params.call === 'unsized' ? t.params.bufferSize : t.params.bufferSize / 2; + // Skip any case that results in an invalid memory reference. + // Split the writes up one per thread. + if (isValidArrayLengthCase(testcase, offset, 0, usedSize)) { + values.push(calculateArrayLength(testcase, offset, 0, usedSize)); + // Setup the offset and size buffer values so they can be indexed by invocation id. + offsets.push(offset); + } + } + + assert(values.length > 0, 'no tests run'); + + const access = testcase.access ?? ''; + const addrOf = access === '' ? '' : '&'; + + const enables = testcase.f16 === true ? 'enable f16;' : ''; + + const wgsl = ` +${enables} + +${kStructDecls} + +fn unsizedLength(p : ptr<${t.params.aspace}, buffer>, gidx : u32) -> u32 { + return arrayLength(${addrOf}bufferView<${testcase.type}>(p, offsets[gidx])${access}); +} + +fn sizedLength(p : ptr<${t.params.aspace}, buffer<${t.params.bufferSize / 2}>>, gidx : u32) -> u32 { + return arrayLength(${addrOf}bufferView<${testcase.type}>(p, offsets[gidx])${access}); +} + +fn sized_indirectLength(p : ptr<${t.params.aspace}, buffer<${ + t.params.bufferSize / 2 + }>>, gidx : u32) -> u32 { + return unsizedLength(p, gidx); +} + +override bufferSize : u32 = 0u; +${decl} +@group(0) @binding(1) var offsets : array; +@group(0) @binding(2) var sizes : array; // unused +@group(0) @binding(3) var out : array; + +override wgx : u32; +@compute @workgroup_size(wgx) +fn main(@builtin(global_invocation_id) gid : vec3u) { + ${t.params.aspace === 'workgroup' ? '_ = dummy[0];' : ''} + + if gid.x >= ${values.length} { + return; + } + + out[gid.x] = ${t.params.call}Length(&v, gid.x); +}`; + + runLengthTest( + t, + wgsl, + t.params.aspace === 'uniform' ? GPUBufferUsage.UNIFORM : GPUBufferUsage.STORAGE, + t.params.bufferSize, + t.params.dynamic_offset, + offsets, + sizes, + values + ); + }); + +g.test('read_layout') + .desc('Test reading memory layout from a bufferView') + .params(u => + u + .combine('case', keysOf(kLayoutCases)) + .beginSubcases() + .combine('aspace', ['storage', 'ro_storage', 'uniform', 'workgroup'] as const) + .combine('offset', kOffsets) + .filter(t => { + const testcase = kLayoutCases[t.case]; + if ((t.offset & (testcase.align - 1)) !== 0) { + return false; + } + return testcase.offset + t.offset < 252; + }) + ) + .fn(t => { + const testcase = kLayoutCases[t.params.case]; + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (t.params.aspace === 'uniform' && testcase.uniformStdLayoutView === true) { + t.skipIfLanguageFeatureNotSupported('uniform_buffer_standard_layout'); + } + if (testcase.f16) { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + + const enables = testcase.f16 === true ? 'enable f16;' : ''; + const v = t.params.aspace === 'workgroup' ? 'wg_var' : 'in'; + + const wgsl = ` +${enables} + +${testcase.decl ?? ''} + +@group(0) @binding(0) +var<${t.params.aspace === 'uniform' ? 'uniform' : 'storage'}${ + t.params.aspace === 'storage' ? ', read_write' : '' + }> in : buffer<256 * 4>; + +@group(0) @binding(1) var out : u32; + +var wg_var : buffer<256 * 4>; + +@compute @workgroup_size(1) +fn main() { + let in_ptr = bufferView>(&in, 0); + let wg_ptr = bufferView>(&wg_var, 0); + for (var i = 0; i < 256 / 4; i++) { + (*wg_ptr)[i] = (*in_ptr)[i]; + } + + workgroupBarrier(); + + let p = bufferView<${testcase.type}>(&${v}, ${t.params.offset}); + out = u32((*p)${testcase.access}); +} +`; + + runReadLayoutTest(t, testcase, wgsl, t.params.aspace, t.params.offset); + }); + +g.test('write_layout') + .desc('Test writing memory layout via a bufferView') + .params(u => + u + .combine('case', keysOf(kLayoutCases)) + .beginSubcases() + .combine('assign', ['let', 'call'] as const) + .combine('aspace', ['storage', 'workgroup'] as const) + .combine('offset', kOffsets) + .filter(t => { + const testcase = kLayoutCases[t.case]; + if ((t.offset & (testcase.align - 1)) !== 0) { + return false; + } + return testcase.offset + t.offset < 252; + }) + ) + .fn(t => { + const testcase = kLayoutCases[t.params.case]; + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (testcase.f16) { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + + const enables = testcase.f16 === true ? 'enable f16;' : ''; + const v = t.params.aspace === 'workgroup' ? 'wg_var' : 'out'; + const call_expr = `bufferView<${testcase.type}>(&${v}, ${t.params.offset})`; + + let post_assign = ''; + if (t.params.aspace === 'workgroup') { + post_assign = ` + workgroupBarrier(); + + let out_ptr = bufferView>(&out, 0); + let wg_ptr = bufferView>(&wg_var, 0); + for (var i = 0; i < 256; i++) { + (*out_ptr)[i] = (*wg_ptr)[i]; + } +`; + } + + const wgsl = ` +${enables} + +${testcase.decl ?? ''} + +@group(0) @binding(0) var in : u32; +@group(0) @binding(1) var out : buffer<256 * 4>; + + +var wg_var : buffer<256 * 4>; + +@compute @workgroup_size(1) +fn main() { + let val = ${testcase.f16 ? 'f16' : testcase.f32 ? 'f32' : 'u32'}(in); + let ptr = ${call_expr}; + (*${t.params.assign === 'call' ? call_expr : 'ptr'})${testcase.access} = val; + + ${post_assign} +} +`; + + runWriteLayoutTest(t, testcase, wgsl, t.params.offset); + }); + +g.test('read') + .desc('Test reading various types from bufferView') + .params(u => + u + .combine('base_type', ['u32', 'i32', 'f32', 'f16'] as const) + .beginSubcases() + .combine('wrap', ['none', 'vector', 'array', 'matrix'] as const) + .combine('width', [1, 2, 3, 4] as const) + .combine('aspace', ['workgroup', 'storage', 'ro_storage', 'uniform'] as const) + .combine('offset', [0, 4, 8, 12, 16, 32, 48, 64] as const) + .filter(t => { + if (t.wrap !== 'none' && t.width === 1) { + return false; + } + if (t.wrap === 'none' && t.width !== 1) { + return false; + } + if (t.wrap === 'matrix' && t.base_type !== 'f32' && t.base_type !== 'f16') { + return false; + } + if (t.aspace === 'uniform' && t.wrap === 'array') { + return false; + } + const ty = Type[t.base_type]; + let align = ty.alignment; + switch (t.wrap) { + case 'vector': + align = Type.vec(t.width, ty).alignment; + break; + case 'array': + align = Type.array(2, ty).alignment; + break; + case 'matrix': + align = Type.mat(t.width, 2, ty).alignment; + break; + case 'none': + break; + } + if (t.offset % align !== 0) { + return false; + } + return true; + }) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (t.params.base_type === 'f16') { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + + const ele_ty = Type[t.params.base_type]; + let ty: Type = ele_ty; + switch (t.params.wrap) { + case 'vector': + ty = Type.vec(t.params.width, ele_ty); + break; + case 'array': + ty = Type.array(t.params.width, ele_ty); + break; + case 'matrix': + ty = Type.mat(t.params.width, 2, ele_ty); + break; + case 'none': + break; + } + + const bufferSize = 128; + const enables = t.params.base_type === 'f16' ? 'enable f16;' : ''; + const v = t.params.aspace === 'workgroup' ? 'wg_var' : 'in'; + const wgsl = ` +${enables} + + +@group(0) @binding(0) +var<${t.params.aspace === 'uniform' ? 'uniform' : 'storage'}${ + t.params.aspace === 'storage' ? ', read_write' : '' + }> in : buffer<${bufferSize}>; + +@group(0) @binding(1) var out : ${ty.toString()}; + +var wg_var : buffer<${bufferSize}>; + +@compute @workgroup_size(1) +fn main() { + let in_ptr = bufferView>(&in, 0); + let wg_ptr = bufferView>(&wg_var, 0); + for (var i = 0; i < 256 / 4; i++) { + (*wg_ptr)[i] = (*in_ptr)[i]; + } + + workgroupBarrier(); + + let p = bufferView<${ty.toString()}>(&${v}, ${t.params.offset}); + out = (*p); +} +`; + + runReadWriteTest(true, t, wgsl, ele_ty, ty, t.params.aspace, t.params.offset, bufferSize); + }); + +g.test('write') + .desc('Test writing various types via bufferView') + .params(u => + u + .combine('base_type', ['u32', 'i32', 'f32', 'f16'] as const) + .beginSubcases() + .combine('wrap', ['none', 'vector', 'array', 'matrix'] as const) + .combine('width', [1, 2, 3, 4] as const) + .combine('aspace', ['workgroup', 'storage'] as const) + .combine('offset', [0, 4, 8, 12, 16, 32, 48, 64] as const) + .combine('assign', ['let', 'call'] as const) + .combine('swizzle', [false, true] as const) + .filter(t => { + if (t.wrap !== 'none' && t.width === 1) { + return false; + } + if (t.wrap === 'none' && t.width !== 1) { + return false; + } + if (t.wrap === 'matrix' && t.base_type !== 'f32' && t.base_type !== 'f16') { + return false; + } + if (t.wrap !== 'vector' && t.swizzle) { + return false; + } + const ty = Type[t.base_type]; + let align = ty.alignment; + switch (t.wrap) { + case 'vector': + align = Type.vec(t.width, ty).alignment; + break; + case 'array': + align = Type.array(2, ty).alignment; + break; + case 'matrix': + align = Type.mat(t.width, 2, ty).alignment; + break; + case 'none': + break; + } + if (t.offset % align !== 0) { + return false; + } + return true; + }) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (t.params.base_type === 'f16') { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + if (t.params.swizzle) { + t.skipIfLanguageFeatureNotSupported('swizzle_assignment'); + } + + let swizzle = ''; + const ele_ty = Type[t.params.base_type]; + let ty: Type = ele_ty; + switch (t.params.wrap) { + case 'vector': + ty = Type.vec(t.params.width, ele_ty); + if (t.params.swizzle) { + swizzle = '.xyzw'.substring(0, 1 + t.params.width); + } + break; + case 'array': + ty = Type.array(t.params.width, ele_ty); + break; + case 'matrix': + ty = Type.mat(t.params.width, 2, ele_ty); + break; + case 'none': + break; + } + + const bufferSize = 128; + const enables = t.params.base_type === 'f16' ? 'enable f16;' : ''; + const v = t.params.aspace === 'workgroup' ? 'wg_var' : 'out'; + + let post_assign = ''; + if (t.params.aspace === 'workgroup') { + post_assign = ` + workgroupBarrier(); + + let out_ptr = bufferView>(&out, 0); + let wg_ptr = bufferView>(&wg_var, 0); + for (var i = 0; i < 256; i++) { + (*out_ptr)[i] = (*wg_ptr)[i]; + } +`; + } + + const call_expr = `bufferView<${ty.toString()}>(&${v}, ${t.params.offset})`; + const wgsl = ` +${enables} + +@group(0) @binding(0) var in : ${ty.toString()}; +@group(0) @binding(1) var out : buffer<${bufferSize}>; + + +var wg_var : buffer<${bufferSize}>; + +@compute @workgroup_size(1) +fn main() { + let ptr = ${call_expr}; + (*${t.params.assign === 'call' ? call_expr : 'ptr'})${swizzle} = in; + + ${post_assign} +} +`; + + runReadWriteTest(false, t, wgsl, ele_ty, ty, t.params.aspace, t.params.offset, bufferSize); + }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/buffer_view_utils.ts b/src/webgpu/shader/execution/expression/call/builtin/buffer_view_utils.ts new file mode 100644 index 000000000000..eca908c1d684 --- /dev/null +++ b/src/webgpu/shader/execution/expression/call/builtin/buffer_view_utils.ts @@ -0,0 +1,957 @@ +import { + iterRange, + typedArrayParam, + typedArrayFromParam, + TypedArrayBufferView, +} from '../../../../../../common/util/util.js'; +import { GPUTest } from '../../../../../gpu_test.js'; +import { + Type, + ScalarType, + ArrayType, + MatrixType, + VectorType, +} from '../../../../../util/conversion.js'; + +export const kBufferSizes = [128, 256, 512, 1024] as const; +export const kOffsets = [0, 1, 2, 3, 4, 8, 12, 16, 32, 100, 156, 480, 768] as const; +export const kSizes = [...kBufferSizes, 32, 48, 64] as const; +export const kCalls = ['unsized', 'sized', 'sized_indirect'] as const; + +interface ArrayLengthType { + type: string; + access?: string; + align: number; + stride: number; + arrayOffset?: number; + f16?: boolean; + uniformStdLayout?: boolean; +} + +export function isValidArrayLengthCase( + type: ArrayLengthType, + offset: number, + size: number, + bufferSize: number +): boolean { + if (size === 0) { + // bufferView case + return offset + (type.arrayOffset ?? 0) + type.stride < bufferSize; + } else { + // bufferArrayView case + if ((type.arrayOffset ?? 0) + type.stride > size) { + return false; + } + return offset + size < bufferSize; + } +} + +export function calculateArrayLength( + type: ArrayLengthType, + offset: number, + size: number, + bufferSize: number +): number { + const alignOffset = offset & ~(type.align - 1); + if (size === 0) { + // bufferView case + return Math.floor((bufferSize - alignOffset - (type.arrayOffset ?? 0)) / type.stride); + } else { + // bufferArrayView case + return Math.floor((size - (type.arrayOffset ?? 0)) / type.stride); + } +} + +export const kStructDecls = ` +struct S { + a: vec4u, + b: u32, +} +struct T { + a : u32, + b : array +} +struct S_arr { + a : vec4u, + b : array, +} +struct T_arr { + a : S, + b : array, +} +`; + +export const kArrayLengthTypes: Record = { + // Scalars and vectors + u32: { + type: 'array', + align: 4, + stride: 4, + uniformStdLayout: true, + }, + vec2f: { + type: 'array', + align: 8, + stride: 8, + uniformStdLayout: true, + }, + vec3i: { + type: 'array', + align: 16, + stride: 16, + }, + vec4u: { + type: 'array', + align: 16, + stride: 16, + }, + f16: { + type: 'array', + align: 2, + stride: 2, + f16: true, + uniformStdLayout: true, + }, + vec2h: { + type: 'array', + align: 4, + stride: 4, + f16: true, + uniformStdLayout: true, + }, + vec3h: { + type: 'array', + align: 8, + stride: 8, + f16: true, + uniformStdLayout: true, + }, + vec4h: { + type: 'array', + align: 8, + stride: 8, + f16: true, + uniformStdLayout: true, + }, + // Matrices + mat2x2f: { + type: 'array', + align: 8, + stride: 16, + }, + mat2x3f: { + type: 'array', + align: 16, + stride: 32, + }, + mat2x4f: { + type: 'array', + align: 16, + stride: 32, + }, + mat4x2f: { + type: 'array', + align: 8, + stride: 32, + }, + mat4x3f: { + type: 'array', + align: 16, + stride: 64, + }, + mat4x4f: { + type: 'array', + align: 16, + stride: 64, + }, + mat3x2h: { + type: 'array', + align: 4, + stride: 12, + f16: true, + }, + mat3x3h: { + type: 'array', + align: 8, + stride: 24, + f16: true, + uniformStdLayout: true, + }, + mat3x4h: { + type: 'array', + align: 8, + stride: 24, + f16: true, + uniformStdLayout: true, + }, + // Structs without arrays + S: { + type: 'array', + align: 16, + stride: 32, + }, + T: { + type: 'array', + align: 4, + stride: 24, + uniformStdLayout: true, + }, + // Sized arrays + array_u32_4: { + type: 'array>', + align: 4, + stride: 16, + uniformStdLayout: true, + }, + array_f16_6: { + type: 'array>', + align: 2, + stride: 12, + f16: true, + uniformStdLayout: true, + }, + array_vec2f_2: { + type: 'array>', + align: 8, + stride: 16, + uniformStdLayout: true, + }, + array_S_2: { + type: 'array>', + align: 16, + stride: 64, + }, + array_T_2: { + type: 'array>', + align: 4, + stride: 48, + }, + // Structs with runtime arrays + S_arr: { + type: 'S_arr', + access: '.b', + align: 16, + stride: 4, + arrayOffset: 16, + }, + T_arr: { + type: 'T_arr', + access: '.b', + align: 16, + stride: 24, + arrayOffset: 32, + }, +}; + +/** + * Run arrayLength tests for bufferView and bufferArrayView + * + * @param t The test + * @param wgsl The shader code. Interface requirements: + * * Overrides: + * * wgx: workgroup size x + * * bufferSize: size of the buffer + * * (0, 0): input buffer with 'usage' usage + * * (0, 1): read-only-storage-buffer for offsets + * * (0, 2): read-only-storage-buffer for sizes + * * (0, 3): storage-buffer for output + * @param usage The test buffer usage + * @param bufferSize The size of the test buffer + * @param dynOffset The size of the dynamic offset for the test buffer + * @param offsets The values for the offset buffer + * @param sizes The values for the size buffer + * @param values The expected results + */ +export function runLengthTest( + t: GPUTest, + wgsl: string, + usage: GPUBufferUsageFlags, + bufferSize: number, + dynOffset: number, + offsets: number[], + sizes: number[], + values: number[] +) { + const fullBufferSize = bufferSize + dynOffset; + const inputBuffer = t.createBufferTracked({ + size: fullBufferSize, + usage, + }); + const offsetBuffer = t.makeBufferWithContents( + new Uint32Array(offsets), + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + ); + const sizeBuffer = t.makeBufferWithContents( + new Uint32Array(sizes), + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + ); + const outputBuffer = t.createBufferTracked({ + size: values.length * 4, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + + const bgLayout = t.device.createBindGroupLayout({ + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: usage === GPUBufferUsage.UNIFORM ? 'uniform' : 'read-only-storage', + hasDynamicOffset: dynOffset !== 0, + }, + }, + { + binding: 1, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'read-only-storage', + }, + }, + { + binding: 2, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'read-only-storage', + }, + }, + { + binding: 3, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'storage', + }, + }, + ], + }); + + // Limit the number of invocations per workgroup to 128 to fit the default compat limits. Launch extra workgroups to fit all the required invocations. + const wgx = Math.min(values.length, 128); + const num_wgs = Math.ceil(values.length / wgx); + + const pipelineLayout = t.device.createPipelineLayout({ + bindGroupLayouts: [bgLayout], + }); + const pipeline = t.device.createComputePipeline({ + layout: pipelineLayout, + compute: { + module: t.device.createShaderModule({ code: wgsl }), + constants: { wgx, bufferSize }, + }, + }); + + const bg = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer: inputBuffer, + size: bufferSize, + }, + }, + { + binding: 1, + resource: { + buffer: offsetBuffer, + }, + }, + { + binding: 2, + resource: { + buffer: sizeBuffer, + }, + }, + { + binding: 3, + resource: { + buffer: outputBuffer, + }, + }, + ], + }); + + const encoder = t.device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup(0, bg, dynOffset === 0 ? [] : [dynOffset]); + pass.dispatchWorkgroups(num_wgs, 1, 1); + pass.end(); + t.queue.submit([encoder.finish()]); + + t.expectGPUBufferValuesEqual(outputBuffer, new Uint32Array(values)); +} + +interface LayoutCase { + type: string; + decl?: string; + access: string; + align: number; + offset: number; + f16?: boolean; + f32?: boolean; + uniformStdLayoutView?: boolean; + uniformStdLayoutArrayView?: boolean; +} + +export const kLayoutCases: Record = { + vec2u_align8: { + type: 'S_vec2u_align', + decl: 'struct S_vec2u_align { x : u32, y : vec2u, }', + access: '.y[1]', + align: 8, + offset: 12, + uniformStdLayoutArrayView: true, + }, + vec3u_align16: { + type: 'S_vec3u_align', + decl: 'struct S_vec3u_align { x : u32, y : vec3u, }', + access: '.y[2]', + align: 16, + offset: 24, + }, + vec4u_align16: { + type: 'S_vec4u_align', + decl: 'struct S_vec4u_align { x : u32, y : vec4u, }', + access: '.y[0]', + align: 16, + offset: 16, + }, + struct_align32: { + type: 'S_align32', + decl: 'struct S_align32 { x : u32, @align(32) y : u32, }', + access: '.y', + align: 32, + offset: 32, + }, + vec2h_align4: { + type: 'S_vec2h_align', + decl: 'struct S_vec2h_align { x : f16, y : vec2h }', + access: '.y[0]', + align: 4, + offset: 4, + f16: true, + uniformStdLayoutArrayView: true, + }, + vec3h_align8: { + type: 'S_vec3h_align', + decl: 'struct S_vec3h_align { x : f16, y : vec3h }', + access: '.y.z', + align: 8, + offset: 12, + f16: true, + uniformStdLayoutArrayView: true, + }, + vec4h_align8: { + type: 'S_vec4h_align', + decl: 'struct S_vec4h_align { x : f16, y : vec4h }', + access: '.y.z', + align: 8, + offset: 12, + f16: true, + uniformStdLayoutArrayView: true, + }, + vec3i_size12: { + type: 'S_vec3i_size', + decl: 'struct S_vec3i_size { x : vec3i, y : u32 }', + access: '.y', + align: 16, + offset: 12, + }, + vec3h_size6: { + type: 'S_vec3h_size', + decl: 'struct S_vec3h_size { x : vec3h, y : f16, z : f16 }', + access: '.z', + align: 8, + offset: 8, + f16: true, + uniformStdLayoutArrayView: true, + }, + size80: { + type: 'S_size80', + decl: 'struct S_size80 { @size(80) x : u32, y : u32 }', + access: '.y', + align: 4, + offset: 80, + uniformStdLayoutArrayView: true, + }, + mat2x2f_align8: { + type: 'S_mat2x2f_align', + decl: 'struct S_mat2x2f_align { x : u32, y : mat2x2f }', + access: '.y[0][0]', + align: 8, + offset: 8, + f32: true, + uniformStdLayoutArrayView: true, + }, + mat3x3f_align16: { + type: 'S_mat3x3f_align', + decl: 'struct S_mat3x3f_align { x : u32, y : mat3x3f }', + access: '.y[0][0]', + align: 16, + offset: 16, + f32: true, + }, + mat4x4f_align16: { + type: 'S_mat4x4f_align', + decl: 'struct S_mat4x4f_align { x : u32, y : mat4x4f }', + access: '.y[1][0]', + align: 16, + offset: 32, + f32: true, + }, + mat3x2h_align4: { + type: 'S_mat3x2h_align', + decl: 'struct S_mat3x2h_align { x : f16, y : mat3x2h }', + access: '.y[0][0]', + align: 4, + offset: 4, + f16: true, + uniformStdLayoutArrayView: true, + }, + mat4x3h_align8: { + type: 'S_mat4x3h_align', + decl: 'struct S_mat4x3h_align { x : f16, y : mat4x3h }', + access: '.y[0][0]', + align: 8, + offset: 8, + f16: true, + uniformStdLayoutArrayView: true, + }, + mat2x4h_align8: { + type: 'S_mat2x4h_align', + decl: 'struct S_mat2x4h_align { x : f16, y : mat2x4h }', + access: '.y[0][0]', + align: 8, + offset: 8, + f16: true, + uniformStdLayoutArrayView: true, + }, + mat2x2f_size: { + type: 'S_mat2x2f_size', + decl: 'struct S_mat2x2f_size { x : mat2x2f, y : u32 }', + access: '.y', + align: 8, + offset: 16, + uniformStdLayoutArrayView: true, + }, + mat3x2f_size: { + type: 'S_mat3x2f_size', + decl: 'struct S_mat3x2f_size { x : mat3x2f, y : u32 }', + access: '.y', + align: 8, + offset: 24, + uniformStdLayoutArrayView: true, + }, + mat2x3f_size: { + type: 'S_mat2x3f_size', + decl: 'struct S_mat2x3f_size { x : mat2x3f, y : u32 }', + access: '.y', + align: 16, + offset: 32, + }, + mat3x3f_size: { + type: 'S_mat3x3f_size', + decl: 'struct S_mat3x3f_size { x : mat3x3f, y : u32 }', + access: '.y', + align: 16, + offset: 48, + }, + mat2x4f_size: { + type: 'S_mat2x4f_size', + decl: 'struct S_mat2x4f_size { x : mat2x4f, y : u32 }', + access: '.y', + align: 16, + offset: 32, + }, + mat3x4f_size: { + type: 'S_mat3x4f_size', + decl: 'struct S_mat3x4f_size { x : mat3x4f, y : u32 }', + access: '.y', + align: 16, + offset: 48, + }, + mat2x2h_size: { + type: 'S_mat2x2h_size', + decl: 'struct S_mat2x2h_size { x : mat2x2h, y : f16 }', + access: '.y', + align: 4, + offset: 8, + f16: true, + uniformStdLayoutArrayView: true, + }, + mat4x2h_size: { + type: 'S_mat4x2h_size', + decl: 'struct S_mat4x2h_size { x : mat4x2h, y : f16 }', + access: '.y', + align: 4, + offset: 16, + f16: true, + uniformStdLayoutArrayView: true, + }, + mat2x3h_size: { + type: 'S_mat2x3h_size', + decl: 'struct S_mat2x3h_size { x : mat2x3h, y : f16 }', + access: '.y', + align: 8, + offset: 16, + f16: true, + uniformStdLayoutArrayView: true, + }, + mat4x3h_size: { + type: 'S_mat4x3h_size', + decl: 'struct S_mat4x3h_size { x : mat4x3h, y : f16 }', + access: '.y', + align: 8, + offset: 32, + f16: true, + uniformStdLayoutArrayView: true, + }, + mat2x4h_size: { + type: 'S_mat2x4h_size', + decl: 'struct S_mat2x4h_size { x : mat2x4h, y : f16 }', + access: '.y', + align: 8, + offset: 16, + f16: true, + uniformStdLayoutArrayView: true, + }, + mat4x4h_size: { + type: 'S_mat4x4h_size', + decl: 'struct S_mat4x4h_size { x : mat4x4h, y : f16 }', + access: '.y', + align: 8, + offset: 32, + f16: true, + uniformStdLayoutArrayView: true, + }, + struct_size_roundup: { + type: 'S_struct_size_roundup', + decl: `struct Inner { x : vec3u } + struct S_struct_size_roundup { x : Inner, y : u32 }`, + access: '.y', + align: 16, + offset: 16, + }, + struct_inner_size: { + type: 'S_struct_inner_size', + decl: `struct Inner { @size(112) x : u32 } + struct S_struct_inner_size { x : Inner, y : u32 }`, + access: '.y', + align: 4, + offset: 112, + uniformStdLayoutArrayView: true, + }, + struct_inner_align: { + type: 'S_struct_inner_align', + decl: `struct Inner { @align(64) x : u32 } + struct S_struct_inner_align { x : Inner, y : u32 }`, + access: '.y', + align: 64, + offset: 64, + }, + struct_inner_size_and_align: { + type: 'S_struct_inner_size_and_align', + decl: `struct Inner { @align(32) @size(33) x : u32 } + struct S_struct_inner_size_and_align { x : Inner, y : u32 }`, + access: '.y', + align: 32, + offset: 64, + }, + struct_override_size: { + type: 'S_struct_override_size', + decl: `struct Inner { @size(32) x : u32 } + struct S_struct_override_size { @size(64) x : Inner, y : u32 }`, + access: '.y', + align: 4, + offset: 64, + uniformStdLayoutArrayView: true, + }, + struct_double_align: { + type: 'S_struct_double_align', + decl: `struct Inner { x : u32, @align(32) y : u32 } + struct S_struct_double_align { x : u32, @align(64) y : Inner }`, + access: '.y.y', + align: 64, + offset: 96, + }, + array_stride_size: { + type: 'array', + decl: 'struct S_stride { @size(16) x : u32 }', + access: '[2].x', + align: 4, + offset: 32, + }, +}; + +/** + * Runs a read test for memory layouts from bufferView and bufferArrayView + * + * @param t The GPUTest + * @param testcase the LayoutCase + * @param wgsl the shader code + * * (0, 0) - uniform or storage buffer + * * (0, 1) - storage buffer + * @param aspace the address space being tested + * @param offset the offset to the view call + */ +export function runReadLayoutTest( + t: GPUTest, + testcase: LayoutCase, + wgsl: string, + aspace: string, + offset: number +) { + let usage = GPUBufferUsage.COPY_DST; + if (aspace === 'uniform') { + usage |= GPUBufferUsage.UNIFORM; + } else { + usage |= GPUBufferUsage.STORAGE; + } + + const kBufferLength = 256; + // Magic number is 42 in various representations. + const inMagicNumber = testcase.f16 ? 0x5140 : testcase.f32 ? 0x42280000 : 42; + const inputBuffer = t.makeBufferWithContents( + new Uint32Array([ + ...iterRange(kBufferLength, x => { + if (x * 4 === testcase.offset + offset) { + return inMagicNumber; + } else { + return 0; + } + }), + ]), + usage + ); + + const outputBuffer = t.createBufferTracked({ + size: 4, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ code: wgsl }), + }, + }); + + const bg = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer: inputBuffer, + }, + }, + { + binding: 1, + resource: { + buffer: outputBuffer, + }, + }, + ], + }); + + 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()]); + + t.expectGPUBufferValuesEqual(outputBuffer, new Uint32Array([42])); +} + +/** + * Runs a write test for memory layout from bufferView and bufferArrayView + * + * @param t The GPUTest + * @param testcase the LayoutCase + * @param wgsl The shader code + * * (0, 0) - a read-only storage buffer + * * (0, 1) - a read-write storage buffer + * @param offset The offset for the view call + */ +export function runWriteLayoutTest(t: GPUTest, testcase: LayoutCase, wgsl: string, offset: number) { + const inputBuffer = t.makeBufferWithContents( + new Uint32Array([42]), + GPUBufferUsage.COPY_SRC | GPUBufferUsage.STORAGE + ); + + const outputBuffer = t.createBufferTracked({ + size: 256 * 4, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ code: wgsl }), + }, + }); + + const bg = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer: inputBuffer, + }, + }, + { + binding: 1, + resource: { + buffer: outputBuffer, + }, + }, + ], + }); + + 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()]); + + // Magic number is 42 in various representations. + const outMagicNumber = testcase.f16 ? 0x5140 : testcase.f32 ? 0x42280000 : 42; + const expect = new Uint32Array([ + ...iterRange(128, x => { + if (x * 4 === testcase.offset + offset) { + return outMagicNumber; + } else { + return 0; + } + }), + ]); + t.expectGPUBufferValuesEqual(outputBuffer, expect); +} + +/** + * Performs a type-based read or write test for bufferView and bufferArrayView + * + * Note: This only covers types with no padding. Padding is covered by the layout tests. + * @param read Test read if true + * @param t The GPUTest + * @param wgsl The shader code + * @param eleTy The base element type + * @param ty The test type + * @param aspace The address space to test + * @param offset The offset for the call + * @param bufferSize The test buffer size + */ +export function runReadWriteTest( + read: boolean, + t: GPUTest, + wgsl: string, + eleTy: ScalarType, + ty: ScalarType | VectorType | MatrixType | ArrayType, + aspace: string, + offset: number, + bufferSize: number +) { + let num_eles = 1; + if (ty instanceof VectorType) { + num_eles *= ty.width; + } else if (ty instanceof ArrayType) { + num_eles *= ty.count; + } else if (ty instanceof MatrixType) { + num_eles *= ty.rows * ty.cols; + } + + const bufferElements = bufferSize / eleTy.size; + + const start = offset / eleTy.size; + const end = start + num_eles; + const ins: number[] = []; + const outs: number[] = []; + for (let i = 0; i < bufferElements; i++) { + if (i >= start && i < end) { + ins.push(i + 4); + outs.push(i + 4); + } else { + if (read) { + ins.push(0); + } else { + outs.push(0); + } + } + } + if (!read && eleTy === Type.f16 && ins.length % 2 === 1) { + ins.push(0); + } + if (read && eleTy === Type.f16 && outs.length % 2 === 1) { + outs.push(0); + } + + const inMap = { + u32: typedArrayParam('Uint32Array', ins), + i32: typedArrayParam('Int32Array', ins), + f32: typedArrayParam('Float32Array', ins), + f16: typedArrayParam('Float16Array', ins), + }; + const inputData: TypedArrayBufferView = typedArrayFromParam( + inMap[eleTy.toString() as 'u32' | 'i32' | 'f32' | 'f16'] + ); + + const outMap = { + u32: typedArrayParam('Uint32Array', outs), + i32: typedArrayParam('Int32Array', outs), + f32: typedArrayParam('Float32Array', outs), + f16: typedArrayParam('Float16Array', outs), + }; + const outputData: TypedArrayBufferView = typedArrayFromParam( + outMap[eleTy.toString() as 'u32' | 'i32' | 'f32' | 'f16'] + ); + + const inputBuffer = t.makeBufferWithContents( + inputData, + GPUBufferUsage.COPY_SRC | + (aspace === 'uniform' ? GPUBufferUsage.UNIFORM : GPUBufferUsage.STORAGE) + ); + const outputBuffer = t.createBufferTracked({ + size: outs.length * eleTy.size, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ code: wgsl }), + }, + }); + + const bg = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer: inputBuffer, + }, + }, + { + binding: 1, + resource: { + buffer: outputBuffer, + }, + }, + ], + }); + + 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()]); + + t.expectGPUBufferValuesEqual(outputBuffer, outputData); +} diff --git a/src/webgpu/shader/validation/decl/let.spec.ts b/src/webgpu/shader/validation/decl/let.spec.ts index 3c1ec06d4bd5..ecae282f2a62 100644 --- a/src/webgpu/shader/validation/decl/let.spec.ts +++ b/src/webgpu/shader/validation/decl/let.spec.ts @@ -10,7 +10,7 @@ export const g = makeTestGroup(ShaderValidationTest); interface Case { code: string; - valid: boolean | 'texture_and_sampler_let'; + valid: boolean | 'texture_and_sampler_let' | 'buffer_view'; decls?: string; } @@ -109,6 +109,36 @@ const kTypeCases: Record = { valid: 'texture_and_sampler_let', decls: `@group(0) @binding(0) var samp_comp : sampler_comparison;`, }, + buffer: { + code: `let x : buffer = b;`, + valid: false, + decls: `@group(0) @binding(0) var b : buffer;`, + }, + buffer_sized: { + code: `let x = buffer<128> = b;`, + valid: false, + decls: `@group(0) @binding(0) var b : buffer<128>;`, + }, + buffer_override_sized: { + code: `let x : buffer = b;`, + valid: false, + decls: `override o = 16u; var b : buffer;`, + }, + ptr_buffer: { + code: `let x : ptr = &b;`, + valid: 'buffer_view', + decls: `@group(0) @binding(0) var b : buffer;`, + }, + ptr_buffer_sized: { + code: `let x : ptr> = &b;`, + valid: 'buffer_view', + decls: `@group(0) @binding(0) var b : buffer<128>;`, + }, + ptr_buffer_override_sized: { + code: `let x : ptr> = &b;`, + valid: 'buffer_view', + decls: `override o = 16u; var b : buffer;`, + }, }; g.test('type') @@ -129,6 +159,8 @@ fn foo() { let expect: boolean = testcase.valid === true; if (testcase.valid === 'texture_and_sampler_let') { expect = t.hasLanguageFeature('texture_and_sampler_let'); + } else if (testcase.valid === 'buffer_view') { + expect = t.hasLanguageFeature('buffer_view'); } t.expectCompileResult(expect, code); }); @@ -178,6 +210,16 @@ const kInitCases: Record = { code: `var x = 1;\nlet y = x << 1;`, valid: true, }, + buffer_mismatch: { + code: `let x : ptr = &b;`, + valid: false, + decls: `@group(0) @binding(0) var b : buffer<128>;`, + }, + buffer_smaller_size: { + code: `let x : ptr> = &b;`, + valid: false, + decls: `@group(0) @binding(0) var b : buffer<128>;`, + }, }; g.test('initializer') diff --git a/src/webgpu/shader/validation/decl/var.spec.ts b/src/webgpu/shader/validation/decl/var.spec.ts index a4afc6ab84e5..1ecbaf60a7b6 100644 --- a/src/webgpu/shader/validation/decl/var.spec.ts +++ b/src/webgpu/shader/validation/decl/var.spec.ts @@ -26,30 +26,35 @@ const kTypes = { isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, i32: { isHostShareable: true, isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, u32: { isHostShareable: true, isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, f32: { isHostShareable: true, isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, f16: { isHostShareable: true, isConstructible: true, isFixedFootprint: true, requiresF16: true, + requiresBufferView: false, }, // Vectors. @@ -58,30 +63,35 @@ const kTypes = { isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, vec3i: { isHostShareable: true, isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, vec4u: { isHostShareable: true, isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, vec2f: { isHostShareable: true, isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, vec3h: { isHostShareable: true, isConstructible: true, isFixedFootprint: true, requiresF16: true, + requiresBufferView: false, }, // Matrices. @@ -90,12 +100,14 @@ const kTypes = { isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, mat3x4h: { isHostShareable: true, isConstructible: true, isFixedFootprint: true, requiresF16: true, + requiresBufferView: false, }, // Atomics. @@ -104,12 +116,14 @@ const kTypes = { isConstructible: false, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, 'atomic': { isHostShareable: true, isConstructible: false, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, // Arrays. @@ -118,36 +132,42 @@ const kTypes = { isConstructible: false, isFixedFootprint: false, requiresF16: false, + requiresBufferView: false, }, 'array, 4>': { isHostShareable: false, isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, 'array': { isHostShareable: true, isConstructible: false, isFixedFootprint: false, requiresF16: false, + requiresBufferView: false, }, 'array': { isHostShareable: true, isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, 'array': { isHostShareable: true, isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, 'array': { isHostShareable: false, isConstructible: false, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, // Structures. @@ -156,36 +176,42 @@ const kTypes = { isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, S_bool: { isHostShareable: false, isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, S_S_bool: { isHostShareable: false, isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, S_array_vec4u: { isHostShareable: true, isConstructible: false, isFixedFootprint: false, requiresF16: false, + requiresBufferView: false, }, S_array_vec4u_4: { isHostShareable: true, isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, S_array_bool_4: { isHostShareable: false, isConstructible: true, isFixedFootprint: true, requiresF16: false, + requiresBufferView: false, }, // Misc. @@ -194,18 +220,44 @@ const kTypes = { isConstructible: false, isFixedFootprint: false, requiresF16: false, + requiresBufferView: false, }, sampler: { isHostShareable: false, isConstructible: false, isFixedFootprint: false, requiresF16: false, + requiresBufferView: false, }, 'texture_2d': { isHostShareable: false, isConstructible: false, isFixedFootprint: false, requiresF16: false, + requiresBufferView: false, + }, + + // Buffers + buffer: { + isHostShareable: true, + isConstructible: false, + isFixedFootprint: false, + requiresF16: false, + requiresBufferView: true, + }, + 'buffer<128>': { + isHostShareable: true, + isConstructible: false, + isFixedFootprint: true, + requiresF16: false, + requiresBufferView: true, + }, + 'buffer': { + isHostShareable: false, + isConstructible: false, + isFixedFootprint: true, + requiresF16: false, + requiresBufferView: true, }, }; @@ -259,7 +311,8 @@ g.test('module_scope_types') break; case 'uniform': decl = '@group(0) @binding(0) var foo : '; - shouldPass = type.isHostShareable && type.isConstructible; + shouldPass = + type.isHostShareable && (type.isConstructible || t.params.type === 'buffer<128>'); break; case 'workgroup': decl = 'var foo : '; @@ -283,6 +336,7 @@ g.test('module_scope_types') ${decl} ${t.params.via_alias ? 'MyType' : t.params.type}; `; + shouldPass &&= !kTypes[t.params.type].requiresBufferView || t.hasLanguageFeature('buffer_view'); t.expectCompileResult(shouldPass, wgsl); }); @@ -332,6 +386,7 @@ g.test('function_scope_types') ${decl} ${t.params.via_alias ? 'MyType' : t.params.type}; }`; + shouldPass &&= !kTypes[t.params.type].requiresBufferView || t.hasLanguageFeature('buffer_view'); t.expectCompileResult(shouldPass, wgsl); }); diff --git a/src/webgpu/shader/validation/expression/call/builtin/bufferArrayView.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/bufferArrayView.spec.ts new file mode 100644 index 000000000000..e78a60e0fb34 --- /dev/null +++ b/src/webgpu/shader/validation/expression/call/builtin/bufferArrayView.spec.ts @@ -0,0 +1,524 @@ +export const description = ` +Validation tests for bufferArrayView +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js'; +import { WGSLLanguageFeature } from '../../../../../capability_info.js'; +import { Type, elementTypeOf, kAllScalarsAndVectors } from '../../../../../util/conversion.js'; +import { ShaderValidationTest } from '../../../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +g.test('must_use') + .desc('Tests that the builtin has the @must_use attribute') + .params(u => u.combine('must_use', [true, false] as const)) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + const wgsl = ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + ${t.params.must_use ? 'let p = ' : ''}bufferArrayView>(&v, 0, 64); +}`; + + t.expectCompileResult(t.params.must_use, wgsl); + }); + +const kTypes = objectsToRecord(kAllScalarsAndVectors); + +g.test('return_type') + .desc('Validates return type') + .params(u => + u + .combine('template_type', keysOf(kTypes)) + .filter(t => { + const type = kTypes[t.template_type]; + const eleType = elementTypeOf(type); + return ( + eleType !== Type.abstractInt && eleType !== Type.abstractFloat && eleType !== Type.bool + ); + }) + .combine('return_type', keysOf(kTypes)) + .filter(t => { + const type = kTypes[t.return_type]; + const eleType = elementTypeOf(type); + return ( + eleType !== Type.abstractInt && eleType !== Type.abstractFloat && eleType !== Type.bool + ); + }) + .beginSubcases() + .combine('ptr', [false, true] as const) + .combine('array', [false, true] as const) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + const template_type = kTypes[t.params.template_type]; + const return_type = kTypes[t.params.return_type]; + let ret_type = return_type.toString(); + let temp_type = template_type.toString(); + if (t.params.array) { + ret_type = `array<${return_type.toString()}>`; + temp_type = `array<${temp_type}>`; + } + if (t.params.ptr) { + ret_type = `ptr`; + } + let enables = ``; + if (template_type.requiresF16() || return_type.requiresF16()) { + enables = `enable f16;`; + } + const wgsl = ` +${enables} +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let res : ${ret_type} = bufferArrayView<${temp_type}>(&v, 0, 128); +}`; + + t.expectCompileResult(t.params.ptr && t.params.array && template_type === return_type, wgsl); + }); + +g.test('buffer_type') + .desc('Validates the buffer parameter type') + .params(u => + u + .combine('type', [ + 'unsized_ro_storage', + 'unsized_storage', + 'sized_ro_storage', + 'sized_storage', + 'sized_uniform', + 'sized_workgroup', + 'override_workgroup', + ] as const) + .beginSubcases() + .combine('ptr', [false, true] as const) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + const wgsl = ` +@group(0) @binding(0) var unsized_ro_storage : buffer; +@group(0) @binding(1) var unsized_storage : buffer; +@group(0) @binding(2) var sized_ro_storage : buffer<128>; +@group(0) @binding(3) var sized_storage : buffer<128>; +@group(0) @binding(4) var sized_uniform : buffer<128>; +var sized_workgroup : buffer<128>; +override o : u32; +var override_workgroup : buffer; + +@compute @workgroup_size(1) +fn main() { + let p = bufferArrayView>(${t.params.ptr ? '&' : ''}${t.params.type}, 0, 128); +}`; + + t.expectCompileResult(t.params.ptr, wgsl); + }); + +g.test('offset_type') + .desc('Validates the offset parameter type') + .params(u => u.combine('type', keysOf(kTypes))) + .fn(t => { + const type = kTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable f16;`; + } + const wgsl = ` +${enables} +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let p = bufferArrayView>(&v, ${type.create(0).wgsl()}, 128); +}`; + + t.expectCompileResult( + type === Type.abstractInt || type === Type.u32 || type === Type.i32, + wgsl + ); + }); + +g.test('size_type') + .desc('Validates the offset parameter type') + .params(u => u.combine('type', keysOf(kTypes))) + .fn(t => { + const type = kTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable f16;`; + } + const wgsl = ` +${enables} +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let p = bufferArrayView>(&v, 0, ${type.create(128).wgsl()}); +}`; + + t.expectCompileResult( + type === Type.abstractInt || type === Type.u32 || type === Type.i32, + wgsl + ); + }); + +interface EarlyEvalCase { + code: string; + valid: boolean | 'pipeline'; + constants?: Record; + ptr_param?: boolean; +} + +const kEarlyEvalCases: Record = { + ro_storage_buffer: { + code: ` +@group(0) @binding(0) var v : buffer<16>; +fn foo() { + let p = bufferArrayView>(&v, 0, 16); +}`, + valid: true, + }, + storage_buffer: { + code: ` +@group(0) @binding(0) var v : buffer<16>; +fn foo() { + let p = bufferArrayView>(&v, 0, 16); +}`, + valid: true, + }, + uniform_buffer: { + code: ` +@group(0) @binding(0) var v : buffer<16>; +fn foo() { + let p = bufferArrayView>(&v, 0, 16); +}`, + valid: true, + }, + workgroup_buffer: { + code: ` +var v : buffer<16>; +fn foo() { + let p = bufferArrayView>(&v, 0, 16); +}`, + valid: true, + }, + const_size_too_small_for_type: { + code: ` +@group(0) @binding(0) var v : buffer; +fn foo() { + let p = bufferArrayView>(&v, 0, 12); +}`, + valid: false, + }, + buffer_too_small_for_const_size: { + code: ` +@group(0) @binding(0) var v : buffer<12>; +fn foo() { + let p = bufferArrayView>(&v, 0, 16); +}`, + valid: false, + }, + buffer_too_small_for_type: { + code: ` +@group(0) @binding(0) var v : buffer<12>; +fn foo() { + let p = bufferArrayView>(&v, 0, 16); +}`, + valid: false, + }, + const_size_too_small_for_const_offset_and_type: { + code: ` +var v : buffer<128>; +fn foo() { + let p = bufferArrayView>(&v, 16, 30); +}`, + valid: false, + }, + buffer_size_too_small_for_const_offset_and_type: { + code: ` +var v : buffer<30>; +fn foo() { + let p = bufferArrayView>(&v, 16, 128); +}`, + valid: false, + }, + override_size_too_small_for_type: { + code: ` +override size : u32; +@group(0) @binding(0) var v : buffer; +fn foo() { + let p = bufferArrayView>(&v, 0, size); +}`, + valid: 'pipeline', + constants: { size: 12 }, + }, + buffer_too_small_for_override_size: { + code: ` +override buffer_size : u32; +override size : u32; +var v : buffer; +fn foo() { + let p = bufferArrayView>(&v, 0, size); +}`, + valid: 'pipeline', + constants: { buffer_size: 12, size: 16 }, + }, + override_buffer_too_small_for_type: { + code: ` +override buffer_size : u32; +override size : u32; +var v : buffer; +fn foo() { + let p = bufferArrayView>(&v, 0, size); +}`, + valid: 'pipeline', + constants: { buffer_size: 12, size: 16 }, + }, + override_size_too_small_for_override_offset_and_type: { + code: ` +override offset : u32; +override size : u32; +@group(0) @binding(0) var v : buffer<128>; +fn foo() { + let p = bufferArrayView>(&v, offset, size); +}`, + valid: 'pipeline', + constants: { offset: 16, size: 30 }, + }, + buffer_size_too_small_for_override_offset_and_type: { + code: ` +override buffer_size : u32; +override offset : u32; +var v : buffer; +fn foo() { + let p = bufferArrayView>(&v, offset, 128); +}`, + valid: 'pipeline', + constants: { buffer_size: 30, offset: 16 }, + }, + const_offset_not_aligned: { + code: ` +@group(0) @binding(0) var v : buffer; +fn foo() { + let p = bufferarrayview>(&v, 12, 128); +}`, + valid: false, + }, + override_offset_not_aligned: { + code: ` +override offset : u32; +@group(0) @binding(0) var v : buffer; +fn foo() { + let p = bufferarrayview>(&v, offset, 128); +}`, + valid: false, + constants: { offset: 12 }, + }, + const_offset_plus_type_size_out_of_range: { + code: ` +@group(0) @binding(0) var v : buffer; +struct S { + a: vec4u, + b: vec4f, +} +fn foo() { + let p = buffer>(&v, 4294967279, 128); +}`, + valid: false, + }, + override_offset_plus_type_size_out_of_range: { + code: ` +override offset : u32; +@group(0) @binding(0) var v : buffer; +struct S { + a: vec4u, + b: vec4f, +} +fn foo() { + let p = buffer>(&v, offset, 128); +}`, + valid: false, + constants: { offset: 4294967279 }, + }, + const_multiple_of_stride: { + code: ` +struct S { + a : vec4u, + b : array, +} +var v : buffer<256>; +fn foo() { + let p = bufferArrayView(&v, 0, 24); +}`, + valid: true, + }, + const_not_multiple_of_stride: { + code: ` +struct S { + a : vec4u, + b : array, +} +var v : buffer<256>; +fn foo() { + let p = bufferArrayView(&v, 0, 26); +}`, + valid: false, + }, + override_not_multiple_of_stride: { + code: ` +override size : u32; +struct S { + a : vec4u, + b : array, +} +var v : buffer<256>; +fn foo() { + let p = bufferArrayView(&v, 0, size); +}`, + valid: 'pipeline', + constants: { size: 26 }, + }, + const_buffer_too_small_through_unsized_function: { + code: ` +@group(0) @binding(0) var buffer<16>; +fn bar(p : ptr) { + let q = bufferArrayView>(p, 0, 32); +} +fn foo() { + bar(&v); +}`, + valid: false, + ptr_param: true, + }, + override_buffer_too_small_through_unsized_function: { + code: ` +override size : u32; +@group(0) @binding(0) var v : buffer<16>; +fn bar(p : ptr) { + let q = bufferArrayView>(p, 0, size); +} +fn foo() { + bar(&v); +}`, + valid: 'pipeline', + constants: { size: 32 }, + ptr_param: true, + }, + sized_buffer_to_smaller_buffer_too_small: { + code: ` +var v : buffer<256>; +fn bar(p : ptr>) { + let q = bufferArrayView>(q, 16, 4); +} +fn foo() { + bar(&v); +}`, + valid: false, + ptr_param: true, + }, +}; + +g.test('early_eval_errors') + .desc('Test shader-creation and pipeline-creation errors') + .params(u => u.combine('case', keysOf(kEarlyEvalCases))) + .fn(t => { + const testcase = kEarlyEvalCases[t.params.case]; + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (testcase.ptr_param === true) { + t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters'); + } + + t.expectCompileResult(testcase.valid !== false, testcase.code); + if (testcase.valid !== false) { + t.expectPipelineResult({ + expectedResult: testcase.valid === true, + code: testcase.code, + constants: testcase.constants, + statements: ['foo();'], + }); + } + }); + +interface LHSCallCase { + code: string; + requires?: WGSLLanguageFeature[]; +} + +const kLHSCallCases: Record = { + u32_array: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + (*bufferArrayView>(&v, 0, 128))[0] = 123u; +}`, + }, + f32_vector_letter: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + (*bufferArrayView>(&v, 0, 128))[0].x = 13.3f; +}`, + }, + f32_struct_access: { + code: ` +struct S { + member_a: f32, +} +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + (*bufferArrayView>(&v, 0, 128))[0].member_a = 13.3f; +}`, + }, + u32_array_ptr_composite: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + bufferArrayView>(&v, 0, 128)[0] = 123u; +}`, + requires: ['pointer_composite_access'], + }, + vec4u_swizzle_assign: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + (*bufferArrayView>(&v, 0, 128))[0].zxy = vec3u(42); +}`, + requires: ['swizzle_assignment'], + }, + vec4u_ptr_composite_swizzle_assign: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + bufferArrayView>(&v, 0, 128)[0].zxy = vec3u(42); +}`, + requires: ['pointer_composite_access', 'swizzle_assignment'], + }, + compound_assign: { + code: ` +var v : buffer<256>; +@compute @workgroup_size(1) +fn main() { + (*bufferArrayView>(&v, 0, 128))[0] += 123u; +}`, + }, +}; + +g.test('lhs_call') + .desc('Validate that bufferView can be on the LHS of an assignment') + .params(u => u.combine('case', keysOf(kLHSCallCases))) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + + const testcase = kLHSCallCases[t.params.case]; + let valid = true; + const features = testcase.requires ?? []; + features.forEach(f => { + valid &&= t.hasLanguageFeature(f); + }); + t.expectCompileResult(valid, testcase.code); + }); diff --git a/src/webgpu/shader/validation/expression/call/builtin/bufferLength.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/bufferLength.spec.ts new file mode 100644 index 000000000000..706fcc1a1be9 --- /dev/null +++ b/src/webgpu/shader/validation/expression/call/builtin/bufferLength.spec.ts @@ -0,0 +1,90 @@ +export const description = ` +Validation tests for bufferLength +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js'; +import { Type, elementTypeOf, kAllScalarsAndVectors } from '../../../../../util/conversion.js'; +import { ShaderValidationTest } from '../../../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +g.test('must_use') + .desc('Tests that the builtin has the @must_use attribute') + .params(u => u.combine('must_use', [true, false] as const)) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + const wgsl = ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + ${t.params.must_use ? '_ = ' : ''}bufferLength(&v); +}`; + + t.expectCompileResult(t.params.must_use, wgsl); + }); + +const kTypes = objectsToRecord(kAllScalarsAndVectors); + +g.test('return_type') + .desc('Validates return type') + .params(u => + u.combine('type', keysOf(kTypes)).filter(t => { + const type = kTypes[t.type]; + const eleType = elementTypeOf(type); + return eleType !== Type.abstractInt && eleType !== Type.abstractFloat; + }) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + const type = kTypes[t.params.type]; + let enables = ``; + if (type.requiresF16()) { + enables = `enable f16;`; + } + const wgsl = ` +${enables} +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let res : ${type.toString()} = bufferLength(&v); +}`; + + t.expectCompileResult(type === Type.u32, wgsl); + }); + +g.test('data_type') + .desc('Validates the input parameter type') + .params(u => + u + .combine('type', [ + 'unsized_ro_storage', + 'unsized_storage', + 'sized_ro_storage', + 'sized_storage', + 'sized_uniform', + 'sized_workgroup', + 'override_workgroup', + ] as const) + .beginSubcases() + .combine('ptr', [false, true] as const) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + const wgsl = ` +@group(0) @binding(0) var unsized_ro_storage : buffer; +@group(0) @binding(1) var unsized_storage : buffer; +@group(0) @binding(2) var sized_ro_storage : buffer<128>; +@group(0) @binding(3) var sized_storage : buffer<128>; +@group(0) @binding(4) var sized_uniform : buffer<128>; +var sized_workgroup : buffer<128>; +override o : u32; +var override_workgroup : buffer; + +@compute @workgroup_size(1) +fn main() { + _ = bufferLength(${t.params.ptr ? '&' : ''}${t.params.type}); +}`; + + t.expectCompileResult(t.params.ptr, wgsl); + }); diff --git a/src/webgpu/shader/validation/expression/call/builtin/bufferView.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/bufferView.spec.ts new file mode 100644 index 000000000000..13c7b771adf6 --- /dev/null +++ b/src/webgpu/shader/validation/expression/call/builtin/bufferView.spec.ts @@ -0,0 +1,435 @@ +export const description = ` +Validation tests for bufferView +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js'; +import { WGSLLanguageFeature } from '../../../../../capability_info.js'; +import { Type, elementTypeOf, kAllScalarsAndVectors } from '../../../../../util/conversion.js'; +import { ShaderValidationTest } from '../../../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +g.test('must_use') + .desc('Tests that the builtin has the @must_use attribute') + .params(u => u.combine('must_use', [true, false] as const)) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + const wgsl = ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + ${t.params.must_use ? 'let p = ' : ''}bufferView(&v, 0); +}`; + + t.expectCompileResult(t.params.must_use, wgsl); + }); + +const kTypes = objectsToRecord(kAllScalarsAndVectors); + +g.test('return_type') + .desc('Validates return type') + .params(u => + u + .combine('template_type', keysOf(kTypes)) + .filter(t => { + const type = kTypes[t.template_type]; + const eleType = elementTypeOf(type); + return ( + eleType !== Type.abstractInt && eleType !== Type.abstractFloat && eleType !== Type.bool + ); + }) + .combine('return_type', keysOf(kTypes)) + .filter(t => { + const type = kTypes[t.return_type]; + const eleType = elementTypeOf(type); + return ( + eleType !== Type.abstractInt && eleType !== Type.abstractFloat && eleType !== Type.bool + ); + }) + .beginSubcases() + .combine('ptr', [false, true] as const) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + const template_type = kTypes[t.params.template_type]; + const return_type = kTypes[t.params.return_type]; + let ret_type = return_type.toString(); + if (t.params.ptr) { + ret_type = `ptr`; + } + let enables = ``; + if (template_type.requiresF16() || return_type.requiresF16()) { + enables = `enable f16;`; + } + const wgsl = ` +${enables} +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let res : ${ret_type} = bufferView<${template_type.toString()}>(&v, 0); +}`; + + t.expectCompileResult(t.params.ptr && template_type === return_type, wgsl); + }); + +g.test('buffer_type') + .desc('Validates the buffer parameter type') + .params(u => + u + .combine('type', [ + 'unsized_ro_storage', + 'unsized_storage', + 'sized_ro_storage', + 'sized_storage', + 'sized_uniform', + 'sized_workgroup', + 'override_workgroup', + ] as const) + .beginSubcases() + .combine('ptr', [false, true] as const) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + const wgsl = ` +@group(0) @binding(0) var unsized_ro_storage : buffer; +@group(0) @binding(1) var unsized_storage : buffer; +@group(0) @binding(2) var sized_ro_storage : buffer<128>; +@group(0) @binding(3) var sized_storage : buffer<128>; +@group(0) @binding(4) var sized_uniform : buffer<128>; +var sized_workgroup : buffer<128>; +override o : u32; +var override_workgroup : buffer; + +@compute @workgroup_size(1) +fn main() { + let p = bufferView(${t.params.ptr ? '&' : ''}${t.params.type}, 0); +}`; + + t.expectCompileResult(t.params.ptr, wgsl); + }); + +g.test('offset_type') + .desc('Validates the offset parameter type') + .params(u => u.combine('type', keysOf(kTypes))) + .fn(t => { + const type = kTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable f16;`; + } + const wgsl = ` +${enables} +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + let p = bufferView(&v, ${type.create(0).wgsl()}); +}`; + + t.expectCompileResult( + type === Type.abstractInt || type === Type.u32 || type === Type.i32, + wgsl + ); + }); + +interface EarlyEvalCase { + code: string; + valid: boolean | 'pipeline'; + constants?: Record; + ptr_param?: boolean; +} + +const kEarlyEvalCases: Record = { + ro_storage_buffer: { + code: ` +@group(0) @binding(0) var v : buffer<16>; +fn foo() { + let p = bufferView(&v, 0); +}`, + valid: true, + }, + storage_buffer: { + code: ` +@group(0) @binding(0) var v : buffer<16>; +fn foo() { + let p = bufferView(&v, 0); +}`, + valid: true, + }, + uniform_buffer: { + code: ` +@group(0) @binding(0) var v : buffer<16>; +fn foo() { + let p = bufferView(&v, 0); +}`, + valid: true, + }, + workgroup_buffer: { + code: ` +var v : buffer<16>; +fn foo() { + let p = bufferView(&v, 0); +}`, + valid: true, + }, + ro_storage_buffer_too_small_for_type: { + code: ` +@group(0) @binding(0) var v : buffer<8>; +fn foo() { + let p = bufferView(&v, 0); +}`, + valid: false, + }, + uniform_buffer_too_small_with_const_offset: { + code: ` +@group(0) @binding(0) var v : buffer<24>; +fn foo() { + let p = bufferView(&v, 16); +}`, + valid: false, + }, + uniform_buffer_too_small_with_override_offset: { + code: ` +@group(0) @binding(0) var v : buffer<24>; +override offset : u32; +fn foo() { + let p = bufferView(&v, offset); +}`, + valid: 'pipeline', + constants: { offset: 16 }, + }, + workgroup_buffer_too_small_with_override: { + code: ` +override size : u32; +var v : buffer; +fn foo() { + let p = bufferView(&v, 0); +}`, + valid: 'pipeline', + constants: { size: 12 }, + }, + storage_buffer_too_small_for_one_array_element: { + code: ` +@group(0) @binding(0) var v : buffer<28>; +struct S { + a: vec4u, + b: vec4f, +} +fn foo() { + let p = bufferView>(&v, 0); +}`, + valid: false, + }, + storage_buffer_out_of_range_u32: { + code: ` +@group(0) @binding(0) var v : buffer; +struct S { + a: vec4u, + b: vec4f, +} +fn foo() { + let p = buffer(&v, 4294967279); +}`, + valid: false, + }, + uniform_buffer_too_small_through_unsized_function: { + code: ` +@group(0) @binding(0) var v : buffer<28>; +fn bar(p : ptr) { + let q = bufferView(p, 16); +} +fn foo() { + bar(&v); +}`, + valid: false, + ptr_param: true, + }, + uniform_buffer_too_small_through_unsized_function_override: { + code: ` +@group(0) @binding(0) var v : buffer<28>; +override offset : u32; +fn bar(p : ptr) { + let q = bufferView(p, offset); +} +fn foo() { + bar(&v); +}`, + valid: 'pipeline', + constants: { offset: 16 }, + ptr_param: true, + }, + workgroup_to_smaller_size_param: { + code: ` +var v : buffer<128>; +fn bar(p : ptr) { + let q = bufferView(p, 0); +} +fn foo() { + bar(&v); +}`, + valid: false, + ptr_param: true, + }, + ro_storage_unsized_buffer_to_smaller_param: { + code: ` +@group(0) @binding(0) var v : buffer; +fn bar(p : ptr) { + let q = bufferView(p, 0); +} +fn foo() { + bar(&v); +}`, + valid: false, + ptr_param: true, + }, + offset_not_aligned: { + code: ` +@group(0) @binding(0) var v : buffer<128>; +fn foo() { + let p = bufferView(&v, 12); +}`, + valid: false, + }, + offset_not_aligned_override: { + code: ` +override offset : u32; +@group(0) @binding(0) var v : buffer<128>; +fn foo() { + let p = bufferView(&v, offset); +}`, + valid: 'pipeline', + constants: { offset: 12 }, + }, + offset_negative: { + code: ` +@group(0) @binding(0) var v : buffer<128>; +fn foo() { + let p = bufferView(&v, -16); +}`, + valid: false, + }, + offset_negative_override: { + code: ` +override offset : u32; +@group(0) @binding(0) var v : buffer<128>; +fn foo() { + let p = bufferView(&v, offset); +}`, + valid: 'pipeline', + constants: { offset: -16 }, + }, +}; + +g.test('early_eval_errors') + .desc('Test shader-creation and pipeline-creation errors') + .params(u => u.combine('case', keysOf(kEarlyEvalCases))) + .fn(t => { + const testcase = kEarlyEvalCases[t.params.case]; + t.skipIfLanguageFeatureNotSupported('buffer_view'); + if (testcase.ptr_param === true) { + t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters'); + } + + t.expectCompileResult(testcase.valid !== false, testcase.code); + if (testcase.valid !== false) { + t.expectPipelineResult({ + expectedResult: testcase.valid === true, + code: testcase.code, + constants: testcase.constants, + statements: ['foo();'], + }); + } + }); + +interface LHSCallCase { + code: string; + requires?: WGSLLanguageFeature[]; +} + +const kLHSCallCases: Record = { + u32: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + *bufferView(&v, 0) = 123u; +}`, + }, + u32_array: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + (*bufferView>(&v, 0))[0] = 123u; +}`, + }, + f32_vector_letter: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + (*bufferView(&v, 0)).x = 13.3f; +}`, + }, + f32_struct_access: { + code: ` +struct S { + member_a: f32, +} +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + (*bufferView(&v, 0)).member_a = 13.3f; +}`, + }, + u32_array_ptr_composite: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + bufferView>(&v, 0)[0] = 123u; +}`, + requires: ['pointer_composite_access'], + }, + vec4u_swizzle_assign: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + (*bufferView(&v, 0)).zxy = vec3u(42); +}`, + requires: ['swizzle_assignment'], + }, + vec4u_ptr_composite_swizzle_assign: { + code: ` +@group(0) @binding(0) var v : buffer; +@compute @workgroup_size(1) +fn main() { + bufferView(&v, 0).zxy = vec3u(42); +}`, + requires: ['pointer_composite_access', 'swizzle_assignment'], + }, + compound_assign: { + code: ` +var v : buffer<256>; +@compute @workgroup_size(1) +fn main() { + *bufferView(&v, 0) += 123u; +}`, + }, +}; + +g.test('lhs_call') + .desc('Validate that bufferView can be on the LHS of an assignment') + .params(u => u.combine('case', keysOf(kLHSCallCases))) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + + const testcase = kLHSCallCases[t.params.case]; + let valid = true; + const features = testcase.requires ?? []; + features.forEach(f => { + valid &&= t.hasLanguageFeature(f); + }); + t.expectCompileResult(valid, testcase.code); + }); diff --git a/src/webgpu/shader/validation/functions/restrictions.spec.ts b/src/webgpu/shader/validation/functions/restrictions.spec.ts index 73c3a376765e..5cec4c3dc21b 100644 --- a/src/webgpu/shader/validation/functions/restrictions.spec.ts +++ b/src/webgpu/shader/validation/functions/restrictions.spec.ts @@ -202,7 +202,7 @@ fn foo() -> ${testcase.name} { interface ParamTypeCase { name: string; - valid: boolean | 'with_unrestricted_pointer_parameters'; + valid: boolean | 'with_unrestricted_pointer_parameters' | 'with_buffer_view'; } const kFunctionParamTypeCases: Record = { @@ -304,6 +304,68 @@ const kFunctionParamTypeCases: Record = { invalid_ptr6: { name: `ptr`, valid: false }, // Can't specify access mode invalid_ptr7: { name: `ptr`, valid: false }, // Invalid store type invalid_ptr8: { name: `ptr`, valid: false }, // non-constructible pointer type + + // Buffers (need buffer_view and unrestricted_pointer_parameters) + ptrBufferUnsized_storage: { + name: `ptr`, + valid: 'with_buffer_view', + }, + ptrBufferSized_storage: { + name: `ptr, read_write>`, + valid: 'with_buffer_view', + }, + ptrBufferSizedSmall_storage: { + name: `ptr, read_write>`, + valid: 'with_buffer_view', + }, + ptrBufferUnsized_ro_storage: { + name: `ptr`, + valid: 'with_buffer_view', + }, + ptrBufferSized_ro_storage: { + name: `ptr>`, + valid: 'with_buffer_view', + }, + ptrBufferSizedSmall_ro_storage: { + name: `ptr>`, + valid: 'with_buffer_view', + }, + ptrBufferUnsized_uniform: { + name: `ptr`, + valid: 'with_buffer_view', + }, + ptrBufferSized_uniform: { + name: `ptr>`, + valid: 'with_buffer_view', + }, + ptrBufferSizedSmall_uniform: { + name: `ptr>`, + valid: 'with_buffer_view', + }, + ptrBufferUnsized_workgroup: { + name: `ptr`, + valid: 'with_buffer_view', + }, + ptrBufferSized_workgroup: { + name: `ptr>`, + valid: 'with_buffer_view', + }, + ptrBufferSizedSmall_workgroup: { + name: `ptr>`, + valid: 'with_buffer_view', + }, + ptrBufferOverrideNoDefault_workgroup: { + name: `ptr>`, + valid: 'with_buffer_view', + }, + ptrBufferOverrideDefault_workgroup: { + name: `ptr>`, + valid: 'with_buffer_view', + }, + ptrBufferOverrideExpr_workgroup: { + name: `ptr>`, + valid: 'with_buffer_view', + }, }; g.test('function_parameter_types') @@ -324,6 +386,10 @@ fn foo(param : ${testcase.name}) { let isValid = testcase.valid; if (isValid === 'with_unrestricted_pointer_parameters') { isValid = t.hasLanguageFeature('unrestricted_pointer_parameters'); + } else if (isValid === 'with_buffer_view') { + isValid = + t.hasLanguageFeature('unrestricted_pointer_parameters') && + t.hasLanguageFeature('buffer_view'); } t.expectCompileResult(isValid, code); @@ -333,6 +399,7 @@ interface ParamValueCase { value: string; matches: string[]; needsUnrestrictedPointerParameters?: boolean; + needsBufferView?: boolean; } const kFunctionParamValueCases: Record = { @@ -519,6 +586,81 @@ const kFunctionParamValueCases: Record = { matches: ['ptrWorkgroupOverrideExpr'], needsUnrestrictedPointerParameters: true, }, + + // Buffer view + ptrStorageBufferUnsized: { + value: `&storage_buffer_unsized`, + matches: ['ptrBufferUnsized_storage'], + needsBufferView: true, + }, + ptrStorageBufferSized: { + value: `&storage_buffer_sized`, + matches: ['ptrBufferUnsized_storage', 'ptrBufferSized_storage', 'ptrBufferSizedSmall_storage'], + needsBufferView: true, + }, + ptrStorageBufferSizedSmall: { + value: `&storage_buffer_sized_small`, + matches: ['ptrBufferUnsized_storage', 'ptrBufferSizedSmall_storage'], + needsBufferView: true, + }, + ptrROStorageBufferUnsized: { + value: `&ro_storage_buffer_unsized`, + matches: ['ptrBufferUnsized_ro_storage'], + needsBufferView: true, + }, + ptrROStorageBufferSized: { + value: `&ro_storage_buffer_sized`, + matches: [ + 'ptrBufferUnsized_ro_storage', + 'ptrBufferSized_ro_storage', + 'ptrBufferSizedSmall_ro_storage', + ], + needsBufferView: true, + }, + ptrROStorageBufferSizedSmall: { + value: `&ro_storage_buffer_sized_small`, + matches: ['ptrBufferUnsized_ro_storage', 'ptrBufferSizedSmall_ro_storage'], + needsBufferView: true, + }, + ptrUniformBufferSized: { + value: `&uniform_buffer_sized`, + matches: ['ptrBufferUnsized_uniform', 'ptrBufferSized_uniform', 'ptrBufferSizedSmall_uniform'], + needsBufferView: true, + }, + ptrUniformBufferSizedSmall: { + value: `&uniform_buffer_sized_small`, + matches: ['ptrBufferUnsized_uniform', 'ptrBufferSizedSmall_uniform'], + needsBufferView: true, + }, + ptrWorkgroupBufferSized: { + value: `&wg_buffer_sized`, + matches: [ + 'ptrBufferUnsized_workgroup', + 'ptrBufferSized_workgroup', + 'ptrBufferSizedSmall_workgroup', + ], + needsBufferView: true, + }, + ptrWorkgroupBufferSizedSmall: { + value: `&wg_buffer_sized_small`, + matches: ['ptrBufferUnsized_workgroup', 'ptrBufferSizedSmall_workgroup'], + needsBufferView: true, + }, + ptrWorkgroupBufferOverrideNoDefault: { + value: `&wg_buffer_override_no_default`, + matches: ['ptrBufferUnsized_workgroup', 'ptrBufferOverrideNoDefault_workgroup'], + needsBufferView: true, + }, + ptrWorkgroupBufferOverrideDefault: { + value: `&wg_buffer_override_default`, + matches: ['ptrBufferUnsized_workgroup', 'ptrBufferOverrideDefault_workgroup'], + needsBufferView: true, + }, + ptrWorkgroupBufferOverrideExpr: { + value: `&wg_buffer_override_expr`, + matches: ['ptrBufferUnsized_workgroup', 'ptrBufferOverrideExpr_workgroup'], + needsBufferView: true, + }, }; function parameterMatches(decl: string, matches: string[]): boolean { @@ -548,6 +690,26 @@ g.test('function_parameter_matching') const param = kFunctionParamTypeCases[t.params.decl]; const arg = kFunctionParamValueCases[t.params.arg]; const enable = param.name === 'f16' ? 'enable f16;' : ''; + let buffer_decls = ``; + if (param.valid === 'with_buffer_view') { + buffer_decls = ` +@group(2) @binding(0) var storage_buffer_unsized : buffer; +@group(2) @binding(1) var storage_buffer_sized : buffer<128>; +@group(2) @binding(2) var storage_buffer_sized_small : buffer<64>; +@group(2) @binding(3) var ro_storage_buffer_unsized : buffer; +@group(2) @binding(4) var ro_storage_buffer_sized : buffer<128>; +@group(2) @binding(5) var ro_storage_buffer_sized_small : buffer<64>; +@group(2) @binding(6) var uniform_buffer_sized : buffer<128>; +@group(2) @binding(7) var uniform_buffer_sized_small : buffer<64>; + +var wg_buffer_sized : buffer<128>; +var wg_buffer_sized_small : buffer<64>; +var wg_buffer_override_no_default : buffer; +var wg_buffer_override_default : buffer; +var wg_buffer_override_expr : buffer; +`; + } + const code = ` ${enable} @@ -574,6 +736,8 @@ var rw_host_shareable : host_shareable; @group(1) @binding(2) var uniform_host_shareable : host_shareable; +${buffer_decls} + fn bar(param : ${param.name}) { } var g_u32 : u32; @@ -642,11 +806,20 @@ fn foo() { (kFunctionParamTypeCases[t.params.decl].valid === 'with_unrestricted_pointer_parameters' || arg.needsUnrestrictedPointerParameters) ?? false; + const needsBufferView = + (kFunctionParamTypeCases[t.params.decl].valid === 'with_buffer_view' || + arg.needsBufferView) ?? + false; let isValid = parameterMatches(t.params.decl, arg.matches); if (isValid && needsUnrestrictedPointerParameters) { isValid = t.hasLanguageFeature('unrestricted_pointer_parameters'); } + if (isValid && needsBufferView) { + isValid = + t.hasLanguageFeature('unrestricted_pointer_parameters') && + t.hasLanguageFeature('buffer_view'); + } t.expectCompileResult(isValid, code); }); diff --git a/src/webgpu/shader/validation/shader_io/locations.spec.ts b/src/webgpu/shader/validation/shader_io/locations.spec.ts index 5641b788cad9..c569e28a890d 100644 --- a/src/webgpu/shader/validation/shader_io/locations.spec.ts +++ b/src/webgpu/shader/validation/shader_io/locations.spec.ts @@ -95,6 +95,8 @@ const kInvalidLocationTypes = new Set([ 'texture_depth_multisampled_2d', 'sampler', 'sampler_comparison', + 'buffer', + 'buffer<16>', ]); g.test('stage_inout') diff --git a/src/webgpu/shader/validation/types/buffer.spec.ts b/src/webgpu/shader/validation/types/buffer.spec.ts new file mode 100644 index 000000000000..fde3145ee472 --- /dev/null +++ b/src/webgpu/shader/validation/types/buffer.spec.ts @@ -0,0 +1,129 @@ +export const description = ` +Validation tests for buffer types +`; + +import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { keysOf } from '../../../../common/util/data_tables.js'; +import { ShaderValidationTest } from '../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +const kParseCases = { + unsized: { + code: `alias T = buffer;`, + valid: true, + }, + literal: { + code: `alias T = buffer<16>;`, + valid: true, + }, + literal_negative: { + code: `alias T = buffer<-1>;`, + valid: false, + }, + literal_zero: { + code: `alias T = buffer<0>;`, + valid: false, + }, + const: { + code: `const x = 16; alias T = buffer;`, + valid: true, + }, + const_negative: { + code: `const x = -1; alias T = buffer;`, + valid: false, + }, + const_zero: { + code: `const x = 0; alias T = buffer;`, + valid: false, + }, + const_functio: { + code: `const x = 16; const y = 32; alias T = buffer;`, + valid: true, + }, + override: { + code: `override x = 16; alias T = buffer;`, + valid: true, + }, + override_negative: { + code: `override x = -1; alias T = buffer;`, + valid: true, + }, + override_zero: { + code: `override x = 0; alias T = buffer;`, + valid: true, + }, + override_function: { + code: `override x = 16; override y = 32; alias T = buffer;`, + valid: true, + }, + empty_template: { + code: `alias T = buffer<>;`, + valid: false, + }, + missing_rparen: { + code: `alias T = buffer<16;`, + valid: false, + }, + subtype: { + code: `alias T = buffer;`, + valid: false, + }, +}; + +g.test('parse') + .desc('Test buffer type parsing') + .params(u => u.combine('case', keysOf(kParseCases))) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + const testcase = kParseCases[t.params.case]; + t.expectCompileResult(testcase.valid, testcase.code); + }); + +g.test('address_space') + .desc('Test buffer type validity for each address space') + .params(u => + u + .combine('case', keysOf(kParseCases)) + .filter(t => { + return kParseCases[t.case].valid; + }) + .beginSubcases() + .combine('aspace', ['function', 'private', 'storage', 'uniform', 'workgroup'] as const) + ) + .fn(t => { + t.skipIfLanguageFeatureNotSupported('buffer_view'); + const testcase = kParseCases[t.params.case]; + + let mvar = ''; + let fvar = ''; + switch (t.params.aspace) { + case 'function': + fvar = `var v : T;`; + break; + case 'private': + case 'workgroup': + mvar = `var<${t.params.aspace}> v : T;`; + break; + case 'storage': + case 'uniform': + mvar = `@group(0) @binding(0) var<${t.params.aspace}> v : T;`; + break; + } + + const wgsl = ` +${testcase.code} +${mvar} +@compute @workgroup_size(1) +fn main() { + ${fvar} +}`; + + let expected = t.params.aspace !== 'function' && t.params.aspace !== 'private'; + if (t.params.case === 'unsized') { + expected = t.params.aspace === 'storage'; + } else if (t.params.case.includes('override')) { + expected = t.params.aspace === 'workgroup'; + } + t.expectCompileResult(expected, wgsl); + });