[Openmp-commits] [openmp] r361638 - Revert "[OPENMP][NVPTX]Fix barriers and parallel level counters, NFC."

[Alexey Bataev via Openmp-commits]

Author: abataev
Date: Fri May 24 07:06:47 2019
New Revision: 361638

URL: http://llvm.org/viewvc/llvm-project?rev=361638&view=rev
Log:
Revert "[OPENMP][NVPTX]Fix barriers and parallel level counters, NFC."

This reverts commit r361421 to split the patch into 3 parts.

Modified:
    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

Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omp_data.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omp_data.cu?rev=361638&r1=361637&r2=361638&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omp_data.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omp_data.cu Fri May 24 07:06:47 2019
@@ -31,8 +31,7 @@ __device__ omptarget_nvptx_SimpleMemoryM
 __device__ __shared__ uint32_t usedMemIdx;
 __device__ __shared__ uint32_t usedSlotIdx;
 
-__device__ __shared__ volatile uint8_t
-    parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
+__device__ __shared__ uint8_t parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
 __device__ __shared__ uint16_t threadLimit;
 __device__ __shared__ uint16_t threadsInTeam;
 __device__ __shared__ uint16_t nThreads;

Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h?rev=361638&r1=361637&r2=361638&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h Fri May 24 07:06:47 2019
@@ -398,7 +398,7 @@ extern __device__ omptarget_nvptx_Simple
     omptarget_nvptx_simpleMemoryManager;
 extern __device__ __shared__ uint32_t usedMemIdx;
 extern __device__ __shared__ uint32_t usedSlotIdx;
-extern __device__ __shared__ volatile uint8_t
+extern __device__ __shared__ uint8_t
     parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
 extern __device__ __shared__ uint16_t threadLimit;
 extern __device__ __shared__ uint16_t threadsInTeam;

Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu?rev=361638&r1=361637&r2=361638&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/sync.cu Fri May 24 07:06:47 2019
@@ -62,8 +62,6 @@ EXTERN void __kmpc_barrier(kmp_Ident *lo
         // 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");
   }
@@ -132,7 +130,7 @@ EXTERN void __kmpc_end_single(kmp_Ident
 
 EXTERN void __kmpc_flush(kmp_Ident *loc) {
   PRINT0(LD_IO, "call kmpc_flush\n");
-  __threadfence();
+  __threadfence_system();
 }
 
 ////////////////////////////////////////////////////////////////////////////////