https://gcc.gnu.org/bugzilla/show_bug.cgi?id=103708
Bug ID: 103708
Summary: [OpenMP][OMP 5.0/5.x] use_device_addr - array sections
not supported
Product: gcc
Version: 12.0
Status: UNCONFIRMED
Keywords: openmp, rejects-valid
Severity: normal
Priority: P3
Component: middle-end
Assignee: unassigned at gcc dot gnu.org
Reporter: burnus at gcc dot gnu.org
CC: jakub at gcc dot gnu.org
Target Milestone: ---
Seemingly missed both when implementing the C/C++ version and when implementing
the Fortran version.
Fails with:
foo.c: In function 'foo':
foo.c:5:44: error: expected ')' before '[' token
5 | #pragma omp target data use_device_addr(a[1:2])
Example (shows partially mapped data):
int
foo ()
{
int a[5];
#pragma omp target enter data map(to:a[1:2])
#pragma omp target data use_device_addr(a[1:2])
{
int *p = &a[1];
// #pragma omp target has_device_addr(a[1:2])
// { /* ... */ }
}
#pragma omp target exit data map(delete:a[1:2])
}
OpenMP specification:
* 5.0: https://www.openmp.org/spec-html/5.0/openmpsu57.html#x83-2580002.12.2
"Array sections privatize the base of the array section and assign the
private copy
to the address of the corresponding array section in the device data
environment."
* 5.1: https://www.openmp.org/spec-html/5.1/openmpsu65.html#x93-1030002.14.2
Same/similar to 5.2
* 5.2: "5.4.10 use_device_addr Clause" [123:23-27]
"The list items that appear in a use_device_addr clause may include array
sections."
C/C++: "If a list item in a use_device_addr clause is an array section that
has a
base pointer, the effect of the clause is to convert the base pointer
to a
pointer that is local to the structured block and that contains the
device
address. This conversion may be elided if the list item was not
already mapped."
NOTES:
* has_device_addr also supports an array section
* Specification: Lacks description for Fortran - especially for
has_device_addr, but
also for use_device_addr. That's tracked as OpenMP Specification Issue 3180.
* has_device_addr: Patch to implement it
https://gcc.gnu.org/pipermail/gcc-patches/2021-November/585361.html
