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

Compare commits

base: Administrator:cuda/final_coverage
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: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

37 Commits
cuda/final ... 5.x

Author SHA1 Message Date
Abhishek Gola
370175b3bd Merge pull request #30624 from sshivensolanki/core-lut-test-bool 
core: remove CV_Bool conversion from LUT tests
 2026-11-10 16:23:21 +05:30
Abhishek Gola
e2eac555df Merge pull request #30724 from pbkx/port-seamless-cloning-mask-preservation-5x 
photo: preserve input masks in seamless cloning APIs
 2026-11-10 16:22:21 +05:30
Abhishek Gola
25029079fb Merge pull request #30732 from sshivensolanki/core-norm-absdiff-64s 
core: fix cv_absdiff overflow for int64 in norm of CV_64S
 2026-11-10 16:20:14 +05:30
Alexander Smorkalov
06ed129364 Merge pull request #30736 from abhishek-gola/fix/east-text-detection-30730 
samples: use fixed ONNX input size for EAST text detection
 2026-11-10 12:26:35 +04:00
Abhishek gola
688c406431 samples: use fixed ONNX input size for EAST text detection 2026-11-10 08:50:59 +05:30
Pratham Kumar
c261e6e4b4 imgproc: vectorize alphaComposite row kernel with universal intrinsics (#30471) 
* imgproc: vectorize alphaComposite row kernel with universal intrinsics

* use CV_ALWAYS_INLINE for alphaComposite SIMD helpers
 2026-11-10 08:49:56 +05:30
Shiven Solanki
f5f76c0786 core: fix cv_absdiff overflow for int64 in norm of CV_64S 
Fixes #30731
 2026-11-10 00:32:38 +08:00
bob
71babee383 Merge pull request #30661 from NUAAwqy:docs/deriv-kernel-size-range 
docs(imgproc): correct the documented Sobel kernel size range - #30661

## Problem

`cv::Sobel` and `cv::getDerivKernels` document the kernel size as

> `@param ksize size of the extended Sobel kernel; it must be 1, 3, 5, or 7.`
> `@param ksize Aperture size. It can be FILTER_SCHARR, 1, 3, 5, or 7.`

but `getDerivKernels` (modules/imgproc/src/deriv.cpp) actually enforces a different constraint:

```cpp
if( _ksize % 2 == 0 || _ksize > 31 )
    CV_Error( cv::Error::StsOutOfRange, "The kernel size must be odd and not larger than 31" );
```

so every odd size up to 31 is accepted, and `FILTER_SCHARR` (-1) selects the 3x3 Scharr kernels.

Reproduced with opencv-python-headless 5.0.0:

```
getDerivKernels(1, 0, 9)   -> OK, 9-tap kernel
getDerivKernels(1, 0, 15)  -> OK, 15-tap kernel
getDerivKernels(1, 0, 31)  -> OK, 31-tap kernel
getDerivKernels(1, 0, 33)  -> error: "The kernel size must be odd and not larger than 31"
getDerivKernels(1, 0, -1)  -> OK (Scharr 3x3)
cv2.Sobel(img, cv2.CV_32F, 1, 0, ksize=9/15)  -> OK
cv2.Sobel(img, cv2.CV_32F, 1, 0, ksize=101)   -> error (out of range)
```

## Change

Documentation-only: the two `@param ksize` entries now state the constraint the
implementation enforces (odd, not larger than 31) while keeping 1, 3, 5, 7 as the typical
choices. No code change.

Signed-off-by: NUAAwqy <162230229@nuaa.edu.cn>
 2026-11-09 15:15:14 +03:00
Andrei Fedorov
2e0c9a44e6 Merge pull request #30663 from intel-staging:andreyfe1/5.x/extend_perf_tests 
Porting of perf tests extension from 4.x to 5.x - #30663

Following PRs are being ported:
* #29908
* #29922 
* #29972

No more changes are introduced.
PR for opencv-extra: https://github.com/opencv/opencv_extra/pull/1458

### Pull Request Readiness Checklist

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

- [x] I agree to contribute to the project under Apache 2 License.
- [x] To the best of my knowledge, the proposed patch is not based on a code under GPL or another license that is incompatible with OpenCV
- [x] The PR is proposed to the proper branch
- [x] There is a reference to the original bug report and related work
- [ ] There is accuracy test, performance test and test data in opencv_extra repository, if applicable
      Patch to opencv_extra has the same branch name.
- [ ] The feature is well documented and sample code can be built with the project CMake
 2026-11-09 11:46:44 +03:00
uwezkhan
dfbf43df29 dnn: check weight, bias and past_state in CausalConvWithState 🤖🤖🤖 (#30628) 
* dnn: check weight, bias and past_state in CausalConvWithState

Check the ONNX shapes of weight [C, 1, k], bias [C] and past_state [B, C, k - 1] against the input in getMemoryShapes, and require them to be float in forward, since the kernel indexes all three by input channel and reads them as float.

* dnn: check CausalConvWithState input types before fp16 conversion, require rank-1 bias
 2026-11-09 12:23:45 +05:30
Shiven Solanki
566df44633 core: support fp8 depths in cv::sum and cv::mean (#30670) 
* core: support fp8 depths in cv::sum

cv::sum and cv::trace threw for CV_8F_E4M3FN and CV_8F_E4M3FNUZ. Add
sum8f and sum8fnuz to the function table, and collect their float totals
in blocks like CV_16F and CV_16BF.

* core: test cv::sum on fp8 past the block size

Add a (1 << 15) + 1 element case to Core_Sum_FP8, so the float total is
moved to the double result and reset between blocks.

* core: support fp8 depths in cv::mean

cv::mean uses the same adding functions as cv::sum, so adding fp8 to that table made it accept these depths. Its blockSum check still left them out, so the float totals were written straight into the double result and the mean came back wrong. Use the block path for CV_8F_E4M3FN and CV_8F_E4M3FNUZ, as cv::sum does, and check cv::mean in Core_Sum_FP8.

* core: test cv::mean with a mask on fp8

* core: split Core_Sum_FP8 into sum, mean and mean_mask tests

Move the array setup into a fixture, and make the length a test parameter, so each case runs on its own.
 2026-11-09 09:41:13 +05:30
pbkx
d1e278028c photo: preserve input masks in seamless cloning APIs 2026-11-08 14:04:26 -08:00
Alexander Smorkalov
ac33ba8f18 Merge pull request #30719 from pbkx:port-tonemap-reinhard-intensity-5x 
photo: preserve TonemapReinhard intensity during processing
 2026-11-08 20:26:08 +03:00
pbkx
5d1f514835 preserve TonemapReinhard intensity during processing 2026-11-08 02:31:03 -08:00
Shiven Solanki
bb8cbcdb0b Merge pull request #30715 from sshivensolanki:core-carttopolar-inplace-5x 
core: fix in-place cartToPolar when only the magnitude overwrites an input (5.x) - #30715

5.x port of #30646, as requested in its review. Fixes #30645 on 5.x.

### Problem

`cv::cartToPolar` gives wrong results on `CV_32F` when only the magnitude output is the same `Mat` as `x` or `y`. The SIMD tail handling repeats the last full block, and `cartToPolar32f_` only skipped that for in-place calls when the angle was aliased.

### Fix

Same as #30646: added `mag == X || mag == Y` to the check, so these calls handle the leftover elements with the scalar code.

### Tests

Same `Core_CartToPolar_inplace_mag.accuracy` test as #30646 (`CV_32F`, lengths 1 to 39, magnitude overwriting `x`, then `y`). On Apple M1 Pro it fails without the fix on 5.x and passes with it. The full `opencv_test_core` passes.

### Pull Request Readiness Checklist

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

- [x] I agree to contribute to the project under Apache 2 License.
- [x] To the best of my knowledge, the proposed patch is not based on a code under GPL or another license that is incompatible with OpenCV
- [x] The PR is proposed to the proper branch
- [x] There is a reference to the original bug report and related work
- [x] There is accuracy test, performance test and test data in opencv_extra repository, if applicable
      Patch to opencv_extra has the same branch name.
- [ ] The feature is well documented and sample code can be built with the project CMake
 2026-11-08 11:47:11 +03:00
pranayr710
1780d20b80 core: vectorize cv::sum for CV_64F (#30582) 
* core: vectorize cv::sum for CV_64F

Sum_SIMD has specializations for the 8-, 16- and 32-bit depths, but not for
CV_64F, so cv::sum and cv::mean on a double matrix fall through to the scalar
unrolled loop in sum_. CV_64FC1 is not in TYPICAL_MAT_TYPES, so no perf test
covered it.

The source type is already the accumulator type, so the specialization needs
neither expansion nor conversion: it loads three vectors at the native width and
adds them. The channel limit matches DEFINE_SUM_SIMD_32 because init_maskbuf can
only build the cn == 4 masks once the vector holds four lanes, which rules the
case out on an SSE2 baseline.

Extend sum_wide with the CV_64F rows so the change is measurable.

* core: reuse DEFINE_SUM_SIMD_16 for CV_64F

Drop the near-duplicate DEFINE_SUM_SIMD_64 and instantiate DEFINE_SUM_SIMD_16
with v_float64 instead, as suggested in review.

The one thing that kept the two apart was the channel guard. DEFINE_SUM_SIMD_16
rejects cn > 4, which is correct only while the vector holds at least four lanes:
init_maskbuf builds the cn == 4 masks under that condition, so with a narrower
vector the reduction runs against all-zero masks and drops the vectorized
elements without reporting anything. Every 32-bit register type has four lanes or
more, so the existing instantiations never saw it; v_float64 has two on an SSE2
baseline.

Widening the guard to the DEFINE_SUM_SIMD_32 form leaves the existing
instantiations unchanged, since 3+(simd_width>=4) is 4 for all of them.
 2026-11-08 08:31:23 +05:30
Alexander Smorkalov
4288cda8af Merge pull request #30711 from Kumataro:trial/donot_convert_from_enum_to_double 
core, videoio: fix C++20 compilation warnings for enum-to-double conversions
 2026-11-07 21:22:46 +03:00
bob
f1a605c037 Merge pull request #30662 from NUAAwqy:docs/medianblur-ksize-1 
docs(imgproc): medianBlur also accepts ksize = 1 - #30662

## Problem

`cv::medianBlur` documents the aperture as

> `@param ksize aperture linear size; it must be odd and greater than 1, for example: 3, 5, 7 ...`

but the implementation only requires an odd size: the aperture check is
`(ksize % 2 == 1) && (_src0.dims() <= 2)`, and `ksize = 1` is accepted and returns the image
unchanged (a 1x1 median is the pixel itself).

Reproduced with opencv-python-headless 5.0.0:

```
cv2.medianBlur(img, 1) -> OK, output is identical to the input
cv2.medianBlur(img, 2) -> error: "(ksize % 2 == 1) && (_src0.dims() <= 2)" in 'medianBlur'
```

## Change

Documentation-only: mention `ksize = 1` so the docs match the implemented behavior.
No code change.

Signed-off-by: NUAAwqy <162230229@nuaa.edu.cn>
 2026-11-07 16:39:27 +03:00
Lazizbek Ergashev
e2c04b9cac Merge pull request #30672 from lazerg:fix/issue-30666-v4l-imdecode-fail 
videoio(v4l): return false when MJPEG frame decoding fails - #30672

### Pull Request Readiness Checklist

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

- [x] I agree to contribute to the project under Apache 2 License.
- [x] To the best of my knowledge, the proposed patch is not based on a code under GPL or another license that is incompatible with OpenCV
- [x] The PR is proposed to the proper branch
- [x] There is a reference to the original bug report and related work
- [ ] There is accuracy test, performance test and test data in opencv_extra repository, if applicable
      Patch to opencv_extra has the same branch name.
- [ ] The feature is well documented and sample code can be built with the project CMake

When `imdecode` fails on a corrupt MJPEG/JPEG buffer it leaves `frame` untouched, so `retrieveFrame()` handed back the previous frame and returned true. Now the frame is released on decode failure and `retrieveFrame()` returns false for an empty frame, so `read()` returns false too.

I don't have a V4L2 device here, so this is only compile-checked on Linux (videoio with V4L and built-in libjpeg-turbo). The existing V4L tests use vivid and have MJPEG disabled, so I didn't add a test.

Fixes #30666
 2026-11-07 16:28:12 +03:00
Kumataro
e027ae3f3b core, videoio: fix C++20 compilation warnings for enum-to-double conversions 
- Replace anonymous enums with `constexpr` variables.
- Add `static_cast<double>()` to enum values in `EXPECT_EQ` assertions.
 2026-11-07 20:13:55 +09:00
Alexander Smorkalov
43f8b1f5da Merge pull request #30701 from Kumataro:bump_release_to_5.1_and_update_version_dropdown 
doc(sphinx): add OpenCV 5.1 to version list and update release to 5.1
 2026-11-06 18:17:58 +03:00
Kumataro
4d0b66ec1d doc(sphinx): add OpenCV 5.1 to version list and update release to 5.1 2026-11-06 17:55:52 +09:00
Alexander Smorkalov
07666386e2 Merge branch 'release_5.1.0' into 5.x 2026-11-05 16:14:14 +03:00
Alexander Smorkalov
e4f78fbcf1 release: OpenCV 5.1.0 2026-11-05 14:18:05 +03:00
Shiven Solanki
545147fdab Merge pull request #30685 from sshivensolanki:core-filestorage-unsigned 
core: keep large CV_32U values when reading with FileStorage - #30685

Fixes #30684

### Problem

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

### Fix

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

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

### Tests

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

### Pull Request Readiness Checklist

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

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

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

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

### Pull Request Readiness Checklist

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

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

## Summary

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

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

## Validation

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

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

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

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

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

### Pull Request Readiness Checklist

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

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

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

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

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

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

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

Outline the temp-bearing loops into `noinline` helpers
(`execFastStripWithTemps` and `execTileWithTemps`) and run the
`nbuffers == 0` tile body directly without `tstoreBuf`. This preserves
4KB L1 fragments when temporaries are present while saving ~8.7KB of
stack space on single-op `broadcastOp` calls.
 2026-11-03 16:17:42 +01:00
Shiven Solanki
482f8f784b core: remove CV_Bool conversion from LUT tests 
The Core_LUT tests converted CV_Bool results to CV_8U before comparing
them, because cv::norm did not support CV_Bool at the time. cv::norm
handles CV_Bool now, so remove the conversion and its TODO.
 2026-11-02 13:09:19 +08:00
Kaustubh Arora
da9b4e5744 dnn: enable the CastLike conformance cases that pass with the FP8/FP4/INT4 support 
34 test_castlike_* cases stayed on the parser denylist after the FP8, FP4
and INT4 cast support landed; they pass on current 5.x for OCV/CPU and
OCV/CPU_FP16. Remove them from the denylist and add them to the OpenVINO
filter as SKIP, like the other CastLike entries there.

The remaining 22 CastLike cases are INT2/UINT2 (not in the ONNX proto
OpenCV ships) and STRING, and stay denylisted.
 2026-11-01 12:38:35 +05:30
190 changed files with 1501 additions and 6751 deletions
Expand all files Collapse all files

1
docs_sphinx/_static/version.js
View File

@@ -12,6 +12,7 @@
* 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.0"
release = "5.1"
# -- Sphinx core ------------------------------------------------------------
extensions = ["myst_parser"]

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

@@ -35,6 +35,14 @@ 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,7 +50,9 @@ 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 :
@@ -265,7 +267,9 @@ 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,6 +321,7 @@ 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_OQ, _Tpvec, suffix) \
OPENCV_HAL_IMPL_AVX_CMP_FLT(v_ne, _CMP_NEQ_UQ, _Tpvec, suffix) \
OPENCV_HAL_IMPL_AVX_CMP_FLT(v_lt, _CMP_LT_OQ, _Tpvec, suffix) \
OPENCV_HAL_IMPL_AVX_CMP_FLT(v_gt, _CMP_GT_OQ, _Tpvec, suffix) \
OPENCV_HAL_IMPL_AVX_CMP_FLT(v_le, _CMP_LE_OQ, _Tpvec, suffix) \

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_OQ, _Tpvec, sufcmp, sufset, tval) \
OPENCV_HAL_IMPL_AVX512_CMP_FLT(v_ne, _CMP_NEQ_UQ, _Tpvec, sufcmp, sufset, tval) \
OPENCV_HAL_IMPL_AVX512_CMP_FLT(v_lt, _CMP_LT_OQ, _Tpvec, sufcmp, sufset, tval) \
OPENCV_HAL_IMPL_AVX512_CMP_FLT(v_gt, _CMP_GT_OQ, _Tpvec, sufcmp, sufset, tval) \
OPENCV_HAL_IMPL_AVX512_CMP_FLT(v_le, _CMP_LE_OQ, _Tpvec, sufcmp, sufset, tval) \

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_cne, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LASX_CMP_FLT(v_eq, xvfcmp_ceq, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LASX_CMP_FLT(v_ne, xvfcmp_cune, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LASX_CMP_FLT(v_lt, xvfcmp_clt, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LASX_CMP_FLT(v_le, xvfcmp_cle, _Tpvec, ssuffix)

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_cne, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LSX_CMP_FLT(v_eq, vfcmp_ceq, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LSX_CMP_FLT(v_ne, vfcmp_cune, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LSX_CMP_FLT(v_lt, vfcmp_clt, _Tpvec, ssuffix) \
OPENCV_HAL_IMPL_LSX_CMP_FLT(v_le, vfcmp_cle, _Tpvec, ssuffix) \

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

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

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 1
#define CV_VERSION_MINOR 2
#define CV_VERSION_REVISION 0
#define CV_VERSION_STATUS "-pre"
#define CV_VERSION_STATUS "-dev"
#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_16SC1, CV_16SC2, CV_16SC3, CV_16SC4, CV_32SC1, CV_32FC1, CV_64FC1)
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)
)
);

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(TYPICAL_MAT_TYPES),
testing::Values(CV_8UC1, CV_8UC4, CV_16UC1, CV_16SC1, CV_32FC1),
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 ),
testing::Values( CV_8UC1, CV_8SC1, CV_16SC1, CV_16UC1, CV_32SC1, CV_32FC1, CV_64FC1 ),
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
#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_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_32FC4))
testing::Values(CV_8UC1, CV_8UC2, CV_8UC3, CV_16UC1, CV_16UC3, CV_32SC1, CV_32SC2, CV_32FC3, 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))
testing::Values(CV_8UC1, CV_8UC2, CV_8UC3, CV_8UC4, CV_16UC3, CV_16UC4, CV_32FC3, CV_32FC4))
)
{
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)
testing::Values((int)NORM_INF, (int)NORM_L1, (int)NORM_L2, (int)NORM_L2SQR)
)
)
{
@@ -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_64FC1),
testing::Values((int)NORM_INF, (int)NORM_L1, (int)NORM_L2)
CV_32FC1, CV_32FC3, CV_64FC1),
testing::Values((int)NORM_INF, (int)NORM_L1, (int)NORM_L2, (int)NORM_L2SQR)
)
)
{
@@ -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_RELATIVE+NORM_INF), (int)(NORM_RELATIVE+NORM_L1), (int)(NORM_RELATIVE+NORM_L2))
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))
)
)
{
@@ -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_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))
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))
)
)
{

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_32FC1
#define TYPICAL_MAT_TYPES_SORT CV_8UC1, CV_16UC1, CV_16SC1, CV_32SC1, CV_32FC1, CV_64FC1
#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,11 +4,12 @@ namespace opencv_test
{
using namespace perf;
// 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.
// The wide 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)))
testing::Values(CV_32SC1, CV_32UC1, CV_32SC3,
CV_64FC1, CV_64FC3)))
{
Size sz = get<0>(GetParam());
int type = get<1>(GetParam());
@@ -23,7 +24,13 @@ PERF_TEST_P(Size_MatType, sum_wide,
SANITY_CHECK_NOTHING();
}
PERF_TEST_P(Size_MatType, sum, TYPICAL_MATS)
// 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 ) ))
{
Size sz = get<0>(GetParam());
int type = get<1>(GetParam());
@@ -38,7 +45,8 @@ PERF_TEST_P(Size_MatType, sum, TYPICAL_MATS)
SANITY_CHECK(s, 1e-6, ERROR_RELATIVE);
}
PERF_TEST_P(Size_MatType, mean, TYPICAL_MATS)
PERF_TEST_P(Size_MatType, mean, testing::Combine( testing::Values( TYPICAL_MAT_SIZES ),
testing::Values( TYPICAL_MAT_TYPES_STAT ) ))
{
Size sz = get<0>(GetParam());
int type = get<1>(GetParam());
@@ -53,7 +61,8 @@ PERF_TEST_P(Size_MatType, mean, TYPICAL_MATS)
SANITY_CHECK(s, 1e-5);
}
PERF_TEST_P(Size_MatType, mean_mask, TYPICAL_MATS)
PERF_TEST_P(Size_MatType, mean_mask, testing::Combine( testing::Values( TYPICAL_MAT_SIZES ),
testing::Values( TYPICAL_MAT_TYPES_STAT_MASK ) ))
{
Size sz = get<0>(GetParam());
int type = get<1>(GetParam());
@@ -69,7 +78,8 @@ PERF_TEST_P(Size_MatType, mean_mask, TYPICAL_MATS)
SANITY_CHECK(s, 5e-5);
}
PERF_TEST_P(Size_MatType, meanStdDev, TYPICAL_MATS)
PERF_TEST_P(Size_MatType, meanStdDev, testing::Combine( testing::Values( TYPICAL_MAT_SIZES ),
testing::Values( TYPICAL_MAT_TYPES_STAT_MASK ) ))
{
Size sz = get<0>(GetParam());
int matType = get<1>(GetParam());
@@ -86,7 +96,8 @@ PERF_TEST_P(Size_MatType, meanStdDev, TYPICAL_MATS)
SANITY_CHECK(dev, 1e-5, ERROR_RELATIVE);
}
PERF_TEST_P(Size_MatType, meanStdDev_mask, TYPICAL_MATS)
PERF_TEST_P(Size_MatType, meanStdDev_mask, testing::Combine( testing::Values( TYPICAL_MAT_SIZES ),
testing::Values( TYPICAL_MAT_TYPES_STAT_MASK ) ))
{
Size sz = get<0>(GetParam());
int matType = get<1>(GetParam());

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

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

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

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

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 || angle == X || angle == Y )
if( i == 0 || mag == X || mag == Y || angle == X || angle == Y )
break;
i = len - VECSZ*2;
}

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

@@ -52,7 +52,8 @@ 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;
bool blockSum = partialSumIsInt || depth == CV_16F || depth == CV_16BF ||
depth == CV_8F_E4M3FN || depth == CV_8F_E4M3FNUZ;
size_t esz = 0, nz0 = 0;
if( blockSum )

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

@@ -73,17 +73,14 @@ enum Kind
KIND_MERGE // 2..4 planes are written into one interleaved stream (e.g. NCHW -> NHWC)
};
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
};
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;
struct Plan
{

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

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

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

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

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

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

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

@@ -17,14 +17,11 @@ namespace cv {
namespace {
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
};
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;
//////////////////////////////// vector traits ////////////////////////////////

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

@@ -152,7 +152,8 @@ 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;
bool blockSum = partialSumIsInt || depth == CV_16F || depth == CV_16BF ||
depth == CV_8F_E4M3FN || depth == CV_8F_E4M3FNUZ;
if( blockSum )
{

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

@@ -124,6 +124,10 @@ 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> \
@@ -134,10 +138,10 @@ template<> struct Sum_SIMD<T, ST> \
} \
int operator ()(const T* src, const uchar* mask, ST* dst, int len, int cn) const \
{ \
if (mask || (cn < 1 || cn > 4)) \
int x = 0, simd_width = VTraits<VecT>::vlanes(); \
if (mask || (cn < 1 || cn > 3+(simd_width>=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; \
@@ -229,6 +233,8 @@ 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
@@ -386,6 +392,12 @@ 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] =
@@ -403,6 +415,8 @@ 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,6 +3183,30 @@ 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.
@@ -3261,6 +3285,52 @@ 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
@@ -4022,6 +4092,35 @@ 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;
@@ -4624,14 +4723,6 @@ 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));
}
@@ -4656,14 +4747,6 @@ 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));
}
@@ -4688,14 +4771,6 @@ 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));
}
@@ -4990,7 +5065,7 @@ TEST(Core_Arithm, inRange_unsupported_depth_30153)
TEST(Core_Arithm, sum_unsupported_depth)
{
for (int depth : {CV_Bool, CV_8F_E4M3FN, CV_8F_E4M3FNUZ})
for (int depth : {CV_Bool})
{
for (int cn : {1, 3})
{
@@ -5129,4 +5204,65 @@ TEST_P(Core_PatchNaNs, replaces_only_nans)
INSTANTIATE_TEST_CASE_P(/**/, Core_PatchNaNs, testing::Values(CV_16F, CV_16BF, CV_32F, CV_64F));
TEST(Core_Multiply, scale_simd_scalar_consistency)
{
const float big = std::ldexp(1.0f, 80); // exact power of 2 (< FLT_MAX = 2^128)
const float small = std::ldexp(1.0f, -80);
// Two contiguous inputs: (a * s) * b = 2^80 vs (a * b) * s = +inf.
{
Mat a(1, 65, CV_32F, Scalar(big));
Mat b(1, 65, CV_32F, Scalar(big));
cv::multiply(a, b, a, small); // in-place -> tail element 64 uses EwMul::scl
EXPECT_EQ(a.at<float>(0), big);
EXPECT_EQ(a.at<float>(64), big);
}
// src1 scalar/column broadcast (s1x == 0): SIMD folds scale into src1, computing (b * s) * a = 2^80;
// without canonicalizing src1 onto src0, the scalar tail would compute (a * s) * b = +inf.
{
Mat a(1, 65, CV_32F, Scalar(big));
Mat b(1, 1, CV_32F, Scalar(small));
cv::multiply(a, b, a, big);
EXPECT_EQ(a.at<float>(0), big);
EXPECT_EQ(a.at<float>(64), big);
}
{
Mat a(1, 65, CV_32F, Scalar(big));
cv::multiply(a, Scalar(small), a, big);
EXPECT_EQ(a.at<float>(0), big);
EXPECT_EQ(a.at<float>(64), big);
}
// src1 short-row broadcast (s1y == 0, width <= 4): short-row SIMD processes full row blocks and
// leaves row 64 for the scalar loop.
{
Mat a(65, 3, CV_32F, Scalar(big));
Mat b(1, 3, CV_32F, Scalar(small));
cv::multiply(a, b, a, big);
EXPECT_EQ(a.at<float>(0, 0), big);
EXPECT_EQ(a.at<float>(64, 0), big);
}
// src1 wide-row broadcast (s1y == 0, width > 4):
{
Mat a(2, 65, CV_32F, Scalar(big));
Mat b(1, 65, CV_32F, Scalar(small));
cv::multiply(a, b, a, big);
EXPECT_EQ(a.at<float>(0, 0), big);
EXPECT_EQ(a.at<float>(1, 64), big);
}
// CV_64F broadcast (covers v_float64 SIMD and scalarBinaryKernel fallback):
{
const double big64 = std::ldexp(1.0, 600);
const double small64 = std::ldexp(1.0, -600);
Mat a(1, 65, CV_64F, Scalar(big64));
Mat b(1, 1, CV_64F, Scalar(small64));
cv::multiply(a, b, a, big64);
EXPECT_EQ(a.at<double>(0), big64);
EXPECT_EQ(a.at<double>(64), big64);
}
}
}} // namespace

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

@@ -1,28 +0,0 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2026, BigVision LLC, all rights reserved.
// Third party copyrights are property of their respective owners.
#ifndef OPENCV_DNN_SRC_CUDA4DNN_KERNELS_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
View File

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

11
modules/dnn/src/cuda4dnn/primitives/normalize_bbox.hpp
View File

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

76
modules/dnn/src/cuda4dnn/primitives/one_hot.hpp
View File

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

32
modules/dnn/src/cuda4dnn/primitives/padding.hpp
View File

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

59
modules/dnn/src/cuda4dnn/primitives/pooling.hpp
View File

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

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