From 724ad2a290af41acaa6a6d390412a6d460638f8c Mon Sep 17 00:00:00 2001 From: dan sinclair Date: Thu, 24 Nov 2022 21:54:00 +0000 Subject: [PATCH] Const eval for `fma` This CL adds const-eval for the `fma` builtin. Bug: tint:1581 Change-Id: Ia4df818fec9d5d969b364b2c165400d787a9e275 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/111584 Commit-Queue: Dan Sinclair Kokoro: Kokoro Reviewed-by: Antonio Maiorano --- src/tint/intrinsics.def | 4 +- src/tint/resolver/const_eval.cc | 27 ++++++++ src/tint/resolver/const_eval.h | 9 +++ src/tint/resolver/const_eval_builtin_test.cc | 25 +++++++ src/tint/resolver/intrinsic_table.inl | 12 ++-- .../tint/builtins/gen/literal/fma/143d5d.wgsl | 43 ++++++++++++ .../literal/fma/143d5d.wgsl.expected.dxc.hlsl | 30 +++++++++ .../literal/fma/143d5d.wgsl.expected.fxc.hlsl | 30 +++++++++ .../gen/literal/fma/143d5d.wgsl.expected.glsl | 49 ++++++++++++++ .../gen/literal/fma/143d5d.wgsl.expected.msl | 33 +++++++++ .../literal/fma/143d5d.wgsl.expected.spvasm | 65 ++++++++++++++++++ .../gen/literal/fma/143d5d.wgsl.expected.wgsl | 19 ++++++ .../tint/builtins/gen/literal/fma/1f5084.wgsl | 43 ++++++++++++ .../literal/fma/1f5084.wgsl.expected.dxc.hlsl | 30 +++++++++ .../literal/fma/1f5084.wgsl.expected.fxc.hlsl | 30 +++++++++ .../gen/literal/fma/1f5084.wgsl.expected.glsl | 49 ++++++++++++++ .../gen/literal/fma/1f5084.wgsl.expected.msl | 33 +++++++++ .../literal/fma/1f5084.wgsl.expected.spvasm | 67 +++++++++++++++++++ .../gen/literal/fma/1f5084.wgsl.expected.wgsl | 19 ++++++ .../literal/fma/26a7a9.wgsl.expected.dxc.hlsl | 2 +- .../literal/fma/26a7a9.wgsl.expected.fxc.hlsl | 2 +- .../gen/literal/fma/26a7a9.wgsl.expected.glsl | 6 +- .../gen/literal/fma/26a7a9.wgsl.expected.msl | 2 +- .../literal/fma/26a7a9.wgsl.expected.spvasm | 37 +++++----- .../tint/builtins/gen/literal/fma/466442.wgsl | 43 ++++++++++++ .../literal/fma/466442.wgsl.expected.dxc.hlsl | 30 +++++++++ .../literal/fma/466442.wgsl.expected.fxc.hlsl | 30 +++++++++ .../gen/literal/fma/466442.wgsl.expected.glsl | 49 ++++++++++++++ .../gen/literal/fma/466442.wgsl.expected.msl | 33 +++++++++ .../literal/fma/466442.wgsl.expected.spvasm | 64 ++++++++++++++++++ .../gen/literal/fma/466442.wgsl.expected.wgsl | 19 ++++++ .../literal/fma/6a3283.wgsl.expected.dxc.hlsl | 2 +- .../literal/fma/6a3283.wgsl.expected.fxc.hlsl | 2 +- .../gen/literal/fma/6a3283.wgsl.expected.glsl | 6 +- .../gen/literal/fma/6a3283.wgsl.expected.msl | 2 +- .../literal/fma/6a3283.wgsl.expected.spvasm | 33 +++++---- .../literal/fma/ab7818.wgsl.expected.dxc.hlsl | 2 +- .../gen/literal/fma/ab7818.wgsl.expected.glsl | 6 +- .../gen/literal/fma/ab7818.wgsl.expected.msl | 2 +- .../literal/fma/ab7818.wgsl.expected.spvasm | 36 +++++----- .../literal/fma/bf21b6.wgsl.expected.dxc.hlsl | 2 +- .../gen/literal/fma/bf21b6.wgsl.expected.glsl | 6 +- .../gen/literal/fma/bf21b6.wgsl.expected.msl | 2 +- .../literal/fma/bf21b6.wgsl.expected.spvasm | 36 +++++----- .../literal/fma/c10ba3.wgsl.expected.dxc.hlsl | 2 +- .../literal/fma/c10ba3.wgsl.expected.fxc.hlsl | 2 +- .../gen/literal/fma/c10ba3.wgsl.expected.glsl | 6 +- .../gen/literal/fma/c10ba3.wgsl.expected.msl | 2 +- .../literal/fma/c10ba3.wgsl.expected.spvasm | 31 +++++---- .../literal/fma/c8abb3.wgsl.expected.dxc.hlsl | 2 +- .../gen/literal/fma/c8abb3.wgsl.expected.glsl | 6 +- .../gen/literal/fma/c8abb3.wgsl.expected.msl | 2 +- .../literal/fma/c8abb3.wgsl.expected.spvasm | 34 +++++----- .../literal/fma/e17c5c.wgsl.expected.dxc.hlsl | 2 +- .../literal/fma/e17c5c.wgsl.expected.fxc.hlsl | 2 +- .../gen/literal/fma/e17c5c.wgsl.expected.glsl | 6 +- .../gen/literal/fma/e17c5c.wgsl.expected.msl | 2 +- .../literal/fma/e17c5c.wgsl.expected.spvasm | 37 +++++----- .../literal/fma/e7abdc.wgsl.expected.dxc.hlsl | 2 +- .../gen/literal/fma/e7abdc.wgsl.expected.glsl | 6 +- .../gen/literal/fma/e7abdc.wgsl.expected.msl | 2 +- .../literal/fma/e7abdc.wgsl.expected.spvasm | 36 +++++----- .../tint/builtins/gen/literal/fma/eb25d7.wgsl | 43 ++++++++++++ .../literal/fma/eb25d7.wgsl.expected.dxc.hlsl | 30 +++++++++ .../literal/fma/eb25d7.wgsl.expected.fxc.hlsl | 30 +++++++++ .../gen/literal/fma/eb25d7.wgsl.expected.glsl | 49 ++++++++++++++ .../gen/literal/fma/eb25d7.wgsl.expected.msl | 33 +++++++++ .../literal/fma/eb25d7.wgsl.expected.spvasm | 67 +++++++++++++++++++ .../gen/literal/fma/eb25d7.wgsl.expected.wgsl | 19 ++++++ test/tint/builtins/gen/var/fma/143d5d.wgsl | 46 +++++++++++++ .../gen/var/fma/143d5d.wgsl.expected.dxc.hlsl | 30 +++++++++ .../gen/var/fma/143d5d.wgsl.expected.fxc.hlsl | 30 +++++++++ .../gen/var/fma/143d5d.wgsl.expected.glsl | 49 ++++++++++++++ .../gen/var/fma/143d5d.wgsl.expected.msl | 33 +++++++++ .../gen/var/fma/143d5d.wgsl.expected.spvasm | 65 ++++++++++++++++++ .../gen/var/fma/143d5d.wgsl.expected.wgsl | 22 ++++++ test/tint/builtins/gen/var/fma/1f5084.wgsl | 46 +++++++++++++ .../gen/var/fma/1f5084.wgsl.expected.dxc.hlsl | 30 +++++++++ .../gen/var/fma/1f5084.wgsl.expected.fxc.hlsl | 30 +++++++++ .../gen/var/fma/1f5084.wgsl.expected.glsl | 49 ++++++++++++++ .../gen/var/fma/1f5084.wgsl.expected.msl | 33 +++++++++ .../gen/var/fma/1f5084.wgsl.expected.spvasm | 67 +++++++++++++++++++ .../gen/var/fma/1f5084.wgsl.expected.wgsl | 22 ++++++ test/tint/builtins/gen/var/fma/466442.wgsl | 46 +++++++++++++ .../gen/var/fma/466442.wgsl.expected.dxc.hlsl | 30 +++++++++ .../gen/var/fma/466442.wgsl.expected.fxc.hlsl | 30 +++++++++ .../gen/var/fma/466442.wgsl.expected.glsl | 49 ++++++++++++++ .../gen/var/fma/466442.wgsl.expected.msl | 33 +++++++++ .../gen/var/fma/466442.wgsl.expected.spvasm | 64 ++++++++++++++++++ .../gen/var/fma/466442.wgsl.expected.wgsl | 22 ++++++ test/tint/builtins/gen/var/fma/eb25d7.wgsl | 46 +++++++++++++ .../gen/var/fma/eb25d7.wgsl.expected.dxc.hlsl | 30 +++++++++ .../gen/var/fma/eb25d7.wgsl.expected.fxc.hlsl | 30 +++++++++ .../gen/var/fma/eb25d7.wgsl.expected.glsl | 49 ++++++++++++++ .../gen/var/fma/eb25d7.wgsl.expected.msl | 33 +++++++++ .../gen/var/fma/eb25d7.wgsl.expected.spvasm | 67 +++++++++++++++++++ .../gen/var/fma/eb25d7.wgsl.expected.wgsl | 22 ++++++ 97 files changed, 2429 insertions(+), 198 deletions(-) create mode 100644 test/tint/builtins/gen/literal/fma/143d5d.wgsl create mode 100644 test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/literal/fma/1f5084.wgsl create mode 100644 test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/literal/fma/466442.wgsl create mode 100644 test/tint/builtins/gen/literal/fma/466442.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/fma/466442.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/fma/466442.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/fma/466442.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/fma/466442.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/fma/466442.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/literal/fma/eb25d7.wgsl create mode 100644 test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/var/fma/143d5d.wgsl create mode 100644 test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/var/fma/1f5084.wgsl create mode 100644 test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/var/fma/466442.wgsl create mode 100644 test/tint/builtins/gen/var/fma/466442.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/var/fma/466442.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/var/fma/466442.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/var/fma/466442.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/var/fma/466442.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/var/fma/466442.wgsl.expected.wgsl create mode 100644 test/tint/builtins/gen/var/fma/eb25d7.wgsl create mode 100644 test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.dxc.hlsl create mode 100644 test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.fxc.hlsl create mode 100644 test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.msl create mode 100644 test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.spvasm create mode 100644 test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.wgsl diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def index 0752d5a1ba..ba8a5810c9 100644 --- a/src/tint/intrinsics.def +++ b/src/tint/intrinsics.def @@ -473,8 +473,8 @@ fn faceForward(vec, vec, vec) -> vec @const fn firstTrailingBit(vec) -> vec @const fn floor(@test_value(1.5) T) -> T @const fn floor(@test_value(1.5) vec) -> vec -fn fma(T, T, T) -> T -fn fma(vec, vec, vec) -> vec +@const fn fma(T, T, T) -> T +@const fn fma(vec, vec, vec) -> vec fn fract(T) -> T fn fract(vec) -> vec @const fn frexp(T) -> __frexp_result diff --git a/src/tint/resolver/const_eval.cc b/src/tint/resolver/const_eval.cc index 424cd8279f..0e6f6824ea 100644 --- a/src/tint/resolver/const_eval.cc +++ b/src/tint/resolver/const_eval.cc @@ -2446,6 +2446,33 @@ ConstEval::Result ConstEval::floor(const sem::Type* ty, return TransformElements(builder, ty, transform, args[0]); } +ConstEval::Result ConstEval::fma(const sem::Type* ty, + utils::VectorRef args, + const Source& source) { + auto transform = [&](const sem::Constant* c1, const sem::Constant* c2, + const sem::Constant* c3) { + auto create = [&](auto e1, auto e2, auto e3) -> ImplResult { + auto err_msg = [&] { + AddNote("when calculating fma", source); + return utils::Failure; + }; + + auto mul = Mul(source, e1, e2); + if (!mul) { + return err_msg(); + } + + auto val = Add(source, mul.Get(), e3); + if (!val) { + return err_msg(); + } + return CreateElement(builder, source, c1->Type(), val.Get()); + }; + return Dispatch_fa_f32_f16(create, c1, c2, c3); + }; + return TransformElements(builder, ty, transform, args[0], args[1], args[2]); +} + ConstEval::Result ConstEval::frexp(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 33ec0a49ef..b65bb2821f 100644 --- a/src/tint/resolver/const_eval.h +++ b/src/tint/resolver/const_eval.h @@ -629,6 +629,15 @@ class ConstEval { utils::VectorRef args, const Source& source); + /// fma 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 fma(const sem::Type* ty, + utils::VectorRef args, + const Source& source); + /// frexp 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 9c5532362d..849f3ef384 100644 --- a/src/tint/resolver/const_eval_builtin_test.cc +++ b/src/tint/resolver/const_eval_builtin_test.cc @@ -1072,6 +1072,31 @@ INSTANTIATE_TEST_SUITE_P( // testing::ValuesIn(Concat(FloorCases(), // FloorCases(), FloorCases())))); + +template +std::vector FmaCases() { + auto error_msg = [](auto a, const char* op, auto b) { + return "12:34 error: " + OverflowErrorMessage(a, op, b) + R"( +12:34 note: when calculating fma)"; + }; + return { + C({T(0), T(0), T(0)}, T(0)), + C({T(1), T(2), T(3)}, T(5)), + C({Vec(T(1), T(2.5), -T(1)), Vec(T(2), T(2.5), T(1)), Vec(T(4), T(3.75), -T(2))}, + Vec(T(6), T(10), -T(3))), + + E({T::Highest(), T::Highest(), T(0)}, error_msg(T::Highest(), "*", T::Highest())), + E({T::Highest(), T(1), T::Highest()}, error_msg(T::Highest(), "+", T::Highest())), + }; +} +INSTANTIATE_TEST_SUITE_P( // + Fma, + ResolverConstEvalBuiltinTest, + testing::Combine(testing::Values(sem::BuiltinType::kFma), + testing::ValuesIn(Concat(FmaCases(), // + FmaCases(), + FmaCases())))); + template std::vector FrexpCases() { using F = T; // fract type diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl index e213f813df..1a2634f585 100644 --- a/src/tint/resolver/intrinsic_table.inl +++ b/src/tint/resolver/intrinsic_table.inl @@ -12510,24 +12510,24 @@ constexpr OverloadInfo kOverloads[] = { /* num parameters */ 3, /* num template types */ 1, /* num template numbers */ 0, - /* template types */ &kTemplateTypes[26], + /* template types */ &kTemplateTypes[23], /* template numbers */ &kTemplateNumbers[10], /* parameters */ &kParameters[462], /* return matcher indices */ &kMatcherIndices[3], /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), - /* const eval */ nullptr, + /* const eval */ &ConstEval::fma, }, { /* [350] */ /* num parameters */ 3, /* num template types */ 1, /* num template numbers */ 1, - /* template types */ &kTemplateTypes[26], + /* template types */ &kTemplateTypes[23], /* template numbers */ &kTemplateNumbers[4], /* parameters */ &kParameters[465], /* return matcher indices */ &kMatcherIndices[30], /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), - /* const eval */ nullptr, + /* const eval */ &ConstEval::fma, }, { /* [351] */ @@ -14245,8 +14245,8 @@ constexpr IntrinsicInfo kBuiltins[] = { }, { /* [38] */ - /* fn fma(T, T, T) -> T */ - /* fn fma(vec, vec, vec) -> vec */ + /* fn fma(T, T, T) -> T */ + /* fn fma(vec, vec, vec) -> vec */ /* num overloads */ 2, /* overloads */ &kOverloads[349], }, diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl b/test/tint/builtins/gen/literal/fma/143d5d.wgsl new file mode 100644 index 0000000000..5a83670083 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/143d5d.wgsl @@ -0,0 +1,43 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +//////////////////////////////////////////////////////////////////////////////// +// File generated by tools/src/cmd/gen +// using the template: +// test/tint/builtins/gen/gen.wgsl.tmpl +// +// Do not modify this file directly +//////////////////////////////////////////////////////////////////////////////// + + +// fn fma(vec<4, fa>, vec<4, fa>, vec<4, fa>) -> vec<4, fa> +fn fma_143d5d() { + var res = fma(vec4(1.), vec4(1.), vec4(1.)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_143d5d(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_143d5d(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_143d5d(); +} diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..6c3ad71f1d --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.dxc.hlsl @@ -0,0 +1,30 @@ +void fma_143d5d() { + float4 res = (2.0f).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_143d5d(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_143d5d(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_143d5d(); + return; +} diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..6c3ad71f1d --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.fxc.hlsl @@ -0,0 +1,30 @@ +void fma_143d5d() { + float4 res = (2.0f).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_143d5d(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_143d5d(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_143d5d(); + return; +} diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.glsl new file mode 100644 index 0000000000..7e6b4b42f2 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.glsl @@ -0,0 +1,49 @@ +#version 310 es + +void fma_143d5d() { + vec4 res = vec4(2.0f); +} + +vec4 vertex_main() { + fma_143d5d(); + return vec4(0.0f); +} + +void main() { + gl_PointSize = 1.0; + vec4 inner_result = vertex_main(); + gl_Position = inner_result; + gl_Position.y = -(gl_Position.y); + gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w); + return; +} +#version 310 es +precision mediump float; + +void fma_143d5d() { + vec4 res = vec4(2.0f); +} + +void fragment_main() { + fma_143d5d(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +void fma_143d5d() { + vec4 res = vec4(2.0f); +} + +void compute_main() { + fma_143d5d(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.msl new file mode 100644 index 0000000000..af723352cb --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; +void fma_143d5d() { + float4 res = float4(2.0f); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner() { + fma_143d5d(); + return float4(0.0f); +} + +vertex tint_symbol vertex_main() { + float4 const inner_result = vertex_main_inner(); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main() { + fma_143d5d(); + return; +} + +kernel void compute_main() { + fma_143d5d(); + return; +} + diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.spvasm new file mode 100644 index 0000000000..fba8c32478 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.spvasm @@ -0,0 +1,65 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 31 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size + OpEntryPoint Fragment %fragment_main "fragment_main" + OpEntryPoint GLCompute %compute_main "compute_main" + OpExecutionMode %fragment_main OriginUpperLeft + OpExecutionMode %compute_main LocalSize 1 1 1 + OpName %value "value" + OpName %vertex_point_size "vertex_point_size" + OpName %fma_143d5d "fma_143d5d" + OpName %res "res" + OpName %vertex_main_inner "vertex_main_inner" + OpName %vertex_main "vertex_main" + OpName %fragment_main "fragment_main" + OpName %compute_main "compute_main" + OpDecorate %value BuiltIn Position + OpDecorate %vertex_point_size BuiltIn PointSize + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%_ptr_Output_v4float = OpTypePointer Output %v4float + %5 = OpConstantNull %v4float + %value = OpVariable %_ptr_Output_v4float Output %5 +%_ptr_Output_float = OpTypePointer Output %float + %8 = OpConstantNull %float +%vertex_point_size = OpVariable %_ptr_Output_float Output %8 + %void = OpTypeVoid + %9 = OpTypeFunction %void + %float_2 = OpConstant %float 2 + %14 = OpConstantComposite %v4float %float_2 %float_2 %float_2 %float_2 +%_ptr_Function_v4float = OpTypePointer Function %v4float + %17 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 + %fma_143d5d = OpFunction %void None %9 + %12 = OpLabel + %res = OpVariable %_ptr_Function_v4float Function %5 + OpStore %res %14 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %17 + %19 = OpLabel + %20 = OpFunctionCall %void %fma_143d5d + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %9 + %22 = OpLabel + %23 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %23 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %9 + %26 = OpLabel + %27 = OpFunctionCall %void %fma_143d5d + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %9 + %29 = OpLabel + %30 = OpFunctionCall %void %fma_143d5d + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.wgsl new file mode 100644 index 0000000000..887606fe1c --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/143d5d.wgsl.expected.wgsl @@ -0,0 +1,19 @@ +fn fma_143d5d() { + var res = fma(vec4(1.0), vec4(1.0), vec4(1.0)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_143d5d(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_143d5d(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_143d5d(); +} diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl b/test/tint/builtins/gen/literal/fma/1f5084.wgsl new file mode 100644 index 0000000000..3453e11a1f --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/1f5084.wgsl @@ -0,0 +1,43 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +//////////////////////////////////////////////////////////////////////////////// +// File generated by tools/src/cmd/gen +// using the template: +// test/tint/builtins/gen/gen.wgsl.tmpl +// +// Do not modify this file directly +//////////////////////////////////////////////////////////////////////////////// + + +// fn fma(vec<2, fa>, vec<2, fa>, vec<2, fa>) -> vec<2, fa> +fn fma_1f5084() { + var res = fma(vec2(1.), vec2(1.), vec2(1.)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_1f5084(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_1f5084(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_1f5084(); +} diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..e7655ab50d --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.dxc.hlsl @@ -0,0 +1,30 @@ +void fma_1f5084() { + float2 res = (2.0f).xx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_1f5084(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_1f5084(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_1f5084(); + return; +} diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..e7655ab50d --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.fxc.hlsl @@ -0,0 +1,30 @@ +void fma_1f5084() { + float2 res = (2.0f).xx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_1f5084(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_1f5084(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_1f5084(); + return; +} diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.glsl new file mode 100644 index 0000000000..215c7401c2 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.glsl @@ -0,0 +1,49 @@ +#version 310 es + +void fma_1f5084() { + vec2 res = vec2(2.0f); +} + +vec4 vertex_main() { + fma_1f5084(); + return vec4(0.0f); +} + +void main() { + gl_PointSize = 1.0; + vec4 inner_result = vertex_main(); + gl_Position = inner_result; + gl_Position.y = -(gl_Position.y); + gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w); + return; +} +#version 310 es +precision mediump float; + +void fma_1f5084() { + vec2 res = vec2(2.0f); +} + +void fragment_main() { + fma_1f5084(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +void fma_1f5084() { + vec2 res = vec2(2.0f); +} + +void compute_main() { + fma_1f5084(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.msl new file mode 100644 index 0000000000..bc3b0f35c6 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; +void fma_1f5084() { + float2 res = float2(2.0f); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner() { + fma_1f5084(); + return float4(0.0f); +} + +vertex tint_symbol vertex_main() { + float4 const inner_result = vertex_main_inner(); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main() { + fma_1f5084(); + return; +} + +kernel void compute_main() { + fma_1f5084(); + return; +} + diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.spvasm new file mode 100644 index 0000000000..3c5ab1b481 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.spvasm @@ -0,0 +1,67 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 33 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size + OpEntryPoint Fragment %fragment_main "fragment_main" + OpEntryPoint GLCompute %compute_main "compute_main" + OpExecutionMode %fragment_main OriginUpperLeft + OpExecutionMode %compute_main LocalSize 1 1 1 + OpName %value "value" + OpName %vertex_point_size "vertex_point_size" + OpName %fma_1f5084 "fma_1f5084" + OpName %res "res" + OpName %vertex_main_inner "vertex_main_inner" + OpName %vertex_main "vertex_main" + OpName %fragment_main "fragment_main" + OpName %compute_main "compute_main" + OpDecorate %value BuiltIn Position + OpDecorate %vertex_point_size BuiltIn PointSize + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%_ptr_Output_v4float = OpTypePointer Output %v4float + %5 = OpConstantNull %v4float + %value = OpVariable %_ptr_Output_v4float Output %5 +%_ptr_Output_float = OpTypePointer Output %float + %8 = OpConstantNull %float +%vertex_point_size = OpVariable %_ptr_Output_float Output %8 + %void = OpTypeVoid + %9 = OpTypeFunction %void + %v2float = OpTypeVector %float 2 + %float_2 = OpConstant %float 2 + %15 = OpConstantComposite %v2float %float_2 %float_2 +%_ptr_Function_v2float = OpTypePointer Function %v2float + %18 = OpConstantNull %v2float + %19 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 + %fma_1f5084 = OpFunction %void None %9 + %12 = OpLabel + %res = OpVariable %_ptr_Function_v2float Function %18 + OpStore %res %15 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %19 + %21 = OpLabel + %22 = OpFunctionCall %void %fma_1f5084 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %9 + %24 = OpLabel + %25 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %25 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %9 + %28 = OpLabel + %29 = OpFunctionCall %void %fma_1f5084 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %9 + %31 = OpLabel + %32 = OpFunctionCall %void %fma_1f5084 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.wgsl new file mode 100644 index 0000000000..116db76657 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/1f5084.wgsl.expected.wgsl @@ -0,0 +1,19 @@ +fn fma_1f5084() { + var res = fma(vec2(1.0), vec2(1.0), vec2(1.0)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_1f5084(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_1f5084(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_1f5084(); +} diff --git a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.dxc.hlsl index d58d1dbb11..7f4c278a1d 100644 --- a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.dxc.hlsl @@ -1,5 +1,5 @@ void fma_26a7a9() { - float2 res = mad((1.0f).xx, (1.0f).xx, (1.0f).xx); + float2 res = (2.0f).xx; } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.fxc.hlsl index d58d1dbb11..7f4c278a1d 100644 --- a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.fxc.hlsl @@ -1,5 +1,5 @@ void fma_26a7a9() { - float2 res = mad((1.0f).xx, (1.0f).xx, (1.0f).xx); + float2 res = (2.0f).xx; } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.glsl index 1905d1615b..ca63145405 100644 --- a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.glsl +++ b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.glsl @@ -1,7 +1,7 @@ #version 310 es void fma_26a7a9() { - vec2 res = ((vec2(1.0f)) * (vec2(1.0f)) + (vec2(1.0f))); + vec2 res = vec2(2.0f); } vec4 vertex_main() { @@ -21,7 +21,7 @@ void main() { precision mediump float; void fma_26a7a9() { - vec2 res = ((vec2(1.0f)) * (vec2(1.0f)) + (vec2(1.0f))); + vec2 res = vec2(2.0f); } void fragment_main() { @@ -35,7 +35,7 @@ void main() { #version 310 es void fma_26a7a9() { - vec2 res = ((vec2(1.0f)) * (vec2(1.0f)) + (vec2(1.0f))); + vec2 res = vec2(2.0f); } void compute_main() { diff --git a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.msl index 98fe6440fe..77a1973e42 100644 --- a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; void fma_26a7a9() { - float2 res = fma(float2(1.0f), float2(1.0f), float2(1.0f)); + float2 res = float2(2.0f); } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.spvasm index 2e2385e2b2..63647d45a7 100644 --- a/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/literal/fma/26a7a9.wgsl.expected.spvasm @@ -1,10 +1,9 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 34 +; Bound: 33 ; 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" @@ -32,37 +31,37 @@ %void = OpTypeVoid %9 = OpTypeFunction %void %v2float = OpTypeVector %float 2 - %float_1 = OpConstant %float 1 - %17 = OpConstantComposite %v2float %float_1 %float_1 + %float_2 = OpConstant %float 2 + %15 = OpConstantComposite %v2float %float_2 %float_2 %_ptr_Function_v2float = OpTypePointer Function %v2float - %20 = OpConstantNull %v2float - %21 = OpTypeFunction %v4float + %18 = OpConstantNull %v2float + %19 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 %fma_26a7a9 = OpFunction %void None %9 %12 = OpLabel - %res = OpVariable %_ptr_Function_v2float Function %20 - %13 = OpExtInst %v2float %15 Fma %17 %17 %17 - OpStore %res %13 + %res = OpVariable %_ptr_Function_v2float Function %18 + OpStore %res %15 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %21 - %23 = OpLabel - %24 = OpFunctionCall %void %fma_26a7a9 +%vertex_main_inner = OpFunction %v4float None %19 + %21 = OpLabel + %22 = OpFunctionCall %void %fma_26a7a9 OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %26 = OpLabel - %27 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %27 + %24 = OpLabel + %25 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %25 OpStore %vertex_point_size %float_1 OpReturn OpFunctionEnd %fragment_main = OpFunction %void None %9 - %29 = OpLabel - %30 = OpFunctionCall %void %fma_26a7a9 + %28 = OpLabel + %29 = OpFunctionCall %void %fma_26a7a9 OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %32 = OpLabel - %33 = OpFunctionCall %void %fma_26a7a9 + %31 = OpLabel + %32 = OpFunctionCall %void %fma_26a7a9 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl b/test/tint/builtins/gen/literal/fma/466442.wgsl new file mode 100644 index 0000000000..e46749a183 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/466442.wgsl @@ -0,0 +1,43 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +//////////////////////////////////////////////////////////////////////////////// +// File generated by tools/src/cmd/gen +// using the template: +// test/tint/builtins/gen/gen.wgsl.tmpl +// +// Do not modify this file directly +//////////////////////////////////////////////////////////////////////////////// + + +// fn fma(fa, fa, fa) -> fa +fn fma_466442() { + var res = fma(1., 1., 1.); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_466442(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_466442(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_466442(); +} diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..c070a74b09 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.dxc.hlsl @@ -0,0 +1,30 @@ +void fma_466442() { + float res = 2.0f; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_466442(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_466442(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_466442(); + return; +} diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..c070a74b09 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.fxc.hlsl @@ -0,0 +1,30 @@ +void fma_466442() { + float res = 2.0f; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_466442(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_466442(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_466442(); + return; +} diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.glsl new file mode 100644 index 0000000000..c094238683 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.glsl @@ -0,0 +1,49 @@ +#version 310 es + +void fma_466442() { + float res = 2.0f; +} + +vec4 vertex_main() { + fma_466442(); + return vec4(0.0f); +} + +void main() { + gl_PointSize = 1.0; + vec4 inner_result = vertex_main(); + gl_Position = inner_result; + gl_Position.y = -(gl_Position.y); + gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w); + return; +} +#version 310 es +precision mediump float; + +void fma_466442() { + float res = 2.0f; +} + +void fragment_main() { + fma_466442(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +void fma_466442() { + float res = 2.0f; +} + +void compute_main() { + fma_466442(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.msl new file mode 100644 index 0000000000..26b2b96a55 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; +void fma_466442() { + float res = 2.0f; +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner() { + fma_466442(); + return float4(0.0f); +} + +vertex tint_symbol vertex_main() { + float4 const inner_result = vertex_main_inner(); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main() { + fma_466442(); + return; +} + +kernel void compute_main() { + fma_466442(); + return; +} + diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.spvasm new file mode 100644 index 0000000000..55e5e1f0b8 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.spvasm @@ -0,0 +1,64 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 30 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size + OpEntryPoint Fragment %fragment_main "fragment_main" + OpEntryPoint GLCompute %compute_main "compute_main" + OpExecutionMode %fragment_main OriginUpperLeft + OpExecutionMode %compute_main LocalSize 1 1 1 + OpName %value "value" + OpName %vertex_point_size "vertex_point_size" + OpName %fma_466442 "fma_466442" + OpName %res "res" + OpName %vertex_main_inner "vertex_main_inner" + OpName %vertex_main "vertex_main" + OpName %fragment_main "fragment_main" + OpName %compute_main "compute_main" + OpDecorate %value BuiltIn Position + OpDecorate %vertex_point_size BuiltIn PointSize + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%_ptr_Output_v4float = OpTypePointer Output %v4float + %5 = OpConstantNull %v4float + %value = OpVariable %_ptr_Output_v4float Output %5 +%_ptr_Output_float = OpTypePointer Output %float + %8 = OpConstantNull %float +%vertex_point_size = OpVariable %_ptr_Output_float Output %8 + %void = OpTypeVoid + %9 = OpTypeFunction %void + %float_2 = OpConstant %float 2 +%_ptr_Function_float = OpTypePointer Function %float + %16 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 + %fma_466442 = OpFunction %void None %9 + %12 = OpLabel + %res = OpVariable %_ptr_Function_float Function %8 + OpStore %res %float_2 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %16 + %18 = OpLabel + %19 = OpFunctionCall %void %fma_466442 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %9 + %21 = OpLabel + %22 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %22 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %9 + %25 = OpLabel + %26 = OpFunctionCall %void %fma_466442 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %9 + %28 = OpLabel + %29 = OpFunctionCall %void %fma_466442 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.wgsl new file mode 100644 index 0000000000..cae52f3716 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/466442.wgsl.expected.wgsl @@ -0,0 +1,19 @@ +fn fma_466442() { + var res = fma(1.0, 1.0, 1.0); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_466442(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_466442(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_466442(); +} diff --git a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.dxc.hlsl index 8959dfa17a..a518076f47 100644 --- a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.dxc.hlsl @@ -1,5 +1,5 @@ void fma_6a3283() { - float4 res = mad((1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx); + float4 res = (2.0f).xxxx; } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.fxc.hlsl index 8959dfa17a..a518076f47 100644 --- a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.fxc.hlsl @@ -1,5 +1,5 @@ void fma_6a3283() { - float4 res = mad((1.0f).xxxx, (1.0f).xxxx, (1.0f).xxxx); + float4 res = (2.0f).xxxx; } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.glsl index 96bc9f232b..f18280d69f 100644 --- a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.glsl +++ b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.glsl @@ -1,7 +1,7 @@ #version 310 es void fma_6a3283() { - vec4 res = ((vec4(1.0f)) * (vec4(1.0f)) + (vec4(1.0f))); + vec4 res = vec4(2.0f); } vec4 vertex_main() { @@ -21,7 +21,7 @@ void main() { precision mediump float; void fma_6a3283() { - vec4 res = ((vec4(1.0f)) * (vec4(1.0f)) + (vec4(1.0f))); + vec4 res = vec4(2.0f); } void fragment_main() { @@ -35,7 +35,7 @@ void main() { #version 310 es void fma_6a3283() { - vec4 res = ((vec4(1.0f)) * (vec4(1.0f)) + (vec4(1.0f))); + vec4 res = vec4(2.0f); } void compute_main() { diff --git a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.msl index 6c675c1166..747f0d64c7 100644 --- a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; void fma_6a3283() { - float4 res = fma(float4(1.0f), float4(1.0f), float4(1.0f)); + float4 res = float4(2.0f); } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.spvasm index 197b06b98a..2f1ceb190a 100644 --- a/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/literal/fma/6a3283.wgsl.expected.spvasm @@ -1,10 +1,9 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 32 +; Bound: 31 ; Schema: 0 OpCapability Shader - %14 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size OpEntryPoint Fragment %fragment_main "fragment_main" @@ -31,36 +30,36 @@ %vertex_point_size = OpVariable %_ptr_Output_float Output %8 %void = OpTypeVoid %9 = OpTypeFunction %void - %float_1 = OpConstant %float 1 - %16 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1 + %float_2 = OpConstant %float 2 + %14 = OpConstantComposite %v4float %float_2 %float_2 %float_2 %float_2 %_ptr_Function_v4float = OpTypePointer Function %v4float - %19 = OpTypeFunction %v4float + %17 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 %fma_6a3283 = OpFunction %void None %9 %12 = OpLabel %res = OpVariable %_ptr_Function_v4float Function %5 - %13 = OpExtInst %v4float %14 Fma %16 %16 %16 - OpStore %res %13 + OpStore %res %14 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %19 - %21 = OpLabel - %22 = OpFunctionCall %void %fma_6a3283 +%vertex_main_inner = OpFunction %v4float None %17 + %19 = OpLabel + %20 = OpFunctionCall %void %fma_6a3283 OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %24 = OpLabel - %25 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %25 + %22 = OpLabel + %23 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %23 OpStore %vertex_point_size %float_1 OpReturn OpFunctionEnd %fragment_main = OpFunction %void None %9 - %27 = OpLabel - %28 = OpFunctionCall %void %fma_6a3283 + %26 = OpLabel + %27 = OpFunctionCall %void %fma_6a3283 OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %30 = OpLabel - %31 = OpFunctionCall %void %fma_6a3283 + %29 = OpLabel + %30 = OpFunctionCall %void %fma_6a3283 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.dxc.hlsl index 3336b31ee2..aaafc9a7c3 100644 --- a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.dxc.hlsl @@ -1,5 +1,5 @@ void fma_ab7818() { - vector res = mad((float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx, (float16_t(1.0h)).xxxx); + vector res = (float16_t(2.0h)).xxxx; } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.glsl index fa4e7607bd..034f24ac07 100644 --- a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.glsl +++ b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.glsl @@ -2,7 +2,7 @@ #extension GL_AMD_gpu_shader_half_float : require void fma_ab7818() { - f16vec4 res = ((f16vec4(1.0hf)) * (f16vec4(1.0hf)) + (f16vec4(1.0hf))); + f16vec4 res = f16vec4(2.0hf); } vec4 vertex_main() { @@ -23,7 +23,7 @@ void main() { precision mediump float; void fma_ab7818() { - f16vec4 res = ((f16vec4(1.0hf)) * (f16vec4(1.0hf)) + (f16vec4(1.0hf))); + f16vec4 res = f16vec4(2.0hf); } void fragment_main() { @@ -38,7 +38,7 @@ void main() { #extension GL_AMD_gpu_shader_half_float : require void fma_ab7818() { - f16vec4 res = ((f16vec4(1.0hf)) * (f16vec4(1.0hf)) + (f16vec4(1.0hf))); + f16vec4 res = f16vec4(2.0hf); } void compute_main() { diff --git a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.msl index eb2484d7ae..77bd22dea9 100644 --- a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; void fma_ab7818() { - half4 res = fma(half4(1.0h), half4(1.0h), half4(1.0h)); + half4 res = half4(2.0h); } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.spvasm index ccf40ea391..b278ae3a34 100644 --- a/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/literal/fma/ab7818.wgsl.expected.spvasm @@ -1,14 +1,13 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 36 +; Bound: 34 ; Schema: 0 OpCapability Shader OpCapability Float16 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" @@ -37,38 +36,37 @@ %9 = OpTypeFunction %void %half = OpTypeFloat 16 %v4half = OpTypeVector %half 4 -%half_0x1p_0 = OpConstant %half 0x1p+0 - %18 = OpConstantComposite %v4half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 +%half_0x1p_1 = OpConstant %half 0x1p+1 + %16 = OpConstantComposite %v4half %half_0x1p_1 %half_0x1p_1 %half_0x1p_1 %half_0x1p_1 %_ptr_Function_v4half = OpTypePointer Function %v4half - %21 = OpConstantNull %v4half - %22 = OpTypeFunction %v4float + %19 = OpConstantNull %v4half + %20 = OpTypeFunction %v4float %float_1 = OpConstant %float 1 %fma_ab7818 = OpFunction %void None %9 %12 = OpLabel - %res = OpVariable %_ptr_Function_v4half Function %21 - %13 = OpExtInst %v4half %16 Fma %18 %18 %18 - OpStore %res %13 + %res = OpVariable %_ptr_Function_v4half Function %19 + OpStore %res %16 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %22 - %24 = OpLabel - %25 = OpFunctionCall %void %fma_ab7818 +%vertex_main_inner = OpFunction %v4float None %20 + %22 = OpLabel + %23 = OpFunctionCall %void %fma_ab7818 OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %27 = OpLabel - %28 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %28 + %25 = OpLabel + %26 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %26 OpStore %vertex_point_size %float_1 OpReturn OpFunctionEnd %fragment_main = OpFunction %void None %9 - %31 = OpLabel - %32 = OpFunctionCall %void %fma_ab7818 + %29 = OpLabel + %30 = OpFunctionCall %void %fma_ab7818 OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %34 = OpLabel - %35 = OpFunctionCall %void %fma_ab7818 + %32 = OpLabel + %33 = OpFunctionCall %void %fma_ab7818 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.dxc.hlsl index 9e48041ae2..846fd3367d 100644 --- a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.dxc.hlsl @@ -1,5 +1,5 @@ void fma_bf21b6() { - vector res = mad((float16_t(1.0h)).xx, (float16_t(1.0h)).xx, (float16_t(1.0h)).xx); + vector res = (float16_t(2.0h)).xx; } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.glsl index 627c7c9ade..fb633c8bc0 100644 --- a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.glsl +++ b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.glsl @@ -2,7 +2,7 @@ #extension GL_AMD_gpu_shader_half_float : require void fma_bf21b6() { - f16vec2 res = ((f16vec2(1.0hf)) * (f16vec2(1.0hf)) + (f16vec2(1.0hf))); + f16vec2 res = f16vec2(2.0hf); } vec4 vertex_main() { @@ -23,7 +23,7 @@ void main() { precision mediump float; void fma_bf21b6() { - f16vec2 res = ((f16vec2(1.0hf)) * (f16vec2(1.0hf)) + (f16vec2(1.0hf))); + f16vec2 res = f16vec2(2.0hf); } void fragment_main() { @@ -38,7 +38,7 @@ void main() { #extension GL_AMD_gpu_shader_half_float : require void fma_bf21b6() { - f16vec2 res = ((f16vec2(1.0hf)) * (f16vec2(1.0hf)) + (f16vec2(1.0hf))); + f16vec2 res = f16vec2(2.0hf); } void compute_main() { diff --git a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.msl index 54c227bc22..69f91587ac 100644 --- a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; void fma_bf21b6() { - half2 res = fma(half2(1.0h), half2(1.0h), half2(1.0h)); + half2 res = half2(2.0h); } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.spvasm index 6b33758ad3..d0c5c44d8a 100644 --- a/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/literal/fma/bf21b6.wgsl.expected.spvasm @@ -1,14 +1,13 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 36 +; Bound: 34 ; Schema: 0 OpCapability Shader OpCapability Float16 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" @@ -37,38 +36,37 @@ %9 = OpTypeFunction %void %half = OpTypeFloat 16 %v2half = OpTypeVector %half 2 -%half_0x1p_0 = OpConstant %half 0x1p+0 - %18 = OpConstantComposite %v2half %half_0x1p_0 %half_0x1p_0 +%half_0x1p_1 = OpConstant %half 0x1p+1 + %16 = OpConstantComposite %v2half %half_0x1p_1 %half_0x1p_1 %_ptr_Function_v2half = OpTypePointer Function %v2half - %21 = OpConstantNull %v2half - %22 = OpTypeFunction %v4float + %19 = OpConstantNull %v2half + %20 = OpTypeFunction %v4float %float_1 = OpConstant %float 1 %fma_bf21b6 = OpFunction %void None %9 %12 = OpLabel - %res = OpVariable %_ptr_Function_v2half Function %21 - %13 = OpExtInst %v2half %16 Fma %18 %18 %18 - OpStore %res %13 + %res = OpVariable %_ptr_Function_v2half Function %19 + OpStore %res %16 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %22 - %24 = OpLabel - %25 = OpFunctionCall %void %fma_bf21b6 +%vertex_main_inner = OpFunction %v4float None %20 + %22 = OpLabel + %23 = OpFunctionCall %void %fma_bf21b6 OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %27 = OpLabel - %28 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %28 + %25 = OpLabel + %26 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %26 OpStore %vertex_point_size %float_1 OpReturn OpFunctionEnd %fragment_main = OpFunction %void None %9 - %31 = OpLabel - %32 = OpFunctionCall %void %fma_bf21b6 + %29 = OpLabel + %30 = OpFunctionCall %void %fma_bf21b6 OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %34 = OpLabel - %35 = OpFunctionCall %void %fma_bf21b6 + %32 = OpLabel + %33 = OpFunctionCall %void %fma_bf21b6 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.dxc.hlsl index 88e5bb2cff..b8ef7d66f1 100644 --- a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.dxc.hlsl @@ -1,5 +1,5 @@ void fma_c10ba3() { - float res = mad(1.0f, 1.0f, 1.0f); + float res = 2.0f; } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.fxc.hlsl index 88e5bb2cff..b8ef7d66f1 100644 --- a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.fxc.hlsl @@ -1,5 +1,5 @@ void fma_c10ba3() { - float res = mad(1.0f, 1.0f, 1.0f); + float res = 2.0f; } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.glsl index 611c0b660c..009c216d1b 100644 --- a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.glsl +++ b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.glsl @@ -1,7 +1,7 @@ #version 310 es void fma_c10ba3() { - float res = ((1.0f) * (1.0f) + (1.0f)); + float res = 2.0f; } vec4 vertex_main() { @@ -21,7 +21,7 @@ void main() { precision mediump float; void fma_c10ba3() { - float res = ((1.0f) * (1.0f) + (1.0f)); + float res = 2.0f; } void fragment_main() { @@ -35,7 +35,7 @@ void main() { #version 310 es void fma_c10ba3() { - float res = ((1.0f) * (1.0f) + (1.0f)); + float res = 2.0f; } void compute_main() { diff --git a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.msl index 9f6001806b..e222ebd86f 100644 --- a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; void fma_c10ba3() { - float res = fma(1.0f, 1.0f, 1.0f); + float res = 2.0f; } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.spvasm index bc9816f1b1..53355676fe 100644 --- a/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/literal/fma/c10ba3.wgsl.expected.spvasm @@ -1,10 +1,9 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 31 +; Bound: 30 ; Schema: 0 OpCapability Shader - %14 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size OpEntryPoint Fragment %fragment_main "fragment_main" @@ -31,35 +30,35 @@ %vertex_point_size = OpVariable %_ptr_Output_float Output %8 %void = OpTypeVoid %9 = OpTypeFunction %void - %float_1 = OpConstant %float 1 + %float_2 = OpConstant %float 2 %_ptr_Function_float = OpTypePointer Function %float - %18 = OpTypeFunction %v4float + %16 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 %fma_c10ba3 = OpFunction %void None %9 %12 = OpLabel %res = OpVariable %_ptr_Function_float Function %8 - %13 = OpExtInst %float %14 Fma %float_1 %float_1 %float_1 - OpStore %res %13 + OpStore %res %float_2 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %18 - %20 = OpLabel - %21 = OpFunctionCall %void %fma_c10ba3 +%vertex_main_inner = OpFunction %v4float None %16 + %18 = OpLabel + %19 = OpFunctionCall %void %fma_c10ba3 OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %23 = OpLabel - %24 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %24 + %21 = OpLabel + %22 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %22 OpStore %vertex_point_size %float_1 OpReturn OpFunctionEnd %fragment_main = OpFunction %void None %9 - %26 = OpLabel - %27 = OpFunctionCall %void %fma_c10ba3 + %25 = OpLabel + %26 = OpFunctionCall %void %fma_c10ba3 OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %29 = OpLabel - %30 = OpFunctionCall %void %fma_c10ba3 + %28 = OpLabel + %29 = OpFunctionCall %void %fma_c10ba3 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.dxc.hlsl index f6e79c58c6..ad0ae1e173 100644 --- a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.dxc.hlsl @@ -1,5 +1,5 @@ void fma_c8abb3() { - float16_t res = mad(float16_t(1.0h), float16_t(1.0h), float16_t(1.0h)); + float16_t res = float16_t(2.0h); } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.glsl index 8540f07399..172c792cc5 100644 --- a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.glsl +++ b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.glsl @@ -2,7 +2,7 @@ #extension GL_AMD_gpu_shader_half_float : require void fma_c8abb3() { - float16_t res = ((1.0hf) * (1.0hf) + (1.0hf)); + float16_t res = 2.0hf; } vec4 vertex_main() { @@ -23,7 +23,7 @@ void main() { precision mediump float; void fma_c8abb3() { - float16_t res = ((1.0hf) * (1.0hf) + (1.0hf)); + float16_t res = 2.0hf; } void fragment_main() { @@ -38,7 +38,7 @@ void main() { #extension GL_AMD_gpu_shader_half_float : require void fma_c8abb3() { - float16_t res = ((1.0hf) * (1.0hf) + (1.0hf)); + float16_t res = 2.0hf; } void compute_main() { diff --git a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.msl index 1152af9b3b..58ba67a58c 100644 --- a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; void fma_c8abb3() { - half res = fma(1.0h, 1.0h, 1.0h); + half res = 2.0h; } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.spvasm index 8bde975503..16b175f5e9 100644 --- a/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/literal/fma/c8abb3.wgsl.expected.spvasm @@ -1,14 +1,13 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 34 +; Bound: 32 ; Schema: 0 OpCapability Shader OpCapability Float16 OpCapability UniformAndStorageBuffer16BitAccess OpCapability StorageBuffer16BitAccess OpCapability StorageInputOutput16 - %15 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size OpEntryPoint Fragment %fragment_main "fragment_main" @@ -36,37 +35,36 @@ %void = OpTypeVoid %9 = OpTypeFunction %void %half = OpTypeFloat 16 -%half_0x1p_0 = OpConstant %half 0x1p+0 +%half_0x1p_1 = OpConstant %half 0x1p+1 %_ptr_Function_half = OpTypePointer Function %half - %19 = OpConstantNull %half - %20 = OpTypeFunction %v4float + %17 = OpConstantNull %half + %18 = OpTypeFunction %v4float %float_1 = OpConstant %float 1 %fma_c8abb3 = OpFunction %void None %9 %12 = OpLabel - %res = OpVariable %_ptr_Function_half Function %19 - %13 = OpExtInst %half %15 Fma %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 - OpStore %res %13 + %res = OpVariable %_ptr_Function_half Function %17 + OpStore %res %half_0x1p_1 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %20 - %22 = OpLabel - %23 = OpFunctionCall %void %fma_c8abb3 +%vertex_main_inner = OpFunction %v4float None %18 + %20 = OpLabel + %21 = OpFunctionCall %void %fma_c8abb3 OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %25 = OpLabel - %26 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %26 + %23 = OpLabel + %24 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %24 OpStore %vertex_point_size %float_1 OpReturn OpFunctionEnd %fragment_main = OpFunction %void None %9 - %29 = OpLabel - %30 = OpFunctionCall %void %fma_c8abb3 + %27 = OpLabel + %28 = OpFunctionCall %void %fma_c8abb3 OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %32 = OpLabel - %33 = OpFunctionCall %void %fma_c8abb3 + %30 = OpLabel + %31 = OpFunctionCall %void %fma_c8abb3 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.dxc.hlsl index 89cd7e83db..ba5da32741 100644 --- a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.dxc.hlsl @@ -1,5 +1,5 @@ void fma_e17c5c() { - float3 res = mad((1.0f).xxx, (1.0f).xxx, (1.0f).xxx); + float3 res = (2.0f).xxx; } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.fxc.hlsl index 89cd7e83db..ba5da32741 100644 --- a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.fxc.hlsl @@ -1,5 +1,5 @@ void fma_e17c5c() { - float3 res = mad((1.0f).xxx, (1.0f).xxx, (1.0f).xxx); + float3 res = (2.0f).xxx; } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.glsl index d1b1d2ee1f..728500552a 100644 --- a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.glsl +++ b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.glsl @@ -1,7 +1,7 @@ #version 310 es void fma_e17c5c() { - vec3 res = ((vec3(1.0f)) * (vec3(1.0f)) + (vec3(1.0f))); + vec3 res = vec3(2.0f); } vec4 vertex_main() { @@ -21,7 +21,7 @@ void main() { precision mediump float; void fma_e17c5c() { - vec3 res = ((vec3(1.0f)) * (vec3(1.0f)) + (vec3(1.0f))); + vec3 res = vec3(2.0f); } void fragment_main() { @@ -35,7 +35,7 @@ void main() { #version 310 es void fma_e17c5c() { - vec3 res = ((vec3(1.0f)) * (vec3(1.0f)) + (vec3(1.0f))); + vec3 res = vec3(2.0f); } void compute_main() { diff --git a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.msl index 469605fc25..24c7af4cd0 100644 --- a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; void fma_e17c5c() { - float3 res = fma(float3(1.0f), float3(1.0f), float3(1.0f)); + float3 res = float3(2.0f); } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.spvasm index f1d385b819..87bcc0e575 100644 --- a/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/literal/fma/e17c5c.wgsl.expected.spvasm @@ -1,10 +1,9 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 34 +; Bound: 33 ; 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" @@ -32,37 +31,37 @@ %void = OpTypeVoid %9 = OpTypeFunction %void %v3float = OpTypeVector %float 3 - %float_1 = OpConstant %float 1 - %17 = OpConstantComposite %v3float %float_1 %float_1 %float_1 + %float_2 = OpConstant %float 2 + %15 = OpConstantComposite %v3float %float_2 %float_2 %float_2 %_ptr_Function_v3float = OpTypePointer Function %v3float - %20 = OpConstantNull %v3float - %21 = OpTypeFunction %v4float + %18 = OpConstantNull %v3float + %19 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 %fma_e17c5c = OpFunction %void None %9 %12 = OpLabel - %res = OpVariable %_ptr_Function_v3float Function %20 - %13 = OpExtInst %v3float %15 Fma %17 %17 %17 - OpStore %res %13 + %res = OpVariable %_ptr_Function_v3float Function %18 + OpStore %res %15 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %21 - %23 = OpLabel - %24 = OpFunctionCall %void %fma_e17c5c +%vertex_main_inner = OpFunction %v4float None %19 + %21 = OpLabel + %22 = OpFunctionCall %void %fma_e17c5c OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %26 = OpLabel - %27 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %27 + %24 = OpLabel + %25 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %25 OpStore %vertex_point_size %float_1 OpReturn OpFunctionEnd %fragment_main = OpFunction %void None %9 - %29 = OpLabel - %30 = OpFunctionCall %void %fma_e17c5c + %28 = OpLabel + %29 = OpFunctionCall %void %fma_e17c5c OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %32 = OpLabel - %33 = OpFunctionCall %void %fma_e17c5c + %31 = OpLabel + %32 = OpFunctionCall %void %fma_e17c5c OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.dxc.hlsl index 4930e9ec54..7229fbbfe6 100644 --- a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.dxc.hlsl @@ -1,5 +1,5 @@ void fma_e7abdc() { - vector res = mad((float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx, (float16_t(1.0h)).xxx); + vector res = (float16_t(2.0h)).xxx; } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.glsl index 182adeec3c..3e06476cb7 100644 --- a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.glsl +++ b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.glsl @@ -2,7 +2,7 @@ #extension GL_AMD_gpu_shader_half_float : require void fma_e7abdc() { - f16vec3 res = ((f16vec3(1.0hf)) * (f16vec3(1.0hf)) + (f16vec3(1.0hf))); + f16vec3 res = f16vec3(2.0hf); } vec4 vertex_main() { @@ -23,7 +23,7 @@ void main() { precision mediump float; void fma_e7abdc() { - f16vec3 res = ((f16vec3(1.0hf)) * (f16vec3(1.0hf)) + (f16vec3(1.0hf))); + f16vec3 res = f16vec3(2.0hf); } void fragment_main() { @@ -38,7 +38,7 @@ void main() { #extension GL_AMD_gpu_shader_half_float : require void fma_e7abdc() { - f16vec3 res = ((f16vec3(1.0hf)) * (f16vec3(1.0hf)) + (f16vec3(1.0hf))); + f16vec3 res = f16vec3(2.0hf); } void compute_main() { diff --git a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.msl index 2360c4d41f..cd1a833459 100644 --- a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; void fma_e7abdc() { - half3 res = fma(half3(1.0h), half3(1.0h), half3(1.0h)); + half3 res = half3(2.0h); } struct tint_symbol { diff --git a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.spvasm index b2d0b4c4d7..0af3de3410 100644 --- a/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/literal/fma/e7abdc.wgsl.expected.spvasm @@ -1,14 +1,13 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 36 +; Bound: 34 ; Schema: 0 OpCapability Shader OpCapability Float16 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" @@ -37,38 +36,37 @@ %9 = OpTypeFunction %void %half = OpTypeFloat 16 %v3half = OpTypeVector %half 3 -%half_0x1p_0 = OpConstant %half 0x1p+0 - %18 = OpConstantComposite %v3half %half_0x1p_0 %half_0x1p_0 %half_0x1p_0 +%half_0x1p_1 = OpConstant %half 0x1p+1 + %16 = OpConstantComposite %v3half %half_0x1p_1 %half_0x1p_1 %half_0x1p_1 %_ptr_Function_v3half = OpTypePointer Function %v3half - %21 = OpConstantNull %v3half - %22 = OpTypeFunction %v4float + %19 = OpConstantNull %v3half + %20 = OpTypeFunction %v4float %float_1 = OpConstant %float 1 %fma_e7abdc = OpFunction %void None %9 %12 = OpLabel - %res = OpVariable %_ptr_Function_v3half Function %21 - %13 = OpExtInst %v3half %16 Fma %18 %18 %18 - OpStore %res %13 + %res = OpVariable %_ptr_Function_v3half Function %19 + OpStore %res %16 OpReturn OpFunctionEnd -%vertex_main_inner = OpFunction %v4float None %22 - %24 = OpLabel - %25 = OpFunctionCall %void %fma_e7abdc +%vertex_main_inner = OpFunction %v4float None %20 + %22 = OpLabel + %23 = OpFunctionCall %void %fma_e7abdc OpReturnValue %5 OpFunctionEnd %vertex_main = OpFunction %void None %9 - %27 = OpLabel - %28 = OpFunctionCall %v4float %vertex_main_inner - OpStore %value %28 + %25 = OpLabel + %26 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %26 OpStore %vertex_point_size %float_1 OpReturn OpFunctionEnd %fragment_main = OpFunction %void None %9 - %31 = OpLabel - %32 = OpFunctionCall %void %fma_e7abdc + %29 = OpLabel + %30 = OpFunctionCall %void %fma_e7abdc OpReturn OpFunctionEnd %compute_main = OpFunction %void None %9 - %34 = OpLabel - %35 = OpFunctionCall %void %fma_e7abdc + %32 = OpLabel + %33 = OpFunctionCall %void %fma_e7abdc OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl new file mode 100644 index 0000000000..e136fa3d11 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl @@ -0,0 +1,43 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +//////////////////////////////////////////////////////////////////////////////// +// File generated by tools/src/cmd/gen +// using the template: +// test/tint/builtins/gen/gen.wgsl.tmpl +// +// Do not modify this file directly +//////////////////////////////////////////////////////////////////////////////// + + +// fn fma(vec<3, fa>, vec<3, fa>, vec<3, fa>) -> vec<3, fa> +fn fma_eb25d7() { + var res = fma(vec3(1.), vec3(1.), vec3(1.)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_eb25d7(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_eb25d7(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_eb25d7(); +} diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..7036fc68eb --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.dxc.hlsl @@ -0,0 +1,30 @@ +void fma_eb25d7() { + float3 res = (2.0f).xxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_eb25d7(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_eb25d7(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_eb25d7(); + return; +} diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..7036fc68eb --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.fxc.hlsl @@ -0,0 +1,30 @@ +void fma_eb25d7() { + float3 res = (2.0f).xxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_eb25d7(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_eb25d7(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_eb25d7(); + return; +} diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.glsl b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.glsl new file mode 100644 index 0000000000..16ddde5d98 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.glsl @@ -0,0 +1,49 @@ +#version 310 es + +void fma_eb25d7() { + vec3 res = vec3(2.0f); +} + +vec4 vertex_main() { + fma_eb25d7(); + return vec4(0.0f); +} + +void main() { + gl_PointSize = 1.0; + vec4 inner_result = vertex_main(); + gl_Position = inner_result; + gl_Position.y = -(gl_Position.y); + gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w); + return; +} +#version 310 es +precision mediump float; + +void fma_eb25d7() { + vec3 res = vec3(2.0f); +} + +void fragment_main() { + fma_eb25d7(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +void fma_eb25d7() { + vec3 res = vec3(2.0f); +} + +void compute_main() { + fma_eb25d7(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.msl b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.msl new file mode 100644 index 0000000000..5d32a93771 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; +void fma_eb25d7() { + float3 res = float3(2.0f); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner() { + fma_eb25d7(); + return float4(0.0f); +} + +vertex tint_symbol vertex_main() { + float4 const inner_result = vertex_main_inner(); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main() { + fma_eb25d7(); + return; +} + +kernel void compute_main() { + fma_eb25d7(); + return; +} + diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.spvasm new file mode 100644 index 0000000000..a3f813e7f3 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.spvasm @@ -0,0 +1,67 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 33 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size + OpEntryPoint Fragment %fragment_main "fragment_main" + OpEntryPoint GLCompute %compute_main "compute_main" + OpExecutionMode %fragment_main OriginUpperLeft + OpExecutionMode %compute_main LocalSize 1 1 1 + OpName %value "value" + OpName %vertex_point_size "vertex_point_size" + OpName %fma_eb25d7 "fma_eb25d7" + OpName %res "res" + OpName %vertex_main_inner "vertex_main_inner" + OpName %vertex_main "vertex_main" + OpName %fragment_main "fragment_main" + OpName %compute_main "compute_main" + OpDecorate %value BuiltIn Position + OpDecorate %vertex_point_size BuiltIn PointSize + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%_ptr_Output_v4float = OpTypePointer Output %v4float + %5 = OpConstantNull %v4float + %value = OpVariable %_ptr_Output_v4float Output %5 +%_ptr_Output_float = OpTypePointer Output %float + %8 = OpConstantNull %float +%vertex_point_size = OpVariable %_ptr_Output_float Output %8 + %void = OpTypeVoid + %9 = OpTypeFunction %void + %v3float = OpTypeVector %float 3 + %float_2 = OpConstant %float 2 + %15 = OpConstantComposite %v3float %float_2 %float_2 %float_2 +%_ptr_Function_v3float = OpTypePointer Function %v3float + %18 = OpConstantNull %v3float + %19 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 + %fma_eb25d7 = OpFunction %void None %9 + %12 = OpLabel + %res = OpVariable %_ptr_Function_v3float Function %18 + OpStore %res %15 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %19 + %21 = OpLabel + %22 = OpFunctionCall %void %fma_eb25d7 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %9 + %24 = OpLabel + %25 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %25 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %9 + %28 = OpLabel + %29 = OpFunctionCall %void %fma_eb25d7 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %9 + %31 = OpLabel + %32 = OpFunctionCall %void %fma_eb25d7 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.wgsl new file mode 100644 index 0000000000..7a380d4563 --- /dev/null +++ b/test/tint/builtins/gen/literal/fma/eb25d7.wgsl.expected.wgsl @@ -0,0 +1,19 @@ +fn fma_eb25d7() { + var res = fma(vec3(1.0), vec3(1.0), vec3(1.0)); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_eb25d7(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_eb25d7(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_eb25d7(); +} diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl b/test/tint/builtins/gen/var/fma/143d5d.wgsl new file mode 100644 index 0000000000..124c6a776d --- /dev/null +++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl @@ -0,0 +1,46 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +//////////////////////////////////////////////////////////////////////////////// +// File generated by tools/src/cmd/gen +// using the template: +// test/tint/builtins/gen/gen.wgsl.tmpl +// +// Do not modify this file directly +//////////////////////////////////////////////////////////////////////////////// + + +// fn fma(vec<4, fa>, vec<4, fa>, vec<4, fa>) -> vec<4, fa> +fn fma_143d5d() { + const arg_0 = vec4(1.); + const arg_1 = vec4(1.); + const arg_2 = vec4(1.); + var res = fma(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_143d5d(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_143d5d(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_143d5d(); +} diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..6c3ad71f1d --- /dev/null +++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.dxc.hlsl @@ -0,0 +1,30 @@ +void fma_143d5d() { + float4 res = (2.0f).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_143d5d(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_143d5d(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_143d5d(); + return; +} diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..6c3ad71f1d --- /dev/null +++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.fxc.hlsl @@ -0,0 +1,30 @@ +void fma_143d5d() { + float4 res = (2.0f).xxxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_143d5d(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_143d5d(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_143d5d(); + return; +} diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.glsl b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.glsl new file mode 100644 index 0000000000..7e6b4b42f2 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.glsl @@ -0,0 +1,49 @@ +#version 310 es + +void fma_143d5d() { + vec4 res = vec4(2.0f); +} + +vec4 vertex_main() { + fma_143d5d(); + return vec4(0.0f); +} + +void main() { + gl_PointSize = 1.0; + vec4 inner_result = vertex_main(); + gl_Position = inner_result; + gl_Position.y = -(gl_Position.y); + gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w); + return; +} +#version 310 es +precision mediump float; + +void fma_143d5d() { + vec4 res = vec4(2.0f); +} + +void fragment_main() { + fma_143d5d(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +void fma_143d5d() { + vec4 res = vec4(2.0f); +} + +void compute_main() { + fma_143d5d(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.msl b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.msl new file mode 100644 index 0000000000..af723352cb --- /dev/null +++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; +void fma_143d5d() { + float4 res = float4(2.0f); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner() { + fma_143d5d(); + return float4(0.0f); +} + +vertex tint_symbol vertex_main() { + float4 const inner_result = vertex_main_inner(); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main() { + fma_143d5d(); + return; +} + +kernel void compute_main() { + fma_143d5d(); + return; +} + diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.spvasm b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.spvasm new file mode 100644 index 0000000000..fba8c32478 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.spvasm @@ -0,0 +1,65 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 31 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size + OpEntryPoint Fragment %fragment_main "fragment_main" + OpEntryPoint GLCompute %compute_main "compute_main" + OpExecutionMode %fragment_main OriginUpperLeft + OpExecutionMode %compute_main LocalSize 1 1 1 + OpName %value "value" + OpName %vertex_point_size "vertex_point_size" + OpName %fma_143d5d "fma_143d5d" + OpName %res "res" + OpName %vertex_main_inner "vertex_main_inner" + OpName %vertex_main "vertex_main" + OpName %fragment_main "fragment_main" + OpName %compute_main "compute_main" + OpDecorate %value BuiltIn Position + OpDecorate %vertex_point_size BuiltIn PointSize + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%_ptr_Output_v4float = OpTypePointer Output %v4float + %5 = OpConstantNull %v4float + %value = OpVariable %_ptr_Output_v4float Output %5 +%_ptr_Output_float = OpTypePointer Output %float + %8 = OpConstantNull %float +%vertex_point_size = OpVariable %_ptr_Output_float Output %8 + %void = OpTypeVoid + %9 = OpTypeFunction %void + %float_2 = OpConstant %float 2 + %14 = OpConstantComposite %v4float %float_2 %float_2 %float_2 %float_2 +%_ptr_Function_v4float = OpTypePointer Function %v4float + %17 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 + %fma_143d5d = OpFunction %void None %9 + %12 = OpLabel + %res = OpVariable %_ptr_Function_v4float Function %5 + OpStore %res %14 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %17 + %19 = OpLabel + %20 = OpFunctionCall %void %fma_143d5d + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %9 + %22 = OpLabel + %23 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %23 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %9 + %26 = OpLabel + %27 = OpFunctionCall %void %fma_143d5d + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %9 + %29 = OpLabel + %30 = OpFunctionCall %void %fma_143d5d + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.wgsl b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.wgsl new file mode 100644 index 0000000000..03e642675c --- /dev/null +++ b/test/tint/builtins/gen/var/fma/143d5d.wgsl.expected.wgsl @@ -0,0 +1,22 @@ +fn fma_143d5d() { + const arg_0 = vec4(1.0); + const arg_1 = vec4(1.0); + const arg_2 = vec4(1.0); + var res = fma(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_143d5d(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_143d5d(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_143d5d(); +} diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl b/test/tint/builtins/gen/var/fma/1f5084.wgsl new file mode 100644 index 0000000000..20fb0932ec --- /dev/null +++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl @@ -0,0 +1,46 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +//////////////////////////////////////////////////////////////////////////////// +// File generated by tools/src/cmd/gen +// using the template: +// test/tint/builtins/gen/gen.wgsl.tmpl +// +// Do not modify this file directly +//////////////////////////////////////////////////////////////////////////////// + + +// fn fma(vec<2, fa>, vec<2, fa>, vec<2, fa>) -> vec<2, fa> +fn fma_1f5084() { + const arg_0 = vec2(1.); + const arg_1 = vec2(1.); + const arg_2 = vec2(1.); + var res = fma(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_1f5084(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_1f5084(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_1f5084(); +} diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..e7655ab50d --- /dev/null +++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.dxc.hlsl @@ -0,0 +1,30 @@ +void fma_1f5084() { + float2 res = (2.0f).xx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_1f5084(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_1f5084(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_1f5084(); + return; +} diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..e7655ab50d --- /dev/null +++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.fxc.hlsl @@ -0,0 +1,30 @@ +void fma_1f5084() { + float2 res = (2.0f).xx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_1f5084(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_1f5084(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_1f5084(); + return; +} diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.glsl b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.glsl new file mode 100644 index 0000000000..215c7401c2 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.glsl @@ -0,0 +1,49 @@ +#version 310 es + +void fma_1f5084() { + vec2 res = vec2(2.0f); +} + +vec4 vertex_main() { + fma_1f5084(); + return vec4(0.0f); +} + +void main() { + gl_PointSize = 1.0; + vec4 inner_result = vertex_main(); + gl_Position = inner_result; + gl_Position.y = -(gl_Position.y); + gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w); + return; +} +#version 310 es +precision mediump float; + +void fma_1f5084() { + vec2 res = vec2(2.0f); +} + +void fragment_main() { + fma_1f5084(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +void fma_1f5084() { + vec2 res = vec2(2.0f); +} + +void compute_main() { + fma_1f5084(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.msl b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.msl new file mode 100644 index 0000000000..bc3b0f35c6 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; +void fma_1f5084() { + float2 res = float2(2.0f); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner() { + fma_1f5084(); + return float4(0.0f); +} + +vertex tint_symbol vertex_main() { + float4 const inner_result = vertex_main_inner(); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main() { + fma_1f5084(); + return; +} + +kernel void compute_main() { + fma_1f5084(); + return; +} + diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.spvasm b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.spvasm new file mode 100644 index 0000000000..3c5ab1b481 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.spvasm @@ -0,0 +1,67 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 33 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size + OpEntryPoint Fragment %fragment_main "fragment_main" + OpEntryPoint GLCompute %compute_main "compute_main" + OpExecutionMode %fragment_main OriginUpperLeft + OpExecutionMode %compute_main LocalSize 1 1 1 + OpName %value "value" + OpName %vertex_point_size "vertex_point_size" + OpName %fma_1f5084 "fma_1f5084" + OpName %res "res" + OpName %vertex_main_inner "vertex_main_inner" + OpName %vertex_main "vertex_main" + OpName %fragment_main "fragment_main" + OpName %compute_main "compute_main" + OpDecorate %value BuiltIn Position + OpDecorate %vertex_point_size BuiltIn PointSize + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%_ptr_Output_v4float = OpTypePointer Output %v4float + %5 = OpConstantNull %v4float + %value = OpVariable %_ptr_Output_v4float Output %5 +%_ptr_Output_float = OpTypePointer Output %float + %8 = OpConstantNull %float +%vertex_point_size = OpVariable %_ptr_Output_float Output %8 + %void = OpTypeVoid + %9 = OpTypeFunction %void + %v2float = OpTypeVector %float 2 + %float_2 = OpConstant %float 2 + %15 = OpConstantComposite %v2float %float_2 %float_2 +%_ptr_Function_v2float = OpTypePointer Function %v2float + %18 = OpConstantNull %v2float + %19 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 + %fma_1f5084 = OpFunction %void None %9 + %12 = OpLabel + %res = OpVariable %_ptr_Function_v2float Function %18 + OpStore %res %15 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %19 + %21 = OpLabel + %22 = OpFunctionCall %void %fma_1f5084 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %9 + %24 = OpLabel + %25 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %25 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %9 + %28 = OpLabel + %29 = OpFunctionCall %void %fma_1f5084 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %9 + %31 = OpLabel + %32 = OpFunctionCall %void %fma_1f5084 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.wgsl b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.wgsl new file mode 100644 index 0000000000..a339a368d2 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/1f5084.wgsl.expected.wgsl @@ -0,0 +1,22 @@ +fn fma_1f5084() { + const arg_0 = vec2(1.0); + const arg_1 = vec2(1.0); + const arg_2 = vec2(1.0); + var res = fma(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_1f5084(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_1f5084(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_1f5084(); +} diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl b/test/tint/builtins/gen/var/fma/466442.wgsl new file mode 100644 index 0000000000..ef43d65498 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/466442.wgsl @@ -0,0 +1,46 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +//////////////////////////////////////////////////////////////////////////////// +// File generated by tools/src/cmd/gen +// using the template: +// test/tint/builtins/gen/gen.wgsl.tmpl +// +// Do not modify this file directly +//////////////////////////////////////////////////////////////////////////////// + + +// fn fma(fa, fa, fa) -> fa +fn fma_466442() { + const arg_0 = 1.; + const arg_1 = 1.; + const arg_2 = 1.; + var res = fma(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_466442(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_466442(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_466442(); +} diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..c070a74b09 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.dxc.hlsl @@ -0,0 +1,30 @@ +void fma_466442() { + float res = 2.0f; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_466442(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_466442(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_466442(); + return; +} diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..c070a74b09 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.fxc.hlsl @@ -0,0 +1,30 @@ +void fma_466442() { + float res = 2.0f; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_466442(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_466442(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_466442(); + return; +} diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl.expected.glsl b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.glsl new file mode 100644 index 0000000000..c094238683 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.glsl @@ -0,0 +1,49 @@ +#version 310 es + +void fma_466442() { + float res = 2.0f; +} + +vec4 vertex_main() { + fma_466442(); + return vec4(0.0f); +} + +void main() { + gl_PointSize = 1.0; + vec4 inner_result = vertex_main(); + gl_Position = inner_result; + gl_Position.y = -(gl_Position.y); + gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w); + return; +} +#version 310 es +precision mediump float; + +void fma_466442() { + float res = 2.0f; +} + +void fragment_main() { + fma_466442(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +void fma_466442() { + float res = 2.0f; +} + +void compute_main() { + fma_466442(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl.expected.msl b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.msl new file mode 100644 index 0000000000..26b2b96a55 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; +void fma_466442() { + float res = 2.0f; +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner() { + fma_466442(); + return float4(0.0f); +} + +vertex tint_symbol vertex_main() { + float4 const inner_result = vertex_main_inner(); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main() { + fma_466442(); + return; +} + +kernel void compute_main() { + fma_466442(); + return; +} + diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl.expected.spvasm b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.spvasm new file mode 100644 index 0000000000..55e5e1f0b8 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.spvasm @@ -0,0 +1,64 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 30 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size + OpEntryPoint Fragment %fragment_main "fragment_main" + OpEntryPoint GLCompute %compute_main "compute_main" + OpExecutionMode %fragment_main OriginUpperLeft + OpExecutionMode %compute_main LocalSize 1 1 1 + OpName %value "value" + OpName %vertex_point_size "vertex_point_size" + OpName %fma_466442 "fma_466442" + OpName %res "res" + OpName %vertex_main_inner "vertex_main_inner" + OpName %vertex_main "vertex_main" + OpName %fragment_main "fragment_main" + OpName %compute_main "compute_main" + OpDecorate %value BuiltIn Position + OpDecorate %vertex_point_size BuiltIn PointSize + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%_ptr_Output_v4float = OpTypePointer Output %v4float + %5 = OpConstantNull %v4float + %value = OpVariable %_ptr_Output_v4float Output %5 +%_ptr_Output_float = OpTypePointer Output %float + %8 = OpConstantNull %float +%vertex_point_size = OpVariable %_ptr_Output_float Output %8 + %void = OpTypeVoid + %9 = OpTypeFunction %void + %float_2 = OpConstant %float 2 +%_ptr_Function_float = OpTypePointer Function %float + %16 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 + %fma_466442 = OpFunction %void None %9 + %12 = OpLabel + %res = OpVariable %_ptr_Function_float Function %8 + OpStore %res %float_2 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %16 + %18 = OpLabel + %19 = OpFunctionCall %void %fma_466442 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %9 + %21 = OpLabel + %22 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %22 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %9 + %25 = OpLabel + %26 = OpFunctionCall %void %fma_466442 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %9 + %28 = OpLabel + %29 = OpFunctionCall %void %fma_466442 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/var/fma/466442.wgsl.expected.wgsl b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.wgsl new file mode 100644 index 0000000000..1ae85a0722 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/466442.wgsl.expected.wgsl @@ -0,0 +1,22 @@ +fn fma_466442() { + const arg_0 = 1.0; + const arg_1 = 1.0; + const arg_2 = 1.0; + var res = fma(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_466442(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_466442(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_466442(); +} diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl b/test/tint/builtins/gen/var/fma/eb25d7.wgsl new file mode 100644 index 0000000000..4f564147d4 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl @@ -0,0 +1,46 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +//////////////////////////////////////////////////////////////////////////////// +// File generated by tools/src/cmd/gen +// using the template: +// test/tint/builtins/gen/gen.wgsl.tmpl +// +// Do not modify this file directly +//////////////////////////////////////////////////////////////////////////////// + + +// fn fma(vec<3, fa>, vec<3, fa>, vec<3, fa>) -> vec<3, fa> +fn fma_eb25d7() { + const arg_0 = vec3(1.); + const arg_1 = vec3(1.); + const arg_2 = vec3(1.); + var res = fma(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_eb25d7(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_eb25d7(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_eb25d7(); +} diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..7036fc68eb --- /dev/null +++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.dxc.hlsl @@ -0,0 +1,30 @@ +void fma_eb25d7() { + float3 res = (2.0f).xxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_eb25d7(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_eb25d7(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_eb25d7(); + return; +} diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..7036fc68eb --- /dev/null +++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.fxc.hlsl @@ -0,0 +1,30 @@ +void fma_eb25d7() { + float3 res = (2.0f).xxx; +} + +struct tint_symbol { + float4 value : SV_Position; +}; + +float4 vertex_main_inner() { + fma_eb25d7(); + return (0.0f).xxxx; +} + +tint_symbol vertex_main() { + const float4 inner_result = vertex_main_inner(); + tint_symbol wrapper_result = (tint_symbol)0; + wrapper_result.value = inner_result; + return wrapper_result; +} + +void fragment_main() { + fma_eb25d7(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + fma_eb25d7(); + return; +} diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.glsl b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.glsl new file mode 100644 index 0000000000..16ddde5d98 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.glsl @@ -0,0 +1,49 @@ +#version 310 es + +void fma_eb25d7() { + vec3 res = vec3(2.0f); +} + +vec4 vertex_main() { + fma_eb25d7(); + return vec4(0.0f); +} + +void main() { + gl_PointSize = 1.0; + vec4 inner_result = vertex_main(); + gl_Position = inner_result; + gl_Position.y = -(gl_Position.y); + gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w); + return; +} +#version 310 es +precision mediump float; + +void fma_eb25d7() { + vec3 res = vec3(2.0f); +} + +void fragment_main() { + fma_eb25d7(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +void fma_eb25d7() { + vec3 res = vec3(2.0f); +} + +void compute_main() { + fma_eb25d7(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.msl b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.msl new file mode 100644 index 0000000000..5d32a93771 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; +void fma_eb25d7() { + float3 res = float3(2.0f); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner() { + fma_eb25d7(); + return float4(0.0f); +} + +vertex tint_symbol vertex_main() { + float4 const inner_result = vertex_main_inner(); + tint_symbol wrapper_result = {}; + wrapper_result.value = inner_result; + return wrapper_result; +} + +fragment void fragment_main() { + fma_eb25d7(); + return; +} + +kernel void compute_main() { + fma_eb25d7(); + return; +} + diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.spvasm b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.spvasm new file mode 100644 index 0000000000..a3f813e7f3 --- /dev/null +++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.spvasm @@ -0,0 +1,67 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 33 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size + OpEntryPoint Fragment %fragment_main "fragment_main" + OpEntryPoint GLCompute %compute_main "compute_main" + OpExecutionMode %fragment_main OriginUpperLeft + OpExecutionMode %compute_main LocalSize 1 1 1 + OpName %value "value" + OpName %vertex_point_size "vertex_point_size" + OpName %fma_eb25d7 "fma_eb25d7" + OpName %res "res" + OpName %vertex_main_inner "vertex_main_inner" + OpName %vertex_main "vertex_main" + OpName %fragment_main "fragment_main" + OpName %compute_main "compute_main" + OpDecorate %value BuiltIn Position + OpDecorate %vertex_point_size BuiltIn PointSize + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 +%_ptr_Output_v4float = OpTypePointer Output %v4float + %5 = OpConstantNull %v4float + %value = OpVariable %_ptr_Output_v4float Output %5 +%_ptr_Output_float = OpTypePointer Output %float + %8 = OpConstantNull %float +%vertex_point_size = OpVariable %_ptr_Output_float Output %8 + %void = OpTypeVoid + %9 = OpTypeFunction %void + %v3float = OpTypeVector %float 3 + %float_2 = OpConstant %float 2 + %15 = OpConstantComposite %v3float %float_2 %float_2 %float_2 +%_ptr_Function_v3float = OpTypePointer Function %v3float + %18 = OpConstantNull %v3float + %19 = OpTypeFunction %v4float + %float_1 = OpConstant %float 1 + %fma_eb25d7 = OpFunction %void None %9 + %12 = OpLabel + %res = OpVariable %_ptr_Function_v3float Function %18 + OpStore %res %15 + OpReturn + OpFunctionEnd +%vertex_main_inner = OpFunction %v4float None %19 + %21 = OpLabel + %22 = OpFunctionCall %void %fma_eb25d7 + OpReturnValue %5 + OpFunctionEnd +%vertex_main = OpFunction %void None %9 + %24 = OpLabel + %25 = OpFunctionCall %v4float %vertex_main_inner + OpStore %value %25 + OpStore %vertex_point_size %float_1 + OpReturn + OpFunctionEnd +%fragment_main = OpFunction %void None %9 + %28 = OpLabel + %29 = OpFunctionCall %void %fma_eb25d7 + OpReturn + OpFunctionEnd +%compute_main = OpFunction %void None %9 + %31 = OpLabel + %32 = OpFunctionCall %void %fma_eb25d7 + OpReturn + OpFunctionEnd diff --git a/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.wgsl b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.wgsl new file mode 100644 index 0000000000..216968e9ca --- /dev/null +++ b/test/tint/builtins/gen/var/fma/eb25d7.wgsl.expected.wgsl @@ -0,0 +1,22 @@ +fn fma_eb25d7() { + const arg_0 = vec3(1.0); + const arg_1 = vec3(1.0); + const arg_2 = vec3(1.0); + var res = fma(arg_0, arg_1, arg_2); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + fma_eb25d7(); + return vec4(); +} + +@fragment +fn fragment_main() { + fma_eb25d7(); +} + +@compute @workgroup_size(1) +fn compute_main() { + fma_eb25d7(); +}