[Openmp-commits] [openmp] r343402 - [libomptarget-nvptx] Align data sharing stack

[Jonas Hahnfeld via Openmp-commits]

Author: hahnfeld
Date: Sun Sep 30 02:23:21 2018
New Revision: 343402

URL: http://llvm.org/viewvc/llvm-project?rev=343402&view=rev
Log:
[libomptarget-nvptx] Align data sharing stack

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.

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

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

Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu?rev=343402&r1=343401&r2=343402&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu Sun Sep 30 02:23:21 2018
@@ -384,6 +384,13 @@ EXTERN void* __kmpc_data_sharing_push_st
     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];

Added: openmp/trunk/libomptarget/deviceRTLs/nvptx/test/data_sharing/alignment.c
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/test/data_sharing/alignment.c?rev=343402&view=auto
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/test/data_sharing/alignment.c (added)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/test/data_sharing/alignment.c Sun Sep 30 02:23:21 2018
@@ -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;
+}