the following code from OpenMP examples 4.0.2(Section 58.3) causes an execution error. I suppose the problem is that the _target_ construct tries to map A instead of using A[0:4].
The testsuit for libgomp contains a similar testcase but it has an additional _map_ clause for A:
#pragma omp target map(p[7:20]) map(A[0:4])
*** Example code
void foo ()
int A, *p;
#pragma omp target data map(A[0:4])
p = &A;
#pragma omp target map(p[7:20])
A = 777;
p = 777;
if (A != 777 || A != 777)
*** Compilation flags
$ x86_64-pc-linux-gnu-gcc -O2 -fopenmp source.c
*** Error message
libgomp: Trying to map into device [0x7fff9538a420..0x7fff9538a498) object when [0x7fff9538a420..0x7fff9538a430) is already mapped
Using built-in specs.
Configured with: ../configure --disable-multilib --enable-languages=c,c++,fortran,lto --build=x86_64-pc-linux-gnu --host=x86_64-pc-linux-gnu --target=x86_64-pc-linux-gnu --enable-offload-targets=x86_64-intelmicemul-linux-gnu --disable-bootstrap
Thread model: posix
gcc version 7.0.0 20160711 (experimental) (GCC)
This is invalid OpenMP.
As you are mentioning A in the target body, but not in its clauses, and as it is not a scalar (new In OpenMP 4.5), it is implicitly map(tofrom:A).
And, as you mapped just a portion of the array earlier, and trying this implicit
map of the whole array which also overlaps the explicit p[7:20] array section, you are violating several restrictions, e.g. OpenMP 4.0:
- List items of map clauses in the same construct must not share original storage.
- If any part of the original storage of a list item has corresponding storage in the enclosing device data environment, all of the original storage must have
corresponding storage in the enclosing device data environment.
The first one that map(p[7:20]) and map(tofrom:A) share original storage, and the second one that A[0:4] is already mapped, and you are mapping whole A.
The fix is simple, add map(A[0:4]) to omp target. It really will not map anything, just bump the reference counts, because it is already mapped, but will avoid the implicit mapping of the whole A.
Jakub, this code is taken verbatim from the openmp-examples document. Should we report it at their github issue tracker? The example in their public git has always been without the A[0:4] map on 'omp target' as quoted by Eugene.
GCC's testcase, on the other hand, had the map clause from day 1. This is a bit strange; assuming GCC didn't change the example code when importing, it suggests that the published openmp-examples text changed the example after GCC imported an earlier version.
I've asked on omp-lang, no responses so far.
I can't find anything in OpenMP 4.0 that would make it valid, in 4.5 there is the mysterious 2.15.5 "When specified, explicit map clauses on target data and target directives determine these attributes." sentence that one could be read either way (but if it allows it, it is IMHO insufficient in saying how it is supposed to work).
The testcase has been changed on the GCC side after lengthy discussion on it back in 2014, the result has been that the testcase is invalid, but there is a desire to make it work somehow in the future. Nothing happened in that area in OpenMP 4.5 though.