These are some unfinished notes on pointer mapping kinds used within GCC/libgomp for offloading target code (for OpenMP/OpenACC).
GOMP_MAP_POINTER
Used for mapping the base pointer of array slices. A simple slice in C might use a pair of mappings, such as:
#pragma acc parallel copyin(arr[0:n])
mapping
decl
size (bias)
GOMP_MAP_TO
*arr
<size of array in bytes>
GOMP_MAP_POINTER
arr
0
The GOMP_MAP_POINTER mapping always appears after a GOMP_MAP_{TO,FROM} mapping.
Inside an offload region, the remapped pointer (arr) is used to locate the target data.
Similar to GOMP_MAP_TO_PSET below, it appears that an additional GOMP_MAP_POINTER can be emitted when we have a reference to an array:
mapping
decl
size (bias)
GOMP_MAP_TO
**arr
<size of array in bytes>
GOMP_MAP_POINTER
*arr
0
GOMP_MAP_POINTER
arr
0
The OpenACC pointer-mapping code expecting array slices to use either 2 (no array descriptor) or 3 (with array descriptor) mappings probably deals with the second GOMP_MAP_POINTER separately. I'm not sure what circumstances we get an array reference but no descriptor (assumed-size arrays?).
The size field for GOMP_MAP_POINTER is a bias, i.e. the offset to the "virtual" zero'th element when the pointed-to data is partially mapped.
The tree dump shows GOMP_MAP_POINTER as "alloc ... [pointer assign ...]".
It appears that GOMP_MAP_POINTER was supposed to be an internal-only mapping kind, but OpenACC support code in libgomp nevertheless expects to see concrete instances of the GOMP_MAP_POINTER mapping kind. That might be a bug.
GOMP_MAP_ALWAYS_POINTER
Like GOMP_MAP_POINTER, appears directly after a data block mapping (TO/FROM/etc.). That is, a GOMP_MAP_ALWAYS_POINTER as clause N will update using the mapped address for clause N-1.
The pointer must be present on the target already, and is rewritten to point to the current host pointer's target data mapping (which is not necessarily the same as the old mapping). Does not affect reference counts. Used for "update" directive.
Using with GOMP_MAP_TO_PSET is likely wrong, since we wouldn't want to update the pointer using the address of the descriptor.
The size field is a bias.
GOMP_MAP_TO_PSET
Used for "pointer sets", e.g. Fortran array descriptors and class descriptors. GOMP_MAP_TO_PSET is followed by one (or more!) GOMP_MAP_POINTER mappings.
Some code in oacc-parallel.c assumes that GOMP_MAP_TO_PSET mappings are always formed in groups of 3.
!$acc parallel copyin(arr(1:n))
mapping
decl
size (bias)
GOMP_MAP_TO
*(arr->data)
<size of array in bytes>
GOMP_MAP_TO_PSET
arr
<array descriptor size>
GOMP_MAP_POINTER
arr
0
!$acc parallel copyin(class)
mapping
decl
size (bias)
GOMP_MAP_TO
class->_data
class->_vptr->size
GOMP_MAP_TO_PSET
class
8 (class descriptor size)
GOMP_MAP_ATTACH_DETACH
class->_data
0
GOMP_MAP_POINTER
class
0
If arr is a reference (e.g. function parameter), we may have an additional GOMP_MAP_POINTER mapping for the reference itself:
mapping
decl
size (bias)
GOMP_MAP_TO
*(*arr->data)
<size of array in bytes>
GOMP_MAP_TO_PSET
*arr
<descriptor size>
GOMP_MAP_POINTER
*arr
0
GOMP_MAP_POINTER
arr
0
The size field is the size of the descriptor.
GOMP_MAP_DECLARE_ALLOCATE/GOMP_MAP_DECLARE_DEALLOCATE
These mappings always form a group of 3, although only two of the mappings appear to be used in oacc-mem.c.
mapping
decl
size (bias)
GOMP_MAP_DECLARE_ALLOCATE
?
?
GOMP_MAP_TO_PSET
?
?
GOMP_MAP_POINTER
?
?
Do these mappings only work with Fortran arrays with descriptors?
GOMP_MAP_ALLOC
Appears similar to GOMP_MAP_POINTER in tree dumps, but is not a pointer mapping. Used for "create" clauses, etc.
GOMP_MAP_FIRSTPRIVATE_POINTER/GOMP_MAP_FIRSTPRIVATE_REFERENCE
Used internally to GCC, not in libgomp. Rewritten into open-coded initialisation of an XYZ(?) mapping in omp-low.c.
Unclear on the precise semantic difference between the pointer and reference versions.
It is not an error if the pointer does not correspond to currently-mapped data on the target. We use -- or perhaps abuse -- this feature during some OpenACC lowering in gimplify.c.
In some cases, GOMP_MAP_FIRSTPRIVATE_POINTER appears to substitute for GOMP_MAP_POINTER. Maybe the firstprivatization is good for efficiency?
Initially used instead of GOMP_MAP_POINTER when firstprivatize_array_bases is true (for...?).
Introduced on gcc-gomp-4_1 branch by:
GOMP_MAP_FIRSTPRIVATE_INT
Introduced at the same time as GOMP_MAP_FIRSTPRIVATE_POINTER. Used to type-pun a scalar value (equal or smaller in size to a pointer).
This is a real mapping kind that is passed through to libgomp.
Used for OpenMP "is_device_ptr".
GOMP_MAP_FIRSTPRIVATE
Unlike GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE}, this one is a real mapping used in libgomp. Not sure if this is used by the OpenACC support, or only OpenMP.
GOMP_MAP_ATTACH_DETACH
Used internally to GCC. Rewritten to GOMP_MAP_ATTACH or GOMP_MAP_DETACH in gimplify.c.
In case of "acc update" directives, GOMP_MAP_ATTACH_DETACH may be rewritten to GOMP_MAP_ALWAYS_POINTER. That means the mapping-ordering restriction for GOMP_MAP_ALWAYS_POINTER applies to GOMP_MAP_ATTACH_DETACH also, at least when used in "acc update" directives.
GOMP_MAP_ATTACH/GOMP_MAP_DETACH
Used for attach/detach operations. Clause size field is a bias. (FIXME I'm not sure if this makes sense, nor if it is applied consistently by the gimplifier! There are cases -- e.g. bare attach/detach clauses -- where the OMP_CLAUSE_SIZE seems to be set to the pointer size, which is wrong).