tint: const eval of cross

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

src/tint/intrinsics.def

@@ -439,7 +439,7 @@ fn cosh<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
 @const fn countOneBits<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
 @const fn countTrailingZeros<T: iu32>(T) -> T
 @const fn countTrailingZeros<N: num, T: iu32>(vec<N, T>) -> vec<N, T>
-fn cross<T: f32_f16>(vec3<T>, vec3<T>) -> vec3<T>
+@const fn cross<T: fa_f32_f16>(vec3<T>, vec3<T>) -> vec3<T>
 fn degrees<T: f32_f16>(T) -> T
 fn degrees<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
 fn determinant<N: num, T: f32_f16>(mat<N, N, T>) -> T

src/tint/resolver/builtin_test.cc

@@ -748,7 +748,7 @@ TEST_F(ResolverBuiltinFloatTest, Cross_Error_NoArgs) {
     EXPECT_EQ(r()->error(), R"(error: no matching call to cross()
 
 1 candidate function:
-  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is f32 or f16
+  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is abstract-float, f32 or f16
 )");
 }
 
@@ -761,7 +761,7 @@ TEST_F(ResolverBuiltinFloatTest, Cross_Error_Scalar) {
     EXPECT_EQ(r()->error(), R"(error: no matching call to cross(f32, f32)
 
 1 candidate function:
-  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is f32 or f16
+  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is abstract-float, f32 or f16
 )");
 }
 
@@ -775,7 +775,7 @@ TEST_F(ResolverBuiltinFloatTest, Cross_Error_Vec3Int) {
               R"(error: no matching call to cross(vec3<i32>, vec3<i32>)
 
 1 candidate function:
-  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is f32 or f16
+  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is abstract-float, f32 or f16
 )");
 }
 
@@ -790,7 +790,7 @@ TEST_F(ResolverBuiltinFloatTest, Cross_Error_Vec4) {
               R"(error: no matching call to cross(vec4<f32>, vec4<f32>)
 
 1 candidate function:
-  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is f32 or f16
+  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is abstract-float, f32 or f16
 )");
 }
 
@@ -806,7 +806,7 @@ TEST_F(ResolverBuiltinFloatTest, Cross_Error_TooManyParams) {
               R"(error: no matching call to cross(vec3<f32>, vec3<f32>, vec3<f32>)
 
 1 candidate function:
-  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is f32 or f16
+  cross(vec3<T>, vec3<T>) -> vec3<T>  where: T is abstract-float, f32 or f16
 )");
 }
 

src/tint/resolver/const_eval.cc

@@ -680,6 +680,33 @@ utils::Result<NumberT> ConstEval::Add(NumberT a, NumberT b) {
     return result;
 }
 
+template <typename NumberT>
+utils::Result<NumberT> ConstEval::Sub(NumberT a, NumberT b) {
+    NumberT result;
+    if constexpr (IsAbstract<NumberT>) {
+        // Check for over/underflow for abstract values
+        if (auto r = CheckedSub(a, b)) {
+            result = r->value;
+        } else {
+            AddError(OverflowErrorMessage(a, "-", b), *current_source);
+            return utils::Failure;
+        }
+    } else {
+        using T = UnwrapNumber<NumberT>;
+        auto sub_values = [](T lhs, T rhs) {
+            if constexpr (std::is_integral_v<T> && std::is_signed_v<T>) {
+                // Ensure no UB for signed overflow
+                using UT = std::make_unsigned_t<T>;
+                return static_cast<T>(static_cast<UT>(lhs) - static_cast<UT>(rhs));
+            } else {
+                return lhs - rhs;
+            }
+        };
+        result = sub_values(a.value, b.value);
+    }
+    return result;
+}
+
 template <typename NumberT>
 utils::Result<NumberT> ConstEval::Mul(NumberT a, NumberT b) {
     using T = UnwrapNumber<NumberT>;
@@ -794,6 +821,23 @@ utils::Result<NumberT> ConstEval::Dot4(NumberT a1,
     return r;
 }
 
+template <typename NumberT>
+utils::Result<NumberT> ConstEval::Det2(NumberT a1, NumberT a2, NumberT b1, NumberT b2) {
+    auto r1 = Mul(a1, b2);
+    if (!r1) {
+        return utils::Failure;
+    }
+    auto r2 = Mul(b1, a2);
+    if (!r2) {
+        return utils::Failure;
+    }
+    auto r = Sub(r1.Get(), r2.Get());
+    if (!r) {
+        return utils::Failure;
+    }
+    return r;
+}
+
 template <typename NumberT>
 utils::Result<NumberT> ConstEval::Clamp(NumberT e, NumberT low, NumberT high) {
     return NumberT{std::min(std::max(e, low), high)};
@@ -817,6 +861,15 @@ auto ConstEval::AddFunc(const sem::Type* elem_ty) {
     };
 }
 
+auto ConstEval::SubFunc(const sem::Type* elem_ty) {
+    return [=](auto a1, auto a2) -> ImplResult {
+        if (auto r = Sub(a1, a2)) {
+            return CreateElement(builder, elem_ty, r.Get());
+        }
+        return utils::Failure;
+    };
+}
+
 auto ConstEval::MulFunc(const sem::Type* elem_ty) {
     return [=](auto a1, auto a2) -> ImplResult {
         if (auto r = Mul(a1, a2)) {
@@ -854,6 +907,15 @@ auto ConstEval::Dot4Func(const sem::Type* elem_ty) {
         };
 }
 
+auto ConstEval::Det2Func(const sem::Type* elem_ty) {
+    return [=](auto a, auto b, auto c, auto d) -> ImplResult {
+        if (auto r = Det2(a, b, c, d)) {
+            return CreateElement(builder, elem_ty, r.Get());
+        }
+        return utils::Failure;
+    };
+}
+
 ConstEval::Result ConstEval::Literal(const sem::Type* ty, const ast::LiteralExpression* literal) {
     return Switch(
         literal,
@@ -1115,34 +1177,9 @@ ConstEval::Result ConstEval::OpPlus(const sem::Type* ty,
 ConstEval::Result ConstEval::OpMinus(const sem::Type* ty,
                                      utils::VectorRef<const sem::Constant*> args,
                                      const Source& source) {
+    TINT_SCOPED_ASSIGNMENT(current_source, &source);
     auto transform = [&](const sem::Constant* c0, const sem::Constant* c1) {
-        auto create = [&](auto i, auto j) -> ImplResult {
+        return Dispatch_fia_fiu32_f16(SubFunc(c0->Type()), c0, c1);
-            using NumberT = decltype(i);
-            NumberT result;
-            if constexpr (IsAbstract<NumberT>) {
-                // Check for over/underflow for abstract values
-                if (auto r = CheckedSub(i, j)) {
-                    result = r->value;
-                } else {
-                    AddError(OverflowErrorMessage(i, "-", j), source);
-                    return utils::Failure;
-                }
-            } else {
-                using T = UnwrapNumber<NumberT>;
-                auto subtract_values = [](T lhs, T rhs) {
-                    if constexpr (std::is_integral_v<T> && std::is_signed_v<T>) {
-                        // Ensure no UB for signed underflow
-                        using UT = std::make_unsigned_t<T>;
-                        return static_cast<T>(static_cast<UT>(lhs) - static_cast<UT>(rhs));
-                    } else {
-                        return lhs - rhs;
-                    }
-                };
-                result = subtract_values(i.value, j.value);
-            }
-            return CreateElement(builder, c0->Type(), result);
-        };
-        return Dispatch_fia_fiu32_f16(create, c0, c1);
     };
 
     return TransformBinaryElements(builder, ty, transform, args[0], args[1]);
@@ -1627,7 +1664,8 @@ ConstEval::Result ConstEval::acos(const sem::Type* ty,
         auto create = [&](auto i) -> ImplResult {
             using NumberT = decltype(i);
             if (i < NumberT(-1.0) || i > NumberT(1.0)) {
-                AddError("acos must be called with a value in the range [-1 .. 1] (inclusive)", source);
+                AddError("acos must be called with a value in the range [-1 .. 1] (inclusive)",
+                         source);
                 return utils::Failure;
             }
             return CreateElement(builder, c0->Type(), NumberT(std::acos(i.value)));
@@ -1656,7 +1694,8 @@ ConstEval::Result ConstEval::asin(const sem::Type* ty,
         auto create = [&](auto i) -> ImplResult {
             using NumberT = decltype(i);
             if (i < NumberT(-1.0) || i > NumberT(1.0)) {
-                AddError("asin must be called with a value in the range [-1 .. 1] (inclusive)", source);
+                AddError("asin must be called with a value in the range [-1 .. 1] (inclusive)",
+                         source);
                 return utils::Failure;
             }
             return CreateElement(builder, c0->Type(), NumberT(std::asin(i.value)));
@@ -1702,7 +1741,8 @@ ConstEval::Result ConstEval::atanh(const sem::Type* ty,
         auto create = [&](auto i) -> ImplResult {
             using NumberT = decltype(i);
             if (i <= NumberT(-1.0) || i >= NumberT(1.0)) {
-                AddError("atanh must be called with a value in the range (-1 .. 1) (exclusive)", source);
+                AddError("atanh must be called with a value in the range (-1 .. 1) (exclusive)",
+                         source);
                 return utils::Failure;
             }
             return CreateElement(builder, c0->Type(), NumberT(std::atanh(i.value)));
@@ -1801,6 +1841,52 @@ ConstEval::Result ConstEval::countTrailingZeros(const sem::Type* ty,
     return TransformElements(builder, ty, transform, args[0]);
 }
 
+ConstEval::Result ConstEval::cross(const sem::Type* ty,
+                                   utils::VectorRef<const sem::Constant*> args,
+                                   const Source& source) {
+    TINT_SCOPED_ASSIGNMENT(current_source, &source);
+    auto* u = args[0];
+    auto* v = args[1];
+    auto* elem_ty = u->Type()->As<sem::Vector>()->type();
+
+    // cross product of a v3 is the determinant of the 3x3 matrix:
+    //
+    // |i   j   k |
+    // |u0  u1  u2|
+    // |v0  v1  v2|
+    //
+    // |u1 u2|i  - |u0 u2|j + |u0 u1|k
+    // |v1 v2|     |v0 v2|    |v0 v1|
+    //
+    // |u1 u2|i  + |v0 v2|j + |u0 u1|k
+    // |v1 v2|     |u0 u2|    |v0 v1|
+
+    auto* u0 = u->Index(0);
+    auto* u1 = u->Index(1);
+    auto* u2 = u->Index(2);
+    auto* v0 = v->Index(0);
+    auto* v1 = v->Index(1);
+    auto* v2 = v->Index(2);
+
+    // auto x = Dispatch_fa_f32_f16(ab_minus_cd_func(elem_ty), u->Index(1), v->Index(2),
+    //                              v->Index(1), u->Index(2));
+    auto x = Dispatch_fa_f32_f16(Det2Func(elem_ty), u1, u2, v1, v2);
+    if (!x) {
+        return utils::Failure;
+    }
+    auto y = Dispatch_fa_f32_f16(Det2Func(elem_ty), v0, v2, u0, u2);
+    if (!y) {
+        return utils::Failure;
+    }
+    auto z = Dispatch_fa_f32_f16(Det2Func(elem_ty), u0, u1, v0, v1);
+    if (!z) {
+        return utils::Failure;
+    }
+
+    return CreateComposite(builder, ty,
+                           utils::Vector<const sem::Constant*, 3>{x.Get(), y.Get(), z.Get()});
+}
+
 ConstEval::Result ConstEval::extractBits(const sem::Type* ty,
                                          utils::VectorRef<const sem::Constant*> args,
                                          const Source& source) {

src/tint/resolver/const_eval.h

@@ -503,6 +503,15 @@ class ConstEval {
                               utils::VectorRef<const sem::Constant*> args,
                               const Source& source);
 
+    /// cross builtin
+    /// @param ty the expression type
+    /// @param args the input arguments
+    /// @param source the source location of the conversion
+    /// @return the result value, or null if the value cannot be calculated
+    Result cross(const sem::Type* ty,
+                 utils::VectorRef<const sem::Constant*> args,
+                 const Source& source);
+
     /// extractBits builtin
     /// @param ty the expression type
     /// @param args the input arguments
@@ -697,6 +706,13 @@ class ConstEval {
     template <typename NumberT>
     utils::Result<NumberT> Add(NumberT a, NumberT b);
 
+    /// Subtracts two Number<T>s
+    /// @param a the lhs number
+    /// @param b the rhs number
+    /// @returns the result number on success, or logs an error and returns Failure
+    template <typename NumberT>
+    utils::Result<NumberT> Sub(NumberT a, NumberT b);
+
     /// Multiplies two Number<T>s
     /// @param a the lhs number
     /// @param b the rhs number
@@ -749,6 +765,14 @@ class ConstEval {
                                 NumberT b3,
                                 NumberT b4);
 
+    /// Returns the determinant of the 2x2 matrix [(a1, a2), (b1, b2)]
+    /// @param a1 component 1 of the first column vector
+    /// @param a2 component 2 of the first column vector
+    /// @param b1 component 1 of the second column vector
+    /// @param b2 component 2 of the second column vector
+    template <typename NumberT>
+    utils::Result<NumberT> Det2(NumberT a1, NumberT a2, NumberT b1, NumberT b2);
+
     /// Clamps e between low and high
     /// @param e the number to clamp
     /// @param low the lower bound
@@ -763,6 +787,12 @@ class ConstEval {
     /// @returns the callable function
     auto AddFunc(const sem::Type* elem_ty);
 
+    /// Returns a callable that calls Sub, and creates a Constant with its result of type `elem_ty`
+    /// if successful, or returns Failure otherwise.
+    /// @param elem_ty the element type of the Constant to create on success
+    /// @returns the callable function
+    auto SubFunc(const sem::Type* elem_ty);
+
     /// Returns a callable that calls Mul, and creates a Constant with its result of type `elem_ty`
     /// if successful, or returns Failure otherwise.
     /// @param elem_ty the element type of the Constant to create on success
@@ -787,6 +817,12 @@ class ConstEval {
     /// @returns the callable function
     auto Dot4Func(const sem::Type* elem_ty);
 
+    /// Returns a callable that calls Det2, and creates a Constant with its result of type `elem_ty`
+    /// if successful, or returns Failure otherwise.
+    /// @param elem_ty the element type of the Constant to create on success
+    /// @returns the callable function
+    auto Det2Func(const sem::Type* elem_ty);
+
     /// Returns a callable that calls Clamp, and creates a Constant with its result of type
     /// `elem_ty` if successful, or returns Failure otherwise.
     /// @param elem_ty the element type of the Constant to create on success

src/tint/resolver/const_eval_builtin_test.cc

@@ -17,6 +17,7 @@
 #include "src/tint/utils/result.h"
 
 using namespace tint::number_suffixes;  // NOLINT
+using ::testing::HasSubstr;
 
 namespace tint::resolver {
 namespace {
@@ -28,8 +29,8 @@ struct Case {
     Case(utils::VectorRef<Types> in_args, Types expected_value)
         : args(std::move(in_args)), expected(Success{std::move(expected_value), false, false}) {}
 
-    Case(utils::VectorRef<Types> in_args, const char* expected_err)
+    Case(utils::VectorRef<Types> in_args, std::string expected_err)
-        : args(std::move(in_args)), expected(Failure{expected_err}) {}
+        : args(std::move(in_args)), expected(Failure{std::move(expected_err)}) {}
 
     /// Expected value may be positive or negative
     Case& PosOrNeg() {
@@ -53,7 +54,7 @@ struct Case {
         bool float_compare = false;
     };
     struct Failure {
-        const char* error = nullptr;
+        std::string error = nullptr;
     };
 
     utils::Vector<Types, 8> args;
@@ -94,17 +95,27 @@ static Case C(std::initializer_list<ScalarTypes> sargs, ScalarTypes sresult) {
 }
 
 /// Creates a Case with Values for args and expected error
-static Case E(std::initializer_list<Types> args, const char* err) {
+static Case E(std::initializer_list<Types> args, std::string err) {
-    return Case{utils::Vector<Types, 8>{args}, err};
+    return Case{utils::Vector<Types, 8>{args}, std::move(err)};
 }
 
 /// Convenience overload that creates an expected-error Case with just scalars
-static Case E(std::initializer_list<ScalarTypes> sargs, const char* err) {
+static Case E(std::initializer_list<ScalarTypes> sargs, std::string err) {
     utils::Vector<Types, 8> args;
     for (auto& sa : sargs) {
         std::visit([&](auto&& v) { return args.Push(Val(v)); }, sa);
     }
-    return Case{std::move(args), err};
+    return Case{std::move(args), std::move(err)};
+}
+
+/// Returns the overflow error message for binary ops
+template <typename NumberT>
+std::string OverflowErrorMessage(NumberT lhs, const char* op, NumberT rhs) {
+    std::stringstream ss;
+    ss << std::setprecision(20);
+    ss << "'" << lhs.value << " " << op << " " << rhs.value << "' cannot be represented as '"
+       << FriendlyName<NumberT>() << "'";
+    return ss.str();
 }
 
 using ResolverConstEvalBuiltinTest = ResolverTestWithParam<std::tuple<sem::BuiltinType, Case>>;
@@ -776,6 +787,130 @@ INSTANTIATE_TEST_SUITE_P(  //
                      testing::ValuesIn(Concat(CountOneBitsCases<i32>(),  //
                                               CountOneBitsCases<u32>()))));
 
+template <typename T, bool finite_only>
+std::vector<Case> CrossCases() {
+    constexpr auto vec_x = [](T v) { return Vec(T(v), T(0), T(0)); };
+    constexpr auto vec_y = [](T v) { return Vec(T(0), T(v), T(0)); };
+    constexpr auto vec_z = [](T v) { return Vec(T(0), T(0), T(v)); };
+
+    const auto zero = Vec(T(0), T(0), T(0));
+    const auto unit_x = vec_x(T(1));
+    const auto unit_y = vec_y(T(1));
+    const auto unit_z = vec_z(T(1));
+    const auto neg_unit_x = vec_x(-T(1));
+    const auto neg_unit_y = vec_y(-T(1));
+    const auto neg_unit_z = vec_z(-T(1));
+    const auto highest_x = vec_x(T::Highest());
+    const auto highest_y = vec_y(T::Highest());
+    const auto highest_z = vec_z(T::Highest());
+    const auto smallest_x = vec_x(T::Smallest());
+    const auto smallest_y = vec_y(T::Smallest());
+    const auto smallest_z = vec_z(T::Smallest());
+    const auto lowest_x = vec_x(T::Lowest());
+    const auto lowest_y = vec_y(T::Lowest());
+    const auto lowest_z = vec_z(T::Lowest());
+    const auto inf_x = vec_x(T::Inf());
+    const auto inf_y = vec_y(T::Inf());
+    const auto inf_z = vec_z(T::Inf());
+    const auto neg_inf_x = vec_x(-T::Inf());
+    const auto neg_inf_y = vec_y(-T::Inf());
+    const auto neg_inf_z = vec_z(-T::Inf());
+
+    std::vector<Case> r = {
+        C({zero, zero}, zero),
+
+        C({unit_x, unit_x}, zero),
+        C({unit_y, unit_y}, zero),
+        C({unit_z, unit_z}, zero),
+
+        C({smallest_x, smallest_x}, zero),
+        C({smallest_y, smallest_y}, zero),
+        C({smallest_z, smallest_z}, zero),
+
+        C({lowest_x, lowest_x}, zero),
+        C({lowest_y, lowest_y}, zero),
+        C({lowest_z, lowest_z}, zero),
+
+        C({highest_x, highest_x}, zero),
+        C({highest_y, highest_y}, zero),
+        C({highest_z, highest_z}, zero),
+
+        C({smallest_x, highest_x}, zero),
+        C({smallest_y, highest_y}, zero),
+        C({smallest_z, highest_z}, zero),
+
+        C({unit_x, neg_unit_x}, zero).PosOrNeg(),
+        C({unit_y, neg_unit_y}, zero).PosOrNeg(),
+        C({unit_z, neg_unit_z}, zero).PosOrNeg(),
+
+        C({unit_x, unit_y}, unit_z),
+        C({unit_y, unit_x}, neg_unit_z),
+
+        C({unit_z, unit_x}, unit_y),
+        C({unit_x, unit_z}, neg_unit_y),
+
+        C({unit_y, unit_z}, unit_x),
+        C({unit_z, unit_y}, neg_unit_x),
+
+        C({vec_x(T(1)), vec_y(T(2))}, vec_z(T(2))),
+        C({vec_y(T(1)), vec_x(T(2))}, vec_z(-T(2))),
+        C({vec_x(T(2)), vec_y(T(3))}, vec_z(T(6))),
+        C({vec_y(T(2)), vec_x(T(3))}, vec_z(-T(6))),
+
+        C({Vec(T(1), T(2), T(3)), Vec(T(1), T(5), T(7))}, Vec(T(-1), T(-4), T(3))),
+        C({Vec(T(33), T(44), T(55)), Vec(T(13), T(42), T(39))}, Vec(T(-594), T(-572), T(814))),
+        C({Vec(T(3.5), T(4), T(5.5)), Vec(T(1), T(4.5), T(3.5))},
+          Vec(T(-10.75), T(-6.75), T(11.75))),
+    };
+
+    ConcatIntoIf<!finite_only>(  //
+        r, std::vector<Case>{
+               C({highest_x, highest_y}, inf_z).PosOrNeg(),  //
+               C({highest_y, highest_x}, inf_z).PosOrNeg(),  //
+               C({highest_z, highest_x}, inf_y).PosOrNeg(),  //
+               C({highest_x, highest_z}, inf_y).PosOrNeg(),  //
+               C({highest_y, highest_z}, inf_x).PosOrNeg(),  //
+               C({highest_z, highest_y}, inf_x).PosOrNeg(),  //
+               C({lowest_x, lowest_y}, inf_z).PosOrNeg(),    //
+               C({lowest_y, lowest_x}, inf_z).PosOrNeg(),    //
+               C({lowest_z, lowest_x}, inf_y).PosOrNeg(),    //
+               C({lowest_x, lowest_z}, inf_y).PosOrNeg(),    //
+               C({lowest_y, lowest_z}, inf_x).PosOrNeg(),    //
+               C({lowest_z, lowest_y}, inf_x).PosOrNeg(),
+           });
+
+    std::string pos_error_msg =
+        "12:34 error: " + OverflowErrorMessage(T::Highest(), "*", T::Highest());
+    std::string neg_error_msg =
+        "12:34 error: " + OverflowErrorMessage(T::Lowest(), "*", T::Lowest());
+
+    ConcatIntoIf<finite_only>(  //
+        r, std::vector<Case>{
+               E({highest_x, highest_y}, pos_error_msg),
+               E({highest_y, highest_x}, pos_error_msg),
+               E({highest_z, highest_x}, pos_error_msg),
+               E({highest_x, highest_z}, pos_error_msg),
+               E({highest_y, highest_z}, pos_error_msg),
+               E({highest_z, highest_y}, pos_error_msg),
+               E({lowest_x, lowest_y}, neg_error_msg),
+               E({lowest_y, lowest_x}, neg_error_msg),
+               E({lowest_z, lowest_x}, neg_error_msg),
+               E({lowest_x, lowest_z}, neg_error_msg),
+               E({lowest_y, lowest_z}, neg_error_msg),
+               E({lowest_z, lowest_y}, neg_error_msg),
+           });
+
+    return r;
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Cross,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kCross),
+                     testing::ValuesIn(Concat(CrossCases<AFloat, true>(),  //
+                                              CrossCases<f32, false>(),
+                                              CrossCases<f32, false>(),  //
+                                              CrossCases<f16, false>()))));
+
 template <typename T>
 std::vector<Case> FirstLeadingBitCases() {
     using B = BitValues<T>;

test/tint/builtins/gen/literal/cross/041cb0.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/cross/041cb0.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/cross/041cb0.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/cross/041cb0.wgsl.expected.msl

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

test/tint/builtins/gen/literal/cross/041cb0.wgsl.expected.spvasm

@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 31
 ; Schema: 0
                OpCapability Shader
-         %15 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -32,37 +31,35 @@
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
     %v3float = OpTypeVector %float 3
-    %float_1 = OpConstant %float 1
+         %14 = OpConstantNull %v3float
-         %17 = OpConstantComposite %v3float %float_1 %float_1 %float_1
 %_ptr_Function_v3float = OpTypePointer Function %v3float
-         %20 = OpConstantNull %v3float
+         %17 = OpTypeFunction %v4float
-         %21 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
 %cross_041cb0 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_v3float Function %20
+        %res = OpVariable %_ptr_Function_v3float Function %14
-         %13 = OpExtInst %v3float %15 Cross %17 %17
+               OpStore %res %14
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
+%vertex_main_inner = OpFunction %v4float None %17
-         %23 = OpLabel
+         %19 = OpLabel
-         %24 = OpFunctionCall %void %cross_041cb0
+         %20 = OpFunctionCall %void %cross_041cb0
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %26 = OpLabel
+         %22 = OpLabel
-         %27 = OpFunctionCall %v4float %vertex_main_inner
+         %23 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %27
+               OpStore %value %23
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
+         %26 = OpLabel
+         %27 = OpFunctionCall %void %cross_041cb0
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
          %29 = OpLabel
          %30 = OpFunctionCall %void %cross_041cb0
                OpReturn
                OpFunctionEnd
-%compute_main = OpFunction %void None %9
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %cross_041cb0
-               OpReturn
-               OpFunctionEnd

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

test/tint/builtins/gen/literal/cross/1d7933.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/cross/1d7933.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/cross/1d7933.wgsl.expected.glsl

@@ -0,0 +1,49 @@
+#version 310 es
+
+void cross_1d7933() {
+  vec3 res = vec3(0.0f);
+}
+
+vec4 vertex_main() {
+  cross_1d7933();
+  return vec4(0.0f);
+}
+
+void main() {
+  gl_PointSize = 1.0;
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+void cross_1d7933() {
+  vec3 res = vec3(0.0f);
+}
+
+void fragment_main() {
+  cross_1d7933();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void cross_1d7933() {
+  vec3 res = vec3(0.0f);
+}
+
+void compute_main() {
+  cross_1d7933();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/literal/cross/1d7933.wgsl.expected.msl

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

test/tint/builtins/gen/literal/cross/1d7933.wgsl.expected.spvasm

@@ -0,0 +1,65 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %cross_1d7933 "cross_1d7933"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+    %v3float = OpTypeVector %float 3
+         %14 = OpConstantNull %v3float
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %17 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%cross_1d7933 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_v3float Function %14
+               OpStore %res %14
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %17
+         %19 = OpLabel
+         %20 = OpFunctionCall %void %cross_1d7933
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %22 = OpLabel
+         %23 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %23
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %26 = OpLabel
+         %27 = OpFunctionCall %void %cross_1d7933
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %cross_1d7933
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/cross/1d7933.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/cross/9857cb.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/cross/9857cb.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/cross/9857cb.wgsl.expected.msl

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

test/tint/builtins/gen/literal/cross/9857cb.wgsl.expected.spvasm

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

test/tint/builtins/gen/var/cross/1d7933.wgsl

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

test/tint/builtins/gen/var/cross/1d7933.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/var/cross/1d7933.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/var/cross/1d7933.wgsl.expected.glsl

@@ -0,0 +1,49 @@
+#version 310 es
+
+void cross_1d7933() {
+  vec3 res = vec3(0.0f);
+}
+
+vec4 vertex_main() {
+  cross_1d7933();
+  return vec4(0.0f);
+}
+
+void main() {
+  gl_PointSize = 1.0;
+  vec4 inner_result = vertex_main();
+  gl_Position = inner_result;
+  gl_Position.y = -(gl_Position.y);
+  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+  return;
+}
+#version 310 es
+precision mediump float;
+
+void cross_1d7933() {
+  vec3 res = vec3(0.0f);
+}
+
+void fragment_main() {
+  cross_1d7933();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+void cross_1d7933() {
+  vec3 res = vec3(0.0f);
+}
+
+void compute_main() {
+  cross_1d7933();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/var/cross/1d7933.wgsl.expected.msl

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

test/tint/builtins/gen/var/cross/1d7933.wgsl.expected.spvasm

@@ -0,0 +1,65 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %value "value"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %cross_1d7933 "cross_1d7933"
+               OpName %res "res"
+               OpName %vertex_main_inner "vertex_main_inner"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %value BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %5 = OpConstantNull %v4float
+      %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+          %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+    %v3float = OpTypeVector %float 3
+         %14 = OpConstantNull %v3float
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %17 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%cross_1d7933 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_v3float Function %14
+               OpStore %res %14
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %17
+         %19 = OpLabel
+         %20 = OpFunctionCall %void %cross_1d7933
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %22 = OpLabel
+         %23 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %23
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %26 = OpLabel
+         %27 = OpFunctionCall %void %cross_1d7933
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %cross_1d7933
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/var/cross/1d7933.wgsl.expected.wgsl

@@ -0,0 +1,21 @@
+fn cross_1d7933() {
+  const arg_0 = vec3(1.0);
+  const arg_1 = vec3(1.0);
+  var res = cross(arg_0, arg_1);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  cross_1d7933();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  cross_1d7933();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  cross_1d7933();
+}