From 72ffdfc170c3f750ac206584249a43e2b83e11e2 Mon Sep 17 00:00:00 2001 From: Abhishek Gola Date: Mon, 8 Jun 2026 18:06:28 +0530 Subject: [PATCH] cuda support and Layer Split + per-op executors --- modules/dnn/include/opencv2/dnn/dnn.hpp | 192 ++++++------- .../dnn/include/opencv2/dnn/layer.details.hpp | 30 ++ modules/dnn/include/opencv2/dnn/layer.hpp | 13 + .../src/cuda4dnn/csl/cudnn/convolution.hpp | 8 +- modules/dnn/src/graph_block_layout.cpp | 13 +- modules/dnn/src/graph_buffer_allocator.cpp | 2 +- modules/dnn/src/graph_const_args.cpp | 6 +- modules/dnn/src/graph_const_fold.cpp | 14 +- modules/dnn/src/graph_fusion_attention.cpp | 64 ++--- modules/dnn/src/graph_fusion_basic.cpp | 32 +-- .../dnn/src/graph_fusion_matmul_to_gemm.cpp | 8 +- modules/dnn/src/graph_fusion_qdq.cpp | 67 ++--- .../src/graph_fusion_reshape_transpose.cpp | 12 +- .../dnn/src/graph_fusion_scale_softmax.cpp | 8 +- modules/dnn/src/graph_fusion_shared_gemm.cpp | 22 +- .../dnn/src/graph_fusion_transpose_matmul.cpp | 8 +- modules/dnn/src/init.cpp | 15 + modules/dnn/src/layer.cpp | 68 +++-- modules/dnn/src/layer_factory.cpp | 66 +++++ modules/dnn/src/layers/batch_norm2_layer.cpp | 24 +- modules/dnn/src/layers/conv2_layer.cpp | 147 ++++++++++ modules/dnn/src/layers/maxpool_layer.cpp | 43 +++ modules/dnn/src/net.cpp | 4 + modules/dnn/src/net_impl.cpp | 30 +- modules/dnn/src/net_impl.hpp | 15 +- modules/dnn/src/net_impl2.cpp | 262 +++++++++++++++--- modules/dnn/src/net_impl_backend.cpp | 23 +- modules/dnn/src/onnx/onnx_importer2.cpp | 11 +- modules/dnn/src/op_cuda.cpp | 60 ++++ modules/dnn/src/tensorflow/tf_importer.cpp | 2 +- modules/dnn/src/tflite/tflite_importer.cpp | 2 +- modules/dnn/test/test_model.cpp | 52 +++- 32 files changed, 995 insertions(+), 328 deletions(-) diff --git a/modules/dnn/include/opencv2/dnn/dnn.hpp b/modules/dnn/include/opencv2/dnn/dnn.hpp index 00f086e516..83a2bbef24 100644 --- a/modules/dnn/include/opencv2/dnn/dnn.hpp +++ b/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 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 blobs; std::vector inputs; std::vector outputs; - void* netimpl; + void* netimpl = nullptr; + + CV_PROP String name; + CV_PROP String type; virtual std::vector >* 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 &inputs, + const int requiredOutputs, + std::vector &outputs, + std::vector &internals) const; + + virtual void getTypes(const std::vector& inputs, + const int requiredOutputs, + const int requiredInternals, + std::vector&outputs, + std::vector&internals) const; + + virtual int getLayouts(const std::vector& actualInputs, + std::vector& desiredInputs, + const int requiredOutputs, + std::vector& outputs) const; + + virtual int64 getFLOPS(const std::vector &inputs, + const std::vector &outputs) const; + + virtual bool updateMemoryShapes(const std::vector &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 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 &inputs, CV_OUT std::vector &outputs, CV_IN_OUT std::vector &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 "%[.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& 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 >& inputs, + const std::vector >& outputs, + void* workspace); /** * @brief "Detaches" all the layers, attached to particular layer. */ virtual void unsetAttached(); - virtual bool getMemoryShapes(const std::vector &inputs, - const int requiredOutputs, - std::vector &outputs, - std::vector &internals) const; - - virtual void getTypes(const std::vector& inputs, - const int requiredOutputs, - const int requiredInternals, - std::vector&outputs, - std::vector&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& actualInputs, - std::vector& desiredInputs, - const int requiredOutputs, - std::vector& outputs) const; - - virtual int64 getFLOPS(const std::vector &inputs, - const std::vector &outputs) const; - - virtual bool updateMemoryShapes(const std::vector &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 ¶ms); //!< Initializes only #name, #type and #blobs fields. - void setParamsFrom(const LayerParams ¶ms); //!< 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& append(Ptr& layer, + virtual const std::vector& append(Ptr& op, const std::vector& outnames=std::vector()) = 0; - virtual Arg append(Ptr& layer, const std::string& outname=std::string()) = 0; + virtual Arg append(Ptr& op, const std::string& outname=std::string()) = 0; virtual std::ostream& dump(std::ostream& strm, int indent, bool comma) = 0; virtual const std::vector& inputs() const = 0; virtual const std::vector& outputs() const = 0; virtual void setOutputs(const std::vector& outputs) = 0; - virtual const std::vector >& prog() const = 0; - virtual void setProg(const std::vector >& newprog) = 0; + virtual const std::vector >& prog() const = 0; + virtual void setProg(const std::vector >& 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. diff --git a/modules/dnn/include/opencv2/dnn/layer.details.hpp b/modules/dnn/include/opencv2/dnn/layer.details.hpp index 1133da562e..0fa36c1264 100644 --- a/modules/dnn/include/opencv2/dnn/layer.details.hpp +++ b/modules/dnn/include/opencv2/dnn/layer.details.hpp @@ -45,6 +45,24 @@ Ptr __LayerStaticRegisterer_func_##type(LayerParams ¶ms) \ { return Ptr(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 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); + +/** @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 create(const Ptr&, 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); + namespace details { template @@ -53,6 +71,18 @@ Ptr _layerDynamicRegisterer(LayerParams ¶ms) return Ptr(LayerClass::create(params)); } +template +Ptr _opDynamicRegisterer(const LayerParams ¶ms) +{ + return Ptr(OpClass::create(params)); +} + +template +Ptr _execDynamicRegisterer(const Ptr& data, void* backendCtx) +{ + return Ptr(ExecClass::create(data, backendCtx)); +} + //allows automatically register created layer on module load time class _LayerStaticRegisterer { diff --git a/modules/dnn/include/opencv2/dnn/layer.hpp b/modules/dnn/include/opencv2/dnn/layer.hpp index a4d167564d..196d86eb20 100644 --- a/modules/dnn/include/opencv2/dnn/layer.hpp +++ b/modules/dnn/include/opencv2/dnn/layer.hpp @@ -76,6 +76,19 @@ public: */ static Ptr createLayerInstance(const String &type, LayerParams& params); + + // Builds the abstract (metadata) node for an operation type. + typedef Ptr(*OpConstructor)(const LayerParams& params); + //! Builds a backend-specific executor from an OpData; returns null Ptr if unsupported. + typedef Ptr(*ExecConstructor)(const Ptr& data, void* backendCtx); + + static void registerOp(const String& type, OpConstructor constructor); + static Ptr createOp(const String& type, const LayerParams& params); + + static void registerExec(const String& type, int backendId, ExecConstructor constructor); + static Ptr createExec(const String& type, int backendId, + const Ptr& data, void* backendCtx); + private: LayerFactory(); }; diff --git a/modules/dnn/src/cuda4dnn/csl/cudnn/convolution.hpp b/modules/dnn/src/cuda4dnn/csl/cudnn/convolution.hpp index 93f3101bf6..ffe46d6757 100644 --- a/modules/dnn/src/cuda4dnn/csl/cudnn/convolution.hpp +++ b/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::value) diff --git a/modules/dnn/src/graph_block_layout.cpp b/modules/dnn/src/graph_block_layout.cpp index 3a90ba794e..f3d1b0bd86 100644 --- a/modules/dnn/src/graph_block_layout.cpp +++ b/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; +using PLayer = Ptr; using PGraph = Ptr; /* Inserts layout conversion operations (if needed) into the model graph and subgraphs. @@ -121,13 +121,15 @@ struct BlockLayoutTransformer std::vector 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& inputs = layer->inputs; const vector& outputs = layer->outputs; size_t ninputs = inputs.size(), noutputs = outputs.size(); std::string op_name = layer->type; std::string name = layer->name; vector* 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++) { diff --git a/modules/dnn/src/graph_buffer_allocator.cpp b/modules/dnn/src/graph_buffer_allocator.cpp index 2c4b5f500e..3a58d89556 100644 --- a/modules/dnn/src/graph_buffer_allocator.cpp +++ b/modules/dnn/src/graph_buffer_allocator.cpp @@ -226,7 +226,7 @@ struct BufferAllocator } } } - const std::vector >& prog = graph->prog(); + const std::vector >& prog = graph->prog(); for (const auto& layer: prog) { bool inplace = false; Arg reuseArg; diff --git a/modules/dnn/src/graph_const_args.cpp b/modules/dnn/src/graph_const_args.cpp index 706bd67538..1bb144b50f 100644 --- a/modules/dnn/src/graph_const_args.cpp +++ b/modules/dnn/src/graph_const_args.cpp @@ -36,14 +36,14 @@ struct ConstArgs void processGraph(Ptr& graph) { - const std::vector >& prog = graph->prog(); + const std::vector >& prog = graph->prog(); size_t i, nops = prog.size(); std::vector removed_args; std::vector saved_tail_inputs; for (i = 0; i < nops; i++) { - const Ptr& layer = prog[i]; - Layer* layer_ptr = const_cast(layer.get()); + const Ptr& layer = prog[i]; + OpData* layer_ptr = const_cast(layer.get()); std::vector >* subgraphs = layer->subgraphs(); if (subgraphs) { for (Ptr& g: *subgraphs) { diff --git a/modules/dnn/src/graph_const_fold.cpp b/modules/dnn/src/graph_const_fold.cpp index 7b999bc64e..442e5f0df8 100644 --- a/modules/dnn/src/graph_const_fold.cpp +++ b/modules/dnn/src/graph_const_fold.cpp @@ -30,7 +30,7 @@ struct ConstFolding netimpl->scratchBufs.clear(); } - Layer* getLayer(std::vector >& newprog, int op_idx) const + OpData* getLayer(std::vector >& 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 >& prog = graph->prog(); + const std::vector >& prog = graph->prog(); size_t i, nops = prog.size(); - std::vector > newprog; + std::vector > newprog; std::vector removed_args; std::vector inpMats, tempMats; std::vector inpTypes, outTypes, tempTypes; std::vector inpShapes, outShapes, tempShapes; for (i = 0; i < nops; i++) { - const Ptr& layer = prog[i]; + const Ptr& layer = prog[i]; std::vector >* subgraphs = layer->subgraphs(); if (subgraphs) { for (Ptr& 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 execLayer = layer.dynamicCast(); + 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]; diff --git a/modules/dnn/src/graph_fusion_attention.cpp b/modules/dnn/src/graph_fusion_attention.cpp index 9f6cb3e1b8..fc5ab02cd9 100644 --- a/modules/dnn/src/graph_fusion_attention.cpp +++ b/modules/dnn/src/graph_fusion_attention.cpp @@ -32,35 +32,35 @@ struct ModelFusionAttention return it->second[0]; } - bool isReshape(const vector>& prog, int idx) const + bool isReshape(const vector>& prog, int idx) const { if (idx < 0 || idx >= (int)prog.size() || !prog[idx]) return false; return dynamic_cast(prog[idx].get()) != nullptr; } - bool isTranspose(const vector>& prog, int idx) const + bool isTranspose(const vector>& prog, int idx) const { if (idx < 0 || idx >= (int)prog.size() || !prog[idx]) return false; return dynamic_cast(prog[idx].get()) != nullptr; } - bool isSoftmax(const vector>& prog, int idx) const + bool isSoftmax(const vector>& prog, int idx) const { if (idx < 0 || idx >= (int)prog.size() || !prog[idx]) return false; return prog[idx]->type == "Softmax"; } - bool isMatMul(const vector>& prog, int idx) const + bool isMatMul(const vector>& prog, int idx) const { if (idx < 0 || idx >= (int)prog.size() || !prog[idx]) return false; return dynamic_cast(prog[idx].get()) != nullptr; } - static bool isProjCandidate(const Ptr& l) + static bool isProjCandidate(const Ptr& l) { if (l->blobs.empty() || l->inputs.size() != 1) return false; if (dynamic_cast(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& l) + static Mat getProjWeight(const Ptr& l) { const Mat& W = l->blobs[0]; GemmLayer* g = dynamic_cast(l.get()); @@ -86,7 +86,7 @@ struct ModelFusionAttention return W; } - bool isScalarBinOp(const vector>& prog, int idx, + bool isScalarBinOp(const vector>& 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>& prog, int idx, float* val) const + bool isScalarMul(const vector>& prog, int idx, float* val) const { return isScalarBinOp(prog, idx, NaryEltwiseLayer::OPERATION::PROD, val); } - bool isScalarDiv(const vector>& prog, int idx, float* val) const + bool isScalarDiv(const vector>& 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>& prog, Arg arg, + bool isRuntimeQKScaleChain(const vector>& prog, Arg arg, std::set& chain_ops) const { const std::vector 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& l = prog[idx]; + const Ptr& l = prog[idx]; if (want == "NaryEltwise") { NaryEltwiseLayer* elt = dynamic_cast(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>& prog, int idx, Arg* out_mask) const + bool isMaskAdd(const vector>& 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>& prog, Arg a) const + int extractConstInt(const vector>& 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>& prog, int concat_idx, + void collectShapeChain(const vector>& prog, int concat_idx, std::set& chain) const { if (concat_idx < 0 || concat_idx >= (int)prog.size() || !prog[concat_idx]) @@ -238,7 +238,7 @@ struct ModelFusionAttention } template - int findMatchingConsumer(const vector>& prog, Arg out, + int findMatchingConsumer(const vector>& prog, Arg out, Pred pred, std::set* extra_shape_ops) const { auto it = consumers_.find(out.idx); @@ -258,7 +258,7 @@ struct ModelFusionAttention return matched; } - int followProjChain(const vector>& prog, + int followProjChain(const vector>& 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(L) != nullptr; }, + [](OpData* L){ return dynamic_cast(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>& prog, int qkv_matmul_idx, + bool tryFuseCombinedQKV(const vector>& prog, int qkv_matmul_idx, std::set& removed_ops, - vector>>& replacements) + vector>>& 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(L) != nullptr; }, + [](OpData* L){ return dynamic_cast(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 attn_layer = LayerFactory::createLayerInstance(attn_params.type, attn_params); + Ptr 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>& prog, Arg arg, + int traceClipBranch(const vector>& 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>& prog, int softmax_idx, + bool tryFuseClipAttention(const vector>& prog, int softmax_idx, std::set& removed_ops, - vector>>& replacements) + vector>>& 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 attn_layer = + Ptr 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) { - const vector>& prog = graph->prog(); + const vector>& 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(L) != nullptr || dynamic_cast(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(L) != nullptr || dynamic_cast(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(L) != nullptr; }, + [](OpData* L){ return dynamic_cast(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 attn_layer = LayerFactory::createLayerInstance( + Ptr attn_layer = LayerFactory::createLayerInstance( attn_params.type, attn_params); CV_Assert(attn_layer); @@ -1157,7 +1157,7 @@ struct ModelFusionAttention } if (modified) { - vector> newprog; + vector> 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 producer_; std::map> consumers_; - vector>> attention_replacements_; + vector>> attention_replacements_; }; void Net::Impl::fuseAttention() diff --git a/modules/dnn/src/graph_fusion_basic.cpp b/modules/dnn/src/graph_fusion_basic.cpp index 2d6bd21c60..7e32685b63 100644 --- a/modules/dnn/src/graph_fusion_basic.cpp +++ b/modules/dnn/src/graph_fusion_basic.cpp @@ -30,7 +30,7 @@ struct ModelFusionBasic } template _LayerType* - getLayer(std::vector >& newprog, int op_idx) const + getLayer(std::vector >& 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 removed_args; bool modified = false; - const std::vector >& prog = graph->prog(); + const std::vector >& prog = graph->prog(); size_t i, nargs = netimpl->args.size(), nops = prog.size(); std::vector producer_of(nargs, -1); - std::vector > newprog; + std::vector > newprog; std::vector fused_inputs; for (i = 0; i < nops; i++) { - const Ptr& layer = prog[i]; + const Ptr& layer = prog[i]; Layer* layer_ptr = (Layer*)layer.get(); int fused_layer_idx = -1; std::vector >* subgraphs = layer->subgraphs(); @@ -74,7 +74,7 @@ struct ModelFusionBasic int conv_layer_idx = producer_of.at(bn_inp.idx); Conv2Layer* conv = getLayer(newprog, conv_layer_idx); if (conv) { - bool ok = conv->fuseBatchNorm(layer); + bool ok = conv->fuseBatchNorm(layer.dynamicCast()); 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(newprog, conv_layer_idx); if (conv) { - bool ok = conv->fuseActivation(layer); + bool ok = conv->fuseActivation(layer.dynamicCast()); 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 gnlayer = GroupNormLayer::create(gnparams); + Ptr 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(); - newprog[reshape2_idx] = Ptr(); - newprog[mul_idx] = Ptr(); + newprog[reshape1_idx] = Ptr(); + newprog[reshape2_idx] = Ptr(); + newprog[mul_idx] = Ptr(); 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) { - const std::vector >& prog = graph->prog(); + const std::vector >& prog = graph->prog(); size_t nops = prog.size(), nargs = netimpl->args.size(); - std::vector > newprog; + std::vector > newprog; newprog.reserve(nops); std::vector producer_of((int)nargs, -1); bool modified = false; for (size_t i = 0; i < nops; i++) { - const Ptr& layer = prog[i]; - Layer* layer_ptr = const_cast(layer.get()); + const Ptr& layer = prog[i]; + Layer* layer_ptr = (Layer*)layer.get(); std::vector >* 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(); + newprog[bn_idx] = Ptr(); modified = true; } } diff --git a/modules/dnn/src/graph_fusion_matmul_to_gemm.cpp b/modules/dnn/src/graph_fusion_matmul_to_gemm.cpp index b3f8ee4096..7eac023f0e 100644 --- a/modules/dnn/src/graph_fusion_matmul_to_gemm.cpp +++ b/modules/dnn/src/graph_fusion_matmul_to_gemm.cpp @@ -43,7 +43,7 @@ struct ModelFusionMatMulToGemm bool fuseGraph(Ptr& graph) { - const vector>& prog = graph->prog(); + const vector>& prog = graph->prog(); size_t nops = prog.size(); bool modified = false; @@ -56,11 +56,11 @@ struct ModelFusionMatMulToGemm } } - vector> newprog = prog; + vector> newprog = prog; bool changed = false; for (size_t i = 0; i < nops; i++) { - const Ptr& layer = newprog[i]; + const Ptr& layer = newprog[i]; if (!layer) continue; MatMulLayer* mm = dynamic_cast(layer.get()); @@ -120,7 +120,7 @@ struct ModelFusionMatMulToGemm gp.blobs.push_back(B); if (have_bias) gp.blobs.push_back(layer->blobs[1]); - Ptr gemm = LayerFactory::createLayerInstance("Gemm", gp); + Ptr gemm = LayerFactory::createLayerInstance("Gemm", gp); if (!gemm) continue; gemm->inputs = layer->inputs; gemm->outputs = layer->outputs; diff --git a/modules/dnn/src/graph_fusion_qdq.cpp b/modules/dnn/src/graph_fusion_qdq.cpp index 667a28a636..4c286b0817 100644 --- a/modules/dnn/src/graph_fusion_qdq.cpp +++ b/modules/dnn/src/graph_fusion_qdq.cpp @@ -32,7 +32,7 @@ struct ModelFusionQDQ } template _LayerType* - getLayer(std::vector >& newprog, int op_idx) const + getLayer(std::vector >& 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 createFusedLayer(const LayerParams& src) const + Ptr createFusedLayer(const LayerParams& src) const { LayerParams params = src; - Ptr layer = LayerFactory::createLayerInstance(params.type, params); + Ptr 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& inputs, const std::vector& producer_of, - std::vector >& newprog, + std::vector >& newprog, Arg& q_data_in, Arg& out_scale, Arg& out_zp, @@ -118,10 +118,10 @@ struct ModelFusionQDQ { vector removed_args; bool modified = false; - const std::vector >& prog = graph->prog(); + const std::vector >& prog = graph->prog(); size_t i, nargs = netimpl->args.size(), nops = prog.size(); std::vector producer_of(nargs, -1); - std::vector > newprog; + std::vector > newprog; std::vector fused_inputs; std::set skip_indices; std::vector override_outputs; @@ -129,7 +129,7 @@ struct ModelFusionQDQ for (i = 0; i < nops; i++) { if (skip_indices.count((int)i)) continue; - const Ptr& layer = prog[i]; + const Ptr& layer = prog[i]; Layer* layer_ptr = (Layer*)layer.get(); int fused_layer_idx = -1; std::vector >* 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(); + newprog[dq_prog_idx] = Ptr(); 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(); + newprog[dq_idx] = Ptr(); 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(); + newprog[dq_idx] = Ptr(); 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(); + newprog[relu_layer_idx2] = Ptr(); 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(); + newprog[dq_prog_idx] = Ptr(); } } 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(); + newprog[dq_bias_idx] = Ptr(); } if (usecounts.at(conv_x.idx) == 1) { removed_args.push_back(conv_x); - newprog[dq_x_idx] = Ptr(); + newprog[dq_x_idx] = Ptr(); } - newprog[dq_w_idx] = Ptr(); + newprog[dq_w_idx] = Ptr(); 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(); - newprog[dq_w_idx] = Ptr(); + newprog[dq_x_idx] = Ptr(); + newprog[dq_w_idx] = Ptr(); 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(); - newprog[dq_x_idx] = Ptr(); - newprog[dq_w_idx] = Ptr(); + newprog[mm2_idx] = Ptr(); + newprog[dq_x_idx] = Ptr(); + newprog[dq_w_idx] = Ptr(); 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(); - newprog[dq_b_idx] = Ptr(); + newprog[dq_a_idx] = Ptr(); + newprog[dq_b_idx] = Ptr(); 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(); + newprog[dq_idx] = Ptr(); 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(); + newprog[dq_idx] = Ptr(); 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(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 activ_layer = layer.dynamicCast(); - 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& 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(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(); + newprog[ql_idx] = Ptr(); modified = true; } @@ -1338,7 +1339,7 @@ struct ModelFusionQDQ if (inp.idx >= 0 && inp.idx < (int)nargs) uc[inp.idx]--; } - layer = Ptr(); + layer = Ptr(); changed = true; } } diff --git a/modules/dnn/src/graph_fusion_reshape_transpose.cpp b/modules/dnn/src/graph_fusion_reshape_transpose.cpp index a0eb6cf62e..15446d222c 100644 --- a/modules/dnn/src/graph_fusion_reshape_transpose.cpp +++ b/modules/dnn/src/graph_fusion_reshape_transpose.cpp @@ -46,7 +46,7 @@ struct ModelFusionReshapeTranspose bool fuseGraph(Ptr& graph) { - const vector>& prog = graph->prog(); + const vector>& prog = graph->prog(); size_t nops = prog.size(); bool modified = false; @@ -72,7 +72,7 @@ struct ModelFusionReshapeTranspose vector dropped(nops, false); for (size_t i = 0; i < nops; i++) { - const Ptr& layer = prog[i]; + const Ptr& 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& pl = prog[prod_idx]; + const Ptr& pl = prog[prod_idx]; TransposeLayer* prevTr = dynamic_cast(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& pl = prog[prod_idx]; + const Ptr& pl = prog[prod_idx]; Reshape2Layer* prevRs = dynamic_cast(pl.get()); Arg prevOut = layer->inputs[0]; bool single_consumer = usecounts[prevOut.idx] == 1 @@ -154,7 +154,7 @@ struct ModelFusionReshapeTranspose } if (modified) { - vector> newprog; + vector> 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>& prog, + void redirectConsumers(const vector>& prog, const vector& dropped, size_t start_idx, Arg from, Arg to) { diff --git a/modules/dnn/src/graph_fusion_scale_softmax.cpp b/modules/dnn/src/graph_fusion_scale_softmax.cpp index 81c841064a..6156f0573a 100644 --- a/modules/dnn/src/graph_fusion_scale_softmax.cpp +++ b/modules/dnn/src/graph_fusion_scale_softmax.cpp @@ -40,7 +40,7 @@ struct ModelFusionScaleSoftmax bool fuseGraph(Ptr& graph) { - const vector>& prog = graph->prog(); + const vector>& prog = graph->prog(); size_t nops = prog.size(); bool modified = false; @@ -70,7 +70,7 @@ struct ModelFusionScaleSoftmax vector dropped(nops, false); for (size_t i = 0; i < nops; i++) { - const Ptr& layer = prog[i]; + const Ptr& layer = prog[i]; if (!layer || dropped[i]) continue; SoftmaxLayer* sm = dynamic_cast(layer.get()); @@ -83,7 +83,7 @@ struct ModelFusionScaleSoftmax int prod_idx = it->second; if (prod_idx < 0 || dropped[prod_idx]) continue; - const Ptr& pl = prog[prod_idx]; + const Ptr& pl = prog[prod_idx]; NaryEltwiseLayer* elt = dynamic_cast(pl.get()); if (!elt) continue; const auto op = elt->op; @@ -123,7 +123,7 @@ struct ModelFusionScaleSoftmax } if (modified) { - vector> newprog; + vector> newprog; newprog.reserve(nops); for (size_t i = 0; i < nops; i++) { if (!dropped[i] && prog[i]) diff --git a/modules/dnn/src/graph_fusion_shared_gemm.cpp b/modules/dnn/src/graph_fusion_shared_gemm.cpp index 0dba14ef19..62824a8dcd 100644 --- a/modules/dnn/src/graph_fusion_shared_gemm.cpp +++ b/modules/dnn/src/graph_fusion_shared_gemm.cpp @@ -35,7 +35,7 @@ using std::string; namespace { -static bool readGemmWeight(const Ptr& l, bool trans_b, Mat& W_out) +static bool readGemmWeight(const Ptr& 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& l, bool trans_b, Mat& W_out) return true; } -static bool readGemmBias(const Ptr& l, Mat& b_out) +static bool readGemmBias(const Ptr& 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>& prog, int idx, GemmInfo& info) const + bool inspectGemm(const vector>& prog, int idx, GemmInfo& info) const { if (idx < 0 || idx >= (int)prog.size() || !prog[idx]) return false; - const Ptr& l = prog[idx]; + const Ptr& l = prog[idx]; GemmLayer* g = dynamic_cast(l.get()); if (!g) return false; @@ -125,7 +125,7 @@ struct ModelFusionSharedGemm bool fuseGraph(Ptr& graph) { - const vector>& prog = graph->prog(); + const vector>& 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 removed_ops; - vector>>> insertions; // (insert_pos, fused-and-slice layers) + vector>>> 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 fused = LayerFactory::createLayerInstance("Gemm", fp); + Ptr 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> slices; + vector> 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 slice = LayerFactory::createLayerInstance("Slice2", sp); + Ptr 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> bundle; + vector> 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> newprog; + vector> newprog; size_t ins_idx = 0; for (size_t i = 0; i < nops; i++) { while (ins_idx < insertions.size() && diff --git a/modules/dnn/src/graph_fusion_transpose_matmul.cpp b/modules/dnn/src/graph_fusion_transpose_matmul.cpp index 0b13cc99bf..1c6ded90f2 100644 --- a/modules/dnn/src/graph_fusion_transpose_matmul.cpp +++ b/modules/dnn/src/graph_fusion_transpose_matmul.cpp @@ -38,7 +38,7 @@ struct ModelFusionTransposeMatMul bool fuseGraph(Ptr& graph) { - const vector>& prog = graph->prog(); + const vector>& prog = graph->prog(); size_t nops = prog.size(); bool modified = false; @@ -68,7 +68,7 @@ struct ModelFusionTransposeMatMul vector dropped(nops, false); for (size_t i = 0; i < nops; i++) { - const Ptr& layer = prog[i]; + const Ptr& layer = prog[i]; if (!layer || dropped[i]) continue; MatMulLayer* mm = dynamic_cast(layer.get()); @@ -83,7 +83,7 @@ struct ModelFusionTransposeMatMul int prod_idx = it->second; if (prod_idx < 0 || dropped[prod_idx]) continue; - const Ptr& pl = prog[prod_idx]; + const Ptr& pl = prog[prod_idx]; TransposeLayer* tr = dynamic_cast(pl.get()); if (!tr || pl->outputs.size() != 1) continue; if (!isLastTwoSwap(tr->perm)) continue; @@ -103,7 +103,7 @@ struct ModelFusionTransposeMatMul } if (modified) { - vector> newprog; + vector> newprog; newprog.reserve(nops); for (size_t i = 0; i < nops; i++) { if (!dropped[i] && prog[i]) diff --git a/modules/dnn/src/init.cpp b/modules/dnn/src/init.cpp index c146dea1b8..0bbf6dde8e 100644 --- a/modules/dnn/src/init.cpp +++ b/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); diff --git a/modules/dnn/src/layer.cpp b/modules/dnn/src/layer.cpp index f783762898..461344efbd 100644 --- a/modules/dnn/src/layer.cpp +++ b/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 Layer::initCann(const std::vector > &inputs bool Layer::setActivation(const Ptr&) { return false; } bool Layer::tryFuse(Ptr&) { return false; } -void Layer::getScaleShift(Mat& scale, Mat& shift) const + +void Layer::forwardCUDA(const std::vector >&, + const std::vector >&, + 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& inputs, std::vector& outputs, std:: Layer::~Layer() {} -bool Layer::getMemoryShapes(const std::vector& inputs, +bool OpData::getMemoryShapes(const std::vector& inputs, const int requiredOutputs, std::vector& outputs, std::vector& internals) const @@ -257,7 +276,7 @@ bool Layer::getMemoryShapes(const std::vector& inputs, return false; } -void Layer::getTypes(const std::vector&inputs, +void OpData::getTypes(const std::vector&inputs, const int requiredOutputs, const int requiredInternals, std::vector&outputs, @@ -265,20 +284,13 @@ void Layer::getTypes(const std::vector&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& actualInputs, +int OpData::getLayouts(const std::vector& actualInputs, std::vector& desiredInputs, const int requiredOutputs, std::vector& outputs) const @@ -288,43 +300,43 @@ int Layer::getLayouts(const std::vector& actualInputs, return 0; } -int64 Layer::getFLOPS(const std::vector&, +int64 OpData::getFLOPS(const std::vector&, const std::vector&) const { return 0; } -bool Layer::updateMemoryShapes(const std::vector& inputs) +bool OpData::updateMemoryShapes(const std::vector& inputs) { return true; } -std::vector >* Layer::subgraphs() const +std::vector >* 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(); diff --git a/modules/dnn/src/layer_factory.cpp b/modules/dnn/src/layer_factory.cpp index e5b835143e..01e46ed4f7 100644 --- a/modules/dnn/src/layer_factory.cpp +++ b/modules/dnn/src/layer_factory.cpp @@ -104,6 +104,72 @@ Ptr LayerFactory::createLayerInstance(const String& type, LayerParams& pa } } +typedef std::map OpFactory_Impl; +typedef std::map > 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 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(); // 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 LayerFactory::createExec(const String& type, int backendId, + const Ptr& 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(); +} + CV__DNN_INLINE_NS_END }} // namespace cv::dnn diff --git a/modules/dnn/src/layers/batch_norm2_layer.cpp b/modules/dnn/src/layers/batch_norm2_layer.cpp index 7d686567cb..bfd194fa8a 100644 --- a/modules/dnn/src/layers/batch_norm2_layer.cpp +++ b/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 initCUDA(void* context_, + const std::vector >&, + const std::vector >&) CV_OVERRIDE + { + auto context = reinterpret_cast(context_); + Mat scale, bias; + getScaleBias(scale, bias); // per-channel FP32 scale and shift + return make_cuda_node( + preferableTarget, std::move(context->stream), scale, bias); + } +#endif + MatShape getOutShape(const MatShape& inpShape) const { return inpShape; diff --git a/modules/dnn/src/layers/conv2_layer.cpp b/modules/dnn/src/layers/conv2_layer.cpp index 7781f788f1..123c407741 100644 --- a/modules/dnn/src/layers/conv2_layer.cpp +++ b/modules/dnn/src/layers/conv2_layer.cpp @@ -11,6 +11,13 @@ #include #include +#include "../op_cuda.hpp" +#include // 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 initCudaConvNode(void* context_, const MatShape& inpShape, + const MatShape& outShape, int targetId) + { + csl::CSLContext context = *reinterpret_cast(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* wp = filters.ptr(); + 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( + targetId, std::move(context.stream), std::move(context.cudnn_handle), + config, filters, biasMat); + } +#endif + std::vector emptyKernelShape; Ptr 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::create(const LayerParams& params) return Ptr(new Conv2LayerImpl(params)); } +#ifdef HAVE_CUDA +class CUDAConv2Layer : public Layer +{ +public: + CUDAConv2Layer(const Ptr& conv_, void* ctx_) : conv(conv_), ctx(ctx_) {} + + static Ptr create(const Ptr& data, void* backendCtx) + { + Ptr conv = data.dynamicCast(); + if (!conv || !backendCtx || !conv->cudaSupported()) + return Ptr(); + Ptr 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 >& inputs, + const std::vector >& outputs, + void* workspace) CV_OVERRIDE + { + CV_Assert(!inputs.empty() && !outputs.empty()); + auto& ws = *reinterpret_cast(workspace); + if (!node) { + auto inW = inputs[0].dynamicCast(); + auto outW = outputs[0].dynamicCast(); + node = conv->initCudaConvNode(ctx, inW->getShape(), outW->getShape(), preferableTarget); + cudaNode = node.dynamicCast(); + CV_Assert(cudaNode); + ws.require(cudaNode->get_workspace_memory_in_bytes()); + } + cudaNode->forward(inputs, outputs, ws); + } + + Ptr conv; + void* ctx; + Ptr node; + Ptr cudaNode; +}; + +void registerConv2CudaBackend() +{ + CV_DNN_REGISTER_EXEC_CLASS(Conv2, DNN_BACKEND_CUDA, CUDAConv2Layer); +} +#endif + }} diff --git a/modules/dnn/src/layers/maxpool_layer.cpp b/modules/dnn/src/layers/maxpool_layer.cpp index bb156d1aba..46016187de 100644 --- a/modules/dnn/src/layers/maxpool_layer.cpp +++ b/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 initCUDA(void* context_, + const std::vector >& inputs, + const std::vector >&) CV_OVERRIDE + { + auto context = reinterpret_cast(context_); + auto inW = inputs[0].dynamicCast(); + 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(preferableTarget, std::move(context->cudnn_handle), config); + } +#endif + virtual int64_t getFLOPS(const std::vector &inputs, const std::vector &outputs) const CV_OVERRIDE { diff --git a/modules/dnn/src/net.cpp b/modules/dnn/src/net.cpp index 4aedd100ff..f0d69feffd 100644 --- a/modules/dnn/src/net.cpp +++ b/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& inputBlobNames) diff --git a/modules/dnn/src/net_impl.cpp b/modules/dnn/src/net_impl.cpp index a88fb66361..fc6f9404f2 100644 --- a/modules/dnn/src/net_impl.cpp +++ b/modules/dnn/src/net_impl.cpp @@ -274,11 +274,11 @@ Ptr Net::Impl::getLayer(int layerId) const CV_Assert(0 <= layerId && layerId < totalLayers); int graph_ofs = 0; for (const Ptr& graph : allgraphs) { - const std::vector >& prog = graph->prog(); + const std::vector >& 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(); 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>& prog = mainGraph->prog(); + const std::vector>& prog = mainGraph->prog(); for (const auto& layer : prog) { bool produces = false; for (const Arg& out : layer->outputs) @@ -2385,8 +2385,8 @@ std::vector Net::Impl::getLayerNames() const if (mainGraph) { res.reserve(totalLayers); for (const Ptr& graph: allgraphs) { - const std::vector >& prog = graph->prog(); - for (const Ptr& layer: prog) + const std::vector >& prog = graph->prog(); + for (const Ptr& layer: prog) res.push_back(layer->name); } } else { @@ -2416,7 +2416,7 @@ std::vector Net::Impl::getUnconnectedOutLayers() const int graph_ofs = 0; for (const auto& graph : allgraphs) { - const std::vector>& prog = graph->prog(); + const std::vector>& 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, return 0; int64 flops = 0; - const std::vector>& prog = graph->prog(); + const std::vector>& prog = graph->prog(); - for (const Ptr& layer : prog) { + for (const Ptr& layer : prog) { if (!layer) continue; @@ -2595,7 +2595,7 @@ int64 Net::Impl::getFLOPS( for (const Ptr& graph : allgraphs) { int progSize = (int)graph->prog().size(); if (localIdx < progSize) { - const Ptr& layer = graph->prog()[localIdx]; + const Ptr& layer = graph->prog()[localIdx]; if (!layer) return 0; @@ -2694,8 +2694,8 @@ void Net::Impl::collectLayerInfo(std::vector& names, std::vector names.reserve(totalLayers); types.reserve(totalLayers); for (const Ptr& graph : allgraphs) { - const std::vector>& prog = graph->prog(); - for (const Ptr& layer : prog) { + const std::vector>& prog = graph->prog(); + for (const Ptr& 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& layersTypes) const if (mainGraph) { std::set layersTypesSet; for (const Ptr& g: allgraphs) { - const std::vector >& prog = g->prog(); - for (const Ptr& layer: prog) { + const std::vector >& prog = g->prog(); + for (const Ptr& 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& g: allgraphs) { - const std::vector >& prog = g->prog(); - for (const Ptr& layer: prog) { + const std::vector >& prog = g->prog(); + for (const Ptr& layer: prog) { if (!layer) continue; if (layer->type == layerType) diff --git a/modules/dnn/src/net_impl.hpp b/modules/dnn/src/net_impl.hpp index eb4df4c5b8..57655ed153 100644 --- a/modules/dnn/src/net_impl.hpp +++ b/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 > argWrappers; + std::vector argWrapperData; TracingMode tracingMode; ProfilingMode profilingMode; std::vector 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, bool useCUDA); +#ifdef HAVE_CUDA + Ptr 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, + void allocateLayerOutputs(const Ptr& layer, const std::vector& inpTypes, const std::vector& inpShapes, std::vector& 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(layer->netimpl); + return reinterpret_cast(op->netimpl); } Net readNetFromONNX2(const String&); diff --git a/modules/dnn/src/net_impl2.cpp b/modules/dnn/src/net_impl2.cpp index e32f82e9b1..2c751058c7 100644 --- a/modules/dnn/src/net_impl2.cpp +++ b/modules/dnn/src/net_impl2.cpp @@ -316,30 +316,30 @@ public: return g; }*/ - virtual const std::vector& append(Ptr& layer, + virtual const std::vector& append(Ptr& op, const std::vector& 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, + virtual Arg append(Ptr& op, const std::string& outname) override { std::vector outnames = {outname}; - const std::vector& outputs = append(layer, outnames); + const std::vector& 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 = prog_[i]; - layer->dump(strm, argindent, i+1 < nlayers); + const Ptr& 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 >& prog() const override { return prog_; } - virtual void setProg(const std::vector >& newprog) override { prog_ = newprog; } + virtual const std::vector >& 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 >& newprog) override + { + prog_ = newprog; + exec_.clear(); + execBackend_.clear(); + inH2D_.clear(); + outD2H_.clear(); + } -protected: Net::Impl* netimpl_; std::string name_; std::vector inputs_; std::vector outputs_; - std::vector > prog_; + std::vector > prog_; + std::vector > exec_; + std::vector execBackend_; + std::vector > inH2D_; + std::vector > outD2H_; }; Ptr 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, bool useCUDA) +{ + GraphImpl* g = static_cast(graph.get()); + const std::vector >& prog = g->prog_; + size_t i, nops = prog.size(); + g->exec_.assign(nops, Ptr()); + g->execBackend_.assign(nops, DNN_BACKEND_OPENCV); + + for (i = 0; i < nops; i++) { + const Ptr& op = prog[i]; + if (!op) + continue; + + // recurse into subgraphs (If/Loop bodies) first + const std::vector >* subs = op->subgraphs(); + if (subs) { + for (const Ptr& sub : *subs) + finalizeGraph(sub, useCUDA); + } + + Ptr 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(); + 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(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& g : allgraphs) + finalizeGraph(g, useCUDA); + useBlockLayout(); + assignBuffers(); + totalLayers = updateGraphOfs(mainGraph, 0, true); + + for (const Ptr& g : allgraphs) + finalizeGraph(g, useCUDA); + + finalized = true; +} + void Net::Impl::allocateLayerOutputs( - const Ptr& layer, + const Ptr& layer, const std::vector& inpTypes, const std::vector& inpShapes, std::vector& 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& perIter, int axis, bool reverse } return stacked; } +#ifdef HAVE_CUDA +Ptr Net::Impl::getCudaArgWrapper(Arg arg, Mat& hostMat) +{ + int idx = arg.idx; + if ((int)argWrappers.size() != (int)args.size()) { + argWrappers.assign(args.size(), Ptr()); + argWrapperData.assign(args.size(), nullptr); + } + Ptr cw = argWrappers[idx].dynamicCast(); + if (!cw || argWrapperData[idx] != (const void*)hostMat.data) { + Ptr w = wrapMat(DNN_BACKEND_CUDA, preferableTarget, hostMat); + cw = w.dynamicCast(); + 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& inputs, const std::vector& outputs, + std::vector& inpMats, std::vector& outMats) +{ + Ptr exec = gimpl->exec_[opidx]; + CV_Assert(exec && netimpl->cudaInfo); + std::vector > 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, InputArrayOfArrays inputs_, OutputArrayOfArrays outputs_, bool isMainGraph) @@ -1184,7 +1323,8 @@ void Net::Impl::forwardGraph(Ptr& 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 >& prog = graph->prog(); + GraphImpl* gimpl = static_cast(graph.get()); + const std::vector >& prog = graph->prog(); size_t i, nops = prog.size(); const std::vector& gr_inputs = graph->inputs(); const std::vector& gr_outputs = graph->outputs(); @@ -1206,15 +1346,29 @@ void Net::Impl::forwardGraph(Ptr& 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 cw = argWrappers[ginp.idx].dynamicCast(); + if (cw) cw->setHostDirty(); + } +#endif } } for (size_t opidx = 0; opidx < nops; opidx++) { - const Ptr& 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& 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& inputs = layer->inputs; - const std::vector& outputs = layer->outputs; + Ptr layer = (opidx < gimpl->exec_.size()) ? gimpl->exec_[opidx] : Ptr(); + if (!layer) + layer = op.dynamicCast(); + 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& inputs = op->inputs; + const std::vector& outputs = op->outputs; size_t ninputs = inputs.size(), noutputs = outputs.size(); inpMats.resize(ninputs); @@ -1233,15 +1387,15 @@ void Net::Impl::forwardGraph(Ptr& 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, InputArrayOfArrays inputs_, timestamp = getTickCount(); - std::vector >* subgraphs = layer->subgraphs(); + std::vector >* 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 cw = argWrappers[inputs[k].idx].dynamicCast(); + 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 cw = argWrappers[outputs[k].idx].dynamicCast(); + if (cw) cw->setHostDirty(); + } + } +#endif + } } else { Ptr iflayer = layer.dynamicCast(); @@ -1425,7 +1606,7 @@ void Net::Impl::forwardGraph(Ptr& 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, 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 cw = argWrappers[out.idx].dynamicCast(); + 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, std::vector& useco CV_Assert(output.idx < (int)usecounts.size()); usecounts[output.idx]++; } - const std::vector >& prog = graph->prog(); - for (const Ptr& layer: prog) { + const std::vector >& prog = graph->prog(); + for (const Ptr& layer: prog) { const std::vector& 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, int currofs, bool ismain) allgraphs.clear(); layerNameToId.clear(); } - const std::vector >& prog = graph->prog(); + const std::vector >& 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 = prog[i]; + const Ptr& layer = prog[i]; layerNameToId.insert(std::make_pair(layer->name, currofs + (int)i)); const std::vector >* subgraphs = layer->subgraphs(); if (subgraphs) { @@ -1718,12 +1906,12 @@ bool Net::Impl::tryInferGraphShapes(const Ptr& graph, if (!graph) return true; - const std::vector >& prog = graph->prog(); + const std::vector >& prog = graph->prog(); std::vector inpShapes, outShapes, tempShapes; std::vector inpTypes, outTypes, tempTypes; - for (const Ptr& layer: prog) { + for (const Ptr& layer: prog) { if (!layer) continue; diff --git a/modules/dnn/src/net_impl_backend.cpp b/modules/dnn/src/net_impl_backend.cpp index 0f176d03d3..beba37ff74 100644 --- a/modules/dnn/src/net_impl_backend.cpp +++ b/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 && diff --git a/modules/dnn/src/onnx/onnx_importer2.cpp b/modules/dnn/src/onnx/onnx_importer2.cpp index 961d207538..df5c671979 100644 --- a/modules/dnn/src/onnx/onnx_importer2.cpp +++ b/modules/dnn/src/onnx/onnx_importer2.cpp @@ -167,7 +167,7 @@ protected: std::string onnxBasePath; Ptr curr_graph; opencv_onnx::GraphProto* curr_graph_proto; - std::vector > curr_prog; + std::vector > curr_prog; std::vector node_inputs, node_outputs; std::string framework_name; @@ -892,7 +892,7 @@ Ptr ONNXImporter2::parseGraph(opencv_onnx::GraphProto* graph_proto, bool opencv_onnx::GraphProto* saved_graph_proto = curr_graph_proto; Ptr saved_graph = curr_graph; - std::vector > saved_prog; + std::vector > saved_prog; curr_graph_proto = graph_proto; std::vector 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& loopLayer = curr_prog.back(); + Ptr& 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& ifLayer = curr_prog.back(); + Ptr& ifLayer = curr_prog.back(); *ifLayer->subgraphs() = thenelse; } diff --git a/modules/dnn/src/op_cuda.cpp b/modules/dnn/src/op_cuda.cpp index 372c36f11a..0f2b6263ec 100644 --- a/modules/dnn/src/op_cuda.cpp +++ b/modules/dnn/src/op_cuda.cpp @@ -8,10 +8,70 @@ #include "op_cuda.hpp" #include "cuda4dnn/init.hpp" #include "net_impl.hpp" +#include namespace cv { namespace dnn { CV__DNN_INLINE_NS_BEGIN +class CUDALegacyExec : public Layer +{ +public: + CUDALegacyExec(const Ptr& impl_, void* ctx_) : impl(impl_), ctx(ctx_) {} + + static Ptr create(const Ptr& data, void* backendCtx) + { + Ptr impl = data.dynamicCast(); + if (!impl || !backendCtx || !impl->supportBackend(DNN_BACKEND_CUDA)) + return Ptr(); // unsupported -> CPU fallback + Ptr 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 >& inputs, + const std::vector >& outputs, + void* workspace) CV_OVERRIDE + { + cuda4dnn::csl::Workspace& ws = *reinterpret_cast(workspace); + if (!node) { + impl->preferableTarget = preferableTarget; // initCUDA may pick FP16/FP32 by target + cuda4dnn::csl::CSLContext context = *reinterpret_cast(ctx); + node = impl->initCUDA(&context, inputs, outputs); + CV_Assert(node); + cudaNode = node.dynamicCast(); + CV_Assert(cudaNode); + ws.require(cudaNode->get_workspace_memory_in_bytes()); + } + cudaNode->forward(inputs, outputs, ws); + } + + Ptr impl; + void* ctx; + Ptr node; + Ptr 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& blobsToKeep_) { diff --git a/modules/dnn/src/tensorflow/tf_importer.cpp b/modules/dnn/src/tensorflow/tf_importer.cpp index 441d73004b..720bc855c7 100644 --- a/modules/dnn/src/tensorflow/tf_importer.cpp +++ b/modules/dnn/src/tensorflow/tf_importer.cpp @@ -576,7 +576,7 @@ protected: std::map layer_id; bool newEngine; - std::vector> curProg; + std::vector> curProg; std::vector> layersOutputs; std::vector modelInputs; std::unordered_map tensorsShape; diff --git a/modules/dnn/src/tflite/tflite_importer.cpp b/modules/dnn/src/tflite/tflite_importer.cpp index 9239edac41..8195c5f87c 100644 --- a/modules/dnn/src/tflite/tflite_importer.cpp +++ b/modules/dnn/src/tflite/tflite_importer.cpp @@ -34,7 +34,7 @@ private: const flatbuffers::Vector >* modelTensors; std::map allTensors; Net& dstNet; - std::vector> curProg; + std::vector> curProg; // This is a vector of pairs (layerId, outputId) where we iterate over // indices from TFLite notation and get created OpenCV layers. diff --git a/modules/dnn/test/test_model.cpp b/modules/dnn/test/test_model.cpp index 81be489d98..47935ae7bd 100644 --- a/modules/dnn/test/test_model.cpp +++ b/modules/dnn/test/test_model.cpp @@ -690,17 +690,31 @@ static void topK(const Mat& probs, std::vector >& result, } } -typedef testing::TestWithParam Reproducibility_ResNet50_ONNX; +// Returns the CPU (OpenCV backend) and CUDA backend/target pairs for benchmarking. +static std::vector > resnet50BackendsAndTargets() +{ + std::vector > 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 > 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 Reproducibility_ResNet50_QDQ_ONNX; TEST_P(Reproducibility_ResNet50_QDQ_ONNX, Accuracy)