cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

mux85
Journeyman III

fixed size array not working when used inside loop

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

0 Likes
35 Replies
mux85
Journeyman III

fixed size array not working when used inside loop

i forgot to say that the array the causes the problem is the one called temp in the code i posted

0 Likes
omkaranathan
Journeyman III

fixed size array not working when used inside loop

mux85,

Could you post the runtime side code too? A compilable test case would make it easier for us to tackle the problem.

 

0 Likes
mux85
Journeyman III

fixed size array not working when used inside loop

i'm using an opencl wrapper for c# (cloo). i am using also other external libraries to capture and write the video frames. if you want i can send the whole project but i'll probably need an email address

0 Likes
omkaranathan
Journeyman III

fixed size array not working when used inside loop

You can send it to streamdeveloper@amd.com

0 Likes
ryta1203
Journeyman III

fixed size array not working when used inside loop

I get a Byte Addressable Stores Invalid in SKA.

Can you just vectorize this code? So that instead of operating on one rgb, you are operating on 4 at a time?

0 Likes
mux85
Journeyman III

fixed size array not working when used inside loop

Originally posted by: ryta1203 I get a Byte Addressable Stores Invalid in SKA.

 

Can you just vectorize this code? So that instead of operating on one rgb, you are operating on 4 at a time?

 

probably your device does not support that extension. i get no such error in ska

0 Likes
ryta1203
Journeyman III

fixed size array not working when used inside loop

Originally posted by: mux85
Originally posted by: ryta1203 I get a Byte Addressable Stores Invalid in SKA.

 

Can you just vectorize this code? So that instead of operating on one rgb, you are operating on 4 at a time?

 

probably your device does not support that extension. i get no such error in ska

I'm using a HD5870. You are using a 5850, so I'm not sure what the diff would be!?

0 Likes
mux85
Journeyman III

fixed size array not working when used inside loop

Originally posted by: omkaranathan You can send it to streamdeveloper@amd.com

 

hi, i've sent a mail with the title of this discussion as subject. let me know you have received it or not. thanks

0 Likes
mux85
Journeyman III

fixed size array not working when used inside loop

Originally posted by: ryta1203
Originally posted by: mux85
Originally posted by: ryta1203 I get a Byte Addressable Stores Invalid in SKA.

 

 

 

Can you just vectorize this code? So that instead of operating on one rgb, you are operating on 4 at a time?

 

 

 

 

probably your device does not support that extension. i get no such error in ska

 

 

I'm using a HD5870. You are using a 5850, so I'm not sure what the diff would be!?

 

in this case you shouldn't have any problem with cl_khr_byte_addressable_store as i do not. in fact i don't think it is at the origin of my problem. other kernels with that extension enabled work well

0 Likes