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

Merge remote-tracking branch 'upstream/3.4' into merge-3.4

Browse Source
This commit is contained in:
Alexander Alekhin
2021-03-27 15:35:16 +00:00
parent ad2f5ccc66 bf03f5fa3a
commit 35eaacd1db
20 changed files with 558 additions and 84 deletions
Download Patch File Download Diff File Expand all files Collapse all files
modules/core/include/opencv2/core/hal/intrin_neon.hpp
+21
View File
@@ -1539,6 +1539,26 @@ OPENCV_HAL_IMPL_NEON_SELECT(v_float32x4, f32, u32)
OPENCV_HAL_IMPL_NEON_SELECT(v_float64x2, f64, u64)
#endif
#if CV_NEON_AARCH64
#define OPENCV_HAL_IMPL_NEON_EXPAND(_Tpvec, _Tpwvec, _Tp, suffix) \
inline void v_expand(const _Tpvec& a, _Tpwvec& b0, _Tpwvec& b1) \
{ \
b0.val = vmovl_##suffix(vget_low_##suffix(a.val)); \
b1.val = vmovl_high_##suffix(a.val); \
} \
inline _Tpwvec v_expand_low(const _Tpvec& a) \
{ \
return _Tpwvec(vmovl_##suffix(vget_low_##suffix(a.val))); \
} \
inline _Tpwvec v_expand_high(const _Tpvec& a) \
{ \
return _Tpwvec(vmovl_high_##suffix(a.val)); \
} \
inline _Tpwvec v_load_expand(const _Tp* ptr) \
{ \
return _Tpwvec(vmovl_##suffix(vld1_##suffix(ptr))); \
}
#else
#define OPENCV_HAL_IMPL_NEON_EXPAND(_Tpvec, _Tpwvec, _Tp, suffix) \
inline void v_expand(const _Tpvec& a, _Tpwvec& b0, _Tpwvec& b1) \
{ \
@@ -1557,6 +1577,7 @@ inline _Tpwvec v_load_expand(const _Tp* ptr) \
{ \
return _Tpwvec(vmovl_##suffix(vld1_##suffix(ptr))); \
}
#endif
OPENCV_HAL_IMPL_NEON_EXPAND(v_uint8x16, v_uint16x8, uchar, u8)
OPENCV_HAL_IMPL_NEON_EXPAND(v_int8x16, v_int16x8, schar, s8)
modules/core/src/directx.cpp
+2 -2
View File
@@ -1050,7 +1050,7 @@ bool ocl_convert_nv12_to_bgr(
k.args(clImageY, clImageUV, clBuffer, step, cols, rows);
size_t globalsize[] = { (size_t)cols, (size_t)rows };
size_t globalsize[] = { (size_t)cols/2, (size_t)rows/2 };
return k.run(2, globalsize, 0, false);
}
@@ -1071,7 +1071,7 @@ bool ocl_convert_bgr_to_nv12(
k.args(clBuffer, step, cols, rows, clImageY, clImageUV);
size_t globalsize[] = { (size_t)cols, (size_t)rows };
size_t globalsize[] = { (size_t)cols/2, (size_t)rows/2 };
return k.run(2, globalsize, 0, false);
}
modules/core/src/opencl/cvtclr_dx.cl
+41 -51
View File
@@ -91,63 +91,50 @@ void YUV2BGR_NV12_8u(
{
int x = get_global_id(0);
int y = get_global_id(1);
// each iteration computes 2*2=4 pixels
int x2 = x*2;
int y2 = y*2;
if (x + 1 < cols)
{
if (y + 1 < rows)
{
__global uchar* pDstRow1 = pBGR + mad24(y, bgrStep, mad24(x, NCHANNELS, 0));
__global uchar* pDstRow2 = pDstRow1 + bgrStep;
if (x2 + 1 < cols) {
if (y2 + 1 < rows) {
__global uchar *pDstRow1 = pBGR + mad24(y2, bgrStep, mad24(x2, NCHANNELS, 0));
__global uchar *pDstRow2 = pDstRow1 + bgrStep;
float4 Y1 = read_imagef(imgY, (int2)(x+0, y+0));
float4 Y2 = read_imagef(imgY, (int2)(x+1, y+0));
float4 Y3 = read_imagef(imgY, (int2)(x+0, y+1));
float4 Y4 = read_imagef(imgY, (int2)(x+1, y+1));
float4 Y1 = read_imagef(imgY, (int2)(x2 + 0, y2 + 0));
float4 Y2 = read_imagef(imgY, (int2)(x2 + 1, y2 + 0));
float4 Y3 = read_imagef(imgY, (int2)(x2 + 0, y2 + 1));
float4 Y4 = read_imagef(imgY, (int2)(x2 + 1, y2 + 1));
float4 Y = (float4)(Y1.x, Y2.x, Y3.x, Y4.x);
float4 UV = read_imagef(imgUV, (int2)(x/2, y/2)) - d2;
float4 UV = read_imagef(imgUV, (int2)(x, y)) - d2;
__constant float* coeffs = c_YUV2RGBCoeffs_420;
__constant float *coeffs = c_YUV2RGBCoeffs_420;
Y1 = max(0.f, Y1 - d1) * coeffs[0];
Y2 = max(0.f, Y2 - d1) * coeffs[0];
Y3 = max(0.f, Y3 - d1) * coeffs[0];
Y4 = max(0.f, Y4 - d1) * coeffs[0];
Y = max(0.f, Y - d1) * coeffs[0];
float ruv = fma(coeffs[4], UV.y, 0.0f);
float guv = fma(coeffs[3], UV.y, fma(coeffs[2], UV.x, 0.0f));
float buv = fma(coeffs[1], UV.x, 0.0f);
float R1 = (Y1.x + ruv) * CV_8U_MAX;
float G1 = (Y1.x + guv) * CV_8U_MAX;
float B1 = (Y1.x + buv) * CV_8U_MAX;
float4 R = (Y + ruv) * CV_8U_MAX;
float4 G = (Y + guv) * CV_8U_MAX;
float4 B = (Y + buv) * CV_8U_MAX;
float R2 = (Y2.x + ruv) * CV_8U_MAX;
float G2 = (Y2.x + guv) * CV_8U_MAX;
float B2 = (Y2.x + buv) * CV_8U_MAX;
pDstRow1[0*NCHANNELS + 0] = convert_uchar_sat(B.x);
pDstRow1[0*NCHANNELS + 1] = convert_uchar_sat(G.x);
pDstRow1[0*NCHANNELS + 2] = convert_uchar_sat(R.x);
float R3 = (Y3.x + ruv) * CV_8U_MAX;
float G3 = (Y3.x + guv) * CV_8U_MAX;
float B3 = (Y3.x + buv) * CV_8U_MAX;
pDstRow1[1*NCHANNELS + 0] = convert_uchar_sat(B.y);
pDstRow1[1*NCHANNELS + 1] = convert_uchar_sat(G.y);
pDstRow1[1*NCHANNELS + 2] = convert_uchar_sat(R.y);
float R4 = (Y4.x + ruv) * CV_8U_MAX;
float G4 = (Y4.x + guv) * CV_8U_MAX;
float B4 = (Y4.x + buv) * CV_8U_MAX;
pDstRow2[0*NCHANNELS + 0] = convert_uchar_sat(B.z);
pDstRow2[0*NCHANNELS + 1] = convert_uchar_sat(G.z);
pDstRow2[0*NCHANNELS + 2] = convert_uchar_sat(R.z);
pDstRow1[0*NCHANNELS + 0] = convert_uchar_sat(B1);
pDstRow1[0*NCHANNELS + 1] = convert_uchar_sat(G1);
pDstRow1[0*NCHANNELS + 2] = convert_uchar_sat(R1);
pDstRow1[1*NCHANNELS + 0] = convert_uchar_sat(B2);
pDstRow1[1*NCHANNELS + 1] = convert_uchar_sat(G2);
pDstRow1[1*NCHANNELS + 2] = convert_uchar_sat(R2);
pDstRow2[0*NCHANNELS + 0] = convert_uchar_sat(B3);
pDstRow2[0*NCHANNELS + 1] = convert_uchar_sat(G3);
pDstRow2[0*NCHANNELS + 2] = convert_uchar_sat(R3);
pDstRow2[1*NCHANNELS + 0] = convert_uchar_sat(B4);
pDstRow2[1*NCHANNELS + 1] = convert_uchar_sat(G4);
pDstRow2[1*NCHANNELS + 2] = convert_uchar_sat(R4);
pDstRow2[1*NCHANNELS + 0] = convert_uchar_sat(B.w);
pDstRow2[1*NCHANNELS + 1] = convert_uchar_sat(G.w);
pDstRow2[1*NCHANNELS + 2] = convert_uchar_sat(R.w);
}
}
}
@@ -172,12 +159,15 @@ void BGR2YUV_NV12_8u(
{
int x = get_global_id(0);
int y = get_global_id(1);
// each iteration computes 2*2=4 pixels
int x2 = x*2;
int y2 = y*2;
if (x < cols)
if (x2 + 1 < cols)
{
if (y < rows)
if (y2 + 1 < rows)
{
__global const uchar* pSrcRow1 = pBGR + mad24(y, bgrStep, mad24(x, NCHANNELS, 0));
__global const uchar* pSrcRow1 = pBGR + mad24(y2, bgrStep, mad24(x2, NCHANNELS, 0));
__global const uchar* pSrcRow2 = pSrcRow1 + bgrStep;
float4 src_pix1 = convert_float4(vload4(0, pSrcRow1 + 0*NCHANNELS)) * CV_8U_SCALE;
@@ -196,12 +186,12 @@ void BGR2YUV_NV12_8u(
UV.x = fma(coeffs[3], src_pix1.z, fma(coeffs[4], src_pix1.y, fma(coeffs[5], src_pix1.x, d2)));
UV.y = fma(coeffs[5], src_pix1.z, fma(coeffs[6], src_pix1.y, fma(coeffs[7], src_pix1.x, d2)));
write_imagef(imgY, (int2)(x+0, y+0), Y1);
write_imagef(imgY, (int2)(x+1, y+0), Y2);
write_imagef(imgY, (int2)(x+0, y+1), Y3);
write_imagef(imgY, (int2)(x+1, y+1), Y4);
write_imagef(imgY, (int2)(x2+0, y2+0), Y1);
write_imagef(imgY, (int2)(x2+1, y2+0), Y2);
write_imagef(imgY, (int2)(x2+0, y2+1), Y3);
write_imagef(imgY, (int2)(x2+1, y2+1), Y4);
write_imagef(imgUV, (int2)((x/2), (y/2)), UV);
write_imagef(imgUV, (int2)(x, y), UV);
}
}
}
modules/core/src/va_intel.cpp
+2 -2
View File
@@ -189,7 +189,7 @@ static bool ocl_convert_nv12_to_bgr(cl_mem clImageY, cl_mem clImageUV, cl_mem cl
k.args(clImageY, clImageUV, clBuffer, step, cols, rows);
size_t globalsize[] = { (size_t)cols, (size_t)rows };
size_t globalsize[] = { (size_t)cols/2, (size_t)rows/2 };
return k.run(2, globalsize, 0, false);
}
@@ -202,7 +202,7 @@ static bool ocl_convert_bgr_to_nv12(cl_mem clBuffer, int step, int cols, int row
k.args(clBuffer, step, cols, rows, clImageY, clImageUV);
size_t globalsize[] = { (size_t)cols, (size_t)rows };
size_t globalsize[] = { (size_t)cols/2, (size_t)rows/2 };
return k.run(2, globalsize, 0, false);
}
#endif // HAVE_VA_INTEL
modules/dnn/include/opencv2/dnn/all_layers.hpp
+1
View File
@@ -364,6 +364,7 @@ CV__DNN_INLINE_NS_BEGIN
* Inner vector has slice ranges for the first number of input dimensions.
*/
std::vector<std::vector<Range> > sliceRanges;
std::vector<std::vector<int> > sliceSteps;
int axis;
int num_split;
modules/dnn/src/layers/slice_layer.cpp
+82 -6
View File
@@ -70,6 +70,7 @@ public:
SliceLayerImpl(const LayerParams& params)
{
setParamsFrom(params);
hasSteps = false;
axis = params.get<int>("axis", 1);
num_split = params.get<int>("num_split", 0);
hasDynamicShapes = params.get<bool>("has_dynamic_shapes", false);
@@ -118,6 +119,22 @@ public:
sliceRanges[0][i].end = end; // We'll finalize a negative value later.
}
}
if (params.has("steps"))
{
const DictValue &steps = params.get("steps");
sliceSteps.resize(1);
sliceSteps[0].resize(steps.size());
for (int i = 0; i < steps.size(); ++i)
{
int step = steps.get<int>(i);
CV_Assert(step >= 1);
if (step > 1)
hasSteps = true;
sliceSteps[0][i] = step;
}
}
}
}
@@ -126,14 +143,17 @@ public:
#ifdef HAVE_DNN_IE_NN_BUILDER_2019
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NN_BUILDER_2019)
return INF_ENGINE_VER_MAJOR_GE(INF_ENGINE_RELEASE_2019R1) &&
sliceRanges.size() == 1 && sliceRanges[0].size() == 4;
sliceRanges.size() == 1 && sliceRanges[0].size() == 4 && !hasSteps;
#endif
#ifdef HAVE_DNN_NGRAPH
if (backendId == DNN_BACKEND_INFERENCE_ENGINE_NGRAPH)
return sliceRanges.size() == 1;
return sliceRanges.size() == 1 && !hasSteps;
#endif
return backendId == DNN_BACKEND_OPENCV ||
backendId == DNN_BACKEND_CUDA;
#ifdef HAVE_CUDA
if (backendId == DNN_BACKEND_CUDA)
return !hasSteps;
#endif
return backendId == DNN_BACKEND_OPENCV;
}
bool getMemoryShapes(const std::vector<MatShape> &inputs,
@@ -154,6 +174,9 @@ public:
{
if (shapesInitialized || inpShape[j] > 0)
outputs[i][j] = normalize_axis_range(sliceRanges[i][j], inpShape[j]).size();
if (!sliceSteps.empty() && (i < sliceSteps.size()) && (j < sliceSteps[i].size()) && (sliceSteps[i][j] > 1))
outputs[i][j] = (outputs[i][j] + sliceSteps[i][j] - 1) / sliceSteps[i][j];
}
}
}
@@ -188,6 +211,7 @@ public:
const MatSize& inpShape = inputs[0].size;
finalSliceRanges = sliceRanges;
if (sliceRanges.empty())
{
// Divide input blob on equal parts by axis.
@@ -220,6 +244,9 @@ public:
}
}
if (!sliceSteps.empty() && sliceSteps[0].size() != inputs[0].dims)
sliceSteps[0].resize(inputs[0].dims, 1);
#if 0
std::cout << "DEBUG: DNN/Slice: " << outputs.size() << " inpShape=" << inpShape << std::endl;
for (int i = 0; i < outputs.size(); ++i)
@@ -427,6 +454,9 @@ public:
{
CV_TRACE_FUNCTION();
if (hasSteps)
return false; // TODO not implemented yet: https://github.com/opencv/opencv/pull/19546
std::vector<UMat> inputs;
std::vector<UMat> outputs;
@@ -485,9 +515,24 @@ public:
const Mat& inpMat = inputs[0];
CV_Assert(outputs.size() == finalSliceRanges.size());
for (size_t i = 0; i < outputs.size(); i++)
if (!hasSteps)
{
inpMat(finalSliceRanges[i]).copyTo(outputs[i]);
for (size_t i = 0; i < outputs.size(); i++)
{
inpMat(finalSliceRanges[i]).copyTo(outputs[i]);
}
}
else
{
int dimsNum = inpMat.dims;
for (size_t i = 0; i < outputs.size(); i++)
{
std::vector<int> inpIdx(dimsNum, 0);
std::vector<int> outIdx(dimsNum, 0);
getSliceRecursive(inpMat, inpIdx, finalSliceRanges[i], sliceSteps[i], 0, dimsNum, outputs[i], outIdx);
}
}
}
@@ -603,11 +648,42 @@ public:
#endif
private:
void getSliceRecursive(const Mat &inpMat, std::vector<int> &inpIdx,
const std::vector<Range> &sliceRanges,
const std::vector<int> &sliceSteps, int dim, int dimsNum,
Mat &outputs, std::vector<int> &outIdx)
{
int begin = sliceRanges[dim].start;
int end = sliceRanges[dim].end;
int step = !sliceSteps.empty() ? sliceSteps[dim] : 1;
const bool is32F = inpMat.depth() == CV_32F;
// TODO optimization is required (for 2D tail case at least)
for (int k = begin, j = 0; k < end; k += step, j++)
{
inpIdx[dim] = k;
outIdx[dim] = j;
if (dim + 1 < dimsNum)
getSliceRecursive(inpMat, inpIdx, sliceRanges, sliceSteps, dim + 1, dimsNum, outputs, outIdx);
else
{
if (is32F)
outputs.at<float>(outIdx.data()) = inpMat.at<float>(inpIdx.data());
else
outputs.at<short>(outIdx.data()) = inpMat.at<short>(inpIdx.data()); // 16F emulation
}
}
}
protected:
// The actual non-negative values determined from @p sliceRanges depends on input size.
std::vector<std::vector<Range> > finalSliceRanges;
bool hasDynamicShapes;
bool shapesInitialized;
bool hasSteps;
};
class CropLayerImpl CV_FINAL : public SliceLayerImpl
modules/dnn/src/ocl4dnn/src/math_functions.cpp
+2 -2
View File
@@ -112,14 +112,14 @@ ocl::Image2D ocl4dnnGEMMCopyBufferToImage(UMat buffer, int offset,
global_copy[0] = padded_width;
global_copy[1] = padded_height;
oclk_gemm_copy
bool res = oclk_gemm_copy
.args(
ocl::KernelArg::PtrReadOnly(buffer),
image, offset,
width, height,
ld)
.run(2, global_copy, NULL, false);
oclk_gemm_copy.run(2, global_copy, NULL, false);
CV_Assert(res);
}
}
modules/dnn/src/onnx/onnx_importer.cpp
+11 -12
View File
@@ -641,20 +641,11 @@ void ONNXImporter::handleNode(const opencv_onnx::NodeProto& node_proto_)
int axis = 0;
std::vector<int> begin;
std::vector<int> end;
std::vector<int> steps;
int inp_size = node_proto.input_size();
if (inp_size == 1)
{
if (layerParams.has("steps"))
{
DictValue steps = layerParams.get("steps");
for (int i = 0; i < steps.size(); ++i)
{
if (steps.get<int>(i) != 1)
CV_Error(Error::StsNotImplemented,
"Slice layer only supports steps = 1");
}
}
if (layerParams.has("axes")) {
DictValue axes = layerParams.get("axes");
for (int i = 1; i < axes.size(); ++i) {
@@ -677,7 +668,7 @@ void ONNXImporter::handleNode(const opencv_onnx::NodeProto& node_proto_)
int finish = ends.get<int>(i);
end.push_back((finish < 0) ? --finish : finish); // numpy doesn't include last dim
}
} else {
} else { // inp_size > 1
CV_Assert(inp_size >= 3);
for (int i = 1; i < inp_size; i++) {
CV_Assert(constBlobs.find(node_proto.input(i)) != constBlobs.end());
@@ -711,6 +702,12 @@ void ONNXImporter::handleNode(const opencv_onnx::NodeProto& node_proto_)
if (inp_size == 5) {
CV_Assert(constBlobs.find(node_proto.input(4)) != constBlobs.end());
Mat step_blob = getBlob(node_proto, 4);
const int* steps_ptr = step_blob.ptr<int>();
if (axis > 0)
steps.resize(axis, 1);
std::copy(steps_ptr, steps_ptr + step_blob.total(), std::back_inserter(steps));
// Very strange application for Slice op with tensor reversing.
// We just workaround it for 2d constants.
@@ -728,13 +725,15 @@ void ONNXImporter::handleNode(const opencv_onnx::NodeProto& node_proto_)
return;
}
}
CV_CheckEQ(countNonZero(step_blob != 1), 0, "Slice layer only supports steps = 1");
}
}
layerParams.set("begin", DictValue::arrayInt(&begin[0], begin.size()));
layerParams.set("end", DictValue::arrayInt(&end[0], end.size()));
layerParams.set("axis", axis);
if (!steps.empty())
layerParams.set("steps", DictValue::arrayInt(&steps[0], steps.size()));
if (constBlobs.find(node_proto.input(0)) != constBlobs.end())
{
Mat inp = getBlob(node_proto, 0);
modules/dnn/test/test_halide_layers.cpp
+11 -1
View File
@@ -258,7 +258,17 @@ TEST_P(LRN, Accuracy)
int sz[] = {1, inChannels, inSize.height, inSize.width};
Mat input(4, &sz[0], CV_32F);
test(lp, input, backendId, targetId);
double l1 = 0.0, lInf = 0.0;
// The OpenCL kernels use the native_ math functions which have
// implementation defined accuracy, so we use relaxed thresholds. See
// https://github.com/opencv/opencv/issues/9821 for more details.
if (targetId == DNN_TARGET_OPENCL)
{
l1 = 0.01;
lInf = 0.01;
}
test(lp, input, backendId, targetId, false, l1, lInf);
}
INSTANTIATE_TEST_CASE_P(Layer_Test_Halide, LRN, Combine(
modules/dnn/test/test_layers.cpp
+11 -2
View File
@@ -169,8 +169,17 @@ TEST_P(Test_Caffe_layers, Softmax)
TEST_P(Test_Caffe_layers, LRN)
{
testLayerUsingCaffeModels("layer_lrn_spatial");
testLayerUsingCaffeModels("layer_lrn_channels");
double l1 = 0.0, lInf = 0.0;
// The OpenCL kernels use the native_ math functions which have
// implementation defined accuracy, so we use relaxed thresholds. See
// https://github.com/opencv/opencv/issues/9821 for more details.
if (target == DNN_TARGET_OPENCL)
{
l1 = 0.01;
lInf = 0.01;
}
testLayerUsingCaffeModels("layer_lrn_spatial", false, true, l1, lInf);
testLayerUsingCaffeModels("layer_lrn_channels", false, true, l1, lInf);
}
TEST_P(Test_Caffe_layers, Convolution)
modules/dnn/test/test_onnx_importer.cpp
+20
View File
@@ -664,6 +664,26 @@ TEST_P(Test_ONNX_layers, Slice)
#endif
}
TEST_P(Test_ONNX_layers, Slice_Steps_2DInput)
{
testONNXModels("slice_opset_11_steps_2d");
}
TEST_P(Test_ONNX_layers, Slice_Steps_3DInput)
{
testONNXModels("slice_opset_11_steps_3d");
}
TEST_P(Test_ONNX_layers, Slice_Steps_4DInput)
{
testONNXModels("slice_opset_11_steps_4d");
}
TEST_P(Test_ONNX_layers, Slice_Steps_5DInput)
{
testONNXModels("slice_opset_11_steps_5d");
}
TEST_P(Test_ONNX_layers, Softmax)
{
testONNXModels("softmax");
modules/dnn/test/test_torch_importer.cpp
+8
View File
@@ -258,6 +258,14 @@ TEST_P(Test_Torch_layers, net_conv_gemm_lrn)
l1 = 0.0042;
lInf = 0.021;
}
// The OpenCL kernels use the native_ math functions which have
// implementation defined accuracy, so we use relaxed thresholds. See
// https://github.com/opencv/opencv/issues/9821 for more details.
else if (target == DNN_TARGET_OPENCL)
{
l1 = 0.02;
lInf = 0.02;
}
runTorchNet("net_conv_gemm_lrn", "", false, true, true, l1, lInf);
}