diff --git a/modules/dnn/include/opencv2/dnn/dnn.hpp b/modules/dnn/include/opencv2/dnn/dnn.hpp index 00f086e516..97aed97a20 100644 --- a/modules/dnn/include/opencv2/dnn/dnn.hpp +++ b/modules/dnn/include/opencv2/dnn/dnn.hpp @@ -262,23 +262,91 @@ 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. + * %LayerInfo 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 %LayerInfo nodes (see Graph::prog()); + * executable, backend-specific instances (Layer subclasses) are constructed from an + * %LayerInfo during Net::finalizeNet(). + * + * Each operation type registers a `static Ptr create(const LayerParams&)` factory + * via @ref CV_DNN_REGISTER_OP_CLASS. */ - class CV_EXPORTS_W Layer : public Algorithm + class CV_EXPORTS_W LayerInfo { public: + LayerInfo(); + explicit LayerInfo(const LayerParams& params); + virtual ~LayerInfo(); + + 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 a LayerInfo node: it + * implements forward() (and finalize()) for a particular backend/target. In the new graph + * engine a %Layer is created from a LayerInfo by the executor factory; shape/type inference + * stays on the LayerInfo 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. + */ + class CV_EXPORTS_W Layer : public LayerInfo + { + public: + /** @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 +409,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. @@ -379,6 +435,20 @@ CV__DNN_INLINE_NS_BEGIN const std::vector>& outputs ); + /** + * @brief Returns a CUDA backend node for the new graph engine (wrapper-free). + * + * Inputs and outputs are device tensors (arrays of cuda::GpuMatND) carrying shape and type; + * only that metadata is needed to build the node, the buffers are filled later by + * forwardCUDA(). The default adapts the wrapper-based initCUDA() so classic-engine ops keep + * working. @p context is a void pointer to a CSLContext object. + */ + virtual Ptr initCUDA( + void *context, + InputArrayOfArrays inputs, + InputArrayOfArrays outputs + ); + /** * @brief Returns a TimVX backend node * @@ -419,102 +489,26 @@ 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. Inputs and outputs are + * device-resident tensors passed as arrays of cuda::GpuMatND; no backend wrappers are + * involved. 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(InputArrayOfArrays inputs, + OutputArrayOfArrays 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 +527,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; + 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..e1807c40fb 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 LayerInfo (metadata node) class for the new graph engine. + * @param type string, containing the operation type name. + * @param class C++ class derived from LayerInfo, 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..b66bb62e3c 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 LayerInfo; 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/cuda4dnn/primitives/activation.hpp b/modules/dnn/src/cuda4dnn/primitives/activation.hpp index c10f9014a5..ecd306e584 100644 --- a/modules/dnn/src/cuda4dnn/primitives/activation.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/activation.hpp @@ -25,17 +25,14 @@ namespace cv { namespace dnn { namespace cuda4dnn { using wrapper_type = GetCUDABackendWrapperType; void forward( - const std::vector>& inputs, - const std::vector>& outputs, + const std::vector& inputs, + const std::vector& outputs, csl::Workspace& workspace) override { for (int i = 0; i < inputs.size(); i++) { - auto input_wrapper = inputs[i].dynamicCast(); - auto input = input_wrapper->getView(); - - auto output_wrapper = outputs[i].dynamicCast(); - auto output = output_wrapper->getSpan(); + auto input = csl::viewOf(inputs[i]); + auto output = csl::spanOf(outputs[i]); static_cast*>(this)->calculate(output, input); } diff --git a/modules/dnn/src/cuda4dnn/primitives/batch_norm.hpp b/modules/dnn/src/cuda4dnn/primitives/batch_norm.hpp index 293811f73d..6180ff0453 100644 --- a/modules/dnn/src/cuda4dnn/primitives/batch_norm.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/batch_norm.hpp @@ -32,17 +32,14 @@ namespace cv { namespace dnn { namespace cuda4dnn { } void forward( - const std::vector>& inputs, - const std::vector>& outputs, + const std::vector& inputs, + const std::vector& outputs, csl::Workspace& workspace) override { CV_Assert(inputs.size() == 1 && outputs.size() == 1); - auto input_wrapper = inputs[0].dynamicCast(); - auto input = input_wrapper->getView(); - - auto output_wrapper = outputs[0].dynamicCast(); - auto output = output_wrapper->getSpan(); + auto input = csl::viewOf(inputs[0]); + auto output = csl::spanOf(outputs[0]); std::size_t inner_size = input.size_range(2, input.rank()); kernels::scaleN_with_biasN(stream, output, input, inner_size, weightsTensor, biasTensor); diff --git a/modules/dnn/src/cuda4dnn/primitives/convolution.hpp b/modules/dnn/src/cuda4dnn/primitives/convolution.hpp index 12cf97404e..a4534a19f5 100644 --- a/modules/dnn/src/cuda4dnn/primitives/convolution.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/convolution.hpp @@ -286,8 +286,8 @@ namespace cv { namespace dnn { namespace cuda4dnn { } void forward( - const std::vector>& inputs, - const std::vector>& outputs, + const std::vector& inputs, + const std::vector& outputs, csl::Workspace& workspace) override { /* input[0] = conv input, input[1] = bias (from fused eltwise layer) */ @@ -296,8 +296,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { csl::WorkspaceAllocator allocator(workspace); - auto input_wrapper = inputs[0].dynamicCast(); - auto input = input_wrapper->getView(); + auto input = csl::viewOf(inputs[0]); if (!transformed_shape.empty()) { @@ -309,8 +308,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { auto conv_scratchpad = allocator.get_instance(); - auto output_wrapper = outputs[0].dynamicCast(); - auto output = output_wrapper->getSpan(); + auto output = csl::spanOf(outputs[0]); if (fusion_location == InternalFusionLocation::CUDNN) { @@ -320,8 +318,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { convoluter.convolve_with_bias_activation(output, input, filtersTensor, biasTensor, conv_scratchpad); else if (fusion_mode == ConvolutionConfiguration::FusionMode::ELTWISE_SUM_THEN_ACTIVATION) { - auto eltwise_wrapper = inputs[1].dynamicCast(); - auto eltwise = eltwise_wrapper->getView(); + auto eltwise = csl::viewOf(inputs[1]); CV_Assert(is_shape_same(eltwise, output)); convoluter.convolve_with_bias_eltwise_activation(output, input, filtersTensor, biasTensor, eltwise, conv_scratchpad); @@ -357,8 +354,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { fusion_mode == ConvolutionConfiguration::FusionMode::ELTWISE_SUM_THEN_ACTIVATION || fusion_mode == ConvolutionConfiguration::FusionMode::ACTIVATION_THEN_ELTWISE_SUM); - auto eltwise_wrapper = inputs[1].dynamicCast(); - auto eltwise = eltwise_wrapper->getView(); + auto eltwise = csl::viewOf(inputs[1]); CV_Assert(is_shape_same(eltwise, output)); std::size_t inner_size = output.size_range(2, output.rank()); @@ -472,8 +468,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { fusion_mode == ConvolutionConfiguration::FusionMode::ELTWISE_SUM_THEN_ACTIVATION || fusion_mode == ConvolutionConfiguration::FusionMode::ACTIVATION_THEN_ELTWISE_SUM); - auto eltwise_wrapper = inputs[1].dynamicCast(); - auto eltwise = eltwise_wrapper->getView(); + auto eltwise = csl::viewOf(inputs[1]); CV_Assert(is_shape_same(eltwise, output)); /* we pass `eltwise` as `bias` (with `inner_size` as one) to bias-activation kernels */ diff --git a/modules/dnn/src/cuda4dnn/primitives/eltwise.hpp b/modules/dnn/src/cuda4dnn/primitives/eltwise.hpp index 634ea2d806..b283401eb0 100644 --- a/modules/dnn/src/cuda4dnn/primitives/eltwise.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/eltwise.hpp @@ -59,8 +59,8 @@ namespace cv { namespace dnn { namespace cuda4dnn { } void forward( - const std::vector>& inputs, - const std::vector>& outputs, + const std::vector& inputs, + const std::vector& outputs, csl::Workspace& workspace) override { CV_Assert(outputs.size() == 1); @@ -68,16 +68,12 @@ namespace cv { namespace dnn { namespace cuda4dnn { CV_Assert(coeffs.size() == 0 || op == EltwiseOpType::SUM); CV_Assert(coeffs.size() == 0 || inputs.size() == coeffs.size()); - auto output_wrapper = outputs[0].dynamicCast(); - auto output = output_wrapper->getSpan(); + auto output = csl::spanOf(outputs[0]); if (inputs.size() == 2) { - auto input_wrapper_x = inputs[0].dynamicCast(); - auto input_x = input_wrapper_x->getView(); - - auto input_wrapper_y = inputs[1].dynamicCast(); - auto input_y = input_wrapper_y->getView(); + auto input_x = csl::viewOf(inputs[0]); + auto input_y = csl::viewOf(inputs[1]); switch (op) { @@ -97,20 +93,17 @@ namespace cv { namespace dnn { namespace cuda4dnn { case EltwiseOpType::POW: kernels::eltwise_pow_2(stream, output, input_x, input_y); break; } } else if (inputs.size() == 1) { - auto input_wrapper_0 = inputs[0].dynamicCast(); - auto input_0 = input_wrapper_0->getView(); + auto input_0 = csl::viewOf(inputs[0]); csl::tensor_ops::copy(stream, output, input_0); } else { - auto input_wrapper_0 = inputs[0].dynamicCast(); - auto input_0 = input_wrapper_0->getView(); + auto input_0 = csl::viewOf(inputs[0]); /* we first make a copy and then apply EltwiseOp cumulatively */ csl::tensor_ops::copy(stream, output, input_0); for (int i = 1; i < inputs.size(); i++) { - auto input_wrapper = inputs[i].dynamicCast(); - auto input = input_wrapper->getView(); + auto input = csl::viewOf(inputs[i]); switch (op) { diff --git a/modules/dnn/src/cuda4dnn/primitives/inner_product.hpp b/modules/dnn/src/cuda4dnn/primitives/inner_product.hpp index a0a00ce125..70f1f79c52 100644 --- a/modules/dnn/src/cuda4dnn/primitives/inner_product.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/inner_product.hpp @@ -43,17 +43,14 @@ namespace cv { namespace dnn { namespace cuda4dnn { } void forward( - const std::vector>& inputs, - const std::vector>& outputs, + const std::vector& inputs, + const std::vector& outputs, csl::Workspace& workspace) override { for (int i = 0; i < inputs.size(); i++) { - auto input_wrapper = inputs[i].dynamicCast(); - auto input = input_wrapper->getView(); - - auto output_wrapper = outputs[i].dynamicCast(); - auto output = output_wrapper->getSpan(); + auto input = csl::viewOf(inputs[i]); + auto output = csl::spanOf(outputs[i]); std::size_t batch_size = input.size_range(0, axis); diff --git a/modules/dnn/src/cuda4dnn/primitives/max_unpooling.hpp b/modules/dnn/src/cuda4dnn/primitives/max_unpooling.hpp index 759cf65b16..e606019a42 100644 --- a/modules/dnn/src/cuda4dnn/primitives/max_unpooling.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/max_unpooling.hpp @@ -111,6 +111,22 @@ namespace cv { namespace dnn { namespace cuda4dnn { ); } + void forward( + const std::vector& inputs, + const std::vector& outputs, + csl::Workspace& workspace) override + { + CV_Assert(inputs.size() == 1 && outputs.size() == 2); + + auto input_data = csl::viewOf(inputs[0]); + auto output_data = csl::spanOf(outputs[0]); + auto output_indices = csl::spanOf(outputs[1]); + + kernels::max_pooling_with_indices( + stream, output_data, output_indices, input_data, window_size, strides, padding_left + ); + } + private: csl::Stream stream; diff --git a/modules/dnn/src/cuda4dnn/primitives/pooling.hpp b/modules/dnn/src/cuda4dnn/primitives/pooling.hpp index bd8a73c2f2..7ab328f1a4 100644 --- a/modules/dnn/src/cuda4dnn/primitives/pooling.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/pooling.hpp @@ -224,14 +224,13 @@ namespace cv { namespace dnn { namespace cuda4dnn { } void forward( - const std::vector>& inputs, - const std::vector>& outputs, + const std::vector& inputs, + const std::vector& outputs, csl::Workspace& workspace) override { CV_Assert(inputs.size() == 1 && outputs.size() == 1); - auto input_wrapper = inputs[0].dynamicCast(); - auto input = input_wrapper->getView(); + auto input = csl::viewOf(inputs[0]); if (!transformedInput.empty()) { @@ -239,8 +238,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { input = csl::TensorView(transformedInput); } - auto output_wrapper = outputs[0].dynamicCast(); - auto output = output_wrapper->getSpan(); + auto output = csl::spanOf(outputs[0]); pooler.pool(input, output); } diff --git a/modules/dnn/src/cuda4dnn/primitives/reshape.hpp b/modules/dnn/src/cuda4dnn/primitives/reshape.hpp index 2cf1d676ea..676a30c76b 100644 --- a/modules/dnn/src/cuda4dnn/primitives/reshape.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/reshape.hpp @@ -23,8 +23,8 @@ namespace cv { namespace dnn { namespace cuda4dnn { ReshapeOp(csl::Stream stream_) : stream(std::move(stream_)) { } void forward( - const std::vector>& inputs, - const std::vector>& outputs, + const std::vector& inputs, + const std::vector& outputs, csl::Workspace& workspace) override { /* sometimes the output shape is passed as extra inputs; hence, >= instead of == */ @@ -32,11 +32,8 @@ namespace cv { namespace dnn { namespace cuda4dnn { for (int i = 0; i < outputs.size(); i++) { - auto input_wrapper = inputs[i].dynamicCast(); - auto input = input_wrapper->getView(); - - auto output_wrapper = outputs[i].dynamicCast(); - auto output = output_wrapper->getSpan(); + auto input = csl::viewOf(inputs[i]); + auto output = csl::spanOf(outputs[i]); if (input.get() != output.get()) { diff --git a/modules/dnn/src/graph_block_layout.cpp b/modules/dnn/src/graph_block_layout.cpp index 3a90ba794e..22834c2909 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..8283c0fae2 100644 --- a/modules/dnn/src/graph_buffer_allocator.cpp +++ b/modules/dnn/src/graph_buffer_allocator.cpp @@ -143,13 +143,13 @@ struct BufferAllocator std::unordered_set bodyDefined; for (Arg ba : body->inputs()) bodyDefined.insert(ba.idx); - for (const Ptr& blayer : body->prog()) { + for (const Ptr& blayer : body->prog()) { if (!blayer) continue; for (Arg bo : blayer->outputs) bodyDefined.insert(bo.idx); } std::unordered_set closureBumped; - for (const Ptr& blayer : body->prog()) { + for (const Ptr& blayer : body->prog()) { if (!blayer) continue; for (Arg bi : blayer->inputs) { if (bi.idx <= 0) continue; @@ -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..fc88010762 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]; + LayerInfo* 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..9979398bb2 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 + LayerInfo* 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..7c7c3a60d0 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; }, + [](LayerInfo* 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; }, + [](LayerInfo* 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) { + [](LayerInfo* 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) { + [](LayerInfo* 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; }, + [](LayerInfo* 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..ef95b05d76 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..059a69ff41 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..454eccb5e2 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..eb5d4e9672 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..f0747e697b 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..bfbb52beab 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..12ec3f207c 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..a30d1c5a8b 100644 --- a/modules/dnn/src/layer.cpp +++ b/modules/dnn/src/layer.cpp @@ -10,46 +10,78 @@ namespace dnn { CV__DNN_INLINE_NS_BEGIN -Layer::Layer() { +LayerInfo::LayerInfo() { netimpl = nullptr; - preferableTarget = DNN_TARGET_CPU; } -Layer::Layer(const LayerParams& params) +LayerInfo::LayerInfo(const LayerParams& params) : blobs(params.blobs) , name(params.name) , type(params.type) { netimpl = nullptr; - preferableTarget = DNN_TARGET_CPU; } -void Layer::setParamsFrom(const LayerParams& params) +LayerInfo::~LayerInfo() {} + +void LayerInfo::setParamsFrom(const LayerParams& params) { blobs = params.blobs; name = params.name; type = params.type; } -int Layer::inputNameToIndex(String) +int LayerInfo::inputNameToIndex(String) { return -1; } -int Layer::outputNameToIndex(const String&) +int LayerInfo::outputNameToIndex(const String&) { return 0; } +Layer::Layer() { + netimpl = nullptr; + preferableTarget = DNN_TARGET_CPU; +} + +Layer::Layer(const LayerParams& params) + : LayerInfo(params) +{ + preferableTarget = DNN_TARGET_CPU; +} + bool Layer::supportBackend(int backendId) { return backendId == DNN_BACKEND_OPENCV; } +Ptr Layer::initCUDA( + void* context, + const std::vector>& inputs, + const std::vector>& outputs) +{ +#ifdef HAVE_CUDA + // Adapt the classic wrapper-based entry point to the array-based one, so ops ported to the + // new graph engine only need to override initCUDA(context, inputs, outputs) with GpuMatND. + std::vector inGpu(inputs.size()), outGpu(outputs.size()); + for (size_t i = 0; i < inputs.size(); i++) + inGpu[i] = inputs[i].dynamicCast()->getDeviceMatND(); + for (size_t i = 0; i < outputs.size(); i++) + outGpu[i] = outputs[i].dynamicCast()->getDeviceMatND(); + return initCUDA(context, inGpu, outGpu); +#else + CV_UNUSED(context); CV_UNUSED(inputs); CV_UNUSED(outputs); + CV_Error(Error::StsNotImplemented, "CUDA pipeline of " + type + " layers is not defined."); + return Ptr(); +#endif +} + Ptr Layer::initCUDA( void*, - const std::vector>&, - const std::vector>&) + InputArrayOfArrays, + InputArrayOfArrays) { CV_Error(Error::StsNotImplemented, "CUDA pipeline of " + type + " layers is not defined."); return Ptr(); @@ -94,13 +126,19 @@ 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(InputArrayOfArrays, OutputArrayOfArrays, void*) +{ + CV_Error(Error::StsNotImplemented, "CUDA forward of " + type + " layers is not defined."); +} + +void LayerInfo::getScaleShift(Mat& scale, Mat& shift) const { scale = Mat(); shift = Mat(); } -void Layer::getScaleZeropoint(float& scale, int& zeropoint) const +void LayerInfo::getScaleZeropoint(float& scale, int& zeropoint) const { scale = 1.f; zeropoint = 0; @@ -247,7 +285,7 @@ void Layer::run(const std::vector& inputs, std::vector& outputs, std:: Layer::~Layer() {} -bool Layer::getMemoryShapes(const std::vector& inputs, +bool LayerInfo::getMemoryShapes(const std::vector& inputs, const int requiredOutputs, std::vector& outputs, std::vector& internals) const @@ -257,7 +295,7 @@ bool Layer::getMemoryShapes(const std::vector& inputs, return false; } -void Layer::getTypes(const std::vector&inputs, +void LayerInfo::getTypes(const std::vector&inputs, const int requiredOutputs, const int requiredInternals, std::vector&outputs, @@ -265,20 +303,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 LayerInfo::getLayouts(const std::vector& actualInputs, std::vector& desiredInputs, const int requiredOutputs, std::vector& outputs) const @@ -288,43 +319,43 @@ int Layer::getLayouts(const std::vector& actualInputs, return 0; } -int64 Layer::getFLOPS(const std::vector&, +int64 LayerInfo::getFLOPS(const std::vector&, const std::vector&) const { return 0; } -bool Layer::updateMemoryShapes(const std::vector& inputs) +bool LayerInfo::updateMemoryShapes(const std::vector& inputs) { return true; } -std::vector >* Layer::subgraphs() const +std::vector >* LayerInfo::subgraphs() const { return nullptr; } -bool Layer::alwaysSupportInplace() const +bool LayerInfo::alwaysSupportInplace() const { return false; } -bool Layer::dynamicOutputShapes() const +bool LayerInfo::dynamicOutputShapes() const { return false; } -bool Layer::isDataShuffling() const +bool LayerInfo::isDataShuffling() const { return false; } -std::ostream& Layer::dumpAttrs(std::ostream& strm, int) const +std::ostream& LayerInfo::dumpAttrs(std::ostream& strm, int) const { return strm; } -std::ostream& Layer::dump(std::ostream& strm, int indent, bool comma) const +std::ostream& LayerInfo::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..7b070998f6 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 LayerInfo 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..92595632af 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_, + InputArrayOfArrays, + InputArrayOfArrays) 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_common.cpp b/modules/dnn/src/layers/conv2_common.cpp index 89d5d3b052..876b185e23 100644 --- a/modules/dnn/src/layers/conv2_common.cpp +++ b/modules/dnn/src/layers/conv2_common.cpp @@ -396,7 +396,7 @@ MatShape deconvInferShape(const MatShape& inpShape, const MatShape& wshape, const std::vector& adjustPads, AutoPadding autoPad) { - bool blockLayout = true; + bool blockLayout = (inpShape.layout == DATA_LAYOUT_BLOCK); int ndims = inpShape.dims; int nspatialdims = ndims - 2 - int(blockLayout); CV_Assert(nspatialdims >= 1); @@ -413,9 +413,13 @@ MatShape deconvInferShape(const MatShape& inpShape, const MatShape& wshape, kshape_[i] = wshape[i + 2]; } - int C0 = inpShape[ndims - 1]; int K_out = ngroups * wshape[1]; - outshape[1] = (K_out + C0 - 1) / C0; + if (blockLayout) { + int C0 = inpShape[ndims - 1]; + outshape[1] = (K_out + C0 - 1) / C0; + } else { + outshape[1] = K_out; + } CV_Assert(strides.empty() || (int)strides.size() == nspatialdims); CV_Assert(dilations.empty() || (int)dilations.size() == nspatialdims); @@ -439,7 +443,7 @@ MatShape deconvInferShape(const MatShape& inpShape, const MatShape& wshape, } outshape[i + 2] = outsz; } - outshape.C = K_out; + outshape.C = blockLayout ? K_out : 0; return outshape; } diff --git a/modules/dnn/src/layers/conv2_layer.cpp b/modules/dnn/src/layers/conv2_layer.cpp index 7781f788f1..337d9e5403 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,54 @@ 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->name = conv->name; + layer->type = conv->type; + layer->inputs = conv->inputs; + layer->outputs = conv->outputs; + return layer; + } + + void forwardCUDA(InputArrayOfArrays inputs_, + OutputArrayOfArrays outputs_, + void* workspace) CV_OVERRIDE + { + std::vector inputs, outputs; + inputs_.getGpuMatNDVector(inputs); + outputs_.getGpuMatNDVector(outputs); + CV_Assert(!inputs.empty() && !outputs.empty()); + + auto& ws = *reinterpret_cast(workspace); + if (!node) { + node = conv->initCudaConvNode(ctx, inputs[0].size, outputs[0].size, 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/elementwise_layers.cpp b/modules/dnn/src/layers/elementwise_layers.cpp index 269b3fe96a..174b4d02f2 100644 --- a/modules/dnn/src/layers/elementwise_layers.cpp +++ b/modules/dnn/src/layers/elementwise_layers.cpp @@ -310,8 +310,8 @@ public: #ifdef HAVE_CUDA Ptr initCUDA( void *context_, - const std::vector>& inputs, - const std::vector>& outputs + InputArrayOfArrays /*inputs*/, + InputArrayOfArrays /*outputs*/ ) override { auto context = reinterpret_cast(context_); diff --git a/modules/dnn/src/layers/flatten_layer.cpp b/modules/dnn/src/layers/flatten_layer.cpp index e628bc3b42..7d29408f0c 100644 --- a/modules/dnn/src/layers/flatten_layer.cpp +++ b/modules/dnn/src/layers/flatten_layer.cpp @@ -281,15 +281,18 @@ public: #ifdef HAVE_CUDA Ptr initCUDA( void *context_, - const std::vector>& inputs, - const std::vector>& outputs + InputArrayOfArrays inputs_, + InputArrayOfArrays outputs ) override { auto context = reinterpret_cast(context_); - if (inputs[0]->getHostMatDepth() == CV_Bool) + std::vector inputs; + inputs_.getGpuMatNDVector(inputs); + int depth = CV_MAT_DEPTH(inputs[0].type()); + if (depth == CV_Bool) return make_cuda_node_bool(std::move(context->stream)); else - return make_cuda_node_with_type(preferableTarget, inputs[0]->getHostMatDepth(), std::move(context->stream)); + return make_cuda_node_with_type(preferableTarget, depth, std::move(context->stream)); } #endif diff --git a/modules/dnn/src/layers/gemm_layer.cpp b/modules/dnn/src/layers/gemm_layer.cpp index 8e4bd9465a..bac4998455 100644 --- a/modules/dnn/src/layers/gemm_layer.cpp +++ b/modules/dnn/src/layers/gemm_layer.cpp @@ -476,18 +476,19 @@ public: #ifdef HAVE_CUDA // Y = A * B + C. B should be guaranteed as two dimensional. Ptr initCUDA(void *context_, - const std::vector>& inputs, - const std::vector>& outputs) CV_OVERRIDE { + InputArrayOfArrays inputs_, + InputArrayOfArrays outputs) CV_OVERRIDE { CV_CheckFalse(trans_a, "DNN/Gemm/Cuda: does not support transA"); CV_CheckTrue(const_B, "DNN/Gemm/Cuda: input B (weight) is required to be constant"); auto context = reinterpret_cast(context_); - auto wrapper_A = inputs[0].dynamicCast(); + std::vector inputs; + inputs_.getGpuMatNDVector(inputs); auto B = blobs[0]; auto C = have_bias && const_C ? blobs[1] : Mat(); // in most cases C is constant if (!trans_b) cv::transpose(B, B); - auto flatten_start_axis = normalize_axis(1, wrapper_A->getRank()); + auto flatten_start_axis = normalize_axis(1, (int)inputs[0].size.size()); return make_cuda_node(preferableTarget, std::move(context->stream), std::move(context->cublas_handle), flatten_start_axis, B, C); } #endif // HAVE_CUDA diff --git a/modules/dnn/src/layers/maxpool_layer.cpp b/modules/dnn/src/layers/maxpool_layer.cpp index bb156d1aba..d4f000ca08 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,49 @@ 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_, + InputArrayOfArrays inputs_, + InputArrayOfArrays) CV_OVERRIDE + { + auto context = reinterpret_cast(context_); + std::vector inputs; + inputs_.getGpuMatNDVector(inputs); + MatShape inShape = inputs[0].size; + 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/layers/nary_eltwise_layers.cpp b/modules/dnn/src/layers/nary_eltwise_layers.cpp index 650036d0c3..b48b11787f 100644 --- a/modules/dnn/src/layers/nary_eltwise_layers.cpp +++ b/modules/dnn/src/layers/nary_eltwise_layers.cpp @@ -1320,11 +1320,13 @@ public: #ifdef HAVE_CUDA Ptr initCUDA( void *context_, - const std::vector>& inputs, - const std::vector>& outputs + InputArrayOfArrays inputs_, + InputArrayOfArrays outputs ) override { auto context = reinterpret_cast(context_); + std::vector inputs; + inputs_.getGpuMatNDVector(inputs); cuda4dnn::EltwiseOpType op_ = cuda4dnn::EltwiseOpType::SUM; switch (op) { @@ -1361,7 +1363,7 @@ public: default: return Ptr(); // return empty cuda_node if the EltwiseOpType is unsupported type. }; - return make_cuda_node_with_type(preferableTarget, inputs[0]->getHostMatDepth(), std::move(context->stream), op_, std::vector()); + return make_cuda_node_with_type(preferableTarget, CV_MAT_DEPTH(inputs[0].type()), std::move(context->stream), op_, std::vector()); } #endif diff --git a/modules/dnn/src/layers/pooling_layer.cpp b/modules/dnn/src/layers/pooling_layer.cpp index 1f3932dc7b..a912eeb791 100644 --- a/modules/dnn/src/layers/pooling_layer.cpp +++ b/modules/dnn/src/layers/pooling_layer.cpp @@ -373,16 +373,18 @@ public: #ifdef HAVE_CUDA Ptr initCUDA( void *context_, - const std::vector>& inputs, - const std::vector>& outputs + InputArrayOfArrays inputs_, + InputArrayOfArrays outputs_ ) override { auto context = reinterpret_cast(context_); if (type == ROI) return make_cuda_node(preferableTarget, std::move(context->stream), spatialScale); - auto input_wrapper = inputs[0].dynamicCast(); - auto input_shape = input_wrapper->getShape(); + std::vector inputs, outputs; + inputs_.getGpuMatNDVector(inputs); + outputs_.getGpuMatNDVector(outputs); + MatShape input_shape = inputs[0].size; /* storing max indices is a special case and we deal with it separately */ if (computeMaxIdx) { @@ -412,13 +414,13 @@ public: config.input_shape.assign(std::begin(input_shape), std::end(input_shape)); - int indicesType = outputs[1]->getHostMatDepth(); + int indicesType = CV_MAT_DEPTH(outputs[1].type()); CV_CheckType(indicesType, indicesType == CV_32S || indicesType == CV_64S, "Unsupported indices type"); if (indicesType == CV_32S) - return make_cuda_node_with_indices(preferableTarget, inputs[0]->getHostMatDepth(), std::move(context->stream), config); + return make_cuda_node_with_indices(preferableTarget, CV_MAT_DEPTH(inputs[0].type()), std::move(context->stream), config); else if (indicesType == CV_64S) - return make_cuda_node_with_indices(preferableTarget, inputs[0]->getHostMatDepth(), std::move(context->stream), config); + return make_cuda_node_with_indices(preferableTarget, CV_MAT_DEPTH(inputs[0].type()), std::move(context->stream), config); CV_Error(Error::BadDepth, "Unsupported indices type"); return Ptr(); 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..796d9a1a2d 100644 --- a/modules/dnn/src/net_impl.cpp +++ b/modules/dnn/src/net_impl.cpp @@ -154,6 +154,9 @@ void Net::Impl::clear() prepared = false; finalizeLayers = true; + finalized = false; + fusedSnapshotValid = false; + fusedSnapshot.clear(); } @@ -274,11 +277,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)); @@ -1419,6 +1422,9 @@ void Net::Impl::getLayerShapes(const ShapesVec& netInputShapes, LayerShapes& shapes) { if (mainGraph) { + // Fusion emits block-layout layers; their TransformLayout conversions are + // only inserted by finalize(), so shape inference must run post-finalize. + finalize(); std::vector shapeCache; std::vector typeCache; CV_Assert(layerId == 0); @@ -1679,7 +1685,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 +2391,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 +2422,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 +2507,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; @@ -2551,6 +2557,7 @@ int64 Net::Impl::getFLOPS(const std::vector& netInputShapes, const std::vector& netInputTypes) /*const*/ { if (mainGraph) { + finalize(); // The new graph engine executes in FP32 on CPU regardless of the requested // target, so FP16 input types (e.g. coming from an OpenCL FP16 target) would be // rejected by Layer::getTypes(). Normalize them to FP32 for shape/FLOPS inference. @@ -2584,6 +2591,7 @@ int64 Net::Impl::getFLOPS( const std::vector& netInputTypes) /*const*/ { if (mainGraph) { + finalize(); std::vector inputTypes = filterFP16InputTypes(netInputTypes); LayerShapes shapes; std::vector shapeCache; @@ -2595,7 +2603,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 +2702,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 +3017,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 +3050,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..c59e5cb290 100644 --- a/modules/dnn/src/net_impl.hpp +++ b/modules/dnn/src/net_impl.hpp @@ -143,6 +143,19 @@ 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 + + // Post-fusion (pre block-layout) snapshot so finalize() can re-run from a clean + // state on a backend/target change; useBlockLayout() is destructive and must + // run after backend assignment (see deviceOp handling in graph_block_layout.cpp). + struct FusedGraphSnapshot { + Ptr graph; + std::vector > prog; + std::vector > inputs; + std::vector > outputs; + }; + bool fusedSnapshotValid = false; + std::vector fusedSnapshot; TracingMode tracingMode; ProfilingMode profilingMode; std::vector dimvalues; @@ -280,6 +293,22 @@ struct Net::Impl : public detail::NetImplBase std::unique_ptr cudaInfo; void initCUDABackend(const std::vector& blobsToKeep_); + + // New graph engine: per-Arg device-resident tensors owned directly by the net (no backend + // wrappers). Sized lazily via GpuMatND::fit() and reused across forwards. Dirty flags track + // which copy (host cv::Mat vs device GpuMatND) is authoritative so transfers happen only at + // CPU<->CUDA boundaries; intermediates stay device-resident across consecutive CUDA ops. + std::vector cudaArgBuffers; + std::vector cudaArgHostDirty; // 1: host copy is authoritative -> needs H2D before device read + std::vector cudaArgDeviceDirty; // 1: device copy is authoritative -> needs D2H before host read + + // Device element type for a host tensor (half for float tensors under the FP16 target). + int cudaDeviceType(const Mat& hostMat) const; + // Returns the device buffer for @p arg, fit() to the host Mat's shape and device type. + cuda::GpuMatND& getCudaArgBuffer(Arg arg, const Mat& hostMat); + void cudaSetHostDirty(Arg arg); // mark host authoritative (e.g. after a CPU op wrote it) + void cudaUploadArg(Arg arg, const Mat& hostMat); // H2D if host dirty + void cudaDownloadArg(Arg arg, Mat& hostMat); // D2H if device dirty #endif #ifdef HAVE_ONNXRUNTIME @@ -420,12 +449,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); + // Save/restore the fused graph so finalize() is re-entrant across backend changes. + void saveFusedSnapshot(); + void restoreFusedSnapshot(); // 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 +558,9 @@ struct Net::Impl : public detail::NetImplBase }; // Net::Impl -inline Net::Impl* getNetImpl(const Layer* layer) +inline Net::Impl* getNetImpl(const LayerInfo* 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..4a504d181d 100644 --- a/modules/dnn/src/net_impl2.cpp +++ b/modules/dnn/src/net_impl2.cpp @@ -8,6 +8,10 @@ #include +#ifdef HAVE_CUDA +#include +#endif + #ifdef HAVE_ONNXRUNTIME #include #endif @@ -316,30 +320,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 +382,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 +404,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 +575,174 @@ 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); + +#ifdef HAVE_CUDA + // Whole-graph CUDA gate: run on CUDA only if *every* op has a CUDA executor; if any op is + // unsupported (or the graph has control-flow subgraphs) run the entire graph on CPU. This + // avoids partial CPU<->CUDA execution and its host/device coherence hazards for now. + std::vector > cudaExecs; + bool graphOnCuda = false; + if (useCUDA && cudaInfo) { + graphOnCuda = true; + cudaExecs.assign(nops, Ptr()); + for (i = 0; i < nops; i++) { + const Ptr& op = prog[i]; + if (!op) + continue; + if (op->subgraphs()) { graphOnCuda = false; break; } // control-flow bodies stay on CPU + Ptr e = LayerFactory::createExec(op->type, DNN_BACKEND_CUDA, op, &cudaInfo->context); + if (!e) { graphOnCuda = false; break; } + cudaExecs[i] = e; + } + if (!graphOnCuda) { + cudaExecs.clear(); + CV_LOG_INFO(NULL, cv::format("DNN/NewEngine: graph '%s' has layer(s) without CUDA support; " + "running the whole graph on CPU", graph->name().c_str())); + } + } +#endif + + 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 + if (graphOnCuda && cudaExecs[i]) { + exec = cudaExecs[i]; + 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::saveFusedSnapshot() +{ + fusedSnapshot.clear(); + for (const Ptr& g : allgraphs) { + FusedGraphSnapshot snap; + snap.graph = g; + const std::vector >& prog = g->prog(); + snap.prog = prog; + snap.inputs.reserve(prog.size()); + snap.outputs.reserve(prog.size()); + for (const Ptr& op : prog) { + snap.inputs.push_back(op ? op->inputs : std::vector()); + snap.outputs.push_back(op ? op->outputs : std::vector()); + } + fusedSnapshot.push_back(std::move(snap)); + } +} + +void Net::Impl::restoreFusedSnapshot() +{ + // Roll the graph back to its post-fusion state: undo the layer-input rewiring + // and remove the TransformLayout ops inserted by a previous useBlockLayout(). + for (const FusedGraphSnapshot& snap : fusedSnapshot) { + for (size_t i = 0; i < snap.prog.size(); i++) { + const Ptr& op = snap.prog[i]; + if (!op) + continue; + op->inputs = snap.inputs[i]; + op->outputs = snap.outputs[i]; + } + snap.graph->setProg(snap.prog); + } + totalLayers = updateGraphOfs(mainGraph, 0, true); +} + +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; + + // Snapshot the fused graph once so finalize() can re-run cleanly on a + // backend/target change (block layout + buffer assignment are destructive). + if (!fusedSnapshotValid) { + saveFusedSnapshot(); + fusedSnapshotValid = true; + } else { + restoreFusedSnapshot(); + } + + bool useCUDA = false; +#ifdef HAVE_CUDA + cudaArgBuffers.clear(); + cudaArgHostDirty.clear(); + cudaArgDeviceDirty.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 +846,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; @@ -1133,6 +1311,125 @@ void Net::Impl::setGraphInput(Ptr& graph, size_t idx, const Mat& m) } } +#ifdef HAVE_CUDA +// cv::cuda::Stream view over the (non-owning) cuda4dnn inference stream, so GpuMatND transfers +// are ordered against the op compute that runs on the same cudaStream_t. +static inline cuda::Stream wrapCudaStream(cuda4dnn::csl::Stream& s) +{ + return cuda::StreamAccessor::wrapStream(s.get()); +} + +// GpuMatND upload/download require at least one dimension (rank-0 shapes trip GpuMatND::setFields). +// View a rank-0 scalar as a 1-element 1D header over the same data; higher-rank Mats pass through. +static inline Mat asAtLeast1D(const Mat& m) +{ + if (m.shape().dims != 0) + return m; + int one = 1; + return Mat(1, &one, m.type(), const_cast(m.data)); +} + +// Device element type for a host tensor: float tensors are stored as half under the FP16 target, +// everything else mirrors the host type. fit() reuses the existing allocation when large enough, +// so buffers persist across forwards. +int Net::Impl::cudaDeviceType(const Mat& hostMat) const +{ + if (preferableTarget == DNN_TARGET_CUDA_FP16 && CV_MAT_DEPTH(hostMat.type()) == CV_32F) + return CV_MAKETYPE(CV_16F, CV_MAT_CN(hostMat.type())); + return hostMat.type(); +} + +cuda::GpuMatND& Net::Impl::getCudaArgBuffer(Arg arg, const Mat& hostMat) +{ + int idx = arg.idx; + if ((int)cudaArgBuffers.size() != (int)args.size()) { + cudaArgBuffers.assign(args.size(), cuda::GpuMatND()); + cudaArgHostDirty.assign(args.size(), 1); + cudaArgDeviceDirty.assign(args.size(), 0); + } + // GpuMatND (and cuda4dnn tensors) require at least one dimension; represent a rank-0 scalar + // tensor as a 1-element 1D tensor so element-wise ops still see the single value. + MatShape shape = hostMat.shape(); + if (shape.dims == 0) + shape = MatShape({1}); + cudaArgBuffers[idx].fit(shape, cudaDeviceType(hostMat)); + return cudaArgBuffers[idx]; +} + +void Net::Impl::cudaSetHostDirty(Arg arg) +{ + int idx = arg.idx; + if (idx >= 0 && idx < (int)cudaArgHostDirty.size()) { + cudaArgHostDirty[idx] = 1; + cudaArgDeviceDirty[idx] = 0; + } +} + +void Net::Impl::cudaUploadArg(Arg arg, const Mat& hostMat) +{ + int idx = arg.idx; + Mat src = asAtLeast1D(hostMat); + cuda::GpuMatND& g = getCudaArgBuffer(arg, src); + if (cudaArgHostDirty[idx]) { + cuda::Stream s = wrapCudaStream(cudaInfo->context.stream); + if (g.type() == src.type()) { + g.upload(src, s); + } else { + // FP32 -> FP16 (device stores half): convert on host, then copy up. + Mat tmp; + src.convertTo(tmp, CV_MAT_DEPTH(g.type())); + g.upload(tmp, s); + } + cudaArgHostDirty[idx] = 0; + cudaArgDeviceDirty[idx] = 0; + } +} + +void Net::Impl::cudaDownloadArg(Arg arg, Mat& hostMat) +{ + int idx = arg.idx; + if (idx < 0 || idx >= (int)cudaArgDeviceDirty.size()) + return; + if (cudaArgDeviceDirty[idx]) { + cuda::GpuMatND& g = cudaArgBuffers[idx]; + cuda::Stream s = wrapCudaStream(cudaInfo->context.stream); + Mat dst = asAtLeast1D(hostMat); // fill the scalar's storage through a 1D header + if (g.type() == dst.type()) { + g.download(dst, s); + cudaInfo->context.stream.synchronize(); // host read follows immediately + } else { + // device stores half: copy down, then convert up to the host FP32 tensor. + Mat tmp; + g.download(tmp, s); + cudaInfo->context.stream.synchronize(); + tmp.convertTo(dst, dst.type()); + } + cudaArgDeviceDirty[idx] = 0; + cudaArgHostDirty[idx] = 0; + } +} + +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 inpG(inputs.size()), outG(outputs.size()); + for (size_t i = 0; i < inputs.size(); i++) { + netimpl->cudaUploadArg(inputs[i], inpMats[i]); // H2D only if host-authoritative + inpG[i] = netimpl->cudaArgBuffers[inputs[i].idx]; + } + for (size_t i = 0; i < outputs.size(); i++) { + outG[i] = netimpl->getCudaArgBuffer(outputs[i], outMats[i]); + int oidx = outputs[i].idx; // op writes the device buffer + netimpl->cudaArgDeviceDirty[oidx] = 1; + netimpl->cudaArgHostDirty[oidx] = 0; + } + exec->forwardCUDA(inpG, outG, &netimpl->cudaInfo->workspace); +} +#endif + // Slice a Scan input at index `idx` along `axis`, removing that axis (contiguous result). static Mat sliceScanAxis(const Mat& m, int axis, int idx) { @@ -1184,7 +1481,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(); @@ -1208,13 +1506,31 @@ void Net::Impl::forwardGraph(Ptr& graph, InputArrayOfArrays inputs_, setGraphInput(graph, i, m); } } +#ifdef HAVE_CUDA + // Graph inputs are host-authoritative and may be updated between forward() calls + // (setInput writes them outside this function, so the loop above is skipped when + // forward() is called with no explicit inputs). Mark their device copies stale each + // forward so the current input is re-uploaded; otherwise a second forward with a + // changed input would read the previous forward's stale device data. + if (cudaInfo) { + for (i = 0; i < n_gr_inputs; i++) + cudaSetHostDirty(gr_inputs[i]); + } +#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); @@ -1225,6 +1541,13 @@ void Net::Impl::forwardGraph(Ptr& graph, InputArrayOfArrays inputs_, for (i = 0; i < ninputs; i++) { Arg inp = inputs[i]; +#ifdef HAVE_CUDA + // CPU op: bring any device-resident input back to host before reading its + // shape/data. This must happen before allocateLayerOutputs(), which may fit() + // an in-place op's output onto this very buffer and rewrite its header. + if (opBackend != DNN_BACKEND_CUDA) + cudaDownloadArg(inp, argTensor(inp)); +#endif const Mat& m = argTensor(inp); inpMats[i] = m; inpTypes[i] = m.type(); @@ -1233,15 +1556,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 +1577,26 @@ 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 + { + // Device-resident inputs were already synced to host in the capture loop above. + 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++) + cudaSetHostDirty(outputs[k]); +#endif + } } else { Ptr iflayer = layer.dynamicCast(); @@ -1425,7 +1763,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 +1846,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 >= 0) { + Mat& t = argTensor(out); + cudaDownloadArg(out, t); + } +#endif const Mat& outm = argTensor(out); if (isMainGraph) { if (outm.size.layout == DATA_LAYOUT_BLOCK) { @@ -1532,8 +1877,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 +1909,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 +2063,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..3b63e59114 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; } @@ -1761,7 +1764,7 @@ void ONNXImporter2::parseScan(LayerParams& layerParams, } addLayer(layerParams, node_proto); - Ptr& scanLayer = curr_prog.back(); + Ptr& scanLayer = curr_prog.back(); *scanLayer->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..bb5eb7839e 100644 --- a/modules/dnn/src/op_cuda.cpp +++ b/modules/dnn/src/op_cuda.cpp @@ -8,10 +8,73 @@ #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->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(InputArrayOfArrays inputs_, + OutputArrayOfArrays outputs_, + void* workspace) CV_OVERRIDE + { + std::vector inputs, outputs; + inputs_.getGpuMatNDVector(inputs); + outputs_.getGpuMatNDVector(outputs); + + 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/op_cuda.hpp b/modules/dnn/src/op_cuda.hpp index ecc8b3c856..2914d40e85 100644 --- a/modules/dnn/src/op_cuda.hpp +++ b/modules/dnn/src/op_cuda.hpp @@ -18,6 +18,7 @@ #include #include +#include #include #include @@ -177,6 +178,22 @@ namespace cv { namespace dnn { if (temp.data != destMat.data) temp.copyTo(destMat); } + + /** @brief builds a read-only TensorView over the device memory of a GpuMatND (no copy) */ + template + TensorView viewOf(const cuda::GpuMatND& g) { + using const_ptr = typename TensorView::const_pointer; + return TensorView(const_ptr(reinterpret_cast(g.getDevicePtr())), + std::begin(g.size), std::end(g.size)); + } + + /** @brief builds a writable TensorSpan over the device memory of a GpuMatND (no copy) */ + template + TensorSpan spanOf(const cuda::GpuMatND& g) { + using ptr = typename TensorSpan::pointer; + return TensorSpan(ptr(reinterpret_cast(g.getDevicePtr())), + std::begin(g.size), std::end(g.size)); + } }} /* namespace cuda4dnn::csl */ /** base class for CUDA operation nodes (for all supported targets) */ @@ -185,10 +202,25 @@ namespace cv { namespace dnn { CUDABackendNode() : BackendNode(DNN_BACKEND_CUDA) { } virtual ~CUDABackendNode() { } + /** classic-engine entry point (wrapper-based). + * + * The default adapts the wrappers to GpuMatND headers and dispatches to the GpuMatND + * overload, so ops ported to the new graph engine only implement the GpuMatND forward. + * Ops not yet ported keep overriding this method directly. + */ virtual void forward( const std::vector>& inputs, const std::vector>& outputs, - cuda4dnn::csl::Workspace& workspace) = 0; + cuda4dnn::csl::Workspace& workspace); + + /** new graph-engine entry point (wrapper-free): operates directly on GpuMatND device tensors */ + virtual void forward( + const std::vector& inputs, + const std::vector& outputs, + cuda4dnn::csl::Workspace& workspace) + { + CV_Error(Error::StsNotImplemented, "GpuMatND CUDA forward is not implemented for this operation"); + } virtual std::size_t get_workspace_memory_in_bytes() const noexcept { return 0; } }; @@ -229,7 +261,7 @@ namespace cv { namespace dnn { template