[Openmp-commits] [PATCH] D52655: [libomptarget-nvptx] Align data sharing stack

[Jonas Hahnfeld via Phabricator via Openmp-commits]

Hahnfeld created this revision.
Hahnfeld added reviewers: gtbercea, grokos, AlexEichenberger.
Herald added a subscriber: openmp-commits.

NVPTX requires addresses of pointer locations to be 8-byte aligned
or there will be an exception during runtime.
This could happen without this patch as shown in the added test:
getId() requires 4 byte of stack and putValueInParallel() uses 16
bytes to store the addresses of the captured variables.


Repository:
  rOMP OpenMP

https://reviews.llvm.org/D52655

Files:
  libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
  libomptarget/deviceRTLs/nvptx/test/data_sharing/alignment.c


Index: libomptarget/deviceRTLs/nvptx/test/data_sharing/alignment.c
===================================================================
--- /dev/null
+++ libomptarget/deviceRTLs/nvptx/test/data_sharing/alignment.c
@@ -0,0 +1,55 @@
+// RUN: %compile-run-and-check
+
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp declare target
+static void putValueInParallel(int *ptr, int value) {
+  #pragma omp parallel
+  {
+    *ptr = value;
+  }
+}
+
+static int getId() {
+  int id;
+  putValueInParallel(&id, omp_get_thread_num());
+  return id;
+}
+#pragma omp end declare target
+
+const int MaxThreads = 1024;
+const int Threads = 64;
+
+int main(int argc, char *argv[]) {
+  int master;
+  int check[MaxThreads];
+  for (int i = 0; i < MaxThreads; i++) {
+    check[i] = 0;
+  }
+
+  #pragma omp target map(master, check[:])
+  {
+    master = getId();
+
+    #pragma omp parallel num_threads(Threads)
+    {
+      check[omp_get_thread_num()] = getId();
+    }
+  }
+
+  // CHECK: master = 0.
+  printf("master = %d.\n", master);
+  // CHECK-NOT: invalid
+  for (int i = 0; i < MaxThreads; i++) {
+    if (i < Threads) {
+      if (check[i] != i) {
+        printf("invalid: check[%d] should be %d, is %d\n", i, i, check[i]);
+      }
+    } else if (check[i] != 0) {
+      printf("invalid: check[%d] should be 0, is %d\n", i, check[i]);
+    }
+  }
+
+  return 0;
+}
Index: libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
===================================================================
--- libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
+++ libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
@@ -389,6 +389,13 @@
     return omptarget_nvptx_SimpleThreadPrivateContext::Allocate(DataSize);
   }
 
+  // Add worst-case padding to DataSize so that future stack allocations are
+  // correctly aligned.
+  const size_t Alignment = 8;
+  if (DataSize % Alignment != 0) {
+    DataSize += (Alignment - DataSize % Alignment);
+  }
+
   // Frame pointer must be visible to all workers in the same warp.
   unsigned WID = getWarpId();
   void *&FrameP = DataSharingState.FramePtr[WID];


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D52655.167480.patch
Type: text/x-patch
Size: 2088 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20180928/0c409255/attachment.bin>