What's wrong with this kernel?

[renegr]

struct RGB16
{
  ushort R, G, B;
};

__kernel void scale2D_3u16(__global ushort* restrict in, __global ushort* restrict out)
{
  int xPos = get_global_id(0);
  int yPos = get_global_id(1);
  int p = yPos*get_global_size(0)+xPos;
  
  __global struct RGB16* pRGBIn = (__global struct RGB16*)in;
  __global struct RGB16* pRGBOut = (__global struct RGB16*)out;
  pRGBOut
 = pRGBIn
;
}

Input data (memory view) is:

0, 0, 0, 1, 1, 1, 2, 2, 2, 3, 3, 3, 4, 4, 4 ...

Using NVidia and AMD I get the same result as output. But using Intel OCL on HD4600 I get:

0, 0, 0, 0, 0, 1, 2, 2, 2, 0, 0, 3, 4, 4, 4, ...

I'm totally confused!

[renegr]

Additional note, using following code works as expected.

  pRGBOut
.R = pRGBIn
.R;
  pRGBOut
.G = pRGBIn
.G;
  pRGBOut
.B = pRGBIn
.B;

[Robert_I_Intel]

The problem is that the way you wrote you code in the first example, two adjacent work items are reading and writing overlapping data, and there is no guarantee of the order in which the work items will write their data out. In the second case you just got lucky :) Technically, you would need barriers between those instructions in the second case to guarantee the results.

To fix, try to make sure that work items read and write non-overlapping pieces of memory.

[renegr]

Thank you Robert,

unfortunately I don't see any overlapping within this code. Maybe I'm totally blind. It is nearly the same code as I use within a "parallel_for()". I could rewrite the code to:
 

__kernel void scale2D_3u16(__global ushort* restrict in, __global ushort* restrict out)
{
  int xPos = get_global_id(0);
  int yPos = get_global_id(1);
  int p = (yPos*get_global_size(0)+xPos)*3;

  out[p+0] = in[p+0];
  out[p+1] = in[p+1];
  out[p+2] = in[p+2];
}

I think this overlapping occurs because the addresses of pRGBOut/pRGBIn are not multiples of 4/8/16. Is this correct?

[Robert_I_Intel · Accepted solution]

Hi Rene,

My bad: there is no issue with your original code - the problem is with the Intel compiler and how it handles structs, so indeed the workaround is

pRGBOut

.R = pRGBIn

.R;

pRGBOut

.G = pRGBIn

.G;

pRGBOut

.B = pRGBIn

.B;

[Robert_I_Intel]

BTW, your last proposal should work as well.