Skip to content

Commit

Permalink
[Find 2.0] Bias for Find 2.0 fusion (#2525)
Browse files Browse the repository at this point in the history
  • Loading branch information
@DrizztDoUrden
DrizztDoUrden authored Nov 30, 2023
1 parent 1a4b476 commit 8605896
Show file tree
Hide file tree
Showing 7 changed files with 102 additions and 24 deletions.
13 changes: 13 additions & 0 deletions include/miopen/miopen.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -5234,6 +5234,9 @@ typedef enum
miopenTensorActivationY = 5,
miopenTensorActivationDX = 6,
miopenTensorActivationDY = 7,
miopenTensorBiasX = 8,
miopenTensorBiasY = 9,
miopenTensorBias = 10,
#endif
} miopenTensorArgumentId_t;

Expand Down Expand Up @@ -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

/** @} */
Expand Down
24 changes: 21 additions & 3 deletions src/api/find2_0_commons.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -28,7 +28,6 @@

#include <miopen/common.hpp>
#include <miopen/errors.hpp>
#include <miopen/handle.hpp>
#include <miopen/logger.hpp>
#include <miopen/problem.hpp>
#include <miopen/search_options.hpp>
Expand Down Expand Up @@ -62,18 +61,34 @@ 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);
}

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<miopen::Problem>(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);
Expand Down Expand Up @@ -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;
}

Expand Down
10 changes: 9 additions & 1 deletion src/include/miopen/problem.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -59,7 +59,13 @@ namespace conv {
struct ProblemDescription;
} // namespace conv

using OperatorDescriptor = boost::variant<ConvolutionDescriptor, ActivationDescriptor>;
struct BiasDescriptor
{
};

// The order of types is important for deserialization and should be preserved between releases.
using OperatorDescriptor =
boost::variant<ConvolutionDescriptor, ActivationDescriptor, BiasDescriptor>;

struct Problem
{
Expand Down Expand Up @@ -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<Problem, FusedProblem>;

Item item;

ProblemContainer() = default;
Expand Down
52 changes: 36 additions & 16 deletions src/problem.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -177,6 +177,9 @@ Problem::FindSolutions(Handle& handle, const FindOptions& options, std::size_t m
},
[&](const ActivationDescriptor& /*op_desc*/) -> std::vector<Solution> {
MIOPEN_THROW(miopenStatusNotImplemented);
},
[&](const BiasDescriptor& /*op_desc*/) -> std::vector<Solution> {
MIOPEN_THROW(miopenStatusNotImplemented);
}),
operator_descriptor);

Expand Down Expand Up @@ -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);
}
Expand All @@ -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{
Expand Down Expand Up @@ -556,34 +564,37 @@ 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);
}

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);
}

Expand Down Expand Up @@ -664,6 +675,10 @@ void FusedProblem::AddProblemToPlan(FusionPlanDescriptor& plan, const Problem& p
plan.AddOp(std::make_shared<ActivFwdFusionOpDescriptor>(activ_desc.GetMode()));
else
plan.AddOp(std::make_shared<ActivBwdFusionOpDescriptor>(activ_desc.GetMode()));
},
[&](const BiasDescriptor&) {
plan.AddOp(std::make_shared<BiasFusionOpDescriptor>(
problem.GetTensorDescriptorChecked(miopenTensorBias, "miopenTensorBias")));
}),
problem.operator_descriptor);
}
Expand Down Expand Up @@ -721,6 +736,11 @@ fusion::FusionInvokeParams FusedProblem::MakeInvokeParams(
std::make_unique<miopen::fusion::ActivationBwdOpInvokeParam>(
y, x, alpha, beta, gamma));
}
},
[&](const BiasDescriptor&) {
const auto bias_ptr = buffers.at(miopenTensorBias);
operator_args.params.emplace_back(
std::make_unique<miopen::fusion::BiasOpInvokeParam>(bias_ptr));
}),
problem.operator_descriptor);
}
Expand Down
9 changes: 7 additions & 2 deletions src/solution.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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());
},
Expand Down Expand Up @@ -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
Expand Down
17 changes: 15 additions & 2 deletions test/gtest/cba_find2.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -87,7 +87,7 @@ struct ConvBiasActivInferFind2Test
// Setup the fusion problem
fused_problem = miopen::FusedProblem{{
MakeConvProblem(),
// MakeBiasProblem(),
MakeBiasProblem(),
MakeActivationProblem(),
}};

Expand Down Expand Up @@ -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{};
Expand All @@ -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);
Expand All @@ -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,
Expand Down
1 change: 1 addition & 0 deletions test/gtest/cba_find2_infer.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down

0 comments on commit 8605896

Please sign in to comment.