On Fri, Jul 17, 2015 at 15:05:59 +0200, Jakub Jelinek wrote:
> As the testcases show, #pragma omp declare target has now a new form (well,
> two; with some issues on it pending), where it is used just as a single
> declarative directive rather than a pair of them and allows marking
> vars and functions by name as "omp declare target" vars/functions (which the
> middle-end etc. already handles), but also "omp declare target link", which
> is a deferred var, that is not initially mapped (on devices without shared
> memory with host), but has to be mapped explicitly.
I don't quite understand how link should work. OpenMP 4.5 says:
"The list items of a link clause are not mapped by the declare target directive.
Instead, their mapping is deferred until they are mapped by target data or
target constructs. They are mapped only for such regions."
But doesn't this mean that the example bellow should work identically
with/without USE_LINK defined? Or is there some difference on other testcases?
int a = 1;
#ifdef USE_LINK
#pragma omp declare target link(a)
#endif
int main ()
{
a = 2;
int res;
#pragma omp target map(to: a) map(from: res)
res = a;
return res;
}
> This patch only marks them with the new attribute, the actual middle-end
> implementation needs to be implemented.
>
> I believe OpenACC has something similar, but no idea if it is already
> implemented.
>
> Anyway, I think the implementation should be that in some pass running on
> the ACCEL_COMPILER side (guarded by separate address space aka non-HSA)
HSA does not define ACCEL_COMPILER, because it uses only one compiler.
> we actually replace the variables with pointers to variables, then need
> to somehow also mark those in the offloading tables, so that the library
I see 2 possible options: use the MSB of the size, or introduce the third field
for flags.
> registers them (the locations of the pointers to the vars), but also marks
> them for special treatment, and then when actually trying to map them
> (or their parts, guess that needs to be discussed) we allocate them or
> whatever is requested and store the device pointer into the corresponding
> variable.
>
> Ilya, Thomas, thoughts on this?
-- Ilya
