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..e5c0a3981eb1
--- /dev/null
+++ b/src/webgpu/shader/validation/decl/immediate.spec.ts
@@ -0,0 +1,293 @@
+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: { enable: ``, prelude: ``, type: `u32` },
+  i32: { enable: ``, prelude: ``, type: `i32` },
+  f32: { enable: ``, prelude: ``, type: `f32` },
+  f16: { enable: `enable f16;`, prelude: ``, type: `f16` },
+  vec2u: { enable: ``, prelude: ``, type: `vec2u` },
+  vec3i: { enable: ``, prelude: ``, type: `vec3i` },
+  vec4f: { enable: ``, prelude: ``, type: `vec4f` },
+  vec3h: { enable: `enable f16;`, prelude: ``, type: `vec3h` },
+  mat2x2f: { enable: ``, prelude: ``, type: `mat2x2f` },
+  struct_numeric: { enable: ``, prelude: `struct S { a : u32, b : vec4f }`, type: `S` },
+} as const;
+
+const kInvalidStoreTypes = {
+  bool: { enable: ``, prelude: ``, type: `bool` },
+  vec2_bool: { enable: ``, prelude: ``, type: `vec2<bool>` },
+  atomic_u32: { enable: ``, prelude: ``, type: `atomic<u32>` },
+  ptr_function_u32: { enable: ``, prelude: ``, type: `ptr<function, u32>` },
+  sampler: { enable: ``, prelude: ``, type: `sampler` },
+  sampler_comparison: { enable: ``, prelude: ``, type: `sampler_comparison` },
+  texture_2d: { enable: ``, prelude: ``, type: `texture_2d<f32>` },
+  runtime_array: { enable: ``, prelude: ``, type: `array<u32>` },
+  fixed_array: { enable: ``, prelude: ``, type: `array<u32, 4>` },
+  struct_runtime_array: { enable: ``, prelude: `struct S { data : array<u32> }`, type: `S` },
+  struct_fixed_array: {
+    enable: ``,
+    prelude: `struct S { data : array<vec4u, 4> }`,
+    type: `S`,
+  },
+} 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 testcase = kValidStoreTypes[t.params.type];
+    if (testcase.enable.includes('f16')) {
+      t.skipIfDeviceDoesNotHaveFeature('shader-f16');
+    }
+    const wgsl = `
+${kImmediateHeader}
+${testcase.enable}
+${testcase.prelude}
+var<immediate> data : ${testcase.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];
+    const wgsl = `
+${kImmediateHeader}
+${testcase.enable}
+${testcase.prelude}
+var<immediate> data : ${testcase.type};
+@compute @workgroup_size(1)
+fn main() {
+  _ = data;
+}`;
+    t.expectCompileResult(false, wgsl);
+  });
+
+g.test('scope')
+  .desc('Validates that immediate variables cannot be declared at function scope.')
+  .params(u => u.combine('addressSpace', ['function', 'immediate'] as const))
+  .fn(t => {
+    skipIfImmediateDataNotSupported(t);
+    const wgsl = `
+${kImmediateHeader}
+@compute @workgroup_size(1)
+fn main() {
+  var<${t.params.addressSpace}> data : u32;
+  _ = data;
+}`;
+    t.expectCompileResult(t.params.addressSpace === 'function', 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<immediate> 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<immediate${suffix}> data : u32;
+@compute @workgroup_size(1)
+fn main() {
+  _ = data;
+}`;
+    t.expectCompileResult(t.params.accessMode === '', wgsl);
+  });
+
+const kEntryPointCases = {
+  one_used: {
+    valid: true,
+    body: `
+var<immediate> a : u32;
+@compute @workgroup_size(1)
+fn main() {
+  _ = a;
+}`,
+  },
+  two_declared_one_used: {
+    valid: true,
+    body: `
+var<immediate> a : u32;
+var<immediate> b : u32;
+@compute @workgroup_size(1)
+fn main() {
+  _ = a;
+}`,
+  },
+  two_entry_points_one_each: {
+    valid: true,
+    body: `
+var<immediate> a : u32;
+var<immediate> 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<immediate> a : u32;
+var<immediate> b : u32;
+@compute @workgroup_size(1)
+fn main() {
+  _ = a + b;
+}`,
+  },
+  one_entry_point_uses_two_through_helper: {
+    valid: false,
+    body: `
+var<immediate> a : u32;
+var<immediate> 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<immediate, u32>;
+var<immediate> data : u32;
+@compute @workgroup_size(1)
+fn main() {
+  let p : P = &data;
+  _ = *p;
+}`,
+  },
+  let_inside_function: {
+    valid: true,
+    needsUnrestrictedPointerParameters: false,
+    body: `
+var<immediate> data : u32;
+@compute @workgroup_size(1)
+fn main() {
+  let p : ptr<immediate, u32> = &data;
+  _ = *p;
+}`,
+  },
+  write_through_pointer: {
+    valid: false,
+    needsUnrestrictedPointerParameters: false,
+    body: `
+var<immediate> data : u32;
+@compute @workgroup_size(1)
+fn main() {
+  let p : ptr<immediate, u32> = &data;
+  *p = 1u;
+}`,
+  },
+  pointer_parameter: {
+    valid: true,
+    needsUnrestrictedPointerParameters: true,
+    body: `
+var<immediate> data : u32;
+fn read_data(p : ptr<immediate, u32>) -> u32 {
+  return *p;
+}
+@compute @workgroup_size(1)
+fn main() {
+  _ = read_data(&data);
+}`,
+  },
+  explicit_read_access: {
+    valid: false,
+    needsUnrestrictedPointerParameters: false,
+    body: `
+alias P = ptr<immediate, u32, read>;`,
+  },
+  explicit_write_access: {
+    valid: false,
+    needsUnrestrictedPointerParameters: false,
+    body: `
+alias P = ptr<immediate, u32, write>;`,
+  },
+  explicit_read_write_access: {
+    valid: false,
+    needsUnrestrictedPointerParameters: false,
+    body: `
+alias P = ptr<immediate, u32, read_write>;`,
+  },
+  missing_store_type: {
+    valid: false,
+    needsUnrestrictedPointerParameters: false,
+    body: `
+alias P = ptr<immediate>;`,
+  },
+} as const;
+
+g.test('pointers')
+  .desc('Validates ptr<immediate> 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/util.ts b/src/webgpu/shader/validation/decl/util.ts
index a2f175a5298c..a5f2759bda96 100644
--- a/src/webgpu/shader/validation/decl/util.ts
+++ b/src/webgpu/shader/validation/decl/util.ts
@@ -1,5 +1,3 @@
-import { getGPU } from '../../../../common/util/navigator_gpu.js';
-import { supportsImmediateData } from '../../../../common/util/util.js';
 import {
   AccessMode,
   AddressSpace,
@@ -20,28 +18,19 @@ export function requiredLanguageFeatureHeader(addressSpace: AddressSpace): strin
 }
 
 type AddressSpaceSupportTest = {
-  readonly rec: Parameters<typeof getGPU>[0];
-  skip(message: string): never;
   skipIfLanguageFeatureNotSupported(
     langFeature: NonNullable<AddressSpaceInfo['wgslLanguageFeature']>
   ): void;
 };
 
 export function skipIfImmediateDataNotSupported(t: AddressSpaceSupportTest): void {
-  if (!supportsImmediateData(getGPU(t.rec))) {
-    t.skip('Immediate data not supported');
-  }
+  t.skipIfLanguageFeatureNotSupported('immediate_address_space');
 }
 
 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);
diff --git a/src/webgpu/shader/validation/decl/var.spec.ts b/src/webgpu/shader/validation/decl/var.spec.ts
index ebb4e75b471e..5d5c091a55cd 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,28 @@ const kTypes = {
   },
 };
 
+const kImmediateTypesWithArray = new Set<string>([
+  'array<vec4<bool>>',
+  'array<vec4<bool>, 4>',
+  'array<vec4u>',
+  'array<vec4u, 4>',
+  'array<vec4u, array_size_const>',
+  'array<vec4u, array_size_override>',
+  'S_array_vec4u',
+  'S_array_vec4u_4',
+  'S_array_bool_4',
+] as const);
+
+function isImmediateStoreType(typeName: keyof typeof kTypes): boolean {
+  const type = kTypes[typeName];
+  return (
+    type.isHostShareable &&
+    type.isConstructible &&
+    type.isFixedFootprint &&
+    !kImmediateTypesWithArray.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 +247,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 +294,16 @@ g.test('module_scope_types')
         decl = 'var<workgroup> foo : ';
         shouldPass = type.isFixedFootprint;
         break;
+      case 'immediate':
+        decl = 'var<immediate> 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 +497,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 +579,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 +625,12 @@ g.test('address_space_access_mode')
       case 'function':
         fdecl = `var<function${suffix}> x : u32;`;
         break;
+      case 'immediate':
+        header = requiredLanguageFeatureHeader('immediate');
+        mdecl = `var<immediate${suffix}> x : u32;`;
+        break;
     }
-    const code = `${mdecl}
+    const code = `${header}${mdecl}
     fn foo() {
       ${fdecl}
     }`;
@@ -601,6 +657,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 +685,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 +708,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 +727,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);