Forum Discussion

Altera_Forum's avatar
Altera_Forum
Icon for Honored Contributor rankHonored Contributor
10 years ago

restrict Keyword & Pointer Aliasing

Hello, I have read over the Altera Optimization Guide and Programming Guide concerning the restrict keyword:

"Insert the restrict keyword in pointer arguments whenever possible. Including the restrict keyword in pointer arguments prevents the Altera Offline Compiler (AOC) from creating unnecessary memory

dependencies between non-conflicting load and store operations."

Based on that, the restrict keyword tells the compiler that you promise not to alias pointer A and B in the following situation:

__kernel void myKernel (__global int * restrict PtrA, __global int * restrict PtrB)

Here both A and B point to global memory, but the programmer is telling the compiler that PtrA and PtrB will never point to the same address.

My problem is that I have two pointers that are derived from the same kernel input pointer, similar too:

__kernel void myKernel (__local int * restrict ptrA, int Offset)

{

for( int i = 0; i < 4; i++ )

ptrA[i] = ptrA[i+Offset];

}

The value of Offset that I pass will guarantee that the read and write will never overlap, but when I compile the kernel, I get a report message like:

| Memory dependency on Load Operation from: | [1]:4 |

| Store Operation | [1]:4 |

So even though I have used the restrict keyword, aoc is still seeing a dependency between the read and write. How can I tell aoc to ignore the dependency in this case?
No RepliesBe the first to reply