Implement const eval of abs

This CL adds const-eval for the `abs` builtin.

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

src/tint/intrinsics.def

@@ -404,8 +404,8 @@ match storage
 ////////////////////////////////////////////////////////////////////////////////
 
 // https://gpuweb.github.io/gpuweb/wgsl/#builtin-functions
-fn abs<T: fiu32_f16>(T) -> T
+@const fn abs<T: fia_fiu32_f16>(T) -> T
-fn abs<N: num, T: fiu32_f16>(vec<N, T>) -> vec<N, T>
+@const fn abs<N: num, T: fia_fiu32_f16>(vec<N, T>) -> vec<N, T>
 fn acos<T: f32_f16>(T) -> T
 fn acos<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
 fn acosh<T: f32_f16>(T) -> T

src/tint/resolver/builtin_test.cc

@@ -61,12 +61,10 @@ inline std::ostream& operator<<(std::ostream& out, BuiltinData data) {
 }
 
 TEST_F(ResolverBuiltinTest, ModuleScopeUsage) {
-    GlobalConst("c", ty.f32(), Call(Source{{12, 34}}, "abs", 1._f));
+    GlobalConst("c", ty.f32(), Call(Source{{12, 34}}, "dpdy", 1._f));
 
     EXPECT_FALSE(r()->Resolve());
 
-    // TODO(crbug.com/tint/1581): Once 'abs' is implemented as @const, this will no longer be an
-    // error.
     EXPECT_EQ(
         r()->error(),
         R"(12:34 error: const initializer requires a const-expression, but expression is a runtime-expression)");

src/tint/resolver/const_eval.cc

@@ -1570,6 +1570,31 @@ ConstEval::Result ConstEval::OpShiftLeft(const sem::Type* ty,
     return r;
 }
 
+ConstEval::Result ConstEval::abs(const sem::Type* ty,
+                                 utils::VectorRef<const sem::Constant*> args,
+                                 const Source&) {
+    auto transform = [&](const sem::Constant* c0) {
+        auto create = [&](auto e) {
+            using NumberT = decltype(e);
+            NumberT result;
+            if constexpr (IsUnsignedIntegral<NumberT>) {
+                result = e;
+            } else if constexpr (IsSignedIntegral<NumberT>) {
+                if (e == NumberT::Lowest()) {
+                    result = e;
+                } else {
+                    result = NumberT{std::abs(e)};
+                }
+            } else {
+                result = NumberT{std::abs(e)};
+            }
+            return CreateElement(builder, c0->Type(), result);
+        };
+        return Dispatch_fia_fiu32_f16(create, c0);
+    };
+    return TransformElements(builder, ty, transform, args[0]);
+}
+
 ConstEval::Result ConstEval::all(const sem::Type* ty,
                                  utils::VectorRef<const sem::Constant*> args,
                                  const Source&) {

src/tint/resolver/const_eval.h

@@ -377,6 +377,15 @@ class ConstEval {
     // Builtins
     ////////////////////////////////////////////////////////////////////////////
 
+    /// abs 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 abs(const sem::Type* ty,
+               utils::VectorRef<const sem::Constant*> args,
+               const Source& source);
+
     /// all builtin
     /// @param ty the expression type
     /// @param args the input arguments

src/tint/resolver/const_eval_builtin_test.cc

@@ -146,6 +146,53 @@ INSTANTIATE_TEST_SUITE_P(  //
                          C({1.0_a, 0_a}, kPiOver2<AFloat>),
                      })));
 
+template <typename T, bool finite_only>
+std::vector<Case> AbsCases() {
+    std::vector<Case> cases = {
+        C({T(0)}, T(0)),
+        C({T(2.0)}, T(2.0)),
+        C({T::Highest()}, T::Highest()),
+
+        // Vector tests
+        C({Vec(T(2.0), T::Highest())}, Vec(T(2.0), T::Highest())),
+    };
+
+    ConcatIntoIf<IsSignedIntegral<T>>(
+        cases,
+        std::vector<Case>{
+            C({Negate(T(0))}, T(0)),
+            C({Negate(T(2.0))}, T(2.0)),
+            // If e is signed and is the largest negative, the result is e
+            C({T::Lowest()}, T::Lowest()),
+
+            // 1 more then min i32
+            C({Negate(T(2147483647))}, T(2147483647)),
+
+            C({Vec(T(0), Negate(T(0)))}, Vec(T(0), T(0))),
+            C({Vec(Negate(T(2.0)), T(2.0), T::Highest())}, Vec(T(2.0), T(2.0), T::Highest())),
+        });
+
+    ConcatIntoIf<!finite_only>(cases, std::vector<Case>{
+                                          C({Negate(T::Inf())}, T::Inf()),
+                                          C({T::Inf()}, T::Inf()),
+                                          C({T::NaN()}, T::NaN()),
+                                          C({Vec(Negate(T::Inf()), T::Inf(), T::NaN())},
+                                            Vec(T::Inf(), T::Inf(), T::NaN())),
+                                      });
+
+    return cases;
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Abs,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kAbs),
+                     testing::ValuesIn(Concat(AbsCases<AInt, false>(),  //
+                                              AbsCases<i32, false>(),
+                                              AbsCases<u32, false>(),
+                                              AbsCases<AFloat, true>(),
+                                              AbsCases<f32, false>(),
+                                              AbsCases<f16, false>()))));
+
 static std::vector<Case> AllCases() {
     return {
         C({Val(true)}, Val(true)),

src/tint/resolver/intrinsic_table.inl

@@ -12945,24 +12945,24 @@ constexpr OverloadInfo kOverloads[] = {
     /* num parameters */ 1,
     /* num template types */ 1,
     /* num template numbers */ 0,
-    /* template types */ &kTemplateTypes[25],
+    /* template types */ &kTemplateTypes[28],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[1018],
     /* return matcher indices */ &kMatcherIndices[1],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::abs,
   },
   {
     /* [385] */
     /* num parameters */ 1,
     /* num template types */ 1,
     /* num template numbers */ 1,
-    /* template types */ &kTemplateTypes[25],
+    /* template types */ &kTemplateTypes[28],
     /* template numbers */ &kTemplateNumbers[6],
     /* parameters */ &kParameters[853],
     /* return matcher indices */ &kMatcherIndices[30],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::abs,
   },
   {
     /* [386] */
@@ -14013,8 +14013,8 @@ constexpr OverloadInfo kOverloads[] = {
 constexpr IntrinsicInfo kBuiltins[] = {
   {
     /* [0] */
-    /* fn abs<T : fiu32_f16>(T) -> T */
+    /* fn abs<T : fia_fiu32_f16>(T) -> T */
-    /* fn abs<N : num, T : fiu32_f16>(vec<N, T>) -> vec<N, T> */
+    /* fn abs<N : num, T : fia_fiu32_f16>(vec<N, T>) -> vec<N, T> */
     /* num overloads */ 2,
     /* overloads */ &kOverloads[384],
   },

test/tint/builtins/gen/literal/abs/002533.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/abs/002533.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/abs/002533.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/abs/002533.wgsl.expected.msl

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

test/tint/builtins/gen/literal/abs/002533.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
  %abs_002533 = OpFunction %void None %9
          %12 = OpLabel
         %res = OpVariable %_ptr_Function_v4float Function %5
-         %13 = OpExtInst %v4float %14 FAbs %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 %abs_002533
+         %20 = OpFunctionCall %void %abs_002533
                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 %abs_002533
+         %26 = OpFunctionCall %void %abs_002533
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %30 = OpLabel
+         %28 = OpLabel
-         %31 = OpFunctionCall %void %abs_002533
+         %29 = OpFunctionCall %void %abs_002533
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/005174.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/abs/005174.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/abs/005174.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/abs/005174.wgsl.expected.msl

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

test/tint/builtins/gen/literal/abs/005174.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
  %abs_005174 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v3float Function %20
+        %res = OpVariable %_ptr_Function_v3float Function %18
-         %13 = OpExtInst %v3float %15 FAbs %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 %abs_005174
+         %22 = OpFunctionCall %void %abs_005174
                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 %abs_005174
+         %28 = OpFunctionCall %void %abs_005174
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
+         %30 = OpLabel
-         %33 = OpFunctionCall %void %abs_005174
+         %31 = OpFunctionCall %void %abs_005174
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/1ce782.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/abs/1ce782.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/abs/1ce782.wgsl.expected.msl

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

test/tint/builtins/gen/literal/abs/1ce782.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 35
+; Bound: 34
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -33,36 +33,36 @@
        %uint = OpTypeInt 32 0
      %v4uint = OpTypeVector %uint 4
      %uint_1 = OpConstant %uint 1
-         %17 = 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
-         %20 = OpConstantNull %v4uint
+         %19 = OpConstantNull %v4uint
-         %21 = OpTypeFunction %v4float
+         %20 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %abs_1ce782 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v4uint Function %20
+        %res = OpVariable %_ptr_Function_v4uint Function %19
-               OpStore %res %17
+               OpStore %res %16
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
+%vertex_main_inner = OpFunction %v4float None %20
-         %23 = OpLabel
+         %22 = OpLabel
-         %24 = OpFunctionCall %void %abs_1ce782
+         %23 = OpFunctionCall %void %abs_1ce782
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %26 = OpLabel
+         %25 = OpLabel
-         %27 = OpFunctionCall %v4float %vertex_main_inner
+         %26 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %27
+               OpStore %value %26
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %30 = OpLabel
+         %29 = OpLabel
-         %31 = OpFunctionCall %void %abs_1ce782
+         %30 = OpFunctionCall %void %abs_1ce782
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %33 = OpLabel
+         %32 = OpLabel
-         %34 = OpFunctionCall %void %abs_1ce782
+         %33 = OpFunctionCall %void %abs_1ce782
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/1e9d53.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/abs/1e9d53.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/abs/1e9d53.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/abs/1e9d53.wgsl.expected.msl

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

test/tint/builtins/gen/literal/abs/1e9d53.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
  %abs_1e9d53 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v2float Function %20
+        %res = OpVariable %_ptr_Function_v2float Function %18
-         %13 = OpExtInst %v2float %15 FAbs %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 %abs_1e9d53
+         %22 = OpFunctionCall %void %abs_1e9d53
                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 %abs_1e9d53
+         %28 = OpFunctionCall %void %abs_1e9d53
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
+         %30 = OpLabel
-         %33 = OpFunctionCall %void %abs_1e9d53
+         %31 = OpFunctionCall %void %abs_1e9d53
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/2f861b.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 abs(vec<3, fa>) -> vec<3, fa>
+fn abs_2f861b() {
+  var res = abs(vec3(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  abs_2f861b();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  abs_2f861b();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  abs_2f861b();
+}

test/tint/builtins/gen/literal/abs/2f861b.wgsl.expected.dxc.hlsl

@@ -0,0 +1,30 @@
+void abs_2f861b() {
+  int3 res = (1).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  abs_2f861b();
+  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() {
+  abs_2f861b();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  abs_2f861b();
+  return;
+}

test/tint/builtins/gen/literal/abs/2f861b.wgsl.expected.fxc.hlsl

@@ -0,0 +1,30 @@
+void abs_2f861b() {
+  int3 res = (1).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  abs_2f861b();
+  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() {
+  abs_2f861b();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  abs_2f861b();
+  return;
+}

test/tint/builtins/gen/literal/abs/2f861b.wgsl.expected.glsl

@@ -0,0 +1,49 @@
+#version 310 es
+
+void abs_2f861b() {
+  ivec3 res = ivec3(1);
+}
+
+vec4 vertex_main() {
+  abs_2f861b();
+  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 abs_2f861b() {
+  ivec3 res = ivec3(1);
+}
+
+void fragment_main() {
+  abs_2f861b();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void abs_2f861b() {
+  ivec3 res = ivec3(1);
+}
+
+void compute_main() {
+  abs_2f861b();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/abs/2f861b.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void abs_2f861b() {
+  int3 res = int3(1);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  abs_2f861b();
+  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() {
+  abs_2f861b();
+  return;
+}
+
+kernel void compute_main() {
+  abs_2f861b();
+  return;
+}
+

test/tint/builtins/gen/literal/abs/2f861b.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 %abs_2f861b "abs_2f861b"
+               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
+ %abs_2f861b = 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 %abs_2f861b
+               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 %abs_2f861b
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %abs_2f861b
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/abs/2f861b.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/abs/421ca3.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/abs/421ca3.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/abs/421ca3.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void abs_421ca3() {
-  half3 res = fabs(half3(0.0h));
+  half3 res = half3(0.0h);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/421ca3.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
-         %16 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -37,36 +36,35 @@
           %9 = OpTypeFunction %void
        %half = OpTypeFloat 16
      %v3half = OpTypeVector %half 3
-         %17 = OpConstantNull %v3half
+         %15 = OpConstantNull %v3half
 %_ptr_Function_v3half = OpTypePointer Function %v3half
-         %20 = OpTypeFunction %v4float
+         %18 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %abs_421ca3 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v3half Function %17
+        %res = OpVariable %_ptr_Function_v3half Function %15
-         %13 = OpExtInst %v3half %16 FAbs %17
+               OpStore %res %15
-               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 %abs_421ca3
+         %21 = OpFunctionCall %void %abs_421ca3
                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 %abs_421ca3
+         %28 = OpFunctionCall %void %abs_421ca3
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
+         %30 = OpLabel
-         %33 = OpFunctionCall %void %abs_421ca3
+         %31 = OpFunctionCall %void %abs_421ca3
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/467cd1.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/abs/467cd1.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/abs/467cd1.wgsl.expected.msl

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

test/tint/builtins/gen/literal/abs/467cd1.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 33
+; Bound: 32
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -33,34 +33,34 @@
        %uint = OpTypeInt 32 0
      %uint_1 = OpConstant %uint 1
 %_ptr_Function_uint = OpTypePointer Function %uint
-         %18 = OpConstantNull %uint
+         %17 = OpConstantNull %uint
-         %19 = OpTypeFunction %v4float
+         %18 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %abs_467cd1 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_uint Function %18
+        %res = OpVariable %_ptr_Function_uint Function %17
                OpStore %res %uint_1
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %19
+%vertex_main_inner = OpFunction %v4float None %18
-         %21 = OpLabel
+         %20 = OpLabel
-         %22 = OpFunctionCall %void %abs_467cd1
+         %21 = OpFunctionCall %void %abs_467cd1
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %24 = OpLabel
+         %23 = OpLabel
-         %25 = OpFunctionCall %v4float %vertex_main_inner
+         %24 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %25
+               OpStore %value %24
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %28 = OpLabel
+         %27 = OpLabel
-         %29 = OpFunctionCall %void %abs_467cd1
+         %28 = OpFunctionCall %void %abs_467cd1
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %31 = OpLabel
+         %30 = OpLabel
-         %32 = OpFunctionCall %void %abs_467cd1
+         %31 = OpFunctionCall %void %abs_467cd1
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/4ad288.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void abs_4ad288() {
-  int res = abs(1);
+  int res = 1;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/4ad288.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void abs_4ad288() {
-  int res = abs(1);
+  int res = 1;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/4ad288.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/abs/4ad288.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void abs_4ad288() {
-  int res = abs(1);
+  int res = 1;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/4ad288.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 @@
         %int = OpTypeInt 32 1
       %int_1 = OpConstant %int 1
 %_ptr_Function_int = OpTypePointer Function %int
-         %19 = OpConstantNull %int
+         %17 = OpConstantNull %int
-         %20 = OpTypeFunction %v4float
+         %18 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %abs_4ad288 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_int Function %19
+        %res = OpVariable %_ptr_Function_int Function %17
-         %13 = OpExtInst %int %15 SAbs %int_1
+               OpStore %res %int_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 %abs_4ad288
+         %21 = OpFunctionCall %void %abs_4ad288
                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 %abs_4ad288
+         %28 = OpFunctionCall %void %abs_4ad288
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
+         %30 = OpLabel
-         %33 = OpFunctionCall %void %abs_4ad288
+         %31 = OpFunctionCall %void %abs_4ad288
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/538d29.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void abs_538d29() {
-  vector<float16_t, 4> res = abs((float16_t(0.0h)).xxxx);
+  vector<float16_t, 4> res = (float16_t(0.0h)).xxxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/538d29.wgsl.expected.glsl

@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void abs_538d29() {
-  f16vec4 res = abs(f16vec4(0.0hf));
+  f16vec4 res = f16vec4(0.0hf);
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@ void main() {
 precision mediump float;
 
 void abs_538d29() {
-  f16vec4 res = abs(f16vec4(0.0hf));
+  f16vec4 res = f16vec4(0.0hf);
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@ void main() {
 #extension GL_AMD_gpu_shader_half_float : require
 
 void abs_538d29() {
-  f16vec4 res = abs(f16vec4(0.0hf));
+  f16vec4 res = f16vec4(0.0hf);
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/abs/538d29.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void abs_538d29() {
-  half4 res = fabs(half4(0.0h));
+  half4 res = half4(0.0h);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/538d29.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
-         %16 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -37,36 +36,35 @@
           %9 = OpTypeFunction %void
        %half = OpTypeFloat 16
      %v4half = OpTypeVector %half 4
-         %17 = OpConstantNull %v4half
+         %15 = OpConstantNull %v4half
 %_ptr_Function_v4half = OpTypePointer Function %v4half
-         %20 = OpTypeFunction %v4float
+         %18 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %abs_538d29 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v4half Function %17
+        %res = OpVariable %_ptr_Function_v4half Function %15
-         %13 = OpExtInst %v4half %16 FAbs %17
+               OpStore %res %15
-               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 %abs_538d29
+         %21 = OpFunctionCall %void %abs_538d29
                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 %abs_538d29
+         %28 = OpFunctionCall %void %abs_538d29
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
+         %30 = OpLabel
-         %33 = OpFunctionCall %void %abs_538d29
+         %31 = OpFunctionCall %void %abs_538d29
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/577d6e.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 abs(vec<2, ia>) -> vec<2, ia>
+fn abs_577d6e() {
+  var res = abs(vec2(1.0));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  abs_577d6e();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  abs_577d6e();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  abs_577d6e();
+}

test/tint/builtins/gen/literal/abs/577d6e.wgsl.expected.dxc.hlsl

@@ -0,0 +1,30 @@
+void abs_577d6e() {
+  float2 res = (1.0f).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  abs_577d6e();
+  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() {
+  abs_577d6e();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  abs_577d6e();
+  return;
+}

test/tint/builtins/gen/literal/abs/577d6e.wgsl.expected.fxc.hlsl

@@ -0,0 +1,30 @@
+void abs_577d6e() {
+  float2 res = (1.0f).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  abs_577d6e();
+  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() {
+  abs_577d6e();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  abs_577d6e();
+  return;
+}

test/tint/builtins/gen/literal/abs/577d6e.wgsl.expected.glsl

@@ -0,0 +1,49 @@
+#version 310 es
+
+void abs_577d6e() {
+  vec2 res = vec2(1.0f);
+}
+
+vec4 vertex_main() {
+  abs_577d6e();
+  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 abs_577d6e() {
+  vec2 res = vec2(1.0f);
+}
+
+void fragment_main() {
+  abs_577d6e();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void abs_577d6e() {
+  vec2 res = vec2(1.0f);
+}
+
+void compute_main() {
+  abs_577d6e();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/abs/577d6e.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void abs_577d6e() {
+  float2 res = float2(1.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  abs_577d6e();
+  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() {
+  abs_577d6e();
+  return;
+}
+
+kernel void compute_main() {
+  abs_577d6e();
+  return;
+}
+

test/tint/builtins/gen/literal/abs/577d6e.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 %abs_577d6e "abs_577d6e"
+               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
+    %v2float = OpTypeVector %float 2
+    %float_1 = OpConstant %float 1
+         %15 = OpConstantComposite %v2float %float_1 %float_1
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+         %18 = OpConstantNull %v2float
+         %19 = OpTypeFunction %v4float
+ %abs_577d6e = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_v2float Function %18
+               OpStore %res %15
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %19
+         %21 = OpLabel
+         %22 = OpFunctionCall %void %abs_577d6e
+               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 %abs_577d6e
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %abs_577d6e
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/abs/577d6e.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/abs/5a8af1.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 abs(ia) -> ia
+fn abs_5a8af1() {
+  var res = abs(1.0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  abs_5a8af1();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  abs_5a8af1();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  abs_5a8af1();
+}

test/tint/builtins/gen/literal/abs/5a8af1.wgsl.expected.dxc.hlsl

@@ -0,0 +1,30 @@
+void abs_5a8af1() {
+  float res = 1.0f;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  abs_5a8af1();
+  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() {
+  abs_5a8af1();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  abs_5a8af1();
+  return;
+}

test/tint/builtins/gen/literal/abs/5a8af1.wgsl.expected.fxc.hlsl

@@ -0,0 +1,30 @@
+void abs_5a8af1() {
+  float res = 1.0f;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  abs_5a8af1();
+  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() {
+  abs_5a8af1();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  abs_5a8af1();
+  return;
+}

test/tint/builtins/gen/literal/abs/5a8af1.wgsl.expected.glsl

@@ -0,0 +1,49 @@
+#version 310 es
+
+void abs_5a8af1() {
+  float res = 1.0f;
+}
+
+vec4 vertex_main() {
+  abs_5a8af1();
+  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 abs_5a8af1() {
+  float res = 1.0f;
+}
+
+void fragment_main() {
+  abs_5a8af1();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void abs_5a8af1() {
+  float res = 1.0f;
+}
+
+void compute_main() {
+  abs_5a8af1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/abs/5a8af1.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void abs_5a8af1() {
+  float res = 1.0f;
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  abs_5a8af1();
+  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() {
+  abs_5a8af1();
+  return;
+}
+
+kernel void compute_main() {
+  abs_5a8af1();
+  return;
+}
+

test/tint/builtins/gen/literal/abs/5a8af1.wgsl.expected.spvasm

@@ -0,0 +1,63 @@
+; 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" %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 %abs_5a8af1 "abs_5a8af1"
+               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
+%_ptr_Function_float = OpTypePointer Function %float
+         %16 = OpTypeFunction %v4float
+ %abs_5a8af1 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_float Function %8
+               OpStore %res %float_1
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %16
+         %18 = OpLabel
+         %19 = OpFunctionCall %void %abs_5a8af1
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %21 = OpLabel
+         %22 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %22
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %24 = OpLabel
+         %25 = OpFunctionCall %void %abs_5a8af1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %abs_5a8af1
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/abs/5a8af1.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/abs/5ad50a.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/abs/5ad50a.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/abs/5ad50a.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/abs/5ad50a.wgsl.expected.msl

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

test/tint/builtins/gen/literal/abs/5ad50a.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
  %abs_5ad50a = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v3int Function %21
+        %res = OpVariable %_ptr_Function_v3int Function %19
-         %13 = OpExtInst %v3int %16 SAbs %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 %abs_5ad50a
+         %23 = OpFunctionCall %void %abs_5ad50a
                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 %abs_5ad50a
+         %30 = OpFunctionCall %void %abs_5ad50a
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %34 = OpLabel
+         %32 = OpLabel
-         %35 = OpFunctionCall %void %abs_5ad50a
+         %33 = OpFunctionCall %void %abs_5ad50a
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/5ae4fe.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/abs/5ae4fe.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/abs/5ae4fe.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void abs_5ae4fe() {
-  half2 res = fabs(half2(0.0h));
+  half2 res = half2(0.0h);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/5ae4fe.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
-         %16 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -37,36 +36,35 @@
           %9 = OpTypeFunction %void
        %half = OpTypeFloat 16
      %v2half = OpTypeVector %half 2
-         %17 = OpConstantNull %v2half
+         %15 = OpConstantNull %v2half
 %_ptr_Function_v2half = OpTypePointer Function %v2half
-         %20 = OpTypeFunction %v4float
+         %18 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %abs_5ae4fe = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v2half Function %17
+        %res = OpVariable %_ptr_Function_v2half Function %15
-         %13 = OpExtInst %v2half %16 FAbs %17
+               OpStore %res %15
-               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 %abs_5ae4fe
+         %21 = OpFunctionCall %void %abs_5ae4fe
                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 %abs_5ae4fe
+         %28 = OpFunctionCall %void %abs_5ae4fe
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
+         %30 = OpLabel
-         %33 = OpFunctionCall %void %abs_5ae4fe
+         %31 = OpFunctionCall %void %abs_5ae4fe
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/7326de.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/abs/7326de.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/abs/7326de.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void abs_7326de() {
-  uint3 res = abs(uint3(1u));
+  uint3 res = uint3(1u);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/7326de.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 35
+; Bound: 34
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -33,36 +33,36 @@
        %uint = OpTypeInt 32 0
      %v3uint = OpTypeVector %uint 3
      %uint_1 = OpConstant %uint 1
-         %17 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
+         %16 = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
 %_ptr_Function_v3uint = OpTypePointer Function %v3uint
-         %20 = OpConstantNull %v3uint
+         %19 = OpConstantNull %v3uint
-         %21 = OpTypeFunction %v4float
+         %20 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %abs_7326de = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v3uint Function %20
+        %res = OpVariable %_ptr_Function_v3uint Function %19
-               OpStore %res %17
+               OpStore %res %16
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
+%vertex_main_inner = OpFunction %v4float None %20
-         %23 = OpLabel
+         %22 = OpLabel
-         %24 = OpFunctionCall %void %abs_7326de
+         %23 = OpFunctionCall %void %abs_7326de
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %26 = OpLabel
+         %25 = OpLabel
-         %27 = OpFunctionCall %v4float %vertex_main_inner
+         %26 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %27
+               OpStore %value %26
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %30 = OpLabel
+         %29 = OpLabel
-         %31 = OpFunctionCall %void %abs_7326de
+         %30 = OpFunctionCall %void %abs_7326de
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %33 = OpLabel
+         %32 = OpLabel
-         %34 = OpFunctionCall %void %abs_7326de
+         %33 = OpFunctionCall %void %abs_7326de
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/7f28e6.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/abs/7f28e6.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/abs/7f28e6.wgsl.expected.msl

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

test/tint/builtins/gen/literal/abs/7f28e6.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 35
+; Bound: 34
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -33,36 +33,36 @@
        %uint = OpTypeInt 32 0
      %v2uint = OpTypeVector %uint 2
      %uint_1 = OpConstant %uint 1
-         %17 = OpConstantComposite %v2uint %uint_1 %uint_1
+         %16 = OpConstantComposite %v2uint %uint_1 %uint_1
 %_ptr_Function_v2uint = OpTypePointer Function %v2uint
-         %20 = OpConstantNull %v2uint
+         %19 = OpConstantNull %v2uint
-         %21 = OpTypeFunction %v4float
+         %20 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %abs_7f28e6 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v2uint Function %20
+        %res = OpVariable %_ptr_Function_v2uint Function %19
-               OpStore %res %17
+               OpStore %res %16
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
+%vertex_main_inner = OpFunction %v4float None %20
-         %23 = OpLabel
+         %22 = OpLabel
-         %24 = OpFunctionCall %void %abs_7f28e6
+         %23 = OpFunctionCall %void %abs_7f28e6
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %26 = OpLabel
+         %25 = OpLabel
-         %27 = OpFunctionCall %v4float %vertex_main_inner
+         %26 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %27
+               OpStore %value %26
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %30 = OpLabel
+         %29 = OpLabel
-         %31 = OpFunctionCall %void %abs_7f28e6
+         %30 = OpFunctionCall %void %abs_7f28e6
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %33 = OpLabel
+         %32 = OpLabel
-         %34 = OpFunctionCall %void %abs_7f28e6
+         %33 = OpFunctionCall %void %abs_7f28e6
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/7faa9e.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void abs_7faa9e() {
-  int2 res = abs((1).xx);
+  int2 res = (1).xx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/7faa9e.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void abs_7faa9e() {
-  int2 res = abs((1).xx);
+  int2 res = (1).xx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/7faa9e.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/abs/7faa9e.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void abs_7faa9e() {
-  int2 res = abs(int2(1));
+  int2 res = int2(1);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/abs/7faa9e.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
       %v2int = OpTypeVector %int 2
       %int_1 = OpConstant %int 1
-         %18 = OpConstantComposite %v2int %int_1 %int_1
+         %16 = OpConstantComposite %v2int %int_1 %int_1
 %_ptr_Function_v2int = OpTypePointer Function %v2int
-         %21 = OpConstantNull %v2int
+         %19 = OpConstantNull %v2int
-         %22 = OpTypeFunction %v4float
+         %20 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %abs_7faa9e = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v2int Function %21
+        %res = OpVariable %_ptr_Function_v2int Function %19
-         %13 = OpExtInst %v2int %16 SAbs %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 %abs_7faa9e
+         %23 = OpFunctionCall %void %abs_7faa9e
                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 %abs_7faa9e
+         %30 = OpFunctionCall %void %abs_7faa9e
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %34 = OpLabel
+         %32 = OpLabel
-         %35 = OpFunctionCall %void %abs_7faa9e
+         %33 = OpFunctionCall %void %abs_7faa9e
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/abs/82ff9d.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 abs(vec<2, fa>) -> vec<2, fa>
+fn abs_82ff9d() {
+  var res = abs(vec2(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  abs_82ff9d();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  abs_82ff9d();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  abs_82ff9d();
+}

test/tint/builtins/gen/literal/abs/82ff9d.wgsl.expected.dxc.hlsl

@@ -0,0 +1,30 @@
+void abs_82ff9d() {
+  int2 res = (1).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  abs_82ff9d();
+  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() {
+  abs_82ff9d();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  abs_82ff9d();
+  return;
+}

test/tint/builtins/gen/literal/abs/82ff9d.wgsl.expected.fxc.hlsl

@@ -0,0 +1,30 @@
+void abs_82ff9d() {
+  int2 res = (1).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  abs_82ff9d();
+  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() {
+  abs_82ff9d();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  abs_82ff9d();
+  return;
+}

test/tint/builtins/gen/literal/abs/82ff9d.wgsl.expected.glsl

@@ -0,0 +1,49 @@
+#version 310 es
+
+void abs_82ff9d() {
+  ivec2 res = ivec2(1);
+}
+
+vec4 vertex_main() {
+  abs_82ff9d();
+  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 abs_82ff9d() {
+  ivec2 res = ivec2(1);
+}
+
+void fragment_main() {
+  abs_82ff9d();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void abs_82ff9d() {
+  ivec2 res = ivec2(1);
+}
+
+void compute_main() {
+  abs_82ff9d();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/abs/82ff9d.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void abs_82ff9d() {
+  int2 res = int2(1);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  abs_82ff9d();
+  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() {
+  abs_82ff9d();
+  return;
+}
+
+kernel void compute_main() {
+  abs_82ff9d();
+  return;
+}
+

test/tint/builtins/gen/literal/abs/82ff9d.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 %abs_82ff9d "abs_82ff9d"
+               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
+ %abs_82ff9d = 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 %abs_82ff9d
+               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 %abs_82ff9d
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %abs_82ff9d
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/abs/82ff9d.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/abs/8ca9b1.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 abs(vec<4, ia>) -> vec<4, ia>
+fn abs_8ca9b1() {
+  var res = abs(vec4(1.0));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  abs_8ca9b1();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  abs_8ca9b1();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  abs_8ca9b1();
+}

test/tint/builtins/gen/literal/abs/8ca9b1.wgsl.expected.dxc.hlsl

@@ -0,0 +1,30 @@
+void abs_8ca9b1() {
+  float4 res = (1.0f).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  abs_8ca9b1();
+  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() {
+  abs_8ca9b1();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  abs_8ca9b1();
+  return;
+}

test/tint/builtins/gen/literal/abs/8ca9b1.wgsl.expected.fxc.hlsl

@@ -0,0 +1,30 @@
+void abs_8ca9b1() {
+  float4 res = (1.0f).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  abs_8ca9b1();
+  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() {
+  abs_8ca9b1();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  abs_8ca9b1();
+  return;
+}

test/tint/builtins/gen/literal/abs/8ca9b1.wgsl.expected.glsl

@@ -0,0 +1,49 @@
+#version 310 es
+
+void abs_8ca9b1() {
+  vec4 res = vec4(1.0f);
+}
+
+vec4 vertex_main() {
+  abs_8ca9b1();
+  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 abs_8ca9b1() {
+  vec4 res = vec4(1.0f);
+}
+
+void fragment_main() {
+  abs_8ca9b1();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void abs_8ca9b1() {
+  vec4 res = vec4(1.0f);
+}
+
+void compute_main() {
+  abs_8ca9b1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/abs/8ca9b1.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void abs_8ca9b1() {
+  float4 res = float4(1.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  abs_8ca9b1();
+  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() {
+  abs_8ca9b1();
+  return;
+}
+
+kernel void compute_main() {
+  abs_8ca9b1();
+  return;
+}
+

test/tint/builtins/gen/literal/abs/8ca9b1.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 %abs_8ca9b1 "abs_8ca9b1"
+               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
+ %abs_8ca9b1 = 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 %abs_8ca9b1
+               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 %abs_8ca9b1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %abs_8ca9b1
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/abs/8ca9b1.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/abs/9c80a6.wgsl.expected.dxc.hlsl

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