buffer_view CTS (#4652)
* Validation of buffer type and related builtin functions * Execution tests for buffer_view
This commit is contained in:
alan-baker
GitHub
parent
70c1848940
commit
af9a578195
19 changed files with 5450 additions and 113 deletions
3
.gitignore
vendored
3
.gitignore
vendored
|
|
@ -11,6 +11,9 @@
|
|||
*.tmp.txt
|
||||
/docs/tsdoc/
|
||||
|
||||
# vim swap files
|
||||
*.swp*
|
||||
|
||||
# `pnpm` lockfile
|
||||
pnpm-lock.yaml
|
||||
|
||||
|
|
|
|||
|
|
@ -373,6 +373,7 @@ interface TypedArrayMap {
|
|||
Uint8ClampedArray: Uint8ClampedArray;
|
||||
Int32Array: Int32Array;
|
||||
Uint32Array: Uint32Array;
|
||||
Float16Array: Float16Array;
|
||||
Float32Array: Float32Array;
|
||||
Float64Array: Float64Array;
|
||||
BigInt64Array: BigInt64Array;
|
||||
|
|
|
|||
220
src/resources/cache/hashes.json
generated
vendored
220
src/resources/cache/hashes.json
generated
vendored
|
|
@ -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": "7fb5b26a",
|
||||
"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": "78c2f699",
|
||||
"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"
|
||||
}
|
||||
|
|
@ -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,428 @@ 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; // Size must be greater than 4
|
||||
binding_type: GPUBufferBindingType;
|
||||
requires?: WGSLLanguageFeature[];
|
||||
}
|
||||
|
||||
const kRequiredSizeCases: Record<string, RequiredSizeCase> = {
|
||||
ro_storage_32bytes: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : array<vec4u, 2>;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
_ = v;
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: 'read-only-storage',
|
||||
},
|
||||
storage_32bytes: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : array<vec4u, 2>;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
_ = v;
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: 'storage',
|
||||
},
|
||||
uniform_32bytes: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<uniform> v : array<vec4u, 2>;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
_ = v;
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: 'uniform',
|
||||
},
|
||||
ro_storage_sized_buffer: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> 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<storage, read_write> 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<uniform> 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<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<array<u32, 16>>(&v, 0);
|
||||
}`,
|
||||
size: 16 * 4,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
ro_storage_unsized_buffer_bufferView2: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<array<u32, 16>>(&v, 56);
|
||||
}`,
|
||||
// Offset does not count toward min binding size.
|
||||
size: 16 * 4,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView3: {
|
||||
code: `
|
||||
struct S {
|
||||
a: vec4u,
|
||||
b: vec2u,
|
||||
}
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let x = 16u;
|
||||
let p = bufferView<S>(&v, x);
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView4: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<array<vec2u>>(&v, 0);
|
||||
}`,
|
||||
size: 8,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView5: {
|
||||
code: `
|
||||
struct S {
|
||||
a: array<u32, 4>
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<vec2u>(&v, 0);
|
||||
let q = bufferView<S>(&v, 0);
|
||||
}`,
|
||||
size: 16,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView6: {
|
||||
code: `
|
||||
struct S {
|
||||
a: array<u32, 4>
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<u32>(&v, 16);
|
||||
let q = bufferView<S>(&v, 0);
|
||||
}`,
|
||||
// Offset does not count toward min binding size.
|
||||
size: 16,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView7: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<vec2i>(&v, 16);
|
||||
}
|
||||
@fragment
|
||||
fn main2() {
|
||||
let q = bufferView<vec4u>(&v, 32);
|
||||
}`,
|
||||
// Offset does not count toward min binding size.
|
||||
size: 8,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView8: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<vec2f>(&v, 16);
|
||||
}
|
||||
fn foo() {
|
||||
let q = bufferView<vec4u>(&v, 32);
|
||||
}`,
|
||||
// Offset does not count toward min binding size.
|
||||
size: 8,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView9: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<u32>(&v, 16);
|
||||
foo(&v);
|
||||
}
|
||||
fn foo(p : ptr<storage, buffer, read_write>) {
|
||||
let q = bufferView<vec4u>(&v, 32);
|
||||
}`,
|
||||
// Offset does not count toward min binding size.
|
||||
size: 16,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view', 'unrestricted_pointer_parameters'],
|
||||
},
|
||||
ro_storage_unsized_buffer_bufferArrayView1: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferArrayView<array<vec2u>>(&v, 0, 32);
|
||||
}`,
|
||||
// Size does not count toward min binding size.
|
||||
size: 8,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
ro_storage_unsized_buffer_bufferArrayView2: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferArrayView<array<vec2f>>(&v, 16, 32);
|
||||
}`,
|
||||
// Offset and size do not count toward min binding size.
|
||||
size: 8,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
ro_storage_unsized_buffer_bufferArrayView3: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let o = 16;
|
||||
let s = 32;
|
||||
let p = bufferArrayView<array<vec2u>>(&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<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let o = 16;
|
||||
let s = 32;
|
||||
let p = bufferArrayView<array<S>>(&v, o, s);
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferArrayView5: {
|
||||
code: `
|
||||
struct S {
|
||||
a : vec4f,
|
||||
b : vec4f,
|
||||
}
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let o = 0u;
|
||||
let s = 64;
|
||||
let p1 = bufferArrayView<array<u32>>(&v, o, s);
|
||||
let p2 = bufferArrayView<array<S>>(&v, o, s);
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferArrayView6: {
|
||||
code: `
|
||||
struct S {
|
||||
a : vec4f,
|
||||
b : vec4f,
|
||||
}
|
||||
var<private> o : i32;
|
||||
var<private> s : i32;
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, o, s);
|
||||
}
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<S>>(&v, o, s);
|
||||
}`,
|
||||
size: 16,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferArrayView7: {
|
||||
code: `
|
||||
struct S {
|
||||
a : vec4f,
|
||||
b : vec4f,
|
||||
}
|
||||
var<private> o : i32;
|
||||
var<private> s : i32;
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, o, s);
|
||||
}
|
||||
@fragment
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<S>>(&v, o, s);
|
||||
}`,
|
||||
size: 16,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferArrayView8: {
|
||||
code: `
|
||||
struct S {
|
||||
a : vec4f,
|
||||
b : vec4f,
|
||||
}
|
||||
var<private> o : i32;
|
||||
var<private> s : i32;
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
bar();
|
||||
foo();
|
||||
}
|
||||
fn bar() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, o, s);
|
||||
}
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<S>>(&v, o, s);
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: '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);
|
||||
}
|
||||
});
|
||||
|
|
|
|||
658
src/webgpu/api/validation/min_buffer_binding_size.spec.ts
Normal file
658
src/webgpu/api/validation/min_buffer_binding_size.spec.ts
Normal file
|
|
@ -0,0 +1,658 @@
|
|||
export const description = `
|
||||
Minimum buffer binding size 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);
|
||||
|
||||
interface RequiredSizeCase {
|
||||
code: string;
|
||||
size: number; // Size must be greater than 4
|
||||
binding_type: GPUBufferBindingType;
|
||||
requires?: WGSLLanguageFeature[];
|
||||
}
|
||||
|
||||
const kRequiredSizeCases: Record<string, RequiredSizeCase> = {
|
||||
ro_storage_32bytes: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : array<vec4u, 2>;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
_ = v;
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: 'read-only-storage',
|
||||
},
|
||||
storage_32bytes: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : array<vec4u, 2>;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
_ = v;
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: 'storage',
|
||||
},
|
||||
uniform_32bytes: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<uniform> v : array<vec4u, 2>;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
_ = v;
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: 'uniform',
|
||||
},
|
||||
ro_storage_sized_buffer: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> 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<storage, read_write> 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<uniform> 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<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<array<u32, 16>>(&v, 0);
|
||||
}`,
|
||||
size: 16 * 4,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
ro_storage_unsized_buffer_bufferView2: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<array<u32, 16>>(&v, 56);
|
||||
}`,
|
||||
// Offset does not count toward min binding size.
|
||||
size: 16 * 4,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView3: {
|
||||
code: `
|
||||
struct S {
|
||||
a: vec4u,
|
||||
b: vec2u,
|
||||
}
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let x = 16u;
|
||||
let p = bufferView<S>(&v, x);
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView4: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<array<vec2u>>(&v, 0);
|
||||
}`,
|
||||
size: 8,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView5: {
|
||||
code: `
|
||||
struct S {
|
||||
a: array<u32, 4>
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<vec2u>(&v, 0);
|
||||
let q = bufferView<S>(&v, 0);
|
||||
}`,
|
||||
size: 16,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView6: {
|
||||
code: `
|
||||
struct S {
|
||||
a: array<u32, 4>
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<u32>(&v, 16);
|
||||
let q = bufferView<S>(&v, 0);
|
||||
}`,
|
||||
// Offset does not count toward min binding size.
|
||||
size: 16,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView7: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<vec2i>(&v, 16);
|
||||
}
|
||||
@fragment
|
||||
fn main2() {
|
||||
let q = bufferView<vec4u>(&v, 32);
|
||||
}`,
|
||||
// Offset does not count toward min binding size.
|
||||
size: 8,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView8: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<vec2f>(&v, 16);
|
||||
}
|
||||
fn foo() {
|
||||
let q = bufferView<vec4u>(&v, 32);
|
||||
}`,
|
||||
// Offset does not count toward min binding size.
|
||||
size: 8,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferView9: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<u32>(&v, 16);
|
||||
foo(&v);
|
||||
}
|
||||
fn foo(p : ptr<storage, buffer, read_write>) {
|
||||
let q = bufferView<vec4u>(&v, 32);
|
||||
}`,
|
||||
// Offset does not count toward min binding size.
|
||||
size: 16,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view', 'unrestricted_pointer_parameters'],
|
||||
},
|
||||
ro_storage_unsized_buffer_bufferArrayView1: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferArrayView<array<vec2u>>(&v, 0, 32);
|
||||
}`,
|
||||
// Size does not count toward min binding size.
|
||||
size: 8,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
ro_storage_unsized_buffer_bufferArrayView2: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferArrayView<array<vec2f>>(&v, 16, 32);
|
||||
}`,
|
||||
// Offset and size do not count toward min binding size.
|
||||
size: 8,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
ro_storage_unsized_buffer_bufferArrayView3: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let o = 16;
|
||||
let s = 32;
|
||||
let p = bufferArrayView<array<vec2u>>(&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<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let o = 16;
|
||||
let s = 32;
|
||||
let p = bufferArrayView<array<S>>(&v, o, s);
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferArrayView5: {
|
||||
code: `
|
||||
struct S {
|
||||
a : vec4f,
|
||||
b : vec4f,
|
||||
}
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let o = 0u;
|
||||
let s = 64;
|
||||
let p1 = bufferArrayView<array<u32>>(&v, o, s);
|
||||
let p2 = bufferArrayView<array<S>>(&v, o, s);
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferArrayView6: {
|
||||
code: `
|
||||
struct S {
|
||||
a : vec4f,
|
||||
b : vec4f,
|
||||
}
|
||||
var<private> o : i32;
|
||||
var<private> s : i32;
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, o, s);
|
||||
}
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<S>>(&v, o, s);
|
||||
}`,
|
||||
size: 16,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferArrayView7: {
|
||||
code: `
|
||||
struct S {
|
||||
a : vec4f,
|
||||
b : vec4f,
|
||||
}
|
||||
var<private> o : i32;
|
||||
var<private> s : i32;
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, o, s);
|
||||
}
|
||||
@fragment
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<S>>(&v, o, s);
|
||||
}`,
|
||||
size: 16,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
storage_unsized_buffer_bufferArrayView8: {
|
||||
code: `
|
||||
struct S {
|
||||
a : vec4f,
|
||||
b : vec4f,
|
||||
}
|
||||
var<private> o : i32;
|
||||
var<private> s : i32;
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
bar();
|
||||
foo();
|
||||
}
|
||||
fn bar() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, o, s);
|
||||
}
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<S>>(&v, o, s);
|
||||
}`,
|
||||
size: 32,
|
||||
binding_type: 'storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
};
|
||||
|
||||
g.test('compute,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);
|
||||
}
|
||||
});
|
||||
|
||||
const kRequiredRenderSizeCases: Record<string, RequiredSizeCase> = {
|
||||
fragment_larger_buffer_view_ro_storage: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> buf : buffer;
|
||||
@group(0) @binding(1) var<uniform> uni : u32;
|
||||
|
||||
@vertex
|
||||
fn vs() -> @builtin(position) vec4f {
|
||||
_ = uni;
|
||||
_ = bufferView<array<u32, 2>>(&buf, 0);
|
||||
return vec4f(0);
|
||||
}
|
||||
|
||||
@fragment
|
||||
fn fs() -> @location(0) vec4f {
|
||||
_ = uni;
|
||||
_ = bufferView<array<u32, 4>>(&buf, 0);
|
||||
return vec4f(0);
|
||||
}
|
||||
`,
|
||||
size: 16,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
vertex_larger_buffer_view_ro_storage: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> buf : buffer;
|
||||
@group(0) @binding(1) var<uniform> uni : u32;
|
||||
|
||||
@vertex
|
||||
fn vs() -> @builtin(position) vec4f {
|
||||
_ = uni;
|
||||
_ = bufferView<array<vec4f, 2>>(&buf, 0);
|
||||
return vec4f(0);
|
||||
}
|
||||
|
||||
@fragment
|
||||
fn fs() -> @location(0) vec4f {
|
||||
_ = uni;
|
||||
_ = bufferView<array<u32, 4>>(&buf, 0);
|
||||
return vec4f(0);
|
||||
}
|
||||
`,
|
||||
size: 32,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
fragment_larger_buffer_array_view_ro_storage: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> buf : buffer;
|
||||
@group(0) @binding(1) var<uniform> uni : u32;
|
||||
|
||||
@vertex
|
||||
fn vs() -> @builtin(position) vec4f {
|
||||
_ = bufferArrayView<array<u32>>(&buf, 0, uni);
|
||||
return vec4f(0);
|
||||
}
|
||||
|
||||
@fragment
|
||||
fn fs() -> @location(0) vec4f {
|
||||
_ = bufferArrayView<array<vec4u>>(&buf, 0, uni);
|
||||
return vec4f(0);
|
||||
}
|
||||
`,
|
||||
size: 16,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
vertex_larger_buffer_array_view_ro_storage: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> buf : buffer;
|
||||
@group(0) @binding(1) var<uniform> uni : u32;
|
||||
|
||||
@vertex
|
||||
fn vs() -> @builtin(position) vec4f {
|
||||
_ = bufferArrayView<array<vec4f>>(&buf, 0, uni);
|
||||
return vec4f(0);
|
||||
}
|
||||
|
||||
@fragment
|
||||
fn fs() -> @location(0) vec4f {
|
||||
_ = bufferArrayView<array<u32>>(&buf, 0, uni);
|
||||
return vec4f(0);
|
||||
}
|
||||
`,
|
||||
size: 16,
|
||||
binding_type: 'read-only-storage',
|
||||
requires: ['buffer_view'],
|
||||
},
|
||||
};
|
||||
|
||||
g.test('render,shader_required_buffer_size')
|
||||
.params(u =>
|
||||
u
|
||||
.combine('case', keysOf(kRequiredRenderSizeCases))
|
||||
.beginSubcases()
|
||||
.combine('valid', [false, true] as const)
|
||||
.combine('layout', ['auto', 'explicit'] as const)
|
||||
)
|
||||
.fn(t => {
|
||||
const testcase = kRequiredRenderSizeCases[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 uniform = t.createBufferTracked({
|
||||
size: 4,
|
||||
usage: GPUBufferUsage.UNIFORM,
|
||||
});
|
||||
|
||||
const bgLayout = t.device.createBindGroupLayout({
|
||||
entries: [
|
||||
{
|
||||
binding: 0,
|
||||
visibility: GPUShaderStage.VERTEX | GPUShaderStage.FRAGMENT,
|
||||
buffer: {
|
||||
type: testcase.binding_type,
|
||||
minBindingSize: 0,
|
||||
},
|
||||
},
|
||||
{
|
||||
binding: 1,
|
||||
visibility: GPUShaderStage.VERTEX | GPUShaderStage.FRAGMENT,
|
||||
buffer: {
|
||||
type: 'uniform',
|
||||
minBindingSize: 4,
|
||||
},
|
||||
},
|
||||
],
|
||||
});
|
||||
const layout = t.device.createPipelineLayout({
|
||||
bindGroupLayouts: [bgLayout],
|
||||
});
|
||||
|
||||
const colorTexture = t.createTextureTracked({
|
||||
size: { width: 1, height: 1 },
|
||||
format: 'rgba8unorm',
|
||||
usage: GPUTextureUsage.RENDER_ATTACHMENT,
|
||||
});
|
||||
|
||||
const shaderModule = t.device.createShaderModule({ code: testcase.code });
|
||||
const pipeline = t.device.createRenderPipeline({
|
||||
layout: t.params.layout === 'auto' ? 'auto' : layout,
|
||||
vertex: { module: shaderModule },
|
||||
fragment: {
|
||||
module: shaderModule,
|
||||
targets: [{ format: 'rgba8unorm' }],
|
||||
},
|
||||
primitive: {
|
||||
topology: 'triangle-list',
|
||||
},
|
||||
});
|
||||
|
||||
if (t.params.layout === 'auto' && !t.params.valid) {
|
||||
// 'auto' layout get minBindingSize from shaders
|
||||
t.expectValidationError(() => {
|
||||
t.device.createBindGroup({
|
||||
layout: pipeline.getBindGroupLayout(0),
|
||||
entries: [
|
||||
{
|
||||
binding: 0,
|
||||
resource: {
|
||||
buffer,
|
||||
},
|
||||
},
|
||||
{
|
||||
binding: 1,
|
||||
resource: {
|
||||
buffer: uniform,
|
||||
},
|
||||
},
|
||||
],
|
||||
}),
|
||||
true;
|
||||
});
|
||||
} else {
|
||||
const bg = t.device.createBindGroup({
|
||||
layout: pipeline.getBindGroupLayout(0),
|
||||
entries: [
|
||||
{
|
||||
binding: 0,
|
||||
resource: {
|
||||
buffer,
|
||||
},
|
||||
},
|
||||
{
|
||||
binding: 1,
|
||||
resource: {
|
||||
buffer: uniform,
|
||||
},
|
||||
},
|
||||
],
|
||||
});
|
||||
|
||||
const encoder = t.device.createCommandEncoder();
|
||||
const pass = encoder.beginRenderPass({
|
||||
colorAttachments: [
|
||||
{
|
||||
view: colorTexture.createView(),
|
||||
clearValue: [0, 0, 0, 0],
|
||||
loadOp: 'clear',
|
||||
storeOp: 'store',
|
||||
},
|
||||
],
|
||||
});
|
||||
pass.setPipeline(pipeline);
|
||||
pass.setBindGroup(0, bg);
|
||||
pass.draw(1), pass.end();
|
||||
|
||||
t.expectValidationError(() => {
|
||||
encoder.finish();
|
||||
}, !t.params.valid);
|
||||
}
|
||||
});
|
||||
|
|
@ -991,6 +991,7 @@ export const kKnownWGSLLanguageFeatures = [
|
|||
'linear_indexing',
|
||||
'texture_formats_tier1',
|
||||
'immediate_address_space',
|
||||
'buffer_view',
|
||||
] as const;
|
||||
|
||||
export type WGSLLanguageFeature = (typeof kKnownWGSLLanguageFeatures)[number];
|
||||
|
|
|
|||
|
|
@ -0,0 +1,618 @@
|
|||
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<workgroup> v : ${buffer_ty};\n@group(0) @binding(0) var<storage> dummy : array<u32>;`;
|
||||
break;
|
||||
case 'uniform':
|
||||
decl = `@group(0) @binding(0) var<uniform> v : ${buffer_ty};`;
|
||||
break;
|
||||
case 'storage':
|
||||
decl = `@group(0) @binding(0) var<storage> 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<storage> offsets : array<u32>;
|
||||
@group(0) @binding(2) var<storage> sizes : array<u32>; // unused
|
||||
@group(0) @binding(3) var<storage, read_write> out : array<u32>;
|
||||
|
||||
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<workgroup> v : ${buffer_ty};\n@group(0) @binding(0) var<storage> dummy : array<u32>;`;
|
||||
break;
|
||||
case 'uniform':
|
||||
decl = `@group(0) @binding(0) var<uniform> v : ${buffer_ty};`;
|
||||
break;
|
||||
case 'storage':
|
||||
decl = `@group(0) @binding(0) var<storage> 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);
|
||||
}
|
||||
|
||||
fn sized_sized_indirectLength(p : ptr<${t.params.aspace}, buffer<${
|
||||
t.params.bufferSize
|
||||
}>>, gidx : u32) -> u32 {
|
||||
return sizedLength(p, gidx);
|
||||
}
|
||||
|
||||
override bufferSize : u32 = 0u;
|
||||
${decl}
|
||||
@group(0) @binding(1) var<storage> offsets : array<u32>;
|
||||
@group(0) @binding(2) var<storage> sizes : array<u32>; // unused
|
||||
@group(0) @binding(3) var<storage, read_write> out : array<u32>;
|
||||
|
||||
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<storage, read_write> out : u32;
|
||||
|
||||
var<workgroup> wg_var : buffer<256 * 4>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let in_ptr = bufferView<array<vec4u, 256 / 4>>(&in, 0);
|
||||
let wg_ptr = bufferView<array<vec4u, 256 / 4>>(&wg_var, 0);
|
||||
for (var i = 0; i < 256 / 4; i++) {
|
||||
(*wg_ptr)[i] = (*in_ptr)[i];
|
||||
}
|
||||
|
||||
workgroupBarrier();
|
||||
|
||||
let p = bufferArrayView<array<${testcase.type}>>(&${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<array<${testcase.type}>>(&${v}, ${t.params.offset}, bufferLength(&${v}) - ${t.params.offset})`;
|
||||
|
||||
let post_assign = '';
|
||||
if (t.params.aspace === 'workgroup') {
|
||||
post_assign = `
|
||||
workgroupBarrier();
|
||||
|
||||
let out_ptr = bufferView<array<u32>>(&out, 0);
|
||||
let wg_ptr = bufferView<array<u32>>(&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<storage> in : u32;
|
||||
@group(0) @binding(1) var<storage, read_write> out : buffer<256 * 4>;
|
||||
|
||||
|
||||
var<workgroup> 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<storage, read_write> out : ${ty.toString()};
|
||||
|
||||
var<workgroup> wg_var : buffer<${bufferSize}>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let in_ptr = bufferView<array<vec4u, ${bufferSize} / 16>>(&in, 0);
|
||||
let wg_ptr = bufferView<array<vec4u, ${bufferSize} / 16>>(&wg_var, 0);
|
||||
for (var i = 0; i < 256 / 4; i++) {
|
||||
(*wg_ptr)[i] = (*in_ptr)[i];
|
||||
}
|
||||
|
||||
workgroupBarrier();
|
||||
|
||||
let p = bufferArrayView<array<${ty.toString()}>>(&${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<array<u32>>(&out, 0);
|
||||
let wg_ptr = bufferView<array<u32>>(&wg_var, 0);
|
||||
for (var i = 0; i < 256; i++) {
|
||||
(*out_ptr)[i] = (*wg_ptr)[i];
|
||||
}
|
||||
`;
|
||||
}
|
||||
|
||||
const call_expr = `bufferArrayView<array<${ty.toString()}>>(&${v}, ${
|
||||
t.params.offset
|
||||
}, bufferLength(&${v}) - ${t.params.offset})`;
|
||||
const wgsl = `
|
||||
${enables}
|
||||
|
||||
@group(0) @binding(0) var<storage> in : ${ty.toString()};
|
||||
@group(0) @binding(1) var<storage, read_write> out : buffer<${bufferSize}>;
|
||||
|
||||
|
||||
var<workgroup> 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);
|
||||
});
|
||||
|
|
@ -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<uniform, buffer>) -> u32 {
|
||||
return bufferLength(p);
|
||||
}
|
||||
fn roStorageUnsized(p : ptr<storage, buffer>) -> u32 {
|
||||
return bufferLength(p);
|
||||
}
|
||||
fn storageUnsized(p : ptr<storage, buffer, read_write>) -> u32 {
|
||||
return bufferLength(p);
|
||||
}
|
||||
fn workgroupUnsized(p : ptr<workgroup, buffer>) -> 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<uniform, buffer<${t.params.size / 2}>>) -> u32 {
|
||||
return bufferLength(p);
|
||||
}
|
||||
fn roStorageSized(p : ptr<storage, buffer<${t.params.size / 2}>>) -> u32 {
|
||||
return bufferLength(p);
|
||||
}
|
||||
fn storageSized(p : ptr<storage, buffer<${t.params.size / 2}>, read_write>) -> u32 {
|
||||
return bufferLength(p);
|
||||
}
|
||||
fn workgroupSized(p : ptr<workgroup, buffer<${t.params.size / 2}>>) -> 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> uniform_buffer : buffer<${t.params.size}>;
|
||||
@group(0) @binding(1) var<storage> ro_storage_buffer : buffer<${t.params.size}>;
|
||||
@group(0) @binding(2) var<storage, read_write> storage_buffer : buffer<${t.params.size}>;
|
||||
var<workgroup> workgroup_buffer : buffer<${t.params.size}>;
|
||||
|
||||
@group(0) @binding(3) var<storage, read_write> out : array<u32, 4>;
|
||||
|
||||
${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<storage, buffer>) -> u32 {
|
||||
return bufferLength(p);
|
||||
}
|
||||
fn storageUnsized(p : ptr<storage, buffer, read_write>) -> u32 {
|
||||
return bufferLength(p);
|
||||
}`;
|
||||
assigns[0] = 'roStorageUnsized(&ro_storage_buffer)';
|
||||
assigns[1] = 'storageUnsized(&storage_buffer)';
|
||||
break;
|
||||
}
|
||||
|
||||
const wgsl = `
|
||||
@group(0) @binding(0) var<storage> ro_storage_buffer : buffer;
|
||||
@group(0) @binding(1) var<storage, read_write> storage_buffer : buffer;
|
||||
|
||||
@group(0) @binding(2) var<storage, read_write> out : array<u32, 4>;
|
||||
|
||||
${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<workgroup> v : ${type};\n@group(0) @binding(1) var<storage> dummy : u32;`;
|
||||
break;
|
||||
case 'uniform':
|
||||
decl = `@group(0) @binding(1) var<uniform> v: ${type};`;
|
||||
break;
|
||||
case 'storage':
|
||||
decl = `@group(0) @binding(1) var<storage, read_write> v : ${type};`;
|
||||
break;
|
||||
case 'ro_storage':
|
||||
decl = `@group(0) @binding(1) var<storage> v : ${type};`;
|
||||
break;
|
||||
}
|
||||
|
||||
const wgsl = `
|
||||
@group(0) @binding(0) var<storage, read_write> 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);
|
||||
});
|
||||
|
|
@ -0,0 +1,584 @@
|
|||
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<workgroup> v : ${buffer_ty};\n@group(0) @binding(0) var<storage> dummy : array<u32>;`;
|
||||
break;
|
||||
case 'uniform':
|
||||
decl = `@group(0) @binding(0) var<uniform> v : ${buffer_ty};`;
|
||||
break;
|
||||
case 'storage':
|
||||
decl = `@group(0) @binding(0) var<storage> 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<storage> offsets : array<u32>;
|
||||
@group(0) @binding(2) var<storage> sizes : array<u32>; // unused
|
||||
@group(0) @binding(3) var<storage, read_write> out : array<u32>;
|
||||
|
||||
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<workgroup> v : ${buffer_ty};\n@group(0) @binding(0) var<storage> dummy : array<u32>;`;
|
||||
break;
|
||||
case 'uniform':
|
||||
decl = `@group(0) @binding(0) var<uniform> v : ${buffer_ty};`;
|
||||
break;
|
||||
case 'storage':
|
||||
decl = `@group(0) @binding(0) var<storage> 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);
|
||||
}
|
||||
|
||||
fn sized_sized_indirectLength(p : ptr<${t.params.aspace}, buffer<${
|
||||
t.params.bufferSize
|
||||
}>>, gidx : u32) -> u32 {
|
||||
return sizedLength(p, gidx);
|
||||
}
|
||||
|
||||
override bufferSize : u32 = 0u;
|
||||
${decl}
|
||||
@group(0) @binding(1) var<storage> offsets : array<u32>;
|
||||
@group(0) @binding(2) var<storage> sizes : array<u32>; // unused
|
||||
@group(0) @binding(3) var<storage, read_write> out : array<u32>;
|
||||
|
||||
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<storage, read_write> out : u32;
|
||||
|
||||
var<workgroup> wg_var : buffer<256 * 4>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let in_ptr = bufferView<array<vec4u, 256 / 4>>(&in, 0);
|
||||
let wg_ptr = bufferView<array<vec4u, 256 / 4>>(&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<array<u32>>(&out, 0);
|
||||
let wg_ptr = bufferView<array<u32>>(&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<storage> in : u32;
|
||||
@group(0) @binding(1) var<storage, read_write> out : buffer<256 * 4>;
|
||||
|
||||
|
||||
var<workgroup> 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<storage, read_write> out : ${ty.toString()};
|
||||
|
||||
var<workgroup> wg_var : buffer<${bufferSize}>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let in_ptr = bufferView<array<vec4u, ${bufferSize} / 16>>(&in, 0);
|
||||
let wg_ptr = bufferView<array<vec4u, ${bufferSize} / 16>>(&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<array<u32>>(&out, 0);
|
||||
let wg_ptr = bufferView<array<u32>>(&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<storage> in : ${ty.toString()};
|
||||
@group(0) @binding(1) var<storage, read_write> out : buffer<${bufferSize}>;
|
||||
|
||||
|
||||
var<workgroup> 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);
|
||||
});
|
||||
|
|
@ -0,0 +1,975 @@
|
|||
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', 'sized_sized_indirect'] as const;
|
||||
|
||||
interface ArrayLengthType {
|
||||
type: string;
|
||||
access?: string;
|
||||
align: number;
|
||||
stride: number;
|
||||
arrayOffset?: number;
|
||||
f16?: boolean;
|
||||
uniformStdLayout?: boolean;
|
||||
}
|
||||
|
||||
/**
|
||||
* Return true if the case is valid
|
||||
*
|
||||
* @param type The testcase
|
||||
* @param offset The offset to the view call
|
||||
* @param size The size to the view call
|
||||
* @param bufferSize The size of the buffer
|
||||
* @returns true if case is valid.
|
||||
*/
|
||||
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;
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Calculates the return length in elements
|
||||
*
|
||||
* @param type The testcase
|
||||
* @param offset The offset to the view call
|
||||
* @param size The size to the view call
|
||||
* @param bufferSize The buffer size
|
||||
* @returns The number of elements in the arrayLength call
|
||||
*/
|
||||
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<u32, 5>
|
||||
}
|
||||
struct S_arr {
|
||||
a : vec4u,
|
||||
b : array<u32>,
|
||||
}
|
||||
struct T_arr {
|
||||
a : S,
|
||||
b : array<T>,
|
||||
}
|
||||
`;
|
||||
|
||||
export const kArrayLengthTypes: Record<string, ArrayLengthType> = {
|
||||
// Scalars and vectors
|
||||
u32: {
|
||||
type: 'array<u32>',
|
||||
align: 4,
|
||||
stride: 4,
|
||||
uniformStdLayout: true,
|
||||
},
|
||||
vec2f: {
|
||||
type: 'array<vec2f>',
|
||||
align: 8,
|
||||
stride: 8,
|
||||
uniformStdLayout: true,
|
||||
},
|
||||
vec3i: {
|
||||
type: 'array<vec3i>',
|
||||
align: 16,
|
||||
stride: 16,
|
||||
},
|
||||
vec4u: {
|
||||
type: 'array<vec4u>',
|
||||
align: 16,
|
||||
stride: 16,
|
||||
},
|
||||
f16: {
|
||||
type: 'array<f16>',
|
||||
align: 2,
|
||||
stride: 2,
|
||||
f16: true,
|
||||
uniformStdLayout: true,
|
||||
},
|
||||
vec2h: {
|
||||
type: 'array<vec2h>',
|
||||
align: 4,
|
||||
stride: 4,
|
||||
f16: true,
|
||||
uniformStdLayout: true,
|
||||
},
|
||||
vec3h: {
|
||||
type: 'array<vec3h>',
|
||||
align: 8,
|
||||
stride: 8,
|
||||
f16: true,
|
||||
uniformStdLayout: true,
|
||||
},
|
||||
vec4h: {
|
||||
type: 'array<vec4h>',
|
||||
align: 8,
|
||||
stride: 8,
|
||||
f16: true,
|
||||
uniformStdLayout: true,
|
||||
},
|
||||
// Matrices
|
||||
mat2x2f: {
|
||||
type: 'array<mat2x2f>',
|
||||
align: 8,
|
||||
stride: 16,
|
||||
},
|
||||
mat2x3f: {
|
||||
type: 'array<mat2x3f>',
|
||||
align: 16,
|
||||
stride: 32,
|
||||
},
|
||||
mat2x4f: {
|
||||
type: 'array<mat2x4f>',
|
||||
align: 16,
|
||||
stride: 32,
|
||||
},
|
||||
mat4x2f: {
|
||||
type: 'array<mat4x2f>',
|
||||
align: 8,
|
||||
stride: 32,
|
||||
},
|
||||
mat4x3f: {
|
||||
type: 'array<mat4x3f>',
|
||||
align: 16,
|
||||
stride: 64,
|
||||
},
|
||||
mat4x4f: {
|
||||
type: 'array<mat4x4f>',
|
||||
align: 16,
|
||||
stride: 64,
|
||||
},
|
||||
mat3x2h: {
|
||||
type: 'array<mat3x2h>',
|
||||
align: 4,
|
||||
stride: 12,
|
||||
f16: true,
|
||||
},
|
||||
mat3x3h: {
|
||||
type: 'array<mat3x3h>',
|
||||
align: 8,
|
||||
stride: 24,
|
||||
f16: true,
|
||||
uniformStdLayout: true,
|
||||
},
|
||||
mat3x4h: {
|
||||
type: 'array<mat3x4h>',
|
||||
align: 8,
|
||||
stride: 24,
|
||||
f16: true,
|
||||
uniformStdLayout: true,
|
||||
},
|
||||
// Structs without arrays
|
||||
S: {
|
||||
type: 'array<S>',
|
||||
align: 16,
|
||||
stride: 32,
|
||||
},
|
||||
T: {
|
||||
type: 'array<T>',
|
||||
align: 4,
|
||||
stride: 24,
|
||||
uniformStdLayout: true,
|
||||
},
|
||||
// Sized arrays
|
||||
array_u32_4: {
|
||||
type: 'array<array<u32, 4>>',
|
||||
align: 4,
|
||||
stride: 16,
|
||||
uniformStdLayout: true,
|
||||
},
|
||||
array_f16_6: {
|
||||
type: 'array<array<f16, 6>>',
|
||||
align: 2,
|
||||
stride: 12,
|
||||
f16: true,
|
||||
uniformStdLayout: true,
|
||||
},
|
||||
array_vec2f_2: {
|
||||
type: 'array<array<vec2f, 2>>',
|
||||
align: 8,
|
||||
stride: 16,
|
||||
uniformStdLayout: true,
|
||||
},
|
||||
array_S_2: {
|
||||
type: 'array<array<S, 2>>',
|
||||
align: 16,
|
||||
stride: 64,
|
||||
},
|
||||
array_T_2: {
|
||||
type: 'array<array<T, 2>>',
|
||||
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<string, LayoutCase> = {
|
||||
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<S_stride, 4>',
|
||||
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);
|
||||
}
|
||||
|
|
@ -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<string, Case> = {
|
|||
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<storage> b : buffer;`,
|
||||
},
|
||||
buffer_sized: {
|
||||
code: `let x : buffer<128> = b;`,
|
||||
valid: false,
|
||||
decls: `@group(0) @binding(0) var<uniform> b : buffer<128>;`,
|
||||
},
|
||||
buffer_override_sized: {
|
||||
code: `let x : buffer<o> = b;`,
|
||||
valid: false,
|
||||
decls: `override o = 16u; var<workgroup> b : buffer<o>;`,
|
||||
},
|
||||
ptr_buffer: {
|
||||
code: `let x : ptr<storage, buffer> = &b;`,
|
||||
valid: 'buffer_view',
|
||||
decls: `@group(0) @binding(0) var<storage> b : buffer;`,
|
||||
},
|
||||
ptr_buffer_sized: {
|
||||
code: `let x : ptr<uniform, buffer<128>> = &b;`,
|
||||
valid: 'buffer_view',
|
||||
decls: `@group(0) @binding(0) var<uniform> b : buffer<128>;`,
|
||||
},
|
||||
ptr_buffer_override_sized: {
|
||||
code: `let x : ptr<workgroup, buffer<o>> = &b;`,
|
||||
valid: 'buffer_view',
|
||||
decls: `override o = 16u; var<workgroup> b : buffer<o>;`,
|
||||
},
|
||||
};
|
||||
|
||||
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,21 @@ const kInitCases: Record<string, Case> = {
|
|||
code: `var x = 1;\nlet y = x << 1;`,
|
||||
valid: true,
|
||||
},
|
||||
buffer_mismatch: {
|
||||
code: `let x : ptr<uniform, buffer> = &b;`,
|
||||
valid: false,
|
||||
decls: `@group(0) @binding(0) var<uniform> b : buffer<128>;`,
|
||||
},
|
||||
buffer_smaller_size: {
|
||||
code: `let x : ptr<uniform, buffer<64>> = &b;`,
|
||||
valid: false,
|
||||
decls: `@group(0) @binding(0) var<uniform> b : buffer<128>;`,
|
||||
},
|
||||
buffer_mismatched_overrides: {
|
||||
code: `let p : ptr<workgroup, buffer<o1>> = &b;`,
|
||||
valid: false,
|
||||
decls: `override o1 : u32;\noverride o2 : u32;\nvar<workgroup> b : buffer<o2>;`,
|
||||
},
|
||||
};
|
||||
|
||||
g.test('initializer')
|
||||
|
|
|
|||
|
|
@ -29,30 +29,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.
|
||||
|
|
@ -61,30 +66,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.
|
||||
|
|
@ -93,12 +103,14 @@ const kTypes = {
|
|||
isConstructible: true,
|
||||
isFixedFootprint: true,
|
||||
requiresF16: false,
|
||||
requiresBufferView: false,
|
||||
},
|
||||
mat3x4h: {
|
||||
isHostShareable: true,
|
||||
isConstructible: true,
|
||||
isFixedFootprint: true,
|
||||
requiresF16: true,
|
||||
requiresBufferView: false,
|
||||
},
|
||||
|
||||
// Atomics.
|
||||
|
|
@ -107,12 +119,14 @@ const kTypes = {
|
|||
isConstructible: false,
|
||||
isFixedFootprint: true,
|
||||
requiresF16: false,
|
||||
requiresBufferView: false,
|
||||
},
|
||||
'atomic<u32>': {
|
||||
isHostShareable: true,
|
||||
isConstructible: false,
|
||||
isFixedFootprint: true,
|
||||
requiresF16: false,
|
||||
requiresBufferView: false,
|
||||
},
|
||||
|
||||
// Arrays.
|
||||
|
|
@ -121,36 +135,42 @@ const kTypes = {
|
|||
isConstructible: false,
|
||||
isFixedFootprint: false,
|
||||
requiresF16: false,
|
||||
requiresBufferView: false,
|
||||
},
|
||||
'array<vec4<bool>, 4>': {
|
||||
isHostShareable: false,
|
||||
isConstructible: true,
|
||||
isFixedFootprint: true,
|
||||
requiresF16: false,
|
||||
requiresBufferView: false,
|
||||
},
|
||||
'array<vec4u>': {
|
||||
isHostShareable: true,
|
||||
isConstructible: false,
|
||||
isFixedFootprint: false,
|
||||
requiresF16: false,
|
||||
requiresBufferView: false,
|
||||
},
|
||||
'array<vec4u, 4>': {
|
||||
isHostShareable: true,
|
||||
isConstructible: true,
|
||||
isFixedFootprint: true,
|
||||
requiresF16: false,
|
||||
requiresBufferView: false,
|
||||
},
|
||||
'array<vec4u, array_size_const>': {
|
||||
isHostShareable: true,
|
||||
isConstructible: true,
|
||||
isFixedFootprint: true,
|
||||
requiresF16: false,
|
||||
requiresBufferView: false,
|
||||
},
|
||||
'array<vec4u, array_size_override>': {
|
||||
isHostShareable: false,
|
||||
isConstructible: false,
|
||||
isFixedFootprint: true,
|
||||
requiresF16: false,
|
||||
requiresBufferView: false,
|
||||
},
|
||||
|
||||
// Structures.
|
||||
|
|
@ -159,36 +179,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.
|
||||
|
|
@ -197,18 +223,44 @@ const kTypes = {
|
|||
isConstructible: false,
|
||||
isFixedFootprint: false,
|
||||
requiresF16: false,
|
||||
requiresBufferView: false,
|
||||
},
|
||||
sampler: {
|
||||
isHostShareable: false,
|
||||
isConstructible: false,
|
||||
isFixedFootprint: false,
|
||||
requiresF16: false,
|
||||
requiresBufferView: false,
|
||||
},
|
||||
'texture_2d<f32>': {
|
||||
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<array_size_override>': {
|
||||
isHostShareable: false,
|
||||
isConstructible: false,
|
||||
isFixedFootprint: true,
|
||||
requiresF16: false,
|
||||
requiresBufferView: true,
|
||||
},
|
||||
};
|
||||
|
||||
|
|
@ -288,7 +340,8 @@ g.test('module_scope_types')
|
|||
break;
|
||||
case 'uniform':
|
||||
decl = '@group(0) @binding(0) var<uniform> foo : ';
|
||||
shouldPass = type.isHostShareable && type.isConstructible;
|
||||
shouldPass =
|
||||
type.isHostShareable && (type.isConstructible || t.params.type === 'buffer<128>');
|
||||
break;
|
||||
case 'workgroup':
|
||||
decl = 'var<workgroup> foo : ';
|
||||
|
|
@ -319,6 +372,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);
|
||||
});
|
||||
|
||||
|
|
@ -368,6 +422,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);
|
||||
});
|
||||
|
||||
|
|
|
|||
|
|
@ -25,6 +25,20 @@ fn main() {
|
|||
t.expectCompileResult(false, code);
|
||||
});
|
||||
|
||||
g.test('buffer_type')
|
||||
.desc('Validates buffers are rejected')
|
||||
.params(u => u.combine('sized', [false, true] as const))
|
||||
.fn(t => {
|
||||
const code = `
|
||||
@group(0) @binding(0) var<storage> b : buffer${t.params.sized ? '<1024>' : ''};
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
_ = arrayLength(&b);
|
||||
}`;
|
||||
|
||||
t.expectCompileResult(false, code);
|
||||
});
|
||||
|
||||
const atomic_types = ['u32', 'i32'].map(j => `atomic<${j}>`);
|
||||
const vec_types = [2, 3, 4]
|
||||
.map(i => ['i32', 'u32', 'f32', 'f16'].map(j => `vec${i}<${j}>`))
|
||||
|
|
|
|||
|
|
@ -0,0 +1,540 @@
|
|||
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<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
${t.params.must_use ? 'let p = ' : ''}bufferArrayView<array<u32>>(&v, 0, 64);
|
||||
}`;
|
||||
|
||||
t.expectCompileResult(t.params.must_use, wgsl);
|
||||
});
|
||||
|
||||
const kTypes = objectsToRecord(kAllScalarsAndVectors);
|
||||
|
||||
g.test('return_store_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<storage, ${ret_type}>`;
|
||||
}
|
||||
let enables = ``;
|
||||
if (template_type.requiresF16() || return_type.requiresF16()) {
|
||||
enables = `enable f16;`;
|
||||
}
|
||||
const wgsl = `
|
||||
${enables}
|
||||
@group(0) @binding(0) var<storage> 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 aspace = t.params.type.substr(t.params.type.lastIndexOf('_') + 1);
|
||||
let access = '';
|
||||
switch (t.params.type) {
|
||||
case 'unsized_ro_storage':
|
||||
case 'sized_ro_storage':
|
||||
access = ', read';
|
||||
break;
|
||||
case 'unsized_storage':
|
||||
case 'sized_storage':
|
||||
access = ', read_write';
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
const wgsl = `
|
||||
@group(0) @binding(0) var<storage> unsized_ro_storage : buffer;
|
||||
@group(0) @binding(1) var<storage, read_write> unsized_storage : buffer;
|
||||
@group(0) @binding(2) var<storage> sized_ro_storage : buffer<128>;
|
||||
@group(0) @binding(3) var<storage, read_write> sized_storage : buffer<128>;
|
||||
@group(0) @binding(4) var<uniform> sized_uniform : buffer<128>;
|
||||
var<workgroup> sized_workgroup : buffer<128>;
|
||||
override o : u32;
|
||||
var<workgroup> override_workgroup : buffer<o>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p : ptr<${aspace}, array<u32>${access}> = bufferArrayView<array<u32>>(${
|
||||
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<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferArrayView<array<u32>>(&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<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferArrayView<array<u32>>(&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<string, GPUPipelineConstantValue>;
|
||||
ptr_param?: boolean;
|
||||
}
|
||||
|
||||
const kEarlyEvalCases: Record<string, EarlyEvalCase> = {
|
||||
ro_storage_buffer: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer<16>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, 0, 16);
|
||||
}`,
|
||||
valid: true,
|
||||
},
|
||||
storage_buffer: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer<16>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, 0, 16);
|
||||
}`,
|
||||
valid: true,
|
||||
},
|
||||
uniform_buffer: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<uniform> v : buffer<16>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, 0, 16);
|
||||
}`,
|
||||
valid: true,
|
||||
},
|
||||
workgroup_buffer: {
|
||||
code: `
|
||||
var<workgroup> v : buffer<16>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, 0, 16);
|
||||
}`,
|
||||
valid: true,
|
||||
},
|
||||
const_size_too_small_for_type: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, 0, 12);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
buffer_too_small_for_const_size: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer<12>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<u32>>(&v, 0, 16);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
buffer_too_small_for_type: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer<12>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, 0, 16);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
const_size_too_small_for_const_offset_and_type: {
|
||||
code: `
|
||||
var<workgroup> v : buffer<128>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, 16, 30);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
buffer_size_too_small_for_const_offset_and_type: {
|
||||
code: `
|
||||
var<workgroup> v : buffer<30>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, 16, 128);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
override_size_too_small_for_type: {
|
||||
code: `
|
||||
override size : u32;
|
||||
@group(0) @binding(0) var<storage> v : buffer;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, 0, size);
|
||||
}`,
|
||||
valid: 'pipeline',
|
||||
constants: { size: 12 },
|
||||
},
|
||||
buffer_too_small_for_override_size: {
|
||||
code: `
|
||||
override buffer_size : u32;
|
||||
override size : u32;
|
||||
var<workgroup> v : buffer<buffer_size>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<u32>>(&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<workgroup> v : buffer<buffer_size>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<vec4u>>(&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<uniform> v : buffer<128>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<vec4u>>(&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<workgroup> v : buffer<buffer_size>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<array<vec4u>>(&v, offset, 128);
|
||||
}`,
|
||||
valid: 'pipeline',
|
||||
constants: { buffer_size: 30, offset: 16 },
|
||||
},
|
||||
const_offset_not_aligned: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
fn foo() {
|
||||
let p = bufferarrayview<array<vec4u>>(&v, 12, 128);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
override_offset_not_aligned: {
|
||||
code: `
|
||||
override offset : u32;
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
fn foo() {
|
||||
let p = bufferarrayview<array<vec4u>>(&v, offset, 128);
|
||||
}`,
|
||||
valid: false,
|
||||
constants: { offset: 12 },
|
||||
},
|
||||
const_offset_plus_type_size_out_of_range: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
struct S {
|
||||
a: vec4u,
|
||||
b: vec4f,
|
||||
}
|
||||
fn foo() {
|
||||
let p = buffer<array<S>>(&v, 4294967279, 128);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
override_offset_plus_type_size_out_of_range: {
|
||||
code: `
|
||||
override offset : u32;
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
struct S {
|
||||
a: vec4u,
|
||||
b: vec4f,
|
||||
}
|
||||
fn foo() {
|
||||
let p = buffer<array<S>>(&v, offset, 128);
|
||||
}`,
|
||||
valid: false,
|
||||
constants: { offset: 4294967279 },
|
||||
},
|
||||
const_multiple_of_stride: {
|
||||
code: `
|
||||
struct S {
|
||||
a : vec4u,
|
||||
b : array<vec2u>,
|
||||
}
|
||||
var<workgroup> v : buffer<256>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<S>(&v, 0, 24);
|
||||
}`,
|
||||
valid: true,
|
||||
},
|
||||
const_not_multiple_of_stride: {
|
||||
code: `
|
||||
struct S {
|
||||
a : vec4u,
|
||||
b : array<vec2u>,
|
||||
}
|
||||
var<workgroup> v : buffer<256>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<S>(&v, 0, 26);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
override_not_multiple_of_stride: {
|
||||
code: `
|
||||
override size : u32;
|
||||
struct S {
|
||||
a : vec4u,
|
||||
b : array<vec2u>,
|
||||
}
|
||||
var<workgroup> v : buffer<256>;
|
||||
fn foo() {
|
||||
let p = bufferArrayView<S>(&v, 0, size);
|
||||
}`,
|
||||
valid: 'pipeline',
|
||||
constants: { size: 26 },
|
||||
},
|
||||
const_buffer_too_small_through_unsized_function: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> buffer<16>;
|
||||
fn bar(p : ptr<storage, buffer>) {
|
||||
let q = bufferArrayView<array<u32>>(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<storage> v : buffer<16>;
|
||||
fn bar(p : ptr<storage, buffer>) {
|
||||
let q = bufferArrayView<array<u32>>(p, 0, size);
|
||||
}
|
||||
fn foo() {
|
||||
bar(&v);
|
||||
}`,
|
||||
valid: 'pipeline',
|
||||
constants: { size: 32 },
|
||||
ptr_param: true,
|
||||
},
|
||||
sized_buffer_to_smaller_buffer_too_small: {
|
||||
code: `
|
||||
var<workgroup> v : buffer<256>;
|
||||
fn bar(p : ptr<workgroup, buffer<16>>) {
|
||||
let q = bufferArrayView<array<u32>>(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<string, LHSCallCase> = {
|
||||
u32_array: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
(*bufferArrayView<array<u32>>(&v, 0, 128))[0] = 123u;
|
||||
}`,
|
||||
},
|
||||
f32_vector_letter: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
(*bufferArrayView<array<vec4f>>(&v, 0, 128))[0].x = 13.3f;
|
||||
}`,
|
||||
},
|
||||
f32_struct_access: {
|
||||
code: `
|
||||
struct S {
|
||||
member_a: f32,
|
||||
}
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
(*bufferArrayView<array<S>>(&v, 0, 128))[0].member_a = 13.3f;
|
||||
}`,
|
||||
},
|
||||
u32_array_ptr_composite: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
bufferArrayView<array<u32>>(&v, 0, 128)[0] = 123u;
|
||||
}`,
|
||||
requires: ['pointer_composite_access'],
|
||||
},
|
||||
vec4u_swizzle_assign: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
(*bufferArrayView<array<vec4u>>(&v, 0, 128))[0].zxy = vec3u(42);
|
||||
}`,
|
||||
requires: ['swizzle_assignment'],
|
||||
},
|
||||
vec4u_ptr_composite_swizzle_assign: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
bufferArrayView<array<vec4u>>(&v, 0, 128)[0].zxy = vec3u(42);
|
||||
}`,
|
||||
requires: ['pointer_composite_access', 'swizzle_assignment'],
|
||||
},
|
||||
compound_assign: {
|
||||
code: `
|
||||
var<workgroup> v : buffer<256>;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
(*bufferArrayView<array<u32>>(&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);
|
||||
});
|
||||
|
|
@ -0,0 +1,115 @@
|
|||
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<storage> 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<storage> 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<storage> unsized_ro_storage : buffer;
|
||||
@group(0) @binding(1) var<storage, read_write> unsized_storage : buffer;
|
||||
@group(0) @binding(2) var<storage> sized_ro_storage : buffer<128>;
|
||||
@group(0) @binding(3) var<storage, read_write> sized_storage : buffer<128>;
|
||||
@group(0) @binding(4) var<uniform> sized_uniform : buffer<128>;
|
||||
var<workgroup> sized_workgroup : buffer<128>;
|
||||
override o : u32;
|
||||
var<workgroup> override_workgroup : buffer<o>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
_ = bufferLength(${t.params.ptr ? '&' : ''}${t.params.type});
|
||||
}`;
|
||||
|
||||
t.expectCompileResult(t.params.ptr, wgsl);
|
||||
});
|
||||
|
||||
const kInvalidCases = {
|
||||
baseline: `&b`,
|
||||
no_args: `bufferLength()`,
|
||||
two_args: `bufferLength(&b, 0)`,
|
||||
scalar: `bufferLength(0u)`,
|
||||
ptr_to_scalar: `bufferLength(&v[0])`,
|
||||
ptr_to_array: `bufferLength(&v)`,
|
||||
};
|
||||
|
||||
g.test('invalid_cases')
|
||||
.desc('Test a smattering of invalid bufferLength calls')
|
||||
.params(u => u.combine('test', keysOf(kInvalidCases)))
|
||||
.fn(t => {
|
||||
t.skipIfLanguageFeatureNotSupported('buffer_view');
|
||||
const code = `
|
||||
var<workgroup> v : array<u32, 4>;
|
||||
var<workgroup> b : buffer<128>;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
_ = ${kInvalidCases[t.params.test]};
|
||||
}`;
|
||||
|
||||
t.expectCompileResult(t.params.test === 'baseline', code);
|
||||
});
|
||||
|
|
@ -0,0 +1,451 @@
|
|||
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<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
${t.params.must_use ? 'let p = ' : ''}bufferView<u32>(&v, 0);
|
||||
}`;
|
||||
|
||||
t.expectCompileResult(t.params.must_use, wgsl);
|
||||
});
|
||||
|
||||
const kTypes = objectsToRecord(kAllScalarsAndVectors);
|
||||
|
||||
g.test('return_store_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<storage, ${ret_type}>`;
|
||||
}
|
||||
let enables = ``;
|
||||
if (template_type.requiresF16() || return_type.requiresF16()) {
|
||||
enables = `enable f16;`;
|
||||
}
|
||||
const wgsl = `
|
||||
${enables}
|
||||
@group(0) @binding(0) var<storage> 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 aspace = t.params.type.substr(t.params.type.lastIndexOf('_') + 1);
|
||||
let access = '';
|
||||
switch (t.params.type) {
|
||||
case 'unsized_ro_storage':
|
||||
case 'sized_ro_storage':
|
||||
access = ', read';
|
||||
break;
|
||||
case 'unsized_storage':
|
||||
case 'sized_storage':
|
||||
access = ', read_write';
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
const wgsl = `
|
||||
@group(0) @binding(0) var<storage> unsized_ro_storage : buffer;
|
||||
@group(0) @binding(1) var<storage, read_write> unsized_storage : buffer;
|
||||
@group(0) @binding(2) var<storage> sized_ro_storage : buffer<128>;
|
||||
@group(0) @binding(3) var<storage, read_write> sized_storage : buffer<128>;
|
||||
@group(0) @binding(4) var<uniform> sized_uniform : buffer<128>;
|
||||
var<workgroup> sized_workgroup : buffer<128>;
|
||||
override o : u32;
|
||||
var<workgroup> override_workgroup : buffer<o>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p : ptr<${aspace}, u32${access}> = bufferView<u32>(${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<storage> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
let p = bufferView<u32>(&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<string, GPUPipelineConstantValue>;
|
||||
ptr_param?: boolean;
|
||||
}
|
||||
|
||||
const kEarlyEvalCases: Record<string, EarlyEvalCase> = {
|
||||
ro_storage_buffer: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer<16>;
|
||||
fn foo() {
|
||||
let p = bufferView<vec4u>(&v, 0);
|
||||
}`,
|
||||
valid: true,
|
||||
},
|
||||
storage_buffer: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer<16>;
|
||||
fn foo() {
|
||||
let p = bufferView<vec4u>(&v, 0);
|
||||
}`,
|
||||
valid: true,
|
||||
},
|
||||
uniform_buffer: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<uniform> v : buffer<16>;
|
||||
fn foo() {
|
||||
let p = bufferView<vec4u>(&v, 0);
|
||||
}`,
|
||||
valid: true,
|
||||
},
|
||||
workgroup_buffer: {
|
||||
code: `
|
||||
var<workgroup> v : buffer<16>;
|
||||
fn foo() {
|
||||
let p = bufferView<vec4u>(&v, 0);
|
||||
}`,
|
||||
valid: true,
|
||||
},
|
||||
ro_storage_buffer_too_small_for_type: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer<8>;
|
||||
fn foo() {
|
||||
let p = bufferView<vec4u>(&v, 0);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
uniform_buffer_too_small_with_const_offset: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<uniform> v : buffer<24>;
|
||||
fn foo() {
|
||||
let p = bufferView<vec4u>(&v, 16);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
uniform_buffer_too_small_with_override_offset: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<uniform> v : buffer<24>;
|
||||
override offset : u32;
|
||||
fn foo() {
|
||||
let p = bufferView<vec4u>(&v, offset);
|
||||
}`,
|
||||
valid: 'pipeline',
|
||||
constants: { offset: 16 },
|
||||
},
|
||||
workgroup_buffer_too_small_with_override: {
|
||||
code: `
|
||||
override size : u32;
|
||||
var<workgroup> v : buffer<size>;
|
||||
fn foo() {
|
||||
let p = bufferView<vec4u>(&v, 0);
|
||||
}`,
|
||||
valid: 'pipeline',
|
||||
constants: { size: 12 },
|
||||
},
|
||||
storage_buffer_too_small_for_one_array_element: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer<28>;
|
||||
struct S {
|
||||
a: vec4u,
|
||||
b: vec4f,
|
||||
}
|
||||
fn foo() {
|
||||
let p = bufferView<array<S>>(&v, 0);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
storage_buffer_out_of_range_u32: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
struct S {
|
||||
a: vec4u,
|
||||
b: vec4f,
|
||||
}
|
||||
fn foo() {
|
||||
let p = buffer<S>(&v, 4294967279);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
uniform_buffer_too_small_through_unsized_function: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<uniform> v : buffer<28>;
|
||||
fn bar(p : ptr<uniform, buffer>) {
|
||||
let q = bufferView<vec4u>(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<uniform> v : buffer<28>;
|
||||
override offset : u32;
|
||||
fn bar(p : ptr<uniform, buffer>) {
|
||||
let q = bufferView<vec4u>(p, offset);
|
||||
}
|
||||
fn foo() {
|
||||
bar(&v);
|
||||
}`,
|
||||
valid: 'pipeline',
|
||||
constants: { offset: 16 },
|
||||
ptr_param: true,
|
||||
},
|
||||
workgroup_to_smaller_size_param: {
|
||||
code: `
|
||||
var<workgroup> v : buffer<128>;
|
||||
fn bar(p : ptr<workgroup, buffer<16>) {
|
||||
let q = bufferView<array<vec2u, 3>(p, 0);
|
||||
}
|
||||
fn foo() {
|
||||
bar(&v);
|
||||
}`,
|
||||
valid: false,
|
||||
ptr_param: true,
|
||||
},
|
||||
ro_storage_unsized_buffer_to_sized_param: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage> v : buffer;
|
||||
fn bar(p : ptr<storage, buffer<16>) {
|
||||
let q = bufferView<array<vec2u, 3>(p, 0);
|
||||
}
|
||||
fn foo() {
|
||||
bar(&v);
|
||||
}`,
|
||||
valid: false,
|
||||
ptr_param: true,
|
||||
},
|
||||
offset_not_aligned: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<uniform> v : buffer<128>;
|
||||
fn foo() {
|
||||
let p = bufferView<vec4u>(&v, 12);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
offset_not_aligned_override: {
|
||||
code: `
|
||||
override offset : u32;
|
||||
@group(0) @binding(0) var<uniform> v : buffer<128>;
|
||||
fn foo() {
|
||||
let p = bufferView<vec4u>(&v, offset);
|
||||
}`,
|
||||
valid: 'pipeline',
|
||||
constants: { offset: 12 },
|
||||
},
|
||||
offset_negative: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<uniform> v : buffer<128>;
|
||||
fn foo() {
|
||||
let p = bufferView<vec4u>(&v, -16);
|
||||
}`,
|
||||
valid: false,
|
||||
},
|
||||
offset_negative_override: {
|
||||
code: `
|
||||
override offset : u32;
|
||||
@group(0) @binding(0) var<uniform> v : buffer<128>;
|
||||
fn foo() {
|
||||
let p = bufferView<vec4u>(&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<string, LHSCallCase> = {
|
||||
u32: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
*bufferView<u32>(&v, 0) = 123u;
|
||||
}`,
|
||||
},
|
||||
u32_array: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
(*bufferView<array<u32, 4>>(&v, 0))[0] = 123u;
|
||||
}`,
|
||||
},
|
||||
f32_vector_letter: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
(*bufferView<vec4f>(&v, 0)).x = 13.3f;
|
||||
}`,
|
||||
},
|
||||
f32_struct_access: {
|
||||
code: `
|
||||
struct S {
|
||||
member_a: f32,
|
||||
}
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
(*bufferView<S>(&v, 0)).member_a = 13.3f;
|
||||
}`,
|
||||
},
|
||||
u32_array_ptr_composite: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
bufferView<array<u32, 4>>(&v, 0)[0] = 123u;
|
||||
}`,
|
||||
requires: ['pointer_composite_access'],
|
||||
},
|
||||
vec4u_swizzle_assign: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
(*bufferView<vec4u>(&v, 0)).zxy = vec3u(42);
|
||||
}`,
|
||||
requires: ['swizzle_assignment'],
|
||||
},
|
||||
vec4u_ptr_composite_swizzle_assign: {
|
||||
code: `
|
||||
@group(0) @binding(0) var<storage, read_write> v : buffer;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
bufferView<vec4u>(&v, 0).zxy = vec3u(42);
|
||||
}`,
|
||||
requires: ['pointer_composite_access', 'swizzle_assignment'],
|
||||
},
|
||||
compound_assign: {
|
||||
code: `
|
||||
var<workgroup> v : buffer<256>;
|
||||
@compute @workgroup_size(1)
|
||||
fn main() {
|
||||
*bufferView<u32>(&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);
|
||||
});
|
||||
|
|
@ -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<string, ParamTypeCase> = {
|
||||
|
|
@ -304,6 +304,68 @@ const kFunctionParamTypeCases: Record<string, ParamTypeCase> = {
|
|||
invalid_ptr6: { name: `ptr<private,u32,read_write>`, valid: false }, // Can't specify access mode
|
||||
invalid_ptr7: { name: `ptr<private,clamp>`, valid: false }, // Invalid store type
|
||||
invalid_ptr8: { name: `ptr<function, texture_external>`, valid: false }, // non-constructible pointer type
|
||||
|
||||
// Buffers (need buffer_view and unrestricted_pointer_parameters)
|
||||
ptrBufferUnsized_storage: {
|
||||
name: `ptr<storage, buffer, read_write>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferSized_storage: {
|
||||
name: `ptr<storage, buffer<128>, read_write>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferSizedSmall_storage: {
|
||||
name: `ptr<storage, buffer<64>, read_write>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferUnsized_ro_storage: {
|
||||
name: `ptr<storage, buffer>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferSized_ro_storage: {
|
||||
name: `ptr<storage, buffer<128>>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferSizedSmall_ro_storage: {
|
||||
name: `ptr<storage, buffer<64>>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferUnsized_uniform: {
|
||||
name: `ptr<uniform, buffer>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferSized_uniform: {
|
||||
name: `ptr<uniform, buffer<128>>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferSizedSmall_uniform: {
|
||||
name: `ptr<uniform, buffer<64>>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferUnsized_workgroup: {
|
||||
name: `ptr<workgroup, buffer>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferSized_workgroup: {
|
||||
name: `ptr<workgroup, buffer<128>>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferSizedSmall_workgroup: {
|
||||
name: `ptr<workgroup, buffer<64>>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferOverrideNoDefault_workgroup: {
|
||||
name: `ptr<workgroup, buffer<override_no_default>>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferOverrideDefault_workgroup: {
|
||||
name: `ptr<workgroup, buffer<override_default>>`,
|
||||
valid: 'with_buffer_view',
|
||||
},
|
||||
ptrBufferOverrideExpr_workgroup: {
|
||||
name: `ptr<workgroup, buffer<override_expr>>`,
|
||||
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<string, ParamValueCase> = {
|
||||
|
|
@ -519,6 +586,81 @@ const kFunctionParamValueCases: Record<string, ParamValueCase> = {
|
|||
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, read_write> storage_buffer_unsized : buffer;
|
||||
@group(2) @binding(1) var<storage, read_write> storage_buffer_sized : buffer<128>;
|
||||
@group(2) @binding(2) var<storage, read_write> storage_buffer_sized_small : buffer<64>;
|
||||
@group(2) @binding(3) var<storage> ro_storage_buffer_unsized : buffer;
|
||||
@group(2) @binding(4) var<storage> ro_storage_buffer_sized : buffer<128>;
|
||||
@group(2) @binding(5) var<storage> ro_storage_buffer_sized_small : buffer<64>;
|
||||
@group(2) @binding(6) var<uniform> uniform_buffer_sized : buffer<128>;
|
||||
@group(2) @binding(7) var<uniform> uniform_buffer_sized_small : buffer<64>;
|
||||
|
||||
var<workgroup> wg_buffer_sized : buffer<128>;
|
||||
var<workgroup> wg_buffer_sized_small : buffer<64>;
|
||||
var<workgroup> wg_buffer_override_no_default : buffer<override_no_default>;
|
||||
var<workgroup> wg_buffer_override_default : buffer<override_default>;
|
||||
var<workgroup> wg_buffer_override_expr : buffer<override_expr>;
|
||||
`;
|
||||
}
|
||||
|
||||
const code = `
|
||||
${enable}
|
||||
|
||||
|
|
@ -574,6 +736,8 @@ var<storage, read_write> rw_host_shareable : host_shareable;
|
|||
@group(1) @binding(2)
|
||||
var<uniform> uniform_host_shareable : host_shareable;
|
||||
|
||||
${buffer_decls}
|
||||
|
||||
fn bar(param : ${param.name}) { }
|
||||
|
||||
var<private> 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);
|
||||
});
|
||||
|
|
|
|||
|
|
@ -95,6 +95,8 @@ const kInvalidLocationTypes = new Set([
|
|||
'texture_depth_multisampled_2d',
|
||||
'sampler',
|
||||
'sampler_comparison',
|
||||
'buffer',
|
||||
'buffer<16>',
|
||||
]);
|
||||
|
||||
g.test('stage_inout')
|
||||
|
|
|
|||
188
src/webgpu/shader/validation/types/buffer.spec.ts
Normal file
188
src/webgpu/shader/validation/types/buffer.spec.ts
Normal file
|
|
@ -0,0 +1,188 @@
|
|||
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<x>;`,
|
||||
valid: true,
|
||||
},
|
||||
const_negative: {
|
||||
code: `const x = -1; alias T = buffer<x>;`,
|
||||
valid: false,
|
||||
},
|
||||
const_zero: {
|
||||
code: `const x = 0; alias T = buffer<x>;`,
|
||||
valid: false,
|
||||
},
|
||||
const_function: {
|
||||
code: `const x = 16; const y = 32; alias T = buffer<min(x, y)>;`,
|
||||
valid: true,
|
||||
},
|
||||
override: {
|
||||
code: `override x = 16; alias T = buffer<x>;`,
|
||||
valid: true,
|
||||
},
|
||||
override_negative: {
|
||||
code: `override x = -1; alias T = buffer<x>;`,
|
||||
valid: true,
|
||||
},
|
||||
override_zero: {
|
||||
code: `override x = 0; alias T = buffer<x>;`,
|
||||
valid: true,
|
||||
},
|
||||
override_function: {
|
||||
code: `override x = 16; override y = 32; alias T = buffer<max(x, y)>;`,
|
||||
valid: true,
|
||||
},
|
||||
empty_template: {
|
||||
code: `alias T = buffer<>;`,
|
||||
valid: false,
|
||||
},
|
||||
missing_rparen: {
|
||||
code: `alias T = buffer<16;`,
|
||||
valid: false,
|
||||
},
|
||||
subtype: {
|
||||
code: `alias T = buffer<u32>;`,
|
||||
valid: false,
|
||||
},
|
||||
pointer_workgroup_unsized_buffer: {
|
||||
code: `alias T = buffer; alias PT = ptr<workgroup, T>;`,
|
||||
valid: true,
|
||||
},
|
||||
pointer_uniform_unsized_buffer: {
|
||||
code: `alias T = buffer; alias PT = ptr<uniform, T>;`,
|
||||
valid: true,
|
||||
},
|
||||
pointer_storage_unsized_buffer: {
|
||||
code: `alias T = buffer; alias PT = ptr<storage, T>;`,
|
||||
valid: true,
|
||||
},
|
||||
pointer_workgroup_sized_buffer: {
|
||||
code: `alias T = buffer<32>; alias PT = ptr<workgroup, T>;`,
|
||||
valid: true,
|
||||
},
|
||||
pointer_uniform_sized_buffer: {
|
||||
code: `alias T = buffer<32>; alias PT = ptr<uniform, T>;`,
|
||||
valid: true,
|
||||
},
|
||||
pointer_storage_sized_buffer: {
|
||||
code: `alias T = buffer<32>; alias PT = ptr<storage, T>;`,
|
||||
valid: true,
|
||||
},
|
||||
pointer_workgroup_override_sized_buffer: {
|
||||
code: `override o : i32; alias T = buffer<o>; alias PT = ptr<workgroup, T>;`,
|
||||
valid: true,
|
||||
},
|
||||
pointer_uniform_override_sized_buffer: {
|
||||
code: `override o : i32; alias T = buffer<o>; alias PT = ptr<uniform, T>;`,
|
||||
valid: false,
|
||||
},
|
||||
pointer_storage_override_sized_buffer: {
|
||||
code: `override o : i32; alias T = buffer<o>; alias PT = ptr<storage, T>;`,
|
||||
valid: false,
|
||||
},
|
||||
buffer_3bytes: {
|
||||
code: `alias T = buffer<3>;`,
|
||||
valid: false,
|
||||
},
|
||||
buffer_2bytes_f16: {
|
||||
code: `enable f16;\nalias T = buffer<2>;`,
|
||||
valid: true,
|
||||
},
|
||||
buffer_2bytes_no_f16: {
|
||||
code: `alias T = buffer<2>;`,
|
||||
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 && !t.case.includes('pointer');
|
||||
})
|
||||
.beginSubcases()
|
||||
.combine('aspace', [
|
||||
'function',
|
||||
'private',
|
||||
'storage',
|
||||
'uniform',
|
||||
'workgroup',
|
||||
'immediate',
|
||||
] 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':
|
||||
case 'immediate':
|
||||
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' &&
|
||||
t.params.aspace !== 'immediate';
|
||||
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);
|
||||
});
|
||||
Loading…
Reference in a new issue