[clang-tools-extra] [OpenMP][libomptarget] Add map checks when running under unified shared memory (PR #69005)

via cfe-commits cfe-commits at lists.llvm.org
Fri Oct 13 13:23:52 PDT 2023


================
@@ -444,6 +486,29 @@ DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool UpdateRefCount,
          LR.TPR.getEntry()->dynRefCountToStr().c_str(), DynRefCountAction,
          LR.TPR.getEntry()->holdRefCountToStr().c_str(), HoldRefCountAction);
     LR.TPR.TargetPointer = (void *)TP;
+
+    // If this entry is not marked as being host pointer (the way the
+    // implementation works today this is never true, mistake?) then we
+    // have to check if this is a host pointer or not. This is a host pointer
+    // if the host address matches the target address.
+    if ((PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) &&
+        !LR.TPR.Flags.IsHostPointer) {
----------------
carlobertolli wrote:

This is odd: you set IsHostPointer on the new entry under unified_shared_memory in getTargetPointer.
I am expecting this to be reflected in getTgtPtrBegin, no? Maybe I am misunderstanding how getTgtPtrBegin is really used. I see you are perplexed by this behavior too...?

Can you write a test that reproduces this? Like

```
#pragma opm requires unified_shared_memory

void foo() {
  int a[10];
  #pragma omp target enter data map(to:a[10])

  #pragma omp target  \\ should this call getTgtPtrBegin?
 {
    a[0] = 1;
  }
}
```


https://github.com/llvm/llvm-project/pull/69005


More information about the cfe-commits mailing list