From 19bd9c495d1f0c5010f0694a062e0558cca81777 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mirza=20Halil=C4=8Devi=C4=87?= <109971222+mirza-halilcevic@users.noreply.github.com> Date: Wed, 22 Nov 2023 19:59:29 +0100 Subject: [PATCH] Add support for the dilations attribute to Pooling ops (#2105) Introduce dilations attribute to pooling operators reference implementation. --- src/include/migraphx/op/pooling.hpp | 51 +- src/include/migraphx/rewrite_pooling.hpp | 1 + src/onnx/parse_pooling.cpp | 18 +- src/rewrite_pooling.cpp | 148 +++++- src/targets/cpu/pooling.cpp | 17 +- .../gpu/include/migraphx/gpu/miopen.hpp | 6 + test/onnx/averagepool_dilate_test.onnx | 17 + test/onnx/gen_onnx.py | 32 ++ test/onnx/maxpool_dilate_test.onnx | 17 + test/onnx/onnx_test.cpp | 77 ++- test/op_shape_test.cpp | 76 ++- test/py/onnx_backend_test.py | 2 + test/ref/pooling.cpp | 326 +++++++++--- test/rewrite_pooling_test.cpp | 492 +++++++++++++++++- test/simplify_qdq_test.cpp | 3 + test/verify/test_avg_pooling_1d.cpp | 2 +- test/verify/test_avg_pooling_3d.cpp | 2 +- test/verify/test_avg_pooling_3d_opt.cpp | 2 +- test/verify/test_avg_pooling_ceil_3d.cpp | 2 +- test/verify/test_avg_pooling_pad.cpp | 2 +- test/verify/test_concat_pooling.cpp | 3 +- test/verify/test_conv_bn_relu_pooling.cpp | 3 +- test/verify/test_conv_bn_relu_pooling2.cpp | 3 +- test/verify/test_max_pooling_ceil_3d.cpp | 2 +- 24 files changed, 1170 insertions(+), 134 deletions(-) create mode 100644 test/onnx/averagepool_dilate_test.onnx create mode 100644 test/onnx/maxpool_dilate_test.onnx diff --git a/src/include/migraphx/op/pooling.hpp b/src/include/migraphx/op/pooling.hpp index 7bfe456f3a0..3037e75752b 100644 --- a/src/include/migraphx/op/pooling.hpp +++ b/src/include/migraphx/op/pooling.hpp @@ -70,7 +70,8 @@ struct pooling // 2 smaller than the input tensor rank (NCHW layout) std::vector lengths = {1, 1}; - // Dilations are not supported at this time. + // Spacing between the elements of the pooling kernel. Must be the same ndim as lengths. + std::vector dilations = {1, 1}; // ceiling mode is a flag affecting output size // or equivalently, placements of the pooling kernel. @@ -99,6 +100,7 @@ struct pooling f(self.padding_mode, "padding_mode"), f(self.stride, "stride"), f(self.lengths, "lengths"), + f(self.dilations, "dilations"), f(self.ceil_mode, "ceil_mode"), f(self.lp_order, "lp_order"), f(self.dyn_global, "dyn_global")); @@ -112,14 +114,17 @@ struct pooling return; if((padding_mode != default_ and padding.size() != stride.size() and (padding.size()) != stride.size() * 2) or - stride.size() != lengths.size()) + stride.size() != lengths.size() or dilations.size() != lengths.size()) { MIGRAPHX_THROW("POOLING: inconsistent attribute sizes"); } - if(std::any_of(lengths.begin(), lengths.end(), [&](auto i) { return (i == 0); }) or - std::any_of(stride.begin(), stride.end(), [&](auto i) { return (i == 0); })) + + const auto is_zero = [](auto el) { return el == 0; }; + if(std::any_of(lengths.begin(), lengths.end(), is_zero) or + std::any_of(stride.begin(), stride.end(), is_zero) or + std::any_of(dilations.begin(), dilations.end(), is_zero)) { - MIGRAPHX_THROW("POOLING: size 0 pooling kernel or stride"); + MIGRAPHX_THROW("POOLING: size 0 pooling kernel or stride or dilations"); } // TODO: update lowering to run the reference @@ -142,6 +147,11 @@ struct pooling value attributes() const { return {{"normalize_padding", "padding"}}; } + inline std::size_t dilate_dim(std::size_t dim, std::size_t dilation) const + { + return 1 + dilation * (dim - 1); + } + std::vector calc_spatial_dim_out(const std::vector& input_lens, std::size_t kdims) const { @@ -151,8 +161,9 @@ struct pooling std::size_t padding_factor = 2 * padding[i]; if(padding.size() == 2 * kdims) padding_factor = padding[i] + padding[i + kdims]; + std::size_t dilated_length = dilate_dim(lengths[i], dilations[i]); std::size_t dim_size; - if(input_lens[i + 2] + padding_factor < lengths[i]) + if(input_lens[i + 2] + padding_factor < dilated_length) { if(padding_mode == default_) MIGRAPHX_THROW("POOLING: not enough padding for the given kernel size"); @@ -162,7 +173,7 @@ struct pooling } else { - dim_size = input_lens[i + 2] + padding_factor - lengths[i]; + dim_size = input_lens[i + 2] + padding_factor - dilated_length; } std::size_t len = (ceil_mode) @@ -331,6 +342,7 @@ struct pooling int start = static_cast(idx_o[dim] * stride[d_2]) - static_cast(padding_vals[d_2]); int end; + std::size_t dilated_kernel_dim = dilate_dim(kernel_dims[d_2], dilations[d_2]); // NOLINT if(count_include_pad and ceil_mode and (mode != pooling_mode::max)) { @@ -340,15 +352,14 @@ struct pooling // padding. Clip out-of-bounds indexes but not padding. // Check if this kernel extends beyond the padding at end of dimension - end = std::min(start + kernel_dims[d_2], + end = std::min(start + dilated_kernel_dim, in_lens[dim] + static_cast(padding_vals[d_2])); } else { // In non-ceiling mode, when // count_include_pad is false, or for max pooling, clip off padding. - end = std::min(start + kernel_dims[d_2], in_lens[dim]); - start = std::max(start, 0); + end = std::min(start + dilated_kernel_dim, in_lens[dim]); } win_start.push_back(start); if(end < start) @@ -366,6 +377,16 @@ struct pooling // for each element in the window... shape_for_each(win_shape, [&](const auto& idx_w) { + // Skip elements that belong to the dilated area + for(size_t axis = 0; axis < idx_w.size(); ++axis) + { + if(idx_w[axis] % dilations[axis]) + { + pool_size -= 1; + return; + } + } + // the coordinates of this element auto idx = idx_o; @@ -390,7 +411,15 @@ struct pooling // this is a padding element. Padding locations // don't contribute to average or max pooling total but can play in // lpnorm pooling. - output_val = op(output_val, 0); + if(mode == pooling_mode::lpnorm) + { + output_val = op(output_val, op.template init()); + } + if(mode == pooling_mode::average) + { + // Ignore padding + pool_size -= 1; + } } }); output[i] = Type(op.final(output_val, pool_size)); diff --git a/src/include/migraphx/rewrite_pooling.hpp b/src/include/migraphx/rewrite_pooling.hpp index a250e255f78..e8102481728 100644 --- a/src/include/migraphx/rewrite_pooling.hpp +++ b/src/include/migraphx/rewrite_pooling.hpp @@ -26,6 +26,7 @@ #include #include +#include namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { diff --git a/src/onnx/parse_pooling.cpp b/src/onnx/parse_pooling.cpp index ade3ebbac99..c61a8844376 100644 --- a/src/onnx/parse_pooling.cpp +++ b/src/onnx/parse_pooling.cpp @@ -91,6 +91,14 @@ struct parse_pooling : op_parser kdims, values["lengths"].size(), "PARSE_POOLING: inconsistent lengths"); } + if(contains(info.attributes, "dilations")) + { + values["dilations"].clear(); + copy(info.attributes["dilations"].ints(), std::back_inserter(values["dilations"])); + check_attr_sizes( + kdims, values["dilations"].size(), "PARSE_POOLING: inconsistent dilations"); + } + // lp_order attribute if(contains(info.attributes, "p")) { @@ -169,10 +177,15 @@ struct parse_pooling : op_parser std::fill_n(values["stride"].begin(), kdims, 1); } + if(values["dilations"].size() != kdims) + { + values["dilations"].resize(kdims); + std::fill_n(values["dilations"].begin(), kdims, 1); + } + // used to calculate the supposed output shape std::vector orig_padding = paddings; - // TODO: add parsing for dilations if(contains(info.attributes, "auto_pad") and to_upper(info.attributes["auto_pad"].s()) != "NOTSET") { @@ -189,11 +202,10 @@ struct parse_pooling : op_parser else { // Calculate auto padding - // dilations (argument 4) not supported; default to all 1's cal_auto_padding_size(info, values, values["lengths"].to_vector(), - std::vector(in_shape.ndim() - 2, 1), + values["dilations"].to_vector(), in_shape.lens(), paddings); values["padding"] = paddings; diff --git a/src/rewrite_pooling.cpp b/src/rewrite_pooling.cpp index b381a8a73f9..18f636e3c67 100644 --- a/src/rewrite_pooling.cpp +++ b/src/rewrite_pooling.cpp @@ -35,6 +35,110 @@ namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { +static void replace_with_reduce(module& m, instruction_ref ins) +{ + auto&& s = ins->inputs().front()->get_shape(); + auto&& op = any_cast(ins->get_operator()); + auto lens = s.lens(); + std::vector axes(lens.size() - 2); + std::iota(axes.begin(), axes.end(), 2); + + // average pooling + if(op.mode == op::pooling_mode::average) + { + m.replace_instruction(ins, make_op("reduce_mean", {{"axes", axes}}), ins->inputs()); + } + // max pooling + else + { + m.replace_instruction(ins, make_op("reduce_max", {{"axes", axes}}), ins->inputs()); + } +} + +static void replace_dilations_with_gather_pooling(module& m, instruction_ref ins) +{ + // TODO remove this when MIOpen supports dilated pooling + auto&& s = ins->inputs().front()->get_shape(); + auto&& op = any_cast(ins->get_operator()); + // Ignore N, C axes + std::vector dims = {s.lens().cbegin() + 2, s.lens().cend()}; + + bool default_padding = + std::all_of(op.padding.cbegin(), op.padding.cend(), [](auto i) { return i == 0; }); + + if(not default_padding) + { + for(size_t idx{0}; idx < op.padding.size(); ++idx) + { + // We need to pad both ends + dims[idx] += op.padding.at(idx) * 2; + } + } + std::vector kernels = op.lengths; + std::vector strides = op.stride; + std::vector dilations = op.dilations; + + std::vector> axis_indices; + axis_indices.resize(dims.size()); + + for(auto idx{0}; idx < dims.size(); ++idx) + { + // Only consider if iw fits into the window + for(size_t stride{0}; stride < dims.at(idx) - dilations.at(idx) * (kernels.at(idx) - 1); + stride += strides.at(idx)) + { + for(size_t step{0}; step < kernels.at(idx); ++step) + { + axis_indices.at(idx).push_back(stride + dilations.at(idx) * step); + } + } + } + + auto elements = ins->inputs().front(); + if(not default_padding) + { + // Pad supports asym, we need to provide both ends + std::vector padding(2 * s.lens().size(), 0); + // Format will be e.g {N, C, P1, P2, N, C, P1, P2} + for(size_t idx{0}; idx < op.padding.size(); ++idx) + { + // Ignore N, C axes + padding.at(2 + idx) = op.padding.at(idx); + padding.at(2 + idx + s.lens().size()) = op.padding.at(idx); + } + + // Default value needed for Max pooling + elements = m.insert_instruction( + ins, + make_op("pad", {{"pads", padding}, {"value", std::numeric_limits::lowest()}}), + elements); + } + + for(auto idx{0}; idx < axis_indices.size(); ++idx) + { + migraphx::shape s_indices{migraphx::shape::int32_type, {axis_indices.at(idx).size()}}; + auto indices = m.add_literal(migraphx::literal{s_indices, axis_indices.at(idx)}); + elements = m.insert_instruction( + ins, make_op("gather", {{"axis", idx + 2 /*ignore N,C*/}}), elements, indices); + } + + // Ignore padding + std::vector new_padding(kernels.size(), 0); + // The kernel window elements are places next to each other. E.g. {x1, y1, x2, y2, ...} + // We need to skip them to not overlap + std::vector new_strides(kernels); + // Ignore dilations + std::vector new_dilations(kernels.size(), 1); + m.replace_instruction(ins, + make_op("pooling", + {{"mode", op.mode}, + {"padding", new_padding}, + {"stride", new_strides}, + {"lengths", kernels}, + {"dilations", new_dilations}}), + elements); +} + void rewrite_pooling::apply(module& m) const { for(auto ins : iterator_for(m)) @@ -43,26 +147,36 @@ void rewrite_pooling::apply(module& m) const continue; if(ins->inputs().empty()) continue; - auto&& s = ins->inputs().front()->get_shape(); - auto&& op = any_cast(ins->get_operator()); - if(not std::all_of(op.padding.begin(), op.padding.end(), [](auto i) { return i == 0; })) - continue; - if(not std::all_of(op.stride.begin(), op.stride.end(), [](auto i) { return i == 1; })) - continue; - auto lens = s.lens(); - if(not std::equal(lens.begin() + 2, lens.end(), op.lengths.begin(), op.lengths.end())) - continue; - std::vector axes(lens.size() - 2); - std::iota(axes.begin(), axes.end(), 2); - // average pooling - if(op.mode == op::pooling_mode::average) + auto&& s = ins->inputs().front()->get_shape(); + auto&& op = any_cast(ins->get_operator()); + bool same_kernel_as_shape = std::equal( + s.lens().cbegin() + 2, s.lens().cend(), op.lengths.cbegin(), op.lengths.cend()); + bool default_strides = + std::all_of(op.stride.cbegin(), op.stride.cend(), [](auto i) { return i == 1; }); + bool default_padding = + std::all_of(op.padding.cbegin(), op.padding.cend(), [](auto i) { return i == 0; }); + bool default_dilations = + std::all_of(op.dilations.cbegin(), op.dilations.cend(), [](auto i) { return i == 1; }); + if(same_kernel_as_shape and default_strides and default_padding and default_dilations) { - m.replace_instruction(ins, make_op("reduce_mean", {{"axes", axes}}), ins->inputs()); + replace_with_reduce(m, ins); } - // max pooling - else + else if(not default_dilations) { - m.replace_instruction(ins, make_op("reduce_max", {{"axes", axes}}), ins->inputs()); + // Dilated AvgPool with padding is not supported + if(not default_padding and op.mode == op::pooling_mode::average) + { + continue; + } + auto size = + std::accumulate(s.lens().cbegin(), s.lens().cend(), 1, std::multiplies()); + // Can't handle too much size because of literal size + if(size > 100000) + { + continue; + } + + replace_dilations_with_gather_pooling(m, ins); } } } diff --git a/src/targets/cpu/pooling.cpp b/src/targets/cpu/pooling.cpp index 28ab11f0ea6..560f2156504 100644 --- a/src/targets/cpu/pooling.cpp +++ b/src/targets/cpu/pooling.cpp @@ -34,23 +34,32 @@ namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { namespace cpu { -struct dnnl_pooling : dnnl_extend_op +struct dnnl_pooling : dnnl_extend_op { std::vector arg_map(int) const { return {MIGRAPHX_DNNL_PREFIX(ARG_SRC)}; } - dnnl::pooling_forward::desc get_desc(const std::unordered_map& m) const + dnnl::pooling_v2_forward::desc + get_desc(const std::unordered_map& m) const { - auto algo = op.mode == op::pooling_mode::max ? dnnl::algorithm::pooling_max - : dnnl::algorithm::pooling_avg; + auto algo = op.mode == op::pooling_mode::max ? dnnl::algorithm::pooling_max + : dnnl::algorithm::pooling_avg; auto kdims = op.kdims(); std::vector padding_l(op.padding.begin(), op.padding.begin() + kdims); std::vector padding_r(op.padding.begin() + kdims, op.padding.end()); + // Note: It is not documented, but the default dilation seems to be 0 instead of 1. + // We need to offset dilations with -1. + std::vector dilations; + std::transform(op.dilations.cbegin(), + op.dilations.cend(), + std::back_inserter(dilations), + [](size_t d) { return d - 1; }); return {dnnl::prop_kind::forward_inference, algo, m.at(MIGRAPHX_DNNL_PREFIX(ARG_SRC)), m.at(MIGRAPHX_DNNL_PREFIX(ARG_DST)), to_dnnl_dims(op.stride), to_dnnl_dims(op.lengths), + to_dnnl_dims(dilations), to_dnnl_dims(padding_l), to_dnnl_dims(padding_r)}; } diff --git a/src/targets/gpu/include/migraphx/gpu/miopen.hpp b/src/targets/gpu/include/migraphx/gpu/miopen.hpp index 8163579eb48..a2a4991e98a 100644 --- a/src/targets/gpu/include/migraphx/gpu/miopen.hpp +++ b/src/targets/gpu/include/migraphx/gpu/miopen.hpp @@ -211,6 +211,12 @@ inline pooling_descriptor make_pooling(const migraphx::op::pooling& op) ss << op.mode; MIGRAPHX_THROW(ss.str()); } + if(not std::all_of( + op.dilations.cbegin(), op.dilations.cend(), [](std::size_t d) { return d == 1; })) + { + MIGRAPHX_THROW("Unsupported dilations for pooling: [" + to_string_range(op.dilations) + + "]"); + } auto p = make_obj(&miopenCreatePoolingDescriptor); int kdims = op.kdims(); diff --git a/test/onnx/averagepool_dilate_test.onnx b/test/onnx/averagepool_dilate_test.onnx new file mode 100644 index 00000000000..02ec87c16dc --- /dev/null +++ b/test/onnx/averagepool_dilate_test.onnx @@ -0,0 +1,17 @@ +averagepool_dilate_test:¦ +Y +xy" AveragePool* + dilations@ * + kernel_shape@ * +pads@@ * +strides@ averagepool_dilate_testZ +x + + + +b +y + + + +B \ No newline at end of file diff --git a/test/onnx/gen_onnx.py b/test/onnx/gen_onnx.py index 9940859156a..141767cfc7f 100644 --- a/test/onnx/gen_onnx.py +++ b/test/onnx/gen_onnx.py @@ -276,6 +276,22 @@ def averagepool_1d_test(): return ([node], [x], [out]) +@onnx_test() +def averagepool_dilate_test(): + x = helper.make_tensor_value_info('x', TensorProto.FLOAT, [1, 4, 3]) + y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [1, 4, 2]) + + node = onnx.helper.make_node('AveragePool', + inputs=['x'], + outputs=['y'], + kernel_shape=[2], + strides=[1], + pads=[1, 1], + dilations=[3]) + + return ([node], [x], [y]) + + @onnx_test() def averagepool_3d_test(): x = helper.make_tensor_value_info('0', TensorProto.FLOAT, [1, 3, 5, 5, 5]) @@ -4882,6 +4898,22 @@ def maxpool_notset_test(): return ([node], [x], [y]) +@onnx_test() +def maxpool_dilate_test(): + x = helper.make_tensor_value_info('x', TensorProto.FLOAT, [1, 4, 3]) + y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [1, 4, 2]) + + node = onnx.helper.make_node('MaxPool', + inputs=['x'], + outputs=['y'], + kernel_shape=[2], + strides=[1], + pads=[1, 1], + dilations=[3]) + + return ([node], [x], [y]) + + @onnx_test() def maxpool_same_upper_test(): x = helper.make_tensor_value_info('x', TensorProto.FLOAT, [1, 1, 5, 5]) diff --git a/test/onnx/maxpool_dilate_test.onnx b/test/onnx/maxpool_dilate_test.onnx new file mode 100644 index 00000000000..017a6a99b37 --- /dev/null +++ b/test/onnx/maxpool_dilate_test.onnx @@ -0,0 +1,17 @@ +maxpool_dilate_test:ž +U +xy"MaxPool* + dilations@ * + kernel_shape@ * +pads@@ * +strides@ maxpool_dilate_testZ +x + + + +b +y + + + +B \ No newline at end of file diff --git a/test/onnx/onnx_test.cpp b/test/onnx/onnx_test.cpp index ab81fff7e62..d5db5ed6622 100644 --- a/test/onnx/onnx_test.cpp +++ b/test/onnx/onnx_test.cpp @@ -296,13 +296,32 @@ TEST_CASE(averagepool_1d_test) {{"mode", migraphx::op::pooling_mode::average}, {"padding", {0, 0}}, {"stride", {1}}, - {"lengths", {3}}}), + {"lengths", {3}}, + {"dilations", {1}}}), l0); auto prog = optimize_onnx("averagepool_1d_test.onnx"); EXPECT(p == prog); } +TEST_CASE(averagepool_dilate_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + auto input = mm->add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 4, 3}}); + mm->add_instruction(migraphx::make_op("pooling", + {{"mode", migraphx::op::pooling_mode::average}, + {"padding", {1, 1}}, + {"stride", {1}}, + {"lengths", {2}}, + {"dilations", {3}}}), + input); + + auto prog = optimize_onnx("averagepool_dilate_test.onnx"); + + EXPECT(p == prog); +} + TEST_CASE(averagepool_3d_test) { migraphx::program p; @@ -312,7 +331,8 @@ TEST_CASE(averagepool_3d_test) {{"mode", migraphx::op::pooling_mode::average}, {"padding", {0, 0, 0, 0, 0, 0}}, {"stride", {1, 1, 1}}, - {"lengths", {3, 3, 3}}}), + {"lengths", {3, 3, 3}}, + {"dilations", {1, 1, 1}}}), l0); auto prog = optimize_onnx("averagepool_3d_test.onnx"); @@ -332,6 +352,7 @@ TEST_CASE(averagepool_dyn_test) {"mode", migraphx::op::pooling_mode::average}, {"stride", {2, 2, 2}}, {"lengths", {3, 3, 3}}, + {"dilations", {1, 1, 1}}, {"padding", {1, 1, 1, 1, 1, 1}}, {"padding_mode", 0}, }), @@ -357,6 +378,7 @@ TEST_CASE(averagepool_dyn_autopad_test) {"mode", migraphx::op::pooling_mode::average}, {"stride", {2, 2, 2}}, {"lengths", {3, 3, 3}}, + {"dilations", {1, 1, 1}}, {"padding", {0, 0, 0, 0, 0, 0}}, {"padding_mode", migraphx::op::padding_mode_t::same_upper}, }), @@ -394,7 +416,8 @@ TEST_CASE(averagepool_notset_test) {{"mode", migraphx::op::pooling_mode::average}, {"padding", {2, 2, 2, 2}}, {"stride", {2, 2}}, - {"lengths", {6, 6}}}), + {"lengths", {6, 6}}, + {"dilations", {1, 1}}}), input); auto ret = mm->add_instruction( migraphx::make_op("slice", {{"axes", {2, 3}}, {"starts", {1, 1}}, {"ends", {2, 2}}}), ins); @@ -415,7 +438,8 @@ TEST_CASE(averagepool_nt_cip_test) {{"mode", migraphx::op::pooling_mode::average}, {"padding", {0, 0, 0, 0}}, {"stride", {2, 2}}, - {"lengths", {6, 6}}}), + {"lengths", {6, 6}}, + {"dilations", {1, 1}}}), ins_pad); mm->add_return({ret}); @@ -437,6 +461,7 @@ TEST_CASE(averagepool_same_lower_test) {"padding", {1, 1, 1, 1}}, {"stride", {1, 1}}, {"lengths", {2, 2}}, + {"dilations", {1, 1}}, {"padding_mode", migraphx::op::padding_mode_t::default_}, }), input); @@ -459,7 +484,8 @@ TEST_CASE(averagepool_sl_cip_test) {{"mode", migraphx::op::pooling_mode::average}, {"padding", {0, 0, 0, 0}}, {"stride", {1, 1}}, - {"lengths", {2, 2}}}), + {"lengths", {2, 2}}, + {"dilations", {1, 1}}}), ins_pad); mm->add_return({ret}); auto prog = migraphx::parse_onnx("averagepool_sl_cip_test.onnx"); @@ -476,7 +502,8 @@ TEST_CASE(averagepool_same_upper_test) {{"mode", migraphx::op::pooling_mode::average}, {"padding", {1, 1, 1, 1}}, {"stride", {1, 1}}, - {"lengths", {2, 2}}}), + {"lengths", {2, 2}}, + {"dilations", {1, 1}}}), input); auto ret = mm->add_instruction( migraphx::make_op("slice", {{"axes", {2, 3}}, {"starts", {1, 1}}, {"ends", {6, 6}}}), ins); @@ -1307,7 +1334,8 @@ TEST_CASE(conv_bn_relu_maxpool_test) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {0, 0, 0, 0}}, {"stride", {2, 2}}, - {"lengths", {2, 2}}}), + {"lengths", {2, 2}}, + {"dilations", {1, 1}}}), l7); auto prog = optimize_onnx("conv_bn_relu_maxpool_test.onnx"); @@ -1505,7 +1533,8 @@ TEST_CASE(conv_relu_maxpool_test) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {0, 0, 0, 0}}, {"stride", {2, 2}}, - {"lengths", {2, 2}}}), + {"lengths", {2, 2}}, + {"dilations", {1, 1}}}), l6); auto prog = optimize_onnx("conv_relu_maxpool_test.onnx"); @@ -1530,7 +1559,8 @@ TEST_CASE(conv_relu_maxpool_x2_test) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {0, 0, 0, 0}}, {"stride", {2, 2}}, - {"lengths", {2, 2}}}), + {"lengths", {2, 2}}, + {"dilations", {1, 1}}}), l6); auto l8 = mm->add_parameter("3", {migraphx::shape::float_type, {1, 5, 5, 5}}); @@ -1546,7 +1576,8 @@ TEST_CASE(conv_relu_maxpool_x2_test) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {0, 0, 0, 0}}, {"stride", {2, 2}}, - {"lengths", {2, 2}}}), + {"lengths", {2, 2}}, + {"dilations", {1, 1}}}), l13); auto prog = optimize_onnx("conv_relu_maxpool_x2_test.onnx"); @@ -4245,6 +4276,7 @@ TEST_CASE(lppool_l1_test) {"padding", {0, 0}}, {"stride", {1}}, {"lengths", {3}}, + {"dilations", {1}}, {"lp_order", 1}}), l0); auto prog = optimize_onnx("lppool_l1_test.onnx"); @@ -4261,6 +4293,7 @@ TEST_CASE(lppool_l2_test) {"padding", {0, 0}}, {"stride", {1}}, {"lengths", {3}}, + {"dilations", {1}}, {"lp_order", 2}}), l0); auto prog = optimize_onnx("lppool_l2_test.onnx"); @@ -4513,7 +4546,8 @@ TEST_CASE(maxpool_notset_test) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {0, 0, 1, 1}}, {"stride", {2, 2}}, - {"lengths", {6, 6}}}), + {"lengths", {6, 6}}, + {"dilations", {1, 1}}}), input); auto prog = optimize_onnx("maxpool_notset_test.onnx"); @@ -4521,6 +4555,24 @@ TEST_CASE(maxpool_notset_test) EXPECT(p == prog); } +TEST_CASE(maxpool_dilate_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + auto input = mm->add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 4, 3}}); + mm->add_instruction(migraphx::make_op("pooling", + {{"mode", migraphx::op::pooling_mode::max}, + {"padding", {1, 1}}, + {"stride", {1}}, + {"lengths", {2}}, + {"dilations", {3}}}), + input); + + auto prog = optimize_onnx("maxpool_dilate_test.onnx"); + + EXPECT(p == prog); +} + TEST_CASE(maxpool_same_upper_test) { migraphx::program p; @@ -4530,7 +4582,8 @@ TEST_CASE(maxpool_same_upper_test) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {0, 0, 1, 1}}, {"stride", {1, 1}}, - {"lengths", {2, 2}}}), + {"lengths", {2, 2}}, + {"dilations", {1, 1}}}), input); auto prog = optimize_onnx("maxpool_same_upper_test.onnx"); diff --git a/test/op_shape_test.cpp b/test/op_shape_test.cpp index b3f688ed4db..aaec91c2818 100644 --- a/test/op_shape_test.cpp +++ b/test/op_shape_test.cpp @@ -2202,7 +2202,8 @@ TEST_CASE(pooling_shape0) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {1}}, {"stride", {0}}, - {"lengths", {1}}}), + {"lengths", {1}}, + {"dilations", {1}}}), input); } @@ -2215,7 +2216,8 @@ TEST_CASE(pooling_shape1) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {0, 0}}, {"stride", {3, 3}}, - {"lengths", {1, 1}}}), + {"lengths", {1, 1}}, + {"dilations", {1, 1}}}), input); } @@ -2229,6 +2231,7 @@ TEST_CASE(pooling_shape2) {"padding", {0, 0}}, {"stride", {3, 3}}, {"lengths", {1, 1}}, + {"dilations", {1, 1}}, {"ceil_mode", true}}), input); } @@ -2243,6 +2246,7 @@ TEST_CASE(pooling_shape3) {"padding", {2, 2}}, {"stride", {3, 3}}, {"lengths", {3, 3}}, + {"dilations", {1, 1}}, {"ceil_mode", true}}), input); } @@ -2254,6 +2258,63 @@ TEST_CASE(pooling_shape4) tiny_input); } +TEST_CASE(pooling_shape5) +{ + migraphx::shape input{migraphx::shape::float_type, {4, 3, 3, 3}}; + migraphx::shape output{migraphx::shape::float_type, {4, 3, 1, 1}}; + expect_shape(output, + migraphx::make_op("pooling", + {{"mode", migraphx::op::pooling_mode::max}, + {"padding", {0, 0}}, + {"stride", {1, 1}}, + {"lengths", {2, 2}}, + {"dilations", {2, 2}}}), + input); +} + +TEST_CASE(pooling_shape6) +{ + migraphx::shape input{migraphx::shape::float_type, {4, 3, 3, 3}}; + migraphx::shape output{migraphx::shape::float_type, {4, 3, 2, 2}}; + expect_shape(output, + migraphx::make_op("pooling", + {{"mode", migraphx::op::pooling_mode::max}, + {"padding", {0, 0}}, + {"stride", {2, 2}}, + {"lengths", {1, 1}}, + {"dilations", {2, 2}}}), + input); +} + +TEST_CASE(pooling_shape7) +{ + migraphx::shape input{migraphx::shape::float_type, {4, 3, 3, 3}}; + migraphx::shape output{migraphx::shape::float_type, {4, 3, 2, 2}}; + expect_shape(output, + migraphx::make_op("pooling", + {{"mode", migraphx::op::pooling_mode::max}, + {"padding", {0, 0}}, + {"stride", {3, 3}}, + {"lengths", {1, 1}}, + {"dilations", {3, 3}}, + {"ceil_mode", true}}), + input); +} + +TEST_CASE(pooling_shape8) +{ + migraphx::shape input{migraphx::shape::float_type, {4, 3, 3, 3}}; + migraphx::shape output{migraphx::shape::float_type, {4, 3, 3, 3}}; + expect_shape(output, + migraphx::make_op("pooling", + {{"mode", migraphx::op::pooling_mode::max}, + {"padding", {2, 2}}, + {"stride", {1, 1}}, + {"lengths", {3, 3}}, + {"dilations", {2, 2}}}), + input); +} + TEST_CASE(pooling_dyn_shape0) { migraphx::shape input{migraphx::shape::float_type, {{1, 4}, {3, 3, {3}}, {3, 3, {3}}, {3, 3}}}; @@ -2261,7 +2322,8 @@ TEST_CASE(pooling_dyn_shape0) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {1}}, {"stride", {0}}, - {"lengths", {1}}}), + {"lengths", {1}}, + {"dilations", {1}}}), input); } @@ -2274,7 +2336,8 @@ TEST_CASE(pooling_dyn_shape1) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {0, 0}}, {"stride", {3, 3}}, - {"lengths", {1, 1}}}), + {"lengths", {1, 1}}, + {"dilations", {1, 1}}}), input); } @@ -2288,6 +2351,7 @@ TEST_CASE(pooling_dyn_shape2) {"padding", {0, 0}}, {"stride", {3, 3}}, {"lengths", {1, 1}}, + {"dilations", {1, 1}}, {"ceil_mode", true}}), input); } @@ -2302,7 +2366,8 @@ TEST_CASE(pooling_dyn_shape3) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {0, 0}}, {"stride", {3, 3}}, - {"lengths", {1, 1}}}), + {"lengths", {1, 1}}, + {"dilations", {1, 1}}}), input); } @@ -2317,6 +2382,7 @@ TEST_CASE(pooling_dyn_shape4) {"padding", {2, 2}}, {"stride", {3, 3}}, {"lengths", {3, 3}}, + {"dilations", {1, 1}}, {"ceil_mode", true}}), input); } diff --git a/test/py/onnx_backend_test.py b/test/py/onnx_backend_test.py index f51e4ce45f4..dc86d5ac3d0 100644 --- a/test/py/onnx_backend_test.py +++ b/test/py/onnx_backend_test.py @@ -576,6 +576,8 @@ def disabled_tests_onnx_1_9_0(backend_test): backend_test.exclude(r'test_gru_batchwise_cpu') backend_test.exclude(r'test_simple_rnn_batchwise_cpu') # from OnnxBackendPyTorchConvertedModelTest + # MaxPool dialtion is partially supported on GPU by a workaround + # But these tests require too large allocations to work properly backend_test.exclude(r'test_MaxPool1d_stride_padding_dilation_cpu') backend_test.exclude(r'test_MaxPool2d_stride_padding_dilation_cpu') diff --git a/test/ref/pooling.cpp b/test/ref/pooling.cpp index a24c5bc6dbc..bf3a01148bc 100644 --- a/test/ref/pooling.cpp +++ b/test/ref/pooling.cpp @@ -35,12 +35,13 @@ TEST_CASE(avgpool_rank3_test) { // 1D case 1, input is 3D migraphx::program p; - auto* mm = p.get_main_module(); - auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; - auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average}; - op.lengths = {2}; - op.padding = {0}; - op.stride = {1}; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average}; + op.lengths = {2}; + op.padding = {0}; + op.stride = {1}; + op.dilations = {1}; std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; auto l0 = mm->add_literal(migraphx::literal{s, data}); @@ -54,6 +55,103 @@ TEST_CASE(avgpool_rank3_test) EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); } +TEST_CASE(avgpool_rank3_dil_test) +{ + // 1D case 1, input is 3D + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average}; + op.lengths = {2}; + op.padding = {0}; + op.stride = {1}; + op.dilations = {2}; + + std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{0.35, 0.15, 0.85, 0.3, 0.1, 0.65}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + +TEST_CASE(avgpool_rank3_dil_test2) +{ + // 1D case 1, input is 3D + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average}; + op.lengths = {2}; + op.padding = {0}; + op.stride = {1}; + op.dilations = {3}; + + std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{0.2, 0.45, 0.35}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + +TEST_CASE(avgpool_rank3_pad_test) +{ + // 1D case 1, input is 3D + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average}; + op.lengths = {2}; + op.padding = {1}; + op.stride = {1}; + op.dilations = {1}; + + std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{ + 0.3, 0.25, 0.3, 0.25, 0.1, 0.8, 0.65, 0.7, 0.5, 0.1, 0.1, 0.4, 0.4, 0.35, 0.6}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + +TEST_CASE(avgpool_rank3_pad_dil_test) +{ + // 1D case 1, input is 3D + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average}; + op.lengths = {2}; + op.padding = {1}; + op.stride = {1}; + op.dilations = {3}; + + std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{0.4, 0.2, 0.2, 0.9, 0.45, 0.5, 0.1, 0.35, 0.7}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + TEST_CASE(avgpool_dyn_test) { // Dynamic input, no padding @@ -65,7 +163,8 @@ TEST_CASE(avgpool_dyn_test) {{"mode", migraphx::op::pooling_mode::average}, {"lengths", {2}}, {"padding", {0}}, - {"stride", {1}}}), + {"stride", {1}}, + {"dilations", {1}}}), x); p.compile(migraphx::make_target("ref")); @@ -82,7 +181,7 @@ TEST_CASE(avgpool_dyn_test) TEST_CASE(avgpool_dyn_pad_test) { - // Dynamic input with explicit padding/ + // Dynamic input with explicit padding migraphx::program p; auto* mm = p.get_main_module(); auto s = migraphx::shape{migraphx::shape::float_type, {{1, 3}, {3, 3}, {4, 4}}}; @@ -91,7 +190,8 @@ TEST_CASE(avgpool_dyn_pad_test) {{"mode", migraphx::op::pooling_mode::average}, {"lengths", {2}}, {"padding", {1}}, - {"stride", {1}}}), + {"stride", {1}}, + {"dilations", {1}}}), x); p.compile(migraphx::make_target("ref")); @@ -158,7 +258,8 @@ TEST_CASE(avgpool_dyn_auto_pad_1d_test) // padding added will be {1, 0} to make output // the same size as input {"padding_mode", migraphx::op::padding_mode_t::same_lower}, - {"stride", {1}}}), + {"stride", {1}}, + {"dilations", {1}}}), x); p.compile(migraphx::make_target("ref")); @@ -171,8 +272,8 @@ TEST_CASE(avgpool_dyn_auto_pad_1d_test) result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); // clang-format off - std::vector gold{0.3, 0.25, 0.3, 0.25, - 0.8, 0.65, 0.7, 0.5, + std::vector gold{0.3, 0.25, 0.3, 0.25, + 0.8, 0.65, 0.7, 0.5, 0.1, 0.4, 0.4, 0.35}; // clang-format on EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); @@ -190,7 +291,8 @@ TEST_CASE(avgpool_dyn_pad_ceil_test) {"lengths", {2, 3}}, {"padding", {1, 2}}, {"ceil_mode", true}, - {"stride", {1, 1}}}), + {"stride", {1, 1}}, + {"dilations", {1, 1}}}), x); p.compile(migraphx::make_target("ref")); @@ -219,12 +321,13 @@ TEST_CASE(avgpool_rank3_stride2_test) { // 1D case 2, stride 2 migraphx::program p; - auto* mm = p.get_main_module(); - auto s = migraphx::shape{migraphx::shape::float_type, {2, 2, 4}}; - auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average}; - op.lengths = {2}; - op.padding = {1}; - op.stride = {2}; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {2, 2, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average}; + op.lengths = {2}; + op.padding = {1}; + op.stride = {2}; + op.dilations = {1}; // clang-format off std::vector data{1.6321, -2.4186, 0.2239, -1.4232, @@ -252,12 +355,13 @@ TEST_CASE(avgpool_rank5_test) { // 3D, input is 5D migraphx::program p; - auto* mm = p.get_main_module(); - auto s = migraphx::shape{migraphx::shape::float_type, {2, 2, 3, 3, 3}}; - auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average}; - op.lengths = {2, 2, 2}; - op.padding = {0, 0, 0}; - op.stride = {1, 1, 1}; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {2, 2, 3, 3, 3}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average}; + op.lengths = {2, 2, 2}; + op.padding = {0, 0, 0}; + op.stride = {1, 1, 1}; + op.dilations = {1, 1, 1}; std::vector data{ -0.179, -1.756, 0.651, 1.955, 1.87, -0.604, 0.247, 0.449, -0.137, 1.187, 1.593, @@ -423,13 +527,14 @@ TEST_CASE(lppool_l1_norm_test) { // L1 norm test migraphx::program p; - auto* mm = p.get_main_module(); - auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; - auto op = migraphx::op::pooling{migraphx::op::pooling_mode::lpnorm}; - op.lengths = {2}; - op.padding = {0}; - op.stride = {1}; - op.lp_order = 1; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::lpnorm}; + op.lengths = {2}; + op.padding = {0}; + op.stride = {1}; + op.dilations = {1}; + op.lp_order = 1; std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; auto l0 = mm->add_literal(migraphx::literal{s, data}); @@ -449,13 +554,14 @@ TEST_CASE(lppool_l1_norm_test) // { // // padding too large for kernel size // migraphx::program p; -// auto* mm = p.get_main_module(); -// auto s = migraphx::shape{migraphx::shape::float_type, {1, 2, 5}}; -// auto op = migraphx::op::pooling{migraphx::op::pooling_mode::lpnorm}; -// op.lengths = {3}; -// op.padding = {2}; -// op.stride = {1}; -// op.lp_order = 1; +// auto* mm = p.get_main_module(); +// auto s = migraphx::shape{migraphx::shape::float_type, {1, 2, 5}}; +// auto op = migraphx::op::pooling{migraphx::op::pooling_mode::lpnorm}; +// op.lengths = {3}; +// op.padding = {2}; +// op.stride = {1}; +// op.dilations = {1}; +// op.lp_order = 1; // std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7}; // auto l0 = mm->add_literal(migraphx::literal{s, data}); @@ -468,13 +574,14 @@ TEST_CASE(lppool_l2_norm_test) { // L2 norm test migraphx::program p; - auto* mm = p.get_main_module(); - auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; - auto op = migraphx::op::pooling{migraphx::op::pooling_mode::lpnorm}; - op.lengths = {2}; - op.padding = {0}; - op.stride = {1}; - op.lp_order = 2; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::lpnorm}; + op.lengths = {2}; + op.padding = {0}; + op.stride = {1}; + op.dilations = {1}; + op.lp_order = 2; std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; auto l0 = mm->add_literal(migraphx::literal{s, data}); @@ -506,7 +613,8 @@ TEST_CASE(lppool_dyn_test) {{"mode", migraphx::op::pooling_mode::lpnorm}, {"lengths", {2}}, {"padding", {0}}, - {"stride", {1}}}), + {"stride", {1}}, + {"dilations", {1}}}), x); p.compile(migraphx::make_target("ref")); @@ -571,7 +679,8 @@ TEST_CASE(maxpool_test) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {0, 0}}, {"stride", {2, 2}}, - {"lengths", {3, 2}}}), + {"lengths", {3, 2}}, + {"dilations", {1, 1}}}), al); p.compile(migraphx::make_target("ref")); auto result = p.eval({}).back(); @@ -599,7 +708,8 @@ TEST_CASE(maxpool_pad_test) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {1, 1}}, {"stride", {2, 2}}, - {"lengths", {3, 2}}}), + {"lengths", {3, 2}}, + {"dilations", {1, 1}}}), al); // * * * * * * * * @@ -620,12 +730,13 @@ TEST_CASE(maxpool_rank3_test0) { // 1D case 1, input is 3D migraphx::program p; - auto* mm = p.get_main_module(); - auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; - auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; - op.lengths = {2}; - op.padding = {0}; - op.stride = {1}; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; + op.lengths = {2}; + op.padding = {0}; + op.stride = {1}; + op.dilations = {1}; std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; auto l0 = mm->add_literal(migraphx::literal{s, data}); @@ -643,12 +754,13 @@ TEST_CASE(maxpool_rank3_test1) { // 1D case 2, input is 3D migraphx::program p; - auto* mm = p.get_main_module(); - auto s = migraphx::shape{migraphx::shape::float_type, {2, 2, 5}}; - auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; - op.lengths = {2}; - op.padding = {0}; - op.stride = {2}; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {2, 2, 5}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; + op.lengths = {2}; + op.padding = {0}; + op.stride = {2}; + op.dilations = {1}; std::vector data{0.4975, -0.1226, -0.0405, -0.2861, -0.1227, -0.6186, -0.9618, 0.6022, -0.1912, 1.1925, 0.5493, 0.1692, -0.8039, -1.0281, @@ -664,6 +776,55 @@ TEST_CASE(maxpool_rank3_test1) EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); } +TEST_CASE(maxpool_rank3_test2) +{ + // 1D case 1, input is 3D + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; + op.lengths = {2}; + op.padding = {0}; + op.stride = {1}; + op.dilations = {2}; + + std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{0.4, 0.2, 0.9, 0.5, 0.1, 0.7}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + +TEST_CASE(maxpool_rank3_test4) +{ + // 1D case 1, input is 3D + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; + op.lengths = {2}; + op.padding = {1}; + op.stride = {1}; + op.dilations = {3}; + + std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{0.4, 0.3, 0.2, 0.9, 0.8, 0.5, 0.1, 0.6, 0.7}; + + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + TEST_CASE(maxpool_rank3_ceil_test) { // 1D case 2, input is 3D, ceil mode @@ -674,6 +835,7 @@ TEST_CASE(maxpool_rank3_ceil_test) op.lengths = {2}; op.padding = {0}; op.stride = {2}; + op.dilations = {1}; op.ceil_mode = true; // clang-format off @@ -702,12 +864,13 @@ TEST_CASE(maxpool_rank5_test) { // 3D, input is 5D migraphx::program p; - auto* mm = p.get_main_module(); - auto s = migraphx::shape{migraphx::shape::float_type, {2, 2, 3, 3, 3}}; - auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; - op.lengths = {2, 2, 2}; - op.padding = {0, 0, 0}; - op.stride = {2, 2, 2}; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {2, 2, 3, 3, 3}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; + op.lengths = {2, 2, 2}; + op.padding = {0, 0, 0}; + op.stride = {2, 2, 2}; + op.dilations = {1, 1, 1}; std::vector data{ -2.8029, 0.5861, 0.7015, 0.1297, -1.44, -1.9472, 0.7812, 2.408, -0.3145, 0.3405, @@ -741,7 +904,8 @@ TEST_CASE(maxpool_dyn_test) {{"mode", migraphx::op::pooling_mode::max}, {"lengths", {2}}, {"padding", {0}}, - {"stride", {1}}}), + {"stride", {1}}, + {"dilations", {1}}}), x); p.compile(migraphx::make_target("ref")); @@ -755,3 +919,29 @@ TEST_CASE(maxpool_dyn_test) std::vector gold{0.3, 0.4, 0.4, 0.8, 0.9, 0.9, 0.7, 0.7, 0.6}; EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); } + +TEST_CASE(maxpool_dyn_test2) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {{1, 4}, {3, 3}, {4, 4}}}; + auto x = mm->add_parameter("X", s); + mm->add_instruction(migraphx::make_op("pooling", + {{"mode", migraphx::op::pooling_mode::max}, + {"lengths", {2}}, + {"padding", {0}}, + {"stride", {1}}, + {"dilations", {2}}}), + x); + p.compile(migraphx::make_target("ref")); + + std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; + migraphx::shape input_fixed_shape{migraphx::shape::float_type, {1, 3, 4}}; + migraphx::parameter_map params; + params["X"] = migraphx::argument(input_fixed_shape, data.data()); + auto result = p.eval(params).back(); + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{0.4, 0.2, 0.9, 0.5, 0.1, 0.7}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} diff --git a/test/rewrite_pooling_test.cpp b/test/rewrite_pooling_test.cpp index b7977cd7a5c..b33386491d6 100644 --- a/test/rewrite_pooling_test.cpp +++ b/test/rewrite_pooling_test.cpp @@ -53,7 +53,8 @@ TEST_CASE(rewrite_pooling_test) {{"mode", mode}, {"padding", {0, 0, 0}}, {"stride", {1, 1, 1}}, - {"lengths", {3, 4, 5}}}), + {"lengths", {3, 4, 5}}, + {"dilations", {1, 1, 1}}}), input); m.add_return({ret}); return m; @@ -80,6 +81,483 @@ TEST_CASE(rewrite_pooling_test) migraphx::make_op("reduce_max", {{"axes", {2, 3, 4}}})); } +TEST_CASE(rewrite_pooling_dialtions_test) +{ + migraphx::shape s{migraphx::shape::float_type, {1, 1, 5, 5}}; + auto pooling_program = [&](const migraphx::op::pooling_mode mode) { + migraphx::module m; + auto input = m.add_parameter("x", s); + auto ret = m.add_instruction(migraphx::make_op("pooling", + {{"mode", mode}, + {"padding", {0, 0}}, + {"stride", {1, 1}}, + {"lengths", {2, 2}}, + {"dilations", {2, 2}}}), + input); + m.add_return({ret}); + return m; + }; + + auto opt_program = [&](const migraphx::op::pooling_mode mode) { + migraphx::module m; + auto input = m.add_parameter("x", s); + std::vector indices{0, 2, 1, 3, 2, 4}; + migraphx::shape s_indices{migraphx::shape::int32_type, {indices.size()}}; + auto i1 = m.add_literal(migraphx::literal{s_indices, indices}); + auto g1 = m.add_instruction(migraphx::make_op("gather", {{"axis", 2}}), input, i1); + auto i2 = m.add_literal(migraphx::literal{s_indices, indices}); + auto g2 = m.add_instruction(migraphx::make_op("gather", {{"axis", 3}}), g1, i2); + auto ret = m.add_instruction(migraphx::make_op("pooling", + {{"mode", mode}, + {"padding", {0, 0}}, + {"stride", {2, 2}}, + {"lengths", {2, 2}}, + {"dilations", {1, 1}}}), + g2); + m.add_return({ret}); + return m; + }; + + auto test_rewrite = [&](const migraphx::op::pooling_mode mode) { + migraphx::module m1 = pooling_program(mode); + migraphx::module m2 = opt_program(mode); + opt_pooling(m1); + EXPECT(m1 == m2); + }; + + test_rewrite(migraphx::op::pooling_mode::average); + test_rewrite(migraphx::op::pooling_mode::max); +} + +TEST_CASE(rewrite_pooling_dialtions_test2) +{ + migraphx::shape s{migraphx::shape::float_type, {1, 1, 5, 5, 5}}; + auto pooling_program = [&](const migraphx::op::pooling_mode mode) { + migraphx::module m; + auto input = m.add_parameter("x", s); + auto ret = m.add_instruction(migraphx::make_op("pooling", + {{"mode", mode}, + {"padding", {0, 0, 0}}, + {"stride", {1, 1, 1}}, + {"lengths", {2, 2, 2}}, + {"dilations", {2, 2, 2}}}), + input); + m.add_return({ret}); + return m; + }; + + auto opt_program = [&](const migraphx::op::pooling_mode mode) { + migraphx::module m; + auto input = m.add_parameter("x", s); + std::vector indices{0, 2, 1, 3, 2, 4}; + migraphx::shape s_indices{migraphx::shape::int32_type, {indices.size()}}; + auto i1 = m.add_literal(migraphx::literal{s_indices, indices}); + auto g1 = m.add_instruction(migraphx::make_op("gather", {{"axis", 2}}), input, i1); + auto i2 = m.add_literal(migraphx::literal{s_indices, indices}); + auto g2 = m.add_instruction(migraphx::make_op("gather", {{"axis", 3}}), g1, i2); + auto i3 = m.add_literal(migraphx::literal{s_indices, indices}); + auto g3 = m.add_instruction(migraphx::make_op("gather", {{"axis", 4}}), g2, i3); + auto ret = m.add_instruction(migraphx::make_op("pooling", + {{"mode", mode}, + {"padding", {0, 0, 0}}, + {"stride", {2, 2, 2}}, + {"lengths", {2, 2, 2}}, + {"dilations", {1, 1, 1}}}), + g3); + m.add_return({ret}); + return m; + }; + + auto test_rewrite = [&](const migraphx::op::pooling_mode mode) { + migraphx::module m1 = pooling_program(mode); + migraphx::module m2 = opt_program(mode); + opt_pooling(m1); + EXPECT(m1 == m2); + }; + + test_rewrite(migraphx::op::pooling_mode::average); + test_rewrite(migraphx::op::pooling_mode::max); +} + +TEST_CASE(rewrite_pooling_dialtions_test3) +{ + migraphx::shape s{migraphx::shape::float_type, {2, 2, 5}}; + auto pooling_program = [&]() { + migraphx::module m; + + auto input = m.add_parameter("x", s); + auto ret = + m.add_instruction(migraphx::make_op("pooling", + {{"mode", migraphx::op::pooling_mode::average}, + {"padding", {1}}, + {"stride", {1}}, + {"lengths", {3}}, + {"dilations", {2}}}), + input); + m.add_return({ret}); + return m; + }; + + migraphx::module m1 = pooling_program(); + migraphx::module m2 = m1; + + opt_pooling(m1); + EXPECT(m1 == m2); +} + +TEST_CASE(rewrite_pooling_dialtions_test4) +{ + migraphx::shape s{migraphx::shape::float_type, {1, 1, 5, 5}}; + auto pooling_program = [&](const migraphx::op::pooling_mode mode) { + migraphx::module m; + auto input = m.add_parameter("x", s); + auto ret = m.add_instruction(migraphx::make_op("pooling", + {{"mode", mode}, + {"padding", {1, 0}}, + {"stride", {1, 3}}, + {"lengths", {3, 1}}, + {"dilations", {1, 2}}}), + input); + m.add_return({ret}); + return m; + }; + + auto opt_program = [&](const migraphx::op::pooling_mode mode) { + migraphx::module m; + auto input = m.add_parameter("x", s); + std::vector col_indices{0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14}; + migraphx::shape s_col_indices{migraphx::shape::int32_type, {col_indices.size()}}; + std::vector row_indices{0, 3}; + migraphx::shape s_row_indices{migraphx::shape::int32_type, {row_indices.size()}}; + auto p = + m.add_instruction(migraphx::make_op("pad", + {{"pads", {0, 0, 1, 0, 0, 0, 1, 0}}, + {"value", std::numeric_limits::lowest()}}), + input); + auto i1 = m.add_literal(migraphx::literal{s_col_indices, col_indices}); + auto g1 = m.add_instruction(migraphx::make_op("gather", {{"axis", 2}}), p, i1); + auto i2 = m.add_literal(migraphx::literal{s_row_indices, row_indices}); + auto g2 = m.add_instruction(migraphx::make_op("gather", {{"axis", 3}}), g1, i2); + auto ret = m.add_instruction(migraphx::make_op("pooling", + {{"mode", mode}, + {"padding", {0, 0}}, + {"stride", {3, 1}}, + {"lengths", {3, 1}}, + {"dilations", {1, 1}}}), + g2); + m.add_return({ret}); + return m; + }; + + auto test_rewrite = [&](const migraphx::op::pooling_mode mode) { + migraphx::module m1 = pooling_program(mode); + migraphx::module m2 = opt_program(mode); + opt_pooling(m1); + EXPECT(m1 == m2); + }; + + // Average won't work because of padding + test_rewrite(migraphx::op::pooling_mode::max); +} + +TEST_CASE(rewrite_pooling_dialtions_test5) +{ + migraphx::shape s{migraphx::shape::float_type, {1, 1, 5, 5}}; + auto pooling_program = [&](const migraphx::op::pooling_mode mode) { + migraphx::module m; + auto input = m.add_parameter("x", s); + auto ret = m.add_instruction(migraphx::make_op("pooling", + {{"mode", mode}, + {"padding", {0, 0}}, + {"stride", {2, 3}}, + {"lengths", {2, 1}}, + {"dilations", {1, 2}}}), + input); + m.add_return({ret}); + return m; + }; + + auto opt_program = [&](const migraphx::op::pooling_mode mode) { + migraphx::module m; + auto input = m.add_parameter("x", s); + std::vector col_indices{0, 1, 2, 3}; + migraphx::shape s_col_indices{migraphx::shape::int32_type, {col_indices.size()}}; + std::vector row_indices{0, 3}; + migraphx::shape s_row_indices{migraphx::shape::int32_type, {row_indices.size()}}; + auto i1 = m.add_literal(migraphx::literal{s_col_indices, col_indices}); + auto g1 = m.add_instruction(migraphx::make_op("gather", {{"axis", 2}}), input, i1); + auto i2 = m.add_literal(migraphx::literal{s_row_indices, row_indices}); + auto g2 = m.add_instruction(migraphx::make_op("gather", {{"axis", 3}}), g1, i2); + auto ret = m.add_instruction(migraphx::make_op("pooling", + {{"mode", mode}, + {"padding", {0, 0}}, + {"stride", {2, 1}}, + {"lengths", {2, 1}}, + {"dilations", {1, 1}}}), + g2); + m.add_return({ret}); + return m; + }; + + auto test_rewrite = [&](const migraphx::op::pooling_mode mode) { + migraphx::module m1 = pooling_program(mode); + migraphx::module m2 = opt_program(mode); + opt_pooling(m1); + EXPECT(m1 == m2); + }; + + test_rewrite(migraphx::op::pooling_mode::average); + test_rewrite(migraphx::op::pooling_mode::max); +} + +TEST_CASE(rewrite_avgpool_rank3_dil_test) +{ + // 1D case 1, input is 3D + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average}; + op.lengths = {2}; + op.padding = {0}; + op.stride = {1}; + op.dilations = {2}; + + std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + opt_pooling(*mm); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{0.35, 0.15, 0.85, 0.3, 0.1, 0.65}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + +TEST_CASE(rewrite_avgpool_rank3_dil_test2) +{ + // 1D case 1, input is 3D + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average}; + op.lengths = {2}; + op.padding = {0}; + op.stride = {1}; + op.dilations = {3}; + + std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + opt_pooling(*mm); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{0.2, 0.45, 0.35}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + +TEST_CASE(rewrite_avgpool_rank4_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 1, 5, 5}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average}; + op.lengths = {2, 1}; + op.padding = {0, 0}; + op.stride = {2, 3}; + op.dilations = {1, 2}; + + std::vector data(25); + std::iota(data.begin(), data.end(), 1); + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + opt_pooling(*mm); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{3.5, 6.5, 13.5, 16.5}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + +TEST_CASE(rewrite_maxpool_rank3_test) +{ + // 1D case 1, input is 3D + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; + op.lengths = {2}; + op.padding = {0}; + op.stride = {1}; + op.dilations = {2}; + + std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + opt_pooling(*mm); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{0.4, 0.2, 0.9, 0.5, 0.1, 0.7}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + +TEST_CASE(rewrite_maxpool_rank3_test2) +{ + // 1D case 1, input is 3D + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; + op.lengths = {2}; + op.padding = {1}; + op.stride = {1}; + op.dilations = {3}; + + std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + opt_pooling(*mm); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{0.4, 0.3, 0.2, 0.9, 0.8, 0.5, 0.1, 0.6, 0.7}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + +TEST_CASE(rewrite_maxpool_rank3_test3) +{ + // 1D case 1, input is 3D + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 3, 4}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; + op.lengths = {3}; + op.padding = {2}; + op.stride = {2}; + op.dilations = {3}; + + std::vector data{0.3, 0.2, 0.4, 0.1, 0.8, 0.5, 0.9, 0.1, 0.1, 0.7, 0.1, 0.6}; + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + opt_pooling(*mm); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{0.2, 0.5, 0.7}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + +TEST_CASE(rewrite_maxpool_rank4_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {1, 1, 5, 5}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; + op.lengths = {3, 1}; + op.padding = {1, 0}; + op.stride = {1, 3}; + op.dilations = {1, 2}; + + std::vector data(25); + std::iota(data.begin(), data.end(), 1); + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + opt_pooling(*mm); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{6, 9, 11, 14, 16, 19, 21, 24, 21, 24}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + +TEST_CASE(maxpool_rank5_test) +{ + // 3D, input is 5D + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {2, 2, 3, 3, 3}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; + op.lengths = {2, 2, 2}; + op.padding = {0, 0, 0}; + op.stride = {1, 1, 1}; + op.dilations = {2, 2, 2}; + + std::vector data{ + -2.8029, 0.5861, 0.7015, 0.1297, -1.44, -1.9472, 0.7812, 2.408, -0.3145, 0.3405, + -0.9146, 0.0624, 1.5064, -0.8345, 1.7977, 1.8949, 1.0073, -0.2102, -0.042, -0.7146, + 0.6227, -0.5263, -2.2598, 0.1713, 0.449, 0.5303, -0.8622, -0.5691, 0.907, -0.0569, + -1.5348, -0.4109, -0.1461, -0.5445, 0.4266, 0.2282, 1.3655, -2.1519, 0.6068, -0.2001, + -0.4702, 0.3864, 1.7083, 0.9096, 0.4286, -1.8866, 0.7034, 0.0293, 1.4587, 0.7672, + -2.8614, 0.8124, -0.053, 1.0449, 0.845, -0.0131, 0.1139, -0.859, -1.2681, -0.6337, + -0.4644, 0.1938, 0.2889, 0.9035, 0.7118, -0.5767, 0.4577, -0.0549, 0.2237, 0.5756, + 0.0677, -0.0223, -0.329, 0.2364, 2.7666, -0.7417, -1.3196, -0.2655, 0.1698, -0.1777, + -0.9427, 2.6859, -0.7501, 0.5175, 1.0029, -2.6436, -0.4388, -1.2348, -0.1539, -0.6229, + -0.4136, 0.5085, 0.4136, -0.6439, -1.1953, -0.406, -0.0195, 0.1869, -0.8664, 1.1364, + 0.5041, 0.0647, 0.1941, -1.0819, -0.4629, -0.5107, 0.3612, -0.3583}; + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + opt_pooling(*mm); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{0.7812, 1.0449, 2.7666, 2.6859}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + +TEST_CASE(maxpool_rank5_test2) +{ + // 3D, input is 5D + migraphx::program p; + auto* mm = p.get_main_module(); + auto s = migraphx::shape{migraphx::shape::float_type, {2, 2, 3, 3, 3}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::max}; + op.lengths = {2, 2, 2}; + op.padding = {2, 2, 2}; + op.stride = {2, 2, 2}; + op.dilations = {3, 3, 3}; + + std::vector data{ + -2.8029, 0.5861, 0.7015, 0.1297, -1.44, -1.9472, 0.7812, 2.408, -0.3145, 0.3405, + -0.9146, 0.0624, 1.5064, -0.8345, 1.7977, 1.8949, 1.0073, -0.2102, -0.042, -0.7146, + 0.6227, -0.5263, -2.2598, 0.1713, 0.449, 0.5303, -0.8622, -0.5691, 0.907, -0.0569, + -1.5348, -0.4109, -0.1461, -0.5445, 0.4266, 0.2282, 1.3655, -2.1519, 0.6068, -0.2001, + -0.4702, 0.3864, 1.7083, 0.9096, 0.4286, -1.8866, 0.7034, 0.0293, 1.4587, 0.7672, + -2.8614, 0.8124, -0.053, 1.0449, 0.845, -0.0131, 0.1139, -0.859, -1.2681, -0.6337, + -0.4644, 0.1938, 0.2889, 0.9035, 0.7118, -0.5767, 0.4577, -0.0549, 0.2237, 0.5756, + 0.0677, -0.0223, -0.329, 0.2364, 2.7666, -0.7417, -1.3196, -0.2655, 0.1698, -0.1777, + -0.9427, 2.6859, -0.7501, 0.5175, 1.0029, -2.6436, -0.4388, -1.2348, -0.1539, -0.6229, + -0.4136, 0.5085, 0.4136, -0.6439, -1.1953, -0.406, -0.0195, 0.1869, -0.8664, 1.1364, + 0.5041, 0.0647, 0.1941, -1.0819, -0.4629, -0.5107, 0.3612, -0.3583}; + auto l0 = mm->add_literal(migraphx::literal{s, data}); + mm->add_instruction(op, l0); + opt_pooling(*mm); + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + std::vector results_vector; + result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); }); + std::vector gold{-0.8345, 1.5064, -0.9146, 0.3405, -1.44, 0.1297, 0.5861, -2.8029, + -0.4702, -0.2001, -2.1519, 1.3655, -0.4109, -1.5348, 0.907, -0.5691, + -0.0549, 0.4577, 0.7118, 0.9035, -1.2681, -0.859, -0.0131, 0.845, + -1.1953, -0.6439, 0.5085, -0.4136, -2.6436, 1.0029, -0.7501, 2.6859}; + EXPECT(migraphx::verify::verify_rms_range(results_vector, gold)); +} + TEST_CASE(rewrite_avepooling_na1_test) { migraphx::shape s{migraphx::shape::float_type, {2, 2, 3, 4, 5}}; @@ -92,7 +570,8 @@ TEST_CASE(rewrite_avepooling_na1_test) {{"mode", migraphx::op::pooling_mode::average}, {"padding", {0, 1, 0}}, {"stride", {1, 1, 1}}, - {"lengths", {3, 4, 5}}}), + {"lengths", {3, 4, 5}}, + {"dilations", {1, 1, 1}}}), input); m.add_return({ret}); return m; @@ -117,7 +596,8 @@ TEST_CASE(rewrite_avepooling_na2_test) {{"mode", migraphx::op::pooling_mode::average}, {"padding", {0, 0, 0}}, {"stride", {1, 2, 1}}, - {"lengths", {3, 4, 5}}}), + {"lengths", {3, 4, 5}}, + {"dilations", {1, 1, 1}}}), input); m.add_return({ret}); return m; @@ -141,7 +621,8 @@ TEST_CASE(rewrite_avepooling_na3_test) {{"mode", migraphx::op::pooling_mode::max}, {"padding", {0, 0, 0}}, {"stride", {1, 1, 1}}, - {"lengths", {3, 3, 5}}}), + {"lengths", {3, 3, 5}}, + {"dilations", {1, 1, 1}}}), input); m.add_return({ret}); return m; @@ -169,7 +650,8 @@ TEST_CASE(literal_rewrite_pooling_test) {{"mode", mode}, {"padding", {0, 0, 0}}, {"stride", {1, 1, 1}}, - {"lengths", {3, 4, 5}}}), + {"lengths", {3, 4, 5}}, + {"dilations", {1, 1, 1}}}), input); mm->add_return({ret}); return p; diff --git a/test/simplify_qdq_test.cpp b/test/simplify_qdq_test.cpp index 3cc4f77ff23..68dc845123a 100644 --- a/test/simplify_qdq_test.cpp +++ b/test/simplify_qdq_test.cpp @@ -788,6 +788,7 @@ TEST_CASE(conv_pooling_dot) {"padding", {0, 0, 0, 0}}, {"stride", {1, 1}}, {"lengths", {7, 7}}, + {"dilations", {1, 1}}, {"ceil_mode", 0}}), a1); auto fl = m1.add_instruction(migraphx::make_op("flatten", {{"axis", 1}}), ap); @@ -835,6 +836,7 @@ TEST_CASE(conv_pooling_dot) {"padding", {0, 0, 0, 0}}, {"stride", {1, 1}}, {"lengths", {7, 7}}, + {"dilations", {1, 1}}, {"ceil_mode", 0}}), a1); auto fl = m2.add_instruction(migraphx::make_op("flatten", {{"axis", 1}}), ap); @@ -896,6 +898,7 @@ TEST_CASE(mobilenet_snippet) {"padding", {0, 0, 0, 0}}, {"stride", {1, 1}}, {"lengths", {7, 7}}, + {"dilations", {1, 1}}, {"ceil_mode", 0}}), d6); auto q3 = add_quantize_op(mm, "quantizelinear", ap, scale, zero); diff --git a/test/verify/test_avg_pooling_1d.cpp b/test/verify/test_avg_pooling_1d.cpp index ef82b86ee60..81e0e03d85c 100644 --- a/test/verify/test_avg_pooling_1d.cpp +++ b/test/verify/test_avg_pooling_1d.cpp @@ -35,7 +35,7 @@ struct test_avg_pooling_1d : verify_program auto* mm = p.get_main_module(); auto input = mm->add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 5}}); - auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average, {0}, {1}, {3}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average, {0}, {1}, {3}, {1}}; mm->add_instruction(op, input); return p; } diff --git a/test/verify/test_avg_pooling_3d.cpp b/test/verify/test_avg_pooling_3d.cpp index 6e83e15148a..cc4bf42faa1 100644 --- a/test/verify/test_avg_pooling_3d.cpp +++ b/test/verify/test_avg_pooling_3d.cpp @@ -36,7 +36,7 @@ struct test_avg_pooling_3d : verify_program auto input = mm->add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 5, 5, 5}}); auto op = migraphx::op::pooling{ - migraphx::op::pooling_mode::average, {1, 1, 1}, {3, 3, 3}, {3, 3, 3}}; + migraphx::op::pooling_mode::average, {1, 1, 1}, {3, 3, 3}, {3, 3, 3}, {1, 1, 1}}; mm->add_instruction(op, input); return p; } diff --git a/test/verify/test_avg_pooling_3d_opt.cpp b/test/verify/test_avg_pooling_3d_opt.cpp index 541d6fcbfe6..34e7122e9c4 100644 --- a/test/verify/test_avg_pooling_3d_opt.cpp +++ b/test/verify/test_avg_pooling_3d_opt.cpp @@ -36,7 +36,7 @@ struct test_avg_pooling_3d_opt : verify_program auto input = mm->add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 2, 3, 3, 3}}); auto op = migraphx::op::pooling{ - migraphx::op::pooling_mode::average, {0, 0, 0}, {1, 1, 1}, {3, 3, 3}}; + migraphx::op::pooling_mode::average, {0, 0, 0}, {1, 1, 1}, {3, 3, 3}, {1, 1, 1}}; mm->add_instruction(op, input); return p; } diff --git a/test/verify/test_avg_pooling_ceil_3d.cpp b/test/verify/test_avg_pooling_ceil_3d.cpp index bfd7ee85b13..ef2518c0639 100644 --- a/test/verify/test_avg_pooling_ceil_3d.cpp +++ b/test/verify/test_avg_pooling_ceil_3d.cpp @@ -37,7 +37,7 @@ struct test_avg_pooling_ceil_3d : verify_program auto input = mm->add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 5, 5, 5}}); auto op = migraphx::op::pooling{ - migraphx::op::pooling_mode::average, {1, 1, 1}, {3, 3, 3}, {3, 3, 3}, true}; + migraphx::op::pooling_mode::average, {1, 1, 1}, {3, 3, 3}, {3, 3, 3}, {1, 1, 1}, true}; mm->add_instruction(op, input); return p; } diff --git a/test/verify/test_avg_pooling_pad.cpp b/test/verify/test_avg_pooling_pad.cpp index 4e7f60d2df0..69a4d522b5e 100644 --- a/test/verify/test_avg_pooling_pad.cpp +++ b/test/verify/test_avg_pooling_pad.cpp @@ -36,7 +36,7 @@ struct test_avg_pooling_pad : verify_program auto* mm = p.get_main_module(); auto input = mm->add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 7}}); - auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average, {2}, {1}, {3}}; + auto op = migraphx::op::pooling{migraphx::op::pooling_mode::average, {2}, {1}, {3}, {1}}; mm->add_instruction(op, input); return p; } diff --git a/test/verify/test_concat_pooling.cpp b/test/verify/test_concat_pooling.cpp index c08f872f626..21cbfddf57a 100644 --- a/test/verify/test_concat_pooling.cpp +++ b/test/verify/test_concat_pooling.cpp @@ -47,7 +47,8 @@ struct test_concat_pooling : verify_program {{"mode", migraphx::op::pooling_mode::average}, {"padding", {0, 0}}, {"stride", {1, 1}}, - {"lengths", {8, 8}}}), + {"lengths", {8, 8}}, + {"dilations", {1, 1}}}), concat_t); mm->add_instruction(migraphx::make_op("relu"), pooling); return p; diff --git a/test/verify/test_conv_bn_relu_pooling.cpp b/test/verify/test_conv_bn_relu_pooling.cpp index e260920811a..4d4c8abb1f4 100644 --- a/test/verify/test_conv_bn_relu_pooling.cpp +++ b/test/verify/test_conv_bn_relu_pooling.cpp @@ -76,7 +76,8 @@ struct test_conv_bn_relu_pooling : verify_program {{"mode", migraphx::op::pooling_mode::average}, {"padding", {1, 1}}, {"stride", {2, 2}}, - {"lengths", {3, 3}}}), + {"lengths", {3, 3}}, + {"dilations", {1, 1}}}), relu); return p; } diff --git a/test/verify/test_conv_bn_relu_pooling2.cpp b/test/verify/test_conv_bn_relu_pooling2.cpp index 3568dab6666..39abacd7c28 100644 --- a/test/verify/test_conv_bn_relu_pooling2.cpp +++ b/test/verify/test_conv_bn_relu_pooling2.cpp @@ -92,7 +92,8 @@ struct test_conv_bn_relu_pooling2 : verify_program {{"mode", migraphx::op::pooling_mode::average}, {"padding", {1, 1}}, {"stride", {2, 2}}, - {"lengths", {3, 3}}}), + {"lengths", {3, 3}}, + {"dilations", {1, 1}}}), relu); return p; } diff --git a/test/verify/test_max_pooling_ceil_3d.cpp b/test/verify/test_max_pooling_ceil_3d.cpp index ff8283eb1dd..239a0806e49 100644 --- a/test/verify/test_max_pooling_ceil_3d.cpp +++ b/test/verify/test_max_pooling_ceil_3d.cpp @@ -36,7 +36,7 @@ struct test_max_pooling_ceil_3d : verify_program auto input = mm->add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 5, 5, 5}}); auto op = migraphx::op::pooling{ - migraphx::op::pooling_mode::max, {1, 1, 1}, {3, 3, 3}, {3, 3, 3}, true}; + migraphx::op::pooling_mode::max, {1, 1, 1}, {3, 3, 3}, {3, 3, 3}, {1, 1, 1}, true}; mm->add_instruction(op, input); return p; }