Hi Thomas,
regarding your TODO in your test case about implicit mapping of variables, I did
some testing.
The 'copy' issue is a general feature and not restricted to no_create.
Additionally, 'int *arr' is not a real array: as the compiler does not know the
size, it cannot distinguish a pointer to a scalar integer from a pointer to an
integer array. – OpenACC and OpenMP map 'int *arr' slightly differently.
* * *
Looking at the spec (thanks Frederik for the help), I read it such that
* For OpenACC 2.6+2.7 in both kernels (2.5.2) and parallel constructs (2.5.1)
[for both, see last paragraph of 'Description']:
– with 'default(none)': nothing is done explicitly.
– with 'default(present)': then scalars = 'copy', arrays/combined types
'present'
- otherwise: arrays/combined types = 'copy' and
parallel: scalars = 'firstprivate'
kernels: scalars = 'copy'
(Per definition, Fortran's allocatable, pointer + character are never a
'scalar'.)
* For OpenMP, implicit mappings is handled similar to 'parallel':
– scalars = firstprivate (unless: 'defaultmap(tofrom:scalars)')
– nonscalars = 'map(tofrom:'
(OpenMP 5 permits more 'defaultmap's and Fortran allocatable/pointer scalars
are
then also 'map(tofrom:' by default; note Fortran's 'character' is not a
'scalar'
per OpenMP terminology.)
For 'int *arr', one has a pointer which can point to a single or multiple
("array") integer – the in C/C++ compiler cannot know, contrary to 'int
arr2[4]'.
Assume now: 'int var, *arr, arr2[4]' (all -fdump-tree-omplower). Result:
(A) OpenACC
oacc_parallel map(tofrom:arr2 [len: 16]) firstprivate(arr) firstprivate(var)
oacc_kernels map(tofrom:arr2 [len: 16]) map(force_tofrom:arr [len: 8])
map(force_tofrom:var [len: 4])
(B) OpenMP
omp target map(tofrom:arr2 [len: 16]) \
map(alloc:MEM[(char *)arr] [len: 0]) map(firstprivate:arr [pointer
assign, bias: 0]) \
firstprivate(var)
Which looks fine – despite the difference between OpenMP and OpenACC.
(OpenACC: Using default(present) also works – giving 'map(force_present:arr2';
as does
default(none) – causing the compiler to complain about unmapped variables.)
* * *
When enclosing this in 'acc data' (or 'omp data target'), the following of
OpenACC applies:
'implicitly determine data attributes for variables that are referenced
in the compute construct that […] do not appear in a data clause on […]
a lexically containing data construct […]".
Testing shows that independent of the used clause, 'copy()' is always done, also
for scalars in 'parallel'.
For OpenMP 4.5, 2.15.5 is a bit unclear whether 'omp data target's map() apply
or not,
but GCC currently ignores them completely and does the normal 'map(tofrom:' +
'firstprivate'
mapping in this case.
Tobias
PS: Your example was:
On 12/3/19 4:16 PM, Thomas Schwinge wrote:
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+#pragma acc data no_create(var, arr[0:N])
+ {
+ devptr[0] = (int *) acc_deviceptr (&var);
+ devptr[1] = (int *) acc_deviceptr (&arr[2]);
+
+#if ACC_MEM_SHARED
+ if (devptr[0] == NULL)
+ __builtin_abort ();
+ if (devptr[1] == NULL)
+ __builtin_abort ();
+#else
+ if (devptr[0] != NULL)
+ __builtin_abort ();
+ if (devptr[1] != NULL)
+ __builtin_abort ();
+#endif
+
+#pragma acc parallel copyout(devptr) // TODO implicit 'copy(var)' -- huh?!
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (devptr[0] != &var)
+ __builtin_abort (); // { dg-xfail-run-if "TODO" { *-*-* } {
"-DACC_MEM_SHARED=0" } }
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+ }
+
