利用OpenCL對OpenCV並行化心得（1）

做了一年多的opencv並行化，總算小有成就了

opencv裡面最重要的一種型別恐怕非8uc1莫屬，灰度圖一般都用這種型別表示，而且大多數演算法只提供這種資料型別的支援，所以8uc1的優化就成了整個工作的重中之重。以matrix add為例，最簡單的kernel的寫法是這樣的

__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 x=get_global_id(0);

int y=get_global_id(1);

if(xdst[mad24(y,dst_step,x)]=src1[mad24(y,src1t_step,x)]+src2[mad24(y,src2_step,x)];

這種寫法非常簡潔，但是不夠高效，沒有充分利用gpu的記憶體頻寬，存在很大的浪費，為了充分利用gpu的記憶體頻寬，至少每次要寫出4byte，所以kernel就變成這樣

__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 x=get_global_id(0)<<2;

int y=get_global_id(1);

if(x*(__global ucha4*)(dst+mad24(y,dst_step,x))=*(__global ucha4*)(src1+mad24(y,src1t_step,x))+*(__global ucha4*)(src2+mad24(y,src2_step,x));

這樣效率會提高四倍左右，但是可以這樣寫的前提是沒有roi問題，如果有roi情況會複雜很多