[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