利用OpenCL对OpenCV并行化心得(2)

上文说到了没有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<cols&&y<rows)

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

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

阅读更多
想对作者说点什么? 我来说一句

没有更多推荐了,返回首页

加入CSDN,享受更精准的内容推荐,与500万程序员共同成长!
关闭
关闭