From f454303f6a2a0c4c1e3d38c88f4a44f72662c332 Mon Sep 17 00:00:00 2001 From: alexlyulkov Date: Tue, 9 Apr 2024 11:20:56 +0300 Subject: [PATCH] Merge pull request #25241 from alexlyulkov:al/int64-padding Added int support to padding layer #25241 Added int32 and int64 support to padding layer (CPU and CUDA). ONNX parser doesn't convert non-zero padding value to float now. ### Pull Request Readiness Checklist See details at https://github.com/opencv/opencv/wiki/How_to_contribute#making-a-good-pull-request - [x] I agree to contribute to the project under Apache 2 License. - [x] To the best of my knowledge, the proposed patch is not based on a code under GPL or another license that is incompatible with OpenCV - [x] The PR is proposed to the proper branch - [ ] There is a reference to the original bug report and related work - [ ] There is accuracy test, performance test and test data in opencv_extra repository, if applicable Patch to opencv_extra has the same branch name. - [x] The feature is well documented and sample code can be built with the project CMake --- modules/dnn/src/cuda/padding.cu | 2 + .../dnn/src/cuda4dnn/primitives/padding.hpp | 2 +- modules/dnn/src/layers/padding_layer.cpp | 26 ++++++- modules/dnn/src/onnx/onnx_importer.cpp | 14 +++- modules/dnn/src/tensorflow/tf_importer.cpp | 2 +- modules/dnn/test/test_int.cpp | 73 +++++++++++++++++++ 6 files changed, 110 insertions(+), 9 deletions(-) diff --git a/modules/dnn/src/cuda/padding.cu b/modules/dnn/src/cuda/padding.cu index fc55ce0633..b08e22894f 100644 --- a/modules/dnn/src/cuda/padding.cu +++ b/modules/dnn/src/cuda/padding.cu @@ -197,5 +197,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { template void copy_with_reflection101(const Stream&, TensorSpan<__half>, TensorView<__half>, std::vector> ranges); #endif template void copy_with_reflection101(const Stream&, TensorSpan, TensorView, std::vector> ranges); + template void copy_with_reflection101(const Stream&, TensorSpan, TensorView, std::vector> ranges); + template void copy_with_reflection101(const Stream&, TensorSpan, TensorView, std::vector> ranges); }}}} /* namespace namespace cv::dnn::cuda4dnn::kernels */ diff --git a/modules/dnn/src/cuda4dnn/primitives/padding.hpp b/modules/dnn/src/cuda4dnn/primitives/padding.hpp index ce2a5c3c47..689436901c 100644 --- a/modules/dnn/src/cuda4dnn/primitives/padding.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/padding.hpp @@ -34,7 +34,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { using wrapper_type = GetCUDABackendWrapperType; /* `ranges` is indexed by axis and contains the range in the output where the input is copied to */ - PaddingOp(csl::Stream stream_, PaddingType type_, T value_, std::vector ranges) + PaddingOp(csl::Stream stream_, PaddingType type_, T value_, const std::vector& ranges) : stream(std::move(stream_)), type{ type_ }, value{ value_ }, dstRanges(std::move(ranges)) { } diff --git a/modules/dnn/src/layers/padding_layer.cpp b/modules/dnn/src/layers/padding_layer.cpp index ee3dac87e2..4e864c1f66 100644 --- a/modules/dnn/src/layers/padding_layer.cpp +++ b/modules/dnn/src/layers/padding_layer.cpp @@ -34,7 +34,7 @@ public: PaddingLayerImpl(const LayerParams ¶ms) { setParamsFrom(params); - paddingValue = params.get("value", 0); + paddingValue = params.get("value", 0); inputDims = params.get("input_dims", -1); paddingType = params.get("type", "constant"); @@ -70,6 +70,23 @@ public: return false; } + void getTypes(const std::vector& inputs, + const int requiredOutputs, + const int requiredInternals, + std::vector& outputs, + std::vector& internals) const CV_OVERRIDE + { + CV_CheckEQ(inputs.size(), 1u, ""); + if (preferableTarget == DNN_TARGET_CUDA_FP16 || preferableTarget == DNN_TARGET_CUDA) + CV_CheckType(inputs[0], inputs[0] == CV_32F || inputs[0] == CV_32S || inputs[0] == CV_64S, ""); + else if (preferableTarget == DNN_TARGET_OPENCL_FP16) + CV_CheckType(inputs[0], inputs[0] == CV_16F || inputs[0] == CV_8S || inputs[0] == CV_32S || inputs[0] == CV_64S, ""); + else + CV_CheckType(inputs[0], inputs[0] == CV_32F || inputs[0] == CV_8S || inputs[0] == CV_32S || inputs[0] == CV_64S, ""); + + outputs.assign(requiredOutputs, inputs[0]); + } + void finalize(InputArrayOfArrays inputs_arr, OutputArrayOfArrays) CV_OVERRIDE { std::vector inputs; @@ -184,7 +201,7 @@ public: else CV_Error(Error::StsNotImplemented, "Unsupported padding mode"); - return make_cuda_node(preferableTarget, std::move(context->stream), ptype, paddingValue, dstRanges); + return make_cuda_node_with_type(preferableTarget, inputs[0]->getHostMatDepth(), std::move(context->stream), ptype, paddingValue, dstRanges); } #endif @@ -248,7 +265,8 @@ public: auto padding_below = std::make_shared(ov::element::i64, ov::Shape{begins.size()}, begins.data()); auto padding_above = std::make_shared(ov::element::i64, ov::Shape{ends.size()}, ends.data()); auto pad_mode = paddingType == "constant" ? ov::op::PadMode::CONSTANT : ov::op::PadMode::REFLECT; // SYMMETRIC - auto arg_pad_value = std::make_shared(ov::element::f32, ov::Shape{}, &paddingValue);; + float paddingValueFloat = paddingValue; + auto arg_pad_value = std::make_shared(ov::element::f32, ov::Shape{}, &paddingValueFloat); auto pad = paddingType == "constant" ? std::make_shared(ieInpNode, padding_below, padding_above, arg_pad_value, pad_mode) : @@ -261,7 +279,7 @@ private: std::vector > paddings; // Pairs pad before, pad after. std::vector dstRanges; int inputDims; - float paddingValue; + double paddingValue; std::string paddingType; }; diff --git a/modules/dnn/src/onnx/onnx_importer.cpp b/modules/dnn/src/onnx/onnx_importer.cpp index eb2d1b3ea5..2e41c81b3a 100644 --- a/modules/dnn/src/onnx/onnx_importer.cpp +++ b/modules/dnn/src/onnx/onnx_importer.cpp @@ -2410,8 +2410,16 @@ void ONNXImporter::parsePad(LayerParams& layerParams, const opencv_onnx::NodePro if (node_proto.input_size() == 3 && !node_proto.input(2).empty()) { Mat value = getBlob(node_proto, 2); - float padValue = (depth == CV_8S) ? (float)value.ptr()[0] : value.ptr()[0]; - layerParams.set("value", padValue); + double padValue = 0; + switch(value.depth()) + { + case CV_32F: padValue = value.ptr()[0]; break; + case CV_32S: padValue = value.ptr()[0]; break; + case CV_64S: padValue = value.ptr()[0]; break; + case CV_8S: padValue = value.ptr()[0]; break; + default: CV_Error(Error::BadDepth, "Unsupported type"); + } + layerParams.set("value", (double)padValue); } } addLayer(layerParams, node_proto); @@ -3403,7 +3411,7 @@ void ONNXImporter::parseQConv(LayerParams& layerParams, const opencv_onnx::NodeP padLp.type = "PaddingInt8"; padLp.set("paddings", DictValue::arrayInt(&paddings[0], paddings.size())); padLp.set("depth", CV_8S); - padLp.set("value", inp_zp); + padLp.set("value", (double)inp_zp); opencv_onnx::NodeProto proto; proto.add_input(node_proto.input(0)); diff --git a/modules/dnn/src/tensorflow/tf_importer.cpp b/modules/dnn/src/tensorflow/tf_importer.cpp index 8fe9d130e0..1eb49a773c 100644 --- a/modules/dnn/src/tensorflow/tf_importer.cpp +++ b/modules/dnn/src/tensorflow/tf_importer.cpp @@ -617,7 +617,7 @@ void TFImporter::setPadding(LayerParams &layerParams, const tensorflow::NodeDef padLp.name = layer.name() + "/pad"; padLp.type = "Padding"; padLp.set("paddings", DictValue::arrayInt(pads, sizeof(pads) / sizeof(pads[0]))); - padLp.set("value", value); + padLp.set("value", (double)value); int id = dstNet.addLayer(padLp.name, padLp.type, padLp); layer_id[padLp.name] = id; diff --git a/modules/dnn/test/test_int.cpp b/modules/dnn/test/test_int.cpp index 4346f4c211..a17b0b4bb1 100644 --- a/modules/dnn/test/test_int.cpp +++ b/modules/dnn/test/test_int.cpp @@ -775,6 +775,79 @@ INSTANTIATE_TEST_CASE_P(/**/, Test_Cast_Int, Combine( dnnBackendsAndTargets() )); +typedef testing::TestWithParam > > Test_Pad_Int; +TEST_P(Test_Pad_Int, random) +{ + int matType = get<0>(GetParam()); + tuple backend_target= get<1>(GetParam()); + Backend backend = get<0>(backend_target); + Target target = get<1>(backend_target); + + std::vector inShape{2, 3, 4, 5}; + int64_t low = 1000000; + Mat input(inShape, matType); + cv::randu(input, low, low + 100); + std::vector paddings{0, 0, 0, 0, 1, 0, 0, 1}; + + Net net; + LayerParams lp; + lp.type = "Padding"; + lp.name = "testLayer"; + lp.set("paddings", DictValue::arrayInt(&paddings[0], paddings.size())); + lp.set("value", 25); + + net.addLayerToPrev(lp.name, lp.type, lp); + + net.setInput(input); + net.setPreferableBackend(backend); + net.setPreferableTarget(target); + + Mat re; + re = net.forward(); + EXPECT_EQ(re.depth(), matType); + EXPECT_EQ(re.size.dims(), 4); + EXPECT_EQ(re.size[0], 2); + EXPECT_EQ(re.size[1], 3); + EXPECT_EQ(re.size[2], 5); + EXPECT_EQ(re.size[3], 6); + + std::vector reIndices(4); + std::vector inIndices(4); + for (int i0 = 0; i0 < re.size[0]; ++i0) + { + reIndices[0] = i0; + inIndices[0] = i0; + for (int i1 = 0; i1 < re.size[1]; ++i1) + { + reIndices[1] = i1; + inIndices[1] = i1; + for (int i2 = 0; i2 < re.size[2]; ++i2) + { + reIndices[2] = i2; + inIndices[2] = i2 - 1; + for (int i3 = 0; i3 < re.size[3]; ++i3) + { + reIndices[3] = i3; + inIndices[3] = i3; + if (i2 < 1 || i3 >= input.size[3]) + { + EXPECT_EQ(getValueAt(re, reIndices.data()), 25l); + } + else + { + EXPECT_EQ(getValueAt(re, reIndices.data()), getValueAt(input, inIndices.data())); + } + } + } + } + } +} + +INSTANTIATE_TEST_CASE_P(/**/, Test_Pad_Int, Combine( + testing::Values(CV_32S, CV_64S), + dnnBackendsAndTargets() +)); + typedef testing::TestWithParam > > Test_Slice_Int; TEST_P(Test_Slice_Int, random) {