[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