[Openmp-commits] [PATCH] D62199: [OPENMP][NVPTX]Fix barriers and parallel level counters, NFC.

[Alexey Bataev via Phabricator via Openmp-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rL361421: [OPENMP][NVPTX]Fix barriers and parallel level counters, NFC. (authored by ABataev, committed by ).
Herald added a project: LLVM.
Herald added a subscriber: llvm-commits.

Repository:
  rL LLVM

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D62199/new/

https://reviews.llvm.org/D62199

Files:
  openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omp_data.cu
  openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
  openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu


Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omp_data.cu
===================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omp_data.cu
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omp_data.cu
@@ -31,7 +31,8 @@
 __device__ __shared__ uint32_t usedMemIdx;
 __device__ __shared__ uint32_t usedSlotIdx;
 
-__device__ __shared__ uint8_t parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
+__device__ __shared__ volatile uint8_t
+    parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
 __device__ __shared__ uint16_t threadLimit;
 __device__ __shared__ uint16_t threadsInTeam;
 __device__ __shared__ uint16_t nThreads;
Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu
===================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu
@@ -62,6 +62,8 @@
         // Barrier #1 is for synchronization among active threads.
         named_sync(L1_BARRIER, threads);
       }
+    } else {
+      __kmpc_flush(loc_ref);
     } // numberOfActiveOMPThreads > 1
     PRINT0(LD_SYNC, "completed kmpc_barrier\n");
   }
@@ -130,7 +132,7 @@
 
 EXTERN void __kmpc_flush(kmp_Ident *loc) {
   PRINT0(LD_IO, "call kmpc_flush\n");
-  __threadfence_system();
+  __threadfence();
 }
 
 ////////////////////////////////////////////////////////////////////////////////
Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
===================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
@@ -398,7 +398,7 @@
     omptarget_nvptx_simpleMemoryManager;
 extern __device__ __shared__ uint32_t usedMemIdx;
 extern __device__ __shared__ uint32_t usedSlotIdx;
-extern __device__ __shared__ uint8_t
+extern __device__ __shared__ volatile uint8_t
     parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
 extern __device__ __shared__ uint16_t threadLimit;
 extern __device__ __shared__ uint16_t threadsInTeam;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D62199.200809.patch
Type: text/x-patch
Size: 2124 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20190522/3aee7b2b/attachment.bin>