From ad3e9194dba941a57468ad25f4c8b47885a41769 Mon Sep 17 00:00:00 2001 From: sanjay Date: Fri, 6 Nov 2026 12:56:53 +0530 Subject: [PATCH] refactored suggested changes --- modules/dnn/src/cuda/argmax.cu | 2 +- modules/dnn/src/cuda/functors.hpp | 11 ++-- modules/dnn/src/cuda/quantize_dequantize.cu | 6 ++- modules/dnn/src/cuda/range.cu | 16 ++++-- modules/dnn/src/cuda/roi_align.cu | 8 ++- modules/dnn/src/cuda/topk.cu | 42 ++++++++++++---- .../dnn/src/cuda4dnn/kernels/roi_align.hpp | 2 +- .../src/cuda4dnn/primitives/layer_norm.hpp | 6 ++- .../dnn/src/cuda4dnn/primitives/one_hot.hpp | 6 +-- .../dnn/src/cuda4dnn/primitives/roi_align.hpp | 1 + modules/dnn/src/layers/conv2_layer.cpp | 6 ++- modules/dnn/src/layers/fusion_cuda.hpp | 6 ++- modules/dnn/src/layers/matmulnbits_layer.cpp | 50 +++++++++++++++++++ modules/dnn/src/layers/range_layer.cpp | 3 -- modules/dnn/src/net_impl2.cpp | 13 +++-- modules/dnn/src/onnx/onnx_importer2.cpp | 11 ++-- modules/dnn/src/op_cuda.cpp | 2 + modules/dnn/test/test_onnx_conformance.cpp | 12 +++++ ...rmance_layer_filter__cuda_denylist.inl.hpp | 10 ++++ 19 files changed, 174 insertions(+), 39 deletions(-) diff --git a/modules/dnn/src/cuda/argmax.cu b/modules/dnn/src/cuda/argmax.cu index 69960376f9..00d9c9d95c 100644 --- a/modules/dnn/src/cuda/argmax.cu +++ b/modules/dnn/src/cuda/argmax.cu @@ -35,7 +35,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { std::int64_t best_index = 0; for (int k = 1; k < axis_size; ++k) { const T value = p_input[base + static_cast(k) * inner_size]; - if constexpr (IsArgMax) { + if (IsArgMax) { if (SelectLast ? value >= best_value : value > best_value) { best_value = value; best_index = k; diff --git a/modules/dnn/src/cuda/functors.hpp b/modules/dnn/src/cuda/functors.hpp index 0c418886e2..2040c637e6 100644 --- a/modules/dnn/src/cuda/functors.hpp +++ b/modules/dnn/src/cuda/functors.hpp @@ -864,12 +864,13 @@ struct PReluFunctor { CUDA4DNN_DEVICE PReluFunctor(const Params& params) { } CUDA4DNN_DEVICE T operator()(T x, T y) { - // an unsigned value is never negative, so PRelu passes it through - if constexpr (std::is_unsigned::value) - return x; - else - return x >= T(0) ? x : T(x * y); + return apply(x, y, std::is_unsigned{}); } + + // an unsigned value is never negative, so PRelu passes it through; an overload + // instead of if constexpr (C++17), which .cu files built as C++14 can't use + CUDA4DNN_DEVICE static T apply(T x, T, std::true_type) { return x; } + CUDA4DNN_DEVICE static T apply(T x, T y, std::false_type) { return x >= T(0) ? x : T(x * y); } }; // bitwise functors are instantiated for integer types only diff --git a/modules/dnn/src/cuda/quantize_dequantize.cu b/modules/dnn/src/cuda/quantize_dequantize.cu index 4fca85b2e0..1a66b28dac 100644 --- a/modules/dnn/src/cuda/quantize_dequantize.cu +++ b/modules/dnn/src/cuda/quantize_dequantize.cu @@ -53,7 +53,11 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { for (auto i : grid_stride_range(n)) { std::size_t a_idx = (i / slice_size) % sz_a; int zp = zero_point ? static_cast(zero_point[a_idx]) : 0; - int rounded = __float2int_rn(input[i] / scale[a_idx]); + // clamp before converting: __float2int_rn turns +Inf or a huge value into + // INT_MIN, which would saturate to qmin instead of qmax + float scaled = input[i] / scale[a_idx]; + scaled = (scaled != scaled) ? 0.f : fminf(fmaxf(scaled, -1024.f), 1024.f); + int rounded = __float2int_rn(scaled); output[i] = saturate_quantized(rounded + zp); } } diff --git a/modules/dnn/src/cuda/range.cu b/modules/dnn/src/cuda/range.cu index e005b4d9a7..7cad7fe4f9 100644 --- a/modules/dnn/src/cuda/range.cu +++ b/modules/dnn/src/cuda/range.cu @@ -30,13 +30,23 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { __device__ inline double as_double(std::int32_t v) { return static_cast(v); } __device__ inline double as_double(std::int64_t v) { return static_cast(v); } + template + __device__ inline T range_value(T start, T delta, std::size_t i) { + return static_cast(as_double(start) + static_cast(i) * as_double(delta)); + } + + // int64 stays in integer arithmetic: double holds integers exactly only up to 2^53 + __device__ inline std::int64_t range_value(std::int64_t start, std::int64_t delta, std::size_t i) { + return start + static_cast(i) * delta; + } + template __global__ void range(Span output, const T* start_p, const T* delta_p) { - const double start = as_double(*start_p); - const double delta = as_double(*delta_p); + const T start = *start_p; + const T delta = *delta_p; for (auto i : grid_stride_range(output.size())) - output[i] = static_cast(start + static_cast(i) * delta); + output[i] = range_value(start, delta, static_cast(i)); } } diff --git a/modules/dnn/src/cuda/roi_align.cu b/modules/dnn/src/cuda/roi_align.cu index ac29d82ac2..b5e51a7a2b 100644 --- a/modules/dnn/src/cuda/roi_align.cu +++ b/modules/dnn/src/cuda/roi_align.cu @@ -75,6 +75,12 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { const float step_x = bin_size_w / gw; const std::int64_t b = read_batch_index(batch_indices, batch_index_element_size, r); + // a kernel can't raise the CPU path's error for an index outside the batch, + // so it writes 0 instead of reading outside the input + if (b < 0 || b >= p.batch_size) { + output[id] = static_cast(0.f); + continue; + } const T* img = input.data().get() + ((b * p.channels) + c) * static_cast(p.height) * p.width; @@ -133,7 +139,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { CV_Assert(batch_index_element_size == sizeof(std::int32_t) || batch_index_element_size == sizeof(std::int64_t)); CV_Assert(params.output_height > 0 && params.output_width > 0); - CV_Assert(params.height > 0 && params.width > 0 && params.channels > 0); + CV_Assert(params.batch_size > 0 && params.height > 0 && params.width > 0 && params.channels > 0); if (output.size() == 0) return; diff --git a/modules/dnn/src/cuda/topk.cu b/modules/dnn/src/cuda/topk.cu index c4304e041f..f85a782ea2 100644 --- a/modules/dnn/src/cuda/topk.cu +++ b/modules/dnn/src/cuda/topk.cu @@ -32,22 +32,41 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { namespace raw { + /* float is exact for __half, float, int8 and uint8; int32/int64 compare in + * their own type, since float holds integers exactly only up to 2^24. */ + template struct TopKCompare { using type = float; }; + template <> struct TopKCompare { using type = std::int32_t; }; + template <> struct TopKCompare { using type = std::int64_t; }; + + /* Starting value that any real element beats: infinity for float types, + * and the lowest or highest integer for int32/int64. */ + template + __device__ inline C topk_worst(bool largest) { + using limits = ::cuda::std::numeric_limits; + if (limits::has_infinity) + return largest ? -limits::infinity() : limits::infinity(); + return largest ? limits::lowest() : limits::max(); + } + + template struct Candidate { - float value; + C value; int index; }; /* Strict weak ordering matching topk2_layer.cpp's comparators: better value * first, and on an exact tie the smaller index. */ - __device__ __forceinline__ bool better_than(const Candidate& a, const Candidate& b, bool largest) { + template + __device__ inline bool better_than(const Candidate& a, const Candidate& b, bool largest) { if (a.value == b.value) return a.index < b.index; return largest ? (a.value > b.value) : (a.value < b.value); } + template struct MergeCandidate { bool largest; - __device__ __forceinline__ Candidate operator()(const Candidate& a, const Candidate& b) const { + __device__ inline Candidate operator()(const Candidate& a, const Candidate& b) const { return better_than(a, b, largest) ? a : b; } }; @@ -56,13 +75,14 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { __global__ void topk(Span values, Span indices, View input, int outer, int dim_axis, int inner, int k, bool largest) { - using BlockReduce = cub::BlockReduce; + using C = typename TopKCompare::type; + using Cand = Candidate; + using BlockReduce = cub::BlockReduce; __shared__ typename BlockReduce::TempStorage reduce_storage; - __shared__ Candidate selected; + __shared__ Cand selected; const int nslices = outer * inner; - const float worst = largest ? -::cuda::std::numeric_limits::infinity() - : ::cuda::std::numeric_limits::infinity(); + const C worst = topk_worst(largest); for (int slice = blockIdx.x; slice < nslices; slice += gridDim.x) { const int b = slice / inner; @@ -72,11 +92,11 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { const std::int64_t out_base = static_cast(b) * k * inner + j; for (int rank = 0; rank < k; ++rank) { - Candidate best{worst, ::cuda::std::numeric_limits::max()}; + Cand best{worst, ::cuda::std::numeric_limits::max()}; for (int u = threadIdx.x; u < dim_axis; u += BlockSize) { - const float v = static_cast(input.data().get()[in_base + static_cast(u) * inner]); - const Candidate cand{v, u}; + const C v = static_cast(input.data().get()[in_base + static_cast(u) * inner]); + const Cand cand{v, u}; /* Everything at or before the previous pick is spent. Ordering * strictly means no element can be chosen twice, and duplicates @@ -88,7 +108,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { best = cand; } - const Candidate winner = BlockReduce(reduce_storage).Reduce(best, MergeCandidate{largest}); + const Cand winner = BlockReduce(reduce_storage).Reduce(best, MergeCandidate{largest}); if (threadIdx.x == 0) { selected = winner; diff --git a/modules/dnn/src/cuda4dnn/kernels/roi_align.hpp b/modules/dnn/src/cuda4dnn/kernels/roi_align.hpp index ac4ccbe937..6c970bdadc 100644 --- a/modules/dnn/src/cuda4dnn/kernels/roi_align.hpp +++ b/modules/dnn/src/cuda4dnn/kernels/roi_align.hpp @@ -15,7 +15,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { struct RoiAlignParams { - int channels{}, height{}, width{}; + int batch_size{}, channels{}, height{}, width{}; int output_height{}, output_width{}; int sampling_ratio{}; float spatial_scale{1.f}; diff --git a/modules/dnn/src/cuda4dnn/primitives/layer_norm.hpp b/modules/dnn/src/cuda4dnn/primitives/layer_norm.hpp index f1b57664ce..84c599eb1d 100644 --- a/modules/dnn/src/cuda4dnn/primitives/layer_norm.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/layer_norm.hpp @@ -77,7 +77,8 @@ namespace cv { namespace dnn { namespace cuda4dnn { auto loops = input.size_range(0, axis); auto norm_size = input.size_range(axis, input.rank()); - if (norm_size == 1) { + // with a size-one axis the result is 0, or the bias when there is one + if (norm_size == 1 && bias.empty()) { kernels::fill(stream, output, 0.f); return; } else { @@ -114,8 +115,9 @@ namespace cv { namespace dnn { namespace cuda4dnn { auto loops = input.size_range(0, axis); auto norm_size = input.size_range(axis, input.rank()); + // with a size-one axis the result is 0, or the bias when there is one; // the optional Mean/InvStdDev outputs still need the statistics - if (norm_size == 1 && outputs.size() == 1) { + if (norm_size == 1 && outputs.size() == 1 && bias.empty()) { kernels::fill(stream, output, 0.f); return; } diff --git a/modules/dnn/src/cuda4dnn/primitives/one_hot.hpp b/modules/dnn/src/cuda4dnn/primitives/one_hot.hpp index ce438f06f2..d6d6660241 100644 --- a/modules/dnn/src/cuda4dnn/primitives/one_hot.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/one_hot.hpp @@ -25,7 +25,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { public: OneHotOp(csl::Stream stream_, int axis_, double off_, double on_, bool runtimeValues_) : stream(std::move(stream_)), axis(axis_), - offValue(static_cast(static_cast(off_))), onValue(static_cast(static_cast(on_))), + offValue(static_cast(off_)), onValue(static_cast(on_)), runtimeValues(runtimeValues_) { } @@ -50,8 +50,8 @@ namespace cv { namespace dnn { namespace cuda4dnn { inputs[2].copyTo(values); CV_Assert(values.total() == 2); values.convertTo(values64, CV_64F); - offValue = static_cast(static_cast(values64.ptr()[0])); - onValue = static_cast(static_cast(values64.ptr()[1])); + offValue = static_cast(values64.ptr()[0]); + onValue = static_cast(values64.ptr()[1]); } std::size_t inner = 1; diff --git a/modules/dnn/src/cuda4dnn/primitives/roi_align.hpp b/modules/dnn/src/cuda4dnn/primitives/roi_align.hpp index 9437c0d2d9..ff145cbaca 100644 --- a/modules/dnn/src/cuda4dnn/primitives/roi_align.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/roi_align.hpp @@ -45,6 +45,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { * per-run fields are filled here and the attribute-derived ones come from * the layer. */ kernels::RoiAlignParams p = params; + p.batch_size = inShape[0]; p.channels = inShape[1]; p.height = inShape[2]; p.width = inShape[3]; diff --git a/modules/dnn/src/layers/conv2_layer.cpp b/modules/dnn/src/layers/conv2_layer.cpp index 1d22247815..7592fa901a 100644 --- a/modules/dnn/src/layers/conv2_layer.cpp +++ b/modules/dnn/src/layers/conv2_layer.cpp @@ -559,7 +559,7 @@ public: static bool cudaPostKernel(ActivationFunc fn) { return fn == getActivationFunc(ACTIV_HARDSWISH) || fn == getActivationFunc(ACTIV_GELU) || - fn == getActivationFunc(ACTIV_GELU_APPROX); + fn == getActivationFunc(ACTIV_GELU_APPROX) || fn == getActivationFunc(ACTIV_EXP); } //! Takes an expression the layer has a kernel for. Only a one-layer chain carries @@ -1281,6 +1281,10 @@ public: post = make_cuda_node(targetId, context.stream); } else if (activationFunc == getActivationFunc(ACTIV_GELU_APPROX)) { post = make_cuda_node(targetId, context.stream); + } else if (activationFunc == getActivationFunc(ACTIV_EXP)) { + const float scale = activParams.size() > 0 ? activParams[0] : 1.f; + const float shift = activParams.size() > 1 ? activParams[1] : 0.f; + post = make_cuda_node(targetId, context.stream, scale, shift); } config.activation_type = ConvolutionConfiguration::ActivationType::IDENTITY; diff --git a/modules/dnn/src/layers/fusion_cuda.hpp b/modules/dnn/src/layers/fusion_cuda.hpp index af8b708400..44ee1b0224 100644 --- a/modules/dnn/src/layers/fusion_cuda.hpp +++ b/modules/dnn/src/layers/fusion_cuda.hpp @@ -45,7 +45,7 @@ inline bool cudaSupportsFusion(const PreparedFusion& fusion, int channels) return fn == getActivationFunc(ACTIV_GELU) || fn == getActivationFunc(ACTIV_GELU_APPROX) || fn == getActivationFunc(ACTIV_RELU) || fn == getActivationFunc(ACTIV_SWISH) || fn == getActivationFunc(ACTIV_SIGMOID) || fn == getActivationFunc(ACTIV_TANH) || - fn == getActivationFunc(ACTIV_MISH); + fn == getActivationFunc(ACTIV_MISH) || fn == getActivationFunc(ACTIV_HARDSIGMOID); if (channels < 0) return true; return (fusion.scaleBuf.empty() || (int)fusion.scaleBuf.total() == channels) && @@ -92,6 +92,10 @@ inline Ptr makeFusedCudaNode(const PreparedFusion& fusion, const Pt postNode = make_cuda_node(target, stream); } else if (fn == getActivationFunc(ACTIV_MISH)) { postNode = make_cuda_node(target, stream); + } else if (fn == getActivationFunc(ACTIV_HARDSIGMOID)) { + const float alpha = fusion.activationParams.size() > 0 ? fusion.activationParams[0] : 0.2f; + const float beta = fusion.activationParams.size() > 1 ? fusion.activationParams[1] : 0.5f; + postNode = make_cuda_node(target, stream, alpha, beta); } else { const float slope = fusion.activationParams.empty() ? 0.f : fusion.activationParams[0]; postNode = make_cuda_node(target, stream, slope); diff --git a/modules/dnn/src/layers/matmulnbits_layer.cpp b/modules/dnn/src/layers/matmulnbits_layer.cpp index aef6747b79..b835f5d436 100644 --- a/modules/dnn/src/layers/matmulnbits_layer.cpp +++ b/modules/dnn/src/layers/matmulnbits_layer.cpp @@ -5,6 +5,12 @@ // Third party copyrights are property of their respective owners. #include "../precomp.hpp" +#include "../net_impl.hpp" +#include "../op_cuda.hpp" +#ifdef HAVE_CUDA +#include "../cuda4dnn/primitives/inner_product.hpp" +using namespace cv::dnn::cuda4dnn; +#endif #include #include @@ -33,6 +39,15 @@ class MatMulNBitsLayerImpl CV_FINAL : public MatMulNBitsLayer { } virtual bool supportBackend(int backendId) CV_OVERRIDE { +#ifdef HAVE_CUDA + if (backendId == DNN_BACKEND_CUDA) { + Net::Impl* netimpl_ = getNetImpl(this); + if (!netimpl_ || inputs.size() != 1 || blobs.size() != 2) + return false; + const int t = netimpl_->argData(inputs[0]).type; + return t < 0 || t == CV_32F; + } +#endif return backendId == DNN_BACKEND_OPENCV; } @@ -168,6 +183,41 @@ class MatMulNBitsLayerImpl CV_FINAL : public MatMulNBitsLayer { default: runImpl<8>(A, B, scales, Y); break; } } + +#ifdef HAVE_CUDA + template + Mat dequantWeights() const { + const int n_blk = (K + block_size - 1) / block_size; + const int blob_bytes = block_size * BITS / 8; + const uchar* bq = blobs[0].ptr(); + const float* sc = blobs[1].ptr(); + Mat W(N, K, CV_32F); + parallel_for_(Range(0, N), [&](const Range& r) { + AutoBuffer wcol(n_blk * block_size); + for (int n = r.start; n < r.end; n++) { + for (int blk = 0; blk < n_blk; blk++) { + const size_t b = static_cast(n) * n_blk + blk; + dequantBlock(bq + b * blob_bytes, blob_bytes, sc[b], wcol.data() + blk * block_size); + } + std::copy(wcol.data(), wcol.data() + K, W.ptr(n)); + } + }); + return W; + } + + Ptr initCUDA(void* context_, InputArrayOfArrays inputs_, InputArrayOfArrays) CV_OVERRIDE { + auto context = reinterpret_cast(context_); + std::vector inps; + inputs_.getUMatVector(inps); + if (dequantW.empty()) + dequantW = bits == 4 ? dequantWeights<4>() : dequantWeights<8>(); + const std::size_t axis = inps[0].dims - 1; + return make_cuda_node(preferableTarget, std::move(context->stream), + std::move(context->cublas_handle), axis, dequantW, Mat()); + } + + Mat dequantW; +#endif }; Ptr MatMulNBitsLayer::create(const LayerParams& params) diff --git a/modules/dnn/src/layers/range_layer.cpp b/modules/dnn/src/layers/range_layer.cpp index 161a312666..20eb570972 100644 --- a/modules/dnn/src/layers/range_layer.cpp +++ b/modules/dnn/src/layers/range_layer.cpp @@ -82,9 +82,6 @@ public: Net::Impl* netimpl_ = getNetImpl(this); if (!netimpl_ || this->inputs.size() != 3) return false; - // a dynamic op on CUDA gets inputs 1.. downloaded before sizing; start must already be on the host - if (!netimpl_->isConstArg(this->inputs[0])) - return false; // an intermediate arg has no type until the first forward; refuse only a known unsupported one const int t = netimpl_->argData(this->inputs[0]).type; return t < 0 || isCudaNodeType(t, false); diff --git a/modules/dnn/src/net_impl2.cpp b/modules/dnn/src/net_impl2.cpp index 7792a5db46..19ba076882 100644 --- a/modules/dnn/src/net_impl2.cpp +++ b/modules/dnn/src/net_impl2.cpp @@ -600,6 +600,13 @@ void Net::Impl::prepareForInference() } } +#ifdef HAVE_CUDA +static bool sizedFromHostInput(const std::string& type, size_t j, size_t ninp) +{ + return j > 0 || ninp == 1 || type == "Range"; +} +#endif + void Net::Impl::finalizeGraph(const Ptr& graph, bool useCUDA) { GraphImpl* g = static_cast(graph.get()); @@ -678,7 +685,7 @@ void Net::Impl::finalizeGraph(const Ptr& graph, bool useCUDA) const bool sizedOnHost = op->dynamicOutputShapes() && hostSizedDynamicOps.count(op->type); for (size_t j = 0; j < ninp; j++) { Arg inp = op->inputs[j]; - if (inp.empty() || !hostArg[inp.idx] || (sizedOnHost && (j > 0 || ninp == 1))) + if (inp.empty() || !hostArg[inp.idx] || (sizedOnHost && sizedFromHostInput(op->type, j, ninp))) continue; unmarkHostOp((size_t)producer[inp.idx]); changed = true; @@ -1858,7 +1865,7 @@ void Net::Impl::forwardGraph(Ptr& graph, InputArrayOfArrays inputs_, outOrigData.clear(); // a dynamic op on CUDA is sized from its value inputs, read on the host: all but the - // first, or the only one of a single-input op such as ConstantOfShape + // first, the only one of a single-input op such as ConstantOfShape, or all of Range's const bool dynamicOnCuda = opBackend == DNN_BACKEND_CUDA && op->dynamicOutputShapes(); for (i = 0; i < ninputs; i++) { @@ -1870,7 +1877,7 @@ void Net::Impl::forwardGraph(Ptr& graph, InputArrayOfArrays inputs_, // A host Shape op reads only the header, which the host shadow already carries. const bool headerOnly = opidx < gimpl->hostShapeOp_.size() && gimpl->hostShapeOp_[opidx] && op->type == "Shape"; - if (!headerOnly && (opBackend != DNN_BACKEND_CUDA || (dynamicOnCuda && (i > 0 || ninputs == 1)))) + if (!headerOnly && (opBackend != DNN_BACKEND_CUDA || (dynamicOnCuda && sizedFromHostInput(op->type, i, ninputs)))) cudaDownloadArg(inp, argTensor(inp)); #endif const Mat& m = argTensor(inp); diff --git a/modules/dnn/src/onnx/onnx_importer2.cpp b/modules/dnn/src/onnx/onnx_importer2.cpp index de2a48ebbd..176bce821f 100644 --- a/modules/dnn/src/onnx/onnx_importer2.cpp +++ b/modules/dnn/src/onnx/onnx_importer2.cpp @@ -2446,9 +2446,11 @@ void ONNXImporter2::parseUpsample(LayerParams& layerParams, const opencv_onnx::N layerParams.set("zoom_factor_y", scales.at(2)); layerParams.set("zoom_factor_x", scales.at(3)); n_inputs = 1; - } else if (layerParams.get("mode", "nearest") != "nearest") { - // runtime scales, or a 4-D tensor whose H and W the output takes (the simplifier's - // resize-to-reference subgraph); Resize2 handles both, the legacy Resize neither on CUDA + } else if (layerParams.get("mode", "nearest") != "nearest" || + (!layerParams.get("align_corners", false) && + !layerParams.get("half_pixel_centers", false))) { + // runtime scales go to Resize2, which can stay on CUDA; nearest only in its plain + // asymmetric form, since that is the one Resize2 can match exactly layerParams.type = "Resize2"; // keep the legacy layer's sampling, which only the two flags selected if (!layerParams.has("coordinate_transformation_mode")) { @@ -2458,6 +2460,9 @@ void ONNXImporter2::parseUpsample(LayerParams& layerParams, const opencv_onnx::N halfPixel ? "half_pixel" : layerParams.get("align_corners", false) ? "align_corners" : "asymmetric"); } + // the legacy nearest path (cv::resize INTER_NEAREST) rounds down + if (layerParams.get("mode", "nearest") == "nearest" && !layerParams.has("nearest_mode")) + layerParams.set("nearest_mode", "floor"); } } replaceLayerParam(layerParams, "mode", "interpolation"); diff --git a/modules/dnn/src/op_cuda.cpp b/modules/dnn/src/op_cuda.cpp index 8f65d50c80..aa41678908 100644 --- a/modules/dnn/src/op_cuda.cpp +++ b/modules/dnn/src/op_cuda.cpp @@ -163,6 +163,8 @@ void registerCudaCommonExecs() CV_DNN_REGISTER_EXEC_CLASS(DepthToSpace, DNN_BACKEND_CUDA, CUDALegacyExec); CV_DNN_REGISTER_EXEC_CLASS(SpaceToDepth, DNN_BACKEND_CUDA, CUDALegacyExec); CV_DNN_REGISTER_EXEC_CLASS(NormalizeBBox, DNN_BACKEND_CUDA, CUDALegacyExec); + CV_DNN_REGISTER_EXEC_CLASS(Normalize, DNN_BACKEND_CUDA, CUDALegacyExec); + CV_DNN_REGISTER_EXEC_CLASS(MatMulNBits, DNN_BACKEND_CUDA, CUDALegacyExec); CV_DNN_REGISTER_EXEC_CLASS(Pad2, DNN_BACKEND_CUDA, CUDALegacyExec); CV_DNN_REGISTER_EXEC_CLASS(Clip, DNN_BACKEND_CUDA, CUDALegacyExec); CV_DNN_REGISTER_EXEC_CLASS(ConvTranspose2, DNN_BACKEND_CUDA, CUDALegacyExec); diff --git a/modules/dnn/test/test_onnx_conformance.cpp b/modules/dnn/test/test_onnx_conformance.cpp index d0c301ab86..047768ca00 100644 --- a/modules/dnn/test/test_onnx_conformance.cpp +++ b/modules/dnn/test/test_onnx_conformance.cpp @@ -1897,6 +1897,18 @@ public: bool checkFallbacks(Net& net) const { + // the check below reads the classic engine's layer list; for the new engine, + // ask which ops it actually placed on the CPU + if (net.getMainGraph()) + { + if (backend != DNN_BACKEND_CUDA) + return false; + const std::vector cpuOps = cudaFallbackOps(net); + for (const std::string& op : cpuOps) + std::cout << "FALLBACK: " << op << " runs on the CPU instead of CUDA" << endl; + return !cpuOps.empty(); + } + // Check if all the layers are supported with current backend and target. // Some layers might be fused so their timings equal to zero. std::vector timings; diff --git a/modules/dnn/test/test_onnx_conformance_layer_filter__cuda_denylist.inl.hpp b/modules/dnn/test/test_onnx_conformance_layer_filter__cuda_denylist.inl.hpp index 86b145cb30..2f58604f13 100644 --- a/modules/dnn/test/test_onnx_conformance_layer_filter__cuda_denylist.inl.hpp +++ b/modules/dnn/test/test_onnx_conformance_layer_filter__cuda_denylist.inl.hpp @@ -6,6 +6,16 @@ "test_max_float64", "test_min_float64", "test_mod_mixed_sign_float64", +// New-engine CUDA: supportBackend() turns down Scatter/ScatterND reductions and blocked Q/DQ, +// so the whole graph would run on the CPU. +"test_scatter_elements_with_reduction_max", +"test_scatter_elements_with_reduction_min", +"test_scatternd_add", +"test_scatternd_max", +"test_scatternd_min", +"test_scatternd_multiply", +"test_dequantizelinear_blocked", +"test_quantizelinear_blocked", // New-engine CUDA: arithmetic ops lack integer (int16/uint16/uint32/uint64) support and // standalone BatchNorm with non-const params; denylisted for now (pass on CPU, wrong on CUDA). "test_add_int16",