diff --git a/modules/ocl/perf/perf_opticalflow.cpp b/modules/ocl/perf/perf_opticalflow.cpp index 936d7a77fc..10763b5b0f 100644 --- a/modules/ocl/perf/perf_opticalflow.cpp +++ b/modules/ocl/perf/perf_opticalflow.cpp @@ -48,8 +48,8 @@ ///////////// PyrLKOpticalFlow //////////////////////// PERFTEST(PyrLKOpticalFlow) { - std::string images1[] = {"rubberwhale1.png", "basketball1.png"}; - std::string images2[] = {"rubberwhale2.png", "basketball2.png"}; + std::string images1[] = {"rubberwhale1.png", "aloeL.jpg"}; + std::string images2[] = {"rubberwhale2.png", "aloeR.jpg"}; for (size_t i = 0; i < sizeof(images1) / sizeof(std::string); i++) { diff --git a/modules/ocl/src/opencl/pyrlk.cl b/modules/ocl/src/opencl/pyrlk.cl index 40a1993952..02cf3afa44 100644 --- a/modules/ocl/src/opencl/pyrlk.cl +++ b/modules/ocl/src/opencl/pyrlk.cl @@ -17,6 +17,7 @@ // @Authors // Dachuan Zhao, dachuan@multicorewareinc.com // Yao Wang, bitwangyaoyao@gmail.com +// Xiaopeng Fu, fuxiaopeng2222@163.com // // Redistribution and use in source and binary forms, with or without modification, // are permitted provided that the following conditions are met: @@ -47,6 +48,7 @@ //#pragma OPENCL EXTENSION cl_amd_printf : enable #define BUFFER 64 +#define BUFFER2 BUFFER>>1 #ifndef WAVE_SIZE #define WAVE_SIZE 1 #endif @@ -58,53 +60,16 @@ void reduce3(float val1, float val2, float val3, __local float* smem1, __local smem3[tid] = val3; barrier(CLK_LOCAL_MEM_FENCE); - if (tid < 32) + for(int i = BUFFER2; i > 0; i >>= 1) { - smem1[tid] += smem1[tid + 32]; - smem2[tid] += smem2[tid + 32]; - smem3[tid] += smem3[tid + 32]; + if(tid < i) + { + smem1[tid] += smem1[tid + i]; + smem2[tid] += smem2[tid + i]; + smem3[tid] += smem3[tid + i]; + } + barrier(CLK_LOCAL_MEM_FENCE); } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 16) - { - smem1[tid] += smem1[tid + 16]; - smem2[tid] += smem2[tid + 16]; - smem3[tid] += smem3[tid + 16]; - } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 8) - { - smem1[tid] += smem1[tid + 8]; - smem2[tid] += smem2[tid + 8]; - smem3[tid] += smem3[tid + 8]; - } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 4) - { - smem1[tid] += smem1[tid + 4]; - smem2[tid] += smem2[tid + 4]; - smem3[tid] += smem3[tid + 4]; - } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 2) - { - smem1[tid] += smem1[tid + 2]; - smem2[tid] += smem2[tid + 2]; - smem3[tid] += smem3[tid + 2]; - } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 1) - { - smem1[BUFFER] = smem1[tid] + smem1[tid + 1]; - smem2[BUFFER] = smem2[tid] + smem2[tid + 1]; - smem3[BUFFER] = smem3[tid] + smem3[tid + 1]; - } - barrier(CLK_LOCAL_MEM_FENCE); } void reduce2(float val1, float val2, volatile __local float* smem1, volatile __local float* smem2, int tid) @@ -113,47 +78,15 @@ void reduce2(float val1, float val2, volatile __local float* smem1, volatile __l smem2[tid] = val2; barrier(CLK_LOCAL_MEM_FENCE); - if (tid < 32) + for(int i = BUFFER2; i > 0; i >>= 1) { - smem1[tid] += smem1[tid + 32]; - smem2[tid] += smem2[tid + 32]; + if(tid < i) + { + smem1[tid] += smem1[tid + i]; + smem2[tid] += smem2[tid + i]; + } + barrier(CLK_LOCAL_MEM_FENCE); } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 16) - { - smem1[tid] += smem1[tid + 16]; - smem2[tid] += smem2[tid + 16]; - } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 8) - { - smem1[tid] += smem1[tid + 8]; - smem2[tid] += smem2[tid + 8]; - } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 4) - { - smem1[tid] += smem1[tid + 4]; - smem2[tid] += smem2[tid + 4]; - } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 2) - { - smem1[tid] += smem1[tid + 2]; - smem2[tid] += smem2[tid + 2]; - } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 1) - { - smem1[BUFFER] = smem1[tid] + smem1[tid + 1]; - smem2[BUFFER] = smem2[tid] + smem2[tid + 1]; - } - barrier(CLK_LOCAL_MEM_FENCE); } void reduce1(float val1, volatile __local float* smem1, int tid) @@ -161,45 +94,18 @@ void reduce1(float val1, volatile __local float* smem1, int tid) smem1[tid] = val1; barrier(CLK_LOCAL_MEM_FENCE); - if (tid < 32) + for(int i = BUFFER2; i > 0; i >>= 1) { - smem1[tid] += smem1[tid + 32]; + if(tid < i) + { + smem1[tid] += smem1[tid + i]; + } + barrier(CLK_LOCAL_MEM_FENCE); } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 16) - { - smem1[tid] += smem1[tid + 16]; - } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 8) - { - smem1[tid] += smem1[tid + 8]; - } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 4) - { - smem1[tid] += smem1[tid + 4]; - } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 2) - { - smem1[tid] += smem1[tid + 2]; - } - barrier(CLK_LOCAL_MEM_FENCE); - - if (tid < 1) - { - smem1[BUFFER] = smem1[tid] + smem1[tid + 1]; - } - barrier(CLK_LOCAL_MEM_FENCE); } #else -void reduce3(float val1, float val2, float val3, -__local volatile float* smem1, __local volatile float* smem2, __local volatile float* smem3, int tid) +void reduce3(float val1, float val2, float val3, + __local volatile float* smem1, __local volatile float* smem2, __local volatile float* smem3, int tid) { smem1[tid] = val1; smem2[tid] = val2; @@ -212,15 +118,19 @@ __local volatile float* smem1, __local volatile float* smem2, __local volatile f smem2[tid] += smem2[tid + 32]; smem3[tid] += smem3[tid + 32]; #if WAVE_SIZE < 32 - } barrier(CLK_LOCAL_MEM_FENCE); - if (tid < 16) { + } + barrier(CLK_LOCAL_MEM_FENCE); + if (tid < 16) + { #endif smem1[tid] += smem1[tid + 16]; smem2[tid] += smem2[tid + 16]; smem3[tid] += smem3[tid + 16]; #if WAVE_SIZE <16 - } barrier(CLK_LOCAL_MEM_FENCE); - if (tid < 8) { + } + barrier(CLK_LOCAL_MEM_FENCE); + if (tid < 8) + { #endif smem1[tid] += smem1[tid + 8]; smem2[tid] += smem2[tid + 8]; @@ -238,6 +148,7 @@ __local volatile float* smem1, __local volatile float* smem2, __local volatile f smem2[tid] += smem2[tid + 1]; smem3[tid] += smem3[tid + 1]; } + barrier(CLK_LOCAL_MEM_FENCE); } void reduce2(float val1, float val2, __local volatile float* smem1, __local volatile float* smem2, int tid) @@ -251,14 +162,18 @@ void reduce2(float val1, float val2, __local volatile float* smem1, __local vola smem1[tid] += smem1[tid + 32]; smem2[tid] += smem2[tid + 32]; #if WAVE_SIZE < 32 - } barrier(CLK_LOCAL_MEM_FENCE); - if (tid < 16) { + } + barrier(CLK_LOCAL_MEM_FENCE); + if (tid < 16) + { #endif smem1[tid] += smem1[tid + 16]; smem2[tid] += smem2[tid + 16]; #if WAVE_SIZE <16 - } barrier(CLK_LOCAL_MEM_FENCE); - if (tid < 8) { + } + barrier(CLK_LOCAL_MEM_FENCE); + if (tid < 8) + { #endif smem1[tid] += smem1[tid + 8]; smem2[tid] += smem2[tid + 8]; @@ -272,6 +187,7 @@ void reduce2(float val1, float val2, __local volatile float* smem1, __local vola smem1[tid] += smem1[tid + 1]; smem2[tid] += smem2[tid + 1]; } + barrier(CLK_LOCAL_MEM_FENCE); } void reduce1(float val1, __local volatile float* smem1, int tid) @@ -283,19 +199,24 @@ void reduce1(float val1, __local volatile float* smem1, int tid) { smem1[tid] += smem1[tid + 32]; #if WAVE_SIZE < 32 - } barrier(CLK_LOCAL_MEM_FENCE); - if (tid < 16) { + } + barrier(CLK_LOCAL_MEM_FENCE); + if (tid < 16) + { #endif smem1[tid] += smem1[tid + 16]; #if WAVE_SIZE <16 - } barrier(CLK_LOCAL_MEM_FENCE); - if (tid < 8) { + } + barrier(CLK_LOCAL_MEM_FENCE); + if (tid < 8) + { #endif smem1[tid] += smem1[tid + 8]; smem1[tid] += smem1[tid + 4]; smem1[tid] += smem1[tid + 2]; smem1[tid] += smem1[tid + 1]; } + barrier(CLK_LOCAL_MEM_FENCE); } #endif @@ -306,106 +227,100 @@ void reduce1(float val1, __local volatile float* smem1, int tid) __constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; void SetPatch(image2d_t I, float x, float y, - float* Pch, float* Dx, float* Dy, - float* A11, float* A12, float* A22) + float* Pch, float* Dx, float* Dy, + float* A11, float* A12, float* A22) { - *Pch = read_imagef(I, sampler, (float2)(x, y)).x; + *Pch = read_imagef(I, sampler, (float2)(x, y)).x; - float dIdx = 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1)).x + 10.0f * read_imagef(I, sampler, (float2)(x + 1, y)).x + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)).x - - (3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)).x + 10.0f * read_imagef(I, sampler, (float2)(x - 1, y)).x + 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1)).x); + float dIdx = 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1)).x + 10.0f * read_imagef(I, sampler, (float2)(x + 1, y)).x + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)).x - + (3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)).x + 10.0f * read_imagef(I, sampler, (float2)(x - 1, y)).x + 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1)).x); - float dIdy = 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1)).x + 10.0f * read_imagef(I, sampler, (float2)(x, y + 1)).x + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)).x - - (3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)).x + 10.0f * read_imagef(I, sampler, (float2)(x, y - 1)).x + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1)).x); + float dIdy = 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1)).x + 10.0f * read_imagef(I, sampler, (float2)(x, y + 1)).x + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)).x - + (3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)).x + 10.0f * read_imagef(I, sampler, (float2)(x, y - 1)).x + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1)).x); - *Dx = dIdx; - *Dy = dIdy; + *Dx = dIdx; + *Dy = dIdy; - *A11 += dIdx * dIdx; - *A12 += dIdx * dIdy; - *A22 += dIdy * dIdy; + *A11 += dIdx * dIdx; + *A12 += dIdx * dIdy; + *A22 += dIdy * dIdy; } void GetPatch(image2d_t J, float x, float y, - float* Pch, float* Dx, float* Dy, - float* b1, float* b2) + float* Pch, float* Dx, float* Dy, + float* b1, float* b2) { - float J_val = read_imagef(J, sampler, (float2)(x, y)).x; - float diff = (J_val - *Pch) * 32.0f; - *b1 += diff**Dx; - *b2 += diff**Dy; + float J_val = read_imagef(J, sampler, (float2)(x, y)).x; + float diff = (J_val - *Pch) * 32.0f; + *b1 += diff**Dx; + *b2 += diff**Dy; } void GetError(image2d_t J, const float x, const float y, const float* Pch, float* errval) { - float diff = read_imagef(J, sampler, (float2)(x,y)).x-*Pch; - *errval += fabs(diff); + float diff = read_imagef(J, sampler, (float2)(x,y)).x-*Pch; + *errval += fabs(diff); } void SetPatch4(image2d_t I, const float x, const float y, - float4* Pch, float4* Dx, float4* Dy, - float* A11, float* A12, float* A22) + float4* Pch, float4* Dx, float4* Dy, + float* A11, float* A12, float* A22) { - *Pch = read_imagef(I, sampler, (float2)(x, y)); + *Pch = read_imagef(I, sampler, (float2)(x, y)); - float4 dIdx = 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x + 1, y)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)) - - (3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x - 1, y)) + 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1))); + float4 dIdx = 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x + 1, y)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)) - + (3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x - 1, y)) + 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1))); - float4 dIdy = 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1)) + 10.0f * read_imagef(I, sampler, (float2)(x, y + 1)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)) - - (3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x, y - 1)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1))); + float4 dIdy = 3.0f * read_imagef(I, sampler, (float2)(x - 1, y + 1)) + 10.0f * read_imagef(I, sampler, (float2)(x, y + 1)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y + 1)) - + (3.0f * read_imagef(I, sampler, (float2)(x - 1, y - 1)) + 10.0f * read_imagef(I, sampler, (float2)(x, y - 1)) + 3.0f * read_imagef(I, sampler, (float2)(x + 1, y - 1))); - *Dx = dIdx; - *Dy = dIdy; - float4 sqIdx = dIdx * dIdx; - *A11 += sqIdx.x + sqIdx.y + sqIdx.z; - sqIdx = dIdx * dIdy; - *A12 += sqIdx.x + sqIdx.y + sqIdx.z; - sqIdx = dIdy * dIdy; - *A22 += sqIdx.x + sqIdx.y + sqIdx.z; + *Dx = dIdx; + *Dy = dIdy; + float4 sqIdx = dIdx * dIdx; + *A11 += sqIdx.x + sqIdx.y + sqIdx.z; + sqIdx = dIdx * dIdy; + *A12 += sqIdx.x + sqIdx.y + sqIdx.z; + sqIdx = dIdy * dIdy; + *A22 += sqIdx.x + sqIdx.y + sqIdx.z; } void GetPatch4(image2d_t J, const float x, const float y, - const float4* Pch, const float4* Dx, const float4* Dy, - float* b1, float* b2) + const float4* Pch, const float4* Dx, const float4* Dy, + float* b1, float* b2) { - float4 J_val = read_imagef(J, sampler, (float2)(x, y)); - float4 diff = (J_val - *Pch) * 32.0f; - float4 xdiff = diff* *Dx; - *b1 += xdiff.x + xdiff.y + xdiff.z; - xdiff = diff* *Dy; - *b2 += xdiff.x + xdiff.y + xdiff.z; + float4 J_val = read_imagef(J, sampler, (float2)(x, y)); + float4 diff = (J_val - *Pch) * 32.0f; + float4 xdiff = diff* *Dx; + *b1 += xdiff.x + xdiff.y + xdiff.z; + xdiff = diff* *Dy; + *b2 += xdiff.x + xdiff.y + xdiff.z; } void GetError4(image2d_t J, const float x, const float y, const float4* Pch, float* errval) { - float4 diff = read_imagef(J, sampler, (float2)(x,y))-*Pch; - *errval += fabs(diff.x) + fabs(diff.y) + fabs(diff.z); + float4 diff = read_imagef(J, sampler, (float2)(x,y))-*Pch; + *errval += fabs(diff.x) + fabs(diff.y) + fabs(diff.z); } #define GRIDSIZE 3 __kernel void lkSparse_C1_D5(image2d_t I, image2d_t 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, int PATCH_X, int PATCH_Y, int cn, int c_winSize_x, int c_winSize_y, int c_iters, char calcErr) + __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, int PATCH_X, int PATCH_Y, int cn, int c_winSize_x, int c_winSize_y, int c_iters, char calcErr) { -#ifdef CPU - __local float smem1[BUFFER+1]; - __local float smem2[BUFFER+1]; - __local float smem3[BUFFER+1]; -#else __local float smem1[BUFFER]; __local float smem2[BUFFER]; __local float smem3[BUFFER]; -#endif - unsigned int xid=get_local_id(0); - unsigned int yid=get_local_id(1); - unsigned int gid=get_group_id(0); - unsigned int xsize=get_local_size(0); - unsigned int ysize=get_local_size(1); - int xBase, yBase, i, j, k; + unsigned int xid=get_local_id(0); + unsigned int yid=get_local_id(1); + unsigned int gid=get_group_id(0); + unsigned int xsize=get_local_size(0); + unsigned int ysize=get_local_size(1); + int xBase, yBase, i, j, k; - float2 c_halfWin = (float2)((c_winSize_x - 1)>>1, (c_winSize_y - 1)>>1); + float2 c_halfWin = (float2)((c_winSize_x - 1)>>1, (c_winSize_y - 1)>>1); const int tid = mad24(yid, xsize, xid); @@ -432,77 +347,71 @@ __kernel void lkSparse_C1_D5(image2d_t I, image2d_t J, float dIdx_patch[GRIDSIZE][GRIDSIZE]; float dIdy_patch[GRIDSIZE][GRIDSIZE]; - yBase=yid; - { - xBase=xid; - SetPatch(I, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, - &I_patch[0][0], &dIdx_patch[0][0], &dIdy_patch[0][0], - &A11, &A12, &A22); + yBase=yid; + { + xBase=xid; + SetPatch(I, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, + &I_patch[0][0], &dIdx_patch[0][0], &dIdy_patch[0][0], + &A11, &A12, &A22); - xBase+=xsize; - SetPatch(I, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, - &I_patch[0][1], &dIdx_patch[0][1], &dIdy_patch[0][1], - &A11, &A12, &A22); + xBase+=xsize; + SetPatch(I, prevPt.x + xBase + 0.5f, prevPt.y + yBase + 0.5f, + &I_patch[0][1], &dIdx_patch[0][1], &dIdy_patch[0][1], + &A11, &A12, &A22); - xBase+=xsize; - if(xBase>1, (c_winSize_y - 1)>>1); + float2 c_halfWin = (float2)((c_winSize_x - 1)>>1, (c_winSize_y - 1)>>1); const int tid = mad24(yid, xsize, xid); @@ -721,7 +615,7 @@ __kernel void lkSparse_C4_D5(image2d_t I, image2d_t J, return; } - nextPt -= c_halfWin; + nextPt -= c_halfWin; // extract the patch from the first image, compute covariation matrix of derivatives @@ -732,80 +626,74 @@ __kernel void lkSparse_C4_D5(image2d_t I, image2d_t J, float4 I_patch[8]; float4 dIdx_patch[8]; float4 dIdy_patch[8]; - float4 I_add,Dx_add,Dy_add; + float4 I_add,Dx_add,Dy_add; - yBase=yid; - { - xBase=xid; - SetPatch4(I, nextPt.x + xBase + 0.5f, nextPt.y + yBase + 0.5f, - &I_patch[0], &dIdx_patch[0], &dIdy_patch[0], - &A11, &A12, &A22); + yBase=yid; + { + xBase=xid; + SetPatch4(I, nextPt.x + xBase + 0.5f, nextPt.y + yBase + 0.5f, + &I_patch[0], &dIdx_patch[0], &dIdy_patch[0], + &A11, &A12, &A22); - xBase+=xsize; - SetPatch4(I, nextPt.x + xBase + 0.5f, nextPt.y + yBase + 0.5f, - &I_patch[1], &dIdx_patch[1], &dIdy_patch[1], - &A11, &A12, &A22); + xBase+=xsize; + SetPatch4(I, nextPt.x + xBase + 0.5f, nextPt.y + yBase + 0.5f, + &I_patch[1], &dIdx_patch[1], &dIdy_patch[1], + &A11, &A12, &A22); - xBase+=xsize; - if(xBase