dnn(cuda): full CUDA coverage for SwinIR, BiRefNet and Depth-Anything

modules/dnn/src/cuda/activations.cu

@@ -347,6 +347,11 @@ template void reciprocal<__half>(const Stream&, Span<__half>, View<__half>);
 
 template void relu<float>(const Stream&, Span<float>, View<float>, float);
 template void clipped_relu<float>(const Stream&, Span<float>, View<float>, float, float);
+// integer Clip (e.g. index arithmetic clamped to a valid range)
+template void clipped_relu<int8_t>(const Stream&, Span<int8_t>, View<int8_t>, int8_t, int8_t);
+template void clipped_relu<uint8_t>(const Stream&, Span<uint8_t>, View<uint8_t>, uint8_t, uint8_t);
+template void clipped_relu<int32_t>(const Stream&, Span<int32_t>, View<int32_t>, int32_t, int32_t);
+template void clipped_relu<int64_t>(const Stream&, Span<int64_t>, View<int64_t>, int64_t, int64_t);
 template void tanh<float>(const Stream&, Span<float>, View<float>);
 template void swish<float>(const Stream&, Span<float>, View<float>);
 template void mish<float>(const Stream&, Span<float>, View<float>);

modules/dnn/src/cuda/cum_scan.cu

@@ -61,13 +61,15 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
                 const std::int64_t target_offset = outer_idx * outer_step_length;
                 const std::int64_t first = target_offset + target_start * inner_size + inner_idx;
 
-                dst[first] = exclusive ? static_cast<T>(is_prod ? W(1) : W(0)) : src[first];
+                // the running value stays in W, so fp16 really accumulates in fp32
+                W acc = exclusive ? (is_prod ? W(1) : W(0)) : static_cast<W>(src[first]);
+                dst[first] = static_cast<T>(acc);
 
                 for (std::int64_t t = target_start + target_delta; t != target_stop; t += target_delta) {
                     const std::int64_t off = target_offset + t * inner_size + inner_idx;
-                    const W prev = static_cast<W>(dst[off - target_step]);
-                    const W cur  = static_cast<W>(src[off - exclusive_delta]);
-                    dst[off] = static_cast<T>(is_prod ? (prev * cur) : (prev + cur));
+                    const W cur = static_cast<W>(src[off - exclusive_delta]);
+                    acc = static_cast<W>(is_prod ? acc * cur : acc + cur);
+                    dst[off] = static_cast<T>(acc);
                 }
             }
         }

modules/dnn/src/cuda/trilu.cu

@@ -30,7 +30,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
 
         template <class T>
         __global__ void trilu(Span<T> output, View<T> input,
-                              int height, int width, int m, int k, bool upper)
+                              int height, int width, int k, bool upper)
         {
             const T* in = input.data().get();
             T* out = output.data().get();
@@ -39,19 +39,8 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
                 const int c = id % width;
                 const int l = (id / width) % height;
 
-                bool zero = false;
-                /* Rows at or past min(height, width) are untouched by the CPU path. */
-                if (l < m) {
-                    if (upper) {
-                        // zeroes columns [0, min(l + k - 1, width - 1)]
-                        const int cmax = device::min<int>(l + k - 1, width - 1);
-                        zero = (c <= cmax);
-                    } else {
-                        // zeroes columns [max(l + k + 1, 0), width - 1]
-                        const int cmin = device::max<int>(l + k + 1, 0);
-                        zero = (c >= cmin);
-                    }
-                }
+                // ONNX keeps c - l >= k (upper) or c - l <= k (lower) on every row, tall matrices included
+                const bool zero = upper ? (c - l < k) : (c - l > k);
 
                 out[id] = zero ? static_cast<T>(0.f) : in[id];
             }
@@ -69,11 +58,9 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
         if (output.size() == 0)
             return;
 
-        const int m = std::min(height, width);
-
         auto kernel = raw::trilu<T>;
         auto policy = make_policy(kernel, output.size(), 0, stream);
-        launch_kernel(kernel, policy, output, input, height, width, m, k, upper);
+        launch_kernel(kernel, policy, output, input, height, width, k, upper);
     }
 
 #if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)

modules/dnn/src/cuda4dnn/kernels/trilu.hpp

@@ -11,12 +11,7 @@
 
 namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
 
-/* ONNX Trilu over the trailing [height, width] of each of `loops` matrices.
- *
- * Mirrors trilu_layer.cpp, including one quirk worth stating: the CPU only rewrites
- * rows below min(height, width) and leaves any row past that as a straight copy. This
- * reproduces that rather than zeroing the full triangle.
- */
+// ONNX Trilu over the trailing [height, width] of each of `loops` matrices
 template <class T>
 void trilu(const csl::Stream& stream,
     csl::Span<T> output, csl::View<T> input,

modules/dnn/src/cuda4dnn/primitives/activation.hpp

@@ -57,8 +57,8 @@ namespace cv { namespace dnn { namespace cuda4dnn {
     template <class T>
     class ClipOp final : public CUDABackendNode {
     public:
-        ClipOp(csl::Stream stream_, float min_, float max_)
-            : stream(std::move(stream_)), min{ static_cast<T>(min_) }, max{ static_cast<T>(max_) } { }
+        ClipOp(csl::Stream stream_, double min_, double max_)
+            : stream(std::move(stream_)), min{ toT(min_) }, max{ toT(max_) } { }
 
         void forward(
             const std::vector<UMat>& inputs,
@@ -73,6 +73,13 @@ namespace cv { namespace dnn { namespace cuda4dnn {
         }
 
     private:
+        // integers convert straight from double so int64 bounds keep their precision; half goes through float
+        static T toT(double v)
+        {
+            using U = typename std::conditional<std::is_integral<T>::value, double, float>::type;
+            return static_cast<T>(static_cast<U>(v));
+        }
+
         csl::Stream stream;
         const T min, max;
     };

modules/dnn/src/cuda4dnn/primitives/conv_transpose2.hpp

@@ -64,18 +64,30 @@ namespace cv { namespace dnn { namespace cuda4dnn {
             scratch_mem_in_bytes = builder.required_workspace_size();
         }
 
+        // constant weights and bias the graph packed into the layer, uploaded once and owned by the node
+        ConvTranspose2Op(csl::Stream stream_, csl::cublas::Handle handle, const ConvTranspose2Configuration& config,
+                         const Mat& weights, const Mat& bias)
+            : ConvTranspose2Op(std::move(stream_), std::move(handle), config)
+        {
+            ownedWeights = upload(weights);
+            if (!bias.empty())
+                ownedBias = upload(bias);
+        }
+
         void forward(
             const std::vector<UMat>& inputs,
             const std::vector<UMat>& outputs,
             csl::Workspace& workspace) override
         {
-            CV_Assert(inputs.size() >= 2 && outputs.size() == 1);
+            const bool owned = !ownedWeights.empty();
+            CV_Assert(inputs.size() >= (owned ? 1u : 2u) && outputs.size() == 1);
 
             auto x = csl::viewOf<T>(inputs[0]);
-            auto w = csl::viewOf<T>(inputs[1]);
+            auto w = csl::viewOf<T>(owned ? ownedWeights : inputs[1]);
             csl::View<T> bias;
-            if (inputs.size() >= 3 && !inputs[2].empty()) {
-                auto b = csl::viewOf<T>(inputs[2]);
+            const UMat& biasMat = owned ? ownedBias : (inputs.size() >= 3 ? inputs[2] : UMat());
+            if (!biasMat.empty()) {
+                auto b = csl::viewOf<T>(biasMat);
                 bias = csl::View<T>(b.get(), b.size());
             }
             auto y = csl::spanOf<T>(outputs[0]);
@@ -111,6 +123,20 @@ namespace cv { namespace dnn { namespace cuda4dnn {
     private:
         std::size_t colsSize() const { return (std::size_t)batch * cout_g * kernel_h * kernel_w * in_h * in_w; }
 
+        // FP32 host data -> device buffer in the node's element type
+        static UMat upload(const Mat& host32)
+        {
+            Mat host;
+            host32.convertTo(host, std::is_same<T, half>::value ? CV_16F : CV_32F);
+            UMat u;
+            u.allocator = cv::cuda::getCudaAllocator();
+            u.fitSameSize(host, host.type());
+            checkCudaAllocated(u);
+            host.copyTo(u);
+            return u;
+        }
+
+        UMat ownedWeights, ownedBias;
         csl::Stream stream;
         csl::cublas::Handle cublasHandle;
         int groups, batch, cin, cin_g, cout, cout_g;

modules/dnn/src/cuda4dnn/primitives/topk.hpp

@@ -9,7 +9,7 @@
 #include "../../op_cuda.hpp"
 
 #include "../csl/stream.hpp"
-#include "../csl/memory.hpp"
+#include "../csl/workspace.hpp"
 
 #include "../kernels/topk.hpp"
 
@@ -25,55 +25,49 @@ namespace cv { namespace dnn { namespace cuda4dnn {
     template <class T>
     class TopKOp final : public CUDABackendNode {
     public:
-        TopKOp(csl::Stream stream_, int axis_, bool largest_)
-            : stream(std::move(stream_)), axis(axis_), largest(largest_) { }
+        TopKOp(csl::Stream stream_, int axis_, bool largest_, const MatShape& inShape, int k_)
+            : stream(std::move(stream_)), largest(largest_), k(k_)
+        {
+            const int a = normalize_axis(axis_, inShape.dims);
+            CV_Assert(a >= 0 && a < inShape.dims);
+            for (int i = 0; i < a; i++) outer *= inShape[i];
+            for (int i = a + 1; i < inShape.dims; i++) inner *= inShape[i];
+            dim_axis = inShape[a];
+            // the sort path's scratch comes from the shared workspace, sized once per shape
+            if (inner == 1 && static_cast<std::int64_t>(k) * dim_axis > sort_threshold)
+                sort_bytes = kernels::topk_sort_workspace<T>(stream, outer, dim_axis, largest);
+        }
 
         void forward(
             const std::vector<UMat>& inputs,
             const std::vector<UMat>& outputs,
             csl::Workspace& workspace) override
         {
-            CV_UNUSED(workspace);
             CV_Assert(inputs.size() >= 1 && outputs.size() == 2);
 
             auto input = csl::viewOf<T>(inputs[0]);
             auto values = csl::spanOf<T>(outputs[0]);
             auto indices = csl::spanOf<std::int64_t>(outputs[1]);
 
-            MatShape inShape = cv::dnn::shape(inputs[0]);
-            MatShape outShape = cv::dnn::shape(outputs[0]);
-            const int rank = inShape.dims;
-            const int a = normalize_axis(axis, rank);
-            CV_Assert(a >= 0 && a < rank);
-
-            int outer = 1, inner = 1;
-            for (int i = 0; i < a; i++) outer *= inShape[i];
-            for (int i = a + 1; i < rank; i++) inner *= inShape[i];
-
-            const int dim_axis = inShape[a];
-            /* K is whatever shape inference settled on; reading it back from the output
-             * avoids duplicating the attribute-vs-input resolution the layer already did. */
-            const int k = outShape[a];
-
-            if (inner == 1 && static_cast<std::int64_t>(k) * dim_axis > sort_threshold) {
-                const std::size_t bytes = kernels::topk_sort_workspace<T>(stream, outer, dim_axis, largest);
-                if (scratch.size() < bytes)
-                    scratch.reset(bytes);
+            if (sort_bytes) {
+                auto scratch = csl::WorkspaceAllocator(workspace).get_span<unsigned char>(sort_bytes);
                 kernels::topk_sort<T>(stream, values, indices, input, outer, dim_axis, k, largest,
-                                      scratch.get().get(), bytes);
+                                      scratch.data().get(), sort_bytes);
                 return;
             }
 
             kernels::topk<T>(stream, values, indices, input, outer, dim_axis, inner, k, largest);
         }
 
+        std::size_t get_workspace_memory_in_bytes() const noexcept override { return sort_bytes; }
+
     private:
         static constexpr std::int64_t sort_threshold = std::int64_t(1) << 16;
 
         csl::Stream stream;
-        int axis;
         bool largest;
-        csl::ManagedPtr<unsigned char> scratch;
+        int k, outer = 1, inner = 1, dim_axis = 1;
+        std::size_t sort_bytes = 0;
     };
 
 }}} /* namespace cv::dnn::cuda4dnn */

modules/dnn/src/layers/clip_layer.cpp

@@ -128,22 +128,25 @@ public:
             for (size_t i = 1; i < inputs.size(); i++)
                 if (!inputs[i].empty() && !netimpl_->isConstArg(inputs[i]))
                     return false;
-            // clipped_relu is float/half only; an intermediate arg has no type until the first forward
+            // an intermediate arg has no type until the first forward; refuse only a known unsupported one
             const int t = netimpl_->argData(inputs[0]).type;
-            return t < 0 || t == CV_32F || t == CV_16F;
+            return t < 0 || isCudaNodeType(t, false);
         }
 #endif
         return backendId == DNN_BACKEND_OPENCV;
     }
 
 #ifdef HAVE_CUDA
-    Ptr<BackendNode> initCUDA(void* context_, InputArrayOfArrays, InputArrayOfArrays) CV_OVERRIDE
+    Ptr<BackendNode> initCUDA(void* context_, InputArrayOfArrays inputs_arr, InputArrayOfArrays) CV_OVERRIDE
     {
         auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
         Net::Impl* netimpl_ = getNetImpl(this);
         CV_Assert(netimpl_);
+        const int type = inputs_arr.type(0);
+        // a half device tensor holds FP32 data, so it takes the float range
+        const int depth = CV_MAT_DEPTH(type) == CV_16F ? CV_32F : CV_MAT_DEPTH(type);
 
-        // same bounds as forward(): const inputs win over attributes, then the float range
+        // same bounds as forward(): const inputs win over attributes, then the type's range
         auto constScalar = [&](size_t i, double& v) {
             if (inputs.size() <= i || inputs[i].empty())
                 return false;
@@ -153,11 +156,11 @@ public:
             v = tmp.at<double>(0);
             return true;
         };
-        double lo = hasMin ? minValue : -FLT_MAX, hi = hasMax ? maxValue : FLT_MAX;
+        double lo = hasMin ? minValue : typeMin(depth), hi = hasMax ? maxValue : typeMax(depth);
         constScalar(1, lo);
         constScalar(2, hi);
         CV_Assert(lo <= hi);
-        return make_cuda_node<cuda4dnn::ClipOp>(preferableTarget, std::move(context->stream), (float)lo, (float)hi);
+        return make_cuda_node_with_type<cuda4dnn::ClipOp>(preferableTarget, type, std::move(context->stream), lo, hi);
     }
 #endif
 

modules/dnn/src/layers/conv2_layer.cpp

@@ -16,6 +16,7 @@
 #include <opencv2/dnn/layer.details.hpp>  // CV_DNN_REGISTER_EXEC_CLASS
 #ifdef HAVE_CUDA
 #include "../cuda4dnn/primitives/convolution.hpp"
+#include "fusion_cuda.hpp"  // FusedPostOpNode and the activation ops
 using namespace cv::dnn::cuda4dnn;
 #endif
 
@@ -554,11 +555,11 @@ public:
                fn == getActivationFunc(ACTIV_MISH) || fn == getActivationFunc(ACTIV_TANH);
     }
 
-    //! On CUDA an activation ConvolutionOp lacks must stay its own op, or the conv loses its CUDA exec.
-    bool targetsCuda() const
+    //! Kernel activations the CUDA path runs as a separate op after the convolution.
+    static bool cudaPostKernel(ActivationFunc fn)
     {
-        Net::Impl* netimpl_ = getNetImpl(this);
-        return netimpl_ && netimpl_->preferableBackend == DNN_BACKEND_CUDA;
+        return fn == getActivationFunc(ACTIV_HARDSWISH) || fn == getActivationFunc(ACTIV_GELU) ||
+               fn == getActivationFunc(ACTIV_GELU_APPROX);
     }
 
     //! Takes an expression the layer has a kernel for. Only a one-layer chain carries
@@ -569,8 +570,6 @@ public:
             return false;
         if (fastActivation != FAST_ACTIV_NONE || activationFunc != nullptr || !activ.empty())
             return false;
-        if (targetsCuda() && !cudaConvKernel(expr->kernel.fn))
-            return false;
 
         activationFunc = expr->kernel.fn;
         activParams.assign(expr->kernel.params, expr->kernel.params + expr->kernel.nparams);
@@ -622,9 +621,6 @@ public:
         int activOperand = -1;
         if (splitActivation(*expr, cur, K, activOperand, act, ap))
             cur = activOperand;
-        // ConvolutionOp has no per-channel PReLU
-        if (act == FAST_ACTIV_PRELU && targetsCuda())
-            return false;
         if (act == FAST_ACTIV_CLIP)
             ap[1] *= postScale;
 
@@ -1218,11 +1214,12 @@ public:
             return false;
         if (!activ.empty())
             return false;
-        // a fused kernel activation runs on CUDA only if ConvolutionOp has the same one
-        if (activationFunc != nullptr && !cudaConvKernel(activationFunc))
+        // a fused kernel activation runs on CUDA inside ConvolutionOp or as a post-op after it
+        if (activationFunc != nullptr && !cudaConvKernel(activationFunc) && !cudaPostKernel(activationFunc))
             return false;
         if (fastActivation != FAST_ACTIV_NONE && fastActivation != FAST_ACTIV_RELU &&
-            fastActivation != FAST_ACTIV_LEAKY_RELU && fastActivation != FAST_ACTIV_CLIP)
+            fastActivation != FAST_ACTIV_LEAKY_RELU && fastActivation != FAST_ACTIV_CLIP &&
+            fastActivation != FAST_ACTIV_PRELU)
             return false;
         return true;
     }
@@ -1273,11 +1270,24 @@ public:
             biasMat = fusedBias;  // already b*scale + bn_bias
         }
 
+        // activations ConvolutionOp lacks run as their own op on the convolution's output
+        Ptr<BackendNode> post;
+        if (fastActivation == FAST_ACTIV_PRELU) {
+            Mat slope(1, (int)activParams.size(), CV_32F, activParams.data());
+            post = make_cuda_node<cuda4dnn::ChannelwiseReLUOp>(targetId, context.stream, slope.clone());
+        } else if (activationFunc == getActivationFunc(ACTIV_HARDSWISH)) {
+            post = make_cuda_node<cuda4dnn::HardSwishOp>(targetId, context.stream);
+        } else if (activationFunc == getActivationFunc(ACTIV_GELU)) {
+            post = make_cuda_node<cuda4dnn::GeluOp>(targetId, context.stream);
+        } else if (activationFunc == getActivationFunc(ACTIV_GELU_APPROX)) {
+            post = make_cuda_node<cuda4dnn::GeluApproximationOp>(targetId, context.stream);
+        }
+
         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 || activationFunc != nullptr;
+        bool hasAct = (fastActivation != FAST_ACTIV_NONE || activationFunc != nullptr) && post.empty();
         if (activationFunc == getActivationFunc(ACTIV_SIGMOID)) {
             config.activation_type = ConvolutionConfiguration::ActivationType::SIGMOID;
         } else if (activationFunc == getActivationFunc(ACTIV_SWISH)) {
@@ -1306,9 +1316,12 @@ public:
         else
             config.fusion_mode = ConvolutionConfiguration::FusionMode::NONE;
 
-        return make_cuda_node<cuda4dnn::ConvolutionOp>(
+        Ptr<BackendNode> convNode = make_cuda_node<cuda4dnn::ConvolutionOp>(
             targetId, std::move(context.stream), std::move(context.cudnn_handle),
             config, filters, biasMat);
+        if (post.empty())
+            return convNode;
+        return makePtr<FusedPostOpNode>(convNode.dynamicCast<CUDABackendNode>(), post.dynamicCast<CUDABackendNode>());
     }
 #endif
 

modules/dnn/src/layers/convtranspose_layer.cpp

@@ -131,14 +131,15 @@ public:
 #ifdef HAVE_CUDA
         if (backendId == DNN_BACKEND_CUDA) {
             Net::Impl* netimpl_ = getNetImpl(this);
-            // ConvTranspose2Op reads weights and bias as device inputs
-            if (!netimpl_ || inputs.size() < 2 || inputs.size() > 3)
+            // weights and bias come as device inputs, or packed into the layer (constant ones)
+            const bool packed = inputs.size() == 1 && !origWeights.empty();
+            if (!netimpl_ || (!packed && (inputs.size() < 2 || inputs.size() > 3)))
                 return false;
             // 2-D and FP32 only; the FP16 target keeps the CPU path
             if (netimpl_->preferableTarget == DNN_TARGET_CUDA_FP16)
                 return false;
             const int xdims = netimpl_->argData(inputs[0]).shape.dims;
-            const int wdims = netimpl_->argData(inputs[1]).shape.dims;
+            const int wdims = packed ? origWeights.dims : netimpl_->argData(inputs[1]).shape.dims;
             if ((xdims >= 0 && xdims != 4) || (wdims >= 0 && wdims != 4))
                 return false;
             // an intermediate arg has no type until the first forward; refuse only a known unsupported one
@@ -159,9 +160,10 @@ public:
         auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
         const int nsd = 2;
 
+        const bool packed = inputs_arr.total() == 1;
         cuda4dnn::ConvTranspose2Configuration config;
         config.input_shape = inputs_arr.shape(0);
-        config.weight_shape = inputs_arr.shape(1);
+        config.weight_shape = packed ? origWeights.shape() : inputs_arr.shape(1);
         config.output_shape = outputs_arr.shape(0);
         config.groups = ngroups;
         // same pads as the CPU path; output_padding only grows the end, which the output shape carries
@@ -171,6 +173,9 @@ public:
             config.dilations.push_back(dilations.empty() ? 1 : dilations[i]);
             config.pads_begin.push_back(pads_resolved.empty() ? 0 : pads_resolved[i]);
         }
+        if (packed)
+            return make_cuda_node<cuda4dnn::ConvTranspose2Op>(preferableTarget, std::move(context->stream),
+                                                              std::move(context->cublas_handle), config, origWeights, bias);
         return make_cuda_node<cuda4dnn::ConvTranspose2Op>(preferableTarget, std::move(context->stream),
                                                           std::move(context->cublas_handle), config);
     }
@@ -190,6 +195,7 @@ public:
             rawWeights.convertTo(wfloat, CV_32F);
         else
             wfloat = rawWeights;
+        origWeights = wfloat.clone();  // the unpacked [Cin, Cout/g, kh, kw] filter the CUDA path reads
         packGemmWeights(wfloat, C0);
         weights.release();
         if (gemmWeights.empty())
@@ -656,6 +662,7 @@ public:
     std::vector<int> emptyKernelShape;
     std::vector<int> explicit_out_shape;
     Mat weights, bias;
+    Mat origWeights;  // unpacked FP32 filter from setWeights(), kept for the CUDA path
     Mat gemmWeights;  // MLAS-packed B per (group, channel chunk), empty without MLAS
     size_t gemmChunkStep = 0;
     int gemmCg = 0, gemmKg = 0, gemmKgp = 0, gemmKC = 0, gemmChunks = 0, gemmKsize = 0, gemmC0 = 0;

modules/dnn/src/layers/topk2_layer.cpp

@@ -84,11 +84,14 @@ public:
 #ifdef HAVE_CUDA
     Ptr<BackendNode> initCUDA(void* context_,
                               InputArrayOfArrays inputs_arr,
-                              InputArrayOfArrays) CV_OVERRIDE
+                              InputArrayOfArrays outputs_arr) CV_OVERRIDE
     {
         auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
+        const MatShape inShape = inputs_arr.shape(0);
+        // K as shape inference resolved it (attribute or constant input)
+        const int k = outputs_arr.shape(0)[normalize_axis(axis, inShape.dims)];
         return make_cuda_node_with_type<cuda4dnn::TopKOp>(preferableTarget, inputs_arr.type(0),
-                                                          std::move(context->stream), axis, largest);
+                                                          std::move(context->stream), axis, largest, inShape, k);
     }
 #endif
 

modules/dnn/src/layers/trilu_layer.cpp

@@ -80,7 +80,6 @@ class TriluLayerImpl CV_FINAL : public TriluLayer {
             const int w = inputs[0].size[shape_input.size() - 1];
             const int h = shape_input.size() >= 2 ? inputs[0].size[shape_input.size() - 2] : 1;
 
-            const int m = std::min(h,w);
             int loops = 1;
             for (int i = 0; i < cdims; ++i)
                 loops *= shape_input[i];
@@ -89,7 +88,7 @@ class TriluLayerImpl CV_FINAL : public TriluLayer {
             uchar *dst = outputs[0].ptr<uchar>();
             auto fn = [&](const Range &r) {
                 for (int i = r.start; i < r.end; i++) {
-                    for(int l=0; l < m; l+=1) {
+                    for(int l=0; l < h; l+=1) {
                         int cmin = upperTri ? 0 : (l + k + 1);
                         cmin = std::max(cmin, 0);
                         const int cmax = upperTri ? min(l + k -1, w-1) : w-1;

modules/dnn/src/net_impl2.cpp

@@ -1687,8 +1687,12 @@ static void forwardOpCUDA(Net::Impl* netimpl, GraphImpl* gimpl, size_t opidx,
     }
     std::vector<UMat> inpG(inputs.size()), outG(outputs.size());
     for (size_t i = 0; i < inputs.size(); i++) {
-        if (inpMats[i].empty())
+        if (inpMats[i].empty()) {
+            // a zero-element tensor still has a shape (Shape reads it); this header allocates nothing
+            if (inpMats[i].dims > 0 && inpMats[i].total() == 0)
+                inpG[i] = UMat(inpMats[i].dims, inpMats[i].size.p, inpMats[i].type());
             continue;
+        }
         netimpl->cudaUploadArg(inputs[i], inpMats[i]);   // H2D only if host-authoritative
         inpG[i] = netimpl->cudaDeviceBuffer(inputs[i]);
     }

modules/dnn/test/test_common.hpp

@@ -15,7 +15,13 @@
 #ifdef HAVE_CUDA
 #include "opencv2/core/cuda.hpp"
 #endif
-#include "../src/dnn_common.hpp"  // cudaFallbackOps()
+
+namespace cv { namespace dnn {
+CV__DNN_INLINE_NS_BEGIN
+// defined in src/net_impl2.cpp and declared in src/dnn_common.hpp, which tests cannot include
+CV_EXPORTS std::vector<std::string> cudaFallbackOps(const Net& net);
+CV__DNN_INLINE_NS_END
+}}
 
 // src/op_inf_engine.hpp
 #define INF_ENGINE_VER_MAJOR_GT(ver) (((INF_ENGINE_RELEASE) / 10000) > ((ver) / 10000))

modules/dnn/test/test_layers.cpp

@@ -2088,6 +2088,41 @@ INSTANTIATE_TEST_CASE_P(/**/, Layer_Test_GlobalPoolingBlock, testing::Combine(
     testing::Bool()),
     GlobalPoolingParamName());
 
+// rows 3 and 4 of a 5x3 matrix lie past min(h, w); ONNX Trilu still masks them
+TEST(Layer_Test_Trilu, tall_matrix)
+{
+    const int sz[] = {2, 5, 3};
+    Mat inp(3, sz, CV_32F);
+    randu(inp, 1, 10);
+
+    for (int upper = 0; upper <= 1; upper++)
+    {
+        for (int k = -2; k <= 1; k++)
+        {
+            LayerParams lp;
+            lp.name = "testTrilu";
+            lp.type = "Trilu";
+            lp.set("upper", upper == 1);
+            Ptr<Layer> layer = LayerFactory::createLayerInstance(lp.type, lp);
+
+            Mat kMat(1, 1, CV_64S);
+            kMat.at<int64>(0, 0) = k;
+            std::vector<Mat> inputs = {inp, kMat}, outputs(1, Mat(3, sz, CV_32F));
+            layer->forward(inputs, outputs, std::vector<Mat>());
+
+            for (int n = 0; n < sz[0]; n++)
+                for (int l = 0; l < sz[1]; l++)
+                    for (int c = 0; c < sz[2]; c++)
+                    {
+                        const bool keep = upper ? (c - l >= k) : (c - l <= k);
+                        const float expected = keep ? inp.ptr<float>(n, l)[c] : 0.f;
+                        EXPECT_EQ(expected, outputs[0].ptr<float>(n, l)[c])
+                            << "upper=" << upper << " k=" << k << " n=" << n << " l=" << l << " c=" << c;
+                    }
+        }
+    }
+}
+
 typedef testing::TestWithParam<tuple<Vec4i, int, tuple<Backend, Target> > > Layer_Test_ShuffleChannel;
 TEST_P(Layer_Test_ShuffleChannel, Accuracy)
 {

modules/dnn/test/test_onnx_importer.cpp

@@ -1155,6 +1155,8 @@ TEST_P(Test_ONNX_layers, Colorization)
         CV_TEST_TAG_DEBUG_VERYLONG
     );
     checkBackend();
+    if (backend == DNN_BACKEND_CUDA)
+        applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA);  // TF32 deconvolution exceeds the FP32 reference bounds
 
     const String model = findDataFile("dnn/colorization_deploy_v2.onnx", false);
     Net net = readNetFromONNX(model);