tint: const eval of determinant builtin

Bug: tint:1581
Change-Id: Ifed8202ba2346eee435ee4e3d0e82ab614a86255
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/111281
Reviewed-by: Ben Clayton <bclayton@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Kokoro: Antonio Maiorano <amaiorano@google.com>

src/tint/intrinsics.def

@@ -442,7 +442,7 @@ fn arrayLength<T, A: access>(ptr<storage, array<T>, A>) -> u32
 @const fn cross<T: fa_f32_f16>(vec3<T>, vec3<T>) -> vec3<T>
 @const fn degrees<T: fa_f32_f16>(T) -> T
 @const fn degrees<N: num, T: fa_f32_f16>(vec<N, T>) -> vec<N, T>
-fn determinant<N: num, T: f32_f16>(mat<N, N, T>) -> T
+@const fn determinant<N: num, T: fa_f32_f16>(mat<N, N, T>) -> T
 fn distance<T: f32_f16>(T, T) -> T
 fn distance<N: num, T: f32_f16>(vec<N, T>, vec<N, T>) -> T
 @const fn dot<N: num, T: fia_fiu32_f16>(vec<N, T>, vec<N, T>) -> T

src/tint/resolver/builtin_test.cc

@@ -1932,7 +1932,7 @@ TEST_F(ResolverBuiltinTest, Determinant_NotSquare) {
     EXPECT_EQ(r()->error(), R"(error: no matching call to determinant(mat2x3<f32>)
 
 1 candidate function:
-  determinant(matNxN<T>) -> T  where: T is f32 or f16
+  determinant(matNxN<T>) -> T  where: T is abstract-float, f32 or f16
 )");
 }
 
@@ -1947,7 +1947,7 @@ TEST_F(ResolverBuiltinTest, Determinant_NotMatrix) {
     EXPECT_EQ(r()->error(), R"(error: no matching call to determinant(f32)
 
 1 candidate function:
-  determinant(matNxN<T>) -> T  where: T is f32 or f16
+  determinant(matNxN<T>) -> T  where: T is abstract-float, f32 or f16
 )");
 }
 

src/tint/resolver/const_eval.cc

@@ -893,15 +893,22 @@ utils::Result<NumberT> ConstEval::Dot4(const Source& source,
 
 template <typename NumberT>
 utils::Result<NumberT> ConstEval::Det2(const Source& source,
-                                       NumberT a1,
+                                       NumberT a,
-                                       NumberT a2,
+                                       NumberT b,
-                                       NumberT b1,
+                                       NumberT c,
-                                       NumberT b2) {
+                                       NumberT d) {
-    auto r1 = Mul(source, a1, b2);
+    // | a c |
+    // | b d |
+    //
+    // =
+    //
+    // a * d - c * b
+
+    auto r1 = Mul(source, a, d);
     if (!r1) {
         return utils::Failure;
     }
-    auto r2 = Mul(source, b1, a2);
+    auto r2 = Mul(source, c, b);
     if (!r2) {
         return utils::Failure;
     }
@@ -912,6 +919,129 @@ utils::Result<NumberT> ConstEval::Det2(const Source& source,
     return r;
 }
 
+template <typename NumberT>
+utils::Result<NumberT> ConstEval::Det3(const Source& source,
+                                       NumberT a,
+                                       NumberT b,
+                                       NumberT c,
+                                       NumberT d,
+                                       NumberT e,
+                                       NumberT f,
+                                       NumberT g,
+                                       NumberT h,
+                                       NumberT i) {
+    // | a d g |
+    // | b e h |
+    // | c f i |
+    //
+    // =
+    //
+    // a | e h | - d | b h | + g | b e |
+    //   | f i |     | c i |     | c f |
+
+    auto det1 = Det2(source, e, f, h, i);
+    if (!det1) {
+        return utils::Failure;
+    }
+    auto a_det1 = Mul(source, a, det1.Get());
+    if (!a_det1) {
+        return utils::Failure;
+    }
+    auto det2 = Det2(source, b, c, h, i);
+    if (!det2) {
+        return utils::Failure;
+    }
+    auto d_det2 = Mul(source, d, det2.Get());
+    if (!d_det2) {
+        return utils::Failure;
+    }
+    auto det3 = Det2(source, b, c, e, f);
+    if (!det3) {
+        return utils::Failure;
+    }
+    auto g_det3 = Mul(source, g, det3.Get());
+    if (!g_det3) {
+        return utils::Failure;
+    }
+    auto r = Sub(source, a_det1.Get(), d_det2.Get());
+    if (!r) {
+        return utils::Failure;
+    }
+    return Add(source, r.Get(), g_det3.Get());
+}
+
+template <typename NumberT>
+utils::Result<NumberT> ConstEval::Det4(const Source& source,
+                                       NumberT a,
+                                       NumberT b,
+                                       NumberT c,
+                                       NumberT d,
+                                       NumberT e,
+                                       NumberT f,
+                                       NumberT g,
+                                       NumberT h,
+                                       NumberT i,
+                                       NumberT j,
+                                       NumberT k,
+                                       NumberT l,
+                                       NumberT m,
+                                       NumberT n,
+                                       NumberT o,
+                                       NumberT p) {
+    // | a e i m |
+    // | b f j n |
+    // | c g k o |
+    // | d h l p |
+    //
+    // =
+    //
+    // a | f j n | - e | b j n | + i | b f n | - m | b f j |
+    //   | g k o |     | c k o |     | c g o |     | c g k |
+    //   | h l p |     | d l p |     | d h p |     | d h l |
+
+    auto det1 = Det3(source, f, g, h, j, k, l, n, o, p);
+    if (!det1) {
+        return utils::Failure;
+    }
+    auto a_det1 = Mul(source, a, det1.Get());
+    if (!a_det1) {
+        return utils::Failure;
+    }
+    auto det2 = Det3(source, b, c, d, j, k, l, n, o, p);
+    if (!det2) {
+        return utils::Failure;
+    }
+    auto e_det2 = Mul(source, e, det2.Get());
+    if (!e_det2) {
+        return utils::Failure;
+    }
+    auto det3 = Det3(source, b, c, d, f, g, h, n, o, p);
+    if (!det3) {
+        return utils::Failure;
+    }
+    auto i_det3 = Mul(source, i, det3.Get());
+    if (!i_det3) {
+        return utils::Failure;
+    }
+    auto det4 = Det3(source, b, c, d, f, g, h, j, k, l);
+    if (!det4) {
+        return utils::Failure;
+    }
+    auto m_det4 = Mul(source, m, det4.Get());
+    if (!m_det4) {
+        return utils::Failure;
+    }
+    auto r = Sub(source, a_det1.Get(), e_det2.Get());
+    if (!r) {
+        return utils::Failure;
+    }
+    r = Add(source, r.Get(), i_det3.Get());
+    if (!r) {
+        return utils::Failure;
+    }
+    return Sub(source, r.Get(), m_det4.Get());
+}
+
 template <typename NumberT>
 utils::Result<NumberT> ConstEval::Sqrt(const Source& source, NumberT v) {
     if (v < NumberT(0)) {
@@ -1044,6 +1174,26 @@ auto ConstEval::Det2Func(const Source& source, const sem::Type* elem_ty) {
     };
 }
 
+auto ConstEval::Det3Func(const Source& source, const sem::Type* elem_ty) {
+    return
+        [=](auto a, auto b, auto c, auto d, auto e, auto f, auto g, auto h, auto i) -> ImplResult {
+            if (auto r = Det3(source, a, b, c, d, e, f, g, h, i)) {
+                return CreateElement(builder, source, elem_ty, r.Get());
+            }
+            return utils::Failure;
+        };
+}
+
+auto ConstEval::Det4Func(const Source& source, const sem::Type* elem_ty) {
+    return [=](auto a, auto b, auto c, auto d, auto e, auto f, auto g, auto h, auto i, auto j,
+               auto k, auto l, auto m, auto n, auto o, auto p) -> ImplResult {
+        if (auto r = Det4(source, a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p)) {
+            return CreateElement(builder, source, elem_ty, r.Get());
+        }
+        return utils::Failure;
+    };
+}
+
 ConstEval::Result ConstEval::Literal(const sem::Type* ty, const ast::LiteralExpression* literal) {
     auto& source = literal->source;
     return Switch(
@@ -2036,6 +2186,41 @@ ConstEval::Result ConstEval::degrees(const sem::Type* ty,
     return TransformElements(builder, ty, transform, args[0]);
 }
 
+ConstEval::Result ConstEval::determinant(const sem::Type* ty,
+                                         utils::VectorRef<const sem::Constant*> args,
+                                         const Source& source) {
+    auto calculate = [&]() -> ImplResult {
+        auto* m = args[0];
+        auto* mat_ty = m->Type()->As<sem::Matrix>();
+        auto me = [&](size_t r, size_t c) { return m->Index(c)->Index(r); };
+        switch (mat_ty->rows()) {
+            case 2:
+                return Dispatch_fa_f32_f16(Det2Func(source, ty),  //
+                                           me(0, 0), me(1, 0),    //
+                                           me(0, 1), me(1, 1));
+
+            case 3:
+                return Dispatch_fa_f32_f16(Det3Func(source, ty),          //
+                                           me(0, 0), me(1, 0), me(2, 0),  //
+                                           me(0, 1), me(1, 1), me(2, 1),  //
+                                           me(0, 2), me(1, 2), me(2, 2));
+
+            case 4:
+                return Dispatch_fa_f32_f16(Det4Func(source, ty),                    //
+                                           me(0, 0), me(1, 0), me(2, 0), me(3, 0),  //
+                                           me(0, 1), me(1, 1), me(2, 1), me(3, 1),  //
+                                           me(0, 2), me(1, 2), me(2, 2), me(3, 2),  //
+                                           me(0, 3), me(1, 3), me(2, 3), me(3, 3));
+        }
+        TINT_ICE(Resolver, builder.Diagnostics()) << "Unexpected number of matrix rows";
+        return utils::Failure;
+    };
+    auto r = calculate();
+    if (!r) {
+        AddNote("when calculating determinant", source);
+    }
+    return r;
+}
 ConstEval::Result ConstEval::dot(const sem::Type*,
                                  utils::VectorRef<const sem::Constant*> args,
                                  const Source& source) {

src/tint/resolver/const_eval.h

@@ -539,6 +539,7 @@ class ConstEval {
                  utils::VectorRef<const sem::Constant*> args,
                  const Source& source);
 
+    /// degrees builtin
     /// @param ty the expression type
     /// @param args the input arguments
     /// @param source the source location of the conversion
@@ -547,6 +548,15 @@ class ConstEval {
                    utils::VectorRef<const sem::Constant*> args,
                    const Source& source);
 
+    /// determinant 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 determinant(const sem::Type* ty,
+                       utils::VectorRef<const sem::Constant*> args,
+                       const Source& source);
+
     /// dot builtin
     /// @param ty the expression type
     /// @param args the input arguments
@@ -1012,18 +1022,87 @@ class ConstEval {
                                 NumberT b3,
                                 NumberT b4);
 
-    /// Returns the determinant of the 2x2 matrix [(a1, a2), (b1, b2)]
+    /// Returns the determinant of the 2x2 matrix:
+    /// | a c |
+    /// | b d |
     /// @param source the source location
-    /// @param a1 component 1 of the first column vector
+    /// @param a component 1 of the first column vector
-    /// @param a2 component 2 of the first column vector
+    /// @param b component 2 of the first column vector
-    /// @param b1 component 1 of the second column vector
+    /// @param c component 1 of the second column vector
-    /// @param b2 component 2 of the second column vector
+    /// @param d component 2 of the second column vector
     template <typename NumberT>
-    utils::Result<NumberT> Det2(const Source& source,
+    utils::Result<NumberT> Det2(const Source& source,  //
-                                NumberT a1,
+                                NumberT a,
-                                NumberT a2,
+                                NumberT b,
-                                NumberT b1,
+                                NumberT c,
-                                NumberT b2);
+                                NumberT d);
+
+    /// Returns the determinant of the 3x3 matrix:
+    /// | a d g |
+    /// | b e h |
+    /// | c f i |
+    /// @param source the source location
+    /// @param a component 1 of the first column vector
+    /// @param b component 2 of the first column vector
+    /// @param c component 3 of the first column vector
+    /// @param d component 1 of the second column vector
+    /// @param e component 2 of the second column vector
+    /// @param f component 3 of the second column vector
+    /// @param g component 1 of the third column vector
+    /// @param h component 2 of the third column vector
+    /// @param i component 3 of the third column vector
+    template <typename NumberT>
+    utils::Result<NumberT> Det3(const Source& source,
+                                NumberT a,
+                                NumberT b,
+                                NumberT c,
+                                NumberT d,
+                                NumberT e,
+                                NumberT f,
+                                NumberT g,
+                                NumberT h,
+                                NumberT i);
+
+    /// Returns the determinant of the 4x4 matrix:
+    /// | a e i m |
+    /// | b f j n |
+    /// | c g k o |
+    /// | d h l p |
+    /// @param source the source location
+    /// @param a component 1 of the first column vector
+    /// @param b component 2 of the first column vector
+    /// @param c component 3 of the first column vector
+    /// @param d component 4 of the first column vector
+    /// @param e component 1 of the second column vector
+    /// @param f component 2 of the second column vector
+    /// @param g component 3 of the second column vector
+    /// @param h component 4 of the second column vector
+    /// @param i component 1 of the third column vector
+    /// @param j component 2 of the third column vector
+    /// @param k component 3 of the third column vector
+    /// @param l component 4 of the third column vector
+    /// @param m component 1 of the fourth column vector
+    /// @param n component 2 of the fourth column vector
+    /// @param o component 3 of the fourth column vector
+    /// @param p component 4 of the fourth column vector
+    template <typename NumberT>
+    utils::Result<NumberT> Det4(const Source& source,
+                                NumberT a,
+                                NumberT b,
+                                NumberT c,
+                                NumberT d,
+                                NumberT e,
+                                NumberT f,
+                                NumberT g,
+                                NumberT h,
+                                NumberT i,
+                                NumberT j,
+                                NumberT k,
+                                NumberT l,
+                                NumberT m,
+                                NumberT n,
+                                NumberT o,
+                                NumberT p);
 
     template <typename NumberT>
     utils::Result<NumberT> Sqrt(const Source& source, NumberT v);
@@ -1093,6 +1172,20 @@ class ConstEval {
     /// @returns the callable function
     auto Det2Func(const Source& source, const sem::Type* elem_ty);
 
+    /// Returns a callable that calls Det3, and creates a Constant with its result of type `elem_ty`
+    /// if successful, or returns Failure otherwise.
+    /// @param source the source location
+    /// @param elem_ty the element type of the Constant to create on success
+    /// @returns the callable function
+    auto Det3Func(const Source& source, const sem::Type* elem_ty);
+
+    /// Returns a callable that calls Det4, and creates a Constant with its result of type `elem_ty`
+    /// if successful, or returns Failure otherwise.
+    /// @param source the source location
+    /// @param elem_ty the element type of the Constant to create on success
+    /// @returns the callable function
+    auto Det4Func(const Source& source, 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 source the source location

src/tint/resolver/const_eval_builtin_test.cc

@@ -853,6 +853,83 @@ INSTANTIATE_TEST_SUITE_P(  //
                                               DotCases<f32>(),     //
                                               DotCases<f16>()))));
 
+template <typename T>
+std::vector<Case> DeterminantCases() {
+    auto error_msg = [](auto a, const char* op, auto b) {
+        return "12:34 error: " + OverflowErrorMessage(a, op, b) + R"(
+12:34 note: when calculating determinant)";
+    };
+
+    auto r = std::vector<Case>{
+        // All zero == 0
+        C({Mat({T(0), T(0)},    //
+               {T(0), T(0)})},  //
+          Val(T(0))),
+
+        C({Mat({T(0), T(0), T(0)},    //
+               {T(0), T(0), T(0)},    //
+               {T(0), T(0), T(0)})},  //
+          Val(T(0))),
+
+        C({Mat({T(0), T(0), T(0), T(0)},    //
+               {T(0), T(0), T(0), T(0)},    //
+               {T(0), T(0), T(0), T(0)},    //
+               {T(0), T(0), T(0), T(0)})},  //
+          Val(T(0))),
+
+        // All same == 0
+        C({Mat({T(42), T(42)},    //
+               {T(42), T(42)})},  //
+          Val(T(0))),
+
+        C({Mat({T(42), T(42), T(42)},    //
+               {T(42), T(42), T(42)},    //
+               {T(42), T(42), T(42)})},  //
+          Val(T(0))),
+
+        C({Mat({T(42), T(42), T(42), T(42)},    //
+               {T(42), T(42), T(42), T(42)},    //
+               {T(42), T(42), T(42), T(42)},    //
+               {T(42), T(42), T(42), T(42)})},  //
+          Val(T(0))),
+
+        // Various values
+        C({Mat({-T(2), T(17)},   //
+               {T(5), T(45)})},  //
+          Val(-T(175))),
+
+        C({Mat({T(4), T(6), -T(13)},    //
+               {T(12), T(5), T(8)},     //
+               {T(9), T(17), T(16)})},  //
+          Val(-T(3011))),
+
+        C({Mat({T(2), T(9), T(8), T(1)},       //
+               {-T(4), T(11), -T(3), T(7)},    //
+               {T(6), T(5), T(12), -T(6)},     //
+               {T(3), -T(10), T(4), -T(7)})},  //
+          Val(T(469))),
+
+        // Overflow during multiply
+        E({Mat({T::Highest(), T(0)},  //
+               {T(0), T(2)})},        //
+          error_msg(T::Highest(), "*", T(2))),
+
+        // Overflow during subtract
+        E({Mat({T::Highest(), T::Lowest()},  //
+               {T(1), T(1)})},               //
+          error_msg(T::Highest(), "-", T::Lowest())),
+    };
+
+    return r;
+}
+INSTANTIATE_TEST_SUITE_P(  //
+    Determinant,
+    ResolverConstEvalBuiltinTest,
+    testing::Combine(testing::Values(sem::BuiltinType::kDeterminant),
+                     testing::ValuesIn(Concat(DeterminantCases<AFloat>(),  //
+                                              DeterminantCases<f32>(),     //
+                                              DeterminantCases<f16>()))));
+
 template <typename T>
 std::vector<Case> FirstLeadingBitCases() {
     using B = BitValues<T>;

src/tint/resolver/const_eval_test.h

@@ -288,6 +288,16 @@ using Types = std::variant<  //
     Value<builder::mat2x2<f32>>,
     Value<builder::mat2x2<f16>>,
 
+    Value<builder::mat3x3<AInt>>,
+    Value<builder::mat3x3<AFloat>>,
+    Value<builder::mat3x3<f32>>,
+    Value<builder::mat3x3<f16>>,
+
+    Value<builder::mat4x4<AInt>>,
+    Value<builder::mat4x4<AFloat>>,
+    Value<builder::mat4x4<f32>>,
+    Value<builder::mat4x4<f16>>,
+
     Value<builder::mat2x3<AInt>>,
     Value<builder::mat2x3<AFloat>>,
     Value<builder::mat2x3<f32>>,

src/tint/resolver/intrinsic_table.inl

@@ -13566,12 +13566,12 @@ constexpr OverloadInfo kOverloads[] = {
     /* num parameters */ 1,
     /* num template types */ 1,
     /* num template numbers */ 1,
-    /* template types */ &kTemplateTypes[26],
+    /* template types */ &kTemplateTypes[23],
     /* template numbers */ &kTemplateNumbers[4],
     /* parameters */ &kParameters[837],
     /* return matcher indices */ &kMatcherIndices[3],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
-    /* const eval */ nullptr,
+    /* const eval */ &ConstEval::determinant,
   },
   {
     /* [438] */
@@ -14124,7 +14124,7 @@ constexpr IntrinsicInfo kBuiltins[] = {
   },
   {
     /* [20] */
-    /* fn determinant<N : num, T : f32_f16>(mat<N, N, T>) -> T */
+    /* fn determinant<N : num, T : fa_f32_f16>(mat<N, N, T>) -> T */
     /* num overloads */ 1,
     /* overloads */ &kOverloads[437],
   },

test/tint/bug/tint/1061.spvasm.expected.dxc.hlsl

@@ -6,7 +6,7 @@ static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f);
 void main_1() {
   float f = 0.0f;
   float4 v = float4(0.0f, 0.0f, 0.0f, 0.0f);
-  f = determinant(float3x3(float3(1.0f, 0.0f, 0.0f), float3(0.0f, 1.0f, 0.0f), float3(0.0f, 0.0f, 1.0f)));
+  f = 1.0f;
   const float x_33 = f;
   const float x_35 = f;
   const float x_37 = f;

test/tint/bug/tint/1061.spvasm.expected.fxc.hlsl

@@ -6,7 +6,7 @@ static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f);
 void main_1() {
   float f = 0.0f;
   float4 v = float4(0.0f, 0.0f, 0.0f, 0.0f);
-  f = determinant(float3x3(float3(1.0f, 0.0f, 0.0f), float3(0.0f, 1.0f, 0.0f), float3(0.0f, 0.0f, 1.0f)));
+  f = 1.0f;
   const float x_33 = f;
   const float x_35 = f;
   const float x_37 = f;

test/tint/bug/tint/1061.spvasm.expected.glsl

@@ -14,7 +14,7 @@ vec4 x_GLF_color = vec4(0.0f, 0.0f, 0.0f, 0.0f);
 void main_1() {
   float f = 0.0f;
   vec4 v = vec4(0.0f, 0.0f, 0.0f, 0.0f);
-  f = determinant(mat3(vec3(1.0f, 0.0f, 0.0f), vec3(0.0f, 1.0f, 0.0f), vec3(0.0f, 0.0f, 1.0f)));
+  f = 1.0f;
   float x_33 = f;
   float x_35 = f;
   float x_37 = f;

test/tint/bug/tint/1061.spvasm.expected.msl

@@ -8,7 +8,7 @@ struct buf0 {
 void main_1(const constant buf0* const tint_symbol_3, thread float4* const tint_symbol_4) {
   float f = 0.0f;
   float4 v = 0.0f;
-  f = determinant(float3x3(float3(1.0f, 0.0f, 0.0f), float3(0.0f, 1.0f, 0.0f), float3(0.0f, 0.0f, 1.0f)));
+  f = 1.0f;
   float const x_33 = f;
   float const x_35 = f;
   float const x_37 = f;

test/tint/bug/tint/1061.spvasm.expected.spvasm

@@ -1,10 +1,10 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 64
+; Bound: 57
 ; Schema: 0
                OpCapability Shader
-         %22 = OpExtInstImport "GLSL.std.450"
+         %27 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Fragment %main "main" %x_GLF_color_1_1
                OpExecutionMode %main OriginUpperLeft
@@ -46,64 +46,57 @@
 %_ptr_Function_float = OpTypePointer Function %float
          %18 = OpConstantNull %float
 %_ptr_Function_v4float = OpTypePointer Function %v4float
-    %v3float = OpTypeVector %float 3
-%mat3v3float = OpTypeMatrix %v3float 3
     %float_1 = OpConstant %float 1
-         %26 = OpConstantComposite %v3float %float_1 %18 %18
-         %27 = OpConstantComposite %v3float %18 %float_1 %18
-         %28 = OpConstantComposite %v3float %18 %18 %float_1
-         %29 = OpConstantComposite %mat3v3float %26 %27 %28
        %uint = OpTypeInt 32 0
      %uint_0 = OpConstant %uint 0
 %_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
 %float_0_100000001 = OpConstant %float 0.100000001
        %bool = OpTypeBool
-         %52 = OpConstantComposite %v4float %float_1 %18 %18 %float_1
+         %45 = OpConstantComposite %v4float %float_1 %18 %18 %float_1
    %main_out = OpTypeStruct %v4float
-         %53 = OpTypeFunction %main_out
+         %46 = OpTypeFunction %main_out
      %main_1 = OpFunction %void None %12
          %15 = OpLabel
           %f = OpVariable %_ptr_Function_float Function %18
           %v = OpVariable %_ptr_Function_v4float Function %5
-         %21 = OpExtInst %float %22 Determinant %29
+               OpStore %f %float_1
-               OpStore %f %21
+         %22 = OpLoad %float %f
-         %30 = OpLoad %float %f
+         %23 = OpLoad %float %f
-         %31 = OpLoad %float %f
+         %24 = OpLoad %float %f
-         %32 = OpLoad %float %f
+         %25 = OpLoad %float %f
-         %33 = OpLoad %float %f
+         %26 = OpExtInst %float %27 Sin %22
-         %34 = OpExtInst %float %22 Sin %30
+         %28 = OpExtInst %float %27 Cos %23
-         %35 = OpExtInst %float %22 Cos %31
+         %29 = OpExtInst %float %27 Exp2 %24
-         %36 = OpExtInst %float %22 Exp2 %32
+         %30 = OpExtInst %float %27 Log %25
-         %37 = OpExtInst %float %22 Log %33
+         %31 = OpCompositeConstruct %v4float %26 %28 %29 %30
-         %38 = OpCompositeConstruct %v4float %34 %35 %36 %37
+               OpStore %v %31
-               OpStore %v %38
+         %32 = OpLoad %v4float %v
-         %39 = OpLoad %v4float %v
+         %36 = OpAccessChain %_ptr_Uniform_v4float %x_7 %uint_0 %uint_0
-         %43 = OpAccessChain %_ptr_Uniform_v4float %x_7 %uint_0 %uint_0
+         %37 = OpLoad %v4float %36
-         %44 = OpLoad %v4float %43
+         %38 = OpExtInst %float %27 Distance %32 %37
-         %45 = OpExtInst %float %22 Distance %39 %44
+         %40 = OpFOrdLessThan %bool %38 %float_0_100000001
-         %47 = OpFOrdLessThan %bool %45 %float_0_100000001
+               OpSelectionMerge %42 None
-               OpSelectionMerge %49 None
+               OpBranchConditional %40 %43 %44
-               OpBranchConditional %47 %50 %51
+         %43 = OpLabel
-         %50 = OpLabel
+               OpStore %x_GLF_color %45
-               OpStore %x_GLF_color %52
+               OpBranch %42
-               OpBranch %49
+         %44 = OpLabel
-         %51 = OpLabel
                OpStore %x_GLF_color %5
-               OpBranch %49
+               OpBranch %42
-         %49 = OpLabel
+         %42 = OpLabel
                OpReturn
                OpFunctionEnd
- %main_inner = OpFunction %main_out None %53
+ %main_inner = OpFunction %main_out None %46
-         %56 = OpLabel
+         %49 = OpLabel
-         %57 = OpFunctionCall %void %main_1
+         %50 = OpFunctionCall %void %main_1
-         %58 = OpLoad %v4float %x_GLF_color
+         %51 = OpLoad %v4float %x_GLF_color
-         %59 = OpCompositeConstruct %main_out %58
+         %52 = OpCompositeConstruct %main_out %51
-               OpReturnValue %59
+               OpReturnValue %52
                OpFunctionEnd
        %main = OpFunction %void None %12
-         %61 = OpLabel
+         %54 = OpLabel
-         %62 = OpFunctionCall %main_out %main_inner
+         %55 = OpFunctionCall %main_out %main_inner
-         %63 = OpCompositeExtract %v4float %62 0
+         %56 = OpCompositeExtract %v4float %55 0
-               OpStore %x_GLF_color_1_1 %63
+               OpStore %x_GLF_color_1_1 %56
                OpReturn
                OpFunctionEnd

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

test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.msl

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

test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.spvasm

@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; 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 %determinant_1bf6e7 "determinant_1bf6e7"
+               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
+%_ptr_Function_float = OpTypePointer Function %float
+         %15 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%determinant_1bf6e7 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_float Function %8
+               OpStore %res %8
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %15
+         %17 = OpLabel
+         %18 = OpFunctionCall %void %determinant_1bf6e7
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %20 = OpLabel
+         %21 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %21
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %24 = OpLabel
+         %25 = OpFunctionCall %void %determinant_1bf6e7
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %determinant_1bf6e7
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void determinant_2b62ba() {
-  float res = determinant(float3x3((1.0f).xxx, (1.0f).xxx, (1.0f).xxx));
+  float res = 0.0f;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void determinant_2b62ba() {
-  float res = determinant(float3x3((1.0f).xxx, (1.0f).xxx, (1.0f).xxx));
+  float res = 0.0f;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void determinant_2b62ba() {
-  float res = determinant(float3x3(float3(1.0f), float3(1.0f), float3(1.0f)));
+  float res = 0.0f;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.spvasm

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

test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void determinant_32bfde() {
-  float16_t res = determinant(matrix<float16_t, 4, 4>((float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx));
+  float16_t res = float16_t(0.0h);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.glsl

@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void determinant_32bfde() {
-  float16_t res = determinant(f16mat4(f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf)));
+  float16_t res = 0.0hf;
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@ void main() {
 precision mediump float;
 
 void determinant_32bfde() {
-  float16_t res = determinant(f16mat4(f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf)));
+  float16_t res = 0.0hf;
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@ void main() {
 #extension GL_AMD_gpu_shader_half_float : require
 
 void determinant_32bfde() {
-  float16_t res = determinant(f16mat4(f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf), f16vec4(1.0hf)));
+  float16_t res = 0.0hf;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void determinant_32bfde() {
-  half res = determinant(half4x4(half4(1.0h), half4(1.0h), half4(1.0h), half4(1.0h)));
+  half res = 0.0h;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.spvasm

@@ -1,14 +1,13 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 38
+; Bound: 31
 ; Schema: 0
                OpCapability Shader
                OpCapability Float16
                OpCapability UniformAndStorageBuffer16BitAccess
                OpCapability StorageBuffer16BitAccess
                OpCapability StorageInputOutput16
-         %15 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -36,41 +35,35 @@
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
        %half = OpTypeFloat 16
-     %v4half = OpTypeVector %half 4
+         %14 = OpConstantNull %half
- %mat4v4half = OpTypeMatrix %v4half 4
-%half_0x1p_0 = OpConstant %half 0x1p+0
-         %19 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
-         %20 = OpConstantComposite %mat4v4half %19 %19 %19 %19
 %_ptr_Function_half = OpTypePointer Function %half
-         %23 = OpConstantNull %half
+         %17 = OpTypeFunction %v4float
-         %24 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
 %determinant_32bfde = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_half Function %23
+        %res = OpVariable %_ptr_Function_half Function %14
-         %13 = OpExtInst %half %15 Determinant %20
+               OpStore %res %14
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %24
+%vertex_main_inner = OpFunction %v4float None %17
-         %26 = OpLabel
+         %19 = OpLabel
-         %27 = OpFunctionCall %void %determinant_32bfde
+         %20 = OpFunctionCall %void %determinant_32bfde
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %29 = OpLabel
+         %22 = OpLabel
-         %30 = OpFunctionCall %v4float %vertex_main_inner
+         %23 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %30
+               OpStore %value %23
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %33 = OpLabel
+         %26 = OpLabel
-         %34 = OpFunctionCall %void %determinant_32bfde
+         %27 = OpFunctionCall %void %determinant_32bfde
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %36 = OpLabel
+         %29 = OpLabel
-         %37 = OpFunctionCall %void %determinant_32bfde
+         %30 = OpFunctionCall %void %determinant_32bfde
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void determinant_a0a87c() {
-  float res = determinant(float4x4((1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx));
+  float res = 0.0f;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void determinant_a0a87c() {
-  float res = determinant(float4x4((1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx));
+  float res = 0.0f;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void determinant_a0a87c() {
-  float res = determinant(float4x4(float4(1.0f), float4(1.0f), float4(1.0f), float4(1.0f)));
+  float res = 0.0f;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.spvasm

@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 29
 ; Schema: 0
                OpCapability Shader
-         %14 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -31,38 +30,34 @@
 %vertex_point_size = OpVariable %_ptr_Output_float Output %8
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
-%mat4v4float = OpTypeMatrix %v4float 4
-    %float_1 = OpConstant %float 1
-         %17 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
-         %18 = OpConstantComposite %mat4v4float %17 %17 %17 %17
 %_ptr_Function_float = OpTypePointer Function %float
-         %21 = OpTypeFunction %v4float
+         %15 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
 %determinant_a0a87c = OpFunction %void None %9
          %12 = OpLabel
         %res = OpVariable %_ptr_Function_float Function %8
-         %13 = OpExtInst %float %14 Determinant %18
+               OpStore %res %8
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
+%vertex_main_inner = OpFunction %v4float None %15
-         %23 = OpLabel
+         %17 = OpLabel
-         %24 = OpFunctionCall %void %determinant_a0a87c
+         %18 = OpFunctionCall %void %determinant_a0a87c
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %26 = OpLabel
+         %20 = OpLabel
-         %27 = OpFunctionCall %v4float %vertex_main_inner
+         %21 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %27
+               OpStore %value %21
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %29 = OpLabel
+         %24 = OpLabel
-         %30 = OpFunctionCall %void %determinant_a0a87c
+         %25 = OpFunctionCall %void %determinant_a0a87c
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %32 = OpLabel
+         %27 = OpLabel
-         %33 = OpFunctionCall %void %determinant_a0a87c
+         %28 = OpFunctionCall %void %determinant_a0a87c
                OpReturn
                OpFunctionEnd

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

test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.msl

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

test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.spvasm

@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; 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 %determinant_c8251d "determinant_c8251d"
+               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
+%_ptr_Function_float = OpTypePointer Function %float
+         %15 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%determinant_c8251d = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_float Function %8
+               OpStore %res %8
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %15
+         %17 = OpLabel
+         %18 = OpFunctionCall %void %determinant_c8251d
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %20 = OpLabel
+         %21 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %21
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %24 = OpLabel
+         %25 = OpFunctionCall %void %determinant_c8251d
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %determinant_c8251d
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.wgsl

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

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

test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.dxc.hlsl

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

test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.fxc.hlsl

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

test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.msl

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

test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.spvasm

@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; 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 %determinant_cefdf3 "determinant_cefdf3"
+               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
+%_ptr_Function_float = OpTypePointer Function %float
+         %15 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
+%determinant_cefdf3 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_float Function %8
+               OpStore %res %8
+               OpReturn
+               OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %15
+         %17 = OpLabel
+         %18 = OpFunctionCall %void %determinant_cefdf3
+               OpReturnValue %5
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %20 = OpLabel
+         %21 = OpFunctionCall %v4float %vertex_main_inner
+               OpStore %value %21
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %24 = OpLabel
+         %25 = OpFunctionCall %void %determinant_cefdf3
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %determinant_cefdf3
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.wgsl

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

test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void determinant_d7c86f() {
-  float16_t res = determinant(matrix<float16_t, 3, 3>((float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx));
+  float16_t res = float16_t(0.0h);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void determinant_d7c86f() {
-  half res = determinant(half3x3(half3(1.0h), half3(1.0h), half3(1.0h)));
+  half res = 0.0h;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.spvasm

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

test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void determinant_e19305() {
-  float res = determinant(float2x2((1.0f).xx, (1.0f).xx));
+  float res = 0.0f;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.fxc.hlsl

@@ -1,5 +1,5 @@
 void determinant_e19305() {
-  float res = determinant(float2x2((1.0f).xx, (1.0f).xx));
+  float res = 0.0f;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.glsl

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

test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void determinant_e19305() {
-  float res = determinant(float2x2(float2(1.0f), float2(1.0f)));
+  float res = 0.0f;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.spvasm

@@ -1,10 +1,9 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 35
+; Bound: 29
 ; Schema: 0
                OpCapability Shader
-         %14 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -31,39 +30,34 @@
 %vertex_point_size = OpVariable %_ptr_Output_float Output %8
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
-    %v2float = OpTypeVector %float 2
-%mat2v2float = OpTypeMatrix %v2float 2
-    %float_1 = OpConstant %float 1
-         %18 = OpConstantComposite %v2float %float_1 %float_1
-         %19 = OpConstantComposite %mat2v2float %18 %18
 %_ptr_Function_float = OpTypePointer Function %float
-         %22 = OpTypeFunction %v4float
+         %15 = OpTypeFunction %v4float
+    %float_1 = OpConstant %float 1
 %determinant_e19305 = OpFunction %void None %9
          %12 = OpLabel
         %res = OpVariable %_ptr_Function_float Function %8
-         %13 = OpExtInst %float %14 Determinant %19
+               OpStore %res %8
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
+%vertex_main_inner = OpFunction %v4float None %15
-         %24 = OpLabel
+         %17 = OpLabel
-         %25 = OpFunctionCall %void %determinant_e19305
+         %18 = OpFunctionCall %void %determinant_e19305
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %27 = OpLabel
+         %20 = OpLabel
-         %28 = OpFunctionCall %v4float %vertex_main_inner
+         %21 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %28
+               OpStore %value %21
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %30 = OpLabel
+         %24 = OpLabel
-         %31 = OpFunctionCall %void %determinant_e19305
+         %25 = OpFunctionCall %void %determinant_e19305
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %33 = OpLabel
+         %27 = OpLabel
-         %34 = OpFunctionCall %void %determinant_e19305
+         %28 = OpFunctionCall %void %determinant_e19305
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.dxc.hlsl

@@ -1,5 +1,5 @@
 void determinant_fc12a5() {
-  float16_t res = determinant(matrix<float16_t, 2, 2>((float16_t(1.0h)).xx, (float16_t(1.0h)).xx));
+  float16_t res = float16_t(0.0h);
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.glsl

@@ -2,7 +2,7 @@
 #extension GL_AMD_gpu_shader_half_float : require
 
 void determinant_fc12a5() {
-  float16_t res = determinant(f16mat2(f16vec2(1.0hf), f16vec2(1.0hf)));
+  float16_t res = 0.0hf;
 }
 
 vec4 vertex_main() {
@@ -23,7 +23,7 @@ void main() {
 precision mediump float;
 
 void determinant_fc12a5() {
-  float16_t res = determinant(f16mat2(f16vec2(1.0hf), f16vec2(1.0hf)));
+  float16_t res = 0.0hf;
 }
 
 void fragment_main() {
@@ -38,7 +38,7 @@ void main() {
 #extension GL_AMD_gpu_shader_half_float : require
 
 void determinant_fc12a5() {
-  float16_t res = determinant(f16mat2(f16vec2(1.0hf), f16vec2(1.0hf)));
+  float16_t res = 0.0hf;
 }
 
 void compute_main() {

test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void determinant_fc12a5() {
-  half res = determinant(half2x2(half2(1.0h), half2(1.0h)));
+  half res = 0.0h;
 }
 
 struct tint_symbol {

test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.spvasm

@@ -1,14 +1,13 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 38
+; Bound: 31
 ; Schema: 0
                OpCapability Shader
                OpCapability Float16
                OpCapability UniformAndStorageBuffer16BitAccess
                OpCapability StorageBuffer16BitAccess
                OpCapability StorageInputOutput16
-         %15 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
                OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -36,41 +35,35 @@
        %void = OpTypeVoid
           %9 = OpTypeFunction %void
        %half = OpTypeFloat 16
-     %v2half = OpTypeVector %half 2
+         %14 = OpConstantNull %half
- %mat2v2half = OpTypeMatrix %v2half 2
-%half_0x1p_0 = OpConstant %half 0x1p+0
-         %19 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
-         %20 = OpConstantComposite %mat2v2half %19 %19
 %_ptr_Function_half = OpTypePointer Function %half
-         %23 = OpConstantNull %half
+         %17 = OpTypeFunction %v4float
-         %24 = OpTypeFunction %v4float
     %float_1 = OpConstant %float 1
 %determinant_fc12a5 = OpFunction %void None %9
          %12 = OpLabel
-        %res = OpVariable %_ptr_Function_half Function %23
+        %res = OpVariable %_ptr_Function_half Function %14
-         %13 = OpExtInst %half %15 Determinant %20
+               OpStore %res %14
-               OpStore %res %13
                OpReturn
                OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %24
+%vertex_main_inner = OpFunction %v4float None %17
-         %26 = OpLabel
+         %19 = OpLabel
-         %27 = OpFunctionCall %void %determinant_fc12a5
+         %20 = OpFunctionCall %void %determinant_fc12a5
                OpReturnValue %5
                OpFunctionEnd
 %vertex_main = OpFunction %void None %9
-         %29 = OpLabel
+         %22 = OpLabel
-         %30 = OpFunctionCall %v4float %vertex_main_inner
+         %23 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %30
+               OpStore %value %23
                OpStore %vertex_point_size %float_1
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %9
-         %33 = OpLabel
+         %26 = OpLabel
-         %34 = OpFunctionCall %void %determinant_fc12a5
+         %27 = OpFunctionCall %void %determinant_fc12a5
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %9
-         %36 = OpLabel
+         %29 = OpLabel
-         %37 = OpFunctionCall %void %determinant_fc12a5
+         %30 = OpFunctionCall %void %determinant_fc12a5
                OpReturn
                OpFunctionEnd

test/tint/builtins/gen/var/determinant/1bf6e7.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 determinant(mat<2, 2, fa>) -> fa
+fn determinant_1bf6e7() {
+  const arg_0 = mat2x2(1., 1., 1., 1.);
+  var res = determinant(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  determinant_1bf6e7();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  determinant_1bf6e7();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  determinant_1bf6e7();
+}

test/tint/builtins/gen/var/determinant/c8251d.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 determinant(mat<3, 3, fa>) -> fa
+fn determinant_c8251d() {
+  const arg_0 = mat3x3(1., 1., 1., 1., 1., 1., 1., 1., 1.);
+  var res = determinant(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  determinant_c8251d();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  determinant_c8251d();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  determinant_c8251d();
+}

test/tint/builtins/gen/var/determinant/cefdf3.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 determinant(mat<4, 4, fa>) -> fa
+fn determinant_cefdf3() {
+  const arg_0 = mat4x4(1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1., 1.);
+  var res = determinant(arg_0);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  determinant_cefdf3();
+  return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+  determinant_cefdf3();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  determinant_cefdf3();
+}