diff --git a/mindspore/ccsrc/device/cpu/kernel/mkldnn/conv2d_cpu_kernel.cc b/mindspore/ccsrc/device/cpu/kernel/mkldnn/conv2d_cpu_kernel.cc index f7527c47507e68a0097b51cbe1584bbf777bfe70..5d63aee6cd23b920bed0eeb20a0a3c44e0c583dd 100644 --- a/mindspore/ccsrc/device/cpu/kernel/mkldnn/conv2d_cpu_kernel.cc +++ b/mindspore/ccsrc/device/cpu/kernel/mkldnn/conv2d_cpu_kernel.cc @@ -35,8 +35,22 @@ void Conv2dCPUKernel::InitKernel(const CNodePtr &kernel_node) { dnnl::memory::desc dst_desc = GetDefaultMemDesc(dst_shape); int kernel_size = SizeToInt(weight_shape[3]); - int stride = AnfAlgo::GetNodeAttr(kernel_node, STRIDE); - int dilation = AnfAlgo::GetNodeAttr(kernel_node, DILATION); + auto stride_ori = AnfAlgo::GetNodeAttr>(kernel_node, STRIDE); + auto dilation_ori = AnfAlgo::GetNodeAttr>(kernel_node, DILATION); + if (stride_ori.size() != 4 || stride_ori[2] != stride_ori[3]) { + MS_LOG(EXCEPTION) << "conv2d only support equal stride, and stride must be 4d!"; + } + if (stride_ori[0] != 1 || stride_ori[1] != 1) { + MS_LOG(EXCEPTION) << "conv2d stride only support 1 in N axis and C axis!"; + } + if (dilation_ori.size() != 4 || dilation_ori[2] != 1 || dilation_ori[3] != 1) { + MS_LOG(EXCEPTION) << "conv2d dilation only support 1, and dilation must be 4d!"; + } + if (dilation_ori[0] != 1 || dilation_ori[1] != 1) { + MS_LOG(EXCEPTION) << "conv2d dilation only support 1 in N axis and C axis!"; + } + int stride = stride_ori[2]; + int dilation = dilation_ori[2]; dnnl::memory::dims strides{stride, stride}; dnnl::memory::dims dilates{dilation - 1, dilation - 1}; diff --git a/mindspore/ccsrc/device/cpu/kernel/mkldnn/conv2d_grad_filter_cpu_kernel.cc b/mindspore/ccsrc/device/cpu/kernel/mkldnn/conv2d_grad_filter_cpu_kernel.cc index f4c0e58350ec98b276842bfd3af47e603e3ae433..1a7c10a531aaf2e51cf434fb06baf9450f8efe91 100644 --- a/mindspore/ccsrc/device/cpu/kernel/mkldnn/conv2d_grad_filter_cpu_kernel.cc +++ b/mindspore/ccsrc/device/cpu/kernel/mkldnn/conv2d_grad_filter_cpu_kernel.cc @@ -35,8 +35,19 @@ void Conv2dGradFilterCPUKernel::InitKernel(const CNodePtr &kernel_node) { dnnl::memory::desc dst_desc = GetDefaultMemDesc(dst_shape); int kernel_size = SizeToInt(weight_shape[3]); - int stride = AnfAlgo::GetNodeAttr(kernel_node, STRIDE); - int dilation = AnfAlgo::GetNodeAttr(kernel_node, DILATION); + auto stride_ori = AnfAlgo::GetNodeAttr>(kernel_node, STRIDE); + auto dilation_ori = AnfAlgo::GetNodeAttr>(kernel_node, DILATION); + if (stride_ori.size() != 2 || stride_ori[0] != stride_ori[1]) { + MS_LOG(EXCEPTION) << "Conv2dGradFilterCPUKernel only support equal stride, and stride must be 2d!"; + } + if (dilation_ori.size() != 4 || dilation_ori[2] != 1 || dilation_ori[3] != 1) { + MS_LOG(EXCEPTION) << "Conv2dGradFilterCPUKernel dilation only support 1, and dilation must be 4d!"; + } + if (dilation_ori[0] != 1 || dilation_ori[1] != 1) { + MS_LOG(EXCEPTION) << "Conv2dGradFilterCPUKernel dilation only support 1 in N axis and C axis!"; + } + int stride = stride_ori[0]; + int dilation = dilation_ori[2]; dnnl::memory::dims strides{stride, stride}; dnnl::memory::dims dilates{dilation - 1, dilation - 1}; diff --git a/mindspore/ccsrc/device/cpu/kernel/mkldnn/conv2d_grad_input_cpu_kernel.cc b/mindspore/ccsrc/device/cpu/kernel/mkldnn/conv2d_grad_input_cpu_kernel.cc index 492e2d6280fac3dcff615019a9a3d732c20b5a51..04dda20acdaf72153e562b49059fb5de4db74658 100644 --- a/mindspore/ccsrc/device/cpu/kernel/mkldnn/conv2d_grad_input_cpu_kernel.cc +++ b/mindspore/ccsrc/device/cpu/kernel/mkldnn/conv2d_grad_input_cpu_kernel.cc @@ -35,8 +35,19 @@ void Conv2dGradInputCPUKernel::InitKernel(const CNodePtr &kernel_node) { dnnl::memory::desc dst_desc = GetDefaultMemDesc(dst_shape); int kernel_size = SizeToInt(weight_shape[3]); - int stride = AnfAlgo::GetNodeAttr(kernel_node, STRIDE); - int dilation = AnfAlgo::GetNodeAttr(kernel_node, DILATION); + auto stride_ori = AnfAlgo::GetNodeAttr>(kernel_node, STRIDE); + auto dilation_ori = AnfAlgo::GetNodeAttr>(kernel_node, DILATION); + if (stride_ori.size() != 2 || stride_ori[0] != stride_ori[1]) { + MS_LOG(EXCEPTION) << "Conv2dGradInputCPUKernel only support equal stride, and stride must be 2d!"; + } + if (dilation_ori.size() != 4 || dilation_ori[2] != 1 || dilation_ori[3] != 1) { + MS_LOG(EXCEPTION) << "Conv2dGradInputCPUKernel dilation only support 1, and dilation must be 4d!"; + } + if (dilation_ori[0] != 1 || dilation_ori[1] != 1) { + MS_LOG(EXCEPTION) << "Conv2dGradInputCPUKernel dilation only support 1 in N axis and C axis!"; + } + int stride = stride_ori[0]; + int dilation = dilation_ori[2]; dnnl::memory::dims strides{stride, stride}; dnnl::memory::dims dilates{dilation - 1, dilation - 1}; std::vector int_padding_l; diff --git a/mindspore/ccsrc/kernel/gpu/nn/conv2d_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/conv2d_gpu_kernel.h index 7a4adff970c86d85087af2e8a4cb5c5f254a20c5..75b2a97cf85c7a12b55dbab90eb80f229af6407e 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/conv2d_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/conv2d_gpu_kernel.h @@ -113,9 +113,24 @@ class Conv2dGpuFwdKernel : public GpuKernel { CHECK_CUDNN_RET_WITH_EXCEPT(cudnnSetConvolutionGroupCount(conv_desc_, group_), "cudnnSetConvGroupCount failed"); pad_height_ = GetAttr(kernel_node, "pad"); pad_width_ = pad_height_; - stride_ = GetAttr(kernel_node, "stride"); - dilation_ = GetAttr(kernel_node, "dilation"); pad_mode_ = GetAttr(kernel_node, "pad_mode"); + auto stride_ori = AnfAlgo::GetNodeAttr>(kernel_node, "stride"); + auto dilation_ori = AnfAlgo::GetNodeAttr>(kernel_node, "dilation"); + if (stride_ori.size() != 4 || stride_ori[2] != stride_ori[3]) { + MS_LOG(EXCEPTION) << "conv2d only support equal stride, and stride must be 4d!"; + } + if (stride_ori[0] != 1 || stride_ori[1] != 1) { + MS_LOG(EXCEPTION) << "conv2d stride only support 1 in N axis and C axis!"; + } + if (dilation_ori.size() != 4 || dilation_ori[2] != dilation_ori[3]) { + MS_LOG(EXCEPTION) << "conv2d only support equal dilation, and dilation must be 4d!"; + } + if (dilation_ori[0] != 1 || dilation_ori[1] != 1) { + MS_LOG(EXCEPTION) << "conv2d dilation only support 1 in N axis and C axis!"; + } + stride_ = stride_ori[2]; + dilation_ = dilation_ori[2]; + cudnnTensorDescriptor_t input_descriptor_real = nullptr; if (pad_mode_ == kSamePadModeUpperCase || pad_mode_ == kSamePadModeLowerCase) { SetPad(in_shape, kernel_node); diff --git a/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_filter_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_filter_gpu_kernel.h index f8afad4f847617ca295492343fb2de34c5fbea13..e481fd448e926e1eb71b5e052a37cb56988ad11e 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_filter_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_filter_gpu_kernel.h @@ -116,9 +116,20 @@ class ConvGradFilterGpuBkwKernel : public GpuKernel { pad_height_ = GetAttr(kernel_node, "pad"); pad_width_ = pad_height_; - stride_ = GetAttr(kernel_node, "stride"); - dilation_ = GetAttr(kernel_node, "dilation"); pad_mode_ = GetAttr(kernel_node, "pad_mode"); + auto stride_ori = AnfAlgo::GetNodeAttr>(kernel_node, "stride"); + auto dilation_ori = AnfAlgo::GetNodeAttr>(kernel_node, "dilation"); + if (stride_ori.size() != 2 || stride_ori[0] != stride_ori[1]) { + MS_LOG(EXCEPTION) << "ConvGradFilterGpuBkwKernel only support equal stride, and stride must be 2d!"; + } + if (dilation_ori.size() != 4 || dilation_ori[2] != dilation_ori[3]) { + MS_LOG(EXCEPTION) << "ConvGradFilterGpuBkwKernel only support equal dilation, and dilation must be 4d!"; + } + if (dilation_ori[0] != 1 || dilation_ori[1] != 1) { + MS_LOG(EXCEPTION) << "ConvGradFilterGpuBkwKernel dilation only support 1 in N axis and C axis!"; + } + stride_ = stride_ori[0]; + dilation_ = dilation_ori[2]; cudnnTensorDescriptor_t x_desc_real = nullptr; if (pad_mode_ == kSamePadModeUpperCase || pad_mode_ == kSamePadModeLowerCase) { SetPad(in_shape, kernel_node); diff --git a/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_input_gpu_kernel.h b/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_input_gpu_kernel.h index be7739981da9afb475c21488d73387eb905aa04d..008abcc65859f6bf5e99c10df0e5446fd84f64df 100644 --- a/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_input_gpu_kernel.h +++ b/mindspore/ccsrc/kernel/gpu/nn/conv2d_grad_input_gpu_kernel.h @@ -117,9 +117,20 @@ class ConvGradInputGpuBkwKernel : public GpuKernel { pad_height_ = GetAttr(kernel_node, "pad"); pad_width_ = pad_height_; - stride_ = GetAttr(kernel_node, "stride"); - dilation_ = GetAttr(kernel_node, "dilation"); pad_mode_ = GetAttr(kernel_node, "pad_mode"); + auto stride_ori = AnfAlgo::GetNodeAttr>(kernel_node, "stride"); + auto dilation_ori = AnfAlgo::GetNodeAttr>(kernel_node, "dilation"); + if (stride_ori.size() != 2 || stride_ori[0] != stride_ori[1]) { + MS_LOG(EXCEPTION) << "ConvGradInputGpuBkwKernel only support equal stride, and stride must be 2d!"; + } + if (dilation_ori.size() != 4 || dilation_ori[2] != dilation_ori[3]) { + MS_LOG(EXCEPTION) << "ConvGradInputGpuBkwKernel only support equal dilation, and dilation must be 4d!"; + } + if (dilation_ori[0] != 1 || dilation_ori[1] != 1) { + MS_LOG(EXCEPTION) << "ConvGradInputGpuBkwKernel dilation only support 1 in N axis and C axis!"; + } + stride_ = stride_ori[0]; + dilation_ = dilation_ori[2]; cudnnTensorDescriptor_t dx_desc_real = nullptr; if (pad_mode_ == kSamePadModeUpperCase || pad_mode_ == kSamePadModeLowerCase) { SetPad(input_shape, kernel_node); diff --git a/mindspore/ccsrc/kernel/tbe/tbe_adapter.cc b/mindspore/ccsrc/kernel/tbe/tbe_adapter.cc index 229a3eb34a72e40867d5f9c85c8941345ece6eb0..481ac755047d5b24704bf423cf2bf71bf8c1f1d1 100644 --- a/mindspore/ccsrc/kernel/tbe/tbe_adapter.cc +++ b/mindspore/ccsrc/kernel/tbe/tbe_adapter.cc @@ -148,9 +148,6 @@ void TbeAdapter::InputOrderPass(const std::string &op_name, std::vector TbeAdapter::build_json_attr_pass_map_ = { - {"Conv2D", TbeAdapter::Conv2DAttrJsonPass}, - {"Conv2DBackpropFilter", TbeAdapter::Conv2DBackpropFilterAttrJsonPass}, - {"Conv2DBackpropInput", TbeAdapter::Conv2DBackpropInputAttrJsonPass}, {"MaximumGrad", TbeAdapter::MaximumGradAttrJsonPass}, {"MinimumGrad", TbeAdapter::MinimumGradAttrJsonPass}, {"Cast", TbeAdapter::CastAttrJsonPass}}; @@ -168,135 +165,6 @@ bool TbeAdapter::RunAttrPass(const mindspore::AnfNodePtr &anf_node, return false; } -void TbeAdapter::Conv2DAttrJsonPass(const mindspore::AnfNodePtr &anf_node, - const std::vector> &op_info_attrs, - nlohmann::json *attrs_json) { - MS_EXCEPTION_IF_NULL(anf_node); - MS_EXCEPTION_IF_NULL(attrs_json); - auto attr_num = op_info_attrs.size(); - auto primitive = AnfAlgo::GetCNodePrimitive(anf_node); - MS_EXCEPTION_IF_NULL(primitive); - for (size_t i = 0; i < attr_num; i++) { - nlohmann::json attr_obj; - MS_EXCEPTION_IF_NULL(op_info_attrs[i]); - std::string attr_name = op_info_attrs[i]->name(); - std::vector attr_value; - if (primitive->GetAttr(attr_name) != nullptr) { - auto value = primitive->GetAttr(attr_name); - int data = GetValue(value); - size_t list_int_size = 0; - if (attr_name == "stride") { - list_int_size = 4; - } else if (attr_name == "dilation") { - list_int_size = 4; - } else if (attr_name == "pad") { - value = primitive->GetAttr("pad_list"); - attr_value = GetValue>(value); - } - for (size_t j = 0; j < list_int_size; j++) { - attr_value.push_back(data); - } - attr_obj["value"] = attr_value; - } else { - attr_obj["value"] = 0; - } - attr_obj["name"] = attr_name; - attr_obj["valid"] = true; - (*attrs_json).push_back(attr_obj); - } - MS_LOG(INFO) << "Conv2DAttrPass done."; -} - -void TbeAdapter::Conv2DBackpropFilterAttrJsonPass( - const mindspore::AnfNodePtr &anf_node, const std::vector> &op_info_attrs, - nlohmann::json *attrs_json) { - MS_EXCEPTION_IF_NULL(anf_node); - MS_EXCEPTION_IF_NULL(attrs_json); - auto attr_num = op_info_attrs.size(); - auto primitive = AnfAlgo::GetCNodePrimitive(anf_node); - MS_EXCEPTION_IF_NULL(primitive); - for (size_t i = 0; i < attr_num; i++) { - nlohmann::json attr_obj; - MS_EXCEPTION_IF_NULL(op_info_attrs[i]); - std::string attr_name = op_info_attrs[i]->name(); - if (primitive->GetAttr(attr_name) != nullptr) { - auto value = primitive->GetAttr(attr_name); - if (attr_name == "pad_mode") { - std::string attr_value = GetValue(value); - (void)transform(attr_value.begin(), attr_value.end(), attr_value.begin(), ::toupper); - attr_obj["value"] = attr_value; - } else if (attr_name == "filter_sizes") { - std::vector attr_value = GetValue>(value); - attr_obj["value"] = attr_value; - } else { - std::vector attr_value; - int data = GetValue(value); - size_t list_int_size = 0; - if (attr_name == "stride") { - list_int_size = 2; - } else if (attr_name == "dilation") { - list_int_size = 4; - } - for (size_t j = 0; j < list_int_size; j++) { - attr_value.push_back(data); - } - attr_obj["value"] = attr_value; - } - attr_obj["valid"] = true; - } else { - attr_obj["valid"] = false; - } - attr_obj["name"] = attr_name; - attrs_json->push_back(attr_obj); - } - MS_LOG(INFO) << "Conv2DBackpropFilterAttrJsonPass done."; -} - -void TbeAdapter::Conv2DBackpropInputAttrJsonPass( - const mindspore::AnfNodePtr &anf_node, const std::vector> &op_info_attrs, - nlohmann::json *attrs_json) { - MS_EXCEPTION_IF_NULL(anf_node); - MS_EXCEPTION_IF_NULL(attrs_json); - auto attr_num = op_info_attrs.size(); - auto primitive = AnfAlgo::GetCNodePrimitive(anf_node); - MS_EXCEPTION_IF_NULL(primitive); - for (size_t i = 0; i < attr_num; i++) { - nlohmann::json attr_obj; - MS_EXCEPTION_IF_NULL(op_info_attrs[i]); - std::string attr_name = op_info_attrs[i]->name(); - if (primitive->GetAttr(attr_name) != nullptr) { - auto value = primitive->GetAttr(attr_name); - if (attr_name == "pad_mode") { - std::string attr_value = GetValue(value); - (void)transform(attr_value.begin(), attr_value.end(), attr_value.begin(), ::toupper); - attr_obj["value"] = attr_value; - } else if (attr_name == "input_sizes") { - std::vector attr_value = GetValue>(value); - attr_obj["value"] = attr_value; - } else { - std::vector attr_value; - int data = GetValue(value); - size_t list_int_size = 0; - if (attr_name == "stride") { - list_int_size = 2; - } else if (attr_name == "dilation") { - list_int_size = 4; - } - for (size_t j = 0; j < list_int_size; j++) { - attr_value.push_back(data); - } - attr_obj["value"] = attr_value; - } - attr_obj["valid"] = true; - } else { - attr_obj["valid"] = false; - } - attr_obj["name"] = attr_name; - attrs_json->push_back(attr_obj); - } - MS_LOG(INFO) << "Conv2DBackpropInputAttrJsonPass done."; -} - void TbeAdapter::MaximumGradAttrJsonPass(const mindspore::AnfNodePtr &anf_node, const std::vector> &op_info_attrs, nlohmann::json *attrs_json) { diff --git a/mindspore/ccsrc/onnx/onnx_exporter.cc b/mindspore/ccsrc/onnx/onnx_exporter.cc index 3bd4a38881da6e55bb90b3ee01fb3662508c7db8..4767d13cd958445f5c73a7bf29f7f748eafd6445 100644 --- a/mindspore/ccsrc/onnx/onnx_exporter.cc +++ b/mindspore/ccsrc/onnx/onnx_exporter.cc @@ -179,7 +179,7 @@ OPERATOR_ONNX_CONVERT_DEFINE(Squeeze, Squeeze, OPERATOR_ONNX_CONVERT_DEFINE( Conv2D, Conv, OpNameInfo() - .Attr("dilation", "dilations", onnx::AttributeProto_AttributeType_INTS, SetAttrValueToProto) + .Attr("dilation", "dilations", onnx::AttributeProto_AttributeType_INTS, SetAttrTupleValueToProto<2>) .Attr("group", "group", onnx::AttributeProto_AttributeType_INT, SetAttrValueToProto) .Attr("kernel_size", "kernel_shape", onnx::AttributeProto_AttributeType_INTS, SetAttrTupleValueToProto) .Attr("pad_mode", "auto_pad", onnx::AttributeProto_AttributeType_STRING, @@ -197,8 +197,7 @@ OPERATOR_ONNX_CONVERT_DEFINE( prim); } }) - .Attr("stride", "strides", onnx::AttributeProto_AttributeType_INTS, SetAttrValueToProto)) - + .Attr("stride", "strides", onnx::AttributeProto_AttributeType_INTS, SetAttrTupleValueToProto<2>)) OPERATOR_ONNX_CONVERT_DEFINE(BiasAdd, Add, OpNameInfo()) OPERATOR_ONNX_CONVERT_DEFINE(MatMul, Gemm, OpNameInfo() diff --git a/mindspore/ccsrc/transform/op_declare.cc b/mindspore/ccsrc/transform/op_declare.cc index 419805c37fd6f204b402998a8cda609d18cf0493..76348799592b158e301f6c40fa125628ec4956b3 100644 --- a/mindspore/ccsrc/transform/op_declare.cc +++ b/mindspore/ccsrc/transform/op_declare.cc @@ -754,9 +754,9 @@ OUTPUT_MAP(MaxPoolGradWithArgmax) = {{0, OUTPUT_DESC(y)}}; // Conv2D INPUT_MAP(Conv2D) = {{1, INPUT_DESC(x)}, {2, INPUT_DESC(filter)}}; ATTR_MAP(Conv2D) = { - {"stride", ATTR_DESC(strides, "pad", AnyTraits>())}, + {"stride", ATTR_DESC(strides, AnyTraits>(), AnyTraits>())}, {"pad_list", ATTR_DESC(pads, AnyTraits>(), AnyTraits>())}, - {"dilation", ATTR_DESC(dilations, "pad", AnyTraits>())}, + {"dilation", ATTR_DESC(dilations, AnyTraits>(), AnyTraits>())}, }; OUTPUT_MAP(Conv2D) = {{0, OUTPUT_DESC(y)}}; @@ -766,8 +766,8 @@ INPUT_ATTR_MAP(Conv2DBackpropInputD) = { {3, ATTR_DESC(input_sizes, AnyTraits>(), AnyTraits>())}}; ATTR_MAP(Conv2DBackpropInputD) = { {"pad_list", ATTR_DESC(pads, AnyTraits>(), AnyTraits>())}, - {"stride", ATTR_DESC(strides, "strides", AnyTraits>())}, - {"dilation", ATTR_DESC(dilations, "pad", AnyTraits>())}, + {"stride", ATTR_DESC(strides, AnyTraits>(), AnyTraits>())}, + {"dilation", ATTR_DESC(dilations, AnyTraits>(), AnyTraits>())}, }; OUTPUT_MAP(Conv2DBackpropInputD) = {{0, OUTPUT_DESC(y)}}; @@ -777,17 +777,17 @@ INPUT_ATTR_MAP(Conv2DBackpropFilterD) = { {3, ATTR_DESC(filter_sizes, AnyTraits>(), AnyTraits>())}}; ATTR_MAP(Conv2DBackpropFilterD) = { {"pad_list", ATTR_DESC(pads, AnyTraits>(), AnyTraits>())}, - {"stride", ATTR_DESC(strides, "strides", AnyTraits>())}, - {"dilation", ATTR_DESC(dilations, "pad", AnyTraits>())}, + {"stride", ATTR_DESC(strides, AnyTraits>(), AnyTraits>())}, + {"dilation", ATTR_DESC(dilations, AnyTraits>(), AnyTraits>())}, }; OUTPUT_MAP(Conv2DBackpropFilterD) = {{0, OUTPUT_DESC(y)}}; // DepthwiseConv2D INPUT_MAP(DepthwiseConv2D) = {{1, INPUT_DESC(x)}, {2, INPUT_DESC(filter)}}; ATTR_MAP(DepthwiseConv2D) = { - {"stride", ATTR_DESC(strides, "pad", AnyTraits>())}, + {"stride", ATTR_DESC(strides, AnyTraits>(), AnyTraits>())}, {"pads", ATTR_DESC(pads, AnyTraits>(), AnyTraits>())}, - {"dilation", ATTR_DESC(dilations, "pad", AnyTraits>())}, + {"dilation", ATTR_DESC(dilations, AnyTraits>(), AnyTraits>())}, {"data_format", ATTR_DESC(data_format, AnyTraits())}, }; OUTPUT_MAP(DepthwiseConv2D) = {{0, OUTPUT_DESC(y)}}; @@ -797,9 +797,9 @@ INPUT_MAP(DepthwiseConv2DBackpropInputD) = {{2, INPUT_DESC(filter)}, {3, INPUT_D INPUT_ATTR_MAP(DepthwiseConv2DBackpropInputD) = { {1, ATTR_DESC(input_size, AnyTraits>(), AnyTraits>())}}; ATTR_MAP(DepthwiseConv2DBackpropInputD) = { - {"stride", ATTR_DESC(strides, "pad", AnyTraits>())}, + {"stride", ATTR_DESC(strides, AnyTraits>(), AnyTraits>())}, {"pads", ATTR_DESC(pads, AnyTraits>(), AnyTraits>())}, - {"dilation", ATTR_DESC(dilations, "pad", AnyTraits>())}, + {"dilation", ATTR_DESC(dilations, AnyTraits>(), AnyTraits>())}, }; OUTPUT_MAP(DepthwiseConv2DBackpropInputD) = {{0, OUTPUT_DESC(input_grad)}}; @@ -808,9 +808,9 @@ INPUT_MAP(DepthwiseConv2DBackpropFilterD) = {{1, INPUT_DESC(input)}, {3, INPUT_D INPUT_ATTR_MAP(DepthwiseConv2DBackpropFilterD) = { {2, ATTR_DESC(filter_size, AnyTraits>(), AnyTraits>())}}; ATTR_MAP(DepthwiseConv2DBackpropFilterD) = { - {"stride", ATTR_DESC(strides, "pad", AnyTraits>())}, + {"stride", ATTR_DESC(strides, AnyTraits>(), AnyTraits>())}, {"pads", ATTR_DESC(pads, AnyTraits>(), AnyTraits>())}, - {"dilation", ATTR_DESC(dilations, "pad", AnyTraits>())}, + {"dilation", ATTR_DESC(dilations, AnyTraits>(), AnyTraits>())}, }; OUTPUT_MAP(DepthwiseConv2DBackpropFilterD) = {{0, OUTPUT_DESC(filter_grad)}}; diff --git a/mindspore/nn/layer/conv.py b/mindspore/nn/layer/conv.py index eb73a9ce5afb3c4ca14ae8cd3ab3344958394f32..fbf6ad2a0c063e033dfafb3739b544ec17e7dafe 100644 --- a/mindspore/nn/layer/conv.py +++ b/mindspore/nn/layer/conv.py @@ -17,7 +17,7 @@ from mindspore import log as logger from mindspore.ops import operations as P from mindspore.common.parameter import Parameter from mindspore.common.initializer import initializer -from mindspore._checkparam import check_bool, twice, check_int_positive, check_int_non_negative, check_int +from mindspore._checkparam import check_bool, twice, check_int_positive, check_int_non_negative from mindspore._extends import cell_attr_register from ..cell import Cell @@ -42,17 +42,23 @@ class _Conv(Cell): self.in_channels = check_int_positive(in_channels) self.out_channels = check_int_positive(out_channels) self.kernel_size = kernel_size - self.stride = check_int_positive(stride) + self.stride = stride self.pad_mode = pad_mode self.padding = check_int_non_negative(padding) - self.dilation = check_int(dilation) + self.dilation = dilation self.group = check_int_positive(group) self.has_bias = has_bias - if (not isinstance(kernel_size, tuple)) or len(kernel_size) != 2 or \ - (not isinstance(kernel_size[0], int)) or (not isinstance(kernel_size[1], int)) or \ - kernel_size[0] < 1 or kernel_size[1] < 1: + if (not isinstance(kernel_size[0], int)) or (not isinstance(kernel_size[1], int)) or \ + kernel_size[0] < 1 or kernel_size[1] < 1: raise ValueError("Attr 'kernel_size' of 'Conv2D' Op passed " + str(self.kernel_size) + ", should be a int or tuple and equal to or greater than 1.") + if (not isinstance(stride[0], int)) or (not isinstance(stride[1], int)) or stride[0] < 1 or stride[1] < 1: + raise ValueError("Attr 'stride' of 'Conv2D' Op passed " + + str(self.stride) + ", should be a int or tuple and equal to or greater than 1.") + if (not isinstance(dilation[0], int)) or (not isinstance(dilation[1], int)) or \ + dilation[0] < 1 or dilation[1] < 1: + raise ValueError("Attr 'dilation' of 'Conv2D' Op passed " + + str(self.dilation) + ", should equal to or greater than 1.") if in_channels % group != 0: raise ValueError("Attr 'in_channels' of 'Conv2D' Op must be divisible by " "attr 'group' of 'Conv2D' Op.") @@ -107,12 +113,13 @@ class Conv2d(_Conv): Args: in_channels (int): The number of input channel :math:`C_{in}`. out_channels (int): The number of output channel :math:`C_{out}`. - kernel_size (Union[int, tuple]): The data type is int or tuple with 2 integers. Specifies the height + kernel_size (Union[int, tuple[int]]): The data type is int or tuple with 2 integers. Specifies the height and width of the 2D convolution window. Single int means the value if for both height and width of the kernel. A tuple of 2 ints means the first value is for the height and the other is for the width of the kernel. - stride (int): Specifies stride for all spatial dimensions with the same value. Value of stride should be - greater or equal to 1 but bounded by the height and width of the input. Default: 1. + stride (Union[int, tuple[int]]): The distance of kernel moving, an int number that represents + the height and width of movement are both strides, or a tuple of two int numbers that + represent height and width of movement respectively. Default: 1. pad_mode (str): Specifies padding mode. The optional values are "same", "valid", "pad". Default: "same". @@ -130,9 +137,11 @@ class Conv2d(_Conv): Tensor borders. `padding` should be greater than or equal to 0. padding (int): Implicit paddings on both sides of the input. Default: 0. - dilation (int): Specifying the dilation rate to use for dilated convolution. If set to be :math:`k > 1`, - there will be :math:`k - 1` pixels skipped for each sampling location. Its value should be greater - or equal to 1 and bounded by the height and width of the input. Default: 1. + dilation (Union[int, tuple[int]]): The data type is int or tuple with 2 integers. Specifies the dilation rate + to use for dilated convolution. If set to be :math:`k > 1`, there will + be :math:`k - 1` pixels skipped for each sampling location. Its value should + be greater or equal to 1 and bounded by the height and width of the + input. Default: 1. group (int): Split filter into groups, `in_ channels` and `out_channels` should be divisible by the number of groups. Default: 1. has_bias (bool): Specifies whether the layer uses a bias vector. Default: False. @@ -172,6 +181,8 @@ class Conv2d(_Conv): weight_init='normal', bias_init='zeros'): kernel_size = twice(kernel_size) + stride = twice(stride) + dilation = twice(dilation) super(Conv2d, self).__init__( in_channels, out_channels, @@ -241,7 +252,9 @@ class Conv2dTranspose(_Conv): and width of the 2D convolution window. Single int means the value is for both height and width of the kernel. A tuple of 2 ints means the first value is for the height and the other is for the width of the kernel. - stride (int): Specifies the same value for all spatial dimensions. Default: 1. + stride (Union[int, tuple[int]]): The distance of kernel moving, an int number that represents + the height and width of movement are both strides, or a tuple of two int numbers that + represent height and width of movement respectively. Default: 1. pad_mode (str): Select the mode of the pad. The optional values are "pad", "same", "valid". Default: "same". @@ -251,8 +264,11 @@ class Conv2dTranspose(_Conv): - valid: Adopted the way of discarding. padding (int): Implicit paddings on both sides of the input. Default: 0. - dilation (int): Specifies the dilation rate to use for dilated - convolution. Default: 1. + dilation (Union[int, tuple[int]]): The data type is int or tuple with 2 integers. Specifies the dilation rate + to use for dilated convolution. If set to be :math:`k > 1`, there will + be :math:`k - 1` pixels skipped for each sampling location. Its value should + be greater or equal to 1 and bounded by the height and width of the + input. Default: 1. group (int): Split filter into groups, `in_channels` and `out_channels` should be divisible by the number of groups. Default: 1. has_bias (bool): Specifies whether the layer uses a bias vector. Default: False. @@ -290,6 +306,8 @@ class Conv2dTranspose(_Conv): weight_init='normal', bias_init='zeros'): kernel_size = twice(kernel_size) + stride = twice(stride) + dilation = twice(dilation) # out_channels and in_channels swap. # cause Conv2DBackpropInput's out_channel refers to Conv2D's out_channel, # then Conv2dTranspose's out_channel refers to Conv2DBackpropInput's in_channel. @@ -333,26 +351,26 @@ class Conv2dTranspose(_Conv): self.conv2d_transpose.set_strategy(strategy) return self - def _deconv_output_length(self, input_length, filter_size): + def _deconv_output_length(self, input_length, filter_size, stride_size, dilation_size): """Calculate the width and height of output.""" length = 0 if self.is_valid: - if filter_size - self.stride > 0: - length = input_length * self.stride + filter_size - self.stride + if filter_size - stride_size > 0: + length = input_length * stride_size + filter_size - stride_size else: - length = input_length * self.stride + length = input_length * stride_size elif self.is_same: - length = input_length * self.stride + length = input_length * stride_size elif self.is_pad: - length = input_length * self.stride - 2 * self.padding + filter_size + \ - (filter_size - 1) * (self.dilation - 1) - self.stride + length = input_length * stride_size - 2 * self.padding + filter_size + \ + (filter_size - 1) * (dilation_size - 1) - stride_size return length def construct(self, x): n, _, h, w = self.shape(x) - h_out = self._deconv_output_length(h, self.kernel_size[0]) - w_out = self._deconv_output_length(w, self.kernel_size[1]) + h_out = self._deconv_output_length(h, self.kernel_size[0], self.stride[0], self.dilation[0]) + w_out = self._deconv_output_length(w, self.kernel_size[1], self.stride[1], self.dilation[1]) if self.has_bias: return self.bias_add(self.conv2d_transpose(x, self.weight, (n, self.out_channels, h_out, w_out)), self.bias) diff --git a/mindspore/ops/_op_impl/tbe/conv2d.py b/mindspore/ops/_op_impl/tbe/conv2d.py index 52a9eac1fad46f2d0abec98655bb271397c08336..da0f34e35921dc03fc51efe0c675cd192880e790 100644 --- a/mindspore/ops/_op_impl/tbe/conv2d.py +++ b/mindspore/ops/_op_impl/tbe/conv2d.py @@ -34,7 +34,7 @@ from mindspore.ops.op_info_register import op_info_register "value": "all" }, { - "name": "pad", + "name": "pad_list", "param_type": "required", "type": "listInt", "value": "all" diff --git a/mindspore/ops/operations/_grad_ops.py b/mindspore/ops/operations/_grad_ops.py index f38044ab6ac07817a8a9ffb3c72991462c965d9d..e03b163225072d99422225b45c0a99875a0b0b4a 100644 --- a/mindspore/ops/operations/_grad_ops.py +++ b/mindspore/ops/operations/_grad_ops.py @@ -119,8 +119,8 @@ class Conv2DBackpropFilter(PrimitiveWithInfer): pad (int): The pad value to fill. Default: 0. mode (int): 0 Math convolutiuon, 1 cross-correlation convolution , 2 deconvolution, 3 depthwise convolution. Default: 1. - stride (int): The stride to apply conv filter. Default: 1. - dilation (int): Specifies the dilation rate to use for dilated convolution. Default: 1. + stride (tuple): The stride to apply conv filter. Default: (1, 1). + dilation (tuple): Specifies the dilation rate to use for dilated convolution. Default: (1, 1, 1, 1). group (int): Splits input into groups. Default: 1. Returns: @@ -135,8 +135,8 @@ class Conv2DBackpropFilter(PrimitiveWithInfer): pad=0, pad_list=(0, 0, 0, 0), mode=1, - stride=1, - dilation=1, + stride=(1, 1), + dilation=(1, 1, 1, 1), group=1): """init Convolution""" self.init_prim_io_names(inputs=['out_backprop', 'input', 'filter_sizes'], outputs=['output']) @@ -146,7 +146,9 @@ class Conv2DBackpropFilter(PrimitiveWithInfer): pad_mode = pad_mode.upper() self.add_prim_attr('pad_mode', pad_mode) self.pad = pad - self.stride = stride + if isinstance(stride, tuple) and len(stride) == 4: + self.stride = (stride[2], stride[3]) + self.add_prim_attr('stride', self.stride) self.dilation = dilation self.group = group self.add_prim_attr('data_format', "NCHW") diff --git a/mindspore/ops/operations/nn_ops.py b/mindspore/ops/operations/nn_ops.py index 83f76455e02b625b5b64acdd5f97e8910b64a325..a7d7357e4688793105ed510044a4492089ef2659 100644 --- a/mindspore/ops/operations/nn_ops.py +++ b/mindspore/ops/operations/nn_ops.py @@ -460,8 +460,8 @@ class Conv2D(PrimitiveWithInfer): 2 deconvolution, 3 depthwise convolution. Default: 1. pad_mode (str): "valid", "same", "pad" the mode to fill padding. Default: "valid". pad (int): The pad value to fill. Default: 0. - stride (int): The stride to apply conv filter. Default: 1. - dilation (int): Specify the space to use between kernel elements. Default: 1. + stride (Union(int, tuple[int])): The stride to apply conv filter. Default: 1. + dilation (Union(int, tuple[int])): Specify the space to use between kernel elements. Default: 1. group (int): Split input into groups. Default: 1. Returns: @@ -488,11 +488,35 @@ class Conv2D(PrimitiveWithInfer): group=1): """init Conv2D""" self.init_prim_io_names(inputs=['x', 'w'], outputs=['output']) - self.kernel_size = kernel_size self.kernel_size = validator.check_type('kernel_size', kernel_size, (int, tuple)) - if isinstance(self.kernel_size, int): - self.kernel_size = (self.kernel_size, self.kernel_size) - validator.check_integer('length of kernel_size', len(self.kernel_size), 2, Rel.GE) + if isinstance(kernel_size, int): + self.kernel_size = (kernel_size, kernel_size) + if len(self.kernel_size) != 2 or (not isinstance(self.kernel_size[0], int)) or \ + (not isinstance(self.kernel_size[1], int)) or \ + self.kernel_size[0] < 1 or self.kernel_size[1] < 1: + raise ValueError(f"The \'kernel_size\' of \'Conv2D\' should be an positive int number or " + f"a tuple of two positive int numbers, but got {kernel_size}") + self.stride = validator.check_type('stride', stride, (int, tuple)) + if isinstance(stride, int): + self.stride = (stride, stride) + if len(self.stride) != 2 or (not isinstance(self.stride[0], int)) or \ + (not isinstance(self.stride[1], int)) or \ + self.stride[0] < 1 or self.stride[1] < 1: + raise ValueError(f"The \'stride\' of \'Conv2D\' should be an positive int number or " + f"a tuple of two positive int numbers, but got {stride}") + self.add_prim_attr('stride', (1, 1, self.stride[0], self.stride[1])) + self.dilation = validator.check_type('dilation', dilation, (tuple, int)) + if isinstance(dilation, int): + self.dilation = (1, 1, dilation, dilation) + elif len(dilation) == 2: + self.dilation = (1, 1, dilation[0], dilation[1]) + if len(self.dilation) != 4 or (not isinstance(self.dilation[0], int) or self.dilation[0] < 1) or \ + (not isinstance(self.dilation[1], int) or self.dilation[1] < 1) or \ + (not isinstance(self.dilation[2], int) or self.dilation[2] < 1) or \ + (not isinstance(self.dilation[3], int) or self.dilation[3] < 1): + raise ValueError(f"The \'dilation\' of \'Conv2D\' should be an positive int number or " + f"a tuple of two or four positive int numbers, but got {dilation}") + self.add_prim_attr('dilation', self.dilation) validator.equal('type of pad', type(pad), 'not bool', not isinstance(pad, bool)) validator.equal('type of pad', type(pad), 'int', isinstance(pad, int)) self.pad_mode = validator.check_string('pad_mode', pad_mode, ['valid', 'same', 'pad']) @@ -504,18 +528,6 @@ class Conv2D(PrimitiveWithInfer): self.add_prim_attr('data_format', "NCHW") self.out_channel = validator.check_integer('out_channel', out_channel, 0, Rel.GT) self.group = validator.check_integer('group', group, 0, Rel.GT) - self.dilation = validator.check_integer('dilation', dilation, 1, Rel.GE) - validator.check_type('kernel_size', kernel_size, [int, tuple]) - if isinstance(kernel_size, int) and kernel_size < 1: - raise ValueError('Attr \'kernel_size\' of \'Conv2D\' Op passed ' - + str(self.kernel_size) + ', should be a int or tuple and equal to or greater than 1.') - if isinstance(kernel_size, tuple) and (len(kernel_size) != 2 or - (not isinstance(kernel_size[0], int)) or - (not isinstance(kernel_size[1], int)) or - kernel_size[0] < 1 or kernel_size[1] < 1): - raise ValueError('Attr \'kernel_size\' of \'Conv2D\' Op passed ' - + str(self.kernel_size) + ', should be a int or tuple and equal to or greater than 1.') - self.stride = validator.check_integer('stride', stride, 1, Rel.GE) def infer_shape(self, x_shape, w_shape): validator.check_integer("weight_shape", len(w_shape), 4, Rel.EQ) @@ -526,29 +538,33 @@ class Conv2D(PrimitiveWithInfer): kernel_size_h = w_shape[2] kernel_size_w = w_shape[3] + stride_h = self.stride[2] + stride_w = self.stride[3] + dilation_h = self.dilation[2] + dilation_w = self.dilation[3] if self.pad_mode == "valid": - h_out = math.ceil((x_shape[2] - self.dilation * (kernel_size_h - 1)) / self.stride) - w_out = math.ceil((x_shape[3] - self.dilation * (kernel_size_w - 1)) / self.stride) + h_out = math.ceil((x_shape[2] - dilation_h * (kernel_size_h - 1)) / stride_h) + w_out = math.ceil((x_shape[3] - dilation_w * (kernel_size_w - 1)) / stride_w) pad_top, pad_bottom, pad_left, pad_right = 0, 0, 0, 0 elif self.pad_mode == "same": - h_out = math.ceil(x_shape[2] / self.stride) - w_out = math.ceil(x_shape[3] / self.stride) + h_out = math.ceil(x_shape[2] / stride_h) + w_out = math.ceil(x_shape[3] / stride_w) - pad_needed_h = max(0, (h_out - 1) * self.stride + self.dilation * (kernel_size_h - 1) + 1 - x_shape[2]) + pad_needed_h = max(0, (h_out - 1) * stride_h + dilation_h * (kernel_size_h - 1) + 1 - x_shape[2]) pad_top = math.floor(pad_needed_h / 2) pad_bottom = pad_needed_h - pad_top - pad_needed_w = max(0, (w_out - 1) * self.stride + self.dilation * (kernel_size_w - 1) + 1 - x_shape[3]) + pad_needed_w = max(0, (w_out - 1) * stride_w + dilation_w * (kernel_size_w - 1) + 1 - x_shape[3]) pad_left = math.floor(pad_needed_w / 2) pad_right = pad_needed_w - pad_left elif self.pad_mode == 'pad': pad_top, pad_bottom, pad_left, pad_right = self.pad, self.pad, self.pad, self.pad - h_out = 1 + (x_shape[2] + 2 * self.pad - kernel_size_h - (kernel_size_h - 1) * (self.dilation - 1)) \ - / self.stride - w_out = 1 + (x_shape[3] + 2 * self.pad - kernel_size_w - (kernel_size_w - 1) * (self.dilation - 1)) \ - / self.stride + h_out = 1 + (x_shape[2] + 2 * self.pad - kernel_size_h - (kernel_size_h - 1) * (dilation_h - 1)) \ + / stride_h + w_out = 1 + (x_shape[3] + 2 * self.pad - kernel_size_w - (kernel_size_w - 1) * (dilation_w - 1)) \ + / stride_w h_out = math.floor(h_out) w_out = math.floor(w_out) @@ -580,19 +596,19 @@ class DepthwiseConv2dNative(PrimitiveWithInfer): Args: channel_multiplier (int): The multipiler for the original output conv. - kernel_size (int or tuple): The size of the conv kernel. + kernel_size (Union[int, tuple[int]]): The size of the conv kernel. mode (int): 0 Math convolution, 1 cross-correlation convolution , 2 deconvolution, 3 depthwise convolution. Default: 3. pad_mode (str): "valid", "same", "pad" the mode to fill padding. Default: "valid". pad (int): The pad value to fill. Default: 0. - stride (int): The stride to apply conv filter. Default: 1. - dilation (int): Specifies the dilation rate to use for dilated convolution. Default: 1. + stride (Union[int, tuple[int]]): The stride to apply conv filter. Default: 1. + dilation (Union[int, tuple[int]]): Specifies the dilation rate to use for dilated convolution. Default: 1. group (int): Splits input into groups. Default: 1. Inputs: - **input** (Tensor) - Tensor of shape :math:`(N, C_{in}, H_{in}, W_{in})`. - **weight** (Tensor) - Set size of kernel is :math:`(K_1, K_2)`, then the shape is - :math:`(C_{out}, C_{in}, K_1, K_2)`. + :math:`(channel_multiplier, C_{in}, K_1, K_2)`. Outputs: Tensor of shape :math:`(N, C_{in} * \text{channel_multiplier}, H_{out}, W_{out})`. @@ -610,15 +626,33 @@ class DepthwiseConv2dNative(PrimitiveWithInfer): group=1): """init DepthwiseConv2dNative""" validator.check_pad_value_by_mode(self.__class__.__name__, pad_mode, pad) - validator.check_type("kernel_size", kernel_size, (int, tuple)) + self.kernel_size = validator.check_type('kernel_size', kernel_size, (int, tuple)) if isinstance(kernel_size, int): - kernel_size = (kernel_size, kernel_size) - if isinstance(kernel_size, tuple) and (len(kernel_size) != 2 or - (not isinstance(kernel_size[0], int)) or - (not isinstance(kernel_size[1], int)) or - kernel_size[0] < 1 or kernel_size[1] < 1): - raise ValueError(f"Attr kernel_size of DepthwiseConv2dNative Op not passed " - f"{kernel_size}, should be a int or tuple and equal to or greater than 1.") + self.kernel_size = (kernel_size, kernel_size) + if len(self.kernel_size) != 2 or (not isinstance(self.kernel_size[0], int)) or \ + (not isinstance(self.kernel_size[1], int)) or \ + self.kernel_size[0] < 1 or self.kernel_size[1] < 1: + raise ValueError(f"The \'kernel_size\' of \'DepthwiseConv2dNative\' should be an positive int number or " + f"a tuple of two positive int numbers, but got {kernel_size}") + self.stride = validator.check_type('stride', stride, (int, tuple)) + if isinstance(stride, int): + self.stride = (stride, stride) + if len(self.stride) != 2 or (not isinstance(self.stride[0], int)) or \ + (not isinstance(self.stride[1], int)) or \ + self.stride[0] < 1 or self.stride[1] < 1: + raise ValueError(f"The \'stride\' of \'DepthwiseConv2dNative\' should be an positive int number or " + f"a tuple of two positive int numbers, but got {stride}") + self.add_prim_attr('stride', (1, 1, self.stride[0], self.stride[1])) + self.dilation = validator.check_type('dilation', dilation, (tuple, int)) + if isinstance(dilation, int): + self.dilation = (dilation, dilation) + if len(self.dilation) != 2 or (not isinstance(self.dilation[0], int)) or \ + (not isinstance(self.dilation[1], int)) or \ + self.dilation[0] < 1 or self.dilation[1] < 1: + raise ValueError(f"The \'dilation\' of \'DepthwiseConv2dNative\' should be an positive int number or " + f"a tuple of two or four positive int numbers, but got {dilation}") + self.add_prim_attr('dilation', (1, 1, self.dilation[0], self.dilation[1])) + validator.equal('type of pad', type(pad), 'not bool', not isinstance(pad, bool)) if pad_mode not in ("same", "valid", "pad"): raise ValueError(f"Attr pad_mode of DepthwiseConv2dNative Op not passed" f"{pad_mode} not in valid, same, pad.") @@ -627,9 +661,6 @@ class DepthwiseConv2dNative(PrimitiveWithInfer): self.add_prim_attr('data_format', "NCHW") self.channel_multiplier = validator.check_integer("channel_multiplier", channel_multiplier, 0, Rel.GT) self.group = validator.check_integer("group", group, 0, Rel.GT) - self.dilation = validator.check_integer("dilation", dilation, 1, Rel.GE) - self.kernel_size = validator.check_value_on_integer("kernel_size", kernel_size, 1, Rel.GE) - self.stride = validator.check_integer("stride", stride, 1, Rel.GE) self.pad = pad def infer_shape(self, x_shape, w_shape): @@ -640,29 +671,33 @@ class DepthwiseConv2dNative(PrimitiveWithInfer): kernel_size_h = w_shape[2] kernel_size_w = w_shape[3] + stride_h = self.stride[2] + stride_w = self.stride[3] + dilation_h = self.dilation[2] + dilation_w = self.dilation[3] if self.pad_mode == "valid": - h_out = math.ceil((x_shape[2] - self.dilation * (kernel_size_h - 1)) / self.stride) - w_out = math.ceil((x_shape[3] - self.dilation * (kernel_size_w - 1)) / self.stride) + h_out = math.ceil((x_shape[2] - dilation_h * (kernel_size_h - 1)) / stride_h) + w_out = math.ceil((x_shape[3] - dilation_w * (kernel_size_w - 1)) / stride_w) pad_top, pad_bottom, pad_left, pad_right = 0, 0, 0, 0 elif self.pad_mode == "same": - h_out = math.ceil(x_shape[2] / self.stride) - w_out = math.ceil(x_shape[3] / self.stride) + h_out = math.ceil(x_shape[2] / stride_h) + w_out = math.ceil(x_shape[3] / stride_w) - pad_needed_h = max(0, (h_out - 1) * self.stride + self.dilation * (kernel_size_h - 1) + 1 - x_shape[2]) + pad_needed_h = max(0, (h_out - 1) * stride_h+ dilation_h * (kernel_size_h - 1) + 1 - x_shape[2]) pad_top = math.floor(pad_needed_h / 2) pad_bottom = pad_needed_h - pad_top - pad_needed_w = max(0, (w_out - 1) * self.stride + self.dilation * (kernel_size_w - 1) + 1 - x_shape[3]) + pad_needed_w = max(0, (w_out - 1) * stride_w + dilation_w * (kernel_size_w - 1) + 1 - x_shape[3]) pad_left = math.floor(pad_needed_w / 2) pad_right = pad_needed_w - pad_left elif self.pad_mode == 'pad': pad_top, pad_bottom, pad_left, pad_right = self.pad, self.pad, self.pad, self.pad - h_out = 1 + (x_shape[2] + 2 * self.pad - kernel_size_h - (kernel_size_h - 1) * (self.dilation - 1)) \ - / self.stride - w_out = 1 + (x_shape[3] + 2 * self.pad - kernel_size_w - (kernel_size_w - 1) * (self.dilation - 1)) \ - / self.stride + h_out = 1 + (x_shape[2] + 2 * self.pad - kernel_size_h - (kernel_size_h - 1) * (dilation_h - 1)) \ + / stride_h + w_out = 1 + (x_shape[3] + 2 * self.pad - kernel_size_w - (kernel_size_w - 1) * (dilation_w - 1)) \ + / stride_w h_out = math.floor(h_out) w_out = math.floor(w_out) else: @@ -715,7 +750,7 @@ class _Pool(PrimitiveWithInfer): (not isinstance(ksize[1], int)) or ksize[0] <= 0 or ksize[1] <= 0): - raise ValueError(f"The 'ksize' passed to operator {self.name} should be an positive int number or" + raise ValueError(f"The 'ksize' passed to operator {self.name} should be an positive int number or " f"a tuple of two positive int numbers, but got {ksize}") self.ksize = (1, 1, ksize[0], ksize[1]) if self.is_maxpoolwithargmax: @@ -731,7 +766,7 @@ class _Pool(PrimitiveWithInfer): (not isinstance(strides[1], int)) or strides[0] <= 0 or strides[1] <= 0): - raise ValueError(f"The 'strides' passed to operator {self.name} should be an positive int number or" + raise ValueError(f"The 'strides' passed to operator {self.name} should be an positive int number or " f"a tuple of two positive int numbers, but got {strides}") self.strides = (1, 1, strides[0], strides[1]) if self.is_maxpoolwithargmax: @@ -853,7 +888,6 @@ class MaxPoolWithArgmax(_Pool): - **output** (Tensor) - Maxpooling result, with shape :math:`(N, C_{out}, H_{out}, W_{out})`. - **mask** (Tensor) - Max values' index represented by the mask. """ - def __init__(self, ksize=1, strides=1, padding="valid"): super(MaxPoolWithArgmax, self).__init__(ksize, strides, padding) self.is_tbe = context.get_context("device_target") == "Ascend" @@ -944,8 +978,8 @@ class Conv2DBackpropInput(PrimitiveWithInfer): pad (int): The pad value to fill. Default: 0. mode (int): 0 Math convolutiuon, 1 cross-correlation convolution , 2 deconvolution, 3 depthwise convolution. Default: 1. - stride (int): The stride to apply conv filter. Default: 1. - dilation (int): Specifies the dilation rate to use for dilated convolution. Default: 1. + stride (Union[int. tuple[int]]): The stride to apply conv filter. Default: 1. + dilation (Union[int. tuple[int]]): Specifies the dilation rate to use for dilated convolution. Default: 1. group (int): Splits input into groups. Default: 1. Returns: @@ -967,25 +1001,41 @@ class Conv2DBackpropInput(PrimitiveWithInfer): self.init_prim_io_names(inputs=['out_backprop', 'filter', 'input_sizes'], outputs=['output']) self.out_channel = validator.check_integer('out_channel', out_channel, 0, Rel.GT) self.kernel_size = validator.check_type('kernel_size', kernel_size, (int, tuple)) - if isinstance(self.kernel_size, int): - if kernel_size < 1: - raise ValueError('Attr \'kernel_size\' of \'Conv2DBackpropInput\' Op passed ' - + str(self.kernel_size) + ', should be a int or tuple and equal to or greater than 1.') - self.kernel_size = (self.kernel_size, self.kernel_size) - elif isinstance(kernel_size, tuple) and (len(kernel_size) != 2 or - (not isinstance(kernel_size[0], int)) or - (not isinstance(kernel_size[1], int)) or - kernel_size[0] < 1 or kernel_size[1] < 1): - raise ValueError('Attr \'kernel_size\' of \'Conv2DBackpropInput\' Op passed ' - + str(self.kernel_size) + ', should be a int or tuple and equal to or greater than 1.') + if isinstance(kernel_size, int): + self.kernel_size = (kernel_size, kernel_size) + if len(self.kernel_size) != 2 or (not isinstance(self.kernel_size[0], int)) or \ + (not isinstance(self.kernel_size[1], int)) or \ + self.kernel_size[0] < 1 or self.kernel_size[1] < 1: + raise ValueError(f"The \'kernel_size\' of \'Conv2DBackpropInput\' should be an positive int number or " + f"a tuple of two positive int numbers, but got {kernel_size}") + self.stride = validator.check_type('stride', stride, (int, tuple)) + if isinstance(stride, int): + self.stride = (stride, stride) + elif isinstance(stride, tuple) and len(stride) == 4: + self.stride = (stride[2], stride[3]) + if len(self.stride) != 2 or (not isinstance(self.stride[0], int)) or (not isinstance(self.stride[1], int)) or \ + self.stride[0] < 1 or self.stride[1] < 1: + raise ValueError(f"The \'stride\' of \'Conv2DBackpropInput\' should be an positive int number or " + f"a tuple of two or four positive int numbers, but got {stride}") + self.add_prim_attr('stride', self.stride) + self.dilation = validator.check_type('dilation', dilation, (tuple, int)) + if isinstance(dilation, int): + self.dilation = (1, 1, dilation, dilation) + elif len(dilation) == 2: + self.dilation = (1, 1, dilation[0], dilation[1]) + if len(self.dilation) != 4 or (not isinstance(self.dilation[0], int) or self.dilation[0] < 1) or \ + (not isinstance(self.dilation[1], int) or self.dilation[1] < 1) or \ + (not isinstance(self.dilation[2], int) or self.dilation[2] < 1) or \ + (not isinstance(self.dilation[3], int) or self.dilation[3] < 1): + raise ValueError(f"The \'dilation\' of \'Conv2DBackpropInput\' should be an positive int number or " + f"a tuple of two or four positive int numbers, but got {dilation}") + self.add_prim_attr('dilation', self.dilation) validator.equal('type of pad', type(pad), 'not bool', not isinstance(pad, bool)) validator.equal('type of pad', type(pad), 'int', isinstance(pad, int)) self.pad_mode = validator.check_string('pad_mode', pad_mode, ['valid', 'same', 'pad']) self.pad = validator.check_pad_value_by_mode(self.__class__.__name__, pad_mode, pad) self.mode = validator.check_integer('mode', mode, 1, Rel.EQ) self.group = validator.check_integer('group', group, 0, Rel.GT) - self.dilation = validator.check_integer('dilation', dilation, 1, Rel.GE) - self.stride = validator.check_integer('stride', stride, 1, Rel.GE) pad_mode = pad_mode.upper() self.add_prim_attr('pad_mode', pad_mode) self.add_prim_attr('data_format', "NCHW") @@ -1004,16 +1054,18 @@ class Conv2DBackpropInput(PrimitiveWithInfer): dout_shape = doutput['shape'] kernel_h = self.kernel_size[0] kernel_w = self.kernel_size[1] + stride_h = self.stride[0] + stride_w = self.stride[1] # default pad mode is valid pad_list = (0, 0, 0, 0) if self.pad_list: pad_list = tuple(self.pad_list) elif self.pad_mode == "SAME": - pad_needed_h = max(0, (dout_shape[2] - 1) * self.stride + kernel_h - x_size_v[2]) + pad_needed_h = max(0, (dout_shape[2] - 1) * stride_h + kernel_h - x_size_v[2]) pad_top = math.floor(pad_needed_h / 2) pad_bottom = pad_needed_h - pad_top - pad_needed_w = max(0, (dout_shape[3] - 1) * self.stride + kernel_w - x_size_v[3]) + pad_needed_w = max(0, (dout_shape[3] - 1) * stride_w + kernel_w - x_size_v[3]) pad_left = math.floor(pad_needed_w / 2) pad_right = pad_needed_w - pad_left pad_list = (pad_top, pad_bottom, pad_left, pad_right) diff --git a/tests/st/ops/cpu/test_conv2d_backprop_filter_op.py b/tests/st/ops/cpu/test_conv2d_backprop_filter_op.py index 75ca91549936bcb154bdaa9422ca009d5bc2e2af..c2f8422e307c99b5eab02124ed9f4a50dc814d60 100644 --- a/tests/st/ops/cpu/test_conv2d_backprop_filter_op.py +++ b/tests/st/ops/cpu/test_conv2d_backprop_filter_op.py @@ -35,8 +35,8 @@ class Net4(nn.Cell): pad_mode="valid", pad=0, mode=1, - stride=1, - dilation=1, + stride=(1, 1), + dilation=(1, 1, 1, 1), group=1) self.w = Parameter(initializer(Tensor(np.array([[[[1, 0, -1], [1, 0, -1], [1, 0, -1]]]]).astype(np.float32)), [1, 1, 3, 3]), name='w') self.x = Parameter(initializer(Tensor(np.array([[[ diff --git a/tests/st/ops/gpu/test_conv2d_backprop_filter_op.py b/tests/st/ops/gpu/test_conv2d_backprop_filter_op.py index 6e2e76cd47cb362133023b516dd1a0f4f1cb12f3..0f66f2fac5fac5bc117535e7c6d425965f5def2c 100644 --- a/tests/st/ops/gpu/test_conv2d_backprop_filter_op.py +++ b/tests/st/ops/gpu/test_conv2d_backprop_filter_op.py @@ -35,8 +35,8 @@ class Conv2dFilter(nn.Cell): pad_mode="valid", pad=0, mode=1, - stride=1, - dilation=1, + stride=(1, 1), + dilation=(1, 1, 1, 1), group=1) self.get_shape = P.Shape() diff --git a/tests/ut/cpp/python_input/gtest_input/pynative/ops_test.py b/tests/ut/cpp/python_input/gtest_input/pynative/ops_test.py index 46c6fdd1cf760457abd631f28d4cf2810d48facf..c7de09dbcbdef93ffc671a4d45b16aac1ba2bb89 100644 --- a/tests/ut/cpp/python_input/gtest_input/pynative/ops_test.py +++ b/tests/ut/cpp/python_input/gtest_input/pynative/ops_test.py @@ -21,17 +21,17 @@ from mindspore.common.tensor import Tensor def im2col(img, filter_h, filter_w, stride=1, pad=0, dilation=1): """Rearranges an image to row vector""" batch_num, channel, height, width = img.shape - out_h = (height + 2*pad - filter_h - (filter_h - 1) * (dilation - 1))//stride + 1 - out_w = (width + 2*pad - filter_w - (filter_w - 1) * (dilation - 1))//stride + 1 + out_h = (height + 2*pad - filter_h - (filter_h - 1) * (dilation[2] - 1))//stride[2] + 1 + out_w = (width + 2*pad - filter_w - (filter_w - 1) * (dilation[3] - 1))//stride[3] + 1 img = np.pad(img, [(0, 0), (0, 0), (pad, pad), (pad, pad)], 'constant') col = np.zeros((batch_num, channel, filter_h, filter_w, out_h, out_w)).astype(img.dtype) for y in range(filter_h): - y_max = y + stride*out_h + y_max = y + stride[2]*out_h for x in range(filter_w): - x_max = x + stride*out_w - col[:, :, y, x, :, :] = img[:, :, y:y_max:stride, x:x_max:stride] + x_max = x + stride[2]*out_w + col[:, :, y, x, :, :] = img[:, :, y:y_max:stride[2], x:x_max:stride[2]] col = col.transpose(0, 4, 5, 1, 2, 3).reshape(batch_num*out_h*out_w, -1) return col @@ -42,8 +42,8 @@ def conv2d(x, weight, bias=None, stride=1, pad=0, """Convolution 2D""" batch_num, _, x_h, x_w = x.shape filter_num, _, filter_h, filter_w = weight.shape - out_h = 1 + int((x_h + 2 * pad - filter_h - (filter_h - 1) * (dilation - 1)) / stride) - out_w = 1 + int((x_w + 2 * pad - filter_w - (filter_w - 1) * (dilation - 1)) / stride) + out_h = 1 + int((x_h + 2 * pad - filter_h - (filter_h - 1) * (dilation[2] - 1)) / stride[2]) + out_w = 1 + int((x_w + 2 * pad - filter_w - (filter_w - 1) * (dilation[3] - 1)) / stride[3]) col = im2col(x, filter_h, filter_w, stride, pad, dilation) col_w = np.reshape(weight, (filter_num, -1)).T out = np.dot(col, col_w) diff --git a/tests/vm_impl/vm_me.py b/tests/vm_impl/vm_me.py index ba51a3b13bcf77116f8586b7537d1d766a0e6e0f..a189aa8b90dfd716bced605b19fd48146e70a3aa 100644 --- a/tests/vm_impl/vm_me.py +++ b/tests/vm_impl/vm_me.py @@ -155,23 +155,35 @@ def batch_norm_grad(dy, x, scale, save_mean, save_inv_variance): def col2im(col, input_shape, filter_h, filter_w, stride=1, pad=0): """Rearranges a row vector to an image.""" - validator.check_integer("stride", stride, 0, Rel.GT) + if isinstance(stride, int): + stride_h = stride + stride_w = stride + elif isinstance(stride, tuple) and len(stride) == 2: + stride_h = stride[0] + stride_w = stride[1] + elif isinstance(stride, tuple) and len(stride) == 3: + stride_h = stride[2] + stride_w = stride[3] + else: + raise ValueError(f"The \'stride\' should be an int number or " + f"a tuple of two or four int numbers, but got {stride}") + batch_num, channel, height, width = input_shape - out_h = (height + 2*pad - filter_h)//stride + 1 - out_w = (width + 2*pad - filter_w)//stride + 1 + out_h = (height + 2*pad - filter_h)//stride_h + 1 + out_w = (width + 2*pad - filter_w)//stride_w + 1 col = col.reshape(batch_num, out_h, out_w, channel, filter_h, filter_w) \ .transpose(0, 3, 4, 5, 1, 2) img = np.zeros((batch_num, channel, - height + 2*pad + stride - 1, - width + 2*pad + stride - 1)) \ + height + 2*pad + stride_h - 1, + width + 2*pad + stride_w - 1)) \ .astype(col.dtype) for y in range(filter_h): - y_max = y + stride*out_h + y_max = y + stride_h*out_h for x in range(filter_w): - x_max = x + stride*out_w - img[:, :, y:y_max:stride, x:x_max:stride] += col[:, :, y, x, :, :] + x_max = x + stride_h*out_w + img[:, :, y:y_max:stride_h, x:x_max:stride_h] += col[:, :, y, x, :, :] return img[:, :, pad:height + pad, pad:width + pad] @@ -205,11 +217,35 @@ def conv2d(x, weight, bias=None, stride=1, pad=0, dilation=1, groups=1, padding_mode='zeros'): """Convolution 2D.""" # pylint: disable=unused-argument - validator.check_integer("stride", stride, 0, Rel.GT) + validator.check_type('stride', stride, (int, tuple)) + if isinstance(stride, int): + stride = (stride, stride) + elif len(stride) == 4: + stride = (stride[2], stride[3]) + if len(stride) != 2 or (not isinstance(stride[0], int)) or \ + (not isinstance(stride[1], int)) or \ + stride[0] < 1 or stride[1] < 1: + raise ValueError(f"The \'stride\' of \'conv2d\' should be an positive int number or " + f"a tuple of two positive int numbers, but got {stride}") + stride_h = stride[0] + stride_w = stride[1] + validator.check_type('dilation', dilation, (int, tuple)) + if isinstance(dilation, int): + dilation = (dilation, dilation) + elif len(dilation) == 4: + dilation = (dilation[2], dilation[3]) + if len(dilation) != 2 or (not isinstance(dilation[0], int)) or \ + (not isinstance(dilation[1], int)) or \ + dilation[0] < 1 or dilation[1] < 1: + raise ValueError(f"The \'dilation\' of \'conv2d\' should be an positive int number or " + f"a tuple of two positive int numbers, but got {dilation}") + dilation_h = dilation[0] + dilation_w = dilation[1] + batch_num, _, x_h, x_w = x.shape filter_num, _, filter_h, filter_w = weight.shape - out_h = 1 + int((x_h + 2 * pad - filter_h - (filter_h - 1) * (dilation - 1)) / stride) - out_w = 1 + int((x_w + 2 * pad - filter_w - (filter_w - 1) * (dilation - 1)) / stride) + out_h = 1 + int((x_h + 2 * pad - filter_h - (filter_h - 1) * (dilation_h - 1)) / stride_h) + out_w = 1 + int((x_w + 2 * pad - filter_w - (filter_w - 1) * (dilation_w - 1)) / stride_w) col = im2col(x, filter_h, filter_w, stride, pad, dilation) col_w = np.reshape(weight, (filter_num, -1)).T out = np.dot(col, col_w) @@ -286,19 +322,43 @@ def flatten_grad(dout, x): def im2col(img, filter_h, filter_w, stride=1, pad=0, dilation=1): """Rearranges an image to row vector.""" - validator.check_integer("stride", stride, 0, Rel.GT) + if isinstance(stride, int): + stride_h = stride + stride_w = stride + elif isinstance(stride, tuple) and len(stride) == 2: + stride_h = stride[0] + stride_w = stride[1] + elif isinstance(stride, tuple) and len(stride) == 3: + stride_h = stride[2] + stride_w = stride[3] + else: + raise ValueError(f"The \'stride\' should be an int number or " + f"a tuple of two or four int numbers, but got {stride}") + if isinstance(dilation, int): + dilation_h = dilation + dilation_w = dilation + elif isinstance(dilation, tuple) and len(dilation) == 2: + dilation_h = dilation[0] + dilation_w = dilation[1] + elif isinstance(dilation, tuple) and len(dilation) == 3: + dilation_h = dilation[2] + dilation_w = dilation[3] + else: + raise ValueError(f"The \'dilation\' should be an int number or " + f"a tuple of two or four int numbers, but got {dilation}") + batch_num, channel, height, width = img.shape - out_h = (height + 2*pad - filter_h- (filter_h - 1) * (dilation - 1))//stride + 1 - out_w = (width + 2*pad - filter_w- (filter_w - 1) * (dilation - 1))//stride + 1 + out_h = (height + 2*pad - filter_h- (filter_h - 1) * (dilation_h - 1))//stride_h + 1 + out_w = (width + 2*pad - filter_w- (filter_w - 1) * (dilation_w - 1))//stride_w + 1 img = np.pad(img, [(0, 0), (0, 0), (pad, pad), (pad, pad)], 'constant') col = np.zeros((batch_num, channel, filter_h, filter_w, out_h, out_w)).astype(img.dtype) for y in range(filter_h): - y_max = y + stride*out_h + y_max = y + stride_h*out_h for x in range(filter_w): - x_max = x + stride*out_w - col[:, :, y, x, :, :] = img[:, :, y:y_max:stride, x:x_max:stride] + x_max = x + stride_h*out_w + col[:, :, y, x, :, :] = img[:, :, y:y_max:stride_h, x:x_max:stride_h] col = col.transpose(0, 4, 5, 1, 2, 3).reshape(batch_num*out_h*out_w, -1) return col