Const eval for `fma`

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

Bug: tint:1581
Change-Id: Ia4df818fec9d5d969b364b2c165400d787a9e275
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/111584
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index 0752d5a..ba8a581 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -473,8 +473,8 @@
 @const fn firstTrailingBit<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
 @const fn floor<T: fa_f32_f16>(@test_value(1.5) T) -> T
 @const fn floor<N: num, T: fa_f32_f16>(@test_value(1.5) vec<N, T>) -> vec<N, T>
-fn fma<T: f32_f16>(T, T, T) -> T
-fn fma<N: num, T: f32_f16>(vec<N, T>, vec<N, T>, vec<N, T>) -> vec<N, T>
+@const fn fma<T: fa_f32_f16>(T, T, T) -> T
+@const fn fma<N: num, T: fa_f32_f16>(vec<N, T>, vec<N, T>, vec<N, T>) -> vec<N, T>
 fn fract<T: f32_f16>(T) -> T
 fn fract<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
 @const fn frexp<T: fa_f32_f16>(T) -> __frexp_result<T>
diff --git a/src/tint/resolver/const_eval.cc b/src/tint/resolver/const_eval.cc
index 424cd82..0e6f682 100644
--- a/src/tint/resolver/const_eval.cc
+++ b/src/tint/resolver/const_eval.cc
@@ -2446,6 +2446,33 @@
     return TransformElements(builder, ty, transform, args[0]);
 }
 
+ConstEval::Result ConstEval::fma(const sem::Type* ty,
+                                 utils::VectorRef<const sem::Constant*> args,
+                                 const Source& source) {
+    auto transform = [&](const sem::Constant* c1, const sem::Constant* c2,
+                         const sem::Constant* c3) {
+        auto create = [&](auto e1, auto e2, auto e3) -> ImplResult {
+            auto err_msg = [&] {
+                AddNote("when calculating fma", source);
+                return utils::Failure;
+            };
+
+            auto mul = Mul(source, e1, e2);
+            if (!mul) {
+                return err_msg();
+            }
+
+            auto val = Add(source, mul.Get(), e3);
+            if (!val) {
+                return err_msg();
+            }
+            return CreateElement(builder, source, c1->Type(), val.Get());
+        };
+        return Dispatch_fa_f32_f16(create, c1, c2, c3);
+    };
+    return TransformElements(builder, ty, transform, args[0], args[1], args[2]);
+}
+
 ConstEval::Result ConstEval::frexp(const sem::Type* ty,
                                    utils::VectorRef<const sem::Constant*> args,
                                    const Source& source) {
diff --git a/src/tint/resolver/const_eval.h b/src/tint/resolver/const_eval.h
index 33ec0a4..b65bb28 100644
--- a/src/tint/resolver/const_eval.h
+++ b/src/tint/resolver/const_eval.h
@@ -629,6 +629,15 @@
                  utils::VectorRef<const sem::Constant*> args,
                  const Source& source);
 
+    /// fma builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location
+    /// @return the result value, or null if the value cannot be calculated
+    Result fma(const sem::Type* ty,
+               utils::VectorRef<const sem::Constant*> args,
+               const Source& source);
+
     /// frexp builtin
     /// @param ty the expression type
     /// @param args the input arguments
diff --git a/src/tint/resolver/const_eval_builtin_test.cc b/src/tint/resolver/const_eval_builtin_test.cc
index 9c55323..849f3ef 100644
--- a/src/tint/resolver/const_eval_builtin_test.cc
+++ b/src/tint/resolver/const_eval_builtin_test.cc
@@ -1072,6 +1072,31 @@
                      testing::ValuesIn(Concat(FloorCases<AFloat>(),  //
                                               FloorCases<f32>(),
                                               FloorCases<f16>()))));
+
+template <typename T>
+std::vector<Case> FmaCases() {
+    auto error_msg = [](auto a, const char* op, auto b) {
+        return "12:34 error: " + OverflowErrorMessage(a, op, b) + R"(
+12:34 note: when calculating fma)";
+    };
+    return {
+        C({T(0), T(0), T(0)}, T(0)),
+        C({T(1), T(2), T(3)}, T(5)),
+        C({Vec(T(1), T(2.5), -T(1)), Vec(T(2), T(2.5), T(1)), Vec(T(4), T(3.75), -T(2))},
+          Vec(T(6), T(10), -T(3))),
+
+        E({T::Highest(), T::Highest(), T(0)}, error_msg(T::Highest(), "*", T::Highest())),
+        E({T::Highest(), T(1), T::Highest()}, error_msg(T::Highest(), "+", T::Highest())),
+    };
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Fma,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kFma),
+                     testing::ValuesIn(Concat(FmaCases<AFloat>(),  //
+                                              FmaCases<f32>(),
+                                              FmaCases<f16>()))));
+
 template <typename T>
 std::vector<Case> FrexpCases() {
     using F = T;                                                         // fract type
diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl
index e213f81..1a2634f 100644
--- a/src/tint/resolver/intrinsic_table.inl
+++ b/src/tint/resolver/intrinsic_table.inl
@@ -12510,24 +12510,24 @@
     /* num parameters */ 3,
     /* num template types */ 1,
     /* num template numbers */ 0,
-    /* template types */ &kTemplateTypes[26],
+    /* template types */ &kTemplateTypes[23],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[462],
     /* return matcher indices */ &kMatcherIndices[3],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::fma,
   },
   {
     /* [350] */
     /* num parameters */ 3,
     /* num template types */ 1,
     /* num template numbers */ 1,
-    /* template types */ &kTemplateTypes[26],
+    /* template types */ &kTemplateTypes[23],
     /* template numbers */ &kTemplateNumbers[4],
     /* parameters */ &kParameters[465],
     /* return matcher indices */ &kMatcherIndices[30],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::fma,
   },
   {
     /* [351] */
@@ -14245,8 +14245,8 @@
   },
   {
     /* [38] */
-    /* fn fma<T : f32_f16>(T, T, T) -> T */
-    /* fn fma<N : num, T : f32_f16>(vec<N, T>, vec<N, T>, vec<N, T>) -> vec<N, T> */
+    /* fn fma<T : fa_f32_f16>(T, T, T) -> T */
+    /* fn fma<N : num, T : fa_f32_f16>(vec<N, T>, vec<N, T>, vec<N, T>) -> vec<N, T> */
     /* num overloads */ 2,
     /* overloads */ &kOverloads[349],
   },
diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl b/test/tint/builtins/gen/literal/fma/143d5d.wgsl
new file mode 100644
index 0000000..5a83670
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/143d5d.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 fma(vec<4, fa>, vec<4, fa>, vec<4, fa>) -> vec<4, fa>
+fn fma_143d5d() {
+  var res = fma(vec4(1.), vec4(1.), vec4(1.));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_143d5d();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_143d5d();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_143d5d();
+}
diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..6c3ad71
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void fma_143d5d() {
+  float4 res = (2.0f).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_143d5d();
+  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() {
+  fma_143d5d();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_143d5d();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..6c3ad71
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void fma_143d5d() {
+  float4 res = (2.0f).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_143d5d();
+  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() {
+  fma_143d5d();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_143d5d();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.glsl
new file mode 100644
index 0000000..7e6b4b4
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void fma_143d5d() {
+  vec4 res = vec4(2.0f);
+}
+
+vec4 vertex_main() {
+  fma_143d5d();
+  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 fma_143d5d() {
+  vec4 res = vec4(2.0f);
+}
+
+void fragment_main() {
+  fma_143d5d();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void fma_143d5d() {
+  vec4 res = vec4(2.0f);
+}
+
+void compute_main() {
+  fma_143d5d();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.msl
new file mode 100644
index 0000000..af72335
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void fma_143d5d() {
+  float4 res = float4(2.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  fma_143d5d();
+  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() {
+  fma_143d5d();
+  return;
+}
+
+kernel void compute_main() {
+  fma_143d5d();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.spvasm
new file mode 100644
index 0000000..fba8c32
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.spvasm
@@ -0,0 +1,65 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %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 %fma_143d5d "fma_143d5d"
+               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_2 = OpConstant %float 2
+         %14 = OpConstantComposite %v4float %float_2 %float_2 %float_2 %float_2
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %17 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %fma_143d5d = 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 %fma_143d5d
+               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
+         %26 = OpLabel
+         %27 = OpFunctionCall %void %fma_143d5d
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %fma_143d5d
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.wgsl
new file mode 100644
index 0000000..887606f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn fma_143d5d() {
+  var res = fma(vec4(1.0), vec4(1.0), vec4(1.0));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_143d5d();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_143d5d();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_143d5d();
+}
diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl b/test/tint/builtins/gen/literal/fma/1f5084.wgsl
new file mode 100644
index 0000000..3453e11
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/1f5084.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 fma(vec<2, fa>, vec<2, fa>, vec<2, fa>) -> vec<2, fa>
+fn fma_1f5084() {
+  var res = fma(vec2(1.), vec2(1.), vec2(1.));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_1f5084();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_1f5084();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_1f5084();
+}
diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..e7655ab
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void fma_1f5084() {
+  float2 res = (2.0f).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_1f5084();
+  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() {
+  fma_1f5084();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_1f5084();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..e7655ab
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void fma_1f5084() {
+  float2 res = (2.0f).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_1f5084();
+  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() {
+  fma_1f5084();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_1f5084();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.glsl
new file mode 100644
index 0000000..215c740
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void fma_1f5084() {
+  vec2 res = vec2(2.0f);
+}
+
+vec4 vertex_main() {
+  fma_1f5084();
+  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 fma_1f5084() {
+  vec2 res = vec2(2.0f);
+}
+
+void fragment_main() {
+  fma_1f5084();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void fma_1f5084() {
+  vec2 res = vec2(2.0f);
+}
+
+void compute_main() {
+  fma_1f5084();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.msl
new file mode 100644
index 0000000..bc3b0f3
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void fma_1f5084() {
+  float2 res = float2(2.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  fma_1f5084();
+  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() {
+  fma_1f5084();
+  return;
+}
+
+kernel void compute_main() {
+  fma_1f5084();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.spvasm
new file mode 100644
index 0000000..3c5ab1b
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.spvasm
@@ -0,0 +1,67 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %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 %fma_1f5084 "fma_1f5084"
+               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_2 = OpConstant %float 2
+         %15 = OpConstantComposite %v2float %float_2 %float_2
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+         %18 = OpConstantNull %v2float
+         %19 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %fma_1f5084 = 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 %fma_1f5084
+               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
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %fma_1f5084
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %fma_1f5084
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.wgsl
new file mode 100644
index 0000000..116db76
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn fma_1f5084() {
+  var res = fma(vec2(1.0), vec2(1.0), vec2(1.0));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_1f5084();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_1f5084();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_1f5084();
+}
diff --git a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.dxc.hlsl
index d58d1db..7f4c278 100644
--- a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void fma_26a7a9() {
-  float2 res = mad((1.0f).xx, (1.0f).xx, (1.0f).xx);
+  float2 res = (2.0f).xx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.fxc.hlsl
index d58d1db..7f4c278 100644
--- a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.fxc.hlsl
@@ -1,5 +1,5 @@
 void fma_26a7a9() {
-  float2 res = mad((1.0f).xx, (1.0f).xx, (1.0f).xx);
+  float2 res = (2.0f).xx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.glsl
index 1905d16..ca63145 100644
--- a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.glsl
@@ -1,7 +1,7 @@
 #version 310 es
 
 void fma_26a7a9() {
-  vec2 res = ((vec2(1.0f)) * (vec2(1.0f)) + (vec2(1.0f)));
+  vec2 res = vec2(2.0f);
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@
 precision mediump float;
 
 void fma_26a7a9() {
-  vec2 res = ((vec2(1.0f)) * (vec2(1.0f)) + (vec2(1.0f)));
+  vec2 res = vec2(2.0f);
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@
 #version 310 es
 
 void fma_26a7a9() {
-  vec2 res = ((vec2(1.0f)) * (vec2(1.0f)) + (vec2(1.0f)));
+  vec2 res = vec2(2.0f);
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.msl
index 98fe644..77a1973 100644
--- a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void fma_26a7a9() {
-  float2 res = fma(float2(1.0f), float2(1.0f), float2(1.0f));
+  float2 res = float2(2.0f);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.spvasm
index 2e2385e..63647d4 100644
--- a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.spvasm
@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 33
 ; 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"
@@ -32,37 +31,37 @@
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
     %v2float = OpTypeVector %float 2
-    %float_1 = OpConstant %float 1
-         %17 = OpConstantComposite %v2float %float_1 %float_1
+    %float_2 = OpConstant %float 2
+         %15 = OpConstantComposite %v2float %float_2 %float_2
 %_ptr_Function_v2float = OpTypePointer Function %v2float
-         %20 = OpConstantNull %v2float
-         %21 = OpTypeFunction %v4float
+         %18 = OpConstantNull %v2float
+         %19 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
  %fma_26a7a9 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v2float Function %20
-         %13 = OpExtInst %v2float %15 Fma %17 %17 %17
-               OpStore %res %13
+        %res = OpVariable %_ptr_Function_v2float Function %18
+               OpStore %res %15
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
-         %23 = OpLabel
-         %24 = OpFunctionCall %void %fma_26a7a9
+%vertex_main_inner = OpFunction %v4float None %19
+         %21 = OpLabel
+         %22 = OpFunctionCall %void %fma_26a7a9
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %26 = OpLabel
-         %27 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %27
+         %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
-         %29 = OpLabel
-         %30 = OpFunctionCall %void %fma_26a7a9
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %fma_26a7a9
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %fma_26a7a9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %fma_26a7a9
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl b/test/tint/builtins/gen/literal/fma/466442.wgsl
new file mode 100644
index 0000000..e46749a
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/466442.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 fma(fa, fa, fa) -> fa
+fn fma_466442() {
+  var res = fma(1., 1., 1.);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_466442();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_466442();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_466442();
+}
diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..c070a74
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void fma_466442() {
+  float res = 2.0f;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_466442();
+  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() {
+  fma_466442();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_466442();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..c070a74
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void fma_466442() {
+  float res = 2.0f;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_466442();
+  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() {
+  fma_466442();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_466442();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.glsl
new file mode 100644
index 0000000..c094238
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void fma_466442() {
+  float res = 2.0f;
+}
+
+vec4 vertex_main() {
+  fma_466442();
+  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 fma_466442() {
+  float res = 2.0f;
+}
+
+void fragment_main() {
+  fma_466442();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void fma_466442() {
+  float res = 2.0f;
+}
+
+void compute_main() {
+  fma_466442();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.msl
new file mode 100644
index 0000000..26b2b96
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void fma_466442() {
+  float res = 2.0f;
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  fma_466442();
+  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() {
+  fma_466442();
+  return;
+}
+
+kernel void compute_main() {
+  fma_466442();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.spvasm
new file mode 100644
index 0000000..55e5e1f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/466442.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 %fma_466442 "fma_466442"
+               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_2 = OpConstant %float 2
+%_ptr_Function_float = OpTypePointer Function %float
+         %16 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %fma_466442 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_float Function %8
+               OpStore %res %float_2
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %16
+         %18 = OpLabel
+         %19 = OpFunctionCall %void %fma_466442
+               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
+         %25 = OpLabel
+         %26 = OpFunctionCall %void %fma_466442
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %fma_466442
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.wgsl
new file mode 100644
index 0000000..cae52f3
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn fma_466442() {
+  var res = fma(1.0, 1.0, 1.0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_466442();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_466442();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_466442();
+}
diff --git a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.dxc.hlsl
index 8959dfa..a518076 100644
--- a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void fma_6a3283() {
-  float4 res = mad((1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx);
+  float4 res = (2.0f).xxxx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.fxc.hlsl
index 8959dfa..a518076 100644
--- a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.fxc.hlsl
@@ -1,5 +1,5 @@
 void fma_6a3283() {
-  float4 res = mad((1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx);
+  float4 res = (2.0f).xxxx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.glsl
index 96bc9f2..f18280d 100644
--- a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.glsl
@@ -1,7 +1,7 @@
 #version 310 es
 
 void fma_6a3283() {
-  vec4 res = ((vec4(1.0f)) * (vec4(1.0f)) + (vec4(1.0f)));
+  vec4 res = vec4(2.0f);
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@
 precision mediump float;
 
 void fma_6a3283() {
-  vec4 res = ((vec4(1.0f)) * (vec4(1.0f)) + (vec4(1.0f)));
+  vec4 res = vec4(2.0f);
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@
 #version 310 es
 
 void fma_6a3283() {
-  vec4 res = ((vec4(1.0f)) * (vec4(1.0f)) + (vec4(1.0f)));
+  vec4 res = vec4(2.0f);
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.msl
index 6c675c1..747f0d6 100644
--- a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void fma_6a3283() {
-  float4 res = fma(float4(1.0f), float4(1.0f), float4(1.0f));
+  float4 res = float4(2.0f);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.spvasm
index 197b06b..2f1ceb1 100644
--- a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.spvasm
@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 32
+; Bound: 31
 ; 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"
@@ -31,36 +30,36 @@
 %vertex_point_size = OpVariable %_ptr_Output_float Output %8
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
-    %float_1 = OpConstant %float 1
-         %16 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+    %float_2 = OpConstant %float 2
+         %14 = OpConstantComposite %v4float %float_2 %float_2 %float_2 %float_2
 %_ptr_Function_v4float = OpTypePointer Function %v4float
-         %19 = OpTypeFunction %v4float
+         %17 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
  %fma_6a3283 = OpFunction %void None %9
          %12 = OpLabel
         %res = OpVariable %_ptr_Function_v4float Function %5
-         %13 = OpExtInst %v4float %14 Fma %16 %16 %16
-               OpStore %res %13
+               OpStore %res %14
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %19
-         %21 = OpLabel
-         %22 = OpFunctionCall %void %fma_6a3283
+%vertex_main_inner = OpFunction %v4float None %17
+         %19 = OpLabel
+         %20 = OpFunctionCall %void %fma_6a3283
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %24 = OpLabel
-         %25 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %25
+         %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
-         %27 = OpLabel
-         %28 = OpFunctionCall %void %fma_6a3283
+         %26 = OpLabel
+         %27 = OpFunctionCall %void %fma_6a3283
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %30 = OpLabel
-         %31 = OpFunctionCall %void %fma_6a3283
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %fma_6a3283
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.dxc.hlsl
index 3336b31..aaafc9a 100644
--- a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void fma_ab7818() {
-  vector<float16_t, 4> res = mad((float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx);
+  vector<float16_t, 4> res = (float16_t(2.0h)).xxxx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.glsl
index fa4e760..034f24a 100644
--- a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.glsl
@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void fma_ab7818() {
-  f16vec4 res = ((f16vec4(1.0hf)) * (f16vec4(1.0hf)) + (f16vec4(1.0hf)));
+  f16vec4 res = f16vec4(2.0hf);
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@
 precision mediump float;
 
 void fma_ab7818() {
-  f16vec4 res = ((f16vec4(1.0hf)) * (f16vec4(1.0hf)) + (f16vec4(1.0hf)));
+  f16vec4 res = f16vec4(2.0hf);
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void fma_ab7818() {
-  f16vec4 res = ((f16vec4(1.0hf)) * (f16vec4(1.0hf)) + (f16vec4(1.0hf)));
+  f16vec4 res = f16vec4(2.0hf);
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.msl
index eb2484d..77bd22d 100644
--- a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void fma_ab7818() {
-  half4 res = fma(half4(1.0h), half4(1.0h), half4(1.0h));
+  half4 res = half4(2.0h);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.spvasm
index ccf40ea..b278ae3 100644
--- a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/fma/ab7818.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"
@@ -37,38 +36,37 @@
           %9 = OpTypeFunction %void
        %half = OpTypeFloat 16
      %v4half = OpTypeVector %half 4
-%half_0x1p_0 = OpConstant %half 0x1p+0
-         %18 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+%half_0x1p_1 = OpConstant %half 0x1p+1
+         %16 = OpConstantComposite %v4half %half_0x1p_1 %half_0x1p_1 %half_0x1p_1 %half_0x1p_1
 %_ptr_Function_v4half = OpTypePointer Function %v4half
-         %21 = OpConstantNull %v4half
-         %22 = OpTypeFunction %v4float
+         %19 = OpConstantNull %v4half
+         %20 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %fma_ab7818 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v4half Function %21
-         %13 = OpExtInst %v4half %16 Fma %18 %18 %18
-               OpStore %res %13
+        %res = OpVariable %_ptr_Function_v4half Function %19
+               OpStore %res %16
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
-         %24 = OpLabel
-         %25 = OpFunctionCall %void %fma_ab7818
+%vertex_main_inner = OpFunction %v4float None %20
+         %22 = OpLabel
+         %23 = OpFunctionCall %void %fma_ab7818
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %27 = OpLabel
-         %28 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %28
+         %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
-         %31 = OpLabel
-         %32 = OpFunctionCall %void %fma_ab7818
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %fma_ab7818
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %fma_ab7818
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %fma_ab7818
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.dxc.hlsl
index 9e48041..846fd33 100644
--- a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void fma_bf21b6() {
-  vector<float16_t, 2> res = mad((float16_t(1.0h)).xx, (float16_t(1.0h)).xx, (float16_t(1.0h)).xx);
+  vector<float16_t, 2> res = (float16_t(2.0h)).xx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.glsl
index 627c7c9..fb633c8 100644
--- a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.glsl
@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void fma_bf21b6() {
-  f16vec2 res = ((f16vec2(1.0hf)) * (f16vec2(1.0hf)) + (f16vec2(1.0hf)));
+  f16vec2 res = f16vec2(2.0hf);
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@
 precision mediump float;
 
 void fma_bf21b6() {
-  f16vec2 res = ((f16vec2(1.0hf)) * (f16vec2(1.0hf)) + (f16vec2(1.0hf)));
+  f16vec2 res = f16vec2(2.0hf);
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void fma_bf21b6() {
-  f16vec2 res = ((f16vec2(1.0hf)) * (f16vec2(1.0hf)) + (f16vec2(1.0hf)));
+  f16vec2 res = f16vec2(2.0hf);
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.msl
index 54c227b..69f9158 100644
--- a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void fma_bf21b6() {
-  half2 res = fma(half2(1.0h), half2(1.0h), half2(1.0h));
+  half2 res = half2(2.0h);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.spvasm
index 6b33758..d0c5c44 100644
--- a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/fma/bf21b6.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"
@@ -37,38 +36,37 @@
           %9 = OpTypeFunction %void
        %half = OpTypeFloat 16
      %v2half = OpTypeVector %half 2
-%half_0x1p_0 = OpConstant %half 0x1p+0
-         %18 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
+%half_0x1p_1 = OpConstant %half 0x1p+1
+         %16 = OpConstantComposite %v2half %half_0x1p_1 %half_0x1p_1
 %_ptr_Function_v2half = OpTypePointer Function %v2half
-         %21 = OpConstantNull %v2half
-         %22 = OpTypeFunction %v4float
+         %19 = OpConstantNull %v2half
+         %20 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %fma_bf21b6 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v2half Function %21
-         %13 = OpExtInst %v2half %16 Fma %18 %18 %18
-               OpStore %res %13
+        %res = OpVariable %_ptr_Function_v2half Function %19
+               OpStore %res %16
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
-         %24 = OpLabel
-         %25 = OpFunctionCall %void %fma_bf21b6
+%vertex_main_inner = OpFunction %v4float None %20
+         %22 = OpLabel
+         %23 = OpFunctionCall %void %fma_bf21b6
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %27 = OpLabel
-         %28 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %28
+         %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
-         %31 = OpLabel
-         %32 = OpFunctionCall %void %fma_bf21b6
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %fma_bf21b6
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %fma_bf21b6
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %fma_bf21b6
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.dxc.hlsl
index 88e5bb2..b8ef7d6 100644
--- a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void fma_c10ba3() {
-  float res = mad(1.0f, 1.0f, 1.0f);
+  float res = 2.0f;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.fxc.hlsl
index 88e5bb2..b8ef7d6 100644
--- a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.fxc.hlsl
@@ -1,5 +1,5 @@
 void fma_c10ba3() {
-  float res = mad(1.0f, 1.0f, 1.0f);
+  float res = 2.0f;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.glsl
index 611c0b6..009c216 100644
--- a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.glsl
@@ -1,7 +1,7 @@
 #version 310 es
 
 void fma_c10ba3() {
-  float res = ((1.0f) * (1.0f) + (1.0f));
+  float res = 2.0f;
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@
 precision mediump float;
 
 void fma_c10ba3() {
-  float res = ((1.0f) * (1.0f) + (1.0f));
+  float res = 2.0f;
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@
 #version 310 es
 
 void fma_c10ba3() {
-  float res = ((1.0f) * (1.0f) + (1.0f));
+  float res = 2.0f;
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.msl
index 9f60018..e222ebd 100644
--- a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void fma_c10ba3() {
-  float res = fma(1.0f, 1.0f, 1.0f);
+  float res = 2.0f;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.spvasm
index bc9816f..5335567 100644
--- a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.spvasm
@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 31
+; 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"
@@ -31,35 +30,35 @@
 %vertex_point_size = OpVariable %_ptr_Output_float Output %8
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
-    %float_1 = OpConstant %float 1
+    %float_2 = OpConstant %float 2
 %_ptr_Function_float = OpTypePointer Function %float
-         %18 = OpTypeFunction %v4float
+         %16 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
  %fma_c10ba3 = OpFunction %void None %9
          %12 = OpLabel
         %res = OpVariable %_ptr_Function_float Function %8
-         %13 = OpExtInst %float %14 Fma %float_1 %float_1 %float_1
-               OpStore %res %13
+               OpStore %res %float_2
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %18
-         %20 = OpLabel
-         %21 = OpFunctionCall %void %fma_c10ba3
+%vertex_main_inner = OpFunction %v4float None %16
+         %18 = OpLabel
+         %19 = OpFunctionCall %void %fma_c10ba3
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %23 = OpLabel
-         %24 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %24
+         %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
-         %26 = OpLabel
-         %27 = OpFunctionCall %void %fma_c10ba3
+         %25 = OpLabel
+         %26 = OpFunctionCall %void %fma_c10ba3
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %29 = OpLabel
-         %30 = OpFunctionCall %void %fma_c10ba3
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %fma_c10ba3
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.dxc.hlsl
index f6e79c5..ad0ae1e 100644
--- a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void fma_c8abb3() {
-  float16_t res = mad(float16_t(1.0h), float16_t(1.0h), float16_t(1.0h));
+  float16_t res = float16_t(2.0h);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.glsl
index 8540f07..172c792 100644
--- a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.glsl
@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void fma_c8abb3() {
-  float16_t res = ((1.0hf) * (1.0hf) + (1.0hf));
+  float16_t res = 2.0hf;
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@
 precision mediump float;
 
 void fma_c8abb3() {
-  float16_t res = ((1.0hf) * (1.0hf) + (1.0hf));
+  float16_t res = 2.0hf;
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void fma_c8abb3() {
-  float16_t res = ((1.0hf) * (1.0hf) + (1.0hf));
+  float16_t res = 2.0hf;
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.msl
index 1152af9..58ba67a 100644
--- a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void fma_c8abb3() {
-  half res = fma(1.0h, 1.0h, 1.0h);
+  half res = 2.0h;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.spvasm
index 8bde975..16b175f 100644
--- a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/fma/c8abb3.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"
@@ -36,37 +35,36 @@
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
        %half = OpTypeFloat 16
-%half_0x1p_0 = OpConstant %half 0x1p+0
+%half_0x1p_1 = OpConstant %half 0x1p+1
 %_ptr_Function_half = OpTypePointer Function %half
-         %19 = OpConstantNull %half
-         %20 = OpTypeFunction %v4float
+         %17 = OpConstantNull %half
+         %18 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %fma_c8abb3 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_half Function %19
-         %13 = OpExtInst %half %15 Fma %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
-               OpStore %res %13
+        %res = OpVariable %_ptr_Function_half Function %17
+               OpStore %res %half_0x1p_1
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
-         %22 = OpLabel
-         %23 = OpFunctionCall %void %fma_c8abb3
+%vertex_main_inner = OpFunction %v4float None %18
+         %20 = OpLabel
+         %21 = OpFunctionCall %void %fma_c8abb3
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %25 = OpLabel
-         %26 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %26
+         %23 = OpLabel
+         %24 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %24
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %29 = OpLabel
-         %30 = OpFunctionCall %void %fma_c8abb3
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %fma_c8abb3
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %fma_c8abb3
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %fma_c8abb3
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.dxc.hlsl
index 89cd7e8..ba5da32 100644
--- a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void fma_e17c5c() {
-  float3 res = mad((1.0f).xxx, (1.0f).xxx, (1.0f).xxx);
+  float3 res = (2.0f).xxx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.fxc.hlsl
index 89cd7e8..ba5da32 100644
--- a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.fxc.hlsl
@@ -1,5 +1,5 @@
 void fma_e17c5c() {
-  float3 res = mad((1.0f).xxx, (1.0f).xxx, (1.0f).xxx);
+  float3 res = (2.0f).xxx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.glsl
index d1b1d2e..7285005 100644
--- a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.glsl
@@ -1,7 +1,7 @@
 #version 310 es
 
 void fma_e17c5c() {
-  vec3 res = ((vec3(1.0f)) * (vec3(1.0f)) + (vec3(1.0f)));
+  vec3 res = vec3(2.0f);
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@
 precision mediump float;
 
 void fma_e17c5c() {
-  vec3 res = ((vec3(1.0f)) * (vec3(1.0f)) + (vec3(1.0f)));
+  vec3 res = vec3(2.0f);
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@
 #version 310 es
 
 void fma_e17c5c() {
-  vec3 res = ((vec3(1.0f)) * (vec3(1.0f)) + (vec3(1.0f)));
+  vec3 res = vec3(2.0f);
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.msl
index 469605f..24c7af4 100644
--- a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void fma_e17c5c() {
-  float3 res = fma(float3(1.0f), float3(1.0f), float3(1.0f));
+  float3 res = float3(2.0f);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.spvasm
index f1d385b..87bcc0e 100644
--- a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.spvasm
@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 33
 ; 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"
@@ -32,37 +31,37 @@
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
     %v3float = OpTypeVector %float 3
-    %float_1 = OpConstant %float 1
-         %17 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+    %float_2 = OpConstant %float 2
+         %15 = OpConstantComposite %v3float %float_2 %float_2 %float_2
 %_ptr_Function_v3float = OpTypePointer Function %v3float
-         %20 = OpConstantNull %v3float
-         %21 = OpTypeFunction %v4float
+         %18 = OpConstantNull %v3float
+         %19 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
  %fma_e17c5c = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v3float Function %20
-         %13 = OpExtInst %v3float %15 Fma %17 %17 %17
-               OpStore %res %13
+        %res = OpVariable %_ptr_Function_v3float Function %18
+               OpStore %res %15
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
-         %23 = OpLabel
-         %24 = OpFunctionCall %void %fma_e17c5c
+%vertex_main_inner = OpFunction %v4float None %19
+         %21 = OpLabel
+         %22 = OpFunctionCall %void %fma_e17c5c
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %26 = OpLabel
-         %27 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %27
+         %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
-         %29 = OpLabel
-         %30 = OpFunctionCall %void %fma_e17c5c
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %fma_e17c5c
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %fma_e17c5c
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %fma_e17c5c
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.dxc.hlsl
index 4930e9e..7229fbb 100644
--- a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void fma_e7abdc() {
-  vector<float16_t, 3> res = mad((float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx);
+  vector<float16_t, 3> res = (float16_t(2.0h)).xxx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.glsl
index 182adee..3e06476 100644
--- a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.glsl
@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void fma_e7abdc() {
-  f16vec3 res = ((f16vec3(1.0hf)) * (f16vec3(1.0hf)) + (f16vec3(1.0hf)));
+  f16vec3 res = f16vec3(2.0hf);
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@
 precision mediump float;
 
 void fma_e7abdc() {
-  f16vec3 res = ((f16vec3(1.0hf)) * (f16vec3(1.0hf)) + (f16vec3(1.0hf)));
+  f16vec3 res = f16vec3(2.0hf);
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void fma_e7abdc() {
-  f16vec3 res = ((f16vec3(1.0hf)) * (f16vec3(1.0hf)) + (f16vec3(1.0hf)));
+  f16vec3 res = f16vec3(2.0hf);
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.msl
index 2360c4d..cd1a833 100644
--- a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void fma_e7abdc() {
-  half3 res = fma(half3(1.0h), half3(1.0h), half3(1.0h));
+  half3 res = half3(2.0h);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.spvasm
index b2d0b4c..0af3de3 100644
--- a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/fma/e7abdc.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"
@@ -37,38 +36,37 @@
           %9 = OpTypeFunction %void
        %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
+%half_0x1p_1 = OpConstant %half 0x1p+1
+         %16 = OpConstantComposite %v3half %half_0x1p_1 %half_0x1p_1 %half_0x1p_1
 %_ptr_Function_v3half = OpTypePointer Function %v3half
-         %21 = OpConstantNull %v3half
-         %22 = OpTypeFunction %v4float
+         %19 = OpConstantNull %v3half
+         %20 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
  %fma_e7abdc = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v3half Function %21
-         %13 = OpExtInst %v3half %16 Fma %18 %18 %18
-               OpStore %res %13
+        %res = OpVariable %_ptr_Function_v3half Function %19
+               OpStore %res %16
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
-         %24 = OpLabel
-         %25 = OpFunctionCall %void %fma_e7abdc
+%vertex_main_inner = OpFunction %v4float None %20
+         %22 = OpLabel
+         %23 = OpFunctionCall %void %fma_e7abdc
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %27 = OpLabel
-         %28 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %28
+         %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
-         %31 = OpLabel
-         %32 = OpFunctionCall %void %fma_e7abdc
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %fma_e7abdc
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %fma_e7abdc
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %fma_e7abdc
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl
new file mode 100644
index 0000000..e136fa3
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/eb25d7.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 fma(vec<3, fa>, vec<3, fa>, vec<3, fa>) -> vec<3, fa>
+fn fma_eb25d7() {
+  var res = fma(vec3(1.), vec3(1.), vec3(1.));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_eb25d7();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_eb25d7();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_eb25d7();
+}
diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..7036fc6
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void fma_eb25d7() {
+  float3 res = (2.0f).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_eb25d7();
+  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() {
+  fma_eb25d7();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_eb25d7();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..7036fc6
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void fma_eb25d7() {
+  float3 res = (2.0f).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_eb25d7();
+  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() {
+  fma_eb25d7();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_eb25d7();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.glsl
new file mode 100644
index 0000000..16ddde5
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void fma_eb25d7() {
+  vec3 res = vec3(2.0f);
+}
+
+vec4 vertex_main() {
+  fma_eb25d7();
+  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 fma_eb25d7() {
+  vec3 res = vec3(2.0f);
+}
+
+void fragment_main() {
+  fma_eb25d7();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void fma_eb25d7() {
+  vec3 res = vec3(2.0f);
+}
+
+void compute_main() {
+  fma_eb25d7();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.msl
new file mode 100644
index 0000000..5d32a93
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void fma_eb25d7() {
+  float3 res = float3(2.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  fma_eb25d7();
+  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() {
+  fma_eb25d7();
+  return;
+}
+
+kernel void compute_main() {
+  fma_eb25d7();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.spvasm
new file mode 100644
index 0000000..a3f813e
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.spvasm
@@ -0,0 +1,67 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %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 %fma_eb25d7 "fma_eb25d7"
+               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_2 = OpConstant %float 2
+         %15 = OpConstantComposite %v3float %float_2 %float_2 %float_2
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %18 = OpConstantNull %v3float
+         %19 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %fma_eb25d7 = 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 %fma_eb25d7
+               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
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %fma_eb25d7
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %fma_eb25d7
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.wgsl
new file mode 100644
index 0000000..7a380d4
--- /dev/null
+++ b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn fma_eb25d7() {
+  var res = fma(vec3(1.0), vec3(1.0), vec3(1.0));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_eb25d7();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_eb25d7();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_eb25d7();
+}
diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl b/test/tint/builtins/gen/var/fma/143d5d.wgsl
new file mode 100644
index 0000000..124c6a7
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl
@@ -0,0 +1,46 @@
+// 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 fma(vec<4, fa>, vec<4, fa>, vec<4, fa>) -> vec<4, fa>
+fn fma_143d5d() {
+  const arg_0 = vec4(1.);
+  const arg_1 = vec4(1.);
+  const arg_2 = vec4(1.);
+  var res = fma(arg_0, arg_1, arg_2);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_143d5d();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_143d5d();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_143d5d();
+}
diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..6c3ad71
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void fma_143d5d() {
+  float4 res = (2.0f).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_143d5d();
+  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() {
+  fma_143d5d();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_143d5d();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..6c3ad71
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void fma_143d5d() {
+  float4 res = (2.0f).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_143d5d();
+  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() {
+  fma_143d5d();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_143d5d();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.glsl b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.glsl
new file mode 100644
index 0000000..7e6b4b4
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void fma_143d5d() {
+  vec4 res = vec4(2.0f);
+}
+
+vec4 vertex_main() {
+  fma_143d5d();
+  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 fma_143d5d() {
+  vec4 res = vec4(2.0f);
+}
+
+void fragment_main() {
+  fma_143d5d();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void fma_143d5d() {
+  vec4 res = vec4(2.0f);
+}
+
+void compute_main() {
+  fma_143d5d();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.msl b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.msl
new file mode 100644
index 0000000..af72335
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void fma_143d5d() {
+  float4 res = float4(2.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  fma_143d5d();
+  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() {
+  fma_143d5d();
+  return;
+}
+
+kernel void compute_main() {
+  fma_143d5d();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.spvasm b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.spvasm
new file mode 100644
index 0000000..fba8c32
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.spvasm
@@ -0,0 +1,65 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %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 %fma_143d5d "fma_143d5d"
+               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_2 = OpConstant %float 2
+         %14 = OpConstantComposite %v4float %float_2 %float_2 %float_2 %float_2
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %17 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %fma_143d5d = 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 %fma_143d5d
+               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
+         %26 = OpLabel
+         %27 = OpFunctionCall %void %fma_143d5d
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %fma_143d5d
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.wgsl b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.wgsl
new file mode 100644
index 0000000..03e6426
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+fn fma_143d5d() {
+  const arg_0 = vec4(1.0);
+  const arg_1 = vec4(1.0);
+  const arg_2 = vec4(1.0);
+  var res = fma(arg_0, arg_1, arg_2);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_143d5d();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_143d5d();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_143d5d();
+}
diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl b/test/tint/builtins/gen/var/fma/1f5084.wgsl
new file mode 100644
index 0000000..20fb093
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl
@@ -0,0 +1,46 @@
+// 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 fma(vec<2, fa>, vec<2, fa>, vec<2, fa>) -> vec<2, fa>
+fn fma_1f5084() {
+  const arg_0 = vec2(1.);
+  const arg_1 = vec2(1.);
+  const arg_2 = vec2(1.);
+  var res = fma(arg_0, arg_1, arg_2);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_1f5084();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_1f5084();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_1f5084();
+}
diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..e7655ab
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void fma_1f5084() {
+  float2 res = (2.0f).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_1f5084();
+  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() {
+  fma_1f5084();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_1f5084();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..e7655ab
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void fma_1f5084() {
+  float2 res = (2.0f).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_1f5084();
+  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() {
+  fma_1f5084();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_1f5084();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.glsl b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.glsl
new file mode 100644
index 0000000..215c740
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void fma_1f5084() {
+  vec2 res = vec2(2.0f);
+}
+
+vec4 vertex_main() {
+  fma_1f5084();
+  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 fma_1f5084() {
+  vec2 res = vec2(2.0f);
+}
+
+void fragment_main() {
+  fma_1f5084();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void fma_1f5084() {
+  vec2 res = vec2(2.0f);
+}
+
+void compute_main() {
+  fma_1f5084();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.msl b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.msl
new file mode 100644
index 0000000..bc3b0f3
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void fma_1f5084() {
+  float2 res = float2(2.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  fma_1f5084();
+  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() {
+  fma_1f5084();
+  return;
+}
+
+kernel void compute_main() {
+  fma_1f5084();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.spvasm b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.spvasm
new file mode 100644
index 0000000..3c5ab1b
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.spvasm
@@ -0,0 +1,67 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %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 %fma_1f5084 "fma_1f5084"
+               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_2 = OpConstant %float 2
+         %15 = OpConstantComposite %v2float %float_2 %float_2
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+         %18 = OpConstantNull %v2float
+         %19 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %fma_1f5084 = 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 %fma_1f5084
+               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
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %fma_1f5084
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %fma_1f5084
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.wgsl b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.wgsl
new file mode 100644
index 0000000..a339a36
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+fn fma_1f5084() {
+  const arg_0 = vec2(1.0);
+  const arg_1 = vec2(1.0);
+  const arg_2 = vec2(1.0);
+  var res = fma(arg_0, arg_1, arg_2);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_1f5084();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_1f5084();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_1f5084();
+}
diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl b/test/tint/builtins/gen/var/fma/466442.wgsl
new file mode 100644
index 0000000..ef43d65
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/466442.wgsl
@@ -0,0 +1,46 @@
+// 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 fma(fa, fa, fa) -> fa
+fn fma_466442() {
+  const arg_0 = 1.;
+  const arg_1 = 1.;
+  const arg_2 = 1.;
+  var res = fma(arg_0, arg_1, arg_2);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_466442();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_466442();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_466442();
+}
diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..c070a74
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void fma_466442() {
+  float res = 2.0f;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_466442();
+  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() {
+  fma_466442();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_466442();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..c070a74
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void fma_466442() {
+  float res = 2.0f;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_466442();
+  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() {
+  fma_466442();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_466442();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl.expected.glsl b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.glsl
new file mode 100644
index 0000000..c094238
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void fma_466442() {
+  float res = 2.0f;
+}
+
+vec4 vertex_main() {
+  fma_466442();
+  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 fma_466442() {
+  float res = 2.0f;
+}
+
+void fragment_main() {
+  fma_466442();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void fma_466442() {
+  float res = 2.0f;
+}
+
+void compute_main() {
+  fma_466442();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl.expected.msl b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.msl
new file mode 100644
index 0000000..26b2b96
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void fma_466442() {
+  float res = 2.0f;
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  fma_466442();
+  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() {
+  fma_466442();
+  return;
+}
+
+kernel void compute_main() {
+  fma_466442();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl.expected.spvasm b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.spvasm
new file mode 100644
index 0000000..55e5e1f
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/466442.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 %fma_466442 "fma_466442"
+               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_2 = OpConstant %float 2
+%_ptr_Function_float = OpTypePointer Function %float
+         %16 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %fma_466442 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_float Function %8
+               OpStore %res %float_2
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %16
+         %18 = OpLabel
+         %19 = OpFunctionCall %void %fma_466442
+               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
+         %25 = OpLabel
+         %26 = OpFunctionCall %void %fma_466442
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %fma_466442
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl.expected.wgsl b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.wgsl
new file mode 100644
index 0000000..1ae85a0
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+fn fma_466442() {
+  const arg_0 = 1.0;
+  const arg_1 = 1.0;
+  const arg_2 = 1.0;
+  var res = fma(arg_0, arg_1, arg_2);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_466442();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_466442();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_466442();
+}
diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl b/test/tint/builtins/gen/var/fma/eb25d7.wgsl
new file mode 100644
index 0000000..4f56414
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl
@@ -0,0 +1,46 @@
+// 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 fma(vec<3, fa>, vec<3, fa>, vec<3, fa>) -> vec<3, fa>
+fn fma_eb25d7() {
+  const arg_0 = vec3(1.);
+  const arg_1 = vec3(1.);
+  const arg_2 = vec3(1.);
+  var res = fma(arg_0, arg_1, arg_2);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_eb25d7();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_eb25d7();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_eb25d7();
+}
diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..7036fc6
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void fma_eb25d7() {
+  float3 res = (2.0f).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_eb25d7();
+  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() {
+  fma_eb25d7();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_eb25d7();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..7036fc6
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void fma_eb25d7() {
+  float3 res = (2.0f).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  fma_eb25d7();
+  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() {
+  fma_eb25d7();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  fma_eb25d7();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.glsl b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.glsl
new file mode 100644
index 0000000..16ddde5
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void fma_eb25d7() {
+  vec3 res = vec3(2.0f);
+}
+
+vec4 vertex_main() {
+  fma_eb25d7();
+  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 fma_eb25d7() {
+  vec3 res = vec3(2.0f);
+}
+
+void fragment_main() {
+  fma_eb25d7();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void fma_eb25d7() {
+  vec3 res = vec3(2.0f);
+}
+
+void compute_main() {
+  fma_eb25d7();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.msl b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.msl
new file mode 100644
index 0000000..5d32a93
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void fma_eb25d7() {
+  float3 res = float3(2.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  fma_eb25d7();
+  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() {
+  fma_eb25d7();
+  return;
+}
+
+kernel void compute_main() {
+  fma_eb25d7();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.spvasm b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.spvasm
new file mode 100644
index 0000000..a3f813e
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.spvasm
@@ -0,0 +1,67 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %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 %fma_eb25d7 "fma_eb25d7"
+               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_2 = OpConstant %float 2
+         %15 = OpConstantComposite %v3float %float_2 %float_2 %float_2
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %18 = OpConstantNull %v3float
+         %19 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+ %fma_eb25d7 = 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 %fma_eb25d7
+               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
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %fma_eb25d7
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %fma_eb25d7
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.wgsl b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.wgsl
new file mode 100644
index 0000000..216968e
--- /dev/null
+++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+fn fma_eb25d7() {
+  const arg_0 = vec3(1.0);
+  const arg_1 = vec3(1.0);
+  const arg_2 = vec3(1.0);
+  var res = fma(arg_0, arg_1, arg_2);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  fma_eb25d7();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  fma_eb25d7();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  fma_eb25d7();
+}