Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
6 changes: 2 additions & 4 deletions projects/miopen/docs/how-to/debug-log.rst
Original file line number Diff line number Diff line change
Expand Up @@ -151,11 +151,9 @@ Some of the solutions have individual controls, which affect both find and immed
* ``MIOPEN_DEBUG_CONV_DIRECT_HIP_FWD11X11`` -- ``ConvHipDirectFwd11x11``
* ``MIOPEN_DEBUG_CONV_DIRECT_OCL_FWDGEN`` -- ``ConvOclDirectFwdGen``
* ``MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD`` -- ``ConvOclDirectFwd``

* ``MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW2`` -- ``ConvOclBwdWrW2<n>`` (where n =
``{1,2,4,8,16}``) and ``ConvOclBwdWrW2NonTunable``
* ``MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1`` -- ``ConvOclDirectFwd1x1``
* ``MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW53`` -- ``ConvOclBwdWrW53``
* ``MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW1X1`` -- ``ConvOclBwdWrW1x1``


* Winograd solutions:

Expand Down
9 changes: 2 additions & 7 deletions projects/miopen/docs/reference/env_variables.rst
Original file line number Diff line number Diff line change
Expand Up @@ -278,8 +278,8 @@ For more information, see :doc:`Logging and debugging <../how-to/debug-log>`.
- | 0: Disable
| 1: Enable

* - | ``MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW2``
| Controls ConvOclBwdWrW2<n> (n={1,2,4,8,16}) and ConvOclBwdWrW2NonTunable solutions.
* - | ``MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD1X1``
| Controls ConvOclDirectFwd1x1 direct solution.
- | 0: Disable
| 1: Enable

Expand All @@ -288,11 +288,6 @@ For more information, see :doc:`Logging and debugging <../how-to/debug-log>`.
- | 0: Disable
| 1: Enable

* - | ``MIOPEN_DEBUG_CONV_DIRECT_OCL_WRW1X1``
| Controls ConvOclBwdWrW1x1 direct solution.
- | 0: Disable
| 1: Enable

* - | ``MIOPEN_DEBUG_CONV_DIRECT_ASM_1X1U_PERF_VALS``
| Controls performance values for ConvAsm1x1U direct solution.
- | 0: Disable
Expand Down
6 changes: 1 addition & 5 deletions projects/miopen/src/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -233,8 +233,7 @@ set( MIOpen_Source
solver/conv/conv_mlir_igemm_wrw.cpp
solver/conv/conv_mlir_igemm_wrw_xdlops.cpp
solver/conv/conv_multipass_wino3x3WrW.cpp
solver/conv/conv_ocl_dir2D_bwdWrW_1x1.cpp
solver/conv/conv_ocl_dir2D_bwdWrW_2.cpp

solver/conv/conv_ocl_dir2D_bwdWrW_53.cpp
solver/conv/conv_hip_dir2D11x11.cpp
solver/conv/conv_ocl_dir2Dfwd.cpp
Expand Down Expand Up @@ -583,11 +582,8 @@ if( MIOPEN_BACKEND MATCHES "OpenCL" OR MIOPEN_BACKEND STREQUAL "HIPOC" OR MIOPEN
kernels/MIOpenIm3d2Col.cpp
kernels/MIOpenCol2Im2d.cpp
kernels/MIOpenCol2Im3d.cpp
kernels/MIOpenConvBwdWrWS2.cl
kernels/MIOpenGroupConvBwdWrWS2.cl
kernels/MIOpenConvBwdWrW_LxG_P53.cl
kernels/MIOpenGroupConvBwdWrW_LxG_P53.cl
kernels/MIOpenConvBwdWrW1x1_PAD_read4.cl
kernels/MIOpenConvFwd_LxL_11.cpp

kernels/MIOpenConvFFT.cpp
Expand Down
5 changes: 0 additions & 5 deletions projects/miopen/src/fin/fin_interface.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -265,11 +265,6 @@ AnySolver<miopen::ExecutionContext, miopen::conv::ProblemDescription>::AnySolver
case 11: SetObject<miopen::solver::conv::ConvOclDirectFwd>(); break;
case 16: SetObject<miopen::solver::conv::ConvAsmBwdWrW3x3>(); break;
case 17: SetObject<miopen::solver::conv::ConvAsmBwdWrW1x1>(); break;
case 18: SetObject<miopen::solver::conv::ConvOclBwdWrW2<1>>(); break;
case 19: SetObject<miopen::solver::conv::ConvOclBwdWrW2<2>>(); break;
case 20: SetObject<miopen::solver::conv::ConvOclBwdWrW2<4>>(); break;
case 21: SetObject<miopen::solver::conv::ConvOclBwdWrW2<8>>(); break;
case 22: SetObject<miopen::solver::conv::ConvOclBwdWrW2<16>>(); break;
case 26: SetObject<miopen::solver::conv::ConvHipImplicitGemmV4R1Fwd>(); break;
case 31: SetObject<miopen::solver::conv::ConvHipImplicitGemmV4R1WrW>(); break;
case 37: SetObject<miopen::solver::conv::ConvBinWinoRxS<3, 2>>(); break;
Expand Down
143 changes: 0 additions & 143 deletions projects/miopen/src/include/miopen/conv/solvers.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -2274,136 +2274,6 @@ struct ConvAsmBwdWrW1x1 final : ConvTunableSolver<PerformanceConfigConvAsmBwdWrW
const PerformanceConfigConvAsmBwdWrW1x1&) const override;
};

/// N_BATCH_LOOPS - {1,2,4,8,16} Num batches processed in single workitem.
/// Required workspace size depends on it. However there is a restriction in the internal
/// Solver API that this shouldn't be so. Therefore the family of Solvers created.
/// Each Solver in the family has constant value of this parameter.
template <int N_BATCH_LOOPS>
struct PerformanceConfigConvOclBwdWrw2
: PerfConfigBase<PerformanceConfigConvOclBwdWrw2<N_BATCH_LOOPS>>
{
// Num waves involved a workgroup.
int n_waves = -1; // {1,2,4,8}
// Num values to read in a workitem (read_unit).
int read_size = -1; // [6..12]
// Num of output channels (top/bottom layer in forward/backward direction)
// that share the same input channel in single workgroup.
// Also represents number of output channels in single tile.
int n_out_channels_per_tile = -1; // {1,2,4,8}
// How many tiles of output channels are processed in a single workgroup?
// n_out_channels_in_lcl * n_out_channels_tiles = total number of
// output channels processed in single workgroup.
int n_out_channels_tiles = -1; // {1,2,4,8}
// Num of output rows processed in a single iteration of loop in a workitem
// (N_ALIGNED_OUT_SCAN_BLK).
int n_out_rows_in_lcl = -1; // [2..11]

PerformanceConfigConvOclBwdWrw2(int nw, int rs, int nocpt, int noct, int noril)
: n_waves(nw),
read_size(rs),
n_out_channels_per_tile(nocpt),
n_out_channels_tiles(noct),
n_out_rows_in_lcl(noril)
{
}
PerformanceConfigConvOclBwdWrw2() {}
PerformanceConfigConvOclBwdWrw2(bool) : PerformanceConfigConvOclBwdWrw2(1, 6, 1, 1, 2) {}
// spare_set is not used in this solver.

template <class Self, class F>
static void Visit(Self&& self, F f)
{
f(self.n_waves, "n_waves");
f(self.read_size, "read_size");
f(self.n_out_channels_per_tile, "n_out_channels_per_tile");
f(self.n_out_channels_tiles, "n_out_channels_tiles");
f(self.n_out_rows_in_lcl, "n_out_rows_in_lcl");
}

int GetNumWaves() const { return n_waves; }
int GetReadSize() const { return read_size; }
int GetNumOutChannelsPerTile() const { return n_out_channels_per_tile; }
int GetNumOutChannelTiles() const { return n_out_channels_tiles; }
int GetNumOutRowsPerIterPerWork() const { return n_out_rows_in_lcl; }

void HeuristicInit(const miopen::conv::ProblemDescription&);
bool IsValidValue() const;
bool SetNextValue(const miopen::conv::ProblemDescription&);
bool IsValid(const ExecutionContext&, const miopen::conv::ProblemDescription&) const;
bool operator==(const PerformanceConfigConvOclBwdWrw2<N_BATCH_LOOPS>& other) const;
};

template <int N_BATCH_LOOPS>
struct ConvOclBwdWrW2 final : ConvTunableSolver<PerformanceConfigConvOclBwdWrw2<N_BATCH_LOOPS>>
{
const std::string& SolverDbId() const override
{
return this->template GetSolverDbId<ConvOclBwdWrW2<N_BATCH_LOOPS>>();
}

PerformanceConfigConvOclBwdWrw2<N_BATCH_LOOPS>
GetDefaultPerformanceConfig(const ExecutionContext&,
const miopen::conv::ProblemDescription&) const override;
bool
IsValidPerformanceConfig(const ExecutionContext&,
const miopen::conv::ProblemDescription&,
const PerformanceConfigConvOclBwdWrw2<N_BATCH_LOOPS>&) const override;
PerformanceConfigConvOclBwdWrw2<N_BATCH_LOOPS>
Search(const ExecutionContext&,
const miopen::conv::ProblemDescription&,
const AnyInvokeParams& invoke_ctx) const override;
bool IsApplicable(const ExecutionContext&,
const miopen::conv::ProblemDescription&) const override;
size_t GetWorkspaceSize(const ExecutionContext&,
const miopen::conv::ProblemDescription&) const override;
bool MayNeedWorkspace() const override { return true; }
ConvSolution GetSolution(const ExecutionContext&,
const miopen::conv::ProblemDescription&,
const PerformanceConfigConvOclBwdWrw2<N_BATCH_LOOPS>&) const override;

protected:
bool IsApplicableBase(const ExecutionContext&, const miopen::conv::ProblemDescription&) const;

friend struct ConvOclBwdWrW2NonTunable;
};

// To suppress misleading clang warnings
#ifndef CONV_OCL_DIR2D_BWDWRW_2_CPP

extern template struct PerformanceConfigConvOclBwdWrw2<1>;
extern template struct PerformanceConfigConvOclBwdWrw2<2>;
extern template struct PerformanceConfigConvOclBwdWrw2<4>;
extern template struct PerformanceConfigConvOclBwdWrw2<8>;
extern template struct PerformanceConfigConvOclBwdWrw2<16>;

extern template struct ConvOclBwdWrW2<1>;
extern template struct ConvOclBwdWrW2<2>;
extern template struct ConvOclBwdWrW2<4>;
extern template struct ConvOclBwdWrW2<8>;
extern template struct ConvOclBwdWrW2<16>;

#endif

/// A separate solver from ConvOclBwdWrW2 to disable auto-tuning for certain configs.
/// Basically, this is *hack* for non-group 3x3 and 1x1 cases.
/// It is assumed that Solutions provided by the ConvOclBwdWrW2 solver
/// would never beat 3x3 and 1x1 assembly WrW kernels, even after tuning.
struct MIOPEN_INTERNALS_EXPORT ConvOclBwdWrW2NonTunable final : ConvSolver
{
const std::string& SolverDbId() const override
{
return GetSolverDbId<ConvOclBwdWrW2NonTunable>();
}

bool IsApplicable(const ExecutionContext&,
const miopen::conv::ProblemDescription&) const override;
size_t GetWorkspaceSize(const ExecutionContext&,
const miopen::conv::ProblemDescription&) const override;
bool MayNeedWorkspace() const override { return true; }
ConvSolution GetSolution(const ExecutionContext&,
const miopen::conv::ProblemDescription&) const override;
};

struct MIOPEN_INTERNALS_EXPORT ConvOclBwdWrW53 final : ConvSolver
{
const std::string& SolverDbId() const override { return GetSolverDbId<ConvOclBwdWrW53>(); }
Expand All @@ -2417,19 +2287,6 @@ struct MIOPEN_INTERNALS_EXPORT ConvOclBwdWrW53 final : ConvSolver
const miopen::conv::ProblemDescription&) const override;
};

struct MIOPEN_INTERNALS_EXPORT ConvOclBwdWrW1x1 final : ConvSolver
{
const std::string& SolverDbId() const override { return GetSolverDbId<ConvOclBwdWrW1x1>(); }

bool IsApplicable(const ExecutionContext&,
const miopen::conv::ProblemDescription&) const override;
size_t GetWorkspaceSize(const ExecutionContext&,
const miopen::conv::ProblemDescription&) const override;
bool MayNeedWorkspace() const override { return true; }
ConvSolution GetSolution(const ExecutionContext&,
const miopen::conv::ProblemDescription&) const override;
};

struct MIOPEN_INTERNALS_EXPORT fft final : ConvSolver
{
const std::string& SolverDbId() const override { return GetSolverDbId<fft>(); }
Expand Down
Loading
Loading