1 Reply Latest reply on Nov 1, 2015 10:19 AM by matszpk

    Images doesn't work with constant const data!

    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.