Chiebot-Mirror
/
opencv_contrib
8
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Adding of OCL & public interface for Pooling layer

Browse Source
pull/707/head
Vitaliy Lyudvichenko 9 years ago
parent 601afeed90
commit 0af7f9ea29
8 changed files with 380 additions and 150 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 21
      modules/dnn/include/opencv2/dnn/all_layers.hpp
  2. 2
      modules/dnn/src/init.cpp
  3. 35
      modules/dnn/src/layers/lrn_layer.cpp
  4. 2
      modules/dnn/src/layers/lrn_layer.hpp
  5. 325
      modules/dnn/src/layers/pooling_layer.cpp
  6. 35
      modules/dnn/src/layers/pooling_layer.hpp
  7. 94
      modules/dnn/src/opencl/pooling.cl
  8. 16
      modules/dnn/test/test_layers.cpp

21
modules/dnn/include/opencv2/dnn/all_layers.hpp
Unescape View File

@ -209,7 +209,7 @@ namespace dnn
{ {
public: public:
Size kernel, pad, stride; Size kernel, stride, pad;
}; };
class CV_EXPORTS_W ConvolutionLayer : public BaseConvolutionLayer class CV_EXPORTS_W ConvolutionLayer : public BaseConvolutionLayer
@ -232,7 +232,7 @@ namespace dnn
{ {
public: public:
enum enum Type
{ {
CHANNEL_NRM, CHANNEL_NRM,
SPATIAL_NRM SPATIAL_NRM
@ -241,9 +241,26 @@ namespace dnn
int size; int size;
double alpha, beta; double alpha, beta;
static Ptr<LRNLayer> create(int type = CHANNEL_NRM, int size = 5, double alpha = 1, double beta = 0.75);
}; };
class CV_EXPORTS_W PoolingLayer : public Layer
{
public:
enum Type
{
MAX,
AVE,
STOCHASTIC
};
int type;
Size kernel, stride, pad;
static Ptr<PoolingLayer> create(int type = MAX, Size kernel = Size(2, 2), Size pad = Size(0, 0), Size stride = Size(1, 1));
};
//! @} //! @}
//! @} //! @}

2
modules/dnn/src/init.cpp
Unescape View File

@ -81,7 +81,7 @@ void initModule()
REG_RUNTIME_LAYER_CLASS(Split, SplitLayer) REG_RUNTIME_LAYER_CLASS(Split, SplitLayer)
REG_RUNTIME_LAYER_CLASS(Reshape, ReshapeLayer) REG_RUNTIME_LAYER_CLASS(Reshape, ReshapeLayer)
REG_STATIC_LAYER_FUNC(Flatten, createFlattenLayer) REG_STATIC_LAYER_FUNC(Flatten, createFlattenLayer)
REG_RUNTIME_LAYER_CLASS(Pooling, PoolingLayerImpl) REG_RUNTIME_LAYER_FUNC(Pooling, createPoolingLayerFromCaffe)
REG_RUNTIME_LAYER_CLASS(MVN, MVNLayer) REG_RUNTIME_LAYER_CLASS(MVN, MVNLayer)
REG_RUNTIME_LAYER_FUNC(LRN, createLRNLayerFromCaffe) REG_RUNTIME_LAYER_FUNC(LRN, createLRNLayerFromCaffe)
REG_RUNTIME_LAYER_CLASS(InnerProduct, FullyConnectedLayer) REG_RUNTIME_LAYER_CLASS(InnerProduct, FullyConnectedLayer)

35
modules/dnn/src/layers/lrn_layer.cpp
Unescape View File

@ -53,17 +53,18 @@ namespace cv
namespace dnn namespace dnn
{ {
LRNLayerImpl::LRNLayerImpl() LRNLayerImpl::LRNLayerImpl(int type_, int size_, double alpha_, double beta_)
{ {
size = 5; type = type_;
alpha = 1; size = size_;
beta = 0.75; alpha = alpha_;
type = CHANNEL_NRM; beta = beta_;
} }
void LRNLayerImpl::allocate(const std::vector<Blob*> &inputs, std::vector<Blob> &outputs) void LRNLayerImpl::allocate(const std::vector<Blob*> &inputs, std::vector<Blob> &outputs)
{ {
CV_Assert(inputs.size() == 1 && inputs[0]->dims() == 4); CV_Assert(inputs.size() == 1 && inputs[0]->dims() == 4);
CV_Assert(type == CHANNEL_NRM || type == SPATIAL_NRM);
useOpenCL = cv::ocl::useOpenCL(); useOpenCL = cv::ocl::useOpenCL();
if (type == SPATIAL_NRM && !useOpenCL) if (type == SPATIAL_NRM && !useOpenCL)
@ -154,6 +155,7 @@ void LRNLayerImpl::channelNoramlization_(Blob &srcBlob, Blob &dstBlob)
bool LRNLayerImpl::channelNoramlization_ocl(const UMat &src, UMat &dst) bool LRNLayerImpl::channelNoramlization_ocl(const UMat &src, UMat &dst)
{ {
#ifdef HAVE_OPENCL
if (src.offset != 0 || dst.offset != 0) //TODO: add offset if (src.offset != 0 || dst.offset != 0) //TODO: add offset
return false; return false;
@ -187,6 +189,9 @@ bool LRNLayerImpl::channelNoramlization_ocl(const UMat &src, UMat &dst)
return false; return false;
return true; return true;
#else
return false;
#endif
} }
void LRNLayerImpl::spatialNormalization(Blob &src, Blob &dst) void LRNLayerImpl::spatialNormalization(Blob &src, Blob &dst)
@ -232,27 +237,31 @@ void LRNLayerImpl::spatialNormalization_(Blob &srcBlob, Blob &dstBlob)
} }
} }
Ptr<Layer> createLRNLayerFromCaffe(LayerParams &params)
Ptr<LRNLayer> LRNLayer::create(int type, int size, double alpha, double beta)
{ {
LRNLayerImpl *l = new LRNLayerImpl(); return Ptr<LRNLayer>(new LRNLayerImpl(type, size, alpha, beta));
}
Ptr<Layer> createLRNLayerFromCaffe(LayerParams &params)
{
int type;
String nrmType = params.get<String>("norm_region", "ACROSS_CHANNELS"); String nrmType = params.get<String>("norm_region", "ACROSS_CHANNELS");
if (nrmType == "ACROSS_CHANNELS") if (nrmType == "ACROSS_CHANNELS")
l->type = LRNLayer::CHANNEL_NRM; type = LRNLayer::CHANNEL_NRM;
else if (nrmType == "WITHIN_CHANNEL") else if (nrmType == "WITHIN_CHANNEL")
l->type = LRNLayer::SPATIAL_NRM; type = LRNLayer::SPATIAL_NRM;
else else
CV_Error(Error::StsBadArg, "Unknown region type \"" + nrmType + "\""); CV_Error(Error::StsBadArg, "Unknown region type \"" + nrmType + "\"");
int size = params.get<int>("local_size", 5); int size = params.get<int>("local_size", 5);
if (size % 2 != 1 || size <= 0) if (size % 2 != 1 || size <= 0)
CV_Error(Error::StsBadArg, "LRN layer supports only positive odd values for local_size"); CV_Error(Error::StsBadArg, "LRN layer supports only positive odd values for local_size");
l->size = size;
l->alpha = params.get<double>("alpha", 1); double alpha = params.get<double>("alpha", 1);
l->beta = params.get<double>("beta", 0.75); double beta = params.get<double>("beta", 0.75);
return Ptr<Layer>(l); return Ptr<Layer>(new LRNLayerImpl(type, size, alpha, beta));
} }
} }

2
modules/dnn/src/layers/lrn_layer.hpp
Unescape View File

@ -64,7 +64,7 @@ namespace dnn
public: public:
LRNLayerImpl(); LRNLayerImpl(int type = CHANNEL_NRM, int size = 5, double alpha = 1, double beta = 0.75);
void allocate(const std::vector<Blob*> &inputs, std::vector<Blob> &outputs); void allocate(const std::vector<Blob*> &inputs, std::vector<Blob> &outputs);
void forward(std::vector<Blob*> &inputs, std::vector<Blob> &outputs); void forward(std::vector<Blob*> &inputs, std::vector<Blob> &outputs);
}; };

325
modules/dnn/src/layers/pooling_layer.cpp
Unescape View File

@ -42,8 +42,10 @@
#include "../precomp.hpp" #include "../precomp.hpp"
#include "layers_common.hpp" #include "layers_common.hpp"
#include "pooling_layer.hpp" #include "pooling_layer.hpp"
#include "opencl_kernels_dnn.hpp"
#include <float.h> #include <float.h>
#include <algorithm> #include <algorithm>
#include <opencv2/core/ocl.hpp>
using std::max; using std::max;
using std::min; using std::min;
@ -53,154 +55,243 @@ namespace dnn
{ {
//TODO: add ceil_mode param //TODO: add ceil_mode param
PoolingLayer::PoolingLayer(LayerParams &params) : Layer(params) PoolingLayerImpl::PoolingLayerImpl()
{ {
if (params.has("pool"))
{
String pool = params.get<String>("pool").toLowerCase();
if (pool == "max")
type = MAX;
else if (pool == "ave")
type = AVE;
else if (pool == "stochastic")
type = STOCHASTIC;
else
CV_Error(cv::Error::StsBadArg, "Unknown pooling type \"" + pool + "\"");
}
else
{
type = MAX;
}
getCaffeConvParams(params, kernel, pad, stride); }
}
void PoolingLayer::allocate(const std::vector<Blob*> &inputs, std::vector<Blob> &outputs) PoolingLayerImpl::PoolingLayerImpl(int type_, Size kernel_, Size pad_, Size stride_)
{ {
CV_Assert(inputs.size() > 0); type = type_;
kernel = kernel_;
pad = pad_;
stride = stride_;
}
inp = inputs[0]->size2(); void PoolingLayerImpl::allocate(const std::vector<Blob*> &inputs, std::vector<Blob> &outputs)
computeOutputShape(inp); {
CV_Assert(inputs.size() > 0);
outputs.resize(inputs.size()); inp = inputs[0]->size2();
for (size_t i = 0; i < inputs.size(); i++) computeOutputShape(inp);
{
CV_Assert(inputs[i]->rows() == inp.height && inputs[i]->cols() == inp.width); useOpenCL = ocl::useOpenCL();
outputs[i].create(BlobShape(inputs[i]->num(), inputs[i]->channels(), out.height, out.width));
} outputs.resize(inputs.size());
for (size_t i = 0; i < inputs.size(); i++)
{
CV_Assert(inputs[i]->rows() == inp.height && inputs[i]->cols() == inp.width);
outputs[i].create(BlobShape(inputs[i]->num(), inputs[i]->channels(), out.height, out.width));
} }
}
void PoolingLayer::forward(std::vector<Blob*> &inputs, std::vector<Blob> &outputs) void PoolingLayerImpl::forward(std::vector<Blob*> &inputs, std::vector<Blob> &outputs)
{
for (size_t ii = 0; ii < inputs.size(); ii++)
{ {
for (size_t ii = 0; ii < inputs.size(); ii++) switch (type)
{ {
switch (type) case MAX:
{ maxPooling(*inputs[ii], outputs[ii]);
case MAX: break;
maxPooling(*inputs[ii], outputs[ii]); case AVE:
break; avePooling(*inputs[ii], outputs[ii]);
case AVE: break;
avePooling(*inputs[ii], outputs[ii]); default:
break; CV_Error(Error::StsNotImplemented, "Not implemented");
default: break;
CV_Error(Error::StsNotImplemented, "Not implemented");
break;
}
} }
} }
}
void PoolingLayerImpl::maxPooling(Blob &src, Blob &dst)
{
if (!useOpenCL)
maxPooling_cpu(src, dst);
else
{
CV_Assert(maxPooling_ocl(src, dst));
}
}
bool PoolingLayerImpl::maxPooling_ocl(Blob &src, Blob &dst)
{
return pooling_ocl("MaxPoolForward", src, dst);
}
void PoolingLayer::maxPooling(Blob &input, Blob &output) void PoolingLayerImpl::avePooling(Blob &src, Blob &dst)
{
if (!useOpenCL)
avePooling_cpu(src, dst);
else
{ {
CV_DbgAssert(output.rows() == out.height && output.cols() == out.width); CV_Assert(avePooling_ocl(src, dst));
}
}
bool PoolingLayerImpl::avePooling_ocl(Blob &src, Blob &dst)
{
return pooling_ocl("AvePoolForward", src, dst);
}
void PoolingLayerImpl::maxPooling_cpu(Blob &src, Blob &dst)
{
CV_DbgAssert(dst.rows() == out.height && dst.cols() == out.width);
for (int n = 0; n < input.num(); ++n) for (int n = 0; n < src.num(); ++n)
{
for (int c = 0; c < src.channels(); ++c)
{ {
for (int c = 0; c < input.channels(); ++c) const float *srcData = src.ptrf(n, c);
{ float *dstData = dst.ptrf(n, c);
float *srcData = input.ptrf(n, c);
float *dstData = output.ptrf(n, c);
for (int ph = 0; ph < out.height; ++ph) for (int ph = 0; ph < out.height; ++ph)
{
for (int pw = 0; pw < out.width; ++pw)
{ {
for (int pw = 0; pw < out.width; ++pw) int hstart = ph * stride.height - pad.height;
{ int wstart = pw * stride.width - pad.width;
int hstart = ph * stride.height - pad.height; int hend = min(hstart + kernel.height, inp.height);
int wstart = pw * stride.width - pad.width; int wend = min(wstart + kernel.width, inp.width);
int hend = min(hstart + kernel.height, inp.height); hstart = max(hstart, 0);
int wend = min(wstart + kernel.width, inp.width); wstart = max(wstart, 0);
hstart = max(hstart, 0); const int poolIndex = ph * out.width + pw;
wstart = max(wstart, 0); float max_val = -FLT_MAX;
const int poolIndex = ph * out.width + pw;
float max_val = -FLT_MAX; for (int h = hstart; h < hend; ++h)
for (int w = wstart; w < wend; ++w)
for (int h = hstart; h < hend; ++h) {
for (int w = wstart; w < wend; ++w) const int index = h * inp.width + w;
{ if (srcData[index] > max_val)
const int index = h * inp.width + w; max_val = srcData[index];
if (srcData[index] > max_val) }
max_val = srcData[index];
} dstData[poolIndex] = max_val;
dstData[poolIndex] = max_val;
}
} }
} }
} }
} }
}
void PoolingLayer::avePooling(Blob &input, Blob &output) #ifdef HAVE_OPENCL
bool PoolingLayerImpl::pooling_ocl(const char *kname, const Blob &src, Blob &dst, Blob *mask)
{
const UMat &srcMat = src.umatRefConst();
UMat &dstMat = dst.umatRef();
CV_Assert(mask == NULL && srcMat.offset == 0 && dstMat.offset == 0);
ocl::Kernel ker(kname, ocl::dnn::pooling_oclsrc, String("-DT=") + ocl::typeToStr(src.type()));
if (ker.empty())
return false;
BlobShape s = src.shape();
size_t nthreads = dst.total();
ker.args((int)nthreads,
ocl::KernelArg::PtrReadOnly(srcMat), s[0], s[1], s[2], s[3],
out.height, out.width, kernel.height, kernel.width,
stride.height, stride.width, pad.height, pad.width,
ocl::KernelArg::PtrWriteOnly(dstMat));
size_t wgSize = ocl::Device::getDefault().maxWorkGroupSize();
if (!ker.run(1, &nthreads, &wgSize, true))
return false;
return true;
}
#else
bool PoolingLayerImpl::pooling_ocl(const char*, const Blob&, Blob&, Blob*)
{
return false;
}
#endif
void PoolingLayerImpl::avePooling_cpu(Blob &src, Blob &dst)
{
for (int n = 0; n < src.num(); ++n)
{ {
for (int n = 0; n < input.num(); ++n) for (int c = 0; c < src.channels(); ++c)
{ {
for (int c = 0; c < input.channels(); ++c) const float *srcData = src.ptrf(n, c);
{ float *dstData = dst.ptrf(n, c);
float *srcData = input.ptrf(n, c);
float *dstData = output.ptrf(n, c);
for (int ph = 0; ph < out.height; ++ph) for (int ph = 0; ph < out.height; ++ph)
{
for (int pw = 0; pw < out.width; ++pw)
{ {
for (int pw = 0; pw < out.width; ++pw) int hstart = ph * stride.height - pad.height;
{ int wstart = pw * stride.width - pad.width;
int hstart = ph * stride.height - pad.height; int hend = min(hstart + kernel.height, inp.height + pad.height);
int wstart = pw * stride.width - pad.width; int wend = min(wstart + kernel.width, inp.width + pad.width);
int hend = min(hstart + kernel.height, inp.height + pad.height); int poolSize = (hend - hstart) * (wend - wstart);
int wend = min(wstart + kernel.width, inp.width + pad.width); hstart = max(hstart, 0);
int poolSize = (hend - hstart) * (wend - wstart); wstart = max(wstart, 0);
hstart = max(hstart, 0); hend = min(hend, inp.height);
wstart = max(wstart, 0); wend = min(wend, inp.width);
hend = min(hend, inp.height);
wend = min(wend, inp.width); dstData[ph * out.width + pw] = 0.f;
dstData[ph * out.width + pw] = 0.f; for (int h = hstart; h < hend; ++h)
for (int w = wstart; w < wend; ++w)
for (int h = hstart; h < hend; ++h) dstData[ph * out.width + pw] += srcData[h * inp.width + w];
for (int w = wstart; w < wend; ++w)
dstData[ph * out.width + pw] += srcData[h * inp.width + w]; dstData[ph * out.width + pw] /= poolSize;
dstData[ph * out.width + pw] /= poolSize;
}
} }
} }
} }
} }
}
void PoolingLayerImpl::computeOutputShape(Size inpSz)
{
//Yeah, something strange Caffe scheme-)
out.height = static_cast<int>(ceil(static_cast<float>(inpSz.height + 2 * pad.height - kernel.height) / stride.height)) + 1;
out.width = static_cast<int>(ceil(static_cast<float>(inpSz.width + 2 * pad.width - kernel.width) / stride.width)) + 1;
void PoolingLayer::computeOutputShape(Size inpSz) if (pad.height || pad.width)
{ {
//Yeah, something strange Caffe scheme-) // If we have padding, ensure that the last pooling starts strictly
out.height = static_cast<int>(ceil(static_cast<float>(inpSz.height + 2 * pad.height - kernel.height) / stride.height)) + 1; // inside the image (instead of at the padding); otherwise clip the last.
out.width = static_cast<int>(ceil(static_cast<float>(inpSz.width + 2 * pad.width - kernel.width) / stride.width)) + 1; if ((out.height - 1) * stride.height >= inpSz.height + pad.height)
--out.height;
if ((out.width - 1) * stride.width >= inpSz.width + pad.width)
--out.width;
CV_Assert((out.height - 1) * stride.height < inpSz.height + pad.height);
CV_Assert((out.width - 1) * stride.width < inpSz.width + pad.width);
}
}
if (pad.height || pad.width) Ptr<PoolingLayer> PoolingLayer::create(int type, Size kernel, Size pad, Size stride)
{ {
// If we have padding, ensure that the last pooling starts strictly return Ptr<PoolingLayer>(new PoolingLayerImpl(type, kernel, pad, stride));
// inside the image (instead of at the padding); otherwise clip the last. }
if ((out.height - 1) * stride.height >= inpSz.height + pad.height)
--out.height; Ptr<Layer> createPoolingLayerFromCaffe(LayerParams &params)
if ((out.width - 1) * stride.width >= inpSz.width + pad.width) {
--out.width; int type;
CV_Assert((out.height - 1) * stride.height < inpSz.height + pad.height); Size kernel, pad, stride;
CV_Assert((out.width - 1) * stride.width < inpSz.width + pad.width);
} if (params.has("pool"))
{
String pool = params.get<String>("pool").toLowerCase();
if (pool == "max")
type = PoolingLayer::MAX;
else if (pool == "ave")
type = PoolingLayer::AVE;
else if (pool == "stochastic")
type = PoolingLayer::STOCHASTIC;
else
CV_Error(Error::StsBadArg, "Unknown pooling type \"" + pool + "\"");
}
else
{
type = PoolingLayer::MAX;
} }
getCaffeConvParams(params, kernel, pad, stride);
return Ptr<Layer>(new PoolingLayerImpl(type, kernel, pad, stride));
}
} }
} }

35
modules/dnn/src/layers/pooling_layer.hpp
Unescape View File

@ -1,4 +1,4 @@
/*M/////////////////////////////////////////////////////////////////////////////////////// /*M///////////////////////////////////////////////////////////////////////////////////////
// //
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
// //
@ -42,33 +42,40 @@
#ifndef __OPENCV_DNN_LAYERS_POOLING_LAYER_HPP__ #ifndef __OPENCV_DNN_LAYERS_POOLING_LAYER_HPP__
#define __OPENCV_DNN_LAYERS_POOLING_LAYER_HPP__ #define __OPENCV_DNN_LAYERS_POOLING_LAYER_HPP__
#include "../precomp.hpp" #include "../precomp.hpp"
#include <opencv2/dnn/all_layers.hpp>
namespace cv namespace cv
{ {
namespace dnn namespace dnn
{ {
class PoolingLayer : public Layer class PoolingLayerImpl : public PoolingLayer
{ {
enum bool useOpenCL;
{
MAX,
AVE,
STOCHASTIC
};
int type;
Size kernel, pad, stride;
Size inp, out; Size inp, out;
void computeOutputShape(Size inpSz); void computeOutputShape(Size inpSz);
void maxPooling(Blob &input, Blob &output);
void avePooling(Blob &input, Blob &output); bool pooling_ocl(const char *kname, const Blob &src, Blob &dst, Blob *mask = NULL);
void maxPooling(Blob &src, Blob &dst);
void maxPooling_cpu(Blob &src, Blob &dst);
bool maxPooling_ocl(Blob &src, Blob &dst);
void avePooling(Blob &src, Blob &dst);
void avePooling_cpu(Blob &src, Blob &dst);
bool avePooling_ocl(Blob &src, Blob &dst);
public: public:
PoolingLayer(LayerParams &params);
PoolingLayerImpl();
PoolingLayerImpl(int type, Size kernel, Size pad, Size stride);
void allocate(const std::vector<Blob*> &inputs, std::vector<Blob> &outputs); void allocate(const std::vector<Blob*> &inputs, std::vector<Blob> &outputs);
void forward(std::vector<Blob*> &inputs, std::vector<Blob> &outputs); void forward(std::vector<Blob*> &inputs, std::vector<Blob> &outputs);
}; };
Ptr<Layer> createPoolingLayerFromCaffe(LayerParams &params);
} }
} }
#endif #endif

94
modules/dnn/src/opencl/pooling.cl
Unescape View File

@ -0,0 +1,94 @@
/*************************************************************************************
* Copyright (c) 2015, Advanced Micro Devices, Inc.
* All rights reserved.
*
* Redistribution and use in source and binary forms, with or without modification,
* are permitted provided that the following conditions are met:
*
* 1. Redistributions of source code must retain the above copyright notice, this
* list of conditions and the following disclaimer.
*
* 2. Redistributions in binary form must reproduce the above copyright notice,
* this list of conditions and the following disclaimer in the documentation and/or
* other materials provided with the distribution.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
* IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
* INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
* BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA,
* OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
* WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
* ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
* POSSIBILITY OF SUCH DAMAGE.
**************************************************************************************/
__kernel void MaxPoolForward(const int nthreads, __global T* bottom_data, const int num, const int channels, const int height, const int width, const int pooled_height, const int pooled_width, const int kernel_h, const int kernel_w, const int stride_h, const int stride_w, const int pad_h, const int pad_w, __global T* top_data
#ifdef MASK
, __global int* mask, __global T* top_mask
#endif
) {
int index = get_global_id(0);
int tmp = get_global_size(0);
for(index; index < nthreads; index += tmp) {
int pw = index % pooled_width;
int ph = (index / pooled_width) % pooled_height;
int c = (index / pooled_width / pooled_height) % channels;
int n = index / pooled_width / pooled_height / channels;
int hstart = ph * stride_h - pad_h;
int wstart = pw * stride_w - pad_w;
const int hend = min(hstart + kernel_h, height);
const int wend = min(wstart + kernel_w, width);
hstart = max(hstart, 0);
wstart = max(wstart, 0);
T maxval = -FLT_MAX;
int maxidx = -1;
bottom_data =
bottom_data + (n * channels + c) * height * width;
for (int h = hstart; h < hend; ++h) {
for (int w = wstart; w < wend; ++w) {
if (bottom_data[h * width + w] > maxval) {
maxidx = h * width + w;
maxval = bottom_data[maxidx];
}
}
}
top_data[index] = maxval;
#ifdef MASK
if (mask) {
mask[index] = maxidx;
} else {
top_mask[index] = maxidx;
}
#endif
}
}
__kernel void AvePoolForward(const int nthreads, __global T* bottom_data, const int num, const int channels, const int height, const int width, const int pooled_height, const int pooled_width, const int kernel_h, const int kernel_w, const int stride_h, const int stride_w, const int pad_h, const int pad_w,__global T* top_data) {
int index = get_global_id(0);
int tmp = get_global_size(0);
for(index; index < nthreads; index+=tmp) {
int pw = index % pooled_width;
int ph = (index / pooled_width) % pooled_height;
int c = (index / pooled_width / pooled_height) % channels;
int n = index / pooled_width / pooled_height / channels; int hstart = ph * stride_h - pad_h; int wstart = pw * stride_w - pad_w;
int hend = min(hstart + kernel_h, height + pad_h);
int wend = min(wstart + kernel_w, width + pad_w);
const int pool_size = (hend - hstart) * (wend - wstart);
hstart = max(hstart, 0);
wstart = max(wstart, 0);
hend = min(hend, height);
wend = min(wend, width);
T aveval = 0;
bottom_data =
bottom_data + (n * channels + c) * height * width;
for (int h = hstart; h < hend; ++h) {
for (int w = wstart; w < wend; ++w) {
aveval += bottom_data[h * width + w];
}
}
top_data[index] = aveval / pool_size;
}
}

16
modules/dnn/test/test_layers.cpp
Unescape View File

@ -142,12 +142,24 @@ TEST(Layer_Test_InnerProduct, Accuracy)
TEST(Layer_Test_Pooling_max, Accuracy) TEST(Layer_Test_Pooling_max, Accuracy)
{ {
testLayerUsingCaffeModels("layer_pooling_max"); OCL_OFF(testLayerUsingCaffeModels("layer_pooling_max"));
OCL_ON();
}
OCL_TEST(Layer_Test_Pooling_max, Accuracy)
{
OCL_ON(testLayerUsingCaffeModels("layer_pooling_max"));
OCL_OFF();
} }
TEST(Layer_Test_Pooling_ave, Accuracy) TEST(Layer_Test_Pooling_ave, Accuracy)
{ {
testLayerUsingCaffeModels("layer_pooling_ave"); OCL_OFF(testLayerUsingCaffeModels("layer_pooling_ave"));
OCL_ON();
}
OCL_TEST(Layer_Test_Pooling_ave, Accuracy)
{
OCL_ON(testLayerUsingCaffeModels("layer_pooling_ave"));
OCL_OFF();
} }
TEST(Layer_Test_MVN, Accuracy) TEST(Layer_Test_MVN, Accuracy)

Write Preview
Loading…
Cancel
Save