From 7a270970dfab89d312af5921333a3e7805ea0169 Mon Sep 17 00:00:00 2001 From: shaoboyan Date: Fri, 5 Jun 2026 14:09:14 +0800 Subject: [PATCH 1/2] Add shared support for WGSL immediate validation This patch adds the shared CTS plumbing needed by the follow-up WGSL `immediate` address space validation tests. It adds `immediate_address_space` to the known WGSL language feature list, models `immediate` in the address space metadata, and treats it as a module-scope, unbound, read-only address space with host-shareable store types. The validation helpers emit the required WGSL `requires` directive and skip immediate tests when immediate data is not exposed by the browser. The skip checks `setImmediates`/`maxImmediateSize` because current Chrome Canary exposes the API before advertising the WGSL feature. This also removes two now-unnecessary address-space type assertions that became redundant once the shared metadata was typed by `AddressSpace`. --- src/webgpu/capability_info.ts | 1 + src/webgpu/shader/types.ts | 39 +++++++++++++++--- src/webgpu/shader/validation/decl/util.ts | 41 ++++++++++++++++++- src/webgpu/shader/validation/decl/var.spec.ts | 6 +-- .../shader/validation/types/pointer.spec.ts | 6 +-- 5 files changed, 77 insertions(+), 16 deletions(-) diff --git a/src/webgpu/capability_info.ts b/src/webgpu/capability_info.ts index 3d7160bb894e..557a0267c419 100644 --- a/src/webgpu/capability_info.ts +++ b/src/webgpu/capability_info.ts @@ -987,6 +987,7 @@ export const kKnownWGSLLanguageFeatures = [ 'swizzle_assignment', 'linear_indexing', 'texture_formats_tier1', + 'immediate_address_space', ] as const; export type WGSLLanguageFeature = (typeof kKnownWGSLLanguageFeatures)[number]; diff --git a/src/webgpu/shader/types.ts b/src/webgpu/shader/types.ts index c7648efdc880..f756703b8cdb 100644 --- a/src/webgpu/shader/types.ts +++ b/src/webgpu/shader/types.ts @@ -1,5 +1,6 @@ import { keysOf } from '../../common/util/data_tables.js'; import { assert } from '../../common/util/util.js'; +import type { WGSLLanguageFeature } from '../capability_info.js'; import { align } from '../util/math.js'; const kDefaultArrayLength = 3; @@ -102,7 +103,14 @@ export const kMatrixContainerTypeLayoutInfo = } } as const; -export type AddressSpace = 'storage' | 'uniform' | 'private' | 'function' | 'workgroup' | 'handle'; +export type AddressSpace = + | 'storage' + | 'uniform' + | 'private' + | 'function' + | 'workgroup' + | 'immediate' + | 'handle'; export type AccessMode = 'read' | 'write' | 'read_write'; export type Scope = 'module' | 'function'; @@ -133,9 +141,12 @@ export type AddressSpaceInfo = { // in the storage address space, must not be specified in the WGSL source. // See ยง13.3 Address Spaces. spellAccessMode: Requirement; + + // WGSL language feature required to use this address space, if any. + wgslLanguageFeature?: WGSLLanguageFeature; }; -export const kAddressSpaceInfo: Record = { +export const kAddressSpaceInfo: Record = { storage: { scope: 'module', binding: true, @@ -171,6 +182,14 @@ export const kAddressSpaceInfo: Record = { accessModes: ['read_write'], spellAccessMode: 'never', }, + immediate: { + scope: 'module', + binding: false, + spell: 'must', + accessModes: ['read'], + spellAccessMode: 'never', + wgslLanguageFeature: 'immediate_address_space', + }, handle: { scope: 'module', binding: true, @@ -237,8 +256,10 @@ export function* generateTypes({ } const scalarType = isAtomic ? `atomic<${baseType}>` : baseType; - // Storage and uniform require host-sharable types. - if (addressSpace === 'storage' || addressSpace === 'uniform') { + // Storage, uniform, and immediate require host-shareable types. + const requiresHostShareable = + addressSpace === 'storage' || addressSpace === 'uniform' || addressSpace === 'immediate'; + if (requiresHostShareable) { assert(isHostSharable(baseType), 'type ' + baseType.toString() + ' is not host sharable'); } @@ -289,6 +310,9 @@ export function* generateTypes({ // Array types if (containerType === 'array') { + if (addressSpace === 'immediate') { + return; + } let arrayElemType: string = scalarType; let arrayElementCount: number = kDefaultArrayLength; let supportsAtomics = scalarInfo.supportsAtomics; @@ -382,8 +406,11 @@ export function* supportedScalarTypes(p: { isAtomic: boolean; addressSpace: stri // Test atomics only on supported scalar types. if (p.isAtomic && !info.supportsAtomics) continue; - // Storage and uniform require host-sharable types. - const isHostShared = p.addressSpace === 'storage' || p.addressSpace === 'uniform'; + // Storage, uniform, and immediate require host-shareable types. + const isHostShared = + p.addressSpace === 'storage' || + p.addressSpace === 'uniform' || + p.addressSpace === 'immediate'; if (isHostShared && info.layout === undefined) continue; yield scalarType; diff --git a/src/webgpu/shader/validation/decl/util.ts b/src/webgpu/shader/validation/decl/util.ts index ab1b08e12a93..a2f175a5298c 100644 --- a/src/webgpu/shader/validation/decl/util.ts +++ b/src/webgpu/shader/validation/decl/util.ts @@ -1,3 +1,5 @@ +import { getGPU } from '../../../../common/util/navigator_gpu.js'; +import { supportsImmediateData } from '../../../../common/util/util.js'; import { AccessMode, AddressSpace, @@ -12,6 +14,40 @@ export type ShaderStage = 'vertex' | 'fragment' | 'compute'; /** The list of all shader stages */ export const kShaderStages = ['vertex', 'fragment', 'compute'] as const; +export function requiredLanguageFeatureHeader(addressSpace: AddressSpace): string { + const feature = kAddressSpaceInfo[addressSpace].wgslLanguageFeature; + return feature === undefined ? '' : `requires ${feature};\n`; +} + +type AddressSpaceSupportTest = { + readonly rec: Parameters[0]; + skip(message: string): never; + skipIfLanguageFeatureNotSupported( + langFeature: NonNullable + ): void; +}; + +export function skipIfImmediateDataNotSupported(t: AddressSpaceSupportTest): void { + if (!supportsImmediateData(getGPU(t.rec))) { + t.skip('Immediate data not supported'); + } +} + +export function skipIfAddressSpaceNotSupported( + t: AddressSpaceSupportTest, + addressSpace: AddressSpace +): void { + if (addressSpace === 'immediate') { + skipIfImmediateDataNotSupported(t); + return; + } + + const feature = kAddressSpaceInfo[addressSpace].wgslLanguageFeature; + if (feature !== undefined) { + t.skipIfLanguageFeatureNotSupported(feature); + } +} + /** * declareEntrypoint emits the WGSL to declare an entry point with the name, stage and body. * The generated function will have an appropriate return type and return statement, so that `body` @@ -107,15 +143,16 @@ export function getVarDeclShader( p.explicitSpace ? p.addressSpace : '', p.explicitAccess ? p.accessMode : '' ); + const header = requiredLanguageFeatureHeader(p.addressSpace); additionalBody = additionalBody ?? ''; switch (info.scope) { case 'module': - return decl + '\n' + declareEntryPoint({ stage: p.stage, body: additionalBody }); + return header + decl + '\n' + declareEntryPoint({ stage: p.stage, body: additionalBody }); case 'function': - return declareEntryPoint({ stage: p.stage, body: decl + '\n' + additionalBody }); + return header + declareEntryPoint({ stage: p.stage, body: decl + '\n' + additionalBody }); } } diff --git a/src/webgpu/shader/validation/decl/var.spec.ts b/src/webgpu/shader/validation/decl/var.spec.ts index a4afc6ab84e5..ebb4e75b471e 100644 --- a/src/webgpu/shader/validation/decl/var.spec.ts +++ b/src/webgpu/shader/validation/decl/var.spec.ts @@ -4,7 +4,7 @@ Validation tests for host-shareable types. import { makeTestGroup } from '../../../../common/framework/test_group.js'; import { keysOf } from '../../../../common/util/data_tables.js'; -import { AddressSpace, kAccessModeInfo, kAddressSpaceInfo } from '../../types.js'; +import { kAccessModeInfo, kAddressSpaceInfo } from '../../types.js'; import { ShaderValidationTest } from '../shader_validation_test.js'; import { @@ -582,9 +582,7 @@ g.test('address_space_access_mode') }); // Address spaces that can hold an i32 variable. -const kNonHandleAddressSpaces = keysOf(kAddressSpaceInfo).filter( - as => as !== 'handle' -) as AddressSpace[]; +const kNonHandleAddressSpaces = keysOf(kAddressSpaceInfo).filter(as => as !== 'handle'); g.test('explicit_access_mode') .desc('Validate uses of an explicit access mode on a var declaration') diff --git a/src/webgpu/shader/validation/types/pointer.spec.ts b/src/webgpu/shader/validation/types/pointer.spec.ts index b42ad4f0d0d2..c787668ad5e4 100644 --- a/src/webgpu/shader/validation/types/pointer.spec.ts +++ b/src/webgpu/shader/validation/types/pointer.spec.ts @@ -2,7 +2,7 @@ export const description = 'Test pointer type validation'; import { makeTestGroup } from '../../../../common/framework/test_group.js'; import { keysOf } from '../../../../common/util/data_tables.js'; -import { AddressSpace, kAccessModeInfo, kAddressSpaceInfo } from '../../types.js'; +import { kAccessModeInfo, kAddressSpaceInfo } from '../../types.js'; import { pointerType, explicitSpaceExpander, @@ -140,9 +140,7 @@ g.test('type') }); // Address spaces that can hold an i32 variable. -const kNonHandleAddressSpaces = keysOf(kAddressSpaceInfo).filter( - as => as !== 'handle' -) as AddressSpace[]; +const kNonHandleAddressSpaces = keysOf(kAddressSpaceInfo).filter(as => as !== 'handle'); g.test('let_ptr_explicit_type_matches_var') .desc( From 91f8171d6368679624266846c747f8ba685b9063 Mon Sep 17 00:00:00 2001 From: shaoboyan Date: Fri, 5 Jun 2026 14:09:21 +0800 Subject: [PATCH 2/2] Add declaration tests for WGSL immediate variables This patch adds a focused declaration suite for `var`. The new tests cover accepted scalar, vector, matrix, and struct store types, rejected store types, module-scope-only declarations, rejection of binding attributes, rejection of explicit access modes, and the one-immediate-variable-per-entry-point rule. It also threads `immediate` through the existing declaration tests so the generic address space coverage checks the same declaration access rules as storage, uniform, private, workgroup, and function variables. The context-dependent resolution table now recognizes `immediate_address_space` as a language feature name. --- .../decl/context_dependent_resolution.spec.ts | 1 + .../shader/validation/decl/immediate.spec.ts | 301 ++++++++++++++++++ src/webgpu/shader/validation/decl/var.spec.ts | 73 ++++- 3 files changed, 370 insertions(+), 5 deletions(-) create mode 100644 src/webgpu/shader/validation/decl/immediate.spec.ts diff --git a/src/webgpu/shader/validation/decl/context_dependent_resolution.spec.ts b/src/webgpu/shader/validation/decl/context_dependent_resolution.spec.ts index 403f4e72b0f4..51a672dc073b 100644 --- a/src/webgpu/shader/validation/decl/context_dependent_resolution.spec.ts +++ b/src/webgpu/shader/validation/decl/context_dependent_resolution.spec.ts @@ -230,6 +230,7 @@ const kLanguageCases = { packed_4x8_integer_dot_product: `requires packed_4x8_integer_dot_product;`, unrestricted_pointer_parameters: `requires unrestricted_pointer_parameters;`, pointer_composite_access: `requires pointer_composite_access;`, + immediate_address_space: `requires immediate_address_space;`, }; g.test('language_names') diff --git a/src/webgpu/shader/validation/decl/immediate.spec.ts b/src/webgpu/shader/validation/decl/immediate.spec.ts new file mode 100644 index 000000000000..63718dd03b04 --- /dev/null +++ b/src/webgpu/shader/validation/decl/immediate.spec.ts @@ -0,0 +1,301 @@ +export const description = ` +Validation tests for the WGSL immediate address space. +`; + +import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { keysOf } from '../../../../common/util/data_tables.js'; +import { ShaderValidationTest } from '../shader_validation_test.js'; + +import { skipIfImmediateDataNotSupported } from './util.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +const kImmediateFeature = 'immediate_address_space' as const; +const kImmediateHeader = `requires ${kImmediateFeature};`; + +const kValidStoreTypes = { + u32: `u32`, + i32: `i32`, + f32: `f32`, + vec2u: `vec2u`, + vec3i: `vec3i`, + vec4f: `vec4f`, + mat2x2f: `mat2x2f`, + struct_numeric: `S`, +} as const; + +const kInvalidStoreTypes = { + bool: { enable: ``, prelude: ``, type: `bool` }, + vec2_bool: { enable: ``, prelude: ``, type: `vec2` }, + atomic_u32: { enable: ``, prelude: ``, type: `atomic` }, + ptr_function_u32: { enable: ``, prelude: ``, type: `ptr` }, + sampler: { enable: ``, prelude: ``, type: `sampler` }, + sampler_comparison: { enable: ``, prelude: ``, type: `sampler_comparison` }, + texture_2d: { enable: ``, prelude: ``, type: `texture_2d` }, + runtime_array: { enable: ``, prelude: ``, type: `array` }, + fixed_array: { enable: ``, prelude: ``, type: `array` }, + struct_runtime_array: { enable: ``, prelude: `struct S { data : array }`, type: `S` }, + struct_fixed_array: { + enable: ``, + prelude: `struct S { data : array }`, + type: `S`, + }, + f16: { enable: `enable f16;`, prelude: ``, type: `f16` }, + vec3h: { enable: `enable f16;`, prelude: ``, type: `vec3h` }, +} as const; + +g.test('store_type,valid') + .desc('Validates immediate store types supported by the current WGSL immediate implementation.') + .params(u => u.combine('type', keysOf(kValidStoreTypes))) + .fn(t => { + skipIfImmediateDataNotSupported(t); + const prelude = t.params.type === 'struct_numeric' ? 'struct S { a : u32, b : vec4f }' : ''; + const wgsl = ` +${kImmediateHeader} +${prelude} +var data : ${kValidStoreTypes[t.params.type]}; +@compute @workgroup_size(1) +fn main() { + _ = data; +}`; + t.expectCompileResult(true, wgsl); + }); + +g.test('store_type,invalid') + .desc('Validates types that cannot be used for immediate variables.') + .params(u => u.combine('type', keysOf(kInvalidStoreTypes))) + .fn(t => { + skipIfImmediateDataNotSupported(t); + const testcase = kInvalidStoreTypes[t.params.type]; + if (testcase.enable.includes('f16')) { + t.skipIfDeviceDoesNotHaveFeature('shader-f16'); + } + const wgsl = ` +${kImmediateHeader} +${testcase.enable} +${testcase.prelude} +var data : ${testcase.type}; +@compute @workgroup_size(1) +fn main() { + _ = data; +}`; + t.expectCompileResult(false, wgsl); + }); + +g.test('scope') + .desc('Validates that immediate variables are module-scope only.') + .params(u => u.combine('scope', ['module', 'function'] as const)) + .fn(t => { + skipIfImmediateDataNotSupported(t); + const wgsl = + t.params.scope === 'module' + ? ` +${kImmediateHeader} +var data : u32; +@compute @workgroup_size(1) +fn main() { + _ = data; +}` + : ` +${kImmediateHeader} +@compute @workgroup_size(1) +fn main() { + var data : u32; + _ = data; +}`; + t.expectCompileResult(t.params.scope === 'module', wgsl); + }); + +g.test('binding_attributes') + .desc('Validates that @group and @binding are not allowed on immediate variables.') + .params(u => + u.combine('group', ['', '@group(0)'] as const).combine('binding', ['', '@binding(0)'] as const) + ) + .fn(t => { + skipIfImmediateDataNotSupported(t); + const wgsl = ` +${kImmediateHeader} +${t.params.group} ${t.params.binding} var data : u32; +@compute @workgroup_size(1) +fn main() { + _ = data; +}`; + t.expectCompileResult(t.params.group === '' && t.params.binding === '', wgsl); + }); + +g.test('access_mode') + .desc('Validates that immediate variables cannot spell an access mode.') + .params(u => u.combine('accessMode', ['', 'read', 'write', 'read_write'] as const)) + .fn(t => { + skipIfImmediateDataNotSupported(t); + const suffix = t.params.accessMode === '' ? '' : `, ${t.params.accessMode}`; + const wgsl = ` +${kImmediateHeader} +var data : u32; +@compute @workgroup_size(1) +fn main() { + _ = data; +}`; + t.expectCompileResult(t.params.accessMode === '', wgsl); + }); + +const kEntryPointCases = { + one_used: { + valid: true, + body: ` +var a : u32; +@compute @workgroup_size(1) +fn main() { + _ = a; +}`, + }, + two_declared_one_used: { + valid: true, + body: ` +var a : u32; +var b : u32; +@compute @workgroup_size(1) +fn main() { + _ = a; +}`, + }, + two_entry_points_one_each: { + valid: true, + body: ` +var a : u32; +var b : u32; +@compute @workgroup_size(1) +fn main_a() { + _ = a; +} +@compute @workgroup_size(1) +fn main_b() { + _ = b; +}`, + }, + one_entry_point_uses_two_directly: { + valid: false, + body: ` +var a : u32; +var b : u32; +@compute @workgroup_size(1) +fn main() { + _ = a + b; +}`, + }, + one_entry_point_uses_two_through_helper: { + valid: false, + body: ` +var a : u32; +var b : u32; +fn read_b() -> u32 { + return b; +} +@compute @workgroup_size(1) +fn main() { + _ = a + read_b(); +}`, + }, +} as const; + +g.test('entry_point_interface') + .desc('Validates one statically used immediate variable per entry point.') + .params(u => u.combine('case', keysOf(kEntryPointCases))) + .fn(t => { + skipIfImmediateDataNotSupported(t); + const testcase = kEntryPointCases[t.params.case]; + t.expectCompileResult(testcase.valid, `${kImmediateHeader}\n${testcase.body}`); + }); + +const kPointerCases = { + alias_module_scope: { + valid: true, + needsUnrestrictedPointerParameters: false, + body: ` +alias P = ptr; +var data : u32; +@compute @workgroup_size(1) +fn main() { + let p : P = &data; + _ = *p; +}`, + }, + let_inside_function: { + valid: true, + needsUnrestrictedPointerParameters: false, + body: ` +var data : u32; +@compute @workgroup_size(1) +fn main() { + let p : ptr = &data; + _ = *p; +}`, + }, + write_through_pointer: { + valid: false, + needsUnrestrictedPointerParameters: false, + body: ` +var data : u32; +@compute @workgroup_size(1) +fn main() { + let p : ptr = &data; + *p = 1u; +}`, + }, + pointer_parameter: { + valid: true, + needsUnrestrictedPointerParameters: true, + body: ` +var data : u32; +fn read_data(p : ptr) -> u32 { + return *p; +} +@compute @workgroup_size(1) +fn main() { + _ = read_data(&data); +}`, + }, + explicit_read_access: { + valid: false, + needsUnrestrictedPointerParameters: false, + body: ` +alias P = ptr;`, + }, + explicit_write_access: { + valid: false, + needsUnrestrictedPointerParameters: false, + body: ` +alias P = ptr;`, + }, + explicit_read_write_access: { + valid: false, + needsUnrestrictedPointerParameters: false, + body: ` +alias P = ptr;`, + }, + missing_store_type: { + valid: false, + needsUnrestrictedPointerParameters: false, + body: ` +alias P = ptr;`, + }, +} as const; + +g.test('pointers') + .desc('Validates ptr type creation, use, access modes, and function parameters.') + .params(u => u.combine('case', keysOf(kPointerCases))) + .fn(t => { + skipIfImmediateDataNotSupported(t); + const testcase = kPointerCases[t.params.case]; + const unrestrictedHeader = + testcase.needsUnrestrictedPointerParameters && + t.hasLanguageFeature('unrestricted_pointer_parameters') + ? 'requires unrestricted_pointer_parameters;\n' + : ''; + const expected = + testcase.valid && + (!testcase.needsUnrestrictedPointerParameters || + t.hasLanguageFeature('unrestricted_pointer_parameters')); + const wgsl = `${kImmediateHeader}\n${unrestrictedHeader}${testcase.body}`; + t.expectCompileResult(expected, wgsl); + }); diff --git a/src/webgpu/shader/validation/decl/var.spec.ts b/src/webgpu/shader/validation/decl/var.spec.ts index ebb4e75b471e..6491fcd14c64 100644 --- a/src/webgpu/shader/validation/decl/var.spec.ts +++ b/src/webgpu/shader/validation/decl/var.spec.ts @@ -14,6 +14,9 @@ import { supportsRead, supportsWrite, ShaderStage, + requiredLanguageFeatureHeader, + skipIfAddressSpaceNotSupported, + skipIfImmediateDataNotSupported, } from './util.js'; export const g = makeTestGroup(ShaderValidationTest); @@ -209,6 +212,31 @@ const kTypes = { }, }; +const kImmediateTypesWithArray = new Set([ + 'array>', + 'array, 4>', + 'array', + 'array', + 'array', + 'array', + 'S_array_vec4u', + 'S_array_vec4u_4', + 'S_array_bool_4', +] as const); + +const kImmediateUnsupportedF16Types = new Set(['f16', 'vec3h', 'mat3x4h'] as const); + +function isImmediateStoreType(typeName: keyof typeof kTypes): boolean { + const type = kTypes[typeName]; + return ( + type.isHostShareable && + type.isConstructible && + type.isFixedFootprint && + !kImmediateTypesWithArray.has(typeName) && + !kImmediateUnsupportedF16Types.has(typeName) + ); +} + g.test('module_scope_types') .desc('Test that only types that are allowed for a given address space are accepted.') .params(u => @@ -222,10 +250,14 @@ g.test('module_scope_types') 'storage_rw', 'uniform', 'workgroup', + 'immediate', ]) .combine('via_alias', [false, true]) ) .fn(t => { + if (t.params.kind === 'immediate') { + skipIfImmediateDataNotSupported(t); + } if (kTypes[t.params.type].requiresF16) { t.skipIfDeviceDoesNotHaveFeature('shader-f16'); } @@ -265,9 +297,16 @@ g.test('module_scope_types') decl = 'var foo : '; shouldPass = type.isFixedFootprint; break; + case 'immediate': + decl = 'var foo : '; + shouldPass = isImmediateStoreType(t.params.type); + break; } - const wgsl = `${type.requiresF16 ? 'enable f16;' : ''} + const featureHeader = + t.params.kind === 'immediate' ? requiredLanguageFeatureHeader('immediate') : ''; + + const wgsl = `${featureHeader}${type.requiresF16 ? 'enable f16;' : ''} const array_size_const = 4; override array_size_override = 4; @@ -461,13 +500,18 @@ g.test('binding_point_on_non_resources') .desc('Test that non-resource variables cannot have either @group or @binding attributes.') .params(u => u - .combine('addrspace', ['private', 'workgroup']) + .combine('addrspace', ['private', 'workgroup', 'immediate'] as const) .combine('group', ['', '@group(0)']) .combine('binding', ['', '@binding(0)']) ) .fn(t => { + if (t.params.addrspace === 'immediate') { + skipIfImmediateDataNotSupported(t); + } const shouldPass = t.params.group === '' && t.params.binding === ''; - const wgsl = `${t.params.group} ${t.params.binding} var<${t.params.addrspace}> foo : i32;`; + const header = + t.params.addrspace === 'immediate' ? requiredLanguageFeatureHeader('immediate') : ''; + const wgsl = `${header}${t.params.group} ${t.params.binding} var<${t.params.addrspace}> foo : i32;`; t.expectCompileResult(shouldPass, wgsl); }); @@ -538,13 +582,24 @@ g.test('address_space_access_mode') .desc('Test that only storage accepts an access mode') .params(u => u - .combine('address_space', ['private', 'storage', 'uniform', 'function', 'workgroup'] as const) + .combine('address_space', [ + 'private', + 'storage', + 'uniform', + 'function', + 'workgroup', + 'immediate', + ] as const) .combine('access_mode', ['', 'read', 'write', 'read_write'] as const) .combine('trailing_comma', [true, false] as const) ) .fn(t => { + if (t.params.address_space === 'immediate') { + skipIfImmediateDataNotSupported(t); + } let fdecl = ``; let mdecl = ``; + let header = ``; // Most address spaces do not accept an access mode, but should accept no // template argument or a trailing comma. let shouldPass = t.params.access_mode === ''; @@ -573,8 +628,12 @@ g.test('address_space_access_mode') case 'function': fdecl = `var x : u32;`; break; + case 'immediate': + header = requiredLanguageFeatureHeader('immediate'); + mdecl = `var x : u32;`; + break; } - const code = `${mdecl} + const code = `${header}${mdecl} fn foo() { ${fdecl} }`; @@ -601,6 +660,7 @@ g.test('explicit_access_mode') .combine('stage', ['compute' as ShaderStage]) // Only need to check compute shaders ) .fn(t => { + skipIfAddressSpaceNotSupported(t, t.params.addressSpace); const prog = getVarDeclShader(t.params); const info = kAddressSpaceInfo[t.params.addressSpace]; @@ -628,6 +688,7 @@ g.test('implicit_access_mode') .combine('stage', ['compute' as ShaderStage]) // Only need to check compute shaders ) .fn(t => { + skipIfAddressSpaceNotSupported(t, t.params.addressSpace); const prog = getVarDeclShader(t.params); // 7.3 var Declarations @@ -650,6 +711,7 @@ g.test('read_access') .combine('stage', ['compute' as ShaderStage]) // Only need to check compute shaders ) .fn(t => { + skipIfAddressSpaceNotSupported(t, t.params.addressSpace); const prog = getVarDeclShader(t.params, 'let copy = x;'); const ok = supportsRead(t.params); t.expectCompileResult(ok, prog); @@ -668,6 +730,7 @@ g.test('write_access') .combine('stage', ['compute' as ShaderStage]) // Only need to check compute shaders ) .fn(t => { + skipIfAddressSpaceNotSupported(t, t.params.addressSpace); const prog = getVarDeclShader(t.params, 'x = 0;'); const ok = supportsWrite(t.params); t.expectCompileResult(ok, prog);