[OpenMP] Add ompx wrappers for "syncthreads"

[tianshilei1992]

openmp/libomptarget/DeviceRTL/src/Synchronization.cpp

602

nfc: empty line between functions

openmp/libomptarget/DeviceRTL/src/Synchronization.cpp

 
 void omp_destroy_lock(omp_lock_t *Lock) { impl::destroyLock(Lock); }
 
 void omp_set_lock(omp_lock_t *Lock) { impl::setLock(Lock); }
 
 void omp_unset_lock(omp_lock_t *Lock) { impl::unsetLock(Lock); }
 
 int omp_test_lock(omp_lock_t *Lock) { return impl::testLock(Lock); }
+
+void ompx_sync_block(int Ordering) {
+  impl::syncThreadsAligned(atomic::OrderingTy(Ordering));
+}
+void ompx_sync_block_acq_rel() {
+  impl::syncThreadsAligned(atomic::OrderingTy::acq_rel);
+}
+void ompx_sync_block_divergent(int Ordering) {
+  impl::syncThreads(atomic::OrderingTy(Ordering));
+}
 } // extern "C"
 
 #pragma omp end declare target

openmp/libomptarget/test/api/ompx_sync.c

+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <ompx.h>
+#include <stdio.h>
+
+void foo(int device) {
+  int X;
+  // clang-format off
+#pragma omp target teams map(from: X) device(device) thread_limit(2) num_teams(1)
+#pragma omp parallel
+  // clang-format on
+  {
+    int tid = ompx_thread_id_x();
+    int bid = ompx_block_id_x();
+    if (tid == 1 && bid == 0) {
+      X = 42;
+      ompx_sync_block_divergent(3);
+    } else {
+      ompx_sync_block_divergent(1);
+    }
+    if (tid == 0 && bid == 0)
+      X++;
+    ompx_sync_block(ompx_seq_cst);
+    if (tid == 1 && bid == 0)
+      X++;
+    ompx_sync_block_acq_rel();
+    if (tid == 0 && bid == 0)
+      X++;
+    ompx_sync_block(ompx_release);
+    if (tid == 0 && bid == 0)
+      X++;
+  }
+  // CHECK: X: 46
+  // CHECK: X: 46
+  printf("X: %i\n", X);
+}
+
+int main() {
+  foo(omp_get_default_device());
+  foo(omp_get_initial_device());
+}

openmp/libomptarget/test/api/ompx_sync.cpp

+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+#include <omp.h>
+#include <ompx.h>
+#include <stdio.h>
+
+void foo(int device) {
+  int X;
+  // clang-format off
+#pragma omp target teams map(from: X) device(device) thread_limit(2) num_teams(1)
+#pragma omp parallel
+  // clang-format on
+  {
+    int tid = ompx::thread_id_x();
+    int bid = ompx::block_id_x();
+    if (tid == 1 && bid == 0) {
+      X = 42;
+      ompx::sync_block_divergent(3);
+    } else {
+      ompx::sync_block_divergent();
+    }
+    if (tid == 0 && bid == 0)
+      X++;
+    ompx::sync_block(ompx::seq_cst);
+    if (tid == 1 && bid == 0)
+      X++;
+    ompx::sync_block();
+    if (tid == 0 && bid == 0)
+      X++;
+    ompx_sync_block(ompx_release);
+    if (tid == 0 && bid == 0)
+      X++;
+  }
+  // CHECK: X: 46
+  // CHECK: X: 46
+  printf("X: %i\n", X);
+}
+
+int main() {
+  foo(omp_get_default_device());
+  foo(omp_get_initial_device());
+}

openmp/runtime/src/include/ompx.h.var

 ///
 ///{
 
 #ifdef __cplusplus
 extern "C" {
 #endif
 
 enum {
+  ompx_relaxed = __ATOMIC_RELAXED,
+  ompx_aquire = __ATOMIC_ACQUIRE,
+  ompx_release = __ATOMIC_RELEASE,
+  ompx_acq_rel = __ATOMIC_ACQ_REL,
+  ompx_seq_cst = __ATOMIC_SEQ_CST,
+};
+
+enum {
   ompx_dim_x = 0,
   ompx_dim_y = 1,
   ompx_dim_z = 2,
 };
 
 /// ompx_{thread,block}_{id,dim}
 ///{
 #pragma omp begin declare variant match(device = {kind(cpu)})
 #define _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C(NAME, VALUE)                     \
   static inline int ompx_##NAME(int Dim) { return VALUE; }
 
 _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C(thread_id,
                                       omp_get_ancestor_thread_num(Dim + 1))
 _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C(thread_dim, omp_get_team_size(Dim + 1))
 _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C(block_id, 0)
 _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C(block_dim, 1)
 #undef _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_C
 ///}
 
+/// ompx_{sync_block}_{,divergent}
+///{
+#define _TGT_KERNEL_LANGUAGE_HOST_IMPL_SYNC_C(RETTY, NAME, ARGS, BODY)         \
+  static inline RETTY ompx_##NAME(ARGS) { BODY; }
+
+_TGT_KERNEL_LANGUAGE_HOST_IMPL_SYNC_C(void, sync_block, int Ordering,
+                                      _Pragma("omp barrier"));
+_TGT_KERNEL_LANGUAGE_HOST_IMPL_SYNC_C(void, sync_block_acq_rel, void,
+                                      ompx_sync_block(ompx_acq_rel));
+_TGT_KERNEL_LANGUAGE_HOST_IMPL_SYNC_C(void, sync_block_divergent, int Ordering,
+                                      ompx_sync_block(Ordering));
+#undef _TGT_KERNEL_LANGUAGE_HOST_IMPL_SYNC_C
+///}
+
 #pragma omp end declare variant
 
+/// ompx_{sync_block}_{,divergent}
+///{
+#define _TGT_KERNEL_LANGUAGE_DECL_SYNC_C(RETTY, NAME, ARGS)         \
+  RETTY ompx_##NAME(ARGS);
+
+_TGT_KERNEL_LANGUAGE_DECL_SYNC_C(void, sync_block, int Ordering);
+_TGT_KERNEL_LANGUAGE_DECL_SYNC_C(void, sync_block_acq_rel, void);
+_TGT_KERNEL_LANGUAGE_DECL_SYNC_C(void, sync_block_divergent, int Ordering);
+#undef _TGT_KERNEL_LANGUAGE_DECL_SYNC_C
+///}
+
 /// ompx_{thread,block}_{id,dim}_{x,y,z}
 ///{
 #define _TGT_KERNEL_LANGUAGE_DECL_GRID_C(NAME)                                 \
   int ompx_##NAME(int Dim);                                                    \
   static inline int ompx_##NAME##_x() { return ompx_##NAME(ompx_dim_x); }      \
   static inline int ompx_##NAME##_y() { return ompx_##NAME(ompx_dim_y); }      \
   static inline int ompx_##NAME##_z() { return ompx_##NAME(ompx_dim_z); }
 

 namespace ompx {
 
 enum {
   dim_x = ompx_dim_x,
   dim_y = ompx_dim_y,
   dim_z = ompx_dim_z,
 };
 
+enum {
+  relaxed = ompx_relaxed ,
+  aquire = ompx_aquire,
+  release = ompx_release,
+  acc_rel = ompx_acq_rel,
+  seq_cst = ompx_seq_cst,
+};
+
 /// ompx::{thread,block}_{id,dim}_{,x,y,z}
 ///{
 #define _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX(NAME)                          \
   static inline int NAME(int Dim) noexcept { return ompx_##NAME(Dim); }        \
   static inline int NAME##_x() noexcept { return NAME(ompx_dim_x); }           \
   static inline int NAME##_y() noexcept { return NAME(ompx_dim_y); }           \
   static inline int NAME##_z() noexcept { return NAME(ompx_dim_z); }
 
 _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX(thread_id)
 _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX(thread_dim)
 _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX(block_id)
 _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX(block_dim)
 #undef _TGT_KERNEL_LANGUAGE_HOST_IMPL_GRID_CXX
 ///}
 
+/// ompx_{sync_block}_{,divergent}
+///{
+#define _TGT_KERNEL_LANGUAGE_HOST_IMPL_SYNC_CXX(RETTY, NAME, ARGS, CALL_ARGS)  \
+  static inline RETTY NAME(ARGS) {               \
+    return ompx_##NAME(CALL_ARGS);                                             \
+  }
+
+_TGT_KERNEL_LANGUAGE_HOST_IMPL_SYNC_CXX(void, sync_block, int Ordering = acc_rel,
+                                        Ordering);
+_TGT_KERNEL_LANGUAGE_HOST_IMPL_SYNC_CXX(void, sync_block_divergent,
+                                        int Ordering = acc_rel, Ordering);
+#undef _TGT_KERNEL_LANGUAGE_HOST_IMPL_SYNC_CXX
+///}
+
 } // namespace ompx
 #endif
 
 ///}
 
 #endif /* __OMPX_H */