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 (); + } +