Commit Graph

60 Commits

Author SHA1 Message Date
jimmylaw21
a7fa1e6f4b
Merge pull request #24610 from jimmylaw21:dnn-onnx-add-group-norm-layer
dnn onnx: add group norm layer #24610

dnn onnx: add group norm layer

Todo:

- [x] speed up by multi-threading
- [x] add perf
- [x] add backend: OpenVINO
- [x] add backend: CUDA
- [x] add backend: OpenCL (no fp16)
- [ ] add backend: CANN

### 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
- [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

Co-authored-by: fengyuentau <yuantao.feng@opencv.org.cn>
2024-01-12 15:13:26 +03:00
Yuantao Feng
e7ccff9805
Merge pull request #24834 from fengyuentau:cuda_naryeltwise_broadcast
dnn (cuda): support broadcasting if a.rank() != b.rank() #24834

Inspired by https://github.com/opencv/opencv/pull/24786. This PR keeps the fusion of `NaryEltwise` and `Concat` while addressed the data missing problem via supporting broadcasting if a.rank() != b.rank().

Resolves https://github.com/opencv/opencv/issues/23977
Resolves https://github.com/opencv/opencv/issues/24606
Resolves https://github.com/opencv/opencv/issues/24635
Resolves https://github.com/opencv/opencv/issues/24721 

### 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
2024-01-11 10:04:46 +03:00
Yuantao Feng
c955564cb3
Merge pull request #24765 from fengyuentau:mod_operator
dnn onnx: add mod #24765

Resolves https://github.com/opencv/opencv/issues/23174

TODO:

- [x] enable some conformance tests
- [x] add backends
    - [x] CANN
    - [x] OpenVINO
    - [x] CUDA

### 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
2024-01-09 19:00:17 +03:00
Yuantao Feng
a2edf4d929
Merge pull request #24647 from fengyuentau:cuda_sub
dnn cuda: support Sub #24647

Related https://github.com/opencv/opencv/issues/24606#issuecomment-1837390257

### 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
2023-12-06 13:46:24 +03:00
Yuantao Feng
d05fb709f9
Merge pull request #24552 from fengyuentau:layernorm_backends
dnn: add openvino, opencl and cuda backends for layer normalization layer #24552

Merge after https://github.com/opencv/opencv/pull/24544.

Todo:

- [x] openvino
- [x] opencl
- [x] cuda

### 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
2023-11-21 15:33:01 +03:00
Yuantao Feng
ee0822dc4d
Merge pull request #24378 from fengyuentau:instance_norm
dnn onnx: add instance norm layer #24378

Resolves https://github.com/opencv/opencv/issues/24377
Relates https://github.com/opencv/opencv/pull/24092#discussion_r1349841644

| Perf | multi-thread | single-thread |
| - | - | - |
| x: [2, 64, 180, 240] | 3.95ms | 11.12ms |

Todo:

- [x] speed up by multi-threading
- [x] add perf
- [x] add backend: OpenVINO
- [x] add backend: CUDA
- [x] add backend: OpenCL (no fp16)
- [ ] add backend: CANN (will be done via https://github.com/opencv/opencv/pull/24462)


### 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

```
force_builders=Linux OpenCL,Win64 OpenCL,Custom
buildworker:Custom=linux-4
build_image:Custom=ubuntu:18.04
modules_filter:Custom=none
disable_ipp:Custom=ON
```
2023-11-07 12:59:10 +03:00
Dmitry Kurtaev
96f23e3da1
Merge pull request #24080 from dkurt:dnn_cuda_layers
Resolve uncovered CUDA dnn layer #24080

### Pull Request Readiness Checklist

* Gelu activation layer on CUDA
* Try to relax GEMM from ONNX

resolves https://github.com/opencv/opencv/issues/24064

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
2023-08-03 09:13:42 +03:00
wanli
46991bcd62 Solve the bug of same shape broadcast with CUDA 2023-05-15 13:55:38 +08:00
rogday
9cd5a0a1e6
Merge pull request #21884 from rogday:cuda_cleanup
Fix CUDA compilation issues and adjust thresholds.

* Fix CUDA compilation issues and adjust thresholds.

* add conformance tests to denylist
2022-04-19 16:40:25 +00:00
zihaomu
e36948cfbc add ONNX OP sign, shrink and reciprocal 2022-04-07 15:32:12 +08:00
luz paz
8e8e4bbabc dnn: fix various dnn related typos
Fixes source comments and documentation related to dnn code.
2022-03-23 18:12:12 -04:00
Smirnov Egor
71a22e45b0 add celu, hardsigmoid, selu, thresholdedrelu layers 2021-12-18 03:19:54 +03:00
Smirnov Egor
1bd382c1d0 Add acos, acosh, asin, asinh, atan, atanh, cos, cosh, erf, hardswish, sin, sinh, softplus, softsign, tan layers 2021-12-17 18:19:40 +03:00
Smirnov Egor
4995aecd62 add alpha parameter to ELU 2021-11-30 14:43:18 +03:00
Supernovae
b594ed99b8
Merge pull request #20933 from shubham-shahh:master
Improved overall readability of the code

* grid_nms.cu: minor fix-ups

* Update grid_stride_range.hpp

* Update tf_importer.cpp
2021-11-28 12:54:29 +00:00
Smirnov Egor
1feb3838b5 add Ceil, Floor, Log, Round, Sqrt, Not, Equal, Less, Greater 2021-10-15 16:02:46 +03:00
Smirnov Egor
9c84749e2c backport YOLOv4x-mish new_coords CUDA implementation 2021-10-08 14:14:49 +03:00
YashasSamaga
505dde09de support broadcasting in eltwise ops 2021-10-04 12:38:45 +05:30
rogday
38b9ec7a18
Merge pull request #20682 from rogday:min
* Add Min layer to CPU, OpenCL, Halide, Inference Engine, NGraph and CUDA

* fix indentation

* add min to fusion and halide tests; fix doc
2021-09-22 15:17:37 +03:00
Alexander Alekhin
c89084e6b7 Merge pull request #19223 from YashasSamaga:cuda4dnn-halfpix-linear-resize 2021-03-30 13:19:41 +00:00
SamFC10
6111935835 Added exp layer 2021-02-20 22:16:00 +05:30
Sergei Slashchinin
ea41f89b40
Merge pull request #19058 from sl-sergei:cuda_1d
Conv1D and Pool1D for CUDA backend

* CUDA-independent changes

* Add Conv1D and Pool1D for CUDA backend

* CUDA-independent changes

* Fix typo

* fix comment

* Update fix

* make changes more correct for pooling layer

* Minor fixes for review

* Split skip blocks
2021-01-21 22:16:56 +00:00
YashasSamaga
8c74d7e4fa add half pixel centers and align corners param 2020-12-27 15:05:39 +05:30
YashasSamaga
f0149cdae2 fix compile-time errors, disable unsupported tests 2020-08-09 14:43:20 +05:30
YashasSamaga
a3106d424b add MVNOp 2020-08-02 12:44:35 +05:30
YashasSamaga
ae293f27cf add DetectionOutputOp 2020-07-29 12:28:00 +05:30
Yashas Samaga B L
d0e6d2438c
Merge pull request #17363 from YashasSamaga:cuda4dnn-eltwise-fusion2
cuda4dnn(conv): fuse eltwise with convolutions

* fuse eltwise with convolutions

* manually rebase to avoid bad git merge
2020-07-09 16:02:21 +03:00
YashasSamaga
cbdaa93e54 reduce slice, concat to copy; enable more concat fusions 2020-07-05 20:52:35 +05:30
YashasSamaga
6573b9ace0 use fp32 mish for fp16 mish 2020-06-22 19:09:36 +05:30
Yashas Samaga B L
9ba5581d17
Merge pull request #17534 from YashasSamaga:cuda4dnn-remove-unused-funcs
cuda4dnn: reduce CUDA version requirements to at least CUDA 9.2

* remove half2 specializations

* do not remove atomicAdd for half in CUDA 10 and below

* remove fp16.hpp
2020-06-17 09:07:52 +00:00
YashasSamaga
87ab4ee567 improve mish performance and accuracy 2020-06-13 16:53:27 +05:30
YashasSamaga
3c35b563d7 add scale_x_y parameter to region 2020-05-10 16:53:28 +05:30
Yashas Samaga B L
d981d04c76
Merge pull request #17200 from YashasSamaga:cuda4dnn-general-opt1
cuda4dnn: optimizations for swish, mish, sigmoid, region, resize based ops, transpose, identity-conv fusion

* bunch of optimizations

* more accurate implementation for mish
2020-05-09 17:20:30 +00:00
Yashas Samaga B L
8808aaccff
Merge pull request #16658 from YashasSamaga:cuda4dnn-refactor-activations
cuda4dnn(activations, eltwise, scale_shift): refactor to reduce code duplication

* refactor activations

* refactor eltwise kernels

* move all functors to functors.hpp

* remove bias1 and scale1 kernels
2020-02-29 11:46:14 +03:00
Alexander Alekhin
2ced568d34 Merge pull request #16220 from YashasSamaga:cuda4dnn-roi-pooling-test_fix-optim 2020-01-29 20:57:15 +00:00
Julien Maille
a696348ec5 FIX: disable dnn cuda input_shortcut on _half for CC<5.3 2020-01-17 14:21:25 +01:00
Yashas Samaga B L
d85e67d3ec Merge pull request #16063 from YashasSamaga:cuda4dnn-shortcut-unequal
support eltwise sum with different number of input channels in CUDA backend

* add shortcut primitive

* add offsets in shortcut kernel

* skip tests involving more than two inputs

* remove redundant modulus operation

* support multiple inputs

* remove whole file indentation

* skip acc in0 trunc test if weighted

* use shortcut iff channels are unequal
2020-01-16 21:54:00 +03:00
Julien
ced3df73da Fix: rsqrt(float) was improperly put in the ifdef for half 2020-01-16 09:21:50 +01:00
YashasSamaga
fd369a5004 fix and optimize ROIPooling 2020-01-15 22:53:48 +05:30
Julien
4e2ef8c8f5 Merge pull request #16218 from JulienMaille:cuda-dnn-for-older-gpus
Enable cuda4dnn on hardware without support for __half

* Enable cuda4dnn on hardware without support for half (ie. compute capability < 5.3)

Update CMakeLists.txt

Lowered minimum CC to 3.0

* UPD: added ifdef on new copy kernel

* added fp16 support detection at runtime

* Clarified #if condition on atomicAdd definition

* More explicit CMake error message
2020-01-15 18:28:37 +03:00
Alexander Alekhin
1f2b2c5242 Merge pull request #16230 from YashasSamaga:cuda4dnn-fp-conversion 2020-01-05 11:59:33 +00:00
YashasSamaga
48eecafc89 simplify code to help MSVC 19.10 and lower 2019-12-30 23:02:17 +05:30
YashasSamaga
01f97f150c perfor fp conversions on GPU 2019-12-30 00:05:39 +05:30
Alexander Alekhin
9ec3d76b21 Merge pull request #16241 from bwignall:typo 2019-12-27 16:18:57 +00:00
Brian Wignall
659ffaddb4 Fix spelling typos 2019-12-26 06:45:03 -05:00
YashasSamaga
16bc505d26 improve reduction logic and add fast transpose kernel 2019-12-24 00:23:45 +05:30
Alexander Alekhin
b8e0898c7c Merge pull request #16082 from YashasSamaga:cuda4dnn-roi-pooling 2019-12-18 14:41:58 +00:00
Yashas Samaga B L
17c485eb03 Merge pull request #16092 from YashasSamaga:cuda4dnn-conv-act-fuse
cuda4dnn: fuse activations with convolutions

* fuse ReLU, ReLU6, TanH, Sigmoid with conv

* fix OpenCL errors

* improve ReLU, add power, swish and mish

* fix missing fusion entries

* fix handling of unsetAttached

* remove whole file indentation

* optimize power = 1.0, use IDENTITY instead of NONE

* handle edge case: change backend and then clear
2019-12-14 22:26:58 +03:00
Yashas Samaga B L
3fddd3bf93 Merge pull request #16069 from YashasSamaga:cuda4dnn-crop_and_resize
add CropAndResize layer for CUDA backend

* add CropAndResize layer

* process multiple channels per iteration
2019-12-09 22:26:58 +03:00
Alexander Alekhin
b505cf84de Merge pull request #16096 from YashasSamaga:cuda4dnn-region-optimize 2019-12-09 14:34:48 +00:00