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
151 changed files with 6642 additions and 950 deletions

View File

@@ -951,7 +951,7 @@ OPENCV_HAL_IMPL_AVX_CMP_OP_64BIT(v_int64x4)
#define OPENCV_HAL_IMPL_AVX_CMP_OP_FLT(_Tpvec, suffix) \
OPENCV_HAL_IMPL_AVX_CMP_FLT(v_eq, _CMP_EQ_OQ, _Tpvec, suffix) \
OPENCV_HAL_IMPL_AVX_CMP_FLT(v_ne, _CMP_NEQ_UQ, _Tpvec, suffix) \
OPENCV_HAL_IMPL_AVX_CMP_FLT(v_ne, _CMP_NEQ_OQ, _Tpvec, suffix) \
OPENCV_HAL_IMPL_AVX_CMP_FLT(v_lt, _CMP_LT_OQ, _Tpvec, suffix) \
OPENCV_HAL_IMPL_AVX_CMP_FLT(v_gt, _CMP_GT_OQ, _Tpvec, suffix) \
OPENCV_HAL_IMPL_AVX_CMP_FLT(v_le, _CMP_LE_OQ, _Tpvec, suffix) \

View File

@@ -888,7 +888,7 @@ OPENCV_HAL_IMPL_AVX512_CMP_OP_INT(v_int64x8, epi64, epi64, (int64)-1)
#define OPENCV_HAL_IMPL_AVX512_CMP_OP_FLT(_Tpvec, sufcmp, sufset, tval) \
OPENCV_HAL_IMPL_AVX512_CMP_FLT(v_eq, _CMP_EQ_OQ, _Tpvec, sufcmp, sufset, tval) \
OPENCV_HAL_IMPL_AVX512_CMP_FLT(v_ne, _CMP_NEQ_UQ, _Tpvec, sufcmp, sufset, tval) \
OPENCV_HAL_IMPL_AVX512_CMP_FLT(v_ne, _CMP_NEQ_OQ, _Tpvec, sufcmp, sufset, tval) \
OPENCV_HAL_IMPL_AVX512_CMP_FLT(v_lt, _CMP_LT_OQ, _Tpvec, sufcmp, sufset, tval) \
OPENCV_HAL_IMPL_AVX512_CMP_FLT(v_gt, _CMP_GT_OQ, _Tpvec, sufcmp, sufset, tval) \
OPENCV_HAL_IMPL_AVX512_CMP_FLT(v_le, _CMP_LE_OQ, _Tpvec, sufcmp, sufset, tval) \

View File

@@ -1026,8 +1026,8 @@ OPENCV_HAL_IMPL_LASX_CMP_OP_64BIT(v_int64x4, d)
{ return _Tpvec(__lasx_##suffix##_##ssuffix(a.val, b.val)); }
#define OPENCV_HAL_IMPL_LASX_CMP_OP_FLT(_Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LASX_CMP_FLT(v_eq, xvfcmp_ceq, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LASX_CMP_FLT(v_ne, xvfcmp_cune, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LASX_CMP_FLT(v_eq, xvfcmp_ceq, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LASX_CMP_FLT(v_ne, xvfcmp_cne, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LASX_CMP_FLT(v_lt, xvfcmp_clt, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LASX_CMP_FLT(v_le, xvfcmp_cle, _Tpvec, ssuffix)

View File

@@ -800,8 +800,8 @@ OPENCV_HAL_IMPL_LSX_CMP_OP_64BIT(v_int64x2, d)
{ return _Tpvec(__lsx_##suffix##_##ssuffix(a.val, b.val)); } \
#define OPENCV_HAL_IMPL_LSX_CMP_OP_FLT(_Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LSX_CMP_FLT(v_eq, vfcmp_ceq, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LSX_CMP_FLT(v_ne, vfcmp_cune, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LSX_CMP_FLT(v_eq, vfcmp_ceq, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LSX_CMP_FLT(v_ne, vfcmp_cne, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LSX_CMP_FLT(v_lt, vfcmp_clt, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LSX_CMP_FLT(v_le, vfcmp_cle, _Tpvec, ssuffix) \

View File

@@ -27,18 +27,14 @@ using namespace ::tbb;
#if TBB_INTERFACE_VERSION >= 8000
static tbb::task_arena& getArena()
{
// Keep the TBB object alive until process termination to avoid static
// destruction during process teardown on Windows.
static tbb::task_arena* tbbArena = new tbb::task_arena(tbb::task_arena::automatic);
return *tbbArena;
static tbb::task_arena tbbArena(tbb::task_arena::automatic);
return tbbArena;
}
#else
static tbb::task_scheduler_init& getScheduler()
{
// Keep the TBB object alive until process termination to avoid static
// destruction during process teardown on Windows.
static tbb::task_scheduler_init* tbbScheduler = new tbb::task_scheduler_init(tbb::task_scheduler_init::deferred);
return *tbbScheduler;
static tbb::task_scheduler_init tbbScheduler(tbb::task_scheduler_init::deferred);
return tbbScheduler;
}
#endif

View File

@@ -8,7 +8,7 @@
#define CV_VERSION_MAJOR 5
#define CV_VERSION_MINOR 1
#define CV_VERSION_REVISION 0
#define CV_VERSION_STATUS ""
#define CV_VERSION_STATUS "-pre"
#define CVAUX_STR_EXP(__A) #__A
#define CVAUX_STR(__A) CVAUX_STR_EXP(__A)

View File

@@ -25,7 +25,6 @@
#include <algorithm>
#include <cmath>
#include <cstring>
#include <type_traits>
namespace cv {
@@ -384,7 +383,7 @@ struct EwMul {
#endif
template<typename V, typename S> static V vec(const V& a, const V& b, const S&) { return vec(a, b); }
template<typename V> static V preproc(const V& a, const V& s) { return v_mul(a, s); }
template<typename W, typename ST> static W scl(W a, W b, ST s) { return a * s * b; }
template<typename W, typename ST> static W scl(W a, W b, ST s) { return a * b * s; }
};
// div has two variants by the COMMON INPUT type (matching cv::'s per-type kernel choice): integer
// inputs guard divide-by-zero -> 0 (cv:: iscalar_div); float inputs do NOT guard (cv:: fscalar_div,
@@ -574,11 +573,6 @@ static int scalarBinaryKernel(const void* src0_, size_t s0y, size_t s0x,
s1y /= sizeof(T);
dsty /= sizeof(Tr);
CV_Assert((s0x|s1x) == 1u || (s0x|s1x) + (size_t)width == 1u);
if constexpr (std::is_same_v<Op, EwMul>) {
if ((s1x == 0 && s0x != 0) || (s1y == 0 && s0y != 0 && s0x == s1x)) {
std::swap(src0_, src1_); std::swap(s0y, s1y); std::swap(s0x, s1x);
}
}
const T* src0 = (const T*)src0_;
const T* src1 = (const T*)src1_;
@@ -922,11 +916,6 @@ static int vecBinaryKernel(const void* src0_, size_t s0y, size_t s0x,
dsty /= sizeof(Tr);
CV_Assert((s0x|s1x) == 1u || (s0x|s1x) + (size_t)width == 1u);
if constexpr (std::is_same_v<Op, EwMul>) {
if ((s1x == 0 && s0x != 0) || (s1y == 0 && s0y != 0 && s0x == s1x)) {
std::swap(src0_, src1_); std::swap(s0y, s1y); std::swap(s0x, s1x);
}
}
const T* src0 = (const T*)src0_;
const T* src1 = (const T*)src1_;

View File

@@ -1231,122 +1231,6 @@ static constexpr int LOCAL_HDRS = 3;
static constexpr int LOCAL_CONSTS = 4;
static constexpr int LOCAL_OPS = 16;
#if defined(__GNUC__)
#define CV_EW_NOINLINE __attribute__((noinline))
#elif defined(_MSC_VER)
#define CV_EW_NOINLINE __declspec(noinline)
#else
#define CV_EW_NOINLINE
#endif
// Outlined from TExpr::exec() so the ~4KB inline scratchBuf is not hoisted into TExpr::exec()'s
// stack frame when execution takes the broadcastOp path instead.
CV_EW_NOINLINE static void execFastStripWithTemps(
const TExpr& expr, const Mat* const* inputs, Mat* outputs, size_t total)
{
const int ninsn = (int)expr.prog.size();
const int totalEsz = expr.bufEszPrefix[expr.nbuffers];
const int wf0 = std::min((int)total, expr.capElems);
const size_t region = alignSize((size_t)wf0, 8);
// Scratch for the temp buffers. AutoBuffer no longer value-inits its tail, so a fresh per-call
// buffer is free (we only WRITE to it); the inline 4KB covers the L1-capped size, heap backs
// the rare larger case. uint64_t is used for memory alignment on the worst case.
AutoBuffer<uint64_t, (4*1024 + 256)/sizeof(uint64_t)> scratchBuf(
divUp((size_t)totalEsz * region, (unsigned)sizeof(uint64_t)));
uchar* scratch = (uchar*)scratchBuf.data();
for (int x0 = 0; x0 < (int)total; x0 += wf0)
{
const int wf = std::min(wf0, (int)total - x0);
auto ptrOf = [&](int s) -> void* {
if (s <= 0) return nullptr;
const TExpr::Arg& ai = expr.arginfo[s];
size_t esz = CV_ELEM_SIZE1(ai.depth);
if (ai.kind == TExpr::INPUT) return (uchar*)inputs[ai.index]->data + (size_t)x0 * esz;
if (ai.kind == TExpr::OUTPUT) return (uchar*)outputs[ai.index].data + (size_t)x0 * esz;
if (ai.kind != TExpr::TEMP) return nullptr; // NONE: moved-from
const int b = expr.bufferOfTemp.empty() ? ai.index : expr.bufferOfTemp[ai.index]; // TEMP
return scratch + (size_t)expr.bufEszPrefix[b] * region; // fragment-local
};
for (int n = 0; n < ninsn; n++)
{
const TExpr::Insn& ins = expr.prog[n];
runInsn(ins, ptrOf(ins.arg0), 0, 1, ptrOf(ins.arg1), 0, 1,
ptrOf(ins.arg2), 0, 1, ptrOf(ins.result), 0, wf, 1);
}
}
}
// Outlined from the broadcastOp body so single-op calls (nbuffers == 0) do not reserve ~4KB of
// unused tstoreBuf stack space in the closure frame.
CV_EW_NOINLINE static void execTileWithTemps(
const EwBody& bc, const TExpr::Arg* arginfo, const BroadcastOp::Tile& tile)
{
using BrSlice = BroadcastOp::Slice;
const TExpr::Insn* insns = bc.prog;
const int* const smap = bc.slotMap;
const signed char* const skind = bc.slotKind;
const int nin = bc.ninsn, nsl = bc.nslots;
const int w = tile.width, h = tile.height;
// Run the program over L1-sized 2D blocks so the intermediates stay in cache. ONE rule covers
// every tile shape: keep the strip as WIDE as fits L1 (long inner loop => good SIMD + hits the
// kernels' width-specific branches), then add as many rows as still fit (bw*bh <= capElems).
// For the dominant 1D tile (w huge, h==1) this is a column strip; for a tall-thin tile
// (w==channels, h huge - e.g. a masked op) it instead splits the LONG axis (height), so the
// temp block stays contiguous and the kernels keep full width. Each temp buffer occupies
// [bufEszPrefix[buf]*region, ...) bytes in tstore; region = alignSize(bw*bh, 8).
const int* const eszPrefix = bc.bufEszPrefix;
const int bw = std::min(w, bc.capElems);
const int bh = std::min(h, std::max(1, bc.capElems / std::max(1, bw)));
const size_t region = alignSize((size_t)bw * bh, 8);
// uint64_t elements (see the note on scratchBuf above): keeps the block-local temp store
// 8-byte aligned, which the 64-bit-lane kernels require on 32-bit ABIs.
AutoBuffer<uint64_t, (4*1024 + 256)/sizeof(uint64_t)> tstoreBuf(
divUp((size_t)eszPrefix[bc.nbuffers] * region, (unsigned)sizeof(uint64_t))); // inline (<= ~4KB)
uchar* tstore = (uchar*)tstoreBuf.data();
AutoBuffer<BrSlice, LOCAL_OPS> args(nsl);
for (int y0 = 0; y0 < h; y0 += bh)
{
const int hf = std::min(bh, h - y0);
for (int x0 = 0; x0 < w; x0 += bw)
{
const int wf = std::min(bw, w - x0);
for (int s = 1; s < nsl; s++)
{
BrSlice& a = args[s];
const int k = skind[s];
if (k == TExpr::NONE) continue;
size_t esz = CV_ELEM_SIZE1(arginfo[s].depth);
if (k == TExpr::TEMP) // contiguous block-local buffer
{
a.ptr = tstore + (size_t)eszPrefix[smap[s]] * region;
a.stepy = (size_t)wf*esz; a.stepx = 1;
}
else // array operand (incl. 0-dim consts): this
// block of the broadcast slice
{
const BrSlice& sl = tile.slices[smap[s]];
a.ptr = (uchar*)sl.ptr +
((size_t)y0 * sl.stepy + (size_t)x0 * sl.stepx) * esz;
a.stepy = sl.stepy*esz; a.stepx = sl.stepx;
}
}
for (int n = 0; n < nin; n++)
{
const TExpr::Insn& ins = insns[n];
const BrSlice& a0 = args[ins.arg0]; const BrSlice& a1 = args[ins.arg1];
const BrSlice& a2 = args[ins.arg2]; const BrSlice& rr = args[ins.result];
runInsn(ins, a0.ptr, a0.stepy, a0.stepx, a1.ptr, a1.stepy, a1.stepx,
a2.ptr, a2.stepy, a2.stepx, (void*)rr.ptr, rr.stepy, wf, hf);
}
}
}
}
#undef CV_EW_NOINLINE
void TExpr::exec(const Mat* const* inputs, Mat* outputs)
{
using BrTile = BroadcastOp::Tile;
@@ -1417,7 +1301,38 @@ void TExpr::exec(const Mat* const* inputs, Mat* outputs)
}
// Temps present: walk the strip in L1-sized fragments so the intermediates stay hot.
execFastStripWithTemps(*this, inputs, outputs, total);
// Byte layout (bufEszPrefix) and the fragment cap (capElems) both come from compile();
// scratch is totalEsz*wf0 (<= ~4KB).
const int totalEsz = bufEszPrefix[nbuffers];
const int wf0 = std::min((int)total, capElems);
const size_t region = alignSize((size_t)wf0, 8);
// Scratch for the temp buffers. AutoBuffer no longer value-inits its tail, so a fresh per-call
// buffer is free (we only WRITE to it); the inline 4KB covers the L1-capped size, heap backs
// the rare larger case. uint64_t is used for memory alignment on the worst case.
AutoBuffer<uint64_t, (4*1024 + 256)/sizeof(uint64_t)> scratchBuf(
divUp((size_t)totalEsz * region, (unsigned)sizeof(uint64_t)));
uchar* scratch = (uchar*)scratchBuf.data();
for (int x0 = 0; x0 < (int)total; x0 += wf0)
{
const int wf = std::min(wf0, (int)total - x0);
auto ptrOf = [&](int s) -> void* {
if (s <= 0) return nullptr;
const Arg& ai = arginfo[s];
size_t esz = CV_ELEM_SIZE1(ai.depth);
if (ai.kind == INPUT) return (uchar*)inputs[ai.index]->data + (size_t)x0 * esz;
if (ai.kind == OUTPUT) return (uchar*)outputs[ai.index].data + (size_t)x0 * esz;
if (ai.kind != TEMP) return nullptr; // NONE: moved-from
const int b = bufferOfTemp.empty() ? ai.index : bufferOfTemp[ai.index]; // TEMP
return scratch + (size_t)bufEszPrefix[b] * region; // fragment-local
};
for (int n = 0; n < ninsn; n++)
{
const Insn& ins = prog[n];
runInsn(ins, ptrOf(ins.arg0), 0, 1, ptrOf(ins.arg1), 0, 1,
ptrOf(ins.arg2), 0, 1, ptrOf(ins.result), 0, wf, 1);
}
}
return;
}
}
@@ -1551,38 +1466,70 @@ void TExpr::exec(const Mat* const* inputs, Mat* outputs)
body.nbuffers = nbuffers;
body.capElems = capElems;
broadcastOp(arr.data(), narr, [&bc=body, this](const BrTile& tile)
broadcastOp(arr.data(), narr, [&](const BrTile& tile)
{
if (bc.nbuffers != 0)
{
execTileWithTemps(bc, arginfo.data(), tile);
return;
}
// No temps: one block over the whole tile, no scratch buffer on the stack, and each
// operand slot is re-pointed once straight at its tile slice.
EwBody& bc = body;
const TExpr::Insn* insns = bc.prog; // hot fields -> locals (registers/stack)
const int* const smap = bc.slotMap;
const signed char* const skind = bc.slotKind;
const int nin = bc.ninsn, nsl = bc.nslots;
const int w = tile.width, h = tile.height;
AutoBuffer<BrSlice, LOCAL_OPS> args(nsl);
for (int s = 1; s < nsl; s++)
{
if (skind[s] == TExpr::NONE) continue;
const BrSlice& sl = tile.slices[smap[s]];
size_t esz = CV_ELEM_SIZE1(arginfo[s].depth);
args[s] = { sl.ptr, sl.stepy * esz, sl.stepx };
}
// Run the program over L1-sized 2D blocks so the intermediates stay in cache. ONE rule covers
// every tile shape: keep the strip as WIDE as fits L1 (long inner loop => good SIMD + hits the
// kernels' width-specific branches), then add as many rows as still fit (bw*bh <= capElems).
// For the dominant 1D tile (w huge, h==1) this is a column strip; for a tall-thin tile
// (w==channels, h huge - e.g. a masked op) it instead splits the LONG axis (height), so the
// temp block stays contiguous and the kernels keep full width. With NO temps capElems==INT_MAX
// => bw=w, bh=h: one block over the whole tile, region temp store is empty (zero bytes), and
// each operand slot is re-pointed once straight at its tile slice. Each temp buffer occupies
// [bufEszPrefix[buf]*region, ...) bytes in tstore; region = bw*bh.
const int* const eszPrefix = bc.bufEszPrefix;
const int bw = std::min(w, bc.capElems);
const int bh = std::min(h, std::max(1, bc.capElems / std::max(1, bw)));
const size_t region = alignSize((size_t)bw * bh, 8);
// uint64_t elements (see the note on scratchBuf above): keeps the block-local temp store
// 8-byte aligned, which the 64-bit-lane kernels require on 32-bit ABIs.
AutoBuffer<uint64_t, (4*1024 + 256)/sizeof(uint64_t)> tstoreBuf(
divUp((size_t)eszPrefix[bc.nbuffers] * region, (unsigned)sizeof(uint64_t))); // inline (<= ~4KB)
uchar* tstore = (uchar*)tstoreBuf.data();
for (int n = 0; n < nin; n++)
AutoBuffer<BrSlice, LOCAL_OPS> args(nsl);
for (int y0 = 0; y0 < h; y0 += bh)
{
const TExpr::Insn& ins = insns[n];
const BrSlice& a0 = args[ins.arg0]; const BrSlice& a1 = args[ins.arg1];
const BrSlice& a2 = args[ins.arg2]; const BrSlice& rr = args[ins.result];
runInsn(ins, a0.ptr, a0.stepy, a0.stepx, a1.ptr, a1.stepy, a1.stepx,
a2.ptr, a2.stepy, a2.stepx, (void*)rr.ptr, rr.stepy, w, h);
const int hf = std::min(bh, h - y0);
for (int x0 = 0; x0 < w; x0 += bw)
{
const int wf = std::min(bw, w - x0);
for (int s = 1; s < nsl; s++)
{
BrSlice& a = args[s];
const int k = skind[s];
size_t esz = CV_ELEM_SIZE1(arginfo[s].depth);
if (k == TExpr::TEMP) // contiguous block-local buffer
{
a.ptr = tstore + (size_t)eszPrefix[smap[s]] * region;
a.stepy = (size_t)wf*esz; a.stepx = 1;
}
else // array operand (incl. 0-dim consts): this
// block of the broadcast slice
{
const BrSlice& sl = tile.slices[smap[s]];
a.ptr = (uchar*)sl.ptr +
((size_t)y0 * sl.stepy + (size_t)x0 * sl.stepx) * esz;
a.stepy = sl.stepy*esz; a.stepx = sl.stepx;
}
}
for (int n = 0; n < nin; n++)
{
const TExpr::Insn& ins = insns[n];
const BrSlice& a0 = args[ins.arg0]; const BrSlice& a1 = args[ins.arg1];
const BrSlice& a2 = args[ins.arg2]; const BrSlice& rr = args[ins.result];
runInsn(ins, a0.ptr, a0.stepy, a0.stepx, a1.ptr, a1.stepy, a1.stepx,
a2.ptr, a2.stepy, a2.stepx, (void*)rr.ptr, rr.stepy, wf, hf);
}
}
}
}, true, nstripes);
}

View File

@@ -465,17 +465,9 @@ namespace {
#if defined HAVE_TBB
#if TBB_INTERFACE_VERSION >= 8000
static tbb::task_arena& getTbbArena() {
// Keep the TBB object alive until process termination to avoid static
// destruction during process teardown on Windows.
CV_SINGLETON_LAZY_INIT_REF(tbb::task_arena, new tbb::task_arena(tbb::task_arena::automatic));
}
static tbb::task_arena tbbArena(tbb::task_arena::automatic);
#else
static tbb::task_scheduler_init& getTbbScheduler() {
// Keep the TBB object alive until process termination to avoid static
// destruction during process teardown on Windows.
CV_SINGLETON_LAZY_INIT_REF(tbb::task_scheduler_init, new tbb::task_scheduler_init(tbb::task_scheduler_init::deferred));
}
static tbb::task_scheduler_init tbbScheduler(tbb::task_scheduler_init::deferred);
#endif
#elif defined HAVE_HPX
// nothing for HPX
@@ -590,7 +582,7 @@ static void parallel_for_impl(const cv::Range& range, const cv::ParallelLoopBody
#if defined HAVE_TBB
#if TBB_INTERFACE_VERSION >= 8000
getTbbArena().execute(pbody);
tbbArena.execute(pbody);
#else
pbody();
#endif
@@ -659,13 +651,13 @@ int getNumThreads(void)
#if defined HAVE_TBB
#if TBB_INTERFACE_VERSION >= 9100
return getTbbArena().max_concurrency();
return tbbArena.max_concurrency();
#elif TBB_INTERFACE_VERSION >= 8000
return numThreads > 0
? numThreads
: tbb::task_scheduler_init::default_num_threads();
#else
return getTbbScheduler().is_active()
return tbbScheduler.is_active()
? numThreads
: tbb::task_scheduler_init::default_num_threads();
#endif
@@ -743,11 +735,11 @@ void setNumThreads( int threads_ )
#ifdef HAVE_TBB
#if TBB_INTERFACE_VERSION >= 8000
if(getTbbArena().is_active()) getTbbArena().terminate();
if(threads > 0) getTbbArena().initialize(threads);
if(tbbArena.is_active()) tbbArena.terminate();
if(threads > 0) tbbArena.initialize(threads);
#else
if(getTbbScheduler().is_active()) getTbbScheduler().terminate();
if(threads > 0) getTbbScheduler().initialize(threads);
if(tbbScheduler.is_active()) tbbScheduler.terminate();
if(threads > 0) tbbScheduler.initialize(threads);
#endif
#elif defined HAVE_HPX

View File

@@ -1837,10 +1837,6 @@ ushort FileStorage::Impl::Base64Decoder::getUInt16() {
return val;
}
unsigned FileStorage::Impl::Base64Decoder::getUInt32() {
return (unsigned)getInt32();
}
int FileStorage::Impl::Base64Decoder::getInt32() {
size_t sz = decoded.size();
if (ofs + 4 > sz && !readMore(4))
@@ -1918,7 +1914,7 @@ char *FileStorage::Impl::parseBase64(char *ptr, int indent, FileNode &collection
ival = (short) base64decoder.getUInt16();
break;
case CV_32U:
ival = base64decoder.getUInt32();
ival = base64decoder.getInt32();
break;
case CV_32S:
ival = base64decoder.getInt32();
@@ -2747,7 +2743,7 @@ FileNodeIterator& FileNodeIterator::readRaw( const String& fmt, void* _data0, si
FileNode node = *(*this);
if( node.isInt() )
{
int64_t ival = (int64_t)node;
int64_t ival = static_cast<int64_t>(elem_size == 8 ? (int64_t)node : (int)node);
switch( elem_type )
{
case CV_8U:

View File

@@ -151,8 +151,6 @@ public:
ushort getUInt16();
unsigned getUInt32();
int getInt32();
int64_t getInt64();

View File

@@ -5129,65 +5129,4 @@ TEST_P(Core_PatchNaNs, replaces_only_nans)
INSTANTIATE_TEST_CASE_P(/**/, Core_PatchNaNs, testing::Values(CV_16F, CV_16BF, CV_32F, CV_64F));
TEST(Core_Multiply, scale_simd_scalar_consistency)
{
const float big = std::ldexp(1.0f, 80); // exact power of 2 (< FLT_MAX = 2^128)
const float small = std::ldexp(1.0f, -80);
// Two contiguous inputs: (a * s) * b = 2^80 vs (a * b) * s = +inf.
{
Mat a(1, 65, CV_32F, Scalar(big));
Mat b(1, 65, CV_32F, Scalar(big));
cv::multiply(a, b, a, small); // in-place -> tail element 64 uses EwMul::scl
EXPECT_EQ(a.at<float>(0), big);
EXPECT_EQ(a.at<float>(64), big);
}
// src1 scalar/column broadcast (s1x == 0): SIMD folds scale into src1, computing (b * s) * a = 2^80;
// without canonicalizing src1 onto src0, the scalar tail would compute (a * s) * b = +inf.
{
Mat a(1, 65, CV_32F, Scalar(big));
Mat b(1, 1, CV_32F, Scalar(small));
cv::multiply(a, b, a, big);
EXPECT_EQ(a.at<float>(0), big);
EXPECT_EQ(a.at<float>(64), big);
}
{
Mat a(1, 65, CV_32F, Scalar(big));
cv::multiply(a, Scalar(small), a, big);
EXPECT_EQ(a.at<float>(0), big);
EXPECT_EQ(a.at<float>(64), big);
}
// src1 short-row broadcast (s1y == 0, width <= 4): short-row SIMD processes full row blocks and
// leaves row 64 for the scalar loop.
{
Mat a(65, 3, CV_32F, Scalar(big));
Mat b(1, 3, CV_32F, Scalar(small));
cv::multiply(a, b, a, big);
EXPECT_EQ(a.at<float>(0, 0), big);
EXPECT_EQ(a.at<float>(64, 0), big);
}
// src1 wide-row broadcast (s1y == 0, width > 4):
{
Mat a(2, 65, CV_32F, Scalar(big));
Mat b(1, 65, CV_32F, Scalar(small));
cv::multiply(a, b, a, big);
EXPECT_EQ(a.at<float>(0, 0), big);
EXPECT_EQ(a.at<float>(1, 64), big);
}
// CV_64F broadcast (covers v_float64 SIMD and scalarBinaryKernel fallback):
{
const double big64 = std::ldexp(1.0, 600);
const double small64 = std::ldexp(1.0, -600);
Mat a(1, 65, CV_64F, Scalar(big64));
Mat b(1, 1, CV_64F, Scalar(small64));
cv::multiply(a, b, a, big64);
EXPECT_EQ(a.at<double>(0), big64);
EXPECT_EQ(a.at<double>(64), big64);
}
}
}} // namespace

View File

@@ -726,38 +726,6 @@ template<typename R> struct TheTest
EXPECT_EQ(dataA[i] < dataB[i], resG[i] != 0);
EXPECT_EQ(dataA[i] <= dataB[i], resH[i] != 0);
}
if (std::numeric_limits<LaneType>::has_quiet_NaN || std::is_same<LaneType, hfloat>::value)
{
const LaneType qnan = (LaneType)std::numeric_limits<float>::quiet_NaN();
dataA[0] = qnan;
if (VTraits<R>::vlanes() > 1)
{
dataA[1] = qnan;
dataB[1] = qnan;
}
if (VTraits<R>::vlanes() > 2)
dataB[2] = qnan;
a = dataA; b = dataB;
resC = (v_eq(a, b));
resD = (v_ne(a, b));
resE = (v_gt(a, b));
resF = (v_ge(a, b));
resG = (v_lt(a, b));
resH = (v_le(a, b));
for (int i = 0; i < VTraits<R>::vlanes(); ++i)
{
SCOPED_TRACE(cv::format("nan i=%d", i));
EXPECT_EQ(dataA[i] == dataB[i], resC.as_int(i) != 0);
EXPECT_EQ(dataA[i] != dataB[i], resD.as_int(i) != 0);
EXPECT_EQ(dataA[i] > dataB[i], resE.as_int(i) != 0);
EXPECT_EQ(dataA[i] >= dataB[i], resF.as_int(i) != 0);
EXPECT_EQ(dataA[i] < dataB[i], resG.as_int(i) != 0);
EXPECT_EQ(dataA[i] <= dataB[i], resH.as_int(i) != 0);
}
}
return *this;
}

View File

@@ -2662,28 +2662,4 @@ TEST(Core_InputOutput, filestorage_read_uppercase_exponent)
}
}
// CV_32U values above INT_MAX must survive saving and loading, in every format
TEST(Core_InputOutput, filestorage_32u_large_values)
{
const unsigned values[] = { 0u, 1u, 2147483647u, 2147483648u, 4294967295u };
const Mat src(1, 5, CV_32U, (void*)values);
for (const char* ext : { ".yml", ".xml", ".json" })
{
for (int base64 : { 0, (int)FileStorage::WRITE_BASE64 })
{
FileStorage writer(ext, FileStorage::WRITE | FileStorage::MEMORY | base64);
writer << "m" << src;
const std::string text = writer.releaseAndGetString();
FileStorage reader(text, FileStorage::READ | FileStorage::MEMORY);
Mat dst;
reader["m"] >> dst;
EXPECT_EQ(CV_32U, dst.type()) << ext << " base64=" << base64;
EXPECT_EQ(0, cv::norm(src, dst, NORM_INF)) << ext << " base64=" << base64;
}
}
}
}} // namespace

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,

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