generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

many bugs fix for intel's HD4000

Browse Source 
from my svn 855
This commit is contained in:
niko
2012-09-21 16:51:45 +08:00
parent f3bc1aede1
commit 0abe7f3196
34 changed files with 1494 additions and 2584 deletions
Download Patch File Download Diff File Expand all files Collapse all files

13
modules/ocl/src/kernels/arithm_addWeighted.cl
View File

@@ -74,8 +74,17 @@ __kernel void addWeighted_D0 (__global uchar *src1,int src1_step,int src1_offset
int dst_end = mad24(y, dst_step, dst_offset + dst_step1);
int dst_index = mad24(y, dst_step, dst_offset + x & (int)0xfffffffc);
uchar4 src1_data = vload4(0, src1 + src1_index);
uchar4 src2_data = vload4(0, src2 + src2_index);
uchar4 src1_data ,src2_data;
src1_data.x= src1_index+0 >= 0 ? src1[src1_index+0] : 0;
src1_data.y= src1_index+1 >= 0 ? src1[src1_index+1] : 0;
src1_data.z= src1_index+2 >= 0 ? src1[src1_index+2] : 0;
src1_data.w= src1_index+3 >= 0 ? src1[src1_index+3] : 0;
src2_data.x= src2_index+0 >= 0 ? src2[src2_index+0] : 0;
src2_data.y= src2_index+1 >= 0 ? src2[src2_index+1] : 0;
src2_data.z= src2_index+2 >= 0 ? src2[src2_index+2] : 0;
src2_data.w= src2_index+3 >= 0 ? src2[src2_index+3] : 0;
uchar4 dst_data = *((__global uchar4 *)(dst + dst_index));
// short4 tmp = convert_short4_sat(src1_data) * alpha + convert_short4_sat(src2_data) * beta + gama;

8
modules/ocl/src/kernels/arithm_div.cl
View File

@@ -48,12 +48,12 @@
typedef double F ;
typedef double4 F4;
#define convert_F4 convert_double4
#define convert_F convert_double
#define convert_F double
#else
typedef float F;
typedef float4 F4;
#define convert_F4 convert_float4
#define convert_F convert_float
#define convert_F float
#endif
uchar round2_uchar(F v){
@@ -229,7 +229,7 @@ __kernel void arithm_div_D4 (__global int *src1, int src1_step, int src1_offset,
int data1 = *((__global int *)((__global char *)src1 + src1_index));
int data2 = *((__global int *)((__global char *)src2 + src2_index));
F tmp = convert_F(data1) * scalar;
F tmp = (convert_F)(data1) * scalar;
int tmp_data = (tmp == 0 || data2 == 0) ? 0 : round2_int(tmp / (convert_F)(data2));
*((__global int *)((__global char *)dst + dst_index)) =tmp_data;
@@ -253,7 +253,7 @@ __kernel void arithm_div_D5 (__global float *src1, int src1_step, int src1_offse
float data1 = *((__global float *)((__global char *)src1 + src1_index));
float data2 = *((__global float *)((__global char *)src2 + src2_index));
F tmp = convert_F(data1) * scalar;
F tmp = (convert_F)(data1) * scalar;
float tmp_data = (tmp == 0 || data2 == 0) ? 0 : convert_float(tmp / (convert_F)(data2));
*((__global float *)((__global char *)dst + dst_index)) = tmp_data;

147
modules/ocl/src/kernels/arithm_minMaxLoc.cl
View File

@@ -46,65 +46,65 @@
/**************************************PUBLICFUNC*************************************/
#if defined (DOUBLE_SUPPORT)
#pragma OPENCL EXTENSION cl_khr_fp64:enable
#define RES_TYPE double8
#define CONVERT_RES_TYPE convert_double8
#define RES_TYPE double4
#define CONVERT_RES_TYPE convert_double4
#else
#define RES_TYPE float8
#define CONVERT_RES_TYPE convert_float8
#define RES_TYPE float4
#define CONVERT_RES_TYPE convert_float4
#endif
#if defined (DEPTH_0)
#define VEC_TYPE uchar8
#define VEC_TYPE_LOC int8
#define CONVERT_TYPE convert_uchar8
#define CONDITION_FUNC(a,b,c) (convert_int8(a) ? b : c)
#define VEC_TYPE uchar4
#define VEC_TYPE_LOC int4
#define CONVERT_TYPE convert_uchar4
#define CONDITION_FUNC(a,b,c) (convert_int4(a) ? b : c)
#define MIN_VAL 0
#define MAX_VAL 255
#endif
#if defined (DEPTH_1)
#define VEC_TYPE char8
#define VEC_TYPE_LOC int8
#define CONVERT_TYPE convert_char8
#define CONDITION_FUNC(a,b,c) (convert_int8(a) ? b : c)
#define VEC_TYPE char4
#define VEC_TYPE_LOC int4
#define CONVERT_TYPE convert_char4
#define CONDITION_FUNC(a,b,c) (convert_int4(a) ? b : c)
#define MIN_VAL -128
#define MAX_VAL 127
#endif
#if defined (DEPTH_2)
#define VEC_TYPE ushort8
#define VEC_TYPE_LOC int8
#define CONVERT_TYPE convert_ushort8
#define CONDITION_FUNC(a,b,c) (convert_int8(a) ? b : c)
#define VEC_TYPE ushort4
#define VEC_TYPE_LOC int4
#define CONVERT_TYPE convert_ushort4
#define CONDITION_FUNC(a,b,c) (convert_int4(a) ? b : c)
#define MIN_VAL 0
#define MAX_VAL 65535
#endif
#if defined (DEPTH_3)
#define VEC_TYPE short8
#define VEC_TYPE_LOC int8
#define CONVERT_TYPE convert_short8
#define CONDITION_FUNC(a,b,c) (convert_int8(a) ? b : c)
#define VEC_TYPE short4
#define VEC_TYPE_LOC int4
#define CONVERT_TYPE convert_short4
#define CONDITION_FUNC(a,b,c) (convert_int4(a) ? b : c)
#define MIN_VAL -32768
#define MAX_VAL 32767
#endif
#if defined (DEPTH_4)
#define VEC_TYPE int8
#define VEC_TYPE_LOC int8
#define CONVERT_TYPE convert_int8
#define VEC_TYPE int4
#define VEC_TYPE_LOC int4
#define CONVERT_TYPE convert_int4
#define CONDITION_FUNC(a,b,c) ((a) ? b : c)
#define MIN_VAL INT_MIN
#define MAX_VAL INT_MAX
#endif
#if defined (DEPTH_5)
#define VEC_TYPE float8
#define VEC_TYPE_LOC float8
#define CONVERT_TYPE convert_float8
#define VEC_TYPE float4
#define VEC_TYPE_LOC float4
#define CONVERT_TYPE convert_float4
#define CONDITION_FUNC(a,b,c) ((a) ? b : c)
#define MIN_VAL (-FLT_MAX)
#define MAX_VAL FLT_MAX
#endif
#if defined (DEPTH_6)
#define VEC_TYPE double8
#define VEC_TYPE_LOC double8
#define CONVERT_TYPE convert_double8
#define VEC_TYPE double4
#define VEC_TYPE_LOC double4
#define CONVERT_TYPE convert_double4
#define CONDITION_FUNC(a,b,c) ((a) ? b : c)
#define MIN_VAL (-DBL_MAX)
#define MAX_VAL DBL_MAX
@@ -122,44 +122,22 @@
#if defined (REPEAT_S3)
#define repeat_s(a) a.s0 = a.s3;a.s1 = a.s3;a.s2 = a.s3;
#endif
#if defined (REPEAT_S4)
#define repeat_s(a) a.s0 = a.s4;a.s1 = a.s4;a.s2 = a.s4;a.s3 = a.s4;
#endif
#if defined (REPEAT_S5)
#define repeat_s(a) a.s0 = a.s5;a.s1 = a.s5;a.s2 = a.s5;a.s3 = a.s5;a.s4 = a.s5;
#endif
#if defined (REPEAT_S6)
#define repeat_s(a) a.s0 = a.s6;a.s1 = a.s6;a.s2 = a.s6;a.s3 = a.s6;a.s4 = a.s6;a.s5 = a.s6;
#endif
#if defined (REPEAT_S7)
#define repeat_s(a) a.s0 = a.s7;a.s1 = a.s7;a.s2 = a.s7;a.s3 = a.s7;a.s4 = a.s7;a.s5 = a.s7;a.s6 = a.s7;
#endif
#if defined (REPEAT_E0)
#define repeat_e(a) a=a;
#endif
#if defined (REPEAT_E1)
#define repeat_e(a) a.s7 = a.s6;
#define repeat_e(a) a.s3 = a.s2;
#endif
#if defined (REPEAT_E2)
#define repeat_e(a) a.s7 = a.s5;a.s6 = a.s5;
#define repeat_e(a) a.s3 = a.s1;a.s2 = a.s1;
#endif
#if defined (REPEAT_E3)
#define repeat_e(a) a.s7 = a.s4;a.s6 = a.s4;a.s5 = a.s4;
#endif
#if defined (REPEAT_E4)
#define repeat_e(a) a.s7 = a.s3;a.s6 = a.s3;a.s5 = a.s3;a.s4 = a.s3;
#endif
#if defined (REPEAT_E5)
#define repeat_e(a) a.s7 = a.s2;a.s6 = a.s2;a.s5 = a.s2;a.s4 = a.s2;a.s3 = a.s2;
#endif
#if defined (REPEAT_E6)
#define repeat_e(a) a.s7 = a.s1;a.s6 = a.s1;a.s5 = a.s1;a.s4 = a.s1;a.s3 = a.s1;a.s2 = a.s1;
#endif
#if defined (REPEAT_E7)
#define repeat_e(a) a.s7 = a.s0;a.s6 = a.s0;a.s5 = a.s0;a.s4 = a.s0;a.s3 = a.s0;a.s2 = a.s0;a.s1 = a.s0;
#define repeat_e(a) a.s3 = a.s0;a.s2 = a.s0;a.s1 = a.s0;
#endif
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics:enable
#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics:enable
@@ -179,8 +157,8 @@ __kernel void arithm_op_minMaxLoc (int cols,int invalid_cols,int offset,int elem
if(id < elemnum)
{
temp = src[idx];
idx_c = idx << 3;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3,idx_c+4,idx_c+5,idx_c+6,idx_c+7);
idx_c = idx << 2;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3);
if(id % cols == 0 )
{
repeat_s(temp);
@@ -203,13 +181,13 @@ __kernel void arithm_op_minMaxLoc (int cols,int invalid_cols,int offset,int elem
minloc = negative;
maxloc = negative;
}
float8 aaa;
float4 aaa;
for(id=id + (groupnum << 8); id < elemnum;id = id + (groupnum << 8))
{
idx = offset + id + (id / cols) * invalid_cols;
temp = src[idx];
idx_c = idx << 3;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3,idx_c+4,idx_c+5,idx_c+6,idx_c+7);
idx_c = idx << 2;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3);
if(id % cols == 0 )
{
repeat_s(temp);
@@ -224,8 +202,8 @@ __kernel void arithm_op_minMaxLoc (int cols,int invalid_cols,int offset,int elem
maxval = max(maxval,temp);
minloc = CONDITION_FUNC(minval == temp, temploc , minloc);
maxloc = CONDITION_FUNC(maxval == temp, temploc , maxloc);
aaa= convert_float8(maxval == temp);
maxloc = convert_int8(aaa) ? temploc : maxloc;
aaa= convert_float4(maxval == temp);
maxloc = convert_int4(aaa) ? temploc : maxloc;
}
if(lid > 127)
{
@@ -278,47 +256,25 @@ __kernel void arithm_op_minMaxLoc (int cols,int invalid_cols,int offset,int elem
#if defined (REPEAT_S3)
#define repeat_ms(a) a.s0 = 0;a.s1 = 0;a.s2 = 0;
#endif
#if defined (REPEAT_S4)
#define repeat_ms(a) a.s0 = 0;a.s1 = 0;a.s2 = 0;a.s3 = 0;
#endif
#if defined (REPEAT_S5)
#define repeat_ms(a) a.s0 = 0;a.s1 = 0;a.s2 = 0;a.s3 = 0;a.s4 = 0;
#endif
#if defined (REPEAT_S6)
#define repeat_ms(a) a.s0 = 0;a.s1 = 0;a.s2 = 0;a.s3 = 0;a.s4 = 0;a.s5 = 0;
#endif
#if defined (REPEAT_S7)
#define repeat_ms(a) a.s0 = 0;a.s1 = 0;a.s2 = 0;a.s3 = 0;a.s4 = 0;a.s5 = 0;a.s6 = 0;
#endif
#if defined (REPEAT_E0)
#define repeat_me(a) a = a;
#endif
#if defined (REPEAT_E1)
#define repeat_me(a) a.s7 = 0;
#define repeat_me(a) a.s3 = 0;
#endif
#if defined (REPEAT_E2)
#define repeat_me(a) a.s7 = 0;a.s6 = 0;
#define repeat_me(a) a.s3 = 0;a.s2 = 0;
#endif
#if defined (REPEAT_E3)
#define repeat_me(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;
#endif
#if defined (REPEAT_E4)
#define repeat_me(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;a.s4 = 0;
#endif
#if defined (REPEAT_E5)
#define repeat_me(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;a.s4 = 0;a.s3 = 0;
#endif
#if defined (REPEAT_E6)
#define repeat_me(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;a.s4 = 0;a.s3 = 0;a.s2 = 0;
#endif
#if defined (REPEAT_E7)
#define repeat_me(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;a.s4 = 0;a.s3 = 0;a.s2 = 0;a.s1 = 0;
#define repeat_me(a) a.s3 = 0;a.s2 = 0;a.s1 = 0;
#endif
/**************************************Array minMaxLoc mask**************************************/
/*
__kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int elemnum,int groupnum,__global VEC_TYPE *src,
int minvalid_cols,int moffset,__global uchar8 *mask,__global RES_TYPE *dst)
int minvalid_cols,int moffset,__global uchar4 *mask,__global RES_TYPE *dst)
{
unsigned int lid = get_local_id(0);
unsigned int gid = get_group_id(0);
@@ -333,8 +289,8 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int
{
temp = src[idx];
m_temp = CONVERT_TYPE(mask[midx]);
int idx_c = idx << 3;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3,idx_c+4,idx_c+5,idx_c+6,idx_c+7);
int idx_c = idx << 2;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3);
if(id % cols == 0 )
{
repeat_ms(m_temp);
@@ -363,8 +319,8 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int
midx = moffset + id + (id / cols) * minvalid_cols;
temp = src[idx];
m_temp = CONVERT_TYPE(mask[midx]);
int idx_c = idx << 3;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3,idx_c+4,idx_c+5,idx_c+6,idx_c+7);
int idx_c = idx << 2;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3);
if(id % cols == 0 )
{
repeat_ms(m_temp);
@@ -421,3 +377,4 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int
}
}
*/

112
modules/ocl/src/kernels/arithm_minMaxLoc_mask.cl
View File

@@ -46,125 +46,101 @@
/**************************************PUBLICFUNC*************************************/
#if defined (DOUBLE_SUPPORT)
#pragma OPENCL EXTENSION cl_khr_fp64:enable
#define RES_TYPE double8
#define CONVERT_RES_TYPE convert_double8
#define RES_TYPE double4
#define CONVERT_RES_TYPE convert_double4
#else
#define RES_TYPE float8
#define CONVERT_RES_TYPE convert_float8
#define RES_TYPE float4
#define CONVERT_RES_TYPE convert_float4
#endif
#if defined (DEPTH_0)
#define TYPE uchar
#define VEC_TYPE uchar8
#define VEC_TYPE_LOC int8
#define CONVERT_TYPE convert_uchar8
#define CONDITION_FUNC(a,b,c) (convert_int8(a) ? b : c)
#define VEC_TYPE uchar4
#define VEC_TYPE_LOC int4
#define CONVERT_TYPE convert_uchar4
#define CONDITION_FUNC(a,b,c) (convert_int4(a) ? b : c)
#define MIN_VAL 0
#define MAX_VAL 255
#endif
#if defined (DEPTH_1)
#define TYPE char
#define VEC_TYPE char8
#define VEC_TYPE_LOC int8
#define CONVERT_TYPE convert_char8
#define CONDITION_FUNC(a,b,c) (convert_int8(a) ? b : c)
#define VEC_TYPE char4
#define VEC_TYPE_LOC int4
#define CONVERT_TYPE convert_char4
#define CONDITION_FUNC(a,b,c) (convert_int4(a) ? b : c)
#define MIN_VAL -128
#define MAX_VAL 127
#endif
#if defined (DEPTH_2)
#define TYPE ushort
#define VEC_TYPE ushort8
#define VEC_TYPE_LOC int8
#define CONVERT_TYPE convert_ushort8
#define CONDITION_FUNC(a,b,c) (convert_int8(a) ? b : c)
#define VEC_TYPE ushort4
#define VEC_TYPE_LOC int4
#define CONVERT_TYPE convert_ushort4
#define CONDITION_FUNC(a,b,c) (convert_int4(a) ? b : c)
#define MIN_VAL 0
#define MAX_VAL 65535
#endif
#if defined (DEPTH_3)
#define TYPE short
#define VEC_TYPE short8
#define VEC_TYPE_LOC int8
#define CONVERT_TYPE convert_short8
#define CONDITION_FUNC(a,b,c) (convert_int8(a) ? b : c)
#define VEC_TYPE short4
#define VEC_TYPE_LOC int4
#define CONVERT_TYPE convert_short4
#define CONDITION_FUNC(a,b,c) (convert_int4(a) ? b : c)
#define MIN_VAL -32768
#define MAX_VAL 32767
#endif
#if defined (DEPTH_4)
#define TYPE int
#define VEC_TYPE int8
#define VEC_TYPE_LOC int8
#define CONVERT_TYPE convert_int8
#define VEC_TYPE int4
#define VEC_TYPE_LOC int4
#define CONVERT_TYPE convert_int4
#define CONDITION_FUNC(a,b,c) ((a) ? b : c)
#define MIN_VAL INT_MIN
#define MAX_VAL INT_MAX
#endif
#if defined (DEPTH_5)
#define TYPE float
#define VEC_TYPE float8
#define VEC_TYPE_LOC float8
#define CONVERT_TYPE convert_float8
#define VEC_TYPE float4
#define VEC_TYPE_LOC float4
#define CONVERT_TYPE convert_float4
#define CONDITION_FUNC(a,b,c) ((a) ? b : c)
#define MIN_VAL (-FLT_MAX)
#define MAX_VAL FLT_MAX
#endif
#if defined (DEPTH_6)
#define TYPE double
#define VEC_TYPE double8
#define VEC_TYPE_LOC double8
#define CONVERT_TYPE convert_double8
#define VEC_TYPE double4
#define VEC_TYPE_LOC double4
#define CONVERT_TYPE convert_double4
#define CONDITION_FUNC(a,b,c) ((a) ? b : c)
#define MIN_VAL (-DBL_MAX)
#define MAX_VAL DBL_MAX
#endif
#if defined (REPEAT_E0)
#define repeat_e(a) a = a;
#define repeat_e(a) a=a;
#endif
#if defined (REPEAT_E1)
#define repeat_e(a) a.s7 = a.s6;
#define repeat_e(a) a.s3 = a.s2;
#endif
#if defined (REPEAT_E2)
#define repeat_e(a) a.s7 = a.s5;a.s6 = a.s5;
#define repeat_e(a) a.s3 = a.s1;a.s2 = a.s1;
#endif
#if defined (REPEAT_E3)
#define repeat_e(a) a.s7 = a.s4;a.s6 = a.s4;a.s5 = a.s4;
#endif
#if defined (REPEAT_E4)
#define repeat_e(a) a.s7 = a.s3;a.s6 = a.s3;a.s5 = a.s3;a.s4 = a.s3;
#endif
#if defined (REPEAT_E5)
#define repeat_e(a) a.s7 = a.s2;a.s6 = a.s2;a.s5 = a.s2;a.s4 = a.s2;a.s3 = a.s2;
#endif
#if defined (REPEAT_E6)
#define repeat_e(a) a.s7 = a.s1;a.s6 = a.s1;a.s5 = a.s1;a.s4 = a.s1;a.s3 = a.s1;a.s2 = a.s1;
#endif
#if defined (REPEAT_E7)
#define repeat_e(a) a.s7 = a.s0;a.s6 = a.s0;a.s5 = a.s0;a.s4 = a.s0;a.s3 = a.s0;a.s2 = a.s0;a.s1 = a.s0;
#define repeat_e(a) a.s3 = a.s0;a.s2 = a.s0;a.s1 = a.s0;
#endif
#if defined (REPEAT_E0)
#define repeat_me(a) a = a;
#endif
#if defined (REPEAT_E1)
#define repeat_me(a) a.s7 = 0;
#define repeat_me(a) a.s3 = 0;
#endif
#if defined (REPEAT_E2)
#define repeat_me(a) a.s7 = 0;a.s6 = 0;
#define repeat_me(a) a.s3 = 0;a.s2 = 0;
#endif
#if defined (REPEAT_E3)
#define repeat_me(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;
#endif
#if defined (REPEAT_E4)
#define repeat_me(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;a.s4 = 0;
#endif
#if defined (REPEAT_E5)
#define repeat_me(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;a.s4 = 0;a.s3 = 0;
#endif
#if defined (REPEAT_E6)
#define repeat_me(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;a.s4 = 0;a.s3 = 0;a.s2 = 0;
#endif
#if defined (REPEAT_E7)
#define repeat_me(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;a.s4 = 0;a.s3 = 0;a.s2 = 0;a.s1 = 0;
#define repeat_me(a) a.s3 = 0;a.s2 = 0;a.s1 = 0;
#endif
/**************************************Array minMaxLoc mask**************************************/
@@ -182,10 +158,10 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int
VEC_TYPE_LOC minloc,maxloc,temploc,negative = -1,one = 1,zero = 0;
if(id < elemnum)
{
temp = vload8(idx, &src[offset]);
m_temp = CONVERT_TYPE(vload8(midx,&mask[moffset]));
int idx_c = (idx << 3) + offset;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3,idx_c+4,idx_c+5,idx_c+6,idx_c+7);
temp = vload4(idx, &src[offset]);
m_temp = CONVERT_TYPE(vload4(midx,&mask[moffset]));
int idx_c = (idx << 2) + offset;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3);
if(id % cols == cols - 1)
{
repeat_me(m_temp);
@@ -207,10 +183,10 @@ __kernel void arithm_op_minMaxLoc_mask (int cols,int invalid_cols,int offset,int
{
idx = id + (id / cols) * invalid_cols;
midx = id + (id / cols) * minvalid_cols;
temp = vload8(idx, &src[offset]);
m_temp = CONVERT_TYPE(vload8(midx,&mask[moffset]));
int idx_c = (idx << 3) + offset;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3,idx_c+4,idx_c+5,idx_c+6,idx_c+7);
temp = vload4(idx, &src[offset]);
m_temp = CONVERT_TYPE(vload4(midx,&mask[moffset]));
int idx_c = (idx << 2) + offset;
temploc = (VEC_TYPE_LOC)(idx_c,idx_c+1,idx_c+2,idx_c+3);
if(id % cols == cols - 1)
{
repeat_me(m_temp);

13
modules/ocl/src/kernels/arithm_mul.cl
View File

@@ -92,8 +92,17 @@ __kernel void arithm_mul_D0 (__global uchar *src1, int src1_step, int src1_offse
int dst_end = mad24(y, dst_step, dst_offset + dst_step1);
int dst_index = mad24(y, dst_step, dst_offset + x & (int)0xfffffffc);
uchar4 src1_data = vload4(0, src1 + src1_index);
uchar4 src2_data = vload4(0, src2 + src2_index);
uchar4 src1_data ,src2_data;
src1_data.x= src1_index+0 >= 0 ? src1[src1_index+0] : 0;
src1_data.y= src1_index+1 >= 0 ? src1[src1_index+1] : 0;
src1_data.z= src1_index+2 >= 0 ? src1[src1_index+2] : 0;
src1_data.w= src1_index+3 >= 0 ? src1[src1_index+3] : 0;
src2_data.x= src2_index+0 >= 0 ? src2[src2_index+0] : 0;
src2_data.y= src2_index+1 >= 0 ? src2[src2_index+1] : 0;
src2_data.z= src2_index+2 >= 0 ? src2[src2_index+2] : 0;
src2_data.w= src2_index+3 >= 0 ? src2[src2_index+3] : 0;
uchar4 dst_data = *((__global uchar4 *)(dst + dst_index));
int4 tmp = convert_int4_sat(src1_data) * convert_int4_sat(src2_data);

15
modules/ocl/src/kernels/arithm_pow.cl
View File

@@ -45,13 +45,19 @@
#if defined (DOUBLE_SUPPORT)
#pragma OPENCL EXTENSION cl_khr_fp64:enable
typedef double F;
typedef double4 F4;
#define convert_F4 convert_double4;
#else
typedef float F;
typedef float4 F4;
#define convert_F4 convert_float4;
#endif
/************************************** pow **************************************/
#if defined (DOUBLE_SUPPORT)
__kernel void arithm_pow_D5 (__global float *src1, int src1_step, int src1_offset,
__global float *dst, int dst_step, int dst_offset,
int rows, int cols, int dst_step1,
double p)
F p)
{
int x = get_global_id(0);
@@ -69,14 +75,12 @@ __kernel void arithm_pow_D5 (__global float *src1, int src1_step, int src1_offse
}
}
#endif
#if defined (DOUBLE_SUPPORT)
__kernel void arithm_pow_D6 (__global double *src1, int src1_step, int src1_offset,
__global double *dst, int dst_step, int dst_offset,
int rows, int cols, int dst_step1,
double p)
F p)
{
int x = get_global_id(0);
@@ -94,4 +98,3 @@ __kernel void arithm_pow_D6 (__global double *src1, int src1_step, int src1_offs
}
#endif

1
modules/ocl/src/kernels/convertC3C4.cl
View File

@@ -123,7 +123,6 @@ __kernel void convertC4C3(__global const GENTYPE4 * restrict src, __global GENTY
int4 outaddr = mul24(id>>2 , 3);
outaddr.y++;
outaddr.z+=2;
//printf("%d ",outaddr.z);
if(outaddr.z <= pixel_end)
{
dst[outaddr.x] = pixel0;

29
modules/ocl/src/kernels/filtering_boxFilter.cl
View File

@@ -238,7 +238,9 @@ __kernel void boxFilter_C4_D0(__global const uchar4 * restrict src, __global uch
int startY = (gY << 1) - anY + src_y_off;
int dst_startX = gX * (THREADS-ksX+1) + dst_x_off;
int dst_startY = (gY << 1) + dst_y_off;
int end_addr = (src_whole_rows-1)*(src_step>>2) + src_whole_cols-4;
//int end_addr = (src_whole_rows-1)*(src_step>>2) + src_whole_cols-4;
int end_addr = src_whole_cols-4;
uint4 data[ksY+1];
__local uint4 temp[2][THREADS];
#ifdef BORDER_CONSTANT
@@ -247,8 +249,13 @@ __kernel void boxFilter_C4_D0(__global const uchar4 * restrict src, __global uch
for(int i=0; i < ksY+1; i++)
{
con = startX+col >= 0 && startX+col < src_whole_cols && startY+i >= 0 && startY+i < src_whole_rows;
int cur_addr = clamp((startY+i)*(src_step>>2)+(startX+col),0,end_addr);
ss = convert_uint4(src[cur_addr]);
//int cur_addr = clamp((startY+i)*(src_step>>2)+(startX+col),0,end_addr);
//ss = convert_uint4(src[cur_addr]);
int cur_col = clamp(startX + col, 0, src_whole_cols);
ss = convert_uint4(src[(startY+i)*(src_step>>2) + cur_col]);
data[i] = con ? ss : 0;
}
#else
@@ -327,8 +334,12 @@ __kernel void boxFilter_C1_D5(__global const float *restrict src, __global float
for(int i=0; i < ksY+1; i++)
{
con = startX+col >= 0 && startX+col < src_whole_cols && startY+i >= 0 && startY+i < src_whole_rows;
int cur_addr = clamp((startY+i)*(src_step>>2)+(startX+col),0,end_addr);
ss = src[cur_addr];
// int cur_addr = clamp((startY+i)*(src_step>>2)+(startX+col),0,end_addr);
// ss = src[cur_addr];
int cur_col = clamp(startX + col, 0, src_whole_cols);
ss = src[(startY+i)*(src_step>>2) + cur_col];
data[i] = con ? ss : 0.f;
}
#else
@@ -407,8 +418,12 @@ __kernel void boxFilter_C4_D5(__global const float4 *restrict src, __global floa
for(int i=0; i < ksY+1; i++)
{
con = startX+col >= 0 && startX+col < src_whole_cols && startY+i >= 0 && startY+i < src_whole_rows;
int cur_addr = clamp((startY+i)*(src_step>>4)+(startX+col),0,end_addr);
ss = src[cur_addr];
//int cur_addr = clamp((startY+i)*(src_step>>4)+(startX+col),0,end_addr);
//ss = src[cur_addr];
int cur_col = clamp(startX + col, 0, src_whole_cols);
ss = src[(startY+i)*(src_step>>4) + cur_col];
data[i] = con ? ss : (float4)(0.0,0.0,0.0,0.0);
}
#else

1331
modules/ocl/src/kernels/img_proc.cl
View File

File diff suppressed because it is too large Load Diff

93
modules/ocl/src/kernels/imgproc_bilateral.cl
View File

@@ -108,71 +108,38 @@ void bilateral4(__global uchar4 *dst,
dst[index_dst] = convert_uchar4_rte(pd);
}
__kernel
void bilateral(__global uchar *dst,
__global uchar *src,
int rows,
int cols,
int channels,
int radius,
int wholerows,
int wholecols,
int src_step,
int dst_step,
int src_offset,
int dst_offset,
__constant float *sigClr,
__constant float *sigSpc)
{
uint lidx = get_local_id(0);
uint lidy = get_local_id(1);
uint gdx = get_global_id(0);
uint gdy = get_global_id(1);
uint gidx = gdx >=cols?cols-1:gdx;
uint gidy = gdy >=rows?rows-1:gdy;
uchar p,q,tmp;
float pf = 0,pq = 0,wt = 0,pd = 0;
int r =radius;
int ij = 0;
int ct = 0;
uint index_src = src_offset + gidy*src_step + gidx;
uint index_dst = dst_offset + gidy*dst_step + gidx;
p = src[index_src];
uint gx,gy;
uint src_index,dst_index;
for(int ii = -r;ii<r+1;ii++)
__kernel void bilateral(__global uchar *dst,
__global const uchar *src,
const int dst_rows,
const int dst_cols,
const int maxk,
const int radius,
const int dst_step,
const int dst_offset,
const int src_step,
const int src_rows,
const int src_cols,
__constant float *color_weight,
__constant float *space_weight,
__constant int *space_ofs)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
if((gidy<dst_rows) && (gidx<dst_cols))
{
for(int jj =-r;jj<r+1;jj++)
{
ij = ii*ii+jj*jj;
if(ij > mul24(radius,radius)) continue;
int src_addr = mad24(gidy+radius,src_step,gidx+radius);
int dst_addr = mad24(gidy,src_step,gidx+dst_offset);
float sum = 0, wsum = 0;
gx = gidx + jj;
gy = gidy + ii;
src_index = src_offset + gy * src_step + gx;
q = src[src_index];
ct = abs(p-q);
wt =sigClr[ct]*sigSpc[(ii+radius)*(2*radius+1)+jj+radius];
pf += q*wt;
pq += wt;
}
int val0 = (int)src[src_addr];
for(int k = 0; k < maxk; k++ )
{
int val = (int)src[src_addr + space_ofs[k]];
float w = space_weight[k]*color_weight[abs(val - val0)];
sum += (float)(val)*w;
wsum += w;
}
dst[dst_addr] = convert_uchar_rtz(sum/wsum+0.5f);
}
pd = pf/pq;
dst[index_dst] = convert_uchar_rte(pd);
}

14
modules/ocl/src/kernels/imgproc_calcHarris.cl
View File

@@ -65,8 +65,8 @@
#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-1+((b_edge)<<1) : (addr))
#endif
#ifdef BORDER_REFLECT_101
//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
#ifdef BORDER_REFLECT101
//BORDER_REFLECT101: gfedcb|abcdefgh|gfedcba
#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i) : (i))
#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr))
#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i) : (i))
@@ -95,6 +95,8 @@ __kernel void calcHarris(__global const float *Dx,__global const float *Dy, __gl
int col = get_local_id(0);
const int gX = get_group_id(0);
const int gY = get_group_id(1);
const int glx = get_global_id(0);
const int gly = get_global_id(1);
int dx_x_off = (dx_offset % dx_step) >> 2;
int dx_y_off = dx_offset / dx_step;
@@ -118,10 +120,10 @@ __kernel void calcHarris(__global const float *Dx,__global const float *Dy, __gl
for(int i=0; i < ksY+1; i++)
{
dx_con = dx_startX+col >= 0 && dx_startX+col < dx_whole_cols && dx_startY+i >= 0 && dx_startY+i < dx_whole_rows;
dx_s = Dx[(dx_startY+i)*(dx_step>>2)+(dx_startX+dx_col)];
dx_s = Dx[(dx_startY+i)*(dx_step>>2)+(dx_startX+col)];
dx_data[i] = dx_con ? dx_s : 0.0;
dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows;
dy_s = Dy[(dy_startY+i)*(dy_step>>2)+(dy_startX+dy_col)];
dy_s = Dy[(dy_startY+i)*(dy_step>>2)+(dy_startX+col)];
dy_data[i] = dy_con ? dy_s : 0.0;
data[0][i] = dx_data[i] * dx_data[i];
data[1][i] = dx_data[i] * dy_data[i];
@@ -144,7 +146,7 @@ __kernel void calcHarris(__global const float *Dx,__global const float *Dy, __gl
dy_selected_row = ADDR_B(dy_startY+i, dy_whole_rows, dy_selected_row);
dy_selected_col = ADDR_L(dy_startX+col, 0, dy_whole_cols);
dy_selected_col = ADDR_R(dy_startX+col, dy_whole_cols, dy_selected_col);
dy_data[i] = Dy[dx_selected_row * (dy_step>>2) + dy_selected_col];
dy_data[i] = Dy[dy_selected_row * (dy_step>>2) + dy_selected_col];
data[0][i] = dx_data[i] * dx_data[i];
data[1][i] = dx_data[i] * dy_data[i];
@@ -176,7 +178,7 @@ __kernel void calcHarris(__global const float *Dx,__global const float *Dy, __gl
{
col += anX;
int posX = dst_startX - dst_x_off + col - anX;
int posY = (gY << 1);
int posY = (gly << 1);
int till = (ksX + 1)%2;
float tmp_sum[6]={ 0.0, 0.0 , 0.0, 0.0, 0.0, 0.0 };
for(int k=0; k<6; k++)

14
modules/ocl/src/kernels/imgproc_calcMinEigenVal.cl
View File

@@ -65,8 +65,8 @@
#define ADDR_B(i, b_edge, addr) ((i) >= (b_edge) ? -(i)-1+((b_edge)<<1) : (addr))
#endif
#ifdef BORDER_REFLECT_101
//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
#ifdef BORDER_REFLECT101
//BORDER_REFLECT101: gfedcb|abcdefgh|gfedcba
#define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i) : (i))
#define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr))
#define ADDR_H(i, t_edge, b_edge) ((i) < (t_edge) ? -(i) : (i))
@@ -95,6 +95,8 @@ __kernel void calcMinEigenVal(__global const float *Dx,__global const float *Dy,
int col = get_local_id(0);
const int gX = get_group_id(0);
const int gY = get_group_id(1);
const int glx = get_global_id(0);
const int gly = get_global_id(1);
int dx_x_off = (dx_offset % dx_step) >> 2;
int dx_y_off = dx_offset / dx_step;
@@ -118,10 +120,10 @@ __kernel void calcMinEigenVal(__global const float *Dx,__global const float *Dy,
for(int i=0; i < ksY+1; i++)
{
dx_con = dx_startX+col >= 0 && dx_startX+col < dx_whole_cols && dx_startY+i >= 0 && dx_startY+i < dx_whole_rows;
dx_s = Dx[(dx_startY+i)*(dx_step>>2)+(dx_startX+dx_col)];
dx_s = Dx[(dx_startY+i)*(dx_step>>2)+(dx_startX+col)];
dx_data[i] = dx_con ? dx_s : 0.0;
dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows;
dy_s = Dy[(dy_startY+i)*(dy_step>>2)+(dy_startX+dy_col)];
dy_s = Dy[(dy_startY+i)*(dy_step>>2)+(dy_startX+col)];
dy_data[i] = dy_con ? dy_s : 0.0;
data[0][i] = dx_data[i] * dx_data[i];
data[1][i] = dx_data[i] * dy_data[i];
@@ -144,7 +146,7 @@ __kernel void calcMinEigenVal(__global const float *Dx,__global const float *Dy,
dy_selected_row = ADDR_B(dy_startY+i, dy_whole_rows, dy_selected_row);
dy_selected_col = ADDR_L(dy_startX+col, 0, dy_whole_cols);
dy_selected_col = ADDR_R(dy_startX+col, dy_whole_cols, dy_selected_col);
dy_data[i] = Dy[dx_selected_row * (dy_step>>2) + dy_selected_col];
dy_data[i] = Dy[dy_selected_row * (dy_step>>2) + dy_selected_col];
data[0][i] = dx_data[i] * dx_data[i];
data[1][i] = dx_data[i] * dy_data[i];
@@ -176,7 +178,7 @@ __kernel void calcMinEigenVal(__global const float *Dx,__global const float *Dy,
{
col += anX;
int posX = dst_startX - dst_x_off + col - anX;
int posY = (gY << 1);
int posY = (gly << 1);
int till = (ksX + 1)%2;
float tmp_sum[6]={ 0.0, 0.0 , 0.0, 0.0, 0.0, 0.0 };
for(int k=0; k<6; k++)

3
modules/ocl/src/kernels/imgproc_canny.cl
View File

@@ -43,7 +43,6 @@
//
//M*/
#pragma OPENCL EXTENSION cl_amd_printf : enable
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
@@ -651,7 +650,7 @@ __kernel
}
__constant int c_dx[8] = {-1, 0, 1, -1, 1, -1, 0, 1};
__constant c_dy[8] = {-1, -1, -1, 0, 0, 1, 1, 1};
__constant int c_dy[8] = {-1, -1, -1, 0, 0, 1, 1, 1};
#define stack_size 512
__kernel

346
modules/ocl/src/kernels/imgproc_copymakeboder.cl
View File

@@ -35,212 +35,166 @@
//
#define get(a,b,c) (( b >= top & b < srcRows+top & a >= left & a < srcCols+left )? c : 8)
__kernel void copyConstBorder_C1_D0(__global uchar * src, __global uchar * dst, int srcOffset, int dstOffset,
int srcCols, int srcRows, int dstCols, int dstRows,
int top, int left, uchar nVal, int srcStep, int dstStep)
{
int idx = get_global_id(0);
int tpr = (dstCols + 3 + (dstOffset&3))>>2;
int dx = ((idx%(tpr))<<2) - (dstOffset&3);
int dy = idx/(tpr);
__global uchar4 * d=(__global uchar4 *)(dst + dstOffset + dy*dstStep + dx);
int start=srcOffset + (dy-top)*srcStep + (dx-left);
uchar8 s=*((__global uchar8 *)(src + ((start>>2)<<2) ));
uchar4 v;
uchar sv[9]={s.s0,s.s1,s.s2,s.s3,s.s4,s.s5,s.s6,s.s7,nVal};
int det=start&3;
v.x=sv[get(dx,dy,det)];
v.y=sv[get(dx+1,dy,det+1)];
v.z=sv[get(dx+2,dy,det+2)];
v.w=sv[get(dx+3,dy,det+3)];
if(dy<dstRows)
{
uchar4 res = *d;
res.x = (dx>=0 && dx<dstCols) ? v.x : res.x;
res.y = (dx+1>=0 && dx+1<dstCols) ? v.y : res.y;
res.z = (dx+2>=0 && dx+2<dstCols) ? v.z : res.z;
res.w = (dx+3>=0 && dx+3<dstCols) ? v.w : res.w;
*d=res;
}
}
#undef get(a,b,c)
#ifdef BORDER_CONSTANT
//BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii
#define ELEM(i,l_edge,r_edge,elem1,elem2) (i)<(l_edge) | (i) >= (r_edge) ? (elem1) : (elem2)
#endif
#define get(a,b,c,d) (( b >= top & b < srcRows+top & a >= left & a < srcCols+left )? c : d)
__kernel void copyConstBorder_C1_D4(__global int * src, __global int * dst, int srcOffset, int dstOffset,
int srcCols, int srcRows, int dstCols, int dstRows,
int top, int left, int nVal, int srcStep, int dstStep)
{
int idx = get_global_id(0);
int tpr = (dstCols + 3)>>2;
int dx = (idx%(tpr))<<2;
int dy = idx/(tpr);
__global int4 * d=(__global int4 *)(dst+dy*dstStep+dx);
int4 s=*((__global int4 *)(src + srcOffset + (dy-top)*srcStep + (dx-left) ));
int4 v;
v.x=get(dx,dy,s.x,nVal);
v.y=get(dx+1,dy,s.y,nVal);
v.z=get(dx+2,dy,s.z,nVal);
v.w=get(dx+3,dy,s.w,nVal);
if(dy<dstRows)
{
int4 res = *d;
v.y = (dx+1<dstCols) ? v.y : res.y;
v.z = (dx+2<dstCols) ? v.z : res.z;
v.w = (dx+3<dstCols) ? v.w : res.w;
*d=v;
}
}
#undef get(a,b,c,d)
#ifdef BORDER_REPLICATE
//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh
#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? (l_edge) : (addr)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (r_edge)-1 : (addr)
#endif
#define get(a,b,c) ( a < srcCols+left ? b : c)
__kernel void copyReplicateBorder_C1_D4(__global int * src, __global int * dst, int srcOffset, int dstOffset,
int srcCols, int srcRows, int dstCols, int dstRows,
int top, int left, int nVal, int srcStep, int dstStep)
{
int idx = get_global_id(0);
int tpr = (dstCols + 3)>>2;
int dx = (idx%(tpr))<<2;
int dy = idx/(tpr);
__global int4 * d=(__global int4 *)(dst + dstOffset + dy*dstStep + dx);
int c=clamp(dx-left,0,srcCols-1);
int4 s=*((__global int4 *)(src + srcOffset + clamp(dy-top,0,srcRows-1) * srcStep + c ));
int sa[4]={s.x,s.y,s.z,s.w};
int4 v;
v.x=get(dx,sa[max(0,(dx-left)-c)],sa[srcCols-1-c]);
v.y=get(dx+1,sa[max(0,(dx+1-left)-c)],sa[srcCols-1-c]);
v.z=get(dx+2,sa[max(0,(dx+2-left)-c)],sa[srcCols-1-c]);
v.w=get(dx+3,sa[max(0,(dx+3-left)-c)],sa[srcCols-1-c]);
if(dy<dstRows)
{
int4 res = *d;
v.y = (dx+1<dstCols) ? v.y : res.y;
v.z = (dx+2<dstCols) ? v.z : res.z;
v.w = (dx+3<dstCols) ? v.w : res.w;
*d=v;
}
}
__kernel void copyReplicateBorder_C1_D0(__global uchar * src, __global uchar * dst, int srcOffset, int dstOffset,
int srcCols, int srcRows, int dstCols, int dstRows,
int top, int left, uchar nVal, int srcStep, int dstStep)
{
int idx = get_global_id(0);
int tpr = (dstCols + 3 + (dstOffset&3))>>2;
int dx = ((idx%(tpr))<<2) - (dstOffset&3);
int dy = idx/(tpr);
__global uchar4 * d=(__global uchar4 *)(dst + dstOffset + dy*dstStep + dx);
int c=clamp(dx-left,0,srcCols-1);
int start= srcOffset + clamp(dy-top,0,srcRows-1) * srcStep + c;
uchar8 s=*((__global uchar8 *)(src + ((start>>2)<<2) ));
uchar4 v;
uchar sa[8]={s.s0,s.s1,s.s2,s.s3,s.s4,s.s5,s.s6,s.s7};
int det=start&3;
v.x=get(dx,sa[max(0,(dx-left)-c)+det],sa[srcCols-1-c+det]);
v.y=get(dx+1,sa[max(0,(dx+1-left)-c)+det],sa[srcCols-1-c+det]);
v.z=get(dx+2,sa[max(0,(dx+2-left)-c)+det],sa[srcCols-1-c+det]);
v.w=get(dx+3,sa[max(0,(dx+3-left)-c)+det],sa[srcCols-1-c+det]);
if(dy<dstRows)
{
uchar4 res = *d;
res.x = (dx>=0 && dx<dstCols) ? v.x : res.x;
res.y = (dx+1>=0 && dx+1<dstCols) ? v.y : res.y;
res.z = (dx+2>=0 && dx+2<dstCols) ? v.z : res.z;
res.w = (dx+3>=0 && dx+3<dstCols) ? v.w : res.w;
*d=res;
}
}
#undef get(a,b,c)
#ifdef BORDER_REFLECT
//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb
#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? -(i)-1 : (addr)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr)
#endif
#ifdef BORDER_REFLECT_101
//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba
#define edge(x,size,rx) rx = abs(x) % ((size<<1)-2); rx = (rx>=size?(size<<1)-2:rx<<1) - rx;
__kernel void copyReflectBorder_C1_D4(__global int * src, __global int * dst, int srcOffset, int dstOffset,
int srcCols, int srcRows, int dstCols, int dstRows,
int top, int left, int nVal, int srcStep, int dstStep)
#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? -(i) : (addr)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr)
#endif
#ifdef BORDER_WRAP
//BORDER_WRAP: cdefgh|abcdefgh|abcdefg
#define ADDR_L(i,l_edge,r_edge,addr) (i) < (l_edge) ? (i)+(r_edge) : (addr)
#define ADDR_R(i,r_edge,addr) (i) >= (r_edge) ? (i)-(r_edge) : (addr)
#endif
__kernel void copymakeborder
(__global const GENTYPE *src,
__global GENTYPE *dst,
const int dst_cols,
const int dst_rows,
const int src_cols,
const int src_rows,
const int src_step_in_pixel,
const int src_offset_in_pixel,
const int dst_step_in_pixel,
const int dst_offset_in_pixel,
const int top,
const int left,
const GENTYPE val
)
{
int idx = get_global_id(0);
int tpr = (dstCols + 3)>>2;
int dx = (idx%(tpr))<<2;
int dy = idx/(tpr);
__global int4 * d=(__global int4 *)(dst + dstOffset + dy*dstStep + dx);
uint4 id;
edge(dx-left,srcCols,id.x);
edge(dx-left+1,srcCols,id.x);
edge(dx-left+2,srcCols,id.x);
edge(dx-left+3,srcCols,id.x);
int start=min(id.x,id.w);
int4 s=*((__global int4 *)(src + srcOffset + clamp(dy-top,0,srcRows-1) * srcStep + start));
int sa[4]={s.x,s.y,s.z,s.w};
int4 v=(int4)(sa[(id.x-start)],sa[(id.y-start)],sa[(id.z-start)],sa[(id.w-start)]);
if(dy<dstRows)
int x = get_global_id(0);
int y = get_global_id(1);
int src_x = x-left;
int src_y = y-top;
int src_addr = mad24(src_y,src_step_in_pixel,src_x+src_offset_in_pixel);
int dst_addr = mad24(y,dst_step_in_pixel,x+dst_offset_in_pixel);
int con = (src_x >= 0) && (src_x < src_cols) && (src_y >= 0) && (src_y < src_rows);
if(con)
{
int4 res = *d;
v.y = (dx+1<dstCols) ? v.y : res.y;
v.z = (dx+2<dstCols) ? v.z : res.z;
v.w = (dx+3<dstCols) ? v.w : res.w;
*d=v;
dst[dst_addr] = src[src_addr];
}
else
{
#ifdef BORDER_CONSTANT
//write the result to dst
if((x<dst_cols) && (y<dst_rows))
{
dst[dst_addr] = val;
}
#else
int s_x,s_y;
//judge if read out of boundary
s_x= ADDR_L(src_x,0,src_cols,src_x);
s_x= ADDR_R(src_x,src_cols,s_x);
s_y= ADDR_L(src_y,0,src_rows,src_y);
s_y= ADDR_R(src_y,src_rows,s_y);
src_addr=mad24(s_y,src_step_in_pixel,s_x+src_offset_in_pixel);
//write the result to dst
if((x<dst_cols) && (y<dst_rows))
{
dst[dst_addr] = src[src_addr];
}
#endif
}
}
__kernel void copyReflectBorder_C1_D0(__global uchar * src, __global uchar * dst, int srcOffset, int dstOffset,
int srcCols, int srcRows, int dstCols, int dstRows,
int top, int left, uchar nVal, int srcStep, int dstStep)
__kernel void copymakeborder_C1_D0
(__global const uchar *src,
__global uchar *dst,
const int dst_cols,
const int dst_rows,
const int src_cols,
const int src_rows,
const int src_step_in_pixel,
const int src_offset_in_pixel,
const int dst_step_in_pixel,
const int dst_offset_in_pixel,
const int top,
const int left,
const uchar val
)
{
int idx = get_global_id(0);
int tpr = (dstCols + 3 + (dstOffset&3))>>2;
int dx = ((idx%(tpr))<<2) - (dstOffset&3);
int dy = idx/(tpr);
__global uchar4 * d=(__global uchar4 *)(dst + dstOffset + dy*dstStep + dx);
uint4 id;
edge(dx-left,srcCols,id.x);
edge(dx-left+1,srcCols,id.x);
edge(dx-left+2,srcCols,id.x);
edge(dx-left+3,srcCols,id.x);
int start=min(id.x,id.w) + srcOffset;
uchar8 s=*((__global uchar8 *)(src + clamp(dy-top,0,srcRows-1) * srcStep + ((start>>2)<<2) ));
uchar sa[8]={s.s0,s.s1,s.s2,s.s3,s.s4,s.s5,s.s6,s.s7};
int det=start&3;
uchar4 v=(uchar4)(sa[(id.x-start)+det],sa[(id.y-start)+det],sa[(id.z-start)+det],sa[(id.w-start)+det]);
if(dy<dstRows)
int x = get_global_id(0)<<2;
int y = get_global_id(1);
int src_x = x-left;
int src_y = y-top;
int src_addr = mad24(src_y,src_step_in_pixel,src_x+src_offset_in_pixel);
int dst_addr = mad24(y,dst_step_in_pixel,x+dst_offset_in_pixel);
int con = (src_x >= 0) && (src_x+3 < src_cols) && (src_y >= 0) && (src_y < src_rows);
if(con)
{
uchar4 res = *d;
res.x = (dx>=0 && dx<dstCols) ? v.x : res.x;
res.y = (dx+1>=0 && dx+1<dstCols) ? v.y : res.y;
res.z = (dx+2>=0 && dx+2<dstCols) ? v.z : res.z;
res.w = (dx+3>=0 && dx+3<dstCols) ? v.w : res.w;
*d=res;
uchar4 tmp = vload4(0,src+src_addr);
*(__global uchar4*)(dst+dst_addr) = tmp;
}
else
{
#ifdef BORDER_CONSTANT
//write the result to dst
if((((src_x<0) && (src_x+3>=0))||(src_x < src_cols) && (src_x+3 >= src_cols)) && (src_y >= 0) && (src_y < src_rows))
{
int4 addr;
uchar4 tmp;
addr.x = ((src_x < 0) || (src_x>= src_cols)) ? 0 : src_addr;
addr.y = ((src_x+1 < 0) || (src_x+1>= src_cols)) ? 0 : (src_addr+1);
addr.z = ((src_x+2 < 0) || (src_x+2>= src_cols)) ? 0 : (src_addr+2);
addr.w = ((src_x+3 < 0) || (src_x+3>= src_cols)) ? 0 : (src_addr+3);
tmp.x = src[addr.x];
tmp.y = src[addr.y];
tmp.z = src[addr.z];
tmp.w = src[addr.w];
tmp.x = (src_x >=0)&&(src_x < src_cols) ? tmp.x : val;
tmp.y = (src_x+1 >=0)&&(src_x +1 < src_cols) ? tmp.y : val;
tmp.z = (src_x+2 >=0)&&(src_x +2 < src_cols) ? tmp.z : val;
tmp.w = (src_x+3 >=0)&&(src_x +3 < src_cols) ? tmp.w : val;
*(__global uchar4*)(dst+dst_addr) = tmp;
}
else if((x<dst_cols) && (y<dst_rows))
{
*(__global uchar4*)(dst+dst_addr) = (uchar4)val;
}
#else
int4 s_x;
int s_y;
//judge if read out of boundary
s_x.x= ADDR_L(src_x,0,src_cols,src_x);
s_x.y= ADDR_L(src_x+1,0,src_cols,src_x+1);
s_x.z= ADDR_L(src_x+2,0,src_cols,src_x+2);
s_x.w= ADDR_L(src_x+3,0,src_cols,src_x+3);
s_x.x= ADDR_R(src_x,src_cols,s_x.x);
s_x.y= ADDR_R(src_x+1,src_cols,s_x.y);
s_x.z= ADDR_R(src_x+2,src_cols,s_x.z);
s_x.w= ADDR_R(src_x+3,src_cols,s_x.w);
s_y= ADDR_L(src_y,0,src_rows,src_y);
s_y= ADDR_R(src_y,src_rows,s_y);
int4 src_addr4=mad24((int4)s_y,(int4)src_step_in_pixel,s_x+(int4)src_offset_in_pixel);
//write the result to dst
if((x<dst_cols) && (y<dst_rows))
{
uchar4 tmp;
tmp.x = src[src_addr4.x];
tmp.y = src[src_addr4.y];
tmp.z = src[src_addr4.z];
tmp.w = src[src_addr4.w];
*(__global uchar4*)(dst+dst_addr) = tmp;
}
#endif
}
}
#undef edge(x,size,rx)

8
modules/ocl/src/kernels/imgproc_histogram.cl
View File

@@ -142,15 +142,17 @@ __kernel void __attribute__((reqd_work_group_size(1,HISTOGRAM256_BIN_COUNT,1)))c
int gy = get_group_id(1);
int gn = get_num_groups(0);
int rowIndex = mad24(gy, gn, gx);
rowIndex &= (PARTIAL_HISTOGRAM256_COUNT - 1);
// rowIndex &= (PARTIAL_HISTOGRAM256_COUNT - 1);
__local int subhist[HISTOGRAM256_BIN_COUNT + 1];
__local int subhist[HISTOGRAM256_LOCAL_MEM_SIZE + 1];
subhist[lidy] = 0;
barrier(CLK_LOCAL_MEM_FENCE);
gidx = ((gidx>left_col) ? (gidx+cols) : gidx);
gidx = ((gidx>=left_col) ? (gidx+cols) : gidx);
int src_index = src_offset + mad24(gidy, src_step, gidx);
barrier(CLK_LOCAL_MEM_FENCE);
int p = (int)src[src_index];
p = gidy >= rows ? HISTOGRAM256_LOCAL_MEM_SIZE : p;
atomic_inc(subhist + p);
barrier(CLK_LOCAL_MEM_FENCE);

4
modules/ocl/src/kernels/imgproc_integral_sum.cl
View File

@@ -56,7 +56,7 @@
#define GET_CONFLICT_OFFSET(lid) ((lid) >> LOG_NUM_BANKS)
kernel void integral_cols(__global uchar4 *src,__global int *sum ,
kernel void integral_sum_cols(__global uchar4 *src,__global int *sum ,
int src_offset,int pre_invalid,int rows,int cols,int src_step,int dst_step)
{
unsigned int lid = get_local_id(0);
@@ -136,7 +136,7 @@ kernel void integral_cols(__global uchar4 *src,__global int *sum ,
}
kernel void integral_rows(__global int4 *srcsum,__global int *sum ,
kernel void integral_sum_rows(__global int4 *srcsum,__global int *sum ,
int rows,int cols,int src_step,int sum_step,
int sum_offset)
{

18
modules/ocl/src/kernels/imgproc_resize.cl
View File

@@ -138,17 +138,14 @@ __kernel void resizeLN_C1_D0(__global uchar * dst, __global uchar const * restri
val2 = mul24(U1 , sdata3) + mul24(U , sdata4);
val = mul24((int4)V1 , val1) + mul24((int4)V , val2);
//__global uchar4* d = (__global uchar4*)(dst + dstoffset_in_pixel + dy * dststep_in_pixel + gx);
//uchar4 dVal = *d;
//int4 con = ( DX >= 0 && DX < dst_cols && dy >= 0 && dy < dst_rows);
val = ((val + (1<<(CAST_BITS-1))) >> CAST_BITS);
//*d = convert_uchar4(con != 0) ? convert_uchar4_sat(val) : dVal;
pos4 = mad24(dy, dststep_in_pixel, gx+dstoffset_in_pixel);
pos4.y++;
pos4.z+=2;
pos4.w+=3;
uchar4 uval = convert_uchar4_sat(val);
int con = (gx >= 0 && gx+3 < dst_cols && dy >= 0 && dy < dst_rows);
int con = (gx >= 0 && gx+3 < dst_cols && dy >= 0 && dy < dst_rows && (dstoffset_in_pixel&3)==0);
if(con)
{
*(__global uchar4*)(dst + pos4.x)=uval;
@@ -167,6 +164,10 @@ __kernel void resizeLN_C1_D0(__global uchar * dst, __global uchar const * restri
{
dst[pos4.z]=uval.z;
}
if(gx+3 >= 0 && gx+3 < dst_cols && dy >= 0 && dy < dst_rows)
{
dst[pos4.w]=uval.w;
}
}
}
@@ -325,8 +326,9 @@ __kernel void resizeNN_C1_D0(__global uchar * dst, __global uchar * src,
pos = mad24(dy, dststep_in_pixel, gx+dstoffset_in_pixel);
pos.y++;
pos.z+=2;
pos.w+=3;
int con = (gx >= 0 && gx+3 < dst_cols && dy >= 0 && dy < dst_rows);
int con = (gx >= 0 && gx+3 < dst_cols && dy >= 0 && dy < dst_rows && (dstoffset_in_pixel&3)==0);
if(con)
{
*(__global uchar4*)(dst + pos.x)=val;
@@ -345,6 +347,10 @@ __kernel void resizeNN_C1_D0(__global uchar * dst, __global uchar * src,
{
dst[pos.z]=val.z;
}
if(gx+3 >= 0 && gx+3 < dst_cols && dy >= 0 && dy < dst_rows)
{
dst[pos.w]=val.w;
}
}
}

873
modules/ocl/src/kernels/imgproc_warpPerspective.cl
View File

File diff suppressed because it is too large Load Diff