Compare commits

...

5 Commits

Author SHA1 Message Date
sanjay
ad3e9194db refactored suggested changes 2026-11-06 12:56:53 +05:30
Savya Sanchi Sharma
f75b1307f9 dnn(cuda): full CUDA coverage for SwinIR, BiRefNet and Depth-Anything 2026-11-04 19:16:17 +05:30
Savya Sanchi Sharma
42f6c98d26 dnn(cuda): Einsum, DropoutMask, DynamicQuantizeLinear and perf fixes 2026-11-04 19:04:57 +05:30
Savya Sanchi Sharma
b4ee7a3366 dnn(cuda): RF-DETR and remaining layer kernels 2026-11-04 19:04:57 +05:30
Savya Sanchi Sharma
ab5554df13 dnn(cuda): CUDA execs for new engine layers and dynamic-shape ops 2026-11-04 19:04:57 +05:30
127 changed files with 6496 additions and 429 deletions

View File

@@ -28,6 +28,10 @@ public:
{
backend = (dnn::Backend)(int)get<0>(GetParam());
target = (dnn::Target)(int)get<1>(GetParam());
// OPENCV_DNN_PERF_PIN_INPUTS=1 pins every test's inputs, to compare with the default unpinned run
static const bool pinAll = utils::getConfigurationParameterBool("OPENCV_DNN_PERF_PIN_INPUTS", false);
if (pinAll)
pinnedInput = "*";
}
void processNet(std::string weights, std::string proto,
@@ -36,10 +40,14 @@ public:
if (!proto.empty())
proto = findDataFile(proto);
net = readNet(weights, proto);
// "*" pins every named input to the shape it is fed
if (pinnedInput == "*")
for (auto& inp : inputs)
net.setInputShape(std::get<1>(inp), shape(std::get<0>(inp)));
// "*" pins every input to the shape it is fed; an unnamed one is the graph input at its position
if (pinnedInput == "*" && net.getMainGraph()) {
const std::vector<Arg>& graphInputs = net.getMainGraph()->inputs();
for (size_t i = 0; i < inputs.size() && i < graphInputs.size(); i++) {
const std::string& name = std::get<1>(inputs[i]);
net.setInputShape(name.empty() ? net.argName(graphInputs[i]) : name, shape(std::get<0>(inputs[i])));
}
}
else if (!pinnedInput.empty())
net.setInputShape(pinnedInput, shape(std::get<0>(inputs[0])));
// Set multiple inputs

View File

@@ -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>);

View File

@@ -0,0 +1,100 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include "types.hpp"
#include "grid_stride_range.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/kernels/argmax.hpp"
#include <cstdint>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
template <class T, bool IsArgMax, bool SelectLast>
__global__ void arg_min_max(View<T> input, Span<std::int64_t> output, size_type num_outputs, int axis_size, int inner_size) {
for (auto id : grid_stride_range(num_outputs)) {
int outer_index = id / inner_size;
int inner_index = id % inner_size;
std::int64_t base = (static_cast<std::int64_t>(outer_index) * axis_size) * inner_size + inner_index;
const T* p_input = input.data().get();
T best_value = p_input[base];
std::int64_t best_index = 0;
for (int k = 1; k < axis_size; ++k) {
const T value = p_input[base + static_cast<std::int64_t>(k) * inner_size];
if (IsArgMax) {
if (SelectLast ? value >= best_value : value > best_value) {
best_value = value;
best_index = k;
}
} else {
if (SelectLast ? value <= best_value : value < best_value) {
best_value = value;
best_index = k;
}
}
}
output[id] = best_index;
}
}
}
template <class T>
void arg_min_max(const Stream& stream,
Span<std::int64_t> output, View<T> input,
int outer_size, int axis_size, int inner_size, bool is_argmax, bool select_last_index)
{
if (axis_size <= 0 || outer_size <= 0 || inner_size <= 0)
return;
const size_type num_outputs = static_cast<size_type>(outer_size) * inner_size;
if (num_outputs == 0)
return;
if (is_argmax) {
if (select_last_index) {
auto kernel = raw::arg_min_max<T, true, true>;
auto policy = make_policy(kernel, num_outputs, 0, stream);
launch_kernel(kernel, policy, input, output, num_outputs, axis_size, inner_size);
} else {
auto kernel = raw::arg_min_max<T, true, false>;
auto policy = make_policy(kernel, num_outputs, 0, stream);
launch_kernel(kernel, policy, input, output, num_outputs, axis_size, inner_size);
}
} else {
if (select_last_index) {
auto kernel = raw::arg_min_max<T, false, true>;
auto policy = make_policy(kernel, num_outputs, 0, stream);
launch_kernel(kernel, policy, input, output, num_outputs, axis_size, inner_size);
} else {
auto kernel = raw::arg_min_max<T, false, false>;
auto policy = make_policy(kernel, num_outputs, 0, stream);
launch_kernel(kernel, policy, input, output, num_outputs, axis_size, inner_size);
}
}
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void arg_min_max(const Stream&, Span<std::int64_t>, View<__half>, int, int, int, bool, bool);
#endif
template void arg_min_max(const Stream&, Span<std::int64_t>, View<float>, int, int, int, bool, bool);
template void arg_min_max(const Stream&, Span<std::int64_t>, View<int8_t>, int, int, int, bool, bool);
template void arg_min_max(const Stream&, Span<std::int64_t>, View<uint8_t>, int, int, int, bool, bool);
template void arg_min_max(const Stream&, Span<std::int64_t>, View<int32_t>, int, int, int, bool, bool);
template void arg_min_max(const Stream&, Span<std::int64_t>, View<int64_t>, int, int, int, bool, bool);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -53,6 +53,44 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
__device__ static std::int64_t apply(float v) { return __float2ll_rz(v); }
};
// 8-bit outputs saturate like the CPU's saturate_cast; integer sources need no rounding
template <class TIn>
struct Converter<std::uint8_t, TIn> {
__device__ static std::uint8_t apply(TIn v) {
const std::int64_t x = static_cast<std::int64_t>(v);
return static_cast<std::uint8_t>(x < 0 ? 0 : (x > 255 ? 255 : x));
}
};
template <class TIn>
struct Converter<std::int8_t, TIn> {
__device__ static std::int8_t apply(TIn v) {
const std::int64_t x = static_cast<std::int64_t>(v);
return static_cast<std::int8_t>(x < -128 ? -128 : (x > 127 ? 127 : x));
}
};
// float truncates toward zero first, as the CPU Cast does; NaN saturates like saturate_cast(INT_MIN)
template <>
struct Converter<std::uint8_t, float> {
__device__ static std::uint8_t apply(float v) {
if (isnan(v))
return 0;
const float t = truncf(v);
return static_cast<std::uint8_t>(t < 0.f ? 0.f : (t > 255.f ? 255.f : t));
}
};
template <>
struct Converter<std::int8_t, float> {
__device__ static std::int8_t apply(float v) {
if (isnan(v))
return -128;
const float t = truncf(v);
return static_cast<std::int8_t>(t < -128.f ? -128.f : (t > 127.f ? 127.f : t));
}
};
template <class TOut, class TIn>
__global__ void cast(Span<TOut> output, View<TIn> input) {
for (auto i : grid_stride_range(output.size()))
@@ -87,4 +125,19 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
template void cast(const Stream&, Span<float>, View<std::int64_t>);
template void cast(const Stream&, Span<float>, View<float>);
#define CV_CUDA_CAST_8BIT(T8) \
template void cast(const Stream&, Span<T8>, View<bool>); \
template void cast(const Stream&, Span<T8>, View<std::int32_t>); \
template void cast(const Stream&, Span<T8>, View<std::int64_t>); \
template void cast(const Stream&, Span<T8>, View<float>); \
template void cast(const Stream&, Span<T8>, View<std::uint8_t>); \
template void cast(const Stream&, Span<T8>, View<std::int8_t>); \
template void cast(const Stream&, Span<bool>, View<T8>); \
template void cast(const Stream&, Span<std::int32_t>, View<T8>); \
template void cast(const Stream&, Span<std::int64_t>, View<T8>); \
template void cast(const Stream&, Span<float>, View<T8>);
CV_CUDA_CAST_8BIT(std::uint8_t)
CV_CUDA_CAST_8BIT(std::int8_t)
#undef CV_CUDA_CAST_8BIT
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -0,0 +1,79 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include "types.hpp"
#include "grid_stride_range.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/kernels/col2im.hpp"
#include <opencv2/core.hpp>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
template <class T>
__global__ void col2im_2d(Span<T> output, View<T> cols, View<T> bias, Col2ImParams p, size_type count)
{
const index_type in_plane = p.in_h * p.in_w;
const index_type out_plane = p.out_h * p.out_w;
for (auto i : grid_stride_range(count)) {
const index_type ox = i % p.out_w;
const index_type oy = (i / p.out_w) % p.out_h;
const index_type c = (i / out_plane) % p.channels;
const index_type n = i / (out_plane * p.channels);
float acc = bias.size() ? static_cast<float>(bias[p.channel_offset + c]) : 0.f;
const index_type batch_base = n * p.channels * p.kernel_h * p.kernel_w * in_plane;
for (index_type ky = 0; ky < p.kernel_h; ky++) {
// input row iy feeds oy when iy * stride == oy + pad - ky * dilation
const index_type ty = oy + p.pad_h - ky * p.dil_h;
if (ty < 0 || ty % p.stride_h != 0)
continue;
const index_type iy = ty / p.stride_h;
if (iy >= p.in_h)
continue;
for (index_type kx = 0; kx < p.kernel_w; kx++) {
const index_type tx = ox + p.pad_w - kx * p.dil_w;
if (tx < 0 || tx % p.stride_w != 0)
continue;
const index_type ix = tx / p.stride_w;
if (ix >= p.in_w)
continue;
const index_type row = (c * p.kernel_h + ky) * p.kernel_w + kx;
acc += static_cast<float>(cols[batch_base + row * in_plane + iy * p.in_w + ix]);
}
}
output[(n * p.total_channels + p.channel_offset + c) * out_plane + oy * p.out_w + ox] = static_cast<T>(acc);
}
}
}
template <class T>
void col2im_2d(const Stream& stream, Span<T> output, View<T> cols, View<T> bias, const Col2ImParams& params)
{
const size_type count = params.batch * params.channels * params.out_h * params.out_w;
if (count == 0)
return;
auto kernel = raw::col2im_2d<T>;
auto policy = make_policy(kernel, count, 0, stream);
launch_kernel(kernel, policy, output, cols, bias, params, count);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void col2im_2d(const Stream&, Span<__half>, View<__half>, View<__half>, const Col2ImParams&);
#endif
template void col2im_2d(const Stream&, Span<float>, View<float>, View<float>, const Col2ImParams&);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -0,0 +1,108 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include "types.hpp"
#include "grid_stride_range.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/kernels/cum_scan.hpp"
#include <opencv2/core.hpp>
#include <cstddef>
#include <cstdint>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
/* fp16 accumulates in fp32; integral types must accumulate in themselves, since
* routing an int64 scan through float would silently lose precision. */
template <class T> struct work_type { using type = T; };
template <> struct work_type<__half> { using type = float; };
template <class T>
__global__ void cum_scan(Span<T> output, View<T> input,
std::int64_t outer_size, std::int64_t target_size, std::int64_t inner_size,
bool is_prod, bool exclusive, bool reverse)
{
using W = typename work_type<T>::type;
const T* src = input.data().get();
T* dst = output.data().get();
const std::int64_t outer_step_length = target_size * inner_size;
/* Same index arithmetic as the CPU scan; target_step is signed so that the
* reverse direction needs no separate code path. */
const std::int64_t target_start = reverse ? target_size - 1 : 0;
const std::int64_t target_stop = reverse ? -1 : target_size;
const std::int64_t target_delta = reverse ? -1 : 1;
const std::int64_t target_step = target_delta * inner_size;
const std::int64_t exclusive_delta = exclusive ? target_step : 0;
const std::int64_t nrows = outer_size * inner_size;
for (auto id : grid_stride_range(static_cast<device::size_type>(nrows))) {
const std::int64_t outer_idx = static_cast<std::int64_t>(id) / inner_size;
const std::int64_t inner_idx = static_cast<std::int64_t>(id) - outer_idx * inner_size;
const std::int64_t target_offset = outer_idx * outer_step_length;
const std::int64_t first = target_offset + target_start * inner_size + inner_idx;
// 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 cur = static_cast<W>(src[off - exclusive_delta]);
acc = static_cast<W>(is_prod ? acc * cur : acc + cur);
dst[off] = static_cast<T>(acc);
}
}
}
}
template <class T>
void cum_scan(const Stream& stream,
Span<T> output, View<T> input,
std::size_t outer_size, std::size_t target_size, std::size_t inner_size,
bool is_prod, bool exclusive, bool reverse)
{
CV_Assert(outer_size > 0 && target_size > 0 && inner_size > 0);
const std::size_t nrows = outer_size * inner_size;
if (nrows == 0)
return;
auto kernel = raw::cum_scan<T>;
auto policy = make_policy(kernel, nrows, 0, stream);
launch_kernel(kernel, policy, output, input,
static_cast<std::int64_t>(outer_size),
static_cast<std::int64_t>(target_size),
static_cast<std::int64_t>(inner_size),
is_prod, exclusive, reverse);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void cum_scan(const Stream&, Span<__half>, View<__half>, std::size_t, std::size_t, std::size_t, bool, bool, bool);
#endif
template void cum_scan(const Stream&, Span<float>, View<float>, std::size_t, std::size_t, std::size_t, bool, bool, bool);
template void cum_scan(const Stream&, Span<int32_t>, View<int32_t>, std::size_t, std::size_t, std::size_t, bool, bool, bool);
template void cum_scan(const Stream&, Span<int64_t>, View<int64_t>, std::size_t, std::size_t, std::size_t, bool, bool, bool);
template void cum_scan(const Stream&, Span<int8_t>, View<int8_t>, std::size_t, std::size_t, std::size_t, bool, bool, bool);
template void cum_scan(const Stream&, Span<uint8_t>, View<uint8_t>, std::size_t, std::size_t, std::size_t, bool, bool, bool);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -355,6 +355,26 @@ void eltwise_pow_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x,
eltwise_op<T, PowFunctor<T>>(stream, output, x, y);
}
template <class T>
void eltwise_prelu_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x, TensorView<T> slope) {
eltwise_op<T, PReluFunctor<T>>(stream, output, x, slope);
}
template <class T>
void eltwise_bitand_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x, TensorView<T> y) {
eltwise_op<T, BitAndFunctor<T>>(stream, output, x, y);
}
template <class T>
void eltwise_bitor_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x, TensorView<T> y) {
eltwise_op<T, BitOrFunctor<T>>(stream, output, x, y);
}
template <class T>
void eltwise_bitxor_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x, TensorView<T> y) {
eltwise_op<T, BitXorFunctor<T>>(stream, output, x, y);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void eltwise_mod_2(const Stream& stream, TensorSpan<__half> output, TensorView<__half> x, TensorView<__half> y);
template void eltwise_fmod_2(const Stream& stream, TensorSpan<__half> output, TensorView<__half> x, TensorView<__half> y);
@@ -422,5 +442,26 @@ void eltwise_pow_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x,
template void eltwise_min_2(const Stream& stream, TensorSpan<int64_t> output, TensorView<int64_t> x, TensorView<int64_t> y);
template void eltwise_pow_2(const Stream& stream, TensorSpan<int64_t> output, TensorView<int64_t> x, TensorView<int64_t> y);
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void eltwise_prelu_2(const Stream&, TensorSpan<__half>, TensorView<__half>, TensorView<__half>);
#endif
template void eltwise_prelu_2(const Stream&, TensorSpan<float>, TensorView<float>, TensorView<float>);
template void eltwise_prelu_2(const Stream&, TensorSpan<int8_t>, TensorView<int8_t>, TensorView<int8_t>);
template void eltwise_prelu_2(const Stream&, TensorSpan<uint8_t>, TensorView<uint8_t>, TensorView<uint8_t>);
template void eltwise_prelu_2(const Stream&, TensorSpan<int32_t>, TensorView<int32_t>, TensorView<int32_t>);
template void eltwise_prelu_2(const Stream&, TensorSpan<int64_t>, TensorView<int64_t>, TensorView<int64_t>);
template void eltwise_bitand_2(const Stream&, TensorSpan<int8_t>, TensorView<int8_t>, TensorView<int8_t>);
template void eltwise_bitand_2(const Stream&, TensorSpan<uint8_t>, TensorView<uint8_t>, TensorView<uint8_t>);
template void eltwise_bitand_2(const Stream&, TensorSpan<int32_t>, TensorView<int32_t>, TensorView<int32_t>);
template void eltwise_bitand_2(const Stream&, TensorSpan<int64_t>, TensorView<int64_t>, TensorView<int64_t>);
template void eltwise_bitor_2(const Stream&, TensorSpan<int8_t>, TensorView<int8_t>, TensorView<int8_t>);
template void eltwise_bitor_2(const Stream&, TensorSpan<uint8_t>, TensorView<uint8_t>, TensorView<uint8_t>);
template void eltwise_bitor_2(const Stream&, TensorSpan<int32_t>, TensorView<int32_t>, TensorView<int32_t>);
template void eltwise_bitor_2(const Stream&, TensorSpan<int64_t>, TensorView<int64_t>, TensorView<int64_t>);
template void eltwise_bitxor_2(const Stream&, TensorSpan<int8_t>, TensorView<int8_t>, TensorView<int8_t>);
template void eltwise_bitxor_2(const Stream&, TensorSpan<uint8_t>, TensorView<uint8_t>, TensorView<uint8_t>);
template void eltwise_bitxor_2(const Stream&, TensorSpan<int32_t>, TensorView<int32_t>, TensorView<int32_t>);
template void eltwise_bitxor_2(const Stream&, TensorSpan<int64_t>, TensorView<int64_t>, TensorView<int64_t>);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -9,6 +9,8 @@
#include "math.hpp"
#include <type_traits>
#include "../cuda4dnn/csl/nvcc_defs.hpp"
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
@@ -851,6 +853,63 @@ struct FModFunctor {
}
};
// ONNX PRelu with a broadcast slope: y is the slope
template <class T>
struct PReluFunctor {
struct Params {
CUDA4DNN_HOST_DEVICE Params() {}
};
CUDA4DNN_DEVICE PReluFunctor() { }
CUDA4DNN_DEVICE PReluFunctor(const Params& params) { }
CUDA4DNN_DEVICE T operator()(T x, T y) {
return apply(x, y, std::is_unsigned<T>{});
}
// 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
template <class T>
struct BitAndFunctor {
struct Params {
CUDA4DNN_HOST_DEVICE Params() {}
};
CUDA4DNN_DEVICE BitAndFunctor() { }
CUDA4DNN_DEVICE BitAndFunctor(const Params& params) { }
CUDA4DNN_DEVICE T operator()(T x, T y) { return x & y; }
};
template <class T>
struct BitOrFunctor {
struct Params {
CUDA4DNN_HOST_DEVICE Params() {}
};
CUDA4DNN_DEVICE BitOrFunctor() { }
CUDA4DNN_DEVICE BitOrFunctor(const Params& params) { }
CUDA4DNN_DEVICE T operator()(T x, T y) { return x | y; }
};
template <class T>
struct BitXorFunctor {
struct Params {
CUDA4DNN_HOST_DEVICE Params() {}
};
CUDA4DNN_DEVICE BitXorFunctor() { }
CUDA4DNN_DEVICE BitXorFunctor(const Params& params) { }
CUDA4DNN_DEVICE T operator()(T x, T y) { return x ^ y; }
};
template <class T>
struct PowFunctor {
struct Params {

View File

@@ -0,0 +1,90 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include "types.hpp"
#include "grid_stride_range.hpp"
#include "execution.hpp"
#include "fast_divmod.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/kernels/gather_elements.hpp"
#include <cstddef>
#include <cstdint>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
__host__ __device__ inline std::int64_t get_element_index_value(const void* index_data, std::size_t index_element_size, std::size_t offset) {
switch (index_element_size) {
case sizeof(std::int32_t):
return *(reinterpret_cast<const std::int32_t*>(index_data) + offset);
case sizeof(std::int64_t):
return *(reinterpret_cast<const std::int64_t*>(index_data) + offset);
default:
break;
}
return 0;
}
template <class T>
__global__ void gather_elements(
Span<T> output, View<T> input,
const void* indices, std::size_t index_element_size,
std::int64_t axis_size_data, std::int64_t inner_size,
fast_divmod axis_inner_size, fast_divmod inner_size_divmod)
{
for (auto id : grid_stride_range(output.size())) {
int outer_index, axis_inner_offset;
axis_inner_size.divmod(id, outer_index, axis_inner_offset);
int axis_index, inner_index;
inner_size_divmod.divmod(axis_inner_offset, axis_index, inner_index);
(void)axis_index;
std::int64_t idx = get_element_index_value(indices, index_element_size, id);
idx = idx < 0 ? idx + axis_size_data : idx;
if (idx < 0 || idx >= axis_size_data) {
output[id] = static_cast<T>(0);
continue;
}
std::int64_t input_index = static_cast<std::int64_t>(outer_index) * axis_size_data * inner_size +
idx * inner_size + inner_index;
output[id] = input.data().get()[input_index];
}
}
}
template <class T>
void gather_elements(const Stream& stream,
Span<T> output, View<T> input,
const void* indices, std::size_t index_element_size,
std::int64_t axis_size_data, std::int64_t inner_size,
fast_divmod axis_inner_size, fast_divmod inner_size_divmod)
{
auto kernel = raw::gather_elements<T>;
auto policy = make_policy(kernel, output.size(), 0, stream);
launch_kernel(kernel, policy, output, input, indices, index_element_size,
axis_size_data, inner_size, axis_inner_size, inner_size_divmod);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void gather_elements(const Stream&, Span<__half>, View<__half>, const void*, std::size_t, std::int64_t, std::int64_t, fast_divmod, fast_divmod);
#endif
template void gather_elements(const Stream&, Span<float>, View<float>, const void*, std::size_t, std::int64_t, std::int64_t, fast_divmod, fast_divmod);
template void gather_elements(const Stream&, Span<int8_t>, View<int8_t>, const void*, std::size_t, std::int64_t, std::int64_t, fast_divmod, fast_divmod);
template void gather_elements(const Stream&, Span<uint8_t>, View<uint8_t>, const void*, std::size_t, std::int64_t, std::int64_t, fast_divmod, fast_divmod);
template void gather_elements(const Stream&, Span<int32_t>, View<int32_t>, const void*, std::size_t, std::int64_t, std::int64_t, fast_divmod, fast_divmod);
template void gather_elements(const Stream&, Span<int64_t>, View<int64_t>, const void*, std::size_t, std::int64_t, std::int64_t, fast_divmod, fast_divmod);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -0,0 +1,118 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include "array.hpp"
#include "types.hpp"
#include "grid_stride_range.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/csl/tensor.hpp"
#include "../cuda4dnn/kernels/gather_nd.hpp"
#include <opencv2/core.hpp>
#include <cstddef>
#include <cstdint>
#include <vector>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
template <class T, class TIdx>
__global__ void gather_nd(
Span<T> output, View<TIdx> indices, View<T> input,
std::int64_t last_index_dimension,
array<std::int64_t, CSL_MAX_TENSOR_RANK> element_counts,
array<std::int64_t, CSL_MAX_TENSOR_RANK> input_dims,
std::int64_t inner_size,
std::int64_t batch_stride,
std::int64_t num_slices_per_batch)
{
for (auto id : grid_stride_range(output.size())) {
const std::int64_t slice = static_cast<std::int64_t>(id) / inner_size;
const std::int64_t inner = static_cast<std::int64_t>(id) - slice * inner_size;
std::int64_t data_offset = 0;
const std::int64_t indices_start = last_index_dimension * slice;
for (std::int64_t j = 0; j < last_index_dimension; ++j) {
std::int64_t index = static_cast<std::int64_t>(indices.data().get()[indices_start + j]);
const std::int64_t dim_value = input_dims[j];
if (index >= 0) {
if (index >= dim_value)
index = dim_value - 1;
} else {
if (index < -dim_value)
index = 0;
else
index += dim_value;
}
data_offset += index * element_counts[j];
}
if (batch_stride > 0)
data_offset += (slice / num_slices_per_batch) * batch_stride;
output[id] = input.data().get()[data_offset + inner];
}
}
}
template <class T, class TIdx>
void gather_nd(const Stream& stream,
Span<T> output, View<TIdx> indices, View<T> input,
std::int64_t last_index_dimension,
const std::vector<std::int64_t>& element_counts,
const std::vector<std::int64_t>& input_dims,
std::int64_t inner_size,
std::int64_t batch_stride,
std::int64_t num_slices_per_batch)
{
CV_Assert(last_index_dimension <= CSL_MAX_TENSOR_RANK);
CV_Assert(element_counts.size() == static_cast<std::size_t>(last_index_dimension));
CV_Assert(input_dims.size() == static_cast<std::size_t>(last_index_dimension));
CV_Assert(inner_size > 0);
CV_Assert(num_slices_per_batch > 0);
if (output.size() == 0)
return;
array<std::int64_t, CSL_MAX_TENSOR_RANK> element_counts_k, input_dims_k;
element_counts_k.assign(std::begin(element_counts), std::end(element_counts));
input_dims_k.assign(std::begin(input_dims), std::end(input_dims));
auto kernel = raw::gather_nd<T, TIdx>;
auto policy = make_policy(kernel, output.size(), 0, stream);
launch_kernel(kernel, policy, output, indices, input,
last_index_dimension, element_counts_k, input_dims_k,
inner_size, batch_stride, num_slices_per_batch);
}
#define CV_GATHER_ND_INST(T) \
template void gather_nd(const Stream&, Span<T>, View<std::int32_t>, View<T>, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::int64_t, std::int64_t, std::int64_t); \
template void gather_nd(const Stream&, Span<T>, View<std::int64_t>, View<T>, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::int64_t, std::int64_t, std::int64_t);
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
CV_GATHER_ND_INST(__half)
#endif
CV_GATHER_ND_INST(float)
CV_GATHER_ND_INST(int8_t)
CV_GATHER_ND_INST(uint8_t)
CV_GATHER_ND_INST(int32_t)
CV_GATHER_ND_INST(int64_t)
#undef CV_GATHER_ND_INST
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -0,0 +1,162 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include "types.hpp"
#include "grid_stride_range.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/kernels/grid_sample.hpp"
#include <opencv2/core.hpp>
#include <cstdint>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
enum { M_NEAREST = 0, M_BILINEAR = 1, M_BICUBIC = 2 };
enum { P_ZEROS = 0, P_BORDER = 1, P_REFLECTION = 2 };
__device__ inline float reflect_coord(float x, int limit, bool align_corners) {
if (limit <= 1) return 0.f;
const float minv = align_corners ? 0.f : -0.5f;
const float maxv = align_corners ? float(limit - 1) : float(limit) - 0.5f;
const float m = maxv - minv;
const float two_m = 2.f * m;
float t = fmodf(x - minv, two_m);
if (t < 0) t += two_m;
if (t > m) t = two_m - t;
return t + minv;
}
__device__ inline void cubic_coeffs(float x, float A, float* c) {
c[0] = ((A * (x + 1.0f) - 5.0f * A) * (x + 1.0f) + 8.0f * A) * (x + 1.0f) - 4.0f * A;
c[1] = ((A + 2.0f) * x - (A + 3.0f)) * x * x + 1.0f;
c[2] = ((A + 2.0f) * (1.0f - x) - (A + 3.0f)) * (1.0f - x) * (1.0f - x) + 1.0f;
c[3] = 1.0f - c[0] - c[1] - c[2];
}
/* The padding mode is decided per layer, so this branch is warp-uniform. */
template <class T>
__device__ inline float fetch(const T* baseNC, int yy, int xx, const GridSampleParams& p) {
int px = xx, py = yy;
if (p.padding == P_BORDER) {
px = static_cast<int>(fminf(float(p.width - 1), fmaxf(0.f, float(px))));
py = static_cast<int>(fminf(float(p.height - 1), fmaxf(0.f, float(py))));
} else if (p.padding == P_REFLECTION) {
px = static_cast<int>(floorf(reflect_coord(float(px), p.width, p.align_corners) + 0.5f));
py = static_cast<int>(floorf(reflect_coord(float(py), p.height, p.align_corners) + 0.5f));
}
if (px < 0 || py < 0 || px >= p.width || py >= p.height) return 0.f;
return static_cast<float>(baseNC[static_cast<std::int64_t>(py) * p.width + px]);
}
template <class T>
__global__ void grid_sample(Span<T> output, View<T> input, View<float> grid, GridSampleParams p)
{
const float delta = p.align_corners ? 1.f : 0.f;
const float xscale = 0.5f * (p.width - delta);
const float yscale = 0.5f * (p.height - delta);
const float xdelta = 0.5f * (p.width - delta) + 0.5f * (delta - 1.f);
const float ydelta = 0.5f * (p.height - delta) + 0.5f * (delta - 1.f);
for (auto id : grid_stride_range(output.size())) {
int rem = id;
const int w = rem % p.output_width; rem /= p.output_width;
const int h = rem % p.output_height; rem /= p.output_height;
const int c = rem % p.channels; rem /= p.channels;
const int n = rem;
const T* baseNC = input.data().get() +
(static_cast<std::int64_t>(n) * p.channels + c) *
static_cast<std::int64_t>(p.height) * p.width;
const std::int64_t g = ((static_cast<std::int64_t>(n) * p.output_height + h) *
p.output_width + w) * 2;
const float nx = grid.data().get()[g + 0];
const float ny = grid.data().get()[g + 1];
const float xf = nx * xscale + xdelta;
const float yf = ny * yscale + ydelta;
float outv = 0.f;
if (p.mode == M_NEAREST) {
outv = fetch<T>(baseNC, __float2int_rn(yf), __float2int_rn(xf), p);
} else if (p.mode == M_BILINEAR) {
const int x0 = static_cast<int>(floorf(xf));
const int y0 = static_cast<int>(floorf(yf));
const float dx = xf - x0, dy = yf - y0;
const float v00 = fetch<T>(baseNC, y0, x0, p);
const float v01 = fetch<T>(baseNC, y0, x0 + 1, p);
const float v10 = fetch<T>(baseNC, y0 + 1, x0, p);
const float v11 = fetch<T>(baseNC, y0 + 1, x0 + 1, p);
const float vx0 = v00 * (1.f - dx) + v01 * dx;
const float vx1 = v10 * (1.f - dx) + v11 * dx;
outv = vx0 * (1.f - dy) + vx1 * dy;
} else {
const int x1 = static_cast<int>(floorf(xf));
const int y1 = static_cast<int>(floorf(yf));
const float tx = xf - x1, ty = yf - y1;
float wx[4], wy[4];
cubic_coeffs(tx, p.cubic_alpha, wx);
cubic_coeffs(ty, p.cubic_alpha, wy);
outv = 0.f;
for (int j = 0; j < 4; ++j) {
float rowv = 0.f;
for (int i = 0; i < 4; ++i)
rowv += fetch<T>(baseNC, y1 - 1 + j, x1 - 1 + i, p) * wx[i];
outv += rowv * wy[j];
}
}
output[id] = static_cast<T>(outv);
}
}
}
template <class T>
void grid_sample(const Stream& stream,
Span<T> output, View<T> input, View<float> grid,
const GridSampleParams& params)
{
CV_Assert(params.channels > 0 && params.height > 0 && params.width > 0);
CV_Assert(params.output_height > 0 && params.output_width > 0);
CV_Assert(params.mode >= 0 && params.mode <= 2);
CV_Assert(params.padding >= 0 && params.padding <= 2);
if (output.size() == 0)
return;
auto kernel = raw::grid_sample<T>;
auto policy = make_policy(kernel, output.size(), 0, stream);
launch_kernel(kernel, policy, output, input, grid, params);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void grid_sample(const Stream&, Span<__half>, View<__half>, View<float>, const GridSampleParams&);
#endif
template void grid_sample(const Stream&, Span<float>, View<float>, View<float>, const GridSampleParams&);
template void grid_sample(const Stream&, Span<int8_t>, View<int8_t>, View<float>, const GridSampleParams&);
template void grid_sample(const Stream&, Span<uint8_t>, View<uint8_t>, View<float>, const GridSampleParams&);
template void grid_sample(const Stream&, Span<int32_t>, View<int32_t>, View<float>, const GridSampleParams&);
template void grid_sample(const Stream&, Span<int64_t>, View<int64_t>, View<float>, const GridSampleParams&);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -34,21 +34,11 @@ namespace raw {
constexpr int MEAN_SQR_SUM_BLOCK_SIZE = 256;
// one block per row: a warp-shuffle and shared-memory reduction, no per-element atomics
template <class T>
__global__ void reduce_mean_sqr_sum(Span<float> means, Span<float> sum_sqrs, View<T> input, size_type inner_size) {
// sums sum and sqr over the block with warp shuffles and shared memory; the totals land in thread 0
__device__ void block_sum2(float& sum, float& sqr) {
__shared__ float warp_sums[MEAN_SQR_SUM_BLOCK_SIZE / GPU_WARP_SIZE];
__shared__ float warp_sqrs[MEAN_SQR_SUM_BLOCK_SIZE / GPU_WARP_SIZE];
const index_type outer_idx = blockIdx.x;
const index_type row_start = outer_idx * inner_size;
float sum = 0.f, sqr = 0.f;
for (index_type i = threadIdx.x; i < inner_size; i += blockDim.x) {
const auto x = static_cast<float>(input[row_start + i]);
sum += x;
sqr += x * x;
}
for (int offset = GPU_WARP_SIZE / 2; offset > 0; offset /= 2) {
sum += WARP_SHFL_XOR(sum, offset);
sqr += WARP_SHFL_XOR(sqr, offset);
@@ -68,10 +58,59 @@ namespace raw {
sum += WARP_SHFL_XOR(sum, offset);
sqr += WARP_SHFL_XOR(sqr, offset);
}
if (lane == 0) {
means[outer_idx] = sum / inner_size;
sum_sqrs[outer_idx] = sqr;
}
}
// one block per row: a warp-shuffle and shared-memory reduction, no per-element atomics
template <class T>
__global__ void reduce_mean_sqr_sum(Span<float> means, Span<float> sum_sqrs, View<T> input, size_type inner_size) {
const index_type outer_idx = blockIdx.x;
const index_type row_start = outer_idx * inner_size;
float sum = 0.f, sqr = 0.f;
for (index_type i = threadIdx.x; i < inner_size; i += blockDim.x) {
const auto x = static_cast<float>(input[row_start + i]);
sum += x;
sqr += x * x;
}
block_sum2(sum, sqr);
if (threadIdx.x == 0) {
means[outer_idx] = sum / inner_size;
sum_sqrs[outer_idx] = sqr;
}
}
// a long row split over `splits` blocks: each writes the partial sums of its chunk
template <class T>
__global__ void reduce_mean_sqr_sum_partial(Span<float> part_sums, Span<float> part_sqrs, View<T> input,
size_type inner_size, size_type chunk, size_type splits) {
const index_type row = blockIdx.x / splits, split = blockIdx.x % splits;
const index_type row_start = row * inner_size;
const index_type begin = split * chunk;
const index_type end = begin + chunk < inner_size ? begin + chunk : inner_size;
float sum = 0.f, sqr = 0.f;
for (index_type i = begin + threadIdx.x; i < end; i += blockDim.x) {
const auto x = static_cast<float>(input[row_start + i]);
sum += x;
sqr += x * x;
}
block_sum2(sum, sqr);
if (threadIdx.x == 0) {
part_sums[row * splits + split] = sum;
part_sqrs[row * splits + split] = sqr;
}
}
// one thread per row adds the partials in split order, so the result does not depend on scheduling
__global__ void reduce_mean_sqr_sum_combine(Span<float> means, Span<float> sum_sqrs, View<float> part_sums,
View<float> part_sqrs, size_type splits, size_type inner_size) {
for (auto row : grid_stride_range(means.size())) {
float sum = 0.f, sqr = 0.f;
for (index_type s = 0; s < splits; s++) {
sum += part_sums[row * splits + s];
sqr += part_sqrs[row * splits + s];
}
means[row] = sum / inner_size;
sum_sqrs[row] = sqr;
}
}
@@ -181,6 +220,56 @@ template void reduce_mean_sqr_sum(const Stream&, Span<float>, Span<float>, View<
#endif
template void reduce_mean_sqr_sum(const Stream&, Span<float>, Span<float>, View<float>, std::size_t);
namespace {
// 64 loads per thread: enough work per block to hide memory latency, small enough to fill the GPU
constexpr std::size_t MEAN_SQR_SUM_CHUNK = raw::MEAN_SQR_SUM_BLOCK_SIZE * 64;
// a row gains from splitting only once it fills two chunks; short rows (LayerNorm) keep one block
std::size_t meanSqrSumSplits(std::size_t inner_size)
{
if (inner_size < 2 * MEAN_SQR_SUM_CHUNK)
return 1;
return (inner_size + MEAN_SQR_SUM_CHUNK - 1) / MEAN_SQR_SUM_CHUNK;
}
}
std::size_t reduce_mean_sqr_sum_scratch(std::size_t rows, std::size_t inner_size)
{
const std::size_t splits = meanSqrSumSplits(inner_size);
return splits > 1 ? 2 * rows * splits : 0;
}
template <class T>
void reduce_mean_sqr_sum(const Stream& stream, Span<float> means, Span<float> sum_sqrs, View<T> input,
std::size_t inner_size, Span<float> scratch)
{
const std::size_t rows = means.size();
const std::size_t splits = meanSqrSumSplits(inner_size);
if (splits == 1 || scratch.size() < 2 * rows * splits) {
reduce_mean_sqr_sum<T>(stream, means, sum_sqrs, input, inner_size);
return;
}
CV_Assert(input.size() / inner_size == rows && sum_sqrs.size() == rows);
const std::size_t chunk = (inner_size + splits - 1) / splits;
Span<float> part_sums(scratch.data(), rows * splits);
Span<float> part_sqrs(scratch.data() + rows * splits, rows * splits);
auto partial = raw::reduce_mean_sqr_sum_partial<T>;
launch_kernel(partial, execution_policy(dim3(static_cast<unsigned int>(rows * splits)),
dim3(raw::MEAN_SQR_SUM_BLOCK_SIZE), 0, stream),
part_sums, part_sqrs, input, inner_size, chunk, splits);
auto combine = raw::reduce_mean_sqr_sum_combine;
auto policy = make_policy(combine, rows, 0, stream);
launch_kernel(combine, policy, means, sum_sqrs, View<float>(part_sums), View<float>(part_sqrs), splits, inner_size);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void reduce_mean_sqr_sum(const Stream&, Span<float>, Span<float>, View<__half>, std::size_t, Span<float>);
#endif
template void reduce_mean_sqr_sum(const Stream&, Span<float>, Span<float>, View<float>, std::size_t, Span<float>);
void compute_normalization_scale(const Stream& stream, Span<float> scale, View<float> means, View<float> sum_sqrs, std::size_t inner_size, float eps)
{
CV_Assert(scale.size() == means.size());

View File

@@ -41,8 +41,9 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
template <class T>
__global__ void reciprocal(Span<T> output, T epsilon) {
// a zero slice scales to 0; in half, eps underflows to 0 and 1/0 would give 0 * inf = NaN
for (auto idx : grid_stride_range(output.size()))
output[idx] = T(1) / (output[idx] + epsilon);
output[idx] = output[idx] == T(0) ? T(0) : T(1) / (output[idx] + epsilon);
}
template <class T>
@@ -60,7 +61,8 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
__global__ void rsqrt(Span<T> output, T epsilon) {
for (auto idx : grid_stride_range(output.size())) {
using device::sqrt;
output[idx] = T(1) / sqrt(output[idx] + epsilon);
// a zero slice scales to 0; in half, eps underflows to 0 and 1/0 would give 0 * inf = NaN
output[idx] = output[idx] == T(0) ? T(0) : T(1) / sqrt(output[idx] + epsilon);
}
}
@@ -115,9 +117,51 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
launch_kernel(scale_kernel, policy, output, input, mid_size * inner_size, inner_size, sums);
}
namespace raw {
template <class T>
__global__ void lrn_across_channels(Span<T> output, View<T> input, size_type channels, size_type plane_size,
int radius, float alpha_scaled, float bias, float beta)
{
for (auto i : grid_stride_range(output.size())) {
const index_type p = i % plane_size;
const index_type c = (i / plane_size) % channels;
const index_type n = i / (plane_size * channels);
const index_type c0 = c - radius < 0 ? 0 : c - radius;
const index_type c1 = c + radius >= channels ? channels - 1 : c + radius;
// same window as the CPU ChannelLRN, accumulated in float
float s = 0.f;
for (index_type k = c0; k <= c1; k++) {
const float v = static_cast<float>(input[(n * channels + k) * plane_size + p]);
s += v * v;
}
const float x = static_cast<float>(input[i]);
output[i] = static_cast<T>(x * powf(bias + alpha_scaled * s, -beta));
}
}
}
template <class T>
void lrn_across_channels(
const Stream& stream,
Span<T> output, View<T> input,
std::size_t batch, std::size_t channels, std::size_t plane_size,
int radius, float alpha_scaled, float bias, float beta)
{
CV_Assert(output.size() == input.size());
CV_Assert(output.size() == batch * channels * plane_size);
if (output.size() == 0)
return;
auto kernel = raw::lrn_across_channels<T>;
auto policy = make_policy(kernel, output.size(), 0, stream);
launch_kernel(kernel, policy, output, input, channels, plane_size, radius, alpha_scaled, bias, beta);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void normalize(const Stream&, Span<__half>, View<__half>, std::size_t, std::size_t, std::size_t, std::size_t, __half, Span<__half>);
template void lrn_across_channels(const Stream&, Span<__half>, View<__half>, std::size_t, std::size_t, std::size_t, int, float, float, float);
#endif
template void lrn_across_channels(const Stream&, Span<float>, View<float>, std::size_t, std::size_t, std::size_t, int, float, float, float);
template void normalize(const Stream&, Span<float>, View<float>, std::size_t, std::size_t, std::size_t, std::size_t, float, Span<float>);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -0,0 +1,78 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include "types.hpp"
#include "grid_stride_range.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/kernels/one_hot.hpp"
#include <opencv2/core.hpp>
#include <cstddef>
#include <cstdint>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
template <class T, class TIdx>
__global__ void one_hot(Span<T> output, View<TIdx> indices, std::int64_t depth, std::int64_t inner,
T off_value, T on_value)
{
for (auto id : grid_stride_range(output.size())) {
const std::int64_t i = id;
const std::int64_t lo = i % inner;
const std::int64_t outer = i / inner;
const std::int64_t d = outer % depth;
const std::int64_t hi = outer / depth;
// only [-depth, depth-1] is valid; anything else matches no d and stays all off_value
std::int64_t v = static_cast<std::int64_t>(indices.data().get()[hi * inner + lo]);
if (v < 0)
v += depth;
output[id] = (v == d) ? on_value : off_value;
}
}
}
template <class T, class TIdx>
void one_hot(const Stream& stream, Span<T> output, View<TIdx> indices,
std::size_t depth, std::size_t inner, T off_value, T on_value)
{
CV_Assert(depth > 0 && inner > 0);
CV_Assert(output.size() == indices.size() * depth);
if (output.size() == 0)
return;
auto kernel = raw::one_hot<T, TIdx>;
auto policy = make_policy(kernel, output.size(), 0, stream);
launch_kernel(kernel, policy, output, indices, static_cast<std::int64_t>(depth),
static_cast<std::int64_t>(inner), off_value, on_value);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void one_hot(const Stream&, Span<__half>, View<int32_t>, std::size_t, std::size_t, __half, __half);
template void one_hot(const Stream&, Span<__half>, View<int64_t>, std::size_t, std::size_t, __half, __half);
#endif
template void one_hot(const Stream&, Span<float>, View<int32_t>, std::size_t, std::size_t, float, float);
template void one_hot(const Stream&, Span<float>, View<int64_t>, std::size_t, std::size_t, float, float);
template void one_hot(const Stream&, Span<int8_t>, View<int32_t>, std::size_t, std::size_t, int8_t, int8_t);
template void one_hot(const Stream&, Span<int8_t>, View<int64_t>, std::size_t, std::size_t, int8_t, int8_t);
template void one_hot(const Stream&, Span<uint8_t>, View<int32_t>, std::size_t, std::size_t, uint8_t, uint8_t);
template void one_hot(const Stream&, Span<uint8_t>, View<int64_t>, std::size_t, std::size_t, uint8_t, uint8_t);
template void one_hot(const Stream&, Span<int32_t>, View<int32_t>, std::size_t, std::size_t, int32_t, int32_t);
template void one_hot(const Stream&, Span<int32_t>, View<int64_t>, std::size_t, std::size_t, int32_t, int32_t);
template void one_hot(const Stream&, Span<int64_t>, View<int32_t>, std::size_t, std::size_t, int64_t, int64_t);
template void one_hot(const Stream&, Span<int64_t>, View<int64_t>, std::size_t, std::size_t, int64_t, int64_t);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -16,6 +16,9 @@
#include "../cuda4dnn/csl/tensor.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/kernels/fill_copy.hpp"
#include "../cuda4dnn/kernels/padding.hpp"
#include <opencv2/core.hpp>
#include <cstddef>
@@ -64,8 +67,114 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
output[i] = input[iidx];
}
}
template <class T, std::size_t Rank>
__global__ void pad_nd(
Span<T> output, array<size_type, Rank> out_strides,
View<T> input, array<size_type, Rank> in_strides, array<index_type, Rank> in_shape,
array<index_type, Rank> pads_begin, int mode, T value)
{
for (auto i : grid_stride_range(output.size())) {
index_type iidx = 0;
bool outside = false;
for (int j = 0; j < Rank; j++) {
const index_type n = in_shape[j];
const index_type out_j = j == 0 ? i / out_strides[0] : (i % out_strides[j - 1]) / out_strides[j];
index_type p = out_j - pads_begin[j];
if (p < 0 || p >= n) {
if (mode == static_cast<int>(kernels::PadMode::CONSTANT)) {
outside = true;
break;
} else if (mode == static_cast<int>(kernels::PadMode::REPLICATE)) {
p = p < 0 ? 0 : n - 1;
} else if (mode == static_cast<int>(kernels::PadMode::WRAP)) {
p = ((p % n) + n) % n;
} else if (n == 1) {
p = 0;
} else {
// reflect101 may bounce more than once when a pad exceeds the axis
while (p < 0 || p >= n)
p = p < 0 ? -p : 2 * n - 2 - p;
}
}
iidx += p * in_strides[j];
}
output[i] = outside ? value : input[iidx];
}
}
}
template <class T, std::size_t Rank> static
void launch_pad_nd(
const Stream& stream,
Span<T> output, const std::vector<std::size_t>& outStride,
View<T> input, const std::vector<std::size_t>& inStride, const std::vector<std::size_t>& inShape,
const std::vector<int>& padsBegin, int mode, T value)
{
CV_Assert(outStride.size() == Rank && inStride.size() == Rank);
CV_Assert(inShape.size() == Rank && padsBegin.size() == Rank);
array<size_type, Rank> outStride_k, inStride_k;
outStride_k.assign(std::begin(outStride), std::end(outStride));
inStride_k.assign(std::begin(inStride), std::end(inStride));
array<index_type, Rank> inShape_k, padsBegin_k;
for (int i = 0; i < Rank; i++) {
inShape_k[i] = inShape[i];
padsBegin_k[i] = padsBegin[i];
}
auto kernel = raw::pad_nd<T, Rank>;
auto policy = make_policy(kernel, output.size(), 0, stream);
launch_kernel(kernel, policy, output, outStride_k, input, inStride_k, inShape_k, padsBegin_k, mode, value);
}
GENERATE_KERNEL_DISPATCHER(pad_nd_dispatcher, launch_pad_nd);
template <class T>
void pad_nd(
const Stream& stream,
TensorSpan<T> output, TensorView<T> input,
const std::vector<int>& pads_begin, PadMode mode, T value)
{
CV_Assert(output.rank() == input.rank());
CV_Assert(output.rank() == pads_begin.size());
if (output.size() == 0)
return;
// an empty input leaves only constant padding
if (input.size() == 0) {
CV_Assert(mode == PadMode::CONSTANT);
fill<T>(stream, output, value);
return;
}
const auto rank = input.rank();
auto inShape = input.shape_as_vector();
auto outShape = output.shape_as_vector();
std::vector<std::size_t> inStride(rank), outStride(rank);
inStride.back() = 1;
outStride.back() = 1;
for (int i = (int)rank - 2; i >= 0; i--) {
inStride[i] = inStride[i + 1] * inShape[i + 1];
outStride[i] = outStride[i + 1] * outShape[i + 1];
}
CV_Assert(1 <= rank && rank <= CSL_MAX_TENSOR_RANK);
pad_nd_dispatcher<T, 1, CSL_MAX_TENSOR_RANK>(rank, stream, output, outStride, input, inStride, inShape,
pads_begin, static_cast<int>(mode), value);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void pad_nd(const Stream&, TensorSpan<__half>, TensorView<__half>, const std::vector<int>&, PadMode, __half);
#endif
template void pad_nd(const Stream&, TensorSpan<float>, TensorView<float>, const std::vector<int>&, PadMode, float);
template void pad_nd(const Stream&, TensorSpan<int8_t>, TensorView<int8_t>, const std::vector<int>&, PadMode, int8_t);
template void pad_nd(const Stream&, TensorSpan<uint8_t>, TensorView<uint8_t>, const std::vector<int>&, PadMode, uint8_t);
template void pad_nd(const Stream&, TensorSpan<int32_t>, TensorView<int32_t>, const std::vector<int>&, PadMode, int32_t);
template void pad_nd(const Stream&, TensorSpan<int64_t>, TensorView<int64_t>, const std::vector<int>&, PadMode, int64_t);
template void pad_nd(const Stream&, TensorSpan<bool>, TensorView<bool>, const std::vector<int>&, PadMode, bool);
template <class T, std::size_t Rank> static
void launch_copy_with_reflection101(
const Stream& stream,

View File

@@ -0,0 +1,177 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
#include <cuda_runtime.h>
#include <cub/block/block_reduce.cuh>
#include "grid_stride_range.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/kernels/quantize_dequantize.hpp"
#include <opencv2/core.hpp>
#include <algorithm>
#include <cstddef>
#include <cstdint>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
template <class T>
__global__ void dequantize_linear(const T* input, const float* scale, const T* zero_point,
float* output, std::size_t sz_a, std::size_t slice_size, std::size_t n)
{
for (auto i : grid_stride_range(n)) {
std::size_t a_idx = (i / slice_size) % sz_a;
T zp = zero_point ? zero_point[a_idx] : T(0);
output[i] = (static_cast<float>(input[i]) - static_cast<float>(zp)) * scale[a_idx];
}
}
template <class T> __device__ __forceinline__ T saturate_quantized(int v);
template <> __device__ __forceinline__ std::int8_t saturate_quantized<std::int8_t>(int v) {
return static_cast<std::int8_t>(::min(::max(v, -128), 127));
}
template <> __device__ __forceinline__ std::uint8_t saturate_quantized<std::uint8_t>(int v) {
return static_cast<std::uint8_t>(::min(::max(v, 0), 255));
}
template <class T>
__global__ void quantize_linear(const float* input, const float* scale, const T* zero_point,
T* output, std::size_t sz_a, std::size_t slice_size, std::size_t n)
{
for (auto i : grid_stride_range(n)) {
std::size_t a_idx = (i / slice_size) % sz_a;
int zp = zero_point ? static_cast<int>(zero_point[a_idx]) : 0;
// 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<T>(rounded + zp);
}
}
// 256 threads, as the module's other block reductions (rms_norm.cu, topk.cu, reduce.cu, mvn.cu)
constexpr int DYNAMIC_QUANTIZE_BLOCK = 256;
struct MinOp { __device__ float operator()(float a, float b) const { return fminf(a, b); } };
struct MaxOp { __device__ float operator()(float a, float b) const { return fmaxf(a, b); } };
// one (min, max) pair per block; zero starts both, since the range must include it
__global__ void dynamic_quantize_minmax(const float* input, size_type n, float* pmin, float* pmax)
{
using BlockReduce = cub::BlockReduce<float, DYNAMIC_QUANTIZE_BLOCK>;
__shared__ typename BlockReduce::TempStorage storage;
float lo = 0.f, hi = 0.f;
for (size_type i = blockIdx.x * DYNAMIC_QUANTIZE_BLOCK + threadIdx.x; i < n; i += gridDim.x * DYNAMIC_QUANTIZE_BLOCK) {
lo = fminf(lo, input[i]);
hi = fmaxf(hi, input[i]);
}
lo = BlockReduce(storage).Reduce(lo, MinOp());
__syncthreads();
hi = BlockReduce(storage).Reduce(hi, MaxOp());
if (threadIdx.x == 0) {
pmin[blockIdx.x] = lo;
pmax[blockIdx.x] = hi;
}
}
// folds the partials and derives the scale and zero point exactly as the CPU layer does
__global__ void dynamic_quantize_params(const float* pmin, const float* pmax, int nparts,
float* scale, std::uint8_t* zero_point)
{
using BlockReduce = cub::BlockReduce<float, DYNAMIC_QUANTIZE_BLOCK>;
__shared__ typename BlockReduce::TempStorage storage;
float lo = 0.f, hi = 0.f;
for (int i = threadIdx.x; i < nparts; i += DYNAMIC_QUANTIZE_BLOCK) {
lo = fminf(lo, pmin[i]);
hi = fmaxf(hi, pmax[i]);
}
lo = BlockReduce(storage).Reduce(lo, MinOp());
__syncthreads();
hi = BlockReduce(storage).Reduce(hi, MaxOp());
if (threadIdx.x == 0) {
const float qmin = 0.f, qmax = 255.f;
float y_scale = (hi - lo) / (qmax - qmin);
if (y_scale == 0.f)
y_scale = 1.f;
const int zp = __float2int_rn(qmin - lo / y_scale);
scale[0] = y_scale;
zero_point[0] = saturate_quantized<std::uint8_t>(zp);
}
}
}
template <class T>
void dequantize_linear(const Stream& stream, Span<float> output, View<T> input,
View<float> scale, View<T> zero_point,
std::size_t sz_a, std::size_t slice_size)
{
CV_Assert(output.size() == input.size());
CV_Assert(zero_point.empty() || zero_point.size() == scale.size());
auto kernel = raw::dequantize_linear<T>;
auto policy = make_policy(kernel, output.size(), 0, stream);
launch_kernel(kernel, policy, input.data().get(), scale.data().get(),
zero_point.empty() ? nullptr : zero_point.data().get(),
output.data().get(), sz_a, slice_size, output.size());
}
template <class T>
void quantize_linear(const Stream& stream, Span<T> output, View<float> input,
View<float> scale, View<T> zero_point,
std::size_t sz_a, std::size_t slice_size)
{
CV_Assert(output.size() == input.size());
CV_Assert(zero_point.empty() || zero_point.size() == scale.size());
auto kernel = raw::quantize_linear<T>;
auto policy = make_policy(kernel, output.size(), 0, stream);
launch_kernel(kernel, policy, input.data().get(), scale.data().get(),
zero_point.empty() ? nullptr : zero_point.data().get(),
output.data().get(), sz_a, slice_size, output.size());
}
std::size_t dynamic_quantize_partials(std::size_t n)
{
const std::size_t block = raw::DYNAMIC_QUANTIZE_BLOCK;
return std::max<std::size_t>(1, std::min<std::size_t>((n + block - 1) / block, block));
}
void dynamic_quantize_scale(const Stream& stream, Span<float> scale, Span<std::uint8_t> zero_point,
View<float> input, Span<float> scratch)
{
CV_Assert(scale.size() == 1 && zero_point.size() == 1);
const std::size_t nparts = dynamic_quantize_partials(input.size());
CV_Assert(scratch.size() >= 2 * nparts);
float* pmin = scratch.data().get();
float* pmax = pmin + nparts;
auto minmax = raw::dynamic_quantize_minmax;
launch_kernel(minmax, execution_policy(nparts, raw::DYNAMIC_QUANTIZE_BLOCK, stream),
input.data().get(), input.size(), pmin, pmax);
auto params = raw::dynamic_quantize_params;
launch_kernel(params, execution_policy(1, raw::DYNAMIC_QUANTIZE_BLOCK, stream),
pmin, pmax, static_cast<int>(nparts), scale.data().get(), zero_point.data().get());
}
template void dequantize_linear<std::int8_t>(const Stream&, Span<float>, View<std::int8_t>, View<float>, View<std::int8_t>, std::size_t, std::size_t);
template void dequantize_linear<std::uint8_t>(const Stream&, Span<float>, View<std::uint8_t>, View<float>, View<std::uint8_t>, std::size_t, std::size_t);
template void quantize_linear<std::int8_t>(const Stream&, Span<std::int8_t>, View<float>, View<float>, View<std::int8_t>, std::size_t, std::size_t);
template void quantize_linear<std::uint8_t>(const Stream&, Span<std::uint8_t>, View<float>, View<float>, View<std::uint8_t>, std::size_t, std::size_t);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -0,0 +1,72 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include "types.hpp"
#include "grid_stride_range.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/kernels/range.hpp"
#include <cstdint>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
__device__ inline double as_double(float v) { return static_cast<double>(v); }
__device__ inline double as_double(__half v) { return static_cast<double>(__half2float(v)); }
__device__ inline double as_double(std::int8_t v) { return static_cast<double>(v); }
__device__ inline double as_double(std::uint8_t v) { return static_cast<double>(v); }
__device__ inline double as_double(std::int32_t v) { return static_cast<double>(v); }
__device__ inline double as_double(std::int64_t v) { return static_cast<double>(v); }
template <class T>
__device__ inline T range_value(T start, T delta, std::size_t i) {
return static_cast<T>(as_double(start) + static_cast<double>(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<std::int64_t>(i) * delta;
}
template <class T>
__global__ void range(Span<T> output, const T* start_p, const T* delta_p)
{
const T start = *start_p;
const T delta = *delta_p;
for (auto i : grid_stride_range(output.size()))
output[i] = range_value(start, delta, static_cast<std::size_t>(i));
}
}
template <class T>
void range(const Stream& stream, Span<T> output, const T* start, const T* delta)
{
if (output.size() == 0)
return;
auto kernel = raw::range<T>;
auto policy = make_policy(kernel, output.size(), 0, stream);
launch_kernel(kernel, policy, output, start, delta);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void range(const Stream&, Span<__half>, const __half*, const __half*);
#endif
template void range(const Stream&, Span<float>, const float*, const float*);
template void range(const Stream&, Span<int8_t>, const int8_t*, const int8_t*);
template void range(const Stream&, Span<uint8_t>, const uint8_t*, const uint8_t*);
template void range(const Stream&, Span<int32_t>, const int32_t*, const int32_t*);
template void range(const Stream&, Span<int64_t>, const int64_t*, const int64_t*);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -242,4 +242,76 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
#endif
template void resize_bilinear<float>(const Stream&, TensorSpan<float>, TensorView<float>, float, float, bool);
namespace raw {
template <class T>
__global__ void resize_separable(Span<T> output, View<T> input, size_type in_h, size_type in_w,
size_type out_h, size_type out_w,
View<int> idx_y, View<float> w_y, View<int> oob_y, int taps_y,
View<int> idx_x, View<float> w_x, View<int> oob_x, int taps_x,
float extrapolation_value)
{
for (auto i : grid_stride_range(output.size())) {
const index_type ox = i % out_w;
const index_type oy = (i / out_w) % out_h;
const index_type plane = i / (out_w * out_h);
if (oob_y[oy] || oob_x[ox]) {
output[i] = static_cast<T>(extrapolation_value);
continue;
}
const index_type in_plane = plane * in_h * in_w;
float acc = 0.f;
for (int ky = 0; ky < taps_y; ky++) {
const int iy = idx_y[oy * taps_y + ky];
if (iy < 0)
continue;
float row = 0.f;
for (int kx = 0; kx < taps_x; kx++) {
const int ix = idx_x[ox * taps_x + kx];
if (ix < 0)
continue;
row += w_x[ox * taps_x + kx] * static_cast<float>(input[in_plane + iy * in_w + ix]);
}
acc += w_y[oy * taps_y + ky] * row;
}
output[i] = static_cast<T>(acc);
}
}
}
template <class T>
void resize_separable(const Stream& stream, TensorSpan<T> output, TensorView<T> input,
View<int> idx_y, View<float> w_y, View<int> oob_y, int taps_y,
View<int> idx_x, View<float> w_x, View<int> oob_x, int taps_x,
float extrapolation_value)
{
CV_Assert(input.rank() == output.rank() && input.rank() >= 2);
const auto out_h = output.get_axis_size(-2), out_w = output.get_axis_size(-1);
const auto in_h = input.get_axis_size(-2), in_w = input.get_axis_size(-1);
CV_Assert(input.size() / (in_h * in_w) == output.size() / (out_h * out_w));
CV_Assert(idx_y.size() == out_h * taps_y && idx_x.size() == out_w * taps_x);
if (output.size() == 0)
return;
auto kernel = raw::resize_separable<T>;
auto policy = make_policy(kernel, output.size(), 0, stream);
launch_kernel(kernel, policy, Span<T>(output.get(), output.size()), View<T>(input.get(), input.size()),
in_h, in_w, out_h, out_w, idx_y, w_y, oob_y, taps_y, idx_x, w_x, oob_x, taps_x,
extrapolation_value);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void resize_separable<__half>(const Stream&, TensorSpan<__half>, TensorView<__half>,
View<int>, View<float>, View<int>, int, View<int>, View<float>, View<int>, int, float);
#endif
template void resize_separable<float>(const Stream&, TensorSpan<float>, TensorView<float>,
View<int>, View<float>, View<int>, int, View<int>, View<float>, View<int>, int, float);
template void resize_separable<int8_t>(const Stream&, TensorSpan<int8_t>, TensorView<int8_t>,
View<int>, View<float>, View<int>, int, View<int>, View<float>, View<int>, int, float);
template void resize_separable<uint8_t>(const Stream&, TensorSpan<uint8_t>, TensorView<uint8_t>,
View<int>, View<float>, View<int>, int, View<int>, View<float>, View<int>, int, float);
template void resize_separable<int32_t>(const Stream&, TensorSpan<int32_t>, TensorView<int32_t>,
View<int>, View<float>, View<int>, int, View<int>, View<float>, View<int>, int, float);
template void resize_separable<int64_t>(const Stream&, TensorSpan<int64_t>, TensorView<int64_t>,
View<int>, View<float>, View<int>, int, View<int>, View<float>, View<int>, int, float);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -0,0 +1,97 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <cub/block/block_reduce.cuh>
#include "types.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/kernels/rms_norm.hpp"
#include <opencv2/core.hpp>
#include <algorithm>
#include <cstddef>
#include <cstdint>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
template <class T, int BlockSize>
__global__ void rms_norm(Span<T> output, View<T> input, View<T> scale,
std::size_t loops, std::size_t norm_size, float epsilon)
{
using BlockReduce = cub::BlockReduce<float, BlockSize>;
__shared__ typename BlockReduce::TempStorage reduce_storage;
__shared__ float inv_stddev;
const T* in = input.data().get();
const T* sc = scale.data().get();
T* out = output.data().get();
for (std::size_t row = blockIdx.x; row < loops; row += gridDim.x) {
const std::int64_t base = static_cast<std::int64_t>(row) * norm_size;
float thread_sum_sq = 0.f;
for (std::size_t j = threadIdx.x; j < norm_size; j += BlockSize) {
const float v = static_cast<float>(in[base + j]);
thread_sum_sq += v * v;
}
const float block_sum_sq = BlockReduce(reduce_storage).Sum(thread_sum_sq);
if (threadIdx.x == 0) {
// recenter=false collapses the variance to the mean square
const float mean_square = block_sum_sq / static_cast<float>(norm_size);
inv_stddev = 1.f / sqrtf(fmaxf(0.f, mean_square) + epsilon);
}
__syncthreads(); // publishes inv_stddev and frees reduce_storage for the next row
for (std::size_t j = threadIdx.x; j < norm_size; j += BlockSize) {
const float v = static_cast<float>(in[base + j]);
const float s = static_cast<float>(sc[j]);
out[base + j] = static_cast<T>(s * v * inv_stddev);
}
__syncthreads();
}
}
}
template <class T>
void rms_norm(const Stream& stream,
Span<T> output, View<T> input, View<T> scale,
std::size_t loops, std::size_t norm_size, float epsilon)
{
CV_Assert(loops > 0 && norm_size > 0);
CV_Assert(scale.size() >= norm_size);
constexpr int block_size = 256;
constexpr int max_blocks = 65535;
const int grid_size = static_cast<int>(std::min<std::size_t>(max_blocks, loops));
raw::rms_norm<T, block_size><<<grid_size, block_size, 0, stream.get()>>>(
output, input, scale, loops, norm_size, epsilon);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void rms_norm(const Stream&, Span<__half>, View<__half>, View<__half>, std::size_t, std::size_t, float);
#endif
template void rms_norm(const Stream&, Span<float>, View<float>, View<float>, std::size_t, std::size_t, float);
template void rms_norm(const Stream&, Span<int8_t>, View<int8_t>, View<int8_t>, std::size_t, std::size_t, float);
template void rms_norm(const Stream&, Span<uint8_t>, View<uint8_t>, View<uint8_t>, std::size_t, std::size_t, float);
template void rms_norm(const Stream&, Span<int32_t>, View<int32_t>, View<int32_t>, std::size_t, std::size_t, float);
template void rms_norm(const Stream&, Span<int64_t>, View<int64_t>, View<int64_t>, std::size_t, std::size_t, float);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -0,0 +1,161 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include "types.hpp"
#include "grid_stride_range.hpp"
#include "math.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/kernels/roi_align.hpp"
#include <opencv2/core.hpp>
#include <cfloat>
#include <cstddef>
#include <cstdint>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
__device__ inline std::int64_t read_batch_index(const void* data, std::size_t element_size, int offset) {
if (element_size == sizeof(std::int32_t))
return *(reinterpret_cast<const std::int32_t*>(data) + offset);
return *(reinterpret_cast<const std::int64_t*>(data) + offset);
}
/* One thread owns one output element. The ROI geometry is recomputed per thread
* rather than cached per ROI (as the CPU path does); it is a handful of flops
* against a memory-bound gather, and it keeps the kernel free of shared state. */
template <class T>
__global__ void roi_align(Span<T> output, View<T> input, View<float> rois,
const void* batch_indices, std::size_t batch_index_element_size,
RoiAlignParams p)
{
for (auto id : grid_stride_range(output.size())) {
int rem = id;
const int pw = rem % p.output_width; rem /= p.output_width;
const int ph = rem % p.output_height; rem /= p.output_height;
const int c = rem % p.channels; rem /= p.channels;
const int r = rem;
const float* roi = rois.data().get() + static_cast<std::int64_t>(r) * 4;
const float x1 = roi[0] * p.spatial_scale - p.offset;
const float y1 = roi[1] * p.spatial_scale - p.offset;
const float x2 = roi[2] * p.spatial_scale - p.offset;
const float y2 = roi[3] * p.spatial_scale - p.offset;
float roi_width = x2 - x1;
float roi_height = y2 - y1;
if (p.clamp_malformed_roi) {
roi_width = fmaxf(roi_width, 1.f);
roi_height = fmaxf(roi_height, 1.f);
}
const float bin_size_w = roi_width / p.output_width;
const float bin_size_h = roi_height / p.output_height;
const int gw = device::max<int>((p.sampling_ratio > 0) ? p.sampling_ratio
: static_cast<int>(ceilf(roi_width / p.output_width)), 1);
const int gh = device::max<int>((p.sampling_ratio > 0) ? p.sampling_ratio
: static_cast<int>(ceilf(roi_height / p.output_height)), 1);
const int sample_count = gw * gh;
const float step_y = bin_size_h / gh;
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<T>(0.f);
continue;
}
const T* img = input.data().get() +
((b * p.channels) + c) * static_cast<std::int64_t>(p.height) * p.width;
float outv = p.max_mode ? -FLT_MAX : 0.f;
for (int s = 0; s < sample_count; ++s) {
const int iy = s / gw;
const int ix = s - iy * gw;
const float yy = y1 + ph * bin_size_h + (iy + 0.5f) * step_y;
const float xx = x1 + pw * bin_size_w + (ix + 0.5f) * step_x;
/* precomputeBilinearSample() returns an all-zero sample when the point
* is fully outside [-1, extent], so it contributes 0 in both modes. */
float v = 0.f;
if (!(yy < -1.f || yy > p.height || xx < -1.f || xx > p.width)) {
const float cy = fminf(fmaxf(yy, 0.f), static_cast<float>(p.height - 1));
const float cx = fminf(fmaxf(xx, 0.f), static_cast<float>(p.width - 1));
const int y_low = static_cast<int>(cy);
const int x_low = static_cast<int>(cx);
const int y_high = device::min<int>(y_low + 1, p.height - 1);
const int x_high = device::min<int>(x_low + 1, p.width - 1);
const float ly = cy - y_low, lx = cx - x_low;
const float hy = 1.f - ly, hx = 1.f - lx;
const std::int64_t row_low = static_cast<std::int64_t>(y_low) * p.width;
const std::int64_t row_high = static_cast<std::int64_t>(y_high) * p.width;
const float v00 = static_cast<float>(img[row_low + x_low]);
const float v01 = static_cast<float>(img[row_low + x_high]);
const float v10 = static_cast<float>(img[row_high + x_low]);
const float v11 = static_cast<float>(img[row_high + x_high]);
v = p.max_mode
? fmaxf(fmaxf(hy * hx * v00, hy * lx * v01),
fmaxf(ly * hx * v10, ly * lx * v11))
: (hy * hx * v00 + hy * lx * v01 + ly * hx * v10 + ly * lx * v11);
}
outv = p.max_mode ? fmaxf(outv, v) : (outv + v);
}
outv = p.max_mode ? ((outv == -FLT_MAX) ? 0.f : outv)
: (outv / static_cast<float>(sample_count));
output[id] = static_cast<T>(outv);
}
}
}
template <class T>
void roi_align(const Stream& stream,
Span<T> output, View<T> input, View<float> rois,
const void* batch_indices, std::size_t batch_index_element_size,
const RoiAlignParams& params)
{
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.batch_size > 0 && params.height > 0 && params.width > 0 && params.channels > 0);
if (output.size() == 0)
return;
auto kernel = raw::roi_align<T>;
auto policy = make_policy(kernel, output.size(), 0, stream);
launch_kernel(kernel, policy, output, input, rois, batch_indices, batch_index_element_size, params);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void roi_align(const Stream&, Span<__half>, View<__half>, View<float>, const void*, std::size_t, const RoiAlignParams&);
#endif
template void roi_align(const Stream&, Span<float>, View<float>, View<float>, const void*, std::size_t, const RoiAlignParams&);
template void roi_align(const Stream&, Span<int8_t>, View<int8_t>, View<float>, const void*, std::size_t, const RoiAlignParams&);
template void roi_align(const Stream&, Span<uint8_t>, View<uint8_t>, View<float>, const void*, std::size_t, const RoiAlignParams&);
template void roi_align(const Stream&, Span<int32_t>, View<int32_t>, View<float>, const void*, std::size_t, const RoiAlignParams&);
template void roi_align(const Stream&, Span<int64_t>, View<int64_t>, View<float>, const void*, std::size_t, const RoiAlignParams&);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -0,0 +1,107 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include "array.hpp"
#include "types.hpp"
#include "grid_stride_range.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/csl/tensor.hpp"
#include "../cuda4dnn/kernels/scatter_elements.hpp"
#include <opencv2/core.hpp>
#include <cstdint>
#include <vector>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
template <class T, class TIdx>
__global__ void scatter_elements(
Span<T> output, View<TIdx> indices, View<T> updates,
int rank, int axis,
array<std::int64_t, CSL_MAX_TENSOR_RANK> indices_dims,
array<std::int64_t, CSL_MAX_TENSOR_RANK> output_strides,
std::int64_t axis_dim_size)
{
for (auto id : grid_stride_range(updates.size())) {
std::int64_t rem = static_cast<std::int64_t>(id);
std::int64_t offset = 0;
for (int d = rank - 1; d >= 0; --d) {
const std::int64_t dim = indices_dims[d];
const std::int64_t coord = rem % dim;
rem /= dim;
std::int64_t use = coord;
if (d == axis) {
std::int64_t idx = static_cast<std::int64_t>(indices.data().get()[id]);
if (idx < 0)
idx += axis_dim_size;
if (idx < 0)
idx = 0;
if (idx >= axis_dim_size)
idx = axis_dim_size - 1;
use = idx;
}
offset += use * output_strides[d];
}
output[offset] = updates.data().get()[id];
}
}
}
template <class T, class TIdx>
void scatter_elements(const Stream& stream,
Span<T> output, View<TIdx> indices, View<T> updates,
int rank, int axis,
const std::vector<std::int64_t>& indices_dims,
const std::vector<std::int64_t>& output_strides,
std::int64_t axis_dim_size)
{
CV_Assert(rank <= CSL_MAX_TENSOR_RANK);
CV_Assert(static_cast<int>(indices_dims.size()) == rank);
CV_Assert(static_cast<int>(output_strides.size()) == rank);
if (updates.size() == 0)
return;
array<std::int64_t, CSL_MAX_TENSOR_RANK> dims_k, strides_k;
dims_k.assign(std::begin(indices_dims), std::end(indices_dims));
strides_k.assign(std::begin(output_strides), std::end(output_strides));
auto kernel = raw::scatter_elements<T, TIdx>;
auto policy = make_policy(kernel, updates.size(), 0, stream);
launch_kernel(kernel, policy, output, indices, updates, rank, axis,
dims_k, strides_k, axis_dim_size);
}
#define CV_SCATTER_ELEM_INST(T) \
template void scatter_elements(const Stream&, Span<T>, View<std::int32_t>, View<T>, int, int, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::int64_t); \
template void scatter_elements(const Stream&, Span<T>, View<std::int64_t>, View<T>, int, int, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::int64_t);
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
CV_SCATTER_ELEM_INST(__half)
#endif
CV_SCATTER_ELEM_INST(float)
CV_SCATTER_ELEM_INST(int8_t)
CV_SCATTER_ELEM_INST(uint8_t)
CV_SCATTER_ELEM_INST(int32_t)
CV_SCATTER_ELEM_INST(int64_t)
#undef CV_SCATTER_ELEM_INST
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -0,0 +1,117 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include "array.hpp"
#include "types.hpp"
#include "grid_stride_range.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/csl/tensor.hpp"
#include "../cuda4dnn/kernels/scatter_nd.hpp"
#include <opencv2/core.hpp>
#include <cstddef>
#include <cstdint>
#include <vector>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
template <class T, class TIdx>
__global__ void scatter_nd(
Span<T> output, View<TIdx> indices, View<T> updates,
size_type num_indices, std::int64_t last_index_dimension,
array<std::int64_t, CSL_MAX_TENSOR_RANK> element_counts,
array<std::int64_t, CSL_MAX_TENSOR_RANK> input_dims,
std::int64_t num_updates_elements)
{
for (auto id : grid_stride_range(num_indices)) {
std::int64_t data_offset = 0;
std::int64_t indices_start = last_index_dimension * id;
for (std::int64_t j = 0; j < last_index_dimension; ++j) {
std::int64_t index = static_cast<std::int64_t>(indices.data().get()[indices_start + j]);
std::int64_t element_count_dim = element_counts[j];
std::int64_t dim_value = input_dims[j];
if (index >= 0) {
if (index >= dim_value) {
index = dim_value - 1;
}
} else {
if (index < -dim_value) {
index = 0;
} else {
index += dim_value;
}
}
data_offset += (index * element_count_dim);
}
const T* updates_data_base = updates.data().get() + num_updates_elements * static_cast<std::int64_t>(id);
T* output_data_base = output.data().get() + data_offset;
for (std::int64_t i = 0; i < num_updates_elements; ++i) {
output_data_base[i] = updates_data_base[i];
}
}
}
}
template <class T, class TIdx>
void scatter_nd(const Stream& stream,
Span<T> output, View<TIdx> indices, View<T> updates,
std::size_t num_indices, std::int64_t last_index_dimension,
const std::vector<std::int64_t>& element_counts,
const std::vector<std::int64_t>& input_dims,
std::size_t num_updates_elements)
{
CV_Assert(last_index_dimension <= CSL_MAX_TENSOR_RANK);
CV_Assert(element_counts.size() == static_cast<std::size_t>(last_index_dimension));
CV_Assert(input_dims.size() == static_cast<std::size_t>(last_index_dimension));
if (num_indices == 0)
return;
array<std::int64_t, CSL_MAX_TENSOR_RANK> element_counts_k, input_dims_k;
element_counts_k.assign(std::begin(element_counts), std::end(element_counts));
input_dims_k.assign(std::begin(input_dims), std::end(input_dims));
auto kernel = raw::scatter_nd<T, TIdx>;
auto policy = make_policy(kernel, num_indices, 0, stream);
launch_kernel(kernel, policy, output, indices, updates,
static_cast<size_type>(num_indices), last_index_dimension,
element_counts_k, input_dims_k,
static_cast<std::int64_t>(num_updates_elements));
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void scatter_nd(const Stream&, Span<__half>, View<std::int32_t>, View<__half>, std::size_t, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::size_t);
template void scatter_nd(const Stream&, Span<__half>, View<std::int64_t>, View<__half>, std::size_t, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::size_t);
#endif
template void scatter_nd(const Stream&, Span<float>, View<std::int32_t>, View<float>, std::size_t, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::size_t);
template void scatter_nd(const Stream&, Span<float>, View<std::int64_t>, View<float>, std::size_t, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::size_t);
template void scatter_nd(const Stream&, Span<int8_t>, View<std::int32_t>, View<int8_t>, std::size_t, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::size_t);
template void scatter_nd(const Stream&, Span<int8_t>, View<std::int64_t>, View<int8_t>, std::size_t, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::size_t);
template void scatter_nd(const Stream&, Span<uint8_t>, View<std::int32_t>, View<uint8_t>, std::size_t, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::size_t);
template void scatter_nd(const Stream&, Span<uint8_t>, View<std::int64_t>, View<uint8_t>, std::size_t, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::size_t);
template void scatter_nd(const Stream&, Span<int32_t>, View<std::int32_t>, View<int32_t>, std::size_t, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::size_t);
template void scatter_nd(const Stream&, Span<int32_t>, View<std::int64_t>, View<int32_t>, std::size_t, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::size_t);
template void scatter_nd(const Stream&, Span<int64_t>, View<std::int32_t>, View<int64_t>, std::size_t, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::size_t);
template void scatter_nd(const Stream&, Span<int64_t>, View<std::int64_t>, View<int64_t>, std::size_t, std::int64_t, const std::vector<std::int64_t>&, const std::vector<std::int64_t>&, std::size_t);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -0,0 +1,287 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <cub/block/block_reduce.cuh>
#include <cub/device/device_radix_sort.cuh>
#include <cub/device/device_segmented_sort.cuh>
#include "types.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/kernels/topk.hpp"
#include <opencv2/core.hpp>
#include <cuda/std/limits>
#include <algorithm>
#include <cstdint>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
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 <class T> struct TopKCompare { using type = float; };
template <> struct TopKCompare<std::int32_t> { using type = std::int32_t; };
template <> struct TopKCompare<std::int64_t> { 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 <class C>
__device__ inline C topk_worst(bool largest) {
using limits = ::cuda::std::numeric_limits<C>;
if (limits::has_infinity)
return largest ? -limits::infinity() : limits::infinity();
return largest ? limits::lowest() : limits::max();
}
template <class C>
struct Candidate {
C value;
int index;
};
/* Strict weak ordering matching topk2_layer.cpp's comparators: better value
* first, and on an exact tie the smaller index. */
template <class C>
__device__ inline bool better_than(const Candidate<C>& a, const Candidate<C>& b, bool largest) {
if (a.value == b.value)
return a.index < b.index;
return largest ? (a.value > b.value) : (a.value < b.value);
}
template <class C>
struct MergeCandidate {
bool largest;
__device__ inline Candidate<C> operator()(const Candidate<C>& a, const Candidate<C>& b) const {
return better_than(a, b, largest) ? a : b;
}
};
template <class T, int BlockSize>
__global__ void topk(Span<T> values, Span<std::int64_t> indices, View<T> input,
int outer, int dim_axis, int inner, int k, bool largest)
{
using C = typename TopKCompare<T>::type;
using Cand = Candidate<C>;
using BlockReduce = cub::BlockReduce<Cand, BlockSize>;
__shared__ typename BlockReduce::TempStorage reduce_storage;
__shared__ Cand selected;
const int nslices = outer * inner;
const C worst = topk_worst<C>(largest);
for (int slice = blockIdx.x; slice < nslices; slice += gridDim.x) {
const int b = slice / inner;
const int j = slice - b * inner;
const std::int64_t in_base = static_cast<std::int64_t>(b) * dim_axis * inner + j;
const std::int64_t out_base = static_cast<std::int64_t>(b) * k * inner + j;
for (int rank = 0; rank < k; ++rank) {
Cand best{worst, ::cuda::std::numeric_limits<int>::max()};
for (int u = threadIdx.x; u < dim_axis; u += BlockSize) {
const C v = static_cast<C>(input.data().get()[in_base + static_cast<std::int64_t>(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
* are still all reachable because index breaks the tie. */
if (rank > 0 && !better_than(selected, cand, largest))
continue;
if (better_than(cand, best, largest))
best = cand;
}
const Cand winner = BlockReduce(reduce_storage).Reduce(best, MergeCandidate<C>{largest});
if (threadIdx.x == 0) {
selected = winner;
const std::int64_t out = out_base + static_cast<std::int64_t>(rank) * inner;
values[out] = static_cast<T>(winner.value);
indices[out] = static_cast<std::int64_t>(winner.index);
}
__syncthreads(); // publishes `selected` and frees reduce_storage for the next rank
}
}
}
__global__ void topk_sort_init(int* ids, int* offsets, int n, int nsegments, int dim_axis)
{
const int stride = gridDim.x * blockDim.x;
for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += stride)
ids[i] = i % dim_axis;
for (int i = blockIdx.x * blockDim.x + threadIdx.x; i <= nsegments; i += stride)
offsets[i] = i * dim_axis;
}
template <class T>
__global__ void topk_sort_gather(Span<T> values, Span<std::int64_t> indices,
const T* keys, const int* ids, int dim_axis, int k)
{
const int m = static_cast<int>(values.size());
const int stride = gridDim.x * blockDim.x;
for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < m; i += stride) {
const int b = i / k;
const int src = b * dim_axis + (i - b * k);
values[i] = keys[src];
indices[i] = static_cast<std::int64_t>(ids[src]);
}
}
}
namespace detail {
struct TopKSortLayout {
std::size_t keys, ids_in, ids_out, offsets, temp, total;
};
inline std::size_t alignTo256(std::size_t bytes)
{
return (bytes + 255) & ~static_cast<std::size_t>(255);
}
template <class T>
TopKSortLayout topkSortLayout(int outer, int dim_axis, bool largest, cudaStream_t stream)
{
const std::int64_t n = static_cast<std::int64_t>(outer) * dim_axis;
std::size_t temp_bytes = 0;
if (outer == 1 && largest)
CUDA4DNN_CHECK_CUDA(cub::DeviceRadixSort::SortPairsDescending(
nullptr, temp_bytes, static_cast<const T*>(nullptr), static_cast<T*>(nullptr),
static_cast<const int*>(nullptr), static_cast<int*>(nullptr), n, 0,
static_cast<int>(sizeof(T) * 8), stream));
else if (outer == 1)
CUDA4DNN_CHECK_CUDA(cub::DeviceRadixSort::SortPairs(
nullptr, temp_bytes, static_cast<const T*>(nullptr), static_cast<T*>(nullptr),
static_cast<const int*>(nullptr), static_cast<int*>(nullptr), n, 0,
static_cast<int>(sizeof(T) * 8), stream));
else if (largest)
CUDA4DNN_CHECK_CUDA(cub::DeviceSegmentedSort::StableSortPairsDescending(
nullptr, temp_bytes, static_cast<const T*>(nullptr), static_cast<T*>(nullptr),
static_cast<const int*>(nullptr), static_cast<int*>(nullptr), n, outer,
static_cast<const int*>(nullptr), static_cast<const int*>(nullptr), stream));
else
CUDA4DNN_CHECK_CUDA(cub::DeviceSegmentedSort::StableSortPairs(
nullptr, temp_bytes, static_cast<const T*>(nullptr), static_cast<T*>(nullptr),
static_cast<const int*>(nullptr), static_cast<int*>(nullptr), n, outer,
static_cast<const int*>(nullptr), static_cast<const int*>(nullptr), stream));
TopKSortLayout layout;
layout.keys = 0;
layout.ids_in = layout.keys + alignTo256(static_cast<std::size_t>(n) * sizeof(T));
layout.ids_out = layout.ids_in + alignTo256(static_cast<std::size_t>(n) * sizeof(int));
layout.offsets = layout.ids_out + alignTo256(static_cast<std::size_t>(n) * sizeof(int));
layout.temp = layout.offsets + alignTo256(static_cast<std::size_t>(outer + 1) * sizeof(int));
layout.total = layout.temp + alignTo256(temp_bytes);
return layout;
}
}
template <class T>
void topk(const Stream& stream,
Span<T> values, Span<std::int64_t> indices, View<T> input,
int outer, int dim_axis, int inner, int k, bool largest)
{
CV_Assert(outer > 0 && dim_axis > 0 && inner > 0);
CV_Assert(k > 0 && k <= dim_axis);
constexpr int block_size = 256;
constexpr int max_blocks = 65535;
const int nslices = outer * inner;
const int grid_size = std::min(max_blocks, nslices);
raw::topk<T, block_size><<<grid_size, block_size, 0, stream.get()>>>(
values, indices, input, outer, dim_axis, inner, k, largest);
}
template <class T>
std::size_t topk_sort_workspace(const Stream& stream, int outer, int dim_axis, bool largest)
{
return detail::topkSortLayout<T>(outer, dim_axis, largest, stream.get()).total;
}
template <class T>
void topk_sort(const Stream& stream,
Span<T> values, Span<std::int64_t> indices, View<T> input,
int outer, int dim_axis, int k, bool largest,
void* workspace, std::size_t workspace_bytes)
{
CV_Assert(outer > 0 && dim_axis > 0);
CV_Assert(k > 0 && k <= dim_axis);
const detail::TopKSortLayout layout = detail::topkSortLayout<T>(outer, dim_axis, largest, stream.get());
CV_Assert(workspace != nullptr && workspace_bytes >= layout.total);
unsigned char* base = static_cast<unsigned char*>(workspace);
T* keys = reinterpret_cast<T*>(base + layout.keys);
int* ids_in = reinterpret_cast<int*>(base + layout.ids_in);
int* ids_out = reinterpret_cast<int*>(base + layout.ids_out);
int* offsets = reinterpret_cast<int*>(base + layout.offsets);
void* temp = base + layout.temp;
std::size_t temp_bytes = layout.total - layout.temp;
const int n = outer * dim_axis;
constexpr int block_size = 256;
constexpr int max_blocks = 4096;
const int init_grid = std::min(max_blocks, (std::max(n, outer + 1) + block_size - 1) / block_size);
raw::topk_sort_init<<<init_grid, block_size, 0, stream.get()>>>(ids_in, offsets, n, outer, dim_axis);
const T* keys_in = input.data().get();
if (outer == 1 && largest)
CUDA4DNN_CHECK_CUDA(cub::DeviceRadixSort::SortPairsDescending(
temp, temp_bytes, keys_in, keys, ids_in, ids_out, n, 0, static_cast<int>(sizeof(T) * 8), stream.get()));
else if (outer == 1)
CUDA4DNN_CHECK_CUDA(cub::DeviceRadixSort::SortPairs(
temp, temp_bytes, keys_in, keys, ids_in, ids_out, n, 0, static_cast<int>(sizeof(T) * 8), stream.get()));
else if (largest)
CUDA4DNN_CHECK_CUDA(cub::DeviceSegmentedSort::StableSortPairsDescending(
temp, temp_bytes, keys_in, keys, ids_in, ids_out, n, outer, offsets, offsets + 1, stream.get()));
else
CUDA4DNN_CHECK_CUDA(cub::DeviceSegmentedSort::StableSortPairs(
temp, temp_bytes, keys_in, keys, ids_in, ids_out, n, outer, offsets, offsets + 1, stream.get()));
const int m = outer * k;
const int gather_grid = std::min(max_blocks, (m + block_size - 1) / block_size);
raw::topk_sort_gather<T><<<gather_grid, block_size, 0, stream.get()>>>(values, indices, keys, ids_out, dim_axis, k);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void topk(const Stream&, Span<__half>, Span<std::int64_t>, View<__half>, int, int, int, int, bool);
#endif
template void topk(const Stream&, Span<float>, Span<std::int64_t>, View<float>, int, int, int, int, bool);
template void topk(const Stream&, Span<int8_t>, Span<std::int64_t>, View<int8_t>, int, int, int, int, bool);
template void topk(const Stream&, Span<uint8_t>, Span<std::int64_t>, View<uint8_t>, int, int, int, int, bool);
template void topk(const Stream&, Span<int32_t>, Span<std::int64_t>, View<int32_t>, int, int, int, int, bool);
template void topk(const Stream&, Span<int64_t>, Span<std::int64_t>, View<int64_t>, int, int, int, int, bool);
#define CUDA4DNN_INSTANTIATE_TOPK_SORT(T) \
template std::size_t topk_sort_workspace<T>(const Stream&, int, int, bool); \
template void topk_sort<T>(const Stream&, Span<T>, Span<std::int64_t>, View<T>, int, int, int, bool, void*, std::size_t);
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
CUDA4DNN_INSTANTIATE_TOPK_SORT(__half)
#endif
CUDA4DNN_INSTANTIATE_TOPK_SORT(float)
CUDA4DNN_INSTANTIATE_TOPK_SORT(int8_t)
CUDA4DNN_INSTANTIATE_TOPK_SORT(uint8_t)
CUDA4DNN_INSTANTIATE_TOPK_SORT(int32_t)
CUDA4DNN_INSTANTIATE_TOPK_SORT(int64_t)
#undef CUDA4DNN_INSTANTIATE_TOPK_SORT
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -0,0 +1,75 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include "types.hpp"
#include "grid_stride_range.hpp"
#include "math.hpp"
#include "execution.hpp"
#include "../cuda4dnn/csl/stream.hpp"
#include "../cuda4dnn/csl/span.hpp"
#include "../cuda4dnn/kernels/trilu.hpp"
#include <opencv2/core.hpp>
#include <algorithm>
#include <cstdint>
using namespace cv::dnn::cuda4dnn::csl;
using namespace cv::dnn::cuda4dnn::csl::device;
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
namespace raw {
template <class T>
__global__ void trilu(Span<T> output, View<T> input,
int height, int width, int k, bool upper)
{
const T* in = input.data().get();
T* out = output.data().get();
for (auto id : grid_stride_range(output.size())) {
const int c = id % width;
const int l = (id / width) % height;
// 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];
}
}
}
template <class T>
void trilu(const Stream& stream,
Span<T> output, View<T> input,
int loops, int height, int width, int k, bool upper)
{
CV_Assert(loops > 0 && height > 0 && width > 0);
CV_UNUSED(loops);
if (output.size() == 0)
return;
auto kernel = raw::trilu<T>;
auto policy = make_policy(kernel, output.size(), 0, stream);
launch_kernel(kernel, policy, output, input, height, width, k, upper);
}
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template void trilu(const Stream&, Span<__half>, View<__half>, int, int, int, int, bool);
#endif
template void trilu(const Stream&, Span<float>, View<float>, int, int, int, int, bool);
template void trilu(const Stream&, Span<int32_t>, View<int32_t>, int, int, int, int, bool);
template void trilu(const Stream&, Span<int64_t>, View<int64_t>, int, int, int, int, bool);
template void trilu(const Stream&, Span<int8_t>, View<int8_t>, int, int, int, int, bool);
template void trilu(const Stream&, Span<uint8_t>, View<uint8_t>, int, int, int, int, bool);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */

View File

@@ -11,7 +11,6 @@
#include "memory.hpp"
#include <opencv2/core.hpp>
#include <opencv2/core/utils/configuration.private.hpp>
#include <cublas_v2.h>

View File

@@ -180,7 +180,10 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace csl { namespace cu
return false;
};
bool found = pick(CUDNN_HEUR_MODE_A, strictFP32) || (strictFP32 && pick(CUDNN_HEUR_MODE_FALLBACK, true));
// mode B ranks engines more accurately than mode A, still deterministically; if it does not
// finalize (e.g. its heuristics library is not loaded), pick() returns false and mode A decides
bool found = pick(CUDNN_HEUR_MODE_B, strictFP32) || pick(CUDNN_HEUR_MODE_A, strictFP32) ||
(strictFP32 && pick(CUDNN_HEUR_MODE_FALLBACK, true));
if (!found && strictFP32 && pick(CUDNN_HEUR_MODE_A, false))
CV_LOG_ONCE_WARNING(NULL, "DNN/CUDA: cuDNN JIT has no exact FP32 engine for some " << what << " layers, using TF32 for them");

View File

@@ -0,0 +1,23 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_ARGMAX_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_ARGMAX_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
#include <cstdint>
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
template <class T>
void arg_min_max(const csl::Stream& stream,
csl::Span<std::int64_t> output, csl::View<T> input,
int outer_size, int axis_size, int inner_size, bool is_argmax, bool select_last_index);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_ARGMAX_HPP */

View File

@@ -0,0 +1,29 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_COL2IM_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_COL2IM_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
struct Col2ImParams {
int batch, channels, total_channels, channel_offset;
int in_h, in_w, out_h, out_w;
int kernel_h, kernel_w, stride_h, stride_w, pad_h, pad_w, dil_h, dil_w;
};
// cols is [batch][channels * kernel_h * kernel_w][in_h * in_w]; each output pixel gathers the
// column entries that land on it, plus bias[channel_offset + c] when bias is not empty
template <class T>
void col2im_2d(const csl::Stream& stream, csl::Span<T> output, csl::View<T> cols, csl::View<T> bias,
const Col2ImParams& params);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_COL2IM_HPP */

View File

@@ -0,0 +1,35 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_CUM_SCAN_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_CUM_SCAN_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
#include <cstddef>
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
/* ONNX CumSum / CumProd over one axis, with the input viewed as
* [outer_size, target_size, inner_size].
*
* `is_prod` picks the operator and the exclusive-mode seed (0 for sum, 1 for product);
* `exclusive` and `reverse` follow the ONNX attributes. Mirrors the scan in
* cumsum_layer.cpp / cumprod_layer.cpp exactly, including their offset arithmetic.
*
* The scan along the axis is inherently sequential, so a thread owns one
* (outer, inner) pair and walks the axis; parallelism comes from outer * inner, which
* is where the width is in practice.
*/
template <class T>
void cum_scan(const csl::Stream& stream,
csl::Span<T> output, csl::View<T> input,
std::size_t outer_size, std::size_t target_size, std::size_t inner_size,
bool is_prod, bool exclusive, bool reverse);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_CUM_SCAN_HPP */

View File

@@ -42,6 +42,19 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
template <class T>
void eltwise_pow_2(const csl::Stream& stream, csl::TensorSpan<T> output, csl::TensorView<T> x, csl::TensorView<T> y);
template <class T>
void eltwise_prelu_2(const csl::Stream& stream, csl::TensorSpan<T> output, csl::TensorView<T> x, csl::TensorView<T> slope);
// integer types only
template <class T>
void eltwise_bitand_2(const csl::Stream& stream, csl::TensorSpan<T> output, csl::TensorView<T> x, csl::TensorView<T> y);
template <class T>
void eltwise_bitor_2(const csl::Stream& stream, csl::TensorSpan<T> output, csl::TensorView<T> x, csl::TensorView<T> y);
template <class T>
void eltwise_bitxor_2(const csl::Stream& stream, csl::TensorSpan<T> output, csl::TensorView<T> x, csl::TensorView<T> y);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_ELTWISE_OPS_HPP */

View File

@@ -0,0 +1,28 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_GATHER_ELEMENTS_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_GATHER_ELEMENTS_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
#include "../../cuda/fast_divmod.hpp"
#include <cstddef>
#include <cstdint>
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
template <class T>
void gather_elements(const csl::Stream& stream,
csl::Span<T> output, csl::View<T> input,
const void* indices, std::size_t index_element_size,
std::int64_t axis_size_data, std::int64_t inner_size,
csl::device::fast_divmod axis_inner_size, csl::device::fast_divmod inner_size_divmod);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_GATHER_ELEMENTS_HPP */

View File

@@ -0,0 +1,30 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_GATHER_ND_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_GATHER_ND_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
#include <cstddef>
#include <cstdint>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
template <class T, class TIdx>
void gather_nd(const csl::Stream& stream,
csl::Span<T> output, csl::View<TIdx> indices, csl::View<T> input,
std::int64_t last_index_dimension,
const std::vector<std::int64_t>& element_counts,
const std::vector<std::int64_t>& input_dims,
std::int64_t inner_size,
std::int64_t batch_stride,
std::int64_t num_slices_per_batch);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_GATHER_ND_HPP */

View File

@@ -0,0 +1,37 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_GRID_SAMPLE_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_GRID_SAMPLE_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
struct GridSampleParams {
int channels{}, height{}, width{};
int output_height{}, output_width{};
int mode{1}; // 0 = nearest, 1 = bilinear, 2 = bicubic
int padding{0}; // 0 = zeros, 1 = border, 2 = reflection
bool align_corners{false};
float cubic_alpha{-0.75f};
};
/* ONNX GridSample over 4-D input. `input` is [N, C, H, W], `grid` is
* [N, H_out, W_out, 2] in CV_32F holding normalised (x, y), and the output is
* [N, C, H_out, W_out].
*
* Mirrors gridSampleComputeRows() in gridsample_layer.cpp. That function has a fast
* path for fully-interior taps and a general `fetch` path; the two agree by
* construction, so this kernel only implements the general one. */
template <class T>
void grid_sample(const csl::Stream& stream,
csl::Span<T> output, csl::View<T> input, csl::View<float> grid,
const GridSampleParams& params);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_GRID_SAMPLE_HPP */

View File

@@ -18,6 +18,14 @@ void reduce_mean(const csl::Stream& stream, csl::Span<float> means, csl::View<T>
template <class T>
void reduce_mean_sqr_sum(const csl::Stream& stream, csl::Span<float> means, csl::Span<float> sum_sqrs, csl::View<T> input, std::size_t inner_size);
// floats of scratch the overload below needs to split rows this long across blocks; 0 when it does not split
std::size_t reduce_mean_sqr_sum_scratch(std::size_t rows, std::size_t inner_size);
// same result, with long rows split across blocks and the partials added in a fixed order
template <class T>
void reduce_mean_sqr_sum(const csl::Stream& stream, csl::Span<float> means, csl::Span<float> sum_sqrs, csl::View<T> input,
std::size_t inner_size, csl::Span<float> scratch);
void compute_normalization_scale(const csl::Stream& stream, csl::Span<float> scale, csl::View<float> means, csl::View<float> sum_sqrs, std::size_t inner_size, float eps);
template <class T>

View File

@@ -19,6 +19,14 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
std::size_t outer_size, std::size_t mid_size, std::size_t inner_size, std::size_t norm, T epsilon,
csl::Span<T> workspace);
// ONNX LRN: y = x * (bias + alpha_scaled * sum of x^2 over channels [c - radius, c + radius])^-beta
template <class T>
void lrn_across_channels(
const csl::Stream& stream,
csl::Span<T> output, csl::View<T> input,
std::size_t batch, std::size_t channels, std::size_t plane_size,
int radius, float alpha_scaled, float bias, float beta);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_NORMALIZE_HPP */

View File

@@ -0,0 +1,22 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_ONE_HOT_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_ONE_HOT_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
#include <cstddef>
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
template <class T, class TIdx>
void one_hot(const csl::Stream& stream, csl::Span<T> output, csl::View<TIdx> indices,
std::size_t depth, std::size_t inner, T off_value, T on_value);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_ONE_HOT_HPP */

View File

@@ -20,6 +20,16 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
csl::TensorSpan<T> output, csl::TensorView<T> input,
std::vector<std::pair<std::size_t, std::size_t>> ranges);
// border rules follow cv::borderInterpolate, as the CPU Pad2 does
enum class PadMode { CONSTANT, REFLECT101, REPLICATE, WRAP };
// pads_begin may be negative (crop); output shape already includes the pads
template <class T>
void pad_nd(
const csl::Stream& stream,
csl::TensorSpan<T> output, csl::TensorView<T> input,
const std::vector<int>& pads_begin, PadMode mode, T value);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_PADDING_HPP */

View File

@@ -0,0 +1,38 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_QUANTIZE_DEQUANTIZE_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_QUANTIZE_DEQUANTIZE_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
#include <cstddef>
#include <cstdint>
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
// Per-tensor (sz_a == 1) and per-axis (sz_a == input.size(axis)) ONNX DequantizeLinear/
// QuantizeLinear, block_size == 0 only. zero_point may be an empty view (implies zero).
template <class T>
void dequantize_linear(const csl::Stream& stream, csl::Span<float> output, csl::View<T> input,
csl::View<float> scale, csl::View<T> zero_point,
std::size_t sz_a, std::size_t slice_size);
template <class T>
void quantize_linear(const csl::Stream& stream, csl::Span<T> output, csl::View<float> input,
csl::View<float> scale, csl::View<T> zero_point,
std::size_t sz_a, std::size_t slice_size);
// Number of (min, max) partial pairs dynamic_quantize_scale() keeps in scratch for n inputs.
std::size_t dynamic_quantize_partials(std::size_t n);
// ONNX DynamicQuantizeLinear y_scale and uint8 y_zero_point from the input range, which is forced
// to include zero; scratch holds 2 * dynamic_quantize_partials(input.size()) floats.
void dynamic_quantize_scale(const csl::Stream& stream, csl::Span<float> scale, csl::Span<std::uint8_t> zero_point,
csl::View<float> input, csl::Span<float> scratch);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_QUANTIZE_DEQUANTIZE_HPP */

View File

@@ -0,0 +1,20 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_RANGE_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_RANGE_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
/* start and delta are read on the device, so no host round trip is needed */
template <class T>
void range(const csl::Stream& stream, csl::Span<T> output, const T* start, const T* delta);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_RANGE_HPP */

View File

@@ -7,9 +7,18 @@
#include "../csl/stream.hpp"
#include "../csl/tensor.hpp"
#include "../csl/span.hpp"
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
// separable gather over host-built per-axis taps (idx < 0 skips a tap); a pixel whose row or
// column is flagged out of bounds takes extrapolation_value (tf_crop_and_resize)
template <class T>
void resize_separable(const csl::Stream& stream, csl::TensorSpan<T> output, csl::TensorView<T> input,
csl::View<int> idx_y, csl::View<float> w_y, csl::View<int> oob_y, int taps_y,
csl::View<int> idx_x, csl::View<float> w_x, csl::View<int> oob_x, int taps_x,
float extrapolation_value);
template <class T>
void resize_nn(const csl::Stream& stream, csl::TensorSpan<T> output, csl::TensorView<T> input, float scale_y, float scale_x, bool round, bool half_pixel_centers);

View File

@@ -0,0 +1,34 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_RMS_NORM_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_RMS_NORM_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
#include <cstddef>
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
/* RMS normalisation: y = scale * x / sqrt(mean(x^2) + epsilon), applied over the last
* `norm_size` elements of each of `loops` rows. `scale` holds `norm_size` elements.
*
* Mirrors fastNormImpl(input, scale, output, eps, axis, recenter=false) in
* cpu_kernels/fast_norm.cpp -- the recenter=false path forces mean to 0, so the
* variance term collapses to the mean square.
*
* Unlike the LayerNorm path this needs no workspace: one block owns one row and
* reduces it in registers plus shared memory, instead of accumulating into a global
* scratch buffer with atomicAdd.
*/
template <class T>
void rms_norm(const csl::Stream& stream,
csl::Span<T> output, csl::View<T> input, csl::View<T> scale,
std::size_t loops, std::size_t norm_size, float epsilon);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_RMS_NORM_HPP */

View File

@@ -0,0 +1,44 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_ROI_ALIGN_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_ROI_ALIGN_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
#include <cstddef>
#include <cstdint>
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
struct RoiAlignParams {
int batch_size{}, channels{}, height{}, width{};
int output_height{}, output_width{};
int sampling_ratio{};
float spatial_scale{1.f};
/* half_pixel subtracts 0.5 from the scaled ROI corners; output_half_pixel does not
* but clamps a degenerate ROI up to one unit. They are the two ONNX
* coordinate_transformation_mode values and they move independently. */
float offset{0.f};
bool clamp_malformed_roi{false};
bool max_mode{false};
};
/* ONNX RoiAlign. `rois` is [num_rois, 4] as (x1, y1, x2, y2) in CV_32F; `batch_indices`
* is [num_rois] of either int32 or int64, selected by `batch_index_element_size`.
* Output is [num_rois, channels, output_height, output_width].
*
* Mirrors the CPU RoiAlignForwardInvoker element for element, including its rule that a
* sample falling fully outside [-1, extent] contributes nothing at all. */
template <class T>
void roi_align(const csl::Stream& stream,
csl::Span<T> output, csl::View<T> input, csl::View<float> rois,
const void* batch_indices, std::size_t batch_index_element_size,
const RoiAlignParams& params);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_ROI_ALIGN_HPP */

View File

@@ -0,0 +1,27 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_SCATTER_ELEMENTS_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_SCATTER_ELEMENTS_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
#include <cstdint>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
template <class T, class TIdx>
void scatter_elements(const csl::Stream& stream,
csl::Span<T> output, csl::View<TIdx> indices, csl::View<T> updates,
int rank, int axis,
const std::vector<std::int64_t>& indices_dims,
const std::vector<std::int64_t>& output_strides,
std::int64_t axis_dim_size);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_SCATTER_ELEMENTS_HPP */

View File

@@ -0,0 +1,28 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_SCATTER_ND_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_SCATTER_ND_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
#include <cstddef>
#include <cstdint>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
template <class T, class TIdx>
void scatter_nd(const csl::Stream& stream,
csl::Span<T> output, csl::View<TIdx> indices, csl::View<T> updates,
std::size_t num_indices, std::int64_t last_index_dimension,
const std::vector<std::int64_t>& element_counts,
const std::vector<std::int64_t>& input_dims,
std::size_t num_updates_elements);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_SCATTER_ND_HPP */

View File

@@ -0,0 +1,44 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_TOPK_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_TOPK_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
#include <cstddef>
#include <cstdint>
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
/* ONNX TopK along one axis. The input is treated as [outer, dim_axis, inner] and the
* outputs as [outer, k, inner], matching TopK2's CPU layout.
*
* Selection is by strict lexicographic order -- value first, then smaller index on a
* tie -- which is what ComparatorGreater/ComparatorLess in topk2_layer.cpp impose. The
* results come out in rank order, so `sorted=1` is satisfied and `sorted=0` (which
* leaves order unspecified) is satisfied too.
*
* One block owns one slice and makes `k` passes over it, so the cost is O(k * dim_axis).
* That suits the small k detection heads use; it is not a general large-k sort.
*/
template <class T>
void topk(const csl::Stream& stream,
csl::Span<T> values, csl::Span<std::int64_t> indices, csl::View<T> input,
int outer, int dim_axis, int inner, int k, bool largest);
template <class T>
std::size_t topk_sort_workspace(const csl::Stream& stream, int outer, int dim_axis, bool largest);
template <class T>
void topk_sort(const csl::Stream& stream,
csl::Span<T> values, csl::Span<std::int64_t> indices, csl::View<T> input,
int outer, int dim_axis, int k, bool largest,
void* workspace, std::size_t workspace_bytes);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_TOPK_HPP */

View File

@@ -0,0 +1,22 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_TRILU_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_TRILU_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
// 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,
int loops, int height, int width, int k, bool upper);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_TRILU_HPP */

View File

@@ -53,6 +53,37 @@ namespace cv { namespace dnn { namespace cuda4dnn {
const T slope;
};
// ONNX Clip: min/max may arrive as extra inputs, so only input 0 is data
template <class T>
class ClipOp final : public CUDABackendNode {
public:
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,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(!inputs.empty() && outputs.size() == 1);
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
kernels::clipped_relu<T>(stream, output, input, min, max);
}
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;
};
template <class T>
class ClippedReLUOp final : public BaseOp<ClippedReLUOp, T> {
public:

View File

@@ -0,0 +1,65 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_ARG_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_ARG_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../kernels/argmax.hpp"
#include <opencv2/core.hpp>
#include <cstdint>
#include <utility>
namespace cv { namespace dnn { namespace cuda4dnn {
template <class T>
class ArgOp final : public CUDABackendNode {
public:
ArgOp(csl::Stream stream_, int axis_, bool is_argmax_, bool select_last_index_)
: stream(std::move(stream_)), axis(axis_), is_argmax(is_argmax_), select_last_index(select_last_index_)
{
}
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() == 1);
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<std::int64_t>(outputs[0]);
MatShape inShape = cv::dnn::shape(inputs[0]);
const int rank = inShape.dims;
const int ax = normalize_axis(axis, rank);
const int axis_size = inShape[ax];
int outer_size = 1;
for (int i = 0; i < ax; i++)
outer_size *= inShape[i];
int inner_size = 1;
for (int i = ax + 1; i < rank; i++)
inner_size *= inShape[i];
kernels::arg_min_max<T>(stream, output, input, outer_size, axis_size, inner_size, is_argmax, select_last_index);
}
private:
csl::Stream stream;
int axis;
bool is_argmax;
bool select_last_index;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_ARG_HPP */

View File

@@ -12,6 +12,7 @@
#include "../csl/tensor.hpp"
#include "../kernels/average_pooling.hpp"
#include "../kernels/reduce.hpp"
#include <opencv2/core.hpp>
@@ -58,11 +59,35 @@ namespace cv { namespace dnn { namespace cuda4dnn {
std::vector<std::int64_t> in_shape(inShape.begin(), inShape.end());
std::vector<std::int64_t> out_shape(outShape.begin(), outShape.end());
// a window over every spatial axis without padding is a per-channel mean, which the
// reduction kernel spreads over a block per channel instead of one thread per output
if (isGlobal(in_shape)) {
std::vector<std::int64_t> axes;
for (std::size_t k = 2; k < in_shape.size(); k++)
axes.push_back(static_cast<std::int64_t>(k));
kernels::reduce_mean<T>(stream, output, input, in_shape, axes);
return;
}
kernels::average_pool<T>(stream, output, input,
in_shape, out_shape, kernel_shape, strides, pads, dilations, count_include_pad);
}
private:
bool isGlobal(const std::vector<std::int64_t>& in_shape) const
{
const std::size_t nspatial = kernel_shape.size();
if (nspatial == 0 || in_shape.size() != nspatial + 2)
return false;
for (std::size_t i = 0; i < nspatial; i++)
if (kernel_shape[i] != in_shape[2 + i] || (!dilations.empty() && dilations[i] != 1))
return false;
for (std::int64_t p : pads)
if (p != 0)
return false;
return true;
}
csl::Stream stream;
std::vector<std::int64_t> kernel_shape, strides, pads, dilations;
bool count_include_pad;

View File

@@ -0,0 +1,150 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_CONV_TRANSPOSE2_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_CONV_TRANSPOSE2_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../csl/cublas.hpp"
#include "../csl/span.hpp"
#include "../csl/tensor.hpp"
#include "../csl/workspace.hpp"
#include "../kernels/col2im.hpp"
#include <opencv2/core.hpp>
#include <cstddef>
#include <vector>
#include <utility>
namespace cv { namespace dnn { namespace cuda4dnn {
struct ConvTranspose2Configuration {
MatShape input_shape, weight_shape, output_shape; // NCHW, [Cin, Cout/group, kh, kw], NCHW
int groups;
std::vector<int> strides, dilations, pads_begin; // two spatial entries each
};
// ConvTranspose without legacy cuDNN: per group, cols = W^T * X (cuBLAS), then col2im adds into Y
template <class T>
class ConvTranspose2Op final : public CUDABackendNode {
public:
ConvTranspose2Op(csl::Stream stream_, csl::cublas::Handle handle, const ConvTranspose2Configuration& config)
: stream(std::move(stream_)), cublasHandle(std::move(handle))
{
const MatShape& in = config.input_shape;
const MatShape& w = config.weight_shape;
const MatShape& out = config.output_shape;
CV_Assert(in.dims == 4 && w.dims == 4 && out.dims == 4);
CV_Assert(config.groups > 0 && in[1] % config.groups == 0 && w[0] == in[1]);
groups = config.groups;
batch = in[0];
cin = in[1];
cin_g = cin / groups;
cout_g = w[1];
cout = out[1];
CV_Assert(cout == cout_g * groups);
in_h = in[2]; in_w = in[3];
out_h = out[2]; out_w = out[3];
kernel_h = w[2]; kernel_w = w[3];
strides = config.strides;
dilations = config.dilations;
pads_begin = config.pads_begin;
// one group's columns for the whole batch; groups reuse it in stream order
csl::WorkspaceBuilder builder;
builder.require<T>(colsSize());
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
{
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>(owned ? ownedWeights : inputs[1]);
csl::View<T> bias;
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]);
auto ws_allocator = csl::WorkspaceAllocator(workspace);
auto cols = ws_allocator.get_span<T>(colsSize());
const std::size_t hw = (std::size_t)in_h * in_w;
const std::size_t m = (std::size_t)cout_g * kernel_h * kernel_w;
for (int g = 0; g < groups; g++) {
// row-major cols(m x hw) = W_g^T(m x cin_g) * X_g(cin_g x hw); in column-major terms cols^T = X_g^T * W_g
csl::cublas::gemmStridedBatched<T>(cublasHandle,
false, true, hw, m, cin_g,
T(1), x.get() + (std::ptrdiff_t)(g * cin_g * hw), hw, (std::size_t)cin * hw,
w.get() + (std::ptrdiff_t)(g * cin_g * m), m, 0,
T(0), cols.data(), hw, m * hw,
batch);
kernels::Col2ImParams p;
p.batch = batch; p.channels = cout_g; p.total_channels = cout; p.channel_offset = g * cout_g;
p.in_h = in_h; p.in_w = in_w; p.out_h = out_h; p.out_w = out_w;
p.kernel_h = kernel_h; p.kernel_w = kernel_w;
p.stride_h = strides[0]; p.stride_w = strides[1];
p.pad_h = pads_begin[0]; p.pad_w = pads_begin[1];
p.dil_h = dilations[0]; p.dil_w = dilations[1];
kernels::col2im_2d<T>(stream, csl::Span<T>(y.get(), y.size()),
csl::View<T>(cols.data(), cols.size()), bias, p);
}
}
std::size_t get_workspace_memory_in_bytes() const noexcept override { return scratch_mem_in_bytes; }
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;
int in_h, in_w, out_h, out_w, kernel_h, kernel_w;
std::vector<int> strides, dilations, pads_begin;
std::size_t scratch_mem_in_bytes;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_CONV_TRANSPOSE2_HPP */

View File

@@ -354,9 +354,12 @@ namespace cv { namespace dnn { namespace cuda4dnn {
builder.require<T>(sz);
}
#if defined(HAVE_CUDNNJIT) && !defined(HAVE_CUDNN)
// NHWC staging copies come from the shared workspace, taken before the conv scratchpad
builder.require<T>(std::accumulate(std::begin(jitInputShapeCL), std::end(jitInputShapeCL), std::size_t(1), std::multiplies<std::size_t>()));
builder.require<T>(std::accumulate(std::begin(jitOutputShapeCL), std::end(jitOutputShapeCL), std::size_t(1), std::multiplies<std::size_t>()));
if (!jitInputShapeCL.empty())
{
// NHWC staging copies come from the shared workspace, taken before the conv scratchpad
builder.require<T>(std::accumulate(std::begin(jitInputShapeCL), std::end(jitInputShapeCL), std::size_t(1), std::multiplies<std::size_t>()));
builder.require<T>(std::accumulate(std::begin(jitOutputShapeCL), std::end(jitOutputShapeCL), std::size_t(1), std::multiplies<std::size_t>()));
}
builder.require(jitConvoluter.get_workspace_size());
#else
builder.require(convoluter.get_workspace_size());
@@ -386,8 +389,12 @@ namespace cv { namespace dnn { namespace cuda4dnn {
}
#if defined(HAVE_CUDNNJIT) && !defined(HAVE_CUDNN)
auto jitInputNHWC = allocator.get_tensor_span<T>(std::begin(jitInputShapeCL), std::end(jitInputShapeCL));
auto jitOutputNHWC = allocator.get_tensor_span<T>(std::begin(jitOutputShapeCL), std::end(jitOutputShapeCL));
csl::TensorSpan<T> jitInputNHWC, jitOutputNHWC;
if (!jitInputShapeCL.empty())
{
jitInputNHWC = allocator.get_tensor_span<T>(std::begin(jitInputShapeCL), std::end(jitInputShapeCL));
jitOutputNHWC = allocator.get_tensor_span<T>(std::begin(jitOutputShapeCL), std::end(jitOutputShapeCL));
}
#endif
auto conv_scratchpad = allocator.get_instance();

View File

@@ -0,0 +1,67 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_CUM_SCAN_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_CUM_SCAN_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../kernels/cum_scan.hpp"
#include <opencv2/core.hpp>
#include <opencv2/dnn/shape_utils.hpp>
#include <cstddef>
#include <utility>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn {
/* Backs CumSum and CumProd. The axis is resolved by the layer at init time -- it can
* arrive as an input tensor, and reading it here would mean a device-to-host stall
* inside forward(). */
template <class T>
class CumScanOp final : public CUDABackendNode {
public:
CumScanOp(csl::Stream stream_, int axis_, bool is_prod_, bool exclusive_, bool reverse_)
: stream(std::move(stream_)), axis(axis_),
is_prod(is_prod_), exclusive(exclusive_), reverse(reverse_) { }
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() == 1);
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
MatShape inShape = cv::dnn::shape(inputs[0]);
const int a = normalize_axis(axis, inShape.dims);
CV_Assert(a >= 0 && a < inShape.dims);
std::size_t outer_size = 1, inner_size = 1;
for (int i = 0; i < a; i++) outer_size *= static_cast<std::size_t>(inShape[i]);
for (int i = a + 1; i < inShape.dims; i++) inner_size *= static_cast<std::size_t>(inShape[i]);
const std::size_t target_size = static_cast<std::size_t>(inShape[a]);
kernels::cum_scan<T>(stream, output, input,
outer_size, target_size, inner_size,
is_prod, exclusive, reverse);
}
private:
csl::Stream stream;
int axis;
bool is_prod, exclusive, reverse;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_CUM_SCAN_HPP */

View File

@@ -40,17 +40,15 @@ namespace cv { namespace dnn { namespace cuda4dnn {
scratch_mem_in_bytes = builder.required_workspace_size();
}
void forward(const std::vector<cv::Ptr<BackendWrapper>> &inputs,
const std::vector<cv::Ptr<BackendWrapper>> &outputs,
void forward(const std::vector<UMat> &inputs,
const std::vector<UMat> &outputs,
csl::Workspace &workspace) override {
CV_CheckEQ(inputs.size(), size_t(1), "DepthSpaceOps: only one input is accepted");
CV_CheckEQ(outputs.size(), size_t(1), "DepthSpaceOps: only one output is accepted");
auto input_wrapper = inputs.front().dynamicCast<wrapper_type>();
auto input = input_wrapper->getView();
auto input = csl::viewOf<T>(inputs.front());
CV_CheckEQ(input.rank(), size_t(4), "DepthSpaceOps: input needs to be 4-dimensional [N, C, H, W]");
auto output_wrapper = outputs.front().dynamicCast<wrapper_type>();
auto output = output_wrapper->getSpan();
auto output = csl::spanOf<T>(outputs.front());
auto ws_allocator = csl::WorkspaceAllocator(workspace);
auto transposed_internal = ws_allocator.get_tensor_span<T>(transposed_internal_shape.begin(), transposed_internal_shape.end());

View File

@@ -0,0 +1,51 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_DROPOUT_MASK_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_DROPOUT_MASK_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../csl/tensor_ops.hpp"
#include "../kernels/fill_copy.hpp"
#include <vector>
#include <utility>
namespace cv { namespace dnn { namespace cuda4dnn {
// Inference-mode Dropout with its mask output(s): y = x, and every mask is all true.
template <class T>
class DropoutMaskOp final : public CUDABackendNode {
public:
DropoutMaskOp(csl::Stream stream_) : stream(std::move(stream_)) { }
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(!inputs.empty() && outputs.size() >= 2);
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
if (input.get() != output.get())
csl::tensor_ops::copy(stream, output, input);
for (std::size_t i = 1; i < outputs.size(); i++)
kernels::fill<bool>(stream, csl::spanOf<bool>(outputs[i]), true);
}
private:
csl::Stream stream;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_DROPOUT_MASK_HPP */

View File

@@ -0,0 +1,108 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_EINSUM_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_EINSUM_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../csl/cublas.hpp"
#include "../csl/tensor.hpp"
#include "../csl/tensor_ops.hpp"
#include "../csl/workspace.hpp"
#include "../kernels/permute.hpp"
#include <cstddef>
#include <utility>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn {
/** A two-operand einsum as permute -> strided batched GEMM -> permute.
*
* A is read as [batch, M, K] (or [batch, K, M] with transA) and B as [batch, K, N]
* (or [batch, N, K] with transB); the GEMM writes [batch, M, N]. An empty perm means
* the operand already has that layout, so it is used in place.
*/
struct EinsumPlan {
std::vector<std::size_t> permA, permB, permC; // permC maps [batch, M, N] to the output
std::vector<std::size_t> shapeA, shapeB, shapeC; // full shapes after permA/permB, and of the GEMM result
bool transA = false, transB = false;
std::size_t batch = 1, M = 1, N = 1, K = 1;
};
template <class T>
class EinsumOp final : public CUDABackendNode {
public:
EinsumOp(csl::Stream stream_, csl::cublas::Handle handle_, EinsumPlan plan_)
: stream(std::move(stream_)), cublasHandle(std::move(handle_)), plan(std::move(plan_))
{
csl::WorkspaceBuilder builder;
if (!plan.permA.empty())
builder.require<T>(plan.batch * plan.M * plan.K);
if (!plan.permB.empty())
builder.require<T>(plan.batch * plan.K * plan.N);
if (!plan.permC.empty())
builder.require<T>(plan.batch * plan.M * plan.N);
scratch_mem_in_bytes = builder.required_workspace_size();
}
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_Assert(inputs.size() == 2 && outputs.size() == 1);
auto allocator = csl::WorkspaceAllocator(workspace);
csl::TensorView<T> a = operand(allocator, csl::viewOf<T>(inputs[0]), plan.permA, plan.shapeA);
csl::TensorView<T> b = operand(allocator, csl::viewOf<T>(inputs[1]), plan.permB, plan.shapeB);
const std::size_t rowsA = plan.transA ? plan.K : plan.M, colsA = plan.transA ? plan.M : plan.K;
const std::size_t rowsB = plan.transB ? plan.N : plan.K, colsB = plan.transB ? plan.K : plan.N;
auto output = csl::spanOf<T>(outputs[0]);
csl::TensorSpan<T> c = plan.permC.empty() ? output
: allocator.get_tensor_span<T>(plan.shapeC.begin(), plan.shapeC.end());
// reshape() cannot raise the rank, so the rank-3 GEMM views are built over the same memory
const std::size_t shapeA3[] = { plan.batch, rowsA, colsA };
const std::size_t shapeB3[] = { plan.batch, rowsB, colsB };
const std::size_t shapeC3[] = { plan.batch, plan.M, plan.N };
csl::TensorView<T> a3(a.get(), std::begin(shapeA3), std::end(shapeA3));
csl::TensorView<T> b3(b.get(), std::begin(shapeB3), std::end(shapeB3));
csl::TensorSpan<T> c3(c.get(), std::begin(shapeC3), std::end(shapeC3));
csl::tensor_ops::gemmStridedBatched<T>(cublasHandle, 0.0, c3, 1.0, plan.transA, a3, plan.transB, b3);
if (!plan.permC.empty())
kernels::permute<T>(stream, output, c, plan.permC);
}
std::size_t get_workspace_memory_in_bytes() const noexcept override { return scratch_mem_in_bytes; }
private:
// the operand itself, or a permuted copy in the workspace
csl::TensorView<T> operand(csl::WorkspaceAllocator& allocator, csl::TensorView<T> input,
const std::vector<std::size_t>& perm, const std::vector<std::size_t>& shape)
{
if (perm.empty())
return input;
auto permuted = allocator.get_tensor_span<T>(shape.begin(), shape.end());
kernels::permute<T>(stream, permuted, input, perm);
return permuted;
}
csl::Stream stream;
csl::cublas::Handle cublasHandle;
EinsumPlan plan;
std::size_t scratch_mem_in_bytes;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_EINSUM_HPP */

View File

@@ -16,6 +16,7 @@
#include <opencv2/core.hpp>
#include <cstddef>
#include <type_traits>
#include <vector>
#include <utility>
@@ -31,6 +32,10 @@ namespace cv { namespace dnn { namespace cuda4dnn {
MOD,
FMOD,
POW,
PRELU,
BIT_AND,
BIT_OR,
BIT_XOR,
};
class EltwiseOpBase : public CUDABackendNode {
@@ -89,6 +94,10 @@ namespace cv { namespace dnn { namespace cuda4dnn {
case EltwiseOpType::MOD: kernels::eltwise_mod_2<T>(stream, output, input_x, input_y); break;
case EltwiseOpType::FMOD: kernels::eltwise_fmod_2<T>(stream, output, input_x, input_y); break;
case EltwiseOpType::POW: kernels::eltwise_pow_2<T>(stream, output, input_x, input_y); break;
case EltwiseOpType::PRELU: kernels::eltwise_prelu_2<T>(stream, output, input_x, input_y); break;
case EltwiseOpType::BIT_AND:
case EltwiseOpType::BIT_OR:
case EltwiseOpType::BIT_XOR: bitwise(output, input_x, input_y); break;
}
} else if (inputs.size() == 1) {
auto input_0 = csl::viewOf<T>(inputs[0]);
@@ -123,10 +132,31 @@ namespace cv { namespace dnn { namespace cuda4dnn {
case EltwiseOpType::MOD: kernels::eltwise_mod_2<T>(stream, output, output, input); break;
case EltwiseOpType::FMOD: kernels::eltwise_fmod_2<T>(stream, output, output, input); break;
case EltwiseOpType::POW: kernels::eltwise_pow_2<T>(stream, output, output, input); break;
case EltwiseOpType::PRELU: CV_Error(Error::StsBadArg, "PRelu takes exactly two inputs");
case EltwiseOpType::BIT_AND:
case EltwiseOpType::BIT_OR:
case EltwiseOpType::BIT_XOR: bitwise(output, output, input); break;
}
}
}
}
private:
// the bitwise kernels exist for integer element types only
void bitwise(csl::TensorSpan<T> output, csl::TensorView<T> x, csl::TensorView<T> y)
{
if constexpr (std::is_integral<T>::value && !std::is_same<T, bool>::value) {
if (op == EltwiseOpType::BIT_AND)
kernels::eltwise_bitand_2<T>(stream, output, x, y);
else if (op == EltwiseOpType::BIT_OR)
kernels::eltwise_bitor_2<T>(stream, output, x, y);
else
kernels::eltwise_bitxor_2<T>(stream, output, x, y);
} else {
CV_UNUSED(output); CV_UNUSED(x); CV_UNUSED(y);
CV_Error(Error::StsNotImplemented, "bitwise ops need integer tensors");
}
}
};
}}} /* namespace cv::dnn::cuda4dnn */

View File

@@ -0,0 +1,71 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_GATHER_ELEMENTS_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_GATHER_ELEMENTS_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../kernels/gather_elements.hpp"
#include <opencv2/core.hpp>
#include <cstddef>
#include <cstdint>
#include <utility>
namespace cv { namespace dnn { namespace cuda4dnn {
template <class T>
class GatherElementsOp final : public CUDABackendNode {
public:
GatherElementsOp(csl::Stream stream_, int axis_)
: stream(std::move(stream_)), axis(axis_)
{
}
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(inputs.size() == 2 && outputs.size() == 1);
auto data = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
MatShape dataShape = cv::dnn::shape(inputs[0]);
MatShape indicesShape = cv::dnn::shape(inputs[1]);
const int dataDims = dataShape.dims;
const int axis_ = normalize_axis(axis, dataDims);
std::int64_t inner_size = 1;
for (int i = axis_ + 1; i < dataDims; i++)
inner_size *= dataShape[i];
const std::int64_t axis_size_data = dataShape[axis_];
const std::int64_t axis_size_indices = indicesShape[axis_];
csl::device::fast_divmod axis_inner_size(static_cast<int>(axis_size_indices * inner_size));
csl::device::fast_divmod inner_size_divmod(static_cast<int>(inner_size));
const std::size_t index_element_size = inputs[1].elemSize();
const void* indices = reinterpret_cast<const uchar*>(inputs[1].u->handle) + inputs[1].offset;
kernels::gather_elements<T>(stream, output, data, indices, index_element_size,
axis_size_data, inner_size, axis_inner_size, inner_size_divmod);
}
private:
csl::Stream stream;
int axis;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_GATHER_ELEMENTS_HPP */

View File

@@ -0,0 +1,104 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_GATHER_ND_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_GATHER_ND_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../kernels/gather_nd.hpp"
#include <opencv2/core.hpp>
#include <cstddef>
#include <cstdint>
#include <utility>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn {
template <class T>
class GatherNDOp final : public CUDABackendNode {
public:
GatherNDOp(csl::Stream stream_, int batch_dims_)
: stream(std::move(stream_)), batch_dims(batch_dims_)
{
}
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(inputs.size() == 2 && outputs.size() == 1);
auto data = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
const MatShape dataShape = cv::dnn::shape(inputs[0]);
const MatShape indicesShape = cv::dnn::shape(inputs[1]);
const int r = dataShape.dims;
const int q = indicesShape.dims;
CV_Assert(r >= 1 && q >= 1);
const std::int64_t last_index_dimension = indicesShape[q - 1];
CV_Assert(last_index_dimension >= 1 && batch_dims + last_index_dimension <= r);
std::vector<std::int64_t> data_strides(r);
data_strides[r - 1] = 1;
for (int i = r - 2; i >= 0; --i)
data_strides[i] = data_strides[i + 1] * dataShape[i + 1];
std::vector<std::int64_t> element_counts(last_index_dimension);
std::vector<std::int64_t> input_dims(last_index_dimension);
for (std::int64_t j = 0; j < last_index_dimension; ++j) {
element_counts[j] = data_strides[batch_dims + j];
input_dims[j] = dataShape[batch_dims + static_cast<int>(j)];
}
std::int64_t indices_total = 1;
for (int i = 0; i < q; i++)
indices_total *= indicesShape[i];
const std::int64_t outer_size = indices_total / last_index_dimension;
// an empty indices or output tensor is a legal no-op, not an error
if (outer_size <= 0 || output.size() == 0)
return;
const std::int64_t inner_size = static_cast<std::int64_t>(output.size()) / outer_size;
std::int64_t num_slices_per_batch = 1;
for (int d = batch_dims; d < q - 1; ++d)
num_slices_per_batch *= indicesShape[d];
if (num_slices_per_batch <= 0)
num_slices_per_batch = 1;
const std::int64_t batch_stride = batch_dims > 0 ? data_strides[batch_dims - 1] : 0;
const int idepth = inputs[1].depth();
if (idepth == CV_32S) {
kernels::gather_nd<T, std::int32_t>(stream, output, csl::viewOf<std::int32_t>(inputs[1]),
data, last_index_dimension, element_counts, input_dims,
inner_size, batch_stride, num_slices_per_batch);
} else if (idepth == CV_64S) {
kernels::gather_nd<T, std::int64_t>(stream, output, csl::viewOf<std::int64_t>(inputs[1]),
data, last_index_dimension, element_counts, input_dims,
inner_size, batch_stride, num_slices_per_batch);
} else {
CV_Error(Error::StsNotImplemented, "GatherND: indices must be CV_32S or CV_64S");
}
}
private:
csl::Stream stream;
int batch_dims;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_GATHER_ND_HPP */

View File

@@ -0,0 +1,61 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_GRID_SAMPLE_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_GRID_SAMPLE_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../kernels/grid_sample.hpp"
#include <opencv2/core.hpp>
#include <utility>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn {
template <class T>
class GridSampleOp final : public CUDABackendNode {
public:
GridSampleOp(csl::Stream stream_, const kernels::GridSampleParams& params_)
: stream(std::move(stream_)), params(params_) { }
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(inputs.size() == 2 && outputs.size() == 1);
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
auto grid = csl::viewOf<float>(inputs[1]);
MatShape inShape = cv::dnn::shape(inputs[0]);
MatShape outShape = cv::dnn::shape(outputs[0]);
CV_Assert(inShape.dims == 4 && outShape.dims == 4);
kernels::GridSampleParams p = params;
p.channels = inShape[1];
p.height = inShape[2];
p.width = inShape[3];
p.output_height = outShape[2];
p.output_width = outShape[3];
kernels::grid_sample<T>(stream, output, input, grid, p);
}
private:
csl::Stream stream;
kernels::GridSampleParams params;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_GRID_SAMPLE_HPP */

View File

@@ -28,27 +28,27 @@ namespace cv { namespace dnn { namespace cuda4dnn {
public:
using wrapper_type = GetCUDABackendWrapperType<T>;
GroupNormOp(csl::Stream stream_, float epsilon_, size_t loops, size_t num_groups)
GroupNormOp(csl::Stream stream_, float epsilon_, size_t loops, size_t num_groups, size_t group_elems)
: stream(std::move(stream_)), epsilon(epsilon_), num_groups(num_groups) {
reduce_scratch = kernels::reduce_mean_sqr_sum_scratch(loops, group_elems);
csl::WorkspaceBuilder builder;
builder.require<float>(loops * num_groups); // mean and stdev for each group
builder.require<float>(loops * num_groups);
if (reduce_scratch)
builder.require<float>(reduce_scratch);
scratch_mem_in_bytes = builder.required_workspace_size();
}
void forward(const std::vector<cv::Ptr<BackendWrapper>>& inputs,
const std::vector<cv::Ptr<BackendWrapper>>& outputs,
void forward(const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override {
auto input_wrapper = inputs[0].dynamicCast<wrapper_type>();
auto scale_wrapper = inputs[1].dynamicCast<wrapper_type>();
auto bias_wrapper = inputs[2].dynamicCast<wrapper_type>();
CV_Assert(inputs.size() == 3 && outputs.size() == 1);
auto input = input_wrapper->getView();
auto scale = scale_wrapper->getView();
auto bias = bias_wrapper->getView();
auto input = csl::viewOf<T>(inputs[0]);
auto scale = csl::viewOf<T>(inputs[1]);
auto bias = csl::viewOf<T>(inputs[2]);
auto output_wrapper = outputs[0].dynamicCast<wrapper_type>();
auto output = output_wrapper->getSpan();
auto output = csl::spanOf<T>(outputs[0]);
auto C = input.get_axis_size(1);
auto loops = input.size_range(0, 2);
@@ -64,8 +64,9 @@ namespace cv { namespace dnn { namespace cuda4dnn {
// reduce_mean_sqr_sum writes every row, so no zero-fill is needed
auto mean = ws_allocator.get_span<float>(loops / group_size);
auto stdev = ws_allocator.get_span<float>(loops / group_size);
auto partials = reduce_scratch ? ws_allocator.get_span<float>(reduce_scratch) : csl::Span<float>();
kernels::reduce_mean_sqr_sum<T>(stream, mean, stdev, input, norm_size * group_size);
kernels::reduce_mean_sqr_sum<T>(stream, mean, stdev, input, norm_size * group_size, partials);
kernels::compute_normalization_scale(stream, stdev, mean, stdev, norm_size * group_size, epsilon);
kernels::normalize_mean_variance_groupwise<T>(stream, output, input, scale, bias, mean, stdev, norm_size, C, num_groups, group_size);
}
@@ -77,6 +78,7 @@ namespace cv { namespace dnn { namespace cuda4dnn {
csl::Stream stream;
float epsilon;
std::size_t num_groups;
std::size_t reduce_scratch;
std::size_t scratch_mem_in_bytes;
};

View File

@@ -28,27 +28,27 @@ namespace cv { namespace dnn { namespace cuda4dnn {
public:
using wrapper_type = GetCUDABackendWrapperType<T>;
InstanceNormOp(csl::Stream stream_, float epsilon_, size_t loops)
InstanceNormOp(csl::Stream stream_, float epsilon_, size_t loops, size_t norm_size)
: stream(std::move(stream_)), epsilon(epsilon_) {
reduce_scratch = kernels::reduce_mean_sqr_sum_scratch(loops, norm_size);
csl::WorkspaceBuilder builder;
builder.require<float>(loops);
builder.require<float>(loops);
if (reduce_scratch)
builder.require<float>(reduce_scratch);
scratch_mem_in_bytes = builder.required_workspace_size();
}
void forward(const std::vector<cv::Ptr<BackendWrapper>>& inputs,
const std::vector<cv::Ptr<BackendWrapper>>& outputs,
void forward(const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override {
auto input_wrapper = inputs[0].dynamicCast<wrapper_type>();
auto scale_wrapper = inputs[1].dynamicCast<wrapper_type>();
auto bias_wrapper = inputs[2].dynamicCast<wrapper_type>();
CV_Assert(inputs.size() == 3 && outputs.size() == 1);
auto input = input_wrapper->getView();
auto scale = scale_wrapper->getView();
auto bias = bias_wrapper->getView();
auto input = csl::viewOf<T>(inputs[0]);
auto scale = csl::viewOf<T>(inputs[1]);
auto bias = csl::viewOf<T>(inputs[2]);
auto output_wrapper = outputs[0].dynamicCast<wrapper_type>();
auto output = output_wrapper->getSpan();
auto output = csl::spanOf<T>(outputs[0]);
auto C = input.get_axis_size(1);
auto loops = input.size_range(0, 2);
@@ -62,8 +62,9 @@ namespace cv { namespace dnn { namespace cuda4dnn {
// reduce_mean_sqr_sum writes every row, so no zero-fill is needed
auto mean = ws_allocator.get_span<float>(loops);
auto stdev = ws_allocator.get_span<float>(loops);
auto partials = reduce_scratch ? ws_allocator.get_span<float>(reduce_scratch) : csl::Span<float>();
kernels::reduce_mean_sqr_sum<T>(stream, mean, stdev, input, norm_size);
kernels::reduce_mean_sqr_sum<T>(stream, mean, stdev, input, norm_size, partials);
kernels::compute_normalization_scale(stream, stdev, mean, stdev, norm_size, epsilon);
kernels::normalize_mean_variance_channelwise<T>(stream, output, input, scale, bias, mean, stdev, norm_size, C);
}
@@ -76,6 +77,7 @@ namespace cv { namespace dnn { namespace cuda4dnn {
float epsilon;
std::size_t reduce_scratch;
std::size_t scratch_mem_in_bytes;
};

View File

@@ -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<T>(stream, output, 0.f);
return;
} else {
@@ -114,7 +115,9 @@ 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;
// the optional Mean/InvStdDev outputs still need the statistics
if (norm_size == 1 && outputs.size() == 1 && bias.empty()) {
kernels::fill<T>(stream, output, 0.f);
return;
}
@@ -131,6 +134,12 @@ namespace cv { namespace dnn { namespace cuda4dnn {
kernels::normalize_mean_variance_layernorm<T>(stream, output, input, scale, bias, mean, inv_stddev, norm_size);
else
kernels::normalize_mean_variance_layernorm<T>(stream, output, input, scale, mean, inv_stddev, norm_size);
// ONNX Mean and InvStdDev are float per row, which is what the scratch holds
if (outputs.size() >= 2 && !outputs[1].empty())
csl::memcpy<float>(csl::spanOf<float>(outputs[1]).get(), mean.data(), loops, stream);
if (outputs.size() >= 3 && !outputs[2].empty())
csl::memcpy<float>(csl::spanOf<float>(outputs[2]).get(), inv_stddev.data(), loops, stream);
}
std::size_t get_workspace_memory_in_bytes() const noexcept override { return scratch_mem_in_bytes; }

View File

@@ -10,8 +10,11 @@
#include "../csl/cudnn.hpp"
#include "../csl/tensor_ops.hpp"
#include "../kernels/normalize.hpp"
#include <cstddef>
#include <utility>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn {
@@ -70,6 +73,38 @@ namespace cv { namespace dnn { namespace cuda4dnn {
std::size_t scratch_mem_in_bytes;
};
// across-channel LRN without legacy cuDNN; matches the CPU ChannelLRN
template <class T>
class LRNAcrossChannelsOp final : public CUDABackendNode {
public:
LRNAcrossChannelsOp(csl::Stream stream_, int local_size, float alpha_scaled_, float beta_, float bias_)
: stream(std::move(stream_)), radius((local_size - 1) / 2), alpha_scaled(alpha_scaled_), beta(beta_), bias(bias_) { }
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(inputs.size() == outputs.size());
for (size_t i = 0; i < inputs.size(); i++) {
auto input = csl::viewOf<T>(inputs[i]);
auto output = csl::spanOf<T>(outputs[i]);
CV_Assert(input.rank() >= 2);
const std::size_t batch = input.get_axis_size(0), channels = input.get_axis_size(1);
const std::size_t plane = input.size_range(2, input.rank());
kernels::lrn_across_channels<T>(stream, csl::Span<T>(output.get(), output.size()),
csl::View<T>(input.get(), input.size()),
batch, channels, plane, radius, alpha_scaled, bias, beta);
}
}
private:
csl::Stream stream;
int radius;
float alpha_scaled, beta, bias;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_LRN_HPP */

View File

@@ -87,17 +87,14 @@ namespace cv { namespace dnn { namespace cuda4dnn {
}
void forward(
const std::vector<cv::Ptr<BackendWrapper>>& inputs,
const std::vector<cv::Ptr<BackendWrapper>>& outputs,
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_Assert(inputs.size() == 1 && outputs.size() == 1);
auto input_wrapper = inputs[0].dynamicCast<wrapper_type>();
auto input = input_wrapper->getView();
auto output_wrapper = outputs[0].dynamicCast<wrapper_type>();
auto output = output_wrapper->getSpan();
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
std::size_t outer_size = input.size_range(0, axis_start);
std::size_t mid_size = input.size_range(axis_start, axis_end);

View File

@@ -0,0 +1,76 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_ONE_HOT_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_ONE_HOT_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../kernels/one_hot.hpp"
#include <opencv2/core.hpp>
#include <cstddef>
#include <vector>
#include <utility>
namespace cv { namespace dnn { namespace cuda4dnn {
template <class T, class TIdx>
class OneHotOp final : public CUDABackendNode {
public:
OneHotOp(csl::Stream stream_, int axis_, double off_, double on_, bool runtimeValues_)
: stream(std::move(stream_)), axis(axis_),
offValue(static_cast<T>(off_)), onValue(static_cast<T>(on_)),
runtimeValues(runtimeValues_)
{
}
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(inputs.size() == 3 && outputs.size() == 1);
MatShape outShape = cv::dnn::shape(outputs[0]);
const int rank = outShape.dims;
const int a = axis < 0 ? axis + rank : axis;
CV_Assert(0 <= a && a < rank);
if (outputs[0].total() == 0)
return;
if (runtimeValues && !inputs[2].empty()) {
Mat values, values64;
inputs[2].copyTo(values);
CV_Assert(values.total() == 2);
values.convertTo(values64, CV_64F);
offValue = static_cast<T>(values64.ptr<double>()[0]);
onValue = static_cast<T>(values64.ptr<double>()[1]);
}
std::size_t inner = 1;
for (int i = a + 1; i < rank; i++)
inner *= static_cast<std::size_t>(outShape[i]);
auto indices = csl::viewOf<TIdx>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
kernels::one_hot<T, TIdx>(stream, output, indices, static_cast<std::size_t>(outShape[a]), inner,
offValue, onValue);
}
private:
csl::Stream stream;
int axis;
T offValue, onValue;
bool runtimeValues;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_ONE_HOT_HPP */

View File

@@ -17,6 +17,7 @@
#include <opencv2/core.hpp>
#include <cstddef>
#include <cstring>
#include <vector>
#include <algorithm>
#include <utility>
@@ -113,6 +114,37 @@ namespace cv { namespace dnn { namespace cuda4dnn {
std::vector<cv::Range> dstRanges;
};
// ONNX Pad: any rank, constant/reflect/edge/wrap, negative pads crop
template <class T>
class Pad2Op final : public CUDABackendNode {
public:
// value_bits holds the pad value already converted to the tensor's element type
Pad2Op(csl::Stream stream_, std::vector<int> pads_begin_, kernels::PadMode mode_, const void* value_bits)
: stream(std::move(stream_)), pads_begin(std::move(pads_begin_)), mode(mode_)
{
std::memcpy(&value, value_bits, sizeof(T));
}
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(!inputs.empty() && outputs.size() == 1);
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
kernels::pad_nd<T>(stream, output, input, pads_begin, mode, value);
}
private:
csl::Stream stream;
std::vector<int> pads_begin;
kernels::PadMode mode;
T value;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_PADDING_HPP */

View File

@@ -17,9 +17,8 @@
#include "../csl/tensor.hpp"
#include "../csl/tensor_ops.hpp"
#if defined(HAVE_CUDNNJIT) && !defined(HAVE_CUDNN)
// max pooling with left padding uses the hand kernel in every build
#include "../kernels/max_unpooling.hpp"
#endif
#include <opencv2/core.hpp>
@@ -181,13 +180,21 @@ namespace cv { namespace dnn { namespace cuda4dnn {
if (std::any_of(std::begin(padding_left), std::end(padding_left), is_not_zero))
{
#if !defined(HAVE_CUDNNJIT) || defined(HAVE_CUDNN)
/* there is padding on the left and we are forced to transform */
auto transformed_input_shape = input_shape;
for (int i = 0; i < rank; i++)
transformed_input_shape[i] += padding_left[i] + padding_right[i];
// the transform pads with zeros, which wins a max over negative borders; the index kernel skips them
if (config.poolMode == PoolingConfiguration::PoolingMode::MAX)
{
useIndexKernel = true;
}
else
{
/* there is padding on the left and we are forced to transform */
auto transformed_input_shape = input_shape;
for (int i = 0; i < rank; i++)
transformed_input_shape[i] += padding_left[i] + padding_right[i];
transformedInput.resize(std::begin(transformed_input_shape), std::end(transformed_input_shape));
inputTransformer = csl::TensorTransform<T>(cudnnHandle, padding_left, padding_right);
transformedInput.resize(std::begin(transformed_input_shape), std::end(transformed_input_shape));
inputTransformer = csl::TensorTransform<T>(cudnnHandle, padding_left, padding_right);
}
#endif
}
}
@@ -237,12 +244,20 @@ namespace cv { namespace dnn { namespace cuda4dnn {
#if defined(HAVE_CUDNNJIT) && !defined(HAVE_CUDNN)
// JIT builds route average pooling to AveragePoolingOp, so only MAX reaches here
CV_Assert(config.poolMode == PoolingConfiguration::PoolingMode::MAX);
poolWindow.assign(std::begin(window_size), std::end(window_size));
poolStrides.assign(std::begin(strides), std::end(strides));
for (std::size_t i = 0; i < window_size.size(); i++)
poolPadsBegin.push_back(common_padding[i + 2] + padding_left[i + 2]);
#else
pooler = csl::Pooling<T>(cudnnHandle, params);
useIndexKernel = true;
#endif
if (useIndexKernel)
{
poolWindow.assign(std::begin(window_size), std::end(window_size));
poolStrides.assign(std::begin(strides), std::end(strides));
for (std::size_t i = 0; i < window_size.size(); i++)
poolPadsBegin.push_back(common_padding[i + 2] + padding_left[i + 2]);
}
#if !defined(HAVE_CUDNNJIT) || defined(HAVE_CUDNN)
else
{
pooler = csl::Pooling<T>(cudnnHandle, params);
}
#endif
}
@@ -256,11 +271,14 @@ namespace cv { namespace dnn { namespace cuda4dnn {
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
#if defined(HAVE_CUDNNJIT) && !defined(HAVE_CUDNN)
CV_UNUSED(workspace);
kernels::max_pooling_with_indices<T, std::int64_t>(stream, output, csl::TensorSpan<std::int64_t>(), input,
poolWindow, poolStrides, poolPadsBegin);
#else
if (useIndexKernel)
{
kernels::max_pooling_with_indices<T, std::int64_t>(stream, output, csl::TensorSpan<std::int64_t>(), input,
poolWindow, poolStrides, poolPadsBegin);
return;
}
#if !defined(HAVE_CUDNNJIT) || defined(HAVE_CUDNN)
if (!transformedInput.empty())
{
inputTransformer.transform(input, transformedInput);
@@ -275,9 +293,10 @@ namespace cv { namespace dnn { namespace cuda4dnn {
csl::Stream stream;
csl::cudnn::Handle cudnnHandle;
#if defined(HAVE_CUDNNJIT) && !defined(HAVE_CUDNN)
// max pooling through the hand kernel: always in JIT builds, for left padding in full cuDNN builds
bool useIndexKernel = false;
std::vector<std::size_t> poolWindow, poolStrides, poolPadsBegin;
#else
#if !defined(HAVE_CUDNNJIT) || defined(HAVE_CUDNN)
csl::Pooling<T> pooler;
csl::Tensor<T> transformedInput;

View File

@@ -0,0 +1,138 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_QUANTIZE_DEQUANTIZE_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_QUANTIZE_DEQUANTIZE_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../csl/workspace.hpp"
#include "../kernels/quantize_dequantize.hpp"
#include <opencv2/dnn/shape_utils.hpp>
#include <cstdint>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn {
// Per-tensor and per-axis only (block_size == 0); gated by the layer's supportBackend.
template <class T>
class DequantizeLinearOp final : public CUDABackendNode {
public:
DequantizeLinearOp(csl::Stream stream_, int axis_) : stream(std::move(stream_)), axis(axis_) { }
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(inputs.size() == 2 || inputs.size() == 3);
CV_Assert(outputs.size() == 1);
auto input = csl::viewOf<T>(inputs[0]);
auto scale = csl::viewOf<float>(inputs[1]);
auto output = csl::spanOf<float>(outputs[0]);
csl::View<T> zero_point = inputs.size() == 3 ? csl::viewOf<T>(inputs[2]) : csl::View<T>();
CV_Assert(zero_point.empty() || zero_point.size() == scale.size());
MatShape inpshape = shape(inputs[0]);
int ndims = inpshape.dims;
int ax = ndims > 0 ? normalize_axis(axis, ndims) : 0;
std::size_t sz_a = scale.size() > 1 ? (std::size_t)inpshape[ax] : 1;
CV_Assert(scale.size() == 1 || scale.size() == sz_a);
std::size_t slice_size = 1;
if (sz_a > 1)
for (int i = ax + 1; i < ndims; i++)
slice_size *= inpshape[i];
kernels::dequantize_linear<T>(stream, output, input, scale, zero_point, sz_a, slice_size);
}
private:
csl::Stream stream;
int axis;
};
template <class T>
class QuantizeLinearOp final : public CUDABackendNode {
public:
QuantizeLinearOp(csl::Stream stream_, int axis_) : stream(std::move(stream_)), axis(axis_) { }
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(inputs.size() == 2 || inputs.size() == 3);
CV_Assert(outputs.size() == 1);
auto input = csl::viewOf<float>(inputs[0]);
auto scale = csl::viewOf<float>(inputs[1]);
auto output = csl::spanOf<T>(outputs[0]);
csl::View<T> zero_point = inputs.size() == 3 ? csl::viewOf<T>(inputs[2]) : csl::View<T>();
CV_Assert(zero_point.empty() || zero_point.size() == scale.size());
MatShape inpshape = shape(inputs[0]);
int ndims = inpshape.dims;
int ax = ndims > 0 ? normalize_axis(axis, ndims) : 0;
std::size_t sz_a = scale.size() > 1 ? (std::size_t)inpshape[ax] : 1;
CV_Assert(scale.size() == 1 || scale.size() == sz_a);
std::size_t slice_size = 1;
if (sz_a > 1)
for (int i = ax + 1; i < ndims; i++)
slice_size *= inpshape[i];
kernels::quantize_linear<T>(stream, output, input, scale, zero_point, sz_a, slice_size);
}
private:
csl::Stream stream;
int axis;
};
// ONNX DynamicQuantizeLinear: float32 x -> uint8 y, scalar y_scale and y_zero_point, all on the device.
class DynamicQuantizeLinearOp final : public CUDABackendNode {
public:
DynamicQuantizeLinearOp(csl::Stream stream_, std::size_t total) : stream(std::move(stream_))
{
partials = kernels::dynamic_quantize_partials(total);
csl::WorkspaceBuilder builder;
builder.require<float>(2 * partials);
scratch_mem_in_bytes = builder.required_workspace_size();
}
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_Assert(inputs.size() == 1 && outputs.size() == 3);
auto input = csl::viewOf<float>(inputs[0]);
auto output = csl::spanOf<std::uint8_t>(outputs[0]);
auto scale = csl::spanOf<float>(outputs[1]);
auto zero_point = csl::spanOf<std::uint8_t>(outputs[2]);
auto allocator = csl::WorkspaceAllocator(workspace);
auto scratch = allocator.get_span<float>(2 * partials);
kernels::dynamic_quantize_scale(stream, scale, zero_point, input, scratch);
// same per-tensor formula as QuantizeLinear, with the scale and zero point read on the device
kernels::quantize_linear<std::uint8_t>(stream, output, input, scale, zero_point, 1, 1);
}
std::size_t get_workspace_memory_in_bytes() const noexcept override { return scratch_mem_in_bytes; }
private:
csl::Stream stream;
std::size_t partials;
std::size_t scratch_mem_in_bytes;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_QUANTIZE_DEQUANTIZE_HPP */

View File

@@ -0,0 +1,52 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_RANGE_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_RANGE_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../kernels/range.hpp"
#include <opencv2/core.hpp>
#include <utility>
namespace cv { namespace dnn { namespace cuda4dnn {
template <class T>
class RangeOp final : public CUDABackendNode {
public:
RangeOp(csl::Stream stream_) : stream(std::move(stream_)) { }
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(inputs.size() == 3 && outputs.size() == 1);
CV_Assert(inputs[0].total() == 1 && inputs[2].total() == 1);
CV_Assert(inputs[0].u && inputs[2].u);
// start/delta are dereferenced inside the kernel: no host round trip, no stall
auto devPtr = [](const UMat& m) {
return reinterpret_cast<const T*>(
reinterpret_cast<const uchar*>(m.u->handle) + m.offset);
};
auto output = csl::spanOf<T>(outputs[0]);
kernels::range<T>(stream, output, devPtr(inputs[0]), devPtr(inputs[2]));
}
private:
csl::Stream stream;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_RANGE_HPP */

View File

@@ -10,8 +10,10 @@
#include "../csl/stream.hpp"
#include "../kernels/resize.hpp"
#include "../csl/span.hpp"
#include <utility>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn {
@@ -40,18 +42,15 @@ namespace cv { namespace dnn { namespace cuda4dnn {
}
void forward(
const std::vector<cv::Ptr<BackendWrapper>>& inputs,
const std::vector<cv::Ptr<BackendWrapper>>& outputs,
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
// sometimes the target shape is taken from the second input; we don't use it however
CV_Assert((inputs.size() == 1 || inputs.size() == 2) && outputs.size() == 1);
// roi/scales/sizes may follow the data input; the output shape already carries them
CV_Assert(inputs.size() >= 1 && outputs.size() == 1);
auto input_wrapper = inputs[0].dynamicCast<wrapper_type>();
auto input = input_wrapper->getView();
auto output_wrapper = outputs[0].dynamicCast<wrapper_type>();
auto output = output_wrapper->getSpan();
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
const auto compute_scale = [this](std::size_t input_size, std::size_t output_size) {
return (align_corners && output_size > 1) ?
@@ -76,6 +75,63 @@ namespace cv { namespace dnn { namespace cuda4dnn {
bool align_corners, half_pixel_centers;
};
// per-axis taps built on the host by the CPU Resize2 builders; idx -1 skips a tap
struct ResizeTables {
int taps_y = 1, taps_x = 1;
std::vector<int> idx_y, idx_x, oob_y, oob_x;
std::vector<float> w_y, w_x;
float extrapolation_value = 0.f;
};
// every ONNX Resize mode as one gather; the tables carry the coordinate and filter math
template <class T>
class ResizeTablesOp final : public CUDABackendNode {
public:
ResizeTablesOp(csl::Stream stream_, const ResizeTables& t)
: stream(std::move(stream_)), taps_y(t.taps_y), taps_x(t.taps_x), extrapolation(t.extrapolation_value)
{
idx_y = upload(t.idx_y); w_y = upload(t.w_y); oob_y = upload(t.oob_y);
idx_x = upload(t.idx_x); w_x = upload(t.w_x); oob_x = upload(t.oob_x);
}
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() == 1);
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
kernels::resize_separable<T>(stream, output, input,
view<int>(idx_y), view<float>(w_y), view<int>(oob_y), taps_y,
view<int>(idx_x), view<float>(w_x), view<int>(oob_x), taps_x, extrapolation);
}
private:
// device copy on the CUDA UMat allocator, like every other device tensor
template <class V>
static UMat upload(const std::vector<V>& v)
{
CV_Assert(!v.empty());
Mat host(1, (int)v.size(), DataType<V>::type, const_cast<V*>(v.data()));
UMat u;
u.allocator = cv::cuda::getCudaAllocator();
u.fitSameSize(host, host.type());
checkCudaAllocated(u);
host.copyTo(u);
return u;
}
template <class V>
static csl::View<V> view(const UMat& u) { return csl::View<V>(csl::viewOf<V>(u).get(), u.total()); }
csl::Stream stream;
int taps_y, taps_x;
float extrapolation;
UMat idx_y, oob_y, idx_x, oob_x, w_y, w_x;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_RESIZE_HPP */

View File

@@ -0,0 +1,61 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_RMS_NORM_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_RMS_NORM_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../kernels/rms_norm.hpp"
#include <opencv2/core.hpp>
#include <opencv2/dnn/shape_utils.hpp>
#include <cstddef>
#include <utility>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn {
template <class T>
class RMSNormOp final : public CUDABackendNode {
public:
RMSNormOp(csl::Stream stream_, int axis_, float epsilon_)
: stream(std::move(stream_)), axis(axis_), epsilon(epsilon_) { }
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(inputs.size() >= 2 && outputs.size() == 1);
auto input = csl::viewOf<T>(inputs[0]);
auto scale = csl::viewOf<T>(inputs[1]);
auto output = csl::spanOf<T>(outputs[0]);
MatShape inShape = cv::dnn::shape(inputs[0]);
const int a = normalize_axis(axis, inShape.dims);
CV_Assert(a >= 0 && a < inShape.dims);
std::size_t loops = 1, norm_size = 1;
for (int i = 0; i < a; i++) loops *= static_cast<std::size_t>(inShape[i]);
for (int i = a; i < inShape.dims; i++) norm_size *= static_cast<std::size_t>(inShape[i]);
kernels::rms_norm<T>(stream, output, input, scale, loops, norm_size, epsilon);
}
private:
csl::Stream stream;
int axis;
float epsilon;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_RMS_NORM_HPP */

View File

@@ -0,0 +1,68 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_ROI_ALIGN_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_ROI_ALIGN_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../kernels/roi_align.hpp"
#include <opencv2/core.hpp>
#include <cstddef>
#include <utility>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn {
template <class T>
class RoiAlignOp final : public CUDABackendNode {
public:
RoiAlignOp(csl::Stream stream_, const kernels::RoiAlignParams& params_)
: stream(std::move(stream_)), params(params_) { }
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(inputs.size() == 3 && outputs.size() == 1);
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
auto rois = csl::viewOf<float>(inputs[1]);
MatShape inShape = cv::dnn::shape(inputs[0]);
CV_Assert(inShape.dims == 4);
/* The spatial extents are only known once shapes are resolved, so the
* 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];
const std::size_t batch_index_element_size = inputs[2].elemSize();
// device pointer, same idiom the gather primitive uses for its index tensor
const void* batch_indices = reinterpret_cast<const uchar*>(inputs[2].u->handle) + inputs[2].offset;
kernels::roi_align<T>(stream, output, input, rois,
batch_indices, batch_index_element_size, p);
}
private:
csl::Stream stream;
kernels::RoiAlignParams params;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_ROI_ALIGN_HPP */

View File

@@ -0,0 +1,87 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_SCATTER_ELEMENTS_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_SCATTER_ELEMENTS_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../kernels/scatter_elements.hpp"
#include "../kernels/fill_copy.hpp"
#include <opencv2/core.hpp>
#include <cstdint>
#include <utility>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn {
template <class T>
class ScatterElementsOp final : public CUDABackendNode {
public:
ScatterElementsOp(csl::Stream stream_, int axis_)
: stream(std::move(stream_)), axis(axis_)
{
}
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(inputs.size() == 3 && outputs.size() == 1);
auto data = csl::viewOf<T>(inputs[0]);
auto updates = csl::viewOf<T>(inputs[2]);
auto output = csl::spanOf<T>(outputs[0]);
// ScatterElements writes a sparse subset, so the untouched elements come from data
if (output.get() != data.get())
kernels::copy<T>(stream, output, data);
const MatShape dataShape = cv::dnn::shape(inputs[0]);
const MatShape idxShape = cv::dnn::shape(inputs[1]);
const int rank = dataShape.dims;
CV_Assert(rank == idxShape.dims && rank >= 1);
const int axis_ = normalize_axis(axis, rank);
std::vector<std::int64_t> output_strides(rank);
output_strides[rank - 1] = 1;
for (int i = rank - 2; i >= 0; --i)
output_strides[i] = output_strides[i + 1] * dataShape[i + 1];
std::vector<std::int64_t> indices_dims(rank);
for (int i = 0; i < rank; i++)
indices_dims[i] = idxShape[i];
const std::int64_t axis_dim_size = dataShape[axis_];
const int idepth = inputs[1].depth();
if (idepth == CV_32S) {
kernels::scatter_elements<T, std::int32_t>(stream, output,
csl::viewOf<std::int32_t>(inputs[1]), updates,
rank, axis_, indices_dims, output_strides, axis_dim_size);
} else if (idepth == CV_64S) {
kernels::scatter_elements<T, std::int64_t>(stream, output,
csl::viewOf<std::int64_t>(inputs[1]), updates,
rank, axis_, indices_dims, output_strides, axis_dim_size);
} else {
CV_Error(Error::StsNotImplemented, "ScatterElements: indices must be CV_32S or CV_64S");
}
}
private:
csl::Stream stream;
int axis;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_SCATTER_ELEMENTS_HPP */

View File

@@ -0,0 +1,88 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_SCATTER_ND_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_SCATTER_ND_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../kernels/scatter_nd.hpp"
#include "../kernels/fill_copy.hpp"
#include <opencv2/core.hpp>
#include <cstddef>
#include <cstdint>
#include <vector>
#include <utility>
namespace cv { namespace dnn { namespace cuda4dnn {
template <class T, class TIdx>
class ScatterNDOp final : public CUDABackendNode {
public:
ScatterNDOp(csl::Stream stream_)
: stream(std::move(stream_))
{
}
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(inputs.size() == 3 && outputs.size() == 1);
if (inputs[0].total() == 0)
return;
auto data = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
kernels::copy<T>(stream, output, data);
if (inputs[1].total() == 0 || inputs[2].total() == 0)
return;
auto indices = csl::viewOf<TIdx>(inputs[1]);
auto updates = csl::viewOf<T>(inputs[2]);
MatShape dataShape = cv::dnn::shape(inputs[0]);
MatShape indShape = cv::dnn::shape(inputs[1]);
const int r = dataShape.dims;
const int q = indShape.dims;
const std::int64_t last_index_dimension = indShape[q - 1];
std::vector<std::int64_t> element_counts(last_index_dimension), input_dims(last_index_dimension);
for (std::int64_t j = 0; j < last_index_dimension; j++) {
std::int64_t stride = 1;
for (int i = static_cast<int>(j) + 1; i < r; i++)
stride *= dataShape[i];
element_counts[j] = stride;
input_dims[j] = dataShape[static_cast<int>(j)];
}
std::size_t num_indices = 1;
for (int i = 0; i < q - 1; i++)
num_indices *= static_cast<std::size_t>(indShape[i]);
std::size_t num_updates_elements = 1;
for (int i = static_cast<int>(last_index_dimension); i < r; i++)
num_updates_elements *= static_cast<std::size_t>(dataShape[i]);
kernels::scatter_nd<T, TIdx>(stream, output, indices, updates,
num_indices, last_index_dimension, element_counts, input_dims, num_updates_elements);
}
private:
csl::Stream stream;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_SCATTER_ND_HPP */

View File

@@ -0,0 +1,47 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_SIZE_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_SIZE_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../kernels/fill_copy.hpp"
#include <opencv2/core.hpp>
#include <cstdint>
#include <utility>
namespace cv { namespace dnn { namespace cuda4dnn {
class SizeOp final : public CUDABackendNode {
public:
SizeOp(csl::Stream stream_) : stream(std::move(stream_)) { }
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() == 1);
auto output = csl::spanOf<std::int64_t>(outputs[0]);
CV_Assert(output.size() == 1);
const std::int64_t n = static_cast<std::int64_t>(inputs[0].total());
kernels::fill<std::int64_t>(stream, output, n);
}
private:
csl::Stream stream;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_SIZE_HPP */

View File

@@ -0,0 +1,75 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_TOPK_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_TOPK_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../csl/workspace.hpp"
#include "../kernels/topk.hpp"
#include <opencv2/core.hpp>
#include <opencv2/dnn/shape_utils.hpp>
#include <cstdint>
#include <utility>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn {
template <class T>
class TopKOp final : public CUDABackendNode {
public:
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_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]);
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.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;
bool largest;
int k, outer = 1, inner = 1, dim_axis = 1;
std::size_t sort_bytes = 0;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_TOPK_HPP */

View File

@@ -0,0 +1,60 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_TRILU_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_TRILU_HPP
#include "../../op_cuda.hpp"
#include "../csl/stream.hpp"
#include "../kernels/trilu.hpp"
#include <opencv2/core.hpp>
#include <utility>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn {
template <class T>
class TriluOp final : public CUDABackendNode {
public:
TriluOp(csl::Stream stream_, int k_, bool upper_)
: stream(std::move(stream_)), k(k_), upper(upper_) { }
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() == 1);
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
MatShape inShape = cv::dnn::shape(inputs[0]);
const int rank = inShape.dims;
CV_Assert(rank >= 1);
const int width = inShape[rank - 1];
const int height = (rank >= 2) ? inShape[rank - 2] : 1;
int loops = 1;
for (int i = 0; i < rank - 2; i++) loops *= inShape[i];
kernels::trilu<T>(stream, output, input, loops, height, width, k, upper);
}
private:
csl::Stream stream;
int k;
bool upper;
};
}}} /* namespace cv::dnn::cuda4dnn */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_TRILU_HPP */

View File

@@ -180,6 +180,9 @@ static inline void prindent(std::ostream& strm, int indent)
strm << ' ';
}
// Test hook: ops a new-engine net set to DNN_BACKEND_CUDA runs on the CPU, as "name (type)", after the first forward()
CV_EXPORTS std::vector<std::string> cudaFallbackOps(const Net& net);
CV__DNN_INLINE_NS_END
namespace accessor {

View File

@@ -4,7 +4,11 @@
#include "../precomp.hpp"
#include "layers_common.hpp"
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#ifdef HAVE_CUDA
#include "../cuda4dnn/primitives/arg.hpp"
#endif
namespace cv { namespace dnn {
@@ -43,9 +47,31 @@ public:
virtual bool supportBackend(int backendId) CV_OVERRIDE
{
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
if (!netimpl_ || this->inputs.size() != 1)
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);
}
#endif
return backendId == DNN_BACKEND_OPENCV;
}
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void* context_,
InputArrayOfArrays inputs_arr,
InputArrayOfArrays) CV_OVERRIDE
{
auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
return make_cuda_node_with_type<cuda4dnn::ArgOp>(preferableTarget, inputs_arr.type(0),
std::move(context->stream), axis,
op == ArgOp::MAX, select_last_index);
}
#endif
void handleKeepDims(MatShape& shape, const int axis_) const
{
if (keepdims)

View File

@@ -40,6 +40,7 @@
//
//M*/
#include "../precomp.hpp"
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#include "../op_inf_engine.hpp"
#include "../ie_ngraph.hpp"
@@ -67,6 +68,17 @@ public:
#ifdef HAVE_INF_ENGINE
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
return true;
#endif
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
// classic-engine nets carry no graph args to inspect
if (!netimpl_ || inputs.empty())
return true;
// 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 || isCudaNodeType(t);
}
#endif
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_CUDA ||
@@ -179,15 +191,16 @@ public:
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(
void *context_,
const std::vector<Ptr<BackendWrapper>>& inputs,
const std::vector<Ptr<BackendWrapper>>& outputs
InputArrayOfArrays inputs_arr,
InputArrayOfArrays
) override
{
auto context = reinterpret_cast<csl::CSLContext*>(context_);
if (inputs[0]->getHostMatDepth() == CV_Bool)
const int depth = inputs_arr.depth(0);
if (depth == CV_Bool)
return make_cuda_node_bool<cuda4dnn::ReshapeOp>(std::move(context->stream));
else
return make_cuda_node_with_type<cuda4dnn::ReshapeOp>(preferableTarget, inputs[0]->getHostMatDepth(), std::move(context->stream));
return make_cuda_node_with_type<cuda4dnn::ReshapeOp>(preferableTarget, depth, std::move(context->stream));
}
#endif
};

View File

@@ -115,9 +115,13 @@ public:
if (inType >= 0 && inType == outType)
return isCudaNodeType(inType);
// an intermediate arg has no type until the first forward; check only the known side(s)
auto castable = [](int t) { return t < 0 || t == CV_Bool || t == CV_32S || t == CV_64S || t == CV_32F; };
auto castable = [](int t) { return t < 0 || t == CV_Bool || t == CV_32S || t == CV_64S || t == CV_32F ||
t == CV_8U || t == CV_8S; };
if (!castable(inType) || !castable(outType))
return false;
// an unknown output type still has to match the static 'to' attribute
if (outType < 0 && toCvDepth_ >= 0 && !castable(resolveStorageDepth(toCvDepth_, false)))
return false;
// the FP16 target stores float as half, which has no CastOp instantiation; an unknown type may be float
return netimpl_->preferableTarget != DNN_TARGET_CUDA_FP16 ||
(inType >= 0 && outType >= 0 && inType != CV_32F && outType != CV_32F);
@@ -147,6 +151,8 @@ public:
case CV_32S: return makeCastFrom<int32_t>(outDepth, std::move(context->stream));
case CV_64S: return makeCastFrom<int64_t>(outDepth, std::move(context->stream));
case CV_32F: return makeCastFrom<float>(outDepth, std::move(context->stream));
case CV_8U: return makeCastFrom<uint8_t>(outDepth, std::move(context->stream));
case CV_8S: return makeCastFrom<int8_t>(outDepth, std::move(context->stream));
}
CV_Error(Error::BadDepth, "Cast: unsupported input type on CUDA");
return Ptr<BackendNode>();
@@ -161,6 +167,8 @@ public:
case CV_32S: return Ptr<BackendNode>(new cuda4dnn::CastOp<int32_t, TIn>(std::move(stream)));
case CV_64S: return Ptr<BackendNode>(new cuda4dnn::CastOp<int64_t, TIn>(std::move(stream)));
case CV_32F: return Ptr<BackendNode>(new cuda4dnn::CastOp<float, TIn>(std::move(stream)));
case CV_8U: return Ptr<BackendNode>(new cuda4dnn::CastOp<uint8_t, TIn>(std::move(stream)));
case CV_8S: return Ptr<BackendNode>(new cuda4dnn::CastOp<int8_t, TIn>(std::move(stream)));
}
CV_Error(Error::BadDepth, "Cast: unsupported output type on CUDA");
return Ptr<BackendNode>();

View File

@@ -24,8 +24,14 @@ static inline void clampFloatChunkDispatch(const float* src, float* dst,
#undef CV_CPU_DISPATCH_MODES_ALL
#include "layers_common.hpp"
#include "../net_impl.hpp"
#include <opencv2/dnn/shape_utils.hpp>
#include <opencv2/core/hal/interface.h>
#ifdef HAVE_CUDA
#include "../op_cuda.hpp"
#include "../cuda4dnn/primitives/activation.hpp"
#endif
#include <limits>
#include <cfloat>
#include <cmath>
@@ -113,9 +119,51 @@ public:
virtual bool supportBackend(int backendId) CV_OVERRIDE
{
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
if (!netimpl_ || inputs.empty() || inputs.size() > 3)
return false;
// initCUDA reads min/max on the host
for (size_t i = 1; i < inputs.size(); i++)
if (!inputs[i].empty() && !netimpl_->isConstArg(inputs[i]))
return false;
// 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 || isCudaNodeType(t, false);
}
#endif
return backendId == DNN_BACKEND_OPENCV;
}
#ifdef HAVE_CUDA
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 type's range
auto constScalar = [&](size_t i, double& v) {
if (inputs.size() <= i || inputs[i].empty())
return false;
Mat tmp;
netimpl_->argTensor(inputs[i]).convertTo(tmp, CV_64F);
CV_Assert(tmp.total() == 1);
v = tmp.at<double>(0);
return true;
};
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_with_type<cuda4dnn::ClipOp>(preferableTarget, type, std::move(context->stream), lo, hi);
}
#endif
// Clip rewrites values in place of the input layout.
virtual bool alwaysSupportInplace() const CV_OVERRIDE { return true; }

View File

@@ -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
@@ -547,6 +548,20 @@ public:
return absorbAffine(expr) || absorbKernel(expr);
}
//! The kernel activations ConvolutionOp applies itself.
static bool cudaConvKernel(ActivationFunc fn)
{
return fn == getActivationFunc(ACTIV_SIGMOID) || fn == getActivationFunc(ACTIV_SWISH) ||
fn == getActivationFunc(ACTIV_MISH) || fn == getActivationFunc(ACTIV_TANH);
}
//! Kernel activations the CUDA path runs as a separate op after the convolution.
static bool cudaPostKernel(ActivationFunc fn)
{
return fn == getActivationFunc(ACTIV_HARDSWISH) || fn == getActivationFunc(ACTIV_GELU) ||
fn == getActivationFunc(ACTIV_GELU_APPROX) || fn == getActivationFunc(ACTIV_EXP);
}
//! Takes an expression the layer has a kernel for. Only a one-layer chain carries
//! one, so this cannot swallow math absorbAffine() just turned down.
bool absorbKernel(const Ptr<AdjacencyGraph>& expr)
@@ -1199,13 +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 &&
activationFunc != getActivationFunc(ACTIV_SIGMOID) && activationFunc != getActivationFunc(ACTIV_SWISH) &&
activationFunc != getActivationFunc(ACTIV_MISH) && activationFunc != getActivationFunc(ACTIV_TANH))
// 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;
}
@@ -1256,11 +1270,28 @@ 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);
} 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<cuda4dnn::ExpOp>(targetId, context.stream, scale, shift);
}
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)) {
@@ -1289,9 +1320,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

View File

@@ -12,6 +12,11 @@
#include "opencv2/core/hal/intrin.hpp"
#include <opencv2/core/utils/configuration.private.hpp>
#ifdef HAVE_CUDA
#include "../op_cuda.hpp"
#include "../cuda4dnn/primitives/conv_transpose2.hpp"
#endif
namespace cv
{
namespace dnn
@@ -96,6 +101,86 @@ public:
return inptype0;
}
// compute actual pads for SAME/VALID auto-padding, or from an explicit output_shape
std::vector<int> resolvePads(const MatShape& inpshape, const MatShape& outshape,
const MatShape& wshape, int nsd) const
{
std::vector<int> pads_resolved = pads;
if (auto_pad != AUTO_PAD_NONE || !explicit_out_shape.empty()) {
pads_resolved.resize(nsd * 2, 0);
for (int i = 0; i < nsd; i++) {
int inpsz = inpshape[2 + i];
int outsz = outshape[2 + i];
int adj_i = adjust_pads.empty() ? 0 : adjust_pads[i];
int stride = strides.empty() ? 1 : strides[i];
int dil = dilations.empty() ? 1 : dilations[i];
int ki = wshape[2 + i];
int total = (inpsz - 1) * stride + dil * (ki - 1) + 1 + adj_i - outsz;
total = std::max(total, 0);
// ONNX splits total_padding evenly; odd pixel goes last for SAME_UPPER.
int pb = (auto_pad == AUTO_PAD_SAME_UPPER) ? total / 2 : total - total / 2;
pads_resolved[i] = pb;
pads_resolved[nsd + i] = total - pb;
}
}
return pads_resolved;
}
virtual bool supportBackend(int backendId) CV_OVERRIDE
{
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
// 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 = 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
for (const Arg& a : inputs) {
const int t = a.empty() ? -1 : netimpl_->argData(a).type;
if (t >= 0 && t != CV_32F)
return false;
}
return true;
}
#endif
return backendId == DNN_BACKEND_OPENCV;
}
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void* context_, InputArrayOfArrays inputs_arr, InputArrayOfArrays outputs_arr) CV_OVERRIDE
{
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 = 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
std::vector<int> pads_resolved = resolvePads(config.input_shape, config.output_shape, config.weight_shape, nsd);
for (int i = 0; i < nsd; i++) {
config.strides.push_back(strides.empty() ? 1 : strides[i]);
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);
}
#endif
virtual void setWeights(InputArray weights_arr, InputArray bias_arr,
int C0, int /*accuracy*/) CV_OVERRIDE
{
@@ -110,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())
@@ -518,26 +604,8 @@ public:
pads, adjust_pads, auto_pad);
applyExplicitOutShape(outshape);
// compute actual pads for SAME/VALID auto-padding, or from an explicit output_shape
int nsd = inpshape.dims - 3;
std::vector<int> pads_resolved = pads;
if (auto_pad != AUTO_PAD_NONE || !explicit_out_shape.empty()) {
pads_resolved.resize(nsd * 2, 0);
for (int i = 0; i < nsd; i++) {
int inpsz = inpshape[2 + i];
int outsz = outshape[2 + i];
int adj_i = adjust_pads.empty() ? 0 : adjust_pads[i];
int stride = strides.empty() ? 1 : strides[i];
int dil = dilations.empty() ? 1 : dilations[i];
int ki = wshape0[2 + i];
int total = (inpsz - 1) * stride + dil * (ki - 1) + 1 + adj_i - outsz;
total = std::max(total, 0);
// ONNX splits total_padding evenly; odd pixel goes last for SAME_UPPER.
int pb = (auto_pad == AUTO_PAD_SAME_UPPER) ? total / 2 : total - total / 2;
pads_resolved[i] = pb;
pads_resolved[nsd + i] = total - pb;
}
}
// block layout: N, C/C0, spatial..., C0
std::vector<int> pads_resolved = resolvePads(inpshape, outshape, wshape0, inpshape.dims - 3);
int outtype = inferType(inptype);
@@ -594,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;

View File

@@ -6,6 +6,11 @@
#include "../precomp.hpp"
#include "layers_common.hpp"
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#ifdef HAVE_CUDA
#include "../cuda4dnn/primitives/cum_scan.hpp"
#endif
#include <opencv2/dnn/shape_utils.hpp>
@@ -29,7 +34,44 @@ public:
setParamsFrom(params);
}
bool supportBackend(int backendId) CV_OVERRIDE { return backendId == DNN_BACKEND_OPENCV; }
bool supportBackend(int backendId) CV_OVERRIDE
{
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
if (!netimpl_ || this->inputs.empty())
return false;
// the axis is baked into the node, so it must be a constant
if (this->inputs.size() > 1 && !netimpl_->isConstArg(this->inputs[1]))
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 || t == CV_32F || t == CV_16F || t == CV_32S || t == CV_64S;
}
#endif
return backendId == DNN_BACKEND_OPENCV;
}
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void* context_,
InputArrayOfArrays inputs_arr,
InputArrayOfArrays) CV_OVERRIDE
{
auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
int axis = axis_raw;
if (this->inputs.size() > 1) {
Net::Impl* netimpl_ = getNetImpl(this);
CV_Assert(netimpl_ && netimpl_->isConstArg(this->inputs[1]));
std::vector<int> av;
tensorToIntVec(netimpl_->argTensor(this->inputs[1]), av);
CV_Assert(av.size() == 1);
axis = av[0];
}
return make_cuda_node_with_type<cuda4dnn::CumScanOp>(preferableTarget, inputs_arr.type(0),
std::move(context->stream), axis, true,
exclusive_raw == 1, reverse_raw == 1);
}
#endif
bool getMemoryShapes(const std::vector<MatShape>& inputs, const int,
std::vector<MatShape>& outputs, std::vector<MatShape>&) const CV_OVERRIDE

View File

@@ -6,6 +6,11 @@
#include "../op_inf_engine.hpp"
#include "../ie_ngraph.hpp"
#include "layers_common.hpp"
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#ifdef HAVE_CUDA
#include "../cuda4dnn/primitives/cum_scan.hpp"
#endif
#include <opencv2/dnn/shape_utils.hpp>
@@ -27,10 +32,44 @@ public:
virtual bool supportBackend(int backendId) CV_OVERRIDE
{
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
if (!netimpl_ || this->inputs.empty())
return false;
// the axis is baked into the node, so it must be a constant
if (this->inputs.size() > 1 && !netimpl_->isConstArg(this->inputs[1]))
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 || t == CV_32F || t == CV_16F || t == CV_32S || t == CV_64S;
}
#endif
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH;
}
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void* context_,
InputArrayOfArrays inputs_arr,
InputArrayOfArrays) CV_OVERRIDE
{
auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
int axis = axis_raw;
if (this->inputs.size() > 1) {
Net::Impl* netimpl_ = getNetImpl(this);
CV_Assert(netimpl_ && netimpl_->isConstArg(this->inputs[1]));
std::vector<int> av;
tensorToIntVec(netimpl_->argTensor(this->inputs[1]), av);
CV_Assert(av.size() == 1);
axis = av[0];
}
return make_cuda_node_with_type<cuda4dnn::CumScanOp>(preferableTarget, inputs_arr.type(0),
std::move(context->stream), axis, false,
exclusive_raw == 1, reverse_raw == 1);
}
#endif
bool getMemoryShapes(const std::vector<MatShape> &inputs,
const int requiredOutputs,
std::vector<MatShape> &outputs,

View File

@@ -18,6 +18,7 @@
// CUDA backend
#ifdef HAVE_CUDA
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#include "../cuda4dnn/primitives/depth_space_ops.hpp"
#endif
@@ -58,6 +59,21 @@ struct DepthSpaceOps {
#endif
}
#ifdef HAVE_CUDA
static bool cudaSupported(const LayerInfo* op) {
Net::Impl* netimpl_ = getNetImpl(op);
// classic-engine nets carry no graph args to inspect
if (!netimpl_ || op->inputs.empty())
return true;
// an intermediate arg has no type or shape until the first forward; refuse only a known unsupported one
const ArgData& inp = netimpl_->argData(op->inputs[0]);
if (inp.shape.dims >= 0 && inp.shape.dims != 4)
return false;
const int t = inp.type;
return t < 0 || isCudaNodeType(t, false);
}
#endif
void cpuCompute(const Mat &input, Mat &output) {
const auto output_shape = shape(output);
Mat tmp;
@@ -136,6 +152,10 @@ public:
}
virtual bool supportBackend(int backendId) CV_OVERRIDE {
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA)
return cudaSupported(this);
#endif
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH ||
backendId == DNN_BACKEND_CUDA ||
@@ -215,12 +235,13 @@ public:
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void *context_,
const std::vector<Ptr<BackendWrapper>>& inputs,
const std::vector<Ptr<BackendWrapper>>& outputs) override {
InputArrayOfArrays inputs_arr,
InputArrayOfArrays) override {
using namespace cv::dnn::cuda4dnn;
auto context = reinterpret_cast<csl::CSLContext*>(context_);
std::vector<size_t> perm(permutation.begin(), permutation.end());
return make_cuda_node<cuda4dnn::DepthSpaceOps>(preferableTarget, std::move(context->stream), internal_shape, perm);
return make_cuda_node_with_type<cuda4dnn::DepthSpaceOps>(preferableTarget, inputs_arr.depth(0),
std::move(context->stream), internal_shape, perm);
}
#endif // HAVE_CUDA
@@ -330,6 +351,10 @@ public:
}
virtual bool supportBackend(int backendId) CV_OVERRIDE {
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA)
return cudaSupported(this);
#endif
// The CANN and TimVX nodes take no mode and always lay the output out as DCR.
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH ||
@@ -404,12 +429,13 @@ public:
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void *context_,
const std::vector<Ptr<BackendWrapper>> &inputs,
const std::vector<Ptr<BackendWrapper>> &outputs) override {
InputArrayOfArrays inputs_arr,
InputArrayOfArrays) override {
using namespace cv::dnn::cuda4dnn;
auto context = reinterpret_cast<csl::CSLContext*>(context_);
std::vector<size_t> perm(permutation.begin(), permutation.end());
return make_cuda_node<cuda4dnn::DepthSpaceOps>(preferableTarget, std::move(context->stream), internal_shape, perm);
return make_cuda_node_with_type<cuda4dnn::DepthSpaceOps>(preferableTarget, inputs_arr.depth(0),
std::move(context->stream), internal_shape, perm);
}
#endif // HAVE_CUDA

View File

@@ -8,6 +8,10 @@
#include "layers_common.hpp"
#include "../net_impl.hpp"
#include "../onnx/onnx_dtype_convert.hpp"
#include "../op_cuda.hpp"
#ifdef HAVE_CUDA
#include "../cuda4dnn/primitives/quantize_dequantize.hpp"
#endif
#if defined(__x86_64__) || defined(_M_X64)
#include <immintrin.h>
@@ -499,6 +503,27 @@ public:
virtual bool supportBackend(int backendId) CV_OVERRIDE
{
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
const size_t ninputs = this->inputs.size();
// the kernel reads the scale and writes the output as float32, which the FP16 target stores as half
if (!netimpl_ || ninputs < 2 || ninputs > 3 || block_size != 0 ||
netimpl_->preferableTarget == DNN_TARGET_CUDA_FP16)
return false;
// an intermediate arg has no type until the first forward; refuse only a known unsupported one
const int xt = netimpl_->argData(this->inputs[0]).type;
const int st = netimpl_->argData(this->inputs[1]).type;
if (st >= 0 && st != CV_32F)
return false;
if (ninputs == 3) {
const int zt = netimpl_->argData(this->inputs[2]).type;
if (xt >= 0 && zt >= 0 && zt != xt)
return false;
}
return xt < 0 || xt == CV_8S || xt == CV_8U;
}
#endif
return backendId == DNN_BACKEND_OPENCV || backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH;
}
@@ -526,6 +551,20 @@ public:
outputs.assign(requiredOutputs, keep);
return 0;
}
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void* context_,
InputArrayOfArrays inputs_arr,
InputArrayOfArrays) CV_OVERRIDE
{
auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
CV_CheckDepthEQ(inputs_arr.depth(1), CV_32F, "DequantizeLinear: CUDA reads the scale as float");
// the kernel is only instantiated for the two 8-bit inputs
if (inputs_arr.depth(0) == CV_8S)
return Ptr<BackendNode>(new cuda4dnn::DequantizeLinearOp<int8_t>(std::move(context->stream), axis));
CV_CheckDepthEQ(inputs_arr.depth(0), CV_8U, "DequantizeLinear: CUDA reads int8 or uint8 only");
return Ptr<BackendNode>(new cuda4dnn::DequantizeLinearOp<uint8_t>(std::move(context->stream), axis));
}
#endif
bool getMemoryShapes(const std::vector<MatShape> &inputs,
const int requiredOutputs,

View File

@@ -6,6 +6,12 @@
#include "../precomp.hpp"
#include "layers_common.hpp"
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#ifdef HAVE_CUDA
#include "../cuda4dnn/primitives/dropout_mask.hpp"
#endif
namespace cv {
namespace dnn {
@@ -20,9 +26,33 @@ public:
bool supportBackend(int backendId) CV_OVERRIDE
{
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
if (!netimpl_ || inputs.empty())
return false;
// 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 || isCudaNodeType(t);
}
#endif
return backendId == DNN_BACKEND_OPENCV;
}
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void* context_,
InputArrayOfArrays inputs_arr,
InputArrayOfArrays outputs_arr) CV_OVERRIDE
{
CV_UNUSED(outputs_arr);
auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
const int depth = inputs_arr.depth(0);
if (depth == CV_Bool)
return make_cuda_node_bool<cuda4dnn::DropoutMaskOp>(std::move(context->stream));
return make_cuda_node_with_type<cuda4dnn::DropoutMaskOp>(preferableTarget, depth, std::move(context->stream));
}
#endif
bool getMemoryShapes(const std::vector<MatShape>& inputs,
const int requiredOutputs,
std::vector<MatShape>& outputs,

View File

@@ -8,12 +8,140 @@
#include "../ie_ngraph.hpp"
#include "layers_common.hpp"
#include "cpu_kernels/fast_gemm.hpp"
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#include <algorithm>
#include <cctype>
#include <map>
#include <set>
#ifdef HAVE_CUDA
#include "../cuda4dnn/primitives/einsum.hpp"
#endif
namespace cv
{
namespace dnn
{
#ifdef HAVE_CUDA
// One label per axis of an einsum term. The axes under "..." get digits, numbered from the
// right so that equal trailing ellipsis axes of both operands share a label.
static bool expandEinsumTerm(const std::string& term, int rank, std::string& labels)
{
const size_t pos = term.find("...");
const std::string left = pos == std::string::npos ? term : term.substr(0, pos);
const std::string right = pos == std::string::npos ? std::string() : term.substr(pos + 3);
const int ellipsis = rank - (int)(left.size() + right.size());
// each ellipsis axis gets one digit label, so at most ten
if (ellipsis < 0 || (pos == std::string::npos && ellipsis != 0) || ellipsis > 10)
return false;
labels = left;
for (int j = 0; j < ellipsis; j++)
labels += char('0' + (ellipsis - 1 - j));
labels += right;
for (char c : left + right)
if (!std::isalpha((unsigned char)c))
return false;
return true;
}
// Plans a two-operand einsum as permute -> strided batched GEMM -> permute. False for what that
// cannot express: a label repeated within a term, a label summed out of one operand only, or
// a shared label whose sizes differ (broadcasting).
static bool planEinsumCUDA(const std::vector<String>& terms, const String& rhs,
const MatShape& shapeA, const MatShape& shapeB, cuda4dnn::EinsumPlan& plan)
{
if (terms.size() != 2 || shapeA.dims < 0 || shapeB.dims < 0)
return false;
auto ellipsisRank = [](const std::string& term, int rank) {
return term.find("...") == std::string::npos ? 0 : rank - (int)(term.size() - 3);
};
const int ellipsisOut = std::max(ellipsisRank(terms[0], shapeA.dims), ellipsisRank(terms[1], shapeB.dims));
const int rankOut = (int)rhs.size() - (rhs.find("...") == std::string::npos ? 0 : 3 - ellipsisOut);
std::string la, lb, lo;
if (!expandEinsumTerm(terms[0], shapeA.dims, la) || !expandEinsumTerm(terms[1], shapeB.dims, lb) ||
!expandEinsumTerm(rhs, rankOut, lo))
return false;
std::map<char, int> size;
for (const std::string& labels : {la, lb, lo})
if (std::set<char>(labels.begin(), labels.end()).size() != labels.size())
return false;
for (size_t i = 0; i < la.size(); i++)
size[la[i]] = shapeA[(int)i];
for (size_t i = 0; i < lb.size(); i++) {
auto it = size.find(lb[i]);
if (it != size.end() && it->second > 0 && shapeB[(int)i] > 0 && it->second != shapeB[(int)i])
return false;
if (it == size.end() || it->second <= 0)
size[lb[i]] = shapeB[(int)i];
}
auto in = [](const std::string& labels, char c) { return labels.find(c) != std::string::npos; };
std::string batch, m, n, k;
for (char c : lo) {
if (!in(la, c) && !in(lb, c))
return false;
(in(la, c) && in(lb, c) ? batch : in(la, c) ? m : n) += c;
}
for (char c : la)
if (!in(lo, c)) {
if (!in(lb, c))
return false;
k += c;
}
for (char c : lb)
if (!in(lo, c) && !in(la, c))
return false;
auto product = [&](const std::string& labels) {
std::size_t p = 1;
for (char c : labels) p *= (std::size_t)std::max(size[c], 0);
return p;
};
auto shapeOf = [&](const std::string& labels) {
std::vector<std::size_t> s;
for (char c : labels) s.push_back((std::size_t)std::max(size[c], 0));
return s;
};
auto positions = [](const std::string& from, const std::string& wanted) {
std::vector<std::size_t> p;
for (char c : wanted) p.push_back(from.find(c));
return p;
};
plan = cuda4dnn::EinsumPlan();
plan.batch = product(batch);
plan.M = product(m);
plan.N = product(n);
plan.K = product(k);
const std::string mk = batch + m + k, kn = batch + k + n, mn = batch + m + n;
if (la != mk) {
if (la == batch + k + m) {
plan.transA = true;
} else {
plan.permA = positions(la, mk);
plan.shapeA = shapeOf(mk);
}
}
if (lb != kn) {
if (lb == batch + n + k) {
plan.transB = true;
} else {
plan.permB = positions(lb, kn);
plan.shapeB = shapeOf(kn);
}
}
if (lo != mn) {
plan.permC = positions(mn, lo);
plan.shapeC = shapeOf(mn);
}
return true;
}
#endif
static bool IsTransposeReshapeForEinsum(const std::vector<size_t>& perm,
const MatShape& input_dims,
MatShape& new_shape) {
@@ -462,10 +590,38 @@ public:
}
virtual bool supportBackend(int backendId) CV_OVERRIDE {
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
if (!netimpl_ || inputs.size() != 2 || !explicitEquation)
return false;
// an intermediate arg has no type until the first forward; refuse only a known unsupported one
for (const Arg& a : inputs) {
const int t = netimpl_->argData(a).type;
if (t >= 0 && t != CV_32F)
return false;
}
cuda4dnn::EinsumPlan plan;
return planEinsumCUDA(lhs_eq_tokens, rhs_eq, netimpl_->argData(inputs[0]).shape,
netimpl_->argData(inputs[1]).shape, plan);
}
#endif
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH;
}
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void* context_, InputArrayOfArrays inputs_arr, InputArrayOfArrays outputs_arr) CV_OVERRIDE
{
CV_UNUSED(outputs_arr);
auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
cuda4dnn::EinsumPlan plan;
CV_Assert(planEinsumCUDA(lhs_eq_tokens, rhs_eq, inputs_arr.shape(0), inputs_arr.shape(1), plan));
return make_cuda_node<cuda4dnn::EinsumOp>(preferableTarget, std::move(context->stream),
std::move(context->cublas_handle), plan);
}
#endif
virtual void getTypes(const std::vector<MatType>& inputs,
const int requiredOutputs,
const int requiredInternals,

View File

@@ -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<BackendNode> makeFusedCudaNode(const PreparedFusion& fusion, const Pt
postNode = make_cuda_node<cuda4dnn::TanHOp>(target, stream);
} else if (fn == getActivationFunc(ACTIV_MISH)) {
postNode = make_cuda_node<cuda4dnn::MishOp>(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<cuda4dnn::HardSigmoidOp>(target, stream, alpha, beta);
} else {
const float slope = fusion.activationParams.empty() ? 0.f : fusion.activationParams[0];
postNode = make_cuda_node<cuda4dnn::ReLUOp>(target, stream, slope);

View File

@@ -1,4 +1,9 @@
#include "../precomp.hpp"
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#ifdef HAVE_CUDA
#include "../cuda4dnn/primitives/gather_nd.hpp"
#endif
#include <opencv2/dnn/shape_utils.hpp>
namespace cv { namespace dnn {
@@ -12,6 +17,35 @@ public:
batch_dims = params.get<int>("batch_dims", 0);
}
virtual bool supportBackend(int backendId) CV_OVERRIDE
{
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
if (!netimpl_ || this->inputs.size() != 2)
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;
const int it = netimpl_->argData(this->inputs[1]).type;
if (it >= 0 && it != CV_32S && it != CV_64S)
return false;
return t < 0 || isCudaNodeType(t, false);
}
#endif
return backendId == DNN_BACKEND_OPENCV;
}
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void* context_,
InputArrayOfArrays inputs_arr,
InputArrayOfArrays) CV_OVERRIDE
{
auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
return make_cuda_node_with_type<cuda4dnn::GatherNDOp>(preferableTarget, inputs_arr.type(0),
std::move(context->stream), batch_dims);
}
#endif
void getTypes(const std::vector<MatType>& inputs,
const int requiredOutputs,
const int requiredInternals,

View File

@@ -3,6 +3,11 @@
// of this distribution and at http://opencv.org/license.html.
#include "../precomp.hpp"
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#ifdef HAVE_CUDA
#include "../cuda4dnn/primitives/gather_elements.hpp"
#endif
#include "../op_inf_engine.hpp"
#include "../ie_ngraph.hpp"
#include <opencv2/dnn/shape_utils.hpp>
@@ -32,10 +37,34 @@ public:
virtual bool supportBackend(int backendId) CV_OVERRIDE
{
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
if (!netimpl_ || this->inputs.size() != 2)
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;
const int it = netimpl_->argData(this->inputs[1]).type;
if (it >= 0 && it != CV_32S && it != CV_64S)
return false;
return t < 0 || isCudaNodeType(t, false);
}
#endif
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH;
}
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void* context_,
InputArrayOfArrays inputs_arr,
InputArrayOfArrays) CV_OVERRIDE
{
auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
return make_cuda_node_with_type<cuda4dnn::GatherElementsOp>(preferableTarget, inputs_arr.type(0),
std::move(context->stream), axis);
}
#endif
virtual bool getMemoryShapes(const std::vector<MatShape> &inputs,
const int requiredOutputs,
std::vector<MatShape> &outputs,

View File

@@ -111,13 +111,15 @@ public:
t = netimpl_->argData(inputs[0]).type;
if (!const_B || trans_a || (t >= 0 && t != CV_32F))
return false;
// InnerProductOp computes A*B + C: no alpha/beta scaling, and C only as a constant length-N row
if (alpha != 1.f || inputs.size() > 1)
// InnerProductOp computes A*B + c; alpha/beta fold into the constant B and C in initCUDA
if (inputs.size() > 1)
return false;
// C must be a constant scalar or length-N row
if (blobs.size() >= 2) {
const Mat& C = blobs[1];
const int N = trans_b ? blobs[0].size[0] : blobs[0].size[1];
if (beta != 1.f || C.dims < 1 || C.total() != (size_t)N || C.size[C.dims - 1] != N)
const bool row = C.dims >= 1 && C.total() == (size_t)N && C.size[C.dims - 1] == N;
if (C.total() != 1 && !row)
return false;
}
return true;
@@ -666,6 +668,19 @@ public:
if (!trans_b)
cv::transpose(B, B);
// alpha*A*B + beta*C == A*(alpha*B) + (beta*C); copy before scaling, blobs stay untouched
if (alpha != 1.f)
B = Mat(B * alpha);
if (!C.empty()) {
const int N = B.size[0];
if (C.total() == 1) {
Mat c32;
C.convertTo(c32, CV_32F);
C = Mat(1, N, CV_32F, Scalar::all(c32.at<float>(0)));
}
if (beta != 1.f)
C = Mat(C * beta);
}
// flatten_a=false keeps A's leading dims as rows, as forward() does: only the last axis is K
auto flatten_start_axis = flatten_a ? normalize_axis(1, inputs[0].dims) : inputs[0].dims - 1;
// the fused activation runs on the same stream, right after the GEMM

View File

@@ -14,6 +14,11 @@
#undef CV_CPU_DISPATCH_MODES_ALL
#include "layers_common.hpp"
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#ifdef HAVE_CUDA
#include "../cuda4dnn/primitives/grid_sample.hpp"
#endif
#include <opencv2/dnn/shape_utils.hpp>
#include <opencv2/core/utility.hpp>
@@ -461,9 +466,44 @@ public:
}
bool supportBackend(int backendId) CV_OVERRIDE {
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
// the kernel reads the grid as float32, which the FP16 target stores as half
if (!netimpl_ || this->inputs.size() != 2 || netimpl_->preferableTarget == DNN_TARGET_CUDA_FP16)
return false;
// the CUDA kernel is 2-D only; a 5-D volumetric sample stays on CPU
const int dims = netimpl_->argData(this->inputs[0]).shape.dims;
if (dims >= 0 && dims != 4)
return false;
// an intermediate arg has no type until the first forward; refuse only a known unsupported one
const int gt = netimpl_->argData(this->inputs[1]).type;
if (gt >= 0 && gt != CV_32F)
return false;
const int t = netimpl_->argData(this->inputs[0]).type;
return t < 0 || t == CV_32F || t == CV_16F;
}
#endif
return backendId == DNN_BACKEND_OPENCV;
}
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void* context_,
InputArrayOfArrays inputs_arr,
InputArrayOfArrays) CV_OVERRIDE
{
auto context = reinterpret_cast<cuda4dnn::csl::CSLContext*>(context_);
CV_CheckDepthEQ(inputs_arr.depth(1), CV_32F, "GridSample: CUDA reads the grid as float");
cuda4dnn::kernels::GridSampleParams params;
params.mode = mode;
params.padding = padding;
params.align_corners = align_corners;
params.cubic_alpha = cubic_alpha;
return make_cuda_node_with_type<cuda4dnn::GridSampleOp>(preferableTarget, inputs_arr.type(0),
std::move(context->stream), params);
}
#endif
bool getMemoryShapes(const std::vector<MatShape>& inputs,
const int requiredOutputs,
std::vector<MatShape>& outputs,

View File

@@ -34,6 +34,24 @@ public:
}
virtual bool supportBackend(int backendId) CV_OVERRIDE {
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
// classic-engine nets carry no graph args to inspect
if (!netimpl_ || this->inputs.empty())
return true;
if (this->inputs.size() != 3)
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;
for (const Arg& a : this->inputs) {
const int ta = netimpl_->argData(a).type;
if ((ta >= 0 && ta != CV_32F && ta != CV_16F) || (ta >= 0 && t >= 0 && ta != t))
return false;
}
return true;
}
#endif
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_CUDA;
}
@@ -175,17 +193,17 @@ public:
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void *context_,
const std::vector<Ptr<BackendWrapper>>& inputs,
const std::vector<Ptr<BackendWrapper>>& outputs) override {
InputArrayOfArrays inputs_arr,
InputArrayOfArrays) CV_OVERRIDE {
auto context = reinterpret_cast<csl::CSLContext*>(context_);
auto input_wrapper = inputs[0].dynamicCast<CUDABackendWrapper>();
auto input_shape = input_wrapper->getShape();
const MatShape input_shape = inputs_arr.shape(0);
size_t N = input_shape[0];
size_t num_groups = this->num_groups;
size_t loops = N * num_groups;
size_t group_elems = static_cast<size_t>(total(input_shape, 1)) / num_groups;
return make_cuda_node<cuda4dnn::GroupNormOp>(preferableTarget, std::move(context->stream), epsilon, loops, num_groups);
return make_cuda_node<cuda4dnn::GroupNormOp>(preferableTarget, std::move(context->stream), epsilon, loops, num_groups, group_elems);
}
#endif // HAVE_CUDA

View File

@@ -15,6 +15,7 @@
#include "../ie_ngraph.hpp"
// CUDA backend
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#ifdef HAVE_CUDA
#include "../cuda4dnn/primitives/instance_norm.hpp"
@@ -42,6 +43,24 @@ public:
#ifdef HAVE_INF_ENGINE
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
return true;
#endif
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
Net::Impl* netimpl_ = getNetImpl(this);
// classic-engine nets carry no graph args to inspect
if (!netimpl_ || this->inputs.empty())
return true;
if (this->inputs.size() != 3)
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;
for (const Arg& a : this->inputs) {
const int ta = netimpl_->argData(a).type;
if ((ta >= 0 && ta != CV_32F && ta != CV_16F) || (ta >= 0 && t >= 0 && ta != t))
return false;
}
return true;
}
#endif
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_CUDA;
@@ -265,15 +284,15 @@ public:
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(void *context_,
const std::vector<Ptr<BackendWrapper>>& inputs,
const std::vector<Ptr<BackendWrapper>>& outputs) override {
InputArrayOfArrays inputs_arr,
InputArrayOfArrays) CV_OVERRIDE {
auto context = reinterpret_cast<csl::CSLContext*>(context_);
auto input_wrapper = inputs[0].dynamicCast<CUDABackendWrapper>();
auto input_shape = input_wrapper->getShape();
const MatShape input_shape = inputs_arr.shape(0);
size_t loops = static_cast<size_t>(total(input_shape, 0, 2));
size_t norm_size = static_cast<size_t>(total(input_shape, 2));
return make_cuda_node<cuda4dnn::InstanceNormOp>(preferableTarget, std::move(context->stream), epsilon, loops);
return make_cuda_node<cuda4dnn::InstanceNormOp>(preferableTarget, std::move(context->stream), epsilon, loops, norm_size);
}
#endif // HAVE_CUDA

View File

@@ -37,10 +37,19 @@ public:
{
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA) {
// LayerNormOp writes Y only; the optional Mean/InvStdDev outputs stay on CPU
// LayerNormOp copies its float Mean/InvStdDev scratch to the optional outputs
Net::Impl* netimpl_ = getNetImpl(this);
if (!netimpl_ || inputs.empty() || outputs.size() != 1)
if (!netimpl_ || inputs.empty() || outputs.empty() || outputs.size() > 3)
return false;
if (outputs.size() > 1) {
if (netimpl_->preferableTarget == DNN_TARGET_CUDA_FP16)
return false;
for (size_t k = 1; k < outputs.size(); k++) {
const int ot = outputs[k].empty() ? -1 : netimpl_->argData(outputs[k]).type;
if (ot >= 0 && ot != CV_32F)
return false;
}
}
// weights are either all constant blobs or all runtime inputs
const bool constWeights = inputs.size() == 1 && (blobs.size() == 1 || blobs.size() == 2);
const bool inputWeights = blobs.empty() && (inputs.size() == 2 || inputs.size() == 3);

View File

@@ -42,6 +42,7 @@
#include "../precomp.hpp"
#include "layers_common.hpp"
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#include "../op_inf_engine.hpp"
#include "../ie_ngraph.hpp"
@@ -99,9 +100,18 @@ public:
virtual bool supportBackend(int backendId) CV_OVERRIDE
{
#if defined(HAVE_CUDNNJIT) && !defined(HAVE_CUDNN)
if (backendId == DNN_BACKEND_CUDA)
return false; // csl::LRN needs the legacy cuDNN API, absent in a JIT-only build
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA && useHandKernelCUDA()) {
// the hand kernel covers across-channel LRN (ONNX); within-channel needs legacy cuDNN
if (type != CHANNEL_NRM)
return false;
Net::Impl* netimpl_ = getNetImpl(this);
if (!netimpl_ || inputs.empty())
return true;
// 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;
}
#endif
#ifdef HAVE_INF_ENGINE
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
@@ -324,14 +334,34 @@ public:
}
#ifdef HAVE_CUDA
// new-engine graphs and JIT-only builds use the hand kernel; classic full-cuDNN keeps csl::LRN
bool useHandKernelCUDA()
{
#if defined(HAVE_CUDNNJIT) && !defined(HAVE_CUDNN)
return true;
#else
return getNetImpl(this) != nullptr && !inputs.empty();
#endif
}
Ptr<BackendNode> initCUDA(
void *context_,
const std::vector<Ptr<BackendWrapper>>& inputs,
const std::vector<Ptr<BackendWrapper>>& outputs
InputArrayOfArrays inputs_arr,
InputArrayOfArrays
) override
{
auto context = reinterpret_cast<csl::CSLContext*>(context_);
if (useHandKernelCUDA()) {
CV_Assert(type == CHANNEL_NRM);
const float alphaScaled = normBySize ? (float)(alpha / size) : (float)alpha;
return make_cuda_node<cuda4dnn::LRNAcrossChannelsOp>(preferableTarget, std::move(context->stream),
size, alphaScaled, (float)beta, (float)bias);
}
#if !defined(HAVE_CUDNNJIT) || defined(HAVE_CUDNN)
std::vector<UMat> inputs;
inputs_arr.getUMatVector(inputs);
cuda4dnn::LRNType type_;
if (type == CHANNEL_NRM)
type_ = cuda4dnn::LRNType::ACROSS_CHANNELS;
@@ -349,17 +379,14 @@ public:
}
std::size_t largestInputSize = 0;
for(auto& wrapper : inputs) {
auto input_wrapper = wrapper.dynamicCast<CUDABackendWrapper>();
auto shape = input_wrapper->getShape();
largestInputSize = std::max<std::size_t>(
largestInputSize,
std::accumulate(std::begin(shape), std::end(shape), 1, std::multiplies<int>())
);
}
for (const UMat& u : inputs)
largestInputSize = std::max<std::size_t>(largestInputSize, u.total());
return make_cuda_node<cuda4dnn::LRNOp>(preferableTarget,
std::move(context->cudnn_handle), type_, size, alphaSize, beta, bias, largestInputSize);
#else
CV_Error(Error::StsNotImplemented, "LRN: within-channel normalization needs the legacy cuDNN API");
#endif
}
#endif

View File

@@ -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 <opencv2/dnn/shape_utils.hpp>
#include <opencv2/core/hal/intrin.hpp>
@@ -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<int BITS>
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 uchar>();
const float* sc = blobs[1].ptr<const float>();
Mat W(N, K, CV_32F);
parallel_for_(Range(0, N), [&](const Range& r) {
AutoBuffer<float, 1024> 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<size_t>(n) * n_blk + blk;
dequantBlock<BITS>(bq + b * blob_bytes, blob_bytes, sc[b], wcol.data() + blk * block_size);
}
std::copy(wcol.data(), wcol.data() + K, W.ptr<float>(n));
}
});
return W;
}
Ptr<BackendNode> initCUDA(void* context_, InputArrayOfArrays inputs_, InputArrayOfArrays) CV_OVERRIDE {
auto context = reinterpret_cast<csl::CSLContext*>(context_);
std::vector<UMat> 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<InnerProductOp>(preferableTarget, std::move(context->stream),
std::move(context->cublas_handle), axis, dequantW, Mat());
}
Mat dequantW;
#endif
};
Ptr<MatMulNBitsLayer> MatMulNBitsLayer::create(const LayerParams& params)

Some files were not shown because too many files have changed in this diff Show More