From 8605896d7b0228ec8da50250d8ded01455b38fd6 Mon Sep 17 00:00:00 2001 From: Vasilii Filippov Date: Thu, 30 Nov 2023 09:47:37 +0100 Subject: [PATCH] [Find 2.0] Bias for Find 2.0 fusion (#2525) --- include/miopen/miopen.h | 13 +++++++++ src/api/find2_0_commons.cpp | 24 ++++++++++++++-- src/include/miopen/problem.hpp | 10 ++++++- src/problem.cpp | 52 +++++++++++++++++++++++----------- src/solution.cpp | 9 ++++-- test/gtest/cba_find2.hpp | 17 +++++++++-- test/gtest/cba_find2_infer.cpp | 1 + 7 files changed, 102 insertions(+), 24 deletions(-) diff --git a/include/miopen/miopen.h b/include/miopen/miopen.h index 37654d4aac..c359c3c5b1 100644 --- a/include/miopen/miopen.h +++ b/include/miopen/miopen.h @@ -5234,6 +5234,9 @@ typedef enum miopenTensorActivationY = 5, miopenTensorActivationDX = 6, miopenTensorActivationDY = 7, + miopenTensorBiasX = 8, + miopenTensorBiasY = 9, + miopenTensorBias = 10, #endif } miopenTensorArgumentId_t; @@ -5507,6 +5510,16 @@ miopenCreateActivationProblem(miopenProblem_t* problem, */ MIOPEN_EXPORT miopenStatus_t miopenFuseProblems(miopenProblem_t problem1, miopenProblem_t problem2); +/*! @brief Initializes a problem object describing an bias operation. + * @note As of now there is no way to actually get any solution for this kind of problems. + * + * @param problem Pointer to the problem to initialize + * @param direction Direction of the operation + * @return miopenStatus_t + */ +MIOPEN_EXPORT miopenStatus_t miopenCreateBiasProblem(miopenProblem_t* problem, + miopenProblemDirection_t direction); + #endif /** @} */ diff --git a/src/api/find2_0_commons.cpp b/src/api/find2_0_commons.cpp index 3f864a930c..5428edf3a0 100644 --- a/src/api/find2_0_commons.cpp +++ b/src/api/find2_0_commons.cpp @@ -28,7 +28,6 @@ #include #include -#include #include #include #include @@ -62,7 +61,7 @@ miopenStatus_t miopenCreateConvProblem(miopenProblem_t* problem, miopenConvolutionDescriptor_t operatorDesc, miopenProblemDirection_t direction) { - MIOPEN_LOG_FUNCTION(problem); + MIOPEN_LOG_FUNCTION(problem, operatorDesc, direction); return MakeProblem(problem, operatorDesc, direction); } @@ -70,10 +69,26 @@ miopenStatus_t miopenCreateActivationProblem(miopenProblem_t* problem, miopenActivationDescriptor_t operatorDesc, miopenProblemDirection_t direction) { - MIOPEN_LOG_FUNCTION(problem); + MIOPEN_LOG_FUNCTION(problem, operatorDesc, direction); return MakeProblem(problem, operatorDesc, direction); } +miopenStatus_t miopenCreateBiasProblem(miopenProblem_t* problem, miopenProblemDirection_t direction) +{ + MIOPEN_LOG_FUNCTION(problem, direction); + + return miopen::try_([&] { + miopen::deref(problem) = new miopen::ProblemContainer(); + auto& container_deref = miopen::deref(*problem); + + container_deref.item = miopen::Problem(); + auto& problem_deref = boost::get(container_deref.item); + + problem_deref.SetOperatorDescriptor(miopen::BiasDescriptor{}); + problem_deref.SetDirection(direction); + }); +} + miopenStatus_t miopenFuseProblems(miopenProblem_t problem1, miopenProblem_t problem2) { MIOPEN_LOG_FUNCTION(problem1, problem2); @@ -245,6 +260,9 @@ inline std::ostream& operator<<(std::ostream& stream, const miopenTensorArgument case miopenTensorActivationDX: stream << "ActivDX"; break; case miopenTensorActivationY: stream << "ActivY"; break; case miopenTensorActivationDY: stream << "ActivDY"; break; + case miopenTensorBias: stream << "Bias"; break; + case miopenTensorBiasX: stream << "BiasX"; break; + case miopenTensorBiasY: stream << "BiasY"; break; case miopenTensorArgumentIdInvalid: stream << "Invalid"; break; } diff --git a/src/include/miopen/problem.hpp b/src/include/miopen/problem.hpp index 6f66a5f9db..2900008785 100644 --- a/src/include/miopen/problem.hpp +++ b/src/include/miopen/problem.hpp @@ -59,7 +59,13 @@ namespace conv { struct ProblemDescription; } // namespace conv -using OperatorDescriptor = boost::variant; +struct BiasDescriptor +{ +}; + +// The order of types is important for deserialization and should be preserved between releases. +using OperatorDescriptor = + boost::variant; struct Problem { @@ -196,7 +202,9 @@ struct FusedProblem struct ProblemContainer : miopenProblem { + // The order of types is important for deserialization and should be preserved between releases. using Item = boost::variant; + Item item; ProblemContainer() = default; diff --git a/src/problem.cpp b/src/problem.cpp index 07778826f1..2ea839391c 100644 --- a/src/problem.cpp +++ b/src/problem.cpp @@ -177,6 +177,9 @@ Problem::FindSolutions(Handle& handle, const FindOptions& options, std::size_t m }, [&](const ActivationDescriptor& /*op_desc*/) -> std::vector { MIOPEN_THROW(miopenStatusNotImplemented); + }, + [&](const BiasDescriptor& /*op_desc*/) -> std::vector { + MIOPEN_THROW(miopenStatusNotImplemented); }), operator_descriptor); @@ -452,7 +455,8 @@ void Problem::LogDriverCommand() const { const auto log_function = boost::hof::match([&](const ConvolutionDescriptor& op_desc) { LogDriverCommand(op_desc); }, - [&](const ActivationDescriptor& op_desc) { LogDriverCommand(op_desc); }); + [&](const ActivationDescriptor& op_desc) { LogDriverCommand(op_desc); }, + [&](const BiasDescriptor&) {}); boost::apply_visitor(log_function, operator_descriptor); } @@ -475,6 +479,10 @@ void Problem::LogDriverCommand(const ActivationDescriptor& descriptor) const miopen::debug::LogCmdActivation(x_desc, descriptor, direction == miopenProblemDirectionForward); } +void to_json(nlohmann::json& json, const BiasDescriptor&) { json = nlohmann::json{}; } + +void from_json(const nlohmann::json&, BiasDescriptor&) {} + void to_json(nlohmann::json& json, const Problem& problem) { json = nlohmann::json{ @@ -556,26 +564,28 @@ void Problem::CalculateOutput() if(!HasInput()) return; - boost::apply_visitor(boost::hof::match( - [&](const ConvolutionDescriptor& conv) { - const auto& in = GetInput(); - conv.GetForwardOutputTensor( - in, - GetTensorDescriptorChecked(miopenTensorConvolutionW, - "miopenTensorConvolutionW"), - in.GetType()); - }, - [&](const ActivationDescriptor&) { - RegisterTensorDescriptor(GetOutputId(), GetInput()); - }), - operator_descriptor); + boost::apply_visitor( + boost::hof::match( + [&](const ConvolutionDescriptor& conv) { + const auto& in = GetInput(); + conv.GetForwardOutputTensor(in, + GetTensorDescriptorChecked(miopenTensorConvolutionW, + "miopenTensorConvolutionW"), + in.GetType()); + }, + [&](const ActivationDescriptor&) { + RegisterTensorDescriptor(GetOutputId(), GetInput()); + }, + [&](const BiasDescriptor&) { RegisterTensorDescriptor(GetOutputId(), GetInput()); }), + operator_descriptor); } miopenTensorArgumentId_t Problem::GetInputId() const { return boost::apply_visitor( boost::hof::match([](const ConvolutionDescriptor&) { return miopenTensorConvolutionX; }, - [](const ActivationDescriptor&) { return miopenTensorActivationX; }), + [](const ActivationDescriptor&) { return miopenTensorActivationX; }, + [](const BiasDescriptor&) { return miopenTensorBiasX; }), operator_descriptor); } @@ -583,7 +593,8 @@ miopenTensorArgumentId_t Problem::GetOutputId() const { return boost::apply_visitor( boost::hof::match([](const ConvolutionDescriptor&) { return miopenTensorConvolutionY; }, - [](const ActivationDescriptor&) { return miopenTensorActivationY; }), + [](const ActivationDescriptor&) { return miopenTensorActivationY; }, + [](const BiasDescriptor&) { return miopenTensorBiasY; }), operator_descriptor); } @@ -664,6 +675,10 @@ void FusedProblem::AddProblemToPlan(FusionPlanDescriptor& plan, const Problem& p plan.AddOp(std::make_shared(activ_desc.GetMode())); else plan.AddOp(std::make_shared(activ_desc.GetMode())); + }, + [&](const BiasDescriptor&) { + plan.AddOp(std::make_shared( + problem.GetTensorDescriptorChecked(miopenTensorBias, "miopenTensorBias"))); }), problem.operator_descriptor); } @@ -721,6 +736,11 @@ fusion::FusionInvokeParams FusedProblem::MakeInvokeParams( std::make_unique( y, x, alpha, beta, gamma)); } + }, + [&](const BiasDescriptor&) { + const auto bias_ptr = buffers.at(miopenTensorBias); + operator_args.params.emplace_back( + std::make_unique(bias_ptr)); }), problem.operator_descriptor); } diff --git a/src/solution.cpp b/src/solution.cpp index e0d658306f..e146191639 100644 --- a/src/solution.cpp +++ b/src/solution.cpp @@ -72,6 +72,9 @@ void Solution::Run(Handle& handle, }, [&](const ActivationDescriptor& /*op_desc*/) { MIOPEN_THROW(miopenStatusNotImplemented); + }, + [&](const BiasDescriptor& /*op_desc*/) { + MIOPEN_THROW(miopenStatusNotImplemented); }), problem_.GetOperatorDescriptor()); }, @@ -109,8 +112,10 @@ void Solution::LogDriverCommand(const ActivationDescriptor& desc) const void Solution::LogDriverCommand(const Problem& problem_) const { - boost::apply_visitor([&](const auto& op_desc) { LogDriverCommand(op_desc); }, - problem_.GetOperatorDescriptor()); + boost::apply_visitor( + boost::hof::match([&](const BiasDescriptor&) { /* \todo: think on how to log bias */ }, + [&](const auto& op_desc) { LogDriverCommand(op_desc); }), + problem_.GetOperatorDescriptor()); } void Solution::LogDriverCommand(const FusedProblem& problem_) const diff --git a/test/gtest/cba_find2.hpp b/test/gtest/cba_find2.hpp index 2ed673e17b..7e2b443d6e 100644 --- a/test/gtest/cba_find2.hpp +++ b/test/gtest/cba_find2.hpp @@ -87,7 +87,7 @@ struct ConvBiasActivInferFind2Test // Setup the fusion problem fused_problem = miopen::FusedProblem{{ MakeConvProblem(), - // MakeBiasProblem(), + MakeBiasProblem(), MakeActivationProblem(), }}; @@ -121,6 +121,14 @@ struct ConvBiasActivInferFind2Test return problem; } + [[nodiscard]] miopen::Problem MakeBiasProblem() const + { + auto problem = miopen::Problem{}; + problem.SetOperatorDescriptor(miopen::BiasDescriptor{}); + problem.RegisterTensorDescriptor(miopenTensorBias, bias.desc); + return problem; + } + [[nodiscard]] miopen::Problem MakeActivationProblem() const { auto problem = miopen::Problem{}; @@ -147,6 +155,11 @@ struct ConvBiasActivInferFind2Test EXPECT_EQ(desc, cfsb::output.desc); return cfsb::out_dev.get(); } + if(id == miopenTensorBias) + { + EXPECT_EQ(desc, bias.desc); + return bias_dev.get(); + } MIOPEN_THROW(miopenStatusInternalError); }, params); @@ -159,7 +172,7 @@ struct ConvBiasActivInferFind2Test cpu_values_calculated = true; cfsb::TearDownConv(); - // cpu_bias_forward(cfsb::ref_out, bias); + cpu_bias_forward(cfsb::ref_out, bias); activationHostInfer(activ_mode, activ_gamma, diff --git a/test/gtest/cba_find2_infer.cpp b/test/gtest/cba_find2_infer.cpp index 2057f86ad0..67da2c2053 100644 --- a/test/gtest/cba_find2_infer.cpp +++ b/test/gtest/cba_find2_infer.cpp @@ -143,6 +143,7 @@ TEST_P(ConvBiasActivFind2InferTestFloatFusionFind, ConvBiasActivFind2Float_testF {miopenTensorConvolutionX, in_dev.get()}, {miopenTensorConvolutionW, wei_dev.get()}, {miopenTensorActivationY, out_dev.get()}, + {miopenTensorBias, bias_dev.get()}, }; for(auto& solution : solutions)