Merge branch 4.x

modules/dnn/include/opencv2/dnn/all_layers.hpp

@@ -1203,6 +1203,12 @@ CV__DNN_INLINE_NS_BEGIN
         static Ptr<SpaceToDepthLayer> create(const LayerParams &params);
     };
 
+    class CV_EXPORTS TopKLayer : public Layer
+    {
+    public:
+        static Ptr<TopKLayer> create(const LayerParams& params);
+    };
+
 //! @}
 //! @}
 CV__DNN_INLINE_NS_END

modules/dnn/misc/js/gen_dict.json

@@ -0,0 +1,12 @@
+{
+    "whitelist":
+    {
+        "dnn_Net": ["setInput", "forward", "setPreferableBackend","getUnconnectedOutLayersNames"],
+        "": ["readNetFromCaffe", "readNetFromTensorflow", "readNetFromTorch", "readNetFromDarknet",
+            "readNetFromONNX", "readNetFromTFLite", "readNet", "blobFromImage"]
+    },
+    "namespace_prefix_override":
+    {
+        "dnn": ""
+    }
+}

modules/dnn/perf/perf_layer.cpp

@@ -1041,4 +1041,67 @@ INSTANTIATE_TEST_CASE_P(/**/, Layer_Elementwise,
                                               /* withWebnn= */           false,
                                               /* withCann= */            false));
 
+struct Layer_TopK : public TestBaseWithParam<tuple<Backend, Target>> {
+    void test_layer(const std::vector<int> &input_shape, const int K, const int axis) {
+        int backend_id = get<0>(GetParam());
+        int target_id = get<1>(GetParam());
+
+        Mat input_data(input_shape, CV_32F);
+        randn(input_data, -1.f, 1.f);
+
+        Net net;
+        LayerParams lp;
+        lp.type = "TopK";
+        lp.name = "testLayer";
+        lp.set("k", K);
+        lp.set("axis", axis);
+        net.addLayerToPrev(lp.name, lp.type, lp);
+
+        // Warmup
+        {
+            net.setInput(input_data);
+            net.setPreferableBackend(backend_id);
+            net.setPreferableTarget(target_id);
+            net.forward();
+        }
+
+        TEST_CYCLE() {
+            net.forward();
+        }
+
+        SANITY_CHECK_NOTHING();
+    }
+
+    std::vector<int> input_shape_2d{1000, 100};
+    std::vector<int> input_shape_3d{100, 100, 100};
+};
+
+PERF_TEST_P_(Layer_TopK, TopK_2D_Axis0) {
+    test_layer(input_shape_2d, input_shape_2d[0] / 2, 0);
+}
+PERF_TEST_P_(Layer_TopK, TopK_2D_Axis0_K5) {
+    test_layer(input_shape_2d, 5, 0);
+}
+PERF_TEST_P_(Layer_TopK, TopK_2D_Axis1) {
+    test_layer(input_shape_2d, input_shape_2d[1] / 2, 1);
+}
+PERF_TEST_P_(Layer_TopK, TopK_3D_Axis0) {
+    test_layer(input_shape_3d, input_shape_3d[0] / 2, 0);
+}
+PERF_TEST_P_(Layer_TopK, TopK_3D_Axis1) {
+    test_layer(input_shape_3d, input_shape_3d[1] / 2, 1);
+}
+PERF_TEST_P_(Layer_TopK, TopK_3D_Axis2) {
+    test_layer(input_shape_3d, input_shape_3d[2] / 2, 2);
+}
+INSTANTIATE_TEST_CASE_P(/**/, Layer_TopK,
+                        dnnBackendsAndTargets(/* withInferenceEngine= */ false,
+                                              /* withHalide= */          false,
+                                              /* withCpuOCV= */          true,
+                                              /* withVkCom= */           false,
+                                              /* withCUDA= */            false,
+                                              /* withNgraph= */          false,
+                                              /* withWebnn= */           false,
+                                              /* withCann= */            false));
+
 } // namespace

modules/dnn/src/cuda4dnn/csl/cublas.hpp

@@ -425,8 +425,8 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace csl { namespace cu
 
         const auto batch_count = static_cast<int>(batchCount);
 
-        AutoBuffer<half> buffer(3 * batch_count);
-        auto A_slices = (half**)(buffer.data());
+        AutoBuffer<half*> buffer(3 * batch_count);
+        auto A_slices = buffer.data();
         auto B_slices = A_slices + batch_count;
         auto C_slices = B_slices + batch_count;
         // collect A, B and C slices
@@ -438,18 +438,18 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace csl { namespace cu
 
         const half **dev_A_slices = 0, **dev_B_slices = 0;
         half **dev_C_slices = 0;
-        cudaMalloc((void**)&dev_A_slices, batch_count * sizeof(half*));
-        cudaMalloc((void**)&dev_B_slices, batch_count * sizeof(half*));
-        cudaMalloc((void**)&dev_C_slices, batch_count * sizeof(half*));
-        cudaMemcpy(dev_A_slices, A_slices, batch_count * sizeof(half*), cudaMemcpyHostToDevice);
-        cudaMemcpy(dev_B_slices, B_slices, batch_count * sizeof(half*), cudaMemcpyHostToDevice);
-        cudaMemcpy(dev_C_slices, C_slices, batch_count * sizeof(half*), cudaMemcpyHostToDevice);
+        CUDA4DNN_CHECK_CUDA(cudaMalloc((void**)&dev_A_slices, batch_count * sizeof(half*)));
+        CUDA4DNN_CHECK_CUDA(cudaMalloc((void**)&dev_B_slices, batch_count * sizeof(half*)));
+        CUDA4DNN_CHECK_CUDA(cudaMalloc((void**)&dev_C_slices, batch_count * sizeof(half*)));
+        CUDA4DNN_CHECK_CUDA(cudaMemcpy(dev_A_slices, A_slices, batch_count * sizeof(half*), cudaMemcpyHostToDevice));
+        CUDA4DNN_CHECK_CUDA(cudaMemcpy(dev_B_slices, B_slices, batch_count * sizeof(half*), cudaMemcpyHostToDevice));
+        CUDA4DNN_CHECK_CUDA(cudaMemcpy(dev_C_slices, C_slices, batch_count * sizeof(half*), cudaMemcpyHostToDevice));
 
         CUDA4DNN_CHECK_CUBLAS(cublasHgemmBatched(handle.get(), opa, opb, iM, iN, iK, &alpha, dev_A_slices, ilda, dev_B_slices, ildb, &beta, dev_C_slices, ildc, batch_count));
 
-        cudaFree(dev_A_slices);
-        cudaFree(dev_B_slices);
-        cudaFree(dev_C_slices);
+        CUDA4DNN_CHECK_CUDA(cudaFree(dev_A_slices));
+        CUDA4DNN_CHECK_CUDA(cudaFree(dev_B_slices));
+        CUDA4DNN_CHECK_CUDA(cudaFree(dev_C_slices));
     }
 
     template <> inline
@@ -475,8 +475,8 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace csl { namespace cu
 
         const auto batch_count = static_cast<int>(batchCount);
 
-        AutoBuffer<float> buffer(3 * batch_count);
-        auto A_slices = (float**)(buffer.data());
+        AutoBuffer<float*> buffer(3 * batch_count);
+        auto A_slices = buffer.data();
         auto B_slices = A_slices + batch_count;
         auto C_slices = B_slices + batch_count;
         // collect A, B and C slices
@@ -488,19 +488,19 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace csl { namespace cu
 
         const float **dev_A_slices = 0, **dev_B_slices = 0;
         float **dev_C_slices = 0;
-        cudaMalloc((void**)&dev_A_slices, batch_count * sizeof(float*));
-        cudaMalloc((void**)&dev_B_slices, batch_count * sizeof(float*));
-        cudaMalloc((void**)&dev_C_slices, batch_count * sizeof(float*));
-        cudaMemcpy(dev_A_slices, A_slices, batch_count * sizeof(float*), cudaMemcpyHostToDevice);
-        cudaMemcpy(dev_B_slices, B_slices, batch_count * sizeof(float*), cudaMemcpyHostToDevice);
-        cudaMemcpy(dev_C_slices, C_slices, batch_count * sizeof(float*), cudaMemcpyHostToDevice);
+        CUDA4DNN_CHECK_CUDA(cudaMalloc((void**)&dev_A_slices, batch_count * sizeof(float*)));
+        CUDA4DNN_CHECK_CUDA(cudaMalloc((void**)&dev_B_slices, batch_count * sizeof(float*)));
+        CUDA4DNN_CHECK_CUDA(cudaMalloc((void**)&dev_C_slices, batch_count * sizeof(float*)));
+        CUDA4DNN_CHECK_CUDA(cudaMemcpy(dev_A_slices, A_slices, batch_count * sizeof(float*), cudaMemcpyHostToDevice));
+        CUDA4DNN_CHECK_CUDA(cudaMemcpy(dev_B_slices, B_slices, batch_count * sizeof(float*), cudaMemcpyHostToDevice));
+        CUDA4DNN_CHECK_CUDA(cudaMemcpy(dev_C_slices, C_slices, batch_count * sizeof(float*), cudaMemcpyHostToDevice));
 
         // cuBLAS is column-major
         CUDA4DNN_CHECK_CUBLAS(cublasSgemmBatched(handle.get(), opa, opb, iM, iN, iK, &alpha, dev_A_slices, ilda, dev_B_slices, ildb, &beta, dev_C_slices, ildc, batch_count));
 
-        cudaFree(dev_A_slices);
-        cudaFree(dev_B_slices);
-        cudaFree(dev_C_slices);
+        CUDA4DNN_CHECK_CUDA(cudaFree(dev_A_slices));
+        CUDA4DNN_CHECK_CUDA(cudaFree(dev_B_slices));
+        CUDA4DNN_CHECK_CUDA(cudaFree(dev_C_slices));
     }
 
 }}}}} /* namespace cv::dnn::cuda4dnn::csl::cublas */

modules/dnn/src/init.cpp

@@ -200,6 +200,7 @@ void initializeLayerFactory()
     CV_DNN_REGISTER_LAYER_CLASS(Scatter,        ScatterLayer);
     CV_DNN_REGISTER_LAYER_CLASS(ScatterND,      ScatterNDLayer);
     CV_DNN_REGISTER_LAYER_CLASS(Tile,           TileLayer);
+    CV_DNN_REGISTER_LAYER_CLASS(TopK,           TopKLayer);
 
     CV_DNN_REGISTER_LAYER_CLASS(Quantize,         QuantizeLayer);
     CV_DNN_REGISTER_LAYER_CLASS(Dequantize,       DequantizeLayer);

modules/dnn/src/layers/einsum_layer.cpp

@@ -459,6 +459,7 @@ public:
     {
         CV_TRACE_FUNCTION();
         CV_TRACE_ARG_VALUE(name, "name", name.c_str());
+        CV_CheckEQ((size_t)inputs_arr.total(), (size_t)numInputs, "Number of inputs in forward and inputs during graph constructions do not match");
 
         if (inputs_arr.depth() == CV_16F)
         {
@@ -541,7 +542,7 @@ public:
                 // Use either the preprocessed inputs (if it is available) or the corresponding raw inputs
                 result = pairwiseOperandProcess(!result.empty() ? result : rawInputs[0],
                                                 !result.empty() ? tmpResult : homogenizedInputDims[0],
-                                                (!preProcessedInputs.empty() && !preProcessedInputs[input].empty()) ? preProcessedInputs[input] : rawInputs[input],
+                                                (!preProcessedInputs[input].empty()) ? preProcessedInputs[input] : rawInputs[input],
                                                 homogenizedInputDims[input],
                                                 reducedDims,
                                                 isFinalPair);
@@ -605,8 +606,8 @@ void LayerEinsumImpl::preProcessInputs(InputArrayOfArrays& inputs_arr)
     std::vector<cv::Mat> inputs;
     inputs_arr.getMatVector(inputs);
 
-    preProcessedInputs.reserve(inputs.size());
-    homogenizedInputDims.reserve(inputs.size());
+    preProcessedInputs.resize(inputs.size());
+    homogenizedInputDims.resize(inputs.size());
 
     int inputIter = 0;
     for(const Mat& input : inputs)
@@ -616,7 +617,7 @@ void LayerEinsumImpl::preProcessInputs(InputArrayOfArrays& inputs_arr)
         // variable to hold processed version of the original input
         MatShape input_dims = shape(input);
         if (input_dims.empty()){
-            homogenizedInputDims.emplace_back(MatShape(numLetterIndices, 1));
+            homogenizedInputDims[inputIter] = MatShape(numLetterIndices, 1);
             ++inputIter;
             continue;
         }
@@ -672,9 +673,9 @@ void LayerEinsumImpl::preProcessInputs(InputArrayOfArrays& inputs_arr)
         {
             preprocessed = preprocessed.reshape(1, homogenizedInputDims_.size(), homogenizedInputDims_.data());
         }
+        preProcessedInputs[inputIter] = preprocessed;
+        homogenizedInputDims[inputIter] = homogenizedInputDims_;
 
-        preProcessedInputs.emplace_back(preprocessed);
-        homogenizedInputDims.emplace_back(homogenizedInputDims_);
         ++inputIter;
     }
 }

modules/dnn/src/layers/elementwise_layers.cpp

@@ -1520,10 +1520,10 @@ struct RoundFunctor : public BaseDefaultFunctor<RoundFunctor>
     inline float calculate(float x) const
     {
         // Rounds to even numbers in halfway cases, so 2.5 -> 2, -2.5 -> -2
-        int old_rounding_direction = std::fegetround();
-        std::fesetround(FE_TONEAREST);
+        int old_rounding_direction = fegetround();
+        fesetround(FE_TONEAREST);
         float y = std::nearbyint(x);
-        std::fesetround(old_rounding_direction);
+        fesetround(old_rounding_direction);
         return y;
     }
 

modules/dnn/src/layers/topk_layer.cpp

@@ -0,0 +1,228 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+
+#include "../precomp.hpp"
+#include "layers_common.hpp"
+
+#include <opencv2/dnn/shape_utils.hpp>
+
+namespace cv { namespace dnn {
+
+namespace {
+
+template<typename T>
+class ComparatorGreater {
+public:
+    ComparatorGreater(const T* data, size_t step)
+        : data_(data), step_(step) {}
+
+    void addOffset(size_t offset) {
+        data_ += offset;
+    }
+
+    void minusOffset(size_t offset) {
+        data_ -= offset;
+    }
+
+    bool operator()(const size_t lhs_idx, const size_t rhs_idx) {
+        T lhs = *(data_ + lhs_idx * step_),
+          rhs = *(data_ + rhs_idx * step_);
+        return (lhs > rhs || (lhs == rhs && lhs_idx < rhs_idx));
+    }
+
+private:
+    const T* data_;
+    size_t step_;
+};
+
+template<typename T>
+class ComparatorLess {
+public:
+    ComparatorLess(const T* data, size_t step)
+        : data_(data), step_(step) {}
+
+    void addOffset(size_t offset) {
+        data_ += offset;
+    }
+
+    void minusOffset(size_t offset) {
+        data_ -= offset;
+    }
+
+    bool operator()(const size_t lhs_idx, const size_t rhs_idx) {
+        T lhs = *(data_ + lhs_idx * step_),
+          rhs = *(data_ + rhs_idx * step_);
+        return (lhs < rhs || (lhs == rhs && lhs_idx < rhs_idx));
+    }
+
+private:
+    const T* data_;
+    size_t step_;
+};
+}
+
+class TopKLayerImpl CV_FINAL : public TopKLayer
+{
+public:
+    TopKLayerImpl(const LayerParams& params)
+    {
+        setParamsFrom(params);
+
+        axis = params.get<int>("axis", -1);
+        largest = params.get<int>("largest", 1) == 1;
+        sorted = params.get<int>("sorted", 1) == 1;
+        CV_CheckTrue(sorted, "TopK: sorted == false is not supported"); // TODO: support sorted
+
+        CV_CheckTrue(params.has("k"), "TopK: parameter k is required but missing");
+        K = params.get<int>("k");
+    }
+
+    virtual bool supportBackend(int backendId) CV_OVERRIDE
+    {
+        return backendId == DNN_BACKEND_OPENCV;
+    }
+
+    virtual bool getMemoryShapes(const std::vector<MatShape> &inputs,
+                                 const int requiredOutputs,
+                                 std::vector<MatShape> &outputs,
+                                 std::vector<MatShape> &internals) const CV_OVERRIDE
+    {
+        const auto &input_shape = inputs.front();
+        int input_dims = input_shape.size();
+
+        // Check if axis is valid
+        CV_CheckGE(axis, -input_dims, "TopK: axis is out of range");
+        CV_CheckLT(axis, input_dims, "TopK: axis is out of range");
+        // Normalize axis
+        int axis_normalized = normalize_axis(axis, input_shape.size());
+
+        // Check if K is in range (0, input_shape[axis])
+        CV_CheckGT(K, 0, "TopK: K needs to be a positive integer");
+        CV_CheckLT(K, input_shape[axis_normalized], "TopK: K is out of range");
+
+        // Assign output shape
+        auto output_shape = input_shape;
+        output_shape[axis_normalized] = K;
+        outputs.assign(1, output_shape);
+        outputs.assign(2, output_shape); // TODO: support indices of type CV_32S on 5.x
+
+        return false;
+    }
+
+    virtual void finalize(InputArrayOfArrays inputs_arr, OutputArrayOfArrays outputs_arr) CV_OVERRIDE {
+        std::vector<Mat> inputs;
+        inputs_arr.getMatVector(inputs);
+
+        // Normalize axis
+        auto input_shape = shape(inputs.front());
+        axis = normalize_axis(axis, input_shape.size());
+    }
+
+    template<class Comparator>
+    void FindTopK(const Mat &input, Mat &output_value, Mat &output_index) {
+        const auto input_shape = shape(input);
+        size_t loops = std::accumulate(input_shape.begin(), input_shape.begin() + axis, 1, std::multiplies<int>());
+        size_t step = std::accumulate(input_shape.begin() + axis + 1, input_shape.end(), 1, std::multiplies<int>());
+        int dim_axis = input_shape[axis];
+        if (loops == 1) {
+            auto worker = [&](const Range &r) {
+                const auto *input_ptr = input.ptr<const float>();   // TODO: support other input type
+                auto *output_value_ptr = output_value.ptr<float>();
+                auto *output_index_ptr = output_index.ptr<float>(); // TODO: use CV_32S on 5.x
+
+                Comparator cmp(input_ptr, step);
+
+                AutoBuffer<int> buffer_index(dim_axis);
+                auto *buffer_index_ptr = buffer_index.data();
+                for (int offset = r.start; offset < r.end; offset++) {
+                    const auto *input_offset_ptr = input_ptr + offset;
+                    cmp.addOffset(offset);
+
+                    std::iota(buffer_index_ptr, buffer_index_ptr + dim_axis, 0);
+                    std::stable_sort(buffer_index_ptr, buffer_index_ptr + dim_axis, cmp);
+
+                    auto *output_value_offset_ptr = output_value_ptr + offset;
+                    auto *output_index_offset_ptr = output_index_ptr + offset;
+                    for (int i = 0; i < K; i++) {
+                        int source_index = buffer_index_ptr[i];
+                        output_value_offset_ptr[i * step] = *(input_offset_ptr + source_index * step);
+                        output_index_offset_ptr[i * step] = source_index;
+                    }
+                    cmp.minusOffset(offset);
+                }
+            };
+            parallel_for_(Range(0, step), worker);
+        } else {
+            auto worker = [&](const Range &r) {
+                const auto *input_ptr = input.ptr<const float>();
+                auto *output_value_ptr = output_value.ptr<float>();
+                auto *output_index_ptr = output_index.ptr<float>();
+
+                Comparator cmp(input_ptr, step);
+
+                AutoBuffer<int> buffer_index(dim_axis);
+                auto *buffer_index_ptr = buffer_index.data();
+                for (int batch_index = r.start; batch_index < r.end; batch_index++) {
+                    for (size_t offset = 0; offset < step; offset++) {
+                        const auto *input_offset_ptr = input_ptr + batch_index * dim_axis * step + offset;
+                        cmp.addOffset(batch_index * dim_axis * step + offset);
+
+                        std::iota(buffer_index_ptr, buffer_index_ptr + dim_axis, 0);
+                        std::stable_sort(buffer_index_ptr, buffer_index_ptr + dim_axis, cmp);
+
+                        auto *output_value_offset_ptr = output_value_ptr + batch_index * K * step + offset;
+                        auto *output_index_offset_ptr = output_index_ptr + batch_index * K * step + offset;
+                        for (int i = 0; i < K; i++) {
+                            int source_index = buffer_index_ptr[i];
+                            output_value_offset_ptr[i * step] = *(input_offset_ptr + source_index * step);
+                            output_index_offset_ptr[i * step] = source_index;
+                        }
+                        cmp.minusOffset(batch_index * dim_axis * step + offset);
+                    }
+                }
+            };
+            parallel_for_(Range(0, loops), worker);
+        }
+    }
+
+    void forward(InputArrayOfArrays inputs_arr, OutputArrayOfArrays outputs_arr, OutputArrayOfArrays internals_arr) CV_OVERRIDE
+    {
+        CV_TRACE_FUNCTION();
+        CV_TRACE_ARG_VALUE(name, "name", name.c_str());
+
+        if (inputs_arr.depth() == CV_16F)
+        {
+            forward_fallback(inputs_arr, outputs_arr, internals_arr);
+            return;
+        }
+
+        std::vector<Mat> inputs, outputs;
+        inputs_arr.getMatVector(inputs);
+        outputs_arr.getMatVector(outputs);
+
+        const auto &input = inputs.front();
+        auto &output_value = outputs.front();
+        auto &output_index = outputs.back();
+
+        if (largest) {
+            FindTopK<ComparatorGreater<float>>(input, output_value, output_index);
+        } else {
+            FindTopK<ComparatorLess<float>>(input, output_value, output_index);
+        }
+    }
+
+private:
+    int axis;
+    bool largest;
+    bool sorted;
+
+    int K; // FIXIT: make it layer input once dynamic shape is supported
+};
+
+Ptr<TopKLayer> TopKLayer::create(const LayerParams& params)
+{
+    return makePtr<TopKLayerImpl>(params);
+}
+
+}} // namespace cv::dnn

modules/dnn/src/onnx/onnx_importer.cpp

@@ -195,6 +195,7 @@ private:
     void parseScatter              (LayerParams& layerParams, const opencv_onnx::NodeProto& node_proto);
     void parseTile                 (LayerParams& layerParams, const opencv_onnx::NodeProto& node_proto);
     void parseLayerNorm            (LayerParams& layerParams, const opencv_onnx::NodeProto& node_proto);
+    void parseTopK                 (LayerParams& LayerParams, const opencv_onnx::NodeProto& node_proto);
     void parseSimpleLayers         (LayerParams& layerParams, const opencv_onnx::NodeProto& node_proto);
     void parseEinsum               (LayerParams& layerParams, const opencv_onnx::NodeProto& node_proto);
 
@@ -3162,6 +3163,21 @@ void ONNXImporter::parseLayerNorm(LayerParams& layerParams, const opencv_onnx::N
     }
 }
 
+void ONNXImporter::parseTopK(LayerParams& layerParams, const opencv_onnx::NodeProto& node_proto)
+{
+    // K needs to be constant in case of being input (since opset 10)
+    if (node_proto.input_size() == 2) {
+        bool K_const = constBlobs.find(node_proto.input(1)) != constBlobs.end();
+        CV_CheckTrue(K_const, "OnnxImporter/TopK: K being non-constant is not supported");
+
+        Mat input_K = getBlob(node_proto, 1);
+        int K = input_K.at<int>(0);
+        layerParams.set("k", K);
+    }
+
+    addLayer(layerParams, node_proto);
+}
+
 void ONNXImporter::parseSimpleLayers(LayerParams& layerParams, const opencv_onnx::NodeProto& node_proto)
 {
     bool is_all_input_const = true;
@@ -3972,6 +3988,7 @@ void ONNXImporter::buildDispatchMap_ONNX_AI(int opset_version)
     dispatch["Tile"] = &ONNXImporter::parseTile;
     dispatch["LayerNormalization"] = &ONNXImporter::parseLayerNorm;
     dispatch["GroupNormalization"] = &ONNXImporter::parseInstanceNormalization;
+    dispatch["TopK"] = &ONNXImporter::parseTopK;
 
     dispatch["Equal"] = dispatch["Greater"] = dispatch["Less"] = dispatch["Pow"] = dispatch["Add"] =
             dispatch["Sub"] = dispatch["Mul"] = dispatch["Div"] = dispatch["GreaterOrEqual"] =

modules/dnn/test/test_onnx_conformance_layer_filter_opencv_ocl_fp16_denylist.inl.hpp

@@ -2,20 +2,8 @@
 "test_dequantizelinear",
 "test_dequantizelinear_axis",
 "test_dequantizelinear_blocked",
-"test_dropout_default_ratio",
-"test_globalmaxpool",
-"test_globalmaxpool_precomputed",
 "test_logsoftmax_large_number",
 "test_logsoftmax_large_number_expanded",
-"test_maxpool_1d_default",
-"test_maxpool_2d_ceil",
-"test_maxpool_2d_default",
-"test_maxpool_2d_pads",
-"test_maxpool_2d_precomputed_pads",
-"test_maxpool_2d_precomputed_same_upper",
-"test_maxpool_2d_precomputed_strides",
-"test_maxpool_2d_same_upper",
-"test_maxpool_2d_strides",
 "test_maxpool_3d_default",
 "test_pow",
 "test_quantizelinear",
@@ -23,12 +11,7 @@
 "test_quantizelinear_blocked",
 "test_softmax_large_number",
 "test_softmax_large_number_expanded",
-"test_split_equal_parts_1d",
-"test_split_equal_parts_2d",
-"test_split_equal_parts_default_axis",
 "test_tan",
-"test_reduce_l2_default_axes_keepdims_example", // Expected: (normL1) <= (l1), actual: 0.00490189 vs 0.004
-"test_reduce_log_sum_exp_default_axes_keepdims_example", // Expected: (normL1) <= (l1), actual: 0.00671387 vs 0.004
 "test_reduce_prod_default_axes_keepdims_example", // Expected: (normL1) <= (l1), actual: inf vs 0.004
 "test_reduce_prod_default_axes_keepdims_random", // Expected: (normL1) <= (l1), actual: 18.6621 vs 0.004, Expected: (normInf) <= (lInf), actual: 18.6621 vs 0.02
 "test_reduce_prod_do_not_keepdims_random", // Expected: (normL1) <= (l1), actual: 0.00436729 vs 0.004, Expected: (normInf) <= (lInf), actual: 0.0201836 vs 0.02
@@ -38,16 +21,3 @@
 "test_reduce_sum_square_do_not_keepdims_random", // Expected: (normL1) <= (l1), actual: 0.010789 vs 0.004, Expected: (normInf) <= (lInf), actual: 0.0290298 vs 0.02
 "test_reduce_sum_square_keepdims_random", // Expected: (normL1) <= (l1), actual: 0.010789 vs 0.004, Expected: (normInf) <= (lInf), actual: 0.0290298 vs 0.02
 "test_reduce_sum_square_negative_axes_keepdims_random", // Expected: (normL1) <= (l1), actual: 0.010789 vs 0.004, Expected: (normInf) <= (lInf), actual: 0.0290298 vs 0.02
-"test_scatter_elements_with_axis",
-"test_scatter_elements_with_duplicate_indices",
-"test_scatter_elements_with_negative_indices",
-"test_scatter_elements_with_reduction_max",
-"test_scatter_elements_with_reduction_min",
-"test_scatter_elements_without_axis",
-"test_scatter_with_axis",
-"test_scatter_without_axis",
-"test_scatternd",
-"test_scatternd_add",
-"test_scatternd_max",
-"test_scatternd_min",
-"test_scatternd_multiply",

modules/dnn/test/test_onnx_conformance_layer_filter_opencv_ocl_fp32_denylist.inl.hpp

@@ -6,16 +6,3 @@
 "test_quantizelinear",
 "test_quantizelinear_axis",
 "test_quantizelinear_blocked",
-"test_scatter_elements_with_axis",
-"test_scatter_elements_with_duplicate_indices",
-"test_scatter_elements_with_negative_indices",
-"test_scatter_elements_with_reduction_max",
-"test_scatter_elements_with_reduction_min",
-"test_scatter_elements_without_axis",
-"test_scatter_with_axis",
-"test_scatter_without_axis",
-"test_scatternd",
-"test_scatternd_add",
-"test_scatternd_max",
-"test_scatternd_min",
-"test_scatternd_multiply",

modules/dnn/test/test_onnx_importer.cpp

@@ -1019,6 +1019,10 @@ TEST_P(Test_ONNX_layers, MatMul_init_bcast)
     testONNXModels("matmul_init_bcast");
 }
 
+TEST_P(Test_ONNX_layers, MatMul_bcast_3dx2d) {
+    testONNXModels("matmul_bcast");
+}
+
 TEST_P(Test_ONNX_layers, MatMulAdd)
 {
 #if defined(INF_ENGINE_RELEASE) && INF_ENGINE_VER_MAJOR_EQ(2022010000)
@@ -3274,6 +3278,40 @@ TEST_P(Test_ONNX_layers, ClipDivSharedConstant) {
     testONNXModels("clip_div_shared_constant");
 }
 
+// Bug: https://github.com/opencv/opencv/issues/26076
+TEST_P(Test_ONNX_layers, DISABLED_TopK) {
+    auto test = [&](const std::string &basename, double l1 = 0, double lInf = 0) {
+        std::string onnxmodel = _tf("models/" + basename + ".onnx", true);
+        Mat input = readTensorFromONNX(_tf("data/input_" + basename + ".pb"));
+        Mat output_ref_val = readTensorFromONNX(_tf("data/output_" + basename + "_0.pb")),
+            output_ref_ind = readTensorFromONNX(_tf("data/output_" + basename + "_1.pb"));
+
+        checkBackend(&input, &output_ref_val);
+        checkBackend(&input, &output_ref_ind);
+        Net net = readNetFromONNX(onnxmodel);
+        net.setPreferableBackend(backend);
+        net.setPreferableTarget(target);
+
+        net.setInput(input);
+        std::vector<Mat> outputs;
+        net.forward(outputs, std::vector<std::string>{"values", "indices"});
+
+        Mat output_res_val = outputs.front(),
+            output_res_ind = outputs.back();
+
+        output_ref_ind.convertTo(output_ref_ind, CV_32F); // TODO: revise this conversion in 5.x
+
+        normAssert(output_ref_val, output_res_val, (basename + " values").c_str(), l1 ? l1 : default_l1, lInf ? lInf : default_lInf);
+        normAssert(output_ref_ind, output_res_ind, (basename + " indices").c_str(), l1 ? l1 : default_l1, lInf ? lInf : default_lInf);
+
+        expectNoFallbacksFromIE(net);
+    };
+
+    test("top_k");
+    test("top_k_negative_axis");
+    test("top_k_smallest");
+}
+
 INSTANTIATE_TEST_CASE_P(/**/, Test_ONNX_nets, dnnBackendsAndTargets());
 
 }} // namespace