diff --git a/modules/nonfree/src/surf.ocl.cpp b/modules/nonfree/src/surf.ocl.cpp index 3d5cb4e083..b1dc166acf 100644 --- a/modules/nonfree/src/surf.ocl.cpp +++ b/modules/nonfree/src/surf.ocl.cpp @@ -55,20 +55,11 @@ namespace cv { namespace ocl { - static const char noImage2dOption[] = "-D DISABLE_IMAGE2D"; - - static bool use_image2d = false; - static void openCLExecuteKernelSURF(Context *clCxt, const cv::ocl::ProgramEntry* source, string kernelName, size_t globalThreads[3], size_t localThreads[3], std::vector< std::pair > &args, int channels, int depth) { char optBuf [100] = {0}; char * optBufPtr = optBuf; - if( !use_image2d ) - { - strcat(optBufPtr, noImage2dOption); - optBufPtr += strlen(noImage2dOption); - } cl_kernel kernel; kernel = openCLGetKernelFromSource(clCxt, source, kernelName, optBufPtr); size_t wave_size = queryWaveFrontSize(kernel); @@ -149,13 +140,10 @@ public: counters.setTo(Scalar::all(0)); integral(img, surf_.sum); - use_image2d = support_image2d(); - if(use_image2d) - { - bindImgTex(img, imgTex); - bindImgTex(surf_.sum, sumTex); - finish(); - } + + bindImgTex(img, imgTex); + bindImgTex(surf_.sum, sumTex); + finish(); maskSumTex = 0; diff --git a/modules/ocl/include/opencv2/ocl/private/util.hpp b/modules/ocl/include/opencv2/ocl/private/util.hpp index bb1a5eeeb7..59e1d05094 100644 --- a/modules/ocl/include/opencv2/ocl/private/util.hpp +++ b/modules/ocl/include/opencv2/ocl/private/util.hpp @@ -165,15 +165,11 @@ private: // for faster clamping, there is no buffer padding for the constructed texture Ptr CV_EXPORTS bindTexturePtr(const oclMat &mat); -// returns whether the current context supports image2d_t format or not -bool CV_EXPORTS support_image2d(Context *clCxt = Context::getContext()); - bool CV_EXPORTS isCpuDevice(); size_t CV_EXPORTS queryWaveFrontSize(cl_kernel kernel); - inline size_t divUp(size_t total, size_t grain) { return (total + grain - 1) / grain; diff --git a/modules/ocl/src/cl_programcache.cpp b/modules/ocl/src/cl_programcache.cpp index ee291c966a..76a115fd70 100644 --- a/modules/ocl/src/cl_programcache.cpp +++ b/modules/ocl/src/cl_programcache.cpp @@ -448,26 +448,30 @@ cl_program ProgramCache::getProgram(const Context *ctx, const cv::ocl::ProgramEn { stringstream src_sign; - src_sign << source->name; - src_sign << getClContext(ctx); - if (NULL != build_options) + if (source->name) { - src_sign << "_" << build_options; - } + src_sign << source->name; + src_sign << getClContext(ctx); + if (NULL != build_options) + { + src_sign << "_" << build_options; + } - { - cv::AutoLock lockCache(mutexCache); - cl_program program = ProgramCache::getProgramCache()->progLookup(src_sign.str()); - if (!!program) { - clRetainProgram(program); - return program; + cv::AutoLock lockCache(mutexCache); + cl_program program = ProgramCache::getProgramCache()->progLookup(src_sign.str()); + if (!!program) + { + clRetainProgram(program); + return program; + } } } cv::AutoLock lockCache(mutexFiles); // second check + if (source->name) { cv::AutoLock lockCache(mutexCache); cl_program program = ProgramCache::getProgramCache()->progLookup(src_sign.str()); @@ -493,6 +497,7 @@ cl_program ProgramCache::getProgram(const Context *ctx, const cv::ocl::ProgramEn cl_program program = programFileCache.getOrBuildProgram(ctx, source, all_build_options); //Cache the binary for future use if build_options is null + if (source->name) { cv::AutoLock lockCache(mutexCache); this->addProgram(src_sign.str(), program); diff --git a/modules/ocl/src/gftt.cpp b/modules/ocl/src/gftt.cpp index 00275a73c2..5c34179c54 100644 --- a/modules/ocl/src/gftt.cpp +++ b/modules/ocl/src/gftt.cpp @@ -202,8 +202,6 @@ void cv::ocl::GoodFeaturesToTrackDetector_OCL::operator ()(const oclMat& image, CV_Assert(qualityLevel > 0 && minDistance >= 0 && maxCorners >= 0); CV_Assert(mask.empty() || (mask.type() == CV_8UC1 && mask.size() == image.size())); - CV_DbgAssert(support_image2d()); - ensureSizeIsEnough(image.size(), CV_32F, eig_); if (useHarrisDetector) diff --git a/modules/ocl/src/mcwutil.cpp b/modules/ocl/src/mcwutil.cpp index 201bc7d818..0a948d198a 100644 --- a/modules/ocl/src/mcwutil.cpp +++ b/modules/ocl/src/mcwutil.cpp @@ -216,41 +216,11 @@ namespace cv { return Ptr(new TextureCL(bindTexture(mat), mat.rows, mat.cols, mat.type())); } + void releaseTexture(cl_mem& texture) { openCLFree(texture); } - - bool support_image2d(Context *clCxt) - { - const cv::ocl::ProgramEntry _kernel = {"test_func", "__kernel void test_func(image2d_t img) {}", NULL}; - static bool _isTested = false; - static bool _support = false; - if(_isTested) - { - return _support; - } - try - { - cv::ocl::openCLGetKernelFromSource(clCxt, &_kernel, "test_func"); - cv::ocl::finish(); - _support = true; - } - catch (const cv::Exception& e) - { - if(e.code == -217) - { - _support = false; - } - else - { - // throw e once again - throw e; - } - } - _isTested = true; - return _support; - } }//namespace ocl }//namespace cv diff --git a/modules/ocl/src/opencl/pyrlk_no_image.cl b/modules/ocl/src/opencl/pyrlk_no_image.cl deleted file mode 100644 index 98a11b5c12..0000000000 --- a/modules/ocl/src/opencl/pyrlk_no_image.cl +++ /dev/null @@ -1,764 +0,0 @@ -/*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) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. -// Third party copyrights are property of their respective owners. -// -// @Authors -// Sen Liu, sen@multicorewareinc.com -// -// 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 oclMaterials 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*/ - -#define BUFFER 256 -void reduce3(float val1, float val2, float val3, __local float *smem1, __local float *smem2, __local float *smem3, int tid) -{ - smem1[tid] = val1; - smem2[tid] = val2; - smem3[tid] = val3; - barrier(CLK_LOCAL_MEM_FENCE); - -#if BUFFER > 128 - - if (tid < 128) - { - smem1[tid] = val1 += smem1[tid + 128]; - smem2[tid] = val2 += smem2[tid + 128]; - smem3[tid] = val3 += smem3[tid + 128]; - } - - barrier(CLK_LOCAL_MEM_FENCE); -#endif - -#if BUFFER > 64 - - if (tid < 64) - { - smem1[tid] = val1 += smem1[tid + 64]; - smem2[tid] = val2 += smem2[tid + 64]; - smem3[tid] = val3 += smem3[tid + 64]; - } - - barrier(CLK_LOCAL_MEM_FENCE); -#endif - - if (tid < 32) - { - smem1[tid] = val1 += smem1[tid + 32]; - smem2[tid] = val2 += smem2[tid + 32]; - smem3[tid] = val3 += smem3[tid + 32]; - } - - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 16) - { - smem1[tid] = val1 += smem1[tid + 16]; - smem2[tid] = val2 += smem2[tid + 16]; - smem3[tid] = val3 += smem3[tid + 16]; - } - - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 8) - { - volatile __local float *vmem1 = smem1; - volatile __local float *vmem2 = smem2; - volatile __local float *vmem3 = smem3; - - vmem1[tid] = val1 += vmem1[tid + 8]; - vmem2[tid] = val2 += vmem2[tid + 8]; - vmem3[tid] = val3 += vmem3[tid + 8]; - - vmem1[tid] = val1 += vmem1[tid + 4]; - vmem2[tid] = val2 += vmem2[tid + 4]; - vmem3[tid] = val3 += vmem3[tid + 4]; - - vmem1[tid] = val1 += vmem1[tid + 2]; - vmem2[tid] = val2 += vmem2[tid + 2]; - vmem3[tid] = val3 += vmem3[tid + 2]; - - vmem1[tid] = val1 += vmem1[tid + 1]; - vmem2[tid] = val2 += vmem2[tid + 1]; - vmem3[tid] = val3 += vmem3[tid + 1]; - } -} - -void reduce2(float val1, float val2, __local float *smem1, __local float *smem2, int tid) -{ - smem1[tid] = val1; - smem2[tid] = val2; - barrier(CLK_LOCAL_MEM_FENCE); - -#if BUFFER > 128 - - if (tid < 128) - { - smem1[tid] = val1 += smem1[tid + 128]; - smem2[tid] = val2 += smem2[tid + 128]; - } - - barrier(CLK_LOCAL_MEM_FENCE); -#endif - -#if BUFFER > 64 - - if (tid < 64) - { - smem1[tid] = val1 += smem1[tid + 64]; - smem2[tid] = val2 += smem2[tid + 64]; - } - - barrier(CLK_LOCAL_MEM_FENCE); -#endif - - if (tid < 32) - { - smem1[tid] = val1 += smem1[tid + 32]; - smem2[tid] = val2 += smem2[tid + 32]; - } - - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 16) - { - smem1[tid] = val1 += smem1[tid + 16]; - smem2[tid] = val2 += smem2[tid + 16]; - } - - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 8) - { - volatile __local float *vmem1 = smem1; - volatile __local float *vmem2 = smem2; - - vmem1[tid] = val1 += vmem1[tid + 8]; - vmem2[tid] = val2 += vmem2[tid + 8]; - - vmem1[tid] = val1 += vmem1[tid + 4]; - vmem2[tid] = val2 += vmem2[tid + 4]; - - vmem1[tid] = val1 += vmem1[tid + 2]; - vmem2[tid] = val2 += vmem2[tid + 2]; - - vmem1[tid] = val1 += vmem1[tid + 1]; - vmem2[tid] = val2 += vmem2[tid + 1]; - } -} - -void reduce1(float val1, __local float *smem1, int tid) -{ - smem1[tid] = val1; - barrier(CLK_LOCAL_MEM_FENCE); - -#if BUFFER > 128 - - if (tid < 128) - { - smem1[tid] = val1 += smem1[tid + 128]; - } - - barrier(CLK_LOCAL_MEM_FENCE); -#endif - -#if BUFFER > 64 - - if (tid < 64) - { - smem1[tid] = val1 += smem1[tid + 64]; - } - - barrier(CLK_LOCAL_MEM_FENCE); -#endif - - if (tid < 32) - { - smem1[tid] = val1 += smem1[tid + 32]; - } - - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 16) - { - volatile __local float *vmem1 = smem1; - - vmem1[tid] = val1 += vmem1[tid + 16]; - } - - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 8) - { - volatile __local float *vmem1 = smem1; - - vmem1[tid] = val1 += vmem1[tid + 8]; - vmem1[tid] = val1 += vmem1[tid + 4]; - vmem1[tid] = val1 += vmem1[tid + 2]; - vmem1[tid] = val1 += vmem1[tid + 1]; - } -} - -#define SCALE (1.0f / (1 << 20)) -#define THRESHOLD 0.01f -#define DIMENSION 21 - -float readImage2Df_C1(__global const float *image, const float x, const float y, const int rows, const int cols, const int elemCntPerRow) -{ - float2 coor = (float2)(x, y); - - int i0 = clamp((int)floor(coor.x), 0, cols - 1); - int j0 = clamp((int)floor(coor.y), 0, rows - 1); - int i1 = clamp((int)floor(coor.x) + 1, 0, cols - 1); - int j1 = clamp((int)floor(coor.y) + 1, 0, rows - 1); - float a = coor.x - floor(coor.x); - float b = coor.y - floor(coor.y); - - return (1 - a) * (1 - b) * image[mad24(j0, elemCntPerRow, i0)] - + a * (1 - b) * image[mad24(j0, elemCntPerRow, i1)] - + (1 - a) * b * image[mad24(j1, elemCntPerRow, i0)] - + a * b * image[mad24(j1, elemCntPerRow, i1)]; -} - -__kernel void lkSparse_C1_D5(__global const float *I, __global const float *J, - __global const float2 *prevPts, int prevPtsStep, __global float2 *nextPts, int nextPtsStep, __global uchar *status, __global float *err, - const int level, const int rows, const int cols, const int elemCntPerRow, - int PATCH_X, int PATCH_Y, int cn, int c_winSize_x, int c_winSize_y, int c_iters, char calcErr) -{ - __local float smem1[BUFFER]; - __local float smem2[BUFFER]; - __local float smem3[BUFFER]; - - float2 c_halfWin = (float2)((c_winSize_x - 1) >> 1, (c_winSize_y - 1) >> 1); - - const int tid = mad24(get_local_id(1), get_local_size(0), get_local_id(0)); - - float2 prevPt = prevPts[get_group_id(0)] * (1.0f / (1 << level)); - - if (prevPt.x < 0 || prevPt.x >= cols || prevPt.y < 0 || prevPt.y >= rows) - { - if (tid == 0 && level == 0) - { - status[get_group_id(0)] = 0; - } - - return; - } - - prevPt -= c_halfWin; - - // extract the patch from the first image, compute covariation matrix of derivatives - - float A11 = 0; - float A12 = 0; - float A22 = 0; - - float I_patch[1][3]; - float dIdx_patch[1][3]; - float dIdy_patch[1][3]; - - for (int yBase = get_local_id(1), i = 0; yBase < c_winSize_y; yBase += get_local_size(1), ++i) - { - for (int xBase = get_local_id(0), j = 0; xBase < c_winSize_x; xBase += get_local_size(0), ++j) - { - float x = (prevPt.x + xBase); - float y = (prevPt.y + yBase); - - I_patch[i][j] = readImage2Df_C1(I, x, y, rows, cols, elemCntPerRow); - float dIdx = 3.0f * readImage2Df_C1(I, x + 1, y - 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C1(I, x + 1, y, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C1(I, x + 1, y + 1, rows, cols, elemCntPerRow) - - (3.0f * readImage2Df_C1(I, x - 1, y - 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C1(I, x - 1, y, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C1(I, x - 1, y + 1, rows, cols, elemCntPerRow)); - - float dIdy = 3.0f * readImage2Df_C1(I, x - 1, y + 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C1(I, x, y + 1, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C1(I, x + 1, y + 1, rows, cols, elemCntPerRow) - - (3.0f * readImage2Df_C1(I, x - 1, y - 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C1(I, x, y - 1, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C1(I, x + 1, y - 1, rows, cols, elemCntPerRow)); - - dIdx_patch[i][j] = dIdx; - dIdy_patch[i][j] = dIdy; - - A11 += dIdx * dIdx; - A12 += dIdx * dIdy; - A22 += dIdy * dIdy; - } - } - - reduce3(A11, A12, A22, smem1, smem2, smem3, tid); - barrier(CLK_LOCAL_MEM_FENCE); - - A11 = smem1[0]; - A12 = smem2[0]; - A22 = smem3[0]; - - float D = A11 * A22 - A12 * A12; - - if (D < 1.192092896e-07f) - { - if (tid == 0 && level == 0) - { - status[get_group_id(0)] = 0; - } - - return; - } - - D = 1.f / D; - - A11 *= D; - A12 *= D; - A22 *= D; - - float2 nextPt = nextPts[get_group_id(0)]; - nextPt = nextPt * 2.0f - c_halfWin; - - for (int k = 0; k < c_iters; ++k) - { - if (nextPt.x < -c_halfWin.x || nextPt.x >= cols || nextPt.y < -c_halfWin.y || nextPt.y >= rows) - { - if (tid == 0 && level == 0) - { - status[get_group_id(0)] = 0; - } - - return; - } - - float b1 = 0; - float b2 = 0; - - for (int y = get_local_id(1), i = 0; y < c_winSize_y; y += get_local_size(1), ++i) - { - for (int x = get_local_id(0), j = 0; x < c_winSize_x; x += get_local_size(0), ++j) - { - float diff = (readImage2Df_C1(J, nextPt.x + x, nextPt.y + y, rows, cols, elemCntPerRow) - I_patch[i][j]) * 32.0f; - - b1 += diff * dIdx_patch[i][j]; - b2 += diff * dIdy_patch[i][j]; - } - } - - reduce2(b1, b2, smem1, smem2, tid); - barrier(CLK_LOCAL_MEM_FENCE); - - b1 = smem1[0]; - b2 = smem2[0]; - - float2 delta; - delta.x = A12 * b2 - A22 * b1; - delta.y = A12 * b1 - A11 * b2; - - nextPt += delta; - - //if (fabs(delta.x) < THRESHOLD && fabs(delta.y) < THRESHOLD) - // break; - } - - float errval = 0.0f; - - if (calcErr) - { - for (int y = get_local_id(1), i = 0; y < c_winSize_y; y += get_local_size(1), ++i) - { - for (int x = get_local_id(0), j = 0; x < c_winSize_x; x += get_local_size(0), ++j) - { - float diff = readImage2Df_C1(J, nextPt.x + x, nextPt.y + y, rows, cols, elemCntPerRow) - I_patch[i][j]; - - errval += fabs(diff); - } - } - - reduce1(errval, smem1, tid); - } - - if (tid == 0) - { - nextPt += c_halfWin; - - nextPts[get_group_id(0)] = nextPt; - - if (calcErr) - { - err[get_group_id(0)] = smem1[0] / (c_winSize_x * c_winSize_y); - } - } -} - -float4 readImage2Df_C4(__global const float4 *image, const float x, const float y, const int rows, const int cols, const int elemCntPerRow) -{ - float2 coor = (float2)(x, y); - - int i0 = clamp((int)floor(coor.x), 0, cols - 1); - int j0 = clamp((int)floor(coor.y), 0, rows - 1); - int i1 = clamp((int)floor(coor.x) + 1, 0, cols - 1); - int j1 = clamp((int)floor(coor.y) + 1, 0, rows - 1); - float a = coor.x - floor(coor.x); - float b = coor.y - floor(coor.y); - - return (1 - a) * (1 - b) * image[mad24(j0, elemCntPerRow, i0)] - + a * (1 - b) * image[mad24(j0, elemCntPerRow, i1)] - + (1 - a) * b * image[mad24(j1, elemCntPerRow, i0)] - + a * b * image[mad24(j1, elemCntPerRow, i1)]; -} - -__kernel void lkSparse_C4_D5(__global const float *I, __global const float *J, - __global const float2 *prevPts, int prevPtsStep, __global float2 *nextPts, int nextPtsStep, __global uchar *status, __global float *err, - const int level, const int rows, const int cols, const int elemCntPerRow, - int PATCH_X, int PATCH_Y, int cn, int c_winSize_x, int c_winSize_y, int c_iters, char calcErr) -{ - __local float smem1[BUFFER]; - __local float smem2[BUFFER]; - __local float smem3[BUFFER]; - - float2 c_halfWin = (float2)((c_winSize_x - 1) >> 1, (c_winSize_y - 1) >> 1); - - const int tid = mad24(get_local_id(1), get_local_size(0), get_local_id(0)); - - float2 prevPt = prevPts[get_group_id(0)] * (1.0f / (1 << level)); - - if (prevPt.x < 0 || prevPt.x >= cols || prevPt.y < 0 || prevPt.y >= rows) - { - if (tid == 0 && level == 0) - { - status[get_group_id(0)] = 0; - } - - return; - } - - prevPt -= c_halfWin; - - // extract the patch from the first image, compute covariation matrix of derivatives - - float A11 = 0; - float A12 = 0; - float A22 = 0; - - float4 I_patch[1][3]; - float4 dIdx_patch[1][3]; - float4 dIdy_patch[1][3]; - - __global float4 *ptrI = (__global float4 *)I; - - for (int yBase = get_local_id(1), i = 0; yBase < c_winSize_y; yBase += get_local_size(1), ++i) - { - for (int xBase = get_local_id(0), j = 0; xBase < c_winSize_x; xBase += get_local_size(0), ++j) - { - float x = (prevPt.x + xBase); - float y = (prevPt.y + yBase); - - I_patch[i][j] = readImage2Df_C4(ptrI, x, y, rows, cols, elemCntPerRow); - - float4 dIdx = 3.0f * readImage2Df_C4(ptrI, x + 1, y - 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C4(ptrI, x + 1, y, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C4(ptrI, x + 1, y + 1, rows, cols, elemCntPerRow) - - (3.0f * readImage2Df_C4(ptrI, x - 1, y - 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C4(ptrI, x - 1, y, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C4(ptrI, x - 1, y + 1, rows, cols, elemCntPerRow)); - - float4 dIdy = 3.0f * readImage2Df_C4(ptrI, x - 1, y + 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C4(ptrI, x, y + 1, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C4(ptrI, x + 1, y + 1, rows, cols, elemCntPerRow) - - (3.0f * readImage2Df_C4(ptrI, x - 1, y - 1, rows, cols, elemCntPerRow) + 10.0f * readImage2Df_C4(ptrI, x, y - 1, rows, cols, elemCntPerRow) + 3.0f * readImage2Df_C4(ptrI, x + 1, y - 1, rows, cols, elemCntPerRow)); - - dIdx_patch[i][j] = dIdx; - dIdy_patch[i][j] = dIdy; - - A11 += (dIdx * dIdx).x + (dIdx * dIdx).y + (dIdx * dIdx).z; - A12 += (dIdx * dIdy).x + (dIdx * dIdy).y + (dIdx * dIdy).z; - A22 += (dIdy * dIdy).x + (dIdy * dIdy).y + (dIdy * dIdy).z; - } - } - - reduce3(A11, A12, A22, smem1, smem2, smem3, tid); - barrier(CLK_LOCAL_MEM_FENCE); - - A11 = smem1[0]; - A12 = smem2[0]; - A22 = smem3[0]; - - float D = A11 * A22 - A12 * A12; - //pD[get_group_id(0)] = D; - - if (D < 1.192092896e-07f) - { - if (tid == 0 && level == 0) - { - status[get_group_id(0)] = 0; - } - - return; - } - - D = 1.f / D; - - A11 *= D; - A12 *= D; - A22 *= D; - - float2 nextPt = nextPts[get_group_id(0)]; - - nextPt = nextPt * 2.0f - c_halfWin; - - __global float4 *ptrJ = (__global float4 *)J; - - for (int k = 0; k < c_iters; ++k) - { - if (nextPt.x < -c_halfWin.x || nextPt.x >= cols || nextPt.y < -c_halfWin.y || nextPt.y >= rows) - { - if (tid == 0 && level == 0) - { - status[get_group_id(0)] = 0; - } - - return; - } - - float b1 = 0; - float b2 = 0; - - for (int y = get_local_id(1), i = 0; y < c_winSize_y; y += get_local_size(1), ++i) - { - for (int x = get_local_id(0), j = 0; x < c_winSize_x; x += get_local_size(0), ++j) - { - float4 diff = (readImage2Df_C4(ptrJ, nextPt.x + x, nextPt.y + y, rows, cols, elemCntPerRow) - I_patch[i][j]) * 32.0f; - - b1 += (diff * dIdx_patch[i][j]).x + (diff * dIdx_patch[i][j]).y + (diff * dIdx_patch[i][j]).z; - b2 += (diff * dIdy_patch[i][j]).x + (diff * dIdy_patch[i][j]).y + (diff * dIdy_patch[i][j]).z; - } - } - - reduce2(b1, b2, smem1, smem2, tid); - barrier(CLK_LOCAL_MEM_FENCE); - - b1 = smem1[0]; - b2 = smem2[0]; - - float2 delta; - delta.x = A12 * b2 - A22 * b1; - delta.y = A12 * b1 - A11 * b2; - - nextPt += delta; - - //if (fabs(delta.x) < THRESHOLD && fabs(delta.y) < THRESHOLD) - // break; - } - - float errval = 0.0f; - - if (calcErr) - { - for (int y = get_local_id(1), i = 0; y < c_winSize_y; y += get_local_size(1), ++i) - { - for (int x = get_local_id(0), j = 0; x < c_winSize_x; x += get_local_size(0), ++j) - { - float4 diff = readImage2Df_C4(ptrJ, nextPt.x + x, nextPt.y + y, rows, cols, elemCntPerRow) - I_patch[i][j]; - - errval += fabs(diff.x) + fabs(diff.y) + fabs(diff.z); - } - } - - reduce1(errval, smem1, tid); - } - - if (tid == 0) - { - nextPt += c_halfWin; - nextPts[get_group_id(0)] = nextPt; - - if (calcErr) - { - err[get_group_id(0)] = smem1[0] / (3 * c_winSize_x * c_winSize_y); - } - } -} - -int readImage2Di_C1(__global const int *image, float2 coor, int2 size, const int elemCntPerRow) -{ - int i = clamp((int)floor(coor.x), 0, size.x - 1); - int j = clamp((int)floor(coor.y), 0, size.y - 1); - return image[mad24(j, elemCntPerRow, i)]; -} - -__kernel void lkDense_C1_D0(__global const int *I, __global const int *J, __global float *u, int uStep, __global float *v, int vStep, __global const float *prevU, int prevUStep, __global const float *prevV, int prevVStep, - const int rows, const int cols, /*__global float* err, int errStep, int cn,*/ - const int elemCntPerRow, int c_winSize_x, int c_winSize_y, int c_iters, char calcErr) -{ - int c_halfWin_x = (c_winSize_x - 1) / 2; - int c_halfWin_y = (c_winSize_y - 1) / 2; - - const int patchWidth = get_local_size(0) + 2 * c_halfWin_x; - const int patchHeight = get_local_size(1) + 2 * c_halfWin_y; - - __local int smem[8192]; - - __local int *I_patch = smem; - __local int *dIdx_patch = I_patch + patchWidth * patchHeight; - __local int *dIdy_patch = dIdx_patch + patchWidth * patchHeight; - - const int xBase = get_group_id(0) * get_local_size(0); - const int yBase = get_group_id(1) * get_local_size(1); - int2 size = (int2)(cols, rows); - - for (int i = get_local_id(1); i < patchHeight; i += get_local_size(1)) - { - for (int j = get_local_id(0); j < patchWidth; j += get_local_size(0)) - { - float x = xBase - c_halfWin_x + j + 0.5f; - float y = yBase - c_halfWin_y + i + 0.5f; - - I_patch[i * patchWidth + j] = readImage2Di_C1(I, (float2)(x, y), size, elemCntPerRow); - - // Sharr Deriv - - dIdx_patch[i * patchWidth + j] = 3 * readImage2Di_C1(I, (float2)(x + 1, y - 1), size, elemCntPerRow) + 10 * readImage2Di_C1(I, (float2)(x + 1, y), size, elemCntPerRow) + 3 * readImage2Di_C1(I, (float2)(x + 1, y + 1), size, elemCntPerRow) - - (3 * readImage2Di_C1(I, (float2)(x - 1, y - 1), size, elemCntPerRow) + 10 * readImage2Di_C1(I, (float2)(x - 1, y), size, elemCntPerRow) + 3 * readImage2Di_C1(I, (float2)(x - 1, y + 1), size, elemCntPerRow)); - - dIdy_patch[i * patchWidth + j] = 3 * readImage2Di_C1(I, (float2)(x - 1, y + 1), size, elemCntPerRow) + 10 * readImage2Di_C1(I, (float2)(x, y + 1), size, elemCntPerRow) + 3 * readImage2Di_C1(I, (float2)(x + 1, y + 1), size, elemCntPerRow) - - (3 * readImage2Di_C1(I, (float2)(x - 1, y - 1), size, elemCntPerRow) + 10 * readImage2Di_C1(I, (float2)(x, y - 1), size, elemCntPerRow) + 3 * readImage2Di_C1(I, (float2)(x + 1, y - 1), size, elemCntPerRow)); - } - } - - barrier(CLK_LOCAL_MEM_FENCE); - - // extract the patch from the first image, compute covariation matrix of derivatives - - const int x = get_global_id(0); - const int y = get_global_id(1); - - if (x >= cols || y >= rows) - { - return; - } - - int A11i = 0; - int A12i = 0; - int A22i = 0; - - for (int i = 0; i < c_winSize_y; ++i) - { - for (int j = 0; j < c_winSize_x; ++j) - { - int dIdx = dIdx_patch[(get_local_id(1) + i) * patchWidth + (get_local_id(0) + j)]; - int dIdy = dIdy_patch[(get_local_id(1) + i) * patchWidth + (get_local_id(0) + j)]; - - A11i += dIdx * dIdx; - A12i += dIdx * dIdy; - A22i += dIdy * dIdy; - } - } - - float A11 = A11i; - float A12 = A12i; - float A22 = A22i; - - float D = A11 * A22 - A12 * A12; - - //if (calcErr && GET_MIN_EIGENVALS) - // (err + y * errStep)[x] = minEig; - - if (D < 1.192092896e-07f) - { - //if (calcErr) - // err(y, x) = 3.402823466e+38f; - - return; - } - - D = 1.f / D; - - A11 *= D; - A12 *= D; - A22 *= D; - - float2 nextPt; - nextPt.x = x + prevU[y / 2 * prevUStep / 4 + x / 2] * 2.0f; - nextPt.y = y + prevV[y / 2 * prevVStep / 4 + x / 2] * 2.0f; - - for (int k = 0; k < c_iters; ++k) - { - if (nextPt.x < 0 || nextPt.x >= cols || nextPt.y < 0 || nextPt.y >= rows) - { - //if (calcErr) - // err(y, x) = 3.402823466e+38f; - - return; - } - - int b1 = 0; - int b2 = 0; - - for (int i = 0; i < c_winSize_y; ++i) - { - for (int j = 0; j < c_winSize_x; ++j) - { - int iI = I_patch[(get_local_id(1) + i) * patchWidth + get_local_id(0) + j]; - int iJ = readImage2Di_C1(J, (float2)(nextPt.x - c_halfWin_x + j + 0.5f, nextPt.y - c_halfWin_y + i + 0.5f), size, elemCntPerRow); - - int diff = (iJ - iI) * 32; - - int dIdx = dIdx_patch[(get_local_id(1) + i) * patchWidth + (get_local_id(0) + j)]; - int dIdy = dIdy_patch[(get_local_id(1) + i) * patchWidth + (get_local_id(0) + j)]; - - b1 += diff * dIdx; - b2 += diff * dIdy; - } - } - - float2 delta; - delta.x = A12 * b2 - A22 * b1; - delta.y = A12 * b1 - A11 * b2; - - nextPt.x += delta.x; - nextPt.y += delta.y; - - if (fabs(delta.x) < 0.01f && fabs(delta.y) < 0.01f) - { - break; - } - } - - u[y * uStep / 4 + x] = nextPt.x - x; - v[y * vStep / 4 + x] = nextPt.y - y; - - if (calcErr) - { - int errval = 0; - - for (int i = 0; i < c_winSize_y; ++i) - { - for (int j = 0; j < c_winSize_x; ++j) - { - int iI = I_patch[(get_local_id(1) + i) * patchWidth + get_local_id(0) + j]; - int iJ = readImage2Di_C1(J, (float2)(nextPt.x - c_halfWin_x + j + 0.5f, nextPt.y - c_halfWin_y + i + 0.5f), size, elemCntPerRow); - - errval += abs(iJ - iI); - } - } - - //err[y * errStep / 4 + x] = static_cast(errval) / (c_winSize_x * c_winSize_y); - } -} diff --git a/modules/ocl/src/pyrlk.cpp b/modules/ocl/src/pyrlk.cpp index bd9b18e8c3..5f76113179 100644 --- a/modules/ocl/src/pyrlk.cpp +++ b/modules/ocl/src/pyrlk.cpp @@ -82,18 +82,16 @@ static void lkSparse_run(oclMat &I, oclMat &J, int level, dim3 patch, Size winSize, int iters) { Context *clCxt = I.clCxt; - int elemCntPerRow = I.step / I.elemSize(); string kernelName = "lkSparse"; - bool isImageSupported = support_image2d(); - size_t localThreads[3] = { 8, isImageSupported ? 8 : 32, 1 }; - size_t globalThreads[3] = { 8 * ptcount, isImageSupported ? 8 : 32, 1}; + size_t localThreads[3] = { 8, 8, 1 }; + size_t globalThreads[3] = { 8 * ptcount, 8, 1}; int cn = I.oclchannels(); char calcErr = level==0?1:0; vector > args; - cl_mem ITex = isImageSupported ? bindTexture(I) : (cl_mem)I.data; - cl_mem JTex = isImageSupported ? bindTexture(J) : (cl_mem)J.data; + cl_mem ITex = bindTexture(I); + cl_mem JTex = bindTexture(J); args.push_back( make_pair( sizeof(cl_mem), (void *)&ITex )); args.push_back( make_pair( sizeof(cl_mem), (void *)&JTex )); @@ -106,8 +104,6 @@ static void lkSparse_run(oclMat &I, oclMat &J, args.push_back( make_pair( sizeof(cl_int), (void *)&level )); args.push_back( make_pair( sizeof(cl_int), (void *)&I.rows )); args.push_back( make_pair( sizeof(cl_int), (void *)&I.cols )); - if (!isImageSupported) - args.push_back( make_pair( sizeof(cl_int), (void *)&elemCntPerRow ) ); args.push_back( make_pair( sizeof(cl_int), (void *)&patch.x )); args.push_back( make_pair( sizeof(cl_int), (void *)&patch.y )); args.push_back( make_pair( sizeof(cl_int), (void *)&cn )); @@ -120,32 +116,23 @@ static void lkSparse_run(oclMat &I, oclMat &J, if (is_cpu) { openCLExecuteKernel(clCxt, &pyrlk, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth(), (char*)" -D CPU"); - releaseTexture(ITex); - releaseTexture(JTex); } else { - if(isImageSupported) - { - stringstream idxStr; - idxStr << kernelName << "_C" << I.oclchannels() << "_D" << I.depth(); - cl_kernel kernel = openCLGetKernelFromSource(clCxt, &pyrlk, idxStr.str()); - int wave_size = (int)queryWaveFrontSize(kernel); - openCLSafeCall(clReleaseKernel(kernel)); - - static char opt[32] = {0}; - sprintf(opt, "-D WAVE_SIZE=%d", wave_size); - - openCLExecuteKernel(clCxt, &pyrlk, kernelName, globalThreads, localThreads, - args, I.oclchannels(), I.depth(), opt); - releaseTexture(ITex); - releaseTexture(JTex); - } - else - { - openCLExecuteKernel(clCxt, &pyrlk_no_image, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth()); - } + stringstream idxStr; + idxStr << kernelName << "_C" << I.oclchannels() << "_D" << I.depth(); + cl_kernel kernel = openCLGetKernelFromSource(clCxt, &pyrlk, idxStr.str()); + int wave_size = (int)queryWaveFrontSize(kernel); + openCLSafeCall(clReleaseKernel(kernel)); + + static char opt[32] = {0}; + sprintf(opt, "-D WAVE_SIZE=%d", wave_size); + + openCLExecuteKernel(clCxt, &pyrlk, kernelName, globalThreads, localThreads, + args, I.oclchannels(), I.depth(), opt); } + releaseTexture(ITex); + releaseTexture(JTex); } void cv::ocl::PyrLKOpticalFlow::sparse(const oclMat &prevImg, const oclMat &nextImg, const oclMat &prevPts, oclMat &nextPts, oclMat &status, oclMat *err) @@ -226,37 +213,19 @@ static void lkDense_run(oclMat &I, oclMat &J, oclMat &u, oclMat &v, oclMat &prevU, oclMat &prevV, oclMat *err, Size winSize, int iters) { Context *clCxt = I.clCxt; - bool isImageSupported = support_image2d(); - int elemCntPerRow = I.step / I.elemSize(); string kernelName = "lkDense"; size_t localThreads[3] = { 16, 16, 1 }; size_t globalThreads[3] = { I.cols, I.rows, 1}; - bool calcErr; - if (err) - { - calcErr = true; - } - else - { - calcErr = false; - } + cl_char calcErr = err ? 1 : 0; cl_mem ITex; cl_mem JTex; - if (isImageSupported) - { - ITex = bindTexture(I); - JTex = bindTexture(J); - } - else - { - ITex = (cl_mem)I.data; - JTex = (cl_mem)J.data; - } + ITex = bindTexture(I); + JTex = bindTexture(J); vector > args; @@ -273,28 +242,15 @@ static void lkDense_run(oclMat &I, oclMat &J, oclMat &u, oclMat &v, args.push_back( make_pair( sizeof(cl_int), (void *)&prevV.step )); args.push_back( make_pair( sizeof(cl_int), (void *)&I.rows )); args.push_back( make_pair( sizeof(cl_int), (void *)&I.cols )); - //args.push_back( make_pair( sizeof(cl_mem), (void *)&(*err).data )); - //args.push_back( make_pair( sizeof(cl_int), (void *)&(*err).step )); - if (!isImageSupported) - { - args.push_back( make_pair( sizeof(cl_int), (void *)&elemCntPerRow ) ); - } args.push_back( make_pair( sizeof(cl_int), (void *)&winSize.width )); args.push_back( make_pair( sizeof(cl_int), (void *)&winSize.height )); args.push_back( make_pair( sizeof(cl_int), (void *)&iters )); args.push_back( make_pair( sizeof(cl_char), (void *)&calcErr )); - if (isImageSupported) - { - openCLExecuteKernel(clCxt, &pyrlk, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth()); + openCLExecuteKernel(clCxt, &pyrlk, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth()); - releaseTexture(ITex); - releaseTexture(JTex); - } - else - { - openCLExecuteKernel(clCxt, &pyrlk_no_image, kernelName, globalThreads, localThreads, args, I.oclchannels(), I.depth()); - } + releaseTexture(ITex); + releaseTexture(JTex); } void cv::ocl::PyrLKOpticalFlow::dense(const oclMat &prevImg, const oclMat &nextImg, oclMat &u, oclMat &v, oclMat *err) diff --git a/modules/ocl/src/tvl1flow.cpp b/modules/ocl/src/tvl1flow.cpp index bbeecb4748..825d932a24 100644 --- a/modules/ocl/src/tvl1flow.cpp +++ b/modules/ocl/src/tvl1flow.cpp @@ -411,9 +411,6 @@ void ocl_tvl1flow::estimateU(oclMat &I1wx, oclMat &I1wy, oclMat &grad, void ocl_tvl1flow::warpBackward(const oclMat &I0, const oclMat &I1, oclMat &I1x, oclMat &I1y, oclMat &u1, oclMat &u2, oclMat &I1w, oclMat &I1wx, oclMat &I1wy, oclMat &grad, oclMat &rho) { Context* clCxt = I0.clCxt; - const bool isImgSupported = support_image2d(clCxt); - - CV_Assert(isImgSupported); int u1ElementSize = u1.elemSize(); int u1Step = u1.step/u1ElementSize;