diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp index 80cb9381ffbb..0255f94599ee 100644 --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -556,22 +556,35 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num, } if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) { - DP("Update pointer (" DPxMOD ") -> [" DPxMOD "]\n", - DPxPTR(PointerTgtPtrBegin), DPxPTR(TgtPtrBegin)); + // Check whether we need to update the pointer on the device + bool UpdateDevPtr = false; + uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase; - void *&TgtPtrBase = AsyncInfo.getVoidPtrLocation(); - TgtPtrBase = (void *)((uint64_t)TgtPtrBegin - Delta); - int rt = Device.submitData(PointerTgtPtrBegin, &TgtPtrBase, - sizeof(void *), AsyncInfo); - if (rt != OFFLOAD_SUCCESS) { - REPORT("Copying data to device failed.\n"); - return OFFLOAD_FAIL; - } - // create shadow pointers for this entry + void *ExpectedTgtPtrBase = (void *)((uint64_t)TgtPtrBegin - Delta); + Device.ShadowMtx.lock(); - Device.ShadowPtrMap[Pointer_HstPtrBegin] = { - HstPtrBase, PointerTgtPtrBegin, TgtPtrBase}; + auto Entry = Device.ShadowPtrMap.find(Pointer_HstPtrBegin); + // If this pointer is not in the map we need to insert it. + if (Entry == Device.ShadowPtrMap.end()) { + // create shadow pointers for this entry + Device.ShadowPtrMap[Pointer_HstPtrBegin] = { + HstPtrBase, PointerTgtPtrBegin, ExpectedTgtPtrBase}; + UpdateDevPtr = true; + } Device.ShadowMtx.unlock(); + + if (UpdateDevPtr) { + DP("Update pointer (" DPxMOD ") -> [" DPxMOD "]\n", + DPxPTR(PointerTgtPtrBegin), DPxPTR(TgtPtrBegin)); + void *&TgtPtrBase = AsyncInfo.getVoidPtrLocation(); + TgtPtrBase = ExpectedTgtPtrBase; + int rt = Device.submitData(PointerTgtPtrBegin, &TgtPtrBase, + sizeof(void *), AsyncInfo); + if (rt != OFFLOAD_SUCCESS) { + REPORT("Copying data to device failed.\n"); + return OFFLOAD_FAIL; + } + } } } diff --git a/openmp/libomptarget/test/mapping/device_ptr_update.c b/openmp/libomptarget/test/mapping/device_ptr_update.c new file mode 100644 index 000000000000..e3dd12f3d434 --- /dev/null +++ b/openmp/libomptarget/test/mapping/device_ptr_update.c @@ -0,0 +1,44 @@ +// RUN: %libomptarget-compile-generic +// RUN: env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefix=DEBUG -check-prefix=CHECK +// REQUIRES: libomptarget-debug + +#include + +struct S { + int *p; +}; + +int main(void) { + int A[10]; + struct S s1; + + s1.p = A; + + // DEBUG: Update pointer ([[DEV_PTR:0x[^ ]+]]) -> {{\[}}[[DEV_OBJ_A:0x[^ ]+]]{{\]}} + #pragma omp target enter data map(alloc : s1.p [0:10]) + + // DEBUG-NOT: Update pointer ([[DEV_PTR]]) -> {{\[}}[[DEV_OBJ_A]]{{\]}} + #pragma omp target map(alloc : s1.p [0:10]) + { + for (int i = 0; i < 10; ++i) + s1.p[i] = i; + } + + #pragma omp target exit data map(from : s1.p [0:10]) + + int fail_A = 0; + for (int i = 0; i < 10; ++i) { + if (A[i] != i) { + fail_A = 1; + break; + } + } + + // CHECK-NOT: Test A failed + if (fail_A) { + printf("Test A failed\n"); + } + + return fail_A; +}