hi, i have a kernel used to compute a background from a video. in this kernel i use an array to store pixels in the same position of different frames. on cpu this kernel works well but on gpu (hd 5850) it doesn't. doing some try simplifing the kernel to find the source of the problem i found out that when i use this array inside a for loop the data aren't written correctly, outside the loop the data are correct (but quite useless). here is the code. thanks in advance for any help
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable #define MAX_FRAMES 64 typedef struct { uchar r; uchar g; uchar b; } rgb; void swap(rgb * a, rgb * b) { rgb t=*a; *a=*b; *b=t; } float rgbToLum(rgb pix) { return 0.3f*pix.r+0.59f*pix.g+0.11f*pix.b; } void sort(rgb * v, uint n) { bool swapped=true; while(swapped) { swapped=false; for(uint i=0; i<n-1; i++) if(rgbToLum(v)>rgbToLum(v[i+1])) { swap(&v,&v[i+1]); swapped=true; } n=n-1; } } rgb median(rgb * v, uint n) { sort(v, n); return v[n/2]; } rgb average(rgb m[2][2]) { rgb a; ushort r=0, g=0, b=0; for(uint i=0; i<2; i++) for(uint j=0; j<2; j++) { r += m.r; g += m.g; b += m.b; } a.r = r/4; a.g = g/4; a.b = b/4; return a; } kernel void BackgroundKernel( global read_only rgb * buf_in, global write_only rgb * buf_bg, read_only uint n) { ushort i = get_global_id(0); ushort j = get_global_id(1); ushort h_bg = get_global_size(0); ushort w_bg = get_global_size(1); ushort h_in = h_bg*2; ushort w_in = w_bg*2; uint frame_size = w_in*h_in; uint pos_in = w_in*i*2+j*2; uint pos_bg = w_bg*i+j; //down-scaling of the frames rgb matr[2][2]; rgb temp[MAX_FRAMES]; for(uint k=0; k<n; k++) { matr[0][0] = buf_in[pos_in]; matr[0][1] = buf_in[pos_in+1]; matr[1][0] = buf_in[pos_in+w_in]; matr[1][1] = buf_in[pos_in+w_in+1]; temp = average(matr); pos_in += frame_size; } //median of the frames used as bg buf_bg[pos_bg] = median(temp, n); }
