[OpenMP][FIX] Do not overalign mapped structures

[jdoerfert]

Summary

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.

[jdoerfert]

Argh, this is not right, it's not what I proposed in the other patch. Will update it later.

[jdoerfert]

Passes the tests now.

[jdoerfert]

Add second test scenario.

[jhuber6]

openmp/libomptarget/src/omptarget.cpp

113

Should try to be consistent with C++ style casts.

[grokos]

openmp/libomptarget/src/omptarget.cpp

77–79

Can you fix this comment as well? Alignment is not necessarily 8 anymore, it could be up to 16.

[pavelkopyl]

Thank you for the patch! I checked it with the D135462 PR that changes alloca alignment - everything works fine.

[pavelkopyl]

openmp/libomptarget/test/mapping/low_alignment.c

11

The only thing is that alignment of the structure 's' may vary depending on many factors. For example, on x64 it's likely 8-aligned. I mean here we check not lowest possible alignment, but rather arbitrary one. Is that OK?

[jdoerfert]

openmp/libomptarget/test/mapping/low_alignment.c

11

The struct (should be) aligned properly for the member that requires the largest alignment. Which is exactly what we are looking for because we might represent only part of the struct and that part might contain this member.

[jhuber6]

Addressing comments

openmp/libomptarget/src/omptarget.cpp

int32_t AsyncInfoTy::runPostProcessing() {
   // procedures might add new procedures themselves.
   const auto PrevBegin = PostProcessingFunctions.begin();
   PostProcessingFunctions.erase(PrevBegin, PrevBegin + Size);
 
   return OFFLOAD_SUCCESS;
 }
 
 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
  *   int b;   // 4-aligned
  *   int *p;  // 8-aligned
  * } s1;
  * ...
  * #pragma omp target map(tofrom: s1.b, s1.p[0:N])

  * address of p will be a misaligned 0x204 (on the host there was no need to add
  * padding between b and p, so p comes exactly 4 bytes after b). If the device
  * kernel tries to access s1.p, a misaligned address error occurs (as reported
  * by the CUDA plugin). By padding the begin address down to a multiple of 8 and
  * extending the size of the allocated chuck accordingly, the chuck on the
  * 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) {
   /*
    * Map global data
    */
   int32_t DeviceId = Device.DeviceID;
   int Rc = OFFLOAD_SUCCESS;
for (int32_t I = 0; I < ArgNum; ++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
     // is a combined entry.
     int64_t Padding = 0;
     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
            "\n",
            Padding, DPxPTR(HstPtrBegin));
         HstPtrBegin = (char *)HstPtrBegin - Padding;
         DataSize += Padding;
       }
if (ArgMappers && ArgMappers[I]) {
         return OFFLOAD_FAIL;
       }
 
       // Skip the rest of this function, continue to the next argument.
       continue;
     }
 
     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
     // is a combined entry.
     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
            "\n",
            Padding, DPxPTR(HstPtrBegin));
         HstPtrBegin = (char *)HstPtrBegin - Padding;
         DataSize += Padding;
       }
struct FirstPrivateArgInfoTy {
     /// Aligned size
     int64_t AlignedSize;
     /// Host pointer name
     map_var_info_t HstPtrName = nullptr;
 
     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
   SmallVector<void *> TgtPtrs;
 
   /// A vector of information of all first-private arguments to be packed
   SmallVector<FirstPrivateArgInfoTy> FirstPrivateArgInfo;
   /// Host buffer for all arguments to be packed

openmp/libomptarget/test/mapping/low_alignment.c

+// 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;
+}