[clang] [llvm] [clang] [OpenMP] New OpenMP 6.0 self_maps clause - CodeGen (PR #134131)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Apr 16 03:25:20 PDT 2025
================
@@ -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
More information about the cfe-commits
mailing list