diff --git a/src/tint/lang/hlsl/writer/raise/builtin_polyfill.cc b/src/tint/lang/hlsl/writer/raise/builtin_polyfill.cc
index b57b16e..e32a896 100644
--- a/src/tint/lang/hlsl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/hlsl/writer/raise/builtin_polyfill.cc
@@ -923,7 +923,8 @@
                                                       tex, args);
             query = b.Load(query);
             if (!swizzle.IsEmpty()) {
-                query = b.Swizzle(ty.vec2<u32>(), query, swizzle);
+                query = b.Swizzle(ty.vec(ty.u32(), static_cast<uint32_t>(swizzle.Length())), query,
+                                  swizzle);
             }
             call->Result(0)->ReplaceAllUsesWith(query->Result(0));
         });
diff --git a/test/tint/builtins/gen/literal/textureDimensions/0890c6.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/literal/textureDimensions/0890c6.wgsl.expected.ir.dxc.hlsl
index de792de..2c40cb4 100644
--- a/test/tint/builtins/gen/literal/textureDimensions/0890c6.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/textureDimensions/0890c6.wgsl.expected.ir.dxc.hlsl
@@ -1,116 +1,46 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<f32>;
-
-fn textureDimensions_0890c6() -> vec3<u32> {
-  var res : vec3<u32> = textureDimensions(arg_0, 1u);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<float4> arg_0 : register(t0, space1);
+uint3 textureDimensions_0890c6() {
+  Texture3D<float4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(1u), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_0890c6();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_0890c6());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_0890c6();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_0890c6());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_0890c6();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :28:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_0890c6 = func():vec3<u32> {
-  $B2: {
-    %4:texture_3d<f32> = load %arg_0
-    %5:u32 = convert 1u
-    %6:ptr<function, vec4<u32>, read_write> = var
-    %7:ptr<function, u32, read_write> = access %6, 0u
-    %8:ptr<function, u32, read_write> = access %6, 1u
-    %9:ptr<function, u32, read_write> = access %6, 2u
-    %10:ptr<function, u32, read_write> = access %6, 3u
-    %11:void = %4.GetDimensions %5, %7, %8, %9, %10
-    %12:vec4<u32> = load %6
-    %13:vec2<u32> = swizzle %12, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-    %15:vec3<u32> = load %res
-    ret %15
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %17:vec3<u32> = call %textureDimensions_0890c6
-    %18:void = %prevent_dce.Store3 0u, %17
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %20:vec3<u32> = call %textureDimensions_0890c6
-    %21:void = %prevent_dce.Store3 0u, %20
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %24:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %24, vec4<f32>(0.0f)
-    %25:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %26:vec3<u32> = call %textureDimensions_0890c6
-    store %25, %26
-    %27:VertexOutput = load %tint_symbol
-    ret %27
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %29:VertexOutput = call %vertex_main_inner
-    %30:vec4<f32> = access %29, 0u
-    %31:vec3<u32> = access %29, 1u
-    %32:vertex_main_outputs = construct %31, %30
-    ret %32
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/literal/textureDimensions/0890c6.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/textureDimensions/0890c6.wgsl.expected.ir.fxc.hlsl
index de792de..2c40cb4 100644
--- a/test/tint/builtins/gen/literal/textureDimensions/0890c6.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/textureDimensions/0890c6.wgsl.expected.ir.fxc.hlsl
@@ -1,116 +1,46 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<f32>;
-
-fn textureDimensions_0890c6() -> vec3<u32> {
-  var res : vec3<u32> = textureDimensions(arg_0, 1u);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<float4> arg_0 : register(t0, space1);
+uint3 textureDimensions_0890c6() {
+  Texture3D<float4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(1u), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_0890c6();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_0890c6());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_0890c6();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_0890c6());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_0890c6();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :28:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_0890c6 = func():vec3<u32> {
-  $B2: {
-    %4:texture_3d<f32> = load %arg_0
-    %5:u32 = convert 1u
-    %6:ptr<function, vec4<u32>, read_write> = var
-    %7:ptr<function, u32, read_write> = access %6, 0u
-    %8:ptr<function, u32, read_write> = access %6, 1u
-    %9:ptr<function, u32, read_write> = access %6, 2u
-    %10:ptr<function, u32, read_write> = access %6, 3u
-    %11:void = %4.GetDimensions %5, %7, %8, %9, %10
-    %12:vec4<u32> = load %6
-    %13:vec2<u32> = swizzle %12, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-    %15:vec3<u32> = load %res
-    ret %15
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %17:vec3<u32> = call %textureDimensions_0890c6
-    %18:void = %prevent_dce.Store3 0u, %17
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %20:vec3<u32> = call %textureDimensions_0890c6
-    %21:void = %prevent_dce.Store3 0u, %20
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %24:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %24, vec4<f32>(0.0f)
-    %25:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %26:vec3<u32> = call %textureDimensions_0890c6
-    store %25, %26
-    %27:VertexOutput = load %tint_symbol
-    ret %27
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %29:VertexOutput = call %vertex_main_inner
-    %30:vec4<f32> = access %29, 0u
-    %31:vec3<u32> = access %29, 1u
-    %32:vertex_main_outputs = construct %31, %30
-    ret %32
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/literal/textureDimensions/1bc428.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/literal/textureDimensions/1bc428.wgsl.expected.ir.dxc.hlsl
index 1d26675..83ebab7 100644
--- a/test/tint/builtins/gen/literal/textureDimensions/1bc428.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/textureDimensions/1bc428.wgsl.expected.ir.dxc.hlsl
@@ -1,116 +1,46 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<f32>;
-
-fn textureDimensions_1bc428() -> vec3<u32> {
-  var res : vec3<u32> = textureDimensions(arg_0, 1i);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<float4> arg_0 : register(t0, space1);
+uint3 textureDimensions_1bc428() {
+  Texture3D<float4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(int(1)), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_1bc428();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_1bc428());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_1bc428();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_1bc428());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_1bc428();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :28:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_1bc428 = func():vec3<u32> {
-  $B2: {
-    %4:texture_3d<f32> = load %arg_0
-    %5:u32 = convert 1i
-    %6:ptr<function, vec4<u32>, read_write> = var
-    %7:ptr<function, u32, read_write> = access %6, 0u
-    %8:ptr<function, u32, read_write> = access %6, 1u
-    %9:ptr<function, u32, read_write> = access %6, 2u
-    %10:ptr<function, u32, read_write> = access %6, 3u
-    %11:void = %4.GetDimensions %5, %7, %8, %9, %10
-    %12:vec4<u32> = load %6
-    %13:vec2<u32> = swizzle %12, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-    %15:vec3<u32> = load %res
-    ret %15
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %17:vec3<u32> = call %textureDimensions_1bc428
-    %18:void = %prevent_dce.Store3 0u, %17
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %20:vec3<u32> = call %textureDimensions_1bc428
-    %21:void = %prevent_dce.Store3 0u, %20
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %24:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %24, vec4<f32>(0.0f)
-    %25:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %26:vec3<u32> = call %textureDimensions_1bc428
-    store %25, %26
-    %27:VertexOutput = load %tint_symbol
-    ret %27
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %29:VertexOutput = call %vertex_main_inner
-    %30:vec4<f32> = access %29, 0u
-    %31:vec3<u32> = access %29, 1u
-    %32:vertex_main_outputs = construct %31, %30
-    ret %32
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/literal/textureDimensions/1bc428.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/textureDimensions/1bc428.wgsl.expected.ir.fxc.hlsl
index 1d26675..83ebab7 100644
--- a/test/tint/builtins/gen/literal/textureDimensions/1bc428.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/textureDimensions/1bc428.wgsl.expected.ir.fxc.hlsl
@@ -1,116 +1,46 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<f32>;
-
-fn textureDimensions_1bc428() -> vec3<u32> {
-  var res : vec3<u32> = textureDimensions(arg_0, 1i);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<float4> arg_0 : register(t0, space1);
+uint3 textureDimensions_1bc428() {
+  Texture3D<float4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(int(1)), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_1bc428();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_1bc428());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_1bc428();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_1bc428());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_1bc428();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :28:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_1bc428 = func():vec3<u32> {
-  $B2: {
-    %4:texture_3d<f32> = load %arg_0
-    %5:u32 = convert 1i
-    %6:ptr<function, vec4<u32>, read_write> = var
-    %7:ptr<function, u32, read_write> = access %6, 0u
-    %8:ptr<function, u32, read_write> = access %6, 1u
-    %9:ptr<function, u32, read_write> = access %6, 2u
-    %10:ptr<function, u32, read_write> = access %6, 3u
-    %11:void = %4.GetDimensions %5, %7, %8, %9, %10
-    %12:vec4<u32> = load %6
-    %13:vec2<u32> = swizzle %12, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-    %15:vec3<u32> = load %res
-    ret %15
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %17:vec3<u32> = call %textureDimensions_1bc428
-    %18:void = %prevent_dce.Store3 0u, %17
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %20:vec3<u32> = call %textureDimensions_1bc428
-    %21:void = %prevent_dce.Store3 0u, %20
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %24:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %24, vec4<f32>(0.0f)
-    %25:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %26:vec3<u32> = call %textureDimensions_1bc428
-    store %25, %26
-    %27:VertexOutput = load %tint_symbol
-    ret %27
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %29:VertexOutput = call %vertex_main_inner
-    %30:vec4<f32> = access %29, 0u
-    %31:vec3<u32> = access %29, 1u
-    %32:vertex_main_outputs = construct %31, %30
-    ret %32
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/literal/textureDimensions/6e6c7a.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/literal/textureDimensions/6e6c7a.wgsl.expected.ir.dxc.hlsl
index ae88f10..e8fb3b6 100644
--- a/test/tint/builtins/gen/literal/textureDimensions/6e6c7a.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/textureDimensions/6e6c7a.wgsl.expected.ir.dxc.hlsl
@@ -1,116 +1,46 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<u32>;
-
-fn textureDimensions_6e6c7a() -> vec3<u32> {
-  var res : vec3<u32> = textureDimensions(arg_0, 1u);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<uint4> arg_0 : register(t0, space1);
+uint3 textureDimensions_6e6c7a() {
+  Texture3D<uint4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(1u), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_6e6c7a();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_6e6c7a());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_6e6c7a();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_6e6c7a());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_6e6c7a();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :28:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<u32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_6e6c7a = func():vec3<u32> {
-  $B2: {
-    %4:texture_3d<u32> = load %arg_0
-    %5:u32 = convert 1u
-    %6:ptr<function, vec4<u32>, read_write> = var
-    %7:ptr<function, u32, read_write> = access %6, 0u
-    %8:ptr<function, u32, read_write> = access %6, 1u
-    %9:ptr<function, u32, read_write> = access %6, 2u
-    %10:ptr<function, u32, read_write> = access %6, 3u
-    %11:void = %4.GetDimensions %5, %7, %8, %9, %10
-    %12:vec4<u32> = load %6
-    %13:vec2<u32> = swizzle %12, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-    %15:vec3<u32> = load %res
-    ret %15
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %17:vec3<u32> = call %textureDimensions_6e6c7a
-    %18:void = %prevent_dce.Store3 0u, %17
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %20:vec3<u32> = call %textureDimensions_6e6c7a
-    %21:void = %prevent_dce.Store3 0u, %20
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %24:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %24, vec4<f32>(0.0f)
-    %25:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %26:vec3<u32> = call %textureDimensions_6e6c7a
-    store %25, %26
-    %27:VertexOutput = load %tint_symbol
-    ret %27
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %29:VertexOutput = call %vertex_main_inner
-    %30:vec4<f32> = access %29, 0u
-    %31:vec3<u32> = access %29, 1u
-    %32:vertex_main_outputs = construct %31, %30
-    ret %32
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/literal/textureDimensions/6e6c7a.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/textureDimensions/6e6c7a.wgsl.expected.ir.fxc.hlsl
index ae88f10..e8fb3b6 100644
--- a/test/tint/builtins/gen/literal/textureDimensions/6e6c7a.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/textureDimensions/6e6c7a.wgsl.expected.ir.fxc.hlsl
@@ -1,116 +1,46 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<u32>;
-
-fn textureDimensions_6e6c7a() -> vec3<u32> {
-  var res : vec3<u32> = textureDimensions(arg_0, 1u);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<uint4> arg_0 : register(t0, space1);
+uint3 textureDimensions_6e6c7a() {
+  Texture3D<uint4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(1u), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_6e6c7a();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_6e6c7a());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_6e6c7a();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_6e6c7a());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_6e6c7a();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :28:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<u32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_6e6c7a = func():vec3<u32> {
-  $B2: {
-    %4:texture_3d<u32> = load %arg_0
-    %5:u32 = convert 1u
-    %6:ptr<function, vec4<u32>, read_write> = var
-    %7:ptr<function, u32, read_write> = access %6, 0u
-    %8:ptr<function, u32, read_write> = access %6, 1u
-    %9:ptr<function, u32, read_write> = access %6, 2u
-    %10:ptr<function, u32, read_write> = access %6, 3u
-    %11:void = %4.GetDimensions %5, %7, %8, %9, %10
-    %12:vec4<u32> = load %6
-    %13:vec2<u32> = swizzle %12, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-    %15:vec3<u32> = load %res
-    ret %15
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %17:vec3<u32> = call %textureDimensions_6e6c7a
-    %18:void = %prevent_dce.Store3 0u, %17
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %20:vec3<u32> = call %textureDimensions_6e6c7a
-    %21:void = %prevent_dce.Store3 0u, %20
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %24:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %24, vec4<f32>(0.0f)
-    %25:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %26:vec3<u32> = call %textureDimensions_6e6c7a
-    store %25, %26
-    %27:VertexOutput = load %tint_symbol
-    ret %27
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %29:VertexOutput = call %vertex_main_inner
-    %30:vec4<f32> = access %29, 0u
-    %31:vec3<u32> = access %29, 1u
-    %32:vertex_main_outputs = construct %31, %30
-    ret %32
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/literal/textureDimensions/756031.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/literal/textureDimensions/756031.wgsl.expected.ir.dxc.hlsl
index 96a7676..93b9cea 100644
--- a/test/tint/builtins/gen/literal/textureDimensions/756031.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/textureDimensions/756031.wgsl.expected.ir.dxc.hlsl
@@ -1,116 +1,46 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<i32>;
-
-fn textureDimensions_756031() -> vec3<u32> {
-  var res : vec3<u32> = textureDimensions(arg_0, 1i);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<int4> arg_0 : register(t0, space1);
+uint3 textureDimensions_756031() {
+  Texture3D<int4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(int(1)), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_756031();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_756031());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_756031();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_756031());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_756031();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :28:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<i32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_756031 = func():vec3<u32> {
-  $B2: {
-    %4:texture_3d<i32> = load %arg_0
-    %5:u32 = convert 1i
-    %6:ptr<function, vec4<u32>, read_write> = var
-    %7:ptr<function, u32, read_write> = access %6, 0u
-    %8:ptr<function, u32, read_write> = access %6, 1u
-    %9:ptr<function, u32, read_write> = access %6, 2u
-    %10:ptr<function, u32, read_write> = access %6, 3u
-    %11:void = %4.GetDimensions %5, %7, %8, %9, %10
-    %12:vec4<u32> = load %6
-    %13:vec2<u32> = swizzle %12, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-    %15:vec3<u32> = load %res
-    ret %15
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %17:vec3<u32> = call %textureDimensions_756031
-    %18:void = %prevent_dce.Store3 0u, %17
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %20:vec3<u32> = call %textureDimensions_756031
-    %21:void = %prevent_dce.Store3 0u, %20
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %24:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %24, vec4<f32>(0.0f)
-    %25:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %26:vec3<u32> = call %textureDimensions_756031
-    store %25, %26
-    %27:VertexOutput = load %tint_symbol
-    ret %27
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %29:VertexOutput = call %vertex_main_inner
-    %30:vec4<f32> = access %29, 0u
-    %31:vec3<u32> = access %29, 1u
-    %32:vertex_main_outputs = construct %31, %30
-    ret %32
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/literal/textureDimensions/756031.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/textureDimensions/756031.wgsl.expected.ir.fxc.hlsl
index 96a7676..93b9cea 100644
--- a/test/tint/builtins/gen/literal/textureDimensions/756031.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/textureDimensions/756031.wgsl.expected.ir.fxc.hlsl
@@ -1,116 +1,46 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<i32>;
-
-fn textureDimensions_756031() -> vec3<u32> {
-  var res : vec3<u32> = textureDimensions(arg_0, 1i);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<int4> arg_0 : register(t0, space1);
+uint3 textureDimensions_756031() {
+  Texture3D<int4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(int(1)), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_756031();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_756031());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_756031();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_756031());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_756031();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :28:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<i32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_756031 = func():vec3<u32> {
-  $B2: {
-    %4:texture_3d<i32> = load %arg_0
-    %5:u32 = convert 1i
-    %6:ptr<function, vec4<u32>, read_write> = var
-    %7:ptr<function, u32, read_write> = access %6, 0u
-    %8:ptr<function, u32, read_write> = access %6, 1u
-    %9:ptr<function, u32, read_write> = access %6, 2u
-    %10:ptr<function, u32, read_write> = access %6, 3u
-    %11:void = %4.GetDimensions %5, %7, %8, %9, %10
-    %12:vec4<u32> = load %6
-    %13:vec2<u32> = swizzle %12, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-    %15:vec3<u32> = load %res
-    ret %15
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %17:vec3<u32> = call %textureDimensions_756031
-    %18:void = %prevent_dce.Store3 0u, %17
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %20:vec3<u32> = call %textureDimensions_756031
-    %21:void = %prevent_dce.Store3 0u, %20
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %24:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %24, vec4<f32>(0.0f)
-    %25:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %26:vec3<u32> = call %textureDimensions_756031
-    store %25, %26
-    %27:VertexOutput = load %tint_symbol
-    ret %27
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %29:VertexOutput = call %vertex_main_inner
-    %30:vec4<f32> = access %29, 0u
-    %31:vec3<u32> = access %29, 1u
-    %32:vertex_main_outputs = construct %31, %30
-    ret %32
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/literal/textureDimensions/c871f3.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/literal/textureDimensions/c871f3.wgsl.expected.ir.dxc.hlsl
index 063ab13..34cda12 100644
--- a/test/tint/builtins/gen/literal/textureDimensions/c871f3.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/textureDimensions/c871f3.wgsl.expected.ir.dxc.hlsl
@@ -1,116 +1,46 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<i32>;
-
-fn textureDimensions_c871f3() -> vec3<u32> {
-  var res : vec3<u32> = textureDimensions(arg_0, 1u);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<int4> arg_0 : register(t0, space1);
+uint3 textureDimensions_c871f3() {
+  Texture3D<int4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(1u), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_c871f3();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_c871f3());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_c871f3();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_c871f3());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_c871f3();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :28:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<i32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_c871f3 = func():vec3<u32> {
-  $B2: {
-    %4:texture_3d<i32> = load %arg_0
-    %5:u32 = convert 1u
-    %6:ptr<function, vec4<u32>, read_write> = var
-    %7:ptr<function, u32, read_write> = access %6, 0u
-    %8:ptr<function, u32, read_write> = access %6, 1u
-    %9:ptr<function, u32, read_write> = access %6, 2u
-    %10:ptr<function, u32, read_write> = access %6, 3u
-    %11:void = %4.GetDimensions %5, %7, %8, %9, %10
-    %12:vec4<u32> = load %6
-    %13:vec2<u32> = swizzle %12, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-    %15:vec3<u32> = load %res
-    ret %15
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %17:vec3<u32> = call %textureDimensions_c871f3
-    %18:void = %prevent_dce.Store3 0u, %17
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %20:vec3<u32> = call %textureDimensions_c871f3
-    %21:void = %prevent_dce.Store3 0u, %20
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %24:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %24, vec4<f32>(0.0f)
-    %25:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %26:vec3<u32> = call %textureDimensions_c871f3
-    store %25, %26
-    %27:VertexOutput = load %tint_symbol
-    ret %27
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %29:VertexOutput = call %vertex_main_inner
-    %30:vec4<f32> = access %29, 0u
-    %31:vec3<u32> = access %29, 1u
-    %32:vertex_main_outputs = construct %31, %30
-    ret %32
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/literal/textureDimensions/c871f3.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/textureDimensions/c871f3.wgsl.expected.ir.fxc.hlsl
index 063ab13..34cda12 100644
--- a/test/tint/builtins/gen/literal/textureDimensions/c871f3.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/textureDimensions/c871f3.wgsl.expected.ir.fxc.hlsl
@@ -1,116 +1,46 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<i32>;
-
-fn textureDimensions_c871f3() -> vec3<u32> {
-  var res : vec3<u32> = textureDimensions(arg_0, 1u);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<int4> arg_0 : register(t0, space1);
+uint3 textureDimensions_c871f3() {
+  Texture3D<int4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(1u), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_c871f3();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_c871f3());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_c871f3();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_c871f3());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_c871f3();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :28:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<i32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_c871f3 = func():vec3<u32> {
-  $B2: {
-    %4:texture_3d<i32> = load %arg_0
-    %5:u32 = convert 1u
-    %6:ptr<function, vec4<u32>, read_write> = var
-    %7:ptr<function, u32, read_write> = access %6, 0u
-    %8:ptr<function, u32, read_write> = access %6, 1u
-    %9:ptr<function, u32, read_write> = access %6, 2u
-    %10:ptr<function, u32, read_write> = access %6, 3u
-    %11:void = %4.GetDimensions %5, %7, %8, %9, %10
-    %12:vec4<u32> = load %6
-    %13:vec2<u32> = swizzle %12, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-    %15:vec3<u32> = load %res
-    ret %15
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %17:vec3<u32> = call %textureDimensions_c871f3
-    %18:void = %prevent_dce.Store3 0u, %17
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %20:vec3<u32> = call %textureDimensions_c871f3
-    %21:void = %prevent_dce.Store3 0u, %20
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %24:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %24, vec4<f32>(0.0f)
-    %25:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %26:vec3<u32> = call %textureDimensions_c871f3
-    store %25, %26
-    %27:VertexOutput = load %tint_symbol
-    ret %27
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %29:VertexOutput = call %vertex_main_inner
-    %30:vec4<f32> = access %29, 0u
-    %31:vec3<u32> = access %29, 1u
-    %32:vertex_main_outputs = construct %31, %30
-    ret %32
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/literal/textureDimensions/e5a203.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/literal/textureDimensions/e5a203.wgsl.expected.ir.dxc.hlsl
index 408ffd5..7099bd1 100644
--- a/test/tint/builtins/gen/literal/textureDimensions/e5a203.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/textureDimensions/e5a203.wgsl.expected.ir.dxc.hlsl
@@ -1,116 +1,46 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<u32>;
-
-fn textureDimensions_e5a203() -> vec3<u32> {
-  var res : vec3<u32> = textureDimensions(arg_0, 1i);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<uint4> arg_0 : register(t0, space1);
+uint3 textureDimensions_e5a203() {
+  Texture3D<uint4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(int(1)), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_e5a203();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_e5a203());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_e5a203();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_e5a203());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_e5a203();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :28:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<u32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_e5a203 = func():vec3<u32> {
-  $B2: {
-    %4:texture_3d<u32> = load %arg_0
-    %5:u32 = convert 1i
-    %6:ptr<function, vec4<u32>, read_write> = var
-    %7:ptr<function, u32, read_write> = access %6, 0u
-    %8:ptr<function, u32, read_write> = access %6, 1u
-    %9:ptr<function, u32, read_write> = access %6, 2u
-    %10:ptr<function, u32, read_write> = access %6, 3u
-    %11:void = %4.GetDimensions %5, %7, %8, %9, %10
-    %12:vec4<u32> = load %6
-    %13:vec2<u32> = swizzle %12, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-    %15:vec3<u32> = load %res
-    ret %15
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %17:vec3<u32> = call %textureDimensions_e5a203
-    %18:void = %prevent_dce.Store3 0u, %17
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %20:vec3<u32> = call %textureDimensions_e5a203
-    %21:void = %prevent_dce.Store3 0u, %20
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %24:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %24, vec4<f32>(0.0f)
-    %25:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %26:vec3<u32> = call %textureDimensions_e5a203
-    store %25, %26
-    %27:VertexOutput = load %tint_symbol
-    ret %27
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %29:VertexOutput = call %vertex_main_inner
-    %30:vec4<f32> = access %29, 0u
-    %31:vec3<u32> = access %29, 1u
-    %32:vertex_main_outputs = construct %31, %30
-    ret %32
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/literal/textureDimensions/e5a203.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/textureDimensions/e5a203.wgsl.expected.ir.fxc.hlsl
index 408ffd5..7099bd1 100644
--- a/test/tint/builtins/gen/literal/textureDimensions/e5a203.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/textureDimensions/e5a203.wgsl.expected.ir.fxc.hlsl
@@ -1,116 +1,46 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<u32>;
-
-fn textureDimensions_e5a203() -> vec3<u32> {
-  var res : vec3<u32> = textureDimensions(arg_0, 1i);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<uint4> arg_0 : register(t0, space1);
+uint3 textureDimensions_e5a203() {
+  Texture3D<uint4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(int(1)), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_e5a203();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_e5a203());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_e5a203();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_e5a203());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_e5a203();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :28:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<u32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_e5a203 = func():vec3<u32> {
-  $B2: {
-    %4:texture_3d<u32> = load %arg_0
-    %5:u32 = convert 1i
-    %6:ptr<function, vec4<u32>, read_write> = var
-    %7:ptr<function, u32, read_write> = access %6, 0u
-    %8:ptr<function, u32, read_write> = access %6, 1u
-    %9:ptr<function, u32, read_write> = access %6, 2u
-    %10:ptr<function, u32, read_write> = access %6, 3u
-    %11:void = %4.GetDimensions %5, %7, %8, %9, %10
-    %12:vec4<u32> = load %6
-    %13:vec2<u32> = swizzle %12, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %13
-    %15:vec3<u32> = load %res
-    ret %15
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %17:vec3<u32> = call %textureDimensions_e5a203
-    %18:void = %prevent_dce.Store3 0u, %17
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %20:vec3<u32> = call %textureDimensions_e5a203
-    %21:void = %prevent_dce.Store3 0u, %20
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %24:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %24, vec4<f32>(0.0f)
-    %25:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %26:vec3<u32> = call %textureDimensions_e5a203
-    store %25, %26
-    %27:VertexOutput = load %tint_symbol
-    ret %27
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %29:VertexOutput = call %vertex_main_inner
-    %30:vec4<f32> = access %29, 0u
-    %31:vec3<u32> = access %29, 1u
-    %32:vertex_main_outputs = construct %31, %30
-    ret %32
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/textureDimensions/0890c6.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/textureDimensions/0890c6.wgsl.expected.ir.dxc.hlsl
index 33670ce..f011d38 100644
--- a/test/tint/builtins/gen/var/textureDimensions/0890c6.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/textureDimensions/0890c6.wgsl.expected.ir.dxc.hlsl
@@ -1,119 +1,47 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<f32>;
-
-fn textureDimensions_0890c6() -> vec3<u32> {
-  var arg_1 = 1u;
-  var res : vec3<u32> = textureDimensions(arg_0, arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<float4> arg_0 : register(t0, space1);
+uint3 textureDimensions_0890c6() {
+  uint arg_1 = 1u;
+  Texture3D<float4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(arg_1), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_0890c6();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_0890c6());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_0890c6();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_0890c6());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_0890c6();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :30:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_0890c6 = func():vec3<u32> {
-  $B2: {
-    %arg_1:ptr<function, u32, read_write> = var, 1u
-    %5:texture_3d<f32> = load %arg_0
-    %6:u32 = load %arg_1
-    %7:u32 = convert %6
-    %8:ptr<function, vec4<u32>, read_write> = var
-    %9:ptr<function, u32, read_write> = access %8, 0u
-    %10:ptr<function, u32, read_write> = access %8, 1u
-    %11:ptr<function, u32, read_write> = access %8, 2u
-    %12:ptr<function, u32, read_write> = access %8, 3u
-    %13:void = %5.GetDimensions %7, %9, %10, %11, %12
-    %14:vec4<u32> = load %8
-    %15:vec2<u32> = swizzle %14, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-    %17:vec3<u32> = load %res
-    ret %17
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %19:vec3<u32> = call %textureDimensions_0890c6
-    %20:void = %prevent_dce.Store3 0u, %19
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %22:vec3<u32> = call %textureDimensions_0890c6
-    %23:void = %prevent_dce.Store3 0u, %22
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %26:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %26, vec4<f32>(0.0f)
-    %27:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %28:vec3<u32> = call %textureDimensions_0890c6
-    store %27, %28
-    %29:VertexOutput = load %tint_symbol
-    ret %29
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %31:VertexOutput = call %vertex_main_inner
-    %32:vec4<f32> = access %31, 0u
-    %33:vec3<u32> = access %31, 1u
-    %34:vertex_main_outputs = construct %33, %32
-    ret %34
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/textureDimensions/0890c6.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/textureDimensions/0890c6.wgsl.expected.ir.fxc.hlsl
index 33670ce..f011d38 100644
--- a/test/tint/builtins/gen/var/textureDimensions/0890c6.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/var/textureDimensions/0890c6.wgsl.expected.ir.fxc.hlsl
@@ -1,119 +1,47 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<f32>;
-
-fn textureDimensions_0890c6() -> vec3<u32> {
-  var arg_1 = 1u;
-  var res : vec3<u32> = textureDimensions(arg_0, arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<float4> arg_0 : register(t0, space1);
+uint3 textureDimensions_0890c6() {
+  uint arg_1 = 1u;
+  Texture3D<float4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(arg_1), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_0890c6();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_0890c6());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_0890c6();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_0890c6());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_0890c6();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :30:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_0890c6 = func():vec3<u32> {
-  $B2: {
-    %arg_1:ptr<function, u32, read_write> = var, 1u
-    %5:texture_3d<f32> = load %arg_0
-    %6:u32 = load %arg_1
-    %7:u32 = convert %6
-    %8:ptr<function, vec4<u32>, read_write> = var
-    %9:ptr<function, u32, read_write> = access %8, 0u
-    %10:ptr<function, u32, read_write> = access %8, 1u
-    %11:ptr<function, u32, read_write> = access %8, 2u
-    %12:ptr<function, u32, read_write> = access %8, 3u
-    %13:void = %5.GetDimensions %7, %9, %10, %11, %12
-    %14:vec4<u32> = load %8
-    %15:vec2<u32> = swizzle %14, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-    %17:vec3<u32> = load %res
-    ret %17
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %19:vec3<u32> = call %textureDimensions_0890c6
-    %20:void = %prevent_dce.Store3 0u, %19
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %22:vec3<u32> = call %textureDimensions_0890c6
-    %23:void = %prevent_dce.Store3 0u, %22
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %26:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %26, vec4<f32>(0.0f)
-    %27:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %28:vec3<u32> = call %textureDimensions_0890c6
-    store %27, %28
-    %29:VertexOutput = load %tint_symbol
-    ret %29
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %31:VertexOutput = call %vertex_main_inner
-    %32:vec4<f32> = access %31, 0u
-    %33:vec3<u32> = access %31, 1u
-    %34:vertex_main_outputs = construct %33, %32
-    ret %34
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/textureDimensions/1bc428.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/textureDimensions/1bc428.wgsl.expected.ir.dxc.hlsl
index 5f17fa6..4970946 100644
--- a/test/tint/builtins/gen/var/textureDimensions/1bc428.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/textureDimensions/1bc428.wgsl.expected.ir.dxc.hlsl
@@ -1,119 +1,47 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<f32>;
-
-fn textureDimensions_1bc428() -> vec3<u32> {
-  var arg_1 = 1i;
-  var res : vec3<u32> = textureDimensions(arg_0, arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<float4> arg_0 : register(t0, space1);
+uint3 textureDimensions_1bc428() {
+  int arg_1 = int(1);
+  Texture3D<float4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(arg_1), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_1bc428();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_1bc428());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_1bc428();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_1bc428());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_1bc428();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :30:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_1bc428 = func():vec3<u32> {
-  $B2: {
-    %arg_1:ptr<function, i32, read_write> = var, 1i
-    %5:texture_3d<f32> = load %arg_0
-    %6:i32 = load %arg_1
-    %7:u32 = convert %6
-    %8:ptr<function, vec4<u32>, read_write> = var
-    %9:ptr<function, u32, read_write> = access %8, 0u
-    %10:ptr<function, u32, read_write> = access %8, 1u
-    %11:ptr<function, u32, read_write> = access %8, 2u
-    %12:ptr<function, u32, read_write> = access %8, 3u
-    %13:void = %5.GetDimensions %7, %9, %10, %11, %12
-    %14:vec4<u32> = load %8
-    %15:vec2<u32> = swizzle %14, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-    %17:vec3<u32> = load %res
-    ret %17
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %19:vec3<u32> = call %textureDimensions_1bc428
-    %20:void = %prevent_dce.Store3 0u, %19
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %22:vec3<u32> = call %textureDimensions_1bc428
-    %23:void = %prevent_dce.Store3 0u, %22
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %26:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %26, vec4<f32>(0.0f)
-    %27:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %28:vec3<u32> = call %textureDimensions_1bc428
-    store %27, %28
-    %29:VertexOutput = load %tint_symbol
-    ret %29
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %31:VertexOutput = call %vertex_main_inner
-    %32:vec4<f32> = access %31, 0u
-    %33:vec3<u32> = access %31, 1u
-    %34:vertex_main_outputs = construct %33, %32
-    ret %34
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/textureDimensions/1bc428.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/textureDimensions/1bc428.wgsl.expected.ir.fxc.hlsl
index 5f17fa6..4970946 100644
--- a/test/tint/builtins/gen/var/textureDimensions/1bc428.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/var/textureDimensions/1bc428.wgsl.expected.ir.fxc.hlsl
@@ -1,119 +1,47 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<f32>;
-
-fn textureDimensions_1bc428() -> vec3<u32> {
-  var arg_1 = 1i;
-  var res : vec3<u32> = textureDimensions(arg_0, arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<float4> arg_0 : register(t0, space1);
+uint3 textureDimensions_1bc428() {
+  int arg_1 = int(1);
+  Texture3D<float4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(arg_1), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_1bc428();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_1bc428());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_1bc428();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_1bc428());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_1bc428();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :30:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_1bc428 = func():vec3<u32> {
-  $B2: {
-    %arg_1:ptr<function, i32, read_write> = var, 1i
-    %5:texture_3d<f32> = load %arg_0
-    %6:i32 = load %arg_1
-    %7:u32 = convert %6
-    %8:ptr<function, vec4<u32>, read_write> = var
-    %9:ptr<function, u32, read_write> = access %8, 0u
-    %10:ptr<function, u32, read_write> = access %8, 1u
-    %11:ptr<function, u32, read_write> = access %8, 2u
-    %12:ptr<function, u32, read_write> = access %8, 3u
-    %13:void = %5.GetDimensions %7, %9, %10, %11, %12
-    %14:vec4<u32> = load %8
-    %15:vec2<u32> = swizzle %14, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-    %17:vec3<u32> = load %res
-    ret %17
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %19:vec3<u32> = call %textureDimensions_1bc428
-    %20:void = %prevent_dce.Store3 0u, %19
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %22:vec3<u32> = call %textureDimensions_1bc428
-    %23:void = %prevent_dce.Store3 0u, %22
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %26:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %26, vec4<f32>(0.0f)
-    %27:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %28:vec3<u32> = call %textureDimensions_1bc428
-    store %27, %28
-    %29:VertexOutput = load %tint_symbol
-    ret %29
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %31:VertexOutput = call %vertex_main_inner
-    %32:vec4<f32> = access %31, 0u
-    %33:vec3<u32> = access %31, 1u
-    %34:vertex_main_outputs = construct %33, %32
-    ret %34
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/textureDimensions/6e6c7a.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/textureDimensions/6e6c7a.wgsl.expected.ir.dxc.hlsl
index c1ddb99..eddecfd 100644
--- a/test/tint/builtins/gen/var/textureDimensions/6e6c7a.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/textureDimensions/6e6c7a.wgsl.expected.ir.dxc.hlsl
@@ -1,119 +1,47 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<u32>;
-
-fn textureDimensions_6e6c7a() -> vec3<u32> {
-  var arg_1 = 1u;
-  var res : vec3<u32> = textureDimensions(arg_0, arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<uint4> arg_0 : register(t0, space1);
+uint3 textureDimensions_6e6c7a() {
+  uint arg_1 = 1u;
+  Texture3D<uint4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(arg_1), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_6e6c7a();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_6e6c7a());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_6e6c7a();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_6e6c7a());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_6e6c7a();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :30:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<u32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_6e6c7a = func():vec3<u32> {
-  $B2: {
-    %arg_1:ptr<function, u32, read_write> = var, 1u
-    %5:texture_3d<u32> = load %arg_0
-    %6:u32 = load %arg_1
-    %7:u32 = convert %6
-    %8:ptr<function, vec4<u32>, read_write> = var
-    %9:ptr<function, u32, read_write> = access %8, 0u
-    %10:ptr<function, u32, read_write> = access %8, 1u
-    %11:ptr<function, u32, read_write> = access %8, 2u
-    %12:ptr<function, u32, read_write> = access %8, 3u
-    %13:void = %5.GetDimensions %7, %9, %10, %11, %12
-    %14:vec4<u32> = load %8
-    %15:vec2<u32> = swizzle %14, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-    %17:vec3<u32> = load %res
-    ret %17
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %19:vec3<u32> = call %textureDimensions_6e6c7a
-    %20:void = %prevent_dce.Store3 0u, %19
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %22:vec3<u32> = call %textureDimensions_6e6c7a
-    %23:void = %prevent_dce.Store3 0u, %22
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %26:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %26, vec4<f32>(0.0f)
-    %27:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %28:vec3<u32> = call %textureDimensions_6e6c7a
-    store %27, %28
-    %29:VertexOutput = load %tint_symbol
-    ret %29
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %31:VertexOutput = call %vertex_main_inner
-    %32:vec4<f32> = access %31, 0u
-    %33:vec3<u32> = access %31, 1u
-    %34:vertex_main_outputs = construct %33, %32
-    ret %34
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/textureDimensions/6e6c7a.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/textureDimensions/6e6c7a.wgsl.expected.ir.fxc.hlsl
index c1ddb99..eddecfd 100644
--- a/test/tint/builtins/gen/var/textureDimensions/6e6c7a.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/var/textureDimensions/6e6c7a.wgsl.expected.ir.fxc.hlsl
@@ -1,119 +1,47 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<u32>;
-
-fn textureDimensions_6e6c7a() -> vec3<u32> {
-  var arg_1 = 1u;
-  var res : vec3<u32> = textureDimensions(arg_0, arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<uint4> arg_0 : register(t0, space1);
+uint3 textureDimensions_6e6c7a() {
+  uint arg_1 = 1u;
+  Texture3D<uint4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(arg_1), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_6e6c7a();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_6e6c7a());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_6e6c7a();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_6e6c7a());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_6e6c7a();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :30:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<u32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_6e6c7a = func():vec3<u32> {
-  $B2: {
-    %arg_1:ptr<function, u32, read_write> = var, 1u
-    %5:texture_3d<u32> = load %arg_0
-    %6:u32 = load %arg_1
-    %7:u32 = convert %6
-    %8:ptr<function, vec4<u32>, read_write> = var
-    %9:ptr<function, u32, read_write> = access %8, 0u
-    %10:ptr<function, u32, read_write> = access %8, 1u
-    %11:ptr<function, u32, read_write> = access %8, 2u
-    %12:ptr<function, u32, read_write> = access %8, 3u
-    %13:void = %5.GetDimensions %7, %9, %10, %11, %12
-    %14:vec4<u32> = load %8
-    %15:vec2<u32> = swizzle %14, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-    %17:vec3<u32> = load %res
-    ret %17
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %19:vec3<u32> = call %textureDimensions_6e6c7a
-    %20:void = %prevent_dce.Store3 0u, %19
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %22:vec3<u32> = call %textureDimensions_6e6c7a
-    %23:void = %prevent_dce.Store3 0u, %22
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %26:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %26, vec4<f32>(0.0f)
-    %27:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %28:vec3<u32> = call %textureDimensions_6e6c7a
-    store %27, %28
-    %29:VertexOutput = load %tint_symbol
-    ret %29
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %31:VertexOutput = call %vertex_main_inner
-    %32:vec4<f32> = access %31, 0u
-    %33:vec3<u32> = access %31, 1u
-    %34:vertex_main_outputs = construct %33, %32
-    ret %34
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/textureDimensions/756031.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/textureDimensions/756031.wgsl.expected.ir.dxc.hlsl
index 7149461..dc43ba9 100644
--- a/test/tint/builtins/gen/var/textureDimensions/756031.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/textureDimensions/756031.wgsl.expected.ir.dxc.hlsl
@@ -1,119 +1,47 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<i32>;
-
-fn textureDimensions_756031() -> vec3<u32> {
-  var arg_1 = 1i;
-  var res : vec3<u32> = textureDimensions(arg_0, arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<int4> arg_0 : register(t0, space1);
+uint3 textureDimensions_756031() {
+  int arg_1 = int(1);
+  Texture3D<int4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(arg_1), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_756031();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_756031());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_756031();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_756031());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_756031();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :30:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<i32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_756031 = func():vec3<u32> {
-  $B2: {
-    %arg_1:ptr<function, i32, read_write> = var, 1i
-    %5:texture_3d<i32> = load %arg_0
-    %6:i32 = load %arg_1
-    %7:u32 = convert %6
-    %8:ptr<function, vec4<u32>, read_write> = var
-    %9:ptr<function, u32, read_write> = access %8, 0u
-    %10:ptr<function, u32, read_write> = access %8, 1u
-    %11:ptr<function, u32, read_write> = access %8, 2u
-    %12:ptr<function, u32, read_write> = access %8, 3u
-    %13:void = %5.GetDimensions %7, %9, %10, %11, %12
-    %14:vec4<u32> = load %8
-    %15:vec2<u32> = swizzle %14, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-    %17:vec3<u32> = load %res
-    ret %17
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %19:vec3<u32> = call %textureDimensions_756031
-    %20:void = %prevent_dce.Store3 0u, %19
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %22:vec3<u32> = call %textureDimensions_756031
-    %23:void = %prevent_dce.Store3 0u, %22
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %26:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %26, vec4<f32>(0.0f)
-    %27:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %28:vec3<u32> = call %textureDimensions_756031
-    store %27, %28
-    %29:VertexOutput = load %tint_symbol
-    ret %29
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %31:VertexOutput = call %vertex_main_inner
-    %32:vec4<f32> = access %31, 0u
-    %33:vec3<u32> = access %31, 1u
-    %34:vertex_main_outputs = construct %33, %32
-    ret %34
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/textureDimensions/756031.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/textureDimensions/756031.wgsl.expected.ir.fxc.hlsl
index 7149461..dc43ba9 100644
--- a/test/tint/builtins/gen/var/textureDimensions/756031.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/var/textureDimensions/756031.wgsl.expected.ir.fxc.hlsl
@@ -1,119 +1,47 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<i32>;
-
-fn textureDimensions_756031() -> vec3<u32> {
-  var arg_1 = 1i;
-  var res : vec3<u32> = textureDimensions(arg_0, arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<int4> arg_0 : register(t0, space1);
+uint3 textureDimensions_756031() {
+  int arg_1 = int(1);
+  Texture3D<int4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(arg_1), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_756031();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_756031());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_756031();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_756031());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_756031();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :30:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<i32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_756031 = func():vec3<u32> {
-  $B2: {
-    %arg_1:ptr<function, i32, read_write> = var, 1i
-    %5:texture_3d<i32> = load %arg_0
-    %6:i32 = load %arg_1
-    %7:u32 = convert %6
-    %8:ptr<function, vec4<u32>, read_write> = var
-    %9:ptr<function, u32, read_write> = access %8, 0u
-    %10:ptr<function, u32, read_write> = access %8, 1u
-    %11:ptr<function, u32, read_write> = access %8, 2u
-    %12:ptr<function, u32, read_write> = access %8, 3u
-    %13:void = %5.GetDimensions %7, %9, %10, %11, %12
-    %14:vec4<u32> = load %8
-    %15:vec2<u32> = swizzle %14, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-    %17:vec3<u32> = load %res
-    ret %17
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %19:vec3<u32> = call %textureDimensions_756031
-    %20:void = %prevent_dce.Store3 0u, %19
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %22:vec3<u32> = call %textureDimensions_756031
-    %23:void = %prevent_dce.Store3 0u, %22
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %26:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %26, vec4<f32>(0.0f)
-    %27:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %28:vec3<u32> = call %textureDimensions_756031
-    store %27, %28
-    %29:VertexOutput = load %tint_symbol
-    ret %29
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %31:VertexOutput = call %vertex_main_inner
-    %32:vec4<f32> = access %31, 0u
-    %33:vec3<u32> = access %31, 1u
-    %34:vertex_main_outputs = construct %33, %32
-    ret %34
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/textureDimensions/c871f3.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/textureDimensions/c871f3.wgsl.expected.ir.dxc.hlsl
index 0cfdb32..03c774d 100644
--- a/test/tint/builtins/gen/var/textureDimensions/c871f3.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/textureDimensions/c871f3.wgsl.expected.ir.dxc.hlsl
@@ -1,119 +1,47 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<i32>;
-
-fn textureDimensions_c871f3() -> vec3<u32> {
-  var arg_1 = 1u;
-  var res : vec3<u32> = textureDimensions(arg_0, arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<int4> arg_0 : register(t0, space1);
+uint3 textureDimensions_c871f3() {
+  uint arg_1 = 1u;
+  Texture3D<int4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(arg_1), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_c871f3();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_c871f3());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_c871f3();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_c871f3());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_c871f3();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :30:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<i32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_c871f3 = func():vec3<u32> {
-  $B2: {
-    %arg_1:ptr<function, u32, read_write> = var, 1u
-    %5:texture_3d<i32> = load %arg_0
-    %6:u32 = load %arg_1
-    %7:u32 = convert %6
-    %8:ptr<function, vec4<u32>, read_write> = var
-    %9:ptr<function, u32, read_write> = access %8, 0u
-    %10:ptr<function, u32, read_write> = access %8, 1u
-    %11:ptr<function, u32, read_write> = access %8, 2u
-    %12:ptr<function, u32, read_write> = access %8, 3u
-    %13:void = %5.GetDimensions %7, %9, %10, %11, %12
-    %14:vec4<u32> = load %8
-    %15:vec2<u32> = swizzle %14, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-    %17:vec3<u32> = load %res
-    ret %17
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %19:vec3<u32> = call %textureDimensions_c871f3
-    %20:void = %prevent_dce.Store3 0u, %19
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %22:vec3<u32> = call %textureDimensions_c871f3
-    %23:void = %prevent_dce.Store3 0u, %22
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %26:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %26, vec4<f32>(0.0f)
-    %27:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %28:vec3<u32> = call %textureDimensions_c871f3
-    store %27, %28
-    %29:VertexOutput = load %tint_symbol
-    ret %29
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %31:VertexOutput = call %vertex_main_inner
-    %32:vec4<f32> = access %31, 0u
-    %33:vec3<u32> = access %31, 1u
-    %34:vertex_main_outputs = construct %33, %32
-    ret %34
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/textureDimensions/c871f3.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/textureDimensions/c871f3.wgsl.expected.ir.fxc.hlsl
index 0cfdb32..03c774d 100644
--- a/test/tint/builtins/gen/var/textureDimensions/c871f3.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/var/textureDimensions/c871f3.wgsl.expected.ir.fxc.hlsl
@@ -1,119 +1,47 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<i32>;
-
-fn textureDimensions_c871f3() -> vec3<u32> {
-  var arg_1 = 1u;
-  var res : vec3<u32> = textureDimensions(arg_0, arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<int4> arg_0 : register(t0, space1);
+uint3 textureDimensions_c871f3() {
+  uint arg_1 = 1u;
+  Texture3D<int4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(arg_1), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_c871f3();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_c871f3());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_c871f3();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_c871f3());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_c871f3();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :30:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<i32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_c871f3 = func():vec3<u32> {
-  $B2: {
-    %arg_1:ptr<function, u32, read_write> = var, 1u
-    %5:texture_3d<i32> = load %arg_0
-    %6:u32 = load %arg_1
-    %7:u32 = convert %6
-    %8:ptr<function, vec4<u32>, read_write> = var
-    %9:ptr<function, u32, read_write> = access %8, 0u
-    %10:ptr<function, u32, read_write> = access %8, 1u
-    %11:ptr<function, u32, read_write> = access %8, 2u
-    %12:ptr<function, u32, read_write> = access %8, 3u
-    %13:void = %5.GetDimensions %7, %9, %10, %11, %12
-    %14:vec4<u32> = load %8
-    %15:vec2<u32> = swizzle %14, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-    %17:vec3<u32> = load %res
-    ret %17
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %19:vec3<u32> = call %textureDimensions_c871f3
-    %20:void = %prevent_dce.Store3 0u, %19
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %22:vec3<u32> = call %textureDimensions_c871f3
-    %23:void = %prevent_dce.Store3 0u, %22
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %26:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %26, vec4<f32>(0.0f)
-    %27:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %28:vec3<u32> = call %textureDimensions_c871f3
-    store %27, %28
-    %29:VertexOutput = load %tint_symbol
-    ret %29
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %31:VertexOutput = call %vertex_main_inner
-    %32:vec4<f32> = access %31, 0u
-    %33:vec3<u32> = access %31, 1u
-    %34:vertex_main_outputs = construct %33, %32
-    ret %34
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/textureDimensions/e5a203.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/textureDimensions/e5a203.wgsl.expected.ir.dxc.hlsl
index c0a859d..b275bf3 100644
--- a/test/tint/builtins/gen/var/textureDimensions/e5a203.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/textureDimensions/e5a203.wgsl.expected.ir.dxc.hlsl
@@ -1,119 +1,47 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<u32>;
-
-fn textureDimensions_e5a203() -> vec3<u32> {
-  var arg_1 = 1i;
-  var res : vec3<u32> = textureDimensions(arg_0, arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<uint4> arg_0 : register(t0, space1);
+uint3 textureDimensions_e5a203() {
+  int arg_1 = int(1);
+  Texture3D<uint4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(arg_1), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_e5a203();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_e5a203());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_e5a203();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_e5a203());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_e5a203();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :30:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<u32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_e5a203 = func():vec3<u32> {
-  $B2: {
-    %arg_1:ptr<function, i32, read_write> = var, 1i
-    %5:texture_3d<u32> = load %arg_0
-    %6:i32 = load %arg_1
-    %7:u32 = convert %6
-    %8:ptr<function, vec4<u32>, read_write> = var
-    %9:ptr<function, u32, read_write> = access %8, 0u
-    %10:ptr<function, u32, read_write> = access %8, 1u
-    %11:ptr<function, u32, read_write> = access %8, 2u
-    %12:ptr<function, u32, read_write> = access %8, 3u
-    %13:void = %5.GetDimensions %7, %9, %10, %11, %12
-    %14:vec4<u32> = load %8
-    %15:vec2<u32> = swizzle %14, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-    %17:vec3<u32> = load %res
-    ret %17
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %19:vec3<u32> = call %textureDimensions_e5a203
-    %20:void = %prevent_dce.Store3 0u, %19
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %22:vec3<u32> = call %textureDimensions_e5a203
-    %23:void = %prevent_dce.Store3 0u, %22
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %26:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %26, vec4<f32>(0.0f)
-    %27:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %28:vec3<u32> = call %textureDimensions_e5a203
-    store %27, %28
-    %29:VertexOutput = load %tint_symbol
-    ret %29
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %31:VertexOutput = call %vertex_main_inner
-    %32:vec4<f32> = access %31, 0u
-    %33:vec3<u32> = access %31, 1u
-    %34:vertex_main_outputs = construct %33, %32
-    ret %34
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/textureDimensions/e5a203.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/textureDimensions/e5a203.wgsl.expected.ir.fxc.hlsl
index c0a859d..b275bf3 100644
--- a/test/tint/builtins/gen/var/textureDimensions/e5a203.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/var/textureDimensions/e5a203.wgsl.expected.ir.fxc.hlsl
@@ -1,119 +1,47 @@
-SKIP: FAILED
+struct VertexOutput {
+  float4 pos;
+  uint3 prevent_dce;
+};
+
+struct vertex_main_outputs {
+  nointerpolation uint3 VertexOutput_prevent_dce : TEXCOORD0;
+  float4 VertexOutput_pos : SV_Position;
+};
 
 
-@group(0) @binding(0) var<storage, read_write> prevent_dce : vec3<u32>;
-
-@group(1) @binding(0) var arg_0 : texture_3d<u32>;
-
-fn textureDimensions_e5a203() -> vec3<u32> {
-  var arg_1 = 1i;
-  var res : vec3<u32> = textureDimensions(arg_0, arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+Texture3D<uint4> arg_0 : register(t0, space1);
+uint3 textureDimensions_e5a203() {
+  int arg_1 = int(1);
+  Texture3D<uint4> v = arg_0;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(arg_1), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 res = v_1.xyz;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = textureDimensions_e5a203();
+void fragment_main() {
+  prevent_dce.Store3(0u, textureDimensions_e5a203());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = textureDimensions_e5a203();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store3(0u, textureDimensions_e5a203());
 }
 
-struct VertexOutput {
-  @builtin(position)
-  pos : vec4<f32>,
-  @location(0) @interpolate(flat)
-  prevent_dce : vec3<u32>,
-}
-
-@vertex
-fn vertex_main() -> VertexOutput {
-  var tint_symbol : VertexOutput;
-  tint_symbol.pos = vec4<f32>();
+VertexOutput vertex_main_inner() {
+  VertexOutput tint_symbol = (VertexOutput)0;
+  tint_symbol.pos = (0.0f).xxxx;
   tint_symbol.prevent_dce = textureDimensions_e5a203();
-  return tint_symbol;
+  VertexOutput v_2 = tint_symbol;
+  return v_2;
 }
 
-Failed to generate: :30:49 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-                                                ^^^
-
-:17:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-VertexOutput = struct @align(16) {
-  pos:vec4<f32> @offset(0)
-  prevent_dce:vec3<u32> @offset(16)
+vertex_main_outputs vertex_main() {
+  VertexOutput v_3 = vertex_main_inner();
+  VertexOutput v_4 = v_3;
+  VertexOutput v_5 = v_3;
+  vertex_main_outputs v_6 = {v_5.prevent_dce, v_4.pos};
+  return v_6;
 }
 
-vertex_main_outputs = struct @align(16) {
-  VertexOutput_prevent_dce:vec3<u32> @offset(0), @location(0), @interpolate(flat)
-  VertexOutput_pos:vec4<f32> @offset(16), @builtin(position)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %arg_0:ptr<handle, texture_3d<u32>, read> = var @binding_point(1, 0)
-}
-
-%textureDimensions_e5a203 = func():vec3<u32> {
-  $B2: {
-    %arg_1:ptr<function, i32, read_write> = var, 1i
-    %5:texture_3d<u32> = load %arg_0
-    %6:i32 = load %arg_1
-    %7:u32 = convert %6
-    %8:ptr<function, vec4<u32>, read_write> = var
-    %9:ptr<function, u32, read_write> = access %8, 0u
-    %10:ptr<function, u32, read_write> = access %8, 1u
-    %11:ptr<function, u32, read_write> = access %8, 2u
-    %12:ptr<function, u32, read_write> = access %8, 3u
-    %13:void = %5.GetDimensions %7, %9, %10, %11, %12
-    %14:vec4<u32> = load %8
-    %15:vec2<u32> = swizzle %14, xyz
-    %res:ptr<function, vec3<u32>, read_write> = var, %15
-    %17:vec3<u32> = load %res
-    ret %17
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %19:vec3<u32> = call %textureDimensions_e5a203
-    %20:void = %prevent_dce.Store3 0u, %19
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %22:vec3<u32> = call %textureDimensions_e5a203
-    %23:void = %prevent_dce.Store3 0u, %22
-    ret
-  }
-}
-%vertex_main_inner = func():VertexOutput {
-  $B5: {
-    %tint_symbol:ptr<function, VertexOutput, read_write> = var
-    %26:ptr<function, vec4<f32>, read_write> = access %tint_symbol, 0u
-    store %26, vec4<f32>(0.0f)
-    %27:ptr<function, vec3<u32>, read_write> = access %tint_symbol, 1u
-    %28:vec3<u32> = call %textureDimensions_e5a203
-    store %27, %28
-    %29:VertexOutput = load %tint_symbol
-    ret %29
-  }
-}
-%vertex_main = @vertex func():vertex_main_outputs {
-  $B6: {
-    %31:VertexOutput = call %vertex_main_inner
-    %32:vec4<f32> = access %31, 0u
-    %33:vec3<u32> = access %31, 1u
-    %34:vertex_main_outputs = construct %33, %32
-    ret %34
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/types/texture/sampled/3d.wgsl.expected.ir.dxc.hlsl b/test/tint/types/texture/sampled/3d.wgsl.expected.ir.dxc.hlsl
index 9cf167f..a372253 100644
--- a/test/tint/types/texture/sampled/3d.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/types/texture/sampled/3d.wgsl.expected.ir.dxc.hlsl
@@ -1,88 +1,20 @@
-SKIP: FAILED
 
-
-@group(0) @binding(0) var t_f : texture_3d<f32>;
-
-@group(0) @binding(1) var t_i : texture_3d<i32>;
-
-@group(0) @binding(2) var t_u : texture_3d<u32>;
-
-@compute @workgroup_size(1)
-fn main() {
-  var fdims = textureDimensions(t_f, 1);
-  var idims = textureDimensions(t_i, 1);
-  var udims = textureDimensions(t_u, 1);
+Texture3D<float4> t_f : register(t0);
+Texture3D<int4> t_i : register(t1);
+Texture3D<uint4> t_u : register(t2);
+[numthreads(1, 1, 1)]
+void main() {
+  Texture3D<float4> v = t_f;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(int(1)), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 fdims = v_1.xyz;
+  Texture3D<int4> v_2 = t_i;
+  uint4 v_3 = (0u).xxxx;
+  v_2.GetDimensions(uint(int(1)), v_3[0u], v_3[1u], v_3[2u], v_3[3u]);
+  uint3 idims = v_3.xyz;
+  Texture3D<uint4> v_4 = t_u;
+  uint4 v_5 = (0u).xxxx;
+  v_4.GetDimensions(uint(int(1)), v_5[0u], v_5[1u], v_5[2u], v_5[3u]);
+  uint3 udims = v_5.xyz;
 }
 
-Failed to generate: :19:51 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %fdims:ptr<function, vec3<u32>, read_write> = var, %14
-                                                  ^^^
-
-:8:3 note: in block
-  $B2: {
-  ^^^
-
-:30:51 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %idims:ptr<function, vec3<u32>, read_write> = var, %25
-                                                  ^^^
-
-:8:3 note: in block
-  $B2: {
-  ^^^
-
-:41:51 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %udims:ptr<function, vec3<u32>, read_write> = var, %36
-                                                  ^^^
-
-:8:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-$B1: {  # root
-  %t_f:ptr<handle, texture_3d<f32>, read> = var @binding_point(0, 0)
-  %t_i:ptr<handle, texture_3d<i32>, read> = var @binding_point(0, 1)
-  %t_u:ptr<handle, texture_3d<u32>, read> = var @binding_point(0, 2)
-}
-
-%main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B2: {
-    %5:texture_3d<f32> = load %t_f
-    %6:u32 = convert 1i
-    %7:ptr<function, vec4<u32>, read_write> = var
-    %8:ptr<function, u32, read_write> = access %7, 0u
-    %9:ptr<function, u32, read_write> = access %7, 1u
-    %10:ptr<function, u32, read_write> = access %7, 2u
-    %11:ptr<function, u32, read_write> = access %7, 3u
-    %12:void = %5.GetDimensions %6, %8, %9, %10, %11
-    %13:vec4<u32> = load %7
-    %14:vec2<u32> = swizzle %13, xyz
-    %fdims:ptr<function, vec3<u32>, read_write> = var, %14
-    %16:texture_3d<i32> = load %t_i
-    %17:u32 = convert 1i
-    %18:ptr<function, vec4<u32>, read_write> = var
-    %19:ptr<function, u32, read_write> = access %18, 0u
-    %20:ptr<function, u32, read_write> = access %18, 1u
-    %21:ptr<function, u32, read_write> = access %18, 2u
-    %22:ptr<function, u32, read_write> = access %18, 3u
-    %23:void = %16.GetDimensions %17, %19, %20, %21, %22
-    %24:vec4<u32> = load %18
-    %25:vec2<u32> = swizzle %24, xyz
-    %idims:ptr<function, vec3<u32>, read_write> = var, %25
-    %27:texture_3d<u32> = load %t_u
-    %28:u32 = convert 1i
-    %29:ptr<function, vec4<u32>, read_write> = var
-    %30:ptr<function, u32, read_write> = access %29, 0u
-    %31:ptr<function, u32, read_write> = access %29, 1u
-    %32:ptr<function, u32, read_write> = access %29, 2u
-    %33:ptr<function, u32, read_write> = access %29, 3u
-    %34:void = %27.GetDimensions %28, %30, %31, %32, %33
-    %35:vec4<u32> = load %29
-    %36:vec2<u32> = swizzle %35, xyz
-    %udims:ptr<function, vec3<u32>, read_write> = var, %36
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/types/texture/sampled/3d.wgsl.expected.ir.fxc.hlsl b/test/tint/types/texture/sampled/3d.wgsl.expected.ir.fxc.hlsl
index 9cf167f..a372253 100644
--- a/test/tint/types/texture/sampled/3d.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/types/texture/sampled/3d.wgsl.expected.ir.fxc.hlsl
@@ -1,88 +1,20 @@
-SKIP: FAILED
 
-
-@group(0) @binding(0) var t_f : texture_3d<f32>;
-
-@group(0) @binding(1) var t_i : texture_3d<i32>;
-
-@group(0) @binding(2) var t_u : texture_3d<u32>;
-
-@compute @workgroup_size(1)
-fn main() {
-  var fdims = textureDimensions(t_f, 1);
-  var idims = textureDimensions(t_i, 1);
-  var udims = textureDimensions(t_u, 1);
+Texture3D<float4> t_f : register(t0);
+Texture3D<int4> t_i : register(t1);
+Texture3D<uint4> t_u : register(t2);
+[numthreads(1, 1, 1)]
+void main() {
+  Texture3D<float4> v = t_f;
+  uint4 v_1 = (0u).xxxx;
+  v.GetDimensions(uint(int(1)), v_1[0u], v_1[1u], v_1[2u], v_1[3u]);
+  uint3 fdims = v_1.xyz;
+  Texture3D<int4> v_2 = t_i;
+  uint4 v_3 = (0u).xxxx;
+  v_2.GetDimensions(uint(int(1)), v_3[0u], v_3[1u], v_3[2u], v_3[3u]);
+  uint3 idims = v_3.xyz;
+  Texture3D<uint4> v_4 = t_u;
+  uint4 v_5 = (0u).xxxx;
+  v_4.GetDimensions(uint(int(1)), v_5[0u], v_5[1u], v_5[2u], v_5[3u]);
+  uint3 udims = v_5.xyz;
 }
 
-Failed to generate: :19:51 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %fdims:ptr<function, vec3<u32>, read_write> = var, %14
-                                                  ^^^
-
-:8:3 note: in block
-  $B2: {
-  ^^^
-
-:30:51 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %idims:ptr<function, vec3<u32>, read_write> = var, %25
-                                                  ^^^
-
-:8:3 note: in block
-  $B2: {
-  ^^^
-
-:41:51 error: var: initializer type 'vec2<u32>' does not match store type 'vec3<u32>'
-    %udims:ptr<function, vec3<u32>, read_write> = var, %36
-                                                  ^^^
-
-:8:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-$B1: {  # root
-  %t_f:ptr<handle, texture_3d<f32>, read> = var @binding_point(0, 0)
-  %t_i:ptr<handle, texture_3d<i32>, read> = var @binding_point(0, 1)
-  %t_u:ptr<handle, texture_3d<u32>, read> = var @binding_point(0, 2)
-}
-
-%main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B2: {
-    %5:texture_3d<f32> = load %t_f
-    %6:u32 = convert 1i
-    %7:ptr<function, vec4<u32>, read_write> = var
-    %8:ptr<function, u32, read_write> = access %7, 0u
-    %9:ptr<function, u32, read_write> = access %7, 1u
-    %10:ptr<function, u32, read_write> = access %7, 2u
-    %11:ptr<function, u32, read_write> = access %7, 3u
-    %12:void = %5.GetDimensions %6, %8, %9, %10, %11
-    %13:vec4<u32> = load %7
-    %14:vec2<u32> = swizzle %13, xyz
-    %fdims:ptr<function, vec3<u32>, read_write> = var, %14
-    %16:texture_3d<i32> = load %t_i
-    %17:u32 = convert 1i
-    %18:ptr<function, vec4<u32>, read_write> = var
-    %19:ptr<function, u32, read_write> = access %18, 0u
-    %20:ptr<function, u32, read_write> = access %18, 1u
-    %21:ptr<function, u32, read_write> = access %18, 2u
-    %22:ptr<function, u32, read_write> = access %18, 3u
-    %23:void = %16.GetDimensions %17, %19, %20, %21, %22
-    %24:vec4<u32> = load %18
-    %25:vec2<u32> = swizzle %24, xyz
-    %idims:ptr<function, vec3<u32>, read_write> = var, %25
-    %27:texture_3d<u32> = load %t_u
-    %28:u32 = convert 1i
-    %29:ptr<function, vec4<u32>, read_write> = var
-    %30:ptr<function, u32, read_write> = access %29, 0u
-    %31:ptr<function, u32, read_write> = access %29, 1u
-    %32:ptr<function, u32, read_write> = access %29, 2u
-    %33:ptr<function, u32, read_write> = access %29, 3u
-    %34:void = %27.GetDimensions %28, %30, %31, %32, %33
-    %35:vec4<u32> = load %29
-    %36:vec2<u32> = swizzle %35, xyz
-    %udims:ptr<function, vec3<u32>, read_write> = var, %36
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1