AnsweredAssumed Answered

Images doesn't work with constant const data!

Question asked by matszpk on Nov 1, 2015
Latest reply on Nov 1, 2015 by matszpk

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[x]; // 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.

Outcomes