"...composable_kernel_onnxruntime.git" did not exist on "d2436a58c4566697abab96b08a016ddaa4159a12"
Commit 376b18af authored by charlie's avatar charlie
Browse files

Initial

parent f2667056
...@@ -36,8 +36,6 @@ namespace op { ...@@ -36,8 +36,6 @@ namespace op {
enum padding_mode_t enum padding_mode_t
{ {
default_, // NOLINT default_, // NOLINT
same,
valid,
same_lower, same_lower,
same_upper same_upper
}; };
......
...@@ -37,13 +37,12 @@ namespace op { ...@@ -37,13 +37,12 @@ namespace op {
struct convolution struct convolution
{ {
std::vector<std::size_t> padding = {0, 0}; std::vector<std::size_t> padding = {};
std::vector<std::size_t> stride = {1, 1}; std::vector<std::size_t> stride = {1, 1};
std::vector<std::size_t> dilation = {1, 1}; std::vector<std::size_t> dilation = {1, 1};
int group = 1; int group = 1;
padding_mode_t padding_mode = default_; padding_mode_t padding_mode = default_;
bool use_dynamic_same_auto_pad = false;
template <class Self, class F> template <class Self, class F>
static auto reflect(Self& self, F f) static auto reflect(Self& self, F f)
...@@ -52,8 +51,7 @@ struct convolution ...@@ -52,8 +51,7 @@ struct convolution
f(self.stride, "stride"), f(self.stride, "stride"),
f(self.dilation, "dilation"), f(self.dilation, "dilation"),
f(self.group, "group"), f(self.group, "group"),
f(self.padding_mode, "padding_mode"), f(self.padding_mode, "padding_mode"));
f(self.use_dynamic_same_auto_pad, "use_dynamic_same_auto_pad"));
} }
std::string name() const { return "convolution"; } std::string name() const { return "convolution"; }
...@@ -93,13 +91,6 @@ struct convolution ...@@ -93,13 +91,6 @@ struct convolution
x_shape.lens().at(1) != (w_shape.lens().at(1) * group)) x_shape.lens().at(1) != (w_shape.lens().at(1) * group))
MIGRAPHX_THROW("CONVOLUTION: mismatched channel numbers"); MIGRAPHX_THROW("CONVOLUTION: mismatched channel numbers");
std::vector<op::padding_mode_t> dyn_pad_modes = {op::padding_mode_t::same_upper,
op::padding_mode_t::same_lower};
if(use_dynamic_same_auto_pad and not contains(dyn_pad_modes, padding_mode))
{
MIGRAPHX_THROW("CONVOLUTION: use_dynamic_same_auto_pad set with invalid padding mode");
}
if(x_shape.dynamic() or w_shape.dynamic()) if(x_shape.dynamic() or w_shape.dynamic())
{ {
return dynamic_compute_shape(x_shape, w_shape); return dynamic_compute_shape(x_shape, w_shape);
...@@ -161,7 +152,7 @@ struct convolution ...@@ -161,7 +152,7 @@ struct convolution
dynamic_shape_push_back(w_shape); dynamic_shape_push_back(w_shape);
const size_t num_spatial_dims = x_shape.max_lens().size() - 2; const size_t num_spatial_dims = x_shape.max_lens().size() - 2;
if(use_dynamic_same_auto_pad) if(padding_mode != op::padding_mode_t::default_)
{ {
for(std::size_t i = 0; i < num_spatial_dims; ++i) for(std::size_t i = 0; i < num_spatial_dims; ++i)
{ {
......
...@@ -37,13 +37,12 @@ namespace op { ...@@ -37,13 +37,12 @@ namespace op {
struct quant_convolution struct quant_convolution
{ {
std::vector<std::size_t> padding = {0, 0}; std::vector<std::size_t> padding = {};
std::vector<std::size_t> stride = {1, 1}; std::vector<std::size_t> stride = {1, 1};
std::vector<std::size_t> dilation = {1, 1}; std::vector<std::size_t> dilation = {1, 1};
padding_mode_t padding_mode = default_; padding_mode_t padding_mode = default_;
int group = 1; int group = 1;
bool use_dynamic_same_auto_pad = false;
template <class Self, class F> template <class Self, class F>
static auto reflect(Self& self, F f) static auto reflect(Self& self, F f)
...@@ -52,8 +51,7 @@ struct quant_convolution ...@@ -52,8 +51,7 @@ struct quant_convolution
f(self.stride, "stride"), f(self.stride, "stride"),
f(self.dilation, "dilation"), f(self.dilation, "dilation"),
f(self.padding_mode, "padding_mode"), f(self.padding_mode, "padding_mode"),
f(self.group, "group"), f(self.group, "group"));
f(self.use_dynamic_same_auto_pad, "use_dynamic_same_auto_pad"));
} }
value attributes() const value attributes() const
......
...@@ -129,7 +129,6 @@ struct parse_convolution : op_parser<parse_convolution> ...@@ -129,7 +129,6 @@ struct parse_convolution : op_parser<parse_convolution>
} }
else else
{ {
values["padding_mode"] = to_value(op::padding_mode_t::same);
// kernel shape will be fixed, so max_lens() == min_len() for kernel lengths // kernel shape will be fixed, so max_lens() == min_len() for kernel lengths
auto weight_lens = weights->get_shape().max_lens(); auto weight_lens = weights->get_shape().max_lens();
std::vector<std::size_t> k_lens(weight_lens.begin() + 2, weight_lens.end()); std::vector<std::size_t> k_lens(weight_lens.begin() + 2, weight_lens.end());
......
...@@ -62,9 +62,9 @@ std::vector<std::size_t> calc_dyn_auto_pad(std::vector<std::size_t> tensor_lens, ...@@ -62,9 +62,9 @@ std::vector<std::size_t> calc_dyn_auto_pad(std::vector<std::size_t> tensor_lens,
padding.resize(2 * k_lens.size()); padding.resize(2 * k_lens.size());
for(std::size_t i = 0; i < padding.size() / 2; i++) for(std::size_t i = 0; i < padding.size() / 2; i++)
{ {
std::ptrdiff_t input_dim = tensor_lens[i]; std::ptrdiff_t input_dim = tensor_lens[i + 2];
std::ptrdiff_t stride = strides[i]; std::ptrdiff_t stride = strides[i];
std::ptrdiff_t weight_dim = k_lens[i]; std::ptrdiff_t weight_dim = k_lens[i + 2];
std::ptrdiff_t dilation = dilations[i]; std::ptrdiff_t dilation = dilations[i];
std::ptrdiff_t output_dim = (input_dim + stride - 1) / stride; // round up result std::ptrdiff_t output_dim = (input_dim + stride - 1) / stride; // round up result
std::ptrdiff_t new_weight_dim = weight_dim + (weight_dim - 1) * (dilation - 1); std::ptrdiff_t new_weight_dim = weight_dim + (weight_dim - 1) * (dilation - 1);
......
...@@ -237,12 +237,12 @@ struct ref_convolution : auto_register_op<ref_convolution<Op>> ...@@ -237,12 +237,12 @@ struct ref_convolution : auto_register_op<ref_convolution<Op>>
argument compute(context&, shape output_shape, std::vector<argument> args) const argument compute(context&, shape output_shape, std::vector<argument> args) const
{ {
std::vector<std::size_t> padding; std::vector<std::size_t> padding;
if(op.use_dynamic_same_auto_pad) if(op.padding_mode != op::padding_mode_t::default_)
{ {
auto input_lens = args[0].get_shape().lens(); auto input_lens = args[0].get_shape().lens();
std::vector<std::size_t> img_lens{input_lens.begin() + 2, input_lens.end()}; std::vector<std::size_t> img_lens{input_lens.begin(), input_lens.end()};
auto weights_lens = args[1].get_shape().lens(); auto weights_lens = args[1].get_shape().lens();
std::vector<std::size_t> k_lens{weights_lens.begin() + 2, weights_lens.end()}; std::vector<std::size_t> k_lens{weights_lens.begin(), weights_lens.end()};
padding = calc_dyn_auto_pad(img_lens, k_lens, op.stride, op.dilation); padding = calc_dyn_auto_pad(img_lens, k_lens, op.stride, op.dilation);
output_shape = output_shape =
compute_padded_shape({args.at(0).get_shape(), args.at(1).get_shape()}, padding); compute_padded_shape({args.at(0).get_shape(), args.at(1).get_shape()}, padding);
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment