Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[MLIR] Implement tuning - step 3: bwd, nonxdlops + xdlops #1152

Merged
merged 1 commit into from
Sep 13, 2021
Merged

[MLIR] Implement tuning - step 3: bwd, nonxdlops + xdlops #1152

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
18 changes: 16 additions & 2 deletions src/include/miopen/solver.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1387,13 +1387,27 @@ struct ConvHipImplicitGemmMlirCppBwd : SolverBase<ConvolutionContext>
struct ConvMlirIgemmBwd : SolverBase<ConvolutionContext>
{
bool IsApplicable(const ConvolutionContext& ctx) const;
ConvSolution GetSolution(const ConvolutionContext& ctx) const;
PerformanceConvMlirIgemm GetPerformanceConfig(const ConvolutionContext& ctx) const;
bool IsValidPerformanceConfig(const ConvolutionContext& ctx,
const PerformanceConvMlirIgemm& config) const;
PerformanceConvMlirIgemm Search(const ConvolutionContext&,
const AnyInvokeParams& invoke_ctx) const;
ConvSolution GetSolution(const ConvolutionContext& ctx,
const PerformanceConvMlirIgemm& config,
bool disableConfigOverrideFromEnv = false) const;
};

struct ConvMlirIgemmBwdXdlops : SolverBase<ConvolutionContext>
{
bool IsApplicable(const ConvolutionContext& ctx) const;
ConvSolution GetSolution(const ConvolutionContext& ctx) const;
PerformanceConvMlirIgemmXdlops GetPerformanceConfig(const ConvolutionContext& ctx) const;
bool IsValidPerformanceConfig(const ConvolutionContext& ctx,
const PerformanceConvMlirIgemmXdlops& config) const;
PerformanceConvMlirIgemmXdlops Search(const ConvolutionContext&,
const AnyInvokeParams& invoke_ctx) const;
ConvSolution GetSolution(const ConvolutionContext& ctx,
const PerformanceConvMlirIgemmXdlops& config,
bool disableConfigOverrideFromEnv = false) const;
};

struct ConvHipImplicitGemmBwdDataV4R1 : SolverBase<ConvolutionContext>
Expand Down
7 changes: 2 additions & 5 deletions src/include/miopen/solver/mlir_common.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,16 +35,13 @@ namespace miopen {
namespace solver {
namespace mlir {

std::string InsertGToLayout(const std::string& layout, char dim);
std::string GetKernelName(const ConvolutionContext& ctx, bool is_xdlops, int kernel_id = 0);

std::string ConstructBuildOptions(const ConvolutionContext& ctx,
const std::string& operation,
const std::string& kernel_name,
bool is_xdlops,
int kernel_id = 0);

std::string ConstructBuildOptions(const ConvolutionContext& ctx,
const std::string& operation,
const std::string& kernel_name,
const std::string& config,
bool is_xdlops,
int kernel_id = 0);
Expand Down
58 changes: 37 additions & 21 deletions src/solver/conv_mlir_igemm_bwd.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,7 @@
#include <miopen/conv/invokers/mlir_impl_gemm.hpp>
#include <miopen/config.h>
#include <miopen/env.hpp>
#include <miopen/generic_search.hpp>
#include <miopen/solver.hpp>
#include <miopen/solver/implicitgemm_util.hpp>
#include <miopen/solver/mlir_common.hpp>
Expand All @@ -36,19 +37,6 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_MLIR_IGEMM_BWD)
namespace miopen {
namespace solver {

namespace {
#if MIOPEN_USE_MLIR
std::string GetKernelName()
{
std::string version = "_v4r1";
std::string direction = "_bwd";
return "mlir_gen_igemm_conv2d" + version + direction;
}

std::string GetOperation() { return "conv2d_bwd_data"; }
#endif
} // Anonymous namespace

bool ConvMlirIgemmBwd::IsApplicable(const ConvolutionContext& ctx) const
{
#if MIOPEN_USE_MLIR
Expand All @@ -59,29 +47,56 @@ bool ConvMlirIgemmBwd::IsApplicable(const ConvolutionContext& ctx) const
if(!IsComposableKernelSupportedHardware(ctx))
return false;

return MiirIsConfigApplicable(
mlir::ConstructBuildOptions(ctx, GetOperation(), GetKernelName(), false));
return MiirIsConfigApplicable(mlir::ConstructBuildOptions(ctx, false));
#else
std::ignore = ctx;
return false;
#endif
}

ConvSolution ConvMlirIgemmBwd::GetSolution(const ConvolutionContext& ctx) const
PerformanceConvMlirIgemm ConvMlirIgemmBwd::GetPerformanceConfig(const ConvolutionContext& ctx) const
{
std::ignore = ctx;
return {};
}
Comment on lines +57 to +61
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This produces invalid performance config, which does not match the specification. Why this is done in GetSoluiton but not in this function?

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ping @jerryyin to check. Sorry to have missed it in last review.

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I agree it break the existing specification, but this is intentional. I will explain why I made this decision in #1154 in case this thread becomes too long. Then based on the discussion result, we can decide what's the best behavior it is for mlir solvers.


bool ConvMlirIgemmBwd::IsValidPerformanceConfig(const ConvolutionContext& ctx,
const PerformanceConvMlirIgemm& config) const
{
MIOPEN_LOG_I("");
return config.IsValid(ctx);
}

PerformanceConvMlirIgemm ConvMlirIgemmBwd::Search(const ConvolutionContext& ctx,
const AnyInvokeParams& invoke_ctx) const
{
return GenericSearch(*this, ctx, invoke_ctx);
}

ConvSolution ConvMlirIgemmBwd::GetSolution(const ConvolutionContext& ctx,
const PerformanceConvMlirIgemm& config,
bool) const
{
#if MIOPEN_USE_MLIR
ConvSolution result;
int kernel_count = MiirGetKernelCount(
mlir::ConstructBuildOptions(ctx, GetOperation(), GetKernelName(), false));
int kernel_count = MiirGetKernelCount(mlir::ConstructBuildOptions(ctx, false));

for(int kernel_id = 0; kernel_id < kernel_count; ++kernel_id)
{
KernelInfo construction_parameters;

construction_parameters.kernel_name = GetKernelName() + std::to_string(kernel_id);
construction_parameters.kernel_name = mlir::GetKernelName(ctx, false, kernel_id);
construction_parameters.kernel_file = construction_parameters.kernel_name + ".mlir";
construction_parameters.comp_options = mlir::ConstructBuildOptions(
ctx, GetOperation(), construction_parameters.kernel_name, false, kernel_id);

if(config == PerformanceConvMlirIgemm())
// At this case, do not pass in the invalid perf config and instead make Miir library to
// do heuristic initialization
construction_parameters.comp_options =
mlir::ConstructBuildOptions(ctx, false, kernel_id);
else
// At this case, Make Miir library to use the valid perf config
construction_parameters.comp_options =
mlir::ConstructBuildOptions(ctx, config.ToString(), false, kernel_id);

size_t local_size = 0;
size_t global_size = 0;
Expand All @@ -100,6 +115,7 @@ ConvSolution ConvMlirIgemmBwd::GetSolution(const ConvolutionContext& ctx) const
return result;
#else
std::ignore = ctx;
std::ignore = config;
return {};
#endif
}
Expand Down
60 changes: 39 additions & 21 deletions src/solver/conv_mlir_igemm_bwd_xdlops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,7 @@
#include <miopen/conv/invokers/mlir_impl_gemm.hpp>
#include <miopen/config.h>
#include <miopen/env.hpp>
#include <miopen/generic_search.hpp>
#include <miopen/mlir_build.hpp>
#include <miopen/solver.hpp>
#include <miopen/solver/implicitgemm_util.hpp>
Expand All @@ -37,19 +38,6 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_MLIR_IGEMM_BWD_XDLOPS)
namespace miopen {
namespace solver {

namespace {
#if MIOPEN_USE_MLIR
std::string GetKernelName()
{
std::string version = "_v4r1";
std::string direction = "_bwd";
return "mlir_gen_igemm_conv2d" + version + direction + "_xdlops";
}

std::string GetOperation() { return "conv2d_bwd_data"; }
#endif
} // Anonymous namespace

bool ConvMlirIgemmBwdXdlops::IsApplicable(const ConvolutionContext& ctx) const
{
#if MIOPEN_USE_MLIR
Expand All @@ -62,29 +50,58 @@ bool ConvMlirIgemmBwdXdlops::IsApplicable(const ConvolutionContext& ctx) const
if(!IsComposableKernelSupportedHardware(ctx))
return false;

return MiirIsConfigApplicable(
mlir::ConstructBuildOptions(ctx, GetOperation(), GetKernelName(), true));
return MiirIsConfigApplicable(mlir::ConstructBuildOptions(ctx, true));
#else
std::ignore = ctx;
return false;
#endif
}

ConvSolution ConvMlirIgemmBwdXdlops::GetSolution(const ConvolutionContext& ctx) const
PerformanceConvMlirIgemmXdlops
ConvMlirIgemmBwdXdlops::GetPerformanceConfig(const ConvolutionContext& ctx) const
{
std::ignore = ctx;
return {};
}

bool ConvMlirIgemmBwdXdlops::IsValidPerformanceConfig(
const ConvolutionContext& ctx, const PerformanceConvMlirIgemmXdlops& config) const
{
MIOPEN_LOG_I("");
return config.IsValid(ctx);
}

PerformanceConvMlirIgemmXdlops
ConvMlirIgemmBwdXdlops::Search(const ConvolutionContext& ctx,
const AnyInvokeParams& invoke_ctx) const
{
return GenericSearch(*this, ctx, invoke_ctx);
}

ConvSolution ConvMlirIgemmBwdXdlops::GetSolution(const ConvolutionContext& ctx,
const PerformanceConvMlirIgemmXdlops& config,
bool) const
{
#if MIOPEN_USE_MLIR
ConvSolution result;
int kernel_count =
MiirGetKernelCount(mlir::ConstructBuildOptions(ctx, GetOperation(), GetKernelName(), true));
int kernel_count = MiirGetKernelCount(mlir::ConstructBuildOptions(ctx, true));

for(int kernel_id = 0; kernel_id < kernel_count; ++kernel_id)
{
KernelInfo construction_parameters;

construction_parameters.kernel_name = GetKernelName() + std::to_string(kernel_id);
construction_parameters.kernel_name = mlir::GetKernelName(ctx, true, kernel_id);
construction_parameters.kernel_file = construction_parameters.kernel_name + ".mlir";
construction_parameters.comp_options = mlir::ConstructBuildOptions(
ctx, GetOperation(), construction_parameters.kernel_name, true, kernel_id);

if(config == PerformanceConvMlirIgemmXdlops())
// At this case, do not pass in the invalid perf config and instead make Miir library to
// do heuristic initialization
construction_parameters.comp_options =
mlir::ConstructBuildOptions(ctx, true, kernel_id);
else
// At this case, Make Miir library to use the valid perf config
construction_parameters.comp_options =
mlir::ConstructBuildOptions(ctx, config.ToString(), true, kernel_id);

size_t local_size = 0;
size_t global_size = 0;
Expand All @@ -103,6 +120,7 @@ ConvSolution ConvMlirIgemmBwdXdlops::GetSolution(const ConvolutionContext& ctx)
return result;
#else
std::ignore = ctx;
std::ignore = config;
return {};
#endif
}
Expand Down
28 changes: 6 additions & 22 deletions src/solver/conv_mlir_igemm_fwd.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,19 +37,6 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_MLIR_IGEMM_FWD)
namespace miopen {
namespace solver {

namespace {
#if MIOPEN_USE_MLIR
std::string GetKernelName()
{
std::string version = "_v4r4";
std::string direction = "_fwd";
return "mlir_gen_igemm_conv2d" + version + direction;
}

std::string GetOperation() { return "conv2d"; }
#endif
} // Anonymous namespace

PerformanceConvMlirIgemm::PerformanceConvMlirIgemm(int BlockSize_,
int GemmMPerBlock_,
int GemmNPerBlock_,
Expand Down Expand Up @@ -98,8 +85,7 @@ bool PerformanceConvMlirIgemm::operator==(const PerformanceConvMlirIgemm& other)
bool PerformanceConvMlirIgemm::IsValid(const ConvolutionContext& ctx) const
{
#if MIOPEN_USE_MLIR
return MiirIsConfigApplicable(
mlir::ConstructBuildOptions(ctx, GetOperation(), GetKernelName(), ToString(), false));
return MiirIsConfigApplicable(mlir::ConstructBuildOptions(ctx, ToString(), false));
#else
std::ignore = ctx;
return false;
Expand Down Expand Up @@ -167,8 +153,7 @@ bool ConvMlirIgemmFwd::IsApplicable(const ConvolutionContext& ctx) const
if(!IsComposableKernelSupportedHardware(ctx))
return false;

return MiirIsConfigApplicable(
mlir::ConstructBuildOptions(ctx, GetOperation(), GetKernelName(), false));
return MiirIsConfigApplicable(mlir::ConstructBuildOptions(ctx, false));
#else
std::ignore = ctx;
return false;
Expand All @@ -183,18 +168,17 @@ ConvSolution ConvMlirIgemmFwd::GetSolution(const ConvolutionContext& ctx,
ConvSolution result;
KernelInfo construction_parameters;

construction_parameters.kernel_name = GetKernelName();
construction_parameters.kernel_name = mlir::GetKernelName(ctx, false);
construction_parameters.kernel_file = construction_parameters.kernel_name + ".mlir";

if(config == PerformanceConvMlirIgemm())
// At this case, do not pass in the invalid perf config and instead make Miir library to do
// heuristic initialization
construction_parameters.comp_options =
mlir::ConstructBuildOptions(ctx, GetOperation(), GetKernelName(), false);
construction_parameters.comp_options = mlir::ConstructBuildOptions(ctx, false);
else
// At this case, Make Miir library to use the valid perf config
construction_parameters.comp_options = mlir::ConstructBuildOptions(
ctx, GetOperation(), GetKernelName(), config.ToString(), false);
construction_parameters.comp_options =
mlir::ConstructBuildOptions(ctx, config.ToString(), false);

size_t local_size = 0;
size_t global_size = 0;
Expand Down
28 changes: 6 additions & 22 deletions src/solver/conv_mlir_igemm_fwd_xdlops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,19 +38,6 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_CONV_MLIR_IGEMM_FWD_XDLOPS)
namespace miopen {
namespace solver {

namespace {
#if MIOPEN_USE_MLIR
std::string GetKernelName()
{
std::string version = "_v4r4";
std::string direction = "_fwd";
return "mlir_gen_igemm_conv2d" + version + direction + "_xdlops";
}

std::string GetOperation() { return "conv2d"; }
#endif
} // Anonymous namespace

bool ConvMlirIgemmFwdXdlops::IsApplicable(const ConvolutionContext& ctx) const
{
#if MIOPEN_USE_MLIR
Expand All @@ -62,8 +49,7 @@ bool ConvMlirIgemmFwdXdlops::IsApplicable(const ConvolutionContext& ctx) const
return false;
if(!IsComposableKernelSupportedHardware(ctx))
return false;
return MiirIsConfigApplicable(
mlir::ConstructBuildOptions(ctx, GetOperation(), GetKernelName(), true));
return MiirIsConfigApplicable(mlir::ConstructBuildOptions(ctx, true));
#else
std::ignore = ctx;
return false;
Expand Down Expand Up @@ -121,8 +107,7 @@ bool PerformanceConvMlirIgemmXdlops::operator==(const PerformanceConvMlirIgemmXd
bool PerformanceConvMlirIgemmXdlops::IsValid(const ConvolutionContext& ctx) const
{
#if MIOPEN_USE_MLIR
bool isValid = MiirIsConfigApplicable(
mlir::ConstructBuildOptions(ctx, GetOperation(), GetKernelName(), ToString(), true));
bool isValid = MiirIsConfigApplicable(mlir::ConstructBuildOptions(ctx, ToString(), true));
return isValid;
#else
std::ignore = ctx;
Expand Down Expand Up @@ -191,18 +176,17 @@ ConvSolution ConvMlirIgemmFwdXdlops::GetSolution(const ConvolutionContext& ctx,
ConvSolution result;
KernelInfo construction_parameters;

construction_parameters.kernel_name = GetKernelName();
construction_parameters.kernel_name = mlir::GetKernelName(ctx, true);
construction_parameters.kernel_file = construction_parameters.kernel_name + ".mlir";

if(config == PerformanceConvMlirIgemmXdlops())
// At this case, do not pass in the invalid perf config and instead make Miir library to do
// heuristic initialization
construction_parameters.comp_options =
mlir::ConstructBuildOptions(ctx, GetOperation(), GetKernelName(), true);
construction_parameters.comp_options = mlir::ConstructBuildOptions(ctx, true);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

⚓

else
// At this case, Make Miir library to use the valid perf config
construction_parameters.comp_options = mlir::ConstructBuildOptions(
ctx, GetOperation(), GetKernelName(), config.ToString(), true);
construction_parameters.comp_options =
mlir::ConstructBuildOptions(ctx, config.ToString(), true);

size_t local_size = 0;
size_t global_size = 0;
Expand Down
Loading