diff options
author | Julian Brown <julian@codesourcery.com> | 2025-04-19 23:41:41 +0000 |
---|---|---|
committer | Sandra Loosemore <sloosemore@baylibre.com> | 2025-05-15 20:25:48 +0000 |
commit | f170e172ab2f559aa765a28af796c1a4897adb00 (patch) | |
tree | daa492336cf1d99ba30624e78720ccbb1d63f1e3 /gcc/c/c-parser.cc | |
parent | 8258d04ad7aaed6b6017da313ca14fcfdc512e7d (diff) | |
download | gcc-f170e172ab2f559aa765a28af796c1a4897adb00.zip gcc-f170e172ab2f559aa765a28af796c1a4897adb00.tar.gz gcc-f170e172ab2f559aa765a28af796c1a4897adb00.tar.bz2 |
OpenACC: Reimplement "inheritance" for lexically-nested offload regions
This patch reimplements "lexical inheritance" for OpenACC offload regions
inside "data" regions, allowing e.g. this to work:
int *ptr;
[...]
#pragma acc data copyin(ptr[10:2])
{
#pragma acc parallel
{ ... }
}
here, the "copyin" is mirrored on the inner "acc parallel" as
"present(ptr[10:2])" -- allowing code within the parallel to use that
section of the array even though the mapping is implicit.
In terms of implementation, this works by expanding mapping nodes for
"acc data" to include pointer mappings that might be needed by inner
offload regions. The resulting mapping group is then copied to the inner
offload region as needed, rewriting the first node to "force_present".
The pointer mapping nodes are then removed from the "acc data" later
during gimplification.
For OpenMP, pointer mapping nodes on equivalent "omp data" regions are
not needed, so remain suppressed during expansion.
gcc/c-family/
* c-omp.cc (c_omp_address_inspector::expand_array_base): Don't omit
pointer nodes for OpenACC.
gcc/
* gimplify.cc (omp_tsort_mark, omp_mapping_group): Move before
gimplify_omp_ctx. Add constructor to omp_mapping_group.
(gimplify_omp_ctx): Add DECL_DATA_CLAUSE field.
(new_omp_context, delete_omp_context): Initialise and free above field.
(omp_gather_mapping_groups_1): Use constructor for omp_mapping_group.
(gimplify_scan_omp_clauses): Record mappings that might be lexically
inherited. Don't remove
GOMP_MAP_FIRSTPRIVATE_POINTER/GOMP_MAP_FIRSTPRIVATE_REFERENCE yet.
(gomp_oacc_needs_data_present): New function.
(gimplify_adjust_omp_clauses_1): Implement lexical inheritance
behaviour for OpenACC.
(gimplify_adjust_omp_clauses): Remove
GOMP_MAP_FIRSTPRIVATE_POINTER/GOMP_MAP_FIRSTPRIVATE_REFERENCE here
instead, after lexical inheritance is done.
gcc/testsuite/
* c-c++-common/goacc/acc-data-chain.c: New test.
* gfortran.dg/goacc/pr70828.f90: Likewise.
* gfortran.dg/goacc/assumed-size.f90: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test.
* testsuite/libgomp.oacc-c-c++-common/pr70828-2.c: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-2.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-3.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-4.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-5.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-6.f90: Likewise.
Diffstat (limited to 'gcc/c/c-parser.cc')
0 files changed, 0 insertions, 0 deletions