利用OpenCL对OpenCV并行化心得（2）

上文说到了没有ROI的情况下怎么优化add，现在看看有roi的情况。 ROI是opencv里面的一个特性，也可以说是图像处理库都有的特性，它的意思是对于整个一幅图像，只处理被ROI框起来的那一块，可以看做是一个mask。如果不注重性能，最简单的方法就是加上一个offse

上文说到了没有ROI的情况下怎么优化add，现在看看有roi的情况。

ROI是opencv里面的一个特性，也可以说是图像处理库都有的特性，它的意思是对于整个一幅图像，只处理被ROI框起来的那一块，可以看做是一个mask。如果不注重性能，最简单的方法就是加上一个offset即可

__kernel void matrix_add(__global uchar* src1,__global uchar* src2, __global uchar* dst, int rows, int cols,int src1_step,int src2_step,int dst_step，int src1_offset, int src2_offset, int dst_offset)

{

int x=get_global_id(0);

int y=get_global_id(1);

if(x

dst[mad24(y,dst_step,x+dst_offset)]=src1[mad24(y,src1t_step,x+src1_offset)]+src2[mad24(y,src2_step,x+src2_offset)];

}

但是在每次读4个点的时候不能这么做，因为这有对齐问题。比如一个矩阵是17列1行，ROI设置的是后16个点，起始地址就不会是4的整数倍，这样在指针强制转换的时候会出现未定义的情况，结果是错的。除此之外，长度不是4的倍数也会出现问题，所以要使用一些技巧。一个办法是可以先强制对齐地址，多读一些数据，判断是否是需要的，不是需要的就扔掉

__kernel void matrix_add (__global uchar *src1, int src1_step, int src1_offset,
__global uchar *src2, int src2_step, int src2_offset,
__global uchar *dst, int dst_step, int dst_offset,
int rows, int cols, int dst_step1)
{
int x = get_global_id(0);
int y = get_global_id(1);

if (x < cols && y < rows)
{
x = x << 2;

#define dst_align (dst_offset & 3)
int src1_index = mad24(y, src1_step, x + src1_offset - dst_align);
int src2_index = mad24(y, src2_step, x + src2_offset - dst_align);

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);

uchar4 src1_data = vload4(0, src1 + src1_index);
uchar4 src2_data = vload4(0, src2 + src2_index);

uchar4 dst_data = *((__global uchar4 *)(dst + dst_index));
short4 tmp = convert_short4_sat(src1_data) + convert_short4_sat(src2_data);
uchar4 tmp_data = convert_uchar4_sat(tmp);

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;

*((__global uchar4 *)(dst + dst_index)) = dst_data;
}
}