I wrote simple test kernel to playing with images with my assembler, and I found bug in AMD Catalyst OpenCL.
If constant const data used with images then write_image doesn't work correctly.
This bug caused when constant data will be saved in global data section (no reduction to constant literals).
This simple example:
constant const uint datax[4] = { 344, 555, 11, 333 };
kernel void imageTransform(uint value1, read_only image2d_t inimg,
write_only image2d_t outimg, global uint* output, uint value2)
{
int x = get_global_id(0)&3;
const int2 pos = (int2)(get_global_id(0), get_global_id(1));
if (pos.x >= get_image_width(inimg) || pos.y >= get_image_height(inimg))
return; // don't touch boundaries
write_imageui(outimg, pos, read_imageui(inimg, pos));
output[0] = value1;
output[1] = value2;
output[x+2] = datax; // use global const data
}
This simple example tested on Linux OpenSUSE 13.2 with AMD Catalyst 15.9 and it causes unexpected results on the output image.
