Update collective APIs naming (#77)
* Update the naming convention for collective APIs to ensure consistency across the interface. * Move all collective API declarations to rocshmem_COLL.hpp * The following APIs were updated as part of this change: - `barrier` - `barrier_all` - `sync` - `sync_all` - `all_to_all` - `broadcast` - `fcollect` - `all_reduce` * Update header file generation code for collective APIs
This commit is contained in:
Avinash Kethineedi authored and 
GitHub committed
1 parent
5e49567b6c
commit
68421895d6
15 files changed
+334
-334
No files matched your search
@@ -490,210 +490,6 @@ __device__ int rocshmem_team_translate_pe(rocshmem_team_t src_team,
|
||||
int src_pe,
|
||||
rocshmem_team_t dest_team);
|
||||
|
||||
/**
|
||||
* @brief perform a collective barrier between all PEs in the system.
|
||||
* The caller is blocked until the barrier is resolved.
|
||||
*
|
||||
* This function must be invoked by a single thread within the PE.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_barrier_all(
|
||||
rocshmem_ctx_t ctx);
|
||||
|
||||
/**
|
||||
* @brief perform a collective barrier between all PEs in the system.
|
||||
* The caller is blocked until the barrier is resolved.
|
||||
*
|
||||
* This function must be called as a wave-front collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_wave_barrier_all(
|
||||
rocshmem_ctx_t ctx);
|
||||
|
||||
/**
|
||||
* @brief perform a collective barrier between all PEs in the system.
|
||||
* The caller is blocked until the barrier is resolved.
|
||||
*
|
||||
* This function must be called as a work-group collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_wg_barrier_all(
|
||||
rocshmem_ctx_t ctx);
|
||||
|
||||
/**
|
||||
* @brief perform a collective barrier between all PEs in the team.
|
||||
* The caller is blocked until the barrier is resolved.
|
||||
*
|
||||
* This function must be invoked by a single thread within the PE.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @param[in] team The team on which to perform barrier synchronization
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ void rocshmem_ctx_barrier(rocshmem_ctx_t ctx, rocshmem_team_t team);
|
||||
|
||||
/**
|
||||
* @brief perform a collective barrier between all PEs in the team.
|
||||
* The caller is blocked until the barrier is resolved.
|
||||
*
|
||||
* This function must be called as a wave-front collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @param[in] team The team on which to perform barrier synchronization
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ void rocshmem_ctx_wave_barrier(rocshmem_ctx_t ctx, rocshmem_team_t team);
|
||||
|
||||
/**
|
||||
* @brief perform a collective barrier between all PEs in the team.
|
||||
* The caller is blocked until the barrier is resolved.
|
||||
*
|
||||
* This function must be called as a work-group collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @param[in] team The team on which to perform barrier synchronization
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ void rocshmem_ctx_wg_barrier(rocshmem_ctx_t ctx, rocshmem_team_t team);
|
||||
|
||||
/**
|
||||
* @brief registers the arrival of a PE at a barrier.
|
||||
* The caller is blocked until the synchronization is resolved.
|
||||
*
|
||||
* In contrast with the shmem_barrier_all routine, shmem_sync_all only ensures
|
||||
* completion and visibility of previously issued memory stores and does not
|
||||
* ensure completion of remote memory updates issued via OpenSHMEM routines.
|
||||
*
|
||||
* This function must be invoked by a single thread within the PE.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_sync_all(rocshmem_ctx_t ctx);
|
||||
|
||||
/**
|
||||
* @brief registers the arrival of a PE at a barrier.
|
||||
* The caller is blocked until the synchronization is resolved.
|
||||
*
|
||||
* In contrast with the shmem_barrier_all routine, shmem_sync_all only ensures
|
||||
* completion and visibility of previously issued memory stores and does not
|
||||
* ensure completion of remote memory updates issued via OpenSHMEM routines.
|
||||
*
|
||||
* This function must be called as a wave-front collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_wave_sync_all(rocshmem_ctx_t ctx);
|
||||
|
||||
/**
|
||||
* @brief registers the arrival of a PE at a barrier.
|
||||
* The caller is blocked until the synchronization is resolved.
|
||||
*
|
||||
* In contrast with the shmem_barrier_all routine, shmem_sync_all only ensures
|
||||
* completion and visibility of previously issued memory stores and does not
|
||||
* ensure completion of remote memory updates issued via OpenSHMEM routines.
|
||||
*
|
||||
* This function must be called as a work-group collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_wg_sync_all(rocshmem_ctx_t ctx);
|
||||
|
||||
/**
|
||||
* @brief registers the arrival of a PE at a barrier.
|
||||
* The caller is blocked until the synchronization is resolved.
|
||||
*
|
||||
* In contrast with the shmem_barrier_all routine, shmem_team_sync only ensures
|
||||
* completion and visibility of previously issued memory stores and does not
|
||||
* ensure completion of remote memory updates issued via OpenSHMEM routines.
|
||||
*
|
||||
* This function must be invoked by a single thread within the PE.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
* @param[in] team Handle of the team being synchronized
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_team_sync(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team);
|
||||
|
||||
/**
|
||||
* @brief registers the arrival of a PE at a barrier.
|
||||
* The caller is blocked until the synchronization is resolved.
|
||||
*
|
||||
* In contrast with the shmem_barrier_all routine, shmem_team_sync only ensures
|
||||
* completion and visibility of previously issued memory stores and does not
|
||||
* ensure completion of remote memory updates issued via OpenSHMEM routines.
|
||||
*
|
||||
* This function must be called as a wave-front collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
* @param[in] team Handle of the team being synchronized
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_wave_team_sync(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team);
|
||||
|
||||
/**
|
||||
* @brief registers the arrival of a PE at a barrier.
|
||||
* The caller is blocked until the synchronization is resolved.
|
||||
*
|
||||
* In contrast with the shmem_barrier_all routine, shmem_team_sync only ensures
|
||||
* completion and visibility of previously issued memory stores and does not
|
||||
* ensure completion of remote memory updates issued via OpenSHMEM routines.
|
||||
*
|
||||
* This function must be called as a work-group collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
* @param[in] team Handle of the team being synchronized
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_wg_team_sync(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team);
|
||||
|
||||
/**
|
||||
* @brief Query a local pointer to a symmetric data object on the
|
||||
* specified \pe . Returns an address that may be used to directly reference
|
||||
* dest on the specified \pe. This address can be accesses with LD/ST ops.
|
||||
*
|
||||
* Can be called per thread with no performance penalty.
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void *rocshmem_ptr(const void *dest, int pe);
|
||||
|
||||
/**
|
||||
* @brief Make all uncacheable GPU data visible to other agents in the sytem.
|
||||
*
|
||||
* This only works for data that was explicitly allocated uncacheable on the
|
||||
* GPU!
|
||||
*
|
||||
* Can be called per thread with no performance penalty.
|
||||
*
|
||||
* @param[in] GPU-side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_threadfence_system(
|
||||
rocshmem_ctx_t ctx);
|
||||
|
||||
|
||||
@@ -41,55 +41,55 @@ namespace rocshmem {
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_float_wg_alltoall(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_float_alltoall_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, float *dest,
|
||||
const float *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_double_wg_alltoall(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_double_alltoall_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, double *dest,
|
||||
const double *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_char_wg_alltoall(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_char_alltoall_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, char *dest,
|
||||
const char *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_schar_wg_alltoall(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_schar_alltoall_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, signed char *dest,
|
||||
const signed char *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_short_wg_alltoall(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_short_alltoall_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest,
|
||||
const short *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_int_wg_alltoall(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_int_alltoall_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest,
|
||||
const int *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_long_wg_alltoall(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_long_alltoall_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest,
|
||||
const long *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_longlong_wg_alltoall(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_longlong_alltoall_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest,
|
||||
const long long *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_uchar_wg_alltoall(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_uchar_alltoall_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned char *dest,
|
||||
const unsigned char *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ushort_wg_alltoall(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ushort_alltoall_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned short *dest,
|
||||
const unsigned short *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_uint_wg_alltoall(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_uint_alltoall_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned int *dest,
|
||||
const unsigned int *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ulong_wg_alltoall(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ulong_alltoall_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned long *dest,
|
||||
const unsigned long *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ulonglong_wg_alltoall(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ulonglong_alltoall_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned long long *dest,
|
||||
const unsigned long long *source, int nelems);
|
||||
|
||||
@@ -116,7 +116,7 @@ __device__ ATTR_NO_INLINE void rocshmem_ctx_ulonglong_wg_alltoall(
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_float_wg_broadcast(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_float_broadcast_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, float *dest,
|
||||
const float *source, int nelems, int pe_root);
|
||||
__host__ void rocshmem_ctx_float_broadcast(
|
||||
@@ -127,7 +127,7 @@ __host__ void rocshmem_ctx_float_broadcast(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, float *dest,
|
||||
const float *source, int nelems, int pe_root);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_double_wg_broadcast(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_double_broadcast_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, double *dest,
|
||||
const double *source, int nelems, int pe_root);
|
||||
__host__ void rocshmem_ctx_double_broadcast(
|
||||
@@ -138,7 +138,7 @@ __host__ void rocshmem_ctx_double_broadcast(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, double *dest,
|
||||
const double *source, int nelems, int pe_root);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_char_wg_broadcast(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_char_broadcast_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, char *dest,
|
||||
const char *source, int nelems, int pe_root);
|
||||
__host__ void rocshmem_ctx_char_broadcast(
|
||||
@@ -149,7 +149,7 @@ __host__ void rocshmem_ctx_char_broadcast(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, char *dest,
|
||||
const char *source, int nelems, int pe_root);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_schar_wg_broadcast(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_schar_broadcast_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, signed char *dest,
|
||||
const signed char *source, int nelems, int pe_root);
|
||||
__host__ void rocshmem_ctx_schar_broadcast(
|
||||
@@ -160,7 +160,7 @@ __host__ void rocshmem_ctx_schar_broadcast(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, signed char *dest,
|
||||
const signed char *source, int nelems, int pe_root);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_short_wg_broadcast(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_short_broadcast_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest,
|
||||
const short *source, int nelems, int pe_root);
|
||||
__host__ void rocshmem_ctx_short_broadcast(
|
||||
@@ -171,7 +171,7 @@ __host__ void rocshmem_ctx_short_broadcast(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest,
|
||||
const short *source, int nelems, int pe_root);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_int_wg_broadcast(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_int_broadcast_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest,
|
||||
const int *source, int nelems, int pe_root);
|
||||
__host__ void rocshmem_ctx_int_broadcast(
|
||||
@@ -182,7 +182,7 @@ __host__ void rocshmem_ctx_int_broadcast(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest,
|
||||
const int *source, int nelems, int pe_root);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_long_wg_broadcast(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_long_broadcast_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest,
|
||||
const long *source, int nelems, int pe_root);
|
||||
__host__ void rocshmem_ctx_long_broadcast(
|
||||
@@ -193,7 +193,7 @@ __host__ void rocshmem_ctx_long_broadcast(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest,
|
||||
const long *source, int nelems, int pe_root);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_longlong_wg_broadcast(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_longlong_broadcast_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest,
|
||||
const long long *source, int nelems, int pe_root);
|
||||
__host__ void rocshmem_ctx_longlong_broadcast(
|
||||
@@ -204,7 +204,7 @@ __host__ void rocshmem_ctx_longlong_broadcast(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest,
|
||||
const long long *source, int nelems, int pe_root);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_uchar_wg_broadcast(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_uchar_broadcast_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned char *dest,
|
||||
const unsigned char *source, int nelems, int pe_root);
|
||||
__host__ void rocshmem_ctx_uchar_broadcast(
|
||||
@@ -215,7 +215,7 @@ __host__ void rocshmem_ctx_uchar_broadcast(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned char *dest,
|
||||
const unsigned char *source, int nelems, int pe_root);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ushort_wg_broadcast(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ushort_broadcast_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned short *dest,
|
||||
const unsigned short *source, int nelems, int pe_root);
|
||||
__host__ void rocshmem_ctx_ushort_broadcast(
|
||||
@@ -226,7 +226,7 @@ __host__ void rocshmem_ctx_ushort_broadcast(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned short *dest,
|
||||
const unsigned short *source, int nelems, int pe_root);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_uint_wg_broadcast(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_uint_broadcast_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned int *dest,
|
||||
const unsigned int *source, int nelems, int pe_root);
|
||||
__host__ void rocshmem_ctx_uint_broadcast(
|
||||
@@ -237,7 +237,7 @@ __host__ void rocshmem_ctx_uint_broadcast(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned int *dest,
|
||||
const unsigned int *source, int nelems, int pe_root);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ulong_wg_broadcast(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ulong_broadcast_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned long *dest,
|
||||
const unsigned long *source, int nelems, int pe_root);
|
||||
__host__ void rocshmem_ctx_ulong_broadcast(
|
||||
@@ -248,7 +248,7 @@ __host__ void rocshmem_ctx_ulong_broadcast(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned long *dest,
|
||||
const unsigned long *source, int nelems, int pe_root);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ulonglong_wg_broadcast(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ulonglong_broadcast_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned long long *dest,
|
||||
const unsigned long long *source, int nelems, int pe_root);
|
||||
__host__ void rocshmem_ctx_ulonglong_broadcast(
|
||||
@@ -276,55 +276,55 @@ __host__ void rocshmem_ctx_ulonglong_broadcast(
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_float_wg_fcollect(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_float_fcollect_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, float *dest,
|
||||
const float *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_double_wg_fcollect(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_double_fcollect_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, double *dest,
|
||||
const double *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_char_wg_fcollect(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_char_fcollect_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, char *dest,
|
||||
const char *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_schar_wg_fcollect(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_schar_fcollect_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, signed char *dest,
|
||||
const signed char *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_short_wg_fcollect(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_short_fcollect_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest,
|
||||
const short *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_int_wg_fcollect(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_int_fcollect_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest,
|
||||
const int *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_long_wg_fcollect(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_long_fcollect_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest,
|
||||
const long *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_longlong_wg_fcollect(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_longlong_fcollect_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest,
|
||||
const long long *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_uchar_wg_fcollect(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_uchar_fcollect_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned char *dest,
|
||||
const unsigned char *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ushort_wg_fcollect(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ushort_fcollect_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned short *dest,
|
||||
const unsigned short *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_uint_wg_fcollect(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_uint_fcollect_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned int *dest,
|
||||
const unsigned int *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ulong_wg_fcollect(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ulong_fcollect_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned long *dest,
|
||||
const unsigned long *source, int nelems);
|
||||
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ulonglong_wg_fcollect(
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_ulonglong_fcollect_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, unsigned long long *dest,
|
||||
const unsigned long long *source, int nelems);
|
||||
|
||||
@@ -345,258 +345,462 @@ __device__ ATTR_NO_INLINE void rocshmem_ctx_ulonglong_wg_fcollect(
|
||||
*
|
||||
* @return int (Zero on successful local completion. Nonzero otherwise.)
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_sum_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_sum_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_short_sum_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_min_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_min_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_short_min_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_max_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_max_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_short_max_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_prod_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_prod_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_short_prod_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_or_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_or_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_short_or_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_and_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_and_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_short_and_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_xor_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_short_xor_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_short_xor_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, short *dest, const short *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_sum_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_sum_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_int_sum_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_min_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_min_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_int_min_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_max_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_max_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_int_max_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_prod_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_prod_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_int_prod_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_or_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_or_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_int_or_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_and_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_and_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_int_and_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_xor_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_int_xor_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_int_xor_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, int *dest, const int *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_sum_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_sum_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_long_sum_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_min_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_min_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_long_min_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_max_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_max_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_long_max_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_prod_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_prod_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_long_prod_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_or_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_or_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_long_or_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_and_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_and_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_long_and_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_xor_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_long_xor_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_long_xor_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long *dest, const long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_sum_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_sum_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_longlong_sum_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_min_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_min_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_longlong_min_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_max_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_max_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_longlong_max_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_prod_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_prod_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_longlong_prod_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_or_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_or_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_longlong_or_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_and_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_and_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_longlong_and_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_xor_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_longlong_xor_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_longlong_xor_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, long long *dest, const long long *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_float_sum_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_float_sum_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, float *dest, const float *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_float_sum_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, float *dest, const float *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_float_min_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_float_min_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, float *dest, const float *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_float_min_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, float *dest, const float *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_float_max_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_float_max_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, float *dest, const float *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_float_max_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, float *dest, const float *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_float_prod_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_float_prod_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, float *dest, const float *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_float_prod_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, float *dest, const float *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_double_sum_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_double_sum_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, double *dest, const double *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_double_sum_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, double *dest, const double *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_double_min_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_double_min_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, double *dest, const double *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_double_min_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, double *dest, const double *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_double_max_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_double_max_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, double *dest, const double *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_double_max_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, double *dest, const double *source,
|
||||
int nreduce);
|
||||
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_double_prod_wg_reduce(
|
||||
__device__ ATTR_NO_INLINE int rocshmem_ctx_double_prod_reduce_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, double *dest, const double *source,
|
||||
int nreduce);
|
||||
__host__ int rocshmem_ctx_double_prod_reduce(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team, double *dest, const double *source,
|
||||
int nreduce);
|
||||
|
||||
/**
|
||||
* @brief perform a collective barrier between all PEs in the system.
|
||||
* The caller is blocked until the barrier is resolved.
|
||||
*
|
||||
* This function must be invoked by a single thread within the PE.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_barrier_all(
|
||||
rocshmem_ctx_t ctx);
|
||||
|
||||
/**
|
||||
* @brief perform a collective barrier between all PEs in the system.
|
||||
* The caller is blocked until the barrier is resolved.
|
||||
*
|
||||
* This function must be called as a wave-front collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_barrier_all_wave(
|
||||
rocshmem_ctx_t ctx);
|
||||
|
||||
/**
|
||||
* @brief perform a collective barrier between all PEs in the system.
|
||||
* The caller is blocked until the barrier is resolved.
|
||||
*
|
||||
* This function must be called as a work-group collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_barrier_all_wg(
|
||||
rocshmem_ctx_t ctx);
|
||||
|
||||
/**
|
||||
* @brief perform a collective barrier between all PEs in the team.
|
||||
* The caller is blocked until the barrier is resolved.
|
||||
*
|
||||
* This function must be invoked by a single thread within the PE.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @param[in] team The team on which to perform barrier synchronization
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ void rocshmem_ctx_barrier(rocshmem_ctx_t ctx, rocshmem_team_t team);
|
||||
|
||||
/**
|
||||
* @brief perform a collective barrier between all PEs in the team.
|
||||
* The caller is blocked until the barrier is resolved.
|
||||
*
|
||||
* This function must be called as a wave-front collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @param[in] team The team on which to perform barrier synchronization
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ void rocshmem_ctx_barrier_wave(rocshmem_ctx_t ctx, rocshmem_team_t team);
|
||||
|
||||
/**
|
||||
* @brief perform a collective barrier between all PEs in the team.
|
||||
* The caller is blocked until the barrier is resolved.
|
||||
*
|
||||
* This function must be called as a work-group collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @param[in] team The team on which to perform barrier synchronization
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ void rocshmem_ctx_barrier_wg(rocshmem_ctx_t ctx, rocshmem_team_t team);
|
||||
|
||||
/**
|
||||
* @brief registers the arrival of a PE at a barrier.
|
||||
* The caller is blocked until the synchronization is resolved.
|
||||
*
|
||||
* In contrast with the shmem_barrier_all routine, shmem_sync_all only ensures
|
||||
* completion and visibility of previously issued memory stores and does not
|
||||
* ensure completion of remote memory updates issued via OpenSHMEM routines.
|
||||
*
|
||||
* This function must be invoked by a single thread within the PE.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_sync_all(rocshmem_ctx_t ctx);
|
||||
|
||||
/**
|
||||
* @brief registers the arrival of a PE at a barrier.
|
||||
* The caller is blocked until the synchronization is resolved.
|
||||
*
|
||||
* In contrast with the shmem_barrier_all routine, shmem_sync_all only ensures
|
||||
* completion and visibility of previously issued memory stores and does not
|
||||
* ensure completion of remote memory updates issued via OpenSHMEM routines.
|
||||
*
|
||||
* This function must be called as a wave-front collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_sync_all_wave(rocshmem_ctx_t ctx);
|
||||
|
||||
/**
|
||||
* @brief registers the arrival of a PE at a barrier.
|
||||
* The caller is blocked until the synchronization is resolved.
|
||||
*
|
||||
* In contrast with the shmem_barrier_all routine, shmem_sync_all only ensures
|
||||
* completion and visibility of previously issued memory stores and does not
|
||||
* ensure completion of remote memory updates issued via OpenSHMEM routines.
|
||||
*
|
||||
* This function must be called as a work-group collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_sync_all_wg(rocshmem_ctx_t ctx);
|
||||
|
||||
/**
|
||||
* @brief registers the arrival of a PE at a barrier.
|
||||
* The caller is blocked until the synchronization is resolved.
|
||||
*
|
||||
* In contrast with the shmem_barrier_all routine, shmem_team_sync only ensures
|
||||
* completion and visibility of previously issued memory stores and does not
|
||||
* ensure completion of remote memory updates issued via OpenSHMEM routines.
|
||||
*
|
||||
* This function must be invoked by a single thread within the PE.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
* @param[in] team Handle of the team being synchronized
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_sync(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team);
|
||||
|
||||
/**
|
||||
* @brief registers the arrival of a PE at a barrier.
|
||||
* The caller is blocked until the synchronization is resolved.
|
||||
*
|
||||
* In contrast with the shmem_barrier_all routine, shmem_team_sync only ensures
|
||||
* completion and visibility of previously issued memory stores and does not
|
||||
* ensure completion of remote memory updates issued via OpenSHMEM routines.
|
||||
*
|
||||
* This function must be called as a wave-front collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
* @param[in] team Handle of the team being synchronized
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_sync_wave(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team);
|
||||
|
||||
/**
|
||||
* @brief registers the arrival of a PE at a barrier.
|
||||
* The caller is blocked until the synchronization is resolved.
|
||||
*
|
||||
* In contrast with the shmem_barrier_all routine, shmem_team_sync only ensures
|
||||
* completion and visibility of previously issued memory stores and does not
|
||||
* ensure completion of remote memory updates issued via OpenSHMEM routines.
|
||||
*
|
||||
* This function must be called as a work-group collective.
|
||||
*
|
||||
* @param[in] handle GPU side handle.
|
||||
* @param[in] team Handle of the team being synchronized
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void rocshmem_ctx_sync_wg(
|
||||
rocshmem_ctx_t ctx, rocshmem_team_t team);
|
||||
|
||||
/**
|
||||
* @brief Query a local pointer to a symmetric data object on the
|
||||
* specified \pe . Returns an address that may be used to directly reference
|
||||
* dest on the specified \pe. This address can be accesses with LD/ST ops.
|
||||
*
|
||||
* Can be called per thread with no performance penalty.
|
||||
*/
|
||||
__device__ ATTR_NO_INLINE void *rocshmem_ptr(const void *dest, int pe);
|
||||
|
||||
/**
|
||||
* @brief Make all uncacheable GPU data visible to other agents in the sytem.
|
||||
*
|
||||
* This only works for data that was explicitly allocated uncacheable on the
|
||||
* GPU!
|
||||
*
|
||||
* Can be called per thread with no performance penalty.
|
||||
*
|
||||
* @param[in] GPU-side handle.
|
||||
*
|
||||
* @return void
|
||||
*/
|
||||
|
||||
} // namespace rocshmem
|
||||
|
||||
|
||||
Reference in new issue
Block a user