Add const-eval for `min` and `max`

This CL adds const-eval for `min` and `max`.

Bug: tint:1581
Change-Id: Ica68ba312f21767c46d57d83570ddc72ee857231
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/110166
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

src/tint/intrinsics.def

@@ -497,10 +497,10 @@ fn log<T: f32_f16>(T) -> T
 fn log<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
 fn log2<T: f32_f16>(T) -> T
 fn log2<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
-fn max<T: fiu32_f16>(T, T) -> T
+@const fn max<T: fia_fiu32_f16>(T, T) -> T
-fn max<N: num, T: fiu32_f16>(vec<N, T>, vec<N, T>) -> vec<N, T>
+@const fn max<N: num, T: fia_fiu32_f16>(vec<N, T>, vec<N, T>) -> vec<N, T>
-fn min<T: fiu32_f16>(T, T) -> T
+@const fn min<T: fia_fiu32_f16>(T, T) -> T
-fn min<N: num, T: fiu32_f16>(vec<N, T>, vec<N, T>) -> vec<N, T>
+@const fn min<N: num, T: fia_fiu32_f16>(vec<N, T>, vec<N, T>) -> vec<N, T>
 fn mix<T: f32_f16>(T, T, T) -> T
 fn mix<N: num, T: f32_f16>(vec<N, T>, vec<N, T>, vec<N, T>) -> vec<N, T>
 fn mix<N: num, T: f32_f16>(vec<N, T>, vec<N, T>, T) -> vec<N, T>

src/tint/resolver/const_eval.cc

@@ -2119,6 +2119,30 @@ ConstEval::Result ConstEval::insertBits(const sem::Type* ty,
     return TransformElements(builder, ty, transform, args[0], args[1]);
 }
 
+ConstEval::Result ConstEval::max(const sem::Type* ty,
+                                 utils::VectorRef<const sem::Constant*> args,
+                                 const Source&) {
+    auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
+        auto create = [&](auto e0, auto e1) {
+            return CreateElement(builder, c0->Type(), decltype(e0)(std::max(e0, e1)));
+        };
+        return Dispatch_fia_fiu32_f16(create, c0, c1);
+    };
+    return TransformElements(builder, ty, transform, args[0], args[1]);
+}
+
+ConstEval::Result ConstEval::min(const sem::Type* ty,
+                                 utils::VectorRef<const sem::Constant*> args,
+                                 const Source&) {
+    auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
+        auto create = [&](auto e0, auto e1) {
+            return CreateElement(builder, c0->Type(), decltype(e0)(std::min(e0, e1)));
+        };
+        return Dispatch_fia_fiu32_f16(create, c0, c1);
+    };
+    return TransformElements(builder, ty, transform, args[0], args[1]);
+}
+
 ConstEval::Result ConstEval::pack2x16float(const sem::Type* ty,
                                            utils::VectorRef<const sem::Constant*> args,
                                            const Source& source) {

src/tint/resolver/const_eval.h

@@ -584,6 +584,24 @@ class ConstEval {
                       utils::VectorRef<const sem::Constant*> args,
                       const Source& source);
 
+    /// max builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    Result max(const sem::Type* ty,
+               utils::VectorRef<const sem::Constant*> args,
+               const Source& source);
+
+    /// min builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    Result min(const sem::Type* ty,  // NOLINT(build/include_what_you_use)  -- confused by min
+               utils::VectorRef<const sem::Constant*> args,
+               const Source& source);
+
     /// pack2x16float builtin
     /// @param ty the expression type
     /// @param args the input arguments

src/tint/resolver/const_eval_builtin_test.cc

@@ -1291,6 +1291,60 @@ INSTANTIATE_TEST_SUITE_P(ExtractBits,
                              std::make_tuple(1, u32::Highest()),  //
                              std::make_tuple(u32::Highest(), u32::Highest())));
 
+template <typename T>
+std::vector<Case> MaxCases() {
+    return {
+        C({T(0), T(0)}, T(0)),
+        C({T(0), T::Highest()}, T::Highest()),
+        C({T::Lowest(), T(0)}, T(0)),
+        C({T::Highest(), T::Lowest()}, T::Highest()),
+        C({T::Highest(), T::Highest()}, T::Highest()),
+        C({T::Lowest(), T::Lowest()}, T::Lowest()),
+
+        // Vector tests
+        C({Vec(T(0), T(0)), Vec(T(0), T(42))}, Vec(T(0), T(42))),
+        C({Vec(T::Lowest(), T(0)), Vec(T(0), T::Lowest())}, Vec(T(0), T(0))),
+        C({Vec(T::Lowest(), T::Highest()), Vec(T::Highest(), T::Lowest())},
+          Vec(T::Highest(), T::Highest())),
+    };
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Max,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kMax),
+                     testing::ValuesIn(Concat(MaxCases<AInt>(),  //
+                                              MaxCases<i32>(),
+                                              MaxCases<u32>(),
+                                              MaxCases<AFloat>(),
+                                              MaxCases<f32>(),
+                                              MaxCases<f16>()))));
+
+template <typename T>
+std::vector<Case> MinCases() {
+    return {C({T(0), T(0)}, T(0)),                //
+            C({T(0), T(42)}, T(0)),               //
+            C({T::Lowest(), T(0)}, T::Lowest()),  //
+            C({T(0), T::Highest()}, T(0)),        //
+            C({T::Highest(), T::Lowest()}, T::Lowest()),
+            C({T::Highest(), T::Highest()}, T::Highest()),
+            C({T::Lowest(), T::Lowest()}, T::Lowest()),
+
+            // Vector tests
+            C({Vec(T(0), T(0)), Vec(T(0), T(42))}, Vec(T(0), T(0))),
+            C({Vec(T::Lowest(), T(0), T(1)), Vec(T(0), T(42), T::Highest())},
+              Vec(T::Lowest(), T(0), T(1)))};
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Min,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kMin),
+                     testing::ValuesIn(Concat(MinCases<AInt>(),  //
+                                              MinCases<i32>(),
+                                              MinCases<u32>(),
+                                              MinCases<AFloat>(),
+                                              MinCases<f32>(),
+                                              MinCases<f16>()))));
+
 std::vector<Case> Pack4x8snormCases() {
     return {
         C({Vec(f32(0), f32(0), f32(0), f32(0))}, Val(u32(0x0000'0000))),

src/tint/resolver/materialize_test.cc

@@ -472,8 +472,13 @@ constexpr Method kSwitchMethods[] = {
 constexpr Method kNoMaterializeMethods[] = {
     Method::kPhonyAssign,  //
     Method::kBinaryOp,
-    // TODO(crbug.com/tint/1504): Enable once "min" supports const evaluation
+};
-    // Method::kBuiltinArg,
+
+/// Methods that do not materialize
+constexpr Method kNoMaterializeScalarVectorMethods[] = {
+    Method::kPhonyAssign,  //
+    Method::kBinaryOp,
+    Method::kBuiltinArg,
 };
 INSTANTIATE_TEST_SUITE_P(
     MaterializeScalar,
@@ -697,6 +702,17 @@ INSTANTIATE_TEST_SUITE_P(NoMaterialize,
                                               Types<AFloatM, AFloatM>(1.0_a, 1.0_a),  //
                                           })));
 
+INSTANTIATE_TEST_SUITE_P(NoMaterializeScalarVector,
+                         MaterializeAbstractNumericToConcreteType,
+                         testing::Combine(testing::Values(Expectation::kNoMaterialize),
+                                          testing::ValuesIn(kNoMaterializeScalarVectorMethods),
+                                          testing::ValuesIn(std::vector<Data>{
+                                              Types<AInt, AInt>(1_a, 1_a),            //
+                                              Types<AIntV, AIntV>(1_a, 1_a),          //
+                                              Types<AFloat, AFloat>(1.0_a, 1.0_a),    //
+                                              Types<AFloatV, AFloatV>(1.0_a, 1.0_a),  //
+                                          })));
+
 INSTANTIATE_TEST_SUITE_P(InvalidConversion,
                          MaterializeAbstractNumericToConcreteType,
                          testing::Combine(testing::Values(Expectation::kInvalidConversion),
@@ -771,9 +787,6 @@ enum class Method {
     // let a = abstract_expr;
     kLet,
 
-    // min(abstract_expr, abstract_expr)
-    kBuiltinArg,
-
     // bitcast<f32>(abstract_expr)
     kBitcastF32Arg,
 
@@ -810,8 +823,6 @@ static std::ostream& operator<<(std::ostream& o, Method m) {
             return o << "var";
         case Method::kLet:
             return o << "let";
-        case Method::kBuiltinArg:
-            return o << "builtin-arg";
         case Method::kBitcastF32Arg:
             return o << "bitcast-f32-arg";
         case Method::kBitcastVec3F32Arg:
@@ -890,10 +901,6 @@ TEST_P(MaterializeAbstractNumericToDefaultType, Test) {
             WrapInFunction(Decl(Let("a", abstract_expr())));
             break;
         }
-        case Method::kBuiltinArg: {
-            WrapInFunction(CallStmt(Call("min", abstract_expr(), abstract_expr())));
-            break;
-        }
         case Method::kBitcastF32Arg: {
             WrapInFunction(Bitcast<f32>(abstract_expr()));
             break;
@@ -949,17 +956,8 @@ TEST_P(MaterializeAbstractNumericToDefaultType, Test) {
         }
         case Expectation::kInvalidConversion: {
             ASSERT_FALSE(r()->Resolve());
-            std::string expect;
+            std::string expect = "error: cannot convert value of type '" + data.abstract_type_name +
-            switch (method) {
+                                 "' to type '" + data.expected_type_name + "'";
-                case Method::kBuiltinArg:
-                    expect = "error: no matching call to min(" + data.abstract_type_name + ", " +
-                             data.abstract_type_name + ")";
-                    break;
-                default:
-                    expect = "error: cannot convert value of type '" + data.abstract_type_name +
-                             "' to type '" + data.expected_type_name + "'";
-                    break;
-            }
             EXPECT_THAT(r()->error(), testing::StartsWith(expect));
             break;
         }
@@ -977,16 +975,17 @@ TEST_P(MaterializeAbstractNumericToDefaultType, Test) {
 constexpr Method kScalarMethods[] = {
     Method::kLet,
     Method::kVar,
-    Method::kBuiltinArg,
     Method::kBitcastF32Arg,
     Method::kTintMaterializeBuiltin,
 };
 
 /// Methods that support vector materialization
 constexpr Method kVectorMethods[] = {
-    Method::kLet,          Method::kVar,
+    Method::kLet,
-    Method::kBuiltinArg,   Method::kBitcastVec3F32Arg,
+    Method::kVar,
-    Method::kRuntimeIndex, Method::kTintMaterializeBuiltin,
+    Method::kBitcastVec3F32Arg,
+    Method::kRuntimeIndex,
+    Method::kTintMaterializeBuiltin,
 };
 
 /// Methods that support matrix materialization

test/tint/builtins/gen/literal/max/067f3a.wgsl

@@ -0,0 +1,43 @@
+// Copyright 2022 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/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn max(vec<2, ia>, vec<2, ia>) -> vec<2, ia>
+fn max_067f3a() {
+  var res = max(vec2(1), vec2(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  max_067f3a();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  max_067f3a();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  max_067f3a();
+}

test/tint/builtins/gen/literal/max/067f3a.wgsl.expected.dxc.hlsl

@@ -0,0 +1,30 @@
+void max_067f3a() {
+  int2 res = (1).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  max_067f3a();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  max_067f3a();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  max_067f3a();
+  return;
+}

test/tint/builtins/gen/literal/max/067f3a.wgsl.expected.fxc.hlsl

@@ -0,0 +1,30 @@
+void max_067f3a() {
+  int2 res = (1).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  max_067f3a();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  max_067f3a();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  max_067f3a();
+  return;
+}

test/tint/builtins/gen/literal/max/067f3a.wgsl.expected.glsl

@@ -0,0 +1,49 @@
+#version 310 es
+
+void max_067f3a() {
+  ivec2 res = ivec2(1);
+}
+
+vec4 vertex_main() {
+  max_067f3a();
+  return vec4(0.0f);
+}
+
+void main() {
+  gl_PointSize = 1.0;
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+void max_067f3a() {
+  ivec2 res = ivec2(1);
+}
+
+void fragment_main() {
+  max_067f3a();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void max_067f3a() {
+  ivec2 res = ivec2(1);
+}
+
+void compute_main() {
+  max_067f3a();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/max/067f3a.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void max_067f3a() {
+  int2 res = int2(1);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  max_067f3a();
+  return float4(0.0f);
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  max_067f3a();
+  return;
+}
+
+kernel void compute_main() {
+  max_067f3a();
+  return;
+}
+

test/tint/builtins/gen/literal/max/067f3a.wgsl.expected.spvasm

@@ -0,0 +1,68 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 34
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %max_067f3a "max_067f3a"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+      %int_1 = OpConstant %int 1
+         %16 = OpConstantComposite %v2int %int_1 %int_1
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+         %19 = OpConstantNull %v2int
+         %20 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %max_067f3a = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_v2int Function %19
+               OpStore %res %16
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %20
+         %22 = OpLabel
+         %23 = OpFunctionCall %void %max_067f3a
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %25 = OpLabel
+         %26 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %26
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %max_067f3a
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %max_067f3a
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/max/067f3a.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn max_067f3a() {
+  var res = max(vec2(1), vec2(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  max_067f3a();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  max_067f3a();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  max_067f3a();
+}

test/tint/builtins/gen/literal/max/0c0aae.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void max_0c0aae() {
-  uint res = max(1u, 1u);
+  uint res = 1u;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/0c0aae.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void max_0c0aae() {
-  uint res = max(1u, 1u);
+  uint res = 1u;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/0c0aae.wgsl.expected.glsl

@@ -1,7 +1,7 @@
 #version 310 es
 
 void max_0c0aae() {
-  uint res = max(1u, 1u);
+  uint res = 1u;
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@ void main() {
 precision mediump float;
 
 void max_0c0aae() {
-  uint res = max(1u, 1u);
+  uint res = 1u;
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@ void main() {
 #version 310 es
 
 void max_0c0aae() {
-  uint res = max(1u, 1u);
+  uint res = 1u;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/max/0c0aae.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void max_0c0aae() {
-  uint res = max(1u, 1u);
+  uint res = 1u;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/0c0aae.wgsl.expected.spvasm

@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 32
 ; Schema: 0
                OpCapability Shader
-         %15 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -34,35 +33,34 @@
        %uint = OpTypeInt 32 0
      %uint_1 = OpConstant %uint 1
 %_ptr_Function_uint = OpTypePointer Function %uint
-         %19 = OpConstantNull %uint
+         %17 = OpConstantNull %uint
-         %20 = OpTypeFunction %v4float
+         %18 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %max_0c0aae = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_uint Function %19
+        %res = OpVariable %_ptr_Function_uint Function %17
-         %13 = OpExtInst %uint %15 UMax %uint_1 %uint_1
+               OpStore %res %uint_1
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
+%vertex_main_inner = OpFunction %v4float None %18
-         %22 = OpLabel
+         %20 = OpLabel
-         %23 = OpFunctionCall %void %max_0c0aae
+         %21 = OpFunctionCall %void %max_0c0aae
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %25 = OpLabel
+         %23 = OpLabel
-         %26 = OpFunctionCall %v4float %vertex_main_inner
+         %24 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %26
+               OpStore %value %24
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %29 = OpLabel
+         %27 = OpLabel
-         %30 = OpFunctionCall %void %max_0c0aae
+         %28 = OpFunctionCall %void %max_0c0aae
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
+         %30 = OpLabel
-         %33 = OpFunctionCall %void %max_0c0aae
+         %31 = OpFunctionCall %void %max_0c0aae
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/max/111ac0.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void max_111ac0() {
-  float16_t res = max(float16_t(1.0h), float16_t(1.0h));
+  float16_t res = float16_t(1.0h);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/111ac0.wgsl.expected.glsl

@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void max_111ac0() {
-  float16_t res = max(1.0hf, 1.0hf);
+  float16_t res = 1.0hf;
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@ void main() {
 precision mediump float;
 
 void max_111ac0() {
-  float16_t res = max(1.0hf, 1.0hf);
+  float16_t res = 1.0hf;
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@ void main() {
 #extension GL_AMD_gpu_shader_half_float : require
 
 void max_111ac0() {
-  float16_t res = max(1.0hf, 1.0hf);
+  float16_t res = 1.0hf;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/max/111ac0.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void max_111ac0() {
-  half res = fmax(1.0h, 1.0h);
+  half res = 1.0h;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/111ac0.wgsl.expected.spvasm

@@ -1,14 +1,13 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 32
 ; Schema: 0
                OpCapability Shader
                OpCapability Float16
                OpCapability UniformAndStorageBuffer16BitAccess
                OpCapability StorageBuffer16BitAccess
                OpCapability StorageInputOutput16
-         %15 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -38,35 +37,34 @@
        %half = OpTypeFloat 16
 %half_0x1p_0 = OpConstant %half 0x1p+0
 %_ptr_Function_half = OpTypePointer Function %half
-         %19 = OpConstantNull %half
+         %17 = OpConstantNull %half
-         %20 = OpTypeFunction %v4float
+         %18 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %max_111ac0 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_half Function %19
+        %res = OpVariable %_ptr_Function_half Function %17
-         %13 = OpExtInst %half %15 NMax %half_0x1p_0 %half_0x1p_0
+               OpStore %res %half_0x1p_0
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
+%vertex_main_inner = OpFunction %v4float None %18
-         %22 = OpLabel
+         %20 = OpLabel
-         %23 = OpFunctionCall %void %max_111ac0
+         %21 = OpFunctionCall %void %max_111ac0
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %25 = OpLabel
+         %23 = OpLabel
-         %26 = OpFunctionCall %v4float %vertex_main_inner
+         %24 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %26
+               OpStore %value %24
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %29 = OpLabel
+         %27 = OpLabel
-         %30 = OpFunctionCall %void %max_111ac0
+         %28 = OpFunctionCall %void %max_111ac0
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
+         %30 = OpLabel
-         %33 = OpFunctionCall %void %max_111ac0
+         %31 = OpFunctionCall %void %max_111ac0
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/max/19070a.wgsl

@@ -0,0 +1,43 @@
+// Copyright 2022 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/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn max(vec<4, ia>, vec<4, ia>) -> vec<4, ia>
+fn max_19070a() {
+  var res = max(vec4(1), vec4(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  max_19070a();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  max_19070a();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  max_19070a();
+}

test/tint/builtins/gen/literal/max/19070a.wgsl.expected.dxc.hlsl

@@ -0,0 +1,30 @@
+void max_19070a() {
+  int4 res = (1).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  max_19070a();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  max_19070a();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  max_19070a();
+  return;
+}

test/tint/builtins/gen/literal/max/19070a.wgsl.expected.fxc.hlsl

@@ -0,0 +1,30 @@
+void max_19070a() {
+  int4 res = (1).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  max_19070a();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  max_19070a();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  max_19070a();
+  return;
+}

test/tint/builtins/gen/literal/max/19070a.wgsl.expected.glsl

@@ -0,0 +1,49 @@
+#version 310 es
+
+void max_19070a() {
+  ivec4 res = ivec4(1);
+}
+
+vec4 vertex_main() {
+  max_19070a();
+  return vec4(0.0f);
+}
+
+void main() {
+  gl_PointSize = 1.0;
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+void max_19070a() {
+  ivec4 res = ivec4(1);
+}
+
+void fragment_main() {
+  max_19070a();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void max_19070a() {
+  ivec4 res = ivec4(1);
+}
+
+void compute_main() {
+  max_19070a();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/max/19070a.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void max_19070a() {
+  int4 res = int4(1);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  max_19070a();
+  return float4(0.0f);
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  max_19070a();
+  return;
+}
+
+kernel void compute_main() {
+  max_19070a();
+  return;
+}
+

test/tint/builtins/gen/literal/max/19070a.wgsl.expected.spvasm

@@ -0,0 +1,68 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 34
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %max_19070a "max_19070a"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v4int = OpTypeVector %int 4
+      %int_1 = OpConstant %int 1
+         %16 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
+%_ptr_Function_v4int = OpTypePointer Function %v4int
+         %19 = OpConstantNull %v4int
+         %20 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %max_19070a = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_v4int Function %19
+               OpStore %res %16
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %20
+         %22 = OpLabel
+         %23 = OpFunctionCall %void %max_19070a
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %25 = OpLabel
+         %26 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %26
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %max_19070a
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %max_19070a
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/max/19070a.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn max_19070a() {
+  var res = max(vec4(1), vec4(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  max_19070a();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  max_19070a();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  max_19070a();
+}

test/tint/builtins/gen/literal/max/25eafe.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void max_25eafe() {
-  int3 res = max((1).xxx, (1).xxx);
+  int3 res = (1).xxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/25eafe.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void max_25eafe() {
-  int3 res = max((1).xxx, (1).xxx);
+  int3 res = (1).xxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/25eafe.wgsl.expected.glsl

@@ -1,7 +1,7 @@
 #version 310 es
 
 void max_25eafe() {
-  ivec3 res = max(ivec3(1), ivec3(1));
+  ivec3 res = ivec3(1);
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@ void main() {
 precision mediump float;
 
 void max_25eafe() {
-  ivec3 res = max(ivec3(1), ivec3(1));
+  ivec3 res = ivec3(1);
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@ void main() {
 #version 310 es
 
 void max_25eafe() {
-  ivec3 res = max(ivec3(1), ivec3(1));
+  ivec3 res = ivec3(1);
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/max/25eafe.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void max_25eafe() {
-  int3 res = max(int3(1), int3(1));
+  int3 res = int3(1);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/25eafe.wgsl.expected.spvasm

@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 34
 ; Schema: 0
                OpCapability Shader
-         %16 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -34,37 +33,36 @@
         %int = OpTypeInt 32 1
       %v3int = OpTypeVector %int 3
       %int_1 = OpConstant %int 1
-         %18 = OpConstantComposite %v3int %int_1 %int_1 %int_1
+         %16 = OpConstantComposite %v3int %int_1 %int_1 %int_1
 %_ptr_Function_v3int = OpTypePointer Function %v3int
-         %21 = OpConstantNull %v3int
+         %19 = OpConstantNull %v3int
-         %22 = OpTypeFunction %v4float
+         %20 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %max_25eafe = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v3int Function %21
+        %res = OpVariable %_ptr_Function_v3int Function %19
-         %13 = OpExtInst %v3int %16 SMax %18 %18
+               OpStore %res %16
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
+%vertex_main_inner = OpFunction %v4float None %20
-         %24 = OpLabel
+         %22 = OpLabel
-         %25 = OpFunctionCall %void %max_25eafe
+         %23 = OpFunctionCall %void %max_25eafe
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %27 = OpLabel
+         %25 = OpLabel
-         %28 = OpFunctionCall %v4float %vertex_main_inner
+         %26 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %28
+               OpStore %value %26
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %31 = OpLabel
+         %29 = OpLabel
-         %32 = OpFunctionCall %void %max_25eafe
+         %30 = OpFunctionCall %void %max_25eafe
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %34 = OpLabel
+         %32 = OpLabel
-         %35 = OpFunctionCall %void %max_25eafe
+         %33 = OpFunctionCall %void %max_25eafe
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/max/320815.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void max_320815() {
-  uint2 res = max((1u).xx, (1u).xx);
+  uint2 res = (1u).xx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/320815.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void max_320815() {
-  uint2 res = max((1u).xx, (1u).xx);
+  uint2 res = (1u).xx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/320815.wgsl.expected.glsl

@@ -1,7 +1,7 @@
 #version 310 es
 
 void max_320815() {
-  uvec2 res = max(uvec2(1u), uvec2(1u));
+  uvec2 res = uvec2(1u);
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@ void main() {
 precision mediump float;
 
 void max_320815() {
-  uvec2 res = max(uvec2(1u), uvec2(1u));
+  uvec2 res = uvec2(1u);
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@ void main() {
 #version 310 es
 
 void max_320815() {
-  uvec2 res = max(uvec2(1u), uvec2(1u));
+  uvec2 res = uvec2(1u);
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/max/320815.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void max_320815() {
-  uint2 res = max(uint2(1u), uint2(1u));
+  uint2 res = uint2(1u);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/320815.wgsl.expected.spvasm

@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 34
 ; Schema: 0
                OpCapability Shader
-         %16 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -34,37 +33,36 @@
        %uint = OpTypeInt 32 0
      %v2uint = OpTypeVector %uint 2
      %uint_1 = OpConstant %uint 1
-         %18 = OpConstantComposite %v2uint %uint_1 %uint_1
+         %16 = OpConstantComposite %v2uint %uint_1 %uint_1
 %_ptr_Function_v2uint = OpTypePointer Function %v2uint
-         %21 = OpConstantNull %v2uint
+         %19 = OpConstantNull %v2uint
-         %22 = OpTypeFunction %v4float
+         %20 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %max_320815 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v2uint Function %21
+        %res = OpVariable %_ptr_Function_v2uint Function %19
-         %13 = OpExtInst %v2uint %16 UMax %18 %18
+               OpStore %res %16
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
+%vertex_main_inner = OpFunction %v4float None %20
-         %24 = OpLabel
+         %22 = OpLabel
-         %25 = OpFunctionCall %void %max_320815
+         %23 = OpFunctionCall %void %max_320815
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %27 = OpLabel
+         %25 = OpLabel
-         %28 = OpFunctionCall %v4float %vertex_main_inner
+         %26 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %28
+               OpStore %value %26
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %31 = OpLabel
+         %29 = OpLabel
-         %32 = OpFunctionCall %void %max_320815
+         %30 = OpFunctionCall %void %max_320815
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %34 = OpLabel
+         %32 = OpLabel
-         %35 = OpFunctionCall %void %max_320815
+         %33 = OpFunctionCall %void %max_320815
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/max/34956e.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void max_34956e() {
-  vector<float16_t, 2> res = max((float16_t(1.0h)).xx, (float16_t(1.0h)).xx);
+  vector<float16_t, 2> res = (float16_t(1.0h)).xx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/34956e.wgsl.expected.glsl

@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void max_34956e() {
-  f16vec2 res = max(f16vec2(1.0hf), f16vec2(1.0hf));
+  f16vec2 res = f16vec2(1.0hf);
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@ void main() {
 precision mediump float;
 
 void max_34956e() {
-  f16vec2 res = max(f16vec2(1.0hf), f16vec2(1.0hf));
+  f16vec2 res = f16vec2(1.0hf);
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@ void main() {
 #extension GL_AMD_gpu_shader_half_float : require
 
 void max_34956e() {
-  f16vec2 res = max(f16vec2(1.0hf), f16vec2(1.0hf));
+  f16vec2 res = f16vec2(1.0hf);
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/max/34956e.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void max_34956e() {
-  half2 res = fmax(half2(1.0h), half2(1.0h));
+  half2 res = half2(1.0h);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/34956e.wgsl.expected.spvasm

@@ -1,14 +1,13 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 34
 ; Schema: 0
                OpCapability Shader
                OpCapability Float16
                OpCapability UniformAndStorageBuffer16BitAccess
                OpCapability StorageBuffer16BitAccess
                OpCapability StorageInputOutput16
-         %16 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -38,37 +37,36 @@
        %half = OpTypeFloat 16
      %v2half = OpTypeVector %half 2
 %half_0x1p_0 = OpConstant %half 0x1p+0
-         %18 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
+         %16 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
 %_ptr_Function_v2half = OpTypePointer Function %v2half
-         %21 = OpConstantNull %v2half
+         %19 = OpConstantNull %v2half
-         %22 = OpTypeFunction %v4float
+         %20 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %max_34956e = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v2half Function %21
+        %res = OpVariable %_ptr_Function_v2half Function %19
-         %13 = OpExtInst %v2half %16 NMax %18 %18
+               OpStore %res %16
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
+%vertex_main_inner = OpFunction %v4float None %20
-         %24 = OpLabel
+         %22 = OpLabel
-         %25 = OpFunctionCall %void %max_34956e
+         %23 = OpFunctionCall %void %max_34956e
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %27 = OpLabel
+         %25 = OpLabel
-         %28 = OpFunctionCall %v4float %vertex_main_inner
+         %26 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %28
+               OpStore %value %26
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %31 = OpLabel
+         %29 = OpLabel
-         %32 = OpFunctionCall %void %max_34956e
+         %30 = OpFunctionCall %void %max_34956e
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %34 = OpLabel
+         %32 = OpLabel
-         %35 = OpFunctionCall %void %max_34956e
+         %33 = OpFunctionCall %void %max_34956e
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/max/445169.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void max_445169() {
-  vector<float16_t, 3> res = max((float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx);
+  vector<float16_t, 3> res = (float16_t(1.0h)).xxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/445169.wgsl.expected.glsl

@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void max_445169() {
-  f16vec3 res = max(f16vec3(1.0hf), f16vec3(1.0hf));
+  f16vec3 res = f16vec3(1.0hf);
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@ void main() {
 precision mediump float;
 
 void max_445169() {
-  f16vec3 res = max(f16vec3(1.0hf), f16vec3(1.0hf));
+  f16vec3 res = f16vec3(1.0hf);
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@ void main() {
 #extension GL_AMD_gpu_shader_half_float : require
 
 void max_445169() {
-  f16vec3 res = max(f16vec3(1.0hf), f16vec3(1.0hf));
+  f16vec3 res = f16vec3(1.0hf);
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/max/445169.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void max_445169() {
-  half3 res = fmax(half3(1.0h), half3(1.0h));
+  half3 res = half3(1.0h);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/445169.wgsl.expected.spvasm

@@ -1,14 +1,13 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 34
 ; Schema: 0
                OpCapability Shader
                OpCapability Float16
                OpCapability UniformAndStorageBuffer16BitAccess
                OpCapability StorageBuffer16BitAccess
                OpCapability StorageInputOutput16
-         %16 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -38,37 +37,36 @@
        %half = OpTypeFloat 16
      %v3half = OpTypeVector %half 3
 %half_0x1p_0 = OpConstant %half 0x1p+0
-         %18 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+         %16 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
 %_ptr_Function_v3half = OpTypePointer Function %v3half
-         %21 = OpConstantNull %v3half
+         %19 = OpConstantNull %v3half
-         %22 = OpTypeFunction %v4float
+         %20 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %max_445169 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v3half Function %21
+        %res = OpVariable %_ptr_Function_v3half Function %19
-         %13 = OpExtInst %v3half %16 NMax %18 %18
+               OpStore %res %16
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
+%vertex_main_inner = OpFunction %v4float None %20
-         %24 = OpLabel
+         %22 = OpLabel
-         %25 = OpFunctionCall %void %max_445169
+         %23 = OpFunctionCall %void %max_445169
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %27 = OpLabel
+         %25 = OpLabel
-         %28 = OpFunctionCall %v4float %vertex_main_inner
+         %26 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %28
+               OpStore %value %26
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %31 = OpLabel
+         %29 = OpLabel
-         %32 = OpFunctionCall %void %max_445169
+         %30 = OpFunctionCall %void %max_445169
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %34 = OpLabel
+         %32 = OpLabel
-         %35 = OpFunctionCall %void %max_445169
+         %33 = OpFunctionCall %void %max_445169
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/max/44a39d.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void max_44a39d() {
-  float res = max(1.0f, 1.0f);
+  float res = 1.0f;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/44a39d.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void max_44a39d() {
-  float res = max(1.0f, 1.0f);
+  float res = 1.0f;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/44a39d.wgsl.expected.glsl

@@ -1,7 +1,7 @@
 #version 310 es
 
 void max_44a39d() {
-  float res = max(1.0f, 1.0f);
+  float res = 1.0f;
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@ void main() {
 precision mediump float;
 
 void max_44a39d() {
-  float res = max(1.0f, 1.0f);
+  float res = 1.0f;
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@ void main() {
 #version 310 es
 
 void max_44a39d() {
-  float res = max(1.0f, 1.0f);
+  float res = 1.0f;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/max/44a39d.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void max_44a39d() {
-  float res = fmax(1.0f, 1.0f);
+  float res = 1.0f;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/44a39d.wgsl.expected.spvasm

@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 31
+; Bound: 29
 ; Schema: 0
                OpCapability Shader
-         %14 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -33,33 +32,32 @@
           %9 = OpTypeFunction %void
     %float_1 = OpConstant %float 1
 %_ptr_Function_float = OpTypePointer Function %float
-         %18 = OpTypeFunction %v4float
+         %16 = OpTypeFunction %v4float
  %max_44a39d = OpFunction %void None %9
          %12 = OpLabel
         %res = OpVariable %_ptr_Function_float Function %8
-         %13 = OpExtInst %float %14 NMax %float_1 %float_1
+               OpStore %res %float_1
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %18
+%vertex_main_inner = OpFunction %v4float None %16
-         %20 = OpLabel
+         %18 = OpLabel
-         %21 = OpFunctionCall %void %max_44a39d
+         %19 = OpFunctionCall %void %max_44a39d
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %23 = OpLabel
+         %21 = OpLabel
-         %24 = OpFunctionCall %v4float %vertex_main_inner
+         %22 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %24
+               OpStore %value %22
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %26 = OpLabel
+         %24 = OpLabel
-         %27 = OpFunctionCall %void %max_44a39d
+         %25 = OpFunctionCall %void %max_44a39d
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %29 = OpLabel
+         %27 = OpLabel
-         %30 = OpFunctionCall %void %max_44a39d
+         %28 = OpFunctionCall %void %max_44a39d
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/max/453e04.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void max_453e04() {
-  uint4 res = max((1u).xxxx, (1u).xxxx);
+  uint4 res = (1u).xxxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/453e04.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void max_453e04() {
-  uint4 res = max((1u).xxxx, (1u).xxxx);
+  uint4 res = (1u).xxxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/453e04.wgsl.expected.glsl

@@ -1,7 +1,7 @@
 #version 310 es
 
 void max_453e04() {
-  uvec4 res = max(uvec4(1u), uvec4(1u));
+  uvec4 res = uvec4(1u);
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@ void main() {
 precision mediump float;
 
 void max_453e04() {
-  uvec4 res = max(uvec4(1u), uvec4(1u));
+  uvec4 res = uvec4(1u);
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@ void main() {
 #version 310 es
 
 void max_453e04() {
-  uvec4 res = max(uvec4(1u), uvec4(1u));
+  uvec4 res = uvec4(1u);
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/max/453e04.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void max_453e04() {
-  uint4 res = max(uint4(1u), uint4(1u));
+  uint4 res = uint4(1u);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/453e04.wgsl.expected.spvasm

@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 34
 ; Schema: 0
                OpCapability Shader
-         %16 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -34,37 +33,36 @@
        %uint = OpTypeInt 32 0
      %v4uint = OpTypeVector %uint 4
      %uint_1 = OpConstant %uint 1
-         %18 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1
+         %16 = OpConstantComposite %v4uint %uint_1 %uint_1 %uint_1 %uint_1
 %_ptr_Function_v4uint = OpTypePointer Function %v4uint
-         %21 = OpConstantNull %v4uint
+         %19 = OpConstantNull %v4uint
-         %22 = OpTypeFunction %v4float
+         %20 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %max_453e04 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v4uint Function %21
+        %res = OpVariable %_ptr_Function_v4uint Function %19
-         %13 = OpExtInst %v4uint %16 UMax %18 %18
+               OpStore %res %16
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
+%vertex_main_inner = OpFunction %v4float None %20
-         %24 = OpLabel
+         %22 = OpLabel
-         %25 = OpFunctionCall %void %max_453e04
+         %23 = OpFunctionCall %void %max_453e04
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %27 = OpLabel
+         %25 = OpLabel
-         %28 = OpFunctionCall %v4float %vertex_main_inner
+         %26 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %28
+               OpStore %value %26
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %31 = OpLabel
+         %29 = OpLabel
-         %32 = OpFunctionCall %void %max_453e04
+         %30 = OpFunctionCall %void %max_453e04
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %34 = OpLabel
+         %32 = OpLabel
-         %35 = OpFunctionCall %void %max_453e04
+         %33 = OpFunctionCall %void %max_453e04
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/max/462050.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void max_462050() {
-  float2 res = max((1.0f).xx, (1.0f).xx);
+  float2 res = (1.0f).xx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/462050.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void max_462050() {
-  float2 res = max((1.0f).xx, (1.0f).xx);
+  float2 res = (1.0f).xx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/462050.wgsl.expected.glsl

@@ -1,7 +1,7 @@
 #version 310 es
 
 void max_462050() {
-  vec2 res = max(vec2(1.0f), vec2(1.0f));
+  vec2 res = vec2(1.0f);
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@ void main() {
 precision mediump float;
 
 void max_462050() {
-  vec2 res = max(vec2(1.0f), vec2(1.0f));
+  vec2 res = vec2(1.0f);
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@ void main() {
 #version 310 es
 
 void max_462050() {
-  vec2 res = max(vec2(1.0f), vec2(1.0f));
+  vec2 res = vec2(1.0f);
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/max/462050.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void max_462050() {
-  float2 res = fmax(float2(1.0f), float2(1.0f));
+  float2 res = float2(1.0f);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/462050.wgsl.expected.spvasm

@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 32
 ; Schema: 0
                OpCapability Shader
-         %15 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -33,36 +32,35 @@
           %9 = OpTypeFunction %void
     %v2float = OpTypeVector %float 2
     %float_1 = OpConstant %float 1
-         %17 = OpConstantComposite %v2float %float_1 %float_1
+         %15 = OpConstantComposite %v2float %float_1 %float_1
 %_ptr_Function_v2float = OpTypePointer Function %v2float
-         %20 = OpConstantNull %v2float
+         %18 = OpConstantNull %v2float
-         %21 = OpTypeFunction %v4float
+         %19 = OpTypeFunction %v4float
  %max_462050 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v2float Function %20
+        %res = OpVariable %_ptr_Function_v2float Function %18
-         %13 = OpExtInst %v2float %15 NMax %17 %17
+               OpStore %res %15
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
+%vertex_main_inner = OpFunction %v4float None %19
-         %23 = OpLabel
+         %21 = OpLabel
-         %24 = OpFunctionCall %void %max_462050
+         %22 = OpFunctionCall %void %max_462050
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %26 = OpLabel
+         %24 = OpLabel
-         %27 = OpFunctionCall %v4float %vertex_main_inner
+         %25 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %27
+               OpStore %value %25
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %29 = OpLabel
+         %27 = OpLabel
-         %30 = OpFunctionCall %void %max_462050
+         %28 = OpFunctionCall %void %max_462050
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
+         %30 = OpLabel
-         %33 = OpFunctionCall %void %max_462050
+         %31 = OpFunctionCall %void %max_462050
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/max/482d23.wgsl

@@ -0,0 +1,43 @@
+// Copyright 2022 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/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn max(vec<3, ia>, vec<3, ia>) -> vec<3, ia>
+fn max_482d23() {
+  var res = max(vec3(1), vec3(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  max_482d23();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  max_482d23();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  max_482d23();
+}

test/tint/builtins/gen/literal/max/482d23.wgsl.expected.dxc.hlsl

@@ -0,0 +1,30 @@
+void max_482d23() {
+  int3 res = (1).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  max_482d23();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  max_482d23();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  max_482d23();
+  return;
+}

test/tint/builtins/gen/literal/max/482d23.wgsl.expected.fxc.hlsl

@@ -0,0 +1,30 @@
+void max_482d23() {
+  int3 res = (1).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  max_482d23();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  max_482d23();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  max_482d23();
+  return;
+}

test/tint/builtins/gen/literal/max/482d23.wgsl.expected.glsl

@@ -0,0 +1,49 @@
+#version 310 es
+
+void max_482d23() {
+  ivec3 res = ivec3(1);
+}
+
+vec4 vertex_main() {
+  max_482d23();
+  return vec4(0.0f);
+}
+
+void main() {
+  gl_PointSize = 1.0;
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+void max_482d23() {
+  ivec3 res = ivec3(1);
+}
+
+void fragment_main() {
+  max_482d23();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void max_482d23() {
+  ivec3 res = ivec3(1);
+}
+
+void compute_main() {
+  max_482d23();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/max/482d23.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void max_482d23() {
+  int3 res = int3(1);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  max_482d23();
+  return float4(0.0f);
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  max_482d23();
+  return;
+}
+
+kernel void compute_main() {
+  max_482d23();
+  return;
+}
+

test/tint/builtins/gen/literal/max/482d23.wgsl.expected.spvasm

@@ -0,0 +1,68 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 34
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %max_482d23 "max_482d23"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+      %int_1 = OpConstant %int 1
+         %16 = OpConstantComposite %v3int %int_1 %int_1 %int_1
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+         %19 = OpConstantNull %v3int
+         %20 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %max_482d23 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_v3int Function %19
+               OpStore %res %16
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %20
+         %22 = OpLabel
+         %23 = OpFunctionCall %void %max_482d23
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %25 = OpLabel
+         %26 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %26
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %max_482d23
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %max_482d23
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/max/482d23.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn max_482d23() {
+  var res = max(vec3(1), vec3(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  max_482d23();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  max_482d23();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  max_482d23();
+}

test/tint/builtins/gen/literal/max/4883ac.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void max_4883ac() {
-  float3 res = max((1.0f).xxx, (1.0f).xxx);
+  float3 res = (1.0f).xxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/4883ac.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void max_4883ac() {
-  float3 res = max((1.0f).xxx, (1.0f).xxx);
+  float3 res = (1.0f).xxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/4883ac.wgsl.expected.glsl

@@ -1,7 +1,7 @@
 #version 310 es
 
 void max_4883ac() {
-  vec3 res = max(vec3(1.0f), vec3(1.0f));
+  vec3 res = vec3(1.0f);
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@ void main() {
 precision mediump float;
 
 void max_4883ac() {
-  vec3 res = max(vec3(1.0f), vec3(1.0f));
+  vec3 res = vec3(1.0f);
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@ void main() {
 #version 310 es
 
 void max_4883ac() {
-  vec3 res = max(vec3(1.0f), vec3(1.0f));
+  vec3 res = vec3(1.0f);
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/max/4883ac.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void max_4883ac() {
-  float3 res = fmax(float3(1.0f), float3(1.0f));
+  float3 res = float3(1.0f);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/4883ac.wgsl.expected.spvasm

@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 32
 ; Schema: 0
                OpCapability Shader
-         %15 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -33,36 +32,35 @@
           %9 = OpTypeFunction %void
     %v3float = OpTypeVector %float 3
     %float_1 = OpConstant %float 1
-         %17 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+         %15 = OpConstantComposite %v3float %float_1 %float_1 %float_1
 %_ptr_Function_v3float = OpTypePointer Function %v3float
-         %20 = OpConstantNull %v3float
+         %18 = OpConstantNull %v3float
-         %21 = OpTypeFunction %v4float
+         %19 = OpTypeFunction %v4float
  %max_4883ac = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v3float Function %20
+        %res = OpVariable %_ptr_Function_v3float Function %18
-         %13 = OpExtInst %v3float %15 NMax %17 %17
+               OpStore %res %15
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
+%vertex_main_inner = OpFunction %v4float None %19
-         %23 = OpLabel
+         %21 = OpLabel
-         %24 = OpFunctionCall %void %max_4883ac
+         %22 = OpFunctionCall %void %max_4883ac
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %26 = OpLabel
+         %24 = OpLabel
-         %27 = OpFunctionCall %v4float %vertex_main_inner
+         %25 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %27
+               OpStore %value %25
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %29 = OpLabel
+         %27 = OpLabel
-         %30 = OpFunctionCall %void %max_4883ac
+         %28 = OpFunctionCall %void %max_4883ac
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
+         %30 = OpLabel
-         %33 = OpFunctionCall %void %max_4883ac
+         %31 = OpFunctionCall %void %max_4883ac
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/max/4bbff2.wgsl

@@ -0,0 +1,43 @@
+// Copyright 2022 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/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn max(vec<4, fa>, vec<4, fa>) -> vec<4, fa>
+fn max_4bbff2() {
+  var res = max(vec4(1.), vec4(1.));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  max_4bbff2();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  max_4bbff2();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  max_4bbff2();
+}

test/tint/builtins/gen/literal/max/4bbff2.wgsl.expected.dxc.hlsl

@@ -0,0 +1,30 @@
+void max_4bbff2() {
+  float4 res = (1.0f).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  max_4bbff2();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  max_4bbff2();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  max_4bbff2();
+  return;
+}

test/tint/builtins/gen/literal/max/4bbff2.wgsl.expected.fxc.hlsl

@@ -0,0 +1,30 @@
+void max_4bbff2() {
+  float4 res = (1.0f).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  max_4bbff2();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  max_4bbff2();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  max_4bbff2();
+  return;
+}

test/tint/builtins/gen/literal/max/4bbff2.wgsl.expected.glsl

@@ -0,0 +1,49 @@
+#version 310 es
+
+void max_4bbff2() {
+  vec4 res = vec4(1.0f);
+}
+
+vec4 vertex_main() {
+  max_4bbff2();
+  return vec4(0.0f);
+}
+
+void main() {
+  gl_PointSize = 1.0;
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+void max_4bbff2() {
+  vec4 res = vec4(1.0f);
+}
+
+void fragment_main() {
+  max_4bbff2();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void max_4bbff2() {
+  vec4 res = vec4(1.0f);
+}
+
+void compute_main() {
+  max_4bbff2();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/max/4bbff2.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void max_4bbff2() {
+  float4 res = float4(1.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  max_4bbff2();
+  return float4(0.0f);
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  max_4bbff2();
+  return;
+}
+
+kernel void compute_main() {
+  max_4bbff2();
+  return;
+}
+

test/tint/builtins/gen/literal/max/4bbff2.wgsl.expected.spvasm

@@ -0,0 +1,64 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 30
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %max_4bbff2 "max_4bbff2"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+         %14 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %17 = OpTypeFunction %v4float
+ %max_4bbff2 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_v4float Function %5
+               OpStore %res %14
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %17
+         %19 = OpLabel
+         %20 = OpFunctionCall %void %max_4bbff2
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %22 = OpLabel
+         %23 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %23
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %25 = OpLabel
+         %26 = OpFunctionCall %void %max_4bbff2
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %max_4bbff2
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/max/4bbff2.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn max_4bbff2() {
+  var res = max(vec4(1.0), vec4(1.0));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  max_4bbff2();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  max_4bbff2();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  max_4bbff2();
+}

test/tint/builtins/gen/literal/max/85e6bc.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void max_85e6bc() {
-  int4 res = max((1).xxxx, (1).xxxx);
+  int4 res = (1).xxxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/85e6bc.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void max_85e6bc() {
-  int4 res = max((1).xxxx, (1).xxxx);
+  int4 res = (1).xxxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/85e6bc.wgsl.expected.glsl

@@ -1,7 +1,7 @@
 #version 310 es
 
 void max_85e6bc() {
-  ivec4 res = max(ivec4(1), ivec4(1));
+  ivec4 res = ivec4(1);
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@ void main() {
 precision mediump float;
 
 void max_85e6bc() {
-  ivec4 res = max(ivec4(1), ivec4(1));
+  ivec4 res = ivec4(1);
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@ void main() {
 #version 310 es
 
 void max_85e6bc() {
-  ivec4 res = max(ivec4(1), ivec4(1));
+  ivec4 res = ivec4(1);
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/max/85e6bc.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void max_85e6bc() {
-  int4 res = max(int4(1), int4(1));
+  int4 res = int4(1);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/85e6bc.wgsl.expected.spvasm

@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 34
 ; Schema: 0
                OpCapability Shader
-         %16 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -34,37 +33,36 @@
         %int = OpTypeInt 32 1
       %v4int = OpTypeVector %int 4
       %int_1 = OpConstant %int 1
-         %18 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
+         %16 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
 %_ptr_Function_v4int = OpTypePointer Function %v4int
-         %21 = OpConstantNull %v4int
+         %19 = OpConstantNull %v4int
-         %22 = OpTypeFunction %v4float
+         %20 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %max_85e6bc = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v4int Function %21
+        %res = OpVariable %_ptr_Function_v4int Function %19
-         %13 = OpExtInst %v4int %16 SMax %18 %18
+               OpStore %res %16
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
+%vertex_main_inner = OpFunction %v4float None %20
-         %24 = OpLabel
+         %22 = OpLabel
-         %25 = OpFunctionCall %void %max_85e6bc
+         %23 = OpFunctionCall %void %max_85e6bc
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %27 = OpLabel
+         %25 = OpLabel
-         %28 = OpFunctionCall %v4float %vertex_main_inner
+         %26 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %28
+               OpStore %value %26
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %31 = OpLabel
+         %29 = OpLabel
-         %32 = OpFunctionCall %void %max_85e6bc
+         %30 = OpFunctionCall %void %max_85e6bc
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %34 = OpLabel
+         %32 = OpLabel
-         %35 = OpFunctionCall %void %max_85e6bc
+         %33 = OpFunctionCall %void %max_85e6bc
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/max/a1b196.wgsl

@@ -0,0 +1,43 @@
+// Copyright 2022 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/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn max(vec<3, fa>, vec<3, fa>) -> vec<3, fa>
+fn max_a1b196() {
+  var res = max(vec3(1.), vec3(1.));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  max_a1b196();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  max_a1b196();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  max_a1b196();
+}

test/tint/builtins/gen/literal/max/a1b196.wgsl.expected.dxc.hlsl

@@ -0,0 +1,30 @@
+void max_a1b196() {
+  float3 res = (1.0f).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  max_a1b196();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  max_a1b196();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  max_a1b196();
+  return;
+}

test/tint/builtins/gen/literal/max/a1b196.wgsl.expected.fxc.hlsl

@@ -0,0 +1,30 @@
+void max_a1b196() {
+  float3 res = (1.0f).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  max_a1b196();
+  return (0.0f).xxxx;
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+void fragment_main() {
+  max_a1b196();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  max_a1b196();
+  return;
+}

test/tint/builtins/gen/literal/max/a1b196.wgsl.expected.glsl

@@ -0,0 +1,49 @@
+#version 310 es
+
+void max_a1b196() {
+  vec3 res = vec3(1.0f);
+}
+
+vec4 vertex_main() {
+  max_a1b196();
+  return vec4(0.0f);
+}
+
+void main() {
+  gl_PointSize = 1.0;
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+void max_a1b196() {
+  vec3 res = vec3(1.0f);
+}
+
+void fragment_main() {
+  max_a1b196();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void max_a1b196() {
+  vec3 res = vec3(1.0f);
+}
+
+void compute_main() {
+  max_a1b196();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/max/a1b196.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void max_a1b196() {
+  float3 res = float3(1.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  max_a1b196();
+  return float4(0.0f);
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
+}
+
+fragment void fragment_main() {
+  max_a1b196();
+  return;
+}
+
+kernel void compute_main() {
+  max_a1b196();
+  return;
+}
+

test/tint/builtins/gen/literal/max/a1b196.wgsl.expected.spvasm

@@ -0,0 +1,66 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 32
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %max_a1b196 "max_a1b196"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+    %v3float = OpTypeVector %float 3
+    %float_1 = OpConstant %float 1
+         %15 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %18 = OpConstantNull %v3float
+         %19 = OpTypeFunction %v4float
+ %max_a1b196 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_v3float Function %18
+               OpStore %res %15
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %19
+         %21 = OpLabel
+         %22 = OpFunctionCall %void %max_a1b196
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %24 = OpLabel
+         %25 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %25
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %max_a1b196
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %max_a1b196
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/max/a1b196.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn max_a1b196() {
+  var res = max(vec3(1.0), vec3(1.0));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  max_a1b196();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  max_a1b196();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  max_a1b196();
+}

test/tint/builtins/gen/literal/max/a93419.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void max_a93419() {
-  float4 res = max((1.0f).xxxx, (1.0f).xxxx);
+  float4 res = (1.0f).xxxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/a93419.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void max_a93419() {
-  float4 res = max((1.0f).xxxx, (1.0f).xxxx);
+  float4 res = (1.0f).xxxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/a93419.wgsl.expected.glsl

@@ -1,7 +1,7 @@
 #version 310 es
 
 void max_a93419() {
-  vec4 res = max(vec4(1.0f), vec4(1.0f));
+  vec4 res = vec4(1.0f);
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@ void main() {
 precision mediump float;
 
 void max_a93419() {
-  vec4 res = max(vec4(1.0f), vec4(1.0f));
+  vec4 res = vec4(1.0f);
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@ void main() {
 #version 310 es
 
 void max_a93419() {
-  vec4 res = max(vec4(1.0f), vec4(1.0f));
+  vec4 res = vec4(1.0f);
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/max/a93419.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void max_a93419() {
-  float4 res = fmax(float4(1.0f), float4(1.0f));
+  float4 res = float4(1.0f);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/a93419.wgsl.expected.spvasm

@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 32
+; Bound: 30
 ; Schema: 0
                OpCapability Shader
-         %14 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -32,35 +31,34 @@
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
     %float_1 = OpConstant %float 1
-         %16 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+         %14 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
 %_ptr_Function_v4float = OpTypePointer Function %v4float
-         %19 = OpTypeFunction %v4float
+         %17 = OpTypeFunction %v4float
  %max_a93419 = OpFunction %void None %9
          %12 = OpLabel
         %res = OpVariable %_ptr_Function_v4float Function %5
-         %13 = OpExtInst %v4float %14 NMax %16 %16
+               OpStore %res %14
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %19
+%vertex_main_inner = OpFunction %v4float None %17
-         %21 = OpLabel
+         %19 = OpLabel
-         %22 = OpFunctionCall %void %max_a93419
+         %20 = OpFunctionCall %void %max_a93419
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %24 = OpLabel
+         %22 = OpLabel
-         %25 = OpFunctionCall %v4float %vertex_main_inner
+         %23 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %25
+               OpStore %value %23
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %27 = OpLabel
+         %25 = OpLabel
-         %28 = OpFunctionCall %void %max_a93419
+         %26 = OpFunctionCall %void %max_a93419
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %30 = OpLabel
+         %28 = OpLabel
-         %31 = OpFunctionCall %void %max_a93419
+         %29 = OpFunctionCall %void %max_a93419
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/max/b1b73a.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void max_b1b73a() {
-  uint3 res = max((1u).xxx, (1u).xxx);
+  uint3 res = (1u).xxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/max/b1b73a.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void max_b1b73a() {
-  uint3 res = max((1u).xxx, (1u).xxx);
+  uint3 res = (1u).xxx;
 }
 
 struct tint_symbol {