From f2cf3c889043ea1e5fb1e25c24b297a5a60bf4f9 Mon Sep 17 00:00:00 2001 From: alexlyulkov Date: Fri, 22 Mar 2024 03:39:42 +0300 Subject: [PATCH] Added int support to flatten, permute, reshape, slice layers (#25236) Co-authored-by: Alexander Lyulkov --- modules/dnn/src/cuda/permute.cu | 4 + modules/dnn/src/cuda/slice.cu | 2 + modules/dnn/src/layers/flatten_layer.cpp | 21 +++- modules/dnn/src/layers/permute_layer.cpp | 4 +- modules/dnn/src/layers/reshape_layer.cpp | 4 +- modules/dnn/src/layers/slice_layer.cpp | 4 +- modules/dnn/test/test_int.cpp | 139 ++++++++++++++++++++++- 7 files changed, 168 insertions(+), 10 deletions(-) diff --git a/modules/dnn/src/cuda/permute.cu b/modules/dnn/src/cuda/permute.cu index 35c95a6737..6510fb84cd 100644 --- a/modules/dnn/src/cuda/permute.cu +++ b/modules/dnn/src/cuda/permute.cu @@ -107,6 +107,8 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { template void transpose(const Stream&, Span<__half>, View<__half>, std::size_t, std::size_t); template void transpose(const Stream&, Span, View, std::size_t, std::size_t); + template void transpose(const Stream&, Span, View, std::size_t, std::size_t); + template void transpose(const Stream&, Span, View, std::size_t, std::size_t); template static void launch_permute_kernel( @@ -284,5 +286,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { template void permute(const Stream&, TensorSpan<__half>, TensorView<__half>, std::vector); #endif template void permute(const Stream&, TensorSpan, TensorView, std::vector); + template void permute(const Stream&, TensorSpan, TensorView, std::vector); + template void permute(const Stream&, TensorSpan, TensorView, std::vector); }}}} /* namespace cv::dnn::cuda4dnn::kernels */ diff --git a/modules/dnn/src/cuda/slice.cu b/modules/dnn/src/cuda/slice.cu index 461e87e549..5465280dae 100644 --- a/modules/dnn/src/cuda/slice.cu +++ b/modules/dnn/src/cuda/slice.cu @@ -199,5 +199,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { template void slice(const Stream&, TensorSpan<__half>, TensorView<__half>, std::vector); #endif template void slice(const Stream&, TensorSpan, TensorView, std::vector); + template void slice(const Stream&, TensorSpan, TensorView, std::vector); + template void slice(const Stream&, TensorSpan, TensorView, std::vector); }}}} /* namespace cv::dnn::cuda4dnn::kernels */ diff --git a/modules/dnn/src/layers/flatten_layer.cpp b/modules/dnn/src/layers/flatten_layer.cpp index d7537fc8fd..f63822b434 100644 --- a/modules/dnn/src/layers/flatten_layer.cpp +++ b/modules/dnn/src/layers/flatten_layer.cpp @@ -118,6 +118,25 @@ public: return true; } + void getTypes(const std::vector& inputs, + const int requiredOutputs, + const int requiredInternals, + std::vector& outputs, + std::vector& internals) const CV_OVERRIDE + { + CV_Assert(inputs.size()); + for (auto input : inputs) + { + if (preferableTarget == DNN_TARGET_OPENCL_FP16) + CV_CheckType(input, input == CV_16F || input == CV_32S || input == CV_64S, ""); + else + CV_CheckType(input, input == CV_32F || input == CV_32S || input == CV_64S, ""); + } + + outputs.assign(requiredOutputs, inputs[0]); + } + + void finalize(InputArrayOfArrays inputs_arr, OutputArrayOfArrays) CV_OVERRIDE { std::vector inputs; @@ -240,7 +259,7 @@ public: ) override { auto context = reinterpret_cast(context_); - return make_cuda_node(preferableTarget, std::move(context->stream)); + return make_cuda_node_with_type(preferableTarget, inputs[0]->getHostMatDepth(), std::move(context->stream)); } #endif diff --git a/modules/dnn/src/layers/permute_layer.cpp b/modules/dnn/src/layers/permute_layer.cpp index ea619512c6..9cb21f7dcf 100644 --- a/modules/dnn/src/layers/permute_layer.cpp +++ b/modules/dnn/src/layers/permute_layer.cpp @@ -188,7 +188,7 @@ public: for (auto input : inputs) { if (preferableTarget == DNN_TARGET_CUDA_FP16 || preferableTarget == DNN_TARGET_CUDA) - CV_CheckTypeEQ(input, CV_32F, "Unsupported type"); + CV_CheckType(input, input == CV_32F || input == CV_32S || input == CV_64S, ""); else if (preferableTarget == DNN_TARGET_OPENCL_FP16) CV_CheckType(input, input == CV_16F || input == CV_8S || input == CV_32S || input == CV_64S, ""); else @@ -521,7 +521,7 @@ public: ) override { auto context = reinterpret_cast(context_); - return make_cuda_node(preferableTarget, std::move(context->stream), _order); + return make_cuda_node_with_type(preferableTarget, inputs[0]->getHostMatDepth(), std::move(context->stream), _order); } #endif diff --git a/modules/dnn/src/layers/reshape_layer.cpp b/modules/dnn/src/layers/reshape_layer.cpp index cdd96d139b..88682a80b1 100644 --- a/modules/dnn/src/layers/reshape_layer.cpp +++ b/modules/dnn/src/layers/reshape_layer.cpp @@ -269,7 +269,7 @@ public: for (auto input : inputs) { if (preferableTarget == DNN_TARGET_CUDA_FP16 || preferableTarget == DNN_TARGET_CUDA) - CV_CheckTypeEQ(input, CV_32F, "Unsupported type"); + CV_CheckType(input, input == CV_32F || input == CV_32S || input == CV_64S, ""); else if (preferableTarget == DNN_TARGET_OPENCL_FP16) CV_CheckType(input, input == CV_16F || input == CV_8S || input == CV_32S || input == CV_64S, ""); else @@ -417,7 +417,7 @@ public: ) override { auto context = reinterpret_cast(context_); - return make_cuda_node(preferableTarget, std::move(context->stream)); + return make_cuda_node_with_type(preferableTarget, inputs[0]->getHostMatDepth(), std::move(context->stream)); } #endif diff --git a/modules/dnn/src/layers/slice_layer.cpp b/modules/dnn/src/layers/slice_layer.cpp index 19b8d1cc43..3888c00791 100644 --- a/modules/dnn/src/layers/slice_layer.cpp +++ b/modules/dnn/src/layers/slice_layer.cpp @@ -288,7 +288,7 @@ public: for (auto input : inputs) { if (preferableTarget == DNN_TARGET_CUDA_FP16 || preferableTarget == DNN_TARGET_CUDA) - CV_CheckEQ(input, CV_32F, "Unsupported type"); + CV_CheckType(input, input == CV_32F || input == CV_32S || input == CV_64S, ""); else if (preferableTarget == DNN_TARGET_OPENCL_FP16) CV_CheckType(input, input == CV_16F || input == CV_8S || input == CV_32S || input == CV_64S, ""); else @@ -827,7 +827,7 @@ public: offsets.push_back(std::move(offsets_i)); } - return make_cuda_node(preferableTarget, std::move(context->stream), std::move(offsets)); + return make_cuda_node_with_type(preferableTarget, inputs[0]->getHostMatDepth(), std::move(context->stream), std::move(offsets)); } #endif diff --git a/modules/dnn/test/test_int.cpp b/modules/dnn/test/test_int.cpp index 283e2f38a9..4ec6ae4ec2 100644 --- a/modules/dnn/test/test_int.cpp +++ b/modules/dnn/test/test_int.cpp @@ -137,9 +137,6 @@ TEST_P(Test_Permute_Int, random) Backend backend = get<0>(backend_target); Target target = get<1>(backend_target); - if(backend == DNN_BACKEND_CUDA) - applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA); - std::vector inShape{2, 3, 4, 5}; int64_t low = matType == CV_64S ? 1000000000000000ll : 1000000000; Mat input(inShape, matType); @@ -363,4 +360,140 @@ INSTANTIATE_TEST_CASE_P(/**/, Test_Cast_Int, Combine( dnnBackendsAndTargets() )); +typedef testing::TestWithParam > > Test_Slice_Int; +TEST_P(Test_Slice_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 inputShape{1, 16, 6, 8}; + std::vector begin{0, 4, 0, 0}; + std::vector end{1, 8, 6, 8}; + int64_t low = matType == CV_64S ? 1000000000000000ll : 1000000000; + Mat input(inputShape, matType); + cv::randu(input, low, low + 100); + + std::vector range(4); + for (int i = 0; i < 4; ++i) + range[i] = Range(begin[i], end[i]); + + Net net; + LayerParams lp; + lp.type = "Slice"; + lp.name = "testLayer"; + lp.set("begin", DictValue::arrayInt(&(begin[0]), 4)); + lp.set("end", DictValue::arrayInt(&(end[0]), 4)); + net.addLayerToPrev(lp.name, lp.type, lp); + + net.setInput(input); + net.setPreferableBackend(backend); + net.setPreferableTarget(target); + Mat out = net.forward(); + + EXPECT_GT(cv::norm(out, NORM_INF), 0); + normAssert(out, input(range)); +} + +INSTANTIATE_TEST_CASE_P(/**/, Test_Slice_Int, Combine( + testing::Values(CV_32S, CV_64S), + dnnBackendsAndTargets() +)); + +typedef testing::TestWithParam > > Test_Reshape_Int; +TEST_P(Test_Reshape_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}; + std::vector outShape{2, 3, 2, 10}; + int64_t low = matType == CV_64S ? 1000000000000000ll : 1000000000; + Mat input(inShape, matType); + cv::randu(input, low, low + 100); + + Net net; + LayerParams lp; + lp.type = "Reshape"; + lp.name = "testLayer"; + lp.set("dim", DictValue::arrayInt(&outShape[0], outShape.size())); + 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], outShape[0]); + EXPECT_EQ(re.size[1], outShape[1]); + EXPECT_EQ(re.size[2], outShape[2]); + EXPECT_EQ(re.size[3], outShape[3]); + + for (int i = 0; i < input.total(); ++i) + { + if (matType == CV_32S) { + EXPECT_EQ(re.ptr()[i], input.ptr()[i]); + } else { + EXPECT_EQ(re.ptr()[i], input.ptr()[i]); + } + } +} + +INSTANTIATE_TEST_CASE_P(/**/, Test_Reshape_Int, Combine( + testing::Values(CV_32S, CV_64S), + dnnBackendsAndTargets() +)); + +typedef testing::TestWithParam > > Test_Flatten_Int; +TEST_P(Test_Flatten_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 = matType == CV_64S ? 1000000000000000ll : 1000000000; + Mat input(inShape, matType); + cv::randu(input, low, low + 100); + + Net net; + LayerParams lp; + lp.type = "Flatten"; + lp.name = "testLayer"; + lp.set("axis", 1); + 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(), 2); + EXPECT_EQ(re.size[0], inShape[0]); + EXPECT_EQ(re.size[1], inShape[1] * inShape[2] * inShape[3]); + + for (int i = 0; i < input.total(); ++i) + { + if (matType == CV_32S) { + EXPECT_EQ(re.ptr()[i], input.ptr()[i]); + } else { + EXPECT_EQ(re.ptr()[i], input.ptr()[i]); + } + } +} + +INSTANTIATE_TEST_CASE_P(/**/, Test_Flatten_Int, Combine( + testing::Values(CV_32S, CV_64S), + dnnBackendsAndTargets() +)); + }} // namespace