diff options
| -rw-r--r-- | openmp/libomptarget/deviceRTLs/nvptx/src/omp_data.cu | 3 | ||||
| -rw-r--r-- | openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h | 2 | ||||
| -rw-r--r-- | openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu | 4 |
3 files changed, 3 insertions, 6 deletions
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/omp_data.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/omp_data.cu index 0cd9b57fd7c..d369da1cb7e 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/omp_data.cu +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/omp_data.cu @@ -31,8 +31,7 @@ __device__ omptarget_nvptx_SimpleMemoryManager __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; diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h index b85d0a750f2..cd51538ad79 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h @@ -398,7 +398,7 @@ extern __device__ omptarget_nvptx_SimpleMemoryManager 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; diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu index 191b046c9f4..d81aa8f0f3a 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu @@ -62,8 +62,6 @@ EXTERN void __kmpc_barrier(kmp_Ident *loc_ref, int32_t tid) { // 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 *loc, int32_t global_tid) { EXTERN void __kmpc_flush(kmp_Ident *loc) { PRINT0(LD_IO, "call kmpc_flush\n"); - __threadfence(); + __threadfence_system(); } //////////////////////////////////////////////////////////////////////////////// |
