diff --git a/openmp/libomptarget/include/device.h b/openmp/libomptarget/include/device.h
--- a/openmp/libomptarget/include/device.h
+++ b/openmp/libomptarget/include/device.h
@@ -245,11 +245,12 @@
unsigned IsContained : 1;
unsigned ExtendsBefore : 1;
unsigned ExtendsAfter : 1;
+ unsigned OnlyBaseFound : 1;
} Flags;
HostDataToTargetListTy::iterator Entry;
- LookupResult() : Flags({0, 0, 0}), Entry() {}
+ LookupResult() : Flags({0, 0, 0, 0}), Entry() {}
};
/// This struct will be returned by \p DeviceTy::getTargetPointer which provides
@@ -333,8 +334,9 @@
bool HasCloseModifier, bool HasPresentModifier,
bool HasHoldModifier, AsyncInfoTy &AsyncInfo);
void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
- TargetPointerResultTy getTgtPtrBegin(void *HstPtrBegin, int64_t Size,
- bool &IsLast, bool UpdateRefCount,
+ TargetPointerResultTy getTgtPtrBegin(void *HstPtrBegin, void *HstPtrBase,
+ int64_t Size, bool &IsLast,
+ bool UpdateRefCount,
bool UseHoldRefCount, bool &IsHostPtr,
bool MustContain = false,
bool ForceDelete = false);
diff --git a/openmp/libomptarget/src/api.cpp b/openmp/libomptarget/src/api.cpp
--- a/openmp/libomptarget/src/api.cpp
+++ b/openmp/libomptarget/src/api.cpp
@@ -108,8 +108,8 @@
bool IsLast; // not used
bool IsHostPtr;
TargetPointerResultTy TPR =
- Device.getTgtPtrBegin(const_cast<void *>(ptr), 0, IsLast,
- /*UpdateRefCount=*/false,
+ Device.getTgtPtrBegin(const_cast<void *>(ptr), const_cast<void *>(ptr), 0,
+ IsLast, /*UpdateRefCount=*/false,
/*UseHoldRefCount=*/false, IsHostPtr);
int rc = (TPR.TargetPointer != NULL);
// Under unified memory the host pointer can be returned by the
diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -169,6 +169,11 @@
hp < HT.HstPtrBegin && (hp + Size) > HT.HstPtrBegin;
// Does it extend beyond the mapped region?
lr.Flags.ExtendsAfter = hp < HT.HstPtrEnd && (hp + Size) > HT.HstPtrEnd;
+ // Did we just find the base? If looking up by base address, the entry can
+ // only be found in the right bin.
+ if (!lr.Flags.ExtendsBefore && !lr.Flags.ExtendsAfter &&
+ hp == HT.HstPtrBase)
+ lr.Flags.OnlyBaseFound = true;
}
if (lr.Flags.ExtendsBefore) {
@@ -204,7 +209,8 @@
// lead to the IsContained flag to be true - then we must ensure that the
// device address is returned even under unified memory conditions.
if (LR.Flags.IsContained ||
- ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && IsImplicit)) {
+ ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && IsImplicit) ||
+ LR.Flags.OnlyBaseFound) {
auto &HT = *LR.Entry;
const char *RefCountAction;
assert(HT.getTotalRefCount() > 0 && "expected existing RefCount > 0");
@@ -219,7 +225,14 @@
}
const char *DynRefCountAction = HasHoldModifier ? "" : RefCountAction;
const char *HoldRefCountAction = HasHoldModifier ? RefCountAction : "";
- uintptr_t Ptr = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
+
+ uintptr_t Ptr = 0;
+ if (LR.Flags.OnlyBaseFound)
+ // Return the implied device base
+ Ptr = HT.TgtPtrBegin - (HT.HstPtrBegin - HT.HstPtrBase);
+ else
+ // Return the device address corresponding to HstPtrBegin
+ Ptr = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
INFO(OMP_INFOTYPE_MAPPING_EXISTS, DeviceID,
"Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
", Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s, Name=%s\n",
@@ -336,9 +349,10 @@
// Return the target pointer begin (where the data will be moved).
// Decrement the reference counter if called from targetDataEnd.
TargetPointerResultTy
-DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
- bool UpdateRefCount, bool UseHoldRefCount,
- bool &IsHostPtr, bool MustContain, bool ForceDelete) {
+DeviceTy::getTgtPtrBegin(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
+ bool &IsLast, bool UpdateRefCount,
+ bool UseHoldRefCount, bool &IsHostPtr,
+ bool MustContain, bool ForceDelete) {
void *TargetPointer = NULL;
bool IsNew = false;
IsHostPtr = false;
@@ -347,7 +361,8 @@
LookupResult lr = lookupMapping(HstPtrBegin, Size);
if (lr.Flags.IsContained ||
- (!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter))) {
+ (!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter)) ||
+ lr.Flags.OnlyBaseFound) {
auto &HT = *lr.Entry;
// We do not zero the total reference count here. deallocTgtPtr does that
// atomically with removing the mapping. Otherwise, before this thread
@@ -376,7 +391,14 @@
}
const char *DynRefCountAction = UseHoldRefCount ? "" : RefCountAction;
const char *HoldRefCountAction = UseHoldRefCount ? RefCountAction : "";
- uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
+
+ uintptr_t tp = 0;
+ if (lr.Flags.OnlyBaseFound)
+ // Return the implied device base
+ tp = HT.TgtPtrBegin - (HT.HstPtrBegin - HT.HstPtrBase);
+ else
+ // Return the device address corresponding to HstPtrBegin
+ tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
INFO(OMP_INFOTYPE_MAPPING_EXISTS, DeviceID,
"Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
"Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s\n",
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -696,6 +696,7 @@
}
void *HstPtrBegin = Args[I];
+ void *HstPtrBase = ArgBases[I];
int64_t DataSize = ArgSizes[I];
// Adjust for proper alignment if this is a combined entry (for structs).
// Look at the next argument - if that is MEMBER_OF this one, then this one
@@ -724,8 +725,8 @@
// If PTR_AND_OBJ, HstPtrBegin is address of pointee
TargetPointerResultTy TPR = Device.getTgtPtrBegin(
- HstPtrBegin, DataSize, IsLast, UpdateRef, HasHoldModifier, IsHostPtr,
- !IsImplicit, ForceDelete);
+ HstPtrBegin, HstPtrBase, DataSize, IsLast, UpdateRef, HasHoldModifier,
+ IsHostPtr, !IsImplicit, ForceDelete);
void *TgtPtrBegin = TPR.TargetPointer;
if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
DP("Mapping does not exist (%s)\n",
@@ -855,7 +856,7 @@
TIMESCOPE_WITH_IDENT(loc);
bool IsLast, IsHostPtr;
TargetPointerResultTy TPR = Device.getTgtPtrBegin(
- HstPtrBegin, ArgSize, IsLast, /*UpdateRefCount=*/false,
+ HstPtrBegin, ArgsBase, ArgSize, IsLast, /*UpdateRefCount=*/false,
/*UseHoldRefCount=*/false, IsHostPtr, /*MustContain=*/true);
void *TgtPtrBegin = TPR.TargetPointer;
if (!TgtPtrBegin) {
@@ -1298,8 +1299,8 @@
void *TgtPtrBegin = (void *)((uintptr_t)TgtPtrBase + Delta);
void *&PointerTgtPtrBegin = AsyncInfo.getVoidPtrLocation();
TargetPointerResultTy TPR = Device.getTgtPtrBegin(
- HstPtrVal, ArgSizes[I], IsLast, /*UpdateRefCount=*/false,
- /*UseHoldRefCount=*/false, IsHostPtr);
+ HstPtrVal, HstPtrBegin, ArgSizes[I], IsLast,
+ /*UpdateRefCount=*/false, /*UseHoldRefCount=*/false, IsHostPtr);
PointerTgtPtrBegin = TPR.TargetPointer;
if (!PointerTgtPtrBegin) {
DP("No lambda captured variable mapped (" DPxMOD ") - ignored\n",
@@ -1355,7 +1356,7 @@
} else {
if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)
HstPtrBase = *reinterpret_cast<void **>(HstPtrBase);
- TPR = Device.getTgtPtrBegin(HstPtrBegin, ArgSizes[I], IsLast,
+ TPR = Device.getTgtPtrBegin(HstPtrBegin, HstPtrBase, ArgSizes[I], IsLast,
/*UpdateRefCount=*/false,
/*UseHoldRefCount=*/false, IsHostPtr);
TgtPtrBegin = TPR.TargetPointer;
diff --git a/openmp/libomptarget/test/mapping/array_section_implicit_capture.c b/openmp/libomptarget/test/mapping/array_section_implicit_capture.c
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/array_section_implicit_capture.c
@@ -0,0 +1,57 @@
+// RUN: %libomptarget-compile-generic -fopenmp-version=51
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 1024
+#define FROM 64
+#define LENGTH 128
+
+int main()
+{
+ float *A = (float *) malloc(N * sizeof(float));
+ float *B = (float *) malloc(N * sizeof(float));
+ float *C = (float *) malloc(N * sizeof(float));
+
+ for (int i = 0; i < N; i++) {
+ C[i] = 0.0;
+ }
+
+ for (int i = 0; i < N; i++) {
+ A[i] = i;
+ B[i] = 2*i;
+ }
+
+ #pragma omp target enter data map(to: A[FROM:LENGTH], B[FROM:LENGTH], C[FROM:LENGTH])
+
+ // A and B have been mapped starting at index FROM, but inside the kernel they
+ // are captured implicitly so the library must look them up using their base
+ // address.
+ #pragma omp target
+ {
+ for (int i = FROM; i < FROM+LENGTH; i++) {
+ C[i] = A[i] + B[i];
+ }
+ }
+
+ #pragma omp target exit data map(from: C[FROM:LENGTH])
+
+ int errors = 0;
+ for (int i = FROM; i < FROM+LENGTH; i++)
+ if (C[i] != A[i] + B[i])
+ ++errors;
+
+ // CHECK: Success
+ if (errors)
+ fprintf(stderr, "Failure\n");
+ else
+ fprintf(stderr, "Success\n");
+
+ free(A);
+ free(B);
+ free(C);
+
+ return 0;
+}