Implement texture_depth_multisampled_2d

Implemented for all readers and writers.

Cleaned up some verbose code in sem::Function and the Inspector in the
process.

Fixed: tint:1032
Change-Id: Ia6f2f59e6d2e511c89160b97be990e8b7c9828d9
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59664
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>

test/BUILD.gn

@@ -162,6 +162,7 @@ tint_unittests_source_set("tint_unittests_core_src") {
     "../src/ast/call_statement_test.cc",
     "../src/ast/case_statement_test.cc",
     "../src/ast/continue_statement_test.cc",
+    "../src/ast/depth_multisampled_texture_test.cc",
     "../src/ast/depth_texture_test.cc",
     "../src/ast/discard_statement_test.cc",
     "../src/ast/else_statement_test.cc",
@@ -262,6 +263,7 @@ tint_unittests_source_set("tint_unittests_core_src") {
     "../src/scope_stack_test.cc",
     "../src/sem/atomic_type_test.cc",
     "../src/sem/bool_type_test.cc",
+    "../src/sem/depth_multisampled_texture_type_test.cc",
     "../src/sem/depth_texture_type_test.cc",
     "../src/sem/external_texture_type_test.cc",
     "../src/sem/f32_type_test.cc",

test/intrinsics/gen/ignore/2a6ac2.wgsl

@@ -0,0 +1,46 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(0)]] var arg_0: texture_depth_multisampled_2d;
+
+// fn ignore(texture_depth_multisampled_2d)
+fn ignore_2a6ac2() {
+  ignore(arg_0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_2a6ac2();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_2a6ac2();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  ignore_2a6ac2();
+}

test/intrinsics/gen/ignore/2a6ac2.wgsl.expected.hlsl

@@ -0,0 +1,26 @@
+Texture2DMS<float4> arg_0 : register(t0, space1);
+
+void ignore_2a6ac2() {
+  arg_0;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  ignore_2a6ac2();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  ignore_2a6ac2();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ignore_2a6ac2();
+  return;
+}

test/intrinsics/gen/ignore/2a6ac2.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void ignore_2a6ac2(depth2d_ms<float, access::read> tint_symbol_2) {
+  (void) tint_symbol_2;
+}
+
+vertex tint_symbol vertex_main(depth2d_ms<float, access::read> tint_symbol_3 [[texture(0)]]) {
+  ignore_2a6ac2(tint_symbol_3);
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main(depth2d_ms<float, access::read> tint_symbol_4 [[texture(0)]]) {
+  ignore_2a6ac2(tint_symbol_4);
+  return;
+}
+
+kernel void compute_main(depth2d_ms<float, access::read> tint_symbol_5 [[texture(0)]]) {
+  ignore_2a6ac2(tint_symbol_5);
+  return;
+}
+

test/intrinsics/gen/ignore/2a6ac2.wgsl.expected.spvasm

@@ -0,0 +1,68 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %arg_0 "arg_0"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %ignore_2a6ac2 "ignore_2a6ac2"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %tint_symbol_1 BuiltIn Position
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+          %7 = OpTypeImage %float 2D 1 0 1 1 Unknown
+%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
+      %arg_0 = OpVariable %_ptr_UniformConstant_7 UniformConstant
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %11 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %11
+       %void = OpTypeVoid
+         %12 = OpTypeFunction %void
+         %18 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%ignore_2a6ac2 = OpFunction %void None %12
+         %15 = OpLabel
+         %17 = OpLoad %7 %arg_0
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %18
+%tint_symbol = OpFunctionParameter %v4float
+         %21 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %12
+         %23 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %25 = OpFunctionCall %void %ignore_2a6ac2
+         %26 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ignore_2a6ac2
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ignore_2a6ac2
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/ignore/2a6ac2.wgsl.expected.wgsl

@@ -0,0 +1,21 @@
+[[group(1), binding(0)]] var arg_0 : texture_depth_multisampled_2d;
+
+fn ignore_2a6ac2() {
+  ignore(arg_0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_2a6ac2();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_2a6ac2();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  ignore_2a6ac2();
+}

test/intrinsics/gen/textureDimensions/f60bdb.wgsl

@@ -0,0 +1,46 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(0)]] var arg_0: texture_depth_multisampled_2d;
+
+// fn textureDimensions(texture: texture_depth_multisampled_2d) -> vec2<i32>
+fn textureDimensions_f60bdb() {
+  var res: vec2<i32> = textureDimensions(arg_0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureDimensions_f60bdb();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureDimensions_f60bdb();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureDimensions_f60bdb();
+}

test/intrinsics/gen/textureDimensions/f60bdb.wgsl.expected.hlsl

@@ -0,0 +1,28 @@
+Texture2DMS<float4> arg_0 : register(t0, space1);
+
+void textureDimensions_f60bdb() {
+  int3 tint_tmp;
+  arg_0.GetDimensions(tint_tmp.x, tint_tmp.y, tint_tmp.z);
+  int2 res = tint_tmp.xy;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  textureDimensions_f60bdb();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  textureDimensions_f60bdb();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureDimensions_f60bdb();
+  return;
+}

test/intrinsics/gen/textureDimensions/f60bdb.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureDimensions_f60bdb(depth2d_ms<float, access::read> tint_symbol_2) {
+  int2 res = int2(tint_symbol_2.get_width(), tint_symbol_2.get_height());
+}
+
+vertex tint_symbol vertex_main(depth2d_ms<float, access::read> tint_symbol_3 [[texture(0)]]) {
+  textureDimensions_f60bdb(tint_symbol_3);
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main(depth2d_ms<float, access::read> tint_symbol_4 [[texture(0)]]) {
+  textureDimensions_f60bdb(tint_symbol_4);
+  return;
+}
+
+kernel void compute_main(depth2d_ms<float, access::read> tint_symbol_5 [[texture(0)]]) {
+  textureDimensions_f60bdb(tint_symbol_5);
+  return;
+}
+

test/intrinsics/gen/textureDimensions/f60bdb.wgsl.expected.spvasm

@@ -0,0 +1,77 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 38
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %arg_0 "arg_0"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %textureDimensions_f60bdb "textureDimensions_f60bdb"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %tint_symbol_1 BuiltIn Position
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+          %7 = OpTypeImage %float 2D 1 0 1 1 Unknown
+%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
+      %arg_0 = OpVariable %_ptr_UniformConstant_7 UniformConstant
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %11 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %11
+       %void = OpTypeVoid
+         %12 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+         %22 = OpConstantNull %v2int
+         %23 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%textureDimensions_f60bdb = OpFunction %void None %12
+         %15 = OpLabel
+        %res = OpVariable %_ptr_Function_v2int Function %22
+         %19 = OpLoad %7 %arg_0
+         %16 = OpImageQuerySize %v2int %19
+               OpStore %res %16
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %23
+%tint_symbol = OpFunctionParameter %v4float
+         %26 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %12
+         %28 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %30 = OpFunctionCall %void %textureDimensions_f60bdb
+         %31 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %33 = OpLabel
+         %34 = OpFunctionCall %void %textureDimensions_f60bdb
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %textureDimensions_f60bdb
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureDimensions/f60bdb.wgsl.expected.wgsl

@@ -0,0 +1,21 @@
+[[group(1), binding(0)]] var arg_0 : texture_depth_multisampled_2d;
+
+fn textureDimensions_f60bdb() {
+  var res : vec2<i32> = textureDimensions(arg_0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureDimensions_f60bdb();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureDimensions_f60bdb();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureDimensions_f60bdb();
+}

test/intrinsics/gen/textureLoad/6273b1.wgsl

@@ -0,0 +1,46 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(0)]] var arg_0: texture_depth_multisampled_2d;
+
+// fn textureLoad(texture: texture_depth_multisampled_2d, coords: vec2<i32>, sample_index: i32) -> f32
+fn textureLoad_6273b1() {
+  var res: f32 = textureLoad(arg_0, vec2<i32>(), 1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureLoad_6273b1();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureLoad_6273b1();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureLoad_6273b1();
+}

test/intrinsics/gen/textureLoad/6273b1.wgsl.expected.hlsl

@@ -0,0 +1,26 @@
+Texture2DMS<float4> arg_0 : register(t0, space1);
+
+void textureLoad_6273b1() {
+  float res = arg_0.Load(int3(0, 0, 0), 1).x;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  textureLoad_6273b1();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  textureLoad_6273b1();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureLoad_6273b1();
+  return;
+}

test/intrinsics/gen/textureLoad/6273b1.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureLoad_6273b1(depth2d_ms<float, access::read> tint_symbol_2) {
+  float res = tint_symbol_2.read(uint2(int2()), 1);
+}
+
+vertex tint_symbol vertex_main(depth2d_ms<float, access::read> tint_symbol_3 [[texture(0)]]) {
+  textureLoad_6273b1(tint_symbol_3);
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main(depth2d_ms<float, access::read> tint_symbol_4 [[texture(0)]]) {
+  textureLoad_6273b1(tint_symbol_4);
+  return;
+}
+
+kernel void compute_main(depth2d_ms<float, access::read> tint_symbol_5 [[texture(0)]]) {
+  textureLoad_6273b1(tint_symbol_5);
+  return;
+}
+

test/intrinsics/gen/textureLoad/6273b1.wgsl.expected.spvasm

@@ -0,0 +1,78 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 40
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %arg_0 "arg_0"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %textureLoad_6273b1 "textureLoad_6273b1"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %tint_symbol_1 BuiltIn Position
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+          %7 = OpTypeImage %float 2D 1 0 1 1 Unknown
+%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
+      %arg_0 = OpVariable %_ptr_UniformConstant_7 UniformConstant
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %11 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %11
+       %void = OpTypeVoid
+         %12 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+         %21 = OpConstantNull %v2int
+      %int_1 = OpConstant %int 1
+%_ptr_Function_float = OpTypePointer Function %float
+         %25 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%textureLoad_6273b1 = OpFunction %void None %12
+         %15 = OpLabel
+        %res = OpVariable %_ptr_Function_float Function %4
+         %18 = OpLoad %7 %arg_0
+         %17 = OpImageFetch %v4float %18 %21 Sample %int_1
+         %16 = OpCompositeExtract %float %17 0
+               OpStore %res %16
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %25
+%tint_symbol = OpFunctionParameter %v4float
+         %28 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %12
+         %30 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %32 = OpFunctionCall %void %textureLoad_6273b1
+         %33 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %textureLoad_6273b1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %38 = OpLabel
+         %39 = OpFunctionCall %void %textureLoad_6273b1
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureLoad/6273b1.wgsl.expected.wgsl

@@ -0,0 +1,21 @@
+[[group(1), binding(0)]] var arg_0 : texture_depth_multisampled_2d;
+
+fn textureLoad_6273b1() {
+  var res : f32 = textureLoad(arg_0, vec2<i32>(), 1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureLoad_6273b1();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureLoad_6273b1();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureLoad_6273b1();
+}

test/intrinsics/gen/textureNumSamples/a3c8a0.wgsl

@@ -0,0 +1,46 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[group(1), binding(0)]] var arg_0: texture_depth_multisampled_2d;
+
+// fn textureNumSamples(texture: texture_depth_multisampled_2d) -> i32
+fn textureNumSamples_a3c8a0() {
+  var res: i32 = textureNumSamples(arg_0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureNumSamples_a3c8a0();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureNumSamples_a3c8a0();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureNumSamples_a3c8a0();
+}

test/intrinsics/gen/textureNumSamples/a3c8a0.wgsl.expected.hlsl

@@ -0,0 +1,28 @@
+Texture2DMS<float4> arg_0 : register(t0, space1);
+
+void textureNumSamples_a3c8a0() {
+  int3 tint_tmp;
+  arg_0.GetDimensions(tint_tmp.x, tint_tmp.y, tint_tmp.z);
+  int res = tint_tmp.z;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  textureNumSamples_a3c8a0();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  textureNumSamples_a3c8a0();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  textureNumSamples_a3c8a0();
+  return;
+}

test/intrinsics/gen/textureNumSamples/a3c8a0.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void textureNumSamples_a3c8a0(depth2d_ms<float, access::read> tint_symbol_2) {
+  int res = int(tint_symbol_2.get_num_samples());
+}
+
+vertex tint_symbol vertex_main(depth2d_ms<float, access::read> tint_symbol_3 [[texture(0)]]) {
+  textureNumSamples_a3c8a0(tint_symbol_3);
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main(depth2d_ms<float, access::read> tint_symbol_4 [[texture(0)]]) {
+  textureNumSamples_a3c8a0(tint_symbol_4);
+  return;
+}
+
+kernel void compute_main(depth2d_ms<float, access::read> tint_symbol_5 [[texture(0)]]) {
+  textureNumSamples_a3c8a0(tint_symbol_5);
+  return;
+}
+

test/intrinsics/gen/textureNumSamples/a3c8a0.wgsl.expected.spvasm

@@ -0,0 +1,76 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 37
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %arg_0 "arg_0"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %textureNumSamples_a3c8a0 "textureNumSamples_a3c8a0"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %tint_symbol_1 BuiltIn Position
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+          %7 = OpTypeImage %float 2D 1 0 1 1 Unknown
+%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
+      %arg_0 = OpVariable %_ptr_UniformConstant_7 UniformConstant
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %11 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %11
+       %void = OpTypeVoid
+         %12 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %21 = OpConstantNull %int
+         %22 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%textureNumSamples_a3c8a0 = OpFunction %void None %12
+         %15 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %21
+         %18 = OpLoad %7 %arg_0
+         %16 = OpImageQuerySamples %int %18
+               OpStore %res %16
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %22
+%tint_symbol = OpFunctionParameter %v4float
+         %25 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %12
+         %27 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %29 = OpFunctionCall %void %textureNumSamples_a3c8a0
+         %30 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %textureNumSamples_a3c8a0
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %textureNumSamples_a3c8a0
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/textureNumSamples/a3c8a0.wgsl.expected.wgsl

@@ -0,0 +1,21 @@
+[[group(1), binding(0)]] var arg_0 : texture_depth_multisampled_2d;
+
+fn textureNumSamples_a3c8a0() {
+  var res : i32 = textureNumSamples(arg_0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  textureNumSamples_a3c8a0();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  textureNumSamples_a3c8a0();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  textureNumSamples_a3c8a0();
+}

test/intrinsics/textureDimensions/depth_ms.spvasm

@@ -0,0 +1,77 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 38
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %arg_0 "arg_0"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %textureDimensions_f60bdb "textureDimensions_f60bdb"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %tint_symbol_1 BuiltIn Position
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+          %7 = OpTypeImage %float 2D 1 0 1 1 Unknown
+%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
+      %arg_0 = OpVariable %_ptr_UniformConstant_7 UniformConstant
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %11 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %11
+       %void = OpTypeVoid
+         %12 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+         %22 = OpConstantNull %v2int
+         %23 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%textureDimensions_f60bdb = OpFunction %void None %12
+         %15 = OpLabel
+        %res = OpVariable %_ptr_Function_v2int Function %22
+         %19 = OpLoad %7 %arg_0
+         %16 = OpImageQuerySize %v2int %19
+               OpStore %res %16
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %23
+%tint_symbol = OpFunctionParameter %v4float
+         %26 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %12
+         %28 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %30 = OpFunctionCall %void %textureDimensions_f60bdb
+         %31 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %33 = OpLabel
+         %34 = OpFunctionCall %void %textureDimensions_f60bdb
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %textureDimensions_f60bdb
+               OpReturn
+               OpFunctionEnd

test/intrinsics/textureDimensions/depth_ms.spvasm.expected.hlsl

@@ -0,0 +1,57 @@
+Texture2DMS<float4> arg_0 : register(t0, space1);
+static float4 tint_symbol_1 = float4(0.0f, 0.0f, 0.0f, 0.0f);
+
+void textureDimensions_f60bdb() {
+  int2 res = int2(0, 0);
+  int3 tint_tmp;
+  arg_0.GetDimensions(tint_tmp.x, tint_tmp.y, tint_tmp.z);
+  const int2 x_16 = int2(tint_tmp.xy);
+  res = x_16;
+  return;
+}
+
+void tint_symbol_2(float4 tint_symbol) {
+  tint_symbol_1 = tint_symbol;
+  return;
+}
+
+void vertex_main_1() {
+  textureDimensions_f60bdb();
+  tint_symbol_2(float4(0.0f, 0.0f, 0.0f, 0.0f));
+  return;
+}
+
+struct vertex_main_out {
+  float4 tint_symbol_1_1;
+};
+struct tint_symbol_3 {
+  float4 tint_symbol_1_1 : SV_Position;
+};
+
+tint_symbol_3 vertex_main() {
+  vertex_main_1();
+  const vertex_main_out tint_symbol_4 = {tint_symbol_1};
+  const tint_symbol_3 tint_symbol_5 = {tint_symbol_4.tint_symbol_1_1};
+  return tint_symbol_5;
+}
+
+void fragment_main_1() {
+  textureDimensions_f60bdb();
+  return;
+}
+
+void fragment_main() {
+  fragment_main_1();
+  return;
+}
+
+void compute_main_1() {
+  textureDimensions_f60bdb();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  compute_main_1();
+  return;
+}

test/intrinsics/textureDimensions/depth_ms.spvasm.expected.msl

@@ -0,0 +1,56 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct vertex_main_out {
+  float4 tint_symbol_1_1;
+};
+struct tint_symbol_3 {
+  float4 tint_symbol_1_1 [[position]];
+};
+
+void textureDimensions_f60bdb(depth2d_ms<float, access::read> tint_symbol_6) {
+  int2 res = int2(0, 0);
+  int2 const x_16 = int2(int2(tint_symbol_6.get_width(), tint_symbol_6.get_height()));
+  res = x_16;
+  return;
+}
+
+void tint_symbol_2(float4 tint_symbol, thread float4* const tint_symbol_7) {
+  *(tint_symbol_7) = tint_symbol;
+  return;
+}
+
+void vertex_main_1(depth2d_ms<float, access::read> tint_symbol_8, thread float4* const tint_symbol_9) {
+  textureDimensions_f60bdb(tint_symbol_8);
+  tint_symbol_2(float4(0.0f, 0.0f, 0.0f, 0.0f), tint_symbol_9);
+  return;
+}
+
+vertex tint_symbol_3 vertex_main(depth2d_ms<float, access::read> tint_symbol_10 [[texture(0)]]) {
+  thread float4 tint_symbol_11 = float4(0.0f, 0.0f, 0.0f, 0.0f);
+  vertex_main_1(tint_symbol_10, &(tint_symbol_11));
+  vertex_main_out const tint_symbol_4 = {.tint_symbol_1_1=tint_symbol_11};
+  tint_symbol_3 const tint_symbol_5 = {.tint_symbol_1_1=tint_symbol_4.tint_symbol_1_1};
+  return tint_symbol_5;
+}
+
+void fragment_main_1(depth2d_ms<float, access::read> tint_symbol_12) {
+  textureDimensions_f60bdb(tint_symbol_12);
+  return;
+}
+
+fragment void fragment_main(depth2d_ms<float, access::read> tint_symbol_13 [[texture(0)]]) {
+  fragment_main_1(tint_symbol_13);
+  return;
+}
+
+void compute_main_1(depth2d_ms<float, access::read> tint_symbol_14) {
+  textureDimensions_f60bdb(tint_symbol_14);
+  return;
+}
+
+kernel void compute_main(depth2d_ms<float, access::read> tint_symbol_15 [[texture(0)]]) {
+  compute_main_1(tint_symbol_15);
+  return;
+}
+

test/intrinsics/textureDimensions/depth_ms.spvasm.expected.spvasm

@@ -0,0 +1,120 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 63
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_4
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %arg_0 "arg_0"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %tint_symbol_4 "tint_symbol_4"
+               OpName %textureDimensions_f60bdb "textureDimensions_f60bdb"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main_1 "vertex_main_1"
+               OpName %vertex_main_out "vertex_main_out"
+               OpMemberName %vertex_main_out 0 "tint_symbol_1_1"
+               OpName %tint_symbol_5 "tint_symbol_5"
+               OpName %tint_symbol_3 "tint_symbol_3"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main_1 "fragment_main_1"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %tint_symbol_4 BuiltIn Position
+               OpMemberDecorate %vertex_main_out 0 Offset 0
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+          %7 = OpTypeImage %float 2D 1 0 1 1 Unknown
+%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
+      %arg_0 = OpVariable %_ptr_UniformConstant_7 UniformConstant
+    %v4float = OpTypeVector %float 4
+    %float_0 = OpConstant %float 0
+         %10 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_0
+%_ptr_Private_v4float = OpTypePointer Private %v4float
+%tint_symbol_1 = OpVariable %_ptr_Private_v4float Private %10
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %15 = OpConstantNull %v4float
+%tint_symbol_4 = OpVariable %_ptr_Output_v4float Output %15
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+      %int_0 = OpConstant %int 0
+         %23 = OpConstantComposite %v2int %int_0 %int_0
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+         %26 = OpConstantNull %v2int
+         %30 = OpTypeFunction %void %v4float
+%vertex_main_out = OpTypeStruct %v4float
+         %38 = OpTypeFunction %void %vertex_main_out
+    %float_1 = OpConstant %float 1
+%textureDimensions_f60bdb = OpFunction %void None %16
+         %19 = OpLabel
+        %res = OpVariable %_ptr_Function_v2int Function %26
+               OpStore %res %23
+         %29 = OpLoad %7 %arg_0
+         %28 = OpImageQuerySize %v2int %29
+               OpStore %res %28
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %30
+%tint_symbol = OpFunctionParameter %v4float
+         %33 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main_1 = OpFunction %void None %16
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %textureDimensions_f60bdb
+         %37 = OpFunctionCall %void %tint_symbol_2 %10
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_5 = OpFunction %void None %38
+%tint_symbol_3 = OpFunctionParameter %vertex_main_out
+         %42 = OpLabel
+         %43 = OpCompositeExtract %v4float %tint_symbol_3 0
+               OpStore %tint_symbol_4 %43
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %16
+         %45 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %47 = OpFunctionCall %void %vertex_main_1
+         %49 = OpLoad %v4float %tint_symbol_1
+         %50 = OpCompositeConstruct %vertex_main_out %49
+         %48 = OpFunctionCall %void %tint_symbol_5 %50
+               OpReturn
+               OpFunctionEnd
+%fragment_main_1 = OpFunction %void None %16
+         %52 = OpLabel
+         %53 = OpFunctionCall %void %textureDimensions_f60bdb
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %16
+         %55 = OpLabel
+         %56 = OpFunctionCall %void %fragment_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %16
+         %58 = OpLabel
+         %59 = OpFunctionCall %void %textureDimensions_f60bdb
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %16
+         %61 = OpLabel
+         %62 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd

test/intrinsics/textureDimensions/depth_ms.spvasm.expected.wgsl

@@ -0,0 +1,52 @@
+[[group(1), binding(0)]] var arg_0 : texture_depth_multisampled_2d;
+
+var<private> tint_symbol_1 : vec4<f32> = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+
+fn textureDimensions_f60bdb() {
+  var res : vec2<i32> = vec2<i32>(0, 0);
+  let x_16 : vec2<i32> = vec2<i32>(textureDimensions(arg_0));
+  res = x_16;
+  return;
+}
+
+fn tint_symbol_2(tint_symbol : vec4<f32>) {
+  tint_symbol_1 = tint_symbol;
+  return;
+}
+
+fn vertex_main_1() {
+  textureDimensions_f60bdb();
+  tint_symbol_2(vec4<f32>(0.0, 0.0, 0.0, 0.0));
+  return;
+}
+
+struct vertex_main_out {
+  [[builtin(position)]]
+  tint_symbol_1_1 : vec4<f32>;
+};
+
+[[stage(vertex)]]
+fn vertex_main() -> vertex_main_out {
+  vertex_main_1();
+  return vertex_main_out(tint_symbol_1);
+}
+
+fn fragment_main_1() {
+  textureDimensions_f60bdb();
+  return;
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  fragment_main_1();
+}
+
+fn compute_main_1() {
+  textureDimensions_f60bdb();
+  return;
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn compute_main() {
+  compute_main_1();
+}

test/intrinsics/textureLoad/depth_ms.spvasm

@@ -0,0 +1,78 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 40
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %arg_0 "arg_0"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %textureLoad_6273b1 "textureLoad_6273b1"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %tint_symbol_1 BuiltIn Position
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+          %7 = OpTypeImage %float 2D 1 0 1 1 Unknown
+%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
+      %arg_0 = OpVariable %_ptr_UniformConstant_7 UniformConstant
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %11 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %11
+       %void = OpTypeVoid
+         %12 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+         %21 = OpConstantNull %v2int
+      %int_1 = OpConstant %int 1
+%_ptr_Function_float = OpTypePointer Function %float
+         %25 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%textureLoad_6273b1 = OpFunction %void None %12
+         %15 = OpLabel
+        %res = OpVariable %_ptr_Function_float Function %4
+         %18 = OpLoad %7 %arg_0
+         %17 = OpImageFetch %v4float %18 %21 Sample %int_1
+         %16 = OpCompositeExtract %float %17 0
+               OpStore %res %16
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %25
+%tint_symbol = OpFunctionParameter %v4float
+         %28 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %12
+         %30 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %32 = OpFunctionCall %void %textureLoad_6273b1
+         %33 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %textureLoad_6273b1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %38 = OpLabel
+         %39 = OpFunctionCall %void %textureLoad_6273b1
+               OpReturn
+               OpFunctionEnd

test/intrinsics/textureLoad/depth_ms.spvasm.expected.hlsl

@@ -0,0 +1,55 @@
+Texture2DMS<float4> arg_0 : register(t0, space1);
+static float4 tint_symbol_1 = float4(0.0f, 0.0f, 0.0f, 0.0f);
+
+void textureLoad_6273b1() {
+  float res = 0.0f;
+  const float4 x_17 = float4(arg_0.Load(int3(0, 0, 0), 1).x, 0.0f, 0.0f, 0.0f);
+  res = x_17.x;
+  return;
+}
+
+void tint_symbol_2(float4 tint_symbol) {
+  tint_symbol_1 = tint_symbol;
+  return;
+}
+
+void vertex_main_1() {
+  textureLoad_6273b1();
+  tint_symbol_2(float4(0.0f, 0.0f, 0.0f, 0.0f));
+  return;
+}
+
+struct vertex_main_out {
+  float4 tint_symbol_1_1;
+};
+struct tint_symbol_3 {
+  float4 tint_symbol_1_1 : SV_Position;
+};
+
+tint_symbol_3 vertex_main() {
+  vertex_main_1();
+  const vertex_main_out tint_symbol_4 = {tint_symbol_1};
+  const tint_symbol_3 tint_symbol_5 = {tint_symbol_4.tint_symbol_1_1};
+  return tint_symbol_5;
+}
+
+void fragment_main_1() {
+  textureLoad_6273b1();
+  return;
+}
+
+void fragment_main() {
+  fragment_main_1();
+  return;
+}
+
+void compute_main_1() {
+  textureLoad_6273b1();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  compute_main_1();
+  return;
+}

test/intrinsics/textureLoad/depth_ms.spvasm.expected.msl

@@ -0,0 +1,56 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct vertex_main_out {
+  float4 tint_symbol_1_1;
+};
+struct tint_symbol_3 {
+  float4 tint_symbol_1_1 [[position]];
+};
+
+void textureLoad_6273b1(depth2d_ms<float, access::read> tint_symbol_6) {
+  float res = 0.0f;
+  float4 const x_17 = float4(tint_symbol_6.read(uint2(int2(0, 0)), 1), 0.0f, 0.0f, 0.0f);
+  res = x_17.x;
+  return;
+}
+
+void tint_symbol_2(float4 tint_symbol, thread float4* const tint_symbol_7) {
+  *(tint_symbol_7) = tint_symbol;
+  return;
+}
+
+void vertex_main_1(depth2d_ms<float, access::read> tint_symbol_8, thread float4* const tint_symbol_9) {
+  textureLoad_6273b1(tint_symbol_8);
+  tint_symbol_2(float4(0.0f, 0.0f, 0.0f, 0.0f), tint_symbol_9);
+  return;
+}
+
+vertex tint_symbol_3 vertex_main(depth2d_ms<float, access::read> tint_symbol_10 [[texture(0)]]) {
+  thread float4 tint_symbol_11 = float4(0.0f, 0.0f, 0.0f, 0.0f);
+  vertex_main_1(tint_symbol_10, &(tint_symbol_11));
+  vertex_main_out const tint_symbol_4 = {.tint_symbol_1_1=tint_symbol_11};
+  tint_symbol_3 const tint_symbol_5 = {.tint_symbol_1_1=tint_symbol_4.tint_symbol_1_1};
+  return tint_symbol_5;
+}
+
+void fragment_main_1(depth2d_ms<float, access::read> tint_symbol_12) {
+  textureLoad_6273b1(tint_symbol_12);
+  return;
+}
+
+fragment void fragment_main(depth2d_ms<float, access::read> tint_symbol_13 [[texture(0)]]) {
+  fragment_main_1(tint_symbol_13);
+  return;
+}
+
+void compute_main_1(depth2d_ms<float, access::read> tint_symbol_14) {
+  textureLoad_6273b1(tint_symbol_14);
+  return;
+}
+
+kernel void compute_main(depth2d_ms<float, access::read> tint_symbol_15 [[texture(0)]]) {
+  compute_main_1(tint_symbol_15);
+  return;
+}
+

test/intrinsics/textureLoad/depth_ms.spvasm.expected.spvasm

@@ -0,0 +1,122 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 65
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_4
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %arg_0 "arg_0"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %tint_symbol_4 "tint_symbol_4"
+               OpName %textureLoad_6273b1 "textureLoad_6273b1"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main_1 "vertex_main_1"
+               OpName %vertex_main_out "vertex_main_out"
+               OpMemberName %vertex_main_out 0 "tint_symbol_1_1"
+               OpName %tint_symbol_5 "tint_symbol_5"
+               OpName %tint_symbol_3 "tint_symbol_3"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main_1 "fragment_main_1"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %tint_symbol_4 BuiltIn Position
+               OpMemberDecorate %vertex_main_out 0 Offset 0
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+          %7 = OpTypeImage %float 2D 1 0 1 1 Unknown
+%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
+      %arg_0 = OpVariable %_ptr_UniformConstant_7 UniformConstant
+    %v4float = OpTypeVector %float 4
+    %float_0 = OpConstant %float 0
+         %10 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_0
+%_ptr_Private_v4float = OpTypePointer Private %v4float
+%tint_symbol_1 = OpVariable %_ptr_Private_v4float Private %10
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %15 = OpConstantNull %v4float
+%tint_symbol_4 = OpVariable %_ptr_Output_v4float Output %15
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+%_ptr_Function_float = OpTypePointer Function %float
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+      %int_0 = OpConstant %int 0
+         %28 = OpConstantComposite %v2int %int_0 %int_0
+      %int_1 = OpConstant %int 1
+         %32 = OpTypeFunction %void %v4float
+%vertex_main_out = OpTypeStruct %v4float
+         %40 = OpTypeFunction %void %vertex_main_out
+    %float_1 = OpConstant %float 1
+%textureLoad_6273b1 = OpFunction %void None %16
+         %19 = OpLabel
+        %res = OpVariable %_ptr_Function_float Function %4
+               OpStore %res %float_0
+         %24 = OpLoad %7 %arg_0
+         %23 = OpImageFetch %v4float %24 %28 Sample %int_1
+         %22 = OpCompositeExtract %float %23 0
+         %30 = OpCompositeConstruct %v4float %22 %float_0 %float_0 %float_0
+         %31 = OpCompositeExtract %float %30 0
+               OpStore %res %31
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %32
+%tint_symbol = OpFunctionParameter %v4float
+         %35 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main_1 = OpFunction %void None %16
+         %37 = OpLabel
+         %38 = OpFunctionCall %void %textureLoad_6273b1
+         %39 = OpFunctionCall %void %tint_symbol_2 %10
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_5 = OpFunction %void None %40
+%tint_symbol_3 = OpFunctionParameter %vertex_main_out
+         %44 = OpLabel
+         %45 = OpCompositeExtract %v4float %tint_symbol_3 0
+               OpStore %tint_symbol_4 %45
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %16
+         %47 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %49 = OpFunctionCall %void %vertex_main_1
+         %51 = OpLoad %v4float %tint_symbol_1
+         %52 = OpCompositeConstruct %vertex_main_out %51
+         %50 = OpFunctionCall %void %tint_symbol_5 %52
+               OpReturn
+               OpFunctionEnd
+%fragment_main_1 = OpFunction %void None %16
+         %54 = OpLabel
+         %55 = OpFunctionCall %void %textureLoad_6273b1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %16
+         %57 = OpLabel
+         %58 = OpFunctionCall %void %fragment_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %16
+         %60 = OpLabel
+         %61 = OpFunctionCall %void %textureLoad_6273b1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %16
+         %63 = OpLabel
+         %64 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd

test/intrinsics/textureLoad/depth_ms.spvasm.expected.wgsl

@@ -0,0 +1,52 @@
+[[group(1), binding(0)]] var arg_0 : texture_depth_multisampled_2d;
+
+var<private> tint_symbol_1 : vec4<f32> = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+
+fn textureLoad_6273b1() {
+  var res : f32 = 0.0;
+  let x_17 : vec4<f32> = vec4<f32>(textureLoad(arg_0, vec2<i32>(0, 0), 1), 0.0, 0.0, 0.0);
+  res = x_17.x;
+  return;
+}
+
+fn tint_symbol_2(tint_symbol : vec4<f32>) {
+  tint_symbol_1 = tint_symbol;
+  return;
+}
+
+fn vertex_main_1() {
+  textureLoad_6273b1();
+  tint_symbol_2(vec4<f32>(0.0, 0.0, 0.0, 0.0));
+  return;
+}
+
+struct vertex_main_out {
+  [[builtin(position)]]
+  tint_symbol_1_1 : vec4<f32>;
+};
+
+[[stage(vertex)]]
+fn vertex_main() -> vertex_main_out {
+  vertex_main_1();
+  return vertex_main_out(tint_symbol_1);
+}
+
+fn fragment_main_1() {
+  textureLoad_6273b1();
+  return;
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  fragment_main_1();
+}
+
+fn compute_main_1() {
+  textureLoad_6273b1();
+  return;
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn compute_main() {
+  compute_main_1();
+}

test/intrinsics/textureNumSamples/depth_ms.spvasm

@@ -0,0 +1,76 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 37
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %arg_0 "arg_0"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %textureNumSamples_a3c8a0 "textureNumSamples_a3c8a0"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %tint_symbol_1 BuiltIn Position
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+          %7 = OpTypeImage %float 2D 1 0 1 1 Unknown
+%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
+      %arg_0 = OpVariable %_ptr_UniformConstant_7 UniformConstant
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %11 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %11
+       %void = OpTypeVoid
+         %12 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %21 = OpConstantNull %int
+         %22 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%textureNumSamples_a3c8a0 = OpFunction %void None %12
+         %15 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %21
+         %18 = OpLoad %7 %arg_0
+         %16 = OpImageQuerySamples %int %18
+               OpStore %res %16
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %22
+%tint_symbol = OpFunctionParameter %v4float
+         %25 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %12
+         %27 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %29 = OpFunctionCall %void %textureNumSamples_a3c8a0
+         %30 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %textureNumSamples_a3c8a0
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %textureNumSamples_a3c8a0
+               OpReturn
+               OpFunctionEnd

test/intrinsics/textureNumSamples/depth_ms.spvasm.expected.hlsl

@@ -0,0 +1,57 @@
+Texture2DMS<float4> arg_0 : register(t0, space1);
+static float4 tint_symbol_1 = float4(0.0f, 0.0f, 0.0f, 0.0f);
+
+void textureNumSamples_a3c8a0() {
+  int res = 0;
+  int3 tint_tmp;
+  arg_0.GetDimensions(tint_tmp.x, tint_tmp.y, tint_tmp.z);
+  const int x_16 = tint_tmp.z;
+  res = x_16;
+  return;
+}
+
+void tint_symbol_2(float4 tint_symbol) {
+  tint_symbol_1 = tint_symbol;
+  return;
+}
+
+void vertex_main_1() {
+  textureNumSamples_a3c8a0();
+  tint_symbol_2(float4(0.0f, 0.0f, 0.0f, 0.0f));
+  return;
+}
+
+struct vertex_main_out {
+  float4 tint_symbol_1_1;
+};
+struct tint_symbol_3 {
+  float4 tint_symbol_1_1 : SV_Position;
+};
+
+tint_symbol_3 vertex_main() {
+  vertex_main_1();
+  const vertex_main_out tint_symbol_4 = {tint_symbol_1};
+  const tint_symbol_3 tint_symbol_5 = {tint_symbol_4.tint_symbol_1_1};
+  return tint_symbol_5;
+}
+
+void fragment_main_1() {
+  textureNumSamples_a3c8a0();
+  return;
+}
+
+void fragment_main() {
+  fragment_main_1();
+  return;
+}
+
+void compute_main_1() {
+  textureNumSamples_a3c8a0();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  compute_main_1();
+  return;
+}

test/intrinsics/textureNumSamples/depth_ms.spvasm.expected.msl

@@ -0,0 +1,56 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct vertex_main_out {
+  float4 tint_symbol_1_1;
+};
+struct tint_symbol_3 {
+  float4 tint_symbol_1_1 [[position]];
+};
+
+void textureNumSamples_a3c8a0(depth2d_ms<float, access::read> tint_symbol_6) {
+  int res = 0;
+  int const x_16 = int(tint_symbol_6.get_num_samples());
+  res = x_16;
+  return;
+}
+
+void tint_symbol_2(float4 tint_symbol, thread float4* const tint_symbol_7) {
+  *(tint_symbol_7) = tint_symbol;
+  return;
+}
+
+void vertex_main_1(depth2d_ms<float, access::read> tint_symbol_8, thread float4* const tint_symbol_9) {
+  textureNumSamples_a3c8a0(tint_symbol_8);
+  tint_symbol_2(float4(0.0f, 0.0f, 0.0f, 0.0f), tint_symbol_9);
+  return;
+}
+
+vertex tint_symbol_3 vertex_main(depth2d_ms<float, access::read> tint_symbol_10 [[texture(0)]]) {
+  thread float4 tint_symbol_11 = float4(0.0f, 0.0f, 0.0f, 0.0f);
+  vertex_main_1(tint_symbol_10, &(tint_symbol_11));
+  vertex_main_out const tint_symbol_4 = {.tint_symbol_1_1=tint_symbol_11};
+  tint_symbol_3 const tint_symbol_5 = {.tint_symbol_1_1=tint_symbol_4.tint_symbol_1_1};
+  return tint_symbol_5;
+}
+
+void fragment_main_1(depth2d_ms<float, access::read> tint_symbol_12) {
+  textureNumSamples_a3c8a0(tint_symbol_12);
+  return;
+}
+
+fragment void fragment_main(depth2d_ms<float, access::read> tint_symbol_13 [[texture(0)]]) {
+  fragment_main_1(tint_symbol_13);
+  return;
+}
+
+void compute_main_1(depth2d_ms<float, access::read> tint_symbol_14) {
+  textureNumSamples_a3c8a0(tint_symbol_14);
+  return;
+}
+
+kernel void compute_main(depth2d_ms<float, access::read> tint_symbol_15 [[texture(0)]]) {
+  compute_main_1(tint_symbol_15);
+  return;
+}
+

test/intrinsics/textureNumSamples/depth_ms.spvasm.expected.spvasm

@@ -0,0 +1,118 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 60
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_4
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %arg_0 "arg_0"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %tint_symbol_4 "tint_symbol_4"
+               OpName %textureNumSamples_a3c8a0 "textureNumSamples_a3c8a0"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main_1 "vertex_main_1"
+               OpName %vertex_main_out "vertex_main_out"
+               OpMemberName %vertex_main_out 0 "tint_symbol_1_1"
+               OpName %tint_symbol_5 "tint_symbol_5"
+               OpName %tint_symbol_3 "tint_symbol_3"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main_1 "fragment_main_1"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %arg_0 DescriptorSet 1
+               OpDecorate %arg_0 Binding 0
+               OpDecorate %tint_symbol_4 BuiltIn Position
+               OpMemberDecorate %vertex_main_out 0 Offset 0
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+          %7 = OpTypeImage %float 2D 1 0 1 1 Unknown
+%_ptr_UniformConstant_7 = OpTypePointer UniformConstant %7
+      %arg_0 = OpVariable %_ptr_UniformConstant_7 UniformConstant
+    %v4float = OpTypeVector %float 4
+    %float_0 = OpConstant %float 0
+         %10 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_0
+%_ptr_Private_v4float = OpTypePointer Private %v4float
+%tint_symbol_1 = OpVariable %_ptr_Private_v4float Private %10
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %15 = OpConstantNull %v4float
+%tint_symbol_4 = OpVariable %_ptr_Output_v4float Output %15
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %24 = OpConstantNull %int
+         %27 = OpTypeFunction %void %v4float
+%vertex_main_out = OpTypeStruct %v4float
+         %35 = OpTypeFunction %void %vertex_main_out
+    %float_1 = OpConstant %float 1
+%textureNumSamples_a3c8a0 = OpFunction %void None %16
+         %19 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %24
+               OpStore %res %int_0
+         %26 = OpLoad %7 %arg_0
+         %25 = OpImageQuerySamples %int %26
+               OpStore %res %25
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %27
+%tint_symbol = OpFunctionParameter %v4float
+         %30 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main_1 = OpFunction %void None %16
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %textureNumSamples_a3c8a0
+         %34 = OpFunctionCall %void %tint_symbol_2 %10
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_5 = OpFunction %void None %35
+%tint_symbol_3 = OpFunctionParameter %vertex_main_out
+         %39 = OpLabel
+         %40 = OpCompositeExtract %v4float %tint_symbol_3 0
+               OpStore %tint_symbol_4 %40
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %16
+         %42 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %44 = OpFunctionCall %void %vertex_main_1
+         %46 = OpLoad %v4float %tint_symbol_1
+         %47 = OpCompositeConstruct %vertex_main_out %46
+         %45 = OpFunctionCall %void %tint_symbol_5 %47
+               OpReturn
+               OpFunctionEnd
+%fragment_main_1 = OpFunction %void None %16
+         %49 = OpLabel
+         %50 = OpFunctionCall %void %textureNumSamples_a3c8a0
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %16
+         %52 = OpLabel
+         %53 = OpFunctionCall %void %fragment_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %16
+         %55 = OpLabel
+         %56 = OpFunctionCall %void %textureNumSamples_a3c8a0
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %16
+         %58 = OpLabel
+         %59 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd

test/intrinsics/textureNumSamples/depth_ms.spvasm.expected.wgsl

@@ -0,0 +1,52 @@
+[[group(1), binding(0)]] var arg_0 : texture_depth_multisampled_2d;
+
+var<private> tint_symbol_1 : vec4<f32> = vec4<f32>(0.0, 0.0, 0.0, 0.0);
+
+fn textureNumSamples_a3c8a0() {
+  var res : i32 = 0;
+  let x_16 : i32 = textureNumSamples(arg_0);
+  res = x_16;
+  return;
+}
+
+fn tint_symbol_2(tint_symbol : vec4<f32>) {
+  tint_symbol_1 = tint_symbol;
+  return;
+}
+
+fn vertex_main_1() {
+  textureNumSamples_a3c8a0();
+  tint_symbol_2(vec4<f32>(0.0, 0.0, 0.0, 0.0));
+  return;
+}
+
+struct vertex_main_out {
+  [[builtin(position)]]
+  tint_symbol_1_1 : vec4<f32>;
+};
+
+[[stage(vertex)]]
+fn vertex_main() -> vertex_main_out {
+  vertex_main_1();
+  return vertex_main_out(tint_symbol_1);
+}
+
+fn fragment_main_1() {
+  textureNumSamples_a3c8a0();
+  return;
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  fragment_main_1();
+}
+
+fn compute_main_1() {
+  textureNumSamples_a3c8a0();
+  return;
+}
+
+[[stage(compute), workgroup_size(1, 1, 1)]]
+fn compute_main() {
+  compute_main_1();
+}