UNSIGNED_INT8 Image2D access impaired
Hi;
I was creating some Image2D manipulation kernels that worked perfectly to filter a given input Bitmap, loading its contents as byte (Unsigned int8), but for some strange reason the read_imageui or write_imageui got broken.
The code worked with previous Catalyst but doesn't work anymore with 10.7. Code works on NVidia hardware too. Interestingly enough, I even managed to get a picture of the result in http://www.cmsoft.com.br/index.php?option=com_content&view=category&layout=blog&id=115&Itemid=172 ....
I would like some advice, please. Kernel code is attached.
__kernel void Border (__read_only image2d_t img1, __write_only image2d_t img2) { const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | //Natural coordinates CLK_ADDRESS_CLAMP | //Clamp to zeros CLK_FILTER_NEAREST; //Don't interpolate int x0 = get_global_id(0); int y0 = get_global_id(1); int2 coord = (int2)(x0+3, y0+3); uint4 centralValue = (uint)49*read_imageui(img1, smp, coord); uint4 val = (uint4)(0,0,0,0); for (int i = 0; i < 7; i++) { for (int j = 0; j < 7; j++) { coord = (int2)(x0+i, y0+j); val += read_imageui(img1, smp, coord); } } centralValue.x = centralValue.x > val.x ? centralValue.x - val.x : val.x - centralValue.x; centralValue.y = centralValue.y > val.y ? centralValue.y - val.y : val.y - centralValue.y; centralValue.z = centralValue.z > val.z ? centralValue.z - val.z : val.z - centralValue.z; //centralValue.w = centralValue.w > val.w ? centralValue.w - val.w : val.w - centralValue.w; centralValue /= (uint)8; centralValue = (uint)255 - centralValue; if (centralValue.x>255) centralValue.x = 255; if (centralValue.y>255) centralValue.y = 255; if (centralValue.z>255) centralValue.z = 255; //if (centralValue.w>255) centralValue.w = 255; coord.x = x0+3;coord.y = y0+3; write_imageui(img2, coord, centralValue); }
