buffer_view CTS (#4652) * Validation of buffer type and related builtin functions * Execution tests for buffer_view

diff --git a/.gitignore b/.gitignore
index 09eb42e..1a359f8 100644
--- a/.gitignore
+++ b/.gitignore

@@ -11,6 +11,9 @@
 *.tmp.txt
 /docs/tsdoc/
 
+# vim swap files
+*.swp*
+
 # `pnpm` lockfile
 pnpm-lock.yaml
 

diff --git a/src/common/util/util.ts b/src/common/util/util.ts
index 540a29f..4e19f9e 100644
--- a/src/common/util/util.ts
+++ b/src/common/util/util.ts

@@ -373,6 +373,7 @@
   Uint8ClampedArray: Uint8ClampedArray;
   Int32Array: Int32Array;
   Uint32Array: Uint32Array;
+  Float16Array: Float16Array;
   Float32Array: Float32Array;
   Float64Array: Float64Array;
   BigInt64Array: BigInt64Array;

diff --git a/src/resources/cache/hashes.json b/src/resources/cache/hashes.json
index 98e9043..c0bfe20 100644
--- a/src/resources/cache/hashes.json
+++ b/src/resources/cache/hashes.json

@@ -1,112 +1,112 @@
 {
-  "webgpu/shader/execution/binary/af_addition.bin": "6e41f81e",
-  "webgpu/shader/execution/binary/af_logical.bin": "254d66ec",
-  "webgpu/shader/execution/binary/af_division.bin": "e15d545b",
-  "webgpu/shader/execution/binary/af_matrix_addition.bin": "4bfc497f",
-  "webgpu/shader/execution/binary/af_matrix_subtraction.bin": "2bbdd5e6",
-  "webgpu/shader/execution/binary/af_multiplication.bin": "23c7c8ad",
-  "webgpu/shader/execution/binary/af_remainder.bin": "9b84ecb1",
-  "webgpu/shader/execution/binary/af_subtraction.bin": "bd4b2ec9",
-  "webgpu/shader/execution/binary/f16_addition.bin": "5fe1fb4e",
-  "webgpu/shader/execution/binary/f16_logical.bin": "225d12ee",
-  "webgpu/shader/execution/binary/f16_division.bin": "dffabf5f",
-  "webgpu/shader/execution/binary/f16_matrix_addition.bin": "acc779ee",
-  "webgpu/shader/execution/binary/f16_matrix_matrix_multiplication.bin": "dcfbfff9",
-  "webgpu/shader/execution/binary/f16_matrix_scalar_multiplication.bin": "4b73a544",
-  "webgpu/shader/execution/binary/f16_matrix_subtraction.bin": "bfe9a25d",
-  "webgpu/shader/execution/binary/f16_matrix_vector_multiplication.bin": "a2222aec",
-  "webgpu/shader/execution/binary/f16_multiplication.bin": "5d5265c9",
-  "webgpu/shader/execution/binary/f16_remainder.bin": "5284a25",
-  "webgpu/shader/execution/binary/f16_subtraction.bin": "b98fc99f",
-  "webgpu/shader/execution/binary/f32_addition.bin": "13d7f835",
-  "webgpu/shader/execution/binary/f32_logical.bin": "c1f1b636",
-  "webgpu/shader/execution/binary/f32_division.bin": "80fee721",
-  "webgpu/shader/execution/binary/f32_matrix_addition.bin": "a0d20c36",
-  "webgpu/shader/execution/binary/f32_matrix_matrix_multiplication.bin": "3cabc829",
-  "webgpu/shader/execution/binary/f32_matrix_scalar_multiplication.bin": "aed8659",
-  "webgpu/shader/execution/binary/f32_matrix_subtraction.bin": "6063545d",
-  "webgpu/shader/execution/binary/f32_matrix_vector_multiplication.bin": "e36324e4",
-  "webgpu/shader/execution/binary/f32_multiplication.bin": "40527d93",
-  "webgpu/shader/execution/binary/f32_remainder.bin": "d8738967",
-  "webgpu/shader/execution/binary/f32_subtraction.bin": "f5453a8",
-  "webgpu/shader/execution/binary/i32_arithmetic.bin": "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"
 }
\ No newline at end of file

diff --git a/src/webgpu/api/validation/dispatch.spec.ts b/src/webgpu/api/validation/dispatch.spec.ts
index 3b06a22..22059201 100644
--- a/src/webgpu/api/validation/dispatch.spec.ts
+++ b/src/webgpu/api/validation/dispatch.spec.ts

@@ -4,6 +4,8 @@
 
 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 @@
     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);
+    }
+  });

diff --git a/src/webgpu/api/validation/min_buffer_binding_size.spec.ts b/src/webgpu/api/validation/min_buffer_binding_size.spec.ts
new file mode 100644
index 0000000..50b19c2
--- /dev/null
+++ b/src/webgpu/api/validation/min_buffer_binding_size.spec.ts

@@ -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);
+    }
+  });

diff --git a/src/webgpu/capability_info.ts b/src/webgpu/capability_info.ts
index 98ec040..871dcad 100644
--- a/src/webgpu/capability_info.ts
+++ b/src/webgpu/capability_info.ts

@@ -991,6 +991,7 @@
   'linear_indexing',
   'texture_formats_tier1',
   'immediate_address_space',
+  'buffer_view',
 ] as const;
 
 export type WGSLLanguageFeature = (typeof kKnownWGSLLanguageFeatures)[number];

diff --git a/src/webgpu/shader/execution/expression/call/builtin/bufferArrayView.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/bufferArrayView.spec.ts
new file mode 100644
index 0000000..a8ec17d
--- /dev/null
+++ b/src/webgpu/shader/execution/expression/call/builtin/bufferArrayView.spec.ts

@@ -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);
+  });

diff --git a/src/webgpu/shader/execution/expression/call/builtin/bufferLength.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/bufferLength.spec.ts
new file mode 100644
index 0000000..57754cd
--- /dev/null
+++ b/src/webgpu/shader/execution/expression/call/builtin/bufferLength.spec.ts

@@ -0,0 +1,485 @@
+export const description = `
+Execution tests for bufferLength
+
+Tests sized and unsized buffers across all address spaces (where applicable).
+Tested against buffers with static and dynamic offsets.
+`;
+
+import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
+import { GPUTest } from '../../../../../gpu_test.js';
+
+export const g = makeTestGroup(GPUTest);
+
+g.test('sized_buffer')
+  .desc('Tests bufferLength directly on sized buffer variables')
+  .params(u =>
+    u
+      .combine('size', [256, 512, 600, 1024, 2048] as const)
+      .beginSubcases()
+      .combine('param', ['none', 'unsized', 'sized'] as const)
+      .combine('padding', [false, true] as const)
+      .combine('dynamic_offset', [false, true] as const)
+      .combine('offset', [0, 256] as const)
+  )
+  .fn(t => {
+    t.skipIfLanguageFeatureNotSupported('buffer_view');
+    if (t.params.param !== 'none') {
+      t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters');
+    }
+
+    let fn_decls = '';
+    const assigns = ['', '', '', ''];
+    switch (t.params.param) {
+      case 'none':
+        assigns[0] = 'bufferLength(&uniform_buffer)';
+        assigns[1] = 'bufferLength(&ro_storage_buffer)';
+        assigns[2] = 'bufferLength(&storage_buffer)';
+        assigns[3] = 'bufferLength(&workgroup_buffer)';
+        break;
+      case 'unsized':
+        fn_decls = `
+fn uniformUnsized(p : ptr<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);
+  });

diff --git a/src/webgpu/shader/execution/expression/call/builtin/bufferView.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/bufferView.spec.ts
new file mode 100644
index 0000000..6c6442d
--- /dev/null
+++ b/src/webgpu/shader/execution/expression/call/builtin/bufferView.spec.ts

@@ -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);
+  });

diff --git a/src/webgpu/shader/execution/expression/call/builtin/buffer_view_utils.ts b/src/webgpu/shader/execution/expression/call/builtin/buffer_view_utils.ts
new file mode 100644
index 0000000..8622b98
--- /dev/null
+++ b/src/webgpu/shader/execution/expression/call/builtin/buffer_view_utils.ts

@@ -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);
+}

diff --git a/src/webgpu/shader/validation/decl/let.spec.ts b/src/webgpu/shader/validation/decl/let.spec.ts
index 3c1ec06..58a86ef 100644
--- a/src/webgpu/shader/validation/decl/let.spec.ts
+++ b/src/webgpu/shader/validation/decl/let.spec.ts

@@ -10,7 +10,7 @@
 
 interface Case {
   code: string;
-  valid: boolean | 'texture_and_sampler_let';
+  valid: boolean | 'texture_and_sampler_let' | 'buffer_view';
   decls?: string;
 }
 
@@ -109,6 +109,36 @@
     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 @@
     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 @@
     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')

diff --git a/src/webgpu/shader/validation/decl/var.spec.ts b/src/webgpu/shader/validation/decl/var.spec.ts
index 5d5c091..7f23f3d 100644
--- a/src/webgpu/shader/validation/decl/var.spec.ts
+++ b/src/webgpu/shader/validation/decl/var.spec.ts

@@ -29,30 +29,35 @@
     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 @@
     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 @@
     isConstructible: true,
     isFixedFootprint: true,
     requiresF16: false,
+    requiresBufferView: false,
   },
   mat3x4h: {
     isHostShareable: true,
     isConstructible: true,
     isFixedFootprint: true,
     requiresF16: true,
+    requiresBufferView: false,
   },
 
   // Atomics.
@@ -107,12 +119,14 @@
     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 @@
     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 @@
     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 @@
     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 @@
         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 @@
     ${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 @@
       ${decl} ${t.params.via_alias ? 'MyType' : t.params.type};
     }`;
 
+    shouldPass &&= !kTypes[t.params.type].requiresBufferView || t.hasLanguageFeature('buffer_view');
     t.expectCompileResult(shouldPass, wgsl);
   });
 

diff --git a/src/webgpu/shader/validation/expression/call/builtin/arrayLength.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/arrayLength.spec.ts
index 59b26b5..879a08b 100644
--- a/src/webgpu/shader/validation/expression/call/builtin/arrayLength.spec.ts
+++ b/src/webgpu/shader/validation/expression/call/builtin/arrayLength.spec.ts

@@ -25,6 +25,20 @@
     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}>`))

diff --git a/src/webgpu/shader/validation/expression/call/builtin/bufferArrayView.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/bufferArrayView.spec.ts
new file mode 100644
index 0000000..aa6ee20
--- /dev/null
+++ b/src/webgpu/shader/validation/expression/call/builtin/bufferArrayView.spec.ts

@@ -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);
+  });

diff --git a/src/webgpu/shader/validation/expression/call/builtin/bufferLength.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/bufferLength.spec.ts
new file mode 100644
index 0000000..be3a3c6
--- /dev/null
+++ b/src/webgpu/shader/validation/expression/call/builtin/bufferLength.spec.ts

@@ -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);
+  });

diff --git a/src/webgpu/shader/validation/expression/call/builtin/bufferView.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/bufferView.spec.ts
new file mode 100644
index 0000000..50dbf12
--- /dev/null
+++ b/src/webgpu/shader/validation/expression/call/builtin/bufferView.spec.ts

@@ -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);
+  });

diff --git a/src/webgpu/shader/validation/functions/restrictions.spec.ts b/src/webgpu/shader/validation/functions/restrictions.spec.ts
index 73c3a37..5cec4c3 100644
--- a/src/webgpu/shader/validation/functions/restrictions.spec.ts
+++ b/src/webgpu/shader/validation/functions/restrictions.spec.ts

@@ -202,7 +202,7 @@
 
 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 @@
   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 @@
     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 @@
   value: string;
   matches: string[];
   needsUnrestrictedPointerParameters?: boolean;
+  needsBufferView?: boolean;
 }
 
 const kFunctionParamValueCases: Record<string, ParamValueCase> = {
@@ -519,6 +586,81 @@
     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 @@
     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 @@
 @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 @@
       (kFunctionParamTypeCases[t.params.decl].valid === 'with_unrestricted_pointer_parameters' ||
         arg.needsUnrestrictedPointerParameters) ??
       false;
+    const needsBufferView =
+      (kFunctionParamTypeCases[t.params.decl].valid === 'with_buffer_view' ||
+        arg.needsBufferView) ??
+      false;
 
     let isValid = parameterMatches(t.params.decl, arg.matches);
     if (isValid && needsUnrestrictedPointerParameters) {
       isValid = t.hasLanguageFeature('unrestricted_pointer_parameters');
     }
+    if (isValid && needsBufferView) {
+      isValid =
+        t.hasLanguageFeature('unrestricted_pointer_parameters') &&
+        t.hasLanguageFeature('buffer_view');
+    }
 
     t.expectCompileResult(isValid, code);
   });

diff --git a/src/webgpu/shader/validation/shader_io/locations.spec.ts b/src/webgpu/shader/validation/shader_io/locations.spec.ts
index 5641b78..c569e28 100644
--- a/src/webgpu/shader/validation/shader_io/locations.spec.ts
+++ b/src/webgpu/shader/validation/shader_io/locations.spec.ts

@@ -95,6 +95,8 @@
   'texture_depth_multisampled_2d',
   'sampler',
   'sampler_comparison',
+  'buffer',
+  'buffer<16>',
 ]);
 
 g.test('stage_inout')

diff --git a/src/webgpu/shader/validation/types/buffer.spec.ts b/src/webgpu/shader/validation/types/buffer.spec.ts
new file mode 100644
index 0000000..ca4be62
--- /dev/null
+++ b/src/webgpu/shader/validation/types/buffer.spec.ts

@@ -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);
+  });