Skip to content

Commit

Permalink
[MLIR] Implement tuning support for xdlops and nonxdlops
Browse files Browse the repository at this point in the history
  • Loading branch information
@jerryyin
jerryyin committed Sep 10, 2021
1 parent 8b2f260 commit a000cfc
Show file tree
Hide file tree
Showing 9 changed files with 162 additions and 140 deletions.
18 changes: 16 additions & 2 deletions src/include/miopen/solver.hpp
View file
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
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
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 {};
}

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

0 comments on commit a000cfc

Please sign in to comment.