tint: Implement const-eval of modf
Bug: tint:1581
Change-Id: I53151ebf43601cd6afcdd2ec91d0ff9c4e650ef3
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/111241
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index 4bb101e..da58aea 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -504,8 +504,8 @@
fn mix<T: f32_f16>(T, T, T) -> T
fn mix<N: num, T: f32_f16>(vec<N, T>, vec<N, T>, vec<N, T>) -> vec<N, T>
fn mix<N: num, T: f32_f16>(vec<N, T>, vec<N, T>, T) -> vec<N, T>
-fn modf<T: f32_f16>(T) -> __modf_result<T>
-fn modf<N: num, T: f32_f16>(vec<N, T>) -> __modf_result_vec<N, T>
+@const fn modf<T: f32_f16>(@test_value(-1.5) T) -> __modf_result<T>
+@const fn modf<N: num, T: f32_f16>(@test_value(-1.5) vec<N, T>) -> __modf_result_vec<N, T>
fn normalize<N: num, T: f32_f16>(vec<N, T>) -> vec<N, T>
@const fn pack2x16float(vec2<f32>) -> u32
@const fn pack2x16snorm(vec2<f32>) -> u32
diff --git a/src/tint/resolver/const_eval.cc b/src/tint/resolver/const_eval.cc
index 6cdf8f4..c5b51c0 100644
--- a/src/tint/resolver/const_eval.cc
+++ b/src/tint/resolver/const_eval.cc
@@ -2266,6 +2266,40 @@
return TransformElements(builder, ty, transform, args[0], args[1]);
}
+ConstEval::Result ConstEval::modf(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source& source) {
+ auto transform_fract = [&](const sem::Constant* c) {
+ auto create = [&](auto e) {
+ return CreateElement(builder, source, c->Type(),
+ decltype(e)(e.value - std::trunc(e.value)));
+ };
+ return Dispatch_fa_f32_f16(create, c);
+ };
+ auto transform_whole = [&](const sem::Constant* c) {
+ auto create = [&](auto e) {
+ return CreateElement(builder, source, c->Type(), decltype(e)(std::trunc(e.value)));
+ };
+ return Dispatch_fa_f32_f16(create, c);
+ };
+
+ utils::Vector<const sem::Constant*, 2> fields;
+
+ if (auto fract = TransformElements(builder, args[0]->Type(), transform_fract, args[0])) {
+ fields.Push(fract.Get());
+ } else {
+ return utils::Failure;
+ }
+
+ if (auto whole = TransformElements(builder, args[0]->Type(), transform_whole, args[0])) {
+ fields.Push(whole.Get());
+ } else {
+ return utils::Failure;
+ }
+
+ return CreateComposite(builder, ty, std::move(fields));
+}
+
ConstEval::Result ConstEval::pack2x16float(const sem::Type* ty,
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 496e138..f290adb 100644
--- a/src/tint/resolver/const_eval.h
+++ b/src/tint/resolver/const_eval.h
@@ -628,6 +628,15 @@
utils::VectorRef<const sem::Constant*> args,
const Source& source);
+ /// modf builtin
+ /// @param ty the expression type
+ /// @param args the input arguments
+ /// @param source the source location
+ /// @return the result value, or null if the value cannot be calculated
+ Result modf(const sem::Type* ty,
+ utils::VectorRef<const sem::Constant*> args,
+ const Source& source);
+
/// pack2x16float builtin
/// @param ty the expression type
/// @param args the input arguments
diff --git a/src/tint/resolver/const_eval_builtin_test.cc b/src/tint/resolver/const_eval_builtin_test.cc
index 13bbac5..eccd711 100644
--- a/src/tint/resolver/const_eval_builtin_test.cc
+++ b/src/tint/resolver/const_eval_builtin_test.cc
@@ -98,6 +98,11 @@
return Case{utils::Vector<Types, 8>{args}, utils::Vector<Types, 2>{std::move(result)}};
}
+/// Creates a Case with Values for args and result
+static Case C(std::initializer_list<Types> args, std::initializer_list<Types> results) {
+ return Case{utils::Vector<Types, 8>{args}, utils::Vector<Types, 2>{results}};
+}
+
/// Convenience overload that creates a Case with just scalars
static Case C(std::initializer_list<ScalarTypes> sargs, ScalarTypes sresult) {
utils::Vector<Types, 8> args;
@@ -109,6 +114,20 @@
return Case{std::move(args), utils::Vector<Types, 2>{std::move(result)}};
}
+/// Creates a Case with Values for args and result
+static Case C(std::initializer_list<ScalarTypes> sargs,
+ std::initializer_list<ScalarTypes> sresults) {
+ utils::Vector<Types, 8> args;
+ for (auto& sa : sargs) {
+ std::visit([&](auto&& v) { return args.Push(Val(v)); }, sa);
+ }
+ utils::Vector<Types, 2> results;
+ for (auto& sa : sresults) {
+ std::visit([&](auto&& v) { return results.Push(Val(v)); }, sa);
+ }
+ return Case{std::move(args), std::move(results)};
+}
+
/// Creates a Case with Values for args and expected error
static Case E(std::initializer_list<Types> args, std::string err) {
return Case{utils::Vector<Types, 8>{args}, std::move(err)};
@@ -1290,6 +1309,38 @@
MinCases<AFloat>(),
MinCases<f32>(),
MinCases<f16>()))));
+template <typename T>
+std::vector<Case> ModfCases() {
+ return {
+ // Scalar tests
+ // in fract whole
+ C({T(0.0)}, {T(0.0), T(0.0)}), //
+ C({T(1.0)}, {T(0.0), T(1.0)}), //
+ C({T(2.0)}, {T(0.0), T(2.0)}), //
+ C({T(1.5)}, {T(0.5), T(1.0)}), //
+ C({T(4.25)}, {T(0.25), T(4.0)}), //
+ C({T(-1.0)}, {T(0.0), T(-1.0)}), //
+ C({T(-2.0)}, {T(0.0), T(-2.0)}), //
+ C({T(-1.5)}, {T(-0.5), T(-1.0)}), //
+ C({T(-4.25)}, {T(-0.25), T(-4.0)}), //
+ C({T::Lowest()}, {T(0.0), T::Lowest()}), //
+ C({T::Highest()}, {T(0.0), T::Highest()}), //
+
+ // Vector tests
+ // in fract whole
+ C({Vec(T(0.0), T(0.0))}, {Vec(T(0.0), T(0.0)), Vec(T(0.0), T(0.0))}),
+ C({Vec(T(1.0), T(2.0))}, {Vec(T(0.0), T(0.0)), Vec(T(1), T(2))}),
+ C({Vec(T(-2.0), T(1.0))}, {Vec(T(0.0), T(0.0)), Vec(T(-2), T(1))}),
+ C({Vec(T(1.5), T(-2.25))}, {Vec(T(0.5), T(-0.25)), Vec(T(1.0), T(-2.0))}),
+ C({Vec(T::Lowest(), T::Highest())}, {Vec(T(0.0), T(0.0)), Vec(T::Lowest(), T::Highest())}),
+ };
+}
+INSTANTIATE_TEST_SUITE_P( //
+ Modf,
+ ResolverConstEvalBuiltinTest,
+ testing::Combine(testing::Values(sem::BuiltinType::kModf),
+ testing::ValuesIn(Concat(ModfCases<f32>(), //
+ ModfCases<f16>()))));
std::vector<Case> Pack4x8snormCases() {
return {
diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl
index 7744960..be0e4bb 100644
--- a/src/tint/resolver/intrinsic_table.inl
+++ b/src/tint/resolver/intrinsic_table.inl
@@ -12851,7 +12851,7 @@
/* parameters */ &kParameters[878],
/* return matcher indices */ &kMatcherIndices[106],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* const eval */ nullptr,
+ /* const eval */ &ConstEval::modf,
},
{
/* [378] */
@@ -12863,7 +12863,7 @@
/* parameters */ &kParameters[879],
/* return matcher indices */ &kMatcherIndices[45],
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* const eval */ nullptr,
+ /* const eval */ &ConstEval::modf,
},
{
/* [379] */
@@ -14351,8 +14351,8 @@
},
{
/* [53] */
- /* fn modf<T : f32_f16>(T) -> __modf_result<T> */
- /* fn modf<N : num, T : f32_f16>(vec<N, T>) -> __modf_result_vec<N, T> */
+ /* fn modf<T : f32_f16>(@test_value(-1.5) T) -> __modf_result<T> */
+ /* fn modf<N : num, T : f32_f16>(@test_value(-1.5) vec<N, T>) -> __modf_result_vec<N, T> */
/* num overloads */ 2,
/* overloads */ &kOverloads[377],
},
diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc
index 0b20c9b..f7d6cb3 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -530,9 +530,13 @@
}
if (auto* str = var->Type()->UnwrapRef()->As<sem::Struct>()) {
- if (!AddressSpaceLayout(str, var->AddressSpace(), str->Declaration()->source, layouts)) {
- AddNote("see declaration of variable", var->Declaration()->source);
- return false;
+ // Check the structure has a declaration. Builtins like modf() and frexp() return untypeable
+ // structures, and so they have no declaration. Just skip validation for these.
+ if (auto* str_decl = str->Declaration()) {
+ if (!AddressSpaceLayout(str, var->AddressSpace(), str_decl->source, layouts)) {
+ AddNote("see declaration of variable", var->Declaration()->source);
+ return false;
+ }
}
} else {
Source source = var->Declaration()->source;
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index e716b1d..86ba657 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -939,9 +939,7 @@
return true;
}
case sem::BuiltinType::kAtomicCompareExchangeWeak: {
- // Emit the builtin return type unique to this overload. This does not
- // exist in the AST, so it will not be generated in Generate().
- if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
+ if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
return false;
}
@@ -2373,10 +2371,12 @@
return true;
},
[&](const sem::Struct* s) {
- if (!EmitType(out, s, ast::AddressSpace::kNone, ast::Access::kUndefined, "")) {
+ if (!EmitStructType(&helpers_, s)) {
return false;
}
+ out << StructName(s);
+
ScopedParen sp(out);
for (size_t i = 0; i < s->Members().size(); i++) {
@@ -2998,6 +2998,11 @@
}
bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
+ auto it = emitted_structs_.emplace(str);
+ if (!it.second) {
+ return true;
+ }
+
auto address_space_uses = str->AddressSpaceUsage();
line(b) << "struct " << StructName(str) << " {";
EmitStructMembers(b, str);
@@ -3007,14 +3012,6 @@
return true;
}
-bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) {
- auto it = emitted_structs_.emplace(str);
- if (!it.second) {
- return true;
- }
- return EmitStructType(buffer, str);
-}
-
bool GeneratorImpl::EmitStructMembers(TextBuffer* b, const sem::Struct* str) {
ScopedIndent si(b);
for (auto* mem : str->Members()) {
diff --git a/src/tint/writer/glsl/generator_impl.h b/src/tint/writer/glsl/generator_impl.h
index c34880d..24d8357 100644
--- a/src/tint/writer/glsl/generator_impl.h
+++ b/src/tint/writer/glsl/generator_impl.h
@@ -430,17 +430,12 @@
ast::AddressSpace address_space,
ast::Access access,
const std::string& name);
- /// Handles generating a structure declaration
+ /// Handles generating a structure declaration. If the structure has already been emitted, then
+ /// this function will simply return `true` without emitting anything.
/// @param buffer the text buffer that the type declaration will be written to
/// @param ty the struct to generate
/// @returns true if the struct is emitted
bool EmitStructType(TextBuffer* buffer, const sem::Struct* ty);
- /// Handles generating a structure declaration only the first time called. Subsequent calls are
- /// a no-op and return true.
- /// @param buffer the text buffer that the type declaration will be written to
- /// @param ty the struct to generate
- /// @returns true if the struct is emitted
- bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty);
/// Handles generating the members of a structure
/// @param buffer the text buffer that the struct members will be written to
/// @param ty the struct to generate
diff --git a/src/tint/writer/glsl/generator_impl_builtin_test.cc b/src/tint/writer/glsl/generator_impl_builtin_test.cc
index 17028ac..7d0a72e 100644
--- a/src/tint/writer/glsl/generator_impl_builtin_test.cc
+++ b/src/tint/writer/glsl/generator_impl_builtin_test.cc
@@ -416,9 +416,9 @@
EXPECT_EQ(out.str(), "((a) * (b) + (c))");
}
-TEST_F(GlslGeneratorImplTest_Builtin, Modf_Scalar_f32) {
- auto* call = Call("modf", 1_f);
- WrapInFunction(CallStmt(call));
+TEST_F(GlslGeneratorImplTest_Builtin, Runtime_Modf_Scalar_f32) {
+ WrapInFunction(Decl(Let("f", Expr(1.5_f))), //
+ Decl(Let("v", Call("modf", "f"))));
GeneratorImpl& gen = SanitizeAndBuild();
@@ -438,7 +438,8 @@
void test_function() {
- tint_modf(1.0f);
+ float f = 1.5f;
+ modf_result v = tint_modf(f);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -449,11 +450,11 @@
)");
}
-TEST_F(GlslGeneratorImplTest_Builtin, Modf_Scalar_f16) {
+TEST_F(GlslGeneratorImplTest_Builtin, Runtime_Modf_Scalar_f16) {
Enable(ast::Extension::kF16);
- auto* call = Call("modf", 1_h);
- WrapInFunction(CallStmt(call));
+ WrapInFunction(Decl(Let("f", Expr(1.5_h))), //
+ Decl(Let("v", Call("modf", "f"))));
GeneratorImpl& gen = SanitizeAndBuild();
@@ -474,7 +475,8 @@
void test_function() {
- tint_modf(1.0hf);
+ float16_t f = 1.5hf;
+ modf_result_f16 v = tint_modf(f);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -485,9 +487,9 @@
)");
}
-TEST_F(GlslGeneratorImplTest_Builtin, Modf_Vector_f32) {
- auto* call = Call("modf", vec3<f32>());
- WrapInFunction(CallStmt(call));
+TEST_F(GlslGeneratorImplTest_Builtin, Runtime_Modf_Vector_f32) {
+ WrapInFunction(Decl(Let("f", vec3<f32>(1.5_f, 2.5_f, 3.5_f))), //
+ Decl(Let("v", Call("modf", "f"))));
GeneratorImpl& gen = SanitizeAndBuild();
@@ -507,7 +509,8 @@
void test_function() {
- tint_modf(vec3(0.0f));
+ vec3 f = vec3(1.5f, 2.5f, 3.5f);
+ modf_result_vec3 v = tint_modf(f);
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -518,11 +521,11 @@
)");
}
-TEST_F(GlslGeneratorImplTest_Builtin, Modf_Vector_f16) {
+TEST_F(GlslGeneratorImplTest_Builtin, Runtime_Modf_Vector_f16) {
Enable(ast::Extension::kF16);
- auto* call = Call("modf", vec3<f16>());
- WrapInFunction(CallStmt(call));
+ WrapInFunction(Decl(Let("f", vec3<f16>(1.5_h, 2.5_h, 3.5_h))), //
+ Decl(Let("v", Call("modf", "f"))));
GeneratorImpl& gen = SanitizeAndBuild();
@@ -543,7 +546,118 @@
void test_function() {
- tint_modf(f16vec3(0.0hf));
+ f16vec3 f = f16vec3(1.5hf, 2.5hf, 3.5hf);
+ modf_result_vec3_f16 v = tint_modf(f);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ test_function();
+ return;
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Builtin, Const_Modf_Scalar_f32) {
+ WrapInFunction(Decl(Let("v", Call("modf", 1.5_f))));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#version 310 es
+
+struct modf_result {
+ float fract;
+ float whole;
+};
+
+
+void test_function() {
+ modf_result v = modf_result(0.5f, 1.0f);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ test_function();
+ return;
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Builtin, Const_Modf_Scalar_f16) {
+ Enable(ast::Extension::kF16);
+
+ WrapInFunction(Decl(Let("v", Call("modf", 1.5_h))));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+struct modf_result_f16 {
+ float16_t fract;
+ float16_t whole;
+};
+
+
+void test_function() {
+ modf_result_f16 v = modf_result_f16(0.5hf, 1.0hf);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ test_function();
+ return;
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Builtin, Const_Modf_Vector_f32) {
+ WrapInFunction(Decl(Let("v", Call("modf", vec3<f32>(1.5_f, 2.5_f, 3.5_f)))));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#version 310 es
+
+struct modf_result_vec3 {
+ vec3 fract;
+ vec3 whole;
+};
+
+
+void test_function() {
+ modf_result_vec3 v = modf_result_vec3(vec3(0.5f), vec3(1.0f, 2.0f, 3.0f));
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ test_function();
+ return;
+}
+)");
+}
+
+TEST_F(GlslGeneratorImplTest_Builtin, Const_Modf_Vector_f16) {
+ Enable(ast::Extension::kF16);
+
+ WrapInFunction(Decl(Let("v", Call("modf", vec3<f16>(1.5_h, 2.5_h, 3.5_h)))));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#version 310 es
+#extension GL_AMD_gpu_shader_half_float : require
+
+struct modf_result_vec3_f16 {
+ f16vec3 fract;
+ f16vec3 whole;
+};
+
+
+void test_function() {
+ modf_result_vec3_f16 v = modf_result_vec3_f16(f16vec3(0.5hf), f16vec3(1.0hf, 2.0hf, 3.0hf));
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index e3abc51..dcad1b5 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -1651,9 +1651,7 @@
return true;
}
case sem::BuiltinType::kAtomicCompareExchangeWeak: {
- // Emit the builtin return type unique to this overload. This does not
- // exist in the AST, so it will not be generated in Generate().
- if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
+ if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
return false;
}
@@ -2506,7 +2504,7 @@
out << "default";
} else {
out << "case ";
- if (!EmitConstant(out, selector->Value())) {
+ if (!EmitConstant(out, selector->Value(), /* is_variable_initializer */ false)) {
return false;
}
}
@@ -2552,7 +2550,13 @@
bool GeneratorImpl::EmitExpression(std::ostream& out, const ast::Expression* expr) {
if (auto* sem = builder_.Sem().Get(expr)) {
if (auto* constant = sem->ConstantValue()) {
- return EmitConstant(out, constant);
+ bool is_variable_initializer = false;
+ if (auto* stmt = sem->Stmt()) {
+ if (auto* decl = As<ast::VariableDeclStatement>(stmt->Declaration())) {
+ is_variable_initializer = decl->variable->initializer == expr;
+ }
+ }
+ return EmitConstant(out, constant, is_variable_initializer);
}
}
return Switch(
@@ -3042,7 +3046,9 @@
return true;
}
-bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constant) {
+bool GeneratorImpl::EmitConstant(std::ostream& out,
+ const sem::Constant* constant,
+ bool is_variable_initializer) {
return Switch(
constant->Type(), //
[&](const sem::Bool*) {
@@ -3072,7 +3078,7 @@
if (constant->AllEqual()) {
{
ScopedParen sp(out);
- if (!EmitConstant(out, constant->Index(0))) {
+ if (!EmitConstant(out, constant->Index(0), is_variable_initializer)) {
return false;
}
}
@@ -3093,7 +3099,7 @@
if (i > 0) {
out << ", ";
}
- if (!EmitConstant(out, constant->Index(i))) {
+ if (!EmitConstant(out, constant->Index(i), is_variable_initializer)) {
return false;
}
}
@@ -3110,7 +3116,7 @@
if (i > 0) {
out << ", ";
}
- if (!EmitConstant(out, constant->Index(i))) {
+ if (!EmitConstant(out, constant->Index(i), is_variable_initializer)) {
return false;
}
}
@@ -3139,7 +3145,7 @@
if (i > 0) {
out << ", ";
}
- if (!EmitConstant(out, constant->Index(i))) {
+ if (!EmitConstant(out, constant->Index(i), is_variable_initializer)) {
return false;
}
}
@@ -3147,25 +3153,45 @@
return true;
},
[&](const sem::Struct* s) {
+ if (!EmitStructType(&helpers_, s)) {
+ return false;
+ }
+
if (constant->AllZero()) {
- out << "(";
- if (!EmitType(out, s, ast::AddressSpace::kNone, ast::Access::kUndefined, "")) {
- return false;
- }
- out << ")0";
+ out << "(" << StructName(s) << ")0";
return true;
}
- out << "{";
- TINT_DEFER(out << "}");
-
- for (size_t i = 0; i < s->Members().size(); i++) {
- if (i > 0) {
- out << ", ";
+ auto emit_member_values = [&](std::ostream& o) {
+ o << "{";
+ for (size_t i = 0; i < s->Members().size(); i++) {
+ if (i > 0) {
+ o << ", ";
+ }
+ if (!EmitConstant(o, constant->Index(i), is_variable_initializer)) {
+ return false;
+ }
}
- if (!EmitConstant(out, constant->Index(i))) {
+ o << "}";
+ return true;
+ };
+
+ if (is_variable_initializer) {
+ if (!emit_member_values(out)) {
return false;
}
+ } else {
+ // HLSL requires structure initializers to be assigned directly to a variable.
+ auto name = UniqueIdentifier("c");
+ {
+ auto decl = line();
+ decl << "const " << StructName(s) << " " << name << " = ";
+ if (!emit_member_values(decl)) {
+ return false;
+ }
+ decl << ";";
+ }
+ out << name;
}
return true;
@@ -3891,6 +3917,11 @@
}
bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
+ auto it = emitted_structs_.emplace(str);
+ if (!it.second) {
+ return true;
+ }
+
line(b) << "struct " << StructName(str) << " {";
{
ScopedIndent si(b);
@@ -3967,14 +3998,6 @@
return true;
}
-bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) {
- auto it = emitted_structs_.emplace(str);
- if (!it.second) {
- return true;
- }
- return EmitStructType(buffer, str);
-}
-
bool GeneratorImpl::EmitUnaryOp(std::ostream& out, const ast::UnaryOpExpression* expr) {
switch (expr->op) {
case ast::UnaryOp::kIndirection:
diff --git a/src/tint/writer/hlsl/generator_impl.h b/src/tint/writer/hlsl/generator_impl.h
index 16b824b..74a172d 100644
--- a/src/tint/writer/hlsl/generator_impl.h
+++ b/src/tint/writer/hlsl/generator_impl.h
@@ -348,8 +348,12 @@
/// Handles a constant value
/// @param out the output stream
/// @param constant the constant value to emit
+ /// @param is_variable_initializer true if the constant is used as the RHS of a variable
+ /// initializer
/// @returns true if the constant value was successfully emitted
- bool EmitConstant(std::ostream& out, const sem::Constant* constant);
+ bool EmitConstant(std::ostream& out,
+ const sem::Constant* constant,
+ bool is_variable_initializer);
/// Handles a literal
/// @param out the output stream
/// @param lit the literal to emit
@@ -420,17 +424,12 @@
ast::AddressSpace address_space,
ast::Access access,
const std::string& name);
- /// Handles generating a structure declaration
+ /// Handles generating a structure declaration. If the structure has already been emitted, then
+ /// this function will simply return `true` without emitting anything.
/// @param buffer the text buffer that the type declaration will be written to
/// @param ty the struct to generate
/// @returns true if the struct is emitted
bool EmitStructType(TextBuffer* buffer, const sem::Struct* ty);
- /// Handles generating a structure declaration only the first time called. Subsequent calls are
- /// a no-op and return true.
- /// @param buffer the text buffer that the type declaration will be written to
- /// @param ty the struct to generate
- /// @returns true if the struct is emitted
- bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty);
/// Handles a unary op expression
/// @param out the output of the expression stream
/// @param expr the expression to emit
diff --git a/src/tint/writer/hlsl/generator_impl_builtin_test.cc b/src/tint/writer/hlsl/generator_impl_builtin_test.cc
index b69e671..3e9b169 100644
--- a/src/tint/writer/hlsl/generator_impl_builtin_test.cc
+++ b/src/tint/writer/hlsl/generator_impl_builtin_test.cc
@@ -381,9 +381,9 @@
EXPECT_EQ(out.str(), "(bool2(true, false) ? b : a)");
}
-TEST_F(HlslGeneratorImplTest_Builtin, Modf_Scalar_f32) {
- auto* call = Call("modf", 1_f);
- WrapInFunction(CallStmt(call));
+TEST_F(HlslGeneratorImplTest_Builtin, Runtime_Modf_Scalar_f32) {
+ WrapInFunction(Decl(Let("f", Expr(1.5_f))), //
+ Decl(Let("v", Call("modf", "f"))));
GeneratorImpl& gen = SanitizeAndBuild();
@@ -400,17 +400,18 @@
[numthreads(1, 1, 1)]
void test_function() {
- tint_modf(1.0f);
+ const float f = 1.5f;
+ const modf_result v = tint_modf(f);
return;
}
)");
}
-TEST_F(HlslGeneratorImplTest_Builtin, Modf_Scalar_f16) {
+TEST_F(HlslGeneratorImplTest_Builtin, Runtime_Modf_Scalar_f16) {
Enable(ast::Extension::kF16);
- auto* call = Call("modf", 1_h);
- WrapInFunction(CallStmt(call));
+ WrapInFunction(Decl(Let("f", Expr(1.5_h))), //
+ Decl(Let("v", Call("modf", "f"))));
GeneratorImpl& gen = SanitizeAndBuild();
@@ -427,15 +428,16 @@
[numthreads(1, 1, 1)]
void test_function() {
- tint_modf(float16_t(1.0h));
+ const float16_t f = float16_t(1.5h);
+ const modf_result_f16 v = tint_modf(f);
return;
}
)");
}
-TEST_F(HlslGeneratorImplTest_Builtin, Modf_Vector_f32) {
- auto* call = Call("modf", vec3<f32>());
- WrapInFunction(CallStmt(call));
+TEST_F(HlslGeneratorImplTest_Builtin, Runtime_Modf_Vector_f32) {
+ WrapInFunction(Decl(Let("f", vec3<f32>(1.5_f, 2.5_f, 3.5_f))), //
+ Decl(Let("v", Call("modf", "f"))));
GeneratorImpl& gen = SanitizeAndBuild();
@@ -452,17 +454,18 @@
[numthreads(1, 1, 1)]
void test_function() {
- tint_modf((0.0f).xxx);
+ const float3 f = float3(1.5f, 2.5f, 3.5f);
+ const modf_result_vec3 v = tint_modf(f);
return;
}
)");
}
-TEST_F(HlslGeneratorImplTest_Builtin, Modf_Vector_f16) {
+TEST_F(HlslGeneratorImplTest_Builtin, Runtime_Modf_Vector_f16) {
Enable(ast::Extension::kF16);
- auto* call = Call("modf", vec3<f16>());
- WrapInFunction(CallStmt(call));
+ WrapInFunction(Decl(Let("f", vec3<f16>(1.5_h, 2.5_h, 3.5_h))), //
+ Decl(Let("v", Call("modf", "f"))));
GeneratorImpl& gen = SanitizeAndBuild();
@@ -479,7 +482,110 @@
[numthreads(1, 1, 1)]
void test_function() {
- tint_modf((float16_t(0.0h)).xxx);
+ const vector<float16_t, 3> f = vector<float16_t, 3>(float16_t(1.5h), float16_t(2.5h), float16_t(3.5h));
+ const modf_result_vec3_f16 v = tint_modf(f);
+ return;
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Builtin, Const_Modf_Scalar_f32) {
+ WrapInFunction(Decl(Let("v", Call("modf", 1.5_f))));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct modf_result {
+ float fract;
+ float whole;
+};
+[numthreads(1, 1, 1)]
+void test_function() {
+ const modf_result v = {0.5f, 1.0f};
+ return;
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Builtin, Const_Modf_Scalar_f16) {
+ Enable(ast::Extension::kF16);
+
+ WrapInFunction(Decl(Let("v", Call("modf", 1.5_h))));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct modf_result_f16 {
+ float16_t fract;
+ float16_t whole;
+};
+[numthreads(1, 1, 1)]
+void test_function() {
+ const modf_result_f16 v = {float16_t(0.5h), float16_t(1.0h)};
+ return;
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Builtin, Const_Modf_Vector_f32) {
+ WrapInFunction(Decl(Let("v", Call("modf", vec3<f32>(1.5_f, 2.5_f, 3.5_f)))));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct modf_result_vec3 {
+ float3 fract;
+ float3 whole;
+};
+[numthreads(1, 1, 1)]
+void test_function() {
+ const modf_result_vec3 v = {(0.5f).xxx, float3(1.0f, 2.0f, 3.0f)};
+ return;
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Builtin, Const_Modf_Vector_f16) {
+ Enable(ast::Extension::kF16);
+
+ WrapInFunction(Decl(Let("v", Call("modf", vec3<f16>(1.5_h, 2.5_h, 3.5_h)))));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct modf_result_vec3_f16 {
+ vector<float16_t, 3> fract;
+ vector<float16_t, 3> whole;
+};
+[numthreads(1, 1, 1)]
+void test_function() {
+ const modf_result_vec3_f16 v = {(float16_t(0.5h)).xxx, vector<float16_t, 3>(float16_t(1.0h), float16_t(2.0h), float16_t(3.0h))};
+ return;
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Builtin, NonInitializer_Modf_Vector_f32) {
+ WrapInFunction(
+ // Declare a variable with the result of a modf call.
+ // This is required to infer the 'var' type.
+ Decl(Var("v", Call("modf", vec3<f32>(1.5_f, 2.5_f, 3.5_f)))),
+ // Now assign 'v' again with another modf call.
+ // This requires generating a temporary variable for the struct initializer.
+ Assign("v", Call("modf", vec3<f32>(4.5_f, 5.5_f, 6.5_f))));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(struct modf_result_vec3 {
+ float3 fract;
+ float3 whole;
+};
+[numthreads(1, 1, 1)]
+void test_function() {
+ modf_result_vec3 v = {(0.5f).xxx, float3(1.0f, 2.0f, 3.0f)};
+ const modf_result_vec3 c = {(0.5f).xxx, float3(4.0f, 5.0f, 6.0f)};
+ v = c;
return;
}
)");
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index 534742e..10efd59 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -907,9 +907,7 @@
auto func = utils::GetOrCreate(
atomicCompareExchangeWeak_, ACEWKeyType{{sc, str}}, [&]() -> std::string {
- // Emit the builtin return type unique to this overload. This does not
- // exist in the AST, so it will not be generated in Generate().
- if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
+ if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
return "";
}
@@ -1748,7 +1746,11 @@
return true;
},
[&](const sem::Struct* s) {
- out << program_->Symbols().NameFor(s->Name()) << "{";
+ if (!EmitStructType(&helpers_, s)) {
+ return false;
+ }
+
+ out << StructName(s) << "{";
TINT_DEFER(out << "}");
if (constant->AllZero()) {
@@ -2747,6 +2749,11 @@
}
bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
+ auto it = emitted_structs_.emplace(str);
+ if (!it.second) {
+ return true;
+ }
+
line(b) << "struct " << StructName(str) << " {";
bool is_host_shareable = str->IsHostShareable();
@@ -2903,14 +2910,6 @@
return true;
}
-bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) {
- auto it = emitted_structs_.emplace(str);
- if (!it.second) {
- return true;
- }
- return EmitStructType(buffer, str);
-}
-
bool GeneratorImpl::EmitUnaryOp(std::ostream& out, const ast::UnaryOpExpression* expr) {
// Handle `-e` when `e` is signed, so that we ensure that if `e` is the
// largest negative value, it returns `e`.
diff --git a/src/tint/writer/msl/generator_impl.h b/src/tint/writer/msl/generator_impl.h
index c3e1ac0..829088d 100644
--- a/src/tint/writer/msl/generator_impl.h
+++ b/src/tint/writer/msl/generator_impl.h
@@ -328,17 +328,12 @@
/// @param sc the address space to generate
/// @returns true if the address space is emitted
bool EmitAddressSpace(std::ostream& out, ast::AddressSpace sc);
- /// Handles generating a struct declaration
+ /// Handles generating a struct declaration. If the structure has already been emitted, then
+ /// this function will simply return `true` without emitting anything.
/// @param buffer the text buffer that the type declaration will be written to
/// @param str the struct to generate
/// @returns true if the struct is emitted
bool EmitStructType(TextBuffer* buffer, const sem::Struct* str);
- /// Handles generating a structure declaration only the first time called. Subsequent calls are
- /// a no-op and return true.
- /// @param buffer the text buffer that the type declaration will be written to
- /// @param ty the struct to generate
- /// @returns true if the struct is emitted
- bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty);
/// Handles a unary op expression
/// @param out the output of the expression stream
/// @param expr the expression to emit
diff --git a/src/tint/writer/msl/generator_impl_builtin_test.cc b/src/tint/writer/msl/generator_impl_builtin_test.cc
index 8974bd9..c611d6b 100644
--- a/src/tint/writer/msl/generator_impl_builtin_test.cc
+++ b/src/tint/writer/msl/generator_impl_builtin_test.cc
@@ -405,9 +405,9 @@
EXPECT_EQ(out.str(), "threadgroup_barrier(mem_flags::mem_threadgroup)");
}
-TEST_F(MslGeneratorImplTest, Modf_Scalar_f32) {
- auto* call = Call("modf", 1_f);
- WrapInFunction(CallStmt(call));
+TEST_F(MslGeneratorImplTest, Runtime_Modf_Scalar_f32) {
+ WrapInFunction(Decl(Let("f", Expr(1.5_f))), //
+ Decl(Let("v", Call("modf", "f"))));
GeneratorImpl& gen = SanitizeAndBuild();
@@ -427,18 +427,19 @@
}
kernel void test_function() {
- tint_modf(1.0f);
+ float const f = 1.5f;
+ modf_result const v = tint_modf(f);
return;
}
)");
}
-TEST_F(MslGeneratorImplTest, Modf_Scalar_f16) {
+TEST_F(MslGeneratorImplTest, Runtime_Modf_Scalar_f16) {
Enable(ast::Extension::kF16);
- auto* call = Call("modf", 1_h);
- WrapInFunction(CallStmt(call));
+ WrapInFunction(Decl(Let("f", Expr(1.5_h))), //
+ Decl(Let("v", Call("modf", "f"))));
GeneratorImpl& gen = SanitizeAndBuild();
@@ -458,16 +459,17 @@
}
kernel void test_function() {
- tint_modf(1.0h);
+ half const f = 1.5h;
+ modf_result_f16 const v = tint_modf(f);
return;
}
)");
}
-TEST_F(MslGeneratorImplTest, Modf_Vector_f32) {
- auto* call = Call("modf", vec3<f32>());
- WrapInFunction(CallStmt(call));
+TEST_F(MslGeneratorImplTest, Runtime_Modf_Vector_f32) {
+ WrapInFunction(Decl(Let("f", vec3<f32>(1.5_f, 2.5_f, 3.5_f))), //
+ Decl(Let("v", Call("modf", "f"))));
GeneratorImpl& gen = SanitizeAndBuild();
@@ -487,18 +489,19 @@
}
kernel void test_function() {
- tint_modf(float3(0.0f));
+ float3 const f = float3(1.5f, 2.5f, 3.5f);
+ modf_result_vec3 const v = tint_modf(f);
return;
}
)");
}
-TEST_F(MslGeneratorImplTest, Modf_Vector_f16) {
+TEST_F(MslGeneratorImplTest, Runtime_Modf_Vector_f16) {
Enable(ast::Extension::kF16);
- auto* call = Call("modf", vec3<f16>());
- WrapInFunction(CallStmt(call));
+ WrapInFunction(Decl(Let("f", vec3<f16>(1.5_h, 2.5_h, 3.5_h))), //
+ Decl(Let("v", Call("modf", "f"))));
GeneratorImpl& gen = SanitizeAndBuild();
@@ -518,7 +521,100 @@
}
kernel void test_function() {
- tint_modf(half3(0.0h));
+ half3 const f = half3(1.5h, 2.5h, 3.5h);
+ modf_result_vec3_f16 const v = tint_modf(f);
+ return;
+}
+
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Const_Modf_Scalar_f32) {
+ WrapInFunction(Decl(Let("v", Call("modf", 1.5_f))));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result {
+ float fract;
+ float whole;
+};
+kernel void test_function() {
+ modf_result const v = modf_result{.fract=0.5f, .whole=1.0f};
+ return;
+}
+
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Const_Modf_Scalar_f16) {
+ Enable(ast::Extension::kF16);
+
+ WrapInFunction(Decl(Let("v", Call("modf", 1.5_h))));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result_f16 {
+ half fract;
+ half whole;
+};
+kernel void test_function() {
+ modf_result_f16 const v = modf_result_f16{.fract=0.5h, .whole=1.0h};
+ return;
+}
+
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Const_Modf_Vector_f32) {
+ WrapInFunction(Decl(Let("v", Call("modf", vec3<f32>(1.5_f, 2.5_f, 3.5_f)))));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result_vec3 {
+ float3 fract;
+ float3 whole;
+};
+kernel void test_function() {
+ modf_result_vec3 const v = modf_result_vec3{.fract=float3(0.5f), .whole=float3(1.0f, 2.0f, 3.0f)};
+ return;
+}
+
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Const_Modf_Vector_f16) {
+ Enable(ast::Extension::kF16);
+
+ WrapInFunction(Decl(Let("v", Call("modf", vec3<f16>(1.5_h, 2.5_h, 3.5_h)))));
+
+ GeneratorImpl& gen = SanitizeAndBuild();
+
+ ASSERT_TRUE(gen.Generate()) << gen.error();
+ EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result_vec3_f16 {
+ half3 fract;
+ half3 whole;
+};
+kernel void test_function() {
+ modf_result_vec3_f16 const v = modf_result_vec3_f16{.fract=half3(0.5h), .whole=half3(1.0h, 2.0h, 3.0h)};
return;
}
diff --git a/src/tint/writer/spirv/builder_builtin_test.cc b/src/tint/writer/spirv/builder_builtin_test.cc
index d00b019..c0995e8 100644
--- a/src/tint/writer/spirv/builder_builtin_test.cc
+++ b/src/tint/writer/spirv/builder_builtin_test.cc
@@ -1630,7 +1630,7 @@
EXPECT_EQ(got, expect);
}
-TEST_F(BuiltinBuilderTest, Call_Modf_f32) {
+TEST_F(BuiltinBuilderTest, Runtime_Call_Modf_f32) {
auto* vec = Var("vec", vec2<f32>(1_f, 2_f));
auto* expr = Call("modf", vec);
Func("a_func", utils::Empty, ty.void_(),
@@ -1682,7 +1682,7 @@
Validate(b);
}
-TEST_F(BuiltinBuilderTest, Call_Modf_f16) {
+TEST_F(BuiltinBuilderTest, Runtime_Call_Modf_f16) {
Enable(ast::Extension::kF16);
auto* vec = Var("vec", vec2<f16>(1_h, 2_h));
@@ -1740,6 +1740,100 @@
Validate(b);
}
+TEST_F(BuiltinBuilderTest, Const_Call_Modf_f32) {
+ auto* expr = Call("modf", vec2<f32>(1_f, 2_f));
+ Func("a_func", utils::Empty, ty.void_(),
+ utils::Vector{
+ CallStmt(expr),
+ },
+ utils::Vector{
+ Stage(ast::PipelineStage::kFragment),
+ });
+
+ spirv::Builder& b = Build();
+
+ ASSERT_TRUE(b.Build()) << b.error();
+ auto got = DumpBuilder(b);
+ auto* expect = R"(OpCapability Shader
+%9 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint Fragment %3 "a_func"
+OpExecutionMode %3 OriginUpperLeft
+OpName %3 "a_func"
+OpName %6 "__modf_result_vec2"
+OpMemberName %6 0 "fract"
+OpMemberName %6 1 "whole"
+OpMemberDecorate %6 0 Offset 0
+OpMemberDecorate %6 1 Offset 8
+%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%8 = OpTypeFloat 32
+%7 = OpTypeVector %8 2
+%6 = OpTypeStruct %7 %7
+%10 = OpConstant %8 1
+%11 = OpConstant %8 2
+%12 = OpConstantComposite %7 %10 %11
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+%5 = OpExtInst %6 %9 ModfStruct %12
+OpReturn
+OpFunctionEnd
+)";
+ EXPECT_EQ(expect, got);
+
+ Validate(b);
+}
+
+TEST_F(BuiltinBuilderTest, Const_Call_Modf_f16) {
+ Enable(ast::Extension::kF16);
+
+ auto* expr = Call("modf", vec2<f16>(1_h, 2_h));
+ Func("a_func", utils::Empty, ty.void_(),
+ utils::Vector{
+ CallStmt(expr),
+ },
+ utils::Vector{
+ Stage(ast::PipelineStage::kFragment),
+ });
+
+ spirv::Builder& b = Build();
+
+ ASSERT_TRUE(b.Build()) << b.error();
+ auto got = DumpBuilder(b);
+ auto* expect = R"(OpCapability Shader
+OpCapability Float16
+OpCapability UniformAndStorageBuffer16BitAccess
+OpCapability StorageBuffer16BitAccess
+OpCapability StorageInputOutput16
+%9 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint Fragment %3 "a_func"
+OpExecutionMode %3 OriginUpperLeft
+OpName %3 "a_func"
+OpName %6 "__modf_result_vec2_f16"
+OpMemberName %6 0 "fract"
+OpMemberName %6 1 "whole"
+OpMemberDecorate %6 0 Offset 0
+OpMemberDecorate %6 1 Offset 4
+%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%8 = OpTypeFloat 16
+%7 = OpTypeVector %8 2
+%6 = OpTypeStruct %7 %7
+%10 = OpConstant %8 0x1p+0
+%11 = OpConstant %8 0x1p+1
+%12 = OpConstantComposite %7 %10 %11
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+%5 = OpExtInst %6 %9 ModfStruct %12
+OpReturn
+OpFunctionEnd
+)";
+ EXPECT_EQ(expect, got);
+
+ Validate(b);
+}
+
TEST_F(BuiltinBuilderTest, Call_Frexp_f32) {
auto* vec = Var("vec", vec2<f32>(1_f, 2_f));
auto* expr = Call("frexp", vec);
diff --git a/test/tint/bug/chromium/1236161.wgsl.expected.dxc.hlsl b/test/tint/bug/chromium/1236161.wgsl.expected.dxc.hlsl
index 1361053..f73a45a 100644
--- a/test/tint/bug/chromium/1236161.wgsl.expected.dxc.hlsl
+++ b/test/tint/bug/chromium/1236161.wgsl.expected.dxc.hlsl
@@ -1,18 +1,8 @@
-struct modf_result {
- float fract;
- float whole;
-};
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
[numthreads(1, 1, 1)]
void unused_entry_point() {
return;
}
void i() {
- const float s = tint_modf(1.0f).whole;
+ const float s = 1.0f;
}
diff --git a/test/tint/bug/chromium/1236161.wgsl.expected.fxc.hlsl b/test/tint/bug/chromium/1236161.wgsl.expected.fxc.hlsl
index 1361053..f73a45a 100644
--- a/test/tint/bug/chromium/1236161.wgsl.expected.fxc.hlsl
+++ b/test/tint/bug/chromium/1236161.wgsl.expected.fxc.hlsl
@@ -1,18 +1,8 @@
-struct modf_result {
- float fract;
- float whole;
-};
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
[numthreads(1, 1, 1)]
void unused_entry_point() {
return;
}
void i() {
- const float s = tint_modf(1.0f).whole;
+ const float s = 1.0f;
}
diff --git a/test/tint/bug/chromium/1236161.wgsl.expected.glsl b/test/tint/bug/chromium/1236161.wgsl.expected.glsl
index 04734ec..f603bf3 100644
--- a/test/tint/bug/chromium/1236161.wgsl.expected.glsl
+++ b/test/tint/bug/chromium/1236161.wgsl.expected.glsl
@@ -1,22 +1,10 @@
#version 310 es
-struct modf_result {
- float fract;
- float whole;
-};
-
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
-
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void unused_entry_point() {
return;
}
void i() {
- float s = tint_modf(1.0f).whole;
+ float s = 1.0f;
}
diff --git a/test/tint/bug/chromium/1236161.wgsl.expected.msl b/test/tint/bug/chromium/1236161.wgsl.expected.msl
index 8d7aa31..6acecbd 100644
--- a/test/tint/bug/chromium/1236161.wgsl.expected.msl
+++ b/test/tint/bug/chromium/1236161.wgsl.expected.msl
@@ -1,18 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-
-struct modf_result {
- float fract;
- float whole;
-};
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void i() {
- float const s = tint_modf(1.0f).whole;
+ float const s = 1.0f;
}
diff --git a/test/tint/bug/chromium/1236161.wgsl.expected.spvasm b/test/tint/bug/chromium/1236161.wgsl.expected.spvasm
index 0d9b5a8..258d3c9 100644
--- a/test/tint/bug/chromium/1236161.wgsl.expected.spvasm
+++ b/test/tint/bug/chromium/1236161.wgsl.expected.spvasm
@@ -1,24 +1,17 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 13
+; Bound: 9
; Schema: 0
OpCapability Shader
- %10 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
OpExecutionMode %unused_entry_point LocalSize 1 1 1
OpName %unused_entry_point "unused_entry_point"
OpName %tint_symbol "tint_symbol"
- OpName %__modf_result "__modf_result"
- OpMemberName %__modf_result 0 "fract"
- OpMemberName %__modf_result 1 "whole"
- OpMemberDecorate %__modf_result 0 Offset 0
- OpMemberDecorate %__modf_result 1 Offset 4
%void = OpTypeVoid
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
-%__modf_result = OpTypeStruct %float %float
%float_1 = OpConstant %float 1
%unused_entry_point = OpFunction %void None %1
%4 = OpLabel
@@ -26,7 +19,5 @@
OpFunctionEnd
%tint_symbol = OpFunction %void None %1
%6 = OpLabel
- %7 = OpExtInst %__modf_result %10 ModfStruct %float_1
- %12 = OpCompositeExtract %float %7 1
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/modf/2d50da.wgsl b/test/tint/builtins/gen/literal/modf/2d50da.wgsl
index 8d3d0dd..e3212fc 100644
--- a/test/tint/builtins/gen/literal/modf/2d50da.wgsl
+++ b/test/tint/builtins/gen/literal/modf/2d50da.wgsl
@@ -23,7 +23,7 @@
// fn modf(vec<2, f32>) -> __modf_result_vec<2, f32>
fn modf_2d50da() {
- var res = modf(vec2<f32>(1.f));
+ var res = modf(vec2<f32>(-1.5f));
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.dxc.hlsl
index 2aba1e7..0dc03d8 100644
--- a/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.dxc.hlsl
@@ -2,14 +2,8 @@
float2 fract;
float2 whole;
};
-modf_result_vec2 tint_modf(float2 param_0) {
- modf_result_vec2 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_2d50da() {
- modf_result_vec2 res = tint_modf((1.0f).xx);
+ modf_result_vec2 res = {(-0.5f).xx, (-1.0f).xx};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.fxc.hlsl
index 2aba1e7..0dc03d8 100644
--- a/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.fxc.hlsl
@@ -2,14 +2,8 @@
float2 fract;
float2 whole;
};
-modf_result_vec2 tint_modf(float2 param_0) {
- modf_result_vec2 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_2d50da() {
- modf_result_vec2 res = tint_modf((1.0f).xx);
+ modf_result_vec2 res = {(-0.5f).xx, (-1.0f).xx};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.glsl b/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.glsl
index 540a580..881c541 100644
--- a/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.glsl
@@ -5,15 +5,9 @@
vec2 whole;
};
-modf_result_vec2 tint_modf(vec2 param_0) {
- modf_result_vec2 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_2d50da() {
- modf_result_vec2 res = tint_modf(vec2(1.0f));
+ modf_result_vec2 res = modf_result_vec2(vec2(-0.5f), vec2(-1.0f));
}
vec4 vertex_main() {
@@ -37,15 +31,9 @@
vec2 whole;
};
-modf_result_vec2 tint_modf(vec2 param_0) {
- modf_result_vec2 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_2d50da() {
- modf_result_vec2 res = tint_modf(vec2(1.0f));
+ modf_result_vec2 res = modf_result_vec2(vec2(-0.5f), vec2(-1.0f));
}
void fragment_main() {
@@ -63,15 +51,9 @@
vec2 whole;
};
-modf_result_vec2 tint_modf(vec2 param_0) {
- modf_result_vec2 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_2d50da() {
- modf_result_vec2 res = tint_modf(vec2(1.0f));
+ modf_result_vec2 res = modf_result_vec2(vec2(-0.5f), vec2(-1.0f));
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.msl b/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.msl
index edccc81..16a1941 100644
--- a/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.msl
@@ -6,14 +6,8 @@
float2 fract;
float2 whole;
};
-modf_result_vec2 tint_modf(float2 param_0) {
- modf_result_vec2 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_2d50da() {
- modf_result_vec2 res = tint_modf(float2(1.0f));
+ modf_result_vec2 res = modf_result_vec2{.fract=float2(-0.5f), .whole=float2(-1.0f)};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.spvasm
index bc75769..84f1689 100644
--- a/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.spvasm
@@ -1,10 +1,9 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 35
+; Bound: 37
; Schema: 0
OpCapability Shader
- %16 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -38,37 +37,40 @@
%9 = OpTypeFunction %void
%v2float = OpTypeVector %float 2
%__modf_result_vec2 = OpTypeStruct %v2float %v2float
- %float_1 = OpConstant %float 1
- %18 = OpConstantComposite %v2float %float_1 %float_1
+ %float_n0_5 = OpConstant %float -0.5
+ %16 = OpConstantComposite %v2float %float_n0_5 %float_n0_5
+ %float_n1 = OpConstant %float -1
+ %18 = OpConstantComposite %v2float %float_n1 %float_n1
+ %19 = OpConstantComposite %__modf_result_vec2 %16 %18
%_ptr_Function___modf_result_vec2 = OpTypePointer Function %__modf_result_vec2
- %21 = OpConstantNull %__modf_result_vec2
- %22 = OpTypeFunction %v4float
+ %22 = OpConstantNull %__modf_result_vec2
+ %23 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
%modf_2d50da = OpFunction %void None %9
%12 = OpLabel
- %res = OpVariable %_ptr_Function___modf_result_vec2 Function %21
- %13 = OpExtInst %__modf_result_vec2 %16 ModfStruct %18
- OpStore %res %13
+ %res = OpVariable %_ptr_Function___modf_result_vec2 Function %22
+ OpStore %res %19
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
- %24 = OpLabel
- %25 = OpFunctionCall %void %modf_2d50da
+%vertex_main_inner = OpFunction %v4float None %23
+ %25 = OpLabel
+ %26 = OpFunctionCall %void %modf_2d50da
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %9
- %27 = OpLabel
- %28 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %28
+ %28 = OpLabel
+ %29 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %29
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %9
- %30 = OpLabel
- %31 = OpFunctionCall %void %modf_2d50da
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %modf_2d50da
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %9
- %33 = OpLabel
- %34 = OpFunctionCall %void %modf_2d50da
+ %35 = OpLabel
+ %36 = OpFunctionCall %void %modf_2d50da
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.wgsl
index d893b25..a267089 100644
--- a/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/modf/2d50da.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
fn modf_2d50da() {
- var res = modf(vec2<f32>(1.0f));
+ var res = modf(vec2<f32>(-(1.5f)));
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/45005f.wgsl b/test/tint/builtins/gen/literal/modf/45005f.wgsl
index fe0b656..2b413ac 100644
--- a/test/tint/builtins/gen/literal/modf/45005f.wgsl
+++ b/test/tint/builtins/gen/literal/modf/45005f.wgsl
@@ -25,7 +25,7 @@
// fn modf(vec<3, f16>) -> __modf_result_vec<3, f16>
fn modf_45005f() {
- var res = modf(vec3<f16>(1.h));
+ var res = modf(vec3<f16>(-1.5h));
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.dxc.hlsl
index e4a0033..79a0841 100644
--- a/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.dxc.hlsl
@@ -2,14 +2,8 @@
vector<float16_t, 3> fract;
vector<float16_t, 3> whole;
};
-modf_result_vec3_f16 tint_modf(vector<float16_t, 3> param_0) {
- modf_result_vec3_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_45005f() {
- modf_result_vec3_f16 res = tint_modf((float16_t(1.0h)).xxx);
+ modf_result_vec3_f16 res = {(float16_t(-0.5h)).xxx, (float16_t(-1.0h)).xxx};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.glsl b/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.glsl
index d402eda..8bbd18f 100644
--- a/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.glsl
@@ -6,15 +6,9 @@
f16vec3 whole;
};
-modf_result_vec3_f16 tint_modf(f16vec3 param_0) {
- modf_result_vec3_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_45005f() {
- modf_result_vec3_f16 res = tint_modf(f16vec3(1.0hf));
+ modf_result_vec3_f16 res = modf_result_vec3_f16(f16vec3(-0.5hf), f16vec3(-1.0hf));
}
vec4 vertex_main() {
@@ -39,15 +33,9 @@
f16vec3 whole;
};
-modf_result_vec3_f16 tint_modf(f16vec3 param_0) {
- modf_result_vec3_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_45005f() {
- modf_result_vec3_f16 res = tint_modf(f16vec3(1.0hf));
+ modf_result_vec3_f16 res = modf_result_vec3_f16(f16vec3(-0.5hf), f16vec3(-1.0hf));
}
void fragment_main() {
@@ -66,15 +54,9 @@
f16vec3 whole;
};
-modf_result_vec3_f16 tint_modf(f16vec3 param_0) {
- modf_result_vec3_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_45005f() {
- modf_result_vec3_f16 res = tint_modf(f16vec3(1.0hf));
+ modf_result_vec3_f16 res = modf_result_vec3_f16(f16vec3(-0.5hf), f16vec3(-1.0hf));
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.msl b/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.msl
index a3071d4..5323690 100644
--- a/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.msl
@@ -6,14 +6,8 @@
half3 fract;
half3 whole;
};
-modf_result_vec3_f16 tint_modf(half3 param_0) {
- modf_result_vec3_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_45005f() {
- modf_result_vec3_f16 res = tint_modf(half3(1.0h));
+ modf_result_vec3_f16 res = modf_result_vec3_f16{.fract=half3(-0.5h), .whole=half3(-1.0h)};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.spvasm
index 63c4bd4..e276d40 100644
--- a/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.spvasm
@@ -1,14 +1,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 37
+; Bound: 38
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
- %17 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -43,38 +42,40 @@
%half = OpTypeFloat 16
%v3half = OpTypeVector %half 3
%__modf_result_vec3_f16 = OpTypeStruct %v3half %v3half
-%half_0x1p_0 = OpConstant %half 0x1p+0
- %19 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+%half_n0x1pn1 = OpConstant %half -0x1p-1
+ %17 = OpConstantComposite %v3half %half_n0x1pn1 %half_n0x1pn1 %half_n0x1pn1
+%half_n0x1p_0 = OpConstant %half -0x1p+0
+ %19 = OpConstantComposite %v3half %half_n0x1p_0 %half_n0x1p_0 %half_n0x1p_0
+ %20 = OpConstantComposite %__modf_result_vec3_f16 %17 %19
%_ptr_Function___modf_result_vec3_f16 = OpTypePointer Function %__modf_result_vec3_f16
- %22 = OpConstantNull %__modf_result_vec3_f16
- %23 = OpTypeFunction %v4float
+ %23 = OpConstantNull %__modf_result_vec3_f16
+ %24 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%modf_45005f = OpFunction %void None %9
%12 = OpLabel
- %res = OpVariable %_ptr_Function___modf_result_vec3_f16 Function %22
- %13 = OpExtInst %__modf_result_vec3_f16 %17 ModfStruct %19
- OpStore %res %13
+ %res = OpVariable %_ptr_Function___modf_result_vec3_f16 Function %23
+ OpStore %res %20
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %23
- %25 = OpLabel
- %26 = OpFunctionCall %void %modf_45005f
+%vertex_main_inner = OpFunction %v4float None %24
+ %26 = OpLabel
+ %27 = OpFunctionCall %void %modf_45005f
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %9
- %28 = OpLabel
- %29 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %29
+ %29 = OpLabel
+ %30 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %30
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %modf_45005f
+ %33 = OpLabel
+ %34 = OpFunctionCall %void %modf_45005f
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %9
- %35 = OpLabel
- %36 = OpFunctionCall %void %modf_45005f
+ %36 = OpLabel
+ %37 = OpFunctionCall %void %modf_45005f
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.wgsl
index 43f497e..5c77370 100644
--- a/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/modf/45005f.wgsl.expected.wgsl
@@ -1,7 +1,7 @@
enable f16;
fn modf_45005f() {
- var res = modf(vec3<f16>(1.0h));
+ var res = modf(vec3<f16>(-(1.5h)));
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/4bfced.wgsl b/test/tint/builtins/gen/literal/modf/4bfced.wgsl
index 1a3fd9a..86e3824 100644
--- a/test/tint/builtins/gen/literal/modf/4bfced.wgsl
+++ b/test/tint/builtins/gen/literal/modf/4bfced.wgsl
@@ -23,7 +23,7 @@
// fn modf(vec<4, f32>) -> __modf_result_vec<4, f32>
fn modf_4bfced() {
- var res = modf(vec4<f32>(1.f));
+ var res = modf(vec4<f32>(-1.5f));
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.dxc.hlsl
index 043bbab..deca14c 100644
--- a/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.dxc.hlsl
@@ -2,14 +2,8 @@
float4 fract;
float4 whole;
};
-modf_result_vec4 tint_modf(float4 param_0) {
- modf_result_vec4 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_4bfced() {
- modf_result_vec4 res = tint_modf((1.0f).xxxx);
+ modf_result_vec4 res = {(-0.5f).xxxx, (-1.0f).xxxx};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.fxc.hlsl
index 043bbab..deca14c 100644
--- a/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.fxc.hlsl
@@ -2,14 +2,8 @@
float4 fract;
float4 whole;
};
-modf_result_vec4 tint_modf(float4 param_0) {
- modf_result_vec4 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_4bfced() {
- modf_result_vec4 res = tint_modf((1.0f).xxxx);
+ modf_result_vec4 res = {(-0.5f).xxxx, (-1.0f).xxxx};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.glsl b/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.glsl
index 7c7e682..6952a13 100644
--- a/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.glsl
@@ -5,15 +5,9 @@
vec4 whole;
};
-modf_result_vec4 tint_modf(vec4 param_0) {
- modf_result_vec4 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_4bfced() {
- modf_result_vec4 res = tint_modf(vec4(1.0f));
+ modf_result_vec4 res = modf_result_vec4(vec4(-0.5f), vec4(-1.0f));
}
vec4 vertex_main() {
@@ -37,15 +31,9 @@
vec4 whole;
};
-modf_result_vec4 tint_modf(vec4 param_0) {
- modf_result_vec4 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_4bfced() {
- modf_result_vec4 res = tint_modf(vec4(1.0f));
+ modf_result_vec4 res = modf_result_vec4(vec4(-0.5f), vec4(-1.0f));
}
void fragment_main() {
@@ -63,15 +51,9 @@
vec4 whole;
};
-modf_result_vec4 tint_modf(vec4 param_0) {
- modf_result_vec4 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_4bfced() {
- modf_result_vec4 res = tint_modf(vec4(1.0f));
+ modf_result_vec4 res = modf_result_vec4(vec4(-0.5f), vec4(-1.0f));
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.msl b/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.msl
index 3b5da9c..9b97487 100644
--- a/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.msl
@@ -6,14 +6,8 @@
float4 fract;
float4 whole;
};
-modf_result_vec4 tint_modf(float4 param_0) {
- modf_result_vec4 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_4bfced() {
- modf_result_vec4 res = tint_modf(float4(1.0f));
+ modf_result_vec4 res = modf_result_vec4{.fract=float4(-0.5f), .whole=float4(-1.0f)};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.spvasm
index 2a58dd4..60d2fc2 100644
--- a/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.spvasm
@@ -1,10 +1,9 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 34
+; Bound: 36
; Schema: 0
OpCapability Shader
- %15 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -37,37 +36,40 @@
%void = OpTypeVoid
%9 = OpTypeFunction %void
%__modf_result_vec4 = OpTypeStruct %v4float %v4float
- %float_1 = OpConstant %float 1
- %17 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+ %float_n0_5 = OpConstant %float -0.5
+ %15 = OpConstantComposite %v4float %float_n0_5 %float_n0_5 %float_n0_5 %float_n0_5
+ %float_n1 = OpConstant %float -1
+ %17 = OpConstantComposite %v4float %float_n1 %float_n1 %float_n1 %float_n1
+ %18 = OpConstantComposite %__modf_result_vec4 %15 %17
%_ptr_Function___modf_result_vec4 = OpTypePointer Function %__modf_result_vec4
- %20 = OpConstantNull %__modf_result_vec4
- %21 = OpTypeFunction %v4float
+ %21 = OpConstantNull %__modf_result_vec4
+ %22 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
%modf_4bfced = OpFunction %void None %9
%12 = OpLabel
- %res = OpVariable %_ptr_Function___modf_result_vec4 Function %20
- %13 = OpExtInst %__modf_result_vec4 %15 ModfStruct %17
- OpStore %res %13
+ %res = OpVariable %_ptr_Function___modf_result_vec4 Function %21
+ OpStore %res %18
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %21
- %23 = OpLabel
- %24 = OpFunctionCall %void %modf_4bfced
+%vertex_main_inner = OpFunction %v4float None %22
+ %24 = OpLabel
+ %25 = OpFunctionCall %void %modf_4bfced
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %9
- %26 = OpLabel
- %27 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %27
+ %27 = OpLabel
+ %28 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %28
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %9
- %29 = OpLabel
- %30 = OpFunctionCall %void %modf_4bfced
+ %31 = OpLabel
+ %32 = OpFunctionCall %void %modf_4bfced
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %modf_4bfced
+ %34 = OpLabel
+ %35 = OpFunctionCall %void %modf_4bfced
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.wgsl
index 0b04b87..2ab4612 100644
--- a/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/modf/4bfced.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
fn modf_4bfced() {
- var res = modf(vec4<f32>(1.0f));
+ var res = modf(vec4<f32>(-(1.5f)));
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/5ea256.wgsl b/test/tint/builtins/gen/literal/modf/5ea256.wgsl
index 673bb25..1d6a103 100644
--- a/test/tint/builtins/gen/literal/modf/5ea256.wgsl
+++ b/test/tint/builtins/gen/literal/modf/5ea256.wgsl
@@ -23,7 +23,7 @@
// fn modf(vec<3, f32>) -> __modf_result_vec<3, f32>
fn modf_5ea256() {
- var res = modf(vec3<f32>(1.f));
+ var res = modf(vec3<f32>(-1.5f));
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.dxc.hlsl
index a2c46f5..e594903 100644
--- a/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.dxc.hlsl
@@ -2,14 +2,8 @@
float3 fract;
float3 whole;
};
-modf_result_vec3 tint_modf(float3 param_0) {
- modf_result_vec3 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_5ea256() {
- modf_result_vec3 res = tint_modf((1.0f).xxx);
+ modf_result_vec3 res = {(-0.5f).xxx, (-1.0f).xxx};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.fxc.hlsl
index a2c46f5..e594903 100644
--- a/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.fxc.hlsl
@@ -2,14 +2,8 @@
float3 fract;
float3 whole;
};
-modf_result_vec3 tint_modf(float3 param_0) {
- modf_result_vec3 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_5ea256() {
- modf_result_vec3 res = tint_modf((1.0f).xxx);
+ modf_result_vec3 res = {(-0.5f).xxx, (-1.0f).xxx};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.glsl b/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.glsl
index 7a727ec..84138c0 100644
--- a/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.glsl
@@ -5,15 +5,9 @@
vec3 whole;
};
-modf_result_vec3 tint_modf(vec3 param_0) {
- modf_result_vec3 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_5ea256() {
- modf_result_vec3 res = tint_modf(vec3(1.0f));
+ modf_result_vec3 res = modf_result_vec3(vec3(-0.5f), vec3(-1.0f));
}
vec4 vertex_main() {
@@ -37,15 +31,9 @@
vec3 whole;
};
-modf_result_vec3 tint_modf(vec3 param_0) {
- modf_result_vec3 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_5ea256() {
- modf_result_vec3 res = tint_modf(vec3(1.0f));
+ modf_result_vec3 res = modf_result_vec3(vec3(-0.5f), vec3(-1.0f));
}
void fragment_main() {
@@ -63,15 +51,9 @@
vec3 whole;
};
-modf_result_vec3 tint_modf(vec3 param_0) {
- modf_result_vec3 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_5ea256() {
- modf_result_vec3 res = tint_modf(vec3(1.0f));
+ modf_result_vec3 res = modf_result_vec3(vec3(-0.5f), vec3(-1.0f));
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.msl b/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.msl
index 33e7fa3..4e73398 100644
--- a/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.msl
@@ -6,14 +6,8 @@
float3 fract;
float3 whole;
};
-modf_result_vec3 tint_modf(float3 param_0) {
- modf_result_vec3 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_5ea256() {
- modf_result_vec3 res = tint_modf(float3(1.0f));
+ modf_result_vec3 res = modf_result_vec3{.fract=float3(-0.5f), .whole=float3(-1.0f)};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.spvasm
index 6762d61..62e6d32 100644
--- a/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.spvasm
@@ -1,10 +1,9 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 35
+; Bound: 37
; Schema: 0
OpCapability Shader
- %16 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -38,37 +37,40 @@
%9 = OpTypeFunction %void
%v3float = OpTypeVector %float 3
%__modf_result_vec3 = OpTypeStruct %v3float %v3float
- %float_1 = OpConstant %float 1
- %18 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+ %float_n0_5 = OpConstant %float -0.5
+ %16 = OpConstantComposite %v3float %float_n0_5 %float_n0_5 %float_n0_5
+ %float_n1 = OpConstant %float -1
+ %18 = OpConstantComposite %v3float %float_n1 %float_n1 %float_n1
+ %19 = OpConstantComposite %__modf_result_vec3 %16 %18
%_ptr_Function___modf_result_vec3 = OpTypePointer Function %__modf_result_vec3
- %21 = OpConstantNull %__modf_result_vec3
- %22 = OpTypeFunction %v4float
+ %22 = OpConstantNull %__modf_result_vec3
+ %23 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
%modf_5ea256 = OpFunction %void None %9
%12 = OpLabel
- %res = OpVariable %_ptr_Function___modf_result_vec3 Function %21
- %13 = OpExtInst %__modf_result_vec3 %16 ModfStruct %18
- OpStore %res %13
+ %res = OpVariable %_ptr_Function___modf_result_vec3 Function %22
+ OpStore %res %19
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %22
- %24 = OpLabel
- %25 = OpFunctionCall %void %modf_5ea256
+%vertex_main_inner = OpFunction %v4float None %23
+ %25 = OpLabel
+ %26 = OpFunctionCall %void %modf_5ea256
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %9
- %27 = OpLabel
- %28 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %28
+ %28 = OpLabel
+ %29 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %29
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %9
- %30 = OpLabel
- %31 = OpFunctionCall %void %modf_5ea256
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %modf_5ea256
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %9
- %33 = OpLabel
- %34 = OpFunctionCall %void %modf_5ea256
+ %35 = OpLabel
+ %36 = OpFunctionCall %void %modf_5ea256
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.wgsl
index ea4142d..8132d91 100644
--- a/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/modf/5ea256.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
fn modf_5ea256() {
- var res = modf(vec3<f32>(1.0f));
+ var res = modf(vec3<f32>(-(1.5f)));
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl
index 591a299..2af78fd 100644
--- a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl
+++ b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl
@@ -25,7 +25,7 @@
// fn modf(f16) -> __modf_result<f16>
fn modf_8dbbbf() {
- var res = modf(1.h);
+ var res = modf(-1.5h);
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.dxc.hlsl
index c84d408..cd06924 100644
--- a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.dxc.hlsl
@@ -2,14 +2,8 @@
float16_t fract;
float16_t whole;
};
-modf_result_f16 tint_modf(float16_t param_0) {
- modf_result_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_8dbbbf() {
- modf_result_f16 res = tint_modf(float16_t(1.0h));
+ modf_result_f16 res = {float16_t(-0.5h), float16_t(-1.0h)};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.glsl b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.glsl
index f9757b9..0741bb8 100644
--- a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.glsl
@@ -6,15 +6,9 @@
float16_t whole;
};
-modf_result_f16 tint_modf(float16_t param_0) {
- modf_result_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_8dbbbf() {
- modf_result_f16 res = tint_modf(1.0hf);
+ modf_result_f16 res = modf_result_f16(-0.5hf, -1.0hf);
}
vec4 vertex_main() {
@@ -39,15 +33,9 @@
float16_t whole;
};
-modf_result_f16 tint_modf(float16_t param_0) {
- modf_result_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_8dbbbf() {
- modf_result_f16 res = tint_modf(1.0hf);
+ modf_result_f16 res = modf_result_f16(-0.5hf, -1.0hf);
}
void fragment_main() {
@@ -66,15 +54,9 @@
float16_t whole;
};
-modf_result_f16 tint_modf(float16_t param_0) {
- modf_result_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_8dbbbf() {
- modf_result_f16 res = tint_modf(1.0hf);
+ modf_result_f16 res = modf_result_f16(-0.5hf, -1.0hf);
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.msl b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.msl
index e9472e7..6d02f05 100644
--- a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.msl
@@ -6,14 +6,8 @@
half fract;
half whole;
};
-modf_result_f16 tint_modf(half param_0) {
- modf_result_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_8dbbbf() {
- modf_result_f16 res = tint_modf(1.0h);
+ modf_result_f16 res = modf_result_f16{.fract=-0.5h, .whole=-1.0h};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.spvasm
index ebb1132..6f758be 100644
--- a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.spvasm
@@ -8,7 +8,6 @@
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
- %16 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -42,7 +41,9 @@
%9 = OpTypeFunction %void
%half = OpTypeFloat 16
%__modf_result_f16 = OpTypeStruct %half %half
-%half_0x1p_0 = OpConstant %half 0x1p+0
+%half_n0x1pn1 = OpConstant %half -0x1p-1
+%half_n0x1p_0 = OpConstant %half -0x1p+0
+ %17 = OpConstantComposite %__modf_result_f16 %half_n0x1pn1 %half_n0x1p_0
%_ptr_Function___modf_result_f16 = OpTypePointer Function %__modf_result_f16
%20 = OpConstantNull %__modf_result_f16
%21 = OpTypeFunction %v4float
@@ -50,8 +51,7 @@
%modf_8dbbbf = OpFunction %void None %9
%12 = OpLabel
%res = OpVariable %_ptr_Function___modf_result_f16 Function %20
- %13 = OpExtInst %__modf_result_f16 %16 ModfStruct %half_0x1p_0
- OpStore %res %13
+ OpStore %res %17
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %21
diff --git a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.wgsl
index c3619a7..d25d72f 100644
--- a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl.expected.wgsl
@@ -1,7 +1,7 @@
enable f16;
fn modf_8dbbbf() {
- var res = modf(1.0h);
+ var res = modf(-(1.5h));
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/995934.wgsl b/test/tint/builtins/gen/literal/modf/995934.wgsl
index bcb6006..59170ae 100644
--- a/test/tint/builtins/gen/literal/modf/995934.wgsl
+++ b/test/tint/builtins/gen/literal/modf/995934.wgsl
@@ -25,7 +25,7 @@
// fn modf(vec<4, f16>) -> __modf_result_vec<4, f16>
fn modf_995934() {
- var res = modf(vec4<f16>(1.h));
+ var res = modf(vec4<f16>(-1.5h));
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.dxc.hlsl
index 24c71de..cc8ede5 100644
--- a/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.dxc.hlsl
@@ -2,14 +2,8 @@
vector<float16_t, 4> fract;
vector<float16_t, 4> whole;
};
-modf_result_vec4_f16 tint_modf(vector<float16_t, 4> param_0) {
- modf_result_vec4_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_995934() {
- modf_result_vec4_f16 res = tint_modf((float16_t(1.0h)).xxxx);
+ modf_result_vec4_f16 res = {(float16_t(-0.5h)).xxxx, (float16_t(-1.0h)).xxxx};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.glsl b/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.glsl
index 065a9e7..cafadca 100644
--- a/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.glsl
@@ -6,15 +6,9 @@
f16vec4 whole;
};
-modf_result_vec4_f16 tint_modf(f16vec4 param_0) {
- modf_result_vec4_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_995934() {
- modf_result_vec4_f16 res = tint_modf(f16vec4(1.0hf));
+ modf_result_vec4_f16 res = modf_result_vec4_f16(f16vec4(-0.5hf), f16vec4(-1.0hf));
}
vec4 vertex_main() {
@@ -39,15 +33,9 @@
f16vec4 whole;
};
-modf_result_vec4_f16 tint_modf(f16vec4 param_0) {
- modf_result_vec4_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_995934() {
- modf_result_vec4_f16 res = tint_modf(f16vec4(1.0hf));
+ modf_result_vec4_f16 res = modf_result_vec4_f16(f16vec4(-0.5hf), f16vec4(-1.0hf));
}
void fragment_main() {
@@ -66,15 +54,9 @@
f16vec4 whole;
};
-modf_result_vec4_f16 tint_modf(f16vec4 param_0) {
- modf_result_vec4_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_995934() {
- modf_result_vec4_f16 res = tint_modf(f16vec4(1.0hf));
+ modf_result_vec4_f16 res = modf_result_vec4_f16(f16vec4(-0.5hf), f16vec4(-1.0hf));
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.msl b/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.msl
index 76d3c04..6371cad 100644
--- a/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.msl
@@ -6,14 +6,8 @@
half4 fract;
half4 whole;
};
-modf_result_vec4_f16 tint_modf(half4 param_0) {
- modf_result_vec4_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_995934() {
- modf_result_vec4_f16 res = tint_modf(half4(1.0h));
+ modf_result_vec4_f16 res = modf_result_vec4_f16{.fract=half4(-0.5h), .whole=half4(-1.0h)};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.spvasm
index b19f599..d67833b 100644
--- a/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.spvasm
@@ -1,14 +1,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 37
+; Bound: 38
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
- %17 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -43,38 +42,40 @@
%half = OpTypeFloat 16
%v4half = OpTypeVector %half 4
%__modf_result_vec4_f16 = OpTypeStruct %v4half %v4half
-%half_0x1p_0 = OpConstant %half 0x1p+0
- %19 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+%half_n0x1pn1 = OpConstant %half -0x1p-1
+ %17 = OpConstantComposite %v4half %half_n0x1pn1 %half_n0x1pn1 %half_n0x1pn1 %half_n0x1pn1
+%half_n0x1p_0 = OpConstant %half -0x1p+0
+ %19 = OpConstantComposite %v4half %half_n0x1p_0 %half_n0x1p_0 %half_n0x1p_0 %half_n0x1p_0
+ %20 = OpConstantComposite %__modf_result_vec4_f16 %17 %19
%_ptr_Function___modf_result_vec4_f16 = OpTypePointer Function %__modf_result_vec4_f16
- %22 = OpConstantNull %__modf_result_vec4_f16
- %23 = OpTypeFunction %v4float
+ %23 = OpConstantNull %__modf_result_vec4_f16
+ %24 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%modf_995934 = OpFunction %void None %9
%12 = OpLabel
- %res = OpVariable %_ptr_Function___modf_result_vec4_f16 Function %22
- %13 = OpExtInst %__modf_result_vec4_f16 %17 ModfStruct %19
- OpStore %res %13
+ %res = OpVariable %_ptr_Function___modf_result_vec4_f16 Function %23
+ OpStore %res %20
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %23
- %25 = OpLabel
- %26 = OpFunctionCall %void %modf_995934
+%vertex_main_inner = OpFunction %v4float None %24
+ %26 = OpLabel
+ %27 = OpFunctionCall %void %modf_995934
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %9
- %28 = OpLabel
- %29 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %29
+ %29 = OpLabel
+ %30 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %30
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %modf_995934
+ %33 = OpLabel
+ %34 = OpFunctionCall %void %modf_995934
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %9
- %35 = OpLabel
- %36 = OpFunctionCall %void %modf_995934
+ %36 = OpLabel
+ %37 = OpFunctionCall %void %modf_995934
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.wgsl
index a73275d..c514967 100644
--- a/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/modf/995934.wgsl.expected.wgsl
@@ -1,7 +1,7 @@
enable f16;
fn modf_995934() {
- var res = modf(vec4<f16>(1.0h));
+ var res = modf(vec4<f16>(-(1.5h)));
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/a545b9.wgsl b/test/tint/builtins/gen/literal/modf/a545b9.wgsl
index 1a66aad..a1c1c6e 100644
--- a/test/tint/builtins/gen/literal/modf/a545b9.wgsl
+++ b/test/tint/builtins/gen/literal/modf/a545b9.wgsl
@@ -25,7 +25,7 @@
// fn modf(vec<2, f16>) -> __modf_result_vec<2, f16>
fn modf_a545b9() {
- var res = modf(vec2<f16>(1.h));
+ var res = modf(vec2<f16>(-1.5h));
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.dxc.hlsl
index 1805568..e3fe286 100644
--- a/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.dxc.hlsl
@@ -2,14 +2,8 @@
vector<float16_t, 2> fract;
vector<float16_t, 2> whole;
};
-modf_result_vec2_f16 tint_modf(vector<float16_t, 2> param_0) {
- modf_result_vec2_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_a545b9() {
- modf_result_vec2_f16 res = tint_modf((float16_t(1.0h)).xx);
+ modf_result_vec2_f16 res = {(float16_t(-0.5h)).xx, (float16_t(-1.0h)).xx};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.glsl b/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.glsl
index 7c2f892..3bcc8fa 100644
--- a/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.glsl
@@ -6,15 +6,9 @@
f16vec2 whole;
};
-modf_result_vec2_f16 tint_modf(f16vec2 param_0) {
- modf_result_vec2_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_a545b9() {
- modf_result_vec2_f16 res = tint_modf(f16vec2(1.0hf));
+ modf_result_vec2_f16 res = modf_result_vec2_f16(f16vec2(-0.5hf), f16vec2(-1.0hf));
}
vec4 vertex_main() {
@@ -39,15 +33,9 @@
f16vec2 whole;
};
-modf_result_vec2_f16 tint_modf(f16vec2 param_0) {
- modf_result_vec2_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_a545b9() {
- modf_result_vec2_f16 res = tint_modf(f16vec2(1.0hf));
+ modf_result_vec2_f16 res = modf_result_vec2_f16(f16vec2(-0.5hf), f16vec2(-1.0hf));
}
void fragment_main() {
@@ -66,15 +54,9 @@
f16vec2 whole;
};
-modf_result_vec2_f16 tint_modf(f16vec2 param_0) {
- modf_result_vec2_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_a545b9() {
- modf_result_vec2_f16 res = tint_modf(f16vec2(1.0hf));
+ modf_result_vec2_f16 res = modf_result_vec2_f16(f16vec2(-0.5hf), f16vec2(-1.0hf));
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.msl b/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.msl
index d7ce71e..18005ff 100644
--- a/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.msl
@@ -6,14 +6,8 @@
half2 fract;
half2 whole;
};
-modf_result_vec2_f16 tint_modf(half2 param_0) {
- modf_result_vec2_f16 result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_a545b9() {
- modf_result_vec2_f16 res = tint_modf(half2(1.0h));
+ modf_result_vec2_f16 res = modf_result_vec2_f16{.fract=half2(-0.5h), .whole=half2(-1.0h)};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.spvasm
index 63d3be1..a1c6140 100644
--- a/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.spvasm
@@ -1,14 +1,13 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 37
+; Bound: 38
; Schema: 0
OpCapability Shader
OpCapability Float16
OpCapability UniformAndStorageBuffer16BitAccess
OpCapability StorageBuffer16BitAccess
OpCapability StorageInputOutput16
- %17 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -43,38 +42,40 @@
%half = OpTypeFloat 16
%v2half = OpTypeVector %half 2
%__modf_result_vec2_f16 = OpTypeStruct %v2half %v2half
-%half_0x1p_0 = OpConstant %half 0x1p+0
- %19 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
+%half_n0x1pn1 = OpConstant %half -0x1p-1
+ %17 = OpConstantComposite %v2half %half_n0x1pn1 %half_n0x1pn1
+%half_n0x1p_0 = OpConstant %half -0x1p+0
+ %19 = OpConstantComposite %v2half %half_n0x1p_0 %half_n0x1p_0
+ %20 = OpConstantComposite %__modf_result_vec2_f16 %17 %19
%_ptr_Function___modf_result_vec2_f16 = OpTypePointer Function %__modf_result_vec2_f16
- %22 = OpConstantNull %__modf_result_vec2_f16
- %23 = OpTypeFunction %v4float
+ %23 = OpConstantNull %__modf_result_vec2_f16
+ %24 = OpTypeFunction %v4float
%float_1 = OpConstant %float 1
%modf_a545b9 = OpFunction %void None %9
%12 = OpLabel
- %res = OpVariable %_ptr_Function___modf_result_vec2_f16 Function %22
- %13 = OpExtInst %__modf_result_vec2_f16 %17 ModfStruct %19
- OpStore %res %13
+ %res = OpVariable %_ptr_Function___modf_result_vec2_f16 Function %23
+ OpStore %res %20
OpReturn
OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %23
- %25 = OpLabel
- %26 = OpFunctionCall %void %modf_a545b9
+%vertex_main_inner = OpFunction %v4float None %24
+ %26 = OpLabel
+ %27 = OpFunctionCall %void %modf_a545b9
OpReturnValue %5
OpFunctionEnd
%vertex_main = OpFunction %void None %9
- %28 = OpLabel
- %29 = OpFunctionCall %v4float %vertex_main_inner
- OpStore %value %29
+ %29 = OpLabel
+ %30 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %30
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %modf_a545b9
+ %33 = OpLabel
+ %34 = OpFunctionCall %void %modf_a545b9
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %9
- %35 = OpLabel
- %36 = OpFunctionCall %void %modf_a545b9
+ %36 = OpLabel
+ %37 = OpFunctionCall %void %modf_a545b9
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.wgsl
index 9d2b453..c7bc52b 100644
--- a/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/modf/a545b9.wgsl.expected.wgsl
@@ -1,7 +1,7 @@
enable f16;
fn modf_a545b9() {
- var res = modf(vec2<f16>(1.0h));
+ var res = modf(vec2<f16>(-(1.5h)));
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl
index d0ffdaa..bdcd0f0 100644
--- a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl
+++ b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl
@@ -23,7 +23,7 @@
// fn modf(f32) -> __modf_result<f32>
fn modf_bbf7f7() {
- var res = modf(1.f);
+ var res = modf(-1.5f);
}
@vertex
diff --git a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.dxc.hlsl
index 020478a..ec4e06b 100644
--- a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.dxc.hlsl
@@ -2,14 +2,8 @@
float fract;
float whole;
};
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_bbf7f7() {
- modf_result res = tint_modf(1.0f);
+ modf_result res = {-0.5f, -1.0f};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.fxc.hlsl
index 020478a..ec4e06b 100644
--- a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.fxc.hlsl
@@ -2,14 +2,8 @@
float fract;
float whole;
};
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_bbf7f7() {
- modf_result res = tint_modf(1.0f);
+ modf_result res = {-0.5f, -1.0f};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.glsl b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.glsl
index 30917ac..177ebe9 100644
--- a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.glsl
@@ -5,15 +5,9 @@
float whole;
};
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_bbf7f7() {
- modf_result res = tint_modf(1.0f);
+ modf_result res = modf_result(-0.5f, -1.0f);
}
vec4 vertex_main() {
@@ -37,15 +31,9 @@
float whole;
};
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_bbf7f7() {
- modf_result res = tint_modf(1.0f);
+ modf_result res = modf_result(-0.5f, -1.0f);
}
void fragment_main() {
@@ -63,15 +51,9 @@
float whole;
};
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_bbf7f7() {
- modf_result res = tint_modf(1.0f);
+ modf_result res = modf_result(-0.5f, -1.0f);
}
void compute_main() {
diff --git a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.msl b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.msl
index 250e830..f5fdb55 100644
--- a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.msl
@@ -6,14 +6,8 @@
float fract;
float whole;
};
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void modf_bbf7f7() {
- modf_result res = tint_modf(1.0f);
+ modf_result res = modf_result{.fract=-0.5f, .whole=-1.0f};
}
struct tint_symbol {
diff --git a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.spvasm
index 21a2358..1f64077 100644
--- a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.spvasm
@@ -1,10 +1,9 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 33
+; Bound: 34
; Schema: 0
OpCapability Shader
- %15 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
OpEntryPoint Fragment %fragment_main "fragment_main"
@@ -37,15 +36,17 @@
%void = OpTypeVoid
%9 = OpTypeFunction %void
%__modf_result = OpTypeStruct %float %float
- %float_1 = OpConstant %float 1
+ %float_n0_5 = OpConstant %float -0.5
+ %float_n1 = OpConstant %float -1
+ %16 = OpConstantComposite %__modf_result %float_n0_5 %float_n1
%_ptr_Function___modf_result = OpTypePointer Function %__modf_result
%19 = OpConstantNull %__modf_result
%20 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
%modf_bbf7f7 = OpFunction %void None %9
%12 = OpLabel
%res = OpVariable %_ptr_Function___modf_result Function %19
- %13 = OpExtInst %__modf_result %15 ModfStruct %float_1
- OpStore %res %13
+ OpStore %res %16
OpReturn
OpFunctionEnd
%vertex_main_inner = OpFunction %v4float None %20
@@ -61,12 +62,12 @@
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %9
- %28 = OpLabel
- %29 = OpFunctionCall %void %modf_bbf7f7
+ %29 = OpLabel
+ %30 = OpFunctionCall %void %modf_bbf7f7
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %modf_bbf7f7
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %modf_bbf7f7
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.wgsl
index 31c86d6..a93f3e3 100644
--- a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
fn modf_bbf7f7() {
- var res = modf(1.0f);
+ var res = modf(-(1.5f));
}
@vertex
diff --git a/test/tint/builtins/gen/var/modf/2d50da.wgsl b/test/tint/builtins/gen/var/modf/2d50da.wgsl
index d8f067b..00be018 100644
--- a/test/tint/builtins/gen/var/modf/2d50da.wgsl
+++ b/test/tint/builtins/gen/var/modf/2d50da.wgsl
@@ -23,7 +23,7 @@
// fn modf(vec<2, f32>) -> __modf_result_vec<2, f32>
fn modf_2d50da() {
- var arg_0 = vec2<f32>(1.f);
+ var arg_0 = vec2<f32>(-1.5f);
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.dxc.hlsl
index bf9728e..a3745ba 100644
--- a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.dxc.hlsl
@@ -9,7 +9,7 @@
}
void modf_2d50da() {
- float2 arg_0 = (1.0f).xx;
+ float2 arg_0 = (-1.5f).xx;
modf_result_vec2 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.fxc.hlsl
index bf9728e..a3745ba 100644
--- a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.fxc.hlsl
@@ -9,7 +9,7 @@
}
void modf_2d50da() {
- float2 arg_0 = (1.0f).xx;
+ float2 arg_0 = (-1.5f).xx;
modf_result_vec2 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.glsl b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.glsl
index a1998b9..d7ecc21 100644
--- a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.glsl
@@ -13,7 +13,7 @@
void modf_2d50da() {
- vec2 arg_0 = vec2(1.0f);
+ vec2 arg_0 = vec2(-1.5f);
modf_result_vec2 res = tint_modf(arg_0);
}
@@ -46,7 +46,7 @@
void modf_2d50da() {
- vec2 arg_0 = vec2(1.0f);
+ vec2 arg_0 = vec2(-1.5f);
modf_result_vec2 res = tint_modf(arg_0);
}
@@ -73,7 +73,7 @@
void modf_2d50da() {
- vec2 arg_0 = vec2(1.0f);
+ vec2 arg_0 = vec2(-1.5f);
modf_result_vec2 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.msl b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.msl
index 7579576..5963525 100644
--- a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.msl
@@ -13,7 +13,7 @@
}
void modf_2d50da() {
- float2 arg_0 = float2(1.0f);
+ float2 arg_0 = float2(-1.5f);
modf_result_vec2 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.spvasm b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.spvasm
index 9e6337b..4fda240 100644
--- a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 39
+; Bound: 40
; Schema: 0
OpCapability Shader
%21 = OpExtInstImport "GLSL.std.450"
@@ -38,14 +38,15 @@
%void = OpTypeVoid
%9 = OpTypeFunction %void
%v2float = OpTypeVector %float 2
- %float_1 = OpConstant %float 1
- %15 = OpConstantComposite %v2float %float_1 %float_1
+ %float_n1_5 = OpConstant %float -1.5
+ %15 = OpConstantComposite %v2float %float_n1_5 %float_n1_5
%_ptr_Function_v2float = OpTypePointer Function %v2float
%18 = OpConstantNull %v2float
%__modf_result_vec2 = OpTypeStruct %v2float %v2float
%_ptr_Function___modf_result_vec2 = OpTypePointer Function %__modf_result_vec2
%25 = OpConstantNull %__modf_result_vec2
%26 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
%modf_2d50da = OpFunction %void None %9
%12 = OpLabel
%arg_0 = OpVariable %_ptr_Function_v2float Function %18
@@ -69,12 +70,12 @@
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %9
- %34 = OpLabel
- %35 = OpFunctionCall %void %modf_2d50da
+ %35 = OpLabel
+ %36 = OpFunctionCall %void %modf_2d50da
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %9
- %37 = OpLabel
- %38 = OpFunctionCall %void %modf_2d50da
+ %38 = OpLabel
+ %39 = OpFunctionCall %void %modf_2d50da
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.wgsl b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.wgsl
index baa5039..1ea4d30 100644
--- a/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/modf/2d50da.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
fn modf_2d50da() {
- var arg_0 = vec2<f32>(1.0f);
+ var arg_0 = vec2<f32>(-(1.5f));
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/45005f.wgsl b/test/tint/builtins/gen/var/modf/45005f.wgsl
index 3764709..8955e23 100644
--- a/test/tint/builtins/gen/var/modf/45005f.wgsl
+++ b/test/tint/builtins/gen/var/modf/45005f.wgsl
@@ -25,7 +25,7 @@
// fn modf(vec<3, f16>) -> __modf_result_vec<3, f16>
fn modf_45005f() {
- var arg_0 = vec3<f16>(1.h);
+ var arg_0 = vec3<f16>(-1.5h);
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.dxc.hlsl
index f9a9e8d..db8c084 100644
--- a/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.dxc.hlsl
@@ -9,7 +9,7 @@
}
void modf_45005f() {
- vector<float16_t, 3> arg_0 = (float16_t(1.0h)).xxx;
+ vector<float16_t, 3> arg_0 = (float16_t(-1.5h)).xxx;
modf_result_vec3_f16 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.glsl b/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.glsl
index 0c6aad6..f019a7d 100644
--- a/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.glsl
@@ -14,7 +14,7 @@
void modf_45005f() {
- f16vec3 arg_0 = f16vec3(1.0hf);
+ f16vec3 arg_0 = f16vec3(-1.5hf);
modf_result_vec3_f16 res = tint_modf(arg_0);
}
@@ -48,7 +48,7 @@
void modf_45005f() {
- f16vec3 arg_0 = f16vec3(1.0hf);
+ f16vec3 arg_0 = f16vec3(-1.5hf);
modf_result_vec3_f16 res = tint_modf(arg_0);
}
@@ -76,7 +76,7 @@
void modf_45005f() {
- f16vec3 arg_0 = f16vec3(1.0hf);
+ f16vec3 arg_0 = f16vec3(-1.5hf);
modf_result_vec3_f16 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.msl b/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.msl
index 106d1b6..5c9713f 100644
--- a/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.msl
@@ -13,7 +13,7 @@
}
void modf_45005f() {
- half3 arg_0 = half3(1.0h);
+ half3 arg_0 = half3(-1.5h);
modf_result_vec3_f16 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.spvasm b/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.spvasm
index 7c7b813..d647058 100644
--- a/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.spvasm
@@ -43,8 +43,8 @@
%9 = OpTypeFunction %void
%half = OpTypeFloat 16
%v3half = OpTypeVector %half 3
-%half_0x1p_0 = OpConstant %half 0x1p+0
- %16 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+%half_n0x1_8p_0 = OpConstant %half -0x1.8p+0
+ %16 = OpConstantComposite %v3half %half_n0x1_8p_0 %half_n0x1_8p_0 %half_n0x1_8p_0
%_ptr_Function_v3half = OpTypePointer Function %v3half
%19 = OpConstantNull %v3half
%__modf_result_vec3_f16 = OpTypeStruct %v3half %v3half
diff --git a/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.wgsl b/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.wgsl
index f9e08f3..96fce1a 100644
--- a/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/modf/45005f.wgsl.expected.wgsl
@@ -1,7 +1,7 @@
enable f16;
fn modf_45005f() {
- var arg_0 = vec3<f16>(1.0h);
+ var arg_0 = vec3<f16>(-(1.5h));
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/4bfced.wgsl b/test/tint/builtins/gen/var/modf/4bfced.wgsl
index 4929415..52581ce 100644
--- a/test/tint/builtins/gen/var/modf/4bfced.wgsl
+++ b/test/tint/builtins/gen/var/modf/4bfced.wgsl
@@ -23,7 +23,7 @@
// fn modf(vec<4, f32>) -> __modf_result_vec<4, f32>
fn modf_4bfced() {
- var arg_0 = vec4<f32>(1.f);
+ var arg_0 = vec4<f32>(-1.5f);
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.dxc.hlsl
index 0cfe790..e084fcc 100644
--- a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.dxc.hlsl
@@ -9,7 +9,7 @@
}
void modf_4bfced() {
- float4 arg_0 = (1.0f).xxxx;
+ float4 arg_0 = (-1.5f).xxxx;
modf_result_vec4 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.fxc.hlsl
index 0cfe790..e084fcc 100644
--- a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.fxc.hlsl
@@ -9,7 +9,7 @@
}
void modf_4bfced() {
- float4 arg_0 = (1.0f).xxxx;
+ float4 arg_0 = (-1.5f).xxxx;
modf_result_vec4 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.glsl b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.glsl
index 5d5615b..5a153b9 100644
--- a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.glsl
@@ -13,7 +13,7 @@
void modf_4bfced() {
- vec4 arg_0 = vec4(1.0f);
+ vec4 arg_0 = vec4(-1.5f);
modf_result_vec4 res = tint_modf(arg_0);
}
@@ -46,7 +46,7 @@
void modf_4bfced() {
- vec4 arg_0 = vec4(1.0f);
+ vec4 arg_0 = vec4(-1.5f);
modf_result_vec4 res = tint_modf(arg_0);
}
@@ -73,7 +73,7 @@
void modf_4bfced() {
- vec4 arg_0 = vec4(1.0f);
+ vec4 arg_0 = vec4(-1.5f);
modf_result_vec4 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.msl b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.msl
index 4bac39e..fe36491 100644
--- a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.msl
@@ -13,7 +13,7 @@
}
void modf_4bfced() {
- float4 arg_0 = float4(1.0f);
+ float4 arg_0 = float4(-1.5f);
modf_result_vec4 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.spvasm b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.spvasm
index 44dedfe..8bdaea4 100644
--- a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 37
+; Bound: 38
; Schema: 0
OpCapability Shader
%19 = OpExtInstImport "GLSL.std.450"
@@ -37,13 +37,14 @@
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%void = OpTypeVoid
%9 = OpTypeFunction %void
- %float_1 = OpConstant %float 1
- %14 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+ %float_n1_5 = OpConstant %float -1.5
+ %14 = OpConstantComposite %v4float %float_n1_5 %float_n1_5 %float_n1_5 %float_n1_5
%_ptr_Function_v4float = OpTypePointer Function %v4float
%__modf_result_vec4 = OpTypeStruct %v4float %v4float
%_ptr_Function___modf_result_vec4 = OpTypePointer Function %__modf_result_vec4
%23 = OpConstantNull %__modf_result_vec4
%24 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
%modf_4bfced = OpFunction %void None %9
%12 = OpLabel
%arg_0 = OpVariable %_ptr_Function_v4float Function %5
@@ -67,12 +68,12 @@
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %9
- %32 = OpLabel
- %33 = OpFunctionCall %void %modf_4bfced
+ %33 = OpLabel
+ %34 = OpFunctionCall %void %modf_4bfced
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %9
- %35 = OpLabel
- %36 = OpFunctionCall %void %modf_4bfced
+ %36 = OpLabel
+ %37 = OpFunctionCall %void %modf_4bfced
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.wgsl b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.wgsl
index cff6f57..cc7e11d 100644
--- a/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/modf/4bfced.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
fn modf_4bfced() {
- var arg_0 = vec4<f32>(1.0f);
+ var arg_0 = vec4<f32>(-(1.5f));
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/5ea256.wgsl b/test/tint/builtins/gen/var/modf/5ea256.wgsl
index f85cfeb..f623263 100644
--- a/test/tint/builtins/gen/var/modf/5ea256.wgsl
+++ b/test/tint/builtins/gen/var/modf/5ea256.wgsl
@@ -23,7 +23,7 @@
// fn modf(vec<3, f32>) -> __modf_result_vec<3, f32>
fn modf_5ea256() {
- var arg_0 = vec3<f32>(1.f);
+ var arg_0 = vec3<f32>(-1.5f);
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.dxc.hlsl
index ffd6a12..96db57b 100644
--- a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.dxc.hlsl
@@ -9,7 +9,7 @@
}
void modf_5ea256() {
- float3 arg_0 = (1.0f).xxx;
+ float3 arg_0 = (-1.5f).xxx;
modf_result_vec3 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.fxc.hlsl
index ffd6a12..96db57b 100644
--- a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.fxc.hlsl
@@ -9,7 +9,7 @@
}
void modf_5ea256() {
- float3 arg_0 = (1.0f).xxx;
+ float3 arg_0 = (-1.5f).xxx;
modf_result_vec3 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.glsl b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.glsl
index 021cf81..9c050b4 100644
--- a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.glsl
@@ -13,7 +13,7 @@
void modf_5ea256() {
- vec3 arg_0 = vec3(1.0f);
+ vec3 arg_0 = vec3(-1.5f);
modf_result_vec3 res = tint_modf(arg_0);
}
@@ -46,7 +46,7 @@
void modf_5ea256() {
- vec3 arg_0 = vec3(1.0f);
+ vec3 arg_0 = vec3(-1.5f);
modf_result_vec3 res = tint_modf(arg_0);
}
@@ -73,7 +73,7 @@
void modf_5ea256() {
- vec3 arg_0 = vec3(1.0f);
+ vec3 arg_0 = vec3(-1.5f);
modf_result_vec3 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.msl b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.msl
index fa5d28a..f5a7183 100644
--- a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.msl
@@ -13,7 +13,7 @@
}
void modf_5ea256() {
- float3 arg_0 = float3(1.0f);
+ float3 arg_0 = float3(-1.5f);
modf_result_vec3 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.spvasm b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.spvasm
index 1587506..42c66e4 100644
--- a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 39
+; Bound: 40
; Schema: 0
OpCapability Shader
%21 = OpExtInstImport "GLSL.std.450"
@@ -38,14 +38,15 @@
%void = OpTypeVoid
%9 = OpTypeFunction %void
%v3float = OpTypeVector %float 3
- %float_1 = OpConstant %float 1
- %15 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+ %float_n1_5 = OpConstant %float -1.5
+ %15 = OpConstantComposite %v3float %float_n1_5 %float_n1_5 %float_n1_5
%_ptr_Function_v3float = OpTypePointer Function %v3float
%18 = OpConstantNull %v3float
%__modf_result_vec3 = OpTypeStruct %v3float %v3float
%_ptr_Function___modf_result_vec3 = OpTypePointer Function %__modf_result_vec3
%25 = OpConstantNull %__modf_result_vec3
%26 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
%modf_5ea256 = OpFunction %void None %9
%12 = OpLabel
%arg_0 = OpVariable %_ptr_Function_v3float Function %18
@@ -69,12 +70,12 @@
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %9
- %34 = OpLabel
- %35 = OpFunctionCall %void %modf_5ea256
+ %35 = OpLabel
+ %36 = OpFunctionCall %void %modf_5ea256
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %9
- %37 = OpLabel
- %38 = OpFunctionCall %void %modf_5ea256
+ %38 = OpLabel
+ %39 = OpFunctionCall %void %modf_5ea256
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.wgsl b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.wgsl
index 5180d14..9d05d93 100644
--- a/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/modf/5ea256.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
fn modf_5ea256() {
- var arg_0 = vec3<f32>(1.0f);
+ var arg_0 = vec3<f32>(-(1.5f));
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl
index f3dea08..3b65908 100644
--- a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl
+++ b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl
@@ -25,7 +25,7 @@
// fn modf(f16) -> __modf_result<f16>
fn modf_8dbbbf() {
- var arg_0 = 1.h;
+ var arg_0 = -1.5h;
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.dxc.hlsl
index 04d940b..ef0eb00 100644
--- a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.dxc.hlsl
@@ -9,7 +9,7 @@
}
void modf_8dbbbf() {
- float16_t arg_0 = float16_t(1.0h);
+ float16_t arg_0 = float16_t(-1.5h);
modf_result_f16 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.glsl b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.glsl
index d6b8976..db8b53a 100644
--- a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.glsl
@@ -14,7 +14,7 @@
void modf_8dbbbf() {
- float16_t arg_0 = 1.0hf;
+ float16_t arg_0 = -1.5hf;
modf_result_f16 res = tint_modf(arg_0);
}
@@ -48,7 +48,7 @@
void modf_8dbbbf() {
- float16_t arg_0 = 1.0hf;
+ float16_t arg_0 = -1.5hf;
modf_result_f16 res = tint_modf(arg_0);
}
@@ -76,7 +76,7 @@
void modf_8dbbbf() {
- float16_t arg_0 = 1.0hf;
+ float16_t arg_0 = -1.5hf;
modf_result_f16 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.msl b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.msl
index 58a30ca..c3b9f88 100644
--- a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.msl
@@ -13,7 +13,7 @@
}
void modf_8dbbbf() {
- half arg_0 = 1.0h;
+ half arg_0 = -1.5h;
modf_result_f16 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.spvasm b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.spvasm
index 28b19b0..5714e2f 100644
--- a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.spvasm
@@ -42,7 +42,7 @@
%void = OpTypeVoid
%9 = OpTypeFunction %void
%half = OpTypeFloat 16
-%half_0x1p_0 = OpConstant %half 0x1p+0
+%half_n0x1_8p_0 = OpConstant %half -0x1.8p+0
%_ptr_Function_half = OpTypePointer Function %half
%17 = OpConstantNull %half
%__modf_result_f16 = OpTypeStruct %half %half
@@ -54,7 +54,7 @@
%12 = OpLabel
%arg_0 = OpVariable %_ptr_Function_half Function %17
%res = OpVariable %_ptr_Function___modf_result_f16 Function %24
- OpStore %arg_0 %half_0x1p_0
+ OpStore %arg_0 %half_n0x1_8p_0
%21 = OpLoad %half %arg_0
%18 = OpExtInst %__modf_result_f16 %20 ModfStruct %21
OpStore %res %18
diff --git a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.wgsl b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.wgsl
index d0815e7..e92ab55 100644
--- a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl.expected.wgsl
@@ -1,7 +1,7 @@
enable f16;
fn modf_8dbbbf() {
- var arg_0 = 1.0h;
+ var arg_0 = -(1.5h);
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/995934.wgsl b/test/tint/builtins/gen/var/modf/995934.wgsl
index 463335e..fa33c70 100644
--- a/test/tint/builtins/gen/var/modf/995934.wgsl
+++ b/test/tint/builtins/gen/var/modf/995934.wgsl
@@ -25,7 +25,7 @@
// fn modf(vec<4, f16>) -> __modf_result_vec<4, f16>
fn modf_995934() {
- var arg_0 = vec4<f16>(1.h);
+ var arg_0 = vec4<f16>(-1.5h);
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/995934.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/modf/995934.wgsl.expected.dxc.hlsl
index 26c46db..82b85d3 100644
--- a/test/tint/builtins/gen/var/modf/995934.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/modf/995934.wgsl.expected.dxc.hlsl
@@ -9,7 +9,7 @@
}
void modf_995934() {
- vector<float16_t, 4> arg_0 = (float16_t(1.0h)).xxxx;
+ vector<float16_t, 4> arg_0 = (float16_t(-1.5h)).xxxx;
modf_result_vec4_f16 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/995934.wgsl.expected.glsl b/test/tint/builtins/gen/var/modf/995934.wgsl.expected.glsl
index 71659e8..49bd43b 100644
--- a/test/tint/builtins/gen/var/modf/995934.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/modf/995934.wgsl.expected.glsl
@@ -14,7 +14,7 @@
void modf_995934() {
- f16vec4 arg_0 = f16vec4(1.0hf);
+ f16vec4 arg_0 = f16vec4(-1.5hf);
modf_result_vec4_f16 res = tint_modf(arg_0);
}
@@ -48,7 +48,7 @@
void modf_995934() {
- f16vec4 arg_0 = f16vec4(1.0hf);
+ f16vec4 arg_0 = f16vec4(-1.5hf);
modf_result_vec4_f16 res = tint_modf(arg_0);
}
@@ -76,7 +76,7 @@
void modf_995934() {
- f16vec4 arg_0 = f16vec4(1.0hf);
+ f16vec4 arg_0 = f16vec4(-1.5hf);
modf_result_vec4_f16 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/995934.wgsl.expected.msl b/test/tint/builtins/gen/var/modf/995934.wgsl.expected.msl
index 1b5ebe7..4e6cf8a 100644
--- a/test/tint/builtins/gen/var/modf/995934.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/modf/995934.wgsl.expected.msl
@@ -13,7 +13,7 @@
}
void modf_995934() {
- half4 arg_0 = half4(1.0h);
+ half4 arg_0 = half4(-1.5h);
modf_result_vec4_f16 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/995934.wgsl.expected.spvasm b/test/tint/builtins/gen/var/modf/995934.wgsl.expected.spvasm
index 94260c897a..440faff 100644
--- a/test/tint/builtins/gen/var/modf/995934.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/modf/995934.wgsl.expected.spvasm
@@ -43,8 +43,8 @@
%9 = OpTypeFunction %void
%half = OpTypeFloat 16
%v4half = OpTypeVector %half 4
-%half_0x1p_0 = OpConstant %half 0x1p+0
- %16 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0
+%half_n0x1_8p_0 = OpConstant %half -0x1.8p+0
+ %16 = OpConstantComposite %v4half %half_n0x1_8p_0 %half_n0x1_8p_0 %half_n0x1_8p_0 %half_n0x1_8p_0
%_ptr_Function_v4half = OpTypePointer Function %v4half
%19 = OpConstantNull %v4half
%__modf_result_vec4_f16 = OpTypeStruct %v4half %v4half
diff --git a/test/tint/builtins/gen/var/modf/995934.wgsl.expected.wgsl b/test/tint/builtins/gen/var/modf/995934.wgsl.expected.wgsl
index 317ae85..e5995a9 100644
--- a/test/tint/builtins/gen/var/modf/995934.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/modf/995934.wgsl.expected.wgsl
@@ -1,7 +1,7 @@
enable f16;
fn modf_995934() {
- var arg_0 = vec4<f16>(1.0h);
+ var arg_0 = vec4<f16>(-(1.5h));
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/a545b9.wgsl b/test/tint/builtins/gen/var/modf/a545b9.wgsl
index 51895fd..c559912 100644
--- a/test/tint/builtins/gen/var/modf/a545b9.wgsl
+++ b/test/tint/builtins/gen/var/modf/a545b9.wgsl
@@ -25,7 +25,7 @@
// fn modf(vec<2, f16>) -> __modf_result_vec<2, f16>
fn modf_a545b9() {
- var arg_0 = vec2<f16>(1.h);
+ var arg_0 = vec2<f16>(-1.5h);
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.dxc.hlsl
index 47cc509..30290bc 100644
--- a/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.dxc.hlsl
@@ -9,7 +9,7 @@
}
void modf_a545b9() {
- vector<float16_t, 2> arg_0 = (float16_t(1.0h)).xx;
+ vector<float16_t, 2> arg_0 = (float16_t(-1.5h)).xx;
modf_result_vec2_f16 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.glsl b/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.glsl
index 86f166d..56974b5 100644
--- a/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.glsl
@@ -14,7 +14,7 @@
void modf_a545b9() {
- f16vec2 arg_0 = f16vec2(1.0hf);
+ f16vec2 arg_0 = f16vec2(-1.5hf);
modf_result_vec2_f16 res = tint_modf(arg_0);
}
@@ -48,7 +48,7 @@
void modf_a545b9() {
- f16vec2 arg_0 = f16vec2(1.0hf);
+ f16vec2 arg_0 = f16vec2(-1.5hf);
modf_result_vec2_f16 res = tint_modf(arg_0);
}
@@ -76,7 +76,7 @@
void modf_a545b9() {
- f16vec2 arg_0 = f16vec2(1.0hf);
+ f16vec2 arg_0 = f16vec2(-1.5hf);
modf_result_vec2_f16 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.msl b/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.msl
index 1cf4d62..2a59af2 100644
--- a/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.msl
@@ -13,7 +13,7 @@
}
void modf_a545b9() {
- half2 arg_0 = half2(1.0h);
+ half2 arg_0 = half2(-1.5h);
modf_result_vec2_f16 res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.spvasm b/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.spvasm
index da154ba..6bd4d3e 100644
--- a/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.spvasm
@@ -43,8 +43,8 @@
%9 = OpTypeFunction %void
%half = OpTypeFloat 16
%v2half = OpTypeVector %half 2
-%half_0x1p_0 = OpConstant %half 0x1p+0
- %16 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0
+%half_n0x1_8p_0 = OpConstant %half -0x1.8p+0
+ %16 = OpConstantComposite %v2half %half_n0x1_8p_0 %half_n0x1_8p_0
%_ptr_Function_v2half = OpTypePointer Function %v2half
%19 = OpConstantNull %v2half
%__modf_result_vec2_f16 = OpTypeStruct %v2half %v2half
diff --git a/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.wgsl b/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.wgsl
index dfb8c47..c6e6265 100644
--- a/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/modf/a545b9.wgsl.expected.wgsl
@@ -1,7 +1,7 @@
enable f16;
fn modf_a545b9() {
- var arg_0 = vec2<f16>(1.0h);
+ var arg_0 = vec2<f16>(-(1.5h));
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl
index 60ab469..7fa9467 100644
--- a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl
+++ b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl
@@ -23,7 +23,7 @@
// fn modf(f32) -> __modf_result<f32>
fn modf_bbf7f7() {
- var arg_0 = 1.f;
+ var arg_0 = -1.5f;
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.dxc.hlsl
index 39adf2d..26d23b2 100644
--- a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.dxc.hlsl
@@ -9,7 +9,7 @@
}
void modf_bbf7f7() {
- float arg_0 = 1.0f;
+ float arg_0 = -1.5f;
modf_result res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.fxc.hlsl
index 39adf2d..26d23b2 100644
--- a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.fxc.hlsl
@@ -9,7 +9,7 @@
}
void modf_bbf7f7() {
- float arg_0 = 1.0f;
+ float arg_0 = -1.5f;
modf_result res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.glsl b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.glsl
index 4907635..12f70df 100644
--- a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.glsl
@@ -13,7 +13,7 @@
void modf_bbf7f7() {
- float arg_0 = 1.0f;
+ float arg_0 = -1.5f;
modf_result res = tint_modf(arg_0);
}
@@ -46,7 +46,7 @@
void modf_bbf7f7() {
- float arg_0 = 1.0f;
+ float arg_0 = -1.5f;
modf_result res = tint_modf(arg_0);
}
@@ -73,7 +73,7 @@
void modf_bbf7f7() {
- float arg_0 = 1.0f;
+ float arg_0 = -1.5f;
modf_result res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.msl b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.msl
index 2ef6251..3e4accd 100644
--- a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.msl
@@ -13,7 +13,7 @@
}
void modf_bbf7f7() {
- float arg_0 = 1.0f;
+ float arg_0 = -1.5f;
modf_result res = tint_modf(arg_0);
}
diff --git a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.spvasm b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.spvasm
index 97fe510..6ffbfde 100644
--- a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 36
+; Bound: 37
; Schema: 0
OpCapability Shader
%18 = OpExtInstImport "GLSL.std.450"
@@ -37,17 +37,18 @@
%vertex_point_size = OpVariable %_ptr_Output_float Output %8
%void = OpTypeVoid
%9 = OpTypeFunction %void
- %float_1 = OpConstant %float 1
+ %float_n1_5 = OpConstant %float -1.5
%_ptr_Function_float = OpTypePointer Function %float
%__modf_result = OpTypeStruct %float %float
%_ptr_Function___modf_result = OpTypePointer Function %__modf_result
%22 = OpConstantNull %__modf_result
%23 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
%modf_bbf7f7 = OpFunction %void None %9
%12 = OpLabel
%arg_0 = OpVariable %_ptr_Function_float Function %8
%res = OpVariable %_ptr_Function___modf_result Function %22
- OpStore %arg_0 %float_1
+ OpStore %arg_0 %float_n1_5
%19 = OpLoad %float %arg_0
%16 = OpExtInst %__modf_result %18 ModfStruct %19
OpStore %res %16
@@ -66,12 +67,12 @@
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %9
- %31 = OpLabel
- %32 = OpFunctionCall %void %modf_bbf7f7
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %modf_bbf7f7
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %9
- %34 = OpLabel
- %35 = OpFunctionCall %void %modf_bbf7f7
+ %35 = OpLabel
+ %36 = OpFunctionCall %void %modf_bbf7f7
OpReturn
OpFunctionEnd
diff --git a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.wgsl b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.wgsl
index 0ebdc4b..418a00d 100644
--- a/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/var/modf/bbf7f7.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
fn modf_bbf7f7() {
- var arg_0 = 1.0f;
+ var arg_0 = -(1.5f);
var res = modf(arg_0);
}
diff --git a/test/tint/builtins/modf.wgsl.expected.dxc.hlsl b/test/tint/builtins/modf.wgsl.expected.dxc.hlsl
index 6ea70f3..2411190 100644
--- a/test/tint/builtins/modf.wgsl.expected.dxc.hlsl
+++ b/test/tint/builtins/modf.wgsl.expected.dxc.hlsl
@@ -2,15 +2,9 @@
float fract;
float whole;
};
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
[numthreads(1, 1, 1)]
void main() {
- const modf_result res = tint_modf(1.230000019f);
+ const modf_result res = {0.230000019f, 1.0f};
const float fract = res.fract;
const float whole = res.whole;
return;
diff --git a/test/tint/builtins/modf.wgsl.expected.fxc.hlsl b/test/tint/builtins/modf.wgsl.expected.fxc.hlsl
index 6ea70f3..2411190 100644
--- a/test/tint/builtins/modf.wgsl.expected.fxc.hlsl
+++ b/test/tint/builtins/modf.wgsl.expected.fxc.hlsl
@@ -2,15 +2,9 @@
float fract;
float whole;
};
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
[numthreads(1, 1, 1)]
void main() {
- const modf_result res = tint_modf(1.230000019f);
+ const modf_result res = {0.230000019f, 1.0f};
const float fract = res.fract;
const float whole = res.whole;
return;
diff --git a/test/tint/builtins/modf.wgsl.expected.glsl b/test/tint/builtins/modf.wgsl.expected.glsl
index 00fa668..e59dfb0 100644
--- a/test/tint/builtins/modf.wgsl.expected.glsl
+++ b/test/tint/builtins/modf.wgsl.expected.glsl
@@ -5,15 +5,9 @@
float whole;
};
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
void tint_symbol() {
- modf_result res = tint_modf(1.230000019f);
+ modf_result res = modf_result(0.230000019f, 1.0f);
float tint_symbol_1 = res.fract;
float whole = res.whole;
}
diff --git a/test/tint/builtins/modf.wgsl.expected.msl b/test/tint/builtins/modf.wgsl.expected.msl
index 108250d..6c73b24 100644
--- a/test/tint/builtins/modf.wgsl.expected.msl
+++ b/test/tint/builtins/modf.wgsl.expected.msl
@@ -6,14 +6,8 @@
float fract;
float whole;
};
-modf_result tint_modf(float param_0) {
- modf_result result;
- result.fract = modf(param_0, result.whole);
- return result;
-}
-
kernel void tint_symbol() {
- modf_result const res = tint_modf(1.230000019f);
+ modf_result const res = modf_result{.fract=0.230000019f, .whole=1.0f};
float const fract = res.fract;
float const whole = res.whole;
return;
diff --git a/test/tint/builtins/modf.wgsl.expected.spvasm b/test/tint/builtins/modf.wgsl.expected.spvasm
index ebbc766..cca23fc 100644
--- a/test/tint/builtins/modf.wgsl.expected.spvasm
+++ b/test/tint/builtins/modf.wgsl.expected.spvasm
@@ -4,7 +4,6 @@
; Bound: 12
; Schema: 0
OpCapability Shader
- %8 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
@@ -18,11 +17,12 @@
%1 = OpTypeFunction %void
%float = OpTypeFloat 32
%__modf_result = OpTypeStruct %float %float
-%float_1_23000002 = OpConstant %float 1.23000002
+%float_0_230000019 = OpConstant %float 0.230000019
+ %float_1 = OpConstant %float 1
+ %9 = OpConstantComposite %__modf_result %float_0_230000019 %float_1
%main = OpFunction %void None %1
%4 = OpLabel
- %5 = OpExtInst %__modf_result %8 ModfStruct %float_1_23000002
- %10 = OpCompositeExtract %float %5 0
- %11 = OpCompositeExtract %float %5 1
+ %10 = OpCompositeExtract %float %9 0
+ %11 = OpCompositeExtract %float %9 1
OpReturn
OpFunctionEnd
diff --git a/test/tint/var/inferred/global.wgsl.expected.dxc.hlsl b/test/tint/var/inferred/global.wgsl.expected.dxc.hlsl
index 09a7a21..ec015d8 100644
--- a/test/tint/var/inferred/global.wgsl.expected.dxc.hlsl
+++ b/test/tint/var/inferred/global.wgsl.expected.dxc.hlsl
@@ -8,7 +8,8 @@
static int3 v4 = (1).xxx;
static uint3 v5 = uint3(1u, 2u, 3u);
static float3 v6 = float3(1.0f, 2.0f, 3.0f);
-static MyStruct v7 = {1.0f};
+const MyStruct c = {1.0f};
+static MyStruct v7 = c;
static float v8[10] = (float[10])0;
static int v9 = 0;
static uint v10 = 0u;
diff --git a/test/tint/var/inferred/global.wgsl.expected.fxc.hlsl b/test/tint/var/inferred/global.wgsl.expected.fxc.hlsl
index 09a7a21..ec015d8 100644
--- a/test/tint/var/inferred/global.wgsl.expected.fxc.hlsl
+++ b/test/tint/var/inferred/global.wgsl.expected.fxc.hlsl
@@ -8,7 +8,8 @@
static int3 v4 = (1).xxx;
static uint3 v5 = uint3(1u, 2u, 3u);
static float3 v6 = float3(1.0f, 2.0f, 3.0f);
-static MyStruct v7 = {1.0f};
+const MyStruct c = {1.0f};
+static MyStruct v7 = c;
static float v8[10] = (float[10])0;
static int v9 = 0;
static uint v10 = 0u;
