Commit Graph

1370 Commits

Author SHA1 Message Date
Vadim Pisarevsky
0739f28e56 Merge pull request #9786 from LaurentBerger:Histo3d 2017-10-10 10:58:34 +00:00
Vadim Pisarevsky
7d55c09a9f Merge pull request #9763 from seiko2plus:addVsxCore 2017-10-10 10:00:31 +00:00
LaurentBerger
752f232335 It's done 2017-10-09 22:25:57 +02:00
Sayed Adel
4b968d1fe2 Added universal intrinsic for VSX 2017-10-09 00:32:41 +00:00
Sayed Adel
d077778074 Added support for VSX 2017-10-09 00:32:29 +00:00
Alexander Alekhin
6be25727ec ocl: refactor program compilation 2017-10-08 19:55:01 +03:00
Jasper Shemilt
0136711cf4 Adds fitEllipseAMS to imgproc: The Approximate Mean Square (AMS) proposed by Taubin 1991.
Adds fitEllipseDirect to imgproc: The Direct least square (Direct) method by Fitzgibbon1999.

New Tests are included for the methods.
fitEllipseAMS Tests
fitEllipseDirect Tests

Comparative examples are added to fitEllipse.cpp in Samples.
2017-10-02 16:38:41 +01: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
Alexander Alekhin
529632f8d0 core: cv::eigenNonSymmetric() via EigenvalueDecomposition 2017-10-01 07:45:32 +00:00
Alexander Alekhin
1283d62e49 ocl: Kernel::runProfiling() 2017-09-19 15:34:35 +03:00
Alexander Alekhin
d9ab31490c ocl: profiling queue 2017-09-19 15:32:15 +03:00
Vadim Pisarevsky
f4136679ea Merge pull request #9551 from ChristofKaufmann:MultiChannelMask 2017-09-18 09:28:34 +00:00
Christof Kaufmann
7ec59fc097 Revert changes of mean and meanStdDev 2017-09-17 21:00:28 +02:00
Alexander Alekhin
9dea296241 Merge pull request #9458 from csukuangfj:fix-doc 2017-09-15 19:35:01 +00:00
Vadim Pisarevsky
5707304777 Merge pull request #9633 from saskatchewancatch:psnr-doc 2017-09-15 12:26:34 +00:00
saskatchewancatch
a90a93b454 i9629 - Added actual documentation for cv::PSNR function 2017-09-14 18:20:00 -06:00
Vadim Pisarevsky
32bb71d686 Merge pull request #9603 from alalek:ocl_device_extensions 2017-09-13 14:43:36 +00:00
Vadim Pisarevsky
f1aa180a40 Merge pull request #9574 from saskatchewancatch:i9482 2017-09-13 13:59:43 +00:00
saskatchewancatch
c9d3c0f206 More whitespace fixes 2017-09-10 19:03:56 -06:00
Alexander Alekhin
169add5aa6 ocl: added cv::ocl::Device::isExtensionSupported() method 2017-09-10 20:32:30 +00:00
Alexander Alekhin
26ad229bd6 Merge pull request #9590 from alalek:ocl_runtime_fix 2017-09-09 19:30:07 +00:00
saskatchewancatch
30ff197f78 Fix whitespace issues 2017-09-08 20:27:03 -06:00
saskatchewancatch
8877e3aedb Adjustmenbts 2017-09-08 19:55:19 -06:00
Alexander Alekhin
576a85368d Merge pull request #9582 from sovrasov:fix_gcc7_macro_warning 2017-09-08 13:31:27 +00:00
Alexander Alekhin
44b75eb116 core: fix OpenCL runtime compilation with HAVE_OPENCL_STATIC 2017-09-08 12:43:33 +03:00
Pavel Vlasov
37ab318657 Compatibility improvement with old IPP versions (tested on 9.0.1);
Manual IPP dispatcher simplification;
2017-09-08 11:08:24 +03:00
saskatchewancatch
fda1e76776 Feedback.
Still need to remove the descriptions of these flags from cv::norm
2017-09-07 22:43:55 -06:00
Vladislav Sovrasov
eeba5c3603 core: fix gcc7 warning about empty VA_ARGS in Assert macro 2017-09-07 16:20:38 +03:00
Vadim Pisarevsky
d25cbaaba8 Merge pull request #9293 from sovrasov:assert_improvement 2017-09-07 11:17:42 +00:00
Maksim Shabunin
c92c99ed0b Enabled forEach for const Mats 2017-09-07 11:35:14 +03:00
saskatchewancatch
570083fb9f i9482:
Removing description for cv::NormTypes that's already in cv::norm

Masking out NORM_TYPE_MASK from docs since it's not intended to be exposed to user
2017-09-06 22:42:01 -06:00
saskatchewancatch
b32685c714 i9482: Improve documentation concerning norm functionality.
Added forkfour Latex command to math js support.

Split cv::norm documentation between the cv::norm and its overload, to make things clearer

Corrected some typos and cleaned up grammar.

Result is clearer documentation for the norms.

Work pending...
2017-09-06 22:35:39 -06:00
Christof Kaufmann
46a668c565 Add multi-channel mask support to mean, meanStdDev and setTo
This adds the possibility to use multi-channel masks for the functions
cv::mean, cv::meanStdDev and the method Mat::setTo. The tests have now a
probability to use multi-channel masks for operations that support them.
This also includes Mat::copyTo, which supported multi-channel masks
before, but there was no test confirming this.
2017-09-04 19:40:27 +02:00
Alexander Alekhin
0451629e22 core(persistence): resolve DMatch/KeyPoint problem 2017-08-31 19:35:48 +03:00
Alexander Alekhin
86b55b3923 core: eliminate CV_ELEM_SIZE() 2017-08-31 19:35:48 +03:00
Alexander Alekhin
7e12c879c2 core: extend traits::Type / traits::Depth for compatible types
DMatch and Keypoint are not compatible types (mixed float/int fields)
2017-08-31 19:35:48 +03:00
Alexander Alekhin
72f789bf34 core: fix type traits 2017-08-31 15:05:46 +03:00
Boris Fomitchev
dde04d5d3e Addressing CUDA9 shfl deprecation 2017-08-29 13:27:21 -07:00
Maarten de Vries
3571e8f263 Restrict std::initializer_list constructors to arithmetic types. 2017-08-29 16:37:20 +02:00
Alexander Alekhin
034aaa7a70 Merge pull request #9465 from tomoaki0705:fixJetsonTK1Build 2017-08-29 11:46:10 +00:00
Tomoaki Teshima
6531fd142c fix build error on Jetson TK1
* guard correctly in header file
  * guard correctly in cmake file
2017-08-29 19:05:13 +09:00
KUANG Fangjun
11fa0094ff Improve the documentation.
Add demo code for cv::reduce, cv::merge and cv::split.
2017-08-28 12:36:23 +02:00
Alexander Alekhin
603fa03ac6 Merge pull request #9441 from wzw-intel:delete_program 2017-08-25 12:03:27 +00:00
Wu Zhiwen
da3da84a20 ocl: Add a function to unload a run-time cached program
This function is the counterpart of "Context::getProg".
With this function, users have chance to unload a program
from global run-time cached programs, and save resource.
2017-08-25 08:42:11 +08:00
Alexander Alekhin
9c14a2f0aa Merge pull request #9395 from lupustr3:pvlasov/icv2017u3_update 2017-08-24 11:48:53 +00:00
Boris Fomitchev
c48807c383 Merge pull request #9418 from borisfom:cuda9
CUDA9 build fixed, added detection (#9418)

* CUDA9 build fixed, added detection

* Replacing deprecated __shfl_xxx with __shfl_sync, fixing bogus CUDA9 warnings
2017-08-24 07:11:44 +00:00
Alexander Alekhin
a893b147dc Merge pull request #9428 from csukuangfj:fix-commandline-parser 2017-08-23 12:48:28 +00:00
Pavel Vlasov
a57718e1ac ICV2017u3 package update;
- Optimizations set change. Now IPP integrations will provide code for SSE42, AVX2 and AVX512 (SKX) CPUs only. For HW below SSE42 IPP code is disabled.
- Performance regressions fixes for IPP code paths;
- cv::boxFilter integration improvement;
- cv::filter2D integration improvement;
2017-08-23 14:24:43 +03:00
KUANG Fangjun
97ec91ad67 fix cv::CommandLineParser.
It should handle bool value not only of "true" but also of "TRUE" and "True".
2017-08-23 11:38:58 +02:00
KUANG Fangjun
336996152a Improve the documentation. 2017-08-20 17:21:39 +02:00