diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index 8744257..e0656cb 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -425,8 +425,8 @@
 @const fn atan2<T: fa_f32_f16, N: num>(vec<N, T>, vec<N, T>) -> vec<N, T>
 @const fn atanh<T: fa_f32_f16>(@test_value(0.5) T) -> T
 @const fn atanh<N: num, T: fa_f32_f16>(@test_value(0.5) vec<N, T>) -> vec<N, T>
-fn ceil<T: f32_f16>(T) -> T
-fn ceil<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
+@const fn ceil<T: fa_f32_f16>(@test_value(1.5) T) -> T
+@const fn ceil<N: num, T: fa_f32_f16>(@test_value(1.5) vec<N, T>) -> vec<N, T>
 @const fn clamp<T: fia_fiu32_f16>(T, T, T) -> T
 @const fn clamp<T: fia_fiu32_f16, N: num>(vec<N, T>, vec<N, T>, vec<N, T>) -> vec<N, T>
 fn cos<T: f32_f16>(T) -> T
diff --git a/src/tint/resolver/const_eval.cc b/src/tint/resolver/const_eval.cc
index 9681961..f32d807 100644
--- a/src/tint/resolver/const_eval.cc
+++ b/src/tint/resolver/const_eval.cc
@@ -1700,6 +1700,18 @@
     return TransformElements(builder, ty, transform, args[0], args[1]);
 }
 
+ConstEval::Result ConstEval::ceil(const sem::Type* ty,
+                                  utils::VectorRef<const sem::Constant*> args,
+                                  const Source&) {
+    auto transform = [&](const sem::Constant* c0) {
+        auto create = [&](auto e) {
+            return CreateElement(builder, c0->Type(), decltype(e)(std::ceil(e)));
+        };
+        return Dispatch_fa_f32_f16(create, c0);
+    };
+    return TransformElements(builder, ty, transform, args[0]);
+}
+
 ConstEval::Result ConstEval::clamp(const sem::Type* ty,
                                    utils::VectorRef<const sem::Constant*> args,
                                    const Source&) {
diff --git a/src/tint/resolver/const_eval.h b/src/tint/resolver/const_eval.h
index 943f0ad..3f83c71 100644
--- a/src/tint/resolver/const_eval.h
+++ b/src/tint/resolver/const_eval.h
@@ -458,6 +458,15 @@
                  utils::VectorRef<const sem::Constant*> args,
                  const Source& source);
 
+    /// ceil builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    Result ceil(const sem::Type* ty,
+                utils::VectorRef<const sem::Constant*> args,
+                const Source& source);
+
     /// clamp 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 0a8978f..4943e3d 100644
--- a/src/tint/resolver/const_eval_builtin_test.cc
+++ b/src/tint/resolver/const_eval_builtin_test.cc
@@ -594,6 +594,37 @@
                                               AsinhCases<f32, false>(),
                                               AsinhCases<f16, false>()))));
 
+template <typename T, bool finite_only>
+std::vector<Case> CeilCases() {
+    std::vector<Case> cases = {
+        C({T(0)}, T(0)),
+        C({-T(0)}, -T(0)),
+        C({-T(1.5)}, -T(1.0)),
+        C({T(1.5)}, T(2.0)),
+        C({T::Lowest()}, T::Lowest()),
+        C({T::Highest()}, T::Highest()),
+
+        C({Vec(T(0), T(1.5), -T(1.5))}, Vec(T(0), T(2.0), -T(1.0))),
+    };
+
+    ConcatIntoIf<!finite_only>(
+        cases, std::vector<Case>{
+                   C({-T::Inf()}, -T::Inf()),
+                   C({T::Inf()}, T::Inf()),
+                   C({T::NaN()}, T::NaN()),
+                   C({Vec(-T::Inf(), T::Inf(), T::NaN())}, Vec(-T::Inf(), T::Inf(), T::NaN())),
+               });
+
+    return cases;
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Ceil,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kCeil),
+                     testing::ValuesIn(Concat(CeilCases<AFloat, true>(),
+                                              CeilCases<f32, false>(),
+                                              CeilCases<f16, false>()))));
+
 template <typename T>
 std::vector<Case> ClampCases() {
     return {
diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl
index c270b5b..ec6c4bf 100644
--- a/src/tint/resolver/intrinsic_table.inl
+++ b/src/tint/resolver/intrinsic_table.inl
@@ -13161,24 +13161,24 @@
     /* num parameters */ 1,
     /* num template types */ 1,
     /* num template numbers */ 0,
-    /* template types */ &kTemplateTypes[22],
+    /* template types */ &kTemplateTypes[23],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[978],
     /* return matcher indices */ &kMatcherIndices[1],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::ceil,
   },
   {
     /* [403] */
     /* num parameters */ 1,
     /* num template types */ 1,
     /* num template numbers */ 1,
-    /* template types */ &kTemplateTypes[22],
+    /* template types */ &kTemplateTypes[23],
     /* template numbers */ &kTemplateNumbers[6],
     /* parameters */ &kParameters[977],
     /* return matcher indices */ &kMatcherIndices[30],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::ceil,
   },
   {
     /* [404] */
@@ -14089,8 +14089,8 @@
   },
   {
     /* [11] */
-    /* fn ceil<T : f32_f16>(T) -> T */
-    /* fn ceil<N : num, T : f32_f16>(vec<N, T>) -> vec<N, T> */
+    /* fn ceil<T : fa_f32_f16>(@test_value(1.5) T) -> T */
+    /* fn ceil<N : num, T : fa_f32_f16>(@test_value(1.5) vec<N, T>) -> vec<N, T> */
     /* num overloads */ 2,
     /* overloads */ &kOverloads[402],
   },
diff --git a/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.dxc.hlsl
index a2d17ac..7bd42e4 100644
--- a/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_09bf52() {
-  vector<float16_t, 3> res = ceil((float16_t(0.0h)).xxx);
+  vector<float16_t, 3> res = (float16_t(0.0h)).xxx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.glsl b/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.glsl
index 5a83910..48a58ee 100644
--- a/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.glsl
@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void ceil_09bf52() {
-  f16vec3 res = ceil(f16vec3(0.0hf));
+  f16vec3 res = f16vec3(0.0hf);
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@
 precision mediump float;
 
 void ceil_09bf52() {
-  f16vec3 res = ceil(f16vec3(0.0hf));
+  f16vec3 res = f16vec3(0.0hf);
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void ceil_09bf52() {
-  f16vec3 res = ceil(f16vec3(0.0hf));
+  f16vec3 res = f16vec3(0.0hf);
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.msl b/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.msl
index 7d0d688..8ff7786 100644
--- a/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ceil_09bf52() {
-  half3 res = ceil(half3(0.0h));
+  half3 res = half3(0.0h);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.spvasm
index 2b87105..926482e 100644
--- a/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/ceil/09bf52.wgsl.expected.spvasm
@@ -1,14 +1,13 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 32
 ; Schema: 0
                OpCapability Shader
                OpCapability Float16
                OpCapability UniformAndStorageBuffer16BitAccess
                OpCapability StorageBuffer16BitAccess
                OpCapability StorageInputOutput16
-         %16 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -37,36 +36,35 @@
           %9 = OpTypeFunction %void
        %half = OpTypeFloat 16
      %v3half = OpTypeVector %half 3
-         %17 = OpConstantNull %v3half
+         %15 = OpConstantNull %v3half
 %_ptr_Function_v3half = OpTypePointer Function %v3half
-         %20 = OpTypeFunction %v4float
+         %18 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
 %ceil_09bf52 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v3half Function %17
-         %13 = OpExtInst %v3half %16 Ceil %17
-               OpStore %res %13
+        %res = OpVariable %_ptr_Function_v3half Function %15
+               OpStore %res %15
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
-         %22 = OpLabel
-         %23 = OpFunctionCall %void %ceil_09bf52
+%vertex_main_inner = OpFunction %v4float None %18
+         %20 = OpLabel
+         %21 = OpFunctionCall %void %ceil_09bf52
                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 %ceil_09bf52
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %ceil_09bf52
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %ceil_09bf52
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %ceil_09bf52
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl b/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl
new file mode 100644
index 0000000..bea2b7e
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/11b1dc.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 ceil(vec<4, fa>) -> vec<4, fa>
+fn ceil_11b1dc() {
+  var res = ceil(vec4(1.5));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_11b1dc();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_11b1dc();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_11b1dc();
+}
diff --git a/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..05f5861
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_11b1dc() {
+  float4 res = (2.0f).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_11b1dc();
+  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() {
+  ceil_11b1dc();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_11b1dc();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..05f5861
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_11b1dc() {
+  float4 res = (2.0f).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_11b1dc();
+  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() {
+  ceil_11b1dc();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_11b1dc();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.glsl b/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.glsl
new file mode 100644
index 0000000..7a6447b
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void ceil_11b1dc() {
+  vec4 res = vec4(2.0f);
+}
+
+vec4 vertex_main() {
+  ceil_11b1dc();
+  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 ceil_11b1dc() {
+  vec4 res = vec4(2.0f);
+}
+
+void fragment_main() {
+  ceil_11b1dc();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void ceil_11b1dc() {
+  vec4 res = vec4(2.0f);
+}
+
+void compute_main() {
+  ceil_11b1dc();
+}
+
+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/ceil/11b1dc.wgsl.expected.msl b/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.msl
new file mode 100644
index 0000000..4da2a34
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void ceil_11b1dc() {
+  float4 res = float4(2.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  ceil_11b1dc();
+  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() {
+  ceil_11b1dc();
+  return;
+}
+
+kernel void compute_main() {
+  ceil_11b1dc();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.spvasm
new file mode 100644
index 0000000..924f17b
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/11b1dc.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 %ceil_11b1dc "ceil_11b1dc"
+               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
+%ceil_11b1dc = 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 %ceil_11b1dc
+               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 %ceil_11b1dc
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %ceil_11b1dc
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.wgsl
new file mode 100644
index 0000000..a5190ae
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/11b1dc.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn ceil_11b1dc() {
+  var res = ceil(vec4(1.5));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_11b1dc();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_11b1dc();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_11b1dc();
+}
diff --git a/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.dxc.hlsl
index 0397f16..7d4e24b 100644
--- a/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_18c240() {
-  vector<float16_t, 2> res = ceil((float16_t(0.0h)).xx);
+  vector<float16_t, 2> res = (float16_t(0.0h)).xx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.glsl b/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.glsl
index b4375a2..b63b07a 100644
--- a/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.glsl
@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void ceil_18c240() {
-  f16vec2 res = ceil(f16vec2(0.0hf));
+  f16vec2 res = f16vec2(0.0hf);
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@
 precision mediump float;
 
 void ceil_18c240() {
-  f16vec2 res = ceil(f16vec2(0.0hf));
+  f16vec2 res = f16vec2(0.0hf);
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void ceil_18c240() {
-  f16vec2 res = ceil(f16vec2(0.0hf));
+  f16vec2 res = f16vec2(0.0hf);
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.msl b/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.msl
index b656d4c..9e480ef 100644
--- a/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ceil_18c240() {
-  half2 res = ceil(half2(0.0h));
+  half2 res = half2(0.0h);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.spvasm
index 757b0b4..7fe9ff4 100644
--- a/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/ceil/18c240.wgsl.expected.spvasm
@@ -1,14 +1,13 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 32
 ; Schema: 0
                OpCapability Shader
                OpCapability Float16
                OpCapability UniformAndStorageBuffer16BitAccess
                OpCapability StorageBuffer16BitAccess
                OpCapability StorageInputOutput16
-         %16 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -37,36 +36,35 @@
           %9 = OpTypeFunction %void
        %half = OpTypeFloat 16
      %v2half = OpTypeVector %half 2
-         %17 = OpConstantNull %v2half
+         %15 = OpConstantNull %v2half
 %_ptr_Function_v2half = OpTypePointer Function %v2half
-         %20 = OpTypeFunction %v4float
+         %18 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
 %ceil_18c240 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v2half Function %17
-         %13 = OpExtInst %v2half %16 Ceil %17
-               OpStore %res %13
+        %res = OpVariable %_ptr_Function_v2half Function %15
+               OpStore %res %15
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
-         %22 = OpLabel
-         %23 = OpFunctionCall %void %ceil_18c240
+%vertex_main_inner = OpFunction %v4float None %18
+         %20 = OpLabel
+         %21 = OpFunctionCall %void %ceil_18c240
                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 %ceil_18c240
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %ceil_18c240
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %ceil_18c240
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %ceil_18c240
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/ceil/32c946.wgsl b/test/tint/builtins/gen/literal/ceil/32c946.wgsl
new file mode 100644
index 0000000..28a3c26
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/32c946.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 ceil(vec<3, fa>) -> vec<3, fa>
+fn ceil_32c946() {
+  var res = ceil(vec3(1.5));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_32c946();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_32c946();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_32c946();
+}
diff --git a/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..fcd472c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_32c946() {
+  float3 res = (2.0f).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_32c946();
+  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() {
+  ceil_32c946();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_32c946();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..fcd472c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_32c946() {
+  float3 res = (2.0f).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_32c946();
+  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() {
+  ceil_32c946();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_32c946();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.glsl b/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.glsl
new file mode 100644
index 0000000..c129c84
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void ceil_32c946() {
+  vec3 res = vec3(2.0f);
+}
+
+vec4 vertex_main() {
+  ceil_32c946();
+  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 ceil_32c946() {
+  vec3 res = vec3(2.0f);
+}
+
+void fragment_main() {
+  ceil_32c946();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void ceil_32c946() {
+  vec3 res = vec3(2.0f);
+}
+
+void compute_main() {
+  ceil_32c946();
+}
+
+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/ceil/32c946.wgsl.expected.msl b/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.msl
new file mode 100644
index 0000000..1096249
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void ceil_32c946() {
+  float3 res = float3(2.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  ceil_32c946();
+  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() {
+  ceil_32c946();
+  return;
+}
+
+kernel void compute_main() {
+  ceil_32c946();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.spvasm
new file mode 100644
index 0000000..2711442
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/32c946.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 %ceil_32c946 "ceil_32c946"
+               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
+%ceil_32c946 = 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 %ceil_32c946
+               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 %ceil_32c946
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ceil_32c946
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.wgsl
new file mode 100644
index 0000000..02d48bd
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/32c946.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn ceil_32c946() {
+  var res = ceil(vec3(1.5));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_32c946();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_32c946();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_32c946();
+}
diff --git a/test/tint/builtins/gen/literal/ceil/34064b.wgsl b/test/tint/builtins/gen/literal/ceil/34064b.wgsl
index 33e3a50..553052f 100644
--- a/test/tint/builtins/gen/literal/ceil/34064b.wgsl
+++ b/test/tint/builtins/gen/literal/ceil/34064b.wgsl
@@ -23,7 +23,7 @@
 
 // fn ceil(vec<3, f32>) -> vec<3, f32>
 fn ceil_34064b() {
-  var res: vec3<f32> = ceil(vec3<f32>(1.f));
+  var res: vec3<f32> = ceil(vec3<f32>(1.5f));
 }
 
 @vertex
diff --git a/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.dxc.hlsl
index 2e8c4ce..9033b8d 100644
--- a/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_34064b() {
-  float3 res = ceil((1.0f).xxx);
+  float3 res = (2.0f).xxx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.fxc.hlsl
index 2e8c4ce..9033b8d 100644
--- a/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.fxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_34064b() {
-  float3 res = ceil((1.0f).xxx);
+  float3 res = (2.0f).xxx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.glsl b/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.glsl
index 35f80ff..ab05efa 100644
--- a/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.glsl
@@ -1,7 +1,7 @@
 #version 310 es
 
 void ceil_34064b() {
-  vec3 res = ceil(vec3(1.0f));
+  vec3 res = vec3(2.0f);
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@
 precision mediump float;
 
 void ceil_34064b() {
-  vec3 res = ceil(vec3(1.0f));
+  vec3 res = vec3(2.0f);
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@
 #version 310 es
 
 void ceil_34064b() {
-  vec3 res = ceil(vec3(1.0f));
+  vec3 res = vec3(2.0f);
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.msl b/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.msl
index 0840f02..2de699a 100644
--- a/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ceil_34064b() {
-  float3 res = ceil(float3(1.0f));
+  float3 res = float3(2.0f);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.spvasm
index 018e998..5538037 100644
--- a/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/ceil/34064b.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
 %ceil_34064b = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v3float Function %20
-         %13 = OpExtInst %v3float %15 Ceil %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 %ceil_34064b
+%vertex_main_inner = OpFunction %v4float None %19
+         %21 = OpLabel
+         %22 = OpFunctionCall %void %ceil_34064b
                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 %ceil_34064b
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ceil_34064b
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %ceil_34064b
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ceil_34064b
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.wgsl
index 9c572bc..c88e88e 100644
--- a/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/ceil/34064b.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
 fn ceil_34064b() {
-  var res : vec3<f32> = ceil(vec3<f32>(1.0f));
+  var res : vec3<f32> = ceil(vec3<f32>(1.5f));
 }
 
 @vertex
diff --git a/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.dxc.hlsl
index 10c1d29..ba26727 100644
--- a/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_4bca2a() {
-  vector<float16_t, 4> res = ceil((float16_t(0.0h)).xxxx);
+  vector<float16_t, 4> res = (float16_t(0.0h)).xxxx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.glsl b/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.glsl
index 4792e79..ff7d3d7 100644
--- a/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.glsl
@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void ceil_4bca2a() {
-  f16vec4 res = ceil(f16vec4(0.0hf));
+  f16vec4 res = f16vec4(0.0hf);
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@
 precision mediump float;
 
 void ceil_4bca2a() {
-  f16vec4 res = ceil(f16vec4(0.0hf));
+  f16vec4 res = f16vec4(0.0hf);
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void ceil_4bca2a() {
-  f16vec4 res = ceil(f16vec4(0.0hf));
+  f16vec4 res = f16vec4(0.0hf);
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.msl b/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.msl
index 63fa1e1..046f92d 100644
--- a/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ceil_4bca2a() {
-  half4 res = ceil(half4(0.0h));
+  half4 res = half4(0.0h);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.spvasm
index 65f288b..abbd917 100644
--- a/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/ceil/4bca2a.wgsl.expected.spvasm
@@ -1,14 +1,13 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 32
 ; Schema: 0
                OpCapability Shader
                OpCapability Float16
                OpCapability UniformAndStorageBuffer16BitAccess
                OpCapability StorageBuffer16BitAccess
                OpCapability StorageInputOutput16
-         %16 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -37,36 +36,35 @@
           %9 = OpTypeFunction %void
        %half = OpTypeFloat 16
      %v4half = OpTypeVector %half 4
-         %17 = OpConstantNull %v4half
+         %15 = OpConstantNull %v4half
 %_ptr_Function_v4half = OpTypePointer Function %v4half
-         %20 = OpTypeFunction %v4float
+         %18 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
 %ceil_4bca2a = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v4half Function %17
-         %13 = OpExtInst %v4half %16 Ceil %17
-               OpStore %res %13
+        %res = OpVariable %_ptr_Function_v4half Function %15
+               OpStore %res %15
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %20
-         %22 = OpLabel
-         %23 = OpFunctionCall %void %ceil_4bca2a
+%vertex_main_inner = OpFunction %v4float None %18
+         %20 = OpLabel
+         %21 = OpFunctionCall %void %ceil_4bca2a
                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 %ceil_4bca2a
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %ceil_4bca2a
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %ceil_4bca2a
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %ceil_4bca2a
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/ceil/678655.wgsl b/test/tint/builtins/gen/literal/ceil/678655.wgsl
index a4d3256..2f81f6d 100644
--- a/test/tint/builtins/gen/literal/ceil/678655.wgsl
+++ b/test/tint/builtins/gen/literal/ceil/678655.wgsl
@@ -23,7 +23,7 @@
 
 // fn ceil(f32) -> f32
 fn ceil_678655() {
-  var res: f32 = ceil(1.f);
+  var res: f32 = ceil(1.5f);
 }
 
 @vertex
diff --git a/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.dxc.hlsl
index 27adeb9..56a79d1 100644
--- a/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_678655() {
-  float res = ceil(1.0f);
+  float res = 2.0f;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.fxc.hlsl
index 27adeb9..56a79d1 100644
--- a/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.fxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_678655() {
-  float res = ceil(1.0f);
+  float res = 2.0f;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.glsl b/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.glsl
index 08ca901..9a15224 100644
--- a/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.glsl
@@ -1,7 +1,7 @@
 #version 310 es
 
 void ceil_678655() {
-  float res = ceil(1.0f);
+  float res = 2.0f;
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@
 precision mediump float;
 
 void ceil_678655() {
-  float res = ceil(1.0f);
+  float res = 2.0f;
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@
 #version 310 es
 
 void ceil_678655() {
-  float res = ceil(1.0f);
+  float res = 2.0f;
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.msl b/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.msl
index 37d1aec..bc40d4f 100644
--- a/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ceil_678655() {
-  float res = ceil(1.0f);
+  float res = 2.0f;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.spvasm
index bb45e0b..81a9a67 100644
--- a/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/ceil/678655.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
 %ceil_678655 = OpFunction %void None %9
          %12 = OpLabel
         %res = OpVariable %_ptr_Function_float Function %8
-         %13 = OpExtInst %float %14 Ceil %float_1
-               OpStore %res %13
+               OpStore %res %float_2
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %18
-         %20 = OpLabel
-         %21 = OpFunctionCall %void %ceil_678655
+%vertex_main_inner = OpFunction %v4float None %16
+         %18 = OpLabel
+         %19 = OpFunctionCall %void %ceil_678655
                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 %ceil_678655
+         %25 = OpLabel
+         %26 = OpFunctionCall %void %ceil_678655
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %29 = OpLabel
-         %30 = OpFunctionCall %void %ceil_678655
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ceil_678655
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.wgsl
index db1c9bb..01d590d 100644
--- a/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/ceil/678655.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
 fn ceil_678655() {
-  var res : f32 = ceil(1.0f);
+  var res : f32 = ceil(1.5f);
 }
 
 @vertex
diff --git a/test/tint/builtins/gen/literal/ceil/96f597.wgsl b/test/tint/builtins/gen/literal/ceil/96f597.wgsl
index ee01d1b..db8fc64 100644
--- a/test/tint/builtins/gen/literal/ceil/96f597.wgsl
+++ b/test/tint/builtins/gen/literal/ceil/96f597.wgsl
@@ -23,7 +23,7 @@
 
 // fn ceil(vec<2, f32>) -> vec<2, f32>
 fn ceil_96f597() {
-  var res: vec2<f32> = ceil(vec2<f32>(1.f));
+  var res: vec2<f32> = ceil(vec2<f32>(1.5f));
 }
 
 @vertex
diff --git a/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.dxc.hlsl
index 01138db..2469cd1 100644
--- a/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_96f597() {
-  float2 res = ceil((1.0f).xx);
+  float2 res = (2.0f).xx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.fxc.hlsl
index 01138db..2469cd1 100644
--- a/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.fxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_96f597() {
-  float2 res = ceil((1.0f).xx);
+  float2 res = (2.0f).xx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.glsl b/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.glsl
index c38418d..9b37453 100644
--- a/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.glsl
@@ -1,7 +1,7 @@
 #version 310 es
 
 void ceil_96f597() {
-  vec2 res = ceil(vec2(1.0f));
+  vec2 res = vec2(2.0f);
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@
 precision mediump float;
 
 void ceil_96f597() {
-  vec2 res = ceil(vec2(1.0f));
+  vec2 res = vec2(2.0f);
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@
 #version 310 es
 
 void ceil_96f597() {
-  vec2 res = ceil(vec2(1.0f));
+  vec2 res = vec2(2.0f);
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.msl b/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.msl
index d0a6c92..2cd2768 100644
--- a/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ceil_96f597() {
-  float2 res = ceil(float2(1.0f));
+  float2 res = float2(2.0f);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.spvasm
index 2eab970..67e52a3 100644
--- a/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/ceil/96f597.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
 %ceil_96f597 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v2float Function %20
-         %13 = OpExtInst %v2float %15 Ceil %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 %ceil_96f597
+%vertex_main_inner = OpFunction %v4float None %19
+         %21 = OpLabel
+         %22 = OpFunctionCall %void %ceil_96f597
                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 %ceil_96f597
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ceil_96f597
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %ceil_96f597
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ceil_96f597
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.wgsl
index fd58dc1..6cd0ba5 100644
--- a/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/ceil/96f597.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
 fn ceil_96f597() {
-  var res : vec2<f32> = ceil(vec2<f32>(1.0f));
+  var res : vec2<f32> = ceil(vec2<f32>(1.5f));
 }
 
 @vertex
diff --git a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl b/test/tint/builtins/gen/literal/ceil/b74c16.wgsl
index 90280c0..c2d9472 100644
--- a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl
+++ b/test/tint/builtins/gen/literal/ceil/b74c16.wgsl
@@ -23,7 +23,7 @@
 
 // fn ceil(vec<4, f32>) -> vec<4, f32>
 fn ceil_b74c16() {
-  var res: vec4<f32> = ceil(vec4<f32>(1.f));
+  var res: vec4<f32> = ceil(vec4<f32>(1.5f));
 }
 
 @vertex
diff --git a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.dxc.hlsl
index 94df849..00beff8 100644
--- a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_b74c16() {
-  float4 res = ceil((1.0f).xxxx);
+  float4 res = (2.0f).xxxx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.fxc.hlsl
index 94df849..00beff8 100644
--- a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.fxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_b74c16() {
-  float4 res = ceil((1.0f).xxxx);
+  float4 res = (2.0f).xxxx;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.glsl b/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.glsl
index d64a737..54c3f79 100644
--- a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.glsl
@@ -1,7 +1,7 @@
 #version 310 es
 
 void ceil_b74c16() {
-  vec4 res = ceil(vec4(1.0f));
+  vec4 res = vec4(2.0f);
 }
 
 vec4 vertex_main() {
@@ -21,7 +21,7 @@
 precision mediump float;
 
 void ceil_b74c16() {
-  vec4 res = ceil(vec4(1.0f));
+  vec4 res = vec4(2.0f);
 }
 
 void fragment_main() {
@@ -35,7 +35,7 @@
 #version 310 es
 
 void ceil_b74c16() {
-  vec4 res = ceil(vec4(1.0f));
+  vec4 res = vec4(2.0f);
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.msl b/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.msl
index 71679e6..c68a4d6 100644
--- a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ceil_b74c16() {
-  float4 res = ceil(float4(1.0f));
+  float4 res = float4(2.0f);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.spvasm
index c673475..4a60c5b 100644
--- a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/ceil/b74c16.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
 %ceil_b74c16 = OpFunction %void None %9
          %12 = OpLabel
         %res = OpVariable %_ptr_Function_v4float Function %5
-         %13 = OpExtInst %v4float %14 Ceil %16
-               OpStore %res %13
+               OpStore %res %14
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %19
-         %21 = OpLabel
-         %22 = OpFunctionCall %void %ceil_b74c16
+%vertex_main_inner = OpFunction %v4float None %17
+         %19 = OpLabel
+         %20 = OpFunctionCall %void %ceil_b74c16
                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 %ceil_b74c16
+         %26 = OpLabel
+         %27 = OpFunctionCall %void %ceil_b74c16
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %30 = OpLabel
-         %31 = OpFunctionCall %void %ceil_b74c16
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %ceil_b74c16
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.wgsl
index d378533..73d40a5 100644
--- a/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/ceil/b74c16.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
 fn ceil_b74c16() {
-  var res : vec4<f32> = ceil(vec4<f32>(1.0f));
+  var res : vec4<f32> = ceil(vec4<f32>(1.5f));
 }
 
 @vertex
diff --git a/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl b/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl
new file mode 100644
index 0000000..a6a02bb
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/bb2ca2.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 ceil(vec<2, fa>) -> vec<2, fa>
+fn ceil_bb2ca2() {
+  var res = ceil(vec2(1.5));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_bb2ca2();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_bb2ca2();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_bb2ca2();
+}
diff --git a/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..ceb4fee
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_bb2ca2() {
+  float2 res = (2.0f).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_bb2ca2();
+  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() {
+  ceil_bb2ca2();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_bb2ca2();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..ceb4fee
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_bb2ca2() {
+  float2 res = (2.0f).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_bb2ca2();
+  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() {
+  ceil_bb2ca2();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_bb2ca2();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.glsl b/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.glsl
new file mode 100644
index 0000000..506d1f0
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void ceil_bb2ca2() {
+  vec2 res = vec2(2.0f);
+}
+
+vec4 vertex_main() {
+  ceil_bb2ca2();
+  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 ceil_bb2ca2() {
+  vec2 res = vec2(2.0f);
+}
+
+void fragment_main() {
+  ceil_bb2ca2();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void ceil_bb2ca2() {
+  vec2 res = vec2(2.0f);
+}
+
+void compute_main() {
+  ceil_bb2ca2();
+}
+
+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/ceil/bb2ca2.wgsl.expected.msl b/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.msl
new file mode 100644
index 0000000..9239b2f
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void ceil_bb2ca2() {
+  float2 res = float2(2.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  ceil_bb2ca2();
+  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() {
+  ceil_bb2ca2();
+  return;
+}
+
+kernel void compute_main() {
+  ceil_bb2ca2();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.spvasm
new file mode 100644
index 0000000..51aa5f7
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/bb2ca2.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 %ceil_bb2ca2 "ceil_bb2ca2"
+               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
+%ceil_bb2ca2 = 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 %ceil_bb2ca2
+               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 %ceil_bb2ca2
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ceil_bb2ca2
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.wgsl
new file mode 100644
index 0000000..f08a92e
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/bb2ca2.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn ceil_bb2ca2() {
+  var res = ceil(vec2(1.5));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_bb2ca2();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_bb2ca2();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_bb2ca2();
+}
diff --git a/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl b/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl
new file mode 100644
index 0000000..c8f8d36
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/e0b70a.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 ceil(fa) -> fa
+fn ceil_e0b70a() {
+  var res = ceil(1.5);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_e0b70a();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_e0b70a();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_e0b70a();
+}
diff --git a/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..8f36f5c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_e0b70a() {
+  float res = 2.0f;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_e0b70a();
+  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() {
+  ceil_e0b70a();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_e0b70a();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..8f36f5c
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_e0b70a() {
+  float res = 2.0f;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_e0b70a();
+  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() {
+  ceil_e0b70a();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_e0b70a();
+  return;
+}
diff --git a/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.glsl b/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.glsl
new file mode 100644
index 0000000..86208fc
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void ceil_e0b70a() {
+  float res = 2.0f;
+}
+
+vec4 vertex_main() {
+  ceil_e0b70a();
+  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 ceil_e0b70a() {
+  float res = 2.0f;
+}
+
+void fragment_main() {
+  ceil_e0b70a();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void ceil_e0b70a() {
+  float res = 2.0f;
+}
+
+void compute_main() {
+  ceil_e0b70a();
+}
+
+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/ceil/e0b70a.wgsl.expected.msl b/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.msl
new file mode 100644
index 0000000..e2956b8
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void ceil_e0b70a() {
+  float res = 2.0f;
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  ceil_e0b70a();
+  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() {
+  ceil_e0b70a();
+  return;
+}
+
+kernel void compute_main() {
+  ceil_e0b70a();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.spvasm
new file mode 100644
index 0000000..ea8a548
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/e0b70a.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 %ceil_e0b70a "ceil_e0b70a"
+               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
+%ceil_e0b70a = 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 %ceil_e0b70a
+               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 %ceil_e0b70a
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ceil_e0b70a
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.wgsl
new file mode 100644
index 0000000..dd1fb8d
--- /dev/null
+++ b/test/tint/builtins/gen/literal/ceil/e0b70a.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn ceil_e0b70a() {
+  var res = ceil(1.5);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_e0b70a();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_e0b70a();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_e0b70a();
+}
diff --git a/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.dxc.hlsl
index b9cee87..1dbabc6 100644
--- a/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_f3f889() {
-  float16_t res = ceil(float16_t(0.0h));
+  float16_t res = float16_t(0.0h);
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.glsl b/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.glsl
index c34111a..dffe515 100644
--- a/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.glsl
@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void ceil_f3f889() {
-  float16_t res = ceil(0.0hf);
+  float16_t res = 0.0hf;
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@
 precision mediump float;
 
 void ceil_f3f889() {
-  float16_t res = ceil(0.0hf);
+  float16_t res = 0.0hf;
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void ceil_f3f889() {
-  float16_t res = ceil(0.0hf);
+  float16_t res = 0.0hf;
 }
 
 void compute_main() {
diff --git a/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.msl b/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.msl
index 2bfecc6..214d570 100644
--- a/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ceil_f3f889() {
-  half res = ceil(0.0h);
+  half res = 0.0h;
 }
 
 struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.spvasm
index b50e1ad..2f13171 100644
--- a/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/ceil/f3f889.wgsl.expected.spvasm
@@ -1,14 +1,13 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 33
+; Bound: 31
 ; 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,36 +35,35 @@
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
        %half = OpTypeFloat 16
-         %16 = OpConstantNull %half
+         %14 = OpConstantNull %half
 %_ptr_Function_half = OpTypePointer Function %half
-         %19 = OpTypeFunction %v4float
+         %17 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
 %ceil_f3f889 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_half Function %16
-         %13 = OpExtInst %half %15 Ceil %16
-               OpStore %res %13
+        %res = OpVariable %_ptr_Function_half Function %14
+               OpStore %res %14
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %19
-         %21 = OpLabel
-         %22 = OpFunctionCall %void %ceil_f3f889
+%vertex_main_inner = OpFunction %v4float None %17
+         %19 = OpLabel
+         %20 = OpFunctionCall %void %ceil_f3f889
                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
-         %28 = OpLabel
-         %29 = OpFunctionCall %void %ceil_f3f889
+         %26 = OpLabel
+         %27 = OpFunctionCall %void %ceil_f3f889
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %31 = OpLabel
-         %32 = OpFunctionCall %void %ceil_f3f889
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %ceil_f3f889
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/ceil/11b1dc.wgsl b/test/tint/builtins/gen/var/ceil/11b1dc.wgsl
new file mode 100644
index 0000000..e87aa07
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/11b1dc.wgsl
@@ -0,0 +1,44 @@
+// 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 ceil(vec<4, fa>) -> vec<4, fa>
+fn ceil_11b1dc() {
+  const arg_0 = vec4(1.5);
+  var res = ceil(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_11b1dc();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_11b1dc();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_11b1dc();
+}
diff --git a/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..05f5861
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_11b1dc() {
+  float4 res = (2.0f).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_11b1dc();
+  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() {
+  ceil_11b1dc();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_11b1dc();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..05f5861
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_11b1dc() {
+  float4 res = (2.0f).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_11b1dc();
+  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() {
+  ceil_11b1dc();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_11b1dc();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.glsl b/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.glsl
new file mode 100644
index 0000000..7a6447b
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void ceil_11b1dc() {
+  vec4 res = vec4(2.0f);
+}
+
+vec4 vertex_main() {
+  ceil_11b1dc();
+  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 ceil_11b1dc() {
+  vec4 res = vec4(2.0f);
+}
+
+void fragment_main() {
+  ceil_11b1dc();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void ceil_11b1dc() {
+  vec4 res = vec4(2.0f);
+}
+
+void compute_main() {
+  ceil_11b1dc();
+}
+
+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/ceil/11b1dc.wgsl.expected.msl b/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.msl
new file mode 100644
index 0000000..4da2a34
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void ceil_11b1dc() {
+  float4 res = float4(2.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  ceil_11b1dc();
+  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() {
+  ceil_11b1dc();
+  return;
+}
+
+kernel void compute_main() {
+  ceil_11b1dc();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.spvasm b/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.spvasm
new file mode 100644
index 0000000..924f17b
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/11b1dc.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 %ceil_11b1dc "ceil_11b1dc"
+               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
+%ceil_11b1dc = 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 %ceil_11b1dc
+               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 %ceil_11b1dc
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %ceil_11b1dc
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.wgsl b/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.wgsl
new file mode 100644
index 0000000..696b29f
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/11b1dc.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+fn ceil_11b1dc() {
+  const arg_0 = vec4(1.5);
+  var res = ceil(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_11b1dc();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_11b1dc();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_11b1dc();
+}
diff --git a/test/tint/builtins/gen/var/ceil/32c946.wgsl b/test/tint/builtins/gen/var/ceil/32c946.wgsl
new file mode 100644
index 0000000..505d32b
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/32c946.wgsl
@@ -0,0 +1,44 @@
+// 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 ceil(vec<3, fa>) -> vec<3, fa>
+fn ceil_32c946() {
+  const arg_0 = vec3(1.5);
+  var res = ceil(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_32c946();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_32c946();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_32c946();
+}
diff --git a/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..fcd472c
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_32c946() {
+  float3 res = (2.0f).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_32c946();
+  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() {
+  ceil_32c946();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_32c946();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..fcd472c
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_32c946() {
+  float3 res = (2.0f).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_32c946();
+  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() {
+  ceil_32c946();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_32c946();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.glsl b/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.glsl
new file mode 100644
index 0000000..c129c84
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void ceil_32c946() {
+  vec3 res = vec3(2.0f);
+}
+
+vec4 vertex_main() {
+  ceil_32c946();
+  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 ceil_32c946() {
+  vec3 res = vec3(2.0f);
+}
+
+void fragment_main() {
+  ceil_32c946();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void ceil_32c946() {
+  vec3 res = vec3(2.0f);
+}
+
+void compute_main() {
+  ceil_32c946();
+}
+
+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/ceil/32c946.wgsl.expected.msl b/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.msl
new file mode 100644
index 0000000..1096249
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void ceil_32c946() {
+  float3 res = float3(2.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  ceil_32c946();
+  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() {
+  ceil_32c946();
+  return;
+}
+
+kernel void compute_main() {
+  ceil_32c946();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.spvasm b/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.spvasm
new file mode 100644
index 0000000..2711442
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/32c946.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 %ceil_32c946 "ceil_32c946"
+               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
+%ceil_32c946 = 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 %ceil_32c946
+               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 %ceil_32c946
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ceil_32c946
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.wgsl b/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.wgsl
new file mode 100644
index 0000000..a357b0b
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/32c946.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+fn ceil_32c946() {
+  const arg_0 = vec3(1.5);
+  var res = ceil(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_32c946();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_32c946();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_32c946();
+}
diff --git a/test/tint/builtins/gen/var/ceil/34064b.wgsl b/test/tint/builtins/gen/var/ceil/34064b.wgsl
index 9d55693..5cdb7e3 100644
--- a/test/tint/builtins/gen/var/ceil/34064b.wgsl
+++ b/test/tint/builtins/gen/var/ceil/34064b.wgsl
@@ -23,7 +23,7 @@
 
 // fn ceil(vec<3, f32>) -> vec<3, f32>
 fn ceil_34064b() {
-  var arg_0 = vec3<f32>(1.f);
+  var arg_0 = vec3<f32>(1.5f);
   var res: vec3<f32> = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.dxc.hlsl
index cf62426..b195a21 100644
--- a/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_34064b() {
-  float3 arg_0 = (1.0f).xxx;
+  float3 arg_0 = (1.5f).xxx;
   float3 res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.fxc.hlsl
index cf62426..b195a21 100644
--- a/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.fxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_34064b() {
-  float3 arg_0 = (1.0f).xxx;
+  float3 arg_0 = (1.5f).xxx;
   float3 res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.glsl b/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.glsl
index 068e52a..c5b7218 100644
--- a/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.glsl
@@ -1,7 +1,7 @@
 #version 310 es
 
 void ceil_34064b() {
-  vec3 arg_0 = vec3(1.0f);
+  vec3 arg_0 = vec3(1.5f);
   vec3 res = ceil(arg_0);
 }
 
@@ -22,7 +22,7 @@
 precision mediump float;
 
 void ceil_34064b() {
-  vec3 arg_0 = vec3(1.0f);
+  vec3 arg_0 = vec3(1.5f);
   vec3 res = ceil(arg_0);
 }
 
@@ -37,7 +37,7 @@
 #version 310 es
 
 void ceil_34064b() {
-  vec3 arg_0 = vec3(1.0f);
+  vec3 arg_0 = vec3(1.5f);
   vec3 res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.msl b/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.msl
index ce94f13..8a2cc0b 100644
--- a/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ceil_34064b() {
-  float3 arg_0 = float3(1.0f);
+  float3 arg_0 = float3(1.5f);
   float3 res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.spvasm b/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.spvasm
index d9674bf..949ebc4 100644
--- a/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 37
 ; Schema: 0
                OpCapability Shader
          %20 = OpExtInstImport "GLSL.std.450"
@@ -33,11 +33,12 @@
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
     %v3float = OpTypeVector %float 3
-    %float_1 = OpConstant %float 1
-         %15 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+  %float_1_5 = OpConstant %float 1.5
+         %15 = OpConstantComposite %v3float %float_1_5 %float_1_5 %float_1_5
 %_ptr_Function_v3float = OpTypePointer Function %v3float
          %18 = OpConstantNull %v3float
          %23 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
 %ceil_34064b = OpFunction %void None %9
          %12 = OpLabel
       %arg_0 = OpVariable %_ptr_Function_v3float Function %18
@@ -61,12 +62,12 @@
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %31 = OpLabel
-         %32 = OpFunctionCall %void %ceil_34064b
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %ceil_34064b
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %ceil_34064b
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %ceil_34064b
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.wgsl b/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.wgsl
index ec9cd28..52f16c4 100644
--- a/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/ceil/34064b.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
 fn ceil_34064b() {
-  var arg_0 = vec3<f32>(1.0f);
+  var arg_0 = vec3<f32>(1.5f);
   var res : vec3<f32> = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/678655.wgsl b/test/tint/builtins/gen/var/ceil/678655.wgsl
index 6906a4d..6446299 100644
--- a/test/tint/builtins/gen/var/ceil/678655.wgsl
+++ b/test/tint/builtins/gen/var/ceil/678655.wgsl
@@ -23,7 +23,7 @@
 
 // fn ceil(f32) -> f32
 fn ceil_678655() {
-  var arg_0 = 1.f;
+  var arg_0 = 1.5f;
   var res: f32 = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.dxc.hlsl
index 61f3dc0..981082d 100644
--- a/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_678655() {
-  float arg_0 = 1.0f;
+  float arg_0 = 1.5f;
   float res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.fxc.hlsl
index 61f3dc0..981082d 100644
--- a/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.fxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_678655() {
-  float arg_0 = 1.0f;
+  float arg_0 = 1.5f;
   float res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.glsl b/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.glsl
index 56f09a6..85fc5c5 100644
--- a/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.glsl
@@ -1,7 +1,7 @@
 #version 310 es
 
 void ceil_678655() {
-  float arg_0 = 1.0f;
+  float arg_0 = 1.5f;
   float res = ceil(arg_0);
 }
 
@@ -22,7 +22,7 @@
 precision mediump float;
 
 void ceil_678655() {
-  float arg_0 = 1.0f;
+  float arg_0 = 1.5f;
   float res = ceil(arg_0);
 }
 
@@ -37,7 +37,7 @@
 #version 310 es
 
 void ceil_678655() {
-  float arg_0 = 1.0f;
+  float arg_0 = 1.5f;
   float res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.msl b/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.msl
index 8714a88..5637519 100644
--- a/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ceil_678655() {
-  float arg_0 = 1.0f;
+  float arg_0 = 1.5f;
   float res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.spvasm b/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.spvasm
index 465ddd4..5b09b5a 100644
--- a/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 33
+; Bound: 34
 ; Schema: 0
                OpCapability Shader
          %17 = OpExtInstImport "GLSL.std.450"
@@ -32,14 +32,15 @@
 %vertex_point_size = OpVariable %_ptr_Output_float Output %8
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
-    %float_1 = OpConstant %float 1
+  %float_1_5 = OpConstant %float 1.5
 %_ptr_Function_float = OpTypePointer Function %float
          %20 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
 %ceil_678655 = OpFunction %void None %9
          %12 = OpLabel
       %arg_0 = OpVariable %_ptr_Function_float Function %8
         %res = OpVariable %_ptr_Function_float Function %8
-               OpStore %arg_0 %float_1
+               OpStore %arg_0 %float_1_5
          %18 = OpLoad %float %arg_0
          %16 = OpExtInst %float %17 Ceil %18
                OpStore %res %16
@@ -58,12 +59,12 @@
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %28 = OpLabel
-         %29 = OpFunctionCall %void %ceil_678655
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %ceil_678655
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %31 = OpLabel
-         %32 = OpFunctionCall %void %ceil_678655
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %ceil_678655
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.wgsl b/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.wgsl
index d3cc9b4..789395f 100644
--- a/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/ceil/678655.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
 fn ceil_678655() {
-  var arg_0 = 1.0f;
+  var arg_0 = 1.5f;
   var res : f32 = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/96f597.wgsl b/test/tint/builtins/gen/var/ceil/96f597.wgsl
index d856956..c149c74 100644
--- a/test/tint/builtins/gen/var/ceil/96f597.wgsl
+++ b/test/tint/builtins/gen/var/ceil/96f597.wgsl
@@ -23,7 +23,7 @@
 
 // fn ceil(vec<2, f32>) -> vec<2, f32>
 fn ceil_96f597() {
-  var arg_0 = vec2<f32>(1.f);
+  var arg_0 = vec2<f32>(1.5f);
   var res: vec2<f32> = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.dxc.hlsl
index 7bdb61b..98e9c51 100644
--- a/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_96f597() {
-  float2 arg_0 = (1.0f).xx;
+  float2 arg_0 = (1.5f).xx;
   float2 res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.fxc.hlsl
index 7bdb61b..98e9c51 100644
--- a/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.fxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_96f597() {
-  float2 arg_0 = (1.0f).xx;
+  float2 arg_0 = (1.5f).xx;
   float2 res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.glsl b/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.glsl
index 3c93695..ceba85e 100644
--- a/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.glsl
@@ -1,7 +1,7 @@
 #version 310 es
 
 void ceil_96f597() {
-  vec2 arg_0 = vec2(1.0f);
+  vec2 arg_0 = vec2(1.5f);
   vec2 res = ceil(arg_0);
 }
 
@@ -22,7 +22,7 @@
 precision mediump float;
 
 void ceil_96f597() {
-  vec2 arg_0 = vec2(1.0f);
+  vec2 arg_0 = vec2(1.5f);
   vec2 res = ceil(arg_0);
 }
 
@@ -37,7 +37,7 @@
 #version 310 es
 
 void ceil_96f597() {
-  vec2 arg_0 = vec2(1.0f);
+  vec2 arg_0 = vec2(1.5f);
   vec2 res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.msl b/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.msl
index e3cb08b..c2e5a5f 100644
--- a/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ceil_96f597() {
-  float2 arg_0 = float2(1.0f);
+  float2 arg_0 = float2(1.5f);
   float2 res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.spvasm b/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.spvasm
index ba9ec75..f825800 100644
--- a/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 37
 ; Schema: 0
                OpCapability Shader
          %20 = OpExtInstImport "GLSL.std.450"
@@ -33,11 +33,12 @@
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
     %v2float = OpTypeVector %float 2
-    %float_1 = OpConstant %float 1
-         %15 = OpConstantComposite %v2float %float_1 %float_1
+  %float_1_5 = OpConstant %float 1.5
+         %15 = OpConstantComposite %v2float %float_1_5 %float_1_5
 %_ptr_Function_v2float = OpTypePointer Function %v2float
          %18 = OpConstantNull %v2float
          %23 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
 %ceil_96f597 = OpFunction %void None %9
          %12 = OpLabel
       %arg_0 = OpVariable %_ptr_Function_v2float Function %18
@@ -61,12 +62,12 @@
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %31 = OpLabel
-         %32 = OpFunctionCall %void %ceil_96f597
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %ceil_96f597
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %ceil_96f597
+         %35 = OpLabel
+         %36 = OpFunctionCall %void %ceil_96f597
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.wgsl b/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.wgsl
index a08c2be..a36992b 100644
--- a/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/ceil/96f597.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
 fn ceil_96f597() {
-  var arg_0 = vec2<f32>(1.0f);
+  var arg_0 = vec2<f32>(1.5f);
   var res : vec2<f32> = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/b74c16.wgsl b/test/tint/builtins/gen/var/ceil/b74c16.wgsl
index 697dab2..b5ff72b 100644
--- a/test/tint/builtins/gen/var/ceil/b74c16.wgsl
+++ b/test/tint/builtins/gen/var/ceil/b74c16.wgsl
@@ -23,7 +23,7 @@
 
 // fn ceil(vec<4, f32>) -> vec<4, f32>
 fn ceil_b74c16() {
-  var arg_0 = vec4<f32>(1.f);
+  var arg_0 = vec4<f32>(1.5f);
   var res: vec4<f32> = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.dxc.hlsl
index f914edf..dedbf8a 100644
--- a/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.dxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_b74c16() {
-  float4 arg_0 = (1.0f).xxxx;
+  float4 arg_0 = (1.5f).xxxx;
   float4 res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.fxc.hlsl
index f914edf..dedbf8a 100644
--- a/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.fxc.hlsl
@@ -1,5 +1,5 @@
 void ceil_b74c16() {
-  float4 arg_0 = (1.0f).xxxx;
+  float4 arg_0 = (1.5f).xxxx;
   float4 res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.glsl b/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.glsl
index f630971..56993a9 100644
--- a/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.glsl
@@ -1,7 +1,7 @@
 #version 310 es
 
 void ceil_b74c16() {
-  vec4 arg_0 = vec4(1.0f);
+  vec4 arg_0 = vec4(1.5f);
   vec4 res = ceil(arg_0);
 }
 
@@ -22,7 +22,7 @@
 precision mediump float;
 
 void ceil_b74c16() {
-  vec4 arg_0 = vec4(1.0f);
+  vec4 arg_0 = vec4(1.5f);
   vec4 res = ceil(arg_0);
 }
 
@@ -37,7 +37,7 @@
 #version 310 es
 
 void ceil_b74c16() {
-  vec4 arg_0 = vec4(1.0f);
+  vec4 arg_0 = vec4(1.5f);
   vec4 res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.msl b/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.msl
index 66c5898..26b12da 100644
--- a/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ceil_b74c16() {
-  float4 arg_0 = float4(1.0f);
+  float4 arg_0 = float4(1.5f);
   float4 res = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.spvasm b/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.spvasm
index 0e581a1..3f86045 100644
--- a/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 35
 ; Schema: 0
                OpCapability Shader
          %18 = OpExtInstImport "GLSL.std.450"
@@ -32,10 +32,11 @@
 %vertex_point_size = OpVariable %_ptr_Output_float Output %8
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
-    %float_1 = OpConstant %float 1
-         %14 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+  %float_1_5 = OpConstant %float 1.5
+         %14 = OpConstantComposite %v4float %float_1_5 %float_1_5 %float_1_5 %float_1_5
 %_ptr_Function_v4float = OpTypePointer Function %v4float
          %21 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
 %ceil_b74c16 = OpFunction %void None %9
          %12 = OpLabel
       %arg_0 = OpVariable %_ptr_Function_v4float Function %5
@@ -59,12 +60,12 @@
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %29 = OpLabel
-         %30 = OpFunctionCall %void %ceil_b74c16
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %ceil_b74c16
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %ceil_b74c16
+         %33 = OpLabel
+         %34 = OpFunctionCall %void %ceil_b74c16
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.wgsl b/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.wgsl
index e2a7c63..5ec0df2 100644
--- a/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/ceil/b74c16.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
 fn ceil_b74c16() {
-  var arg_0 = vec4<f32>(1.0f);
+  var arg_0 = vec4<f32>(1.5f);
   var res : vec4<f32> = ceil(arg_0);
 }
 
diff --git a/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl b/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl
new file mode 100644
index 0000000..c9ea5c5
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl
@@ -0,0 +1,44 @@
+// 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 ceil(vec<2, fa>) -> vec<2, fa>
+fn ceil_bb2ca2() {
+  const arg_0 = vec2(1.5);
+  var res = ceil(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_bb2ca2();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_bb2ca2();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_bb2ca2();
+}
diff --git a/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..ceb4fee
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_bb2ca2() {
+  float2 res = (2.0f).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_bb2ca2();
+  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() {
+  ceil_bb2ca2();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_bb2ca2();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..ceb4fee
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_bb2ca2() {
+  float2 res = (2.0f).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_bb2ca2();
+  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() {
+  ceil_bb2ca2();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_bb2ca2();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.glsl b/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.glsl
new file mode 100644
index 0000000..506d1f0
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void ceil_bb2ca2() {
+  vec2 res = vec2(2.0f);
+}
+
+vec4 vertex_main() {
+  ceil_bb2ca2();
+  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 ceil_bb2ca2() {
+  vec2 res = vec2(2.0f);
+}
+
+void fragment_main() {
+  ceil_bb2ca2();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void ceil_bb2ca2() {
+  vec2 res = vec2(2.0f);
+}
+
+void compute_main() {
+  ceil_bb2ca2();
+}
+
+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/ceil/bb2ca2.wgsl.expected.msl b/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.msl
new file mode 100644
index 0000000..9239b2f
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void ceil_bb2ca2() {
+  float2 res = float2(2.0f);
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  ceil_bb2ca2();
+  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() {
+  ceil_bb2ca2();
+  return;
+}
+
+kernel void compute_main() {
+  ceil_bb2ca2();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.spvasm b/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.spvasm
new file mode 100644
index 0000000..51aa5f7
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/bb2ca2.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 %ceil_bb2ca2 "ceil_bb2ca2"
+               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
+%ceil_bb2ca2 = 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 %ceil_bb2ca2
+               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 %ceil_bb2ca2
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ceil_bb2ca2
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.wgsl b/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.wgsl
new file mode 100644
index 0000000..4a4d7e9
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/bb2ca2.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+fn ceil_bb2ca2() {
+  const arg_0 = vec2(1.5);
+  var res = ceil(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_bb2ca2();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_bb2ca2();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_bb2ca2();
+}
diff --git a/test/tint/builtins/gen/var/ceil/e0b70a.wgsl b/test/tint/builtins/gen/var/ceil/e0b70a.wgsl
new file mode 100644
index 0000000..af56279
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/e0b70a.wgsl
@@ -0,0 +1,44 @@
+// 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 ceil(fa) -> fa
+fn ceil_e0b70a() {
+  const arg_0 = 1.5;
+  var res = ceil(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_e0b70a();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_e0b70a();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_e0b70a();
+}
diff --git a/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..8f36f5c
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_e0b70a() {
+  float res = 2.0f;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_e0b70a();
+  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() {
+  ceil_e0b70a();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_e0b70a();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..8f36f5c
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+void ceil_e0b70a() {
+  float res = 2.0f;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ceil_e0b70a();
+  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() {
+  ceil_e0b70a();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ceil_e0b70a();
+  return;
+}
diff --git a/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.glsl b/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.glsl
new file mode 100644
index 0000000..86208fc
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void ceil_e0b70a() {
+  float res = 2.0f;
+}
+
+vec4 vertex_main() {
+  ceil_e0b70a();
+  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 ceil_e0b70a() {
+  float res = 2.0f;
+}
+
+void fragment_main() {
+  ceil_e0b70a();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void ceil_e0b70a() {
+  float res = 2.0f;
+}
+
+void compute_main() {
+  ceil_e0b70a();
+}
+
+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/ceil/e0b70a.wgsl.expected.msl b/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.msl
new file mode 100644
index 0000000..e2956b8
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void ceil_e0b70a() {
+  float res = 2.0f;
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+float4 vertex_main_inner() {
+  ceil_e0b70a();
+  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() {
+  ceil_e0b70a();
+  return;
+}
+
+kernel void compute_main() {
+  ceil_e0b70a();
+  return;
+}
+
diff --git a/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.spvasm b/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.spvasm
new file mode 100644
index 0000000..ea8a548
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/e0b70a.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 %ceil_e0b70a "ceil_e0b70a"
+               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
+%ceil_e0b70a = 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 %ceil_e0b70a
+               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 %ceil_e0b70a
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ceil_e0b70a
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.wgsl b/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.wgsl
new file mode 100644
index 0000000..64a34d7
--- /dev/null
+++ b/test/tint/builtins/gen/var/ceil/e0b70a.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+fn ceil_e0b70a() {
+  const arg_0 = 1.5;
+  var res = ceil(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ceil_e0b70a();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ceil_e0b70a();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ceil_e0b70a();
+}
diff --git a/webgpu-cts/expectations.txt b/webgpu-cts/expectations.txt
index 09f45ce..527e5df 100644
--- a/webgpu-cts/expectations.txt
+++ b/webgpu-cts/expectations.txt
@@ -403,10 +403,6 @@
 crbug.com/tint/1613 webgpu:shader,execution,expression,call,builtin,atanh:f32:inputSource="const";vectorize=2 [ Failure ]
 crbug.com/tint/1613 webgpu:shader,execution,expression,call,builtin,atanh:f32:inputSource="const";vectorize=3 [ Failure ]
 crbug.com/tint/1613 webgpu:shader,execution,expression,call,builtin,atanh:f32:inputSource="const";vectorize=4 [ Failure ]
-crbug.com/tint/1613 webgpu:shader,execution,expression,call,builtin,ceil:f32:inputSource="const";vectorize="_undef_" [ Failure ]
-crbug.com/tint/1613 webgpu:shader,execution,expression,call,builtin,ceil:f32:inputSource="const";vectorize=2 [ Failure ]
-crbug.com/tint/1613 webgpu:shader,execution,expression,call,builtin,ceil:f32:inputSource="const";vectorize=3 [ Failure ]
-crbug.com/tint/1613 webgpu:shader,execution,expression,call,builtin,ceil:f32:inputSource="const";vectorize=4 [ Failure ]
 crbug.com/tint/1613 webgpu:shader,execution,expression,call,builtin,cos:f32:inputSource="const";vectorize="_undef_" [ Failure ]
 crbug.com/tint/1613 webgpu:shader,execution,expression,call,builtin,cos:f32:inputSource="const";vectorize=2 [ Failure ]
 crbug.com/tint/1613 webgpu:shader,execution,expression,call,builtin,cos:f32:inputSource="const";vectorize=3 [ Failure ]