Merge pull request #707 from ludv1x:dnn
Vadim Pisarevsky
9 years ago
commit
9a342b5187
66 changed files with 5066 additions and 1307 deletions
@ -0,0 +1,97 @@ |
||||
#COPYRIGHT |
||||
# |
||||
#All contributions by the University of California: |
||||
#Copyright (c) 2014, 2015, The Regents of the University of California (Regents) |
||||
#All rights reserved. |
||||
# |
||||
#All other contributions: |
||||
#Copyright (c) 2014, 2015, the respective contributors |
||||
#All rights reserved. |
||||
# |
||||
#Caffe uses a shared copyright model: each contributor holds copyright over |
||||
#their contributions to Caffe. The project versioning records all such |
||||
#contribution and copyright details. If a contributor wants to further mark |
||||
#their specific copyright on a particular contribution, they should indicate |
||||
#their copyright solely in the commit message of the change when it is |
||||
#committed. |
||||
# |
||||
#LICENSE |
||||
# |
||||
#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 OWNER 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. |
||||
# |
||||
#CONTRIBUTION AGREEMENT |
||||
# |
||||
#By contributing to the BVLC/caffe repository through pull-request, comment, |
||||
#or otherwise, the contributor releases their content to the |
||||
#license and copyright terms herein. |
||||
|
||||
|
||||
# Find the Atlas (and Lapack) libraries |
||||
# |
||||
# The following variables are optionally searched for defaults |
||||
# Atlas_ROOT_DIR: Base directory where all Atlas components are found |
||||
# |
||||
# The following are set after configuration is done: |
||||
# Atlas_FOUND |
||||
# Atlas_INCLUDE_DIRS |
||||
# Atlas_LIBRARIES |
||||
# Atlas_LIBRARYRARY_DIRS |
||||
|
||||
set(Atlas_INCLUDE_SEARCH_PATHS |
||||
/usr/include/atlas |
||||
/usr/include/atlas-base |
||||
$ENV{Atlas_ROOT_DIR} |
||||
$ENV{Atlas_ROOT_DIR}/include |
||||
) |
||||
|
||||
set(Atlas_LIB_SEARCH_PATHS |
||||
/usr/lib/atlas |
||||
/usr/lib/atlas-base |
||||
$ENV{Atlas_ROOT_DIR} |
||||
$ENV{Atlas_ROOT_DIR}/lib |
||||
) |
||||
|
||||
find_path(Atlas_CBLAS_INCLUDE_DIR NAMES cblas.h PATHS ${Atlas_INCLUDE_SEARCH_PATHS}) |
||||
find_path(Atlas_CLAPACK_INCLUDE_DIR NAMES clapack.h PATHS ${Atlas_INCLUDE_SEARCH_PATHS}) |
||||
|
||||
find_library(Atlas_CBLAS_LIBRARY NAMES ptcblas_r ptcblas cblas_r cblas PATHS ${Atlas_LIB_SEARCH_PATHS}) |
||||
find_library(Atlas_BLAS_LIBRARY NAMES atlas_r atlas PATHS ${Atlas_LIB_SEARCH_PATHS}) |
||||
find_library(Atlas_LAPACK_LIBRARY NAMES alapack_r alapack lapack_atlas PATHS ${Atlas_LIB_SEARCH_PATHS}) |
||||
|
||||
set(LOOKED_FOR |
||||
Atlas_CBLAS_INCLUDE_DIR |
||||
Atlas_CLAPACK_INCLUDE_DIR |
||||
|
||||
Atlas_CBLAS_LIBRARY |
||||
Atlas_BLAS_LIBRARY |
||||
Atlas_LAPACK_LIBRARY |
||||
) |
||||
|
||||
include(FindPackageHandleStandardArgs) |
||||
find_package_handle_standard_args(Atlas DEFAULT_MSG ${LOOKED_FOR}) |
||||
|
||||
if(ATLAS_FOUND) |
||||
set(Atlas_INCLUDE_DIR ${Atlas_CBLAS_INCLUDE_DIR} ${Atlas_CLAPACK_INCLUDE_DIR}) |
||||
set(Atlas_LIBRARIES ${Atlas_LAPACK_LIBRARY} ${Atlas_CBLAS_LIBRARY} ${Atlas_BLAS_LIBRARY}) |
||||
mark_as_advanced(${LOOKED_FOR}) |
||||
|
||||
message(STATUS "Found Atlas (include: ${Atlas_CBLAS_INCLUDE_DIR}, library: ${Atlas_BLAS_LIBRARY})") |
||||
endif(ATLAS_FOUND) |
||||
@ -0,0 +1,106 @@ |
||||
#COPYRIGHT |
||||
# |
||||
#All contributions by the University of California: |
||||
#Copyright (c) 2014, 2015, The Regents of the University of California (Regents) |
||||
#All rights reserved. |
||||
# |
||||
#All other contributions: |
||||
#Copyright (c) 2014, 2015, the respective contributors |
||||
#All rights reserved. |
||||
# |
||||
#Caffe uses a shared copyright model: each contributor holds copyright over |
||||
#their contributions to Caffe. The project versioning records all such |
||||
#contribution and copyright details. If a contributor wants to further mark |
||||
#their specific copyright on a particular contribution, they should indicate |
||||
#their copyright solely in the commit message of the change when it is |
||||
#committed. |
||||
# |
||||
#LICENSE |
||||
# |
||||
#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 OWNER 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. |
||||
# |
||||
#CONTRIBUTION AGREEMENT |
||||
# |
||||
#By contributing to the BVLC/caffe repository through pull-request, comment, |
||||
#or otherwise, the contributor releases their content to the |
||||
#license and copyright terms herein. |
||||
|
||||
SET(Open_BLAS_INCLUDE_SEARCH_PATHS |
||||
/usr/include |
||||
/usr/include/openblas |
||||
/usr/include/openblas-base |
||||
/usr/local/include |
||||
/usr/local/include/openblas |
||||
/usr/local/include/openblas-base |
||||
/opt/OpenBLAS/include |
||||
$ENV{OpenBLAS_HOME} |
||||
$ENV{OpenBLAS_HOME}/include |
||||
) |
||||
|
||||
SET(Open_BLAS_LIB_SEARCH_PATHS |
||||
/lib/ |
||||
/lib/openblas-base |
||||
/lib64/ |
||||
/usr/lib |
||||
/usr/lib/openblas-base |
||||
/usr/lib64 |
||||
/usr/local/lib |
||||
/usr/local/lib64 |
||||
/opt/OpenBLAS/lib |
||||
$ENV{OpenBLAS}cd |
||||
$ENV{OpenBLAS}/lib |
||||
$ENV{OpenBLAS_HOME} |
||||
$ENV{OpenBLAS_HOME}/lib |
||||
) |
||||
|
||||
FIND_PATH(OpenBLAS_INCLUDE_DIR NAMES cblas.h PATHS ${Open_BLAS_INCLUDE_SEARCH_PATHS}) |
||||
FIND_LIBRARY(OpenBLAS_LIB NAMES openblas PATHS ${Open_BLAS_LIB_SEARCH_PATHS}) |
||||
|
||||
SET(OpenBLAS_FOUND ON) |
||||
|
||||
# Check include files |
||||
IF(NOT OpenBLAS_INCLUDE_DIR) |
||||
SET(OpenBLAS_FOUND OFF) |
||||
MESSAGE(STATUS "Could not find OpenBLAS include. Turning OpenBLAS_FOUND off") |
||||
ENDIF() |
||||
|
||||
# Check libraries |
||||
IF(NOT OpenBLAS_LIB) |
||||
SET(OpenBLAS_FOUND OFF) |
||||
MESSAGE(STATUS "Could not find OpenBLAS lib. Turning OpenBLAS_FOUND off") |
||||
ENDIF() |
||||
|
||||
IF (OpenBLAS_FOUND) |
||||
IF (NOT OpenBLAS_FIND_QUIETLY) |
||||
MESSAGE(STATUS "Found OpenBLAS libraries: ${OpenBLAS_LIB}") |
||||
MESSAGE(STATUS "Found OpenBLAS include: ${OpenBLAS_INCLUDE_DIR}") |
||||
ENDIF (NOT OpenBLAS_FIND_QUIETLY) |
||||
ELSE (OpenBLAS_FOUND) |
||||
IF (OpenBLAS_FIND_REQUIRED) |
||||
MESSAGE(FATAL_ERROR "Could not find OpenBLAS") |
||||
ENDIF (OpenBLAS_FIND_REQUIRED) |
||||
ENDIF (OpenBLAS_FOUND) |
||||
|
||||
MARK_AS_ADVANCED( |
||||
OpenBLAS_INCLUDE_DIR |
||||
OpenBLAS_LIB |
||||
OpenBLAS |
||||
) |
||||
@ -0,0 +1,60 @@ |
||||
macro(_find_file_in_dirs VAR NAME DIRS) |
||||
find_path(${VAR} ${NAME} ${DIRS} NO_DEFAULT_PATH) |
||||
set(${VAR} ${${VAR}}/${NAME}) |
||||
unset(${VAR} CACHE) |
||||
endmacro() |
||||
|
||||
if(${the_module}_WITH_BLAS) |
||||
set(_bp ${the_module}_BLAS) #prefix for blas variables |
||||
set(BLAS_CBLAS_H "cblas.h") |
||||
set(HAVE_BLAS "") |
||||
|
||||
if(NOT HAVE_BLAS) #check custom BLAS from user input |
||||
if(${_bp}_INCLUDE_DIR AND ${_bp}_LIBRARIES AND ${_bp}_CBLAS_H) |
||||
set(HAVE_BLAS "Custom") |
||||
endif() |
||||
endif() |
||||
if(NOT HAVE_BLAS) |
||||
include(cmake/OpenCVFindMKL.cmake) |
||||
if(MKL_FOUND) |
||||
set(BLAS_INCLUDE_DIR ${MKL_INCLUDE_DIRS}) |
||||
set(BLAS_LIBRARIES ${MKL_LIBRARIES} ) |
||||
set(BLAS_CBLAS_H "mkl_cblas.h" ) |
||||
set(HAVE_BLAS "MKL") |
||||
endif() |
||||
endif() |
||||
if(NOT HAVE_BLAS) |
||||
include(cmake/FindOpenBLAS.cmake) |
||||
if(OpenBLAS_FOUND) |
||||
set(BLAS_INCLUDE_DIR ${OpenBLAS_INCLUDE_DIR} ) |
||||
set(BLAS_LIBRARIES ${OpenBLAS_LIB} ) |
||||
set(HAVE_BLAS "OpenBLAS") |
||||
endif() |
||||
endif() |
||||
if(NOT HAVE_BLAS AND UNIX) |
||||
include(cmake/FindAtlas.cmake) |
||||
if(ATLAS_FOUND) |
||||
set(BLAS_INCLUDE_DIR ${Atlas_INCLUDE_DIR}) |
||||
set(BLAS_LIBRARIES ${Atlas_LIBRARIES} ) |
||||
set(HAVE_BLAS "Atlas") |
||||
endif() |
||||
endif() |
||||
|
||||
if(NOT HAVE_BLAS OR NOT (HAVE_BLAS STREQUAL "Custom")) |
||||
set(${_bp}_INCLUDE_DIR ${BLAS_INCLUDE_DIR} CACHE PATH "Path to BLAS include dir" FORCE) |
||||
set(${_bp}_CBLAS_H ${BLAS_CBLAS_H} CACHE STRING "Alternative name of cblas.h" FORCE) |
||||
set(${_bp}_LIBRARIES ${BLAS_LIBRARIES} CACHE FILEPATH "Path to BLAS libraries that will be linked with ${the_module} module" FORCE) |
||||
set(${_bp}_BINARIES ${BLAS_BINARIES} CACHE FILEPATH "Path to BLAS binaries (.so, .dll) that will be installed with ${the_module} module" FORCE) |
||||
endif() |
||||
|
||||
if(HAVE_BLAS) #adding proxy cblas.h header |
||||
_find_file_in_dirs(CBLAS_H_PATH ${${_bp}_CBLAS_H} ${${_bp}_INCLUDE_DIR}) |
||||
if(NOT CBLAS_H_PATH) |
||||
message(WARNING "CBLAS header '${${_bp}_CBLAS_H}' not found into '${${_bp}_INCLUDE_DIR}'") |
||||
endif() |
||||
|
||||
set(CBLAS_H_PROXY_PATH ${CMAKE_CURRENT_BINARY_DIR}/opencv_cblas.hpp) |
||||
set(_include_str "\#include \"${CBLAS_H_PATH}\"") |
||||
file(WRITE ${CBLAS_H_PROXY_PATH} ${_include_str}) |
||||
endif() |
||||
endif() |
||||
@ -0,0 +1,123 @@ |
||||
# |
||||
# The script to detect Intel(R) Math Kernel Library (MKL) |
||||
# installation/package |
||||
# |
||||
# Parameters: |
||||
# MKL_WITH_TBB |
||||
# |
||||
# On return this will define: |
||||
# |
||||
# HAVE_MKL - True if Intel IPP found |
||||
# MKL_ROOT_DIR - root of IPP installation |
||||
# MKL_INCLUDE_DIRS - IPP include folder |
||||
# MKL_LIBRARIES - IPP libraries that are used by OpenCV |
||||
# |
||||
|
||||
macro(mkl_fail) |
||||
set(HAVE_MKL OFF CACHE BOOL "True if MKL found") |
||||
set(MKL_ROOT_DIR ${MKL_ROOT_DIR} CACHE PATH "Path to MKL directory") |
||||
unset(MKL_INCLUDE_DIRS CACHE) |
||||
unset(MKL_LIBRARIES CACHE) |
||||
endmacro() |
||||
|
||||
macro(get_mkl_version VERSION_FILE) |
||||
# read MKL version info from file |
||||
file(STRINGS ${VERSION_FILE} STR1 REGEX "__INTEL_MKL__") |
||||
file(STRINGS ${VERSION_FILE} STR2 REGEX "__INTEL_MKL_MINOR__") |
||||
file(STRINGS ${VERSION_FILE} STR3 REGEX "__INTEL_MKL_UPDATE__") |
||||
#file(STRINGS ${VERSION_FILE} STR4 REGEX "INTEL_MKL_VERSION") |
||||
|
||||
# extract info and assign to variables |
||||
string(REGEX MATCHALL "[0-9]+" MKL_VERSION_MAJOR ${STR1}) |
||||
string(REGEX MATCHALL "[0-9]+" MKL_VERSION_MINOR ${STR2}) |
||||
string(REGEX MATCHALL "[0-9]+" MKL_VERSION_UPDATE ${STR3}) |
||||
set(MKL_VERSION_STR "${MKL_VERSION_MAJOR}.${MKL_VERSION_MINOR}.${MKL_VERSION_UPDATE}" CACHE STRING "MKL version" FORCE) |
||||
endmacro() |
||||
|
||||
|
||||
if(NOT DEFINED MKL_USE_MULTITHREAD) |
||||
OCV_OPTION(MKL_WITH_TBB "Use MKL with TBB multithreading" OFF)#ON IF WITH_TBB) |
||||
OCV_OPTION(MKL_WITH_OPENMP "Use MKL with OpenMP multithreading" OFF)#ON IF WITH_OPENMP) |
||||
endif() |
||||
|
||||
#check current MKL_ROOT_DIR |
||||
if(NOT MKL_ROOT_DIR OR NOT EXISTS ${MKL_ROOT_DIR}/include/mkl.h) |
||||
set(MKLROOT_PATHS ${MKL_ROOT_DIR}) |
||||
if(DEFINED $ENV{MKLROOT}) |
||||
list(APPEND MKLROOT_PATHS $ENV{MKLROOT}) |
||||
endif() |
||||
if(WIN32) |
||||
set(ProgramFilesx86 "ProgramFiles(x86)") |
||||
list(APPEND MKLROOT_PATHS $ENV{${ProgramFilesx86}}/IntelSWTools/compilers_and_libraries/windows/mkl) |
||||
endif() |
||||
if(UNIX) |
||||
list(APPEND MKLROOT_PATHS "/opt/intel/mkl") |
||||
endif() |
||||
|
||||
find_path(MKL_ROOT_DIR include/mkl.h PATHS ${MKLROOT_PATHS}) |
||||
endif() |
||||
|
||||
if(NOT MKL_ROOT_DIR) |
||||
mkl_fail() |
||||
return() |
||||
endif() |
||||
|
||||
set(MKL_INCLUDE_DIRS ${MKL_ROOT_DIR}/include) |
||||
set(MKL_INCLUDE_HEADERS ${MKL_INCLUDE_DIRS}/mkl.h ${MKL_INCLUDE_DIRS}/mkl_version.h) |
||||
|
||||
#determine arch |
||||
if(CMAKE_CXX_SIZEOF_DATA_PTR EQUAL 8) |
||||
set(MKL_X64 1) |
||||
set(MKL_ARCH "intel64") |
||||
|
||||
include(CheckTypeSize) |
||||
CHECK_TYPE_SIZE(int _sizeof_int) |
||||
if (_sizeof_int EQUAL 4) |
||||
set(MKL_LP64 "lp64") |
||||
else() |
||||
set(MKL_LP64 "ilp64") |
||||
endif() |
||||
else() |
||||
set(MKL_ARCH "ia32") |
||||
endif() |
||||
|
||||
if(MSVC) |
||||
set(MKL_EXT ".lib") |
||||
set(MKL_PRE "") |
||||
else() |
||||
set(MKL_EXT ".a") |
||||
set(MKL_PRE "lib") |
||||
endif() |
||||
|
||||
set(MKL_LIB_DIR ${MKL_ROOT_DIR}/lib/${MKL_ARCH}) |
||||
set(MKL_LIBRARIES ${MKL_LIB_DIR}/${MKL_PRE}mkl_core${MKL_EXT} ${MKL_LIB_DIR}/${MKL_PRE}mkl_intel_${MKL_LP64}${MKL_EXT}) |
||||
|
||||
if(MKL_WITH_TBB) |
||||
list(APPEND MKL_LIBRARIES ${MKL_LIB_DIR}/${MKL_PRE}mkl_tbb_thread${MKL_EXT}) |
||||
list(APPEND MKL_LIBRARIES ${MKL_ROOT_DIR}/../tbb/lib/${MKL_ARCH}/tbb${MKL_EXT}) |
||||
elseif(MKL_WITH_OPENMP) |
||||
message(FATAL_ERROR "Multithreaded MKL is not supported yet") |
||||
else() |
||||
list(APPEND MKL_LIBRARIES ${MKL_LIB_DIR}/${MKL_PRE}mkl_sequential${MKL_EXT}) |
||||
endif() |
||||
|
||||
include(FindPackageHandleStandardArgs) |
||||
find_package_handle_standard_args(MKL MKL_INCLUDE_HEADERS MKL_LIBRARIES) |
||||
|
||||
if(MKL_FOUND) |
||||
get_mkl_version(${MKL_INCLUDE_DIRS}/mkl_version.h) |
||||
message(STATUS "Found MKL ${MKL_VERSION_STR} at: ${MKL_ROOT_DIR}") |
||||
|
||||
set(HAVE_MKL ON CACHE BOOL "True if MKL found") |
||||
set(MKL_ROOT_DIR ${MKL_ROOT_DIR} CACHE PATH "Path to MKL directory") |
||||
set(MKL_INCLUDE_DIRS ${MKL_INCLUDE_DIRS} CACHE PATH "Path to MKL include directory") |
||||
if(NOT UNIX) |
||||
set(MKL_LIBRARIES ${MKL_LIBRARIES} CACHE FILEPATH "MKL libarries") |
||||
else() |
||||
#it's ugly but helps to avoid cyclic lib problem |
||||
set(MKL_LIBRARIES ${MKL_LIBRARIES} ${MKL_LIBRARIES} ${MKL_LIBRARIES} "-lpthread" "-lm" "-ldl") |
||||
set(MKL_LIBRARIES ${MKL_LIBRARIES} CACHE STRING "MKL libarries") |
||||
endif() |
||||
else() |
||||
|
||||
endif() |
||||
@ -0,0 +1,371 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's 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.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// 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 Intel Corporation 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.
|
||||
//
|
||||
//M*/
|
||||
|
||||
#ifndef __OPENCV_DNN_DNN_ALL_LAYERS_HPP__ |
||||
#define __OPENCV_DNN_DNN_ALL_LAYERS_HPP__ |
||||
#include <opencv2/dnn.hpp> |
||||
|
||||
namespace cv |
||||
{ |
||||
namespace dnn |
||||
{ |
||||
//! @addtogroup dnn
|
||||
//! @{
|
||||
|
||||
/** @defgroup dnnLayerList Partial List of Implemented Layers
|
||||
@{ |
||||
This subsection of dnn module contains information about bult-in layers and their descriptions. |
||||
|
||||
Classes listed here, in fact, provides C++ API for creating intances of bult-in layers. |
||||
In addition to this way of layers instantiation, there is a more common factory API (see @ref dnnLayerFactory), it allows to create layers dynamically (by name) and register new ones. |
||||
You can use both API, but factory API is less convinient for native C++ programming and basically designed for use inside importers (see @ref Importer, @ref createCaffeImporter(), @ref createTorchImporter()). |
||||
|
||||
Bult-in layers partially reproduce functionality of corresponding Caffe and Torch7 layers. |
||||
In partuclar, the following layers and Caffe @ref Importer were tested to reproduce <a href="http://caffe.berkeleyvision.org/tutorial/layers.html">Caffe</a> functionality: |
||||
- Convolution |
||||
- Deconvolution |
||||
- Pooling |
||||
- InnerProduct |
||||
- TanH, ReLU, Sigmoid, BNLL, Power, AbsVal |
||||
- Softmax |
||||
- Reshape, Flatten, Slice, Split |
||||
- LRN |
||||
- MVN |
||||
- Dropout (since it does nothing on forward pass -)) |
||||
*/ |
||||
|
||||
//! LSTM recurrent layer
|
||||
class CV_EXPORTS_W LSTMLayer : public Layer |
||||
{ |
||||
public: |
||||
/** Creates instance of LSTM layer */ |
||||
static Ptr<LSTMLayer> create(); |
||||
|
||||
/** Set trained weights for LSTM layer.
|
||||
LSTM behavior on each step is defined by current input, previous output, previous cell state and learned weights. |
||||
|
||||
Let @f$x_t@f$ be current input, @f$h_t@f$ be current output, @f$c_t@f$ be current state. |
||||
Than current output and current cell state is computed as follows: |
||||
@f{eqnarray*}{ |
||||
h_t &= o_t \odot tanh(c_t), \\
|
||||
c_t &= f_t \odot c_{t-1} + i_t \odot g_t, \\
|
||||
@f} |
||||
where @f$\odot@f$ is per-element multiply operation and @f$i_t, f_t, o_t, g_t@f$ is internal gates that are computed using learned wights. |
||||
|
||||
Gates are computed as follows: |
||||
@f{eqnarray*}{ |
||||
i_t &= sigmoid&(W_{xi} x_t + W_{hi} h_{t-1} + b_i), \\
|
||||
f_t &= sigmoid&(W_{xf} x_t + W_{hf} h_{t-1} + b_f), \\
|
||||
o_t &= sigmoid&(W_{xo} x_t + W_{ho} h_{t-1} + b_o), \\
|
||||
g_t &= tanh &(W_{xg} x_t + W_{hg} h_{t-1} + b_g), \\
|
||||
@f} |
||||
where @f$W_{x?}@f$, @f$W_{h?}@f$ and @f$b_{?}@f$ are learned weights represented as matrices: |
||||
@f$W_{x?} \in R^{N_h \times N_x}@f$, @f$W_{h?} \in R^{N_h \times N_h}@f$, @f$b_? \in R^{N_h}@f$. |
||||
|
||||
For simplicity and performance purposes we use @f$ W_x = [W_{xi}; W_{xf}; W_{xo}, W_{xg}] @f$ |
||||
(i.e. @f$W_x@f$ is vertical contacentaion of @f$ W_{x?} @f$), @f$ W_x \in R^{4N_h \times N_x} @f$. |
||||
The same for @f$ W_h = [W_{hi}; W_{hf}; W_{ho}, W_{hg}], W_h \in R^{4N_h \times N_h} @f$ |
||||
and for @f$ b = [b_i; b_f, b_o, b_g]@f$, @f$b \in R^{4N_h} @f$. |
||||
|
||||
@param Wh is matrix defining how previous output is transformed to internal gates (i.e. according to abovemtioned notation is @f$ W_h @f$) |
||||
@param Wx is matrix defining how current input is transformed to internal gates (i.e. according to abovemtioned notation is @f$ W_x @f$) |
||||
@param b is bias vector (i.e. according to abovemtioned notation is @f$ b @f$) |
||||
*/ |
||||
virtual void setWeights(const Blob &Wh, const Blob &Wx, const Blob &b) = 0; |
||||
|
||||
/** @brief Specifies shape of output blob which will be [[`T`], `N`] + @p outTailShape.
|
||||
* @details If this parameter is empty or unset then @p outTailShape = [`Wh`.size(0)] will be used, |
||||
* where `Wh` is parameter from setWeights(). |
||||
*/ |
||||
virtual void setOutShape(const BlobShape &outTailShape = BlobShape::empty()) = 0; |
||||
|
||||
/** @brief Set @f$ h_{t-1} @f$ value that will be used in next forward() calls.
|
||||
* @details By-default @f$ h_{t-1} @f$ is inited by zeros and updated after each forward() call. |
||||
*/ |
||||
virtual void setH(const Blob &H) = 0; |
||||
/** @brief Returns current @f$ h_{t-1} @f$ value (deep copy). */ |
||||
virtual Blob getH() const = 0; |
||||
|
||||
/** @brief Set @f$ c_{t-1} @f$ value that will be used in next forward() calls.
|
||||
* @details By-default @f$ c_{t-1} @f$ is inited by zeros and updated after each forward() call. |
||||
*/ |
||||
virtual void setC(const Blob &C) = 0; |
||||
/** @brief Returns current @f$ c_{t-1} @f$ value (deep copy). */ |
||||
virtual Blob getC() const = 0; |
||||
|
||||
/** @brief Specifies either interpet first dimension of input blob as timestamp dimenion either as sample.
|
||||
* |
||||
* If flag is set to true then shape of input blob will be interpeted as [`T`, `N`, `[data dims]`] where `T` specifies number of timpestamps, `N` is number of independent streams. |
||||
* In this case each forward() call will iterate through `T` timestamps and update layer's state `T` times. |
||||
* |
||||
* If flag is set to false then shape of input blob will be interpeted as [`N`, `[data dims]`]. |
||||
* In this case each forward() call will make one iteration and produce one timestamp with shape [`N`, `[out dims]`]. |
||||
*/ |
||||
virtual void setUseTimstampsDim(bool use = true) = 0; |
||||
|
||||
/** @brief If this flag is set to true then layer will produce @f$ c_t @f$ as second output.
|
||||
* @details Shape of the second output is the same as first output. |
||||
*/ |
||||
virtual void setProduceCellOutput(bool produce = false) = 0; |
||||
|
||||
/** In common case it use single input with @f$x_t@f$ values to compute output(s) @f$h_t@f$ (and @f$c_t@f$).
|
||||
* @param input should contain packed values @f$x_t@f$ |
||||
* @param output contains computed outputs: @f$h_t@f$ (and @f$c_t@f$ if setProduceCellOutput() flag was set to true). |
||||
* |
||||
* If setUseTimstampsDim() is set to true then @p input[0] should has at least two dimensions with the following shape: [`T`, `N`, `[data dims]`], |
||||
* where `T` specifies number of timpestamps, `N` is number of independent streams (i.e. @f$ x_{t_0 + t}^{stream} @f$ is stored inside @p input[0][t, stream, ...]). |
||||
* |
||||
* If setUseTimstampsDim() is set to fase then @p input[0] should contain single timestamp, its shape should has form [`N`, `[data dims]`] with at least one dimension. |
||||
* (i.e. @f$ x_{t}^{stream} @f$ is stored inside @p input[0][stream, ...]). |
||||
*/ |
||||
void forward(std::vector<Blob*> &input, std::vector<Blob> &output); |
||||
|
||||
int inputNameToIndex(String inputName); |
||||
|
||||
int outputNameToIndex(String outputName); |
||||
}; |
||||
|
||||
//! Classical recurrent layer
|
||||
class CV_EXPORTS_W RNNLayer : public Layer |
||||
{ |
||||
public: |
||||
/** Creates instance of RNNLayer */ |
||||
static Ptr<RNNLayer> create(); |
||||
|
||||
/** Setups learned weights.
|
||||
|
||||
Recurrent-layer behavior on each step is defined by current input @f$ x_t @f$, previous state @f$ h_t @f$ and learned weights as follows: |
||||
@f{eqnarray*}{ |
||||
h_t &= tanh&(W_{hh} h_{t-1} + W_{xh} x_t + b_h), \\
|
||||
o_t &= tanh&(W_{ho} h_t + b_o), |
||||
@f} |
||||
|
||||
@param Wxh is @f$ W_{xh} @f$ matrix |
||||
@param bh is @f$ b_{h} @f$ vector |
||||
@param Whh is @f$ W_{hh} @f$ matrix |
||||
@param Who is @f$ W_{xo} @f$ matrix |
||||
@param bo is @f$ b_{o} @f$ vector |
||||
*/ |
||||
virtual void setWeights(const Blob &Wxh, const Blob &bh, const Blob &Whh, const Blob &Who, const Blob &bo) = 0; |
||||
|
||||
/** @brief If this flag is set to true then layer will produce @f$ h_t @f$ as second output.
|
||||
* @details Shape of the second output is the same as first output. |
||||
*/ |
||||
virtual void setProduceHiddenOutput(bool produce = false) = 0; |
||||
|
||||
/** Accepts two inputs @f$x_t@f$ and @f$h_{t-1}@f$ and compute two outputs @f$o_t@f$ and @f$h_t@f$.
|
||||
|
||||
@param input should contain packed input @f$x_t@f$. |
||||
@param output should contain output @f$o_t@f$ (and @f$h_t@f$ if setProduceHiddenOutput() is set to true). |
||||
|
||||
@p input[0] should have shape [`T`, `N`, `data_dims`] where `T` and `N` is number of timestamps and number of independent samples of @f$x_t@f$ respectively. |
||||
|
||||
@p output[0] will have shape [`T`, `N`, @f$N_o@f$], where @f$N_o@f$ is number of rows in @f$ W_{xo} @f$ matrix. |
||||
|
||||
If setProduceHiddenOutput() is set to true then @p output[1] will contain a Blob with shape [`T`, `N`, @f$N_h@f$], where @f$N_h@f$ is number of rows in @f$ W_{hh} @f$ matrix. |
||||
*/ |
||||
void forward(std::vector<Blob*> &input, std::vector<Blob> &output); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W BaseConvolutionLayer : public Layer |
||||
{ |
||||
public: |
||||
|
||||
Size kernel, stride, pad; |
||||
}; |
||||
|
||||
class CV_EXPORTS_W ConvolutionLayer : public BaseConvolutionLayer |
||||
{ |
||||
public: |
||||
|
||||
static Ptr<BaseConvolutionLayer> create(Size kernel = Size(3, 3), Size stride = Size(1, 1), Size pad = Size(0, 0)); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W DeconvolutionLayer : public BaseConvolutionLayer |
||||
{ |
||||
public: |
||||
|
||||
static Ptr<BaseConvolutionLayer> create(Size kernel = Size(3, 3), Size stride = Size(1, 1), Size pad = Size(0, 0)); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W LRNLayer : public Layer |
||||
{ |
||||
public: |
||||
|
||||
enum Type |
||||
{ |
||||
CHANNEL_NRM, |
||||
SPATIAL_NRM |
||||
}; |
||||
int type; |
||||
|
||||
int size; |
||||
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 stride = Size(1, 1), Size pad = Size(0, 0)); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W SoftmaxLayer : public Layer |
||||
{ |
||||
public: |
||||
|
||||
static Ptr<SoftmaxLayer> create(int axis = 1); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W InnerProductLayer : public Layer |
||||
{ |
||||
public: |
||||
int axis; |
||||
|
||||
static Ptr<InnerProductLayer> create(int axis = 1); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W MVNLayer : public Layer |
||||
{ |
||||
public: |
||||
double eps; |
||||
bool normVariance, acrossChannels; |
||||
|
||||
static Ptr<MVNLayer> create(bool normVariance = true, bool acrossChannels = false, double eps = 1e-9); |
||||
}; |
||||
|
||||
/* Reshaping */ |
||||
|
||||
class CV_EXPORTS_W ReshapeLayer : public Layer |
||||
{ |
||||
public: |
||||
BlobShape newShapeDesc; |
||||
Range newShapeRange; |
||||
|
||||
static Ptr<ReshapeLayer> create(const BlobShape &newShape, Range applyingRange = Range::all()); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W ConcatLayer : public Layer |
||||
{ |
||||
public: |
||||
int axis; |
||||
|
||||
static Ptr<ConcatLayer> create(int axis = 1); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W SplitLayer : public Layer |
||||
{ |
||||
public: |
||||
int outputsCount; //!< Number of copies that will be produced (is ignored when negative).
|
||||
|
||||
static Ptr<SplitLayer> create(int outputsCount = -1); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W SliceLayer : public Layer |
||||
{ |
||||
public: |
||||
int axis; |
||||
std::vector<int> sliceIndices; |
||||
|
||||
static Ptr<SliceLayer> create(int axis); |
||||
static Ptr<SliceLayer> create(int axis, const std::vector<int> &sliceIndices); |
||||
}; |
||||
|
||||
/* Activations */ |
||||
|
||||
class CV_EXPORTS_W ReLULayer : public Layer |
||||
{ |
||||
public: |
||||
double negativeSlope; |
||||
|
||||
static Ptr<ReLULayer> create(double negativeSlope = 0); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W TanHLayer : public Layer |
||||
{ |
||||
public: |
||||
static Ptr<TanHLayer> create(); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W SigmoidLayer : public Layer |
||||
{ |
||||
public: |
||||
static Ptr<SigmoidLayer> create(); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W BNLLLayer : public Layer |
||||
{ |
||||
public: |
||||
static Ptr<BNLLLayer> create(); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W AbsLayer : public Layer |
||||
{ |
||||
public: |
||||
static Ptr<AbsLayer> create(); |
||||
}; |
||||
|
||||
class CV_EXPORTS_W PowerLayer : public Layer |
||||
{ |
||||
public: |
||||
double power, scale, shift; |
||||
|
||||
static Ptr<PowerLayer> create(double power = 1, double scale = 1, double shift = 0); |
||||
}; |
||||
|
||||
//! @}
|
||||
//! @}
|
||||
|
||||
} |
||||
} |
||||
#endif |
||||
@ -0,0 +1,137 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's 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.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// 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 Intel Corporation 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.
|
||||
//
|
||||
//M*/
|
||||
|
||||
#ifndef __OPENCV_DNN_DNN_SHAPE_UTILS_HPP__ |
||||
#define __OPENCV_DNN_DNN_SHAPE_UTILS_HPP__ |
||||
|
||||
#include <opencv2/core.hpp> |
||||
#include <ostream> |
||||
|
||||
namespace cv { |
||||
namespace dnn { |
||||
|
||||
//Useful shortcut
|
||||
typedef BlobShape Shape; |
||||
|
||||
inline std::ostream &operator<< (std::ostream &s, cv::Range &r) |
||||
{ |
||||
return s << "[" << r.start << ", " << r.end << ")"; |
||||
} |
||||
|
||||
//Reshaping
|
||||
//TODO: add -1 specifier for automatic size inferring
|
||||
|
||||
template<typename Mat> |
||||
void reshape(Mat &m, const BlobShape &shape) |
||||
{ |
||||
m = m.reshape(1, shape.dims(), shape.ptr()); |
||||
} |
||||
|
||||
template<typename Mat> |
||||
Mat reshaped(const Mat &m, const BlobShape &shape) |
||||
{ |
||||
return m.reshape(1, shape.dims(), shape.ptr()); |
||||
} |
||||
|
||||
|
||||
//Slicing
|
||||
|
||||
struct _Range : public cv::Range |
||||
{ |
||||
_Range(const Range &r) : cv::Range(r) {} |
||||
_Range(int start, int size = 1) : cv::Range(start, start + size) {} |
||||
}; |
||||
|
||||
template<typename Mat> |
||||
Mat slice(const Mat &m, const _Range &r0) |
||||
{ |
||||
//CV_Assert(m.dims >= 1);
|
||||
cv::AutoBuffer<cv::Range, 4> ranges(m.dims); |
||||
for (int i = 1; i < m.dims; i++) |
||||
ranges[i] = Range::all(); |
||||
ranges[0] = r0; |
||||
return m(&ranges[0]); |
||||
} |
||||
|
||||
template<typename Mat> |
||||
Mat slice(const Mat &m, const _Range &r0, const _Range &r1) |
||||
{ |
||||
CV_Assert(m.dims >= 2); |
||||
cv::AutoBuffer<cv::Range, 4> ranges(m.dims); |
||||
for (int i = 2; i < m.dims; i++) |
||||
ranges[i] = Range::all(); |
||||
ranges[0] = r0; |
||||
ranges[1] = r1; |
||||
return m(&ranges[0]); |
||||
} |
||||
|
||||
template<typename Mat> |
||||
Mat slice(const Mat &m, const _Range &r0, const _Range &r1, const _Range &r2) |
||||
{ |
||||
CV_Assert(m.dims <= 3); |
||||
cv::AutoBuffer<cv::Range, 4> ranges(m.dims); |
||||
for (int i = 3; i < m.dims; i++) |
||||
ranges[i] = Range::all(); |
||||
ranges[0] = r0; |
||||
ranges[1] = r1; |
||||
ranges[2] = r2; |
||||
return m(&ranges[0]); |
||||
} |
||||
|
||||
template<typename Mat> |
||||
Mat slice(const Mat &m, const _Range &r0, const _Range &r1, const _Range &r2, const _Range &r3) |
||||
{ |
||||
CV_Assert(m.dims <= 4); |
||||
cv::AutoBuffer<cv::Range, 4> ranges(m.dims); |
||||
for (int i = 4; i < m.dims; i++) |
||||
ranges[i] = Range::all(); |
||||
ranges[0] = r0; |
||||
ranges[1] = r1; |
||||
ranges[2] = r2; |
||||
ranges[3] = r3; |
||||
return m(&ranges[0]); |
||||
} |
||||
|
||||
BlobShape computeShapeByReshapeMask(const BlobShape &srcShape, const BlobShape &maskShape, Range srcRange = Range::all()); |
||||
|
||||
} |
||||
} |
||||
#endif |
||||
@ -0,0 +1,80 @@ |
||||
#include "perf_precomp.hpp" |
||||
|
||||
namespace cvtest |
||||
{ |
||||
|
||||
using std::tr1::tuple; |
||||
using std::tr1::get; |
||||
using std::tr1::make_tuple; |
||||
using std::make_pair; |
||||
using namespace perf; |
||||
using namespace testing; |
||||
using namespace cv; |
||||
using namespace cv::dnn; |
||||
|
||||
enum {STRIDE_OFF = 1, STRIDE_ON = 2}; |
||||
CV_ENUM(StrideSize, STRIDE_OFF, STRIDE_ON); |
||||
|
||||
enum {GROUP_OFF = 1, GROUP_2 = 2}; |
||||
CV_ENUM(GroupSize, GROUP_OFF, GROUP_2); |
||||
|
||||
//Squared Size
|
||||
#define SSZ(n) cv::Size(n, n) |
||||
|
||||
typedef std::pair<BlobShape, int> InpShapeNumOut; |
||||
typedef tuple<Size, InpShapeNumOut, GroupSize, StrideSize> ConvParam; //kernel_size, inp shape, groups, stride
|
||||
typedef TestBaseWithParam<ConvParam> ConvolutionPerfTest; |
||||
|
||||
PERF_TEST_P( ConvolutionPerfTest, perf, Combine( |
||||
Values(Size(1, 1), Size(3, 3), Size(5, 5), Size(11, 11)), |
||||
Values(make_pair(BlobShape(1, 4, 224, 224), 64), |
||||
make_pair(BlobShape(1, 64, 112, 122), 128), |
||||
make_pair(BlobShape(1, 256, 28, 28), 512)), |
||||
GroupSize::all(), |
||||
StrideSize::all()) |
||||
) |
||||
{ |
||||
RNG rng(0); |
||||
|
||||
ConvParam params = GetParam(); |
||||
int ksz = get<0>(params).width; |
||||
BlobShape inpShape = get<1>(params).first; |
||||
int outCn = get<1>(params).second; |
||||
int groups = get<2>(params); |
||||
int stride = (ksz >= 11) ? 4 : (int)get<3>(params); |
||||
|
||||
int inpCn = inpShape[1]; |
||||
Blob wgtBlob(BlobShape(outCn, inpCn/groups, ksz, ksz)), biasBlob(BlobShape(outCn, 1, 1, 1)); |
||||
Blob inpBlob(inpShape); |
||||
rng.fill(biasBlob.matRef(), RNG::UNIFORM, -1, +1); |
||||
rng.fill(wgtBlob.matRef(), RNG::UNIFORM, -1, +1); |
||||
rng.fill(inpBlob.matRef(), RNG::UNIFORM, -1, +1); |
||||
|
||||
LayerParams lp; |
||||
lp.set("num_output", outCn); |
||||
lp.set("group", groups); |
||||
lp.set("stride", stride); |
||||
lp.set("kernel_size", ksz); |
||||
lp.blobs.reserve(2); |
||||
lp.blobs.push_back(wgtBlob); |
||||
lp.blobs.push_back(biasBlob); |
||||
|
||||
std::vector<Blob*> inpBlobs(1, &inpBlob); |
||||
std::vector<Blob> outBlobs; |
||||
|
||||
cv::setNumThreads(cv::getNumberOfCPUs()); |
||||
|
||||
Ptr<Layer> layer = cv::dnn::LayerFactory::createLayerInstance("Convolution", lp); |
||||
layer->allocate(inpBlobs, outBlobs); |
||||
|
||||
declare.in(inpBlob.matRef(), wgtBlob.matRef(), WARMUP_RNG).out(outBlobs[0].matRef()).tbb_threads(cv::getNumThreads()); |
||||
|
||||
TEST_CYCLE_N(10) |
||||
{ |
||||
layer->forward(inpBlobs, outBlobs); |
||||
} |
||||
|
||||
SANITY_CHECK_NOTHING(); |
||||
} |
||||
|
||||
} |
||||
@ -0,0 +1,3 @@ |
||||
#include "perf_precomp.hpp" |
||||
|
||||
CV_PERF_TEST_MAIN(dnn) |
||||
@ -0,0 +1,17 @@ |
||||
#ifdef __GNUC__ |
||||
# pragma GCC diagnostic ignored "-Wmissing-declarations" |
||||
# if defined __clang__ || defined __APPLE__ |
||||
# pragma GCC diagnostic ignored "-Wmissing-prototypes" |
||||
# pragma GCC diagnostic ignored "-Wextra" |
||||
# endif |
||||
#endif |
||||
|
||||
#ifndef __OPENCV_PERF_PRECOMP_HPP__ |
||||
#define __OPENCV_PERF_PRECOMP_HPP__ |
||||
|
||||
#include <opencv2/ts.hpp> |
||||
#include <opencv2/imgproc.hpp> |
||||
#include <opencv2/highgui.hpp> |
||||
#include <opencv2/dnn.hpp> |
||||
|
||||
#endif |
||||
@ -0,0 +1 @@ |
||||
*.caffemodel |
||||
@ -0,0 +1,294 @@ |
||||
#include "../precomp.hpp" |
||||
#include "layer_loaders.hpp" |
||||
#include <opencv2/dnn/shape_utils.hpp> |
||||
#include <climits> |
||||
|
||||
namespace cv |
||||
{ |
||||
namespace dnn |
||||
{ |
||||
|
||||
//Utils
|
||||
|
||||
//Extracts params used into Conv, Deconv and Pooling layers
|
||||
static void getCaffeConvParams(LayerParams ¶ms, Size &kernel, Size &pad, Size &stride) |
||||
{ |
||||
if (params.has("kernel_h") && params.has("kernel_w")) |
||||
{ |
||||
kernel.height = params.get<int>("kernel_h"); |
||||
kernel.width = params.get<int>("kernel_w"); |
||||
} |
||||
else if (params.has("kernel_size")) |
||||
{ |
||||
kernel.height = kernel.width = params.get<int>("kernel_size"); |
||||
} |
||||
else |
||||
{ |
||||
CV_Error(Error::StsBadArg, "kernel_size (or kernel_h and kernel_w) not specified"); |
||||
} |
||||
CV_Assert(kernel.height > 0 && kernel.width > 0); |
||||
|
||||
if (params.has("pad_h") && params.has("pad_w")) |
||||
{ |
||||
pad.height = params.get<int>("pad_h"); |
||||
pad.width = params.get<int>("pad_w"); |
||||
} |
||||
else |
||||
{ |
||||
pad.height = pad.width = params.get<int>("pad", 0); |
||||
} |
||||
CV_Assert(pad.height >= 0 && pad.width >= 0); |
||||
|
||||
if (params.has("stride_h") && params.has("stride_w")) |
||||
{ |
||||
stride.height = params.get<int>("stride_h"); |
||||
stride.width = params.get<int>("stride_w"); |
||||
} |
||||
else |
||||
{ |
||||
stride.height = stride.width = params.get<int>("stride", 1); |
||||
} |
||||
CV_Assert(stride.height > 0 && stride.width > 0); |
||||
} |
||||
|
||||
//Layers
|
||||
|
||||
//Convolution and Deconvolution
|
||||
static void initConvDeconvLayerFromCaffe(Ptr<BaseConvolutionLayer> l, LayerParams ¶ms) |
||||
{ |
||||
l->setParamsFrom(params); |
||||
getCaffeConvParams(params, l->kernel, l->pad, l->stride); |
||||
|
||||
bool bias = params.get<bool>("bias_term", true); |
||||
int numOutput = params.get<int>("num_output"); |
||||
int group = params.get<int>("group", 1); |
||||
|
||||
CV_Assert(numOutput % group == 0); |
||||
CV_Assert((bias && l->blobs.size() == 2) || (!bias && l->blobs.size() == 1)); |
||||
} |
||||
|
||||
template<> |
||||
Ptr<Layer> createLayerFromCaffe<ConvolutionLayer>(LayerParams ¶ms) |
||||
{ |
||||
Ptr<BaseConvolutionLayer> l = ConvolutionLayer::create(); |
||||
initConvDeconvLayerFromCaffe(l, params); |
||||
return Ptr<Layer>(l); |
||||
} |
||||
|
||||
template<> |
||||
Ptr<Layer> createLayerFromCaffe<DeconvolutionLayer>(LayerParams ¶ms) |
||||
{ |
||||
Ptr<BaseConvolutionLayer> l = DeconvolutionLayer::create(); |
||||
initConvDeconvLayerFromCaffe(l, params); |
||||
return Ptr<Layer>(l); |
||||
} |
||||
|
||||
template<> |
||||
Ptr<Layer> createLayerFromCaffe<PoolingLayer>(LayerParams ¶ms) |
||||
{ |
||||
int type; |
||||
Size kernel, stride, pad; |
||||
|
||||
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>(PoolingLayer::create(type, kernel, stride, pad)); |
||||
} |
||||
|
||||
template<> |
||||
Ptr<Layer> createLayerFromCaffe<SoftmaxLayer>(LayerParams ¶ms) |
||||
{ |
||||
int axis = params.get<int>("axis", 1); |
||||
return Ptr<Layer>(SoftmaxLayer::create(axis)); |
||||
} |
||||
|
||||
template<> //InnerProduct specialization
|
||||
Ptr<Layer> createLayerFromCaffe<InnerProductLayer>(LayerParams ¶ms) |
||||
{ |
||||
const std::vector<Blob> &blobs = params.blobs; |
||||
CV_Assert(1 <= blobs.size() && blobs.size() <= 2); |
||||
|
||||
int numOutputs = params.get<int>("num_output"); |
||||
int innerSize = (int)blobs[0].total() / numOutputs; |
||||
bool bias = params.get<bool>("bias_term", true); |
||||
int axis = params.get<int>("axis", 1); |
||||
|
||||
CV_Assert(blobs[0].dims() >= 2 && (size_t)(innerSize * numOutputs) == blobs[0].total()); |
||||
CV_Assert(!bias || (blobs.size() == 2 && (size_t)numOutputs == blobs[1].total())); |
||||
|
||||
Ptr<InnerProductLayer> l = InnerProductLayer::create(axis); |
||||
l->setParamsFrom(params); |
||||
l->blobs[0].reshape(Shape(numOutputs, innerSize)); |
||||
if (bias) |
||||
l->blobs[1].reshape(Shape(1, numOutputs)); |
||||
|
||||
return Ptr<Layer>(l); |
||||
} |
||||
|
||||
template<> //LRNLayer specialization
|
||||
Ptr<Layer> createLayerFromCaffe<LRNLayer>(LayerParams& params) |
||||
{ |
||||
int type; |
||||
String nrmType = params.get<String>("norm_region", "ACROSS_CHANNELS"); |
||||
if (nrmType == "ACROSS_CHANNELS") |
||||
type = LRNLayer::CHANNEL_NRM; |
||||
else if (nrmType == "WITHIN_CHANNEL") |
||||
type = LRNLayer::SPATIAL_NRM; |
||||
else |
||||
CV_Error(Error::StsBadArg, "Unknown region type \"" + nrmType + "\""); |
||||
|
||||
int size = params.get<int>("local_size", 5); |
||||
if (size % 2 != 1 || size <= 0) |
||||
CV_Error(Error::StsBadArg, "LRN layer supports only positive odd values for local_size"); |
||||
|
||||
double alpha = params.get<double>("alpha", 1); |
||||
double beta = params.get<double>("beta", 0.75); |
||||
|
||||
return Ptr<Layer>(LRNLayer::create(type, size, alpha, beta)); |
||||
} |
||||
|
||||
template<> |
||||
Ptr<Layer> createLayerFromCaffe<MVNLayer>(LayerParams ¶ms) |
||||
{ |
||||
return Ptr<Layer>(MVNLayer::create( |
||||
params.get<bool>("normalize_variance", true), |
||||
params.get<bool>("across_channels", false), |
||||
params.get<double>("eps", 1e-9) |
||||
)); |
||||
} |
||||
|
||||
/* Reshape layers */ |
||||
|
||||
template<> |
||||
Ptr<Layer> createLayerFromCaffe<ReshapeLayer>(LayerParams ¶ms) |
||||
{ |
||||
int axis = params.get<int>("axis", 0); |
||||
int numAxes = params.get<int>("num_axes", -1); |
||||
CV_Assert(numAxes >= -1); |
||||
Range applyingRange = (numAxes == -1) ? Range(axis, INT_MAX) : Range(axis, axis + numAxes); |
||||
|
||||
Shape newShape; |
||||
if (params.has("dim")) |
||||
{ |
||||
const DictValue ¶mShape = params.get("dim"); |
||||
newShape = Shape::all(paramShape.size()); |
||||
for (int i = 0; i < paramShape.size(); i++) |
||||
newShape[i] = paramShape.get<int>(i); |
||||
} |
||||
else |
||||
newShape = Shape::all(0); |
||||
|
||||
return Ptr<Layer>(ReshapeLayer::create(newShape, applyingRange)); |
||||
} |
||||
|
||||
Ptr<Layer> createFlattenLayerFromCaffe(LayerParams&) |
||||
{ |
||||
return Ptr<Layer>(ReshapeLayer::create(Shape(0, -1))); |
||||
} |
||||
|
||||
template<> |
||||
Ptr<Layer> createLayerFromCaffe<ConcatLayer>(LayerParams& params) |
||||
{ |
||||
return Ptr<Layer>(ConcatLayer::create(params.get<int>("axis", 1))); |
||||
} |
||||
|
||||
template<> |
||||
Ptr<Layer> createLayerFromCaffe<SplitLayer>(LayerParams ¶ms) |
||||
{ |
||||
int outputsCount; |
||||
|
||||
//TODO: maybe "top_count" param is useless because it can be determined by output connections number
|
||||
if (params.has("top_count")) |
||||
{ |
||||
outputsCount = params.get<int>("top_count"); |
||||
CV_Assert(outputsCount >= 0); |
||||
} |
||||
else |
||||
{ |
||||
outputsCount = -1; |
||||
} |
||||
|
||||
return Ptr<Layer>(SplitLayer::create(outputsCount)); |
||||
} |
||||
|
||||
template<> |
||||
Ptr<Layer> createLayerFromCaffe<SliceLayer>(LayerParams& params) |
||||
{ |
||||
int axis = params.get<int>("axis", 1); |
||||
|
||||
if (!params.has("slice_point")) |
||||
{ |
||||
return Ptr<Layer>(SliceLayer::create(axis)); |
||||
} |
||||
else |
||||
{ |
||||
const DictValue &indicesValue = params.get("slice_point"); |
||||
std::vector<int> sliceIndices(indicesValue.size()); |
||||
for (int i = 0; i < indicesValue.size(); i++) |
||||
sliceIndices[i] = indicesValue.get<int>(i); |
||||
|
||||
return Ptr<Layer>(SliceLayer::create(axis, sliceIndices)); |
||||
} |
||||
} |
||||
|
||||
/* Activation layers */ |
||||
|
||||
template <typename ActivationLayer> //Intended for parameters-free activations
|
||||
Ptr<Layer> createLayerFromCaffe(LayerParams&) |
||||
{ |
||||
return Ptr<Layer>(ActivationLayer::create()); |
||||
} |
||||
|
||||
template<> //ReLU specialization
|
||||
Ptr<Layer> createLayerFromCaffe<ReLULayer>(LayerParams& params) |
||||
{ |
||||
float negative_slope = params.get<float>("negative_slope", 0.f); |
||||
return Ptr<Layer>(ReLULayer::create(negative_slope)); |
||||
} |
||||
|
||||
template<> //Power specialization
|
||||
Ptr<Layer> createLayerFromCaffe<PowerLayer>(LayerParams& params) |
||||
{ |
||||
float power = params.get<float>("power", 1.0f); |
||||
float scale = params.get<float>("scale", 1.0f); |
||||
float shift = params.get<float>("shift", 0.0f); |
||||
return Ptr<Layer>(PowerLayer::create(power, scale, shift)); |
||||
} |
||||
|
||||
//Explicit instantiation
|
||||
template Ptr<Layer> createLayerFromCaffe<ConvolutionLayer>(LayerParams&); |
||||
template Ptr<Layer> createLayerFromCaffe<DeconvolutionLayer>(LayerParams&); |
||||
template Ptr<Layer> createLayerFromCaffe<SoftmaxLayer>(LayerParams&); |
||||
template Ptr<Layer> createLayerFromCaffe<InnerProductLayer>(LayerParams&); |
||||
template Ptr<Layer> createLayerFromCaffe<LRNLayer>(LayerParams&); |
||||
template Ptr<Layer> createLayerFromCaffe<MVNLayer>(LayerParams&); |
||||
|
||||
template Ptr<Layer> createLayerFromCaffe<ConcatLayer>(LayerParams&); |
||||
template Ptr<Layer> createLayerFromCaffe<SliceLayer>(LayerParams&); |
||||
template Ptr<Layer> createLayerFromCaffe<SplitLayer>(LayerParams&); |
||||
|
||||
template Ptr<Layer> createLayerFromCaffe<ReLULayer>(LayerParams&); |
||||
template Ptr<Layer> createLayerFromCaffe<SigmoidLayer>(LayerParams&); |
||||
template Ptr<Layer> createLayerFromCaffe<TanHLayer>(LayerParams&); |
||||
template Ptr<Layer> createLayerFromCaffe<AbsLayer>(LayerParams&); |
||||
template Ptr<Layer> createLayerFromCaffe<BNLLLayer>(LayerParams&); |
||||
template Ptr<Layer> createLayerFromCaffe<PowerLayer>(LayerParams&); |
||||
|
||||
} |
||||
} |
||||
@ -0,0 +1,60 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's 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.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// 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 Intel Corporation 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.
|
||||
//
|
||||
//M*/
|
||||
|
||||
#ifndef __OPENCV_DNN_CAFFE_LAYER_LOADERS_HPP__ |
||||
#define __OPENCV_DNN_CAFFE_LAYER_LOADERS_HPP__ |
||||
|
||||
#include <opencv2/dnn/all_layers.hpp> |
||||
|
||||
namespace cv |
||||
{ |
||||
namespace dnn |
||||
{ |
||||
|
||||
//Common template for Caffe layer loaders
|
||||
template <typename PublicLayer> |
||||
Ptr<Layer> createLayerFromCaffe(LayerParams&); |
||||
|
||||
Ptr<Layer> createFlattenLayerFromCaffe(LayerParams&); |
||||
|
||||
} |
||||
} |
||||
#endif |
||||
@ -0,0 +1,46 @@ |
||||
#include "../precomp.hpp" |
||||
#include "elementwise_layers.hpp" |
||||
|
||||
namespace cv |
||||
{ |
||||
namespace dnn |
||||
{ |
||||
|
||||
#define ACTIVATION_CREATOR_FOR(_Layer, _Functor, ...) \ |
||||
Ptr<_Layer> _Layer::create() { \
|
||||
return return Ptr<_Layer>( new ElementWiseLayer<_Functor>(_Functor()) ); } |
||||
|
||||
|
||||
Ptr<ReLULayer> ReLULayer::create(double negativeSlope) |
||||
{ |
||||
return Ptr<ReLULayer>(new ElementWiseLayer<ReLUFunctor>(ReLUFunctor(negativeSlope))); |
||||
} |
||||
|
||||
Ptr<TanHLayer> TanHLayer::create() |
||||
{ |
||||
return Ptr<TanHLayer>(new ElementWiseLayer<TanHFunctor>()); |
||||
} |
||||
|
||||
Ptr<SigmoidLayer> SigmoidLayer::create() |
||||
{ |
||||
return Ptr<SigmoidLayer>(new ElementWiseLayer<SigmoidFunctor>()); |
||||
} |
||||
|
||||
Ptr<AbsLayer> AbsLayer::create() |
||||
{ |
||||
return Ptr<AbsLayer>(new ElementWiseLayer<AbsValFunctor>()); |
||||
} |
||||
|
||||
Ptr<BNLLLayer> BNLLLayer::create() |
||||
{ |
||||
return Ptr<BNLLLayer>(new ElementWiseLayer<BNLLFunctor>()); |
||||
} |
||||
|
||||
Ptr<PowerLayer> PowerLayer::create(double power /*= 1*/, double scale /*= 1*/, double shift /*= 0*/) |
||||
{ |
||||
const PowerFunctor f(power, scale, shift); |
||||
return Ptr<PowerLayer>(new ElementWiseLayer<PowerFunctor>(f)); |
||||
} |
||||
|
||||
} |
||||
} |
||||
@ -0,0 +1,95 @@ |
||||
#include "op_blas.hpp" |
||||
|
||||
#if HAVE_CBLAS |
||||
#include "opencv_cblas.hpp" |
||||
#endif |
||||
|
||||
#include <iostream> |
||||
|
||||
namespace cv |
||||
{ |
||||
namespace dnn |
||||
{ |
||||
|
||||
void gemm(InputArray A, InputArray B, double alpha, InputOutputArray C, double beta, int flags) |
||||
{ |
||||
if (C.isMat()) |
||||
gemmCPU(A.getMat(), B.getMat(), alpha, C.getMatRef(), beta, flags); |
||||
else |
||||
{ |
||||
cv::gemm(A, B, alpha, (beta == 0) ? noArray() : C, beta, C, flags); |
||||
} |
||||
} |
||||
|
||||
inline void SwapRowCols(const Mat &A, int &rows, int &cols, bool isTrans) |
||||
{ |
||||
CV_DbgAssert(A.dims == 2); |
||||
rows = (isTrans) ? A.cols : A.rows; |
||||
cols = (isTrans) ? A.rows : A.cols; |
||||
} |
||||
|
||||
void gemmCPU(const Mat &A, const Mat &B, double alpha, Mat &C, double beta, int flags /*= 0*/) |
||||
{ |
||||
#if HAVE_CBLAS |
||||
bool transA = static_cast<bool>(flags & GEMM_1_T); |
||||
bool transB = static_cast<bool>(flags & GEMM_2_T); |
||||
bool transC = static_cast<bool>(flags & GEMM_3_T); |
||||
|
||||
int Arows, Acols, Brows, Bcols, Crows, Ccols; |
||||
SwapRowCols(A, Arows, Acols, transA); |
||||
SwapRowCols(B, Brows, Bcols, transB); |
||||
SwapRowCols(C, Crows, Ccols, transC); |
||||
|
||||
CV_Assert(!(flags & GEMM_3_T)); |
||||
CV_Assert(Acols == Brows && Arows == Crows && Bcols == Ccols); |
||||
CV_Assert(A.isContinuous() && B.isContinuous() && C.isContinuous()); |
||||
CV_Assert(A.type() == B.type() && B.type() == C.type()); |
||||
CV_Assert(A.data != C.data && B.data != C.data); |
||||
|
||||
if (C.type() == CV_32F) |
||||
{ |
||||
cblas_sgemm(CblasRowMajor, transA ? CblasTrans : CblasNoTrans, transB ? CblasTrans : CblasNoTrans, |
||||
Arows, Bcols, Acols, |
||||
(float)alpha, A.ptr<float>(), A.cols, |
||||
B.ptr<float>(), B.cols, |
||||
(float)beta, C.ptr<float>(), C.cols); |
||||
} |
||||
else if (C.type() == CV_64F) |
||||
{ |
||||
//TODO: Should be tested
|
||||
cblas_dgemm(CblasRowMajor, transA ? CblasTrans : CblasNoTrans, transB ? CblasTrans : CblasNoTrans, |
||||
Arows, Bcols, Acols, |
||||
alpha, A.ptr<double>(), A.cols, |
||||
B.ptr<double>(), B.cols, |
||||
beta, C.ptr<double>(), C.cols); |
||||
} |
||||
else |
||||
{ |
||||
CV_Error(Error::BadDepth, "Only floating point types are supported"); |
||||
} |
||||
#else |
||||
cv::gemm(A, B, alpha, C, beta, C, flags); |
||||
#endif |
||||
} |
||||
|
||||
int getBlasThreads() |
||||
{ |
||||
#ifdef OPENBLAS_VERSION |
||||
return openblas_get_num_threads(); |
||||
#else |
||||
return 1; |
||||
#endif |
||||
} |
||||
|
||||
void setBlasThreads(int numThreads) |
||||
{ |
||||
#ifdef OPENBLAS_VERSION |
||||
openblas_set_num_threads(numThreads); |
||||
goto_set_num_threads(numThreads); |
||||
#else |
||||
(void)numThreads; //suppress compilers' warning
|
||||
#endif |
||||
} |
||||
|
||||
} |
||||
} |
||||
@ -0,0 +1,231 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's 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.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// 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 Intel Corporation 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.
|
||||
//
|
||||
//M*/
|
||||
|
||||
#ifndef __OPENCV_DNN_LAYERS_IM2COL_HPP__ |
||||
#define __OPENCV_DNN_LAYERS_IM2COL_HPP__ |
||||
#include "../precomp.hpp" |
||||
#include <iostream> |
||||
|
||||
namespace cv |
||||
{ |
||||
namespace dnn |
||||
{ |
||||
|
||||
template <typename Dtype> |
||||
class im2col_CpuPBody : public cv::ParallelLoopBody |
||||
{ |
||||
const Dtype* data_im; |
||||
int channels, height, width; |
||||
int kernel_h, kernel_w; |
||||
int pad_h, pad_w; |
||||
int stride_h, stride_w; |
||||
Dtype* data_col; |
||||
int height_col, width_col, channels_col; |
||||
|
||||
im2col_CpuPBody() {} |
||||
public: |
||||
|
||||
static void run(const Dtype* data_im, |
||||
int channels, int height, int width, |
||||
int kernel_h, int kernel_w, |
||||
int pad_h, int pad_w, |
||||
int stride_h, int stride_w, |
||||
Dtype* data_col) |
||||
{ |
||||
im2col_CpuPBody<Dtype> t; |
||||
t.data_im = data_im; |
||||
t.data_col = data_col; |
||||
t.channels = channels; t.height = height; t.width = width; |
||||
t.kernel_h = kernel_h; t.kernel_w = kernel_w; |
||||
t.pad_h = pad_h; t.pad_w = pad_w; |
||||
t.stride_h = stride_h; t.stride_w = stride_w; |
||||
t.height_col = (height + 2 * pad_h - kernel_h) / stride_h + 1; |
||||
t.width_col = (width + 2 * pad_w - kernel_w) / stride_w + 1; |
||||
t.channels_col = channels * kernel_h * kernel_w; |
||||
|
||||
cv::parallel_for_(Range(0, t.channels_col), t); |
||||
} |
||||
|
||||
virtual void operator ()(const Range &r) const |
||||
{ |
||||
for (int c = r.start; c < r.end; ++c) { |
||||
int w_offset = c % kernel_w; |
||||
int h_offset = (c / kernel_w) % kernel_h; |
||||
int c_im = c / kernel_h / kernel_w; |
||||
for (int h = 0; h < height_col; ++h) { |
||||
for (int w = 0; w < width_col; ++w) { |
||||
int h_pad = h * stride_h - pad_h + h_offset; |
||||
int w_pad = w * stride_w - pad_w + w_offset; |
||||
if (h_pad >= 0 && h_pad < height && w_pad >= 0 && w_pad < width) |
||||
data_col[(c * height_col + h) * width_col + w] = |
||||
data_im[(c_im * height + h_pad) * width + w_pad]; |
||||
else |
||||
data_col[(c * height_col + h) * width_col + w] = 0; |
||||
} |
||||
} |
||||
} |
||||
} |
||||
}; |
||||
|
||||
template <typename Dtype> |
||||
class col2im_CpuPBody : public cv::ParallelLoopBody |
||||
{ |
||||
const Dtype* data_col; |
||||
int channels, height, width; |
||||
int kernel_h, kernel_w; |
||||
int pad_h, pad_w; |
||||
int stride_h, stride_w; |
||||
Dtype* data_im; |
||||
int height_col, width_col; |
||||
|
||||
col2im_CpuPBody() {} |
||||
|
||||
public: |
||||
|
||||
static void run(const Dtype* data_col, |
||||
int channels, int height, int width, |
||||
int kernel_h, int kernel_w, |
||||
int pad_h, int pad_w, |
||||
int stride_h, int stride_w, |
||||
Dtype* data_im) |
||||
{ |
||||
//TODO: single-threaded version switch
|
||||
|
||||
col2im_CpuPBody t; |
||||
t.data_col = data_col; |
||||
t.data_im = data_im; |
||||
t.channels = channels; t.height = height; t.width = width; |
||||
t.kernel_h = kernel_h; t.kernel_w = kernel_w; |
||||
t.pad_h = pad_h; t.pad_w = pad_w; |
||||
t.stride_h = stride_h; t.stride_w = stride_w; |
||||
t.height_col = (height + 2 * pad_h - kernel_h) / stride_h + 1; |
||||
t.width_col = (width + 2 * pad_w - kernel_w) / stride_w + 1; |
||||
int img_total = channels * height * width; |
||||
|
||||
cv::parallel_for_(Range(0, img_total), t); |
||||
} |
||||
|
||||
virtual void operator ()(const Range &r) const |
||||
{ |
||||
for (int index = r.start; index < r.end; index++) |
||||
{ |
||||
Dtype val = 0; |
||||
int w = index % width + pad_w; |
||||
int h = (index / width) % height + pad_h; |
||||
int c = index / (width * height); |
||||
|
||||
// compute the start and end of the output
|
||||
int w_col_start = (w < kernel_w) ? 0 : (w - kernel_w) / stride_w + 1; |
||||
int w_col_end = std::min(w / stride_w + 1, width_col); |
||||
int h_col_start = (h < kernel_h) ? 0 : (h - kernel_h) / stride_h + 1; |
||||
int h_col_end = std::min(h / stride_h + 1, height_col); |
||||
|
||||
// equivalent implementation
|
||||
int offset = |
||||
(c * kernel_h * kernel_w + h * kernel_w + w) * height_col * width_col; |
||||
int coeff_h_col = (1 - stride_h * kernel_w * height_col) * width_col; |
||||
int coeff_w_col = (1 - stride_w * height_col * width_col); |
||||
for (int h_col = h_col_start; h_col < h_col_end; ++h_col) { |
||||
for (int w_col = w_col_start; w_col < w_col_end; ++w_col) { |
||||
val += data_col[offset + h_col * coeff_h_col + w_col * coeff_w_col]; |
||||
} |
||||
} |
||||
data_im[index] = val; |
||||
} |
||||
} |
||||
}; |
||||
|
||||
//single-threaded version
|
||||
template <typename Dtype> |
||||
void col2im_cpu(const Dtype* data_col, |
||||
int channels, int height, int width, |
||||
int kernel_h, int kernel_w, |
||||
int pad_h, int pad_w, |
||||
int stride_h, int stride_w, |
||||
Dtype* data_im) |
||||
{ |
||||
int height_col = (height + 2 * pad_h - kernel_h) / stride_h + 1; |
||||
int width_col = (width + 2 * pad_w - kernel_w) / stride_w + 1; |
||||
int channels_col = channels * kernel_h * kernel_w; |
||||
|
||||
std::memset(data_im, 0, height * width * channels * sizeof(Dtype)); |
||||
|
||||
for (int c = 0; c < channels_col; ++c) |
||||
{ |
||||
int w_offset = c % kernel_w; |
||||
int h_offset = (c / kernel_w) % kernel_h; |
||||
int c_im = c / kernel_h / kernel_w; |
||||
|
||||
for (int h = 0; h < height_col; ++h) |
||||
{ |
||||
for (int w = 0; w < width_col; ++w) |
||||
{ |
||||
int h_pad = h * stride_h - pad_h + h_offset; |
||||
int w_pad = w * stride_w - pad_w + w_offset; |
||||
|
||||
if (h_pad >= 0 && h_pad < height && w_pad >= 0 && w_pad < width) |
||||
data_im[(c_im * height + h_pad) * width + w_pad] += |
||||
data_col[(c * height_col + h) * width_col + w]; |
||||
} |
||||
} |
||||
} |
||||
} |
||||
|
||||
#ifdef HAVE_OPENCL |
||||
bool im2col_ocl(const UMat &img, |
||||
int channels, int height, int width, |
||||
int kernel_h, int kernel_w, |
||||
int pad_h, int pad_w, |
||||
int stride_h, int stride_w, |
||||
UMat &col); |
||||
|
||||
bool col2im_ocl(const UMat &col, |
||||
int channels, int height, int width, |
||||
int kernel_h, int kernel_w, |
||||
int pad_h, int pad_w, |
||||
int stride_h, int stride_w, |
||||
UMat &img); |
||||
#endif |
||||
|
||||
} |
||||
} |
||||
|
||||
#endif |
||||
@ -0,0 +1,440 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's 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.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// 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 Intel Corporation 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.
|
||||
//
|
||||
//M*/
|
||||
|
||||
#include "../precomp.hpp" |
||||
#include "recurrent_layers.hpp" |
||||
#include "op_blas.hpp" |
||||
#include <iostream> |
||||
#include <cmath> |
||||
|
||||
namespace cv |
||||
{ |
||||
namespace dnn |
||||
{ |
||||
|
||||
template<typename Dtype> |
||||
static void tanh(const Mat &src, Mat &dst) |
||||
{ |
||||
MatConstIterator_<Dtype> itSrc = src.begin<Dtype>(); |
||||
MatIterator_<Dtype> itDst = dst.begin<Dtype>(); |
||||
|
||||
for (; itSrc != src.end<Dtype>(); itSrc++, itDst++) |
||||
*itDst = std::tanh(*itSrc); |
||||
} |
||||
|
||||
static void tanh(const Mat &src, Mat &dst) |
||||
{ |
||||
dst.create(src.dims, (const int*)src.size, src.type()); |
||||
|
||||
if (src.type() == CV_32F) |
||||
tanh<float>(src, dst); |
||||
else if (src.type() == CV_64F) |
||||
tanh<double>(src, dst); |
||||
else |
||||
CV_Error(Error::StsUnsupportedFormat, "Function supports only floating point types"); |
||||
} |
||||
|
||||
static void sigmoid(const Mat &src, Mat &dst) |
||||
{ |
||||
cv::exp(-src, dst); |
||||
cv::pow(1 + dst, -1, dst); |
||||
} |
||||
|
||||
class LSTMLayerImpl : public LSTMLayer |
||||
{ |
||||
int numOut, numTimeStamps, numSamples, numInp; |
||||
Mat hInternal, cInternal; |
||||
Mat gates, dummyOnes; |
||||
int dtype; |
||||
bool allocated; |
||||
|
||||
BlobShape outTailShape; //shape of single output sample
|
||||
BlobShape outTsMatShape, outTsShape; //shape of N output samples
|
||||
BlobShape outResShape; //shape of T timestamps and N output samples
|
||||
|
||||
bool useTimestampDim; |
||||
bool produceCellOutput; |
||||
|
||||
public: |
||||
|
||||
LSTMLayerImpl() |
||||
{ |
||||
type = "LSTM"; |
||||
useTimestampDim = true; |
||||
produceCellOutput = false; |
||||
allocated = false; |
||||
outTailShape = BlobShape::empty(); |
||||
} |
||||
|
||||
void setUseTimstampsDim(bool use) |
||||
{ |
||||
CV_Assert(!allocated); |
||||
useTimestampDim = use; |
||||
} |
||||
|
||||
void setProduceCellOutput(bool produce) |
||||
{ |
||||
CV_Assert(!allocated); |
||||
produceCellOutput = produce; |
||||
} |
||||
|
||||
void setC(const Blob &C) |
||||
{ |
||||
CV_Assert(cInternal.empty() || C.total() == cInternal.total()); |
||||
if (!cInternal.empty()) |
||||
C.reshaped(BlobShape::like(cInternal)).matRefConst().copyTo(cInternal); |
||||
else |
||||
C.matRefConst().copyTo(cInternal); |
||||
} |
||||
|
||||
void setH(const Blob &H) |
||||
{ |
||||
CV_Assert(hInternal.empty() || H.total() == hInternal.total()); |
||||
if (!hInternal.empty()) |
||||
H.reshaped(BlobShape::like(hInternal)).matRefConst().copyTo(hInternal); |
||||
else |
||||
H.matRefConst().copyTo(hInternal); |
||||
} |
||||
|
||||
Blob getC() const |
||||
{ |
||||
CV_Assert(!cInternal.empty()); |
||||
|
||||
//TODO: add convinient Mat -> Blob constructor
|
||||
Blob res(outTsShape, cInternal.type()); |
||||
res.fill(res.shape(), res.type(), cInternal.data); |
||||
return res; |
||||
} |
||||
|
||||
Blob getH() const |
||||
{ |
||||
CV_Assert(!hInternal.empty()); |
||||
|
||||
Blob res(outTsShape, hInternal.type()); |
||||
res.fill(res.shape(), res.type(), hInternal.data); |
||||
return res; |
||||
} |
||||
|
||||
void setOutShape(const BlobShape &outTailShape_) |
||||
{ |
||||
CV_Assert(!allocated || outTailShape_.total() == outTailShape.total()); |
||||
outTailShape = outTailShape_; |
||||
} |
||||
|
||||
void setWeights(const Blob &Wh, const Blob &Wx, const Blob &bias) |
||||
{ |
||||
CV_Assert(Wh.dims() == 2 && Wx.dims() == 2); |
||||
CV_Assert(Wh.size(0) == Wx.size(0)); |
||||
CV_Assert(Wh.size(0) == 4*Wh.size(1)); |
||||
CV_Assert(Wh.size(0) == (int)bias.total()); |
||||
CV_Assert(Wh.type() == Wx.type() && Wx.type() == bias.type()); |
||||
|
||||
blobs.resize(3); |
||||
blobs[0] = Wh; |
||||
blobs[1] = Wx; |
||||
blobs[2] = bias; |
||||
blobs[2].reshape(BlobShape(1, (int)bias.total())); |
||||
} |
||||
|
||||
void allocate(const std::vector<Blob*> &input, std::vector<Blob> &output) |
||||
{ |
||||
CV_Assert(blobs.size() == 3); |
||||
CV_Assert(input.size() == 1); |
||||
|
||||
Blob &Wh = blobs[0], &Wx = blobs[1]; |
||||
numOut = Wh.size(1); |
||||
numInp = Wx.size(1); |
||||
|
||||
if (!outTailShape.isEmpty()) |
||||
CV_Assert(outTailShape.total() == numOut); |
||||
else |
||||
outTailShape = BlobShape(numOut); |
||||
|
||||
if (useTimestampDim) |
||||
{ |
||||
CV_Assert(input[0]->dims() >= 2 && (int)input[0]->total(2) == numInp); |
||||
numTimeStamps = input[0]->size(0); |
||||
numSamples = input[0]->size(1); |
||||
outResShape = BlobShape(numTimeStamps, numSamples) + outTailShape; |
||||
} |
||||
else |
||||
{ |
||||
CV_Assert(input[0]->dims() >= 1 && (int)input[0]->total(1) == numInp); |
||||
numTimeStamps = 1; |
||||
numSamples = input[0]->size(0); |
||||
outResShape = BlobShape(numSamples) + outTailShape; |
||||
} |
||||
outTsMatShape = BlobShape(numSamples, numOut); |
||||
outTsShape = BlobShape(numSamples) + outTailShape; |
||||
|
||||
dtype = input[0]->type(); |
||||
CV_Assert(dtype == CV_32F || dtype == CV_64F); |
||||
CV_Assert(Wh.type() == dtype); |
||||
|
||||
output.resize( (produceCellOutput) ? 2 : 1 ); |
||||
output[0].create(outResShape, dtype); |
||||
if (produceCellOutput) |
||||
output[1].create(outResShape, dtype); |
||||
|
||||
if (hInternal.empty()) |
||||
{ |
||||
hInternal.create(outTsMatShape.dims(), outTsMatShape.ptr(), dtype); |
||||
hInternal.setTo(0); |
||||
} |
||||
else |
||||
{ |
||||
CV_Assert((int)hInternal.total() == numSamples*numOut); |
||||
hInternal = hInternal.reshape(1, outTsMatShape.dims(), outTsMatShape.ptr()); |
||||
} |
||||
|
||||
if (cInternal.empty()) |
||||
{ |
||||
cInternal.create(outTsMatShape.dims(), outTsMatShape.ptr(), dtype); |
||||
cInternal.setTo(0); |
||||
} |
||||
else |
||||
{ |
||||
CV_Assert((int)cInternal.total() == numSamples*numOut); |
||||
cInternal = cInternal.reshape(1, outTsMatShape.dims(), outTsMatShape.ptr()); |
||||
} |
||||
|
||||
gates.create(numSamples, 4*numOut, dtype); |
||||
|
||||
dummyOnes.create(numSamples, 1, dtype); |
||||
dummyOnes.setTo(1); |
||||
|
||||
allocated = true; |
||||
} |
||||
|
||||
void forward(std::vector<Blob*> &input, std::vector<Blob> &output) |
||||
{ |
||||
const Mat &Wh = blobs[0].matRefConst(); |
||||
const Mat &Wx = blobs[1].matRefConst(); |
||||
const Mat &bias = blobs[2].matRefConst(); |
||||
|
||||
int numSamplesTotal = numTimeStamps*numSamples; |
||||
Mat xTs = input[0]->reshaped(BlobShape(numSamplesTotal, numInp)).matRefConst(); |
||||
|
||||
BlobShape outMatShape(numSamplesTotal, numOut); |
||||
Mat hOutTs = output[0].reshaped(outMatShape).matRef(); |
||||
Mat cOutTs = (produceCellOutput) ? output[1].reshaped(outMatShape).matRef() : Mat(); |
||||
|
||||
for (int ts = 0; ts < numTimeStamps; ts++) |
||||
{ |
||||
Range curRowRange(ts*numSamples, (ts + 1)*numSamples); |
||||
Mat xCurr = xTs.rowRange(curRowRange); |
||||
|
||||
gemmCPU(xCurr, Wx, 1, gates, 0, GEMM_2_T); // Wx * x_t
|
||||
gemmCPU(hInternal, Wh, 1, gates, 1, GEMM_2_T); //+Wh * h_{t-1}
|
||||
gemmCPU(dummyOnes, bias, 1, gates, 1); //+b
|
||||
|
||||
Mat getesIFO = gates.colRange(0, 3*numOut); |
||||
Mat gateI = gates.colRange(0*numOut, 1*numOut); |
||||
Mat gateF = gates.colRange(1*numOut, 2*numOut); |
||||
Mat gateO = gates.colRange(2*numOut, 3*numOut); |
||||
Mat gateG = gates.colRange(3*numOut, 4*numOut); |
||||
|
||||
sigmoid(getesIFO, getesIFO); |
||||
tanh(gateG, gateG); |
||||
|
||||
//compute c_t
|
||||
cv::multiply(gateF, cInternal, gateF); // f_t (*) c_{t-1}
|
||||
cv::multiply(gateI, gateG, gateI); // i_t (*) g_t
|
||||
cv::add(gateF, gateI, cInternal); // c_t = f_t (*) c_{t-1} + i_t (*) g_t
|
||||
|
||||
//compute h_t
|
||||
tanh(cInternal, hInternal); |
||||
cv::multiply(gateO, hInternal, hInternal); |
||||
|
||||
//save results in output blobs
|
||||
hInternal.copyTo(hOutTs.rowRange(curRowRange)); |
||||
if (produceCellOutput) |
||||
cInternal.copyTo(cOutTs.rowRange(curRowRange)); |
||||
} |
||||
} |
||||
}; |
||||
|
||||
Ptr<LSTMLayer> LSTMLayer::create() |
||||
{ |
||||
return Ptr<LSTMLayer>(new LSTMLayerImpl()); |
||||
} |
||||
|
||||
void LSTMLayer::forward(std::vector<Blob*>&, std::vector<Blob>&) |
||||
{ |
||||
CV_Error(Error::StsInternal, "This function should be unreached"); |
||||
} |
||||
|
||||
int LSTMLayer::inputNameToIndex(String inputName) |
||||
{ |
||||
if (inputName.toLowerCase() == "x") |
||||
return 0; |
||||
return -1; |
||||
} |
||||
|
||||
int LSTMLayer::outputNameToIndex(String outputName) |
||||
{ |
||||
if (outputName.toLowerCase() == "h") |
||||
return 0; |
||||
else if (outputName.toLowerCase() == "c") |
||||
return 1; |
||||
return -1; |
||||
} |
||||
|
||||
|
||||
class RNNLayerImpl : public RNNLayer |
||||
{ |
||||
int numX, numH, numO; |
||||
int numSamples, numTimestamps, numSamplesTotal; |
||||
int dtype; |
||||
Mat Whh, Wxh, bh; |
||||
Mat Who, bo; |
||||
Mat hCurr, hPrev, dummyBiasOnes; |
||||
bool produceH; |
||||
|
||||
public: |
||||
|
||||
RNNLayerImpl() |
||||
{ |
||||
type = "RNN"; |
||||
produceH = false; |
||||
} |
||||
|
||||
void setProduceHiddenOutput(bool produce = false) |
||||
{ |
||||
produceH = produce; |
||||
} |
||||
|
||||
void setWeights(const Blob &W_xh, const Blob &b_h, const Blob &W_hh, const Blob &W_ho, const Blob &b_o) |
||||
{ |
||||
CV_Assert(W_hh.dims() == 2 && W_xh.dims() == 2); |
||||
CV_Assert(W_hh.size(0) == W_xh.size(0) && W_hh.size(0) == W_hh.size(1) && (int)b_h.total() == W_xh.size(0)); |
||||
CV_Assert(W_ho.size(0) == (int)b_o.total()); |
||||
CV_Assert(W_ho.size(1) == W_hh.size(1)); |
||||
|
||||
blobs.resize(5); |
||||
blobs[0] = W_xh; |
||||
blobs[1] = b_h; |
||||
blobs[2] = W_hh; |
||||
blobs[3] = W_ho; |
||||
blobs[4] = b_o; |
||||
} |
||||
|
||||
void allocate(const std::vector<Blob*> &input, std::vector<Blob> &output) |
||||
{ |
||||
CV_Assert(input.size() >= 1 && input.size() <= 2); |
||||
|
||||
Wxh = blobs[0].matRefConst(); |
||||
bh = blobs[1].matRefConst(); |
||||
Whh = blobs[2].matRefConst(); |
||||
Who = blobs[3].matRefConst(); |
||||
bo = blobs[4].matRefConst(); |
||||
|
||||
numH = Wxh.rows; |
||||
numX = Wxh.cols; |
||||
numO = Who.rows; |
||||
|
||||
CV_Assert(input[0]->dims() >= 2); |
||||
CV_Assert((int)input[0]->total(2) == numX); |
||||
CV_Assert(input[0]->type() == CV_32F || input[0]->type() == CV_64F); |
||||
dtype = input[0]->type(); |
||||
numTimestamps = input[0]->size(0); |
||||
numSamples = input[0]->size(1); |
||||
numSamplesTotal = numTimestamps * numSamples; |
||||
|
||||
hCurr.create(numSamples, numH, dtype); |
||||
hPrev.create(numSamples, numH, dtype); |
||||
hPrev.setTo(0); |
||||
|
||||
dummyBiasOnes.create(numSamples, 1, dtype); |
||||
dummyBiasOnes.setTo(1); |
||||
bh = bh.reshape(1, 1); //is 1 x numH Mat
|
||||
bo = bo.reshape(1, 1); //is 1 x numO Mat
|
||||
|
||||
reshapeOutput(output); |
||||
} |
||||
|
||||
void reshapeOutput(std::vector<Blob> &output) |
||||
{ |
||||
output.resize((produceH) ? 2 : 1); |
||||
output[0].create(BlobShape(numTimestamps, numSamples, numO), dtype); |
||||
if (produceH) |
||||
output[1].create(BlobShape(numTimestamps, numSamples, numH), dtype); |
||||
} |
||||
|
||||
void forward(std::vector<Blob*> &input, std::vector<Blob> &output) |
||||
{ |
||||
Mat xTs = input[0]->reshaped(BlobShape(numSamplesTotal, numX)).matRefConst(); |
||||
Mat oTs = output[0].reshaped(BlobShape(numSamplesTotal, numO)).matRef(); |
||||
Mat hTs = (produceH) ? output[1].reshaped(BlobShape(numSamplesTotal, numH)).matRef() : Mat(); |
||||
|
||||
for (int ts = 0; ts < numTimestamps; ts++) |
||||
{ |
||||
Range curRowRange = Range(ts * numSamples, (ts + 1) * numSamples); |
||||
Mat xCurr = xTs.rowRange(curRowRange); |
||||
|
||||
gemmCPU(hPrev, Whh, 1, hCurr, 0, GEMM_2_T); // W_{hh} * h_{prev}
|
||||
gemmCPU(xCurr, Wxh, 1, hCurr, 1, GEMM_2_T); //+W_{xh} * x_{curr}
|
||||
gemmCPU(dummyBiasOnes, bh, 1, hCurr, 1); //+bh
|
||||
tanh(hCurr, hPrev); |
||||
|
||||
Mat oCurr = oTs.rowRange(curRowRange); |
||||
gemmCPU(hPrev, Who, 1, oCurr, 0, GEMM_2_T); // W_{ho} * h_{prev}
|
||||
gemmCPU(dummyBiasOnes, bo, 1, oCurr, 1); //+b_o
|
||||
tanh(oCurr, oCurr); |
||||
|
||||
if (produceH) |
||||
hPrev.copyTo(hTs.rowRange(curRowRange)); |
||||
} |
||||
} |
||||
}; |
||||
|
||||
void RNNLayer::forward(std::vector<Blob*>&, std::vector<Blob>&) |
||||
{ |
||||
CV_Error(Error::StsInternal, "This function should be unreached"); |
||||
} |
||||
|
||||
CV_EXPORTS_W Ptr<RNNLayer> RNNLayer::create() |
||||
{ |
||||
return Ptr<RNNLayer>(new RNNLayerImpl()); |
||||
} |
||||
|
||||
} |
||||
} |
||||
@ -0,0 +1,54 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistribution's of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistribution's 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.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// 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 Intel Corporation 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.
|
||||
//
|
||||
//M*/
|
||||
|
||||
#ifndef __OPENCV_DNN_LAYERS_RECURRENT_LAYERS_HPP__ |
||||
#define __OPENCV_DNN_LAYERS_RECURRENT_LAYERS_HPP__ |
||||
#include "../precomp.hpp" |
||||
#include <opencv2/dnn/all_layers.hpp> |
||||
|
||||
namespace cv |
||||
{ |
||||
namespace dnn |
||||
{ |
||||
|
||||
} |
||||
} |
||||
#endif |
||||
@ -0,0 +1,44 @@ |
||||
__kernel void ReLUForward(const int count, __global const T* in, __global T* out |
||||
#ifndef RELU_NO_SLOPE |
||||
, T negative_slope |
||||
#endif |
||||
) { |
||||
int index = get_global_id(0); |
||||
if(index < count) |
||||
#ifndef RELU_NO_SLOPE |
||||
out[index] = in[index] > 0 ? in[index] : in[index] * negative_slope; |
||||
#else |
||||
out[index] = in[index] > 0 ? in[index] : 0; |
||||
#endif |
||||
} |
||||
|
||||
__kernel void TanHForward(const int count, __global T* in, __global T* out) { |
||||
int index = get_global_id(0); |
||||
if(index < count) |
||||
out[index] = tanh(in[index]); |
||||
} |
||||
|
||||
__kernel void SigmoidForward(const int count, __global const T* in, __global T* out) { |
||||
int index = get_global_id(0); |
||||
if(index < count) |
||||
out[index] = 1. / (1. + exp(-in[index])); |
||||
} |
||||
|
||||
__kernel void BNLLForward(const int n, __global const T* in, __global T* out) { |
||||
int index = get_global_id(0); |
||||
if (index < n) { |
||||
out[index] = in[index] > 0 ? in[index] + log(1. + exp(-in[index])) : log(1. + exp(in[index])); |
||||
} |
||||
} |
||||
|
||||
__kernel void AbsValForward(const int n, __global const T* in, __global T* out) { |
||||
int index = get_global_id(0); |
||||
if (index < n) |
||||
out[index] = fabs(in[index]); |
||||
} |
||||
|
||||
__kernel void PowForward(const int n, __global const T* in, __global T* out, const T power, const T scale, const T shift) { |
||||
int index = get_global_id(0); |
||||
if (index < n) |
||||
out[index] = pow(shift + scale * in[index], power); |
||||
} |
||||
@ -0,0 +1,62 @@ |
||||
/************************************************************************************* |
||||
* 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 col2im(const int n, __global const T* data_col, const int col_offset, |
||||
const int height, const int width, const int channels, |
||||
const int patch_h, const int patch_w, |
||||
const int pad_h, const int pad_w, |
||||
const int stride_h, const int stride_w, |
||||
const int height_col, const int width_col, |
||||
__global T* data_im, const int img_offset) |
||||
{ |
||||
data_col = data_col + col_offset; |
||||
data_im = data_im + img_offset; |
||||
int index = get_global_id(0); |
||||
if(index < n) { |
||||
T val = 0; |
||||
int w = index % width + pad_w; |
||||
int h = (index / width) % height + pad_h; |
||||
int c = index / (width * height); |
||||
|
||||
// compute the start and end of the output |
||||
int w_col_start = (w < patch_w) ? 0 : (w - patch_w) / stride_w + 1; |
||||
int w_col_end = min(w / stride_w + 1, width_col); |
||||
int h_col_start = (h < patch_h) ? 0 : (h - patch_h) / stride_h + 1; |
||||
int h_col_end = min(h / stride_h + 1, height_col); |
||||
|
||||
// equivalent implementation |
||||
int offset = |
||||
(c * patch_h * patch_w + h * patch_w + w) * height_col * width_col; |
||||
int coeff_h_col = (1 - stride_h * patch_w * height_col) * width_col; |
||||
int coeff_w_col = (1 - stride_w * height_col * width_col); |
||||
for (int h_col = h_col_start; h_col < h_col_end; ++h_col) { |
||||
for (int w_col = w_col_start; w_col < w_col_end; ++w_col) { |
||||
val += data_col[offset + h_col * coeff_h_col + w_col * coeff_w_col]; |
||||
} |
||||
} |
||||
data_im[index] = val; |
||||
} |
||||
} |
||||
@ -0,0 +1,76 @@ |
||||
/************************************************************************************* |
||||
* 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 LRNComputeOutput(const int nthreads, __global T* in, __global T* scale, const T negative_beta, __global T* out) { |
||||
int index = get_global_id(0); |
||||
int tmp = get_global_size(0); |
||||
for(index; index < nthreads; index += tmp) |
||||
out[index] = in[index] * pow(scale[index], negative_beta); |
||||
} |
||||
|
||||
__kernel void LRNFillScale(const int nthreads, __global T* in, const int num, const int channels, const int height, const int width, const int size, const T alpha_over_size, const T k, __global T* scale) { |
||||
int index = get_global_id(0); |
||||
int tmp = get_global_size(0); |
||||
for(index; index < nthreads; index += tmp) { |
||||
// find out the local offset |
||||
const int w = index % width; |
||||
const int h = (index / width) % height; |
||||
const int n = index / width / height; |
||||
const int offset = (n * channels * height + h) * width + w; |
||||
const int step = height * width; |
||||
in = in + offset; |
||||
scale = scale + offset; |
||||
int head = 0; |
||||
const int pre_pad = (size - 1) / 2; |
||||
const int post_pad = size - pre_pad - 1; |
||||
T accum_scale = 0; |
||||
// fill the scale at [n, :, h, w] |
||||
// accumulate values |
||||
while (head < post_pad && head < channels) { |
||||
accum_scale += in[head * step] * in[head * step]; |
||||
++head; |
||||
} |
||||
// both add and subtract |
||||
while (head < channels) { |
||||
accum_scale += in[head * step] * in[head * step]; |
||||
if (head - size >= 0) { |
||||
accum_scale -= in[(head - size) * step] |
||||
* in[(head - size) * step]; |
||||
} |
||||
scale[(head - post_pad) * step] = k + accum_scale * alpha_over_size; |
||||
++head; |
||||
} |
||||
// subtract only |
||||
while (head < channels + post_pad) { |
||||
if (head - size >= 0) { |
||||
accum_scale -= in[(head - size) * step] |
||||
* in[(head - size) * step]; |
||||
} |
||||
scale[(head - post_pad) * step] = k + accum_scale * alpha_over_size; |
||||
++head; |
||||
} |
||||
} |
||||
} |
||||
@ -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; |
||||
} |
||||
|
||||
} |
||||
@ -0,0 +1,75 @@ |
||||
/************************************************************************************* |
||||
* 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 kernel_channel_max(const int num, const int channels, |
||||
const int spatial_dim, __global const T* data, __global T* out) { |
||||
int index = get_global_id(0); |
||||
if(index < num * spatial_dim) { |
||||
int n = index / spatial_dim; |
||||
int s = index % spatial_dim; |
||||
T maxval = -FLT_MAX; |
||||
for (int c = 0; c < channels; ++c) { |
||||
maxval = max(data[(n * channels + c) * spatial_dim + s], maxval); |
||||
} |
||||
out[index] = maxval; |
||||
} |
||||
} |
||||
|
||||
__kernel void kernel_channel_subtract(const int count, |
||||
const int num, const int channels, |
||||
const int spatial_dim, __global const T* channel_max, __global T* data) { |
||||
int index = get_global_id(0); |
||||
if(index < count) { |
||||
int n = index / channels / spatial_dim; |
||||
int s = index % spatial_dim; |
||||
data[index] -= channel_max[n * spatial_dim + s]; |
||||
} |
||||
} |
||||
|
||||
__kernel void kernel_channel_sum(const int num, const int channels, |
||||
const int spatial_dim, __global const T* data, __global T* channel_sum) { |
||||
int index = get_global_id(0); |
||||
if(index < num * spatial_dim) { |
||||
int n = index / spatial_dim; |
||||
int s = index % spatial_dim; |
||||
T sum = 0; |
||||
for (int c = 0; c < channels; ++c) { |
||||
sum += data[(n * channels + c) * spatial_dim + s]; |
||||
} |
||||
channel_sum[index] = sum; |
||||
} |
||||
} |
||||
|
||||
__kernel void kernel_channel_div(const int count, |
||||
const int num, const int channels, |
||||
const int spatial_dim, __global const T* channel_sum, __global T* data) { |
||||
int index = get_global_id(0); |
||||
if(index < count) { |
||||
int n = index / channels / spatial_dim; |
||||
int s = index % spatial_dim; |
||||
data[index] /= channel_sum[n * spatial_dim + s]; |
||||
} |
||||
} |
||||
@ -1,3 +1,31 @@ |
||||
#include "test_precomp.hpp" |
||||
|
||||
CV_TEST_MAIN("") |
||||
|
||||
namespace cvtest |
||||
{ |
||||
|
||||
using namespace cv; |
||||
using namespace cv::dnn; |
||||
|
||||
TEST(BlobShape_SimpleConstr, Regression) |
||||
{ |
||||
BlobShape sd; |
||||
|
||||
BlobShape s1(0); |
||||
EXPECT_EQ(s1.dims(), 1); |
||||
EXPECT_EQ(s1[0], 0); |
||||
|
||||
BlobShape s2(0, 0); |
||||
EXPECT_EQ(s2.dims(), 2); |
||||
EXPECT_EQ(s2[0], 0); |
||||
EXPECT_EQ(s2[1], 0); |
||||
} |
||||
|
||||
TEST(BlobShape_EmptyFill, Regression) |
||||
{ |
||||
BlobShape s(10, (int*)NULL); |
||||
EXPECT_EQ(s.dims(), 10); |
||||
} |
||||
|
||||
} |
||||
|
||||
@ -0,0 +1 @@ |
||||
*.caffemodel |
||||
Loading…
Reference in new issue