[Openmp-commits] [openmp] 434992c - [OpenMP][FIX] Do not overalign mapped structures

[Joseph Huber via Openmp-commits]

Author: Johannes Doerfert
Date: 2023-02-03T07:57:16-06:00
New Revision: 434992c96ed1c41316e029877bdb516fddeddb16

URL: https://github.com/llvm/llvm-project/commit/434992c96ed1c41316e029877bdb516fddeddb16
DIFF: https://github.com/llvm/llvm-project/commit/434992c96ed1c41316e029877bdb516fddeddb16.diff

LOG: [OpenMP][FIX] Do not overalign mapped structures

While we potentially need to align partially mapped structs more than
the first member, we do not need to align past the struct itself. This
prevents us from moving the base pointer past the struct beginning too.

See https://reviews.llvm.org/D142508 for a discussion.

Reviewed By: pavelkopyl, grokos, jhuber6

Differential Revision: https://reviews.llvm.org/D142586

Added: 
    openmp/libomptarget/test/mapping/low_alignment.c

Modified: 
    openmp/libomptarget/src/omptarget.cpp

Removed: 
    


################################################################################
diff  --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index f3a570f89692e..194c414863903 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -75,8 +75,8 @@ int32_t AsyncInfoTy::runPostProcessing() {
 
 bool AsyncInfoTy::isQueueEmpty() const { return AsyncInfo.Queue == nullptr; }
 
-/* All begin addresses for partially mapped structs must be 8-aligned in order
- * to ensure proper alignment of members. E.g.
+/* All begin addresses for partially mapped structs must be aligned, up to 16,
+ * in order to ensure proper alignment of members. E.g.
  *
  * struct S {
  *   int a;   // 4-aligned
@@ -105,7 +105,14 @@ bool AsyncInfoTy::isQueueEmpty() const { return AsyncInfo.Queue == nullptr; }
  * device will start at 0x200 with the padding (4 bytes), then &s1.b=0x204 and
  * &s1.p=0x208, as they should be to satisfy the alignment requirements.
  */
-static const int64_t Alignment = 8;
+static const int64_t MaxAlignment = 16;
+
+/// Return the alignment requirement of partially mapped structs, see
+/// MaxAlignment above.
+static int64_t getPartialStructRequiredAlignment(void *HstPtrBase) {
+  auto BaseAlignment = reinterpret_cast<uintptr_t>(HstPtrBase) % MaxAlignment;
+  return BaseAlignment == 0 ? MaxAlignment : BaseAlignment;
+}
 
 /// Map global data and execute pending ctors
 static int initLibrary(DeviceTy &Device) {
@@ -585,6 +592,7 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
     const int NextI = I + 1;
     if (getParentIndex(ArgTypes[I]) < 0 && NextI < ArgNum &&
         getParentIndex(ArgTypes[NextI]) == I) {
+      int64_t Alignment = getPartialStructRequiredAlignment(HstPtrBase);
       Padding = (int64_t)HstPtrBegin % Alignment;
       if (Padding) {
         DP("Using a padding of %" PRId64 " bytes for begin address " DPxMOD
@@ -932,6 +940,7 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
     }
 
     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
@@ -939,6 +948,7 @@ int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
     const int NextI = I + 1;
     if (getParentIndex(ArgTypes[I]) < 0 && NextI < ArgNum &&
         getParentIndex(ArgTypes[NextI]) == I) {
+      int64_t Alignment = getPartialStructRequiredAlignment(HstPtrBase);
       int64_t Padding = (int64_t)HstPtrBegin % Alignment;
       if (Padding) {
         DP("Using a Padding of %" PRId64 " bytes for begin address " DPxMOD
@@ -1293,8 +1303,8 @@ class PrivateArgumentManagerTy {
     FirstPrivateArgInfoTy(int Index, void *HstPtr, int64_t Size,
                           const map_var_info_t HstPtrName = nullptr)
         : Index(Index), HstPtrBegin(reinterpret_cast<char *>(HstPtr)),
-          HstPtrEnd(HstPtrBegin + Size), AlignedSize(Size + Size % Alignment),
-          HstPtrName(HstPtrName) {}
+          HstPtrEnd(HstPtrBegin + Size),
+          AlignedSize(Size + Size % MaxAlignment), HstPtrName(HstPtrName) {}
   };
 
   /// A vector of target pointers for all private arguments

diff  --git a/openmp/libomptarget/test/mapping/low_alignment.c b/openmp/libomptarget/test/mapping/low_alignment.c
new file mode 100644
index 0000000000000..615a5a9c31112
--- /dev/null
+++ b/openmp/libomptarget/test/mapping/low_alignment.c
@@ -0,0 +1,49 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  struct S {
+    int i;
+    int j;
+  } s;
+  s.i = 20;
+  s.j = 30;
+#pragma omp target data map(tofrom : s)
+  {
+#pragma omp target map(from : s.i, s.j)
+    {
+      s.i = 21;
+      s.j = 31;
+    }
+  }
+  if (s.i == 21 && s.j == 31)
+    printf("PASS 1\n");
+  // CHECK: PASS 1
+
+  struct T {
+    int a;
+    int b;
+    int c;
+    int d;
+    int i;
+    int j;
+  } t;
+  t.a = 10;
+  t.i = 20;
+  t.j = 30;
+#pragma omp target data map(from : t.i, t.j)
+  {
+#pragma omp target map(from : t.a)
+    {
+      t.a = 11;
+      t.i = 21;
+      t.j = 31;
+    }
+  }
+  if (t.a == 11 && t.i == 21 && t.j == 31)
+    printf("PASS 2\n");
+  // CHECK: PASS 2
+  return 0;
+}