[Openmp-commits] [openmp] c1a6fe1 - [libomptarget] Implement pointer lookup as 5.1 spec.
Ye Luo via Openmp-commits
openmp-commits at lists.llvm.org
Thu Apr 7 21:01:38 PDT 2022
Author: Ye Luo
Date: 2022-04-07T23:01:25-05:00
New Revision: c1a6fe196d8c23b6f614d6d94f313fcb23da7d85
URL: https://github.com/llvm/llvm-project/commit/c1a6fe196d8c23b6f614d6d94f313fcb23da7d85
DIFF: https://github.com/llvm/llvm-project/commit/c1a6fe196d8c23b6f614d6d94f313fcb23da7d85.diff
LOG: [libomptarget] Implement pointer lookup as 5.1 spec.
As described in 5.1 spec
2.21.7.2 Pointer Initialization for Device Data Environments
Reviewed By: RaviNarayanaswamy
Differential Revision: https://reviews.llvm.org/D123093
Added:
openmp/libomptarget/test/mapping/array_section_implicit_capture.c
openmp/libomptarget/test/mapping/array_section_use_device_ptr.c
Modified:
openmp/libomptarget/src/device.cpp
Removed:
################################################################################
diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
index 328d04cb1f8cb..ff9da120e69f1 100644
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -151,36 +151,58 @@ LookupResult DeviceTy::lookupMapping(HDTTMapAccessorTy &HDTTMap,
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 @@ TargetPointerResultTy DeviceTy::getTargetPointer(
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
index 0000000000000..04f0f1167ae02
--- /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 <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])
+#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
index 0000000000000..ca63eb9cd98e6
--- /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 <stdio.h>
+#include <stdlib.h>
+
+#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;
+}
More information about the Openmp-commits
mailing list