Hi Omkar,
The Ouput buffer is NULL(all zeros) from kernel in GPU seems some issue in using __local wherein data from local mem is not copied to output buffer or vise versa.
this kernel is part of a big project .guess may not be feasable to send the complete code and setup.
I have also used the following code before barrier() compared to previous in kernel but output is still NULL ...is there any atomic issues like calculate,read,write and so on.I am cluless why it is not working in GPU.
local_temp[ly][lx * 3 + 0] = output[(gy * 1280 * 3 * WG_HEIGHT_SIZE ) + (gx * 3 * WG_WIDTH_SIZE ) + ((ly * 3 * 1280) + lx * 3 + 0) ];
local_temp[ly][lx * 3 + 1] = output[(gy * 1280 * 3 * WG_HEIGHT_SIZE ) + (gx * 3 * WG_WIDTH_SIZE ) + ((ly * 3 * 1280) + lx * 3 + 1) ];
local_temp[ly][lx * 3 + 2] = output[(gy * 1280 * 3 * WG_HEIGHT_SIZE ) + (gx * 3 * WG_WIDTH_SIZE ) + ((ly * 3 * 1280) + lx * 3 + 2) ];
Actaully I am trying to modify the below kernel in terms of __local usage and for understanding the concept so that i can apply __local mem for performance in other part of program.
**********************************************
__kernel void kernel(__global unsigned char * output,
const int rc,const int gc,const int bc ,
const int ravg,const int gavg,const int bavg,
const int ra,const int ga,const int ba)
{
int B[3];
uint j;
uint tx = get_global_id(0) * 3;
B[0] = ((output[tx] - rc) * ravg) + ra;
B[1] = ((output[tx+1] - gc) * gavg) + ga;
B[2] = ((output[tx+2] - bc) * bavg) + ba;
for(j=0;j<3;j++)
{
if(B < 0) B = 0;
if(B > 255) B = 255;
output[tx+j] = B;
}
}
***********************************
Thanks
Pavan
