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

[Jonas Hahnfeld via Phabricator via Openmp-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rL343402: [libomptarget-nvptx] Align data sharing stack (authored by Hahnfeld, committed by ).
Herald added a subscriber: llvm-commits.

Changed prior to commit:
  https://reviews.llvm.org/D52655?vs=167480&id=167638#toc

Repository:
  rL LLVM

https://reviews.llvm.org/D52655

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


Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
===================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
@@ -384,6 +384,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];
Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/test/data_sharing/alignment.c
===================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/test/data_sharing/alignment.c
+++ openmp/trunk/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;
+}


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D52655.167638.patch
Type: text/x-patch
Size: 2216 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20180930/f712b2e2/attachment.bin>