potok/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2026-07-21 19:33:03 +04:00
Code Packages Releases Wiki Activity

Merge pull request #25900 from fengyuentau:dnn/nary_elementwise_multi_thread

Browse Source 
dnn: merge #25630 to 5.x #25900

Sync changes from https://github.com/opencv/opencv/pull/25630 to 5.x.

### 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
This commit is contained in:
Yuantao Feng
2024-07-15 13:57:50 +08:00
committed by GitHub
parent 9aa5f3f1db
commit 420663498f
13 changed files with 519 additions and 332 deletions
Download Patch File Download Diff File Expand all files Collapse all files
modules/dnn/src/cuda/eltwise_ops.cu
+11
View File
@@ -350,6 +350,11 @@ void eltwise_fmod_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x,
eltwise_op<T, FModFunctor<T>>(stream, output, x, y);
}
template <class T>
void eltwise_pow_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x, TensorView<T> y) {
eltwise_op<T, PowFunctor<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);
@@ -360,6 +365,7 @@ void eltwise_fmod_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x,
template void eltwise_sum_2(const Stream& stream, TensorSpan<__half> output, TensorView<__half> x, TensorView<__half> y);
template void eltwise_max_2(const Stream& stream, TensorSpan<__half> output, TensorView<__half> x, TensorView<__half> y);
template void eltwise_min_2(const Stream& stream, TensorSpan<__half> output, TensorView<__half> x, TensorView<__half> y);
template void eltwise_pow_2(const Stream& stream, TensorSpan<__half> output, TensorView<__half> x, TensorView<__half> y);
#endif
template void eltwise_mod_2(const Stream& stream, TensorSpan<float> output, TensorView<float> x, TensorView<float> y);
template void eltwise_fmod_2(const Stream& stream, TensorSpan<float> output, TensorView<float> x, TensorView<float> y);
@@ -370,6 +376,7 @@ void eltwise_fmod_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x,
template void eltwise_sum_2(const Stream& stream, TensorSpan<float> output, TensorView<float> x, TensorView<float> y);
template void eltwise_max_2(const Stream& stream, TensorSpan<float> output, TensorView<float> x, TensorView<float> y);
template void eltwise_min_2(const Stream& stream, TensorSpan<float> output, TensorView<float> x, TensorView<float> y);
template void eltwise_pow_2(const Stream& stream, TensorSpan<float> output, TensorView<float> x, TensorView<float> y);
template void eltwise_mod_2(const Stream& stream, TensorSpan<int8_t> output, TensorView<int8_t> x, TensorView<int8_t> y);
template void eltwise_fmod_2(const Stream& stream, TensorSpan<int8_t> output, TensorView<int8_t> x, TensorView<int8_t> y);
@@ -380,6 +387,7 @@ void eltwise_fmod_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x,
template void eltwise_sum_2(const Stream& stream, TensorSpan<int8_t> output, TensorView<int8_t> x, TensorView<int8_t> y);
template void eltwise_max_2(const Stream& stream, TensorSpan<int8_t> output, TensorView<int8_t> x, TensorView<int8_t> y);
template void eltwise_min_2(const Stream& stream, TensorSpan<int8_t> output, TensorView<int8_t> x, TensorView<int8_t> y);
template void eltwise_pow_2(const Stream& stream, TensorSpan<int8_t> output, TensorView<int8_t> x, TensorView<int8_t> y);
template void eltwise_mod_2(const Stream& stream, TensorSpan<uint8_t> output, TensorView<uint8_t> x, TensorView<uint8_t> y);
template void eltwise_fmod_2(const Stream& stream, TensorSpan<uint8_t> output, TensorView<uint8_t> x, TensorView<uint8_t> y);
@@ -390,6 +398,7 @@ void eltwise_fmod_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x,
template void eltwise_sum_2(const Stream& stream, TensorSpan<uint8_t> output, TensorView<uint8_t> x, TensorView<uint8_t> y);
template void eltwise_max_2(const Stream& stream, TensorSpan<uint8_t> output, TensorView<uint8_t> x, TensorView<uint8_t> y);
template void eltwise_min_2(const Stream& stream, TensorSpan<uint8_t> output, TensorView<uint8_t> x, TensorView<uint8_t> y);
template void eltwise_pow_2(const Stream& stream, TensorSpan<uint8_t> output, TensorView<uint8_t> x, TensorView<uint8_t> y);
template void eltwise_mod_2(const Stream& stream, TensorSpan<int32_t> output, TensorView<int32_t> x, TensorView<int32_t> y);
template void eltwise_fmod_2(const Stream& stream, TensorSpan<int32_t> output, TensorView<int32_t> x, TensorView<int32_t> y);
@@ -400,6 +409,7 @@ void eltwise_fmod_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x,
template void eltwise_sum_2(const Stream& stream, TensorSpan<int32_t> output, TensorView<int32_t> x, TensorView<int32_t> y);
template void eltwise_max_2(const Stream& stream, TensorSpan<int32_t> output, TensorView<int32_t> x, TensorView<int32_t> y);
template void eltwise_min_2(const Stream& stream, TensorSpan<int32_t> output, TensorView<int32_t> x, TensorView<int32_t> y);
template void eltwise_pow_2(const Stream& stream, TensorSpan<int32_t> output, TensorView<int32_t> x, TensorView<int32_t> y);
template void eltwise_mod_2(const Stream& stream, TensorSpan<int64_t> output, TensorView<int64_t> x, TensorView<int64_t> y);
template void eltwise_fmod_2(const Stream& stream, TensorSpan<int64_t> output, TensorView<int64_t> x, TensorView<int64_t> y);
@@ -410,6 +420,7 @@ void eltwise_fmod_2(const Stream& stream, TensorSpan<T> output, TensorView<T> x,
template void eltwise_sum_2(const Stream& stream, TensorSpan<int64_t> output, TensorView<int64_t> x, TensorView<int64_t> y);
template void eltwise_max_2(const Stream& stream, TensorSpan<int64_t> output, TensorView<int64_t> x, TensorView<int64_t> y);
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);
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
modules/dnn/src/cuda/functors.hpp
+15
View File
@@ -833,6 +833,21 @@ struct FModFunctor {
}
};
template <class T>
struct PowFunctor {
struct Params {
CUDA4DNN_HOST_DEVICE Params() {}
};
CUDA4DNN_DEVICE PowFunctor() { }
CUDA4DNN_DEVICE PowFunctor(const Params& params) { }
CUDA4DNN_DEVICE T operator()(T x, T y) {
using csl::device::pow;
return pow(x, y);
}
};
}}}} /* namespace cv::dnn::cuda4dnn::kernels */
#endif /* OPENCV_DNN_SRC_CUDA_FUNCTORS_HPP */
modules/dnn/src/cuda/math.hpp
+2 -1
View File
@@ -90,7 +90,8 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace csl { namespace de
template <> inline __device__ float tanh(float val) { return tanhf(val); }
template <> inline __device__ double tanh(double val) { return ::tanh(val); }
template <class T> __device__ T pow(T val, T exp);
// caution: overflow warning casting from double to low-bit types
template <class T> __device__ T pow(T val, T exp) { return T(::pow(double(val), double(exp))); }
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
template <> inline __device__ __half pow(__half val, __half exp) { return powf(val, exp); }
#endif
modules/dnn/src/cuda4dnn/kernels/eltwise_ops.hpp
+3
View File
@@ -39,6 +39,9 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
template <class T>
void eltwise_fmod_2(const csl::Stream& stream, csl::TensorSpan<T> output, csl::TensorView<T> x, csl::TensorView<T> y);
template <class T>
void eltwise_pow_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 */
modules/dnn/src/cuda4dnn/primitives/eltwise.hpp
+8 -4
View File
@@ -30,6 +30,7 @@ namespace cv { namespace dnn { namespace cuda4dnn {
SUB,
MOD,
FMOD,
POW,
};
class EltwiseOpBase : public CUDABackendNode {
@@ -62,7 +63,6 @@ namespace cv { namespace dnn { namespace cuda4dnn {
const std::vector<cv::Ptr<BackendWrapper>>& outputs,
csl::Workspace& workspace) override
{
CV_Assert(inputs.size() >= 2);
CV_Assert(outputs.size() == 1);
CV_Assert(coeffs.size() == 0 || op == EltwiseOpType::SUM);
@@ -94,10 +94,13 @@ namespace cv { namespace dnn { namespace cuda4dnn {
case EltwiseOpType::SUB: kernels::eltwise_sub_2<T>(stream, output, input_x, input_y); break;
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;
}
}
else
{
} else if (inputs.size() == 1) {
auto input_wrapper_0 = inputs[0].dynamicCast<wrapper_type>();
auto input_0 = input_wrapper_0->getView();
csl::tensor_ops::copy(stream, output, input_0);
} else {
auto input_wrapper_0 = inputs[0].dynamicCast<wrapper_type>();
auto input_0 = input_wrapper_0->getView();
@@ -128,6 +131,7 @@ namespace cv { namespace dnn { namespace cuda4dnn {
case EltwiseOpType::SUB: kernels::eltwise_sub_2<T>(stream, output, output, input); break;
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;
}
}
}
modules/dnn/src/layers/nary_eltwise_layers.cpp
+433 -301
View File
@@ -44,13 +44,11 @@ public:
std::vector<int> all_ndims;
std::vector<std::vector<int>> orig_shapes;
std::vector<std::vector<size_t>> orig_steps;
std::vector<char*> ptrs;
std::vector<std::vector<int>> shapes;
std::vector<std::vector<size_t>> steps;
std::vector<size_t> elemsize;
NaryEltwiseHelper() {
}
NaryEltwiseHelper() {}
void init(const std::vector<Mat>& inputs, const std::vector<Mat>& outputs)
{
@@ -59,7 +57,6 @@ public:
all_ndims.clear();
orig_shapes.clear();
orig_steps.clear();
ptrs.clear();
shapes.clear();
steps.clear();
elemsize.clear();
@@ -81,7 +78,6 @@ public:
shapes = std::vector<std::vector<int>>(narrays, std::vector<int>(max_ndims, 0));
steps = std::vector<std::vector<size_t>>(narrays, std::vector<size_t>(max_ndims, 0));
ptrs = std::vector<char*>(narrays, nullptr);
for(i = 0; i <= ninputs; i++) {
all_ndims.push_back(i == 0 ? out_ndims : inp_ndims[i-1]);
@@ -279,7 +275,7 @@ public:
#ifdef HAVE_VULKAN
if (backendId == DNN_BACKEND_VKCOM)
return op == OPERATION::ADD || op == OPERATION::PROD || op == OPERATION::SUB ||
op == OPERATION::DIV ;
op == OPERATION::DIV;
#endif
if (backendId == DNN_BACKEND_CUDA) {
@@ -325,7 +321,7 @@ public:
outputs_arr.getMatVector(outputs);
helper.init(inputs, outputs);
CV_Assert(helper.prepare_for_broadcast_op());
CV_CheckTrue(helper.prepare_for_broadcast_op(), "NaryEltwiseLayer: Preparation for broadcasting failed");
}
bool getMemoryShapes(const std::vector<MatShape> &inputs,
@@ -333,8 +329,12 @@ public:
std::vector<MatShape> &outputs,
std::vector<MatShape> &internals) const CV_OVERRIDE
{
MatShape outShape = findCommonShape(inputs);
outputs.assign(1, outShape);
if (inputs.size() == 1) {
outputs.assign(1, inputs.front());
} else {
MatShape outShape = findCommonShape(inputs);
outputs.assign(1, outShape);
}
return false;
}
@@ -360,6 +360,16 @@ public:
return;
}
if (op == OPERATION::POW) {
/*
First input: exponent of Type T;
Second input: power of the exponent of Type T1;
Output: same type T as first input's.
*/
outputs.assign(1, inputs.front());
return;
}
CV_Assert(inputs.size());
for (auto input : inputs)
{
@@ -378,175 +388,232 @@ public:
template <typename T, typename RESULT_T, typename Functor>
void binary_forward_impl(
int ndims, const std::vector<int>& shape,
const char* data1, const std::vector<size_t>& step1,
const char* data2, const std::vector<size_t>& step2,
char* data, const std::vector<size_t>& step,
const Functor& op)
{
void binary_forward_impl(const Functor& op, int ndims, const std::vector<int>& shape,
const char* data1, const std::vector<size_t>& step1,
const char* data2, const std::vector<size_t>& step2,
char* data, const std::vector<size_t>& step, size_t block_size) {
size_t dp1 = 0, dp2 = 0, dp = 0;
int k, n1 = 1, n2 = 1;
size_t inplane_step1 = 0, inplane_step2 = 0, inplane_step = 0;
size_t plane_idx, nplanes = 1;
int plane_size = 1;
size_t nplanes = 1;
if (ndims >= 1) {
dp1 = step1[ndims-1]/sizeof(T);
dp2 = step2[ndims-1]/sizeof(T);
dp = step[ndims-1]/sizeof(RESULT_T);
n1 = shape[ndims-1];
dp1 = step1.back() / sizeof(T);
dp2 = step2.back() / sizeof(T);
dp = step.back() / sizeof(RESULT_T);
plane_size = shape.back();
if (ndims >= 2) {
inplane_step1 = step1[ndims-2];
inplane_step2 = step2[ndims-2];
inplane_step = step[ndims-2];
n2 = shape[ndims-2];
for (k = 0; k < ndims-2; k++) nplanes *= shape[k];
nplanes = std::accumulate(shape.begin(), shape.end() - 1, 1, std::multiplies<size_t>());
}
}
for (plane_idx = 0; plane_idx < nplanes; plane_idx++) {
const char* ptr1_ = data1;
const char* ptr2_ = data2;
char* ptr_ = data;
size_t idx = plane_idx;
for (k = ndims-3; k >= 0; k--) {
size_t next_idx = idx/shape[k];
int i_k = (int)(idx - next_idx*shape[k]);
ptr1_ += i_k*step1[k];
ptr2_ += i_k*step2[k];
ptr_ += i_k*step[k];
idx = next_idx;
}
for (int i2 = 0; i2 < n2; i2++, ptr1_ += inplane_step1,
ptr2_ += inplane_step2,
ptr_ += inplane_step)
{
const T* ptr1 = (const T*)ptr1_;
const T* ptr2 = (const T*)ptr2_;
RESULT_T* ptr = (RESULT_T*)ptr_;
if (nplanes == 1) { // parallelize within the plane
const T* ptr1 = (const T*)data1;
const T* ptr2 = (const T*)data2;
RESULT_T* ptr = (RESULT_T*)data;
auto worker = [&](const Range &r) {
if (dp1 == 1 && dp2 == 1 && dp == 1) {
for(int i1 = 0; i1 < n1; i1++)
ptr[i1] = op(ptr1[i1], ptr2[i1]);
for(int i = r.start; i < r.end; i++) {
ptr[i] = op(ptr1[i], ptr2[i]);
}
} else if (dp1 == 1 && dp2 == 0 && dp == 1){
T x2 = *ptr2;
for(int i1 = 0; i1 < n1; i1++)
ptr[i1] = op(ptr1[i1], x2);
for(int i = r.start; i < r.end; i++) {
ptr[i] = op(ptr1[i], x2);
}
} else if (dp1 == 0 && dp2 == 1 && dp == 1){
T x1 = *ptr1;
for(int i1 = 0; i1 < n1; i1++)
ptr[i1] = op(x1, ptr2[i1]);
for(int i = r.start; i < r.end; i++) {
ptr[i] = op(x1, ptr2[i]);
}
} else {
for(int i1 = 0; i1 < n1; i1++, ptr1 += dp1, ptr2 += dp2, ptr += dp)
for(int i = r.start; i < r.end; i++, ptr1 += dp1, ptr2 += dp2, ptr += dp) {
*ptr = op(*ptr1, *ptr2);
}
}
}
};
double nstripes = plane_size * (1.0 / double(block_size));
parallel_for_(Range(0, plane_size), worker, nstripes);
} else { // parallelize across planes
auto worker = [&](const Range &r) {
for (int plane_idx = r.start; plane_idx < r.end; plane_idx++) {
const char* ptr1_ = data1;
const char* ptr2_ = data2;
char* ptr_ = data;
size_t idx = plane_idx;
for (int k = ndims - 2; k >= 0; k--) {
size_t next_idx = idx / shape[k];
size_t i_k = (int)(idx - next_idx * shape[k]);
ptr1_ += i_k * step1[k];
ptr2_ += i_k * step2[k];
ptr_ += i_k * step[k];
idx = next_idx;
}
const T* ptr1 = (const T*)ptr1_;
const T* ptr2 = (const T*)ptr2_;
RESULT_T* ptr = (RESULT_T*)ptr_;
if (dp1 == 1 && dp2 == 1 && dp == 1) {
for(int i = 0; i < plane_size; i++) {
ptr[i] = op(ptr1[i], ptr2[i]);
}
} else if (dp1 == 1 && dp2 == 0 && dp == 1){
T x2 = *ptr2;
for(int i = 0; i < plane_size; i++) {
ptr[i] = op(ptr1[i], x2);
}
} else if (dp1 == 0 && dp2 == 1 && dp == 1){
T x1 = *ptr1;
for(int i = 0; i < plane_size; i++) {
ptr[i] = op(x1, ptr2[i]);
}
} else {
for(int i = 0; i < plane_size; i++, ptr1 += dp1, ptr2 += dp2, ptr += dp) {
*ptr = op(*ptr1, *ptr2);
}
}
}
};
double nstripes = nplanes * (1.0 / double(block_size));
parallel_for_(Range(0, nplanes), worker, nstripes);
}
}
/*
Elementwise binary operator (like +, -, x, /, etc.) which takes two operands
*/
template <typename T, typename RESULT_T, typename Functor>
void binary_forward(const Functor& f, const std::vector<Mat>& inputs, std::vector<Mat>& outputs)
{
void binary_forward(const Functor& f, const std::vector<Mat>& inputs, std::vector<Mat>& outputs, size_t block_size = 6e6) {
const Mat& a = inputs[0];
const Mat& b = inputs[1];
Mat& out = outputs[0];
CV_Assert(helper.shapes.size() == 3 && helper.steps.size() == 3);
binary_forward_impl<T, RESULT_T, Functor>(
helper.max_ndims, helper.shapes[0], a.ptr<char>(), helper.steps[1],
b.ptr<char>(), helper.steps[2], out.ptr<char>(), helper.steps[0],
f);
binary_forward_impl<T, RESULT_T, Functor>(f, helper.max_ndims, helper.shapes[0], a.ptr<char>(), helper.steps[1],
b.ptr<char>(), helper.steps[2], out.ptr<char>(), helper.steps[0], block_size);
}
template<typename T, typename Functor>
void nary_forward_impl(
const Functor& f, const T scale, int ninputs, int ndims, const std::vector<int>& shape,
const char** inp, char* out,
const std::vector<std::vector<size_t>>& steps, std::vector<char*>& ptrs)
{
void nary_forward_impl(const Functor& op, const T scale, int ninputs, int ndims, const std::vector<int>& shape,
const char** inp, char* out, const std::vector<std::vector<size_t>>& steps, size_t block_size) {
CV_Assert(ndims >= 2);
size_t dp = steps[0][ndims-1]/sizeof(T);
size_t dp1 = steps[1][ndims-1]/sizeof(T);
size_t dp2 = steps[2][ndims-1]/sizeof(T);
size_t dp = steps[0].back() / sizeof(T);
size_t dp1 = steps[1].back() / sizeof(T);
size_t dp2 = steps[2].back() / sizeof(T);
enum { BLOCK_SIZE = 1024 };
T blck[BLOCK_SIZE];
int plane_size = shape.back();
int nplanes = std::accumulate(shape.begin(), shape.end() - 1, 1, std::multiplies<int>());
int k, i, di1=0, n1 = shape[ndims-1], n2 = shape[ndims-2];
int second = ninputs == 1 ? 1 : 2;
size_t plane_idx, nplanes = 1;
for (k = 0; k < ndims-2; k++) nplanes *= shape[k];
for (plane_idx = 0; plane_idx < nplanes; plane_idx++) {
if (nplanes == 1) { // parallelize within the plane
AutoBuffer<char> buf_ptrs(steps.size());
auto ptrs = (char**)buf_ptrs.data();
ptrs[0] = out;
for (i = 0; i < ninputs; i++) ptrs[i+1] = (char*)inp[i];
size_t idx = plane_idx;
for (k = ndims-3; k >= 0; k--) {
size_t next_idx = idx/shape[k];
int i_k = (int)(idx - next_idx*shape[k]);
for (i = 0; i < ninputs; i++)
ptrs[i] += i_k*steps[i][k];
idx = next_idx;
for (int i = 0; i < ninputs; i++) {
ptrs[i+1] = (char*)inp[i];
}
for (int i2 = 0; i2 < n2; i2++)
{
const T* ptr1 = (const T*)(ptrs[1] + steps[1][ndims-2]*i2);
const T* ptr2 = (const T*)(ptrs[second] + steps[second][ndims-2]*i2);
T* ptr = (T*)(ptrs[0] + steps[0][ndims-2]*i2);
if (ninputs <= 2) {
if (dp1 == 1 && dp2 == 1) {
for (int i1 = 0; i1 < n1; i1++)
ptr[i1] = saturate_cast<T>(f(ptr1[i1], ptr2[i1])*scale);
} else {
for(int i1 = 0; i1 < n1; i1++, ptr1 += dp1, ptr2 += dp2, ptr += dp)
*ptr = saturate_cast<T>(f(*ptr1, *ptr2)*scale);
auto worker = [&](const Range &r) {
const T* ptr1 = (const T*)(ptrs[1]);
const T* ptr2 = (const T*)(ptrs[2]);
T* ptr = (T*)(ptrs[0]);
if (dp == 1 && dp1 == 1 && dp2 == 1) {
for (int i = r.start; i < r.end; i++) {
ptr[i] = op(ptr1[i], ptr2[i]);
}
for (int j = 2; j < ninputs; j++) {
size_t dpj = steps[j + 1].back() / sizeof(T);
const T* ptrj = (const T*)(ptrs[j + 1]);
if (dpj == 1) {
for (int i = r.start; i < r.end; i++) {
ptr[i] = saturate_cast<T>(op(ptr[i], ptrj[i]) * scale);
}
} else {
for (int i = r.start; i < r.end; i++, ptrj += dpj) {
ptr[i] = saturate_cast<T>(op(ptr[i], *ptrj) * scale);
}
}
}
} else {
for (int i1 = 0; i1 < n1; i1 += di1, ptr += di1) {
di1 = BLOCK_SIZE < n1-i1 ? BLOCK_SIZE : n1-i1;
if (dp1 == 1 && dp2 == 1) {
for (int j = 0; j < di1; j++)
blck[j] = f(ptr1[j], ptr2[j]);
ptr1 += di1;
ptr2 += di1;
} else {
for(int j = 0; j < di1; j++, ptr1 += dp1, ptr2 += dp2)
blck[j] = f(*ptr1, *ptr2);
auto *tmp = ptr;
for (int i = r.start; i < r.end; i++, ptr += dp, ptr1 += dp1, ptr2 += dp2) {
*ptr = op(*ptr1, *ptr2);
}
ptr = tmp;
for (int j = 2; j < ninputs; j++) {
size_t dpj = steps[j + 1].back() / sizeof(T);
const T* ptr_j = (const T*)(ptrs[j + 1]);
for (int i = r.start; i < r.end; i++, ptr += dp, ptr_j += dpj) {
*ptr = saturate_cast<T>(op(*ptr, *ptr_j) * scale);
}
for(i = 2; i < ninputs; i++) {
int dp_i = steps[i+1][ndims-1]/sizeof(T);
const T* ptr_i = (const T*)(ptrs[i+1] +
steps[i+1][ndims-2]*i2) + i1*dp_i;
if (dp_i == 1) {
if (i < ninputs-1) {
for (int j = 0; j < di1; j++)
blck[j] = f(blck[j], ptr_i[j]);
} else {
for (int j = 0; j < di1; j++)
ptr[j] = saturate_cast<T>(f(blck[j], ptr_i[j]) * scale);
}
}
};
double nstripes = plane_size * (1.0 / double(block_size));
parallel_for_(Range(0, plane_size), worker, nstripes);
} else { // parallelize across the plane
auto worker = [&](const Range &r) {
AutoBuffer<char> buf_ptrs(steps.size());
auto ptrs = (char**)buf_ptrs.data();
for (int plane_idx = r.start; plane_idx < r.end; plane_idx++) {
ptrs[0] = out;
for (int i = 0; i < ninputs; i++) ptrs[i+1] = (char*)inp[i];
size_t idx = plane_idx;
for (int k = ndims - 2; k >= 0; k--) {
size_t next_idx = idx / shape[k];
int i_k = (int)(idx - next_idx * shape[k]);
for (int i = 0; i <= ninputs; i++) {
ptrs[i] += i_k * steps[i][k];
}
idx = next_idx;
}
const T* ptr1 = (const T*)(ptrs[1]);
const T* ptr2 = (const T*)(ptrs[2]);
T* ptr = (T*)(ptrs[0]);
if (dp == 1 && dp1 == 1 && dp2 == 1) {
for (int i = 0; i < plane_size; i++) {
ptr[i] = saturate_cast<T>(op(ptr1[i], ptr2[i]) * scale);
}
for (int j = 2; j < ninputs; j++) {
size_t dpj = steps[j + 1].back() / sizeof(T);
const T* ptrj = (const T*)(ptrs[j + 1]);
if (dpj == 1) {
for (int i = 0; i < plane_size; i++) {
ptr[i] = op(ptr[i], saturate_cast<T>(ptrj[i] * scale));
}
} else {
if (i < ninputs-1) {
for (int j = 0; j < di1; j++, ptr_i += dp_i)
blck[j] = f(blck[j], *ptr_i);
} else {
for (int j = 0; j < di1; j++, ptr_i += dp_i)
ptr[j] = saturate_cast<T>(f(blck[j], *ptr_i) * scale);
for (int i = 0; i < plane_size; i++, ptrj += dpj) {
ptr[i] = op(ptr[i], saturate_cast<T>(*ptrj * scale));
}
}
}
} else {
auto *tmp = ptr;
for (int i = 0; i < plane_size; i++, ptr += dp, ptr1 += dp1, ptr2 += dp2) {
*ptr = saturate_cast<T>(op(*ptr1, *ptr2) * scale);
}
ptr = tmp;
for (int j = 2; j < ninputs; j++) {
size_t dpj = steps[j + 1].back() / sizeof(T);
const T* ptrj = (const T*)(ptrs[j + 1]);
for (int i = 0; i < plane_size; i++, ptr += dp, ptrj += dpj) {
*ptr = op(*ptr, saturate_cast<T>(*ptrj * scale));
}
}
}
}
}
};
double nstripes = nplanes * (1.0 / double(block_size));
parallel_for_(Range(0, nplanes), worker, nstripes);
}
}
/*
Elementwise nary operator (like sum, mean, etc.) which takes at least one operand
*/
template <typename T, typename Functor>
void nary_forward(
const Functor& f, T scale,
const std::vector<Mat>& inputs, std::vector<Mat>& outputs
)
{
void nary_forward(const Functor& f, T scale,
const std::vector<Mat>& inputs, std::vector<Mat>& outputs,
size_t block_size = 6e6) {
// collect all input info
std::vector<const char*> v_inp;
std::transform(inputs.begin(), inputs.end(), std::back_inserter(v_inp), [] (const Mat& m) { return m.template ptr<const char>(); });
@@ -555,13 +622,14 @@ public:
// collect output info
char* out = outputs[0].ptr<char>();
nary_forward_impl<T>(
f, scale, helper.ninputs, helper.max_ndims, helper.shapes[0], inp, out, helper.steps, helper.ptrs);
nary_forward_impl<T, Functor>(f, scale, helper.ninputs, helper.max_ndims, helper.shapes[0], inp, out, helper.steps, block_size);
}
/*
Elementwise ternary operator (like where) which takes three operands
*/
template <typename T_INP1, typename T_INP2, typename T_INP3, typename T_OUT, typename Functor>
void trinary_forward(const Functor& f, const std::vector<Mat>& inputs, std::vector<Mat>& outputs)
{
void ternary_forward(const Functor& f, const std::vector<Mat>& inputs, std::vector<Mat>& outputs, size_t block_size = 6e6) {
const Mat& a = inputs[0];
const Mat& b = inputs[1];
const Mat& c = inputs[2];
@@ -569,69 +637,112 @@ public:
CV_Assert(helper.shapes.size() == 4 && helper.steps.size() == 4);
trinary_forward_impl<T_INP1, T_INP2, T_INP3, T_OUT, Functor>(
helper.max_ndims, helper.shapes[0], a.ptr<char>(), helper.steps[1], b.ptr<char>(), helper.steps[2],
c.ptr<char>(), helper.steps[3], out.ptr<char>(), helper.steps[0],
f);
ternary_forward_impl<T_INP1, T_INP2, T_INP3, T_OUT, Functor>(f, helper.max_ndims, helper.shapes[0],
a.ptr<char>(), helper.steps[1],
b.ptr<char>(), helper.steps[2],
c.ptr<char>(), helper.steps[3],
out.ptr<char>(), helper.steps[0], block_size);
}
template <typename T_INP1, typename T_INP2, typename T_INP3, typename T_OUT, typename Functor>
void trinary_forward_impl(
int ndims, const std::vector<int>& shape,
void ternary_forward_impl(
const Functor& op, int ndims, const std::vector<int>& shape,
const char* data1, const std::vector<size_t>& step1,
const char* data2, const std::vector<size_t>& step2,
const char* data3, const std::vector<size_t>& step3,
char* data, const std::vector<size_t>& step,
const Functor& op)
{
assert(ndims >= 2);
size_t dp1 = step1[ndims-1]/sizeof(T_INP1);
size_t dp2 = step2[ndims-1]/sizeof(T_INP2);
size_t dp3 = step3[ndims-1]/sizeof(T_INP3);
size_t dp = step[ndims-1]/sizeof(T_OUT);
int k, n1 = shape[ndims-1], n2 = shape[ndims-2];
size_t plane_idx, nplanes = 1;
for (k = 0; k < ndims-2; k++) nplanes *= shape[k];
char* data, const std::vector<size_t>& step, size_t block_size) {
CV_Assert(ndims >= 2);
size_t dp1 = step1.back() / sizeof(T_INP1);
size_t dp2 = step2.back() / sizeof(T_INP2);
size_t dp3 = step3.back() / sizeof(T_INP3);
size_t dp = step.back() / sizeof(T_OUT);
int plane_size = shape.back();
int nplanes = std::accumulate(shape.begin(), shape.end() - 1, 1, std::multiplies<int>());
for (plane_idx = 0; plane_idx < nplanes; plane_idx++)
{
const char* ptr1_ = data1;
const char* ptr2_ = data2;
const char* ptr3_ = data3;
char* ptr_ = data;
size_t idx = plane_idx;
for (k = ndims-3; k >= 0; k--)
{
size_t next_idx = idx/shape[k];
int i_k = (int)(idx - next_idx*shape[k]);
ptr1_ += i_k*step1[k];
ptr2_ += i_k*step2[k];
ptr3_ += i_k*step3[k];
ptr_ += i_k*step[k];
idx = next_idx;
}
for (int i2 = 0; i2 < n2; i2++, ptr1_ += step1[ndims-2],
ptr2_ += step2[ndims-2],
ptr3_ += step3[ndims-2],
ptr_ += step[ndims-2])
{
const T_INP1* ptr1 = (const T_INP1*)ptr1_;
const T_INP2* ptr2 = (const T_INP2*)ptr2_;
const T_INP3* ptr3 = (const T_INP3*)ptr3_;
T_OUT* ptr = (T_OUT*)ptr_;
if (dp1 == 1 && dp2 == 1 && dp3 == 1 && dp == 1)
{
for(int i1 = 0; i1 < n1; i1++)
ptr[i1] = op(ptr1[i1], ptr2[i1], ptr3[i1]);
}
else
{
for(int i1 = 0; i1 < n1; i1++, ptr1 += dp1, ptr2 += dp2, ptr3 += dp3, ptr += dp)
if (nplanes == 1) { // parallelize within the plane
const auto *ptr1 = (const T_INP1*)data1;
const auto *ptr2 = (const T_INP2*)data2;
const auto *ptr3 = (const T_INP3*)data3;
auto* ptr = (T_OUT*)data;
auto worker = [&](const Range &r) {
if (dp1 == 1 && dp2 == 1 && dp3 == 1 && dp == 1) {
for (int i = r.start; i < r.end; i++) {
ptr[i] = op(ptr1[i], ptr2[i], ptr3[i]);
}
} else if (dp1 == 0 && dp2 == 1 && dp3 == 1 && dp == 1){
auto x1 = *ptr1;
for (int i = r.start; i < r.end; i++) {
ptr[i] = op(x1, ptr2[i], ptr3[i]);
}
} else if (dp1 == 1 && dp2 == 0 && dp3 == 1 && dp == 1){
auto x2 = *ptr2;
for (int i = r.start; i < r.end; i++) {
ptr[i] = op(ptr1[i], x2, ptr3[i]);
}
} else if (dp1 == 1 && dp2 == 1 && dp3 == 1 && dp == 1) {
auto x3 = *ptr3;
for (int i = r.start; i < r.end; i++) {
ptr[i] = op(ptr1[i], ptr2[i], x3);
}
} else {
for(int i = r.start; i < r.end; i++, ptr1 += dp1, ptr2 += dp2, ptr3 += dp3, ptr += dp) {
*ptr = op(*ptr1, *ptr2, *ptr3);
}
}
}
};
double nstripes = plane_size * (1.0 / double(block_size));
parallel_for_(Range(0, plane_size), worker, nstripes);
} else { // parallelize across planes
auto worker = [&](const Range &r) {
for (int plane_idx = r.start; plane_idx < r.end; plane_idx++) {
const char* ptr1_ = data1;
const char* ptr2_ = data2;
const char* ptr3_ = data3;
char* ptr_ = data;
size_t idx = plane_idx;
for (int k = ndims - 2; k >= 0; k--)
{
size_t next_idx = idx / shape[k];
int i_k = (int)(idx - next_idx * shape[k]);
ptr1_ += i_k * step1[k];
ptr2_ += i_k * step2[k];
ptr3_ += i_k * step3[k];
ptr_ += i_k * step[k];
idx = next_idx;
}
const auto *ptr1 = (const T_INP1*)ptr1_;
const auto *ptr2 = (const T_INP2*)ptr2_;
const auto *ptr3 = (const T_INP3*)ptr3_;
auto* ptr = (T_OUT*)ptr_;
if (dp1 == 1 && dp2 == 1 && dp3 == 1 && dp == 1) {
for (int i = 0; i < plane_size; i++) {
ptr[i] = op(ptr1[i], ptr2[i], ptr3[i]);
}
} else if (dp1 == 0 && dp2 == 1 && dp3 == 1 && dp == 1){
auto x1 = *ptr1;
for (int i = 0; i < plane_size; i++) {
ptr[i] = op(x1, ptr2[i], ptr3[i]);
}
} else if (dp1 == 1 && dp2 == 0 && dp3 == 1 && dp == 1){
auto x2 = *ptr2;
for (int i = 0; i < plane_size; i++) {
ptr[i] = op(ptr1[i], x2, ptr3[i]);
}
} else if (dp1 == 1 && dp2 == 1 && dp3 == 1 && dp == 1) {
auto x3 = *ptr3;
for (int i = 0; i < plane_size; i++) {
ptr[i] = op(ptr1[i], ptr2[i], x3);
}
} else {
for(int i = 0; i < plane_size; i++, ptr1 += dp1, ptr2 += dp2, ptr3 += dp3, ptr += dp) {
*ptr = op(*ptr1, *ptr2, *ptr3);
}
}
}
};
double nstripes = nplanes * (1.0 / double(block_size));
parallel_for_(Range(0, nplanes), worker, nstripes);
}
}
@@ -650,125 +761,133 @@ public:
inputs_arr.getMatVector(inputs);
outputs_arr.getMatVector(outputs);
// TODO: assert types
typeDispatch(inputs.back().type(), inputs.size(), inputs, outputs);
if (inputs.size() == 1) {
inputs[0].copyTo(outputs[0]);
return;
}
int type_for_dispatch = op == OPERATION::WHERE ? outputs.front().type() : inputs.front().type();
typeDispatch(type_for_dispatch, inputs.size(), inputs, outputs);
}
template<typename T, typename... Args>
inline void opDispatch(size_t ninputs, Args&&... args)
{
switch (op)
{
case OPERATION::EQUAL:
{
auto equal = [](const T &a, const T &b) { return a == b; };
binary_forward<T, bool>(equal, std::forward<Args>(args)...);
break;
if (ninputs == 2) { // Operators that take two operands
switch (op) {
case OPERATION::EQUAL: {
auto equal = [](const T &a, const T &b) { return a == b; };
binary_forward<T, bool>(equal, std::forward<Args>(args)...);
break;
}
case OPERATION::GREATER: {
auto greater = [](const T &a, const T &b) { return a > b; };
binary_forward<T, bool>(greater, std::forward<Args>(args)...);
break;
}
case OPERATION::GREATER_EQUAL: {
auto greater_equal = [](const T &a, const T &b) { return a >= b; };
binary_forward<T, bool>(greater_equal, std::forward<Args>(args)...);
break;
}
case OPERATION::LESS: {
auto less = [](const T &a, const T &b) { return a < b; };
binary_forward<T, bool>(less, std::forward<Args>(args)...);
break;
}
case OPERATION::LESS_EQUAL: {
auto less_equal = [](const T &a, const T &b) { return a <= b; };
binary_forward<T, bool>(less_equal, std::forward<Args>(args)...);
break;
}
case OPERATION::POW: {
auto pow = [] (const T& a, const T& b) { return std::pow(a, b); };
binary_forward<T, T>(pow, std::forward<Args>(args)..., 1e5);
break;
}
case OPERATION::BITSHIFT: {
auto bitshift = [] (const uint8_t &a, const uint8_t &b) { return a << b; };
binary_forward<T, T>(bitshift, std::forward<Args>(args)...);
break;
}
case OPERATION::MAX: {
auto max = [](const T &a, const T &b) { return std::max(a, b); };
binary_forward<T, T>(max, std::forward<Args>(args)...);
break;
}
case OPERATION::MEAN: {
auto mean = [](const T &a, const T &b) { return (a + b) / T{2}; };
binary_forward<T, T>(mean, std::forward<Args>(args)...);
break;
}
case OPERATION::MIN: {
auto min = [](const T &a, const T &b) { return std::min(a, b); };
binary_forward<T, T>(min, std::forward<Args>(args)...);
break;
}
case OPERATION::MOD: {
auto mod = [] (const T &a, const T &b) { return static_cast<T>(_mod(int(a), int(b))); };
binary_forward<T, T>(mod, std::forward<Args>(args)...);
break;
}
case OPERATION::FMOD: {
auto fmod = [](const T &a, const T &b) { return std::fmod(a, b); };
binary_forward<T, T>(fmod, std::forward<Args>(args)...);
break;
}
case OPERATION::PROD: {
auto prod = [](const T &a, const T &b) { return a * b; };
binary_forward<T, T>(prod, std::forward<Args>(args)...);
break;
}
case OPERATION::SUB: {
auto sub = [](const T &a, const T &b) { return a - b; };
binary_forward<T, T>(sub, std::forward<Args>(args)...);
break;
}
case OPERATION::ADD:
case OPERATION::SUM: {
auto sum = [](const T &a, const T &b) { return a + b; };
binary_forward<T, T>(sum, std::forward<Args>(args)...);
break;
}
case OPERATION::DIV: {
auto div = [](const T &a, const T &b) { return a / b; };
binary_forward<T, T>(div, std::forward<Args>(args)...);
break;
}
default: CV_Error(Error::StsBadArg, "Unsupported operation");
}
case OPERATION::GREATER:
} else if (ninputs == 3 && op == OPERATION::WHERE) { // Operators that take three operands
auto where = [](const T &a, const T &b, const T &c) { return a ? b : c; };
ternary_forward<bool, T, T, T>(where, std::forward<Args>(args)...);
} else { // Operators that can take multiple (>= 3) operands
switch (op)
{
auto greater = [](const T &a, const T &b) { return a > b; };
binary_forward<T, bool>(greater, std::forward<Args>(args)...);
break;
case OPERATION::MAX: {
auto max = [](const T &a, const T &b) { return std::max(a, b); };
nary_forward<T>(max, T{1}, std::forward<Args>(args)...);
break;
}
case OPERATION::MEAN: {
// Sum up inputs and then calculate mean by scale = 1 / ninputs
auto sum = [](const T &a, const T &b) { return a + b; };
nary_forward<T>(sum, T{1} / ninputs, std::forward<Args>(args)...);
break;
}
case OPERATION::MIN: {
auto min = [](const T &a, const T &b) { return std::min(a, b); };
nary_forward<T>(min, T{1}, std::forward<Args>(args)...);
break;
}
case OPERATION::SUM: {
auto sum = [](const T &a, const T &b) { return a + b; };
nary_forward<T>(sum, T{1}, std::forward<Args>(args)...);
break;
}
default:
CV_Error(Error::StsBadArg, "Unsupported operation.");
}
case OPERATION::GREATER_EQUAL:
{
auto greater_equal = [](const T &a, const T &b) { return a >= b; };
binary_forward<T, bool>(greater_equal, std::forward<Args>(args)...);
break;
}
case OPERATION::LESS:
{
auto less = [](const T &a, const T &b) { return a < b; };
binary_forward<T, bool>(less, std::forward<Args>(args)...);
break;
}
case OPERATION::LESS_EQUAL:
{
auto less_equal = [](const T &a, const T &b) { return a <= b; };
binary_forward<T, bool>(less_equal, std::forward<Args>(args)...);
break;
}
case OPERATION::POW:
{
auto pow = [] (const T& a, const T& b) { return std::pow(a, b); };
binary_forward<T, T>(pow, std::forward<Args>(args)...);
break;
}
case OPERATION::BITSHIFT:
{
auto bitshift = [] (const uint8_t &a, const uint8_t &b) { return a << b; };
binary_forward<T, T>(bitshift, std::forward<Args>(args)...);
break;
}
case OPERATION::MAX:
{
auto max = [](const T &a, const T &b) { return std::max(a, b); };
nary_forward<T>(max, T{1}, std::forward<Args>(args)...);
break;
}
case OPERATION::MEAN:
{
auto mean = [](const T &a, const T &b) { return (a + b) / T{2}; };
nary_forward<T>(mean, T{1} / ninputs, std::forward<Args>(args)...);
break;
}
case OPERATION::MIN:
{
auto min = [](const T &a, const T &b) { return std::min(a, b); };
nary_forward<T>(min, T{1}, std::forward<Args>(args)...);
break;
}
case OPERATION::MOD:
{
auto mod = [] (const T &a, const T &b) { return static_cast<T>(_mod(int(a), int(b))); };
binary_forward<T, T>(mod, std::forward<Args>(args)...);
break;
}
case OPERATION::FMOD:
{
auto fmod = [](const T &a, const T &b) { return std::fmod(a, b); };
binary_forward<T, T>(fmod, std::forward<Args>(args)...);
break;
}
case OPERATION::PROD:
{
auto prod = [](const T &a, const T &b) { return a * b; };
binary_forward<T, T>(prod, std::forward<Args>(args)...);
break;
}
case OPERATION::SUB:
{
auto sub = [](const T &a, const T &b) { return a - b; };
binary_forward<T, T>(sub, std::forward<Args>(args)...);
break;
}
case OPERATION::SUM:
{
auto sum = [](const T &a, const T &b) { return a + b; };
nary_forward<T>(sum, T{1}, std::forward<Args>(args)...);
break;
}
case OPERATION::ADD:
{
auto add = [](const T &a, const T &b) { return a + b; };
binary_forward<T, T>(add, std::forward<Args>(args)...);
break;
}
case OPERATION::DIV:
{
auto div = [](const T &a, const T &b) { return a / b; };
binary_forward<T, T>(div, std::forward<Args>(args)...);
break;
}
case OPERATION::WHERE:
{
auto op_where = [](const bool &a, const T &b, const T &c) { return a ? b : c; };
trinary_forward<bool, T, T, T>(op_where, std::forward<Args>(args)...);
break;
}
default:
CV_Error(Error::StsBadArg, "Unsupported operation.");
};
}
@@ -868,6 +987,9 @@ public:
case OPERATION::FMOD:
op_ = cuda4dnn::EltwiseOpType::FMOD;
break;
case OPERATION::POW:
op_ = cuda4dnn::EltwiseOpType::POW;
break;
default: return Ptr<BackendNode>(); // return empty cuda_node if the EltwiseOpType is unsupported type.
};
@@ -932,6 +1054,16 @@ public:
#ifdef HAVE_DNN_NGRAPH
virtual Ptr<BackendNode> initNgraph(const std::vector<Ptr<BackendWrapper> >& inputs, const std::vector<Ptr<BackendNode> >& nodes) CV_OVERRIDE
{
// In case only one input
if (inputs.size() == 1) {
auto &ieInpNode = nodes[0].dynamicCast<InfEngineNgraphNode>()->node;
ngraph::OutputVector inp{ieInpNode};
auto blank = std::make_shared<ov::op::v0::Concat>(inp, 0);
return Ptr<BackendNode>(new InfEngineNgraphNode(blank));
}
// TODO: Support multiple (>=3) inputs
if (op == OPERATION::WHERE)
CV_CheckEQ(inputs.size(), 3u, "");
else
modules/dnn/src/onnx/onnx_importer.cpp
+1 -9
View File
@@ -2858,14 +2858,6 @@ void ONNXImporter::parseElementWise(LayerParams& layerParams, const opencv_onnx:
};
}
// element-wise layers that can have >=1 inputs but actually have one input
if (node_proto.input_size() == 1 && (op_type == "max" || op_type == "min" || op_type == "mean" || op_type == "sum"))
{
layerParams.type = "Identity";
addLayer(layerParams, node_proto);
return;
}
auto pre_broadcast_transform = [](Mat& t, int t_real_ndims) {
if (t.dims == 2 && t_real_ndims == 1 && t.size[1] == 1)
transpose(t, t);
@@ -3971,7 +3963,7 @@ void ONNXImporter::buildDispatchMap_ONNX_AI(int opset_version)
dispatch["Sub"] = dispatch["Mul"] = dispatch["Div"] = dispatch["GreaterOrEqual"] =
dispatch["LessOrEqual"] = dispatch["Mod"] = dispatch["And"] = dispatch["Or"] = dispatch["Xor"] = &ONNXImporter::parseElementWise;
dispatch["Sum"] = dispatch["Min"] = dispatch["Max"] = &ONNXImporter::parseElementWise;
dispatch["Sum"] = dispatch["Min"] = dispatch["Max"] = dispatch["Mean"] = &ONNXImporter::parseElementWise;
dispatch["Where"] = &ONNXImporter::parseElementWise;
dispatch["Range"] = &ONNXImporter::parseRange;
dispatch["Einsum"] = &ONNXImporter::parseEinsum;
modules/dnn/test/test_onnx_conformance.cpp
+10 -1
View File
@@ -967,6 +967,7 @@ public:
#endif
#ifdef HAVE_CUDA
static std::set<std::string> cuda_deny_list;
static std::set<std::string> cuda_fp16_deny_list;
#endif
Test_ONNX_conformance()
@@ -1046,6 +1047,9 @@ public:
cuda_deny_list = {
#include "test_onnx_conformance_layer_filter__cuda_denylist.inl.hpp"
};
cuda_fp16_deny_list = {
#include "test_onnx_conformance_layer_filter__cuda_fp16_denylist.inl.hpp"
};
#endif
}
@@ -1062,6 +1066,7 @@ std::set<std::string> Test_ONNX_conformance::vulkan_deny_list;
#endif
#ifdef HAVE_CUDA
std::set<std::string> Test_ONNX_conformance::cuda_deny_list;
std::set<std::string> Test_ONNX_conformance::cuda_fp16_deny_list;
#endif
TEST_P(Test_ONNX_conformance, Layer_Test)
@@ -1121,10 +1126,14 @@ TEST_P(Test_ONNX_conformance, Layer_Test)
#ifdef HAVE_CUDA
else if (backend == DNN_BACKEND_CUDA)
{
if (cuda_deny_list.find(name) != cuda_deny_list.end())
if (target == DNN_TARGET_CUDA && cuda_deny_list.find(name) != cuda_deny_list.end())
{
applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA, CV_TEST_TAG_DNN_SKIP_ONNX_CONFORMANCE);
}
if (target == DNN_TARGET_CUDA_FP16 && cuda_fp16_deny_list.find(name) != cuda_fp16_deny_list.end())
{
applyTestTag(CV_TEST_TAG_DNN_SKIP_CUDA_FP16, CV_TEST_TAG_DNN_SKIP_ONNX_CONFORMANCE);
}
}
#endif
else
modules/dnn/test/test_onnx_conformance_layer_filter__cuda_denylist.inl.hpp
-13
View File
@@ -73,21 +73,9 @@
"test_maxunpool_export_with_output_shape",
"test_mul_bcast",
"test_mul_uint8",
"test_reduce_prod_default_axes_keepdims_example", // FP16 only
"test_reduce_prod_default_axes_keepdims_random", // FP16 only
"test_reduce_prod_do_not_keepdims_random", // FP16 only
"test_reduce_prod_keepdims_random", // FP16 only
"test_reduce_prod_negative_axes_keepdims_random", // FP16 only
"test_reduce_sum_square_default_axes_keepdims_random", // FP16 only
"test_reduce_sum_square_do_not_keepdims_random", // FP16 only
"test_reduce_sum_square_keepdims_random", // FP16 only
"test_reduce_sum_square_negative_axes_keepdims_random", // FP16 only
"test_softmax_default_axis",
"test_softmax_large_number", // FP16 only
"test_softmax_large_number_expanded", // FP16 only
"test_sub_bcast",
"test_sub_uint8",
"test_tan", // FP16 only
"test_upsample_nearest",
"test_scatter_elements_with_axis",
"test_scatter_elements_with_duplicate_indices",
@@ -103,4 +91,3 @@
"test_scatternd_min",
"test_scatternd_multiply",
"test_nllloss_NCd1d2d3d4d5_none_no_weight_expanded", // crash: https://github.com/opencv/opencv/issues/25471
"test_pow" // accuracy issue with FP16
modules/dnn/test/test_onnx_conformance_layer_filter__cuda_fp16_denylist.inl.hpp
+21
View File
@@ -0,0 +1,21 @@
"test_basic_conv_with_padding", // (assert failed) !blobs.empty() in initCUDA
"test_basic_conv_without_padding", // (assert failed) !blobs.empty() in initCUDA
"test_conv_with_autopad_same", // (assert failed) !blobs.empty() in initCUDA
"test_conv_with_strides_and_asymmetric_padding", // (assert failed) !blobs.empty() in initCUDA
"test_conv_with_strides_no_padding", // (assert failed) !blobs.empty() in initCUDA
"test_conv_with_strides_padding", // (assert failed) !blobs.empty() in initCUDA
"test_dropout_default_ratio",
"test_logsoftmax_large_number", // fp16 accuracy issue
"test_logsoftmax_large_number_expanded", // fp16 accuracy issue
"test_maxpool_with_argmax_2d_precomputed_pads", // assertion failed mat.type() == CV_32F
"test_nllloss_NCd1d2d3d4d5_none_no_weight_expanded", // crash: https://github.com/opencv/opencv/issues/25471
"test_reduce_prod_default_axes_keepdims_example", // fallback to cpu, accuracy
"test_reduce_prod_default_axes_keepdims_random", // fallback to cpu, accuracy
"test_reduce_sum_square_default_axes_keepdims_random", // fallback to cpu, accuracy
"test_reduce_sum_square_do_not_keepdims_random", // fallback to cpu, accuracy
"test_reduce_sum_square_keepdims_random", // fallback to cpu, accuracy
"test_reduce_sum_square_negative_axes_keepdims_random", // fallback to cpu, accuracy
"test_pow", // fp16 accuracy issue
"test_softmax_large_number", // fp16 accuracy issue
"test_softmax_large_number_expanded", // fp16 accuracy issue
"test_tan", // fp16 accuracy issue
modules/dnn/test/test_onnx_conformance_layer_filter__openvino.inl.hpp
+11 -3
View File
@@ -86,7 +86,11 @@ CASE(test_adam)
CASE(test_adam_multiple)
// no filter
CASE(test_add)
// no filter
if (target == DNN_TARGET_OPENCL)
{
default_l1 = 0.00024; // Expected: (normL1) <= (l1), actual: 0.000234754 vs 1e-05
default_lInf = 0.0011; // Expected: (normInf) <= (lInf), actual: 0.00106502 vs 0.0001
}
CASE(test_add_bcast)
#if SKIP_SET_1
SKIP;
@@ -1044,7 +1048,11 @@ CASE(test_momentum)
CASE(test_momentum_multiple)
// no filter
CASE(test_mul)
// no filter
if (target == DNN_TARGET_OPENCL)
{
default_l1 = 0.00024; // Expected: (normL1) <= (l1), actual: 0.00023824 vs 1e-05
default_lInf = 0.0015; // Expected: (normInf) <= (lInf), actual: 0.00145674 vs 0.0001
}
CASE(test_mul_bcast)
#if SKIP_SET_1
SKIP;
@@ -1194,7 +1202,7 @@ CASE(test_or_bcast4v3d)
CASE(test_or_bcast4v4d)
// no filter
CASE(test_pow)
// no filter
SKIP_OPENCL_FP16;
CASE(test_pow_bcast_array)
// no filter
CASE(test_pow_bcast_scalar)
modules/dnn/test/test_onnx_conformance_layer_filter__vulkan_denylist.inl.hpp
+3
View File
@@ -68,6 +68,9 @@
"test_maxunpool_export_with_output_shape",
"test_maxunpool_export_without_output_shape",
"test_mul_uint8",
"test_pow_types_float32_int32", // vulkan backend does not take tensor other than float32 data type
"test_pow_types_float32_int64", // vulkan backend does not take tensor other than float32 data type
"test_pow_types_int", // vulkan backend does not take tensor other than float32 data type
"test_softmax_default_axis",
"test_sub_bcast",
"test_sub_uint8",
modules/dnn/test/test_onnx_conformance_layer_filter_opencv_ocl_fp16_denylist.inl.hpp
+1
View File
@@ -14,6 +14,7 @@
"test_maxpool_2d_same_upper",
"test_maxpool_2d_strides",
"test_maxpool_3d_default",
"test_pow", // fp16 accuracy issue
"test_softmax_large_number",
"test_softmax_large_number_expanded",
"test_split_equal_parts_1d",