Implement ignore() intrinsic

Fixed: tint:871
Change-Id: I01e275686094611d67a54735593320bce16705c3
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53681
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>

src/intrinsics.def

@@ -313,6 +313,7 @@ fn fwidthCoarse(f32) -> f32
 fn fwidthCoarse<N: num>(vec<N, f32>) -> vec<N, f32>
 fn fwidthFine(f32) -> f32
 fn fwidthFine<N: num>(vec<N, f32>) -> vec<N, f32>
+fn ignore<T>(T)
 fn inverseSqrt(f32) -> f32
 fn inverseSqrt<N: num>(vec<N, f32>) -> vec<N, f32>
 fn isFinite(f32) -> bool

src/sem/intrinsic_type.cc

@@ -127,6 +127,9 @@ IntrinsicType ParseIntrinsicType(const std::string& name) {
   if (name == "fwidthFine") {
     return IntrinsicType::kFwidthFine;
   }
+  if (name == "ignore") {
+    return IntrinsicType::kIgnore;
+  }
   if (name == "inverseSqrt") {
     return IntrinsicType::kInverseSqrt;
   }
@@ -353,6 +356,8 @@ const char* str(IntrinsicType i) {
       return "fwidthCoarse";
     case IntrinsicType::kFwidthFine:
       return "fwidthFine";
+    case IntrinsicType::kIgnore:
+      return "ignore";
     case IntrinsicType::kInverseSqrt:
       return "inverseSqrt";
     case IntrinsicType::kIsFinite:

src/sem/intrinsic_type.h

@@ -66,6 +66,7 @@ enum class IntrinsicType {
   kFwidth,
   kFwidthCoarse,
   kFwidthFine,
+  kIgnore,
   kInverseSqrt,
   kIsFinite,
   kIsInf,

src/writer/hlsl/generator_impl.cc

@@ -591,6 +591,8 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
     } else if (intrinsic->Type() == sem::IntrinsicType::kIsNormal) {
       diagnostics_.add_error("is_normal not supported in HLSL backend yet");
       return false;
+    } else if (intrinsic->Type() == sem::IntrinsicType::kIgnore) {
+      return EmitExpression(pre, out, expr->params()[0]);
     } else if (intrinsic->IsDataPacking()) {
       return EmitDataPackingCall(pre, out, expr, intrinsic);
     } else if (intrinsic->IsDataUnpacking()) {

src/writer/hlsl/generator_impl_intrinsic_test.cc

@@ -474,6 +474,32 @@ void main() {
   Validate();
 }
 
+TEST_F(HlslGeneratorImplTest_Intrinsic, Ignore) {
+  Func("f", {Param("a", ty.i32()), Param("b", ty.i32()), Param("c", ty.i32())},
+       ty.i32(), {Return(Mul(Add("a", "b"), "c"))});
+
+  Func("main", {}, ty.void_(),
+       {create<ast::CallStatement>(Call("ignore", Call("f", 1, 2, 3)))},
+       {Stage(ast::PipelineStage::kCompute)});
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate(out)) << gen.error();
+  EXPECT_EQ(result(), R"(int f(int a, int b, int c) {
+  return ((a + b) * c);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  f(1, 2, 3);
+  return;
+}
+
+)");
+
+  Validate();
+}
+
 }  // namespace
 }  // namespace hlsl
 }  // namespace writer

src/writer/msl/generator_impl.cc

@@ -354,58 +354,7 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
   auto* ident = expr->func();
   auto* call = program_->Sem().Get(expr);
   if (auto* intrinsic = call->Target()->As<sem::Intrinsic>()) {
-    if (intrinsic->IsTexture()) {
-      return EmitTextureCall(expr, intrinsic);
-    }
-    if (intrinsic->Type() == sem::IntrinsicType::kPack2x16float ||
-        intrinsic->Type() == sem::IntrinsicType::kUnpack2x16float) {
-      make_indent();
-      if (intrinsic->Type() == sem::IntrinsicType::kPack2x16float) {
-        out_ << "as_type<uint>(half2(";
-      } else {
-        out_ << "float2(as_type<half2>(";
-      }
-      if (!EmitExpression(expr->params()[0])) {
-        return false;
-      }
-      out_ << "))";
-      return true;
-    }
-    // TODO(crbug.com/tint/661): Combine sequential barriers to a single
-    // instruction.
-    if (intrinsic->Type() == sem::IntrinsicType::kStorageBarrier) {
-      make_indent();
-      out_ << "threadgroup_barrier(mem_flags::mem_device)";
-      return true;
-    }
-    if (intrinsic->Type() == sem::IntrinsicType::kWorkgroupBarrier) {
-      make_indent();
-      out_ << "threadgroup_barrier(mem_flags::mem_threadgroup)";
-      return true;
-    }
-    auto name = generate_builtin_name(intrinsic);
-    if (name.empty()) {
-      return false;
-    }
-
-    make_indent();
-    out_ << name << "(";
-
-    bool first = true;
-    const auto& params = expr->params();
-    for (auto* param : params) {
-      if (!first) {
-        out_ << ", ";
-      }
-      first = false;
-
-      if (!EmitExpression(param)) {
-        return false;
-      }
-    }
-
-    out_ << ")";
-    return true;
+    return EmitIntrinsicCall(expr, intrinsic);
   }
 
   auto name = program_->Symbols().NameFor(ident->symbol());
@@ -492,6 +441,75 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
   return true;
 }
 
+bool GeneratorImpl::EmitIntrinsicCall(ast::CallExpression* expr,
+                                      const sem::Intrinsic* intrinsic) {
+  if (intrinsic->IsTexture()) {
+    return EmitTextureCall(expr, intrinsic);
+  }
+
+  switch (intrinsic->Type()) {
+    case sem::IntrinsicType::kPack2x16float:
+    case sem::IntrinsicType::kUnpack2x16float: {
+      make_indent();
+      if (intrinsic->Type() == sem::IntrinsicType::kPack2x16float) {
+        out_ << "as_type<uint>(half2(";
+      } else {
+        out_ << "float2(as_type<half2>(";
+      }
+      if (!EmitExpression(expr->params()[0])) {
+        return false;
+      }
+      out_ << "))";
+      return true;
+    }
+    // TODO(crbug.com/tint/661): Combine sequential barriers to a single
+    // instruction.
+    case sem::IntrinsicType::kStorageBarrier: {
+      make_indent();
+      out_ << "threadgroup_barrier(mem_flags::mem_device)";
+      return true;
+    }
+    case sem::IntrinsicType::kWorkgroupBarrier: {
+      make_indent();
+      out_ << "threadgroup_barrier(mem_flags::mem_threadgroup)";
+      return true;
+    }
+    case sem::IntrinsicType::kIgnore: {
+      out_ << "(void) ";
+      if (!EmitExpression(expr->params()[0])) {
+        return false;
+      }
+      return true;
+    }
+    default:
+      break;
+  }
+
+  auto name = generate_builtin_name(intrinsic);
+  if (name.empty()) {
+    return false;
+  }
+
+  make_indent();
+  out_ << name << "(";
+
+  bool first = true;
+  const auto& params = expr->params();
+  for (auto* param : params) {
+    if (!first) {
+      out_ << ", ";
+    }
+    first = false;
+
+    if (!EmitExpression(param)) {
+      return false;
+    }
+  }
+
+  out_ << ")";
+  return true;
+}
+
 bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr,
                                     const sem::Intrinsic* intrinsic) {
   using Usage = sem::ParameterUsage;

src/writer/msl/generator_impl.h

@@ -101,6 +101,12 @@ class GeneratorImpl : public TextGenerator {
   /// @param expr the call expression
   /// @returns true if the call expression is emitted
   bool EmitCall(ast::CallExpression* expr);
+  /// Handles generating an intrinsic call expression
+  /// @param expr the call expression
+  /// @param intrinsic the intrinsic being called
+  /// @returns true if the call expression is emitted
+  bool EmitIntrinsicCall(ast::CallExpression* expr,
+                         const sem::Intrinsic* intrinsic);
   /// Handles generating a call to a texture function (`textureSample`,
   /// `textureSampleGrad`, etc)
   /// @param expr the call expression

src/writer/msl/generator_impl_intrinsic_test.cc

@@ -338,6 +338,34 @@ TEST_F(MslGeneratorImplTest, Unpack2x16Float) {
   EXPECT_EQ(gen.result(), "  float2(as_type<half2>(p1))");
 }
 
+TEST_F(MslGeneratorImplTest, Ignore) {
+  Func("f", {Param("a", ty.i32()), Param("b", ty.i32()), Param("c", ty.i32())},
+       ty.i32(), {Return(Mul(Add("a", "b"), "c"))});
+
+  Func("main", {}, ty.void_(),
+       {create<ast::CallStatement>(Call("ignore", Call("f", 1, 2, 3)))},
+       {Stage(ast::PipelineStage::kCompute)});
+
+  GeneratorImpl& gen = Build();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+int f(int a, int b, int c) {
+  return ((a + b) * c);
+}
+
+kernel void main() {
+  (void) f(1, 2, 3);
+  return;
+}
+
+)");
+
+  Validate();
+}
+
 }  // namespace
 }  // namespace msl
 }  // namespace writer

src/writer/spirv/builder.cc

@@ -2237,6 +2237,14 @@ uint32_t Builder::GenerateIntrinsic(ast::CallExpression* call,
     case IntrinsicType::kFwidthFine:
       op = spv::Op::OpFwidthFine;
       break;
+    case IntrinsicType::kIgnore:
+      // Evaluate the single argument, return the non-zero result_id which isn't
+      // associated with any op (ignore returns void, so this cannot be used in
+      // an expression).
+      if (!get_param_as_value_id(0)) {
+        return 0;
+      }
+      return result_id;
     case IntrinsicType::kIsInf:
       op = spv::Op::OpIsInf;
       break;

src/writer/spirv/builder_intrinsic_test.cc

@@ -1818,6 +1818,43 @@ TEST_F(IntrinsicBuilderTest, Call_StorageBarrier) {
   Validate(b);
 }
 
+TEST_F(IntrinsicBuilderTest, Call_Ignore) {
+  Func("f", {Param("a", ty.i32()), Param("b", ty.i32()), Param("c", ty.i32())},
+       ty.i32(), {Return(Mul(Add("a", "b"), "c"))});
+
+  Func("main", {}, ty.void_(),
+       {
+           create<ast::CallStatement>(Call("ignore", Call("f", 1, 2, 3))),
+       },
+       {
+           Stage(ast::PipelineStage::kCompute),
+       });
+
+  spirv::Builder& b = Build();
+
+  ASSERT_TRUE(b.Build()) << b.error();
+
+  ASSERT_EQ(b.functions().size(), 2u);
+
+  auto* expected_types = R"(%2 = OpTypeInt 32 1
+%1 = OpTypeFunction %2 %2 %2 %2
+%11 = OpTypeVoid
+%10 = OpTypeFunction %11
+%16 = OpConstant %2 1
+%17 = OpConstant %2 2
+%18 = OpConstant %2 3
+)";
+  auto got_types = DumpInstructions(b.types());
+  EXPECT_EQ(expected_types, got_types);
+
+  auto* expected_instructions = R"(%15 = OpFunctionCall %2 %3 %16 %17 %18
+)";
+  auto got_instructions = DumpInstructions(b.functions()[1].instructions());
+  EXPECT_EQ(expected_instructions, got_instructions);
+
+  Validate(b);
+}
+
 }  // namespace
 }  // namespace spirv
 }  // namespace writer

test/intrinsics/gen/ignore/5016e5.wgsl

@@ -0,0 +1,44 @@
+// 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: sampler;
+fn ignore_5016e5() {
+  ignore(arg_0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_5016e5();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_5016e5();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_5016e5();
+}

test/intrinsics/gen/ignore/5016e5.wgsl.expected.hlsl

@@ -0,0 +1,27 @@
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+SamplerState arg_0 : register(s0, space1);
+
+void ignore_5016e5() {
+  arg_0;
+}
+
+tint_symbol vertex_main() {
+  ignore_5016e5();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  ignore_5016e5();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ignore_5016e5();
+  return;
+}
+

test/intrinsics/gen/ignore/5016e5.wgsl.expected.msl

@@ -0,0 +1,37 @@
+SKIP: FAILED
+
+
+
+Validation Failure:
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void ignore_5016e5() {
+  (void) arg_0;
+}
+
+vertex tint_symbol vertex_main() {
+  ignore_5016e5();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  ignore_5016e5();
+  return;
+}
+
+kernel void compute_main() {
+  ignore_5016e5();
+  return;
+}
+
+
+tint_nr1Dd1.metal:9:10: error: use of undeclared identifier 'arg_0'
+  (void) arg_0;
+         ^
+1 error generated.

test/intrinsics/gen/ignore/5016e5.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_5016e5 "ignore_5016e5"
+               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 = OpTypeSampler
+%_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_5016e5 = 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_5016e5
+         %26 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ignore_5016e5
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ignore_5016e5
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/ignore/5016e5.wgsl.expected.wgsl

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

test/intrinsics/gen/ignore/509355.wgsl

@@ -0,0 +1,44 @@
+// 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_cube;
+fn ignore_509355() {
+  ignore(arg_0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_509355();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_509355();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_509355();
+}

test/intrinsics/gen/ignore/509355.wgsl.expected.hlsl

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

test/intrinsics/gen/ignore/509355.wgsl.expected.msl

@@ -0,0 +1,37 @@
+SKIP: FAILED
+
+
+
+Validation Failure:
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void ignore_509355() {
+  (void) arg_0;
+}
+
+vertex tint_symbol vertex_main() {
+  ignore_509355();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  ignore_509355();
+  return;
+}
+
+kernel void compute_main() {
+  ignore_509355();
+  return;
+}
+
+
+tint_BUnzrQ.metal:9:10: error: use of undeclared identifier 'arg_0'
+  (void) arg_0;
+         ^
+1 error generated.

test/intrinsics/gen/ignore/509355.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_509355 "ignore_509355"
+               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 Cube 1 0 0 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_509355 = 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_509355
+         %26 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ignore_509355
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ignore_509355
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/ignore/509355.wgsl.expected.wgsl

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

test/intrinsics/gen/ignore/51aeb7.wgsl

@@ -0,0 +1,43 @@
+// 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
+////////////////////////////////////////////////////////////////////////////////
+
+fn ignore_51aeb7() {
+  ignore(1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_51aeb7();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_51aeb7();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_51aeb7();
+}

test/intrinsics/gen/ignore/51aeb7.wgsl.expected.hlsl

@@ -0,0 +1,25 @@
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+void ignore_51aeb7() {
+  1;
+}
+
+tint_symbol vertex_main() {
+  ignore_51aeb7();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  ignore_51aeb7();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ignore_51aeb7();
+  return;
+}
+

test/intrinsics/gen/ignore/51aeb7.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void ignore_51aeb7() {
+  (void) 1;
+}
+
+vertex tint_symbol vertex_main() {
+  ignore_51aeb7();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  ignore_51aeb7();
+  return;
+}
+
+kernel void compute_main() {
+  ignore_51aeb7();
+  return;
+}
+

test/intrinsics/gen/ignore/51aeb7.wgsl.expected.spvasm

@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; 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 %tint_symbol_1 "tint_symbol_1"
+               OpName %ignore_51aeb7 "ignore_51aeb7"
+               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 %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
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+         %16 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%ignore_51aeb7 = OpFunction %void None %9
+         %12 = OpLabel
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %16
+%tint_symbol = OpFunctionParameter %v4float
+         %19 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %21 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %23 = OpFunctionCall %void %ignore_51aeb7
+         %24 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %26 = OpLabel
+         %27 = OpFunctionCall %void %ignore_51aeb7
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %ignore_51aeb7
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/ignore/51aeb7.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn ignore_51aeb7() {
+  ignore(1);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_51aeb7();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_51aeb7();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_51aeb7();
+}

test/intrinsics/gen/ignore/5c9edf.wgsl

@@ -0,0 +1,44 @@
+// 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_external;
+fn ignore_5c9edf() {
+  ignore(arg_0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_5c9edf();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_5c9edf();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_5c9edf();
+}

test/intrinsics/gen/ignore/5c9edf.wgsl.expected.hlsl

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

test/intrinsics/gen/ignore/5c9edf.wgsl.expected.msl

@@ -0,0 +1,37 @@
+SKIP: FAILED
+
+
+
+Validation Failure:
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void ignore_5c9edf() {
+  (void) arg_0;
+}
+
+vertex tint_symbol vertex_main() {
+  ignore_5c9edf();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  ignore_5c9edf();
+  return;
+}
+
+kernel void compute_main() {
+  ignore_5c9edf();
+  return;
+}
+
+
+tint_520czX.metal:9:10: error: use of undeclared identifier 'arg_0'
+  (void) arg_0;
+         ^
+1 error generated.

test/intrinsics/gen/ignore/5c9edf.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_5c9edf "ignore_5c9edf"
+               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 0 0 0 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_5c9edf = 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_5c9edf
+         %26 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ignore_5c9edf
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ignore_5c9edf
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/ignore/5c9edf.wgsl.expected.wgsl

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

test/intrinsics/gen/ignore/6698df.wgsl

@@ -0,0 +1,43 @@
+// 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
+////////////////////////////////////////////////////////////////////////////////
+
+fn ignore_6698df() {
+  ignore(1u);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_6698df();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_6698df();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_6698df();
+}

test/intrinsics/gen/ignore/6698df.wgsl.expected.hlsl

@@ -0,0 +1,25 @@
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+void ignore_6698df() {
+  1u;
+}
+
+tint_symbol vertex_main() {
+  ignore_6698df();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  ignore_6698df();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ignore_6698df();
+  return;
+}
+

test/intrinsics/gen/ignore/6698df.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void ignore_6698df() {
+  (void) 1u;
+}
+
+vertex tint_symbol vertex_main() {
+  ignore_6698df();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  ignore_6698df();
+  return;
+}
+
+kernel void compute_main() {
+  ignore_6698df();
+  return;
+}
+

test/intrinsics/gen/ignore/6698df.wgsl.expected.spvasm

@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; 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 %tint_symbol_1 "tint_symbol_1"
+               OpName %ignore_6698df "ignore_6698df"
+               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 %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
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+         %16 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%ignore_6698df = OpFunction %void None %9
+         %12 = OpLabel
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %16
+%tint_symbol = OpFunctionParameter %v4float
+         %19 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %21 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %23 = OpFunctionCall %void %ignore_6698df
+         %24 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %26 = OpLabel
+         %27 = OpFunctionCall %void %ignore_6698df
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %ignore_6698df
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/ignore/6698df.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn ignore_6698df() {
+  ignore(1u);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_6698df();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_6698df();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_6698df();
+}

test/intrinsics/gen/ignore/ad88be.wgsl

@@ -0,0 +1,44 @@
+// 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_cube_array;
+fn ignore_ad88be() {
+  ignore(arg_0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_ad88be();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_ad88be();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_ad88be();
+}

test/intrinsics/gen/ignore/ad88be.wgsl.expected.hlsl

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

test/intrinsics/gen/ignore/ad88be.wgsl.expected.msl

@@ -0,0 +1,37 @@
+SKIP: FAILED
+
+
+
+Validation Failure:
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void ignore_ad88be() {
+  (void) arg_0;
+}
+
+vertex tint_symbol vertex_main() {
+  ignore_ad88be();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  ignore_ad88be();
+  return;
+}
+
+kernel void compute_main() {
+  ignore_ad88be();
+  return;
+}
+
+
+tint_BhQwjP.metal:9:10: error: use of undeclared identifier 'arg_0'
+  (void) arg_0;
+         ^
+1 error generated.

test/intrinsics/gen/ignore/ad88be.wgsl.expected.spvasm

@@ -0,0 +1,69 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+               OpCapability Shader
+               OpCapability SampledCubeArray
+               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_ad88be "ignore_ad88be"
+               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 Cube 1 1 0 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_ad88be = 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_ad88be
+         %26 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ignore_ad88be
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ignore_ad88be
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/ignore/ad88be.wgsl.expected.wgsl

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

test/intrinsics/gen/ignore/b469af.wgsl

@@ -0,0 +1,44 @@
+// 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: sampler_comparison;
+fn ignore_b469af() {
+  ignore(arg_0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_b469af();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_b469af();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_b469af();
+}

test/intrinsics/gen/ignore/b469af.wgsl.expected.hlsl

@@ -0,0 +1,27 @@
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+SamplerComparisonState arg_0 : register(s0, space1);
+
+void ignore_b469af() {
+  arg_0;
+}
+
+tint_symbol vertex_main() {
+  ignore_b469af();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  ignore_b469af();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ignore_b469af();
+  return;
+}
+

test/intrinsics/gen/ignore/b469af.wgsl.expected.msl

@@ -0,0 +1,37 @@
+SKIP: FAILED
+
+
+
+Validation Failure:
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void ignore_b469af() {
+  (void) arg_0;
+}
+
+vertex tint_symbol vertex_main() {
+  ignore_b469af();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  ignore_b469af();
+  return;
+}
+
+kernel void compute_main() {
+  ignore_b469af();
+  return;
+}
+
+
+tint_vztXdF.metal:9:10: error: use of undeclared identifier 'arg_0'
+  (void) arg_0;
+         ^
+1 error generated.

test/intrinsics/gen/ignore/b469af.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_b469af "ignore_b469af"
+               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 = OpTypeSampler
+%_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_b469af = 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_b469af
+         %26 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ignore_b469af
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ignore_b469af
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/ignore/b469af.wgsl.expected.wgsl

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

test/intrinsics/gen/ignore/c8a0ee.wgsl

@@ -0,0 +1,44 @@
+// 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_2d_array;
+fn ignore_c8a0ee() {
+  ignore(arg_0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_c8a0ee();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_c8a0ee();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_c8a0ee();
+}

test/intrinsics/gen/ignore/c8a0ee.wgsl.expected.hlsl

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

test/intrinsics/gen/ignore/c8a0ee.wgsl.expected.msl

@@ -0,0 +1,37 @@
+SKIP: FAILED
+
+
+
+Validation Failure:
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void ignore_c8a0ee() {
+  (void) arg_0;
+}
+
+vertex tint_symbol vertex_main() {
+  ignore_c8a0ee();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  ignore_c8a0ee();
+  return;
+}
+
+kernel void compute_main() {
+  ignore_c8a0ee();
+  return;
+}
+
+
+tint_77yvjP.metal:9:10: error: use of undeclared identifier 'arg_0'
+  (void) arg_0;
+         ^
+1 error generated.

test/intrinsics/gen/ignore/c8a0ee.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_c8a0ee "ignore_c8a0ee"
+               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 1 0 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_c8a0ee = 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_c8a0ee
+         %26 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ignore_c8a0ee
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ignore_c8a0ee
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/ignore/c8a0ee.wgsl.expected.wgsl

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

test/intrinsics/gen/ignore/d91a2f.wgsl

@@ -0,0 +1,43 @@
+// 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
+////////////////////////////////////////////////////////////////////////////////
+
+fn ignore_d91a2f() {
+  ignore(1.0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_d91a2f();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_d91a2f();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_d91a2f();
+}

test/intrinsics/gen/ignore/d91a2f.wgsl.expected.hlsl

@@ -0,0 +1,25 @@
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+void ignore_d91a2f() {
+  1.0f;
+}
+
+tint_symbol vertex_main() {
+  ignore_d91a2f();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  ignore_d91a2f();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ignore_d91a2f();
+  return;
+}
+

test/intrinsics/gen/ignore/d91a2f.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void ignore_d91a2f() {
+  (void) 1.0f;
+}
+
+vertex tint_symbol vertex_main() {
+  ignore_d91a2f();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  ignore_d91a2f();
+  return;
+}
+
+kernel void compute_main() {
+  ignore_d91a2f();
+  return;
+}
+

test/intrinsics/gen/ignore/d91a2f.wgsl.expected.spvasm

@@ -0,0 +1,61 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; 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 %tint_symbol_1 "tint_symbol_1"
+               OpName %ignore_d91a2f "ignore_d91a2f"
+               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 %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
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+         %15 = OpTypeFunction %void %v4float
+%ignore_d91a2f = OpFunction %void None %9
+         %12 = OpLabel
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %15
+%tint_symbol = OpFunctionParameter %v4float
+         %18 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %20 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %21 = OpFunctionCall %void %ignore_d91a2f
+         %22 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %24 = OpLabel
+         %25 = OpFunctionCall %void %ignore_d91a2f
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %ignore_d91a2f
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/ignore/d91a2f.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn ignore_d91a2f() {
+  ignore(1.0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_d91a2f();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_d91a2f();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_d91a2f();
+}

test/intrinsics/gen/ignore/e0187b.wgsl

@@ -0,0 +1,44 @@
+// 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_2d;
+fn ignore_e0187b() {
+  ignore(arg_0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_e0187b();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_e0187b();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_e0187b();
+}

test/intrinsics/gen/ignore/e0187b.wgsl.expected.hlsl

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

test/intrinsics/gen/ignore/e0187b.wgsl.expected.msl

@@ -0,0 +1,37 @@
+SKIP: FAILED
+
+
+
+Validation Failure:
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void ignore_e0187b() {
+  (void) arg_0;
+}
+
+vertex tint_symbol vertex_main() {
+  ignore_e0187b();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  ignore_e0187b();
+  return;
+}
+
+kernel void compute_main() {
+  ignore_e0187b();
+  return;
+}
+
+
+tint_gjL9b8.metal:9:10: error: use of undeclared identifier 'arg_0'
+  (void) arg_0;
+         ^
+1 error generated.

test/intrinsics/gen/ignore/e0187b.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_e0187b "ignore_e0187b"
+               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 0 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_e0187b = 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_e0187b
+         %26 = OpFunctionCall %void %tint_symbol_2 %11
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %12
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ignore_e0187b
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %12
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ignore_e0187b
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/ignore/e0187b.wgsl.expected.wgsl

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

test/intrinsics/gen/ignore/f414a6.wgsl

@@ -0,0 +1,43 @@
+// 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
+////////////////////////////////////////////////////////////////////////////////
+
+fn ignore_f414a6() {
+  ignore(bool());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_f414a6();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_f414a6();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_f414a6();
+}

test/intrinsics/gen/ignore/f414a6.wgsl.expected.hlsl

@@ -0,0 +1,25 @@
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+void ignore_f414a6() {
+  false;
+}
+
+tint_symbol vertex_main() {
+  ignore_f414a6();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  ignore_f414a6();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ignore_f414a6();
+  return;
+}
+

test/intrinsics/gen/ignore/f414a6.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void ignore_f414a6() {
+  (void) bool();
+}
+
+vertex tint_symbol vertex_main() {
+  ignore_f414a6();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  ignore_f414a6();
+  return;
+}
+
+kernel void compute_main() {
+  ignore_f414a6();
+  return;
+}
+

test/intrinsics/gen/ignore/f414a6.wgsl.expected.spvasm

@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; 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 %tint_symbol_1 "tint_symbol_1"
+               OpName %ignore_f414a6 "ignore_f414a6"
+               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 %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
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+       %bool = OpTypeBool
+         %15 = OpConstantNull %bool
+         %16 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%ignore_f414a6 = OpFunction %void None %9
+         %12 = OpLabel
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %16
+%tint_symbol = OpFunctionParameter %v4float
+         %19 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %21 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %23 = OpFunctionCall %void %ignore_f414a6
+         %24 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %26 = OpLabel
+         %27 = OpFunctionCall %void %ignore_f414a6
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %ignore_f414a6
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/ignore/f414a6.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn ignore_f414a6() {
+  ignore(bool());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  ignore_f414a6();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  ignore_f414a6();
+}
+
+[[stage(compute)]]
+fn compute_main() {
+  ignore_f414a6();
+}

test/intrinsics/ignore.wgsl

@@ -0,0 +1,8 @@
+fn f(a: i32, b: i32, c: i32) -> i32 {
+    return a * b + c;
+}
+
+[[stage(compute)]]
+fn main() {
+    ignore(f(1, 2, 3));
+}

test/intrinsics/ignore.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+int f(int a, int b, int c) {
+  return ((a * b) + c);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  f(1, 2, 3);
+  return;
+}
+

test/intrinsics/ignore.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int f(int a, int b, int c) {
+  return ((a * b) + c);
+}
+
+kernel void tint_symbol() {
+  (void) f(1, 2, 3);
+  return;
+}
+

test/intrinsics/ignore.wgsl.expected.spvasm

@@ -0,0 +1,35 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 19
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %f "f"
+               OpName %a "a"
+               OpName %b "b"
+               OpName %c "c"
+               OpName %main "main"
+        %int = OpTypeInt 32 1
+          %1 = OpTypeFunction %int %int %int %int
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void
+      %int_1 = OpConstant %int 1
+      %int_2 = OpConstant %int 2
+      %int_3 = OpConstant %int 3
+          %f = OpFunction %int None %1
+          %a = OpFunctionParameter %int
+          %b = OpFunctionParameter %int
+          %c = OpFunctionParameter %int
+          %7 = OpLabel
+          %8 = OpIMul %int %a %b
+          %9 = OpIAdd %int %8 %c
+               OpReturnValue %9
+               OpFunctionEnd
+       %main = OpFunction %void None %10
+         %13 = OpLabel
+         %15 = OpFunctionCall %int %f %int_1 %int_2 %int_3
+               OpReturn
+               OpFunctionEnd

test/intrinsics/ignore.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+fn f(a : i32, b : i32, c : i32) -> i32 {
+  return ((a * b) + c);
+}
+
+[[stage(compute)]]
+fn main() {
+  ignore(f(1, 2, 3));
+}