- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
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!
- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
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;
Link Copied
- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
Additional note, using following code works as expected.
pRGBOut.R = pRGBIn
.R; pRGBOut
.G = pRGBIn
.G; pRGBOut
.B = pRGBIn
.B;
- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
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.
- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
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?
- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
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;
- Mark as New
- Bookmark
- Subscribe
- Mute
- Subscribe to RSS Feed
- Permalink
- Report Inappropriate Content
BTW, your last proposal should work as well.

- Subscribe to RSS Feed
- Mark Topic as New
- Mark Topic as Read
- Float this Topic for Current User
- Bookmark
- Subscribe
- Printer Friendly Page