tint: Implement runtime quantizeToF16()

Fixed: tint:991
Fixed: tint:1741
Change-Id: I55dbabce6d28adf5abb710dc1e3e879c5aaa6be5
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/108140
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>

src/tint/intrinsics.def

@@ -514,6 +514,8 @@ fn pack4x8snorm(vec4<f32>) -> u32
 fn pack4x8unorm(vec4<f32>) -> u32
 fn pow<T: f32_f16>(T, T) -> T
 fn pow<N: num, T: f32_f16>(vec<N, T>, vec<N, T>) -> vec<N, T>
+fn quantizeToF16(f32) -> f32
+fn quantizeToF16<N: num>(vec<N, f32>) -> vec<N, f32>
 fn radians<T: f32_f16>(T) -> T
 fn radians<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
 fn reflect<N: num, T: f32_f16>(vec<N, T>, vec<N, T>) -> vec<N, T>

src/tint/reader/spirv/function_conversion_test.cc

@@ -615,7 +615,6 @@ OpFunctionEnd
 // TODO(dneto): OpSConvert // only if multiple widths
 // TODO(dneto): OpUConvert // only if multiple widths
 // TODO(dneto): OpFConvert // only if multiple widths
-// TODO(dneto): OpQuantizeToF16 // only if f16 supported
 // TODO(dneto): OpSatConvertSToU // Kernel (OpenCL), not in WebGPU
 // TODO(dneto): OpSatConvertUToS // Kernel (OpenCL), not in WebGPU
 

src/tint/sem/builtin_type.cc

@@ -210,6 +210,9 @@ BuiltinType ParseBuiltinType(const std::string& name) {
     if (name == "pow") {
         return BuiltinType::kPow;
     }
+    if (name == "quantizeToF16") {
+        return BuiltinType::kQuantizeToF16;
+    }
     if (name == "radians") {
         return BuiltinType::kRadians;
     }
@@ -492,6 +495,8 @@ const char* str(BuiltinType i) {
             return "pack4x8unorm";
         case BuiltinType::kPow:
             return "pow";
+        case BuiltinType::kQuantizeToF16:
+            return "quantizeToF16";
         case BuiltinType::kRadians:
             return "radians";
         case BuiltinType::kReflect:

src/tint/sem/builtin_type.h

@@ -92,6 +92,7 @@ enum class BuiltinType {
     kPack4X8Snorm,
     kPack4X8Unorm,
     kPow,
+    kQuantizeToF16,
     kRadians,
     kReflect,
     kRefract,

src/tint/transform/builtin_polyfill.cc

@@ -556,6 +556,27 @@ struct BuiltinPolyfill::State {
         return name;
     }
 
+    /// Builds the polyfill function for the `quantizeToF16` builtin, by replacing the vector form
+    /// with scalar calls.
+    /// @param vec the vector type
+    /// @return the polyfill function name
+    Symbol quantizeToF16(const sem::Vector* vec) {
+        auto name = b.Symbols().New("tint_quantizeToF16");
+        utils::Vector<const ast::Expression*, 4> args;
+        for (uint32_t i = 0; i < vec->Width(); i++) {
+            args.Push(b.Call("quantizeToF16", b.IndexAccessor("v", u32(i))));
+        }
+        b.Func(name,
+               utils::Vector{
+                   b.Param("v", T(vec)),
+               },
+               T(vec),
+               utils::Vector{
+                   b.Return(b.Construct(T(vec), std::move(args))),
+               });
+        return name;
+    }
+
   private:
     /// @returns the AST type for the given sem type
     const ast::Type* T(const sem::Type* ty) const { return CreateASTTypeFor(ctx, ty); }
@@ -659,6 +680,13 @@ bool BuiltinPolyfill::ShouldRun(const Program* program, const DataMap& data) con
                                 }
                             }
                             break;
+                        case sem::BuiltinType::kQuantizeToF16:
+                            if (builtins.quantize_to_vec_f16) {
+                                if (builtin->ReturnType()->Is<sem::Vector>()) {
+                                    return true;
+                                }
+                            }
+                            break;
                         default:
                             break;
                     }
@@ -778,6 +806,15 @@ void BuiltinPolyfill::Run(CloneContext& ctx, const DataMap& data, DataMap&) cons
                             }
                         }
                         break;
+                    case sem::BuiltinType::kQuantizeToF16:
+                        if (builtins.quantize_to_vec_f16) {
+                            if (auto* vec = builtin->ReturnType()->As<sem::Vector>()) {
+                                polyfill = utils::GetOrCreate(polyfills, builtin,
+                                                              [&] { return s.quantizeToF16(vec); });
+                            }
+                        }
+                        break;
+
                     default:
                         break;
                 }

src/tint/transform/builtin_polyfill.h

@@ -65,6 +65,9 @@ class BuiltinPolyfill final : public Castable<BuiltinPolyfill, Transform> {
         bool saturate = false;
         /// Should `textureSampleBaseClampToEdge()` be polyfilled for texture_2d<f32> textures?
         bool texture_sample_base_clamp_to_edge_2d_f32 = false;
+        /// Should the vector form of `quantizeToF16()` be polyfilled with a scalar implementation?
+         /// See crbug.com/tint/1741
+        bool quantize_to_vec_f16 = false;
     };
 
     /// Config is consumed by the BuiltinPolyfill transform.

src/tint/transform/builtin_polyfill_test.cc

@@ -1996,5 +1996,112 @@ fn f() {
     EXPECT_EQ(expect, str(got));
 }
 
+////////////////////////////////////////////////////////////////////////////////
+// quantizeToF16
+////////////////////////////////////////////////////////////////////////////////
+DataMap polyfillQuantizeToF16_2d_f32() {
+    BuiltinPolyfill::Builtins builtins;
+    builtins.quantize_to_vec_f16 = true;
+    DataMap data;
+    data.Add<BuiltinPolyfill::Config>(builtins);
+    return data;
+}
+
+TEST_F(BuiltinPolyfillTest, ShouldRunQuantizeToF16_Scalar) {
+    auto* src = R"(
+fn f() {
+  let v = 0.5;
+  quantizeToF16(0.5);
+}
+)";
+
+    EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src));
+    EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src, polyfillQuantizeToF16_2d_f32()));
+}
+
+TEST_F(BuiltinPolyfillTest, ShouldRunQuantizeToF16_Vector) {
+    auto* src = R"(
+fn f() {
+  let v = 0.5;
+  quantizeToF16(vec2(v));
+}
+)";
+
+    EXPECT_FALSE(ShouldRun<BuiltinPolyfill>(src));
+    EXPECT_TRUE(ShouldRun<BuiltinPolyfill>(src, polyfillQuantizeToF16_2d_f32()));
+}
+
+TEST_F(BuiltinPolyfillTest, QuantizeToF16_Vec2) {
+    auto* src = R"(
+fn f() {
+  let v = 0.5;
+  quantizeToF16(vec2(v));
+}
+)";
+
+    auto* expect = R"(
+fn tint_quantizeToF16(v : vec2<f32>) -> vec2<f32> {
+  return vec2<f32>(quantizeToF16(v[0u]), quantizeToF16(v[1u]));
+}
+
+fn f() {
+  let v = 0.5;
+  tint_quantizeToF16(vec2(v));
+}
+)";
+
+    auto got = Run<BuiltinPolyfill>(src, polyfillQuantizeToF16_2d_f32());
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, QuantizeToF16_Vec3) {
+    auto* src = R"(
+fn f() {
+  let v = 0.5;
+  quantizeToF16(vec3(v));
+}
+)";
+
+    auto* expect = R"(
+fn tint_quantizeToF16(v : vec3<f32>) -> vec3<f32> {
+  return vec3<f32>(quantizeToF16(v[0u]), quantizeToF16(v[1u]), quantizeToF16(v[2u]));
+}
+
+fn f() {
+  let v = 0.5;
+  tint_quantizeToF16(vec3(v));
+}
+)";
+
+    auto got = Run<BuiltinPolyfill>(src, polyfillQuantizeToF16_2d_f32());
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(BuiltinPolyfillTest, QuantizeToF16_Vec4) {
+    auto* src = R"(
+fn f() {
+  let v = 0.5;
+  quantizeToF16(vec4(v));
+}
+)";
+
+    auto* expect = R"(
+fn tint_quantizeToF16(v : vec4<f32>) -> vec4<f32> {
+  return vec4<f32>(quantizeToF16(v[0u]), quantizeToF16(v[1u]), quantizeToF16(v[2u]), quantizeToF16(v[3u]));
+}
+
+fn f() {
+  let v = 0.5;
+  tint_quantizeToF16(vec4(v));
+}
+)";
+
+    auto got = Run<BuiltinPolyfill>(src, polyfillQuantizeToF16_2d_f32());
+
+    EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace tint::transform

src/tint/writer/glsl/generator_impl.cc

@@ -796,6 +796,9 @@ bool GeneratorImpl::EmitBuiltinCall(std::ostream& out,
     if (builtin->Type() == sem::BuiltinType::kRadians) {
         return EmitRadiansCall(out, expr, builtin);
     }
+    if (builtin->Type() == sem::BuiltinType::kQuantizeToF16) {
+        return EmitQuantizeToF16Call(out, expr, builtin);
+    }
     if (builtin->Type() == sem::BuiltinType::kArrayLength) {
         return EmitArrayLength(out, expr);
     }
@@ -1287,6 +1290,38 @@ bool GeneratorImpl::EmitRadiansCall(std::ostream& out,
                              });
 }
 
+bool GeneratorImpl::EmitQuantizeToF16Call(std::ostream& out,
+                                          const ast::CallExpression* expr,
+                                          const sem::Builtin* builtin) {
+    // Emulate by casting to f16 and back again.
+    return CallBuiltinHelper(
+        out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
+            const auto v = params[0];
+            if (auto* vec = builtin->ReturnType()->As<sem::Vector>()) {
+                switch (vec->Width()) {
+                    case 2: {
+                        line(b) << "return unpackHalf2x16(packHalf2x16(" << v << "));";
+                        return true;
+                    }
+                    case 3: {
+                        line(b) << "return vec3(";
+                        line(b) << "  unpackHalf2x16(packHalf2x16(" << v << ".xy)),";
+                        line(b) << "  unpackHalf2x16(packHalf2x16(" << v << ".zz)).x);";
+                        return true;
+                    }
+                    default: {
+                        line(b) << "return vec4(";
+                        line(b) << "  unpackHalf2x16(packHalf2x16(" << v << ".xy)),";
+                        line(b) << "  unpackHalf2x16(packHalf2x16(" << v << ".zw)));";
+                        return true;
+                    }
+                }
+            }
+            line(b) << "return unpackHalf2x16(packHalf2x16(vec2(" << v << "))).x;";
+            return true;
+        });
+}
+
 bool GeneratorImpl::EmitBarrierCall(std::ostream& out, const sem::Builtin* builtin) {
     // TODO(crbug.com/tint/661): Combine sequential barriers to a single
     // instruction.

src/tint/writer/glsl/generator_impl.h

@@ -270,6 +270,14 @@ class GeneratorImpl : public TextGenerator {
     bool EmitRadiansCall(std::ostream& out,
                          const ast::CallExpression* expr,
                          const sem::Builtin* builtin);
+    /// Handles generating a call to the `quantizeToF16()` intrinsic
+    /// @param out the output of the expression stream
+    /// @param expr the call expression
+    /// @param builtin the semantic information for the builtin
+    /// @returns true if the call expression is emitted
+    bool EmitQuantizeToF16Call(std::ostream& out,
+                               const ast::CallExpression* expr,
+                               const sem::Builtin* builtin);
     /// Handles a case statement
     /// @param stmt the statement
     /// @returns true if the statement was emitted successfully

src/tint/writer/glsl/generator_impl_builtin_test.cc

@@ -1266,5 +1266,118 @@ void main() {
 )");
 }
 
+TEST_F(GlslGeneratorImplTest_Builtin, QuantizeToF16_Scalar) {
+    GlobalVar("v", Expr(2_f), ast::AddressSpace::kPrivate);
+    WrapInFunction(Call("quantizeToF16", "v"));
+
+    GeneratorImpl& gen = SanitizeAndBuild();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+    EXPECT_EQ(gen.result(), R"(#version 310 es
+
+float tint_quantizeToF16(float param_0) {
+  return unpackHalf2x16(packHalf2x16(vec2(param_0))).x;
+}
+
+
+float v = 2.0f;
+void test_function() {
+  float tint_symbol = tint_quantizeToF16(v);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  test_function();
+  return;
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Builtin, QuantizeToF16_Vec2) {
+    GlobalVar("v", vec2<f32>(2_f), ast::AddressSpace::kPrivate);
+    WrapInFunction(Call("quantizeToF16", "v"));
+
+    GeneratorImpl& gen = SanitizeAndBuild();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+    EXPECT_EQ(gen.result(), R"(#version 310 es
+
+vec2 tint_quantizeToF16(vec2 param_0) {
+  return unpackHalf2x16(packHalf2x16(param_0));
+}
+
+
+vec2 v = vec2(2.0f);
+void test_function() {
+  vec2 tint_symbol = tint_quantizeToF16(v);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  test_function();
+  return;
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Builtin, QuantizeToF16_Vec3) {
+    GlobalVar("v", vec3<f32>(2_f), ast::AddressSpace::kPrivate);
+    WrapInFunction(Call("quantizeToF16", "v"));
+
+    GeneratorImpl& gen = SanitizeAndBuild();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+    EXPECT_EQ(gen.result(), R"(#version 310 es
+
+vec3 tint_quantizeToF16(vec3 param_0) {
+  return vec3(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zz)).x);
+}
+
+
+vec3 v = vec3(2.0f);
+void test_function() {
+  vec3 tint_symbol = tint_quantizeToF16(v);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  test_function();
+  return;
+}
+)");
+}
+
+
+TEST_F(GlslGeneratorImplTest_Builtin, QuantizeToF16_Vec4) {
+    GlobalVar("v", vec4<f32>(2_f), ast::AddressSpace::kPrivate);
+    WrapInFunction(Call("quantizeToF16", "v"));
+
+    GeneratorImpl& gen = SanitizeAndBuild();
+
+    ASSERT_TRUE(gen.Generate()) << gen.error();
+    EXPECT_EQ(gen.result(), R"(#version 310 es
+
+vec4 tint_quantizeToF16(vec4 param_0) {
+  return vec4(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zw)));
+}
+
+
+vec4 v = vec4(2.0f);
+void test_function() {
+  vec4 tint_symbol = tint_quantizeToF16(v);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  test_function();
+  return;
+}
+)");
+}
+
 }  // namespace
 }  // namespace tint::writer::glsl

src/tint/writer/hlsl/generator_impl.cc

@@ -1047,6 +1047,9 @@ bool GeneratorImpl::EmitBuiltinCall(std::ostream& out,
     if (type == sem::BuiltinType::kRadians) {
         return EmitRadiansCall(out, expr, builtin);
     }
+    if (type == sem::BuiltinType::kQuantizeToF16) {
+        return EmitQuantizeToF16Call(out, expr, builtin);
+    }
     if (builtin->IsDataPacking()) {
         return EmitDataPackingCall(out, expr, builtin);
     }
@@ -1940,6 +1943,22 @@ bool GeneratorImpl::EmitRadiansCall(std::ostream& out,
                              });
 }
 
+bool GeneratorImpl::EmitQuantizeToF16Call(std::ostream& out,
+                                          const ast::CallExpression* expr,
+                                          const sem::Builtin* builtin) {
+    // Emulate by casting to min16float and back again.
+    std::string width;
+    if (auto* vec = builtin->ReturnType()->As<sem::Vector>()) {
+        width = std::to_string(vec->Width());
+    }
+    out << "float" << width << "(min16float" << width << "(";
+    if (!EmitExpression(out, expr->args[0])) {
+        return false;
+    }
+    out << "))";
+    return true;
+}
+
 bool GeneratorImpl::EmitDataPackingCall(std::ostream& out,
                                         const ast::CallExpression* expr,
                                         const sem::Builtin* builtin) {

src/tint/writer/hlsl/generator_impl.h

@@ -267,6 +267,14 @@ class GeneratorImpl : public TextGenerator {
     bool EmitDataUnpackingCall(std::ostream& out,
                                const ast::CallExpression* expr,
                                const sem::Builtin* builtin);
+    /// Handles generating a call to the `quantizeToF16()` intrinsic
+    /// @param out the output of the expression stream
+    /// @param expr the call expression
+    /// @param builtin the semantic information for the builtin
+    /// @returns true if the call expression is emitted
+    bool EmitQuantizeToF16Call(std::ostream& out,
+                               const ast::CallExpression* expr,
+                               const sem::Builtin* builtin);
     /// Handles generating a call to DP4a builtins (dot4I8Packed and dot4U8Packed)
     /// @param out the output of the expression stream
     /// @param expr the call expression

src/tint/writer/hlsl/generator_impl_import_test.cc

@@ -268,5 +268,31 @@ TEST_F(HlslGeneratorImplTest_Import, HlslImportData_Determinant) {
     EXPECT_EQ(out.str(), std::string("determinant(var)"));
 }
 
+TEST_F(HlslGeneratorImplTest_Import, HlslImportData_QuantizeToF16_Scalar) {
+    GlobalVar("v", Expr(2_f), ast::AddressSpace::kPrivate);
+
+    auto* expr = Call("quantizeToF16", "v");
+    WrapInFunction(expr);
+
+    GeneratorImpl& gen = Build();
+
+    std::stringstream out;
+    ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+    EXPECT_EQ(out.str(), std::string("float(min16float(v))"));
+}
+
+TEST_F(HlslGeneratorImplTest_Import, HlslImportData_QuantizeToF16_Vector) {
+    GlobalVar("v", vec3<f32>(2_f), ast::AddressSpace::kPrivate);
+
+    auto* expr = Call("quantizeToF16", "v");
+    WrapInFunction(expr);
+
+    GeneratorImpl& gen = Build();
+
+    std::stringstream out;
+    ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+    EXPECT_EQ(out.str(), std::string("float3(min16float3(v))"));
+}
+
 }  // namespace
 }  // namespace tint::writer::hlsl

src/tint/writer/msl/generator_impl.cc

@@ -691,6 +691,18 @@ bool GeneratorImpl::EmitBuiltinCall(std::ostream& out,
             out << "))";
             return true;
         }
+        case sem::BuiltinType::kQuantizeToF16: {
+            std::string width = "";
+            if (auto* vec = builtin->ReturnType()->As<sem::Vector>()) {
+                width = std::to_string(vec->Width());
+            }
+            out << "float" << width << "(half" << width << "(";
+            if (!EmitExpression(out, expr->args[0])) {
+                return false;
+            }
+            out << "))";
+            return true;
+        }
         // TODO(crbug.com/tint/661): Combine sequential barriers to a single
         // instruction.
         case sem::BuiltinType::kStorageBarrier: {

src/tint/writer/msl/generator_impl_import_test.cc

@@ -247,5 +247,31 @@ TEST_F(MslGeneratorImplTest, MslImportData_Determinant) {
     EXPECT_EQ(out.str(), std::string("determinant(var)"));
 }
 
+TEST_F(MslGeneratorImplTest, MslImportData_QuantizeToF16_Scalar) {
+    GlobalVar("v", Expr(2_f), ast::AddressSpace::kPrivate);
+
+    auto* expr = Call("quantizeToF16", "v");
+    WrapInFunction(expr);
+
+    GeneratorImpl& gen = Build();
+
+    std::stringstream out;
+    ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+    EXPECT_EQ(out.str(), "float(half(v))");
+}
+
+TEST_F(MslGeneratorImplTest, MslImportData_QuantizeToF16_Vector) {
+    GlobalVar("v", vec3<f32>(2_f), ast::AddressSpace::kPrivate);
+
+    auto* expr = Call("quantizeToF16", "v");
+    WrapInFunction(expr);
+
+    GeneratorImpl& gen = Build();
+
+    std::stringstream out;
+    ASSERT_TRUE(gen.EmitCall(out, expr)) << gen.error();
+    EXPECT_EQ(out.str(), "float3(half3(v))");
+}
+
 }  // namespace
 }  // namespace tint::writer::msl

src/tint/writer/spirv/builder.cc

@@ -2505,6 +2505,9 @@ uint32_t Builder::GenerateBuiltinCall(const sem::Call* call, const sem::Builtin*
             }
             return result_id;
         }
+        case BuiltinType::kQuantizeToF16:
+            op = spv::Op::OpQuantizeToF16;
+            break;
         case BuiltinType::kReverseBits:
             op = spv::Op::OpBitReverse;
             break;

src/tint/writer/spirv/builder_builtin_test.cc

@@ -1854,6 +1854,105 @@ OpFunctionEnd
     Validate(b);
 }
 
+TEST_F(BuiltinBuilderTest, Call_QuantizeToF16_Scalar) {
+    GlobalVar("v", Expr(2_f), ast::AddressSpace::kPrivate);
+
+    Func("a_func", utils::Empty, ty.void_(),
+         utils::Vector{
+             Decl(Let("l", Call("quantizeToF16", "v"))),
+         },
+         utils::Vector{
+             Stage(ast::PipelineStage::kFragment),
+         });
+
+    spirv::Builder& b = SanitizeAndBuild();
+
+    ASSERT_TRUE(b.Build()) << b.error();
+    auto got = DumpBuilder(b);
+    auto* expect = R"(OpCapability Shader
+OpMemoryModel Logical GLSL450
+OpEntryPoint Fragment %7 "a_func"
+OpExecutionMode %7 OriginUpperLeft
+OpName %3 "v"
+OpName %7 "a_func"
+%1 = OpTypeFloat 32
+%2 = OpConstant %1 2
+%4 = OpTypePointer Private %1
+%3 = OpVariable %4 Private %2
+%6 = OpTypeVoid
+%5 = OpTypeFunction %6
+%7 = OpFunction %6 None %5
+%8 = OpLabel
+%10 = OpLoad %1 %3
+%9 = OpQuantizeToF16 %1 %10
+OpReturn
+OpFunctionEnd
+)";
+    EXPECT_EQ(expect, got);
+
+    Validate(b);
+}
+
+TEST_F(BuiltinBuilderTest, Call_QuantizeToF16_Vector) {
+    GlobalVar("v", vec3<f32>(2_f), ast::AddressSpace::kPrivate);
+
+    Func("a_func", utils::Empty, ty.void_(),
+         utils::Vector{
+             Decl(Let("l", Call("quantizeToF16", "v"))),
+         },
+         utils::Vector{
+             Stage(ast::PipelineStage::kFragment),
+         });
+
+    spirv::Builder& b = SanitizeAndBuild();
+
+    ASSERT_TRUE(b.Build()) << b.error();
+    auto got = DumpBuilder(b);
+    auto* expect = R"(OpCapability Shader
+OpMemoryModel Logical GLSL450
+OpEntryPoint Fragment %24 "a_func"
+OpExecutionMode %24 OriginUpperLeft
+OpName %5 "v"
+OpName %8 "tint_quantizeToF16"
+OpName %9 "v_1"
+OpName %24 "a_func"
+%2 = OpTypeFloat 32
+%1 = OpTypeVector %2 3
+%3 = OpConstant %2 2
+%4 = OpConstantComposite %1 %3 %3 %3
+%6 = OpTypePointer Private %1
+%5 = OpVariable %6 Private %4
+%7 = OpTypeFunction %1 %1
+%12 = OpTypeInt 32 0
+%13 = OpConstantNull %12
+%16 = OpConstant %12 1
+%19 = OpConstant %12 2
+%23 = OpTypeVoid
+%22 = OpTypeFunction %23
+%8 = OpFunction %1 None %7
+%9 = OpFunctionParameter %1
+%10 = OpLabel
+%14 = OpCompositeExtract %2 %9 0
+%11 = OpQuantizeToF16 %2 %14
+%17 = OpCompositeExtract %2 %9 1
+%15 = OpQuantizeToF16 %2 %17
+%20 = OpCompositeExtract %2 %9 2
+%18 = OpQuantizeToF16 %2 %20
+%21 = OpCompositeConstruct %1 %11 %15 %18
+OpReturnValue %21
+OpFunctionEnd
+%24 = OpFunction %23 None %22
+%25 = OpLabel
+%27 = OpLoad %1 %5
+%26 = OpFunctionCall %1 %8 %27
+OpReturn
+OpFunctionEnd
+)";
+    EXPECT_EQ(expect, got);
+
+    Validate(b);
+}
+
 }  // namespace float_builtin_tests
 
 // Tests for Numeric builtins with all integer parameter

src/tint/writer/spirv/generator_impl.cc

@@ -61,6 +61,7 @@ SanitizedResult Sanitize(const Program* in, const Options& options) {
         polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
         polyfills.saturate = true;
         polyfills.texture_sample_base_clamp_to_edge_2d_f32 = true;
+        polyfills.quantize_to_vec_f16 = true;  // crbug.com/tint/1741
         data.Add<transform::BuiltinPolyfill::Config>(polyfills);
         manager.Add<transform::BuiltinPolyfill>();
     }

test/tint/builtins/gen/literal/quantizeToF16/12e50e.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 quantizeToF16(f32) -> f32
+fn quantizeToF16_12e50e() {
+  var res: f32 = quantizeToF16(1.f);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  quantizeToF16_12e50e();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  quantizeToF16_12e50e();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  quantizeToF16_12e50e();
+}

test/tint/builtins/gen/literal/quantizeToF16/12e50e.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/quantizeToF16/12e50e.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/quantizeToF16/12e50e.wgsl.expected.glsl

@@ -0,0 +1,64 @@
+#version 310 es
+
+float tint_quantizeToF16(float param_0) {
+  return unpackHalf2x16(packHalf2x16(vec2(param_0))).x;
+}
+
+
+void quantizeToF16_12e50e() {
+  float res = tint_quantizeToF16(1.0f);
+}
+
+vec4 vertex_main() {
+  quantizeToF16_12e50e();
+  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;
+
+float tint_quantizeToF16(float param_0) {
+  return unpackHalf2x16(packHalf2x16(vec2(param_0))).x;
+}
+
+
+void quantizeToF16_12e50e() {
+  float res = tint_quantizeToF16(1.0f);
+}
+
+void fragment_main() {
+  quantizeToF16_12e50e();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+float tint_quantizeToF16(float param_0) {
+  return unpackHalf2x16(packHalf2x16(vec2(param_0))).x;
+}
+
+
+void quantizeToF16_12e50e() {
+  float res = tint_quantizeToF16(1.0f);
+}
+
+void compute_main() {
+  quantizeToF16_12e50e();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/quantizeToF16/12e50e.wgsl.expected.msl

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

test/tint/builtins/gen/literal/quantizeToF16/12e50e.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 %quantizeToF16_12e50e "quantizeToF16_12e50e"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+%_ptr_Function_float = OpTypePointer Function %float
+         %17 = OpTypeFunction %v4float
+%quantizeToF16_12e50e = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_float Function %8
+         %13 = OpQuantizeToF16 %float %float_1
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %17
+         %19 = OpLabel
+         %20 = OpFunctionCall %void %quantizeToF16_12e50e
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %22 = OpLabel
+         %23 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %23
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %25 = OpLabel
+         %26 = OpFunctionCall %void %quantizeToF16_12e50e
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %quantizeToF16_12e50e
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/quantizeToF16/12e50e.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/literal/quantizeToF16/2cddf3.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/quantizeToF16/2cddf3.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/quantizeToF16/2cddf3.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/quantizeToF16/2cddf3.wgsl.expected.msl

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

test/tint/builtins/gen/literal/quantizeToF16/2cddf3.wgsl.expected.spvasm

@@ -0,0 +1,83 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 45
+; 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 %tint_quantizeToF16 "tint_quantizeToF16"
+               OpName %v "v"
+               OpName %quantizeToF16_2cddf3 "quantizeToF16_2cddf3"
+               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
+    %v2float = OpTypeVector %float 2
+          %9 = OpTypeFunction %v2float %v2float
+       %uint = OpTypeInt 32 0
+         %16 = OpConstantNull %uint
+     %uint_1 = OpConstant %uint 1
+       %void = OpTypeVoid
+         %22 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+         %28 = OpConstantComposite %v2float %float_1 %float_1
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+         %31 = OpConstantNull %v2float
+         %32 = OpTypeFunction %v4float
+%tint_quantizeToF16 = OpFunction %v2float None %9
+          %v = OpFunctionParameter %v2float
+         %13 = OpLabel
+         %17 = OpCompositeExtract %float %v 0
+         %14 = OpQuantizeToF16 %float %17
+         %20 = OpCompositeExtract %float %v 1
+         %18 = OpQuantizeToF16 %float %20
+         %21 = OpCompositeConstruct %v2float %14 %18
+               OpReturnValue %21
+               OpFunctionEnd
+%quantizeToF16_2cddf3 = OpFunction %void None %22
+         %25 = OpLabel
+        %res = OpVariable %_ptr_Function_v2float Function %31
+         %26 = OpFunctionCall %v2float %tint_quantizeToF16 %28
+               OpStore %res %26
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %32
+         %34 = OpLabel
+         %35 = OpFunctionCall %void %quantizeToF16_2cddf3
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %22
+         %37 = OpLabel
+         %38 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %38
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %22
+         %40 = OpLabel
+         %41 = OpFunctionCall %void %quantizeToF16_2cddf3
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %22
+         %43 = OpLabel
+         %44 = OpFunctionCall %void %quantizeToF16_2cddf3
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/quantizeToF16/2cddf3.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/literal/quantizeToF16/cba294.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/quantizeToF16/cba294.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/quantizeToF16/cba294.wgsl.expected.glsl

@@ -0,0 +1,70 @@
+#version 310 es
+
+vec4 tint_quantizeToF16(vec4 param_0) {
+  return vec4(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zw)));
+}
+
+
+void quantizeToF16_cba294() {
+  vec4 res = tint_quantizeToF16(vec4(1.0f));
+}
+
+vec4 vertex_main() {
+  quantizeToF16_cba294();
+  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;
+
+vec4 tint_quantizeToF16(vec4 param_0) {
+  return vec4(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zw)));
+}
+
+
+void quantizeToF16_cba294() {
+  vec4 res = tint_quantizeToF16(vec4(1.0f));
+}
+
+void fragment_main() {
+  quantizeToF16_cba294();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+vec4 tint_quantizeToF16(vec4 param_0) {
+  return vec4(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zw)));
+}
+
+
+void quantizeToF16_cba294() {
+  vec4 res = tint_quantizeToF16(vec4(1.0f));
+}
+
+void compute_main() {
+  quantizeToF16_cba294();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/quantizeToF16/cba294.wgsl.expected.msl

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

test/tint/builtins/gen/literal/quantizeToF16/cba294.wgsl.expected.spvasm

@@ -0,0 +1,87 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 49
+; 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 %tint_quantizeToF16 "tint_quantizeToF16"
+               OpName %v "v"
+               OpName %quantizeToF16_cba294 "quantizeToF16_cba294"
+               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
+          %9 = OpTypeFunction %v4float %v4float
+       %uint = OpTypeInt 32 0
+         %15 = OpConstantNull %uint
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+       %void = OpTypeVoid
+         %27 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+         %33 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %36 = OpTypeFunction %v4float
+%tint_quantizeToF16 = OpFunction %v4float None %9
+          %v = OpFunctionParameter %v4float
+         %12 = OpLabel
+         %16 = OpCompositeExtract %float %v 0
+         %13 = OpQuantizeToF16 %float %16
+         %19 = OpCompositeExtract %float %v 1
+         %17 = OpQuantizeToF16 %float %19
+         %22 = OpCompositeExtract %float %v 2
+         %20 = OpQuantizeToF16 %float %22
+         %25 = OpCompositeExtract %float %v 3
+         %23 = OpQuantizeToF16 %float %25
+         %26 = OpCompositeConstruct %v4float %13 %17 %20 %23
+               OpReturnValue %26
+               OpFunctionEnd
+%quantizeToF16_cba294 = OpFunction %void None %27
+         %30 = OpLabel
+        %res = OpVariable %_ptr_Function_v4float Function %5
+         %31 = OpFunctionCall %v4float %tint_quantizeToF16 %33
+               OpStore %res %31
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %36
+         %38 = OpLabel
+         %39 = OpFunctionCall %void %quantizeToF16_cba294
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %27
+         %41 = OpLabel
+         %42 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %42
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %27
+         %44 = OpLabel
+         %45 = OpFunctionCall %void %quantizeToF16_cba294
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %27
+         %47 = OpLabel
+         %48 = OpFunctionCall %void %quantizeToF16_cba294
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/quantizeToF16/cba294.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/literal/quantizeToF16/e8fd14.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/quantizeToF16/e8fd14.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/quantizeToF16/e8fd14.wgsl.expected.glsl

@@ -0,0 +1,70 @@
+#version 310 es
+
+vec3 tint_quantizeToF16(vec3 param_0) {
+  return vec3(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zz)).x);
+}
+
+
+void quantizeToF16_e8fd14() {
+  vec3 res = tint_quantizeToF16(vec3(1.0f));
+}
+
+vec4 vertex_main() {
+  quantizeToF16_e8fd14();
+  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;
+
+vec3 tint_quantizeToF16(vec3 param_0) {
+  return vec3(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zz)).x);
+}
+
+
+void quantizeToF16_e8fd14() {
+  vec3 res = tint_quantizeToF16(vec3(1.0f));
+}
+
+void fragment_main() {
+  quantizeToF16_e8fd14();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+vec3 tint_quantizeToF16(vec3 param_0) {
+  return vec3(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zz)).x);
+}
+
+
+void quantizeToF16_e8fd14() {
+  vec3 res = tint_quantizeToF16(vec3(1.0f));
+}
+
+void compute_main() {
+  quantizeToF16_e8fd14();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/quantizeToF16/e8fd14.wgsl.expected.msl

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

test/tint/builtins/gen/literal/quantizeToF16/e8fd14.wgsl.expected.spvasm

@@ -0,0 +1,86 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; 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 %tint_quantizeToF16 "tint_quantizeToF16"
+               OpName %v "v"
+               OpName %quantizeToF16_e8fd14 "quantizeToF16_e8fd14"
+               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
+    %v3float = OpTypeVector %float 3
+          %9 = OpTypeFunction %v3float %v3float
+       %uint = OpTypeInt 32 0
+         %16 = OpConstantNull %uint
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+       %void = OpTypeVoid
+         %25 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+         %31 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %34 = OpConstantNull %v3float
+         %35 = OpTypeFunction %v4float
+%tint_quantizeToF16 = OpFunction %v3float None %9
+          %v = OpFunctionParameter %v3float
+         %13 = OpLabel
+         %17 = OpCompositeExtract %float %v 0
+         %14 = OpQuantizeToF16 %float %17
+         %20 = OpCompositeExtract %float %v 1
+         %18 = OpQuantizeToF16 %float %20
+         %23 = OpCompositeExtract %float %v 2
+         %21 = OpQuantizeToF16 %float %23
+         %24 = OpCompositeConstruct %v3float %14 %18 %21
+               OpReturnValue %24
+               OpFunctionEnd
+%quantizeToF16_e8fd14 = OpFunction %void None %25
+         %28 = OpLabel
+        %res = OpVariable %_ptr_Function_v3float Function %34
+         %29 = OpFunctionCall %v3float %tint_quantizeToF16 %31
+               OpStore %res %29
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %35
+         %37 = OpLabel
+         %38 = OpFunctionCall %void %quantizeToF16_e8fd14
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %25
+         %40 = OpLabel
+         %41 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %41
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %25
+         %43 = OpLabel
+         %44 = OpFunctionCall %void %quantizeToF16_e8fd14
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %25
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %quantizeToF16_e8fd14
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/quantizeToF16/e8fd14.wgsl.expected.wgsl

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

test/tint/builtins/gen/var/quantizeToF16/12e50e.wgsl

@@ -0,0 +1,44 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn quantizeToF16(f32) -> f32
+fn quantizeToF16_12e50e() {
+  var arg_0 = 1.f;
+  var res: f32 = quantizeToF16(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  quantizeToF16_12e50e();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  quantizeToF16_12e50e();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  quantizeToF16_12e50e();
+}

test/tint/builtins/gen/var/quantizeToF16/12e50e.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/var/quantizeToF16/12e50e.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/var/quantizeToF16/12e50e.wgsl.expected.glsl

@@ -0,0 +1,67 @@
+#version 310 es
+
+float tint_quantizeToF16(float param_0) {
+  return unpackHalf2x16(packHalf2x16(vec2(param_0))).x;
+}
+
+
+void quantizeToF16_12e50e() {
+  float arg_0 = 1.0f;
+  float res = tint_quantizeToF16(arg_0);
+}
+
+vec4 vertex_main() {
+  quantizeToF16_12e50e();
+  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;
+
+float tint_quantizeToF16(float param_0) {
+  return unpackHalf2x16(packHalf2x16(vec2(param_0))).x;
+}
+
+
+void quantizeToF16_12e50e() {
+  float arg_0 = 1.0f;
+  float res = tint_quantizeToF16(arg_0);
+}
+
+void fragment_main() {
+  quantizeToF16_12e50e();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+float tint_quantizeToF16(float param_0) {
+  return unpackHalf2x16(packHalf2x16(vec2(param_0))).x;
+}
+
+
+void quantizeToF16_12e50e() {
+  float arg_0 = 1.0f;
+  float res = tint_quantizeToF16(arg_0);
+}
+
+void compute_main() {
+  quantizeToF16_12e50e();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/var/quantizeToF16/12e50e.wgsl.expected.msl

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

test/tint/builtins/gen/var/quantizeToF16/12e50e.wgsl.expected.spvasm

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

test/tint/builtins/gen/var/quantizeToF16/12e50e.wgsl.expected.wgsl

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

test/tint/builtins/gen/var/quantizeToF16/2cddf3.wgsl

@@ -0,0 +1,44 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn quantizeToF16(vec<2, f32>) -> vec<2, f32>
+fn quantizeToF16_2cddf3() {
+  var arg_0 = vec2<f32>(1.f);
+  var res: vec2<f32> = quantizeToF16(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  quantizeToF16_2cddf3();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  quantizeToF16_2cddf3();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  quantizeToF16_2cddf3();
+}

test/tint/builtins/gen/var/quantizeToF16/2cddf3.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/var/quantizeToF16/2cddf3.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/var/quantizeToF16/2cddf3.wgsl.expected.glsl

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

test/tint/builtins/gen/var/quantizeToF16/2cddf3.wgsl.expected.msl

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

test/tint/builtins/gen/var/quantizeToF16/2cddf3.wgsl.expected.spvasm

@@ -0,0 +1,87 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 47
+; 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 %tint_quantizeToF16 "tint_quantizeToF16"
+               OpName %v "v"
+               OpName %quantizeToF16_2cddf3 "quantizeToF16_2cddf3"
+               OpName %arg_0 "arg_0"
+               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
+    %v2float = OpTypeVector %float 2
+          %9 = OpTypeFunction %v2float %v2float
+       %uint = OpTypeInt 32 0
+         %16 = OpConstantNull %uint
+     %uint_1 = OpConstant %uint 1
+       %void = OpTypeVoid
+         %22 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+         %27 = OpConstantComposite %v2float %float_1 %float_1
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+         %30 = OpConstantNull %v2float
+         %34 = OpTypeFunction %v4float
+%tint_quantizeToF16 = OpFunction %v2float None %9
+          %v = OpFunctionParameter %v2float
+         %13 = OpLabel
+         %17 = OpCompositeExtract %float %v 0
+         %14 = OpQuantizeToF16 %float %17
+         %20 = OpCompositeExtract %float %v 1
+         %18 = OpQuantizeToF16 %float %20
+         %21 = OpCompositeConstruct %v2float %14 %18
+               OpReturnValue %21
+               OpFunctionEnd
+%quantizeToF16_2cddf3 = OpFunction %void None %22
+         %25 = OpLabel
+      %arg_0 = OpVariable %_ptr_Function_v2float Function %30
+        %res = OpVariable %_ptr_Function_v2float Function %30
+               OpStore %arg_0 %27
+         %32 = OpLoad %v2float %arg_0
+         %31 = OpFunctionCall %v2float %tint_quantizeToF16 %32
+               OpStore %res %31
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %34
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %quantizeToF16_2cddf3
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %22
+         %39 = OpLabel
+         %40 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %40
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %22
+         %42 = OpLabel
+         %43 = OpFunctionCall %void %quantizeToF16_2cddf3
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %22
+         %45 = OpLabel
+         %46 = OpFunctionCall %void %quantizeToF16_2cddf3
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/var/quantizeToF16/2cddf3.wgsl.expected.wgsl

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

test/tint/builtins/gen/var/quantizeToF16/cba294.wgsl

@@ -0,0 +1,44 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn quantizeToF16(vec<4, f32>) -> vec<4, f32>
+fn quantizeToF16_cba294() {
+  var arg_0 = vec4<f32>(1.f);
+  var res: vec4<f32> = quantizeToF16(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  quantizeToF16_cba294();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  quantizeToF16_cba294();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  quantizeToF16_cba294();
+}

test/tint/builtins/gen/var/quantizeToF16/cba294.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/var/quantizeToF16/cba294.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/var/quantizeToF16/cba294.wgsl.expected.glsl

@@ -0,0 +1,73 @@
+#version 310 es
+
+vec4 tint_quantizeToF16(vec4 param_0) {
+  return vec4(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zw)));
+}
+
+
+void quantizeToF16_cba294() {
+  vec4 arg_0 = vec4(1.0f);
+  vec4 res = tint_quantizeToF16(arg_0);
+}
+
+vec4 vertex_main() {
+  quantizeToF16_cba294();
+  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;
+
+vec4 tint_quantizeToF16(vec4 param_0) {
+  return vec4(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zw)));
+}
+
+
+void quantizeToF16_cba294() {
+  vec4 arg_0 = vec4(1.0f);
+  vec4 res = tint_quantizeToF16(arg_0);
+}
+
+void fragment_main() {
+  quantizeToF16_cba294();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+vec4 tint_quantizeToF16(vec4 param_0) {
+  return vec4(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zw)));
+}
+
+
+void quantizeToF16_cba294() {
+  vec4 arg_0 = vec4(1.0f);
+  vec4 res = tint_quantizeToF16(arg_0);
+}
+
+void compute_main() {
+  quantizeToF16_cba294();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/var/quantizeToF16/cba294.wgsl.expected.msl

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

test/tint/builtins/gen/var/quantizeToF16/cba294.wgsl.expected.spvasm

@@ -0,0 +1,91 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 51
+; 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 %tint_quantizeToF16 "tint_quantizeToF16"
+               OpName %v "v"
+               OpName %quantizeToF16_cba294 "quantizeToF16_cba294"
+               OpName %arg_0 "arg_0"
+               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
+          %9 = OpTypeFunction %v4float %v4float
+       %uint = OpTypeInt 32 0
+         %15 = OpConstantNull %uint
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+       %void = OpTypeVoid
+         %27 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+         %32 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+         %38 = OpTypeFunction %v4float
+%tint_quantizeToF16 = OpFunction %v4float None %9
+          %v = OpFunctionParameter %v4float
+         %12 = OpLabel
+         %16 = OpCompositeExtract %float %v 0
+         %13 = OpQuantizeToF16 %float %16
+         %19 = OpCompositeExtract %float %v 1
+         %17 = OpQuantizeToF16 %float %19
+         %22 = OpCompositeExtract %float %v 2
+         %20 = OpQuantizeToF16 %float %22
+         %25 = OpCompositeExtract %float %v 3
+         %23 = OpQuantizeToF16 %float %25
+         %26 = OpCompositeConstruct %v4float %13 %17 %20 %23
+               OpReturnValue %26
+               OpFunctionEnd
+%quantizeToF16_cba294 = OpFunction %void None %27
+         %30 = OpLabel
+      %arg_0 = OpVariable %_ptr_Function_v4float Function %5
+        %res = OpVariable %_ptr_Function_v4float Function %5
+               OpStore %arg_0 %32
+         %36 = OpLoad %v4float %arg_0
+         %35 = OpFunctionCall %v4float %tint_quantizeToF16 %36
+               OpStore %res %35
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %38
+         %40 = OpLabel
+         %41 = OpFunctionCall %void %quantizeToF16_cba294
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %27
+         %43 = OpLabel
+         %44 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %44
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %27
+         %46 = OpLabel
+         %47 = OpFunctionCall %void %quantizeToF16_cba294
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %27
+         %49 = OpLabel
+         %50 = OpFunctionCall %void %quantizeToF16_cba294
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/var/quantizeToF16/cba294.wgsl.expected.wgsl

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

test/tint/builtins/gen/var/quantizeToF16/e8fd14.wgsl

@@ -0,0 +1,44 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn quantizeToF16(vec<3, f32>) -> vec<3, f32>
+fn quantizeToF16_e8fd14() {
+  var arg_0 = vec3<f32>(1.f);
+  var res: vec3<f32> = quantizeToF16(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  quantizeToF16_e8fd14();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  quantizeToF16_e8fd14();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  quantizeToF16_e8fd14();
+}

test/tint/builtins/gen/var/quantizeToF16/e8fd14.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/var/quantizeToF16/e8fd14.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/var/quantizeToF16/e8fd14.wgsl.expected.glsl

@@ -0,0 +1,73 @@
+#version 310 es
+
+vec3 tint_quantizeToF16(vec3 param_0) {
+  return vec3(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zz)).x);
+}
+
+
+void quantizeToF16_e8fd14() {
+  vec3 arg_0 = vec3(1.0f);
+  vec3 res = tint_quantizeToF16(arg_0);
+}
+
+vec4 vertex_main() {
+  quantizeToF16_e8fd14();
+  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;
+
+vec3 tint_quantizeToF16(vec3 param_0) {
+  return vec3(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zz)).x);
+}
+
+
+void quantizeToF16_e8fd14() {
+  vec3 arg_0 = vec3(1.0f);
+  vec3 res = tint_quantizeToF16(arg_0);
+}
+
+void fragment_main() {
+  quantizeToF16_e8fd14();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+vec3 tint_quantizeToF16(vec3 param_0) {
+  return vec3(
+    unpackHalf2x16(packHalf2x16(param_0.xy)),
+    unpackHalf2x16(packHalf2x16(param_0.zz)).x);
+}
+
+
+void quantizeToF16_e8fd14() {
+  vec3 arg_0 = vec3(1.0f);
+  vec3 res = tint_quantizeToF16(arg_0);
+}
+
+void compute_main() {
+  quantizeToF16_e8fd14();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/var/quantizeToF16/e8fd14.wgsl.expected.msl

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

test/tint/builtins/gen/var/quantizeToF16/e8fd14.wgsl.expected.spvasm

@@ -0,0 +1,90 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 50
+; 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 %tint_quantizeToF16 "tint_quantizeToF16"
+               OpName %v "v"
+               OpName %quantizeToF16_e8fd14 "quantizeToF16_e8fd14"
+               OpName %arg_0 "arg_0"
+               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
+    %v3float = OpTypeVector %float 3
+          %9 = OpTypeFunction %v3float %v3float
+       %uint = OpTypeInt 32 0
+         %16 = OpConstantNull %uint
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+       %void = OpTypeVoid
+         %25 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+         %30 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %33 = OpConstantNull %v3float
+         %37 = OpTypeFunction %v4float
+%tint_quantizeToF16 = OpFunction %v3float None %9
+          %v = OpFunctionParameter %v3float
+         %13 = OpLabel
+         %17 = OpCompositeExtract %float %v 0
+         %14 = OpQuantizeToF16 %float %17
+         %20 = OpCompositeExtract %float %v 1
+         %18 = OpQuantizeToF16 %float %20
+         %23 = OpCompositeExtract %float %v 2
+         %21 = OpQuantizeToF16 %float %23
+         %24 = OpCompositeConstruct %v3float %14 %18 %21
+               OpReturnValue %24
+               OpFunctionEnd
+%quantizeToF16_e8fd14 = OpFunction %void None %25
+         %28 = OpLabel
+      %arg_0 = OpVariable %_ptr_Function_v3float Function %33
+        %res = OpVariable %_ptr_Function_v3float Function %33
+               OpStore %arg_0 %30
+         %35 = OpLoad %v3float %arg_0
+         %34 = OpFunctionCall %v3float %tint_quantizeToF16 %35
+               OpStore %res %34
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %37
+         %39 = OpLabel
+         %40 = OpFunctionCall %void %quantizeToF16_e8fd14
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %25
+         %42 = OpLabel
+         %43 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %43
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %25
+         %45 = OpLabel
+         %46 = OpFunctionCall %void %quantizeToF16_e8fd14
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %25
+         %48 = OpLabel
+         %49 = OpFunctionCall %void %quantizeToF16_e8fd14
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/var/quantizeToF16/e8fd14.wgsl.expected.wgsl

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