Compare commits

..

13 Commits

Author SHA1 Message Date
Alexander Smorkalov
e4f78fbcf1 release: OpenCV 5.1.0 2026-11-05 14:18:05 +03:00
Shiven Solanki
545147fdab Merge pull request #30685 from sshivensolanki:core-filestorage-unsigned
core: keep large CV_32U values when reading with FileStorage - #30685

Fixes #30684

### Problem

When a `CV_32U` `Mat` was loaded with `FileStorage`, values above `INT_MAX` came back as `0`, in YAML, XML and JSON, with and without base64. The saved file had the correct values, so only reading was affected.

### Fix

- `FileNodeIterator::readRaw` read elements of up to 4 bytes through `(int)node`, so these values turned negative and the `CV_32U` case clamped them to 0. It now always reads the node as `int64_t`. Each case still converts the value to its own type.
- The base64 path stored `getInt32()` as signed, so `CV_32U` now uses a new `Base64Decoder::getUInt32()`, like `getUInt8()` and `getUInt16()`.

`CV_64U` values above `INT64_MAX` are also lost in the text formats, but that has a different cause, in the parsers' `strtoll`, so it is not part of this PR.

### Tests

New `Core_InputOutput.filestorage_32u_large_values` saves values on both sides of `INT_MAX` and loads them back, in YAML, XML and JSON, with and without base64. All 6 combinations fail without the fix, and the test passes with it on Apple M1 Pro.

### Pull Request Readiness Checklist

See details at https://github.com/opencv/opencv/wiki/How_to_contribute#making-a-good-pull-request

- [x] I agree to contribute to the project under Apache 2 License.
- [x] To the best of my knowledge, the proposed patch is not based on a code under GPL or another license that is incompatible with OpenCV
- [x] The PR is proposed to the proper branch
- [x] There is a reference to the original bug report and related work
- [x] There is accuracy test, performance test and test data in opencv_extra repository, if applicable
      Patch to opencv_extra has the same branch name.
- [ ] The feature is well documented and sample code can be built with the project CMake
2026-11-05 14:17:21 +03:00
Vincent Rabaud
3b42da9cda Merge pull request #30690 from vrabaud:imgwarp
imgproc: handle empty input image in findContours and findTRUContours - #30690

Since cv::copyMakeBorder now asserts that the source matrix is non-empty, calling cv::findContours or cv::findTRUContours on an empty image threw an exception instead of returning an empty contour list.

Guard against empty input images by clearing the output contours (and hierarchy) and returning early, matching cv::findContoursLinkRuns.

### Pull Request Readiness Checklist

See details at https://github.com/opencv/opencv/wiki/How_to_contribute#making-a-good-pull-request

- [x] I agree to contribute to the project under Apache 2 License.
- [x] To the best of my knowledge, the proposed patch is not based on a code under GPL or another license that is incompatible with OpenCV
- [x] The PR is proposed to the proper branch
- [x] There is a reference to the original bug report and related work
- [x] There is accuracy test, performance test and test data in opencv_extra repository, if applicable
      Patch to opencv_extra has the same branch name.
- [x] The feature is well documented and sample code can be built with the project CMake
2026-11-05 14:15:49 +03:00
Alexander Smorkalov
4f4d6125f7 Merge pull request #30686 from Kaustubh1235:dnn-conformance-castlike
dnn: enable the CastLike conformance cases that pass with the FP8/FP4/INT4 support
2026-11-05 10:55:59 +03:00
tanveer-arch
1727c38423 Merge pull request #30072 from tanveer-arch:investigate-30067-tbb-lifetime
core: avoid TBB static destruction during Windows teardown - #30072

## Summary

Avoid destruction of the TBB task arena/scheduler during C++ static teardown.

On Windows with the UCRT and TBB backend, process teardown can involve
interaction between TBB cleanup and the Windows loader/CRT shutdown path.
Keeping the OpenCV-owned TBB arena/scheduler alive until process termination
avoids invoking its destructor during static destruction.

## Validation

- Reproduced the reported teardown deadlock with the MSYS2 packaged
  OpenCV 5.0.0 environment.
- Built OpenCV 5.x locally with TBB enabled.
- The patched local build completed 2,001 stress-test iterations without
  reproducing the deadlock.
- `git diff --check` passes.

The local 5.x baseline did not reproduce the issue during the available
baseline stress run, so the stress result is considered supporting evidence
rather than a definitive reproduction/fix confirmation.

Related issue: #30067
2026-11-05 10:26:42 +03:00
Vincent Rabaud
7c1e0d5d50 Merge pull request #30678 from vrabaud:imgwarp
compilation fix: force 256-bit SIMD width under AVX2 in imgwarp - #30678

warp_kernels.simd.hpp and warp_common.vector.hpp only implement up to 256-bit SIMD and mix explicit v256_load* calls with width-generic v_* types. Force CV__SIMD_FORCE_WIDTH=256 when CV_AVX2 is enabled so builds with AVX-512 in CPU_BASELINE (-DCPU_BASELINE=AVX512_SKX) compile cleanly.

The following now compiles:
cmake -B build -DCPU_BASELINE=AVX512_SKX
cmake --build build --target opencv_imgproc --parallel

### Pull Request Readiness Checklist

See details at https://github.com/opencv/opencv/wiki/How_to_contribute#making-a-good-pull-request

- [x] I agree to contribute to the project under Apache 2 License.
- [x] To the best of my knowledge, the proposed patch is not based on a code under GPL or another license that is incompatible with OpenCV
- [x] The PR is proposed to the proper branch
- [ ] There is a reference to the original bug report and related work
- [ ] There is accuracy test, performance test and test data in opencv_extra repository, if applicable
      Patch to opencv_extra has the same branch name.
- [ ] The feature is well documented and sample code can be built with the project CMake
2026-11-04 16:52:24 +03:00
Vladimir Iglovikov
eb47f1f73b imgproc: vectorize float32 bilinear interpolation across channels 🤖🤖🤖 (#30650)
Vectorize the contiguous channel dimension of the float32 bilinear inlier loop for C>4. Coordinates, interpolation tables, border handling, and scalar tails keep their existing routes; C1–C4 kernels are unchanged. https://github.com/opencv/opencv/pull/30650
2026-11-04 19:10:20 +05:30
Abhishek Gola
fca1d938b0 Merge pull request #30674 from vrabaud/NaN
fix NaN inequality semantics in v_ne for AVX, AVX-512, LSX, and LASX
2026-11-04 18:06:20 +05:30
Vincent Rabaud
5bbf4b4b93 fix NaN inequality semantics in v_ne for AVX, AVX-512, LSX, and LASX
Ordered not-equal predicates (_CMP_NEQ_OQ in AVX/AVX-512 and
vfcmp_cne/xvfcmp_cne in LoongArch LSX/LASX) return false when either
operand is NaN, violating IEEE-754 != semantics and diverging from C++,
SSE (_mm_cmpneq_ps/_pd), NEON, MSA, VSX, WASM, and RVV implementations
(as well as callers like intrin_math.hpp that use v_ne(x, x) to detect
NaNs).

Switch floating-point v_ne to unordered not-equal (_CMP_NEQ_UQ and
vfcmp_cune/xvfcmp_cune) and add NaN coverage to test_cmp().
2026-11-04 10:17:08 +01:00
Abhishek Gola
2da0a498da outline 4KB temp scratch buffers in TExpr::exec (#30660)
Even after 0e20693a8d reduced the inline temp scratch buffers from 16KB
to 4KB, Clang hoists the fast-strip `scratchBuf` (~4.3KB) into
`TExpr::exec()`'s prologue when taking the `broadcastOp` path, while the
`broadcastOp` lambda unconditionally reserves another ~4.3KB `tstoreBuf`
on the stack even when `nbuffers == 0`. Together with `arithm_op()` and
`BroadcastOp::run()`, a single `cv::subtract` or `cv::compare` consumes
nearly 19KB of stack space on the calling thread, overflowing 64KB
fiber/coroutine stacks under deep call chains (e.g. `calibrateCameraRO`).

Outline the temp-bearing loops into `noinline` helpers
(`execFastStripWithTemps` and `execTileWithTemps`) and run the
`nbuffers == 0` tile body directly without `tstoreBuf`. This preserves
4KB L1 fragments when temporaries are present while saving ~8.7KB of
stack space on single-op `broadcastOp` calls.
2026-11-04 14:18:06 +05:30
Vincent Rabaud
e40204af2b Have EwMul::scl match the SIMD implementation in vecBinaryKernel (#30643)
* match SIMD and scalar scale association in EwMul

Restore `EwMul::scl` from `(a * b) * s` to `(a * s) * b` so the scalar
path matches SIMD `vec(preproc(a, s), b)` and avoids premature `float`
overflow on `a * b`.

Canonicalize single-sided `src1` broadcasts (`s1x == 0` or `s1y == 0`)
onto `src0` for `EwMul` in `vecBinaryKernel` and `scalarBinaryKernel`,
preserving loop-hoisted `b * s` in SIMD while matching the `(b * s) * a`
grouping in the scalar tails.

* Add more tests
2026-11-04 12:52:29 +05:30
Vincent Rabaud
d4a4497e38 outline 4KB temp scratch buffers in TExpr::exec
Even after 0e20693a8d reduced the inline temp scratch buffers from 16KB
to 4KB, Clang hoists the fast-strip `scratchBuf` (~4.3KB) into
`TExpr::exec()`'s prologue when taking the `broadcastOp` path, while the
`broadcastOp` lambda unconditionally reserves another ~4.3KB `tstoreBuf`
on the stack even when `nbuffers == 0`. Together with `arithm_op()` and
`BroadcastOp::run()`, a single `cv::subtract` or `cv::compare` consumes
nearly 19KB of stack space on the calling thread, overflowing 64KB
fiber/coroutine stacks under deep call chains (e.g. `calibrateCameraRO`).

Outline the temp-bearing loops into `noinline` helpers
(`execFastStripWithTemps` and `execTileWithTemps`) and run the
`nbuffers == 0` tile body directly without `tstoreBuf`. This preserves
4KB L1 fragments when temporaries are present while saving ~8.7KB of
stack space on single-op `broadcastOp` calls.
2026-11-03 16:17:42 +01:00
Kaustubh Arora
da9b4e5744 dnn: enable the CastLike conformance cases that pass with the FP8/FP4/INT4 support
34 test_castlike_* cases stayed on the parser denylist after the FP8, FP4
and INT4 cast support landed; they pass on current 5.x for OCV/CPU and
OCV/CPU_FP16. Remove them from the denylist and add them to the OpenVINO
filter as SKIP, like the other CastLike entries there.

The remaining 22 CastLike cases are INT2/UINT2 (not in the ONNX proto
OpenCV ships) and STRING, and stay denylisted.
2026-11-01 12:38:35 +05:30
151 changed files with 949 additions and 6641 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_OQ, _Tpvec, suffix) \
OPENCV_HAL_IMPL_AVX_CMP_FLT(v_ne, _CMP_NEQ_UQ, _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_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_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_cne, _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_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_cne, _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_lt, vfcmp_clt, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LSX_CMP_FLT(v_le, vfcmp_cle, _Tpvec, ssuffix) \

View File

@@ -27,14 +27,18 @@ using namespace ::tbb;
#if TBB_INTERFACE_VERSION >= 8000
static tbb::task_arena& getArena()
{
static tbb::task_arena tbbArena(tbb::task_arena::automatic);
return tbbArena;
// 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;
}
#else
static tbb::task_scheduler_init& getScheduler()
{
static tbb::task_scheduler_init tbbScheduler(tbb::task_scheduler_init::deferred);
return tbbScheduler;
// 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;
}
#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 "-pre"
#define CV_VERSION_STATUS ""
#define CVAUX_STR_EXP(__A) #__A
#define CVAUX_STR(__A) CVAUX_STR_EXP(__A)

View File

@@ -25,6 +25,7 @@
#include <algorithm>
#include <cmath>
#include <cstring>
#include <type_traits>
namespace cv {
@@ -383,7 +384,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 * b * s; }
template<typename W, typename ST> static W scl(W a, W b, ST s) { return a * s * b; }
};
// 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,
@@ -573,6 +574,11 @@ 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_;
@@ -916,6 +922,11 @@ 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,6 +1231,122 @@ 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;
@@ -1301,38 +1417,7 @@ void TExpr::exec(const Mat* const* inputs, Mat* outputs)
}
// Temps present: walk the strip in L1-sized fragments so the intermediates stay hot.
// 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);
}
}
execFastStripWithTemps(*this, inputs, outputs, total);
return;
}
}
@@ -1466,70 +1551,38 @@ void TExpr::exec(const Mat* const* inputs, Mat* outputs)
body.nbuffers = nbuffers;
body.capElems = capElems;
broadcastOp(arr.data(), narr, [&](const BrTile& tile)
broadcastOp(arr.data(), narr, [&bc=body, this](const BrTile& tile)
{
EwBody& bc = body;
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.
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;
// 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();
AutoBuffer<BrSlice, LOCAL_OPS> args(nsl);
for (int y0 = 0; y0 < h; y0 += bh)
for (int s = 1; s < nsl; s++)
{
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;
}
}
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 };
}
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);
}
}
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, w, h);
}
}, true, nstripes);
}

View File

@@ -465,9 +465,17 @@ namespace {
#if defined HAVE_TBB
#if TBB_INTERFACE_VERSION >= 8000
static tbb::task_arena tbbArena(tbb::task_arena::automatic);
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));
}
#else
static tbb::task_scheduler_init tbbScheduler(tbb::task_scheduler_init::deferred);
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));
}
#endif
#elif defined HAVE_HPX
// nothing for HPX
@@ -582,7 +590,7 @@ static void parallel_for_impl(const cv::Range& range, const cv::ParallelLoopBody
#if defined HAVE_TBB
#if TBB_INTERFACE_VERSION >= 8000
tbbArena.execute(pbody);
getTbbArena().execute(pbody);
#else
pbody();
#endif
@@ -651,13 +659,13 @@ int getNumThreads(void)
#if defined HAVE_TBB
#if TBB_INTERFACE_VERSION >= 9100
return tbbArena.max_concurrency();
return getTbbArena().max_concurrency();
#elif TBB_INTERFACE_VERSION >= 8000
return numThreads > 0
? numThreads
: tbb::task_scheduler_init::default_num_threads();
#else
return tbbScheduler.is_active()
return getTbbScheduler().is_active()
? numThreads
: tbb::task_scheduler_init::default_num_threads();
#endif
@@ -735,11 +743,11 @@ void setNumThreads( int threads_ )
#ifdef HAVE_TBB
#if TBB_INTERFACE_VERSION >= 8000
if(tbbArena.is_active()) tbbArena.terminate();
if(threads > 0) tbbArena.initialize(threads);
if(getTbbArena().is_active()) getTbbArena().terminate();
if(threads > 0) getTbbArena().initialize(threads);
#else
if(tbbScheduler.is_active()) tbbScheduler.terminate();
if(threads > 0) tbbScheduler.initialize(threads);
if(getTbbScheduler().is_active()) getTbbScheduler().terminate();
if(threads > 0) getTbbScheduler().initialize(threads);
#endif
#elif defined HAVE_HPX

View File

@@ -1837,6 +1837,10 @@ 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))
@@ -1914,7 +1918,7 @@ char *FileStorage::Impl::parseBase64(char *ptr, int indent, FileNode &collection
ival = (short) base64decoder.getUInt16();
break;
case CV_32U:
ival = base64decoder.getInt32();
ival = base64decoder.getUInt32();
break;
case CV_32S:
ival = base64decoder.getInt32();
@@ -2743,7 +2747,7 @@ FileNodeIterator& FileNodeIterator::readRaw( const String& fmt, void* _data0, si
FileNode node = *(*this);
if( node.isInt() )
{
int64_t ival = static_cast<int64_t>(elem_size == 8 ? (int64_t)node : (int)node);
int64_t ival = (int64_t)node;
switch( elem_type )
{
case CV_8U:

View File

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

View File

@@ -5129,4 +5129,65 @@ 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,6 +726,38 @@ 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,4 +2662,28 @@ 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,10 +28,6 @@ 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,
@@ -40,14 +36,10 @@ public:
if (!proto.empty())
proto = findDataFile(proto);
net = readNet(weights, proto);
// "*" 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])));
}
}
// "*" 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)));
else if (!pinnedInput.empty())
net.setInputShape(pinnedInput, shape(std::get<0>(inputs[0])));
// Set multiple inputs

View File

@@ -347,11 +347,6 @@ 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

@@ -1,100 +0,0 @@
// 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,44 +53,6 @@ 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()))
@@ -125,19 +87,4 @@ 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

@@ -1,79 +0,0 @@
// 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

@@ -1,108 +0,0 @@
// 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,26 +355,6 @@ 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);
@@ -442,26 +422,5 @@ void eltwise_bitxor_2(const Stream& stream, TensorSpan<T> output, TensorView<T>
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,8 +9,6 @@
#include "math.hpp"
#include <type_traits>
#include "../cuda4dnn/csl/nvcc_defs.hpp"
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
@@ -853,63 +851,6 @@ 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

@@ -1,90 +0,0 @@
// 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

@@ -1,118 +0,0 @@
// 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

@@ -1,162 +0,0 @@
// 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,11 +34,21 @@ namespace raw {
constexpr int MEAN_SQR_SUM_BLOCK_SIZE = 256;
// 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) {
// 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) {
__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);
@@ -58,59 +68,10 @@ namespace raw {
sum += WARP_SHFL_XOR(sum, offset);
sqr += WARP_SHFL_XOR(sqr, offset);
}
}
}
// 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];
if (lane == 0) {
means[outer_idx] = sum / inner_size;
sum_sqrs[outer_idx] = sqr;
}
means[row] = sum / inner_size;
sum_sqrs[row] = sqr;
}
}
@@ -220,56 +181,6 @@ 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,9 +41,8 @@ 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] = output[idx] == T(0) ? T(0) : T(1) / (output[idx] + epsilon);
output[idx] = T(1) / (output[idx] + epsilon);
}
template <class T>
@@ -61,8 +60,7 @@ 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;
// 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);
output[idx] = T(1) / sqrt(output[idx] + epsilon);
}
}
@@ -117,51 +115,9 @@ 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

@@ -1,78 +0,0 @@
// 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,9 +16,6 @@
#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>
@@ -67,114 +64,8 @@ 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

@@ -1,177 +0,0 @@
// 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

@@ -1,72 +0,0 @@
// 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,76 +242,4 @@ 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

@@ -1,97 +0,0 @@
// 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

@@ -1,161 +0,0 @@
// 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

@@ -1,107 +0,0 @@
// 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

@@ -1,117 +0,0 @@
// 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

@@ -1,287 +0,0 @@
// 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

@@ -1,75 +0,0 @@
// 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,6 +11,7 @@
#include "memory.hpp"
#include <opencv2/core.hpp>
#include <opencv2/core/utils/configuration.private.hpp>
#include <cublas_v2.h>

View File

@@ -180,10 +180,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace csl { namespace cu
return false;
};
// 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));
bool found = 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

@@ -1,23 +0,0 @@
// 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

@@ -1,29 +0,0 @@
// 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

@@ -1,35 +0,0 @@
// 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,19 +42,6 @@ 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

@@ -1,28 +0,0 @@
// 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

@@ -1,30 +0,0 @@
// 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

@@ -1,37 +0,0 @@
// 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,14 +18,6 @@ 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,14 +19,6 @@ 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

@@ -1,22 +0,0 @@
// 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,16 +20,6 @@ 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

@@ -1,38 +0,0 @@
// 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

@@ -1,20 +0,0 @@
// 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,18 +7,9 @@
#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

@@ -1,34 +0,0 @@
// 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

@@ -1,44 +0,0 @@
// 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

@@ -1,27 +0,0 @@
// 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

@@ -1,28 +0,0 @@
// 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

@@ -1,44 +0,0 @@
// 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

@@ -1,22 +0,0 @@
// 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,37 +53,6 @@ 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

@@ -1,65 +0,0 @@
// 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,7 +12,6 @@
#include "../csl/tensor.hpp"
#include "../kernels/average_pooling.hpp"
#include "../kernels/reduce.hpp"
#include <opencv2/core.hpp>
@@ -59,35 +58,11 @@ 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

@@ -1,150 +0,0 @@
// 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,12 +354,9 @@ namespace cv { namespace dnn { namespace cuda4dnn {
builder.require<T>(sz);
}
#if defined(HAVE_CUDNNJIT) && !defined(HAVE_CUDNN)
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>()));
}
// 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());
@@ -389,12 +386,8 @@ namespace cv { namespace dnn { namespace cuda4dnn {
}
#if defined(HAVE_CUDNNJIT) && !defined(HAVE_CUDNN)
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));
}
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));
#endif
auto conv_scratchpad = allocator.get_instance();

View File

@@ -1,67 +0,0 @@
// 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,15 +40,17 @@ namespace cv { namespace dnn { namespace cuda4dnn {
scratch_mem_in_bytes = builder.required_workspace_size();
}
void forward(const std::vector<UMat> &inputs,
const std::vector<UMat> &outputs,
void forward(const std::vector<cv::Ptr<BackendWrapper>> &inputs,
const std::vector<cv::Ptr<BackendWrapper>> &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 = csl::viewOf<T>(inputs.front());
auto input_wrapper = inputs.front().dynamicCast<wrapper_type>();
auto input = input_wrapper->getView();
CV_CheckEQ(input.rank(), size_t(4), "DepthSpaceOps: input needs to be 4-dimensional [N, C, H, W]");
auto output = csl::spanOf<T>(outputs.front());
auto output_wrapper = outputs.front().dynamicCast<wrapper_type>();
auto output = output_wrapper->getSpan();
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

@@ -1,51 +0,0 @@
// 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

@@ -1,108 +0,0 @@
// 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,7 +16,6 @@
#include <opencv2/core.hpp>
#include <cstddef>
#include <type_traits>
#include <vector>
#include <utility>
@@ -32,10 +31,6 @@ namespace cv { namespace dnn { namespace cuda4dnn {
MOD,
FMOD,
POW,
PRELU,
BIT_AND,
BIT_OR,
BIT_XOR,
};
class EltwiseOpBase : public CUDABackendNode {
@@ -94,10 +89,6 @@ 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]);
@@ -132,31 +123,10 @@ 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

@@ -1,71 +0,0 @@
// 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

@@ -1,104 +0,0 @@
// 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

@@ -1,61 +0,0 @@
// 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, size_t group_elems)
GroupNormOp(csl::Stream stream_, float epsilon_, size_t loops, size_t num_groups)
: 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<UMat>& inputs,
const std::vector<UMat>& outputs,
void forward(const std::vector<cv::Ptr<BackendWrapper>>& inputs,
const std::vector<cv::Ptr<BackendWrapper>>& outputs,
csl::Workspace& workspace) override {
CV_Assert(inputs.size() == 3 && outputs.size() == 1);
auto input_wrapper = inputs[0].dynamicCast<wrapper_type>();
auto scale_wrapper = inputs[1].dynamicCast<wrapper_type>();
auto bias_wrapper = inputs[2].dynamicCast<wrapper_type>();
auto input = csl::viewOf<T>(inputs[0]);
auto scale = csl::viewOf<T>(inputs[1]);
auto bias = csl::viewOf<T>(inputs[2]);
auto input = input_wrapper->getView();
auto scale = scale_wrapper->getView();
auto bias = bias_wrapper->getView();
auto output = csl::spanOf<T>(outputs[0]);
auto output_wrapper = outputs[0].dynamicCast<wrapper_type>();
auto output = output_wrapper->getSpan();
auto C = input.get_axis_size(1);
auto loops = input.size_range(0, 2);
@@ -64,9 +64,8 @@ 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, partials);
kernels::reduce_mean_sqr_sum<T>(stream, mean, stdev, input, norm_size * group_size);
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);
}
@@ -78,7 +77,6 @@ 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, size_t norm_size)
InstanceNormOp(csl::Stream stream_, float epsilon_, size_t loops)
: 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<UMat>& inputs,
const std::vector<UMat>& outputs,
void forward(const std::vector<cv::Ptr<BackendWrapper>>& inputs,
const std::vector<cv::Ptr<BackendWrapper>>& outputs,
csl::Workspace& workspace) override {
CV_Assert(inputs.size() == 3 && outputs.size() == 1);
auto input_wrapper = inputs[0].dynamicCast<wrapper_type>();
auto scale_wrapper = inputs[1].dynamicCast<wrapper_type>();
auto bias_wrapper = inputs[2].dynamicCast<wrapper_type>();
auto input = csl::viewOf<T>(inputs[0]);
auto scale = csl::viewOf<T>(inputs[1]);
auto bias = csl::viewOf<T>(inputs[2]);
auto input = input_wrapper->getView();
auto scale = scale_wrapper->getView();
auto bias = bias_wrapper->getView();
auto output = csl::spanOf<T>(outputs[0]);
auto output_wrapper = outputs[0].dynamicCast<wrapper_type>();
auto output = output_wrapper->getSpan();
auto C = input.get_axis_size(1);
auto loops = input.size_range(0, 2);
@@ -62,9 +62,8 @@ 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, partials);
kernels::reduce_mean_sqr_sum<T>(stream, mean, stdev, input, norm_size);
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);
}
@@ -77,7 +76,6 @@ namespace cv { namespace dnn { namespace cuda4dnn {
float epsilon;
std::size_t reduce_scratch;
std::size_t scratch_mem_in_bytes;
};

View File

@@ -77,8 +77,7 @@ namespace cv { namespace dnn { namespace cuda4dnn {
auto loops = input.size_range(0, axis);
auto norm_size = input.size_range(axis, input.rank());
// with a size-one axis the result is 0, or the bias when there is one
if (norm_size == 1 && bias.empty()) {
if (norm_size == 1) {
kernels::fill<T>(stream, output, 0.f);
return;
} else {
@@ -115,9 +114,7 @@ namespace cv { namespace dnn { namespace cuda4dnn {
auto loops = input.size_range(0, axis);
auto norm_size = input.size_range(axis, input.rank());
// with a size-one axis the result is 0, or the bias when there is one;
// the optional Mean/InvStdDev outputs still need the statistics
if (norm_size == 1 && outputs.size() == 1 && bias.empty()) {
if (norm_size == 1) {
kernels::fill<T>(stream, output, 0.f);
return;
}
@@ -134,12 +131,6 @@ 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,11 +10,8 @@
#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 {
@@ -73,38 +70,6 @@ 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,14 +87,17 @@ namespace cv { namespace dnn { namespace cuda4dnn {
}
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
const std::vector<cv::Ptr<BackendWrapper>>& inputs,
const std::vector<cv::Ptr<BackendWrapper>>& outputs,
csl::Workspace& workspace) override
{
CV_Assert(inputs.size() == 1 && outputs.size() == 1);
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
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();
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

@@ -1,76 +0,0 @@
// 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,7 +17,6 @@
#include <opencv2/core.hpp>
#include <cstddef>
#include <cstring>
#include <vector>
#include <algorithm>
#include <utility>
@@ -114,37 +113,6 @@ 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,8 +17,9 @@
#include "../csl/tensor.hpp"
#include "../csl/tensor_ops.hpp"
// max pooling with left padding uses the hand kernel in every build
#if defined(HAVE_CUDNNJIT) && !defined(HAVE_CUDNN)
#include "../kernels/max_unpooling.hpp"
#endif
#include <opencv2/core.hpp>
@@ -180,21 +181,13 @@ 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)
// 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];
/* 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
}
}
@@ -244,20 +237,12 @@ 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);
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);
}
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);
#endif
}
@@ -271,14 +256,11 @@ 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);
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)
kernels::max_pooling_with_indices<T, std::int64_t>(stream, output, csl::TensorSpan<std::int64_t>(), input,
poolWindow, poolStrides, poolPadsBegin);
#else
if (!transformedInput.empty())
{
inputTransformer.transform(input, transformedInput);
@@ -293,10 +275,9 @@ namespace cv { namespace dnn { namespace cuda4dnn {
csl::Stream stream;
csl::cudnn::Handle cudnnHandle;
// max pooling through the hand kernel: always in JIT builds, for left padding in full cuDNN builds
bool useIndexKernel = false;
#if defined(HAVE_CUDNNJIT) && !defined(HAVE_CUDNN)
std::vector<std::size_t> poolWindow, poolStrides, poolPadsBegin;
#if !defined(HAVE_CUDNNJIT) || defined(HAVE_CUDNN)
#else
csl::Pooling<T> pooler;
csl::Tensor<T> transformedInput;

View File

@@ -1,138 +0,0 @@
// 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

@@ -1,52 +0,0 @@
// 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,10 +10,8 @@
#include "../csl/stream.hpp"
#include "../kernels/resize.hpp"
#include "../csl/span.hpp"
#include <utility>
#include <vector>
namespace cv { namespace dnn { namespace cuda4dnn {
@@ -42,15 +40,18 @@ namespace cv { namespace dnn { namespace cuda4dnn {
}
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
const std::vector<cv::Ptr<BackendWrapper>>& inputs,
const std::vector<cv::Ptr<BackendWrapper>>& outputs,
csl::Workspace& workspace) override
{
// roi/scales/sizes may follow the data input; the output shape already carries them
CV_Assert(inputs.size() >= 1 && outputs.size() == 1);
// 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);
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
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();
const auto compute_scale = [this](std::size_t input_size, std::size_t output_size) {
return (align_corners && output_size > 1) ?
@@ -75,63 +76,6 @@ 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

@@ -1,61 +0,0 @@
// 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

@@ -1,68 +0,0 @@
// 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

@@ -1,87 +0,0 @@
// 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

@@ -1,88 +0,0 @@
// 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

@@ -1,47 +0,0 @@
// 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

@@ -1,75 +0,0 @@
// 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

@@ -1,60 +0,0 @@
// 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,9 +180,6 @@ 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,11 +4,7 @@
#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 {
@@ -47,31 +43,9 @@ 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,7 +40,6 @@
//
//M*/
#include "../precomp.hpp"
#include "../net_impl.hpp"
#include "../op_cuda.hpp"
#include "../op_inf_engine.hpp"
#include "../ie_ngraph.hpp"
@@ -68,17 +67,6 @@ 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 ||
@@ -191,16 +179,15 @@ public:
#ifdef HAVE_CUDA
Ptr<BackendNode> initCUDA(
void *context_,
InputArrayOfArrays inputs_arr,
InputArrayOfArrays
const std::vector<Ptr<BackendWrapper>>& inputs,
const std::vector<Ptr<BackendWrapper>>& outputs
) override
{
auto context = reinterpret_cast<csl::CSLContext*>(context_);
const int depth = inputs_arr.depth(0);
if (depth == CV_Bool)
if (inputs[0]->getHostMatDepth() == CV_Bool)
return make_cuda_node_bool<cuda4dnn::ReshapeOp>(std::move(context->stream));
else
return make_cuda_node_with_type<cuda4dnn::ReshapeOp>(preferableTarget, depth, std::move(context->stream));
return make_cuda_node_with_type<cuda4dnn::ReshapeOp>(preferableTarget, inputs[0]->getHostMatDepth(), std::move(context->stream));
}
#endif
};

View File

@@ -115,13 +115,9 @@ 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 ||
t == CV_8U || t == CV_8S; };
auto castable = [](int t) { return t < 0 || t == CV_Bool || t == CV_32S || t == CV_64S || t == CV_32F; };
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);
@@ -151,8 +147,6 @@ 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>();
@@ -167,8 +161,6 @@ 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,14 +24,8 @@ 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>
@@ -119,51 +113,9 @@ 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,7 +16,6 @@
#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
@@ -548,20 +547,6 @@ 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)
@@ -1214,12 +1199,13 @@ public:
return false;
if (!activ.empty())
return false;
// a fused kernel activation runs on CUDA inside ConvolutionOp or as a post-op after it
if (activationFunc != nullptr && !cudaConvKernel(activationFunc) && !cudaPostKernel(activationFunc))
// 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))
return false;
if (fastActivation != FAST_ACTIV_NONE && fastActivation != FAST_ACTIV_RELU &&
fastActivation != FAST_ACTIV_LEAKY_RELU && fastActivation != FAST_ACTIV_CLIP &&
fastActivation != FAST_ACTIV_PRELU)
fastActivation != FAST_ACTIV_LEAKY_RELU && fastActivation != FAST_ACTIV_CLIP)
return false;
return true;
}
@@ -1270,28 +1256,11 @@ 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) && post.empty();
bool hasAct = fastActivation != FAST_ACTIV_NONE || activationFunc != nullptr;
if (activationFunc == getActivationFunc(ACTIV_SIGMOID)) {
config.activation_type = ConvolutionConfiguration::ActivationType::SIGMOID;
} else if (activationFunc == getActivationFunc(ACTIV_SWISH)) {
@@ -1320,12 +1289,9 @@ public:
else
config.fusion_mode = ConvolutionConfiguration::FusionMode::NONE;
Ptr<BackendNode> convNode = make_cuda_node<cuda4dnn::ConvolutionOp>(
return 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,11 +12,6 @@
#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
@@ -101,86 +96,6 @@ 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
{
@@ -195,7 +110,6 @@ 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())
@@ -604,8 +518,26 @@ public:
pads, adjust_pads, auto_pad);
applyExplicitOutShape(outshape);
// block layout: N, C/C0, spatial..., C0
std::vector<int> pads_resolved = resolvePads(inpshape, outshape, wshape0, inpshape.dims - 3);
// 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;
}
}
int outtype = inferType(inptype);
@@ -662,7 +594,6 @@ 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,11 +6,6 @@
#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>
@@ -34,44 +29,7 @@ public:
setParamsFrom(params);
}
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 supportBackend(int backendId) CV_OVERRIDE { return backendId == DNN_BACKEND_OPENCV; }
bool getMemoryShapes(const std::vector<MatShape>& inputs, const int,
std::vector<MatShape>& outputs, std::vector<MatShape>&) const CV_OVERRIDE

View File

@@ -6,11 +6,6 @@
#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>
@@ -32,44 +27,10 @@ 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