diff --git a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h --- a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h +++ b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h @@ -42,10 +42,6 @@ #define WARPSIZE 64 -// The named barrier for active parallel threads of a team in an L1 parallel -// region to synchronize with each other. -#define L1_BARRIER (1) - // Maximum number of preallocated arguments to an outlined parallel/simd // function. Anything more requires dynamic memory allocation. #define MAX_SHARED_ARGS 20 @@ -113,10 +109,9 @@ // AMDGCN doesn't need to sync threads in a warp } -INLINE void __kmpc_impl_named_sync(int barrier, uint32_t num_threads) { - // we have protected the master warp from releasing from its barrier - // due to a full workgroup barrier in the middle of a work function. - // So it is ok to issue a full workgroup barrier here. +INLINE void __kmpc_impl_named_sync(uint32_t num_threads) { + (void)num_threads; + // TODO: Implement on top of __SHARED__ __builtin_amdgcn_s_barrier(); } diff --git a/openmp/libomptarget/deviceRTLs/common/src/sync.cu b/openmp/libomptarget/deviceRTLs/common/src/sync.cu --- a/openmp/libomptarget/deviceRTLs/common/src/sync.cu +++ b/openmp/libomptarget/deviceRTLs/common/src/sync.cu @@ -60,8 +60,7 @@ PRINT(LD_SYNC, "call kmpc_barrier with %d omp threads, sync parameter %d\n", (int)numberOfActiveOMPThreads, (int)threads); - // Barrier #1 is for synchronization among active threads. - __kmpc_impl_named_sync(L1_BARRIER, threads); + __kmpc_impl_named_sync(threads); } } else { // Still need to flush the memory per the standard. diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h --- a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h @@ -37,10 +37,6 @@ #define WARPSIZE 32 -// The named barrier for active parallel threads of a team in an L1 parallel -// region to synchronize with each other. -#define L1_BARRIER (1) - // Maximum number of preallocated arguments to an outlined parallel/simd function. // Anything more requires dynamic memory allocation. #define MAX_SHARED_ARGS 20 @@ -187,7 +183,10 @@ #endif // CUDA_VERSION } -INLINE void __kmpc_impl_named_sync(int barrier, uint32_t num_threads) { +INLINE void __kmpc_impl_named_sync(uint32_t num_threads) { + // The named barrier for active parallel threads of a team in an L1 parallel + // region to synchronize with each other. + int barrier = 1; asm volatile("bar.sync %0, %1;" : : "r"(barrier), "r"(num_threads)