- 05 10月, 2017 12 次提交
- 04 10月, 2017 3 次提交
-
-
由 Alexander Alekhin 提交于
-
由 Bisaloo 提交于
-
由 Peter J. Stieber 提交于
-
- 03 10月, 2017 1 次提交
-
-
由 Dmitry Kurtaev 提交于
-
- 02 10月, 2017 11 次提交
-
-
由 Alexander Alekhin 提交于
-
由 Vadim Pisarevsky 提交于
-
由 Alexander Alekhin 提交于
-
由 pengli 提交于
add libdnn acceleration to dnn module (#9114) * import libdnn code Signed-off-by: NLi Peng <peng.li@intel.com> * add convolution layer ocl acceleration Signed-off-by: NLi Peng <peng.li@intel.com> * add pooling layer ocl acceleration Signed-off-by: NLi Peng <peng.li@intel.com> * add softmax layer ocl acceleration Signed-off-by: NLi Peng <peng.li@intel.com> * add lrn layer ocl acceleration Signed-off-by: NLi Peng <peng.li@intel.com> * add innerproduct layer ocl acceleration Signed-off-by: NLi Peng <peng.li@intel.com> * add HAVE_OPENCL macro Signed-off-by: NLi Peng <peng.li@intel.com> * fix for convolution ocl Signed-off-by: NLi Peng <peng.li@intel.com> * enable getUMat() for multi-dimension Mat Signed-off-by: NLi Peng <peng.li@intel.com> * use getUMat for ocl acceleration Signed-off-by: NLi Peng <peng.li@intel.com> * use CV_OCL_RUN macro Signed-off-by: NLi 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: NLi Peng <peng.li@intel.com> * fix innerproduct accuracy test Signed-off-by: NLi Peng <peng.li@intel.com> * remove trailing space Signed-off-by: NLi 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: NLi Peng <peng.li@intel.com> * Fix compile failure Signed-off-by: NLi Peng <peng.li@intel.com> * check env flag for auto tuning Signed-off-by: NLi Peng <peng.li@intel.com> * switch to new ocl kernels for softmax layer Signed-off-by: NLi 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: NLi Peng <peng.li@intel.com> * fallback to cpu path for fc layer with multi output Signed-off-by: NLi Peng <peng.li@intel.com> * update output message Signed-off-by: NLi Peng <peng.li@intel.com> * update fully connected layer fallback to gemm API if libdnn return false Signed-off-by: NLi Peng <peng.li@intel.com> * Add ReLU OCL implementation * disable layer fusion for now Signed-off-by: NLi Peng <peng.li@intel.com> * Add OCL implementation for concat layer Signed-off-by: NWu Zhiwen <zhiwen.wu@intel.com> * libdnn: update license and copyrights Also refine libdnn coding style Signed-off-by: NWu Zhiwen <zhiwen.wu@intel.com> Signed-off-by: NLi 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: NLi Peng <peng.li@intel.com> * update ReLU layer ocl path Signed-off-by: NLi 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: NLi Peng <peng.li@intel.com> * Add cl_event based timer for cv::ocl * Rename libdnn to ocl4dnn Signed-off-by: NLi Peng <peng.li@intel.com> Signed-off-by: Nwzw <zhiwen.wu@intel.com> * use UMat for ocl4dnn internal buffer Remove allocateMemory which use clCreateBuffer directly Signed-off-by: NLi Peng <peng.li@intel.com> Signed-off-by: Nwzw <zhiwen.wu@intel.com> * enable buffer gemm in ocl4dnn innerproduct Signed-off-by: NLi Peng <peng.li@intel.com> * replace int_tp globally for ocl4dnn kernels. Signed-off-by: Nwzw <zhiwen.wu@intel.com> Signed-off-by: NLi Peng <peng.li@intel.com> * create UMat for layer params Signed-off-by: NLi Peng <peng.li@intel.com> * update sign ocl kernel Signed-off-by: NLi Peng <peng.li@intel.com> * update image based gemm of inner product layer Signed-off-by: NLi Peng <peng.li@intel.com> * remove buffer gemm of inner product layer call cv::gemm API instead Signed-off-by: NLi Peng <peng.li@intel.com> * change ocl4dnn forward parameter to UMat Signed-off-by: NLi 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: NLi Peng <peng.li@intel.com> * Refine gemm like convolution kernel. Signed-off-by: NLi 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: NZhigang Gong <zhigang.gong@intel.com> * replace cl_mem with umat in ocl4dnn convolution Signed-off-by: NLi Peng <peng.li@intel.com> * remove weight_image_ of ocl4dnn inner product Actually it is unused in the computation Signed-off-by: NLi 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: NLi Peng <peng.li@intel.com> * add build option for log softmax Signed-off-by: NLi Peng <peng.li@intel.com> * remove unused ocl kernels in ocl4dnn Signed-off-by: NLi Peng <peng.li@intel.com> * replace ocl4dnnSet with opencv setTo Signed-off-by: NLi Peng <peng.li@intel.com> * replace ALIGN with cv::alignSize Signed-off-by: NLi Peng <peng.li@intel.com> * check kernel build options Signed-off-by: NLi 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: NLi Peng <peng.li@intel.com> * remove unused ctx_id Signed-off-by: NLi Peng <peng.li@intel.com> * change clEnqueueNDRangeKernel to kernel.run() Signed-off-by: NLi Peng <peng.li@intel.com> * change cl_mem to UMat in image based gemm Signed-off-by: NLi Peng <peng.li@intel.com> * check intel subgroup support for lrn and pooling layer Signed-off-by: NLi Peng <peng.li@intel.com> * Fix convolution bug if group is greater than 1 Signed-off-by: NLi Peng <peng.li@intel.com> * Set default layer preferableTarget to be DNN_TARGET_CPU Signed-off-by: NLi Peng <peng.li@intel.com> * Add ocl perf test for convolution Signed-off-by: NLi Peng <peng.li@intel.com> * Add more ocl accuracy test Signed-off-by: NLi Peng <peng.li@intel.com> * replace cl_image with ocl::Image2D Signed-off-by: NLi Peng <peng.li@intel.com> * Fix build failure in elementwise layer Signed-off-by: NLi Peng <peng.li@intel.com> * use getUMat() to get blob data Signed-off-by: NLi Peng <peng.li@intel.com> * replace cl_mem handle with ocl::KernelArg Signed-off-by: NLi 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: NLi Peng <peng.li@intel.com> * dnn(ocl4dnn): fix bias bug for gemm like kernel * dnn(ocl4dnn): wrap cl_mem into UMat Signed-off-by: NLi 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: NLi 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: NLi 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: NLi Peng <peng.li@intel.com>
-
由 Pavel Rojtberg 提交于
-
由 Vadim Pisarevsky 提交于
-
由 Alexander Alekhin 提交于
doc: fix typo in py_tutorials
-
由 Vadim Pisarevsky 提交于
-
由 Alexander Alekhin 提交于
-
由 Pranit Bauva 提交于
-
由 Suleyman TURKMEN 提交于
-
- 01 10月, 2017 4 次提交
-
-
由 berak 提交于
-
由 Alexander Alekhin 提交于
-
由 Alexander Alekhin 提交于
-
由 Suleyman TURKMEN 提交于
-
- 29 9月, 2017 9 次提交
-
-
由 Alexander Alekhin 提交于
-
由 Alexander Alekhin 提交于
-
由 Alexander Alekhin 提交于
-
由 Alexander Alekhin 提交于
-
由 Alexander Alekhin 提交于
-
由 Alexander Alekhin 提交于
-
由 Tomoaki Teshima 提交于
* remove raw NEON/SSE2 implementation as much as possible * replace them to universal intrinsic in InRange/Compare/AddWeighted
-
由 bradford barr 提交于
OpenCV fails to detect tbb on a debug build if the platform has only installed debug libraries. This PR adds an additional check to the tbb detect logic for systems that only install tbb debug and not both tbb debug and release.
-
由 Vadim Pisarevsky 提交于
-