From 05c8daac420965194364a6647bfadf5322c284c2 Mon Sep 17 00:00:00 2001 From: Antonio Maiorano Date: Wed, 23 Nov 2022 19:16:15 +0000 Subject: [PATCH] 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 Reviewed-by: Dan Sinclair Commit-Queue: Antonio Maiorano Kokoro: Kokoro Kokoro: Antonio Maiorano --- src/tint/intrinsics.def | 2 +- src/tint/resolver/builtin_test.cc | 4 +- src/tint/resolver/const_eval.cc | 197 +++++++++++++++++- src/tint/resolver/const_eval.h | 113 +++++++++- src/tint/resolver/const_eval_builtin_test.cc | 77 +++++++ src/tint/resolver/const_eval_test.h | 10 + src/tint/resolver/intrinsic_table.inl | 6 +- .../bug/tint/1061.spvasm.expected.dxc.hlsl | 2 +- .../bug/tint/1061.spvasm.expected.fxc.hlsl | 2 +- test/tint/bug/tint/1061.spvasm.expected.glsl | 2 +- test/tint/bug/tint/1061.spvasm.expected.msl | 2 +- .../tint/bug/tint/1061.spvasm.expected.spvasm | 83 ++++---- .../gen/literal/determinant/1bf6e7.wgsl | 43 ++++ .../determinant/1bf6e7.wgsl.expected.dxc.hlsl | 30 +++ .../determinant/1bf6e7.wgsl.expected.fxc.hlsl | 30 +++ .../determinant/1bf6e7.wgsl.expected.glsl | 49 +++++ .../determinant/1bf6e7.wgsl.expected.msl | 33 +++ .../determinant/1bf6e7.wgsl.expected.spvasm | 63 ++++++ .../determinant/1bf6e7.wgsl.expected.wgsl | 19 ++ .../determinant/2b62ba.wgsl.expected.dxc.hlsl | 2 +- .../determinant/2b62ba.wgsl.expected.fxc.hlsl | 2 +- .../determinant/2b62ba.wgsl.expected.glsl | 6 +- .../determinant/2b62ba.wgsl.expected.msl | 2 +- .../determinant/2b62ba.wgsl.expected.spvasm | 34 ++- .../determinant/32bfde.wgsl.expected.dxc.hlsl | 2 +- .../determinant/32bfde.wgsl.expected.glsl | 6 +- .../determinant/32bfde.wgsl.expected.msl | 2 +- .../determinant/32bfde.wgsl.expected.spvasm | 37 ++-- .../determinant/a0a87c.wgsl.expected.dxc.hlsl | 2 +- .../determinant/a0a87c.wgsl.expected.fxc.hlsl | 2 +- .../determinant/a0a87c.wgsl.expected.glsl | 6 +- .../determinant/a0a87c.wgsl.expected.msl | 2 +- .../determinant/a0a87c.wgsl.expected.spvasm | 33 ++- .../gen/literal/determinant/c8251d.wgsl | 43 ++++ .../determinant/c8251d.wgsl.expected.dxc.hlsl | 30 +++ .../determinant/c8251d.wgsl.expected.fxc.hlsl | 30 +++ .../determinant/c8251d.wgsl.expected.glsl | 49 +++++ .../determinant/c8251d.wgsl.expected.msl | 33 +++ .../determinant/c8251d.wgsl.expected.spvasm | 63 ++++++ .../determinant/c8251d.wgsl.expected.wgsl | 19 ++ .../gen/literal/determinant/cefdf3.wgsl | 43 ++++ .../determinant/cefdf3.wgsl.expected.dxc.hlsl | 30 +++ .../determinant/cefdf3.wgsl.expected.fxc.hlsl | 30 +++ .../determinant/cefdf3.wgsl.expected.glsl | 49 +++++ .../determinant/cefdf3.wgsl.expected.msl | 33 +++ .../determinant/cefdf3.wgsl.expected.spvasm | 63 ++++++ .../determinant/cefdf3.wgsl.expected.wgsl | 19 ++ .../determinant/d7c86f.wgsl.expected.dxc.hlsl | 2 +- .../determinant/d7c86f.wgsl.expected.glsl | 6 +- .../determinant/d7c86f.wgsl.expected.msl | 2 +- .../determinant/d7c86f.wgsl.expected.spvasm | 37 ++-- .../determinant/e19305.wgsl.expected.dxc.hlsl | 2 +- .../determinant/e19305.wgsl.expected.fxc.hlsl | 2 +- .../determinant/e19305.wgsl.expected.glsl | 6 +- .../determinant/e19305.wgsl.expected.msl | 2 +- .../determinant/e19305.wgsl.expected.spvasm | 34 ++- .../determinant/fc12a5.wgsl.expected.dxc.hlsl | 2 +- .../determinant/fc12a5.wgsl.expected.glsl | 6 +- .../determinant/fc12a5.wgsl.expected.msl | 2 +- .../determinant/fc12a5.wgsl.expected.spvasm | 37 ++-- .../builtins/gen/var/determinant/1bf6e7.wgsl | 44 ++++ .../builtins/gen/var/determinant/c8251d.wgsl | 44 ++++ .../builtins/gen/var/determinant/cefdf3.wgsl | 44 ++++ 63 files changed, 1482 insertions(+), 229 deletions(-) create mode 100644 test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl create mode 100644 test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/literal/determinant/c8251d.wgsl create mode 100644 test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/literal/determinant/cefdf3.wgsl create mode 100644 test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/var/determinant/1bf6e7.wgsl create mode 100644 test/tint/builtins/gen/var/determinant/c8251d.wgsl create mode 100644 test/tint/builtins/gen/var/determinant/cefdf3.wgsl diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def index 6c17ec6111..57eb10d7cc 100644 --- a/src/tint/intrinsics.def +++ b/src/tint/intrinsics.def @@ -442,7 +442,7 @@ fn arrayLength(ptr, A>) -> u32 @const fn cross(vec3, vec3) -> vec3 @const fn degrees(T) -> T @const fn degrees(vec) -> vec -fn determinant(mat) -> T +@const fn determinant(mat) -> T fn distance(T, T) -> T fn distance(vec, vec) -> T @const fn dot(vec, vec) -> T diff --git a/src/tint/resolver/builtin_test.cc b/src/tint/resolver/builtin_test.cc index 07ade34acb..44f4bc5475 100644 --- a/src/tint/resolver/builtin_test.cc +++ b/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) 1 candidate function: - determinant(matNxN) -> T where: T is f32 or f16 + determinant(matNxN) -> 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 where: T is f32 or f16 + determinant(matNxN) -> T where: T is abstract-float, f32 or f16 )"); } diff --git a/src/tint/resolver/const_eval.cc b/src/tint/resolver/const_eval.cc index b0595f2c15..2019f5955e 100644 --- a/src/tint/resolver/const_eval.cc +++ b/src/tint/resolver/const_eval.cc @@ -893,15 +893,22 @@ utils::Result ConstEval::Dot4(const Source& source, template utils::Result ConstEval::Det2(const Source& source, - NumberT a1, - NumberT a2, - NumberT b1, - NumberT b2) { - auto r1 = Mul(source, a1, b2); + NumberT a, + NumberT b, + NumberT c, + NumberT d) { + // | 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 ConstEval::Det2(const Source& source, return r; } +template +utils::Result 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 +utils::Result 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 utils::Result 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 args, + const Source& source) { + auto calculate = [&]() -> ImplResult { + auto* m = args[0]; + auto* mat_ty = m->Type()->As(); + 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 args, const Source& source) { diff --git a/src/tint/resolver/const_eval.h b/src/tint/resolver/const_eval.h index 60407297c7..fd253b4fb6 100644 --- a/src/tint/resolver/const_eval.h +++ b/src/tint/resolver/const_eval.h @@ -539,6 +539,7 @@ class ConstEval { utils::VectorRef 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 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 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 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 + /// @param a component 1 of the first column vector + /// @param b component 2 of the first column vector + /// @param c component 1 of the second column vector + /// @param d component 2 of the second column vector template - utils::Result Det2(const Source& source, - NumberT a1, - NumberT a2, - NumberT b1, - NumberT b2); + utils::Result Det2(const Source& source, // + NumberT a, + NumberT b, + NumberT c, + 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 + utils::Result 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 + utils::Result 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 utils::Result 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 diff --git a/src/tint/resolver/const_eval_builtin_test.cc b/src/tint/resolver/const_eval_builtin_test.cc index c2880779ce..1bf742311c 100644 --- a/src/tint/resolver/const_eval_builtin_test.cc +++ b/src/tint/resolver/const_eval_builtin_test.cc @@ -853,6 +853,83 @@ INSTANTIATE_TEST_SUITE_P( // DotCases(), // DotCases())))); +template +std::vector 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{ + // 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(), // + DeterminantCases(), // + DeterminantCases())))); + template std::vector FirstLeadingBitCases() { using B = BitValues; diff --git a/src/tint/resolver/const_eval_test.h b/src/tint/resolver/const_eval_test.h index 8975752285..148531d163 100644 --- a/src/tint/resolver/const_eval_test.h +++ b/src/tint/resolver/const_eval_test.h @@ -288,6 +288,16 @@ using Types = std::variant< // Value>, Value>, + Value>, + Value>, + Value>, + Value>, + + Value>, + Value>, + Value>, + Value>, + Value>, Value>, Value>, diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl index ee79ab1363..aff12e4c7f 100644 --- a/src/tint/resolver/intrinsic_table.inl +++ b/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(mat) -> T */ + /* fn determinant(mat) -> T */ /* num overloads */ 1, /* overloads */ &kOverloads[437], }, diff --git a/test/tint/bug/tint/1061.spvasm.expected.dxc.hlsl b/test/tint/bug/tint/1061.spvasm.expected.dxc.hlsl index e8ae5f0fb2..319b7c3e92 100644 --- a/test/tint/bug/tint/1061.spvasm.expected.dxc.hlsl +++ b/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; diff --git a/test/tint/bug/tint/1061.spvasm.expected.fxc.hlsl b/test/tint/bug/tint/1061.spvasm.expected.fxc.hlsl index e8ae5f0fb2..319b7c3e92 100644 --- a/test/tint/bug/tint/1061.spvasm.expected.fxc.hlsl +++ b/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; diff --git a/test/tint/bug/tint/1061.spvasm.expected.glsl b/test/tint/bug/tint/1061.spvasm.expected.glsl index c8518585bb..0a1208992a 100644 --- a/test/tint/bug/tint/1061.spvasm.expected.glsl +++ b/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; diff --git a/test/tint/bug/tint/1061.spvasm.expected.msl b/test/tint/bug/tint/1061.spvasm.expected.msl index f0697f8d1e..2f9478fd67 100644 --- a/test/tint/bug/tint/1061.spvasm.expected.msl +++ b/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; diff --git a/test/tint/bug/tint/1061.spvasm.expected.spvasm b/test/tint/bug/tint/1061.spvasm.expected.spvasm index 946f03cced..78fc8aae80 100644 --- a/test/tint/bug/tint/1061.spvasm.expected.spvasm +++ b/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 %21 - %30 = OpLoad %float %f - %31 = OpLoad %float %f - %32 = OpLoad %float %f - %33 = OpLoad %float %f - %34 = OpExtInst %float %22 Sin %30 - %35 = OpExtInst %float %22 Cos %31 - %36 = OpExtInst %float %22 Exp2 %32 - %37 = OpExtInst %float %22 Log %33 - %38 = OpCompositeConstruct %v4float %34 %35 %36 %37 - OpStore %v %38 - %39 = OpLoad %v4float %v - %43 = OpAccessChain %_ptr_Uniform_v4float %x_7 %uint_0 %uint_0 - %44 = OpLoad %v4float %43 - %45 = OpExtInst %float %22 Distance %39 %44 - %47 = OpFOrdLessThan %bool %45 %float_0_100000001 - OpSelectionMerge %49 None - OpBranchConditional %47 %50 %51 - %50 = OpLabel - OpStore %x_GLF_color %52 - OpBranch %49 - %51 = OpLabel + OpStore %f %float_1 + %22 = OpLoad %float %f + %23 = OpLoad %float %f + %24 = OpLoad %float %f + %25 = OpLoad %float %f + %26 = OpExtInst %float %27 Sin %22 + %28 = OpExtInst %float %27 Cos %23 + %29 = OpExtInst %float %27 Exp2 %24 + %30 = OpExtInst %float %27 Log %25 + %31 = OpCompositeConstruct %v4float %26 %28 %29 %30 + OpStore %v %31 + %32 = OpLoad %v4float %v + %36 = OpAccessChain %_ptr_Uniform_v4float %x_7 %uint_0 %uint_0 + %37 = OpLoad %v4float %36 + %38 = OpExtInst %float %27 Distance %32 %37 + %40 = OpFOrdLessThan %bool %38 %float_0_100000001 + OpSelectionMerge %42 None + OpBranchConditional %40 %43 %44 + %43 = OpLabel + OpStore %x_GLF_color %45 + OpBranch %42 + %44 = OpLabel OpStore %x_GLF_color %5 - OpBranch %49 - %49 = OpLabel + OpBranch %42 + %42 = OpLabel OpReturn OpFunctionEnd - %main_inner = OpFunction %main_out None %53 - %56 = OpLabel - %57 = OpFunctionCall %void %main_1 - %58 = OpLoad %v4float %x_GLF_color - %59 = OpCompositeConstruct %main_out %58 - OpReturnValue %59 + %main_inner = OpFunction %main_out None %46 + %49 = OpLabel + %50 = OpFunctionCall %void %main_1 + %51 = OpLoad %v4float %x_GLF_color + %52 = OpCompositeConstruct %main_out %51 + OpReturnValue %52 OpFunctionEnd %main = OpFunction %void None %12 - %61 = OpLabel - %62 = OpFunctionCall %main_out %main_inner - %63 = OpCompositeExtract %v4float %62 0 - OpStore %x_GLF_color_1_1 %63 + %54 = OpLabel + %55 = OpFunctionCall %main_out %main_inner + %56 = OpCompositeExtract %v4float %55 0 + OpStore %x_GLF_color_1_1 %56 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl b/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl new file mode 100644 index 0000000000..e00541848d --- /dev/null +++ b/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 { + determinant_1bf6e7(); + return vec4(); +} + +@fragment +fn fragment_main() { + determinant_1bf6e7(); +} + +@compute @workgroup_size(1) +fn compute_main() { + determinant_1bf6e7(); +} diff --git a/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..3202766bf5 --- /dev/null +++ b/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; +} diff --git a/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..3202766bf5 --- /dev/null +++ b/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; +} diff --git a/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.glsl b/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.glsl new file mode 100644 index 0000000000..d5b27a98e7 --- /dev/null +++ b/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; +} diff --git a/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.msl b/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.msl new file mode 100644 index 0000000000..3e3494cb51 --- /dev/null +++ b/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +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; +} + diff --git a/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.spvasm new file mode 100644 index 0000000000..b5f4f68b56 --- /dev/null +++ b/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 diff --git a/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/determinant/1bf6e7.wgsl.expected.wgsl new file mode 100644 index 0000000000..0aaf417049 --- /dev/null +++ b/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 { + determinant_1bf6e7(); + return vec4(); +} + +@fragment +fn fragment_main() { + determinant_1bf6e7(); +} + +@compute @workgroup_size(1) +fn compute_main() { + determinant_1bf6e7(); +} diff --git a/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.dxc.hlsl index c9482d7f19..30e80d3aec 100644 --- a/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.dxc.hlsl +++ b/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 { diff --git a/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.fxc.hlsl index c9482d7f19..30e80d3aec 100644 --- a/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.fxc.hlsl +++ b/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 { diff --git a/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.glsl b/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.glsl index d801fb8be7..41fce044c1 100644 --- a/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.glsl +++ b/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() { diff --git a/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.msl b/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.msl index 725abfc8cc..56313b0fec 100644 --- a/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.msl +++ b/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 { diff --git a/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.spvasm index 9673dd3c0d..c7425e839a 100644 --- a/test/tint/builtins/gen/literal/determinant/2b62ba.wgsl.expected.spvasm +++ b/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 %13 + OpStore %res %8 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %22 - %24 = OpLabel - %25 = OpFunctionCall %void %determinant_2b62ba +%vertex_main_inner = OpFunction %v4float None %15 + %17 = OpLabel + %18 = OpFunctionCall %void %determinant_2b62ba OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %27 = OpLabel - %28 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %28 + %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 - %30 = OpLabel - %31 = OpFunctionCall %void %determinant_2b62ba + %24 = OpLabel + %25 = OpFunctionCall %void %determinant_2b62ba OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %33 = OpLabel - %34 = OpFunctionCall %void %determinant_2b62ba + %27 = OpLabel + %28 = OpFunctionCall %void %determinant_2b62ba OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.dxc.hlsl index bae350d219..e602aa22d0 100644 --- a/test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.dxc.hlsl +++ b/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(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 { diff --git a/test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.glsl b/test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.glsl index f94a854c71..2a0aa4fac9 100644 --- a/test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.glsl +++ b/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() { diff --git a/test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.msl b/test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.msl index f416fbf507..c50561de64 100644 --- a/test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.msl +++ b/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 { diff --git a/test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.spvasm index ffe5e4d393..d2fb249d64 100644 --- a/test/tint/builtins/gen/literal/determinant/32bfde.wgsl.expected.spvasm +++ b/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 - %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 + %14 = OpConstantNull %half %_ptr_Function_half = OpTypePointer Function %half - %23 = OpConstantNull %half - %24 = OpTypeFunction %v4float + %17 = OpTypeFunction %v4float %float_1 = OpConstant %float 1 %determinant_32bfde = OpFunction %void None %9 %12 = OpLabel - %res = OpVariable %_ptr_Function_half Function %23 - %13 = OpExtInst %half %15 Determinant %20 - OpStore %res %13 + %res = OpVariable %_ptr_Function_half Function %14 + OpStore %res %14 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %24 - %26 = OpLabel - %27 = OpFunctionCall %void %determinant_32bfde +%vertex_main_inner = OpFunction %v4float None %17 + %19 = OpLabel + %20 = OpFunctionCall %void %determinant_32bfde OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %29 = OpLabel - %30 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %30 + %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 - %33 = OpLabel - %34 = OpFunctionCall %void %determinant_32bfde + %26 = OpLabel + %27 = OpFunctionCall %void %determinant_32bfde OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %36 = OpLabel - %37 = OpFunctionCall %void %determinant_32bfde + %29 = OpLabel + %30 = OpFunctionCall %void %determinant_32bfde OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.dxc.hlsl index 0857fc602e..78a081c8bb 100644 --- a/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.dxc.hlsl +++ b/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 { diff --git a/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.fxc.hlsl index 0857fc602e..78a081c8bb 100644 --- a/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.fxc.hlsl +++ b/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 { diff --git a/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.glsl b/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.glsl index 3d93e51b2c..de14daf8a8 100644 --- a/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.glsl +++ b/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() { diff --git a/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.msl b/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.msl index 9ee1b47ef2..bd6b0a6a0d 100644 --- a/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.msl +++ b/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 { diff --git a/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.spvasm index 57d1c71295..4d1af3a4ed 100644 --- a/test/tint/builtins/gen/literal/determinant/a0a87c.wgsl.expected.spvasm +++ b/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 %13 + OpStore %res %8 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %21 - %23 = OpLabel - %24 = OpFunctionCall %void %determinant_a0a87c +%vertex_main_inner = OpFunction %v4float None %15 + %17 = OpLabel + %18 = OpFunctionCall %void %determinant_a0a87c OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %26 = OpLabel - %27 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %27 + %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 - %29 = OpLabel - %30 = OpFunctionCall %void %determinant_a0a87c + %24 = OpLabel + %25 = OpFunctionCall %void %determinant_a0a87c OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %32 = OpLabel - %33 = OpFunctionCall %void %determinant_a0a87c + %27 = OpLabel + %28 = OpFunctionCall %void %determinant_a0a87c OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/determinant/c8251d.wgsl b/test/tint/builtins/gen/literal/determinant/c8251d.wgsl new file mode 100644 index 0000000000..d38a575a19 --- /dev/null +++ b/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 { + determinant_c8251d(); + return vec4(); +} + +@fragment +fn fragment_main() { + determinant_c8251d(); +} + +@compute @workgroup_size(1) +fn compute_main() { + determinant_c8251d(); +} diff --git a/test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..8415cb1500 --- /dev/null +++ b/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; +} diff --git a/test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..8415cb1500 --- /dev/null +++ b/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; +} diff --git a/test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.glsl b/test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.glsl new file mode 100644 index 0000000000..fb8d99d904 --- /dev/null +++ b/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; +} diff --git a/test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.msl b/test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.msl new file mode 100644 index 0000000000..7132d6f6c4 --- /dev/null +++ b/test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +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; +} + diff --git a/test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.spvasm new file mode 100644 index 0000000000..23a1e6dc03 --- /dev/null +++ b/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 diff --git a/test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/determinant/c8251d.wgsl.expected.wgsl new file mode 100644 index 0000000000..ce0cf9c6ab --- /dev/null +++ b/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 { + determinant_c8251d(); + return vec4(); +} + +@fragment +fn fragment_main() { + determinant_c8251d(); +} + +@compute @workgroup_size(1) +fn compute_main() { + determinant_c8251d(); +} diff --git a/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl b/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl new file mode 100644 index 0000000000..8563a81ab0 --- /dev/null +++ b/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 { + determinant_cefdf3(); + return vec4(); +} + +@fragment +fn fragment_main() { + determinant_cefdf3(); +} + +@compute @workgroup_size(1) +fn compute_main() { + determinant_cefdf3(); +} diff --git a/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..ff286087a0 --- /dev/null +++ b/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; +} diff --git a/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..ff286087a0 --- /dev/null +++ b/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; +} diff --git a/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.glsl b/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.glsl new file mode 100644 index 0000000000..488869b846 --- /dev/null +++ b/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; +} diff --git a/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.msl b/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.msl new file mode 100644 index 0000000000..f687f255ae --- /dev/null +++ b/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +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; +} + diff --git a/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.spvasm new file mode 100644 index 0000000000..39854c3075 --- /dev/null +++ b/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 diff --git a/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/determinant/cefdf3.wgsl.expected.wgsl new file mode 100644 index 0000000000..f3b7dd713b --- /dev/null +++ b/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 { + determinant_cefdf3(); + return vec4(); +} + +@fragment +fn fragment_main() { + determinant_cefdf3(); +} + +@compute @workgroup_size(1) +fn compute_main() { + determinant_cefdf3(); +} diff --git a/test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.dxc.hlsl index 083d524dfc..ecaeb94301 100644 --- a/test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.dxc.hlsl +++ b/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(1.0h)).xxx, (float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx)); + float16_t res = float16_t(0.0h); } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.glsl b/test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.glsl index 064604cedf..2c159929f4 100644 --- a/test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.glsl +++ b/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() { diff --git a/test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.msl b/test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.msl index 17c01fbbd4..2855ff5ab9 100644 --- a/test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.msl +++ b/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 { diff --git a/test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.spvasm index c4cb0a3b00..132435283b 100644 --- a/test/tint/builtins/gen/literal/determinant/d7c86f.wgsl.expected.spvasm +++ b/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 - %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 + %14 = OpConstantNull %half %_ptr_Function_half = OpTypePointer Function %half - %23 = OpConstantNull %half - %24 = OpTypeFunction %v4float + %17 = OpTypeFunction %v4float %float_1 = OpConstant %float 1 %determinant_d7c86f = OpFunction %void None %9 %12 = OpLabel - %res = OpVariable %_ptr_Function_half Function %23 - %13 = OpExtInst %half %15 Determinant %20 - OpStore %res %13 + %res = OpVariable %_ptr_Function_half Function %14 + OpStore %res %14 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %24 - %26 = OpLabel - %27 = OpFunctionCall %void %determinant_d7c86f +%vertex_main_inner = OpFunction %v4float None %17 + %19 = OpLabel + %20 = OpFunctionCall %void %determinant_d7c86f OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %29 = OpLabel - %30 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %30 + %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 - %33 = OpLabel - %34 = OpFunctionCall %void %determinant_d7c86f + %26 = OpLabel + %27 = OpFunctionCall %void %determinant_d7c86f OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %36 = OpLabel - %37 = OpFunctionCall %void %determinant_d7c86f + %29 = OpLabel + %30 = OpFunctionCall %void %determinant_d7c86f OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.dxc.hlsl index 0c372a3a13..de9399c76c 100644 --- a/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.dxc.hlsl +++ b/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 { diff --git a/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.fxc.hlsl index 0c372a3a13..de9399c76c 100644 --- a/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.fxc.hlsl +++ b/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 { diff --git a/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.glsl b/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.glsl index cccb7b6faf..0a2470132a 100644 --- a/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.glsl +++ b/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() { diff --git a/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.msl b/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.msl index 47b8bed88e..c592c006fc 100644 --- a/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.msl +++ b/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 { diff --git a/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.spvasm index 441f89e453..4a86d63117 100644 --- a/test/tint/builtins/gen/literal/determinant/e19305.wgsl.expected.spvasm +++ b/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 %13 + OpStore %res %8 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %22 - %24 = OpLabel - %25 = OpFunctionCall %void %determinant_e19305 +%vertex_main_inner = OpFunction %v4float None %15 + %17 = OpLabel + %18 = OpFunctionCall %void %determinant_e19305 OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %27 = OpLabel - %28 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %28 + %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 - %30 = OpLabel - %31 = OpFunctionCall %void %determinant_e19305 + %24 = OpLabel + %25 = OpFunctionCall %void %determinant_e19305 OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %33 = OpLabel - %34 = OpFunctionCall %void %determinant_e19305 + %27 = OpLabel + %28 = OpFunctionCall %void %determinant_e19305 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.dxc.hlsl index 698a14c9bb..b2730ff0a5 100644 --- a/test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.dxc.hlsl +++ b/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(1.0h)).xx, (float16_t(1.0h)).xx)); + float16_t res = float16_t(0.0h); } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.glsl b/test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.glsl index b647997531..eb4ac54e9e 100644 --- a/test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.glsl +++ b/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() { diff --git a/test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.msl b/test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.msl index a639c13a93..51ce8b3cd8 100644 --- a/test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.msl +++ b/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 { diff --git a/test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.spvasm index 28542451af..9a8aa11103 100644 --- a/test/tint/builtins/gen/literal/determinant/fc12a5.wgsl.expected.spvasm +++ b/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 - %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 + %14 = OpConstantNull %half %_ptr_Function_half = OpTypePointer Function %half - %23 = OpConstantNull %half - %24 = OpTypeFunction %v4float + %17 = OpTypeFunction %v4float %float_1 = OpConstant %float 1 %determinant_fc12a5 = OpFunction %void None %9 %12 = OpLabel - %res = OpVariable %_ptr_Function_half Function %23 - %13 = OpExtInst %half %15 Determinant %20 - OpStore %res %13 + %res = OpVariable %_ptr_Function_half Function %14 + OpStore %res %14 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %24 - %26 = OpLabel - %27 = OpFunctionCall %void %determinant_fc12a5 +%vertex_main_inner = OpFunction %v4float None %17 + %19 = OpLabel + %20 = OpFunctionCall %void %determinant_fc12a5 OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %29 = OpLabel - %30 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %30 + %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 - %33 = OpLabel - %34 = OpFunctionCall %void %determinant_fc12a5 + %26 = OpLabel + %27 = OpFunctionCall %void %determinant_fc12a5 OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %36 = OpLabel - %37 = OpFunctionCall %void %determinant_fc12a5 + %29 = OpLabel + %30 = OpFunctionCall %void %determinant_fc12a5 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/var/determinant/1bf6e7.wgsl b/test/tint/builtins/gen/var/determinant/1bf6e7.wgsl new file mode 100644 index 0000000000..a10c782d61 --- /dev/null +++ b/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 { + determinant_1bf6e7(); + return vec4(); +} + +@fragment +fn fragment_main() { + determinant_1bf6e7(); +} + +@compute @workgroup_size(1) +fn compute_main() { + determinant_1bf6e7(); +} diff --git a/test/tint/builtins/gen/var/determinant/c8251d.wgsl b/test/tint/builtins/gen/var/determinant/c8251d.wgsl new file mode 100644 index 0000000000..63743dd59f --- /dev/null +++ b/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 { + determinant_c8251d(); + return vec4(); +} + +@fragment +fn fragment_main() { + determinant_c8251d(); +} + +@compute @workgroup_size(1) +fn compute_main() { + determinant_c8251d(); +} diff --git a/test/tint/builtins/gen/var/determinant/cefdf3.wgsl b/test/tint/builtins/gen/var/determinant/cefdf3.wgsl new file mode 100644 index 0000000000..3a38f3dcd1 --- /dev/null +++ b/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 { + determinant_cefdf3(); + return vec4(); +} + +@fragment +fn fragment_main() { + determinant_cefdf3(); +} + +@compute @workgroup_size(1) +fn compute_main() { + determinant_cefdf3(); +}