From 329dfd7cbd2df6025c697e3f00acbf91f65bd0fd Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Wed, 23 Nov 2022 00:05:05 +0000 Subject: [PATCH] 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 Kokoro: Kokoro Commit-Queue: Ben Clayton --- src/tint/intrinsics.def | 4 +- src/tint/resolver/const_eval.cc | 34 ++++ src/tint/resolver/const_eval.h | 9 ++ src/tint/resolver/const_eval_builtin_test.cc | 51 ++++++ src/tint/resolver/intrinsic_table.inl | 8 +- src/tint/resolver/validator.cc | 10 +- src/tint/writer/glsl/generator_impl.cc | 21 ++- src/tint/writer/glsl/generator_impl.h | 9 +- .../glsl/generator_impl_builtin_test.cc | 146 ++++++++++++++++-- src/tint/writer/hlsl/generator_impl.cc | 83 ++++++---- src/tint/writer/hlsl/generator_impl.h | 15 +- .../hlsl/generator_impl_builtin_test.cc | 138 +++++++++++++++-- src/tint/writer/msl/generator_impl.cc | 23 ++- src/tint/writer/msl/generator_impl.h | 9 +- .../writer/msl/generator_impl_builtin_test.cc | 128 +++++++++++++-- src/tint/writer/spirv/builder_builtin_test.cc | 98 +++++++++++- .../chromium/1236161.wgsl.expected.dxc.hlsl | 12 +- .../chromium/1236161.wgsl.expected.fxc.hlsl | 12 +- .../bug/chromium/1236161.wgsl.expected.glsl | 14 +- .../bug/chromium/1236161.wgsl.expected.msl | 13 +- .../bug/chromium/1236161.wgsl.expected.spvasm | 11 +- .../builtins/gen/literal/modf/2d50da.wgsl | 2 +- .../modf/2d50da.wgsl.expected.dxc.hlsl | 8 +- .../modf/2d50da.wgsl.expected.fxc.hlsl | 8 +- .../literal/modf/2d50da.wgsl.expected.glsl | 24 +-- .../gen/literal/modf/2d50da.wgsl.expected.msl | 8 +- .../literal/modf/2d50da.wgsl.expected.spvasm | 40 ++--- .../literal/modf/2d50da.wgsl.expected.wgsl | 2 +- .../builtins/gen/literal/modf/45005f.wgsl | 2 +- .../modf/45005f.wgsl.expected.dxc.hlsl | 8 +- .../literal/modf/45005f.wgsl.expected.glsl | 24 +-- .../gen/literal/modf/45005f.wgsl.expected.msl | 8 +- .../literal/modf/45005f.wgsl.expected.spvasm | 39 ++--- .../literal/modf/45005f.wgsl.expected.wgsl | 2 +- .../builtins/gen/literal/modf/4bfced.wgsl | 2 +- .../modf/4bfced.wgsl.expected.dxc.hlsl | 8 +- .../modf/4bfced.wgsl.expected.fxc.hlsl | 8 +- .../literal/modf/4bfced.wgsl.expected.glsl | 24 +-- .../gen/literal/modf/4bfced.wgsl.expected.msl | 8 +- .../literal/modf/4bfced.wgsl.expected.spvasm | 40 ++--- .../literal/modf/4bfced.wgsl.expected.wgsl | 2 +- .../builtins/gen/literal/modf/5ea256.wgsl | 2 +- .../modf/5ea256.wgsl.expected.dxc.hlsl | 8 +- .../modf/5ea256.wgsl.expected.fxc.hlsl | 8 +- .../literal/modf/5ea256.wgsl.expected.glsl | 24 +-- .../gen/literal/modf/5ea256.wgsl.expected.msl | 8 +- .../literal/modf/5ea256.wgsl.expected.spvasm | 40 ++--- .../literal/modf/5ea256.wgsl.expected.wgsl | 2 +- .../builtins/gen/literal/modf/8dbbbf.wgsl | 2 +- .../modf/8dbbbf.wgsl.expected.dxc.hlsl | 8 +- .../literal/modf/8dbbbf.wgsl.expected.glsl | 24 +-- .../gen/literal/modf/8dbbbf.wgsl.expected.msl | 8 +- .../literal/modf/8dbbbf.wgsl.expected.spvasm | 8 +- .../literal/modf/8dbbbf.wgsl.expected.wgsl | 2 +- .../builtins/gen/literal/modf/995934.wgsl | 2 +- .../modf/995934.wgsl.expected.dxc.hlsl | 8 +- .../literal/modf/995934.wgsl.expected.glsl | 24 +-- .../gen/literal/modf/995934.wgsl.expected.msl | 8 +- .../literal/modf/995934.wgsl.expected.spvasm | 39 ++--- .../literal/modf/995934.wgsl.expected.wgsl | 2 +- .../builtins/gen/literal/modf/a545b9.wgsl | 2 +- .../modf/a545b9.wgsl.expected.dxc.hlsl | 8 +- .../literal/modf/a545b9.wgsl.expected.glsl | 24 +-- .../gen/literal/modf/a545b9.wgsl.expected.msl | 8 +- .../literal/modf/a545b9.wgsl.expected.spvasm | 39 ++--- .../literal/modf/a545b9.wgsl.expected.wgsl | 2 +- .../builtins/gen/literal/modf/bbf7f7.wgsl | 2 +- .../modf/bbf7f7.wgsl.expected.dxc.hlsl | 8 +- .../modf/bbf7f7.wgsl.expected.fxc.hlsl | 8 +- .../literal/modf/bbf7f7.wgsl.expected.glsl | 24 +-- .../gen/literal/modf/bbf7f7.wgsl.expected.msl | 8 +- .../literal/modf/bbf7f7.wgsl.expected.spvasm | 19 +-- .../literal/modf/bbf7f7.wgsl.expected.wgsl | 2 +- test/tint/builtins/gen/var/modf/2d50da.wgsl | 2 +- .../var/modf/2d50da.wgsl.expected.dxc.hlsl | 2 +- .../var/modf/2d50da.wgsl.expected.fxc.hlsl | 2 +- .../gen/var/modf/2d50da.wgsl.expected.glsl | 6 +- .../gen/var/modf/2d50da.wgsl.expected.msl | 2 +- .../gen/var/modf/2d50da.wgsl.expected.spvasm | 15 +- .../gen/var/modf/2d50da.wgsl.expected.wgsl | 2 +- test/tint/builtins/gen/var/modf/45005f.wgsl | 2 +- .../var/modf/45005f.wgsl.expected.dxc.hlsl | 2 +- .../gen/var/modf/45005f.wgsl.expected.glsl | 6 +- .../gen/var/modf/45005f.wgsl.expected.msl | 2 +- .../gen/var/modf/45005f.wgsl.expected.spvasm | 4 +- .../gen/var/modf/45005f.wgsl.expected.wgsl | 2 +- test/tint/builtins/gen/var/modf/4bfced.wgsl | 2 +- .../var/modf/4bfced.wgsl.expected.dxc.hlsl | 2 +- .../var/modf/4bfced.wgsl.expected.fxc.hlsl | 2 +- .../gen/var/modf/4bfced.wgsl.expected.glsl | 6 +- .../gen/var/modf/4bfced.wgsl.expected.msl | 2 +- .../gen/var/modf/4bfced.wgsl.expected.spvasm | 15 +- .../gen/var/modf/4bfced.wgsl.expected.wgsl | 2 +- test/tint/builtins/gen/var/modf/5ea256.wgsl | 2 +- .../var/modf/5ea256.wgsl.expected.dxc.hlsl | 2 +- .../var/modf/5ea256.wgsl.expected.fxc.hlsl | 2 +- .../gen/var/modf/5ea256.wgsl.expected.glsl | 6 +- .../gen/var/modf/5ea256.wgsl.expected.msl | 2 +- .../gen/var/modf/5ea256.wgsl.expected.spvasm | 15 +- .../gen/var/modf/5ea256.wgsl.expected.wgsl | 2 +- test/tint/builtins/gen/var/modf/8dbbbf.wgsl | 2 +- .../var/modf/8dbbbf.wgsl.expected.dxc.hlsl | 2 +- .../gen/var/modf/8dbbbf.wgsl.expected.glsl | 6 +- .../gen/var/modf/8dbbbf.wgsl.expected.msl | 2 +- .../gen/var/modf/8dbbbf.wgsl.expected.spvasm | 4 +- .../gen/var/modf/8dbbbf.wgsl.expected.wgsl | 2 +- test/tint/builtins/gen/var/modf/995934.wgsl | 2 +- .../var/modf/995934.wgsl.expected.dxc.hlsl | 2 +- .../gen/var/modf/995934.wgsl.expected.glsl | 6 +- .../gen/var/modf/995934.wgsl.expected.msl | 2 +- .../gen/var/modf/995934.wgsl.expected.spvasm | 4 +- .../gen/var/modf/995934.wgsl.expected.wgsl | 2 +- test/tint/builtins/gen/var/modf/a545b9.wgsl | 2 +- .../var/modf/a545b9.wgsl.expected.dxc.hlsl | 2 +- .../gen/var/modf/a545b9.wgsl.expected.glsl | 6 +- .../gen/var/modf/a545b9.wgsl.expected.msl | 2 +- .../gen/var/modf/a545b9.wgsl.expected.spvasm | 4 +- .../gen/var/modf/a545b9.wgsl.expected.wgsl | 2 +- test/tint/builtins/gen/var/modf/bbf7f7.wgsl | 2 +- .../var/modf/bbf7f7.wgsl.expected.dxc.hlsl | 2 +- .../var/modf/bbf7f7.wgsl.expected.fxc.hlsl | 2 +- .../gen/var/modf/bbf7f7.wgsl.expected.glsl | 6 +- .../gen/var/modf/bbf7f7.wgsl.expected.msl | 2 +- .../gen/var/modf/bbf7f7.wgsl.expected.spvasm | 15 +- .../gen/var/modf/bbf7f7.wgsl.expected.wgsl | 2 +- .../tint/builtins/modf.wgsl.expected.dxc.hlsl | 8 +- .../tint/builtins/modf.wgsl.expected.fxc.hlsl | 8 +- test/tint/builtins/modf.wgsl.expected.glsl | 8 +- test/tint/builtins/modf.wgsl.expected.msl | 8 +- test/tint/builtins/modf.wgsl.expected.spvasm | 10 +- .../inferred/global.wgsl.expected.dxc.hlsl | 3 +- .../inferred/global.wgsl.expected.fxc.hlsl | 3 +- 132 files changed, 966 insertions(+), 774 deletions(-) diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def index 4bb101ee89..da58aeaed6 100644 --- a/src/tint/intrinsics.def +++ b/src/tint/intrinsics.def @@ -504,8 +504,8 @@ fn log2(vec) -> vec fn mix(T, T, T) -> T fn mix(vec, vec, vec) -> vec fn mix(vec, vec, T) -> vec -fn modf(T) -> __modf_result -fn modf(vec) -> __modf_result_vec +@const fn modf(@test_value(-1.5) T) -> __modf_result +@const fn modf(@test_value(-1.5) vec) -> __modf_result_vec fn normalize(vec) -> vec @const fn pack2x16float(vec2) -> u32 @const fn pack2x16snorm(vec2) -> u32 diff --git a/src/tint/resolver/const_eval.cc b/src/tint/resolver/const_eval.cc index 6cdf8f4e16..c5b51c0a38 100644 --- a/src/tint/resolver/const_eval.cc +++ b/src/tint/resolver/const_eval.cc @@ -2266,6 +2266,40 @@ ConstEval::Result ConstEval::min(const sem::Type* ty, return TransformElements(builder, ty, transform, args[0], args[1]); } +ConstEval::Result ConstEval::modf(const sem::Type* ty, + utils::VectorRef 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 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 args, const Source& source) { diff --git a/src/tint/resolver/const_eval.h b/src/tint/resolver/const_eval.h index 496e138ecd..f290adb89c 100644 --- a/src/tint/resolver/const_eval.h +++ b/src/tint/resolver/const_eval.h @@ -628,6 +628,15 @@ class ConstEval { utils::VectorRef 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 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 13bbac50ca..eccd711684 100644 --- a/src/tint/resolver/const_eval_builtin_test.cc +++ b/src/tint/resolver/const_eval_builtin_test.cc @@ -98,6 +98,11 @@ static Case C(std::initializer_list args, Types result) { return Case{utils::Vector{args}, utils::Vector{std::move(result)}}; } +/// Creates a Case with Values for args and result +static Case C(std::initializer_list args, std::initializer_list results) { + return Case{utils::Vector{args}, utils::Vector{results}}; +} + /// Convenience overload that creates a Case with just scalars static Case C(std::initializer_list sargs, ScalarTypes sresult) { utils::Vector args; @@ -109,6 +114,20 @@ static Case C(std::initializer_list sargs, ScalarTypes sresult) { return Case{std::move(args), utils::Vector{std::move(result)}}; } +/// Creates a Case with Values for args and result +static Case C(std::initializer_list sargs, + std::initializer_list sresults) { + utils::Vector args; + for (auto& sa : sargs) { + std::visit([&](auto&& v) { return args.Push(Val(v)); }, sa); + } + utils::Vector 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 args, std::string err) { return Case{utils::Vector{args}, std::move(err)}; @@ -1290,6 +1309,38 @@ INSTANTIATE_TEST_SUITE_P( // MinCases(), MinCases(), MinCases())))); +template +std::vector 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(), // + ModfCases())))); std::vector Pack4x8snormCases() { return { diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl index 774496018f..be0e4bb729 100644 --- a/src/tint/resolver/intrinsic_table.inl +++ b/src/tint/resolver/intrinsic_table.inl @@ -12851,7 +12851,7 @@ constexpr OverloadInfo kOverloads[] = { /* 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 @@ constexpr OverloadInfo kOverloads[] = { /* 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 @@ constexpr IntrinsicInfo kBuiltins[] = { }, { /* [53] */ - /* fn modf(T) -> __modf_result */ - /* fn modf(vec) -> __modf_result_vec */ + /* fn modf(@test_value(-1.5) T) -> __modf_result */ + /* fn modf(@test_value(-1.5) vec) -> __modf_result_vec */ /* num overloads */ 2, /* overloads */ &kOverloads[377], }, diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc index 0b20c9b31c..f7d6cb3a84 100644 --- a/src/tint/resolver/validator.cc +++ b/src/tint/resolver/validator.cc @@ -530,9 +530,13 @@ bool Validator::AddressSpaceLayout(const sem::Variable* var, } if (auto* str = var->Type()->UnwrapRef()->As()) { - 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 e716b1d3c7..86ba6579d2 100644 --- a/src/tint/writer/glsl/generator_impl.cc +++ b/src/tint/writer/glsl/generator_impl.cc @@ -939,9 +939,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out, 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())) { + if (!EmitStructType(&helpers_, builtin->ReturnType()->As())) { return false; } @@ -2373,10 +2371,12 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan 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::EmitTypeAndName(std::ostream& out, } 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 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) { 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 c34880df10..24d835741c 100644 --- a/src/tint/writer/glsl/generator_impl.h +++ b/src/tint/writer/glsl/generator_impl.h @@ -430,17 +430,12 @@ class GeneratorImpl : public TextGenerator { 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 17028aca8b..7d0a72e44e 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 @@ TEST_F(GlslGeneratorImplTest_Builtin, FMA_f16) { 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 @@ modf_result tint_modf(float param_0) { 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 @@ void main() { )"); } -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 @@ modf_result_f16 tint_modf(float16_t param_0) { 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 @@ void main() { )"); } -TEST_F(GlslGeneratorImplTest_Builtin, Modf_Vector_f32) { - auto* call = Call("modf", vec3()); - WrapInFunction(CallStmt(call)); +TEST_F(GlslGeneratorImplTest_Builtin, Runtime_Modf_Vector_f32) { + WrapInFunction(Decl(Let("f", vec3(1.5_f, 2.5_f, 3.5_f))), // + Decl(Let("v", Call("modf", "f")))); GeneratorImpl& gen = SanitizeAndBuild(); @@ -507,7 +509,8 @@ modf_result_vec3 tint_modf(vec3 param_0) { 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 @@ void main() { )"); } -TEST_F(GlslGeneratorImplTest_Builtin, Modf_Vector_f16) { +TEST_F(GlslGeneratorImplTest_Builtin, Runtime_Modf_Vector_f16) { Enable(ast::Extension::kF16); - auto* call = Call("modf", vec3()); - WrapInFunction(CallStmt(call)); + WrapInFunction(Decl(Let("f", vec3(1.5_h, 2.5_h, 3.5_h))), // + Decl(Let("v", Call("modf", "f")))); GeneratorImpl& gen = SanitizeAndBuild(); @@ -543,7 +546,118 @@ modf_result_vec3_f16 tint_modf(f16vec3 param_0) { 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(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(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 e3abc5175e..dcad1b5d1c 100644 --- a/src/tint/writer/hlsl/generator_impl.cc +++ b/src/tint/writer/hlsl/generator_impl.cc @@ -1651,9 +1651,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out, 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())) { + if (!EmitStructType(&helpers_, builtin->ReturnType()->As())) { return false; } @@ -2506,7 +2504,7 @@ bool GeneratorImpl::EmitCase(const ast::SwitchStatement* s, size_t case_idx) { 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::EmitDiscard(const ast::DiscardStatement*) { 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(stmt->Declaration())) { + is_variable_initializer = decl->variable->initializer == expr; + } + } + return EmitConstant(out, constant, is_variable_initializer); } } return Switch( @@ -3042,7 +3046,9 @@ bool GeneratorImpl::EmitEntryPointFunction(const ast::Function* func) { 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 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan 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 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan 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 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan 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 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan 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 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan 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::EmitTypeAndName(std::ostream& out, } 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 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) { 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 16b824b483..74a172d0e0 100644 --- a/src/tint/writer/hlsl/generator_impl.h +++ b/src/tint/writer/hlsl/generator_impl.h @@ -348,8 +348,12 @@ class GeneratorImpl : public TextGenerator { /// 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 @@ class GeneratorImpl : public TextGenerator { 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 b69e67148e..3e9b16973a 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 @@ TEST_F(HlslGeneratorImplTest_Builtin, Select_Vector) { 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 @@ modf_result tint_modf(float param_0) { [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 @@ modf_result_f16 tint_modf(float16_t param_0) { [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()); - WrapInFunction(CallStmt(call)); +TEST_F(HlslGeneratorImplTest_Builtin, Runtime_Modf_Vector_f32) { + WrapInFunction(Decl(Let("f", vec3(1.5_f, 2.5_f, 3.5_f))), // + Decl(Let("v", Call("modf", "f")))); GeneratorImpl& gen = SanitizeAndBuild(); @@ -452,17 +454,18 @@ modf_result_vec3 tint_modf(float3 param_0) { [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()); - WrapInFunction(CallStmt(call)); + WrapInFunction(Decl(Let("f", vec3(1.5_h, 2.5_h, 3.5_h))), // + Decl(Let("v", Call("modf", "f")))); GeneratorImpl& gen = SanitizeAndBuild(); @@ -479,7 +482,110 @@ modf_result_vec3_f16 tint_modf(vector param_0) { [numthreads(1, 1, 1)] void test_function() { - tint_modf((float16_t(0.0h)).xxx); + const vector f = vector(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(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(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 fract; + vector whole; +}; +[numthreads(1, 1, 1)] +void test_function() { + const modf_result_vec3_f16 v = {(float16_t(0.5h)).xxx, vector(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(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(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 534742e674..10efd595dd 100644 --- a/src/tint/writer/msl/generator_impl.cc +++ b/src/tint/writer/msl/generator_impl.cc @@ -907,9 +907,7 @@ bool GeneratorImpl::EmitAtomicCall(std::ostream& out, 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())) { + if (!EmitStructType(&helpers_, builtin->ReturnType()->As())) { return ""; } @@ -1748,7 +1746,11 @@ bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constan 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::EmitAddressSpace(std::ostream& out, ast::AddressSpace sc) { } 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 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) { 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 c3e1ac0486..829088dc37 100644 --- a/src/tint/writer/msl/generator_impl.h +++ b/src/tint/writer/msl/generator_impl.h @@ -328,17 +328,12 @@ class GeneratorImpl : public TextGenerator { /// @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 8974bd9bf9..c611d6b368 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 @@ TEST_F(MslGeneratorImplTest, WorkgroupBarrier) { 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 @@ modf_result tint_modf(float param_0) { } 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 @@ modf_result_f16 tint_modf(half param_0) { } 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()); - WrapInFunction(CallStmt(call)); +TEST_F(MslGeneratorImplTest, Runtime_Modf_Vector_f32) { + WrapInFunction(Decl(Let("f", vec3(1.5_f, 2.5_f, 3.5_f))), // + Decl(Let("v", Call("modf", "f")))); GeneratorImpl& gen = SanitizeAndBuild(); @@ -487,18 +489,19 @@ modf_result_vec3 tint_modf(float3 param_0) { } 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()); - WrapInFunction(CallStmt(call)); + WrapInFunction(Decl(Let("f", vec3(1.5_h, 2.5_h, 3.5_h))), // + Decl(Let("v", Call("modf", "f")))); GeneratorImpl& gen = SanitizeAndBuild(); @@ -518,7 +521,100 @@ modf_result_vec3_f16 tint_modf(half3 param_0) { } 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 + +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 + +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(1.5_f, 2.5_f, 3.5_f))))); + + GeneratorImpl& gen = SanitizeAndBuild(); + + ASSERT_TRUE(gen.Generate()) << gen.error(); + EXPECT_EQ(gen.result(), R"(#include + +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(1.5_h, 2.5_h, 3.5_h))))); + + GeneratorImpl& gen = SanitizeAndBuild(); + + ASSERT_TRUE(gen.Generate()) << gen.error(); + EXPECT_EQ(gen.result(), R"(#include + +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 d00b019cbe..c0995e840d 100644 --- a/src/tint/writer/spirv/builder_builtin_test.cc +++ b/src/tint/writer/spirv/builder_builtin_test.cc @@ -1630,7 +1630,7 @@ OpFunctionEnd EXPECT_EQ(got, expect); } -TEST_F(BuiltinBuilderTest, Call_Modf_f32) { +TEST_F(BuiltinBuilderTest, Runtime_Call_Modf_f32) { auto* vec = Var("vec", vec2(1_f, 2_f)); auto* expr = Call("modf", vec); Func("a_func", utils::Empty, ty.void_(), @@ -1682,7 +1682,7 @@ OpFunctionEnd Validate(b); } -TEST_F(BuiltinBuilderTest, Call_Modf_f16) { +TEST_F(BuiltinBuilderTest, Runtime_Call_Modf_f16) { Enable(ast::Extension::kF16); auto* vec = Var("vec", vec2(1_h, 2_h)); @@ -1740,6 +1740,100 @@ OpFunctionEnd Validate(b); } +TEST_F(BuiltinBuilderTest, Const_Call_Modf_f32) { + auto* expr = Call("modf", vec2(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(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(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 136105307d..f73a45a5b8 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 136105307d..f73a45a5b8 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 04734ec407..f603bf359d 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 8d7aa31d32..6acecbdca9 100644 --- a/test/tint/bug/chromium/1236161.wgsl.expected.msl +++ b/test/tint/bug/chromium/1236161.wgsl.expected.msl @@ -1,18 +1,7 @@ #include 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 0d9b5a8368..258d3c9b91 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 8d3d0ddfd7..e3212fc017 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(1.f)); + var res = modf(vec2(-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 2aba1e7707..0dc03d8f5e 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 @@ struct modf_result_vec2 { 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 2aba1e7707..0dc03d8f5e 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 @@ struct modf_result_vec2 { 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 540a580081..881c541e29 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 @@ struct modf_result_vec2 { 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 @@ struct modf_result_vec2 { 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 @@ struct modf_result_vec2 { 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 edccc81725..16a1941c03 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 @@ struct modf_result_vec2 { 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 bc7576901c..84f16891a6 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 d893b25e80..a267089435 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(1.0f)); + var res = modf(vec2(-(1.5f))); } @vertex diff --git a/test/tint/builtins/gen/literal/modf/45005f.wgsl b/test/tint/builtins/gen/literal/modf/45005f.wgsl index fe0b65699a..2b413ac53b 100644 --- a/test/tint/builtins/gen/literal/modf/45005f.wgsl +++ b/test/tint/builtins/gen/literal/modf/45005f.wgsl @@ -25,7 +25,7 @@ enable f16; // fn modf(vec<3, f16>) -> __modf_result_vec<3, f16> fn modf_45005f() { - var res = modf(vec3(1.h)); + var res = modf(vec3(-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 e4a00330ea..79a0841eb1 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 @@ struct modf_result_vec3_f16 { vector fract; vector whole; }; -modf_result_vec3_f16 tint_modf(vector 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 d402eda656..8bbd18f764 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 @@ struct modf_result_vec3_f16 { 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 @@ struct modf_result_vec3_f16 { 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 @@ struct modf_result_vec3_f16 { 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 a3071d4f5a..5323690a18 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 @@ struct modf_result_vec3_f16 { 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 63c4bd4558..e276d40d4d 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 43f497ee30..5c77370bcc 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(1.0h)); + var res = modf(vec3(-(1.5h))); } @vertex diff --git a/test/tint/builtins/gen/literal/modf/4bfced.wgsl b/test/tint/builtins/gen/literal/modf/4bfced.wgsl index 1a3fd9a1b2..86e38243e1 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(1.f)); + var res = modf(vec4(-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 043bbab60d..deca14c4c0 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 @@ struct modf_result_vec4 { 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 043bbab60d..deca14c4c0 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 @@ struct modf_result_vec4 { 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 7c7e682147..6952a13caa 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 @@ struct modf_result_vec4 { 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 @@ struct modf_result_vec4 { 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 @@ struct modf_result_vec4 { 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 3b5da9c052..9b97487998 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 @@ struct modf_result_vec4 { 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 2a58dd4174..60d2fc255b 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 0b04b87797..2ab46123e5 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(1.0f)); + var res = modf(vec4(-(1.5f))); } @vertex diff --git a/test/tint/builtins/gen/literal/modf/5ea256.wgsl b/test/tint/builtins/gen/literal/modf/5ea256.wgsl index 673bb2517e..1d6a103062 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(1.f)); + var res = modf(vec3(-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 a2c46f5a60..e5949037c8 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 @@ struct modf_result_vec3 { 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 a2c46f5a60..e5949037c8 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 @@ struct modf_result_vec3 { 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 7a727ecf19..84138c0066 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 @@ struct modf_result_vec3 { 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 @@ struct modf_result_vec3 { 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 @@ struct modf_result_vec3 { 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 33e7fa36e6..4e73398b9c 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 @@ struct modf_result_vec3 { 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 6762d611b1..62e6d3208b 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 ea4142d8c6..8132d91bae 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(1.0f)); + var res = modf(vec3(-(1.5f))); } @vertex diff --git a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl index 591a299b75..2af78fda33 100644 --- a/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl +++ b/test/tint/builtins/gen/literal/modf/8dbbbf.wgsl @@ -25,7 +25,7 @@ enable f16; // fn modf(f16) -> __modf_result 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 c84d408785..cd06924560 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 @@ struct modf_result_f16 { 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 f9757b9291..0741bb82d3 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 @@ struct modf_result_f16 { 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 @@ struct modf_result_f16 { 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 @@ struct modf_result_f16 { 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 e9472e7514..6d02f05c82 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 @@ struct modf_result_f16 { 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 ebb1132497..6f758bebbd 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 c3619a7613..d25d72fe41 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 bcb60068c9..59170ae5b9 100644 --- a/test/tint/builtins/gen/literal/modf/995934.wgsl +++ b/test/tint/builtins/gen/literal/modf/995934.wgsl @@ -25,7 +25,7 @@ enable f16; // fn modf(vec<4, f16>) -> __modf_result_vec<4, f16> fn modf_995934() { - var res = modf(vec4(1.h)); + var res = modf(vec4(-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 24c71de828..cc8ede54a2 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 @@ struct modf_result_vec4_f16 { vector fract; vector whole; }; -modf_result_vec4_f16 tint_modf(vector 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 065a9e74ef..cafadcaf68 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 @@ struct modf_result_vec4_f16 { 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 @@ struct modf_result_vec4_f16 { 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 @@ struct modf_result_vec4_f16 { 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 76d3c047e6..6371cad5fe 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 @@ struct modf_result_vec4_f16 { 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 b19f599e2e..d67833bcac 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 a73275d387..c514967ab2 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(1.0h)); + var res = modf(vec4(-(1.5h))); } @vertex diff --git a/test/tint/builtins/gen/literal/modf/a545b9.wgsl b/test/tint/builtins/gen/literal/modf/a545b9.wgsl index 1a66aad4b9..a1c1c6efec 100644 --- a/test/tint/builtins/gen/literal/modf/a545b9.wgsl +++ b/test/tint/builtins/gen/literal/modf/a545b9.wgsl @@ -25,7 +25,7 @@ enable f16; // fn modf(vec<2, f16>) -> __modf_result_vec<2, f16> fn modf_a545b9() { - var res = modf(vec2(1.h)); + var res = modf(vec2(-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 1805568cef..e3fe28651f 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 @@ struct modf_result_vec2_f16 { vector fract; vector whole; }; -modf_result_vec2_f16 tint_modf(vector 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 7c2f892c9e..3bcc8faf73 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 @@ struct modf_result_vec2_f16 { 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 @@ struct modf_result_vec2_f16 { 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 @@ struct modf_result_vec2_f16 { 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 d7ce71e2ba..18005ff5b4 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 @@ struct modf_result_vec2_f16 { 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 63d3be1e4e..a1c614024d 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 9d2b453af8..c7bc52b775 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(1.0h)); + var res = modf(vec2(-(1.5h))); } @vertex diff --git a/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl b/test/tint/builtins/gen/literal/modf/bbf7f7.wgsl index d0ffdaa96d..bdcd0f0b99 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 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 020478a8f8..ec4e06b8ba 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 @@ 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 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 020478a8f8..ec4e06b8ba 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 @@ 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 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 30917acac3..177ebe943d 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 @@ struct modf_result { 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 @@ struct modf_result { 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 @@ struct modf_result { 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 250e830d22..f5fdb5557f 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 @@ 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 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 21a235880f..1f64077942 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 31c86d6146..a93f3e3d20 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 d8f067b9b6..00be018fbb 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(1.f); + var arg_0 = vec2(-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 bf9728ea66..a3745baca5 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 @@ modf_result_vec2 tint_modf(float2 param_0) { } 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 bf9728ea66..a3745baca5 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 @@ modf_result_vec2 tint_modf(float2 param_0) { } 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 a1998b9e45..d7ecc21516 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 @@ modf_result_vec2 tint_modf(vec2 param_0) { 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 @@ modf_result_vec2 tint_modf(vec2 param_0) { 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 @@ modf_result_vec2 tint_modf(vec2 param_0) { 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 75795761e1..5963525f7e 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 @@ modf_result_vec2 tint_modf(float2 param_0) { } 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 9e6337b8ba..4fda240124 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 baa503959e..1ea4d30196 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(1.0f); + var arg_0 = vec2(-(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 37647095ae..8955e2318b 100644 --- a/test/tint/builtins/gen/var/modf/45005f.wgsl +++ b/test/tint/builtins/gen/var/modf/45005f.wgsl @@ -25,7 +25,7 @@ enable f16; // fn modf(vec<3, f16>) -> __modf_result_vec<3, f16> fn modf_45005f() { - var arg_0 = vec3(1.h); + var arg_0 = vec3(-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 f9a9e8d0eb..db8c08486e 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 @@ modf_result_vec3_f16 tint_modf(vector param_0) { } void modf_45005f() { - vector arg_0 = (float16_t(1.0h)).xxx; + vector 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 0c6aad6456..f019a7d1cd 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 @@ modf_result_vec3_f16 tint_modf(f16vec3 param_0) { 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 @@ modf_result_vec3_f16 tint_modf(f16vec3 param_0) { 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 @@ modf_result_vec3_f16 tint_modf(f16vec3 param_0) { 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 106d1b673d..5c9713fa52 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 @@ modf_result_vec3_f16 tint_modf(half3 param_0) { } 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 7c7b81375d..d647058406 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 f9e08f3e48..96fce1ae6d 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(1.0h); + var arg_0 = vec3(-(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 4929415af6..52581ce7c3 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(1.f); + var arg_0 = vec4(-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 0cfe790dc1..e084fcc42c 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 @@ modf_result_vec4 tint_modf(float4 param_0) { } 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 0cfe790dc1..e084fcc42c 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 @@ modf_result_vec4 tint_modf(float4 param_0) { } 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 5d5615b40f..5a153b99c1 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 @@ modf_result_vec4 tint_modf(vec4 param_0) { 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 @@ modf_result_vec4 tint_modf(vec4 param_0) { 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 @@ modf_result_vec4 tint_modf(vec4 param_0) { 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 4bac39e4e9..fe36491366 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 @@ modf_result_vec4 tint_modf(float4 param_0) { } 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 44dedfe0c6..8bdaea472a 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 cff6f57977..cc7e11d579 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(1.0f); + var arg_0 = vec4(-(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 f85cfeb5af..f623263bfe 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(1.f); + var arg_0 = vec3(-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 ffd6a125b3..96db57b752 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 @@ modf_result_vec3 tint_modf(float3 param_0) { } 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 ffd6a125b3..96db57b752 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 @@ modf_result_vec3 tint_modf(float3 param_0) { } 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 021cf81383..9c050b48ce 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 @@ modf_result_vec3 tint_modf(vec3 param_0) { 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 @@ modf_result_vec3 tint_modf(vec3 param_0) { 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 @@ modf_result_vec3 tint_modf(vec3 param_0) { 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 fa5d28af5a..f5a71837c3 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 @@ modf_result_vec3 tint_modf(float3 param_0) { } 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 1587506574..42c66e481a 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 5180d1456a..9d05d93eab 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(1.0f); + var arg_0 = vec3(-(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 f3dea08ec9..3b65908b73 100644 --- a/test/tint/builtins/gen/var/modf/8dbbbf.wgsl +++ b/test/tint/builtins/gen/var/modf/8dbbbf.wgsl @@ -25,7 +25,7 @@ enable f16; // fn modf(f16) -> __modf_result 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 04d940bfde..ef0eb008fc 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 @@ modf_result_f16 tint_modf(float16_t param_0) { } 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 d6b89764a4..db8b53af10 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 @@ modf_result_f16 tint_modf(float16_t param_0) { 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 @@ modf_result_f16 tint_modf(float16_t param_0) { 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 @@ modf_result_f16 tint_modf(float16_t param_0) { 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 58a30cade0..c3b9f88bbe 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 @@ modf_result_f16 tint_modf(half param_0) { } 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 28b19b092d..5714e2fe39 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 d0815e7d1f..e92ab557d6 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 463335efcc..fa33c7011d 100644 --- a/test/tint/builtins/gen/var/modf/995934.wgsl +++ b/test/tint/builtins/gen/var/modf/995934.wgsl @@ -25,7 +25,7 @@ enable f16; // fn modf(vec<4, f16>) -> __modf_result_vec<4, f16> fn modf_995934() { - var arg_0 = vec4(1.h); + var arg_0 = vec4(-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 26c46db522..82b85d3014 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 @@ modf_result_vec4_f16 tint_modf(vector param_0) { } void modf_995934() { - vector arg_0 = (float16_t(1.0h)).xxxx; + vector 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 71659e8717..49bd43b3be 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 @@ modf_result_vec4_f16 tint_modf(f16vec4 param_0) { 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 @@ modf_result_vec4_f16 tint_modf(f16vec4 param_0) { 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 @@ modf_result_vec4_f16 tint_modf(f16vec4 param_0) { 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 1b5ebe731f..4e6cf8aff8 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 @@ modf_result_vec4_f16 tint_modf(half4 param_0) { } 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..440faff511 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 317ae858e1..e5995a9989 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(1.0h); + var arg_0 = vec4(-(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 51895fd8a1..c559912793 100644 --- a/test/tint/builtins/gen/var/modf/a545b9.wgsl +++ b/test/tint/builtins/gen/var/modf/a545b9.wgsl @@ -25,7 +25,7 @@ enable f16; // fn modf(vec<2, f16>) -> __modf_result_vec<2, f16> fn modf_a545b9() { - var arg_0 = vec2(1.h); + var arg_0 = vec2(-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 47cc509cc5..30290bcf11 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 @@ modf_result_vec2_f16 tint_modf(vector param_0) { } void modf_a545b9() { - vector arg_0 = (float16_t(1.0h)).xx; + vector 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 86f166d4c9..56974b58e4 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 @@ modf_result_vec2_f16 tint_modf(f16vec2 param_0) { 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 @@ modf_result_vec2_f16 tint_modf(f16vec2 param_0) { 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 @@ modf_result_vec2_f16 tint_modf(f16vec2 param_0) { 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 1cf4d62063..2a59af29f5 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 @@ modf_result_vec2_f16 tint_modf(half2 param_0) { } 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 da154baf2c..6bd4d3e9ab 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 dfb8c476e9..c6e6265426 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(1.0h); + var arg_0 = vec2(-(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 60ab469def..7fa9467a74 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 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 39adf2d7b4..26d23b2288 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 @@ modf_result tint_modf(float param_0) { } 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 39adf2d7b4..26d23b2288 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 @@ modf_result tint_modf(float param_0) { } 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 4907635438..12f70df10a 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 @@ modf_result tint_modf(float param_0) { void modf_bbf7f7() { - float arg_0 = 1.0f; + float arg_0 = -1.5f; modf_result res = tint_modf(arg_0); } @@ -46,7 +46,7 @@ modf_result tint_modf(float param_0) { void modf_bbf7f7() { - float arg_0 = 1.0f; + float arg_0 = -1.5f; modf_result res = tint_modf(arg_0); } @@ -73,7 +73,7 @@ modf_result tint_modf(float param_0) { 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 2ef62514a5..3e4accd0e9 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 @@ modf_result tint_modf(float param_0) { } 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 97fe5104d6..6ffbfde31f 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 0ebdc4bb05..418a00ddd8 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 6ea70f3eed..2411190418 100644 --- a/test/tint/builtins/modf.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/modf.wgsl.expected.dxc.hlsl @@ -2,15 +2,9 @@ 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 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 6ea70f3eed..2411190418 100644 --- a/test/tint/builtins/modf.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/modf.wgsl.expected.fxc.hlsl @@ -2,15 +2,9 @@ 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 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 00fa668bcb..e59dfb0dc4 100644 --- a/test/tint/builtins/modf.wgsl.expected.glsl +++ b/test/tint/builtins/modf.wgsl.expected.glsl @@ -5,15 +5,9 @@ struct modf_result { 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 108250d0ae..6c73b2445a 100644 --- a/test/tint/builtins/modf.wgsl.expected.msl +++ b/test/tint/builtins/modf.wgsl.expected.msl @@ -6,14 +6,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; -} - 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 ebbc7669d3..cca23fccb2 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 09a7a21949..ec015d8786 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 float v3 = 1.0f; 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 09a7a21949..ec015d8786 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 float v3 = 1.0f; 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;