================
@@ -236,6 +236,8 @@ enum class OpenMPOffloadMappingFlags : uint64_t {
// dynamic.
// This is an OpenMP extension for the sake of OpenACC support.
OMP_MAP_OMPX_HOLD = 0x2000,
+ /// Self directs mapping without creating a separate device copy.
+ OMP_MAP_SELF = 0x4000,
----------------
Ritanya-B-Bharadwaj wrote:
We need `OMP_MAP_SELF` because the runtime still needs to track a self-mapped
variable in the present table, even if no new device copy is created. This is
important when a pointer (like firstprivate) refers to that variable — the
runtime must handle pointer translations and aliasing correctly. Here is an
example:
```
#include <stdio.h>
#include <assert.h>
#define N 1000
int main() {
int aaa[N];
#pragma omp target data map(self: aaa) // map aaa with self map
{
for (int i = 0 ; i < N ; i++) {
// Update host copy of aaa on host
aaa[i] = i + 1;
}
int *p = aaa;
#pragma omp target teams distribute parallel for // p is treated as
firstprivate
for (int i = 0 ; i < N ; i++) {
// Update host copy of aaa on device, since it should be self mapped
p[i] = p[i] - 1; // without self map of aaa, this may not work
}
}
for (int i = 0 ; i < N ; i++) {
assert( aaa[i] == i );
}
printf("PASS\n");
return 0;
}
```
https://github.com/llvm/llvm-project/pull/134131
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
