Commit Graph

227 Commits

Author SHA1 Message Date
Alexander Alekhin
cec700525c core(ocl): fix deadlock in UMatDataAutoLock
UMatData locks are not mapped on real locks (they are mapped to some "pre-initialized" pool).

Concurrent execution of these statements may lead to deadlock:
- a.copyTo(b) from thread 1
- c.copyTo(d) from thread 2
where:
- 'a' and 'd' are mapped to single lock "A".
- 'b' and 'c' are mapped to single lock "B".

Workaround is to process locks with strict order.
2018-01-16 17:33:06 +03:00
Maksim Shabunin
594a93316c Fixed concurrent OpenCL cache folder name generation 2018-01-12 19:03:16 +03:00
Alexander Alekhin
534645a12f ocl: workaround option to disable usage of buffer "Rect" operations 2017-12-22 13:05:03 +03:00
Jiri Horner
3dbf392d48 fix build with intrinsics enabled
* since #10231 opencv with instrumentation does not build
2017-12-17 20:23:15 +01:00
Tomoaki Teshima
267c5a747b suppress warnings on OpenCL build
* stop re-enabling the warning C4127
  * disabling is done in CMakeLists.txt
2017-12-13 15:07:51 +09:00
Vadim Pisarevsky
9fa505027a Merge pull request #10263 from mshabunin:embedded-build 2017-12-11 12:42:45 +00:00
Maksim Shabunin
7349b8f5ce Build for embedded systems 2017-12-11 13:27:37 +03:00
Alexander Alekhin
a82d2363f4 ocl: refactor Program API
- don't store ProgramSource in compiled Programs (resolved problem with "source" buffers lifetime)
- completelly remove Program::read/write methods implementation:
  - replaced with method to query RAW OpenCL binary without any "custom" data
- deprecate Program::getPrefix() methods
2017-12-05 22:25:14 +03:00
Alexander Alekhin
13c4a02157 ocl: low-level API to support OpenCL binary programs 2017-12-05 22:25:14 +03:00
Vadim Pisarevsky
5ce38e516e Merge pull request #10223 from vpisarev:ocl_mac_fixes
* fixed OpenCL functions on Mac, so that the tests pass

* fixed compile warnings; temporarily disabled OCL branch of TV L1 optical flow on mac

* fixed other few warnings on macos
2017-12-05 13:32:28 +03:00
Alexander Alekhin
0595ab3eef ocl: fix usage of invalid OpenCL cache on mixed 64/32-bit platforms
Observed during launch of 32/64-bit applications on Windows.
Added '32-bit' prefix for 32-bit OpenCL devices. No prefix on 64-bit configurations.
2017-12-01 14:20:18 +03:00
Vadim Pisarevsky
f5dba12762 Merge pull request #10180 from alalek:ocl_avoid_unnecessary_initialization 2017-11-29 11:42:22 +00:00
Alexander Alekhin
0ed3209b00 ocl: avoid unnecessary loading/initializing OpenCL subsystem
If there are no OpenCL/UMat methods calls from application.

OpenCL subsystem is initialized:
- haveOpenCL() is called from application
- useOpenCL() is called from application
- access to OpenCL allocator: UMat is created (empty UMat is ignored) or UMat <-> Mat conversions are called

Don't call OpenCL functions if OPENCV_OPENCL_RUNTIME=disabled
(independent from OpenCL linkage type)
2017-11-28 14:02:42 +03:00
Alexander Alekhin
c4b158ff91 Merge pull request #10167 from alalek:ocl_fix_issue_contrib1467 2017-11-27 11:05:07 +00:00
Alexander Alekhin
92b35e6758 ocl: fix null pointer access crash 2017-11-27 12:43:29 +03:00
Alexander Alekhin
b6abf0d3f9 ocl: drop obsolete cache directories after upgrade of OpenCL driver
Entries with the same platform name, the same device name and with different driver versions
are assumed obsolete.
2017-11-24 17:02:28 +03:00
Alexander Alekhin
8e6280fc8e ocl: binary program cache 2017-11-22 12:56:38 +03:00
Alexander Alekhin
9c4f0a984f ocl: drop CV_OclDbgAssert 2017-11-03 13:31:37 +03:00
Alexander Alekhin
8fb48c09f7 ocl: improve debug information 2017-11-03 13:31:37 +03:00
Alexander Alekhin
7809c4156f core(ocl): workaround CL_OUT_OF_RESOURCES error
Flush deallocation queue before calling map/unmap
2017-10-30 17:54:56 +03:00
Alexander Alekhin
d96cac1341 ocl: cache program build failures
To prevent unnecessary compiler invocations
2017-10-23 13:46:56 +03:00
Alexander Alekhin
185faf99bd ocl: simplify ocl::Timer interface 2017-10-18 16:01:21 +03:00
Vadim Pisarevsky
44699c59b3 Merge pull request #9799 from alalek:ocl_program 2017-10-09 12:43:46 +00:00
Wu Zhiwen
dbe9ee0924 ocl: simplify ocl::Timer
Use clFinish to gurantee commands completed, instead of waiting for events.

Signed-off-by: Wu Zhiwen <zhiwen.wu@intel.com>
2017-10-09 13:48:38 +08:00
Alexander Alekhin
6be25727ec ocl: refactor program compilation 2017-10-08 19:55:01 +03:00
Alexander Alekhin
04b4495493 ocl: define ProgramSource before program
no changes in code
2017-10-08 19:55:01 +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
Alexander Alekhin
1f18282f04 build: fix MSVS2017 warning 2017-10-01 10:48:35 +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
Alexander Alekhin
2e68f89225 ocl: update kernel global size adjustment
Prevents 10000x1 => 10000x8 transformation after getContinuousSize() call
2017-09-18 19:04:46 +03:00
Alexander Alekhin
169add5aa6 ocl: added cv::ocl::Device::isExtensionSupported() method 2017-09-10 20:32:30 +00: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
Dmitry Kurtaev
41519d3ac0 Fixed some OpenCL interface bugs 2017-08-09 11:54:55 +03:00
Alexander Alekhin
16fb74425e ocl: fix program cache key 2017-07-31 17:24:08 +03:00
Alexander Alekhin
da8dbf6cf5 ocl: async cl_buffer cleanup queue (for event callback) 2017-07-07 13:41:20 +03:00
Alexander Alekhin
daee982106 ocl: rework events handling with clSetEventCallback 2017-07-06 13:25:32 +03:00
Alexander Alekhin
006966e629 trace: initial support for code trace 2017-06-26 17:07:13 +03:00
Maksim Shabunin
b04ed5956e Fixed several issues found by static analysis in core module 2017-05-23 12:35:31 +03:00
Vadim Pisarevsky
925594d1e3 Merge pull request #7894 from alalek:ocl_program 2017-05-03 13:48:58 +00:00
insoow
2922738b6d Merge pull request #8104 from insoow:master
Gemm kernels for Intel GPU (#8104)

* Fix an issue with Kernel object reset release when consecutive Kernel::run calls

Kernel::run launch OCL gpu kernels and set a event callback function
to decreate the ref count of UMat or remove UMat when the lauched workloads
are completed. However, for some OCL kernels requires multiple call of
Kernel::run function with some kernel parameter changes (e.g., input
and output buffer offset) to get the final computation result.
In the case, the current implementation requires unnecessary
synchronization and cleanupMat.

This fix requires the user to specify whether there will be more work or not.
If there is no remaining computation, the Kernel::run will reset the
kernel object

Signed-off-by: Woo, Insoo <insoo.woo@intel.com>

* GEMM kernel optimization for Intel GEN

The optimized kernels uses cl_intel_subgroups extension for better
performance.

Note: This optimized kernels will be part of ISAAC in a code generation
way under MIT license.

Signed-off-by: Woo, Insoo <insoo.woo@intel.com>

* Fix API compatibility error

This patch fixes a OCV API compatibility error. The error was reported
due to the interface changes of Kernel::run. To resolve the issue,
An overloaded function of Kernel::run is added. It take a flag indicating
whether there are more work to be done with the kernel object without
releasing resources related to it.

Signed-off-by: Woo, Insoo <insoo.woo@intel.com>

* Renaming intel_gpu_gemm.cpp to intel_gpu_gemm.inl.hpp

Signed-off-by: Woo, Insoo <insoo.woo@intel.com>

* Revert "Fix API compatibility error"

This reverts commit 2ef427db91.

Conflicts:
	modules/core/src/intel_gpu_gemm.inl.hpp

* Revert "Fix an issue with Kernel object reset release when consecutive Kernel::run calls"

This reverts commit cc7f9f5469.

* Fix the case of uninitialization D

When C is null and beta is non-zero, D is used without initialization.
This resloves the issue

Signed-off-by: Woo, Insoo <insoo.woo@intel.com>

* fix potential output error due to 0 * nan

Signed-off-by: Woo, Insoo <insoo.woo@intel.com>

* whitespace fix, eliminate non-ASCII symbols

* fix build warning
2017-04-19 12:57:54 +03:00
KUANG, Fangjun
debc1c4c95 fix an error while setting kernel argument for a 3-D matrix. 2017-03-12 18:29:49 +01:00
Vadim Pisarevsky
ca0b6fbb95 Merge pull request #8203 from alalek:ocl_kernelarg_validate 2017-02-16 14:58:25 +00:00
Jinay Patel
f710d96602 Update ocl.cpp function Device::OpenCLVersion()
Instead of CL_DEVICE_EXTENSIONS in Querying OpenCLVersion, it should be CL_DEVICE_VERSION flag which needs to be passed to the function.
2017-02-16 15:20:57 +05:30
Fangjun KUANG
5e4dfb9546 Change the argument type to cl_platform_info.
Although both `cl_platform_info` and `cl_device_info` are defined as macro `cl_uint`, it needs to use `cl_platform_info` to get
the platform information.
2017-02-15 16:44:40 +01:00
Alexander Alekhin
4c7aa8645a ocl: validate arguments in KernelArgs constructor
- don't use undefined flag=0. It should be CONSTANT instead.
- don't allow 'UMat* m=NULL' argument (except LOCAL/CONSTANT flags).
  This case is not handled well to provide NULL __global pointers.
  It is better to use '-D' macro defines instead (at least for performance)
2017-02-14 16:10:32 +03:00
Alexander Alekhin
c3e409f950 ocl: update program cache defaults 2016-12-22 04:52:43 +03:00
Alexander Alekhin
7a95e654eb ocl: update compiled programs
- minimize library initialization time (lazy calculations of program hash)
- LRU cache of in-memory compiled programs
2016-12-19 17:17:20 +03:00
Alexander Alekhin
05d187ec90 ocl: move deprecated code from ocl.cpp 2016-12-19 00:34:50 +03:00