tint: const eval of ldexp builtin

Bug: tint:1581
Change-Id: Ib21717065041b65a637f4d73ce0088544b1fce0d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/114321
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

src/tint/intrinsics.def

@@ -226,6 +226,7 @@ match fia_fiu32: fa | ia | f32 | i32 | u32
 match fa_f32: fa | f32
 match fa_f32_f16: fa | f32 | f16
 match ia_iu32: ia | i32 | u32
+match ia_i32: ia | i32
 match fiu32_f16: f32 | i32 | u32 | f16
 match fiu32: f32 | i32 | u32
 match fi32_f16: f32 | i32 | f16
@@ -527,8 +528,8 @@ fn dot4U8Packed(u32, u32) -> u32
 @const fn insertBits<N: num, T: iu32>(vec<N, T>, vec<N, T>, u32, u32) -> vec<N, T>
 @const fn inverseSqrt<T: fa_f32_f16>(T) -> T
 @const fn inverseSqrt<N: num, T: fa_f32_f16>(vec<N, T>) -> vec<N, T>
-fn ldexp<T: f32_f16>(T, i32) -> T
+@const fn ldexp<T: fa_f32_f16, U: ia_i32>(T, U) -> T
-fn ldexp<N: num, T: f32_f16>(vec<N, T>, vec<N, i32>) -> vec<N, T>
+@const fn ldexp<N: num, T: fa_f32_f16, U: ia_i32>(vec<N, T>, vec<N, U>) -> vec<N, T>
 @const fn length<T: fa_f32_f16>(@test_value(0.0) T) -> T
 @const fn length<N: num, T: fa_f32_f16>(@test_value(0.0) vec<N, T>) -> T
 @const fn log<T: fa_f32_f16>(T) -> T

src/tint/resolver/const_eval.cc

@@ -2654,6 +2654,48 @@ ConstEval::Result ConstEval::inverseSqrt(const type::Type* ty,
     return TransformElements(builder, ty, transform, args[0]);
 }
 
+ConstEval::Result ConstEval::ldexp(const type::Type* ty,
+                                   utils::VectorRef<const constant::Value*> args,
+                                   const Source& source) {
+    auto transform = [&](const constant::Value* c1, size_t index) {
+        auto create = [&](auto e1) -> ConstEval::Result {
+            using E1Type = decltype(e1);
+            // If e1 is AFloat, then e2 is AInt, otherwise it's i32
+            using E2Type = std::conditional_t<std::is_same_v<E1Type, AFloat>, AInt, i32>;
+
+            E2Type e2;
+            auto* c2 = args[1];
+            if (c2->Type()->Is<type::Vector>()) {
+                e2 = c2->Index(index)->ValueAs<E2Type>();
+            } else {
+                e2 = c2->ValueAs<E2Type>();
+            }
+
+            E2Type bias;
+            if constexpr (std::is_same_v<E1Type, f16>) {
+                bias = 15;
+            } else if constexpr (std::is_same_v<E1Type, f32>) {
+                bias = 127;
+            } else {
+                bias = 1023;
+            }
+
+            if (e2 > bias + 1) {
+                AddError("e2 must be less than or equal to " + std::to_string(bias + 1), source);
+                return utils::Failure;
+            }
+
+            auto target_ty = type::Type::DeepestElementOf(ty);
+
+            auto r = std::ldexp(e1, static_cast<int>(e2));
+            return CreateScalar(builder, source, target_ty, E1Type{r});
+        };
+        return Dispatch_fa_f32_f16(create, c1);
+    };
+
+    return TransformElements(builder, ty, transform, args[0]);
+}
+
 ConstEval::Result ConstEval::length(const type::Type* ty,
                                     utils::VectorRef<const constant::Value*> args,
                                     const Source& source) {

src/tint/resolver/const_eval.h

@@ -725,6 +725,15 @@ class ConstEval {
                        utils::VectorRef<const constant::Value*> args,
                        const Source& source);
 
+    /// ldexp builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location
+    /// @return the result value, or null if the value cannot be calculated
+    Result ldexp(const type::Type* ty,
+                 utils::VectorRef<const constant::Value*> args,
+                 const Source& source);
+
     /// length builtin
     /// @param ty the expression type
     /// @param args the input arguments

src/tint/resolver/const_eval_builtin_test.cc

@@ -1550,6 +1550,71 @@ INSTANTIATE_TEST_SUITE_P(  //
                      testing::ValuesIn(Concat(ExtractBitsCases<i32>(),  //
                                               ExtractBitsCases<u32>()))));
 
+template <typename T>
+std::vector<Case> LdexpCases() {
+    using T2 = std::conditional_t<std::is_same_v<T, AFloat>, AInt, i32>;
+    T2 bias;
+    if constexpr (std::is_same_v<T, f16>) {
+        bias = 15;
+    } else if constexpr (std::is_same_v<T, f32>) {
+        bias = 127;
+    } else {
+        bias = 1023;
+    }
+
+    auto compute = [](T e1, T2 e2) { return T{std::ldexp(e1.value, static_cast<int>(e2.value))}; };
+
+    auto r = std::vector<Case>{
+        C({T(0), T2(0)}, T(0)),          //
+        C({T(7), T2(4)}, T(112)),        //
+        C({T(7), T2(5)}, T(224)),        //
+        C({T(7), T2(6)}, T(448)),        //
+        C({T(7), T2(-4)}, T(0.4375)),    //
+        C({T(7), T2(-5)}, T(0.21875)),   //
+        C({T(7), T2(-6)}, T(0.109375)),  //
+        // With bias exponent
+        C({T(0), T2(bias)}, T(0)),                           //
+        C({T(0), T2(bias + 1)}, T(0)),                       //
+        C({T(1), T2(bias)}, compute(T(1), T2(bias))),        //
+        C({T(0.5), T2(bias)}, compute(T(0.5), T2(bias))),    //
+        C({T(0.25), T2(bias)}, compute(T(0.25), T2(bias))),  //
+        // The result may be zero if e2 + bias ≤ 0.
+        C({T(0), T2(-bias)}, T(0)),      //
+        C({T(0), T2(-bias - 1)}, T(0)),  //
+        C({T(0), T2(-bias - 2)}, T(0)),  //
+
+        // Vector tests
+        C({Vec(T(0), T(7), T(7)), Vec(T2(0), T2(4), T2(-4))}, Vec(T(0), T(112), T(0.4375))),
+    };
+
+    std::string e2_too_large_error_msg =
+        "12:34 error: e2 must be less than or equal to " + std::to_string(bias + 1);
+    auto val_overflow_error_msg = [](auto val) {
+        return "12:34 error: " + OverflowErrorMessage(val, FriendlyName<T>());
+    };
+    ConcatInto(r, std::vector<Case>{
+                      // e2 is > bias + 1
+                      E({T(0), T2(bias + 2)}, e2_too_large_error_msg),
+                      E({T(0), T2(bias + 1000)}, e2_too_large_error_msg),
+                      E({T(0), T2::Highest()}, e2_too_large_error_msg),
+                      // Result is inf
+                      E({T(1), T2(bias + 1)}, val_overflow_error_msg(T::Inf())),
+                      E({T(2), T2(bias + 1)}, val_overflow_error_msg(T::Inf())),
+                      E({T::Highest(), T2(bias + 1)}, val_overflow_error_msg(T::Inf())),
+                      E({T(-1), T2(bias + 1)}, val_overflow_error_msg(-T::Inf())),
+                      E({T(-2), T2(bias + 1)}, val_overflow_error_msg(-T::Inf())),
+                      E({T::Lowest(), T2(bias + 1)}, val_overflow_error_msg(-T::Inf())),
+                  });
+    return r;
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Ldexp,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kLdexp),
+                     testing::ValuesIn(Concat(LdexpCases<AFloat>(),  //
+                                              LdexpCases<f32>(),
+                                              LdexpCases<f16>()))));
+
 template <typename T>
 std::vector<Case> LengthCases() {
     const auto kSqrtOfHighest = T(std::sqrt(T::Highest()));

test/tint/builtins/gen/literal/ldexp/082c1f.wgsl

@@ -0,0 +1,45 @@
+// 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
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable f16;
+
+// fn ldexp(f16, ia) -> f16
+fn ldexp_082c1f() {
+  var res: f16 = ldexp(1.h, 1);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ldexp_082c1f();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ldexp_082c1f();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ldexp_082c1f();
+}

test/tint/builtins/gen/literal/ldexp/082c1f.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/082c1f.wgsl.expected.fxc.hlsl

@@ -0,0 +1,36 @@
+SKIP: FAILED
+
+void ldexp_082c1f() {
+  float16_t res = float16_t(2.0h);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ldexp_082c1f();
+  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() {
+  ldexp_082c1f();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ldexp_082c1f();
+  return;
+}
+FXC validation failure:
+C:\src\dawn\test\tint\Shader@0x000001E1B8560100(2,3-11): error X3000: unrecognized identifier 'float16_t'
+C:\src\dawn\test\tint\Shader@0x000001E1B8560100(2,13-15): error X3000: unrecognized identifier 'res'
+

test/tint/builtins/gen/literal/ldexp/082c1f.wgsl.expected.glsl

@@ -0,0 +1,52 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void ldexp_082c1f() {
+  float16_t res = 2.0hf;
+}
+
+vec4 vertex_main() {
+  ldexp_082c1f();
+  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
+#extension GL_AMD_gpu_shader_half_float : require
+precision mediump float;
+
+void ldexp_082c1f() {
+  float16_t res = 2.0hf;
+}
+
+void fragment_main() {
+  ldexp_082c1f();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void ldexp_082c1f() {
+  float16_t res = 2.0hf;
+}
+
+void compute_main() {
+  ldexp_082c1f();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/ldexp/082c1f.wgsl.expected.msl

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

test/tint/builtins/gen/literal/ldexp/082c1f.wgsl.expected.spvasm

@@ -0,0 +1,70 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 32
+; Schema: 0
+               OpCapability Shader
+               OpCapability Float16
+               OpCapability UniformAndStorageBuffer16BitAccess
+               OpCapability StorageBuffer16BitAccess
+               OpCapability StorageInputOutput16
+               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 %ldexp_082c1f "ldexp_082c1f"
+               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
+       %half = OpTypeFloat 16
+%half_0x1p_1 = OpConstant %half 0x1p+1
+%_ptr_Function_half = OpTypePointer Function %half
+         %17 = OpConstantNull %half
+         %18 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%ldexp_082c1f = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_half Function %17
+               OpStore %res %half_0x1p_1
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %18
+         %20 = OpLabel
+         %21 = OpFunctionCall %void %ldexp_082c1f
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %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
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %ldexp_082c1f
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %ldexp_082c1f
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/ldexp/082c1f.wgsl.expected.wgsl

@@ -0,0 +1,21 @@
+enable f16;
+
+fn ldexp_082c1f() {
+  var res : f16 = ldexp(1.0h, 1);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ldexp_082c1f();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ldexp_082c1f();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ldexp_082c1f();
+}

test/tint/builtins/gen/literal/ldexp/217a31.wgsl

@@ -0,0 +1,45 @@
+// 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
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable f16;
+
+// fn ldexp(vec<2, f16>, vec<2, ia>) -> vec<2, f16>
+fn ldexp_217a31() {
+  var res: vec2<f16> = ldexp(vec2<f16>(1.h), vec2(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ldexp_217a31();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ldexp_217a31();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ldexp_217a31();
+}

test/tint/builtins/gen/literal/ldexp/217a31.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/217a31.wgsl.expected.fxc.hlsl

@@ -0,0 +1,35 @@
+SKIP: FAILED
+
+void ldexp_217a31() {
+  vector<float16_t, 2> res = (float16_t(2.0h)).xx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ldexp_217a31();
+  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() {
+  ldexp_217a31();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ldexp_217a31();
+  return;
+}
+FXC validation failure:
+C:\src\dawn\test\tint\Shader@0x00000132D8550FB0(2,10-18): error X3000: syntax error: unexpected token 'float16_t'
+

test/tint/builtins/gen/literal/ldexp/217a31.wgsl.expected.glsl

@@ -0,0 +1,52 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void ldexp_217a31() {
+  f16vec2 res = f16vec2(2.0hf);
+}
+
+vec4 vertex_main() {
+  ldexp_217a31();
+  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
+#extension GL_AMD_gpu_shader_half_float : require
+precision mediump float;
+
+void ldexp_217a31() {
+  f16vec2 res = f16vec2(2.0hf);
+}
+
+void fragment_main() {
+  ldexp_217a31();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void ldexp_217a31() {
+  f16vec2 res = f16vec2(2.0hf);
+}
+
+void compute_main() {
+  ldexp_217a31();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/ldexp/217a31.wgsl.expected.msl

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

test/tint/builtins/gen/literal/ldexp/217a31.wgsl.expected.spvasm

@@ -0,0 +1,72 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 34
+; Schema: 0
+               OpCapability Shader
+               OpCapability Float16
+               OpCapability UniformAndStorageBuffer16BitAccess
+               OpCapability StorageBuffer16BitAccess
+               OpCapability StorageInputOutput16
+               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 %ldexp_217a31 "ldexp_217a31"
+               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
+       %half = OpTypeFloat 16
+     %v2half = OpTypeVector %half 2
+%half_0x1p_1 = OpConstant %half 0x1p+1
+         %16 = OpConstantComposite %v2half %half_0x1p_1 %half_0x1p_1
+%_ptr_Function_v2half = OpTypePointer Function %v2half
+         %19 = OpConstantNull %v2half
+         %20 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%ldexp_217a31 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_v2half Function %19
+               OpStore %res %16
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %20
+         %22 = OpLabel
+         %23 = OpFunctionCall %void %ldexp_217a31
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %25 = OpLabel
+         %26 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %26
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %ldexp_217a31
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %ldexp_217a31
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/ldexp/217a31.wgsl.expected.wgsl

@@ -0,0 +1,21 @@
+enable f16;
+
+fn ldexp_217a31() {
+  var res : vec2<f16> = ldexp(vec2<f16>(1.0h), vec2(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ldexp_217a31();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ldexp_217a31();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ldexp_217a31();
+}

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

test/tint/builtins/gen/literal/ldexp/2bfc68.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/2bfc68.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/2bfc68.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/ldexp/2bfc68.wgsl.expected.msl

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

test/tint/builtins/gen/literal/ldexp/2bfc68.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 %ldexp_2bfc68 "ldexp_2bfc68"
+               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
+%ldexp_2bfc68 = 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 %ldexp_2bfc68
+               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 %ldexp_2bfc68
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ldexp_2bfc68
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/ldexp/2bfc68.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/literal/ldexp/2c6370.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/2c6370.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/2c6370.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/ldexp/2c6370.wgsl.expected.msl

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

test/tint/builtins/gen/literal/ldexp/2c6370.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 %ldexp_2c6370 "ldexp_2c6370"
+               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
+%ldexp_2c6370 = 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 %ldexp_2c6370
+               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 %ldexp_2c6370
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ldexp_2c6370
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/ldexp/2c6370.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/ldexp/376938.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 ldexp(vec<4, fa>, vec<4, i32>) -> vec<4, fa>
+fn ldexp_376938() {
+  var res = ldexp(vec4(1.), vec4<i32>(1i));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ldexp_376938();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ldexp_376938();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ldexp_376938();
+}

test/tint/builtins/gen/literal/ldexp/376938.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/376938.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/376938.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/ldexp/376938.wgsl.expected.msl

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

test/tint/builtins/gen/literal/ldexp/376938.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 %ldexp_376938 "ldexp_376938"
+               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
+%ldexp_376938 = 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 %ldexp_376938
+               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 %ldexp_376938
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %ldexp_376938
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/ldexp/376938.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/ldexp/3d90b4.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/3d90b4.wgsl.expected.fxc.hlsl

@@ -1,7 +1,7 @@
 SKIP: FAILED
 
 void ldexp_3d90b4() {
-  vector<float16_t, 2> res = ldexp((float16_t(0.0h)).xx, (1).xx);
+  vector<float16_t, 2> res = (float16_t(2.0h)).xx;
 }
 
 struct tint_symbol {
@@ -30,3 +30,6 @@ void compute_main() {
   ldexp_3d90b4();
   return;
 }
+FXC validation failure:
+C:\src\dawn\test\tint\Shader@0x000001C8EE651B30(2,10-18): error X3000: syntax error: unexpected token 'float16_t'
+

test/tint/builtins/gen/literal/ldexp/3d90b4.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/ldexp/3d90b4.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ldexp_3d90b4() {
-  half2 res = ldexp(half2(1.0h), int2(1));
+  half2 res = half2(2.0h);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/ldexp/3d90b4.wgsl.expected.spvasm

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

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

test/tint/builtins/gen/literal/ldexp/4a3ad9.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/4a3ad9.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/4a3ad9.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/ldexp/4a3ad9.wgsl.expected.msl

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

test/tint/builtins/gen/literal/ldexp/4a3ad9.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 %ldexp_4a3ad9 "ldexp_4a3ad9"
+               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
+%ldexp_4a3ad9 = 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 %ldexp_4a3ad9
+               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 %ldexp_4a3ad9
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ldexp_4a3ad9
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/ldexp/4a3ad9.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/ldexp/593ff3.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 ldexp(vec<3, f32>, vec<3, ia>) -> vec<3, f32>
+fn ldexp_593ff3() {
+  var res: vec3<f32> = ldexp(vec3<f32>(1.f), vec3(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ldexp_593ff3();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ldexp_593ff3();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ldexp_593ff3();
+}

test/tint/builtins/gen/literal/ldexp/593ff3.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/593ff3.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/593ff3.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/ldexp/593ff3.wgsl.expected.msl

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

test/tint/builtins/gen/literal/ldexp/593ff3.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 %ldexp_593ff3 "ldexp_593ff3"
+               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
+%ldexp_593ff3 = 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 %ldexp_593ff3
+               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 %ldexp_593ff3
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %ldexp_593ff3
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/ldexp/593ff3.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/ldexp/624e0c.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void ldexp_624e0c() {
-  float16_t res = ldexp(float16_t(1.0h), 1);
+  float16_t res = float16_t(2.0h);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/ldexp/624e0c.wgsl.expected.fxc.hlsl

@@ -1,7 +1,7 @@
 SKIP: FAILED
 
 void ldexp_624e0c() {
-  float16_t res = ldexp(float16_t(0.0h), 1);
+  float16_t res = float16_t(2.0h);
 }
 
 struct tint_symbol {
@@ -30,3 +30,7 @@ void compute_main() {
   ldexp_624e0c();
   return;
 }
+FXC validation failure:
+C:\src\dawn\test\tint\Shader@0x00000246DEBE4390(2,3-11): error X3000: unrecognized identifier 'float16_t'
+C:\src\dawn\test\tint\Shader@0x00000246DEBE4390(2,13-15): error X3000: unrecognized identifier 'res'
+

test/tint/builtins/gen/literal/ldexp/624e0c.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/ldexp/624e0c.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ldexp_624e0c() {
-  half res = ldexp(1.0h, 1);
+  half res = 2.0h;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/ldexp/624e0c.wgsl.expected.spvasm

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

test/tint/builtins/gen/literal/ldexp/65a7bd.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 ldexp(vec<4, f32>, vec<4, ia>) -> vec<4, f32>
+fn ldexp_65a7bd() {
+  var res: vec4<f32> = ldexp(vec4<f32>(1.f), vec4(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ldexp_65a7bd();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ldexp_65a7bd();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ldexp_65a7bd();
+}

test/tint/builtins/gen/literal/ldexp/65a7bd.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/65a7bd.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/65a7bd.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/ldexp/65a7bd.wgsl.expected.msl

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

test/tint/builtins/gen/literal/ldexp/65a7bd.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 %ldexp_65a7bd "ldexp_65a7bd"
+               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
+%ldexp_65a7bd = 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 %ldexp_65a7bd
+               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 %ldexp_65a7bd
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %ldexp_65a7bd
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/ldexp/65a7bd.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/ldexp/71ebe3.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 ldexp(fa, i32) -> fa
+fn ldexp_71ebe3() {
+  var res = ldexp(1., 1i);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ldexp_71ebe3();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ldexp_71ebe3();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ldexp_71ebe3();
+}

test/tint/builtins/gen/literal/ldexp/71ebe3.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/71ebe3.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/71ebe3.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/ldexp/71ebe3.wgsl.expected.msl

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

test/tint/builtins/gen/literal/ldexp/71ebe3.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 %ldexp_71ebe3 "ldexp_71ebe3"
+               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
+%ldexp_71ebe3 = 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 %ldexp_71ebe3
+               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 %ldexp_71ebe3
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %ldexp_71ebe3
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/ldexp/71ebe3.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/ldexp/7485ce.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/7485ce.wgsl.expected.fxc.hlsl

@@ -1,7 +1,7 @@
 SKIP: FAILED
 
 void ldexp_7485ce() {
-  vector<float16_t, 3> res = ldexp((float16_t(0.0h)).xxx, (1).xxx);
+  vector<float16_t, 3> res = (float16_t(2.0h)).xxx;
 }
 
 struct tint_symbol {
@@ -30,3 +30,6 @@ void compute_main() {
   ldexp_7485ce();
   return;
 }
+FXC validation failure:
+C:\src\dawn\test\tint\Shader@0x000001A317BC93B0(2,10-18): error X3000: syntax error: unexpected token 'float16_t'
+

test/tint/builtins/gen/literal/ldexp/7485ce.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/ldexp/7485ce.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ldexp_7485ce() {
-  half3 res = ldexp(half3(1.0h), int3(1));
+  half3 res = half3(2.0h);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/ldexp/7485ce.wgsl.expected.spvasm

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

test/tint/builtins/gen/literal/ldexp/7fa13c.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void ldexp_7fa13c() {
-  vector<float16_t, 4> res = ldexp((float16_t(1.0h)).xxxx, (1).xxxx);
+  vector<float16_t, 4> res = (float16_t(2.0h)).xxxx;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/ldexp/7fa13c.wgsl.expected.fxc.hlsl

@@ -1,7 +1,7 @@
 SKIP: FAILED
 
 void ldexp_7fa13c() {
-  vector<float16_t, 4> res = ldexp((float16_t(0.0h)).xxxx, (1).xxxx);
+  vector<float16_t, 4> res = (float16_t(2.0h)).xxxx;
 }
 
 struct tint_symbol {
@@ -30,3 +30,6 @@ void compute_main() {
   ldexp_7fa13c();
   return;
 }
+FXC validation failure:
+C:\src\dawn\test\tint\Shader@0x000001D4BE9AE370(2,10-18): error X3000: syntax error: unexpected token 'float16_t'
+

test/tint/builtins/gen/literal/ldexp/7fa13c.wgsl.expected.glsl

@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void ldexp_7fa13c() {
-  f16vec4 res = ldexp(f16vec4(1.0hf), ivec4(1));
+  f16vec4 res = f16vec4(2.0hf);
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@ void main() {
 precision mediump float;
 
 void ldexp_7fa13c() {
-  f16vec4 res = ldexp(f16vec4(1.0hf), ivec4(1));
+  f16vec4 res = f16vec4(2.0hf);
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@ void main() {
 #extension GL_AMD_gpu_shader_half_float : require
 
 void ldexp_7fa13c() {
-  f16vec4 res = ldexp(f16vec4(1.0hf), ivec4(1));
+  f16vec4 res = f16vec4(2.0hf);
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/ldexp/7fa13c.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void ldexp_7fa13c() {
-  half4 res = ldexp(half4(1.0h), int4(1));
+  half4 res = half4(2.0h);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/ldexp/7fa13c.wgsl.expected.spvasm

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

test/tint/builtins/gen/literal/ldexp/8a0c2f.wgsl

@@ -0,0 +1,45 @@
+// 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
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable f16;
+
+// fn ldexp(vec<4, f16>, vec<4, ia>) -> vec<4, f16>
+fn ldexp_8a0c2f() {
+  var res: vec4<f16> = ldexp(vec4<f16>(1.h), vec4(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ldexp_8a0c2f();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ldexp_8a0c2f();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ldexp_8a0c2f();
+}

test/tint/builtins/gen/literal/ldexp/8a0c2f.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/8a0c2f.wgsl.expected.fxc.hlsl

@@ -0,0 +1,35 @@
+SKIP: FAILED
+
+void ldexp_8a0c2f() {
+  vector<float16_t, 4> res = (float16_t(2.0h)).xxxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ldexp_8a0c2f();
+  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() {
+  ldexp_8a0c2f();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ldexp_8a0c2f();
+  return;
+}
+FXC validation failure:
+C:\src\dawn\test\tint\Shader@0x0000021161096930(2,10-18): error X3000: syntax error: unexpected token 'float16_t'
+

test/tint/builtins/gen/literal/ldexp/8a0c2f.wgsl.expected.glsl

@@ -0,0 +1,52 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void ldexp_8a0c2f() {
+  f16vec4 res = f16vec4(2.0hf);
+}
+
+vec4 vertex_main() {
+  ldexp_8a0c2f();
+  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
+#extension GL_AMD_gpu_shader_half_float : require
+precision mediump float;
+
+void ldexp_8a0c2f() {
+  f16vec4 res = f16vec4(2.0hf);
+}
+
+void fragment_main() {
+  ldexp_8a0c2f();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void ldexp_8a0c2f() {
+  f16vec4 res = f16vec4(2.0hf);
+}
+
+void compute_main() {
+  ldexp_8a0c2f();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/ldexp/8a0c2f.wgsl.expected.msl

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

test/tint/builtins/gen/literal/ldexp/8a0c2f.wgsl.expected.spvasm

@@ -0,0 +1,72 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 34
+; Schema: 0
+               OpCapability Shader
+               OpCapability Float16
+               OpCapability UniformAndStorageBuffer16BitAccess
+               OpCapability StorageBuffer16BitAccess
+               OpCapability StorageInputOutput16
+               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 %ldexp_8a0c2f "ldexp_8a0c2f"
+               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
+       %half = OpTypeFloat 16
+     %v4half = OpTypeVector %half 4
+%half_0x1p_1 = OpConstant %half 0x1p+1
+         %16 = OpConstantComposite %v4half %half_0x1p_1 %half_0x1p_1 %half_0x1p_1 %half_0x1p_1
+%_ptr_Function_v4half = OpTypePointer Function %v4half
+         %19 = OpConstantNull %v4half
+         %20 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%ldexp_8a0c2f = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_v4half Function %19
+               OpStore %res %16
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %20
+         %22 = OpLabel
+         %23 = OpFunctionCall %void %ldexp_8a0c2f
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %25 = OpLabel
+         %26 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %26
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %ldexp_8a0c2f
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %ldexp_8a0c2f
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/ldexp/8a0c2f.wgsl.expected.wgsl

@@ -0,0 +1,21 @@
+enable f16;
+
+fn ldexp_8a0c2f() {
+  var res : vec4<f16> = ldexp(vec4<f16>(1.0h), vec4(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ldexp_8a0c2f();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ldexp_8a0c2f();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ldexp_8a0c2f();
+}

test/tint/builtins/gen/literal/ldexp/8e43e9.wgsl

@@ -0,0 +1,45 @@
+// 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
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable f16;
+
+// fn ldexp(vec<3, f16>, vec<3, ia>) -> vec<3, f16>
+fn ldexp_8e43e9() {
+  var res: vec3<f16> = ldexp(vec3<f16>(1.h), vec3(1));
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  ldexp_8e43e9();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  ldexp_8e43e9();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  ldexp_8e43e9();
+}

test/tint/builtins/gen/literal/ldexp/8e43e9.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/ldexp/8e43e9.wgsl.expected.fxc.hlsl

@@ -0,0 +1,35 @@
+SKIP: FAILED
+
+void ldexp_8e43e9() {
+  vector<float16_t, 3> res = (float16_t(2.0h)).xxx;
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+  ldexp_8e43e9();
+  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() {
+  ldexp_8e43e9();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  ldexp_8e43e9();
+  return;
+}
+FXC validation failure:
+C:\src\dawn\test\tint\Shader@0x00000172CFF90020(2,10-18): error X3000: syntax error: unexpected token 'float16_t'
+

test/tint/builtins/gen/literal/ldexp/8e43e9.wgsl.expected.glsl

@@ -0,0 +1,52 @@
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void ldexp_8e43e9() {
+  f16vec3 res = f16vec3(2.0hf);
+}
+
+vec4 vertex_main() {
+  ldexp_8e43e9();
+  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
+#extension GL_AMD_gpu_shader_half_float : require
+precision mediump float;
+
+void ldexp_8e43e9() {
+  f16vec3 res = f16vec3(2.0hf);
+}
+
+void fragment_main() {
+  ldexp_8e43e9();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+void ldexp_8e43e9() {
+  f16vec3 res = f16vec3(2.0hf);
+}
+
+void compute_main() {
+  ldexp_8e43e9();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/ldexp/8e43e9.wgsl.expected.msl

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