Skip to content

Commit 7707c9b

Browse files
pli2-intelpli2-intel
committed
reorg layer ocl implementation
Signed-off-by: Li Peng <peng.li@intel.com>
1 parent 85b1c40 commit 7707c9b

File tree

2 files changed

+107
-0
lines changed

2 files changed

+107
-0
lines changed

‎modules/dnn/src/layers/reorg_layer.cpp

Lines changed: 44 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -44,6 +44,7 @@
4444
#include <opencv2/dnn/shape_utils.hpp>
4545
#include <opencv2/dnn/all_layers.hpp>
4646
#include <iostream>
47+
#include "opencl_kernels_dnn.hpp"
4748

4849
namespace cv
4950
{
@@ -86,11 +87,54 @@ class ReorgLayerImpl : public ReorgLayer
8687
return backendId == DNN_BACKEND_DEFAULT;
8788
}
8889

90+
#ifdef HAVE_OPENCL
91+
bool forward_ocl(InputArrayOfArrays inps, OutputArrayOfArrays outs, OutputArrayOfArrays internals)
92+
{
93+
std::vector<UMat> inputs;
94+
std::vector<UMat> outputs;
95+
96+
inps.getUMatVector(inputs);
97+
outs.getUMatVector(outputs);
98+
String buildopt = String("-DDtype=") + ocl::typeToStr(inputs[0].type()) + String(" ");
99+
100+
for (size_t i = 0; i < inputs.size(); i++)
101+
{
102+
ocl::Kernel kernel("reorg", ocl::dnn::reorg_oclsrc, buildopt);
103+
if (kernel.empty())
104+
return false;
105+
106+
UMat& srcBlob = inputs[i];
107+
UMat& dstBlob = outputs[0];
108+
int channels = srcBlob.size[1];
109+
int height = srcBlob.size[2];
110+
int width = srcBlob.size[3];
111+
size_t nthreads = channels * height * width;
112+
113+
kernel.set(0, (int)nthreads);
114+
kernel.set(1, ocl::KernelArg::PtrReadOnly(srcBlob));
115+
kernel.set(2, (int)channels);
116+
kernel.set(3, (int)height);
117+
kernel.set(4, (int)width);
118+
kernel.set(5, (int)reorgStride);
119+
kernel.set(6, ocl::KernelArg::PtrWriteOnly(dstBlob));
120+
121+
if (!kernel.run(1, &nthreads, NULL, false))
122+
return false;
123+
}
124+
125+
return true;
126+
}
127+
#endif
128+
89129
void forward(InputArrayOfArrays inputs_arr, OutputArrayOfArrays outputs_arr, OutputArrayOfArrays internals_arr)
90130
{
91131
CV_TRACE_FUNCTION();
92132
CV_TRACE_ARG_VALUE(name, "name", name.c_str());
93133

134+
CV_OCL_RUN((preferableTarget == DNN_TARGET_OPENCL) &&
135+
OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()),
136+
forward_ocl(inputs_arr, outputs_arr, internals_arr))
137+
94138
Layer::forward_fallback(inputs_arr, outputs_arr, internals_arr);
95139
}
96140

‎modules/dnn/src/opencl/reorg.cl

Lines changed: 63 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,63 @@
1+
/*M///////////////////////////////////////////////////////////////////////////////////////
2+
//
3+
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
4+
//
5+
// By downloading, copying, installing or using the software you agree to this license.
6+
// If you do not agree to this license, do not download, install,
7+
// copy or use the software.
8+
//
9+
//
10+
// License Agreement
11+
// For Open Source Computer Vision Library
12+
//
13+
// Copyright (c) 2016-2017 Fabian David Tschopp, all rights reserved.
14+
// Third party copyrights are property of their respective owners.
15+
//
16+
// Redistribution and use in source and binary forms, with or without modification,
17+
// are permitted provided that the following conditions are met:
18+
//
19+
// * Redistribution's of source code must retain the above copyright notice,
20+
// this list of conditions and the following disclaimer.
21+
//
22+
// * Redistribution's in binary form must reproduce the above copyright notice,
23+
// this list of conditions and the following disclaimer in the documentation
24+
// and/or other materials provided with the distribution.
25+
//
26+
// * The name of the copyright holders may not be used to endorse or promote products
27+
// derived from this software without specific prior written permission.
28+
//
29+
// This software is provided by the copyright holders and contributors "as is" and
30+
// any express or implied warranties, including, but not limited to, the implied
31+
// warranties of merchantability and fitness for a particular purpose are disclaimed.
32+
// In no event shall the Intel Corporation or contributors be liable for any direct,
33+
// indirect, incidental, special, exemplary, or consequential damages
34+
// (including, but not limited to, procurement of substitute goods or services;
35+
// loss of use, data, or profits; or business interruption) however caused
36+
// and on any theory of liability, whether in contract, strict liability,
37+
// or tort (including negligence or otherwise) arising in any way out of
38+
// the use of this software, even if advised of the possibility of such damage.
39+
//
40+
//M*/
41+
42+
__kernel void reorg(const int count,
43+
__global const Dtype* src,
44+
const int channels,
45+
const int height,
46+
const int width,
47+
const int reorgStride,
48+
__global Dtype* dst)
49+
{
50+
for (int index = get_global_id(0); index < count; index += get_global_size(0))
51+
{
52+
int k = index / (height * width);
53+
int j = (index - (k * height * width)) / width;
54+
int i = (index - (k * height * width)) % width;
55+
int out_c = channels / (reorgStride*reorgStride);
56+
int c2 = k % out_c;
57+
int offset = k / out_c;
58+
int w2 = i*reorgStride + offset % reorgStride;
59+
int h2 = j*reorgStride + offset / reorgStride;
60+
int in_index = w2 + width*reorgStride*(h2 + height*reorgStride*c2);
61+
dst[index] = src[in_index];
62+
}
63+
}

0 commit comments

Comments
 (0)