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>
diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index 6c17ec6..57eb10d 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -442,7 +442,7 @@
@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
diff --git a/src/tint/resolver/builtin_test.cc b/src/tint/resolver/builtin_test.cc
index 07ade34..44f4bc5 100644
--- a/src/tint/resolver/builtin_test.cc
+++ b/src/tint/resolver/builtin_test.cc
@@ -1932,7 +1932,7 @@
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 @@
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
)");
}
diff --git a/src/tint/resolver/const_eval.cc b/src/tint/resolver/const_eval.cc
index b0595f2..2019f59 100644
--- a/src/tint/resolver/const_eval.cc
+++ b/src/tint/resolver/const_eval.cc
@@ -893,15 +893,22 @@
template <typename NumberT>
utils::Result<NumberT> 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;
}
@@ -913,6 +920,129 @@
}
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)) {
AddError("sqrt must be called with a value >= 0", source);
@@ -1044,6 +1174,26 @@
};
}
+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 @@
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) {
diff --git a/src/tint/resolver/const_eval.h b/src/tint/resolver/const_eval.h
index 6040729..fd253b4 100644
--- a/src/tint/resolver/const_eval.h
+++ b/src/tint/resolver/const_eval.h
@@ -539,6 +539,7 @@
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 @@
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 @@
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 <typename NumberT>
- utils::Result<NumberT> Det2(const Source& source,
- NumberT a1,
- NumberT a2,
- NumberT b1,
- NumberT b2);
+ utils::Result<NumberT> 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 <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 @@
/// @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 c288077..1bf7423 100644
--- a/src/tint/resolver/const_eval_builtin_test.cc
+++ b/src/tint/resolver/const_eval_builtin_test.cc
@@ -854,6 +854,83 @@
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>;
auto r = std::vector<Case>{
diff --git a/src/tint/resolver/const_eval_test.h b/src/tint/resolver/const_eval_test.h
index 8975752..148531d 100644
--- a/src/tint/resolver/const_eval_test.h
+++ b/src/tint/resolver/const_eval_test.h
@@ -288,6 +288,16 @@
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>>,
diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl
index ee79ab1..aff12e4 100644
--- a/src/tint/resolver/intrinsic_table.inl
+++ b/src/tint/resolver/intrinsic_table.inl
@@ -13566,12 +13566,12 @@
/* 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 @@
},
{
/* [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],
},
diff --git a/test/tint/bug/tint/1061.spvasm.expected.dxc.hlsl b/test/tint/bug/tint/1061.spvasm.expected.dxc.hlsl
index e8ae5f0..319b7c3 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 @@
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 e8ae5f0..319b7c3 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 @@
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 c851858..0a12089 100644
--- a/test/tint/bug/tint/1061.spvasm.expected.glsl
+++ b/test/tint/bug/tint/1061.spvasm.expected.glsl
@@ -14,7 +14,7 @@
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 f0697f8..2f9478f 100644
--- a/test/tint/bug/tint/1061.spvasm.expected.msl
+++ b/test/tint/bug/tint/1061.spvasm.expected.msl
@@ -8,7 +8,7 @@
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 946f03c..78fc8aa 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 0000000..e005418
--- /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<f32> {
+ determinant_1bf6e7();
+ return vec4<f32>();
+}
+
+@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 0000000..3202766
--- /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 0000000..3202766
--- /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 0000000..d5b27a9
--- /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 0000000..3e3494c
--- /dev/null
+++ b/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;
+}
+
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 0000000..b5f4f68
--- /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 0000000..0aaf417
--- /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<f32> {
+ determinant_1bf6e7();
+ return vec4<f32>();
+}
+
+@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 c9482d7..30e80d3 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 c9482d7..30e80d3 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 d801fb8..41fce04 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 @@
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 @@
#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 725abfc..56313b0 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 9673dd3..c7425e8 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 bae350d..e602aa2 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, 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 {
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 f94a854..2a0aa4f 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 @@
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 @@
#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 f416fbf..c50561d 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 ffe5e4d..d2fb249 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 0857fc6..78a081c 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 0857fc6..78a081c 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 3d93e51..de14daf 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 @@
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 @@
#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 9ee1b47..bd6b0a6 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 57d1c71..4d1af3a 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 0000000..d38a575
--- /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<f32> {
+ determinant_c8251d();
+ return vec4<f32>();
+}
+
+@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 0000000..8415cb1
--- /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 0000000..8415cb1
--- /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 0000000..fb8d99d
--- /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 0000000..7132d6f
--- /dev/null
+++ b/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;
+}
+
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 0000000..23a1e6d
--- /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 0000000..ce0cf9c
--- /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<f32> {
+ determinant_c8251d();
+ return vec4<f32>();
+}
+
+@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 0000000..8563a81
--- /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<f32> {
+ determinant_cefdf3();
+ return vec4<f32>();
+}
+
+@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 0000000..ff28608
--- /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 0000000..ff28608
--- /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 0000000..488869b
--- /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 0000000..f687f25
--- /dev/null
+++ b/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;
+}
+
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 0000000..39854c3
--- /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 0000000..f3b7dd7
--- /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<f32> {
+ determinant_cefdf3();
+ return vec4<f32>();
+}
+
+@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 083d524..ecaeb94 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, 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 {
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 064604c..2c159929 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 @@
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 @@
#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 17c01fb..2855ff5 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 c4cb0a3..1324352 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 0c372a3..de9399c 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 0c372a3..de9399c 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 cccb7b6..0a24701 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 @@
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 @@
#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 47b8bed8..c592c00 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 441f89e..4a86d631 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 698a14c..b2730ff 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, 2, 2>((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 b647997..eb4ac54 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 @@
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 @@
#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 a639c13..51ce8b3 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 2854245..9a8aa11 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 0000000..a10c782
--- /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<f32> {
+ determinant_1bf6e7();
+ return vec4<f32>();
+}
+
+@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 0000000..63743dd
--- /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<f32> {
+ determinant_c8251d();
+ return vec4<f32>();
+}
+
+@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 0000000..3a38f3d
--- /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<f32> {
+ determinant_cefdf3();
+ return vec4<f32>();
+}
+
+@fragment
+fn fragment_main() {
+ determinant_cefdf3();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ determinant_cefdf3();
+}
