Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

fixed ocl::bitwise_not in case of CV_8UC1, CV_8UC2

Browse Source
pull/1548/head
Ilya Lavrenov 12 years ago
parent 75d89f4331
commit a613a1a1da
1 changed files with 26 additions and 43 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 69
      modules/ocl/src/opencl/arithm_bitwise_not.cl

69
modules/ocl/src/opencl/arithm_bitwise_not.cl
Unescape View File

@ -42,6 +42,7 @@
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#if defined (DOUBLE_SUPPORT)
#ifdef cl_khr_fp64
#pragma OPENCL EXTENSION cl_khr_fp64:enable
@ -49,10 +50,11 @@
#pragma OPENCL EXTENSION cl_amd_fp64:enable
#endif
#endif
//////////////////////////////////////////////////////////////////////////////////////////////////////
////////////////////////////////////////////BITWISE_NOT////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////////////////////////////////
////////////////////////////////////////////BITWISE_NOT////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////////////////////////////////
/**************************************bitwise_not without mask**************************************/
__kernel void arithm_bitwise_not_D0 (__global uchar *src1, int src1_step, int src1_offset,
__global uchar *dst, int dst_step, int dst_offset,
int rows, int cols, int dst_step1)
@ -63,35 +65,22 @@ __kernel void arithm_bitwise_not_D0 (__global uchar *src1, int src1_step, int sr
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
#define dst_align (dst_offset & 3)
int src1_index = mad24(y, src1_step, x + src1_offset - dst_align);
int src1_index = mad24(y, src1_step, x + src1_offset);
int dst_start = mad24(y, dst_step, dst_offset);
int dst_end = mad24(y, dst_step, dst_offset + dst_step1);
int dst_index = mad24(y, dst_step, dst_offset + x & (int)0xfffffffc);
int src1_index_fix = src1_index < 0 ? 0 : src1_index;
uchar4 src1_data = vload4(0, src1 + src1_index_fix);
uchar4 dst_data = *((__global uchar4 *)(dst + dst_index));
uchar4 tmp_data = ~ src1_data;
/* if(src1_index < 0)
{
uchar4 tmp;
tmp.xyzw = (src1_index == -2) ? src1_data.zwxy:src1_data.yzwx;
src1_data.xyzw = (src1_index == -1) ? src1_data.wxyz:tmp.xyzw;
}
*/
dst_data.x = ((dst_index + 0 >= dst_start) && (dst_index + 0 < dst_end)) ? tmp_data.x : dst_data.x;
dst_data.y = ((dst_index + 1 >= dst_start) && (dst_index + 1 < dst_end)) ? tmp_data.y : dst_data.y;
dst_data.z = ((dst_index + 2 >= dst_start) && (dst_index + 2 < dst_end)) ? tmp_data.z : dst_data.z;
dst_data.w = ((dst_index + 3 >= dst_start) && (dst_index + 3 < dst_end)) ? tmp_data.w : dst_data.w;
int dst_index = mad24(y, dst_step, dst_offset + x);
*((__global uchar4 *)(dst + dst_index)) = dst_data;
uchar4 src1_data = vload4(0, src1 + src1_index);
uchar4 dst_data = vload4(0, dst + dst_index);
uchar4 tmp_data = ~src1_data;
dst_data.x = dst_index + 0 < dst_end ? tmp_data.x : dst_data.x;
dst_data.y = dst_index + 1 < dst_end ? tmp_data.y : dst_data.y;
dst_data.z = dst_index + 2 < dst_end ? tmp_data.z : dst_data.z;
dst_data.w = dst_index + 3 < dst_end ? tmp_data.w : dst_data.w;
vstore4(dst_data, 0, dst + dst_index);
}
}
@ -106,28 +95,22 @@ __kernel void arithm_bitwise_not_D1 (__global char *src1, int src1_step, int src
if (x < cols && y < rows)
{
x = x << 2;
#ifdef dst_align
#undef dst_align
#endif
#define dst_align (dst_offset & 3)
int src1_index = mad24(y, src1_step, x + src1_offset - dst_align);
int src1_index = mad24(y, src1_step, x + src1_offset);
int dst_start = mad24(y, dst_step, dst_offset);
int dst_end = mad24(y, dst_step, dst_offset + dst_step1);
int dst_index = mad24(y, dst_step, dst_offset + x & (int)0xfffffffc);
int dst_index = mad24(y, dst_step, dst_offset + x);
char4 src1_data = vload4(0, src1 + src1_index);
char4 dst_data = vload4(0, dst + dst_index);
char4 tmp_data = ~src1_data;
char4 dst_data = *((__global char4 *)(dst + dst_index));
char4 tmp_data = ~ src1_data;
dst_data.x = ((dst_index + 0 >= dst_start) && (dst_index + 0 < dst_end)) ? tmp_data.x : dst_data.x;
dst_data.y = ((dst_index + 1 >= dst_start) && (dst_index + 1 < dst_end)) ? tmp_data.y : dst_data.y;
dst_data.z = ((dst_index + 2 >= dst_start) && (dst_index + 2 < dst_end)) ? tmp_data.z : dst_data.z;
dst_data.w = ((dst_index + 3 >= dst_start) && (dst_index + 3 < dst_end)) ? tmp_data.w : dst_data.w;
dst_data.x = dst_index + 0 < dst_end ? tmp_data.x : dst_data.x;
dst_data.y = dst_index + 1 < dst_end ? tmp_data.y : dst_data.y;
dst_data.z = dst_index + 2 < dst_end ? tmp_data.z : dst_data.z;
dst_data.w = dst_index + 3 < dst_end ? tmp_data.w : dst_data.w;
*((__global char4 *)(dst + dst_index)) = dst_data;
vstore4(dst_data, 0, dst + dst_index);
}
}

Write Preview
Loading…
Cancel
Save