diff --git a/dnn/src/rocm/miopen_wrapper.cpp b/dnn/src/rocm/miopen_wrapper.cpp index 3a48a466..1b82405f 100644 --- a/dnn/src/rocm/miopen_wrapper.cpp +++ b/dnn/src/rocm/miopen_wrapper.cpp @@ -9,6 +9,7 @@ * "AS IS" BASIS, WITHOUT ARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. */ #include "hcc_detail/hcc_defs_prologue.h" +#include "megdnn/opr_param_defs.h" #include "src/rocm/miopen_wrapper.h" #include "src/common/utils.h" @@ -114,6 +115,9 @@ void PoolingDesc::set(const param::Pooling& param) { case param::Pooling::Mode::AVERAGE_COUNT_EXCLUDE_PADDING: mode = miopenPoolingAverage; break; + case param::Pooling::Mode::AVERAGE: + mode = miopenPoolingAverageInclusive; + break; default: megdnn_throw(megdnn_mangle("Unsupported pooling mode for miopen")); } diff --git a/dnn/test/rocm/pooling.cpp b/dnn/test/rocm/pooling.cpp index d60d35a5..08cee8ce 100644 --- a/dnn/test/rocm/pooling.cpp +++ b/dnn/test/rocm/pooling.cpp @@ -31,10 +31,6 @@ TEST_F(ROCM, POOLING_FORWARD) { for (auto format : {Format::NCHW}) for (auto&& arg : args) { auto param = arg.param; - if (param.mode == param::Pooling::Mode::AVERAGE) { - param.mode = - param::Pooling::Mode::AVERAGE_COUNT_EXCLUDE_PADDING; - } auto src = arg.ishape; param.format = format; Checker checker(handle_rocm()); @@ -53,11 +49,6 @@ TEST_F(ROCM, POOLING_BACKWARD) { TensorLayout ilayout = TensorLayout(arg.ishape, dtype::Float32()); TensorLayout olayout; - auto& param = arg.param; - if (param.mode == param::Pooling::Mode::AVERAGE) { - param.mode = param::Pooling::Mode::AVERAGE_COUNT_EXCLUDE_PADDING; - } - auto constraint = [this, arg](CheckerHelper::TensorValueArray& tensors_orig) { megdnn_assert(tensors_orig.size() == 4); @@ -141,7 +132,7 @@ TEST_F(ROCM, POOLING_FWD_BENCHMARK) { benchmarker.set_param(param); size_t OH = infer_conv_shape(IH, FH, SH, PH); size_t OW = infer_conv_shape(IW, FW, SW, PW); - // warm up + // warm up benchmarker.execs({{N, IC, IH, IW}, {N, IC, OH, OW}}); // do actual benchmark auto time_ms = benchmarker.execs({{N, IC, IH, IW}, {N, IC, OH, OW}});