[Openmp-commits] [PATCH] D88474: [nfc][libomptarget] Drop parameter to named_sync

[Jon Chesterfield via Phabricator via Openmp-commits]

JonChesterfield created this revision.
JonChesterfield added reviewers: jdoerfert, ABataev, grokos, jhuber6.
Herald added subscribers: openmp-commits, aaron.ballman, tpr, jvesely.
Herald added a project: OpenMP.
JonChesterfield requested review of this revision.
Herald added a subscriber: sstefan1.

[nfc][libomptarget] Drop parameter to named_sync

named_sync has one call site (in sync.cu) where it always passed L1_BARRIER.
Folding this into the call site and dropping the macro is a simplification.

amdgpu doesn't have ptx' bar.sync instruction. A correct implementation of
__kmpc_impl_named_sync in terms of shared memory is much easier if it can
assume that the barrier argument is this constant. Said implementation is left
for a second patch.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D88474

Files:
  openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
  openmp/libomptarget/deviceRTLs/common/src/sync.cu
  openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h


Index: openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
===================================================================
--- openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
+++ 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)
Index: openmp/libomptarget/deviceRTLs/common/src/sync.cu
===================================================================
--- openmp/libomptarget/deviceRTLs/common/src/sync.cu
+++ 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.
Index: openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
===================================================================
--- openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
+++ 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();
 }
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D88474.294921.patch
Type: text/x-patch
Size: 2818 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20200929/ee45ca1f/attachment.bin>