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 @@ -151,36 +151,58 @@ return lr; auto upper = HDTTMap->upper_bound(hp); - // check the left bin - if (upper != HDTTMap->begin()) { - lr.Entry = std::prev(upper)->HDTT; - auto &HT = *lr.Entry; - // Is it contained? - lr.Flags.IsContained = hp >= HT.HstPtrBegin && hp < HT.HstPtrEnd && - (hp + Size) <= HT.HstPtrEnd; - // Does it extend beyond the mapped region? - lr.Flags.ExtendsAfter = hp < HT.HstPtrEnd && (hp + Size) > HT.HstPtrEnd; - } - // check the right bin - if (!(lr.Flags.IsContained || lr.Flags.ExtendsAfter) && - upper != HDTTMap->end()) { - lr.Entry = upper->HDTT; - auto &HT = *lr.Entry; - // Does it extend into an already mapped region? - lr.Flags.ExtendsBefore = - hp < HT.HstPtrBegin && (hp + Size) > HT.HstPtrBegin; - // Does it extend beyond the mapped region? - lr.Flags.ExtendsAfter = hp < HT.HstPtrEnd && (hp + Size) > HT.HstPtrEnd; - } + if (Size == 0) { + // specification v5.1 Pointer Initialization for Device Data Environments + // upper_bound satisfies + // std::prev(upper)->HDTT.HstPtrBegin <= hp < upper->HDTT.HstPtrBegin + if (upper != HDTTMap->begin()) { + lr.Entry = std::prev(upper)->HDTT; + auto &HT = *lr.Entry; + // the left side of extended address range is satisified. + // hp >= HT.HstPtrBegin || hp >= HT.HstPtrBase + lr.Flags.IsContained = hp < HT.HstPtrEnd || hp < HT.HstPtrBase; + } - if (lr.Flags.ExtendsBefore) { - DP("WARNING: Pointer is not mapped but section extends into already " - "mapped data\n"); - } - if (lr.Flags.ExtendsAfter) { - DP("WARNING: Pointer is already mapped but section extends beyond mapped " - "region\n"); + if (!lr.Flags.IsContained && upper != HDTTMap->end()) { + lr.Entry = upper->HDTT; + auto &HT = *lr.Entry; + // the right side of extended address range is satisified. + // hp < HT.HstPtrEnd || hp < HT.HstPtrBase + lr.Flags.IsContained = hp >= HT.HstPtrBase; + } + } else { + // check the left bin + if (upper != HDTTMap->begin()) { + lr.Entry = std::prev(upper)->HDTT; + auto &HT = *lr.Entry; + // Is it contained? + lr.Flags.IsContained = hp >= HT.HstPtrBegin && hp < HT.HstPtrEnd && + (hp + Size) <= HT.HstPtrEnd; + // Does it extend beyond the mapped region? + lr.Flags.ExtendsAfter = hp < HT.HstPtrEnd && (hp + Size) > HT.HstPtrEnd; + } + + // check the right bin + if (!(lr.Flags.IsContained || lr.Flags.ExtendsAfter) && + upper != HDTTMap->end()) { + lr.Entry = upper->HDTT; + auto &HT = *lr.Entry; + // Does it extend into an already mapped region? + lr.Flags.ExtendsBefore = + hp < HT.HstPtrBegin && (hp + Size) > HT.HstPtrBegin; + // Does it extend beyond the mapped region? + lr.Flags.ExtendsAfter = hp < HT.HstPtrEnd && (hp + Size) > HT.HstPtrEnd; + } + + if (lr.Flags.ExtendsBefore) { + DP("WARNING: Pointer is not mapped but section extends into already " + "mapped data\n"); + } + if (lr.Flags.ExtendsAfter) { + DP("WARNING: Pointer is already mapped but section extends beyond mapped " + "region\n"); + } } return lr; @@ -275,10 +297,10 @@ HstPtrName)) .first->HDTT; INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID, - "Creating new map entry with " - "HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", Size=%ld, " + "Creating new map entry with HstPtrBase= " DPxMOD + ", HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", Size=%ld, " "DynRefCount=%s, HoldRefCount=%s, Name=%s\n", - DPxPTR(HstPtrBegin), DPxPTR(Ptr), Size, + DPxPTR(HstPtrBase), DPxPTR(HstPtrBegin), DPxPTR(Ptr), Size, Entry->dynRefCountToStr().c_str(), Entry->holdRefCountToStr().c_str(), (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown"); TargetPointer = (void *)Ptr; 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,58 @@ +// RUN: %libomptarget-compile-generic -fopenmp-version=51 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +#include +#include + +#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]) +#pragma omp target enter data map(alloc : C [FROM:LENGTH]) + +// A, B and C 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]) +#pragma omp target exit data map(delete : A [FROM:LENGTH], B [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; +} diff --git a/openmp/libomptarget/test/mapping/array_section_use_device_ptr.c b/openmp/libomptarget/test/mapping/array_section_use_device_ptr.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/mapping/array_section_use_device_ptr.c @@ -0,0 +1,35 @@ +// RUN: %libomptarget-compile-generic -fopenmp-version=51 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +#include +#include + +#define N 1024 +#define FROM 64 +#define LENGTH 128 + +int main() { + float *A = (float *)malloc(N * sizeof(float)); + +#pragma omp target enter data map(to : A [FROM:LENGTH]) + + // A, has been mapped starting at index FROM, but inside the use_device_ptr + // clause it is captured by base so the library must look it up using the + // base address. + + float *A_dev = NULL; +#pragma omp target data use_device_ptr(A) + { A_dev = A; } +#pragma omp target exit data map(delete : A [FROM:LENGTH]) + + // CHECK: Success + if (A_dev == NULL || A_dev == A) + fprintf(stderr, "Failure\n"); + else + fprintf(stderr, "Success\n"); + + free(A); + + return 0; +}