Skip to content

Commit 7b7033a

Browse files
pli2-intelpli2-intel
committed
permute layer ocl implementation
Signed-off-by: Li Peng <peng.li@intel.com>
1 parent a3ec2ac commit 7b7033a

File tree

2 files changed

+126
-0
lines changed

2 files changed

+126
-0
lines changed

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

Lines changed: 59 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -44,6 +44,7 @@
4444
#include "layers_common.hpp"
4545
#include <float.h>
4646
#include <algorithm>
47+
#include "opencl_kernels_dnn.hpp"
4748

4849
namespace cv
4950
{
@@ -173,6 +174,24 @@ class PermuteLayerImpl : public PermuteLayer
173174
CV_Assert((int)_numAxes == inp0.dims);
174175

175176
computeStrides(shape(*inputs[0]), shape(outputs[0]));
177+
178+
#ifdef HAVE_OPENCL
179+
if (uorder.empty())
180+
{
181+
std::vector<int> orderVec(_order.begin(), _order.end());;
182+
Mat morder(1, orderVec.size(), CV_32SC1, &orderVec[0]);
183+
184+
std::vector<int> oldStrideVec(_oldStride.begin(), _oldStride.end());
185+
Mat mold_stride(1, _oldStride.size(), CV_32SC1, &oldStrideVec[0]);
186+
187+
std::vector<int> newStrideVec(_newStride.begin(), _newStride.end());
188+
Mat mnew_stride(1, newStrideVec.size(), CV_32SC1, &newStrideVec[0]);
189+
190+
morder.copyTo(uorder);
191+
mold_stride.copyTo(uold_stride);
192+
mnew_stride.copyTo(unew_stride);
193+
}
194+
#endif
176195
}
177196

178197
class PermuteInvoker : public ParallelLoopBody
@@ -247,11 +266,47 @@ class PermuteLayerImpl : public PermuteLayer
247266
}
248267
};
249268

269+
#ifdef HAVE_OPENCL
270+
bool forward_ocl(InputArrayOfArrays inps, OutputArrayOfArrays outs, OutputArrayOfArrays internals)
271+
{
272+
std::vector<UMat> inputs;
273+
std::vector<UMat> outputs;
274+
275+
inps.getUMatVector(inputs);
276+
outs.getUMatVector(outputs);
277+
278+
if (!_needsPermute)
279+
return false;
280+
281+
for (size_t i = 0; i < inputs.size(); i++)
282+
{
283+
ocl::Kernel kernel("permute", ocl::dnn::permute_oclsrc);
284+
285+
kernel.set(0, (int)_count);
286+
kernel.set(1, ocl::KernelArg::PtrReadOnly(inputs[i]));
287+
kernel.set(2, ocl::KernelArg::PtrReadOnly(uorder));
288+
kernel.set(3, ocl::KernelArg::PtrReadOnly(uold_stride));
289+
kernel.set(4, ocl::KernelArg::PtrReadOnly(unew_stride));
290+
kernel.set(5, (int)_numAxes);
291+
kernel.set(6, ocl::KernelArg::PtrWriteOnly(outputs[i]));
292+
293+
if (!kernel.run(1, &_count, NULL, false))
294+
return false;
295+
}
296+
297+
return true;
298+
}
299+
#endif
300+
250301
void forward(InputArrayOfArrays inputs_arr, OutputArrayOfArrays outputs_arr, OutputArrayOfArrays internals_arr)
251302
{
252303
CV_TRACE_FUNCTION();
253304
CV_TRACE_ARG_VALUE(name, "name", name.c_str());
254305

306+
CV_OCL_RUN((preferableTarget == DNN_TARGET_OPENCL) &&
307+
OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel()),
308+
forward_ocl(inputs_arr, outputs_arr, internals_arr))
309+
255310
Layer::forward_fallback(inputs_arr, outputs_arr, internals_arr);
256311
}
257312

@@ -325,6 +380,10 @@ class PermuteLayerImpl : public PermuteLayer
325380
std::vector<size_t> _newStride;
326381
bool _needsPermute;
327382

383+
#ifdef HAVE_OPENCL
384+
UMat uorder, uold_stride, unew_stride;
385+
#endif
386+
328387
size_t _numAxes;
329388
};
330389

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

Lines changed: 67 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,67 @@
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) 2017, Intel Corporation, all rights reserved.
14+
// Copyright (c) 2016-2017 Fabian David Tschopp, all rights reserved.
15+
// Third party copyrights are property of their respective owners.
16+
//
17+
// Redistribution and use in source and binary forms, with or without modification,
18+
// are permitted provided that the following conditions are met:
19+
//
20+
// * Redistribution's of source code must retain the above copyright notice,
21+
// this list of conditions and the following disclaimer.
22+
//
23+
// * Redistribution's in binary form must reproduce the above copyright notice,
24+
// this list of conditions and the following disclaimer in the documentation
25+
// and/or other materials provided with the distribution.
26+
//
27+
// * The name of the copyright holders may not be used to endorse or promote products
28+
// derived from this software without specific prior written permission.
29+
//
30+
// This software is provided by the copyright holders and contributors "as is" and
31+
// any express or implied warranties, including, but not limited to, the implied
32+
// warranties of merchantability and fitness for a particular purpose are disclaimed.
33+
// In no event shall the Intel Corporation or contributors be liable for any direct,
34+
// indirect, incidental, special, exemplary, or consequential damages
35+
// (including, but not limited to, procurement of substitute goods or services;
36+
// loss of use, data, or profits; or business interruption) however caused
37+
// and on any theory of liability, whether in contract, strict liability,
38+
// or tort (including negligence or otherwise) arising in any way out of
39+
// the use of this software, even if advised of the possibility of such damage.
40+
//
41+
//M*/
42+
43+
#define Dtype float
44+
45+
__kernel void permute(const int nthreads,
46+
__global Dtype* bottom_data,
47+
global int* permute_order,
48+
global int* oldStride,
49+
global int* newStride,
50+
const int num_axes,
51+
__global Dtype* top_data)
52+
{
53+
for (int i = get_global_id(0); i < nthreads; i += get_global_size(0))
54+
{
55+
int oldPosition = 0;
56+
int newPosition = i;
57+
58+
for (int j = 0; j < num_axes; ++j)
59+
{
60+
int order = permute_order[j];
61+
oldPosition += (newPosition / newStride[j]) * oldStride[order];
62+
newPosition %= newStride[j];
63+
}
64+
65+
top_data[i] = bottom_data[oldPosition];
66+
}
67+
}

0 commit comments

Comments
 (0)