Administrator/opencv-MIRROR
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2026-11-11 02:59:00 -06:00
Code Issues Packages Projects Releases Wiki Activity

Compare commits

base: Administrator:5.x
Branches Tags
Administrator:4.x Administrator:5.x Administrator:cuda/final_coverage Administrator:copilot/fix-windows-actions-job Administrator:4.13.0-python Administrator:3.4 Administrator:2.4
Administrator:5.1.0 Administrator:4.14.0 Administrator:5.0.0 Administrator:4.13.0 Administrator:4.12.0 Administrator:4.11.0 Administrator:5.0.0-alpha Administrator:4.10.0-kleidicv Administrator:4.10.0 Administrator:4.9.0 Administrator:4.8.1 Administrator:4.8.0 Administrator:3.4.20 Administrator:4.7.0 Administrator:3.4.19 Administrator:4.6.0 Administrator:3.4.18 Administrator:4.5.5-openvino-2022.1.0 Administrator:4.5.5 Administrator:3.4.17 Administrator:4.5.3-openvino-2021.4.2 Administrator:4.5.4 Administrator:3.4.16 Administrator:4.5.3-openvino-2021.4.1 Administrator:4.5.3 Administrator:3.4.15 Administrator:4.5.3-openvino Administrator:4.5.2 Administrator:3.4.14 Administrator:4.5.2-openvino Administrator:4.5.1 Administrator:3.4.13 Administrator:4.5.1-openvino Administrator:4.5.0 Administrator:3.4.12 Administrator:4.5.0-openvino Administrator:4.4.0 Administrator:3.4.11 Administrator:4.4.0-openvino Administrator:4.3.0-openvino-2020.3.0 Administrator:4.3.0 Administrator:3.4.10 Administrator:4.3.0-openvino Administrator:4.2.0-openvino Administrator:4.2.0 Administrator:3.4.9 Administrator:4.1.2 Administrator:3.4.8 Administrator:4.1.2-openvino Administrator:4.1.1 Administrator:3.4.7 Administrator:4.1.1-openvino Administrator:4.1.0 Administrator:3.4.6 Administrator:4.1.0-openvino Administrator:4.0.1 Administrator:3.4.5 Administrator:4.0.1-openvino Administrator:4.0.0 Administrator:3.4.4 Administrator:4.0.0-openvino Administrator:4.0.0-rc Administrator:4.0.0-beta Administrator:4.0.0-alpha Administrator:3.4.3 Administrator:3.4.3-openvino Administrator:3.4.2 Administrator:2.4.13.7 Administrator:3.4.2-openvino Administrator:3.4.1-cvsdk Administrator:3.4.1 Administrator:2.4.13.6 Administrator:3.4.0 Administrator:2.4.13.5 Administrator:3.4.0-rc Administrator:3.3.1 Administrator:2.4.13.4 Administrator:3.3.1-cvsdk Administrator:3.3.0 Administrator:3.3.0-cvsdk Administrator:2.4.13.3 Administrator:3.3.0-rc Administrator:3.2.0 Administrator:3.2.0-rc Administrator:2.4.13.2 Administrator:2.4.13.1 Administrator:2.4.13 Administrator:3.1.0 Administrator:2.4.12.3 Administrator:2.4.12.2 Administrator:2.4.12.1 Administrator:2.4.12 Administrator:3.0.0 Administrator:2.4.10.4 Administrator:2.4.10.3 Administrator:3.0.0-rc1 Administrator:2.4.10.2 Administrator:2.4.11 Administrator:3.0.0-beta Administrator:2.4.10.1 Administrator:2.4.10 Administrator:2.4.9.1 Administrator:3.0.0-alpha Administrator:2.4.8.3 Administrator:2.4.9 Administrator:2.4.8.2 Administrator:3.0-ocl-tp2 Administrator:2.4.8.1 Administrator:3.0-ocl-tech-preview Administrator:2.4.8 Administrator:2.4.7.2 Administrator:2.4.7.1 Administrator:2.4.7 Administrator:2.4.7-rc1 Administrator:2.4.6.2r3 Administrator:2.4.6.2r2 Administrator:2.4.6.2 Administrator:2.4.6.2-rc1 Administrator:2.4.6.1 Administrator:2.4.6 Administrator:2.4.5 Administrator:2.4.4 Administrator:2.4.4-beta Administrator:2.4.3.2 Administrator:2.4.3.1 Administrator:2.4.3 Administrator:2.4.3-rc Administrator:2.4.2 Administrator:2.4.1 Administrator:2.4.0 Administrator:2.3.1 Administrator:2.3.0 Administrator:2.2
..
compare: Administrator:cuda/final_coverage
Branches Tags
Administrator:5.x Administrator:4.x Administrator:cuda/final_coverage Administrator:copilot/fix-windows-actions-job Administrator:4.13.0-python Administrator:3.4 Administrator:2.4
Administrator:5.1.0 Administrator:4.14.0 Administrator:5.0.0 Administrator:4.13.0 Administrator:4.12.0 Administrator:4.11.0 Administrator:5.0.0-alpha Administrator:4.10.0-kleidicv Administrator:4.10.0 Administrator:4.9.0 Administrator:4.8.1 Administrator:4.8.0 Administrator:3.4.20 Administrator:4.7.0 Administrator:3.4.19 Administrator:4.6.0 Administrator:3.4.18 Administrator:4.5.5-openvino-2022.1.0 Administrator:4.5.5 Administrator:3.4.17 Administrator:4.5.3-openvino-2021.4.2 Administrator:4.5.4 Administrator:3.4.16 Administrator:4.5.3-openvino-2021.4.1 Administrator:4.5.3 Administrator:3.4.15 Administrator:4.5.3-openvino Administrator:4.5.2 Administrator:3.4.14 Administrator:4.5.2-openvino Administrator:4.5.1 Administrator:3.4.13 Administrator:4.5.1-openvino Administrator:4.5.0 Administrator:3.4.12 Administrator:4.5.0-openvino Administrator:4.4.0 Administrator:3.4.11 Administrator:4.4.0-openvino Administrator:4.3.0-openvino-2020.3.0 Administrator:4.3.0 Administrator:3.4.10 Administrator:4.3.0-openvino Administrator:4.2.0-openvino Administrator:4.2.0 Administrator:3.4.9 Administrator:4.1.2 Administrator:3.4.8 Administrator:4.1.2-openvino Administrator:4.1.1 Administrator:3.4.7 Administrator:4.1.1-openvino Administrator:4.1.0 Administrator:3.4.6 Administrator:4.1.0-openvino Administrator:4.0.1 Administrator:3.4.5 Administrator:4.0.1-openvino Administrator:4.0.0 Administrator:3.4.4 Administrator:4.0.0-openvino Administrator:4.0.0-rc Administrator:4.0.0-beta Administrator:4.0.0-alpha Administrator:3.4.3 Administrator:3.4.3-openvino Administrator:3.4.2 Administrator:2.4.13.7 Administrator:3.4.2-openvino Administrator:3.4.1-cvsdk Administrator:3.4.1 Administrator:2.4.13.6 Administrator:3.4.0 Administrator:2.4.13.5 Administrator:3.4.0-rc Administrator:3.3.1 Administrator:2.4.13.4 Administrator:3.3.1-cvsdk Administrator:3.3.0 Administrator:3.3.0-cvsdk Administrator:2.4.13.3 Administrator:3.3.0-rc Administrator:3.2.0 Administrator:3.2.0-rc Administrator:2.4.13.2 Administrator:2.4.13.1 Administrator:2.4.13 Administrator:3.1.0 Administrator:2.4.12.3 Administrator:2.4.12.2 Administrator:2.4.12.1 Administrator:2.4.12 Administrator:3.0.0 Administrator:2.4.10.4 Administrator:2.4.10.3 Administrator:3.0.0-rc1 Administrator:2.4.10.2 Administrator:2.4.11 Administrator:3.0.0-beta Administrator:2.4.10.1 Administrator:2.4.10 Administrator:2.4.9.1 Administrator:3.0.0-alpha Administrator:2.4.8.3 Administrator:2.4.9 Administrator:2.4.8.2 Administrator:3.0-ocl-tp2 Administrator:2.4.8.1 Administrator:3.0-ocl-tech-preview Administrator:2.4.8 Administrator:2.4.7.2 Administrator:2.4.7.1 Administrator:2.4.7 Administrator:2.4.7-rc1 Administrator:2.4.6.2r3 Administrator:2.4.6.2r2 Administrator:2.4.6.2 Administrator:2.4.6.2-rc1 Administrator:2.4.6.1 Administrator:2.4.6 Administrator:2.4.5 Administrator:2.4.4 Administrator:2.4.4-beta Administrator:2.4.3.2 Administrator:2.4.3.1 Administrator:2.4.3 Administrator:2.4.3-rc Administrator:2.4.2 Administrator:2.4.1 Administrator:2.4.0 Administrator:2.3.1 Administrator:2.3.0 Administrator:2.2

5 Commits
5.x ... cuda/final

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

2
doc/tutorials/app/raster_io_gdal.markdown
View File

@@ -71,7 +71,7 @@ Mercator.
One easy method to find the corner coordinates of an image is to use the command-line tool gdalinfo.
For imagery which is ortho-rectified and contains the projection information, you can use the [USGS
EarthExplorer](https://earthexplorer.usgs.gov).
EarthExplorer](http://http://earthexplorer.usgs.gov).
@code{.bash}
\f$> gdalinfo N37W123.hgt

2
doc/tutorials/dnn/dnn_text_spotting/dnn_text_spotting.markdown
View File

@@ -50,7 +50,7 @@ In the current version, @ref cv::dnn::TextDetectionModel supports these algorith
- and use @ref cv::dnn::TextDetectionModel_EAST with "EAST" models
The following provided pretrained models are variants of DB (w/o deformable convolution),
and the performance can be referred to the Table.1 in the [paper](https://arxiv.org/abs/1911.08947).
and the performance can be referred to the Table.1 in the [paper]((https://arxiv.org/abs/1911.08947)).
For more information, please refer to the [official code](https://github.com/MhLiao/DB)
---

4
doc/tutorials/introduction/windows_install/windows_install.markdown
View File

@@ -158,7 +158,7 @@ of them, you need to download and install them on your system.
- The latest [CUDA Toolkit](http://developer.nvidia.com/cuda-downloads) will allow you to use the power lying inside your GPU. This will
drastically improve performance for some algorithms (e.g the HOG descriptor). Getting more and
more of our algorithms to work on the GPUs is a constant effort of the OpenCV team.
- [OpenEXR](https://github.com/AcademySoftwareFoundation/openexr/releases) source files are required for the library to work with this high dynamic range (HDR)
- [OpenEXR](http://www.openexr.com/downloads.html) source files are required for the library to work with this high dynamic range (HDR)
image file format.
- The OpenNI Framework contains a set of open source APIs that provide support for natural interaction with devices via methods such as voice command recognition, hand gestures, and body
motion tracking. Prebuilt binaries can be found [here](http://structure.io/openni). The source code of [OpenNI](https://github.com/OpenNI/OpenNI) and [OpenNI2](https://github.com/OpenNI/OpenNI2) are also available on Github.
@@ -212,7 +212,7 @@ libraries). If you do not need the support for some of these, you can just freel
-# In case of the [Eigen](http://eigen.tuxfamily.org/index.php?title=Main_Page#Download) library it is again a case of download and extract to the
`D:/OpenCV/dep` directory.
-# Same as above with [OpenEXR](https://github.com/AcademySoftwareFoundation/openexr/releases).
-# Same as above with [OpenEXR](http://www.openexr.com/downloads.html).
-# For the [OpenNI Framework](http://www.openni.org/) you need to install both the [development
build](http://www.openni.org/downloadfiles/opennimodules/openni-binaries/21-stable) and the
[PrimeSensor

1
docs_sphinx/_static/version.js
View File

@@ -12,7 +12,6 @@
* the dropdown on every version's pages at once. Do not maintain a second
* list anywhere. */
window.OPENCV_DOC_VERSIONS = [
['5.1', '/5.1'],
['4.13.0', '/4.13.0'],
['4.12.0', '/4.12.0'],
['4.11.0', '/4.11.0'],

2
docs_sphinx/conf.py
View File

@@ -24,7 +24,7 @@ from conf_helpers.postprocess import _inline_coll_graphs_on_finish
# -- Project ----------------------------------------------------------------
project = "OpenCV"
author = "OpenCV Team"
release = "5.1"
release = "5.0"
# -- Sphinx core ------------------------------------------------------------
extensions = ["myst_parser"]

8
hal/ipp/include/ipp_hal_core.hpp
View File

@@ -35,14 +35,6 @@ int ipp_hal_minMaxIdxMaskStep(const uchar* src_data, size_t src_step, int width,
# define IPP_DISABLE_NORM_INF_16U_C1MR 0
#endif
// segmentation fault in ippiNorm[Diff]_Inf_32f_C3CMR with the macOS ippicv (2021.9.1);
// the newer Linux/Windows ippicv (2026.0.0) is not affected
#if defined(__APPLE__)
# define IPP_DISABLE_NORM_INF_32F_C3CMR 1
# else
# define IPP_DISABLE_NORM_INF_32F_C3CMR 0
#endif
int ipp_hal_norm(const uchar* src, size_t src_step, const uchar* mask, size_t mask_step,
int width, int height, int type, int norm_type, double* result);

4
hal/ipp/src/norm_ipp.cpp
View File

@@ -50,9 +50,7 @@ int ipp_hal_norm(const uchar* src, size_t src_step, const uchar* mask, size_t ma
norm_type == cv::NORM_INF ?
(type == CV_8UC3 ? (ippiMaskNormFuncC3)ippiNorm_Inf_8u_C3CMR :
type == CV_16UC3 ? (ippiMaskNormFuncC3)ippiNorm_Inf_16u_C3CMR :
#if (!IPP_DISABLE_NORM_INF_32F_C3CMR)
type == CV_32FC3 ? (ippiMaskNormFuncC3)ippiNorm_Inf_32f_C3CMR :
#endif
0) :
norm_type == cv::NORM_L1 ?
(type == CV_8UC3 ? (ippiMaskNormFuncC3)ippiNorm_L1_8u_C3CMR :
@@ -267,9 +265,7 @@ int ipp_hal_normDiff(const uchar* src1, size_t src1_step, const uchar* src2, siz
norm_type == cv::NORM_INF ?
(type == CV_8UC3 ? (ippiMaskNormDiffFuncC3)ippiNormDiff_Inf_8u_C3CMR :
type == CV_16UC3 ? (ippiMaskNormDiffFuncC3)ippiNormDiff_Inf_16u_C3CMR :
#if (!IPP_DISABLE_NORM_INF_32F_C3CMR)
type == CV_32FC3 ? (ippiMaskNormDiffFuncC3)ippiNormDiff_Inf_32f_C3CMR :
#endif
0) :
norm_type == cv::NORM_L1 ?
(type == CV_8UC3 ? (ippiMaskNormDiffFuncC3)ippiNormDiff_L1_8u_C3CMR :

1
modules/core/include/opencv2/core/base.hpp
View File

@@ -321,7 +321,6 @@ inline int cv_absdiff(short x, short y) { return (int)std::abs((int)x - (int)y);
inline unsigned cv_absdiff(int x, int y) { return (unsigned)std::max(x, y) - (unsigned)std::min(x, y); }
inline unsigned cv_absdiff(unsigned x, unsigned y) { return std::max(x, y) - std::min(x, y); }
inline uint64 cv_absdiff(uint64 x, uint64 y) { return std::max(x, y) - std::min(x, y); }
inline uint64 cv_absdiff(int64 x, int64 y) { return (uint64)std::max(x, y) - (uint64)std::min(x, y); }
inline float cv_absdiff(hfloat x, hfloat y) { return std::abs((float)x - (float)y); }
inline float cv_absdiff(bfloat x, bfloat y) { return std::abs((float)x - (float)y); }
inline float cv_absdiff(fp8_t x, fp8_t y) { return std::abs((float)x - (float)y); }

2
modules/core/include/opencv2/core/hal/intrin_avx.hpp
View File

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

2
modules/core/include/opencv2/core/hal/intrin_avx512.hpp
View File

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

4
modules/core/include/opencv2/core/hal/intrin_lasx.hpp
View File

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

4
modules/core/include/opencv2/core/hal/intrin_lsx.hpp
View File

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

12
modules/core/include/opencv2/core/parallel/backend/parallel_for.tbb.hpp
View File

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

4
modules/core/include/opencv2/core/version.hpp
View File

@@ -6,9 +6,9 @@
#define OPENCV_VERSION_HPP
#define CV_VERSION_MAJOR 5
#define CV_VERSION_MINOR 2
#define CV_VERSION_MINOR 1
#define CV_VERSION_REVISION 0
#define CV_VERSION_STATUS "-dev"
#define CV_VERSION_STATUS "-pre"
#define CVAUX_STR_EXP(__A) #__A
#define CVAUX_STR(__A) CVAUX_STR_EXP(__A)

2
modules/core/perf/perf_arithm.cpp
View File

@@ -528,7 +528,7 @@ PERF_TEST_P_(BinaryOpTest, transpose_generic_move_tail_order)
INSTANTIATE_TEST_CASE_P(/*nothing*/ , BinaryOpTest,
testing::Combine(
testing::Values(szVGA, sz720p, sz1080p, Size(127, 61)), // 127x61 guards per-call overhead
testing::Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_8SC1, CV_16UC1, CV_16SC1, CV_16SC2, CV_16SC3, CV_16SC4, CV_32SC1, CV_32FC1, CV_64FC1)
testing::Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_8SC1, CV_16SC1, CV_16SC2, CV_16SC3, CV_16SC4, CV_32SC1, CV_32FC1, CV_64FC1)
)
);

2
modules/core/perf/perf_compare.cpp
View File

@@ -37,7 +37,7 @@ PERF_TEST_P( Size_MatType_CmpType, compare,
PERF_TEST_P( Size_MatType_CmpType, compareScalar,
testing::Combine(
testing::Values(TYPICAL_MAT_SIZES),
testing::Values(CV_8UC1, CV_8UC4, CV_16UC1, CV_16SC1, CV_32FC1),
testing::Values(TYPICAL_MAT_TYPES),
CmpType::all()
)
)

2
modules/core/perf/perf_dot.cpp
View File

@@ -9,7 +9,7 @@ typedef TestBaseWithParam<MatType_Length_t> MatType_Length;
PERF_TEST_P( MatType_Length, dot,
testing::Combine(
testing::Values( CV_8UC1, CV_8SC1, CV_16SC1, CV_16UC1, CV_32SC1, CV_32FC1, CV_64FC1 ),
testing::Values( CV_8UC1, CV_8SC1, CV_16SC1, CV_16UC1, CV_32SC1, CV_32FC1 ),
testing::Values( 32, 64, 128, 256, 512, 1024 )
))
{

2
modules/core/perf/perf_flip.cpp
View File

@@ -15,7 +15,7 @@ enum
};
#define FLIP_SIZES szQVGA, szVGA, sz1080p
#define FLIP_TYPES CV_8UC1, CV_8UC2, CV_8UC3, CV_8UC4, CV_8SC1, CV_16SC1, CV_16SC2, CV_16SC3, CV_16SC4, CV_32SC1, CV_32FC1, CV_32FC3, CV_32FC4
#define FLIP_TYPES CV_8UC1, CV_8UC2, CV_8UC3, CV_8UC4, CV_8SC1, CV_16SC1, CV_16SC2, CV_16SC3, CV_16SC4, CV_32SC1, CV_32FC1
#define FLIP_CODES FLIP_X, FLIP_Y, FLIP_XY
CV_FLAGS(FlipCode, FLIP_X, FLIP_Y, FLIP_XY);

4
modules/core/perf/perf_mat.cpp
View File

@@ -99,7 +99,7 @@ PERF_TEST_P(Size_MatType, Mat_Clone_Roi,
PERF_TEST_P(Size_MatType, Mat_CopyToWithMask,
testing::Combine(testing::Values(::perf::sz1080p, ::perf::szODD),
testing::Values(CV_8UC1, CV_8UC2, CV_8UC3, CV_16UC1, CV_16UC3, CV_32SC1, CV_32SC2, CV_32FC3, CV_32FC4))
testing::Values(CV_8UC1, CV_8UC2, CV_8UC3, CV_16UC1, CV_16UC3, CV_32SC1, CV_32SC2, CV_32FC4))
)
{
const Size_MatType_t params = GetParam();
@@ -119,7 +119,7 @@ PERF_TEST_P(Size_MatType, Mat_CopyToWithMask,
PERF_TEST_P(Size_MatType, Mat_SetToWithMask,
testing::Combine(testing::Values(TYPICAL_MAT_SIZES),
testing::Values(CV_8UC1, CV_8UC2, CV_8UC3, CV_8UC4, CV_16UC3, CV_16UC4, CV_32FC3, CV_32FC4))
testing::Values(CV_8UC1, CV_8UC2))
)
{
const Size_MatType_t params = GetParam();

12
modules/core/perf/perf_norm.cpp
View File

@@ -15,7 +15,7 @@ PERF_TEST_P(Size_MatType_NormType, norm,
testing::Combine(
testing::Values(TYPICAL_MAT_SIZES),
testing::Values(CV_8UC1, CV_8UC4, CV_8SC1, CV_16UC1, CV_16SC1, CV_32SC1, CV_32FC1, CV_64FC1),
testing::Values((int)NORM_INF, (int)NORM_L1, (int)NORM_L2, (int)NORM_L2SQR)
testing::Values((int)NORM_INF, (int)NORM_L1, (int)NORM_L2)
)
)
{
@@ -39,8 +39,8 @@ PERF_TEST_P(Size_MatType_NormType, norm_mask,
testing::Values(CV_8UC1, CV_8UC3, CV_8UC4,
CV_8SC1, CV_8SC3, CV_16UC1, CV_16UC3,
CV_16SC1, CV_16SC3, CV_32SC1, CV_32SC3,
CV_32FC1, CV_32FC3, CV_64FC1),
testing::Values((int)NORM_INF, (int)NORM_L1, (int)NORM_L2, (int)NORM_L2SQR)
CV_32FC1, CV_64FC1),
testing::Values((int)NORM_INF, (int)NORM_L1, (int)NORM_L2)
)
)
{
@@ -63,7 +63,7 @@ PERF_TEST_P(Size_MatType_NormType, norm2,
testing::Combine(
testing::Values(TYPICAL_MAT_SIZES),
testing::Values(CV_8UC1, CV_8UC4, CV_8SC1, CV_16UC1, CV_16SC1, CV_32SC1, CV_32FC1, CV_64FC1),
testing::Values((int)NORM_INF, (int)NORM_L1, (int)NORM_L2, (int)NORM_L2SQR, (int)(NORM_RELATIVE+NORM_INF), (int)(NORM_RELATIVE+NORM_L1), (int)(NORM_RELATIVE+NORM_L2), (int)(NORM_RELATIVE+NORM_L2SQR))
testing::Values((int)NORM_INF, (int)NORM_L1, (int)NORM_L2, (int)(NORM_RELATIVE+NORM_INF), (int)(NORM_RELATIVE+NORM_L1), (int)(NORM_RELATIVE+NORM_L2))
)
)
{
@@ -85,8 +85,8 @@ PERF_TEST_P(Size_MatType_NormType, norm2,
PERF_TEST_P(Size_MatType_NormType, norm2_mask,
testing::Combine(
testing::Values(TYPICAL_MAT_SIZES),
testing::Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_8SC1, CV_16UC1, CV_16UC3, CV_16SC1, CV_32SC1, CV_32FC1, CV_32FC3, CV_64FC1),
testing::Values((int)NORM_INF, (int)NORM_L1, (int)NORM_L2, (int)NORM_L2SQR, (int)(NORM_RELATIVE|NORM_INF), (int)(NORM_RELATIVE|NORM_L1), (int)(NORM_RELATIVE|NORM_L2), (int)(NORM_RELATIVE|NORM_L2SQR))
testing::Values(CV_8UC1, CV_8UC4, CV_8SC1, CV_16UC1, CV_16SC1, CV_32SC1, CV_32FC1, CV_64FC1),
testing::Values((int)NORM_INF, (int)NORM_L1, (int)NORM_L2, (int)(NORM_RELATIVE|NORM_INF), (int)(NORM_RELATIVE|NORM_L1), (int)(NORM_RELATIVE|NORM_L2))
)
)
{

2
modules/core/perf/perf_sort.cpp
View File

@@ -5,7 +5,7 @@ namespace opencv_test
using namespace perf;
#define TYPICAL_MAT_SIZES_SORT TYPICAL_MAT_SIZES
#define TYPICAL_MAT_TYPES_SORT CV_8UC1, CV_16UC1, CV_16SC1, CV_32SC1, CV_32FC1, CV_64FC1
#define TYPICAL_MAT_TYPES_SORT CV_8UC1, CV_16UC1, CV_32FC1
#define SORT_TYPES SORT_EVERY_ROW | SORT_ASCENDING, SORT_EVERY_ROW | SORT_DESCENDING
#define TYPICAL_MATS_SORT testing::Combine( testing::Values(TYPICAL_MAT_SIZES_SORT), testing::Values(TYPICAL_MAT_TYPES_SORT), testing::Values(SORT_TYPES) )

27
modules/core/perf/perf_stat.cpp
View File

@@ -4,12 +4,11 @@ namespace opencv_test
{
using namespace perf;
// The wide depths are not in TYPICAL_MAT_TYPES, so they were not measured; they are the ones
// whose scalar accumulation step this touches.
// The 32-bit integer depths are not in TYPICAL_MAT_TYPES, so they were not measured; they are the
// ones whose scalar accumulation step this touches.
PERF_TEST_P(Size_MatType, sum_wide,
testing::Combine(testing::Values(TYPICAL_MAT_SIZES),
testing::Values(CV_32SC1, CV_32UC1, CV_32SC3,
CV_64FC1, CV_64FC3)))
testing::Values(CV_32SC1, CV_32UC1, CV_32SC3)))
{
Size sz = get<0>(GetParam());
int type = get<1>(GetParam());
@@ -24,13 +23,7 @@ PERF_TEST_P(Size_MatType, sum_wide,
SANITY_CHECK_NOTHING();
}
// Masked variants: 8U/16U/32F x C1/C3/C4.
#define TYPICAL_MAT_TYPES_STAT_MASK CV_8UC1, CV_8UC3, CV_8UC4, CV_16UC1, CV_16UC3, CV_16UC4, CV_32FC1, CV_32FC3, CV_32FC4
// sum/mean/meanStdDev additionally cover 16S.
#define TYPICAL_MAT_TYPES_STAT TYPICAL_MAT_TYPES_STAT_MASK, CV_16SC1, CV_16SC3, CV_16SC4
PERF_TEST_P(Size_MatType, sum, testing::Combine( testing::Values( TYPICAL_MAT_SIZES ),
testing::Values( TYPICAL_MAT_TYPES_STAT ) ))
PERF_TEST_P(Size_MatType, sum, TYPICAL_MATS)
{
Size sz = get<0>(GetParam());
int type = get<1>(GetParam());
@@ -45,8 +38,7 @@ PERF_TEST_P(Size_MatType, sum, testing::Combine( testing::Values( TYPICAL_MAT_SI
SANITY_CHECK(s, 1e-6, ERROR_RELATIVE);
}
PERF_TEST_P(Size_MatType, mean, testing::Combine( testing::Values( TYPICAL_MAT_SIZES ),
testing::Values( TYPICAL_MAT_TYPES_STAT ) ))
PERF_TEST_P(Size_MatType, mean, TYPICAL_MATS)
{
Size sz = get<0>(GetParam());
int type = get<1>(GetParam());
@@ -61,8 +53,7 @@ PERF_TEST_P(Size_MatType, mean, testing::Combine( testing::Values( TYPICAL_MAT_S
SANITY_CHECK(s, 1e-5);
}
PERF_TEST_P(Size_MatType, mean_mask, testing::Combine( testing::Values( TYPICAL_MAT_SIZES ),
testing::Values( TYPICAL_MAT_TYPES_STAT_MASK ) ))
PERF_TEST_P(Size_MatType, mean_mask, TYPICAL_MATS)
{
Size sz = get<0>(GetParam());
int type = get<1>(GetParam());
@@ -78,8 +69,7 @@ PERF_TEST_P(Size_MatType, mean_mask, testing::Combine( testing::Values( TYPICAL_
SANITY_CHECK(s, 5e-5);
}
PERF_TEST_P(Size_MatType, meanStdDev, testing::Combine( testing::Values( TYPICAL_MAT_SIZES ),
testing::Values( TYPICAL_MAT_TYPES_STAT_MASK ) ))
PERF_TEST_P(Size_MatType, meanStdDev, TYPICAL_MATS)
{
Size sz = get<0>(GetParam());
int matType = get<1>(GetParam());
@@ -96,8 +86,7 @@ PERF_TEST_P(Size_MatType, meanStdDev, testing::Combine( testing::Values( TYPICAL
SANITY_CHECK(dev, 1e-5, ERROR_RELATIVE);
}
PERF_TEST_P(Size_MatType, meanStdDev_mask, testing::Combine( testing::Values( TYPICAL_MAT_SIZES ),
testing::Values( TYPICAL_MAT_TYPES_STAT_MASK ) ))
PERF_TEST_P(Size_MatType, meanStdDev_mask, TYPICAL_MATS)
{
Size sz = get<0>(GetParam());
int matType = get<1>(GetParam());

13
modules/core/src/arithm.simd.hpp
View File

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

227
modules/core/src/arithm_expr.cpp
View File

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

2
modules/core/src/mathfuncs_core.simd.hpp
View File

@@ -115,7 +115,7 @@ static void cartToPolar32f_(const float *X, const float *Y, float *mag, float *a
// if it's inplace operation, we cannot repeatedly process
// the tail for the second time, so we have to use the
// scalar code
if( i == 0 || mag == X || mag == Y || angle == X || angle == Y )
if( i == 0 || angle == X || angle == Y )
break;
i = len - VECSZ*2;
}

3
modules/core/src/mean.dispatch.cpp
View File

@@ -52,8 +52,7 @@ static Scalar mean_(const Mat& src, const Mat& mask)
int _buf[CV_CN_MAX];
int* buf = (int*)&s[0];
bool partialSumIsInt = depth < CV_32S;
bool blockSum = partialSumIsInt || depth == CV_16F || depth == CV_16BF ||
depth == CV_8F_E4M3FN || depth == CV_8F_E4M3FNUZ;
bool blockSum = partialSumIsInt || depth == CV_16F || depth == CV_16BF;
size_t esz = 0, nz0 = 0;
if( blockSum )

19
modules/core/src/nd_copy.cpp
View File

@@ -73,14 +73,17 @@ enum Kind
KIND_MERGE // 2..4 planes are written into one interleaved stream (e.g. NCHW -> NHWC)
};
constexpr int TRANSPOSE_TILE = 64; // in elements
constexpr int SPLIT_CHUNK = 1024; // in elements
constexpr int ROW_CHUNK_BYTES = 64*1024;
constexpr double GATHER_BLOCK_BYTES = 8*1024;
// costs are bytes weighted by the kernel; the copies are memory-bound,
// so small stripes only add overhead
constexpr double PARALLEL_MIN_COST = 1024*1024;
constexpr double STRIPE_COST = 1024*1024;
enum
{
TRANSPOSE_TILE = 64, // in elements
SPLIT_CHUNK = 1024, // in elements
ROW_CHUNK_BYTES = 64*1024,
GATHER_BLOCK_BYTES = 8*1024,
// costs are bytes weighted by the kernel; the copies are memory-bound,
// so small stripes only add overhead
PARALLEL_MIN_COST = 1024*1024,
STRIPE_COST = 1024*1024
};
struct Plan
{

62
modules/core/src/parallel.cpp
View File

@@ -454,41 +454,7 @@ namespace {
static inline void pplParallelFor(int first, int last, const Body& body)
{
#if defined(_M_ARM64) || defined(_M_ARM64EC)
// static_partitioner just spawns one task per worker; the atomic loop rebalances.
if (first >= last)
return;
const int total = last - first;
// cv::getNumThreads() returns 0 for WINRT, which also reaches this function.
const unsigned vps = Concurrency::CurrentScheduler::Get()->GetNumberOfVirtualProcessors();
const unsigned workers = std::min(std::max(1u, vps), (unsigned)total);
// total<=workers means one unit each, so skip the contended claim loop below.
if (total <= (int)workers)
{
Concurrency::parallel_for(first, last, body, Concurrency::static_partitioner());
return;
}
const int remaining_multiplier = std::min(total,
(int)std::max(std::min(100u, workers * 4), workers * 2)); // experimental value
std::atomic<int> current_task(0);
const auto worker = [&](int)
{
for (;;)
{
int chunk_size = std::max(1, (total - current_task) / remaining_multiplier);
int id = current_task.fetch_add(chunk_size, std::memory_order_relaxed);
if (id >= total)
break;
int end_id = std::min(total, id + chunk_size);
for (int i = first + id; i < first + end_id; ++i)
body(i);
}
};
Concurrency::parallel_for(0, (int)workers, worker, Concurrency::static_partitioner());
CV_Assert(current_task >= total);
Concurrency::parallel_for(first, last, body, Concurrency::static_partitioner());
#else
Concurrency::parallel_for(first, last, body);
#endif
@@ -499,17 +465,9 @@ namespace {
#if defined HAVE_TBB
#if TBB_INTERFACE_VERSION >= 8000
static tbb::task_arena& getTbbArena() {
// Keep the TBB object alive until process termination to avoid static
// destruction during process teardown on Windows.
CV_SINGLETON_LAZY_INIT_REF(tbb::task_arena, new tbb::task_arena(tbb::task_arena::automatic));
}
static tbb::task_arena tbbArena(tbb::task_arena::automatic);
#else
static tbb::task_scheduler_init& getTbbScheduler() {
// Keep the TBB object alive until process termination to avoid static
// destruction during process teardown on Windows.
CV_SINGLETON_LAZY_INIT_REF(tbb::task_scheduler_init, new tbb::task_scheduler_init(tbb::task_scheduler_init::deferred));
}
static tbb::task_scheduler_init tbbScheduler(tbb::task_scheduler_init::deferred);
#endif
#elif defined HAVE_HPX
// nothing for HPX
@@ -624,7 +582,7 @@ static void parallel_for_impl(const cv::Range& range, const cv::ParallelLoopBody
#if defined HAVE_TBB
#if TBB_INTERFACE_VERSION >= 8000
getTbbArena().execute(pbody);
tbbArena.execute(pbody);
#else
pbody();
#endif
@@ -693,13 +651,13 @@ int getNumThreads(void)
#if defined HAVE_TBB
#if TBB_INTERFACE_VERSION >= 9100
return getTbbArena().max_concurrency();
return tbbArena.max_concurrency();
#elif TBB_INTERFACE_VERSION >= 8000
return numThreads > 0
? numThreads
: tbb::task_scheduler_init::default_num_threads();
#else
return getTbbScheduler().is_active()
return tbbScheduler.is_active()
? numThreads
: tbb::task_scheduler_init::default_num_threads();
#endif
@@ -777,11 +735,11 @@ void setNumThreads( int threads_ )
#ifdef HAVE_TBB
#if TBB_INTERFACE_VERSION >= 8000
if(getTbbArena().is_active()) getTbbArena().terminate();
if(threads > 0) getTbbArena().initialize(threads);
if(tbbArena.is_active()) tbbArena.terminate();
if(threads > 0) tbbArena.initialize(threads);
#else
if(getTbbScheduler().is_active()) getTbbScheduler().terminate();
if(threads > 0) getTbbScheduler().initialize(threads);
if(tbbScheduler.is_active()) tbbScheduler.terminate();
if(threads > 0) tbbScheduler.initialize(threads);
#endif
#elif defined HAVE_HPX

8
modules/core/src/persistence.cpp
View File

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

2
modules/core/src/persistence_impl.hpp
View File

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

13
modules/core/src/reduce_nd.cpp
View File

@@ -17,11 +17,14 @@ namespace cv {
namespace {
constexpr int RND_MAX_DIMS = CV_MAX_DIM + 1; // + the channel axis
constexpr int RND_BLOCK = 1024; // elements converted per call when the source needs conversion
constexpr int RND_VCHUNK = 1024; // innermost-axis chunk in the vertical kernel
constexpr double RND_SERIAL_WORK = 64*1024; // below this (source elements) everything runs in one thread
constexpr double RND_STRIPE_WORK = 64*1024;
enum
{
RND_MAX_DIMS = CV_MAX_DIM + 1, // + the channel axis
RND_BLOCK = 1024, // elements converted per call when the source needs conversion
RND_VCHUNK = 1024, // innermost-axis chunk in the vertical kernel
RND_SERIAL_WORK = 64*1024, // below this (source elements) everything runs in one thread
RND_STRIPE_WORK = 64*1024
};
//////////////////////////////// vector traits ////////////////////////////////

3
modules/core/src/sum.dispatch.cpp
View File

@@ -152,8 +152,7 @@ static Scalar sum_(const Mat& src)
int* buf = (int*)&_res[0];
size_t esz = 0;
bool partialSumIsInt = depth < CV_32S;
bool blockSum = partialSumIsInt || depth == CV_16F || depth == CV_16BF ||
depth == CV_8F_E4M3FN || depth == CV_8F_E4M3FNUZ;
bool blockSum = partialSumIsInt || depth == CV_16F || depth == CV_16BF;
if( blockSum )
{

18
modules/core/src/sum.simd.hpp
View File

@@ -124,10 +124,6 @@ template<> struct Sum_SIMD<T, ST> \
ST maskbuf[VTraits<VecT>::max_nlanes*9]; \
};
// The channel limit is 4 only once the vector holds at least four lanes: init_maskbuf builds
// the cn == 4 masks under that condition, so a narrower vector would reduce against all-zero
// masks and silently drop the vectorized elements. Every 32-bit register type has four lanes
// or more; v_float64 on an SSE2 baseline has two.
#undef DEFINE_SUM_SIMD_16
#define DEFINE_SUM_SIMD_16(T, ST, iST, VecT, load_op) \
template<> struct Sum_SIMD<T, ST> \
@@ -138,10 +134,10 @@ template<> struct Sum_SIMD<T, ST> \
} \
int operator ()(const T* src, const uchar* mask, ST* dst, int len, int cn) const \
{ \
int x = 0, simd_width = VTraits<VecT>::vlanes(); \
if (mask || (cn < 1 || cn > 3+(simd_width>=4))) \
if (mask || (cn < 1 || cn > 4)) \
return 0; \
len *= cn; \
int x = 0, simd_width = VTraits<VecT>::vlanes(); \
VecT m0 = vx_load(maskbuf), m1, m2, m3, m4, m5, m6, m7, m8; \
if (cn == 1) { \
m1 = m2 = m3 = m4 = m5 = m6 = m7 = m8 = m0; \
@@ -233,8 +229,6 @@ template<> struct Sum_SIMD<T, ST> \
DEFINE_SUM_SIMD_32(int, double, int64, v_float64)
DEFINE_SUM_SIMD_32(float, double, int64, v_float64)
DEFINE_SUM_SIMD_16(double, double, int64, v_float64, vx_load)
#endif
#endif
@@ -392,12 +386,6 @@ static int sum16f( const hfloat* src, const uchar* mask, float* dst, int len, in
static int sum16bf( const bfloat* src, const uchar* mask, float* dst, int len, int cn )
{ CV_INSTRUMENT_REGION(); return sum_<bfloat, float, float>(src, mask, dst, len, cn); }
static int sum8f( const fp8_t* src, const uchar* mask, float* dst, int len, int cn )
{ CV_INSTRUMENT_REGION(); return sum_<fp8_t, float, float>(src, mask, dst, len, cn); }
static int sum8fnuz( const fp8a_t* src, const uchar* mask, float* dst, int len, int cn )
{ CV_INSTRUMENT_REGION(); return sum_<fp8a_t, float, float>(src, mask, dst, len, cn); }
SumFunc getSumFunc(int depth)
{
static SumFunc sumTab[CV_DEPTH_MAX] =
@@ -415,8 +403,6 @@ SumFunc getSumFunc(int depth)
sumWrap<uint64, double, sum64u>,
sumWrap<int64, double, sum64s>,
sumWrap<unsigned, double, sum32u>,
sumWrap<fp8_t, float, sum8f>,
sumWrap<fp8a_t, float, sum8fnuz>,
0
};

186
modules/core/test/test_arithm.cpp
View File

@@ -3183,30 +3183,6 @@ TEST(Core_CvAbs, lowest_signed_values)
EXPECT_EQ(1ull, cv_abs((int64)-1));
}
// cv_absdiff() on int64 must return the exact distance, even when it does not fit in int64.
TEST(Core_CvAbsDiff, int64_full_range)
{
EXPECT_EQ((uint64)ULLONG_MAX, cv_absdiff((int64)LLONG_MAX, (int64)LLONG_MIN));
EXPECT_EQ((uint64)ULLONG_MAX, cv_absdiff((int64)LLONG_MIN, (int64)LLONG_MAX));
EXPECT_EQ(3ull, cv_absdiff((int64)-1, (int64)2));
}
// norm(a, b) on CV_64S must not overflow when |a - b| is larger than INT64_MAX.
TEST(Core_Norm, diff_64S_beyond_int64_range)
{
for (int len : {1, 7, 64})
{
Mat a(1, len, CV_64S, Scalar(5e18));
Mat b(1, len, CV_64S, Scalar(-5e18));
Mat mask(1, len, CV_8U, Scalar(255));
EXPECT_EQ(1e19, cv::norm(a, b, NORM_INF)) << "length " << len;
EXPECT_EQ(1e19, cv::norm(a, b, NORM_INF, mask)) << "length " << len << ", masked";
EXPECT_EQ(len * 1e19, cv::norm(a, b, NORM_L1)) << "length " << len;
EXPECT_EQ(len * 1e19, cv::norm(a, b, NORM_L1, mask)) << "length " << len << ", masked";
}
}
// An L1 norm is a sum of magnitudes: never negative, and never dependent on the array length.
// Checked over both signed integer depths, masked and unmasked, at lengths straddling the vector
// width so the body, the tail and the two together are covered.
@@ -3285,52 +3261,6 @@ TEST(Core_Norm, NORM_L2SQR_16SC4_large)
EXPECT_EQ(expected, cv::norm(src, NORM_L2SQR));
}
// sum and mean on the fp8 depths, compared with the same values in CV_32F
class Core_Sum_FP8 : public testing::TestWithParam< tuple<int, int, int> >
{
public:
int len;
Mat a32f, a;
void SetUp()
{
const int depth = get<0>(GetParam());
const int cn = get<1>(GetParam());
len = get<2>(GetParam());
const float vals[] = { 1.f, -2.f, 3.5f, 0.5f, 7.f, -0.25f }; // exact in both fp8 formats
a32f.create(1, len, CV_MAKETYPE(CV_32F, cn));
for (int i = 0; i < len * cn; i++)
a32f.ptr<float>()[i] = vals[i % 6];
a32f.convertTo(a, depth);
}
};
TEST_P(Core_Sum_FP8, sum)
{
EXPECT_EQ(cv::sum(a32f), cv::sum(a));
}
TEST_P(Core_Sum_FP8, mean)
{
EXPECT_EQ(cv::mean(a32f), cv::mean(a));
}
TEST_P(Core_Sum_FP8, mean_mask)
{
// every third element is left out, so the mean divides by the selected count
Mat mask(1, len, CV_8U);
for (int i = 0; i < len; i++)
mask.at<uchar>(i) = i % 3 != 0;
EXPECT_EQ(cv::mean(a32f, mask), cv::mean(a, mask));
}
// (1 << 15) + 1 needs a second block, so the float total gets moved to the result and reset
INSTANTIATE_TEST_CASE_P(/**/, Core_Sum_FP8, testing::Combine(
testing::Values(CV_8F_E4M3FN, CV_8F_E4M3FNUZ),
testing::Range(1, 5),
testing::Values(37, (1 << 15) + 1)));
// cv::sum() returns doubles, so a CV_32U or CV_32S array whose total exceeds the 32-bit range must
// still come back exact. The scalar step of the accumulation adds four elements to each other in
// the source type before reaching the double accumulator, so large values used to wrap - silently
@@ -4092,35 +4022,6 @@ INSTANTIATE_TEST_CASE_P(Core_CartPolar, Core_CartToPolar_inplace,
)
);
typedef testing::TestWithParam<int> Core_CartToPolar_inplace_mag;
// only the magnitude overwrites an input, so the last elements must not be processed twice.
// CV_64F is not tested: it works on separate float buffers, so it can't hit this.
TEST_P(Core_CartToPolar_inplace_mag, accuracy)
{
const int len = GetParam();
cv::Mat x(1, len, CV_32F), y(1, len, CV_32F);
cv::randu(x, cv::Scalar::all(-1000), cv::Scalar::all(1000));
cv::randu(y, cv::Scalar::all(-1000), cv::Scalar::all(1000));
cv::Mat expectedMag, expectedAngle;
cv::cartToPolar(x, y, expectedMag, expectedAngle);
// magnitude into x
cv::Mat x1 = x.clone(), angle1;
cv::cartToPolar(x1, y, x1, angle1);
EXPECT_MAT_NEAR(expectedMag, x1, 1e-2);
EXPECT_MAT_NEAR(expectedAngle, angle1, 1e-2);
// magnitude into y
cv::Mat y1 = y.clone(), angle2;
cv::cartToPolar(x, y1, y1, angle2);
EXPECT_MAT_NEAR(expectedMag, y1, 1e-2);
EXPECT_MAT_NEAR(expectedAngle, angle2, 1e-2);
}
INSTANTIATE_TEST_CASE_P(Core_CartPolar, Core_CartToPolar_inplace_mag, testing::Range(1, 40));
PARAM_TEST_CASE(Core_PolarToCart_inplace, int, bool, bool)
{
int depth;
@@ -4723,6 +4624,14 @@ TEST_P(Core_LUT, accuracy)
cv::Mat gt = reference(input, table);
ASSERT_FALSE(gt.empty());
// Force convert to 8U as CV_Bool is not supported in cv::norm for now
// TODO: Remove conversion after cv::norm fix
if (value_type == CV_Bool)
{
output.convertTo(output, CV_8U);
gt.convertTo(gt, CV_8U);
}
ASSERT_EQ(0, cv::norm(output, gt, cv::NORM_INF));
}
@@ -4747,6 +4656,14 @@ TEST_P(Core_LUT, accuracy_multi)
cv::Mat gt = reference<3>(input, table);
ASSERT_FALSE(gt.empty());
// Force convert to 8U as CV_Bool is not supported in cv::norm for now
// TODO: Remove conversion after cv::norm fix
if (value_type == CV_Bool)
{
output.convertTo(output, CV_8U);
gt.convertTo(gt, CV_8U);
}
ASSERT_EQ(0, cv::norm(output, gt, cv::NORM_INF));
}
@@ -4771,6 +4688,14 @@ TEST_P(Core_LUT, accuracy_multi2)
cv::Mat gt = reference<3, true>(input, table);
ASSERT_FALSE(gt.empty());
// Force convert to 8U as CV_Bool is not supported in cv::norm for now
// TODO: Remove conversion after cv::norm fix
if (value_type == CV_Bool)
{
output.convertTo(output, CV_8U);
gt.convertTo(gt, CV_8U);
}
ASSERT_EQ(0, cv::norm(output, gt, cv::NORM_INF));
}
@@ -5065,7 +4990,7 @@ TEST(Core_Arithm, inRange_unsupported_depth_30153)
TEST(Core_Arithm, sum_unsupported_depth)
{
for (int depth : {CV_Bool})
for (int depth : {CV_Bool, CV_8F_E4M3FN, CV_8F_E4M3FNUZ})
{
for (int cn : {1, 3})
{
@@ -5204,65 +5129,4 @@ TEST_P(Core_PatchNaNs, replaces_only_nans)
INSTANTIATE_TEST_CASE_P(/**/, Core_PatchNaNs, testing::Values(CV_16F, CV_16BF, CV_32F, CV_64F));
TEST(Core_Multiply, scale_simd_scalar_consistency)
{
const float big = std::ldexp(1.0f, 80); // exact power of 2 (< FLT_MAX = 2^128)
const float small = std::ldexp(1.0f, -80);
// Two contiguous inputs: (a * s) * b = 2^80 vs (a * b) * s = +inf.
{
Mat a(1, 65, CV_32F, Scalar(big));
Mat b(1, 65, CV_32F, Scalar(big));
cv::multiply(a, b, a, small); // in-place -> tail element 64 uses EwMul::scl
EXPECT_EQ(a.at<float>(0), big);
EXPECT_EQ(a.at<float>(64), big);
}
// src1 scalar/column broadcast (s1x == 0): SIMD folds scale into src1, computing (b * s) * a = 2^80;
// without canonicalizing src1 onto src0, the scalar tail would compute (a * s) * b = +inf.
{
Mat a(1, 65, CV_32F, Scalar(big));
Mat b(1, 1, CV_32F, Scalar(small));
cv::multiply(a, b, a, big);
EXPECT_EQ(a.at<float>(0), big);
EXPECT_EQ(a.at<float>(64), big);
}
{
Mat a(1, 65, CV_32F, Scalar(big));
cv::multiply(a, Scalar(small), a, big);
EXPECT_EQ(a.at<float>(0), big);
EXPECT_EQ(a.at<float>(64), big);
}
// src1 short-row broadcast (s1y == 0, width <= 4): short-row SIMD processes full row blocks and
// leaves row 64 for the scalar loop.
{
Mat a(65, 3, CV_32F, Scalar(big));
Mat b(1, 3, CV_32F, Scalar(small));
cv::multiply(a, b, a, big);
EXPECT_EQ(a.at<float>(0, 0), big);
EXPECT_EQ(a.at<float>(64, 0), big);
}
// src1 wide-row broadcast (s1y == 0, width > 4):
{
Mat a(2, 65, CV_32F, Scalar(big));
Mat b(1, 65, CV_32F, Scalar(small));
cv::multiply(a, b, a, big);
EXPECT_EQ(a.at<float>(0, 0), big);
EXPECT_EQ(a.at<float>(1, 64), big);
}
// CV_64F broadcast (covers v_float64 SIMD and scalarBinaryKernel fallback):
{
const double big64 = std::ldexp(1.0, 600);
const double small64 = std::ldexp(1.0, -600);
Mat a(1, 65, CV_64F, Scalar(big64));
Mat b(1, 1, CV_64F, Scalar(small64));
cv::multiply(a, b, a, big64);
EXPECT_EQ(a.at<double>(0), big64);
EXPECT_EQ(a.at<double>(64), big64);
}
}
}} // namespace

32
modules/core/test/test_intrin_utils.hpp
View File

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

24
modules/core/test/test_io.cpp
View File

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

28
modules/core/test/test_misc.cpp
View File

@@ -441,34 +441,6 @@ TEST(Core_Parallel, propagate_fp_denormals_ignore_hint)
}
}
class CountingParallelLoopBody : public cv::ParallelLoopBody
{
public:
explicit CountingParallelLoopBody(std::vector<int>& hits) : hits_(hits) {}
void operator()(const cv::Range& r) const CV_OVERRIDE
{
for (int i = r.start; i < r.end; i++)
CV_XADD(&hits_[i], 1);
}
std::vector<int>& hits_;
};
typedef testing::TestWithParam<int> Core_Parallel_IndexCoverage;
TEST_P(Core_Parallel_IndexCoverage, every_index_exactly_once)
{
const int total = GetParam();
std::vector<int> hits(total, 0);
parallel_for_(cv::Range(0, total), CountingParallelLoopBody(hits));
for (int i = 0; i < total; i++)
ASSERT_EQ(1, hits[i]) << "index " << i << " of " << total;
}
// Brackets pplParallelFor's total<=workers fast path, and goes well past it to also
// exercise the shrinking-chunk claim loop.
INSTANTIATE_TEST_CASE_P(/**/, Core_Parallel_IndexCoverage,
testing::Values(1, 2, 4, 8, 16, 17, 32, 64, 100, 300, 100000));
TEST(Core_Version, consistency)
{
// this test verifies that OpenCV version loaded in runtime

16
modules/dnn/perf/perf_net.cpp
View File

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

5
modules/dnn/src/cuda/activations.cu
View File

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

100
modules/dnn/src/cuda/argmax.cu Normal file
View File

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

53
modules/dnn/src/cuda/cast.cu
View File

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

79
modules/dnn/src/cuda/col2im.cu Normal file
View File

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

108
modules/dnn/src/cuda/cum_scan.cu Normal file
View File

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

41
modules/dnn/src/cuda/eltwise_ops.cu
View File

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

59
modules/dnn/src/cuda/functors.hpp
View File

@@ -9,6 +9,8 @@
#include "math.hpp"
#include <type_traits>
#include "../cuda4dnn/csl/nvcc_defs.hpp"
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
@@ -851,6 +853,63 @@ struct FModFunctor {
}
};
// ONNX PRelu with a broadcast slope: y is the slope
template <class T>
struct PReluFunctor {
struct Params {
CUDA4DNN_HOST_DEVICE Params() {}
};
CUDA4DNN_DEVICE PReluFunctor() { }
CUDA4DNN_DEVICE PReluFunctor(const Params& params) { }
CUDA4DNN_DEVICE T operator()(T x, T y) {
return apply(x, y, std::is_unsigned<T>{});
}
// an unsigned value is never negative, so PRelu passes it through; an overload
// instead of if constexpr (C++17), which .cu files built as C++14 can't use
CUDA4DNN_DEVICE static T apply(T x, T, std::true_type) { return x; }
CUDA4DNN_DEVICE static T apply(T x, T y, std::false_type) { return x >= T(0) ? x : T(x * y); }
};
// bitwise functors are instantiated for integer types only
template <class T>
struct BitAndFunctor {
struct Params {
CUDA4DNN_HOST_DEVICE Params() {}
};
CUDA4DNN_DEVICE BitAndFunctor() { }
CUDA4DNN_DEVICE BitAndFunctor(const Params& params) { }
CUDA4DNN_DEVICE T operator()(T x, T y) { return x & y; }
};
template <class T>
struct BitOrFunctor {
struct Params {
CUDA4DNN_HOST_DEVICE Params() {}
};
CUDA4DNN_DEVICE BitOrFunctor() { }
CUDA4DNN_DEVICE BitOrFunctor(const Params& params) { }
CUDA4DNN_DEVICE T operator()(T x, T y) { return x | y; }
};
template <class T>
struct BitXorFunctor {
struct Params {
CUDA4DNN_HOST_DEVICE Params() {}
};
CUDA4DNN_DEVICE BitXorFunctor() { }
CUDA4DNN_DEVICE BitXorFunctor(const Params& params) { }
CUDA4DNN_DEVICE T operator()(T x, T y) { return x ^ y; }
};
template <class T>
struct PowFunctor {
struct Params {

90
modules/dnn/src/cuda/gather_elements.cu Normal file
View File

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

118
modules/dnn/src/cuda/gather_nd.cu Normal file
View File

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

162
modules/dnn/src/cuda/grid_sample.cu Normal file
View File

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

119
modules/dnn/src/cuda/mvn.cu
View File

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

48
modules/dnn/src/cuda/normalize.cu
View File

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

78
modules/dnn/src/cuda/one_hot.cu Normal file
View File

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

109
modules/dnn/src/cuda/padding.cu
View File

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

177
modules/dnn/src/cuda/quantize_dequantize.cu Normal file
View File

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

72
modules/dnn/src/cuda/range.cu Normal file
View File

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

72
modules/dnn/src/cuda/resize.cu
View File

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

97
modules/dnn/src/cuda/rms_norm.cu Normal file
View File

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

161
modules/dnn/src/cuda/roi_align.cu Normal file
View File

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

107
modules/dnn/src/cuda/scatter_elements.cu Normal file
View File

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

117
modules/dnn/src/cuda/scatter_nd.cu Normal file
View File

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

287
modules/dnn/src/cuda/topk.cu Normal file
View File

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

75
modules/dnn/src/cuda/trilu.cu Normal file
View File

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

1
modules/dnn/src/cuda4dnn/csl/cublas.hpp
View File

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

5
modules/dnn/src/cuda4dnn/csl/cudnn/graph.hpp
View File

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

23
modules/dnn/src/cuda4dnn/kernels/argmax.hpp Normal file
View File

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

29
modules/dnn/src/cuda4dnn/kernels/col2im.hpp Normal file
View File

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

35
modules/dnn/src/cuda4dnn/kernels/cum_scan.hpp Normal file
View File

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

13
modules/dnn/src/cuda4dnn/kernels/eltwise_ops.hpp
View File

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

28
modules/dnn/src/cuda4dnn/kernels/gather_elements.hpp Normal file
View File

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

30
modules/dnn/src/cuda4dnn/kernels/gather_nd.hpp Normal file
View File

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

37
modules/dnn/src/cuda4dnn/kernels/grid_sample.hpp Normal file
View File

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

8
modules/dnn/src/cuda4dnn/kernels/mvn.hpp
View File

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

8
modules/dnn/src/cuda4dnn/kernels/normalize.hpp
View File

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

22
modules/dnn/src/cuda4dnn/kernels/one_hot.hpp Normal file
View File

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

10
modules/dnn/src/cuda4dnn/kernels/padding.hpp
View File

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

38
modules/dnn/src/cuda4dnn/kernels/quantize_dequantize.hpp Normal file
View File

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

20
modules/dnn/src/cuda4dnn/kernels/range.hpp Normal file
View File

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

9
modules/dnn/src/cuda4dnn/kernels/resize.hpp
View File

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

34
modules/dnn/src/cuda4dnn/kernels/rms_norm.hpp Normal file
View File

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

44
modules/dnn/src/cuda4dnn/kernels/roi_align.hpp Normal file
View File

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

27
modules/dnn/src/cuda4dnn/kernels/scatter_elements.hpp Normal file
View File

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

28
modules/dnn/src/cuda4dnn/kernels/scatter_nd.hpp Normal file
View File

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

44
modules/dnn/src/cuda4dnn/kernels/topk.hpp Normal file
View File

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

22
modules/dnn/src/cuda4dnn/kernels/trilu.hpp Normal file
View File

@@ -0,0 +1,22 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_TRILU_HPP
#define OPENCV_DNN_SRC_CUDA4DNN_KERNELS_TRILU_HPP
#include "../csl/stream.hpp"
#include "../csl/span.hpp"
namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
// ONNX Trilu over the trailing [height, width] of each of `loops` matrices
template <class T>
void trilu(const csl::Stream& stream,
csl::Span<T> output, csl::View<T> input,
int loops, int height, int width, int k, bool upper);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_TRILU_HPP */

31
modules/dnn/src/cuda4dnn/primitives/activation.hpp
View File

@@ -53,6 +53,37 @@ namespace cv { namespace dnn { namespace cuda4dnn {
const T slope;
};
// ONNX Clip: min/max may arrive as extra inputs, so only input 0 is data
template <class T>
class ClipOp final : public CUDABackendNode {
public:
ClipOp(csl::Stream stream_, double min_, double max_)
: stream(std::move(stream_)), min{ toT(min_) }, max{ toT(max_) } { }
void forward(
const std::vector<UMat>& inputs,
const std::vector<UMat>& outputs,
csl::Workspace& workspace) override
{
CV_UNUSED(workspace);
CV_Assert(!inputs.empty() && outputs.size() == 1);
auto input = csl::viewOf<T>(inputs[0]);
auto output = csl::spanOf<T>(outputs[0]);
kernels::clipped_relu<T>(stream, output, input, min, max);
}
private:
// integers convert straight from double so int64 bounds keep their precision; half goes through float
static T toT(double v)
{
using U = typename std::conditional<std::is_integral<T>::value, double, float>::type;
return static_cast<T>(static_cast<U>(v));
}
csl::Stream stream;
const T min, max;
};
template <class T>
class ClippedReLUOp final : public BaseOp<ClippedReLUOp, T> {
public:

65
modules/dnn/src/cuda4dnn/primitives/arg.hpp Normal file
View File

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

25
modules/dnn/src/cuda4dnn/primitives/average_pooling.hpp
View File

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

150
modules/dnn/src/cuda4dnn/primitives/conv_transpose2.hpp Normal file
View File

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

17
modules/dnn/src/cuda4dnn/primitives/convolution.hpp
View File

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

67
modules/dnn/src/cuda4dnn/primitives/cum_scan.hpp Normal file
View File

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

10
modules/dnn/src/cuda4dnn/primitives/depth_space_ops.hpp
View File

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

51
modules/dnn/src/cuda4dnn/primitives/dropout_mask.hpp Normal file
View File

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

108
modules/dnn/src/cuda4dnn/primitives/einsum.hpp Normal file
View File

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

30
modules/dnn/src/cuda4dnn/primitives/eltwise.hpp
View File

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

71
modules/dnn/src/cuda4dnn/primitives/gather_elements.hpp Normal file
View File

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

104
modules/dnn/src/cuda4dnn/primitives/gather_nd.hpp Normal file
View File

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

61
modules/dnn/src/cuda4dnn/primitives/grid_sample.hpp Normal file
View File

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

26
modules/dnn/src/cuda4dnn/primitives/group_norm.hpp
View File

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

26
modules/dnn/src/cuda4dnn/primitives/instance_norm.hpp
View File

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

13
modules/dnn/src/cuda4dnn/primitives/layer_norm.hpp
View File

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

35
modules/dnn/src/cuda4dnn/primitives/lrn.hpp
View File

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

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