76 short8 gx = (short8)0;
79 short8 gy = (short8)0;
83 uchar16 temp = vload16(0,
offset(&src, -1, -1));
84 short8 left = convert_short8(temp.s01234567);
85 short8 middle = convert_short8(temp.s12345678);
86 short8 right = convert_short8(temp.s23456789);
88 gx += left * (short8)(-1);
89 gx += right * (short8)(+1);
92 gy += left * (short8)(-1);
93 gy += middle * (short8)(-2);
94 gy += right * (short8)(-1);
98 temp = vload16(0,
offset(&src, -1, 0));
99 left = convert_short8(temp.s01234567);
100 right = convert_short8(temp.s23456789);
102 gx += left * (short8)(-2);
103 gx += right * (short8)(+2);
107 temp = vload16(0,
offset(&src, -1, 1));
108 left = convert_short8(temp.s01234567);
109 middle = convert_short8(temp.s12345678);
110 right = convert_short8(temp.s23456789);
112 gx += left * (short8)(-1);
113 gx += right * (short8)(+1);
116 gy += left * (short8)(+1);
117 gy += middle * (short8)(+2);
118 gy += right * (short8)(+1);
123 vstore8(gx, 0, ((__global
short *)dst_gx.
ptr));
126 vstore8(gy, 0, ((__global
short *)dst_gy.
ptr));
156 const short left1_coeff_gx,
157 const short left2_coeff_gx,
158 const short middle_coeff_gx,
159 const short right1_coeff_gx,
160 const short right2_coeff_gx,
161 const short left1_coeff_gy,
162 const short left2_coeff_gy,
163 const short middle_coeff_gy,
164 const short right1_coeff_gy,
165 const short right2_coeff_gy)
167 uchar16 temp = vload16(0,
offset(src, -2, 0));
172 val = convert_short8(temp.s01234567);
173 gx += val * (short8)left1_coeff_gx;
174 gy += val * (short8)left1_coeff_gy;
176 val = convert_short8(temp.s12345678);
177 gx += val * (short8)left2_coeff_gx;
178 gy += val * (short8)left2_coeff_gy;
180 val = convert_short8(temp.s23456789);
181 gx += val * (short8)middle_coeff_gx;
182 gy += val * (short8)middle_coeff_gy;
184 val = convert_short8(temp.s3456789a);
185 gx += val * (short8)right1_coeff_gx;
186 gy += val * (short8)right1_coeff_gy;
188 val = convert_short8(temp.s456789ab);
189 gx += val * (short8)right2_coeff_gx;
190 gy += val * (short8)right2_coeff_gy;
192 return (short16)(gx, gy);
208 const short up1_coeff,
209 const short up2_coeff,
210 const short middle_coeff,
211 const short down1_coeff,
212 const short down2_coeff)
215 short8 out = (short8)0;
217 val = vload8(0, (__global
short *)
offset(src, 0, -2));
218 out += val * (short8)up1_coeff;
220 val = vload8(0, (__global
short *)
offset(src, 0, -1));
221 out += val * (short8)up2_coeff;
223 val = vload8(0, (__global
short *)
offset(src, 0, 0));
224 out += val * (short8)middle_coeff;
226 val = vload8(0, (__global
short *)
offset(src, 0, 1));
227 out += val * (short8)down1_coeff;
229 val = vload8(0, (__global
short *)
offset(src, 0, 2));
230 out += val * (short8)down2_coeff;
232 return (short8)(out);
286 vstore8(gx_gy.s01234567, 0, ((__global
short *)dst_gx.
ptr));
289 vstore8(gx_gy.s89ABCDEF, 0, ((__global
short *)dst_gy.
ptr));
348 vstore8(gx, 0, ((__global
short *)dst_gx.
ptr));
353 vstore8(gy, 0, ((__global
short *)dst_gy.
ptr));
384 #define SOBEL1x1_HOR(src, gx, gy, idx) \ 386 int8 val = convert_int8(vload8(0, offset(src, idx - 3, 0))); \ 387 gx += val * X##idx; \ 388 gy += val * Y##idx; \ 392 #define SOBEL1x1_VERT(src, g, direction, idx) \ 394 int8 val = vload8(0, (__global int *)offset(src, 0, idx - 3)); \ 395 g += val * (int8)direction##idx; \ 399 #define SOBEL1x7(ptr, gx, gy) \ 400 SOBEL1x1_HOR(ptr, gx, gy, 0) \ 401 SOBEL1x1_HOR(ptr, gx, gy, 1) \ 402 SOBEL1x1_HOR(ptr, gx, gy, 2) \ 403 SOBEL1x1_HOR(ptr, gx, gy, 3) \ 404 SOBEL1x1_HOR(ptr, gx, gy, 4) \ 405 SOBEL1x1_HOR(ptr, gx, gy, 5) \ 406 SOBEL1x1_HOR(ptr, gx, gy, 6) 409 #define SOBEL7x1(ptr, g, direction) \ 410 SOBEL1x1_VERT(ptr, g, direction, 0) \ 411 SOBEL1x1_VERT(ptr, g, direction, 1) \ 412 SOBEL1x1_VERT(ptr, g, direction, 2) \ 413 SOBEL1x1_VERT(ptr, g, direction, 3) \ 414 SOBEL1x1_VERT(ptr, g, direction, 4) \ 415 SOBEL1x1_VERT(ptr, g, direction, 5) \ 416 SOBEL1x1_VERT(ptr, g, direction, 6) 468 vstore8(gx, 0, ((__global
int *)dst_gx.
ptr));
471 vstore8(gy, 0, ((__global
int *)dst_gy.
ptr));
530 vstore8(gx, 0, (__global
int *)dst_gx.
ptr);
535 vstore8(gy, 0, (__global
int *)dst_gy.
ptr);
__kernel void sobel_separable7x1(__global uchar *src_x_ptr, uint src_x_stride_x, uint src_x_step_x, uint src_x_stride_y, uint src_x_step_y, uint src_x_offset_first_element_in_bytes, __global uchar *dst_gx_ptr, uint dst_gx_stride_x, uint dst_gx_step_x, uint dst_gx_stride_y, uint dst_gx_step_y, uint dst_gx_offset_first_element_in_bytes, __global uchar *src_y_ptr, uint src_y_stride_x, uint src_y_step_x, uint src_y_stride_y, uint src_y_step_y, uint src_y_offset_first_element_in_bytes, __global uchar *dst_gy_ptr, uint dst_gy_stride_x, uint dst_gy_step_x, uint dst_gy_stride_y, uint dst_gy_step_y, uint dst_gy_offset_first_element_in_bytes, int dummy)
Apply a 7x1 convolution matrix to two single channel S16 input temporary images and output two single...
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
#define CONVERT_TO_IMAGE_STRUCT(name)
__kernel void sobel_separable1x5(__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global uchar *dst_gx_ptr, uint dst_gx_stride_x, uint dst_gx_step_x, uint dst_gx_stride_y, uint dst_gx_step_y, uint dst_gx_offset_first_element_in_bytes, __global uchar *dst_gy_ptr, uint dst_gy_stride_x, uint dst_gy_step_x, uint dst_gy_stride_y, uint dst_gy_step_y, uint dst_gy_offset_first_element_in_bytes)
Apply a 1x5 sobel matrix to a single channel U8 input image and output two temporary channel S16 imag...
#define IMAGE_DECLARATION(name)
SimpleTensor< float > src
__kernel void sobel_separable1x7(__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global uchar *dst_gx_ptr, uint dst_gx_stride_x, uint dst_gx_step_x, uint dst_gx_stride_y, uint dst_gx_step_y, uint dst_gx_offset_first_element_in_bytes, __global uchar *dst_gy_ptr, uint dst_gy_stride_x, uint dst_gy_step_x, uint dst_gy_stride_y, uint dst_gy_step_y, uint dst_gy_offset_first_element_in_bytes)
Apply a 1x7 sobel matrix to a single channel U8 input image and output two temporary channel S16 imag...
#define SOBEL7x1(ptr, g, direction)
#define SOBEL1x7(ptr, gx, gy)
Structure to hold Image information.
__global uchar * ptr
Pointer to the starting postion of the buffer.
short16 sobel1x5(Image *src, const short left1_coeff_gx, const short left2_coeff_gx, const short middle_coeff_gx, const short right1_coeff_gx, const short right2_coeff_gx, const short left1_coeff_gy, const short left2_coeff_gy, const short middle_coeff_gy, const short right1_coeff_gy, const short right2_coeff_gy)
Compute a 1D horizontal sobel filter 1x5 for 8 bytes assuming the input is made of 1 channel of 1 byt...
short8 sobel5x1(Image *src, const short up1_coeff, const short up2_coeff, const short middle_coeff, const short down1_coeff, const short down2_coeff)
Compute a 1D vertical sobel filter 5x1 for 8 bytes assuming the input is made of 1 channel of 1 byte ...
__kernel void sobel3x3(__global uchar *src_ptr, uint src_stride_x, uint src_step_x, uint src_stride_y, uint src_step_y, uint src_offset_first_element_in_bytes, __global uchar *dst_gx_ptr, uint dst_gx_stride_x, uint dst_gx_step_x, uint dst_gx_stride_y, uint dst_gx_step_y, uint dst_gx_offset_first_element_in_bytes, __global uchar *dst_gy_ptr, uint dst_gy_stride_x, uint dst_gy_step_x, uint dst_gy_stride_y, uint dst_gy_step_y, uint dst_gy_offset_first_element_in_bytes)
This OpenCL kernel that computes a Sobel3x3 filter.
__kernel void sobel_separable5x1(__global uchar *src_x_ptr, uint src_x_stride_x, uint src_x_step_x, uint src_x_stride_y, uint src_x_step_y, uint src_x_offset_first_element_in_bytes, __global uchar *dst_gx_ptr, uint dst_gx_stride_x, uint dst_gx_step_x, uint dst_gx_stride_y, uint dst_gx_step_y, uint dst_gx_offset_first_element_in_bytes, __global uchar *src_y_ptr, uint src_y_stride_x, uint src_y_step_x, uint src_y_stride_y, uint src_y_step_y, uint src_y_offset_first_element_in_bytes, __global uchar *dst_gy_ptr, uint dst_gy_stride_x, uint dst_gy_step_x, uint dst_gy_stride_y, uint dst_gy_step_y, uint dst_gy_offset_first_element_in_bytes, int dummy)
Apply a 5x1 convolution matrix to two single channel S16 input temporary images and output two single...