cuda support and Layer Split + per-op executors

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

@@ -262,23 +262,93 @@ CV__DNN_INLINE_NS_BEGIN
     class CV_EXPORTS Graph;
     class CV_EXPORTS ActivationLayer;
 
-    /** @brief This interface class allows to build new Layers - are building blocks of networks.
+    /** @brief Backend-independent description of a graph operation (node).
      *
-     * Each class, derived from Layer, must implement forward() method to compute outputs.
-     * Also before using the new layer into networks you must register your layer by using one of @ref dnnLayerFactory "LayerFactory" macros.
+     * %OpData carries everything needed to reason about an operation *without* executing it:
+     * its parameters (#blobs and type-specific fields of derived classes), graph wiring
+     * (#inputs / #outputs as Arg indices) and shape/type/layout inference. The new DNN graph
+     * engine stores a topologically sorted sequence of %OpData nodes (see Graph::prog());
+     * executable, backend-specific instances (Layer subclasses) are constructed from an
+     * %OpData during Net::finalizeNet().
+     *
+     * Each operation type registers a `static Ptr<OpData> create(const LayerParams&)` factory
+     * via @ref CV_DNN_REGISTER_OP_CLASS_STATIC.
      */
-    class CV_EXPORTS_W Layer : public Algorithm
+    class CV_EXPORTS_W OpData : public Algorithm
     {
     public:
+        OpData();
+        explicit OpData(const LayerParams& params);
+        virtual ~OpData();
+
+        void setParamsFrom(const LayerParams& params);
 
         //! List of learned parameters must be stored here to allow read them by using Net::getParam().
         CV_PROP_RW std::vector<Mat> blobs;
         std::vector<Arg> inputs;
         std::vector<Arg> outputs;
-        void* netimpl;
+        void* netimpl = nullptr;
+
+        CV_PROP String name;
+        CV_PROP String type;
 
         virtual std::vector<Ptr<Graph> >* subgraphs() const;
 
+        virtual int inputNameToIndex(String inputName);  // FIXIT const
+        CV_WRAP virtual int outputNameToIndex(const String& outputName);  // FIXIT const
+
+        virtual bool getMemoryShapes(const std::vector<MatShape> &inputs,
+                                     const int requiredOutputs,
+                                     std::vector<MatShape> &outputs,
+                                     std::vector<MatShape> &internals) const;
+
+        virtual void getTypes(const std::vector<MatType>& inputs,
+                              const int requiredOutputs,
+                              const int requiredInternals,
+                              std::vector<MatType>&outputs,
+                              std::vector<MatType>&internals) const;
+
+        virtual int getLayouts(const std::vector<DataLayout>& actualInputs,
+                               std::vector<DataLayout>& desiredInputs,
+                               const int requiredOutputs,
+                               std::vector<DataLayout>& outputs) const;
+
+        virtual int64 getFLOPS(const std::vector<MatShape> &inputs,
+                               const std::vector<MatShape> &outputs) const;
+
+        virtual bool updateMemoryShapes(const std::vector<MatShape> &inputs);
+
+        virtual bool alwaysSupportInplace() const;
+
+        virtual bool dynamicOutputShapes() const;
+
+        virtual bool isDataShuffling() const;
+
+        virtual void getScaleShift(Mat& scale, Mat& shift) const;
+
+        virtual void getScaleZeropoint(float& scale, int& zeropoint) const;
+
+        virtual std::ostream& dumpAttrs(std::ostream& strm, int indent) const;
+
+        virtual std::ostream& dump(std::ostream& strm, int indent, bool comma) const;
+    };
+
+    /** @brief This interface class allows to build new Layers - are building blocks of networks.
+     *
+     * A %Layer is the *executable*, backend-specific counterpart of an OpData node: it
+     * implements forward() (and finalize()) for a particular backend/target. In the new graph
+     * engine a %Layer is created from an OpData (held in #data) by Net::finalizeNet(); its
+     * inference methods (getMemoryShapes() etc., inherited from OpData) delegate to #data.
+     *
+     * Each class, derived from Layer, must implement forward() method to compute outputs.
+     * Also before using the new layer into networks you must register your layer by using one of @ref dnnLayerFactory "LayerFactory" macros.
+     */
+    class CV_EXPORTS_W Layer : public OpData
+    {
+    public:
+
+        Ptr<OpData> data;
+
         /** @brief Computes and sets internal parameters according to inputs, outputs and blobs.
          *  @deprecated Use Layer::finalize(InputArrayOfArrays, OutputArrayOfArrays) instead
          *  @param[in]  input  vector of already allocated input blobs
@@ -341,18 +411,6 @@ CV__DNN_INLINE_NS_BEGIN
         CV_DEPRECATED CV_WRAP void run(const std::vector<Mat> &inputs, CV_OUT std::vector<Mat> &outputs,
                                        CV_IN_OUT std::vector<Mat> &internals);
 
-        /** @brief Returns index of input blob into the input array.
-         *  @param inputName label of input blob
-         *
-         * Each layer input and output can be labeled to easily identify them using "%<layer_name%>[.output_name]" notation.
-         * This method maps label of input blob to its index into input vector.
-         */
-        virtual int inputNameToIndex(String inputName);  // FIXIT const
-        /** @brief Returns index of output blob in output array.
-         *  @see inputNameToIndex()
-         */
-        CV_WRAP virtual int outputNameToIndex(const String& outputName);  // FIXIT const
-
         /**
          * @brief Ask layer if it support specific backend for doing computations.
          * @param[in] backendId computation backend identifier.
@@ -419,102 +477,25 @@ CV__DNN_INLINE_NS_BEGIN
         virtual bool tryFuse(Ptr<Layer>& top);
 
         /**
-         * @brief Returns parameters of layers with channel-wise multiplication and addition.
-         * @param[out] scale Channel-wise multipliers. Total number of values should
-         *                   be equal to number of channels.
-         * @param[out] shift Channel-wise offsets. Total number of values should
-         *                   be equal to number of channels.
+         * @brief Executes the operation on the CUDA backend (new graph engine).
          *
-         * Some layers can fuse their transformations with further layers.
-         * In example, convolution + batch normalization. This way base layer
-         * use weights from layer after it. Fused layer is skipped.
-         * By default, @p scale and @p shift are empty that means layer has no
-         * element-wise multiplications or additions.
+         * Called by the engine for nodes assigned to DNN_BACKEND_CUDA. The default
+         * implementation raises an error. @p workspace is an opaque pointer to a
+         * cuda4dnn::csl::Workspace (kept void* to avoid leaking internal CUDA types).
          */
-        virtual void getScaleShift(Mat& scale, Mat& shift) const;
-
-        /**
-         * @brief Returns scale and zeropoint of layers
-         * @param[out] scale Output scale
-         * @param[out] zeropoint Output zeropoint
-         *
-         * By default, @p scale is 1 and @p zeropoint is 0.
-         */
-        virtual void getScaleZeropoint(float& scale, int& zeropoint) const;
-
+        virtual void forwardCUDA(const std::vector<Ptr<BackendWrapper> >& inputs,
+                                 const std::vector<Ptr<BackendWrapper> >& outputs,
+                                 void* workspace);
 
         /**
          * @brief "Detaches" all the layers, attached to particular layer.
          */
         virtual void unsetAttached();
 
-        virtual bool getMemoryShapes(const std::vector<MatShape> &inputs,
-                                     const int requiredOutputs,
-                                     std::vector<MatShape> &outputs,
-                                     std::vector<MatShape> &internals) const;
-
-        virtual void getTypes(const std::vector<MatType>& inputs,
-                              const int requiredOutputs,
-                              const int requiredInternals,
-                              std::vector<MatType>&outputs,
-                              std::vector<MatType>&internals) const;
-
-        // this is the method for Layer to express its attitude to the block layout
-        // or any other special form of layout. It takes
-        // layouts of the inputs and should return the desired layouts of
-        // inputs, as well as layouts of the outputs.
-        // By default, no mater what the actual inputs' layouts are,
-        // the desired inputs as well as outputs will get 'Unknown' layout values.
-        // It means that the layer can only handle non-block layout
-        // (depending on the model format, e.g. NCHW for ONNX or NHWC for TFLite)
-        // and will return tensors with non-block layout as well.
-        // Some layers could override this default behaviour:
-        // a) if they _can_ process block-layout data, like element-wise operations, or
-        // b) if they _need_ block-layout data, like convolution
-        virtual int getLayouts(const std::vector<DataLayout>& actualInputs,
-                                std::vector<DataLayout>& desiredInputs,
-                                const int requiredOutputs,
-                                std::vector<DataLayout>& outputs) const;
-
-        virtual int64 getFLOPS(const std::vector<MatShape> &inputs,
-                               const std::vector<MatShape> &outputs) const;
-
-        virtual bool updateMemoryShapes(const std::vector<MatShape> &inputs);
-
-        // returns true if the operation takes a single input and can always be performed in-place,
-        // assuming that the input is contiguous.
-        // Examples of such operations are: Reshape, Flatten, Squeeze, Unsqueeze,
-        // as well many unary element-wise operations (ReLU, Tanh, ...)
-        virtual bool alwaysSupportInplace() const;
-
-        // returns false if the shape of Layer outputs is defined only by the shapes of inputs.
-        // Sometimes the shape depends on the content of the input(s), then the method should return true.
-        // In such a rare case forward() method should take care of proper allocation of the output tensors.
-        // On the other hand, when this method returns false, the engine takes care of proper allocation of the outputs,
-        // so that forward() can assume that the outputs are already allocated.
-        virtual bool dynamicOutputShapes() const;
-
-        // returns true if the layer only rearranges data without changing values.
-        // Examples: Flatten, Reshape, Transpose, Permute, Squeeze, Unsqueeze,
-        // Concat, Split, Slice, Tile, MaxPool.
-        // Used by QDQ fusion to elide redundant dequantize-quantize pairs
-        // when the scale and zero point are the same.
-        virtual bool isDataShuffling() const;
-
-        // dumps attributes of the layer (e.g. strides, dilations in Convolution, MaxPool)
-        virtual std::ostream& dumpAttrs(std::ostream& strm, int indent) const;
-
-        // dumps information about the layer. The default implementation is usually good enough,
-        // just override dumpAttrs().
-        virtual std::ostream& dump(std::ostream& strm, int indent, bool comma) const;
-
-        CV_PROP String name; //!< Name of the layer instance, can be used for logging or other internal purposes.
-        CV_PROP String type; //!< Type name which was used for creating layer by layer factory.
         CV_PROP int preferableTarget; //!< prefer target for layer forwarding
 
         Layer();
         explicit Layer(const LayerParams &params);      //!< Initializes only #name, #type and #blobs fields.
-        void setParamsFrom(const LayerParams &params);  //!< Initializes only #name, #type and #blobs fields.
         virtual ~Layer();
     };
 
@@ -533,15 +514,16 @@ CV__DNN_INLINE_NS_BEGIN
         virtual bool empty() const = 0;
         virtual void clear() = 0;
         virtual std::string name() const = 0;
-        virtual const std::vector<Arg>& append(Ptr<Layer>& layer,
+        virtual const std::vector<Arg>& append(Ptr<OpData>& op,
                     const std::vector<std::string>& outnames=std::vector<std::string>()) = 0;
-        virtual Arg append(Ptr<Layer>& layer, const std::string& outname=std::string()) = 0;
+        virtual Arg append(Ptr<OpData>& op, const std::string& outname=std::string()) = 0;
         virtual std::ostream& dump(std::ostream& strm, int indent, bool comma) = 0;
         virtual const std::vector<Arg>& inputs() const = 0;
         virtual const std::vector<Arg>& outputs() const = 0;
         virtual void setOutputs(const std::vector<Arg>& outputs) = 0;
-        virtual const std::vector<Ptr<Layer> >& prog() const = 0;
-        virtual void setProg(const std::vector<Ptr<Layer> >& newprog) = 0;
+        virtual const std::vector<Ptr<OpData> >& prog() const = 0;
+        virtual void setProg(const std::vector<Ptr<OpData> >& newprog) = 0;
+https://github.com/opencv/opencv/pull/29018        virtual int opBackend(int opidx) const = 0;
     };
 
     /** @brief This class allows to create and manipulate comprehensive artificial neural networks.

modules/dnn/include/opencv2/dnn/layer.details.hpp

@@ -45,6 +45,24 @@ Ptr<Layer> __LayerStaticRegisterer_func_##type(LayerParams &params) \
     { return Ptr<Layer>(new class(params)); }                       \
 static cv::dnn::details::_LayerStaticRegisterer __LayerStaticRegisterer_##type(#type, __LayerStaticRegisterer_func_##type);
 
+/** @brief Registers an OpData (metadata node) class for the new graph engine.
+ *  @param type string, containing the operation type name.
+ *  @param class C++ class derived from OpData, providing `static Ptr<OpData> create(const LayerParams&)`.
+ *  @details This macro must be placed inside the function code (e.g. initializeLayerFactory()).
+ */
+#define CV_DNN_REGISTER_OP_CLASS(type, class) \
+    cv::dnn::LayerFactory::registerOp(#type, cv::dnn::details::_opDynamicRegisterer<class>);
+
+/** @brief Registers a backend executor class for the new graph engine.
+ *  @param type string, containing the operation type name.
+ *  @param backendId backend id the executor targets (e.g. DNN_BACKEND_OPENCV, DNN_BACKEND_CUDA).
+ *  @param class C++ class derived from Layer, providing
+ *         `static Ptr<Layer> create(const Ptr<OpData>&, void* backendCtx)` (null Ptr if unsupported).
+ *  @details This macro must be placed inside the function code.
+ */
+#define CV_DNN_REGISTER_EXEC_CLASS(type, backendId, class) \
+    cv::dnn::LayerFactory::registerExec(#type, backendId, cv::dnn::details::_execDynamicRegisterer<class>);
+
 namespace details {
 
 template<typename LayerClass>
@@ -53,6 +71,18 @@ Ptr<Layer> _layerDynamicRegisterer(LayerParams &params)
     return Ptr<Layer>(LayerClass::create(params));
 }
 
+template<typename OpClass>
+Ptr<OpData> _opDynamicRegisterer(const LayerParams &params)
+{
+    return Ptr<OpData>(OpClass::create(params));
+}
+
+template<typename ExecClass>
+Ptr<Layer> _execDynamicRegisterer(const Ptr<OpData>& data, void* backendCtx)
+{
+    return Ptr<Layer>(ExecClass::create(data, backendCtx));
+}
+
 //allows automatically register created layer on module load time
 class _LayerStaticRegisterer
 {

modules/dnn/include/opencv2/dnn/layer.hpp

@@ -76,6 +76,19 @@ public:
      */
     static Ptr<Layer> createLayerInstance(const String &type, LayerParams& params);
 
+
+    // Builds the abstract (metadata) node for an operation type.
+    typedef Ptr<OpData>(*OpConstructor)(const LayerParams& params);
+    //! Builds a backend-specific executor from an OpData; returns null Ptr if unsupported.
+    typedef Ptr<Layer>(*ExecConstructor)(const Ptr<OpData>& data, void* backendCtx);
+
+    static void registerOp(const String& type, OpConstructor constructor);
+    static Ptr<OpData> createOp(const String& type, const LayerParams& params);
+
+    static void registerExec(const String& type, int backendId, ExecConstructor constructor);
+    static Ptr<Layer> createExec(const String& type, int backendId,
+                                 const Ptr<OpData>& data, void* backendCtx);
+
 private:
     LayerFactory();
 };

modules/dnn/src/cuda4dnn/csl/cudnn/convolution.hpp

@@ -227,11 +227,11 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace csl { namespace cu
                 CUDA4DNN_CHECK_CUDNN(cudnnSetConvolutionGroupCount(descriptor, group_count));
 
 #if CUDNN_MAJOR >= 8
-                /* cuDNN 7 and below use FMA math by default. cuDNN 8 includes TF32 Tensor Ops
-                 * in the default setting. TF32 convolutions have lower precision than FP32.
-                 * Hence, we set the math type to CUDNN_FMA_MATH to reproduce old behavior.
+                /* cuDNN 8 default math includes TF32 Tensor Ops for FP32 convolutions (Ampere+),
+                 * giving near-FP16 throughput at slightly reduced precision. ONNXRuntime enables
+                 * this by default; we follow suit. (Was CUDNN_FMA_MATH to force exact FP32.)
                  */
-                CUDA4DNN_CHECK_CUDNN(cudnnSetConvolutionMathType(descriptor, CUDNN_FMA_MATH));
+                CUDA4DNN_CHECK_CUDNN(cudnnSetConvolutionMathType(descriptor, CUDNN_DEFAULT_MATH));
 #endif
 
                 if (std::is_same<T, half>::value)

modules/dnn/src/graph_block_layout.cpp

@@ -11,7 +11,7 @@ CV__DNN_INLINE_NS_BEGIN
 using std::vector;
 using std::string;
 
-using PLayer = Ptr<Layer>;
+using PLayer = Ptr<OpData>;
 using PGraph = Ptr<Graph>;
 
 /* Inserts layout conversion operations (if needed) into the model graph and subgraphs.
@@ -121,13 +121,15 @@ struct BlockLayoutTransformer
         std::vector<DataLayout> inputLayoutsOrig, inputLayoutsNew, outputLayouts;
         size_t nchanges = 0;
 
-        for (const PLayer& layer: currProg) {
+        for (size_t opidx = 0; opidx < currProg.size(); opidx++) {
+            const PLayer& layer = currProg[opidx];
             const vector<Arg>& inputs = layer->inputs;
             const vector<Arg>& outputs = layer->outputs;
             size_t ninputs = inputs.size(), noutputs = outputs.size();
             std::string op_name = layer->type;
             std::string name = layer->name;
             vector<PGraph>* subgraphs = layer->subgraphs();
+\            bool deviceOp = g->opBackend((int)opidx) != DNN_BACKEND_OPENCV;
             //std::cout << "name: " << name << ", op_name: " << op_name << ", inp0 layout: " << layoutToString(layouts[inputs[0].idx]) << "\n";
 
             if (subgraphs) {
@@ -164,6 +166,13 @@ struct BlockLayoutTransformer
             CV_Assert(inputLayoutsNew.size() == ninputs);
             CV_Assert(outputLayouts.size() == noutputs);
 
+            if (deviceOp) {
+                for (size_t i = 0; i < ninputs; i++)
+                    inputLayoutsNew[i] = defaultLayout;
+                for (size_t i = 0; i < noutputs; i++)
+                    outputLayouts[i] = defaultLayout;
+            }
+
             newInputs.clear();
             bool changedInputs = false;
             for (size_t i = 0; i < ninputs; i++) {

modules/dnn/src/graph_buffer_allocator.cpp

@@ -226,7 +226,7 @@ struct BufferAllocator
                 }
             }
         }
-        const std::vector<Ptr<Layer> >& prog = graph->prog();
+        const std::vector<Ptr<OpData> >& prog = graph->prog();
         for (const auto& layer: prog) {
             bool inplace = false;
             Arg reuseArg;

modules/dnn/src/graph_const_args.cpp

@@ -36,14 +36,14 @@ struct ConstArgs
 
     void processGraph(Ptr<Graph>& graph)
     {
-        const std::vector<Ptr<Layer> >& prog = graph->prog();
+        const std::vector<Ptr<OpData> >& prog = graph->prog();
         size_t i, nops = prog.size();
         std::vector<Arg> removed_args;
         std::vector<Arg> saved_tail_inputs;
 
         for (i = 0; i < nops; i++) {
-            const Ptr<Layer>& layer = prog[i];
-            Layer* layer_ptr = const_cast<Layer*>(layer.get());
+            const Ptr<OpData>& layer = prog[i];
+            OpData* layer_ptr = const_cast<OpData*>(layer.get());
             std::vector<Ptr<Graph> >* subgraphs = layer->subgraphs();
             if (subgraphs) {
                 for (Ptr<Graph>& g: *subgraphs) {

modules/dnn/src/graph_const_fold.cpp

@@ -30,7 +30,7 @@ struct ConstFolding
         netimpl->scratchBufs.clear();
     }
 
-    Layer* getLayer(std::vector<Ptr<Layer> >& newprog, int op_idx) const
+    OpData* getLayer(std::vector<Ptr<OpData> >& newprog, int op_idx) const
     {
         return op_idx >= 0 ? newprog.at(op_idx).get() : 0;
     }
@@ -47,16 +47,16 @@ struct ConstFolding
     {
         netimpl->scratchBufs.clear();
         bool modified = false;
-        const std::vector<Ptr<Layer> >& prog = graph->prog();
+        const std::vector<Ptr<OpData> >& prog = graph->prog();
         size_t i, nops = prog.size();
-        std::vector<Ptr<Layer> > newprog;
+        std::vector<Ptr<OpData> > newprog;
         std::vector<Arg> removed_args;
         std::vector<Mat> inpMats, tempMats;
         std::vector<int> inpTypes, outTypes, tempTypes;
         std::vector<MatShape> inpShapes, outShapes, tempShapes;
 
         for (i = 0; i < nops; i++) {
-            const Ptr<Layer>& layer = prog[i];
+            const Ptr<OpData>& layer = prog[i];
             std::vector<Ptr<Graph> >* subgraphs = layer->subgraphs();
             if (subgraphs) {
                 for (Ptr<Graph>& g: *subgraphs) {
@@ -98,8 +98,10 @@ struct ConstFolding
                     netimpl->allocateLayerOutputs(layer, inpTypes, inpShapes, outTypes,
                                                   outShapes, outOrigData, outMats, tempTypes, tempShapes, tempMats,
                                                   netimpl->scratchBufs, false);
-                layer->finalize(inpMats, outMats);
-                layer->forward(inpMats, outMats, tempMats);
+                Ptr<Layer> execLayer = layer.dynamicCast<Layer>();
+                CV_Assert(execLayer);  // const-folded ops are CPU-executable (monolithic) layers
+                execLayer->finalize(inpMats, outMats);
+                execLayer->forward(inpMats, outMats, tempMats);
                 CV_Assert(outMats.size() == noutputs);
                 for (j = 0; j < noutputs; j++) {
                     Arg out = outputs[j];

modules/dnn/src/graph_fusion_attention.cpp

@@ -32,35 +32,35 @@ struct ModelFusionAttention
         return it->second[0];
     }
 
-    bool isReshape(const vector<Ptr<Layer>>& prog, int idx) const
+    bool isReshape(const vector<Ptr<OpData>>& prog, int idx) const
     {
         if (idx < 0 || idx >= (int)prog.size() || !prog[idx])
             return false;
         return dynamic_cast<Reshape2Layer*>(prog[idx].get()) != nullptr;
     }
 
-    bool isTranspose(const vector<Ptr<Layer>>& prog, int idx) const
+    bool isTranspose(const vector<Ptr<OpData>>& prog, int idx) const
     {
         if (idx < 0 || idx >= (int)prog.size() || !prog[idx])
             return false;
         return dynamic_cast<TransposeLayer*>(prog[idx].get()) != nullptr;
     }
 
-    bool isSoftmax(const vector<Ptr<Layer>>& prog, int idx) const
+    bool isSoftmax(const vector<Ptr<OpData>>& prog, int idx) const
     {
         if (idx < 0 || idx >= (int)prog.size() || !prog[idx])
             return false;
         return prog[idx]->type == "Softmax";
     }
 
-    bool isMatMul(const vector<Ptr<Layer>>& prog, int idx) const
+    bool isMatMul(const vector<Ptr<OpData>>& prog, int idx) const
     {
         if (idx < 0 || idx >= (int)prog.size() || !prog[idx])
             return false;
         return dynamic_cast<MatMulLayer*>(prog[idx].get()) != nullptr;
     }
 
-    static bool isProjCandidate(const Ptr<Layer>& l)
+    static bool isProjCandidate(const Ptr<OpData>& l)
     {
         if (l->blobs.empty() || l->inputs.size() != 1) return false;
         if (dynamic_cast<MatMulLayer*>(l.get()))
@@ -74,7 +74,7 @@ struct ModelFusionAttention
 
     // Returns the projection weight in [K, N] (input_hidden, output_hidden)
     // layout, transposing if the source is a Gemm with trans_b.
-    static Mat getProjWeight(const Ptr<Layer>& l)
+    static Mat getProjWeight(const Ptr<OpData>& l)
     {
         const Mat& W = l->blobs[0];
         GemmLayer* g = dynamic_cast<GemmLayer*>(l.get());
@@ -86,7 +86,7 @@ struct ModelFusionAttention
         return W;
     }
 
-    bool isScalarBinOp(const vector<Ptr<Layer>>& prog, int idx,
+    bool isScalarBinOp(const vector<Ptr<OpData>>& prog, int idx,
                        NaryEltwiseLayer::OPERATION op, float* val) const
     {
         if (idx < 0 || idx >= (int)prog.size() || !prog[idx])
@@ -109,18 +109,18 @@ struct ModelFusionAttention
         return false;
     }
 
-    bool isScalarMul(const vector<Ptr<Layer>>& prog, int idx, float* val) const
+    bool isScalarMul(const vector<Ptr<OpData>>& prog, int idx, float* val) const
     {
         return isScalarBinOp(prog, idx, NaryEltwiseLayer::OPERATION::PROD, val);
     }
 
-    bool isScalarDiv(const vector<Ptr<Layer>>& prog, int idx, float* val) const
+    bool isScalarDiv(const vector<Ptr<OpData>>& prog, int idx, float* val) const
     {
         return isScalarBinOp(prog, idx, NaryEltwiseLayer::OPERATION::DIV, val);
     }
 
     // True if `arg` is produced by the dynamic scale chain Sqrt<-Cast<-Div(1,.)<-Sqrt<-Cast<-Slice<-Shape; visited ops are appended to `chain_ops`.
-    bool isRuntimeQKScaleChain(const vector<Ptr<Layer>>& prog, Arg arg,
+    bool isRuntimeQKScaleChain(const vector<Ptr<OpData>>& prog, Arg arg,
                                 std::set<int>& chain_ops) const
     {
         const std::vector<std::string> expected = {
@@ -133,7 +133,7 @@ struct ModelFusionAttention
             if (it == producer_.end()) return false;
             int idx = it->second;
             if (idx < 0 || idx >= (int)prog.size() || !prog[idx]) return false;
-            const Ptr<Layer>& l = prog[idx];
+            const Ptr<OpData>& l = prog[idx];
             if (want == "NaryEltwise") {
                 NaryEltwiseLayer* elt = dynamic_cast<NaryEltwiseLayer*>(l.get());
                 if (!elt || elt->op != NaryEltwiseLayer::OPERATION::DIV) return false;
@@ -171,7 +171,7 @@ struct ModelFusionAttention
 
     // Accept Add op with exactly two inputs; identify the non-constant runtime
     // input (the mask tensor). Returns false if the Add doesn't match.
-    bool isMaskAdd(const vector<Ptr<Layer>>& prog, int idx, Arg* out_mask) const
+    bool isMaskAdd(const vector<Ptr<OpData>>& prog, int idx, Arg* out_mask) const
     {
         if (idx < 0 || idx >= (int)prog.size() || !prog[idx])
             return false;
@@ -186,7 +186,7 @@ struct ModelFusionAttention
 
     // Extract a scalar integer from a const-valued arg, possibly wrapped in an
     // Unsqueeze of a scalar const. Returns -1 if extraction fails.
-    int extractConstInt(const vector<Ptr<Layer>>& prog, Arg a) const
+    int extractConstInt(const vector<Ptr<OpData>>& prog, Arg a) const
     {
         auto readScalar = [&](Arg x) -> int {
             if (!netimpl->isConstArg(x)) return -1;
@@ -207,7 +207,7 @@ struct ModelFusionAttention
         return -1;
     }
 
-    void collectShapeChain(const vector<Ptr<Layer>>& prog, int concat_idx,
+    void collectShapeChain(const vector<Ptr<OpData>>& prog, int concat_idx,
                            std::set<int>& chain) const
     {
         if (concat_idx < 0 || concat_idx >= (int)prog.size() || !prog[concat_idx])
@@ -238,7 +238,7 @@ struct ModelFusionAttention
     }
 
     template <class Pred>
-    int findMatchingConsumer(const vector<Ptr<Layer>>& prog, Arg out,
+    int findMatchingConsumer(const vector<Ptr<OpData>>& prog, Arg out,
                              Pred pred, std::set<int>* extra_shape_ops) const
     {
         auto it = consumers_.find(out.idx);
@@ -258,7 +258,7 @@ struct ModelFusionAttention
         return matched;
     }
 
-    int followProjChain(const vector<Ptr<Layer>>& prog,
+    int followProjChain(const vector<Ptr<OpData>>& prog,
                         int proj_matmul_idx,
                         int* out_reshape_idx,
                         int* out_num_heads,
@@ -268,7 +268,7 @@ struct ModelFusionAttention
         if (proj_matmul_idx < 0) return -1;
         Arg proj_out = prog[proj_matmul_idx]->outputs[0];
         int reshape_idx = findMatchingConsumer(prog, proj_out,
-            [](Layer* L){ return dynamic_cast<Reshape2Layer*>(L) != nullptr; },
+            [](OpData* L){ return dynamic_cast<Reshape2Layer*>(L) != nullptr; },
             extra_ops_to_remove);
         if (!isReshape(prog, reshape_idx)) return -1;
 
@@ -311,9 +311,9 @@ struct ModelFusionAttention
 
     // Combined-QKV attention: QKV proj -> Reshape ->
     // Transpose -> 3 Gathers -> QK^T -> Softmax(no mask) -> *V.
-    bool tryFuseCombinedQKV(const vector<Ptr<Layer>>& prog, int qkv_matmul_idx,
+    bool tryFuseCombinedQKV(const vector<Ptr<OpData>>& prog, int qkv_matmul_idx,
                             std::set<int>& removed_ops,
-                            vector<std::pair<int, Ptr<Layer>>>& replacements)
+                            vector<std::pair<int, Ptr<OpData>>>& replacements)
     {
         if (qkv_matmul_idx < 0 || qkv_matmul_idx >= (int)prog.size() || !prog[qkv_matmul_idx])
             return false;
@@ -456,7 +456,7 @@ struct ModelFusionAttention
         int out_trans_idx = singleConsumer(prog[av_matmul_idx]->outputs[0]);
         if (!isTranspose(prog, out_trans_idx)) return false;
         int out_reshape_idx = findMatchingConsumer(prog, prog[out_trans_idx]->outputs[0],
-            [](Layer* L){ return dynamic_cast<Reshape2Layer*>(L) != nullptr; },
+            [](OpData* L){ return dynamic_cast<Reshape2Layer*>(L) != nullptr; },
             &extra_ops);
         if (!isReshape(prog, out_reshape_idx)) return false;
 
@@ -487,7 +487,7 @@ struct ModelFusionAttention
         attn_params.blobs.push_back(W_qkv);
         if (has_bias) attn_params.blobs.push_back(bias_qkv);
 
-        Ptr<Layer> attn_layer = LayerFactory::createLayerInstance(attn_params.type, attn_params);
+        Ptr<OpData> attn_layer = LayerFactory::createLayerInstance(attn_params.type, attn_params);
         CV_Assert(attn_layer);
         Arg shared_input = prog[qkv_matmul_idx]->inputs[0];
         attn_layer->inputs  = { shared_input };
@@ -514,7 +514,7 @@ struct ModelFusionAttention
 
     // CLIP-branch trace: arg -> [Transpose3D(K^T)] -> Reshape3D -> Transpose ->
     // Reshape4D -> [Mul(Q scale)] -> [Add(bias)] -> proj_MatMul.
-    int traceClipBranch(const vector<Ptr<Layer>>& prog, Arg arg,
+    int traceClipBranch(const vector<Ptr<OpData>>& prog, Arg arg,
                         bool is_q_branch, bool is_k_branch,
                         Mat& out_W, Mat& out_bias, int& out_num_heads,
                         float& out_q_scale,
@@ -660,9 +660,9 @@ struct ModelFusionAttention
 
     // 3 separate q/k/v projections, Q scaled, K^T at
     // the QK^T matmul, output reshaped+transposed back to (B,S,H*D).
-    bool tryFuseClipAttention(const vector<Ptr<Layer>>& prog, int softmax_idx,
+    bool tryFuseClipAttention(const vector<Ptr<OpData>>& prog, int softmax_idx,
                               std::set<int>& removed_ops,
-                              vector<std::pair<int, Ptr<Layer>>>& replacements)
+                              vector<std::pair<int, Ptr<OpData>>>& replacements)
     {
         if (softmax_idx < 0 || softmax_idx >= (int)prog.size() || !prog[softmax_idx])
             return false;
@@ -782,7 +782,7 @@ struct ModelFusionAttention
         attn_params.blobs.push_back(W_qkv);
         if (!bias_qkv.empty()) attn_params.blobs.push_back(bias_qkv);
 
-        Ptr<Layer> attn_layer =
+        Ptr<OpData> attn_layer =
             LayerFactory::createLayerInstance(attn_params.type, attn_params);
         if (!attn_layer) return false;
         attn_layer->inputs  = { shared_input };
@@ -804,7 +804,7 @@ struct ModelFusionAttention
 
     bool fuseGraph(Ptr<Graph>& graph)
     {
-        const vector<Ptr<Layer>>& prog = graph->prog();
+        const vector<Ptr<OpData>>& prog = graph->prog();
         size_t nops = prog.size();
 
         producer_.clear();
@@ -885,7 +885,7 @@ struct ModelFusionAttention
                 Arg k_tr_out = prog[transpose_idx[k_slot]]->outputs[0];
                 // Tolerate a Shape consumer alongside the Mul/MatMul: the runtime-scale chain (Shape->Slice->Cast->Sqrt...) branches off the Q/K transpose.
                 int k_next = findMatchingConsumer(prog, k_tr_out,
-                    [](Layer* L) {
+                    [](OpData* L) {
                         return dynamic_cast<NaryEltwiseLayer*>(L) != nullptr ||
                                dynamic_cast<MatMulLayer*>(L) != nullptr;
                     },
@@ -939,7 +939,7 @@ struct ModelFusionAttention
 
                     if (vit_style) {
                         int q_next = findMatchingConsumer(prog, q_tr_out,
-                            [](Layer* L) {
+                            [](OpData* L) {
                                 return dynamic_cast<NaryEltwiseLayer*>(L) != nullptr ||
                                        dynamic_cast<MatMulLayer*>(L) != nullptr;
                             },
@@ -1020,7 +1020,7 @@ struct ModelFusionAttention
 
                     Arg out_tr_out = prog[out_transpose_idx]->outputs[0];
                     int out_reshape_idx = findMatchingConsumer(prog, out_tr_out,
-                        [](Layer* L){ return dynamic_cast<Reshape2Layer*>(L) != nullptr; },
+                        [](OpData* L){ return dynamic_cast<Reshape2Layer*>(L) != nullptr; },
                         &extra_ops);
                     if (!isReshape(prog, out_reshape_idx)) continue;
 
@@ -1094,7 +1094,7 @@ struct ModelFusionAttention
                     if (has_bias)
                         attn_params.blobs.push_back(bias_qkv);
 
-                    Ptr<Layer> attn_layer = LayerFactory::createLayerInstance(
+                    Ptr<OpData> attn_layer = LayerFactory::createLayerInstance(
                         attn_params.type, attn_params);
                     CV_Assert(attn_layer);
 
@@ -1157,7 +1157,7 @@ struct ModelFusionAttention
         }
 
         if (modified) {
-            vector<Ptr<Layer>> newprog;
+            vector<Ptr<OpData>> newprog;
             std::sort(attention_replacements_.begin(), attention_replacements_.end(),
                       [](auto& a, auto& b) { return a.first < b.first; });
 
@@ -1183,7 +1183,7 @@ struct ModelFusionAttention
 private:
     std::map<int, int> producer_;
     std::map<int, vector<int>> consumers_;
-    vector<std::pair<int, Ptr<Layer>>> attention_replacements_;
+    vector<std::pair<int, Ptr<OpData>>> attention_replacements_;
 };
 
 void Net::Impl::fuseAttention()

modules/dnn/src/graph_fusion_basic.cpp

@@ -30,7 +30,7 @@ struct ModelFusionBasic
     }
 
     template<typename _LayerType> _LayerType*
-    getLayer(std::vector<Ptr<Layer> >& newprog, int op_idx) const
+    getLayer(std::vector<Ptr<OpData> >& newprog, int op_idx) const
     {
         return op_idx >= 0 ? dynamic_cast<_LayerType*>(newprog.at(op_idx).get()) : 0;
     }
@@ -39,14 +39,14 @@ struct ModelFusionBasic
     {
         vector<Arg> removed_args;
         bool modified = false;
-        const std::vector<Ptr<Layer> >& prog = graph->prog();
+        const std::vector<Ptr<OpData> >& prog = graph->prog();
         size_t i, nargs = netimpl->args.size(), nops = prog.size();
         std::vector<int> producer_of(nargs, -1);
-        std::vector<Ptr<Layer> > newprog;
+        std::vector<Ptr<OpData> > newprog;
         std::vector<Arg> fused_inputs;
 
         for (i = 0; i < nops; i++) {
-            const Ptr<Layer>& layer = prog[i];
+            const Ptr<OpData>& layer = prog[i];
             Layer* layer_ptr = (Layer*)layer.get();
             int fused_layer_idx = -1;
             std::vector<Ptr<Graph> >* subgraphs = layer->subgraphs();
@@ -74,7 +74,7 @@ struct ModelFusionBasic
                     int conv_layer_idx = producer_of.at(bn_inp.idx);
                     Conv2Layer* conv = getLayer<Conv2Layer>(newprog, conv_layer_idx);
                     if (conv) {
-                        bool ok = conv->fuseBatchNorm(layer);
+                        bool ok = conv->fuseBatchNorm(layer.dynamicCast<Layer>());
                         if (ok) {
                             fused_layer_idx = conv_layer_idx;
                             removed_args.push_back(bn_inp);
@@ -122,7 +122,7 @@ struct ModelFusionBasic
                     int conv_layer_idx = producer_of.at(activ_inp.idx);
                     Conv2Layer* conv = getLayer<Conv2Layer>(newprog, conv_layer_idx);
                     if (conv) {
-                        bool ok = conv->fuseActivation(layer);
+                        bool ok = conv->fuseActivation(layer.dynamicCast<Layer>());
                         if (ok) {
                             fused_layer_idx = conv_layer_idx;
                             removed_args.push_back(activ_inp);
@@ -209,7 +209,7 @@ struct ModelFusionBasic
                                                         gnparams.type = "GroupNormalization";
                                                         gnparams.set("epsilon", instnorm->epsilon);
                                                         gnparams.set("num_groups", num_groups);
-                                                        Ptr<Layer> gnlayer = GroupNormLayer::create(gnparams);
+                                                        Ptr<OpData> gnlayer = GroupNormLayer::create(gnparams);
                                                         gnlayer->netimpl = netimpl;
                                                         gnlayer->inputs = {orig_inp, mul_scale_arg, add_bias_arg};
                                                         newprog[instnorm_idx] = gnlayer;
@@ -219,9 +219,9 @@ struct ModelFusionBasic
                                                     removed_args.push_back(reshape2_inp);
                                                     removed_args.push_back(reshape2_out);
                                                     removed_args.push_back(mul_out);
-                                                    newprog[reshape1_idx] = Ptr<Layer>();
-                                                    newprog[reshape2_idx] = Ptr<Layer>();
-                                                    newprog[mul_idx] = Ptr<Layer>();
+                                                    newprog[reshape1_idx] = Ptr<OpData>();
+                                                    newprog[reshape2_idx] = Ptr<OpData>();
+                                                    newprog[mul_idx] = Ptr<OpData>();
                                                     break;
                                                 }
                                             }
@@ -237,7 +237,7 @@ struct ModelFusionBasic
 
             if (fused_layer_idx >= 0) {
                 modified = true;
-                Layer* fused_layer = newprog[fused_layer_idx];
+                Layer* fused_layer = (Layer*)newprog[fused_layer_idx].get();
                 fused_layer->outputs = outputs;
                 for (Arg new_out: outputs)
                     producer_of[new_out.idx] = fused_layer_idx;
@@ -293,16 +293,16 @@ struct FuseBNPass
 
     void fuseGraph(Ptr<Graph>& graph)
     {
-        const std::vector<Ptr<Layer> >& prog = graph->prog();
+        const std::vector<Ptr<OpData> >& prog = graph->prog();
         size_t nops = prog.size(), nargs = netimpl->args.size();
-        std::vector<Ptr<Layer> > newprog;
+        std::vector<Ptr<OpData> > newprog;
         newprog.reserve(nops);
         std::vector<int> producer_of((int)nargs, -1);
         bool modified = false;
 
         for (size_t i = 0; i < nops; i++) {
-            const Ptr<Layer>& layer = prog[i];
-            Layer* layer_ptr = const_cast<Layer*>(layer.get());
+            const Ptr<OpData>& layer = prog[i];
+            Layer* layer_ptr = (Layer*)layer.get();
 
             std::vector<Ptr<Graph> >* subgraphs = layer->subgraphs();
             if (subgraphs)
@@ -324,7 +324,7 @@ struct FuseBNPass
                         usecounts[conv_inp0.idx] = 0;
                         if (bn_inp0.idx >= 0)
                             usecounts[bn_inp0.idx]++;
-                        newprog[bn_idx] = Ptr<Layer>();
+                        newprog[bn_idx] = Ptr<OpData>();
                         modified = true;
                     }
                 }

modules/dnn/src/graph_fusion_matmul_to_gemm.cpp

@@ -43,7 +43,7 @@ struct ModelFusionMatMulToGemm
 
     bool fuseGraph(Ptr<Graph>& graph)
     {
-        const vector<Ptr<Layer>>& prog = graph->prog();
+        const vector<Ptr<OpData>>& prog = graph->prog();
         size_t nops = prog.size();
         bool modified = false;
 
@@ -56,11 +56,11 @@ struct ModelFusionMatMulToGemm
             }
         }
 
-        vector<Ptr<Layer>> newprog = prog;
+        vector<Ptr<OpData>> newprog = prog;
         bool changed = false;
 
         for (size_t i = 0; i < nops; i++) {
-            const Ptr<Layer>& layer = newprog[i];
+            const Ptr<OpData>& layer = newprog[i];
             if (!layer) continue;
 
             MatMulLayer* mm = dynamic_cast<MatMulLayer*>(layer.get());
@@ -120,7 +120,7 @@ struct ModelFusionMatMulToGemm
             gp.blobs.push_back(B);
             if (have_bias) gp.blobs.push_back(layer->blobs[1]);
 
-            Ptr<Layer> gemm = LayerFactory::createLayerInstance("Gemm", gp);
+            Ptr<OpData> gemm = LayerFactory::createLayerInstance("Gemm", gp);
             if (!gemm) continue;
             gemm->inputs  = layer->inputs;
             gemm->outputs = layer->outputs;

modules/dnn/src/graph_fusion_qdq.cpp

@@ -32,7 +32,7 @@ struct ModelFusionQDQ
     }
 
     template<typename _LayerType> _LayerType*
-    getLayer(std::vector<Ptr<Layer> >& newprog, int op_idx) const
+    getLayer(std::vector<Ptr<OpData> >& newprog, int op_idx) const
     {
         return op_idx >= 0 ? dynamic_cast<_LayerType*>(newprog.at(op_idx).get()) : 0;
     }
@@ -48,10 +48,10 @@ struct ModelFusionQDQ
         return params;
     }
 
-    Ptr<Layer> createFusedLayer(const LayerParams& src) const
+    Ptr<OpData> createFusedLayer(const LayerParams& src) const
     {
         LayerParams params = src;
-        Ptr<Layer> layer = LayerFactory::createLayerInstance(params.type, params);
+        Ptr<OpData> layer = LayerFactory::createLayerInstance(params.type, params);
         if (!layer.empty())
             layer->netimpl = netimpl;
         return layer;
@@ -94,7 +94,7 @@ struct ModelFusionQDQ
                               size_t ninputs,
                               const std::vector<Arg>& inputs,
                               const std::vector<int>& producer_of,
-                              std::vector<Ptr<Layer> >& newprog,
+                              std::vector<Ptr<OpData> >& newprog,
                               Arg& q_data_in,
                               Arg& out_scale,
                               Arg& out_zp,
@@ -118,10 +118,10 @@ struct ModelFusionQDQ
     {
         vector<Arg> removed_args;
         bool modified = false;
-        const std::vector<Ptr<Layer> >& prog = graph->prog();
+        const std::vector<Ptr<OpData> >& prog = graph->prog();
         size_t i, nargs = netimpl->args.size(), nops = prog.size();
         std::vector<int> producer_of(nargs, -1);
-        std::vector<Ptr<Layer> > newprog;
+        std::vector<Ptr<OpData> > newprog;
         std::vector<Arg> fused_inputs;
         std::set<int> skip_indices;
         std::vector<Arg> override_outputs;
@@ -129,7 +129,7 @@ struct ModelFusionQDQ
 
         for (i = 0; i < nops; i++) {
             if (skip_indices.count((int)i)) continue;
-            const Ptr<Layer>& layer = prog[i];
+            const Ptr<OpData>& layer = prog[i];
             Layer* layer_ptr = (Layer*)layer.get();
             int fused_layer_idx = -1;
             std::vector<Ptr<Graph> >* subgraphs = layer->subgraphs();
@@ -219,7 +219,7 @@ struct ModelFusionQDQ
                                 removed_args.push_back(add_inp);
 
                             for (int dq_prog_idx : dq_prog_indices)
-                                newprog[dq_prog_idx] = Ptr<Layer>();
+                                newprog[dq_prog_idx] = Ptr<OpData>();
 
                             break;
                         }
@@ -289,7 +289,7 @@ struct ModelFusionQDQ
                             fused_inputs.assign(1, dq->inputs[0]);
                             removed_args.push_back(q_data_in);
                             removed_args.push_back(relu_in);
-                            newprog[dq_idx] = Ptr<Layer>();
+                            newprog[dq_idx] = Ptr<OpData>();
                             break;
                         }
                     }
@@ -363,7 +363,7 @@ struct ModelFusionQDQ
                             fused_inputs.assign(1, dq->inputs[0]);
                             removed_args.push_back(q_data_in);
                             removed_args.push_back(sig_in);
-                            newprog[dq_idx] = Ptr<Layer>();
+                            newprog[dq_idx] = Ptr<OpData>();
                             break;
                         }
                     }
@@ -450,7 +450,7 @@ struct ModelFusionQDQ
                                         }
                                         int prod_idx = producer_of.at(add2->inputs[k].idx);
                                         Layer* prod = prod_idx >= 0 && !newprog[prod_idx].empty()
-                                            ? newprog[prod_idx].get() : nullptr;
+                                            ? (Layer*)newprog[prod_idx].get() : nullptr;
                                         if (!prod || !getInt8OutputParams(prod, in_scales2[k], in_zps2[k]))
                                             elt_in_int82 = false;
                                     }
@@ -479,7 +479,7 @@ struct ModelFusionQDQ
                                         if (relu_out_uc <= 1) {
                                             fused_layer_idx = add_idx2;
                                             newprog[add_idx2] = eltInt8;
-                                            newprog[relu_layer_idx2] = Ptr<Layer>();
+                                            newprog[relu_layer_idx2] = Ptr<OpData>();
                                             removed_args.push_back(q_inp);   // relu_out
                                             removed_args.push_back(relu_in2); // add_out
                                             for (size_t dk = 0; dk < add2->inputs.size(); dk++) {
@@ -487,7 +487,7 @@ struct ModelFusionQDQ
                                             }
                                             for (int dq_prog_idx : dq_prog_indices2) {
                                                 if (dq_prog_idx >= 0)
-                                                    newprog[dq_prog_idx] = Ptr<Layer>();
+                                                    newprog[dq_prog_idx] = Ptr<OpData>();
                                             }
                                         } else {
                                             int new_idx = (int)newprog.size();
@@ -665,13 +665,13 @@ struct ModelFusionQDQ
                                     if (conv->inputs.size() == 3) {
                                         removed_args.push_back(conv->inputs[2]);
                                         if (dq_bias_idx >= 0)
-                                            newprog[dq_bias_idx] = Ptr<Layer>();
+                                            newprog[dq_bias_idx] = Ptr<OpData>();
                                     }
                                     if (usecounts.at(conv_x.idx) == 1) {
                                         removed_args.push_back(conv_x);
-                                        newprog[dq_x_idx] = Ptr<Layer>();
+                                        newprog[dq_x_idx] = Ptr<OpData>();
                                     }
-                                    newprog[dq_w_idx] = Ptr<Layer>();
+                                    newprog[dq_w_idx] = Ptr<OpData>();
                                     break;
                                 }
                             }
@@ -770,8 +770,8 @@ struct ModelFusionQDQ
                                     removed_args.push_back(q_data_in);
                                     removed_args.push_back(mm_x);
                                     removed_args.push_back(mm_w);
-                                    newprog[dq_x_idx] = Ptr<Layer>();
-                                    newprog[dq_w_idx] = Ptr<Layer>();
+                                    newprog[dq_x_idx] = Ptr<OpData>();
+                                    newprog[dq_w_idx] = Ptr<OpData>();
                                     break;
                                 }
                             }
@@ -896,9 +896,9 @@ struct ModelFusionQDQ
                                             removed_args.push_back(add_bias->inputs[mm_inp_k]); // matmul out
                                             removed_args.push_back(mm_x);
                                             removed_args.push_back(mm_w);
-                                            newprog[mm2_idx] = Ptr<Layer>();
-                                            newprog[dq_x_idx] = Ptr<Layer>();
-                                            newprog[dq_w_idx] = Ptr<Layer>();
+                                            newprog[mm2_idx] = Ptr<OpData>();
+                                            newprog[dq_x_idx] = Ptr<OpData>();
+                                            newprog[dq_w_idx] = Ptr<OpData>();
                                             break;
                                         }
                                     }
@@ -1026,8 +1026,8 @@ struct ModelFusionQDQ
                                             removed_args.push_back(q_data_in);
                                             removed_args.push_back(gemm_a);
                                             removed_args.push_back(gemm_b);
-                                            newprog[dq_a_idx] = Ptr<Layer>();
-                                            newprog[dq_b_idx] = Ptr<Layer>();
+                                            newprog[dq_a_idx] = Ptr<OpData>();
+                                            newprog[dq_b_idx] = Ptr<OpData>();
                                             break;
                                         }
                                     }
@@ -1089,7 +1089,7 @@ struct ModelFusionQDQ
                                 fused_inputs.assign(1, dq->inputs[0]);
                                 removed_args.push_back(q_data_in);
                                 removed_args.push_back(pool_in);
-                                newprog[dq_idx] = Ptr<Layer>();
+                                newprog[dq_idx] = Ptr<OpData>();
                                 break;
                             }
                         }
@@ -1107,7 +1107,7 @@ struct ModelFusionQDQ
                         Arg ql_zp = inputs[2];
                         int shuffle_idx = producer_of.at(ql_data.idx);
                         Layer* shuffle_layer = (shuffle_idx >= 0 && !newprog[shuffle_idx].empty())
-                            ? newprog[shuffle_idx].get() : nullptr;
+                            ? (Layer*)newprog[shuffle_idx].get() : nullptr;
 
                         bool is_shuffle = shuffle_layer && shuffle_layer->isDataShuffling();
 
@@ -1142,7 +1142,7 @@ struct ModelFusionQDQ
                                         fused_inputs.push_back(shuffle_layer->inputs[si]);
                                     removed_args.push_back(ql_data);
                                     removed_args.push_back(shuffle_inp);
-                                    newprog[dq_idx] = Ptr<Layer>();
+                                    newprog[dq_idx] = Ptr<OpData>();
                                     break;
                                 }
                             }
@@ -1156,14 +1156,15 @@ struct ModelFusionQDQ
                     Arg activ_inp = inputs[0];
                     int producer_idx = producer_of.at(activ_inp.idx);
                     if (producer_idx >= 0 && !newprog[producer_idx].empty()) {
-                        Layer* producer_layer = newprog[producer_idx].get();
+                        Layer* producer_layer = dynamic_cast<Layer*>(newprog[producer_idx].get());
                         float prod_sc = 0.f;
                         int prod_zp = 0;
-                        if (getInt8OutputParams(producer_layer, prod_sc, prod_zp) &&
+                        if (producer_layer &&
+                            getInt8OutputParams(producer_layer, prod_sc, prod_zp) &&
                             prod_sc == activ_int8->input_sc &&
                             prod_zp == activ_int8->input_zp) {
                             Ptr<ActivationLayer> activ_layer = layer.dynamicCast<ActivationLayer>();
-                            if (newprog[producer_idx]->setActivation(activ_layer)) {
+                            if (producer_layer->setActivation(activ_layer)) {
                                 setInt8OutputParams(producer_layer,
                                                     activ_int8->output_sc,
                                                     activ_int8->output_zp);
@@ -1179,7 +1180,7 @@ struct ModelFusionQDQ
 
             if (fused_layer_idx >= 0) {
                 modified = true;
-                Layer* fused_layer = newprog[fused_layer_idx];
+                Layer* fused_layer = (Layer*)newprog[fused_layer_idx].get();
                 const std::vector<Arg>& final_outputs = override_outputs.empty() ? outputs : override_outputs;
                 fused_layer->outputs = final_outputs;
                 if (!fused_inputs.empty())
@@ -1224,7 +1225,7 @@ struct ModelFusionQDQ
             while (cur >= 0 && !newprog[cur].empty()) {
                 ql = dynamic_cast<QuantizeLinearLayer*>(newprog[cur].get());
                 if (ql) break;
-                Layer* l = newprog[cur].get();
+                Layer* l = (Layer*)newprog[cur].get();
                 if (l->inputs.empty()) { ql = nullptr; break; }
                 cur = producer_of.at(l->inputs[0].idx);
             }
@@ -1303,7 +1304,7 @@ struct ModelFusionQDQ
 
             conv->float_input = true;
             conv->inputs[0] = ql_data;
-            newprog[ql_idx] = Ptr<Layer>();
+            newprog[ql_idx] = Ptr<OpData>();
             modified = true;
         }
 
@@ -1338,7 +1339,7 @@ struct ModelFusionQDQ
                             if (inp.idx >= 0 && inp.idx < (int)nargs)
                                 uc[inp.idx]--;
                         }
-                        layer = Ptr<Layer>();
+                        layer = Ptr<OpData>();
                         changed = true;
                     }
                 }

modules/dnn/src/graph_fusion_reshape_transpose.cpp

@@ -46,7 +46,7 @@ struct ModelFusionReshapeTranspose
 
     bool fuseGraph(Ptr<Graph>& graph)
     {
-        const vector<Ptr<Layer>>& prog = graph->prog();
+        const vector<Ptr<OpData>>& prog = graph->prog();
         size_t nops = prog.size();
         bool modified = false;
 
@@ -72,7 +72,7 @@ struct ModelFusionReshapeTranspose
         vector<bool> dropped(nops, false);
 
         for (size_t i = 0; i < nops; i++) {
-            const Ptr<Layer>& layer = prog[i];
+            const Ptr<OpData>& layer = prog[i];
             if (!layer || dropped[i]) continue;
             if (layer->inputs.empty() || layer->outputs.empty()) continue;
 
@@ -95,7 +95,7 @@ struct ModelFusionReshapeTranspose
                 if (it != producer.end()) {
                     int prod_idx = it->second;
                     if (prod_idx >= 0 && !dropped[prod_idx]) {
-                        const Ptr<Layer>& pl = prog[prod_idx];
+                        const Ptr<OpData>& pl = prog[prod_idx];
                         TransposeLayer* prevTr = dynamic_cast<TransposeLayer*>(pl.get());
                         Arg prevOut = layer->inputs[0];
                         bool single_consumer = usecounts[prevOut.idx] == 1
@@ -135,7 +135,7 @@ struct ModelFusionReshapeTranspose
                 if (it != producer.end()) {
                     int prod_idx = it->second;
                     if (prod_idx >= 0 && !dropped[prod_idx]) {
-                        const Ptr<Layer>& pl = prog[prod_idx];
+                        const Ptr<OpData>& pl = prog[prod_idx];
                         Reshape2Layer* prevRs = dynamic_cast<Reshape2Layer*>(pl.get());
                         Arg prevOut = layer->inputs[0];
                         bool single_consumer = usecounts[prevOut.idx] == 1
@@ -154,7 +154,7 @@ struct ModelFusionReshapeTranspose
         }
 
         if (modified) {
-            vector<Ptr<Layer>> newprog;
+            vector<Ptr<OpData>> newprog;
             newprog.reserve(nops);
             for (size_t i = 0; i < nops; i++) {
                 if (!dropped[i] && prog[i])
@@ -166,7 +166,7 @@ struct ModelFusionReshapeTranspose
         return modified;
     }
 
-    void redirectConsumers(const vector<Ptr<Layer>>& prog,
+    void redirectConsumers(const vector<Ptr<OpData>>& prog,
                            const vector<bool>& dropped,
                            size_t start_idx, Arg from, Arg to)
     {

modules/dnn/src/graph_fusion_scale_softmax.cpp

@@ -40,7 +40,7 @@ struct ModelFusionScaleSoftmax
 
     bool fuseGraph(Ptr<Graph>& graph)
     {
-        const vector<Ptr<Layer>>& prog = graph->prog();
+        const vector<Ptr<OpData>>& prog = graph->prog();
         size_t nops = prog.size();
         bool modified = false;
 
@@ -70,7 +70,7 @@ struct ModelFusionScaleSoftmax
         vector<bool> dropped(nops, false);
 
         for (size_t i = 0; i < nops; i++) {
-            const Ptr<Layer>& layer = prog[i];
+            const Ptr<OpData>& layer = prog[i];
             if (!layer || dropped[i]) continue;
 
             SoftmaxLayer* sm = dynamic_cast<SoftmaxLayer*>(layer.get());
@@ -83,7 +83,7 @@ struct ModelFusionScaleSoftmax
             int prod_idx = it->second;
             if (prod_idx < 0 || dropped[prod_idx]) continue;
 
-            const Ptr<Layer>& pl = prog[prod_idx];
+            const Ptr<OpData>& pl = prog[prod_idx];
             NaryEltwiseLayer* elt = dynamic_cast<NaryEltwiseLayer*>(pl.get());
             if (!elt) continue;
             const auto op = elt->op;
@@ -123,7 +123,7 @@ struct ModelFusionScaleSoftmax
         }
 
         if (modified) {
-            vector<Ptr<Layer>> newprog;
+            vector<Ptr<OpData>> newprog;
             newprog.reserve(nops);
             for (size_t i = 0; i < nops; i++) {
                 if (!dropped[i] && prog[i])

modules/dnn/src/graph_fusion_shared_gemm.cpp

@@ -35,7 +35,7 @@ using std::string;
 
 namespace {
 
-static bool readGemmWeight(const Ptr<Layer>& l, bool trans_b, Mat& W_out)
+static bool readGemmWeight(const Ptr<OpData>& l, bool trans_b, Mat& W_out)
 {
     if (l->blobs.empty()) return false;
     const Mat& W = l->blobs[0];
@@ -48,7 +48,7 @@ static bool readGemmWeight(const Ptr<Layer>& l, bool trans_b, Mat& W_out)
     return true;
 }
 
-static bool readGemmBias(const Ptr<Layer>& l, Mat& b_out)
+static bool readGemmBias(const Ptr<OpData>& l, Mat& b_out)
 {
     if (l->blobs.size() < 2) { b_out.release(); return true; }
     const Mat& b = l->blobs[1];
@@ -76,10 +76,10 @@ struct ModelFusionSharedGemm
         bool   flatten_a = true;
     };
 
-    bool inspectGemm(const vector<Ptr<Layer>>& prog, int idx, GemmInfo& info) const
+    bool inspectGemm(const vector<Ptr<OpData>>& prog, int idx, GemmInfo& info) const
     {
         if (idx < 0 || idx >= (int)prog.size() || !prog[idx]) return false;
-        const Ptr<Layer>& l = prog[idx];
+        const Ptr<OpData>& l = prog[idx];
         GemmLayer* g = dynamic_cast<GemmLayer*>(l.get());
         if (!g) return false;
 
@@ -125,7 +125,7 @@ struct ModelFusionSharedGemm
 
     bool fuseGraph(Ptr<Graph>& graph)
     {
-        const vector<Ptr<Layer>>& prog = graph->prog();
+        const vector<Ptr<OpData>>& prog = graph->prog();
         size_t nops = prog.size();
 
         struct Key { int input_idx; int trans_b; int K; };
@@ -145,7 +145,7 @@ struct ModelFusionSharedGemm
 
         bool modified = false;
         std::set<int> removed_ops;
-        vector<std::pair<int, vector<Ptr<Layer>>>> insertions;  // (insert_pos, fused-and-slice layers)
+        vector<std::pair<int, vector<Ptr<OpData>>>> insertions;  // (insert_pos, fused-and-slice layers)
 
         for (auto& group : groups) {
             auto infos = group.second;
@@ -226,7 +226,7 @@ struct ModelFusionSharedGemm
             fp.blobs.push_back(W_concat);
             if (all_have_bias) fp.blobs.push_back(b_concat);
 
-            Ptr<Layer> fused = LayerFactory::createLayerInstance("Gemm", fp);
+            Ptr<OpData> fused = LayerFactory::createLayerInstance("Gemm", fp);
             if (!fused) continue;
 
             string fused_out_name = fp.name + "_out";
@@ -235,7 +235,7 @@ struct ModelFusionSharedGemm
             fused->outputs = { fused_out_arg };
             fused->netimpl = netimpl;
 
-            vector<Ptr<Layer>> slices;
+            vector<Ptr<OpData>> slices;
             int col_cursor = 0;
             for (size_t s = 0; s < infos.size(); s++) {
                 int N_s = infos[s].N;
@@ -257,7 +257,7 @@ struct ModelFusionSharedGemm
                 Arg axes_arg   = netimpl->newConstArg(sp.name + "_axes",   axes);
                 Arg steps_arg  = netimpl->newConstArg(sp.name + "_steps",  steps);
 
-                Ptr<Layer> slice = LayerFactory::createLayerInstance("Slice2", sp);
+                Ptr<OpData> slice = LayerFactory::createLayerInstance("Slice2", sp);
                 if (!slice) { uniform = false; break; }
                 slice->inputs  = { fused_out_arg, starts_arg, ends_arg, axes_arg, steps_arg };
                 slice->outputs = prog[infos[s].layer_idx]->outputs;
@@ -267,7 +267,7 @@ struct ModelFusionSharedGemm
             if (!uniform) continue;
 
             for (auto& info : infos) removed_ops.insert(info.layer_idx);
-            vector<Ptr<Layer>> bundle;
+            vector<Ptr<OpData>> bundle;
             bundle.push_back(fused);
             for (auto& s : slices) bundle.push_back(s);
             insertions.emplace_back(insert_pos, std::move(bundle));
@@ -279,7 +279,7 @@ struct ModelFusionSharedGemm
         std::sort(insertions.begin(), insertions.end(),
                   [](auto& a, auto& b) { return a.first < b.first; });
 
-        vector<Ptr<Layer>> newprog;
+        vector<Ptr<OpData>> newprog;
         size_t ins_idx = 0;
         for (size_t i = 0; i < nops; i++) {
             while (ins_idx < insertions.size() &&

modules/dnn/src/graph_fusion_transpose_matmul.cpp

@@ -38,7 +38,7 @@ struct ModelFusionTransposeMatMul
 
     bool fuseGraph(Ptr<Graph>& graph)
     {
-        const vector<Ptr<Layer>>& prog = graph->prog();
+        const vector<Ptr<OpData>>& prog = graph->prog();
         size_t nops = prog.size();
         bool modified = false;
 
@@ -68,7 +68,7 @@ struct ModelFusionTransposeMatMul
         vector<bool> dropped(nops, false);
 
         for (size_t i = 0; i < nops; i++) {
-            const Ptr<Layer>& layer = prog[i];
+            const Ptr<OpData>& layer = prog[i];
             if (!layer || dropped[i]) continue;
 
             MatMulLayer* mm = dynamic_cast<MatMulLayer*>(layer.get());
@@ -83,7 +83,7 @@ struct ModelFusionTransposeMatMul
                 int prod_idx = it->second;
                 if (prod_idx < 0 || dropped[prod_idx]) continue;
 
-                const Ptr<Layer>& pl = prog[prod_idx];
+                const Ptr<OpData>& pl = prog[prod_idx];
                 TransposeLayer* tr = dynamic_cast<TransposeLayer*>(pl.get());
                 if (!tr || pl->outputs.size() != 1) continue;
                 if (!isLastTwoSwap(tr->perm)) continue;
@@ -103,7 +103,7 @@ struct ModelFusionTransposeMatMul
         }
 
         if (modified) {
-            vector<Ptr<Layer>> newprog;
+            vector<Ptr<OpData>> newprog;
             newprog.reserve(nops);
             for (size_t i = 0; i < nops; i++) {
                 if (!dropped[i] && prog[i])

modules/dnn/src/init.cpp

@@ -48,6 +48,11 @@
 
 namespace cv {
 namespace dnn {
+
+#ifdef HAVE_CUDA
+void registerConv2CudaBackend();  // defined in layers/conv2_layer.cpp (plain cv::dnn namespace)
+#endif
+
 CV__DNN_INLINE_NS_BEGIN
 
 static Mutex* __initialization_mutex = NULL;
@@ -76,6 +81,10 @@ public:
 } // namespace
 #endif
 
+#ifdef HAVE_CUDA
+void registerCudaCommonExecs();  // op_cuda.cpp (inline namespace)
+#endif
+
 void initializeLayerFactory()
 {
     CV_TRACE_FUNCTION();
@@ -84,6 +93,12 @@ void initializeLayerFactory()
     static ProtobufShutdown protobufShutdown; CV_UNUSED(protobufShutdown);
 #endif
 
+#ifdef HAVE_CUDA
+    // New graph engine: per-op CUDA executors.
+    registerConv2CudaBackend();
+    registerCudaCommonExecs();
+#endif
+
     CV_DNN_REGISTER_LAYER_CLASS(If,             IfLayer);
     CV_DNN_REGISTER_LAYER_CLASS(Loop,           LoopLayer);
     CV_DNN_REGISTER_LAYER_CLASS(Scan,           ScanLayer);

modules/dnn/src/layer.cpp

@@ -10,37 +10,48 @@ namespace dnn {
 CV__DNN_INLINE_NS_BEGIN
 
 
-Layer::Layer() {
+OpData::OpData() {
     netimpl = nullptr;
-    preferableTarget = DNN_TARGET_CPU;
 }
 
-Layer::Layer(const LayerParams& params)
+OpData::OpData(const LayerParams& params)
     : blobs(params.blobs)
     , name(params.name)
     , type(params.type)
 {
     netimpl = nullptr;
-    preferableTarget = DNN_TARGET_CPU;
 }
 
-void Layer::setParamsFrom(const LayerParams& params)
+OpData::~OpData() {}
+
+void OpData::setParamsFrom(const LayerParams& params)
 {
     blobs = params.blobs;
     name = params.name;
     type = params.type;
 }
 
-int Layer::inputNameToIndex(String)
+int OpData::inputNameToIndex(String)
 {
     return -1;
 }
 
-int Layer::outputNameToIndex(const String&)
+int OpData::outputNameToIndex(const String&)
 {
     return 0;
 }
 
+Layer::Layer() {
+    netimpl = nullptr;
+    preferableTarget = DNN_TARGET_CPU;
+}
+
+Layer::Layer(const LayerParams& params)
+    : OpData(params)
+{
+    preferableTarget = DNN_TARGET_CPU;
+}
+
 bool Layer::supportBackend(int backendId)
 {
     return backendId == DNN_BACKEND_OPENCV;
@@ -94,13 +105,21 @@ Ptr<BackendNode> Layer::initCann(const std::vector<Ptr<BackendWrapper> > &inputs
 
 bool Layer::setActivation(const Ptr<ActivationLayer>&) { return false; }
 bool Layer::tryFuse(Ptr<Layer>&) { return false; }
-void Layer::getScaleShift(Mat& scale, Mat& shift) const
+
+void Layer::forwardCUDA(const std::vector<Ptr<BackendWrapper> >&,
+                        const std::vector<Ptr<BackendWrapper> >&,
+                        void*)
+{
+    CV_Error(Error::StsNotImplemented, "CUDA forward of " + type + " layers is not defined.");
+}
+
+void OpData::getScaleShift(Mat& scale, Mat& shift) const
 {
     scale = Mat();
     shift = Mat();
 }
 
-void Layer::getScaleZeropoint(float& scale, int& zeropoint) const
+void OpData::getScaleZeropoint(float& scale, int& zeropoint) const
 {
     scale = 1.f;
     zeropoint = 0;
@@ -247,7 +266,7 @@ void Layer::run(const std::vector<Mat>& inputs, std::vector<Mat>& outputs, std::
 
 Layer::~Layer() {}
 
-bool Layer::getMemoryShapes(const std::vector<MatShape>& inputs,
+bool OpData::getMemoryShapes(const std::vector<MatShape>& inputs,
         const int requiredOutputs,
         std::vector<MatShape>& outputs,
         std::vector<MatShape>& internals) const
@@ -257,7 +276,7 @@ bool Layer::getMemoryShapes(const std::vector<MatShape>& inputs,
     return false;
 }
 
-void Layer::getTypes(const std::vector<MatType>&inputs,
+void OpData::getTypes(const std::vector<MatType>&inputs,
                      const int requiredOutputs,
                      const int requiredInternals,
                      std::vector<MatType>&outputs,
@@ -265,20 +284,13 @@ void Layer::getTypes(const std::vector<MatType>&inputs,
 {
     CV_Assert(inputs.size());
     for (auto input : inputs)
-    {
-        if (preferableTarget == DNN_TARGET_CUDA_FP16 || preferableTarget == DNN_TARGET_CUDA)
-            CV_CheckTypeEQ(input, CV_32F, "");
-        else if (preferableTarget == DNN_TARGET_OPENCL_FP16)
-            CV_CheckType(input, input == CV_16F || input == CV_8S || input == CV_8U || input == CV_64F || input == CV_64S, "");
-        else
-            CV_CheckType(input, input == CV_32F || input == CV_64F || input == CV_8S || input == CV_8U || input == CV_64S, "");
-    }
+        CV_CheckType(input, input == CV_32F || input == CV_64F || input == CV_8S || input == CV_8U || input == CV_64S, "");
 
     outputs.assign(requiredOutputs, inputs[0]);
     internals.assign(requiredInternals, inputs[0]);
 }
 
-int Layer::getLayouts(const std::vector<DataLayout>& actualInputs,
+int OpData::getLayouts(const std::vector<DataLayout>& actualInputs,
                        std::vector<DataLayout>& desiredInputs,
                        const int requiredOutputs,
                        std::vector<DataLayout>& outputs) const
@@ -288,43 +300,43 @@ int Layer::getLayouts(const std::vector<DataLayout>& actualInputs,
     return 0;
 }
 
-int64 Layer::getFLOPS(const std::vector<MatShape>&,
+int64 OpData::getFLOPS(const std::vector<MatShape>&,
                       const std::vector<MatShape>&) const
 {
     return 0;
 }
 
-bool Layer::updateMemoryShapes(const std::vector<MatShape>& inputs)
+bool OpData::updateMemoryShapes(const std::vector<MatShape>& inputs)
 {
     return true;
 }
 
-std::vector<Ptr<Graph> >* Layer::subgraphs() const
+std::vector<Ptr<Graph> >* OpData::subgraphs() const
 {
     return nullptr;
 }
 
-bool Layer::alwaysSupportInplace() const
+bool OpData::alwaysSupportInplace() const
 {
     return false;
 }
 
-bool Layer::dynamicOutputShapes() const
+bool OpData::dynamicOutputShapes() const
 {
     return false;
 }
 
-bool Layer::isDataShuffling() const
+bool OpData::isDataShuffling() const
 {
     return false;
 }
 
-std::ostream& Layer::dumpAttrs(std::ostream& strm, int) const
+std::ostream& OpData::dumpAttrs(std::ostream& strm, int) const
 {
     return strm;
 }
 
-std::ostream& Layer::dump(std::ostream& strm, int indent, bool comma) const
+std::ostream& OpData::dump(std::ostream& strm, int indent, bool comma) const
 {
     CV_Assert(netimpl);
     size_t ninputs = inputs.size();

modules/dnn/src/layer_factory.cpp

@@ -104,6 +104,72 @@ Ptr<Layer> LayerFactory::createLayerInstance(const String& type, LayerParams& pa
     }
 }
 
+typedef std::map<std::string, LayerFactory::OpConstructor> OpFactory_Impl;
+typedef std::map<std::string, std::map<int, LayerFactory::ExecConstructor> > ExecFactory_Impl;
+
+static OpFactory_Impl& getOpFactoryImpl()
+{
+    static OpFactory_Impl impl;
+    return impl;
+}
+
+static ExecFactory_Impl& getExecFactoryImpl()
+{
+    static ExecFactory_Impl impl;
+    return impl;
+}
+
+void LayerFactory::registerOp(const String& type, OpConstructor constructor)
+{
+    CV_TRACE_FUNCTION();
+    CV_TRACE_ARG_VALUE(type, "type", type.c_str());
+    CV_Assert(constructor);
+    cv::AutoLock lock(getLayerFactoryMutex());
+    getOpFactoryImpl()[type] = constructor;  // last registration wins
+}
+
+Ptr<OpData> LayerFactory::createOp(const String& type, const LayerParams& params)
+{
+    CV_TRACE_FUNCTION();
+    CV_TRACE_ARG_VALUE(type, "type", type.c_str());
+    cv::AutoLock lock(getLayerFactoryMutex());
+    OpFactory_Impl& impl = getOpFactoryImpl();
+    OpFactory_Impl::const_iterator it = impl.find(type);
+    if (it != impl.end())
+        return it->second(params);
+    return Ptr<OpData>();  // NULL: no OpData constructor for this type yet
+}
+
+void LayerFactory::registerExec(const String& type, int backendId, ExecConstructor constructor)
+{
+    CV_TRACE_FUNCTION();
+    CV_TRACE_ARG_VALUE(type, "type", type.c_str());
+    CV_Assert(constructor);
+    cv::AutoLock lock(getLayerFactoryMutex());
+    getExecFactoryImpl()[type][backendId] = constructor;
+}
+
+Ptr<Layer> LayerFactory::createExec(const String& type, int backendId,
+                                    const Ptr<OpData>& data, void* backendCtx)
+{
+    CV_TRACE_FUNCTION();
+    CV_TRACE_ARG_VALUE(type, "type", type.c_str());
+    ExecConstructor ctor = nullptr;
+    {
+        cv::AutoLock lock(getLayerFactoryMutex());
+        ExecFactory_Impl& impl = getExecFactoryImpl();
+        ExecFactory_Impl::const_iterator it = impl.find(type);
+        if (it != impl.end()) {
+            auto bit = it->second.find(backendId);
+            if (bit != it->second.end())
+                ctor = bit->second;
+        }
+    }
+    if (ctor)
+        return ctor(data, backendCtx);
+    return Ptr<Layer>();
+}
+
 
 CV__DNN_INLINE_NS_END
 }}  // namespace cv::dnn

modules/dnn/src/layers/batch_norm2_layer.cpp

@@ -7,6 +7,10 @@
 #include "../precomp.hpp"
 #include "layers_common.hpp"
 #include "../net_impl.hpp"
+#include "../op_cuda.hpp"
+#ifdef HAVE_CUDA
+#include "../cuda4dnn/primitives/batch_norm.hpp"
+#endif
 
 namespace cv {
 namespace dnn {
@@ -355,9 +359,27 @@ public:
 
     bool supportBackend(int backendId) CV_OVERRIDE
     {
-        return backendId == DNN_BACKEND_OPENCV;
+        return backendId == DNN_BACKEND_OPENCV
+#ifdef HAVE_CUDA
+               || backendId == DNN_BACKEND_CUDA
+#endif
+               ;
     }
 
+#ifdef HAVE_CUDA
+    // Channel-wise scale+shift on CUDA; reused by the new graph engine via CUDALegacyExec.
+    Ptr<BackendNode> initCUDA(void* context_,
+                              const std::vector<Ptr<BackendWrapper> >&,
+                              const std::vector<Ptr<BackendWrapper> >&) CV_OVERRIDE
+    {
+        auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
+        Mat scale, bias;
+        getScaleBias(scale, bias);  // per-channel FP32 scale and shift
+        return make_cuda_node<cuda4dnn::BatchNormOp>(
+            preferableTarget, std::move(context->stream), scale, bias);
+    }
+#endif
+
     MatShape getOutShape(const MatShape& inpShape) const
     {
         return inpShape;

modules/dnn/src/layers/conv2_layer.cpp

@@ -11,6 +11,13 @@
 #include <algorithm>
 #include <cstring>
 
+#include "../op_cuda.hpp"
+#include <opencv2/dnn/layer.details.hpp>  // CV_DNN_REGISTER_EXEC_CLASS
+#ifdef HAVE_CUDA
+#include "../cuda4dnn/primitives/convolution.hpp"
+using namespace cv::dnn::cuda4dnn;
+#endif
+
 namespace cv
 {
 namespace dnn
@@ -115,6 +122,10 @@ public:
         int wtype = accuracy < 0 ? CV_32F : accuracy;
 
         wshape0 = weights_.shape();
+#ifdef HAVE_CUDA
+        // Retain the original NCHW filter for the CUDA (cuDNN) path.
+        weights_.convertTo(origWeights, CV_32F);
+#endif
         bool depthwise = ngroups == wshape0[0] && wshape0[1] == 1;
 
         if (depthwise) {
@@ -660,9 +671,96 @@ public:
         }
     }
 
+#ifdef HAVE_CUDA
+    bool cudaSupported() const
+    {
+        if (origWeights.empty() || wshape0.dims != 4)  // [Cout, Cin/group, kh, kw] (2D conv)
+            return false;
+        if (auto_pad != AUTO_PAD_NONE && auto_pad != AUTO_PAD_VALID)
+            return false;
+        if (activationFunc != nullptr || !activ.empty())
+            return false;
+        if (fastActivation != FAST_ACTIV_NONE && fastActivation != FAST_ACTIV_RELU &&
+            fastActivation != FAST_ACTIV_LEAKY_RELU && fastActivation != FAST_ACTIV_CLIP)
+            return false;
+        return true;
+    }
+
+    Ptr<BackendNode> initCudaConvNode(void* context_, const MatShape& inpShape,
+                                      const MatShape& outShape, int targetId)
+    {
+        csl::CSLContext context = *reinterpret_cast<csl::CSLContext*>(context_);
+        const int nspatial = wshape0.dims - 2;
+
+        ConvolutionConfiguration config;
+        for (int i = 0; i < nspatial; i++) {
+            config.kernel_size.push_back((size_t)wshape0[2 + i]);
+            config.strides.push_back(strides.empty() ? 1 : (size_t)strides[i]);
+            config.dilations.push_back(dilations.empty() ? 1 : (size_t)dilations[i]);
+        }
+        if (auto_pad == AUTO_PAD_VALID) {
+            config.padMode = ConvolutionConfiguration::PaddingMode::VALID;
+        } else {
+            config.padMode = ConvolutionConfiguration::PaddingMode::MANUAL;
+            for (int i = 0; i < nspatial; i++) {
+                config.pads_begin.push_back(pads.empty() ? 0 : (size_t)pads[i]);
+                config.pads_end.push_back(pads.empty() ? 0 : (size_t)pads[i + nspatial]);
+            }
+        }
+        config.input_shape.assign(inpShape.begin(), inpShape.end());
+        config.output_shape.assign(outShape.begin(), outShape.end());
+        config.groups = (size_t)ngroups;
+
+        Mat filters = origWeights, biasMat = bias;
+        if (fusedBatchNorm) {
+            filters = origWeights.clone();
+            const int Cout = wshape0[0];
+            const size_t inner = filters.total() / (size_t)Cout;
+            const float* sc = fusedScale.ptr<float>();
+            float* wp = filters.ptr<float>();
+            for (int co = 0; co < Cout; co++) {
+                float s = sc[co];
+                for (size_t k = 0; k < inner; k++)
+                    wp[co * inner + k] *= s;
+            }
+            biasMat = fusedBias;  // already b*scale + bn_bias
+        }
+
+        config.activation_type = ConvolutionConfiguration::ActivationType::IDENTITY;
+        config.relu_negative_slope = 0.f;
+        config.crelu_floor = 0.f; config.crelu_ceil = 0.f;
+        config.power_exp = 1.f; config.power_scale = 1.f; config.power_shift = 0.f;
+        bool hasAct = fastActivation != FAST_ACTIV_NONE;
+        if (fastActivation == FAST_ACTIV_RELU) {
+            config.activation_type = ConvolutionConfiguration::ActivationType::RELU;
+        } else if (fastActivation == FAST_ACTIV_LEAKY_RELU) {
+            config.activation_type = ConvolutionConfiguration::ActivationType::RELU;
+            config.relu_negative_slope = activParams.empty() ? 0.f : activParams[0];
+        } else if (fastActivation == FAST_ACTIV_CLIP) {
+            config.activation_type = ConvolutionConfiguration::ActivationType::CLIPPED_RELU;
+            config.crelu_floor = activParams.size() > 0 ? activParams[0] : 0.f;
+            config.crelu_ceil  = activParams.size() > 1 ? activParams[1] : 6.f;
+        }
+
+        if (addResidual && hasAct)
+            config.fusion_mode = ConvolutionConfiguration::FusionMode::ELTWISE_SUM_THEN_ACTIVATION;
+        else if (addResidual)
+            config.fusion_mode = ConvolutionConfiguration::FusionMode::ELTWISE_SUM;
+        else if (hasAct)
+            config.fusion_mode = ConvolutionConfiguration::FusionMode::ACTIVATION;
+        else
+            config.fusion_mode = ConvolutionConfiguration::FusionMode::NONE;
+
+        return make_cuda_node<cuda4dnn::ConvolutionOp>(
+            targetId, std::move(context.stream), std::move(context.cudnn_handle),
+            config, filters, biasMat);
+    }
+#endif
+
     std::vector<int> emptyKernelShape;
     Ptr<Layer> activ, batchNorm;
     Mat weights, bias, fusedScale, fusedBias;
+    Mat origWeights;  // original NCHW filter (FP32), kept for the CUDA path
     MatShape wshape0, prevInpshape;
     ConvState cs;
     bool fusedBatchNorm;
@@ -684,4 +782,53 @@ Ptr<Conv2Layer> Conv2Layer::create(const LayerParams& params)
     return Ptr<Conv2Layer>(new Conv2LayerImpl(params));
 }
 
+#ifdef HAVE_CUDA
+class CUDAConv2Layer : public Layer
+{
+public:
+    CUDAConv2Layer(const Ptr<Conv2LayerImpl>& conv_, void* ctx_) : conv(conv_), ctx(ctx_) {}
+
+    static Ptr<Layer> create(const Ptr<OpData>& data, void* backendCtx)
+    {
+        Ptr<Conv2LayerImpl> conv = data.dynamicCast<Conv2LayerImpl>();
+        if (!conv || !backendCtx || !conv->cudaSupported())
+            return Ptr<Layer>();
+        Ptr<CUDAConv2Layer> layer(new CUDAConv2Layer(conv, backendCtx));
+        layer->data = data;
+        layer->name = conv->name;
+        layer->type = conv->type;
+        layer->inputs = conv->inputs;
+        layer->outputs = conv->outputs;
+        return layer;
+    }
+
+    void forwardCUDA(const std::vector<Ptr<BackendWrapper> >& inputs,
+                     const std::vector<Ptr<BackendWrapper> >& outputs,
+                     void* workspace) CV_OVERRIDE
+    {
+        CV_Assert(!inputs.empty() && !outputs.empty());
+        auto& ws = *reinterpret_cast<cuda4dnn::csl::Workspace*>(workspace);
+        if (!node) {
+            auto inW = inputs[0].dynamicCast<CUDABackendWrapper>();
+            auto outW = outputs[0].dynamicCast<CUDABackendWrapper>();
+            node = conv->initCudaConvNode(ctx, inW->getShape(), outW->getShape(), preferableTarget);
+            cudaNode = node.dynamicCast<CUDABackendNode>();
+            CV_Assert(cudaNode);
+            ws.require(cudaNode->get_workspace_memory_in_bytes());
+        }
+        cudaNode->forward(inputs, outputs, ws);
+    }
+
+    Ptr<Conv2LayerImpl> conv;
+    void* ctx;
+    Ptr<BackendNode> node;
+    Ptr<CUDABackendNode> cudaNode;
+};
+
+void registerConv2CudaBackend()
+{
+    CV_DNN_REGISTER_EXEC_CLASS(Conv2, DNN_BACKEND_CUDA, CUDAConv2Layer);
+}
+#endif
+
 }}

modules/dnn/src/layers/maxpool_layer.cpp

@@ -7,6 +7,10 @@
 #include "../net_impl.hpp"
 #include "conv2_common.hpp"
 #include "opencv2/core/hal/intrin.hpp"
+#include "../op_cuda.hpp"
+#ifdef HAVE_CUDA
+#include "../cuda4dnn/primitives/pooling.hpp"
+#endif
 
 namespace cv
 {
@@ -476,9 +480,48 @@ public:
 
     virtual bool supportBackend(int backendId) CV_OVERRIDE
     {
+#ifdef HAVE_CUDA
+        if (backendId == DNN_BACKEND_CUDA) {
+            if (kernel_shape.size() != 2 || outputs.size() != 1)
+                return false;
+            for (int d : dilations) if (d != 1) return false;
+            return auto_pad == AUTO_PAD_NONE || auto_pad == AUTO_PAD_VALID;
+        }
+#endif
         return backendId == DNN_BACKEND_OPENCV;
     }
 
+#ifdef HAVE_CUDA
+    Ptr<BackendNode> initCUDA(void* context_,
+                              const std::vector<Ptr<BackendWrapper> >& inputs,
+                              const std::vector<Ptr<BackendWrapper> >&) CV_OVERRIDE
+    {
+        auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
+        auto inW = inputs[0].dynamicCast<CUDABackendWrapper>();
+        MatShape inShape = inW->getShape();
+        const int nspatial = (int)kernel_shape.size();
+
+        cuda4dnn::PoolingConfiguration config;
+        config.poolMode = cuda4dnn::PoolingConfiguration::PoolingMode::MAX;
+        config.window_size.assign(kernel_shape.begin(), kernel_shape.end());
+        for (int i = 0; i < nspatial; i++)
+            config.strides.push_back(strides.empty() ? 1 : (size_t)strides[i]);
+        config.padMode = (auto_pad == AUTO_PAD_VALID)
+            ? cuda4dnn::PoolingConfiguration::PaddingMode::VALID
+            : cuda4dnn::PoolingConfiguration::PaddingMode::MANUAL;
+        if (config.padMode == cuda4dnn::PoolingConfiguration::PaddingMode::MANUAL) {
+            for (int i = 0; i < nspatial; i++) {
+                config.pads_begin.push_back(pads.empty() ? 0 : (size_t)pads[i]);
+                config.pads_end.push_back(pads.empty() ? 0 : (size_t)pads[i + nspatial]);
+            }
+        }
+        config.roundMode = ceil_mode ? cuda4dnn::PoolingConfiguration::RoundingMode::CEIL
+                                     : cuda4dnn::PoolingConfiguration::RoundingMode::FLOOR;
+        config.input_shape.assign(inShape.begin(), inShape.end());
+        return make_cuda_node<cuda4dnn::PoolingOp>(preferableTarget, std::move(context->cudnn_handle), config);
+    }
+#endif
+
     virtual int64_t getFLOPS(const std::vector<MatShape> &inputs,
                            const std::vector<MatShape> &outputs) const CV_OVERRIDE
     {

modules/dnn/src/net.cpp

@@ -142,6 +142,10 @@ void Net::finalizeNet()
         return;
     }
 #endif
+    // New graph engine: explicitly select per-op executors for the chosen backend/target now,
+    // so the first forward() isn't slowed by it.
+    if (impl->mainGraph)
+        impl->finalize();
 }
 
 void Net::setInputsNames(const std::vector<String>& inputBlobNames)

modules/dnn/src/net_impl.cpp

@@ -274,11 +274,11 @@ Ptr<Layer> Net::Impl::getLayer(int layerId) const
         CV_Assert(0 <= layerId && layerId < totalLayers);
         int graph_ofs = 0;
         for (const Ptr<Graph>& graph : allgraphs) {
-            const std::vector<Ptr<Layer> >& prog = graph->prog();
+            const std::vector<Ptr<OpData> >& prog = graph->prog();
             int nops = (int)prog.size();
             CV_Assert(layerId >= graph_ofs);
             if (layerId < graph_ofs + nops)
-                return prog[layerId - graph_ofs];
+                return prog[layerId - graph_ofs].dynamicCast<Layer>();
             graph_ofs += nops;
         }
         CV_Error_(Error::StsObjectNotFound, ("layer #%d is not found", layerId));
@@ -1679,7 +1679,7 @@ void Net::Impl::setParam(const std::string& outputTensorName, int numParam, cons
         }
 
         int targetIdx = (int)it->second;
-        const std::vector<Ptr<Layer>>& prog = mainGraph->prog();
+        const std::vector<Ptr<OpData>>& prog = mainGraph->prog();
         for (const auto& layer : prog) {
             bool produces = false;
             for (const Arg& out : layer->outputs)
@@ -2385,8 +2385,8 @@ std::vector<String> Net::Impl::getLayerNames() const
     if (mainGraph) {
         res.reserve(totalLayers);
         for (const Ptr<Graph>& graph: allgraphs) {
-            const std::vector<Ptr<Layer> >& prog = graph->prog();
-            for (const Ptr<Layer>& layer: prog)
+            const std::vector<Ptr<OpData> >& prog = graph->prog();
+            for (const Ptr<OpData>& layer: prog)
                 res.push_back(layer->name);
         }
     } else {
@@ -2416,7 +2416,7 @@ std::vector<int> Net::Impl::getUnconnectedOutLayers() const
 
         int graph_ofs = 0;
         for (const auto& graph : allgraphs) {
-            const std::vector<Ptr<Layer>>& prog = graph->prog();
+            const std::vector<Ptr<OpData>>& prog = graph->prog();
             for (int i = 0; i < (int)prog.size(); i++) {
                 for (const auto& layerOut : prog[i]->outputs) {
                     if (outArgIdxs.count(layerOut.idx)) {
@@ -2501,9 +2501,9 @@ int64 Net::Impl::getFLOPSGraph(const Ptr<Graph>& graph,
         return 0;
 
     int64 flops = 0;
-    const std::vector<Ptr<Layer>>& prog = graph->prog();
+    const std::vector<Ptr<OpData>>& prog = graph->prog();
 
-    for (const Ptr<Layer>& layer : prog) {
+    for (const Ptr<OpData>& layer : prog) {
         if (!layer)
             continue;
 
@@ -2595,7 +2595,7 @@ int64 Net::Impl::getFLOPS(
         for (const Ptr<Graph>& graph : allgraphs) {
             int progSize = (int)graph->prog().size();
             if (localIdx < progSize) {
-                const Ptr<Layer>& layer = graph->prog()[localIdx];
+                const Ptr<OpData>& layer = graph->prog()[localIdx];
                 if (!layer)
                     return 0;
 
@@ -2694,8 +2694,8 @@ void Net::Impl::collectLayerInfo(std::vector<String>& names, std::vector<String>
         names.reserve(totalLayers);
         types.reserve(totalLayers);
         for (const Ptr<Graph>& graph : allgraphs) {
-            const std::vector<Ptr<Layer>>& prog = graph->prog();
-            for (const Ptr<Layer>& layer : prog) {
+            const std::vector<Ptr<OpData>>& prog = graph->prog();
+            for (const Ptr<OpData>& layer : prog) {
                 names.push_back(layer ? layer->name : "null");
                 types.push_back(layer ? layer->type : "null");
             }
@@ -3009,8 +3009,8 @@ void Net::Impl::getLayerTypes(std::vector<String>& layersTypes) const
     if (mainGraph) {
         std::set<std::string> layersTypesSet;
         for (const Ptr<Graph>& g: allgraphs) {
-            const std::vector<Ptr<Layer> >& prog = g->prog();
-            for (const Ptr<Layer>& layer: prog) {
+            const std::vector<Ptr<OpData> >& prog = g->prog();
+            for (const Ptr<OpData>& layer: prog) {
                 if (!layer)
                     continue;
                 layersTypesSet.insert(layer->type);
@@ -3042,8 +3042,8 @@ int Net::Impl::getLayersCount(const String& layerType) const
     if (mainGraph) {
         int count = 0;
         for (const Ptr<Graph>& g: allgraphs) {
-            const std::vector<Ptr<Layer> >& prog = g->prog();
-            for (const Ptr<Layer>& layer: prog) {
+            const std::vector<Ptr<OpData> >& prog = g->prog();
+            for (const Ptr<OpData>& layer: prog) {
                 if (!layer)
                     continue;
                 if (layer->type == layerType)

modules/dnn/src/net_impl.hpp

@@ -143,6 +143,9 @@ struct Net::Impl : public detail::NetImplBase
     bool enableFP16, haveFP16;
     bool prepared; // need to rerun graph transformations/optimizations
     bool finalizeLayers; // need to initialize each layer
+    bool finalized = false; // executors have been selected for the current backend/target
+    std::vector<Ptr<BackendWrapper> > argWrappers;
+    std::vector<const void*> argWrapperData;
     TracingMode tracingMode;
     ProfilingMode profilingMode;
     std::vector<int64_t> dimvalues;
@@ -420,12 +423,18 @@ struct Net::Impl : public detail::NetImplBase
     int findDim(const std::string& name, bool insert=false);
 
     void prepareForInference();
+    void finalize();
+    // Selects executors for a single graph (recursing into subgraphs).
+    void finalizeGraph(const Ptr<Graph>& graph, bool useCUDA);
+#ifdef HAVE_CUDA
+    Ptr<BackendWrapper> getCudaArgWrapper(Arg arg, Mat& hostMat);
+#endif
 
     // pre-allocates memory for output tensors.
     // if useBufferPool==true, the method uses 'buffers'
     // for outputs (according to bufidxs)
     // instead of allocating fresh outputs
-    void allocateLayerOutputs(const Ptr<Layer>& layer,
+    void allocateLayerOutputs(const Ptr<OpData>& layer,
                               const std::vector<int>& inpTypes,
                               const std::vector<MatShape>& inpShapes,
                               std::vector<int>& outTypes,
@@ -523,9 +532,9 @@ struct Net::Impl : public detail::NetImplBase
 
 };  // Net::Impl
 
-inline Net::Impl* getNetImpl(const Layer* layer)
+inline Net::Impl* getNetImpl(const OpData* op)
 {
-    return reinterpret_cast<Net::Impl*>(layer->netimpl);
+    return reinterpret_cast<Net::Impl*>(op->netimpl);
 }
 
 Net readNetFromONNX2(const String&);

modules/dnn/src/net_impl2.cpp

@@ -316,30 +316,30 @@ public:
         return g;
     }*/
 
-    virtual const std::vector<Arg>& append(Ptr<Layer>& layer,
+    virtual const std::vector<Arg>& append(Ptr<OpData>& op,
                 const std::vector<std::string>& outnames) override
     {
-        CV_Assert(layer);
+        CV_Assert(op);
         int i, noutputs = (int)outnames.size();
-        //CV_Assert(layer->minNumOutputs() <= noutputs && noutputs <= layer->maxNumOutputs());
+        //CV_Assert(op->minNumOutputs() <= noutputs && noutputs <= op->maxNumOutputs());
 
-        layer->outputs.resize(noutputs);
+        op->outputs.resize(noutputs);
         for (i = 0; i < noutputs; i++) {
             Arg outarg = netimpl_->getArg(outnames[i]);
             ArgKind kind = netimpl_->argKind(outarg);
             CV_Assert(kind == DNN_ARG_TEMP || kind == DNN_ARG_OUTPUT);
-            layer->outputs[i] = outarg;
+            op->outputs[i] = outarg;
         }
 
-        prog_.push_back(layer);
-        return layer->outputs;
+        prog_.push_back(op);
+        return op->outputs;
     }
 
-    virtual Arg append(Ptr<Layer>& layer,
+    virtual Arg append(Ptr<OpData>& op,
                const std::string& outname) override
     {
         std::vector<std::string> outnames = {outname};
-        const std::vector<Arg>& outputs = append(layer, outnames);
+        const std::vector<Arg>& outputs = append(op, outnames);
         CV_Assert(outputs.size() == 1);
         return outputs[0];
     }
@@ -378,8 +378,8 @@ public:
         for (size_t i = 0; i < nlayers; i++) {
             prindent(strm, argindent);
             strm << "// op #" << i << "\n";
-            const Ptr<Layer>& layer = prog_[i];
-            layer->dump(strm, argindent, i+1 < nlayers);
+            const Ptr<OpData>& op = prog_[i];
+            op->dump(strm, argindent, i+1 < nlayers);
         }
         prindent(strm, subindent);
         strm << "]\n";
@@ -400,15 +400,30 @@ public:
         netimpl_->checkArgs(outputs);
         outputs_ = outputs;
     }
-    virtual const std::vector<Ptr<Layer> >& prog() const override { return prog_; }
-    virtual void setProg(const std::vector<Ptr<Layer> >& newprog) override { prog_ = newprog; }
+    virtual const std::vector<Ptr<OpData> >& prog() const override { return prog_; }
+    virtual int opBackend(int opidx) const override
+    {
+        return (opidx >= 0 && opidx < (int)execBackend_.size()) ? execBackend_[opidx]
+                                                                : DNN_BACKEND_OPENCV;
+    }
+    virtual void setProg(const std::vector<Ptr<OpData> >& newprog) override
+    {
+        prog_ = newprog;
+        exec_.clear();
+        execBackend_.clear();
+        inH2D_.clear();
+        outD2H_.clear();
+    }
 
-protected:
     Net::Impl* netimpl_;
     std::string name_;
     std::vector<Arg> inputs_;
     std::vector<Arg> outputs_;
-    std::vector<Ptr<Layer> > prog_;
+    std::vector<Ptr<OpData> > prog_;
+    std::vector<Ptr<Layer> > exec_;
+    std::vector<int> execBackend_;
+    std::vector<std::vector<uchar> > inH2D_;
+    std::vector<std::vector<uchar> > outD2H_;
 };
 
 Ptr<Graph> Graph::create(void* netimpl, const std::string& name,
@@ -556,16 +571,106 @@ void Net::Impl::prepareForInference()
         fuseTransposeMatMul();
         fuseScaleSoftmax();
         fuseBasic();
-        useBlockLayout();
-        assignBuffers();
         totalLayers = updateGraphOfs(mainGraph, 0, true);
         prepared = true;
         finalizeLayers = true;
     }
 }
 
+void Net::Impl::finalizeGraph(const Ptr<Graph>& graph, bool useCUDA)
+{
+    GraphImpl* g = static_cast<GraphImpl*>(graph.get());
+    const std::vector<Ptr<OpData> >& prog = g->prog_;
+    size_t i, nops = prog.size();
+    g->exec_.assign(nops, Ptr<Layer>());
+    g->execBackend_.assign(nops, DNN_BACKEND_OPENCV);
+
+    for (i = 0; i < nops; i++) {
+        const Ptr<OpData>& op = prog[i];
+        if (!op)
+            continue;
+
+        // recurse into subgraphs (If/Loop bodies) first
+        const std::vector<Ptr<Graph> >* subs = op->subgraphs();
+        if (subs) {
+            for (const Ptr<Graph>& sub : *subs)
+                finalizeGraph(sub, useCUDA);
+        }
+
+        Ptr<Layer> exec;
+        int backend = DNN_BACKEND_OPENCV;
+#ifdef HAVE_CUDA
+        // Try the optimized backend first; its create() returns null if the op is unsupported.
+        if (useCUDA && cudaInfo && !subs) {
+            exec = LayerFactory::createExec(op->type, DNN_BACKEND_CUDA, op, &cudaInfo->context);
+            if (exec) {
+                exec->preferableTarget = preferableTarget;
+                backend = DNN_BACKEND_CUDA;
+            }
+        }
+#endif
+        if (!exec) {
+            exec = LayerFactory::createExec(op->type, DNN_BACKEND_OPENCV, op, nullptr);
+            if (!exec)
+                exec = op.dynamicCast<Layer>();
+            backend = DNN_BACKEND_OPENCV;
+        }
+        CV_Assert(exec);
+        g->exec_[i] = exec;
+        g->execBackend_[i] = backend;
+        CV_LOG_INFO(NULL, cv::format("DNN/NewEngine: finalize op #%zu '%s' (%s) -> %s",
+                    i, op->name.c_str(), op->type.c_str(),
+                    backend == DNN_BACKEND_CUDA ? "CUDA" : "CPU"));
+    }
+}
+
+void Net::Impl::finalize()
+{
+#ifdef HAVE_ONNXRUNTIME
+    if (ort_session)
+        return;  // ONNX Runtime manages its own execution session
+#endif
+    if (!mainGraph)
+        return;
+    if (!prepared)
+        prepareForInference();
+    if (finalized)
+        return;
+
+    bool useCUDA = false;
+#ifdef HAVE_CUDA
+    argWrappers.clear();
+    argWrapperData.clear();
+    if (preferableBackend == DNN_BACKEND_CUDA && haveCUDA()) {
+        useCUDA = true;
+        if (!cudaInfo) {
+            cuda4dnn::csl::CSLContext context;
+            context.stream = cuda4dnn::csl::Stream(true);
+            context.cublas_handle = cuda4dnn::csl::cublas::Handle(context.stream);
+            context.cudnn_handle = cuda4dnn::csl::cudnn::Handle(context.stream);
+            auto d2h_stream = cuda4dnn::csl::Stream(true);
+            cudaInfo = std::unique_ptr<CudaInfo_t>(new CudaInfo_t(std::move(context), std::move(d2h_stream)));
+        }
+    }
+#endif
+
+    CV_LOG_INFO(NULL, cv::format("DNN/NewEngine: finalize() backend=%d target=%d useCUDA=%d over %zu graph(s)",
+                preferableBackend, preferableTarget, (int)useCUDA, allgraphs.size()));
+
+    for (const Ptr<Graph>& g : allgraphs)
+        finalizeGraph(g, useCUDA);
+    useBlockLayout();
+    assignBuffers();
+    totalLayers = updateGraphOfs(mainGraph, 0, true);
+
+    for (const Ptr<Graph>& g : allgraphs)
+        finalizeGraph(g, useCUDA);
+
+    finalized = true;
+}
+
 void Net::Impl::allocateLayerOutputs(
-                          const Ptr<Layer>& layer,
+                          const Ptr<OpData>& layer,
                           const std::vector<int>& inpTypes,
                           const std::vector<MatShape>& inpShapes,
                           std::vector<int>& outTypes,
@@ -669,6 +774,7 @@ void Net::Impl::forwardMainGraph(InputArrayOfArrays inputs, OutputArrayOfArrays
     if (!mainGraph) {
         CV_Error(Error::StsNullPtr, "the model was not loaded");
     }
+    finalize();  // select per-op executors for the chosen backend/target (idempotent)
     // ************ uncomment one of the lines below for debugging **********
     //tracingMode = DNN_TRACE_OP;
     //tracingMode = DNN_TRACE_ALL;
@@ -1175,6 +1281,39 @@ static Mat stackScanAxis(const std::vector<Mat>& perIter, int axis, bool reverse
     }
     return stacked;
 }
+#ifdef HAVE_CUDA
+Ptr<BackendWrapper> Net::Impl::getCudaArgWrapper(Arg arg, Mat& hostMat)
+{
+    int idx = arg.idx;
+    if ((int)argWrappers.size() != (int)args.size()) {
+        argWrappers.assign(args.size(), Ptr<BackendWrapper>());
+        argWrapperData.assign(args.size(), nullptr);
+    }
+    Ptr<CUDABackendWrapper> cw = argWrappers[idx].dynamicCast<CUDABackendWrapper>();
+    if (!cw || argWrapperData[idx] != (const void*)hostMat.data) {
+        Ptr<BackendWrapper> w = wrapMat(DNN_BACKEND_CUDA, preferableTarget, hostMat);
+        cw = w.dynamicCast<CUDABackendWrapper>();
+        cw->setStream(cudaInfo->context.stream, cudaInfo->d2h_stream);
+        argWrappers[idx] = w;
+        argWrapperData[idx] = (const void*)hostMat.data;
+    }
+    return argWrappers[idx];
+}
+
+static void forwardOpCUDA(Net::Impl* netimpl, GraphImpl* gimpl, size_t opidx,
+                          const std::vector<Arg>& inputs, const std::vector<Arg>& outputs,
+                          std::vector<Mat>& inpMats, std::vector<Mat>& outMats)
+{
+    Ptr<Layer> exec = gimpl->exec_[opidx];
+    CV_Assert(exec && netimpl->cudaInfo);
+    std::vector<Ptr<BackendWrapper> > inpWrappers(inputs.size()), outWrappers(outputs.size());
+    for (size_t i = 0; i < inputs.size(); i++)
+        inpWrappers[i] = netimpl->getCudaArgWrapper(inputs[i], inpMats[i]);
+    for (size_t i = 0; i < outputs.size(); i++)
+        outWrappers[i] = netimpl->getCudaArgWrapper(outputs[i], outMats[i]);
+    exec->forwardCUDA(inpWrappers, outWrappers, &netimpl->cudaInfo->workspace);
+}
+#endif
 
 void Net::Impl::forwardGraph(Ptr<Graph>& graph, InputArrayOfArrays inputs_,
                              OutputArrayOfArrays outputs_, bool isMainGraph)
@@ -1184,7 +1323,8 @@ void Net::Impl::forwardGraph(Ptr<Graph>& graph, InputArrayOfArrays inputs_,
         CV_Error_(Error::StsObjectNotFound, ("graph '%s' does not belong to the model", graph->name().c_str()));
     }
     std::ostream& strm_ = dump_strm ? *dump_strm : std::cout;
-    const std::vector<Ptr<Layer> >& prog = graph->prog();
+    GraphImpl* gimpl = static_cast<GraphImpl*>(graph.get());
+    const std::vector<Ptr<OpData> >& prog = graph->prog();
     size_t i, nops = prog.size();
     const std::vector<Arg>& gr_inputs = graph->inputs();
     const std::vector<Arg>& gr_outputs = graph->outputs();
@@ -1206,15 +1346,29 @@ void Net::Impl::forwardGraph(Ptr<Graph>& graph, InputArrayOfArrays inputs_,
         for (i = 0; i < n_gr_inputs; i++) {
             Mat m = inputs_.getMat((int)i);
             setGraphInput(graph, i, m);
+#ifdef HAVE_CUDA
+            Arg ginp = gr_inputs[i];
+            if (ginp.idx < (int)argWrappers.size()) {
+                Ptr<CUDABackendWrapper> cw = argWrappers[ginp.idx].dynamicCast<CUDABackendWrapper>();
+                if (cw) cw->setHostDirty();
+            }
+#endif
         }
     }
 
     for (size_t opidx = 0; opidx < nops; opidx++) {
-        const Ptr<Layer>& layer = prog.at(opidx);
-        if (!layer) // in theory we shouldn't have any 'nops' at this stage, but just in case we skip them.
+        const Ptr<OpData>& op = prog.at(opidx);
+        if (!op) // in theory we shouldn't have any 'nops' at this stage, but just in case we skip them.
             continue;
-        const std::vector<Arg>& inputs = layer->inputs;
-        const std::vector<Arg>& outputs = layer->outputs;
+        Ptr<Layer> layer = (opidx < gimpl->exec_.size()) ? gimpl->exec_[opidx] : Ptr<Layer>();
+        if (!layer)
+            layer = op.dynamicCast<Layer>();
+        CV_Assert(layer);
+        int opBackend = (opidx < gimpl->execBackend_.size()) ? gimpl->execBackend_[opidx]
+                                                             : DNN_BACKEND_OPENCV;
+        CV_UNUSED(opBackend);  // only consumed by the HAVE_CUDA dispatch below
+        const std::vector<Arg>& inputs = op->inputs;
+        const std::vector<Arg>& outputs = op->outputs;
         size_t ninputs = inputs.size(), noutputs = outputs.size();
 
         inpMats.resize(ninputs);
@@ -1233,15 +1387,15 @@ void Net::Impl::forwardGraph(Ptr<Graph>& graph, InputArrayOfArrays inputs_,
 
         if (tracingMode != DNN_TRACE_NONE) {
             strm_ << "-----------\n";
-            strm_ << "'" << graph->name() << "' [" << opidx << "/" << nops << "]. " << layer->type << " node: " << layer->name << "\n";
+            strm_ << "'" << graph->name() << "' [" << opidx << "/" << nops << "]. " << op->type << " node: " << op->name << "\n";
             for (i = 0; i < ninputs; i++) {
                 Arg inp = inputs[i];
                 traceArg(strm_, "Input", i, inp, false);
             }
         }
-        bool dynamicOutShapes = layer->dynamicOutputShapes();
+        bool dynamicOutShapes = op->dynamicOutputShapes();
         if (!dynamicOutShapes) {
-            allocateLayerOutputs(layer, inpTypes, inpShapes, outTypes, outShapes, outOrigData, outMats,
+            allocateLayerOutputs(op, inpTypes, inpShapes, outTypes, outShapes, outOrigData, outMats,
                                  tempTypes, tempShapes, tempMats, scratchBufs, true);
         } else {
             outMats.resize(noutputs);
@@ -1254,11 +1408,38 @@ void Net::Impl::forwardGraph(Ptr<Graph>& graph, InputArrayOfArrays inputs_,
 
         timestamp = getTickCount();
 
-        std::vector<Ptr<Graph> >* subgraphs = layer->subgraphs();
+        std::vector<Ptr<Graph> >* subgraphs = op->subgraphs();
         if (!subgraphs) {
-            if (finalizeLayers)
-                layer->finalize(inpMats, outMats);
-            layer->forward(inpMats, outMats, tempMats);
+#ifdef HAVE_CUDA
+            if (opBackend == DNN_BACKEND_CUDA) {
+                if (finalizeLayers)
+                    layer->finalize(inpMats, outMats);
+                forwardOpCUDA(this, gimpl, opidx, inputs, outputs, inpMats, outMats);
+            } else
+#endif
+            {
+#ifdef HAVE_CUDA
+                // CPU op: bring any device-resident inputs back to host before reading them.
+                for (size_t k = 0; k < ninputs; k++) {
+                    if (inputs[k].idx < (int)argWrappers.size()) {
+                        Ptr<CUDABackendWrapper> cw = argWrappers[inputs[k].idx].dynamicCast<CUDABackendWrapper>();
+                        if (cw) { cw->copyToHost(); inpMats[k] = argTensor(inputs[k]); }
+                    }
+                }
+#endif
+                if (finalizeLayers)
+                    layer->finalize(inpMats, outMats);
+                layer->forward(inpMats, outMats, tempMats);
+#ifdef HAVE_CUDA
+                // CPU produced fresh host data; invalidate any stale device copy of its outputs.
+                for (size_t k = 0; k < noutputs; k++) {
+                    if (outputs[k].idx < (int)argWrappers.size()) {
+                        Ptr<CUDABackendWrapper> cw = argWrappers[outputs[k].idx].dynamicCast<CUDABackendWrapper>();
+                        if (cw) cw->setHostDirty();
+                    }
+                }
+#endif
+            }
         }
         else {
             Ptr<IfLayer> iflayer = layer.dynamicCast<IfLayer>();
@@ -1425,7 +1606,7 @@ void Net::Impl::forwardGraph(Ptr<Graph>& graph, InputArrayOfArrays inputs_,
             }
             else {
                 CV_Error_(Error::StsNotImplemented,
-                          ("unknown layer type '%s' with subgraphs", layer->type.c_str()));
+                          ("unknown layer type '%s' with subgraphs", op->type.c_str()));
             }
         }
         CV_Assert(outMats.size() == noutputs);
@@ -1508,6 +1689,13 @@ void Net::Impl::forwardGraph(Ptr<Graph>& graph, InputArrayOfArrays inputs_,
     outputsVec.resize(n_gr_outputs);
     for (i = 0; i < n_gr_outputs; i++) {
         Arg out = gr_outputs[i];
+#ifdef HAVE_CUDA
+        // A graph output produced on the device must be brought back to host before it is read.
+        if (out.idx < (int)argWrappers.size()) {
+            Ptr<CUDABackendWrapper> cw = argWrappers[out.idx].dynamicCast<CUDABackendWrapper>();
+            if (cw) cw->copyToHost();
+        }
+#endif
         const Mat& outm = argTensor(out);
         if (isMainGraph) {
             if (outm.size.layout == DATA_LAYOUT_BLOCK) {
@@ -1532,8 +1720,8 @@ void Net::Impl::updateUseCounts(const Ptr<Graph>& graph, std::vector<int>& useco
         CV_Assert(output.idx < (int)usecounts.size());
         usecounts[output.idx]++;
     }
-    const std::vector<Ptr<Layer> >& prog = graph->prog();
-    for (const Ptr<Layer>& layer: prog) {
+    const std::vector<Ptr<OpData> >& prog = graph->prog();
+    for (const Ptr<OpData>& layer: prog) {
         const std::vector<Arg>& inputs = layer->inputs;
         for (const Arg& input: inputs) {
             CV_Assert(input.idx < (int)usecounts.size());
@@ -1564,14 +1752,14 @@ int Net::Impl::updateGraphOfs(const Ptr<Graph>& graph, int currofs, bool ismain)
         allgraphs.clear();
         layerNameToId.clear();
     }
-    const std::vector<Ptr<Layer> >& prog = graph->prog();
+    const std::vector<Ptr<OpData> >& prog = graph->prog();
     size_t i, nops = prog.size();
     int subgraph_ofs = currofs + (int)nops;
     std::string name = graph->name();
     graphofs.insert(std::make_pair(name, currofs));
     allgraphs.push_back(graph);
     for (i = 0; i < nops; i++) {
-        const Ptr<Layer>& layer = prog[i];
+        const Ptr<OpData>& layer = prog[i];
         layerNameToId.insert(std::make_pair(layer->name, currofs + (int)i));
         const std::vector<Ptr<Graph> >* subgraphs = layer->subgraphs();
         if (subgraphs) {
@@ -1718,12 +1906,12 @@ bool Net::Impl::tryInferGraphShapes(const Ptr<Graph>& graph,
     if (!graph)
         return true;
 
-    const std::vector<Ptr<Layer> >& prog = graph->prog();
+    const std::vector<Ptr<OpData> >& prog = graph->prog();
 
     std::vector<MatShape> inpShapes, outShapes, tempShapes;
     std::vector<int> inpTypes, outTypes, tempTypes;
 
-    for (const Ptr<Layer>& layer: prog) {
+    for (const Ptr<OpData>& layer: prog) {
         if (!layer)
             continue;
 

modules/dnn/src/net_impl_backend.cpp

@@ -299,6 +299,18 @@ void Net::Impl::setPreferableBackend(Net& net, int backendId)
 
     if (mainGraph)
     {
+        if (backendId == DNN_BACKEND_OPENCV
+#ifdef HAVE_CUDA
+            || backendId == DNN_BACKEND_CUDA
+#endif
+            )
+        {
+            if (preferableBackend != backendId) {
+                preferableBackend = backendId;
+                finalized = false;  // re-select per-op executors on next finalize()
+            }
+            return;
+        }
         CV_LOG_WARNING(NULL, "Back-ends are not supported by the new graph engine for now");
         preferableBackend = backendId;
         return;
@@ -347,10 +359,17 @@ void Net::Impl::setPreferableTarget(int targetId)
 
     if (mainGraph)
     {
-        if (targetId != DNN_TARGET_CPU)
+#ifdef HAVE_CUDA
+        if (targetId == DNN_TARGET_CPU || IS_DNN_CUDA_TARGET(targetId))
         {
-            CV_LOG_WARNING(NULL, "Targets are not supported by the new graph engine for now");
+            if (preferableTarget != targetId) {
+                preferableTarget = targetId;
+                finalized = false;  // re-select per-op executors on next finalize()
+            }
+            return;
         }
+#endif
+        CV_LOG_WARNING(NULL, "Targets are not supported by the new graph engine for now");
         return;
     }
     if (netWasQuantized && targetId != DNN_TARGET_CPU &&

modules/dnn/src/onnx/onnx_importer2.cpp

@@ -167,7 +167,7 @@ protected:
     std::string onnxBasePath;
     Ptr<Graph> curr_graph;
     opencv_onnx::GraphProto* curr_graph_proto;
-    std::vector<Ptr<Layer> > curr_prog;
+    std::vector<Ptr<OpData> > curr_prog;
     std::vector<Arg> node_inputs, node_outputs;
 
     std::string framework_name;
@@ -892,7 +892,7 @@ Ptr<Graph> ONNXImporter2::parseGraph(opencv_onnx::GraphProto* graph_proto, bool
 
     opencv_onnx::GraphProto* saved_graph_proto = curr_graph_proto;
     Ptr<Graph> saved_graph = curr_graph;
-    std::vector<Ptr<Layer> > saved_prog;
+    std::vector<Ptr<OpData> > saved_prog;
 
     curr_graph_proto = graph_proto;
     std::vector<Arg> inputs, outputs;
@@ -1504,9 +1504,12 @@ void ONNXImporter2::parseGemm(LayerParams& layerParams, const opencv_onnx::NodeP
     if (net.isConstArg(node_inputs[1]) && (n_inputs == 2 || net.isConstArg(node_inputs[2]))) {
         Mat B = net.argTensor(node_inputs[1]);
         layerParams.blobs.push_back(B);
+        layerParams.set("constB", true);  // weight folded into blobs[0] (enables CUDA InnerProduct)
         if (n_inputs > 2) {
             Mat bias = net.argTensor(node_inputs[2]);
             layerParams.blobs.push_back(bias);
+            layerParams.set("have_bias", true);
+            layerParams.set("constC", true);
         }
         n_inputs = 1;
     }
@@ -1721,7 +1724,7 @@ void ONNXImporter2::parseLoop(LayerParams& layerParams,
 
     CV_Assert(!subgraphs[0].empty());
 
-    Ptr<Layer>& loopLayer = curr_prog.back();
+    Ptr<OpData>& loopLayer = curr_prog.back();
     *loopLayer->subgraphs() = subgraphs;
 }
 
@@ -1786,7 +1789,7 @@ void ONNXImporter2::parseIf(LayerParams& layerParams,
 
     CV_Assert_N(!thenelse[0].empty(), !thenelse[1].empty());
 
-    Ptr<Layer>& ifLayer = curr_prog.back();
+    Ptr<OpData>& ifLayer = curr_prog.back();
     *ifLayer->subgraphs() = thenelse;
 }
 

modules/dnn/src/op_cuda.cpp

@@ -8,10 +8,70 @@
 #include "op_cuda.hpp"
 #include "cuda4dnn/init.hpp"
 #include "net_impl.hpp"
+#include <opencv2/dnn/layer.details.hpp>
 
 namespace cv { namespace dnn {
 CV__DNN_INLINE_NS_BEGIN
 
+class CUDALegacyExec : public Layer
+{
+public:
+    CUDALegacyExec(const Ptr<Layer>& impl_, void* ctx_) : impl(impl_), ctx(ctx_) {}
+
+    static Ptr<Layer> create(const Ptr<OpData>& data, void* backendCtx)
+    {
+        Ptr<Layer> impl = data.dynamicCast<Layer>();
+        if (!impl || !backendCtx || !impl->supportBackend(DNN_BACKEND_CUDA))
+            return Ptr<Layer>();  // unsupported -> CPU fallback
+        Ptr<CUDALegacyExec> e(new CUDALegacyExec(impl, backendCtx));
+        e->data = data;
+        e->name = impl->name;
+        e->type = impl->type;
+        e->inputs = impl->inputs;
+        e->outputs = impl->outputs;
+        return e;
+    }
+
+    void finalize(InputArrayOfArrays inputs, OutputArrayOfArrays outputs) CV_OVERRIDE
+    {
+        impl->finalize(inputs, outputs);
+    }
+
+    void forwardCUDA(const std::vector<Ptr<BackendWrapper> >& inputs,
+                     const std::vector<Ptr<BackendWrapper> >& outputs,
+                     void* workspace) CV_OVERRIDE
+    {
+        cuda4dnn::csl::Workspace& ws = *reinterpret_cast<cuda4dnn::csl::Workspace*>(workspace);
+        if (!node) {
+            impl->preferableTarget = preferableTarget;  // initCUDA may pick FP16/FP32 by target
+            cuda4dnn::csl::CSLContext context = *reinterpret_cast<cuda4dnn::csl::CSLContext*>(ctx);
+            node = impl->initCUDA(&context, inputs, outputs);
+            CV_Assert(node);
+            cudaNode = node.dynamicCast<CUDABackendNode>();
+            CV_Assert(cudaNode);
+            ws.require(cudaNode->get_workspace_memory_in_bytes());
+        }
+        cudaNode->forward(inputs, outputs, ws);
+    }
+
+    Ptr<Layer> impl;
+    void* ctx;
+    Ptr<BackendNode> node;
+    Ptr<CUDABackendNode> cudaNode;
+};
+
+void registerCudaCommonExecs()
+{
+    CV_DNN_REGISTER_EXEC_CLASS(ReLU,        DNN_BACKEND_CUDA, CUDALegacyExec);
+    CV_DNN_REGISTER_EXEC_CLASS(ReLU6,       DNN_BACKEND_CUDA, CUDALegacyExec);
+    CV_DNN_REGISTER_EXEC_CLASS(NaryEltwise, DNN_BACKEND_CUDA, CUDALegacyExec);
+    CV_DNN_REGISTER_EXEC_CLASS(Flatten,     DNN_BACKEND_CUDA, CUDALegacyExec);
+    CV_DNN_REGISTER_EXEC_CLASS(BatchNorm2,  DNN_BACKEND_CUDA, CUDALegacyExec);
+    CV_DNN_REGISTER_EXEC_CLASS(MaxPool,     DNN_BACKEND_CUDA, CUDALegacyExec);
+    CV_DNN_REGISTER_EXEC_CLASS(Gemm,        DNN_BACKEND_CUDA, CUDALegacyExec);
+    CV_DNN_REGISTER_EXEC_CLASS(Pooling,     DNN_BACKEND_CUDA, CUDALegacyExec);  // GlobalAveragePool/GlobalMaxPool
+}
+
 
 void Net::Impl::initCUDABackend(const std::vector<LayerPin>& blobsToKeep_)
 {

modules/dnn/src/tensorflow/tf_importer.cpp

@@ -576,7 +576,7 @@ protected:
     std::map<String, int> layer_id;
 
     bool newEngine;
-    std::vector<Ptr<Layer>> curProg;
+    std::vector<Ptr<OpData>> curProg;
     std::vector<std::vector<std::string>> layersOutputs;
     std::vector<Arg> modelInputs;
     std::unordered_map<std::string, MatShape> tensorsShape;

modules/dnn/src/tflite/tflite_importer.cpp

@@ -34,7 +34,7 @@ private:
     const flatbuffers::Vector<flatbuffers::Offset<opencv_tflite::Tensor> >* modelTensors;
     std::map<int, Mat> allTensors;
     Net& dstNet;
-    std::vector<Ptr<Layer>> curProg;
+    std::vector<Ptr<OpData>> curProg;
 
     // This is a vector of pairs (layerId, outputId) where we iterate over
     // indices from TFLite notation and get created OpenCV layers.

modules/dnn/test/test_model.cpp

@@ -690,17 +690,31 @@ static void topK(const Mat& probs, std::vector<std::pair<int, float> >& result,
     }
 }
 
-typedef testing::TestWithParam<Target> Reproducibility_ResNet50_ONNX;
+// Returns the CPU (OpenCV backend) and CUDA backend/target pairs for benchmarking.
+static std::vector<tuple<Backend, Target> > resnet50BackendsAndTargets()
+{
+    std::vector<tuple<Backend, Target> > targets;
+    targets.push_back(make_tuple(DNN_BACKEND_OPENCV, DNN_TARGET_CPU));
+#ifdef HAVE_CUDA
+    for (auto target : getAvailableTargets(DNN_BACKEND_CUDA))
+        targets.push_back(make_tuple(DNN_BACKEND_CUDA, target));
+#endif
+    return targets;
+}
+
+typedef testing::TestWithParam<tuple<Backend, Target> > Reproducibility_ResNet50_ONNX;
 TEST_P(Reproducibility_ResNet50_ONNX, Accuracy)
 {
-    Target targetId = GetParam();
+    Backend backendId = get<0>(GetParam());
+    Target targetId = get<1>(GetParam());
     applyTestTag(targetId == DNN_TARGET_CPU ? CV_TEST_TAG_MEMORY_512MB : CV_TEST_TAG_MEMORY_1GB);
-    ASSERT_TRUE(ocl::useOpenCL() || targetId == DNN_TARGET_CPU || targetId == DNN_TARGET_CPU_FP16);
+    ASSERT_TRUE(ocl::useOpenCL() || targetId == DNN_TARGET_CPU || targetId == DNN_TARGET_CPU_FP16
+                || backendId == DNN_BACKEND_CUDA);
 
     std::string modelname = _tf("onnx/models/resnet50v1.onnx", false);
     Net net = readNetFromONNX(modelname);
 
-    net.setPreferableBackend(DNN_BACKEND_OPENCV);
+    net.setPreferableBackend(backendId);
     net.setPreferableTarget(targetId);
 
     if (targetId == DNN_TARGET_CPU_FP16)
@@ -738,9 +752,37 @@ TEST_P(Reproducibility_ResNet50_ONNX, Accuracy)
     for (int i = 0; i < K; i++) {
         EXPECT_NEAR(ref[i].second, res[i].second, eps);
     }
+
+    // Benchmark: warmup runs followed by timed runs, reporting avg/min/max forward time.
+    const int numWarmup = 5;
+    const int numRuns = 30;
+    for (int i = 0; i < numWarmup; i++)
+    {
+        net.setInput(input);
+        net.forward();
+    }
+
+    double timeMin = DBL_MAX, timeMax = 0.0, timeSum = 0.0;
+    for (int i = 0; i < numRuns; i++)
+    {
+        net.setInput(input);
+        TickMeter tm;
+        tm.start();
+        net.forward();
+        tm.stop();
+        double t = tm.getTimeMilli();
+        timeSum += t;
+        timeMin = std::min(timeMin, t);
+        timeMax = std::max(timeMax, t);
+    }
+
+    std::cout << "[ BENCHMARK ] ResNet50 ONNX (backend=" << backendId << ", target=" << targetId << ") over "
+              << numRuns << " runs: avg=" << (timeSum / numRuns) << " ms"
+              << ", min=" << timeMin << " ms"
+              << ", max=" << timeMax << " ms" << std::endl;
 }
 INSTANTIATE_TEST_CASE_P(/**/, Reproducibility_ResNet50_ONNX,
-                        testing::ValuesIn(getAvailableTargets(DNN_BACKEND_OPENCV)));
+                        testing::ValuesIn(resnet50BackendsAndTargets()));
 
 typedef testing::TestWithParam<Target> Reproducibility_ResNet50_QDQ_ONNX;
 TEST_P(Reproducibility_ResNet50_QDQ_ONNX, Accuracy)