diff options
| author | Julian Brown <julian@codesourcery.com> | 2025-04-20 00:00:54 +0000 |
|---|---|---|
| committer | Sandra Loosemore <sloosemore@baylibre.com> | 2025-05-15 20:25:48 +0000 |
| commit | 87548cbb46416486888daad14c606dffa8e58204 (patch) | |
| tree | 8d3af0c26d366243ac9f670e94fab6a30be67449 /gcc/c/c-parser.cc | |
| parent | 693be70f589682c0c65ce92c4602d481aa7f16e3 (diff) | |
| download | gcc-87548cbb46416486888daad14c606dffa8e58204.zip gcc-87548cbb46416486888daad14c606dffa8e58204.tar.gz gcc-87548cbb46416486888daad14c606dffa8e58204.tar.bz2 | |
OpenACC: Improve implicit mapping for non-lexically nested offload regions
This patch enables use of the OMP_CLAUSE_RUNTIME_IMPLICIT_P flag for
OpenACC.
This allows code like this to work correctly:
int arr[100];
[...]
#pragma acc enter data copyin(arr[20:10])
/* No explicit mapping of 'arr' here. */
#pragma acc parallel
{ /* use of arr[20:10]... */ }
#pragma acc exit data copyout(arr[20:10])
Otherwise, the implicit "copy" ("present_or_copy") on the parallel
corresponds to the whole array, and that fails at runtime when the
subarray is mapped.
The numbering of the GOMP_MAP_IMPLICIT bit clashes with the OpenACC
"non-contiguous" dynamic array support, so the GOMP_MAP_NONCONTIG_ARRAY_P
macro has been adjusted to account for that.
gcc/
* gimplify.cc (gimplify_adjust_omp_clauses_1): Set
OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P for OpenACC also.
gcc/testsuite/
* c-c++-common/goacc/combined-reduction.c: Adjust scan output.
* c-c++-common/goacc/implied-copy-1.c: Likewise.
* c-c++-common/goacc/reduction-1.c: Adjust patterns.
* c-c++-common/goacc/reduction-2.c: Likewise.
* c-c++-common/goacc/reduction-3.c: Likewise.
* c-c++-common/goacc/reduction-4.c: Likewise.
* c-c++-common/goacc/reduction-10.c: Likewise.
* gfortran.dg/goacc/common-block-3.f90: Likewise.
* gfortran.dg/goacc/implied-copy-1.f90: Likewise.
* gfortran.dg/goacc/loop-tree-1.f90: Likewise.
* gfortran.dg/goacc/private-explicit-kernels-1.f95: Likewise.
* gfortran.dg/goacc/private-predetermined-kernels-1.f95: Likewise.
include/
* gomp-constants.h (GOMP_MAP_NONCONTIG_ARRAY_P): Tweak condition.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c: New test.
Co-Authored-By: Thomas Schwinge <tschwinge@baylibre.com>
Co-Authored-By: Sandra Loosemore <sloosemore@baylibre.com>
Diffstat (limited to 'gcc/c/c-parser.cc')
0 files changed, 0 insertions, 0 deletions