Add host API for enqueuing barrier on given stream (#274)

* add host API for enqueuing barrier on given stream

[ROCm/rocshmem commit: a44b581997]

projects/rocshmem/include/rocshmem/rocshmem.hpp

@@ -341,6 +341,13 @@ __host__ void rocshmem_quiet();
  */
 __host__ void rocshmem_barrier_all();
 
+/**
+ * @brief enqueues a collective barrier on given stream.
+ *
+ * @return void
+ */
+__host__ void rocshmem_barrier_all_on_stream(hipStream_t stream);
+
 /**
  * @brief registers the arrival of a PE at a barrier.
  * The caller is blocked until the synchronization is resolved.

projects/rocshmem/include/rocshmem/rocshmem_COLL.hpp

@@ -599,6 +599,14 @@ __host__ int rocshmem_ctx_double_prod_reduce(
     rocshmem_ctx_t ctx, rocshmem_team_t team, double *dest, const double *source,
     int nreduce);
 
+/**
+ * @brief kernel for performing a barrier synchronization.
+ * Caller enqueues the kernel on given stream
+ *
+ * @return void
+ */
+__global__ ATTR_NO_INLINE void rocshmem_barrier_all_kernel();
+
 /**
  * @brief perform a collective barrier between all PEs in the system.
  * The caller is blocked until the barrier is resolved.