Commit Graph

600 Commits

Author SHA1 Message Date
Dmitry Kurtaev
2a21c10837 Fix TensorFlow split layer 2017-10-02 22:44:42 +03:00
pengli
e340ff9c3a Merge pull request #9114 from pengli:dnn_rebase
add libdnn acceleration to dnn module  (#9114)

* import libdnn code

Signed-off-by: Li Peng <peng.li@intel.com>

* add convolution layer ocl acceleration

Signed-off-by: Li Peng <peng.li@intel.com>

* add pooling layer ocl acceleration

Signed-off-by: Li Peng <peng.li@intel.com>

* add softmax layer ocl acceleration

Signed-off-by: Li Peng <peng.li@intel.com>

* add lrn layer ocl acceleration

Signed-off-by: Li Peng <peng.li@intel.com>

* add innerproduct layer ocl acceleration

Signed-off-by: Li Peng <peng.li@intel.com>

* add HAVE_OPENCL macro

Signed-off-by: Li Peng <peng.li@intel.com>

* fix for convolution ocl

Signed-off-by: Li Peng <peng.li@intel.com>

* enable getUMat() for multi-dimension Mat

Signed-off-by: Li Peng <peng.li@intel.com>

* use getUMat for ocl acceleration

Signed-off-by: Li Peng <peng.li@intel.com>

* use CV_OCL_RUN macro

Signed-off-by: Li Peng <peng.li@intel.com>

* set OPENCL target when it is available

and disable fuseLayer for OCL target for the time being

Signed-off-by: Li Peng <peng.li@intel.com>

* fix innerproduct accuracy test

Signed-off-by: Li Peng <peng.li@intel.com>

* remove trailing space

Signed-off-by: Li Peng <peng.li@intel.com>

* Fixed tensorflow demo bug.

Root cause is that tensorflow has different algorithm with libdnn
to calculate convolution output dimension.

libdnn don't calculate output dimension anymore and just use one
passed in by config.

* split gemm ocl file

split it into gemm_buffer.cl and gemm_image.cl

Signed-off-by: Li Peng <peng.li@intel.com>

* Fix compile failure

Signed-off-by: Li Peng <peng.li@intel.com>

* check env flag for auto tuning

Signed-off-by: Li Peng <peng.li@intel.com>

* switch to new ocl kernels for softmax layer

Signed-off-by: Li Peng <peng.li@intel.com>

* update softmax layer

on some platform subgroup extension may not work well,
fallback to non subgroup ocl acceleration.

Signed-off-by: Li Peng <peng.li@intel.com>

* fallback to cpu path for fc layer with multi output

Signed-off-by: Li Peng <peng.li@intel.com>

* update output message

Signed-off-by: Li Peng <peng.li@intel.com>

* update fully connected layer

fallback to gemm API if libdnn return false

Signed-off-by: Li Peng <peng.li@intel.com>

* Add ReLU OCL implementation

* disable layer fusion for now

Signed-off-by: Li Peng <peng.li@intel.com>

* Add OCL implementation for concat layer

Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com>

* libdnn: update license and copyrights

Also refine libdnn coding style

Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com>
Signed-off-by: Li Peng <peng.li@intel.com>

* DNN: Don't link OpenCL library explicitly

* DNN: Make default preferableTarget to DNN_TARGET_CPU

User should set it to DNN_TARGET_OPENCL explicitly if want to
use OpenCL acceleration.

Also don't fusion when using DNN_TARGET_OPENCL

* DNN: refine coding style

* Add getOpenCLErrorString

* DNN: Use int32_t/uint32_t instread of alias

* Use namespace ocl4dnn to include libdnn things

* remove extra copyTo in softmax ocl path

Signed-off-by: Li Peng <peng.li@intel.com>

* update ReLU layer ocl path

Signed-off-by: Li Peng <peng.li@intel.com>

* Add prefer target property for layer class

It is used to indicate the target for layer forwarding,
either the default CPU target or OCL target.

Signed-off-by: Li Peng <peng.li@intel.com>

* Add cl_event based timer for cv::ocl

* Rename libdnn to ocl4dnn

Signed-off-by: Li Peng <peng.li@intel.com>
Signed-off-by: wzw <zhiwen.wu@intel.com>

* use UMat for ocl4dnn internal buffer

Remove allocateMemory which use clCreateBuffer directly

Signed-off-by: Li Peng <peng.li@intel.com>
Signed-off-by: wzw <zhiwen.wu@intel.com>

* enable buffer gemm in ocl4dnn innerproduct

Signed-off-by: Li Peng <peng.li@intel.com>

* replace int_tp globally for ocl4dnn kernels.

Signed-off-by: wzw <zhiwen.wu@intel.com>
Signed-off-by: Li Peng <peng.li@intel.com>

* create UMat for layer params

Signed-off-by: Li Peng <peng.li@intel.com>

* update sign ocl kernel

Signed-off-by: Li Peng <peng.li@intel.com>

* update image based gemm of inner product layer

Signed-off-by: Li Peng <peng.li@intel.com>

* remove buffer gemm of inner product layer

call cv::gemm API instead

Signed-off-by: Li Peng <peng.li@intel.com>

* change ocl4dnn forward parameter to UMat

Signed-off-by: Li Peng <peng.li@intel.com>

* Refine auto-tuning mechanism.

- Use OPENCV_OCL4DNN_KERNEL_CONFIG_PATH to set cache directory
  for fine-tuned kernel configuration.
  e.g. export OPENCV_OCL4DNN_KERNEL_CONFIG_PATH=/home/tmp,
  the cache directory will be /home/tmp/spatialkernels/ on Linux.

- Define environment OPENCV_OCL4DNN_ENABLE_AUTO_TUNING to enable
  auto-tuning.

- OPENCV_OPENCL_ENABLE_PROFILING is only used to enable profiling
  for OpenCL command queue. This fix basic kernel get wrong running
  time, i.e. 0ms.

- If creating cache directory failed, disable auto-tuning.

* Detect and create cache dir on windows

Signed-off-by: Li Peng <peng.li@intel.com>

* Refine gemm like convolution kernel.

Signed-off-by: Li Peng <peng.li@intel.com>

* Fix redundant swizzleWeights calling when use cached kernel config.

* Fix "out of resource" bug when auto-tuning too many kernels.

* replace cl_mem with UMat in ocl4dnnConvSpatial class

* OCL4DNN: reduce the tuning kernel candidate.

This patch could reduce 75% of the tuning candidates with less
than 2% performance impact for the final result.

Signed-off-by: Zhigang Gong <zhigang.gong@intel.com>

* replace cl_mem with umat in ocl4dnn convolution

Signed-off-by: Li Peng <peng.li@intel.com>

* remove weight_image_ of ocl4dnn inner product

Actually it is unused in the computation

Signed-off-by: Li Peng <peng.li@intel.com>

* Various fixes for ocl4dnn

1. OCL_PERFORMANCE_CHECK(ocl::Device::getDefault().isIntel())
2. Ptr<OCL4DNNInnerProduct<float> > innerProductOp
3. Code comments cleanup
4. ignore check on OCL cpu device

Signed-off-by: Li Peng <peng.li@intel.com>

* add build option for log softmax

Signed-off-by: Li Peng <peng.li@intel.com>

* remove unused ocl kernels in ocl4dnn

Signed-off-by: Li Peng <peng.li@intel.com>

* replace ocl4dnnSet with opencv setTo

Signed-off-by: Li Peng <peng.li@intel.com>

* replace ALIGN with cv::alignSize

Signed-off-by: Li Peng <peng.li@intel.com>

* check kernel build options

Signed-off-by: Li Peng <peng.li@intel.com>

* Handle program compilation fail properly.

* Use std::numeric_limits<float>::infinity() for large float number

* check ocl4dnn kernel compilation result

Signed-off-by: Li Peng <peng.li@intel.com>

* remove unused ctx_id

Signed-off-by: Li Peng <peng.li@intel.com>

* change clEnqueueNDRangeKernel to kernel.run()

Signed-off-by: Li Peng <peng.li@intel.com>

* change cl_mem to UMat in image based gemm

Signed-off-by: Li Peng <peng.li@intel.com>

* check intel subgroup support for lrn and pooling layer

Signed-off-by: Li Peng <peng.li@intel.com>

* Fix convolution bug if group is greater than 1

Signed-off-by: Li Peng <peng.li@intel.com>

* Set default layer preferableTarget to be DNN_TARGET_CPU

Signed-off-by: Li Peng <peng.li@intel.com>

* Add ocl perf test for convolution

Signed-off-by: Li Peng <peng.li@intel.com>

* Add more ocl accuracy test

Signed-off-by: Li Peng <peng.li@intel.com>

* replace cl_image with ocl::Image2D

Signed-off-by: Li Peng <peng.li@intel.com>

* Fix build failure in elementwise layer

Signed-off-by: Li Peng <peng.li@intel.com>

* use getUMat() to get blob data

Signed-off-by: Li Peng <peng.li@intel.com>

* replace cl_mem handle with ocl::KernelArg

Signed-off-by: Li Peng <peng.li@intel.com>

* dnn(build): don't use C++11, OPENCL_LIBRARIES fix

* dnn(ocl4dnn): remove unused OpenCL kernels

* dnn(ocl4dnn): extract OpenCL code into .cl files

* dnn(ocl4dnn): refine auto-tuning

Defaultly disable auto-tuning, set OPENCV_OCL4DNN_ENABLE_AUTO_TUNING
environment variable to enable it.

Use a set of pre-tuned configs as default config if auto-tuning is disabled.
These configs are tuned for Intel GPU with 48/72 EUs, and for googlenet,
AlexNet, ResNet-50

If default config is not suitable, use the first available kernel config
from the candidates. Candidate priority from high to low is gemm like kernel,
IDLF kernel, basick kernel.

* dnn(ocl4dnn): pooling doesn't use OpenCL subgroups

* dnn(ocl4dnn): fix perf test

OpenCV has default 3sec time limit for each performance test.
Warmup OpenCL backend outside of perf measurement loop.

* use ocl::KernelArg as much as possible

Signed-off-by: Li Peng <peng.li@intel.com>

* dnn(ocl4dnn): fix bias bug for gemm like kernel

* dnn(ocl4dnn): wrap cl_mem into UMat

Signed-off-by: Li Peng <peng.li@intel.com>

* dnn(ocl4dnn): Refine signature of kernel config

- Use more readable string as signture of kernel config
- Don't count device name and vendor in signature string
- Default kernel configurations are tuned for Intel GPU with
  24/48/72 EUs, and for googlenet, AlexNet, ResNet-50 net model.

* dnn(ocl4dnn): swap width/height in configuration

* dnn(ocl4dnn): enable configs for Intel OpenCL runtime only

* core: make configuration helper functions accessible from non-core modules

* dnn(ocl4dnn): update kernel auto-tuning behavior

Avoid unwanted creation of directories

* dnn(ocl4dnn): simplify kernel to workaround OpenCL compiler crash

* dnn(ocl4dnn): remove redundant code

* dnn(ocl4dnn): Add more clear message for simd size dismatch.

* dnn(ocl4dnn): add const to const argument

Signed-off-by: Li Peng <peng.li@intel.com>

* dnn(ocl4dnn): force compiler use a specific SIMD size for IDLF kernel

* dnn(ocl4dnn): drop unused tuneLocalSize()

* dnn(ocl4dnn): specify OpenCL queue for Timer and convolve() method

* dnn(ocl4dnn): sanitize file names used for cache

* dnn(perf): enable Network tests with OpenCL

* dnn(ocl4dnn/conv): drop computeGlobalSize()

* dnn(ocl4dnn/conv): drop unused fields

* dnn(ocl4dnn/conv): simplify ctor

* dnn(ocl4dnn/conv): refactor kernelConfig localSize=NULL

* dnn(ocl4dnn/conv): drop unsupported double / untested half types

* dnn(ocl4dnn/conv): drop unused variable

* dnn(ocl4dnn/conv): alignSize/divUp

* dnn(ocl4dnn/conv): use enum values

* dnn(ocl4dnn): drop unused innerproduct variable

Signed-off-by: Li Peng <peng.li@intel.com>

* dnn(ocl4dnn): add an generic function to check cl option support

* dnn(ocl4dnn): run softmax subgroup version kernel first

Signed-off-by: Li Peng <peng.li@intel.com>
2017-10-02 15:38:00 +03:00
Vadim Pisarevsky
5e93c82023 Merge pull request #9491 from dkurt:tf_lstm 2017-09-28 21:04:06 +00:00
Vadim Pisarevsky
68cc2e292d Merge pull request #9734 from dkurt:fix_deconv_layer_kernel_layout 2017-09-28 11:42:57 +00:00
Vadim Pisarevsky
45365e4df1 Merge pull request #9691 from dkurt:padding_layer_refactoring 2017-09-28 11:34:28 +00:00
Dmitry Kurtaev
6e593cd1f0 Swap dimensions of deconvolution kernel 2017-09-27 22:38:34 +03:00
Dmitry Kurtaev
84cec17913 LSTM layer for TensorFlow importer 2017-09-26 12:59:36 +03:00
Alexander Alekhin
aea25e7f90 Merge pull request #9676 from jrobble:fix_caffe_swaprb 2017-09-25 16:05:11 +00:00
jrobble
c67ad49378 Set swapRB to false in Caffe tests and examples. 2017-09-22 09:58:48 -04:00
Dmitry Kurtaev
222149b9c6 Refactored Padding layer 2017-09-22 12:39:00 +03:00
Dmitry Kurtaev
17a85b16fc Remove reorder_dims attribute of Reshape layer 2017-09-21 16:42:03 +03:00
Dmitry Kurtaev
bdd8cc697a Import wrapped Dropout subgraphs from TensorFlow 2017-09-20 13:30:25 +03:00
Alexander Alekhin
6a5298a532 Merge pull request #9664 from tomoaki0705:fixBuildVS2013 2017-09-19 12:25:18 +00:00
Tomoaki Teshima
48c4bea57c fix build on Visual Studio 2013 and earlier 2017-09-19 19:08:23 +09:00
Vadim Pisarevsky
f7df5dd32c Merge pull request #9305 from dkurt:public_dnn_importer_is_deprecated 2017-09-18 09:35:35 +00:00
Vadim Pisarevsky
e012ccda4a Merge pull request #9517 from dkurt:tf_mobilenet 2017-09-18 09:31:19 +00:00
Vadim Pisarevsky
3358b8910b Merge pull request #9591 from dkurt:feature_dnn_caffe_importer_fp16 2017-09-18 09:26:23 +00:00
Dmitry Kurtaev
bd8e6b7e14 Make external cv::dnn::Importer usage is deprecated 2017-09-18 08:52:36 +03:00
Vadim Pisarevsky
4196543cd5 Merge pull request #9313 from dkurt:dnn_perf_test 2017-09-16 19:39:23 +00:00
Dmitry Kurtaev
d891e9b1d8 Layers for MobileNet from TensorFlow 2017-09-15 20:17:30 +03:00
Dmitry Kurtaev
8646d5fb84 FP16 Caffe models import and export 2017-09-15 18:06:34 +03:00
Vadim Pisarevsky
a391871cce Merge pull request #9301 from dkurt:dnn_npy 2017-09-15 12:42:02 +00:00
Vadim Pisarevsky
41b23fde9f Merge pull request #9524 from dkurt:dnn_torch_openface 2017-09-15 12:38:12 +00:00
Vadim Pisarevsky
cb0d695984 Merge pull request #9616 from dkurt:feature_dnn_tf_importer_atrous_conv 2017-09-15 10:24:05 +00:00
Dmitry Kurtaev
7dc6b1d7d4 Layers for OpenFace face recognition network 2017-09-14 09:11:31 +03:00
Dmitry Kurtaev
58b890b9f7 Dilated convolution import from TensorFlow 2017-09-13 18:44:14 +03:00
Dmitry Kurtaev
6639348c5e Disable SSD on CPU Halide test 2017-09-13 18:19:54 +03:00
Vadim Pisarevsky
8c7f19850f Merge pull request #9576 from dkurt:feature_dnn_tf_importer_fp16 2017-09-13 13:57:38 +00:00
Dmitry Kurtaev
ce41a15437 Import and convert FP16 weights from TensorFlow 2017-09-12 09:50:55 +03:00
Dmitry Kurtaev
cad7c4d51d MobileNet-SSD and VGG-SSD topologies in Halide 2017-09-08 09:55:53 +03:00
Dmitry Kurtaev
5c43a394c5 Added performance test for Caffe framework 2017-08-27 19:40:58 +03:00
Dmitry Kurtaev
e74c475162 Removed 3rdparty serialized NumPy arrays parser 2017-08-22 18:14:32 +03:00
Dmitry Kurtaev
54f0616a13 Deconvolution layer from TensorFlow 2017-08-21 21:38:07 +03:00
dkurt
339793143c Unit tests for TensorFlow importer 2017-08-03 11:29:48 +03:00
Aleksandr Rybnikov
8d6b8b45b6 Added ELU and test for it 2017-08-02 11:13:59 +03:00
Tomoaki Teshima
e63d628677 remove some rand functions
* make test more reproducible
2017-07-20 22:43:11 +09:00
Aleksandr Rybnikov
7d1140340e Rewrote googlenet tests 2017-07-18 18:49:14 +03:00
Alexander Alekhin
dc59476c34 Merge pull request #9078 from arrybn:resnet_squeezenet_tests
Added tests for ResNet-50 and SqueezeNet v1.1
2017-07-14 16:50:25 +00:00
Vadim Pisarevsky
0488d9bdb2 optimize out scaleLayer & concatLayer whenever possible
fixed problem in concat layer by disabling memory re-use in layers with multiple inputs

trying to fix the tests when Halide is used to run deep nets

another attempt to fix Halide tests

see if the Halide tests will pass with concat layer fusion turned off

trying to fix failures in halide tests; another try

one more experiment to make halide_concat & halide_enet tests pass

continue attempts to fix halide tests

moving on

uncomment parallel concat layer

seemingly fixed failures in Halide tests and re-enabled concat layer fusion; thanks to dkurt for the patch
2017-07-14 18:30:53 +03:00
dkurt
3203635765 Eltwise layer fixes 2017-07-10 12:58:11 +03:00
Aleksandr Rybnikov
d3c14ce3ea Added tests for ResNet-50 and SqueezeNet v1.1 2017-07-04 12:19:27 +03:00
Alexander Alekhin
9aaf6a66a3 Merge pull request #9073 from berak:dnn_fix_tutorial_links 2017-07-03 16:11:10 +00:00
Alexander Alekhin
32d1eb36df dnn: fix links 2017-07-03 12:44:27 +03:00
Alexander Alekhin
72a765d766 dnn: fix compilation of Halide tests 2017-06-29 16:36:34 +03:00
Vadim Pisarevsky
bbb14d3746 Merge pull request #9003 from dkurt:halide_bug_fixes 2017-06-28 08:48:27 +00:00
Vadim Pisarevsky
8b3d6603d5 another round of dnn optimization (#9011)
* another round of dnn optimization:
* increased malloc alignment across OpenCV from 16 to 64 bytes to make it AVX2 and even AVX-512 friendly
* improved SIMD optimization of pooling layer, optimized average pooling
* cleaned up convolution layer implementation
* made activation layer "attacheable" to all other layers, including fully connected and addition layer.
* fixed bug in the fusion algorithm: "LayerData::consumers" should not be cleared, because it desctibes the topology.
* greatly optimized permutation layer, which improved SSD performance
* parallelized element-wise binary/ternary/... ops (sum, prod, max)

* also, added missing copyrights to many of the layer implementation files

* temporarily disabled (again) the check for intermediate blobs consistency; fixed warnings from various builders
2017-06-28 11:15:22 +03:00
dkurt
121789f78e Fixed some bugs from Halide tests 2017-06-27 14:52:46 +03:00
Alexander Alekhin
ee54bafe6b dnn: fix dnn python test files 2017-06-26 15:13:00 +03:00
Alexander Alekhin
382e38941c dnn: fix documentation links 2017-06-26 15:10:50 +03:00
Alexander Alekhin
93729784bb dnn: move module from opencv_contrib
e6f63c7a38/modules/dnn
2017-06-26 13:41:51 +03:00