利用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     {
        x = x
        #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_data.y = ((dst_index + 1 >= dst_start) && (dst_index + 1         dst_data.z = ((dst_index + 2 >= dst_start) && (dst_index + 2         dst_data.w = ((dst_index + 3 >= dst_start) && (dst_index + 3
        *((__global uchar4 *)(dst + dst_index)) = dst_data;
    }
}

由于为了处理对齐多出来很多额外的操作，所以速度并不是理论上的4倍，而是2倍多一点