diff --git a/include/roc_shmem/roc_shmem.hpp b/include/roc_shmem/roc_shmem.hpp index 7fa8e605a1..c06d18a6e3 100644 --- a/include/roc_shmem/roc_shmem.hpp +++ b/include/roc_shmem/roc_shmem.hpp @@ -1153,82 +1153,82 @@ __device__ ATTR_NO_INLINE void roc_shmem_threadfence_system(); */ #define WAIT_UNTIL_API_GEN(T, TNAME) \ __device__ void roc_shmem_##TNAME##_wait_until(T *ivars, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val); \ __device__ size_t roc_shmem_##TNAME##_wait_until_any(T *ivars, \ size_t nelems, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val); \ __device__ void roc_shmem_##TNAME##_wait_until_all(T *ivars, \ size_t nelems, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val); \ __device__ size_t roc_shmem_##TNAME##_wait_until_some(T *ivars, \ size_t nelems, \ size_t* indices, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val); \ __device__ size_t roc_shmem_##TNAME##_wait_until_any_vector(T *ivars, \ size_t nelems, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T* vals); \ __device__ void roc_shmem_##TNAME##_wait_until_all_vector(T *ivars, \ size_t nelems, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T* vals); \ __device__ size_t roc_shmem_##TNAME##_wait_until_some_vector(T *ivars, \ size_t nelems, \ size_t* indices, \ const int* status, \ - roc_shmem_cmps cmp,\ + int cmp, \ T* vals); \ __host__ void roc_shmem_##TNAME##_wait_until(T *ivars, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val); \ __host__ size_t roc_shmem_##TNAME##_wait_until_any(T *ivars, \ size_t nelems, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val); \ __host__ void roc_shmem_##TNAME##_wait_until_all(T *ivars, \ size_t nelems, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val); \ __host__ size_t roc_shmem_##TNAME##_wait_until_some(T *ivars, \ size_t nelems, \ size_t* indices, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val); \ __host__ size_t roc_shmem_##TNAME##_wait_until_any_vector(T *ivars, \ size_t nelems, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T* vals); \ __host__ void roc_shmem_##TNAME##_wait_until_all_vector(T *ivars, \ size_t nelems, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T* vals); \ __host__ size_t roc_shmem_##TNAME##_wait_until_some_vector(T *ivars, \ size_t nelems, \ size_t* indices, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T* vals); /* * MACRO DECLARE SHMEM_TEST APIs */ #define TEST_API_GEN(T, TNAME) \ - __device__ int roc_shmem_##TNAME##_test(T *ivars, roc_shmem_cmps cmp, T val); \ - __host__ int roc_shmem_##TNAME##_test(T *ivars, roc_shmem_cmps cmp, T val); + __device__ int roc_shmem_##TNAME##_test(T *ivars, int cmp, T val); \ + __host__ int roc_shmem_##TNAME##_test(T *ivars, int cmp, T val); /** * @name SHMEM_REDUCTIONS diff --git a/src/context.hpp b/src/context.hpp index d2cd759bff..0db4e99b23 100644 --- a/src/context.hpp +++ b/src/context.hpp @@ -74,42 +74,42 @@ class Context { ***************************** DEVICE METHODS ***************************** *************************************************************************/ template - __device__ void wait_until(T *ivars, roc_shmem_cmps cmp, T val); + __device__ void wait_until(T *ivars, int cmp, T val); template __device__ void wait_until_all(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T val); + int cmp, T val); template __device__ size_t wait_until_any(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T val); + int cmp, T val); template __device__ size_t wait_until_some(T *ivars, size_t nelems, size_t* indices, const int *status, - roc_shmem_cmps cmp, T val); + int cmp, T val); template __device__ void wait_until_all_vector(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals); + int cmp, T* vals); template __device__ size_t wait_until_any_vector(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals); + int cmp, T* vals); template __device__ size_t wait_until_some_vector(T *ivars, size_t nelems, size_t* indices, const int *status, - roc_shmem_cmps cmp, T* vals); + int cmp, T* vals); template - __device__ int test(T *ivars, roc_shmem_cmps cmp, T val); + __device__ int test(T *ivars, int cmp, T val); __device__ void threadfence_system(); @@ -365,42 +365,42 @@ class Context { int nreduce); template - __host__ void wait_until(T *ivars, roc_shmem_cmps cmp, T val); + __host__ void wait_until(T *ivars, int cmp, T val); template __host__ void wait_until_all(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T val); + int cmp, T val); template __host__ size_t wait_until_any(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T val); + int cmp, T val); template __host__ size_t wait_until_some(T *ivars, size_t nelems, size_t* indices, const int *status, - roc_shmem_cmps cmp, T val); + int cmp, T val); template __host__ void wait_until_all_vector(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals); + int cmp, T* vals); template __host__ size_t wait_until_any_vector(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals); + int cmp, T* vals); template __host__ size_t wait_until_some_vector(T *ivars, size_t nelems, size_t* indices, const int *status, - roc_shmem_cmps cmp, T* vals); + int cmp, T* vals); template - __host__ int test(T *ivars, roc_shmem_cmps cmp, T val); + __host__ int test(T *ivars, int cmp, T val); public: /** diff --git a/src/context_tmpl_device.hpp b/src/context_tmpl_device.hpp index 680a0eda3e..d0310d476a 100644 --- a/src/context_tmpl_device.hpp +++ b/src/context_tmpl_device.hpp @@ -199,7 +199,7 @@ __device__ void Context::broadcast(T *dest, const T *source, int nelems, } template -__device__ __forceinline__ void Context::wait_until(T *ivars, roc_shmem_cmps cmp, +__device__ __forceinline__ void Context::wait_until(T *ivars, int cmp, T val) { while (!test(ivars, cmp, val)) { } @@ -221,7 +221,7 @@ template __device__ __forceinline__ size_t Context::wait_until_any(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T val) { + int cmp, T val) { // zero nelems error condition if (!nelems) { return SIZE_MAX; @@ -251,7 +251,7 @@ template __device__ __forceinline__ void Context::wait_until_all(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T val) { + int cmp, T val) { // zero nelems error condition if (!nelems) { return; @@ -278,7 +278,7 @@ __device__ __forceinline__ size_t Context::wait_until_some(T *ivars, size_t nelems, size_t* indices, const int *status, - roc_shmem_cmps cmp, T val) { + int cmp, T val) { // zero nelems error condition if (!nelems) { return 0; @@ -313,7 +313,7 @@ template __device__ __forceinline__ void Context::wait_until_all_vector(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { ; } @@ -321,7 +321,7 @@ template __device__ __forceinline__ size_t Context::wait_until_any_vector(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { return 0; } @@ -330,12 +330,12 @@ __device__ __forceinline__ size_t Context::wait_until_some_vector(T *ivars, size_t nelems, size_t* indices, const int *status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { return 0; } template -__device__ __forceinline__ int Context::test(T *ivars, roc_shmem_cmps cmp, +__device__ __forceinline__ int Context::test(T *ivars, int cmp, T val) { int ret = 0; volatile T *vol_ivars = reinterpret_cast(ivars); diff --git a/src/context_tmpl_host.hpp b/src/context_tmpl_host.hpp index 12c225fd0f..a4c7929501 100644 --- a/src/context_tmpl_host.hpp +++ b/src/context_tmpl_host.hpp @@ -234,7 +234,7 @@ __host__ void Context::to_all(roc_shmem_team_t team, T *dest, const T *source, } template -__host__ void Context::wait_until(T *ivars, roc_shmem_cmps cmp, T val) { +__host__ void Context::wait_until(T *ivars, int cmp, T val) { ctxHostStats.incStat(NUM_HOST_WAIT_UNTIL); HOST_DISPATCH(wait_until(ivars, cmp, val)); @@ -243,7 +243,7 @@ __host__ void Context::wait_until(T *ivars, roc_shmem_cmps cmp, T val) { template __host__ size_t Context::wait_until_any(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T val) { + int cmp, T val) { ctxHostStats.incStat(NUM_HOST_WAIT_UNTIL_ANY); return HOST_DISPATCH(wait_until_any(ivars, nelems, status, cmp, val)); @@ -252,7 +252,7 @@ __host__ size_t Context::wait_until_any(T *ivars, size_t nelems, template __host__ void Context::wait_until_all(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T val) { + int cmp, T val) { ctxHostStats.incStat(NUM_HOST_WAIT_UNTIL_ALL); HOST_DISPATCH(wait_until_all(ivars, nelems, status, cmp, val)); @@ -262,7 +262,7 @@ template __host__ size_t Context::wait_until_some(T *ivars, size_t nelems, size_t* indices, const int* status, - roc_shmem_cmps cmp, T val) { + int cmp, T val) { ctxHostStats.incStat(NUM_HOST_WAIT_UNTIL_SOME); HOST_DISPATCH_RET(wait_until_some(ivars, nelems, indices, status, cmp, val)); @@ -271,7 +271,7 @@ __host__ size_t Context::wait_until_some(T *ivars, size_t nelems, template __host__ void Context::wait_until_all_vector(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { ctxHostStats.incStat(NUM_HOST_WAIT_UNTIL_ALL_VECTOR); HOST_DISPATCH(wait_until_all_vector(ivars, nelems, status, cmp, vals)); @@ -280,7 +280,7 @@ __host__ void Context::wait_until_all_vector(T *ivars, size_t nelems, template __host__ size_t Context::wait_until_any_vector(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { ctxHostStats.incStat(NUM_HOST_WAIT_UNTIL_ANY_VECTOR); HOST_DISPATCH_RET(wait_until_any_vector(ivars, nelems, status, cmp, vals)); @@ -290,14 +290,14 @@ template __host__ size_t Context::wait_until_some_vector(T *ivars, size_t nelems, size_t* indices, const int *status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { ctxHostStats.incStat(NUM_HOST_WAIT_UNTIL_SOME_VECTOR); HOST_DISPATCH_RET(wait_until_some_vector(ivars, nelems, indices, status, cmp, vals)); } template -__host__ int Context::test(T *ivars, roc_shmem_cmps cmp, T val) { +__host__ int Context::test(T *ivars, int cmp, T val) { ctxHostStats.incStat(NUM_HOST_TEST); HOST_DISPATCH_RET(test(ivars, cmp, val)); diff --git a/src/host/host.hpp b/src/host/host.hpp index b155b87a3e..4cd2b2e393 100644 --- a/src/host/host.hpp +++ b/src/host/host.hpp @@ -211,41 +211,41 @@ class HostInterface { int nreduce); template - __host__ void wait_until(T *ivars, roc_shmem_cmps cmp, T val, + __host__ void wait_until(T *ivars, int cmp, T val, WindowInfo* window_info); template __host__ void wait_until_all(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T val, + int cmp, T val, WindowInfo* window_info); template __host__ size_t wait_until_any(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T val, + int cmp, T val, WindowInfo* window_info); template __host__ size_t wait_until_some(T *ivars, size_t nelems, size_t* indices, - const int* status, roc_shmem_cmps cmp, T val, + const int* status, int cmp, T val, WindowInfo* window_info); template __host__ void wait_until_all_vector(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T* vals, + int cmp, T* vals, WindowInfo* window_info); template __host__ size_t wait_until_any_vector(T *ivars, size_t nelems, - const int* status, roc_shmem_cmps cmp, + const int* status, int cmp, T* vals, WindowInfo* window_info); template __host__ size_t wait_until_some_vector(T *ivars, size_t nelems, size_t* indices, - const int* status, roc_shmem_cmps cmp, + const int* status, int cmp, T* vals, WindowInfo* window_info); template - __host__ int test(T *ivars, roc_shmem_cmps cmp, T val, WindowInfo* window_info); + __host__ int test(T *ivars, int cmp, T val, WindowInfo* window_info); #ifndef USE_COHERENT_HEAP __host__ void create_hdp_window(); @@ -295,11 +295,11 @@ class HostInterface { __host__ MPI_Datatype get_mpi_type(); template - __host__ int compare(roc_shmem_cmps cmp, T input_val, T target_val); + __host__ int compare(int cmp, T input_val, T target_val); template __host__ int test_and_compare(MPI_Aint offset, MPI_Datatype mpi_type, - roc_shmem_cmps cmp, T val, MPI_Win win); + int cmp, T val, MPI_Win win); template __host__ void to_all_internal(MPI_Comm mpi_comm, T* dest, const T* source, diff --git a/src/host/host_templates.hpp b/src/host/host_templates.hpp index 3a84862a00..41ff7cb094 100644 --- a/src/host/host_templates.hpp +++ b/src/host/host_templates.hpp @@ -392,7 +392,7 @@ __host__ void HostInterface::to_all(roc_shmem_team_t team, T* dest, } template -__host__ inline int HostInterface::compare(roc_shmem_cmps cmp, T input_val, +__host__ inline int HostInterface::compare(int cmp, T input_val, T target_val) { int cond_satisfied{0}; @@ -426,7 +426,7 @@ __host__ inline int HostInterface::compare(roc_shmem_cmps cmp, T input_val, template __host__ inline int HostInterface::test_and_compare(MPI_Aint offset, MPI_Datatype mpi_type, - roc_shmem_cmps cmp, T val, + int cmp, T val, MPI_Win win) { T fetched_val{}; @@ -446,7 +446,7 @@ __host__ inline int HostInterface::test_and_compare(MPI_Aint offset, } template -__host__ void HostInterface::wait_until(T *ivars, roc_shmem_cmps cmp, T val, +__host__ void HostInterface::wait_until(T *ivars, int cmp, T val, WindowInfo* window_info) { DPRINTF("Function: host_wait_until\n"); @@ -502,7 +502,7 @@ __host__ size_t status_entry(size_t nelems, template __host__ size_t HostInterface::wait_until_any(T* ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T val, + int cmp, T val, WindowInfo* window_info) { DPRINTF("Function: host_wait_until_any\n"); @@ -534,7 +534,7 @@ __host__ size_t HostInterface::wait_until_any(T* ivars, size_t nelems, template __host__ void HostInterface::wait_until_all(T* ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T val, + int cmp, T val, WindowInfo* window_info) { DPRINTF("Function: host_wait_until_all\n"); @@ -563,7 +563,7 @@ template __host__ size_t HostInterface::wait_until_some(T* ivars, size_t nelems, size_t* indices, const int *status, - roc_shmem_cmps cmp, T val, + int cmp, T val, WindowInfo* window_info) { DPRINTF("Function: host_wait_until_some\n"); @@ -600,7 +600,7 @@ __host__ size_t HostInterface::wait_until_some(T* ivars, size_t nelems, template __host__ void HostInterface::wait_until_all_vector(T* ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals, + int cmp, T* vals, WindowInfo* window_info) { DPRINTF("Function: host_wait_until_all_vector\n"); } @@ -608,7 +608,7 @@ __host__ void HostInterface::wait_until_all_vector(T* ivars, size_t nelems, template __host__ size_t HostInterface::wait_until_any_vector(T* ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals, + int cmp, T* vals, WindowInfo* window_info) { DPRINTF("Function: host_wait_until_any_vector\n"); return 0; @@ -618,14 +618,14 @@ template __host__ size_t HostInterface::wait_until_some_vector(T* ivars, size_t nelems, size_t* indices, const int *status, - roc_shmem_cmps cmp, T* vals, + int cmp, T* vals, WindowInfo* window_info) { DPRINTF("Function: host_wait_until_some_vector\n"); return 0; } template -__host__ int HostInterface::test(T* ivars, roc_shmem_cmps cmp, T val, +__host__ int HostInterface::test(T* ivars, int cmp, T val, WindowInfo* window_info) { DPRINTF("Function: host_test\n"); diff --git a/src/ipc/context_ipc_host.hpp b/src/ipc/context_ipc_host.hpp index 0089ca8f75..4c2e626d75 100644 --- a/src/ipc/context_ipc_host.hpp +++ b/src/ipc/context_ipc_host.hpp @@ -99,42 +99,42 @@ class IPCHostContext : public Context { int nreduce); template - __host__ void wait_until(T *ivars, roc_shmem_cmps cmp, T val); + __host__ void wait_until(T *ivars, int cmp, T val); template __host__ size_t wait_until_any(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T val); + int cmp, T val); template __host__ void wait_until_all(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T val); + int cmp, T val); template __host__ size_t wait_until_some(T *ivars, size_t nelems, size_t* indices, const int *status, - roc_shmem_cmps cmp, T val); + int cmp, T val); template __host__ void wait_until_all_vector(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals); + int cmp, T* vals); template __host__ size_t wait_until_any_vector(T *ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals); + int cmp, T* vals); template __host__ size_t wait_until_some_vector(T *ivars, size_t nelems, size_t* indices, const int *status, - roc_shmem_cmps cmp, T* vals); + int cmp, T* vals); template - __host__ int test(T *ivars, roc_shmem_cmps cmp, T val); + __host__ int test(T *ivars, int cmp, T val); public: /* Pointer to the backend's host interface */ diff --git a/src/ipc/context_ipc_tmpl_host.hpp b/src/ipc/context_ipc_tmpl_host.hpp index 1c74c80fdf..d78593b888 100644 --- a/src/ipc/context_ipc_tmpl_host.hpp +++ b/src/ipc/context_ipc_tmpl_host.hpp @@ -115,21 +115,21 @@ __host__ void IPCHostContext::to_all(roc_shmem_team_t team, T *dest, } template -__host__ void IPCHostContext::wait_until(T *ivars, roc_shmem_cmps cmp, T val) { +__host__ void IPCHostContext::wait_until(T *ivars, int cmp, T val) { host_interface->wait_until(ivars, cmp, val, context_window_info); } template __host__ void IPCHostContext::wait_until_all(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T val) { + int cmp, T val) { host_interface->wait_until_all(ivars, nelems, status, cmp, val, context_window_info); } template __host__ size_t IPCHostContext::wait_until_any(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T val) { + int cmp, T val) { return host_interface->wait_until_any(ivars, nelems, status, cmp, val, context_window_info); } @@ -137,21 +137,21 @@ template __host__ size_t IPCHostContext::wait_until_some(T *ivars, size_t nelems, size_t* indices, const int* status, - roc_shmem_cmps cmp, T val) { + int cmp, T val) { return host_interface->wait_until_some(ivars, nelems, indices, status, cmp, val, context_window_info); } template __host__ void IPCHostContext::wait_until_all_vector(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { host_interface->wait_until_all_vector(ivars, nelems, status, cmp, vals, context_window_info); } template __host__ size_t IPCHostContext::wait_until_any_vector(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { return host_interface->wait_until_any_vector(ivars, nelems, status, cmp, vals, context_window_info); } @@ -159,12 +159,12 @@ template __host__ size_t IPCHostContext::wait_until_some_vector(T *ivars, size_t nelems, size_t* indices, const int* status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { return host_interface->wait_until_some_vector(ivars, nelems, indices, status, cmp, vals, context_window_info); } template -__host__ int IPCHostContext::test(T *ivars, roc_shmem_cmps cmp, T val) { +__host__ int IPCHostContext::test(T *ivars, int cmp, T val) { return host_interface->test(ivars, cmp, val, context_window_info); } diff --git a/src/roc_shmem.cpp b/src/roc_shmem.cpp index bcef7c0db4..23e2b25a7b 100644 --- a/src/roc_shmem.cpp +++ b/src/roc_shmem.cpp @@ -744,7 +744,7 @@ __host__ void roc_shmem_to_all([[maybe_unused]] roc_shmem_ctx_t ctx, } template -__host__ void roc_shmem_wait_until(T *ivars, roc_shmem_cmps cmp, T val) { +__host__ void roc_shmem_wait_until(T *ivars, int cmp, T val) { DPRINTF("Host function: roc_shmem_wait_until\n"); get_internal_ctx(ROC_SHMEM_HOST_CTX_DEFAULT)->wait_until(ivars, cmp, val); @@ -752,7 +752,7 @@ __host__ void roc_shmem_wait_until(T *ivars, roc_shmem_cmps cmp, T val) { template __host__ void roc_shmem_wait_until_all(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T val) { + int cmp, T val) { DPRINTF("Host function: roc_shmem_wait_until_all\n"); get_internal_ctx(ROC_SHMEM_HOST_CTX_DEFAULT)->wait_until_all(ivars, @@ -761,7 +761,7 @@ __host__ void roc_shmem_wait_until_all(T *ivars, size_t nelems, const int* statu template __host__ size_t roc_shmem_wait_until_any(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T val) { + int cmp, T val) { DPRINTF("Host function: roc_shmem_wait_until_any\n"); return get_internal_ctx(ROC_SHMEM_HOST_CTX_DEFAULT)->wait_until_any(ivars, @@ -770,7 +770,7 @@ __host__ size_t roc_shmem_wait_until_any(T *ivars, size_t nelems, const int* sta template __host__ size_t roc_shmem_wait_until_some(T *ivars, size_t nelems, size_t* indices, - const int* status, roc_shmem_cmps cmp, + const int* status, int cmp, T val) { DPRINTF("Host function: roc_shmem_wait_until_some\n"); @@ -780,7 +780,7 @@ __host__ size_t roc_shmem_wait_until_some(T *ivars, size_t nelems, size_t* indic template __host__ size_t roc_shmem_wait_until_any_vector(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { DPRINTF("Host function: roc_shmem_wait_until_any_vector\n"); return get_internal_ctx(ROC_SHMEM_HOST_CTX_DEFAULT)->wait_until_any_vector(ivars, @@ -789,7 +789,7 @@ __host__ size_t roc_shmem_wait_until_any_vector(T *ivars, size_t nelems, const i template __host__ void roc_shmem_wait_until_all_vector(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { DPRINTF("Host function: roc_shmem_wait_until_all_vector\n"); get_internal_ctx(ROC_SHMEM_HOST_CTX_DEFAULT)->wait_until_all_vector(ivars, @@ -800,7 +800,7 @@ template __host__ size_t roc_shmem_wait_until_some_vector(T *ivars, size_t nelems, size_t* indices, const int* status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { DPRINTF("Host function: roc_shmem_wait_until_some_vector\n"); return get_internal_ctx(ROC_SHMEM_HOST_CTX_DEFAULT)->wait_until_some_vector(ivars, @@ -808,7 +808,7 @@ __host__ size_t roc_shmem_wait_until_some_vector(T *ivars, size_t nelems, } template -__host__ int roc_shmem_test(T *ivars, roc_shmem_cmps cmp, T val) { +__host__ int roc_shmem_test(T *ivars, int cmp, T val) { DPRINTF("Host function: roc_shmem_testl\n"); return get_internal_ctx(ROC_SHMEM_HOST_CTX_DEFAULT)->test(ivars, cmp, val); @@ -940,31 +940,31 @@ __host__ int roc_shmem_test(T *ivars, roc_shmem_cmps cmp, T val) { * Declare templates for the wait types */ #define WAIT_GEN(T) \ - template __host__ void roc_shmem_wait_until(T *ivars, roc_shmem_cmps cmp, \ + template __host__ void roc_shmem_wait_until(T *ivars, int cmp, \ T val); \ - template __host__ int roc_shmem_test(T *ivars, roc_shmem_cmps cmp, T val);\ - template __host__ void Context::wait_until(T *ivars, roc_shmem_cmps cmp, \ + template __host__ int roc_shmem_test(T *ivars, int cmp, T val); \ + template __host__ void Context::wait_until(T *ivars, int cmp, \ T val); \ template __host__ size_t roc_shmem_wait_until_any(T *ivars, \ size_t nelems, const int* status, \ - roc_shmem_cmps cmp, T val); \ + int cmp, T val); \ template __host__ void roc_shmem_wait_until_all(T *ivars, \ size_t nelems, const int* status, \ - roc_shmem_cmps cmp, T val); \ + int cmp, T val); \ template __host__ size_t roc_shmem_wait_until_some(T *ivars, size_t nelems,\ size_t* indices, const int* status, \ - roc_shmem_cmps cmp, T val); \ + int cmp, T val); \ template __host__ size_t roc_shmem_wait_until_any_vector(T *ivars, \ size_t nelems, const int* status, \ - roc_shmem_cmps cmp, T* vals); \ + int cmp, T* vals); \ template __host__ void roc_shmem_wait_until_all_vector(T *ivars, \ size_t nelems, const int* status, \ - roc_shmem_cmps cmp, T* vals); \ + int cmp, T* vals); \ template __host__ size_t roc_shmem_wait_until_some_vector(T *ivars, \ size_t nelems, size_t* indices, \ - const int* status, roc_shmem_cmps cmp, \ + const int* status, int cmp, \ T* vals); \ - template __host__ int Context::test(T *ivars, roc_shmem_cmps cmp, T val); + template __host__ int Context::test(T *ivars, int cmp, T val); /** * Define APIs to call the template functions @@ -1165,33 +1165,33 @@ __host__ int roc_shmem_test(T *ivars, roc_shmem_cmps cmp, T val) { } #define WAIT_DEF_GEN(T, TNAME) \ - __host__ void roc_shmem_##TNAME##_wait_until(T *ivars, roc_shmem_cmps cmp, \ + __host__ void roc_shmem_##TNAME##_wait_until(T *ivars, int cmp, \ T val) { \ roc_shmem_wait_until(ivars, cmp, val); \ } \ __host__ size_t roc_shmem_##TNAME##_wait_until_any(T *ivars, size_t nelems,\ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val) { \ return roc_shmem_wait_until_any(ivars, nelems, status, cmp, val); \ } \ __host__ void roc_shmem_##TNAME##_wait_until_all(T *ivars, size_t nelems, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val) { \ roc_shmem_wait_until_all(ivars, nelems, status, cmp, val); \ } \ __host__ size_t roc_shmem_##TNAME##_wait_until_some(T *ivars, size_t nelems, \ size_t* indices, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val) { \ return roc_shmem_wait_until_some(ivars, nelems, indices, status, cmp, val); \ } \ __host__ size_t roc_shmem_##TNAME##_wait_until_any_vector(T *ivars, \ size_t nelems, \ const int* status, \ - roc_shmem_cmps cmp,\ + int cmp, \ T* vals) { \ return roc_shmem_wait_until_any_vector(ivars, nelems, status, cmp, \ vals); \ @@ -1199,7 +1199,7 @@ __host__ int roc_shmem_test(T *ivars, roc_shmem_cmps cmp, T val) { __host__ void roc_shmem_##TNAME##_wait_until_all_vector(T *ivars, \ size_t nelems, \ const int* status, \ - roc_shmem_cmps cmp,\ + int cmp, \ T* vals) { \ roc_shmem_wait_until_all_vector(ivars, nelems, status, cmp, vals); \ } \ @@ -1207,12 +1207,12 @@ __host__ int roc_shmem_test(T *ivars, roc_shmem_cmps cmp, T val) { size_t nelems, \ size_t* indices, \ const int* status,\ - roc_shmem_cmps cmp,\ + int cmp, \ T* vals) { \ return roc_shmem_wait_until_some_vector(ivars, nelems, indices, \ status, cmp, vals); \ } \ - __host__ int roc_shmem_##TNAME##_test(T *ivars, roc_shmem_cmps cmp, T val) { \ + __host__ int roc_shmem_##TNAME##_test(T *ivars, int cmp, T val) { \ return roc_shmem_test(ivars, cmp, val); \ } diff --git a/src/roc_shmem_gpu.cpp b/src/roc_shmem_gpu.cpp index 9cc723c221..9a3a79f6e5 100644 --- a/src/roc_shmem_gpu.cpp +++ b/src/roc_shmem_gpu.cpp @@ -488,7 +488,7 @@ __device__ void roc_shmem_wg_fcollect(roc_shmem_ctx_t ctx, } template -__device__ void roc_shmem_wait_until(T *ivars, roc_shmem_cmps cmp, T val) { +__device__ void roc_shmem_wait_until(T *ivars, int cmp, T val) { GPU_DPRINTF("Function: roc_shmem_wait_until\n"); Context *ctx_internal = get_internal_ctx(ROC_SHMEM_CTX_DEFAULT); @@ -498,7 +498,7 @@ __device__ void roc_shmem_wait_until(T *ivars, roc_shmem_cmps cmp, T val) { template __device__ void roc_shmem_wait_until_all(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T val) { + int cmp, T val) { GPU_DPRINTF("Function: roc_shmem_wait_until_all\n"); Context *ctx_internal = get_internal_ctx(ROC_SHMEM_CTX_DEFAULT); @@ -508,7 +508,7 @@ __device__ void roc_shmem_wait_until_all(T *ivars, size_t nelems, const int* sta template __device__ size_t roc_shmem_wait_until_any(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T val) { + int cmp, T val) { GPU_DPRINTF("Function: roc_shmem_wait_until_any\n"); Context *ctx_internal = get_internal_ctx(ROC_SHMEM_CTX_DEFAULT); @@ -518,7 +518,7 @@ __device__ size_t roc_shmem_wait_until_any(T *ivars, size_t nelems, const int* s template __device__ size_t roc_shmem_wait_until_some(T *ivars, size_t nelems, size_t* indices, - const int* status, roc_shmem_cmps cmp, + const int* status, int cmp, T val) { DPRINTF("Function: roc_shmem_wait_until_some\n"); @@ -529,7 +529,7 @@ __device__ size_t roc_shmem_wait_until_some(T *ivars, size_t nelems, size_t* ind template __device__ size_t roc_shmem_wait_until_any_vector(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { DPRINTF("Function: roc_shmem_wait_until_any_vector\n"); Context *ctx_internal = get_internal_ctx(ROC_SHMEM_CTX_DEFAULT); @@ -539,7 +539,7 @@ __device__ size_t roc_shmem_wait_until_any_vector(T *ivars, size_t nelems, const template __device__ void roc_shmem_wait_until_all_vector(T *ivars, size_t nelems, const int* status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { DPRINTF("Function: roc_shmem_wait_until_all_vector\n"); Context *ctx_internal = get_internal_ctx(ROC_SHMEM_CTX_DEFAULT); @@ -551,7 +551,7 @@ template __device__ size_t roc_shmem_wait_until_some_vector(T *ivars, size_t nelems, size_t* indices, const int* status, - roc_shmem_cmps cmp, T* vals) { + int cmp, T* vals) { DPRINTF("Function: roc_shmem_wait_until_some_vector\n"); Context *ctx_internal = get_internal_ctx(ROC_SHMEM_CTX_DEFAULT); @@ -560,7 +560,7 @@ __device__ size_t roc_shmem_wait_until_some_vector(T *ivars, size_t nelems, } template -__device__ int roc_shmem_test(T *ivars, roc_shmem_cmps cmp, T val) { +__device__ int roc_shmem_test(T *ivars, int cmp, T val) { GPU_DPRINTF("Function: roc_shmem_testl\n"); Context *ctx_internal = get_internal_ctx(ROC_SHMEM_CTX_DEFAULT); @@ -1031,52 +1031,52 @@ __device__ int roc_shmem_team_translate_pe(roc_shmem_team_t src_team, */ #define WAIT_GEN(T) \ template __device__ void roc_shmem_wait_until(T *ivars, \ - roc_shmem_cmps cmp, T val); \ + int cmp, T val); \ template __device__ size_t roc_shmem_wait_until_any(T *ivars, \ size_t nelems, const int* status, \ - roc_shmem_cmps cmp, T val); \ + int cmp, T val); \ template __device__ void roc_shmem_wait_until_all(T *ivars, \ size_t nelems, const int* status, \ - roc_shmem_cmps cmp, T val); \ + int cmp, T val); \ template __device__ size_t roc_shmem_wait_until_some(T *ivars, \ size_t nelems, size_t* indices, \ const int* status, \ - roc_shmem_cmps cmp, T val); \ + int cmp, T val); \ template __device__ size_t roc_shmem_wait_until_any_vector(T *ivars, \ size_t nelems, const int* status, \ - roc_shmem_cmps cmp, T* vals); \ + int cmp, T* vals); \ template __device__ void roc_shmem_wait_until_all_vector(T *ivars, \ size_t nelems, const int* status, \ - roc_shmem_cmps cmp, T* vals); \ + int cmp, T* vals); \ template __device__ size_t roc_shmem_wait_until_some_vector(T *ivars, \ size_t nelems, size_t* indices, \ - const int* status, roc_shmem_cmps cmp, \ + const int* status, int cmp, \ T* vals); \ - template __device__ int roc_shmem_test(T *ivars, roc_shmem_cmps cmp, \ + template __device__ int roc_shmem_test(T *ivars, int cmp, \ T val); \ - template __device__ void Context::wait_until(T *ivars, roc_shmem_cmps cmp,\ + template __device__ void Context::wait_until(T *ivars, int cmp, \ T val); \ template __device__ size_t Context::wait_until_any(T *ivars, \ size_t nelems, const int* status, \ - roc_shmem_cmps cmp, T val); \ + int cmp, T val); \ template __device__ void Context::wait_until_all(T *ivars, \ size_t nelems, const int* status, \ - roc_shmem_cmps cmp, T val); \ + int cmp, T val); \ template __device__ size_t Context::wait_until_some(T *ivars, \ size_t nelems, \ size_t* indices, const int* status, \ - roc_shmem_cmps cmp, T val); \ + int cmp, T val); \ template __device__ size_t Context::wait_until_any_vector(T *ivars, \ size_t nelems, const int* status, \ - roc_shmem_cmps cmp, T* vals); \ + int cmp, T* vals); \ template __device__ void Context::wait_until_all_vector(T *ivars, \ size_t nelems, const int* status, \ - roc_shmem_cmps cmp, T* vals); \ + int cmp, T* vals); \ template __device__ size_t Context::wait_until_some_vector(T *ivars, \ size_t nelems, size_t* indices, \ - const int* status, roc_shmem_cmps cmp, \ + const int* status, int cmp, \ T* vals); \ - template __device__ int Context::test(T *ivars, roc_shmem_cmps cmp, T val); + template __device__ int Context::test(T *ivars, int cmp, T val); #define ARITH_REDUCTION_GEN(T) \ REDUCTION_GEN(T, ROC_SHMEM_SUM) \ @@ -1371,19 +1371,19 @@ __device__ int roc_shmem_team_translate_pe(roc_shmem_team_t src_team, } #define WAIT_DEF_GEN(T, TNAME) \ - __device__ void roc_shmem_##TNAME##_wait_until(T *ivars, roc_shmem_cmps cmp, \ + __device__ void roc_shmem_##TNAME##_wait_until(T *ivars, int cmp, \ T val) { \ roc_shmem_wait_until(ivars, cmp, val); \ } \ __device__ size_t roc_shmem_##TNAME##_wait_until_any(T *ivars, size_t nelems,\ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val) { \ return roc_shmem_wait_until_any(ivars, nelems, status, cmp, val); \ } \ __device__ void roc_shmem_##TNAME##_wait_until_all(T *ivars, size_t nelems,\ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val) { \ roc_shmem_wait_until_all(ivars, nelems, status, cmp, val); \ } \ @@ -1391,7 +1391,7 @@ __device__ int roc_shmem_team_translate_pe(roc_shmem_team_t src_team, size_t nelems, \ size_t* indices, \ const int* status, \ - roc_shmem_cmps cmp, \ + int cmp, \ T val) { \ return roc_shmem_wait_until_some(ivars, nelems, indices, status, cmp, \ val); \ @@ -1399,7 +1399,7 @@ __device__ int roc_shmem_team_translate_pe(roc_shmem_team_t src_team, __device__ size_t roc_shmem_##TNAME##_wait_until_any_vector(T *ivars, \ size_t nelems, \ const int* status, \ - roc_shmem_cmps cmp,\ + int cmp, \ T* vals) { \ return roc_shmem_wait_until_any_vector(ivars, nelems, status, cmp, \ vals); \ @@ -1407,7 +1407,7 @@ __device__ int roc_shmem_team_translate_pe(roc_shmem_team_t src_team, __device__ void roc_shmem_##TNAME##_wait_until_all_vector(T *ivars, \ size_t nelems, \ const int* status, \ - roc_shmem_cmps cmp,\ + int cmp, \ T* vals) { \ roc_shmem_wait_until_all_vector(ivars, nelems, status, cmp, vals); \ } \ @@ -1415,12 +1415,12 @@ __device__ int roc_shmem_team_translate_pe(roc_shmem_team_t src_team, size_t nelems, \ size_t* indices, \ const int* status,\ - roc_shmem_cmps cmp,\ + int cmp, \ T* vals) { \ return roc_shmem_wait_until_some_vector(ivars, nelems, indices, \ status, cmp, vals); \ } \ - __device__ int roc_shmem_##TNAME##_test(T *ivars, roc_shmem_cmps cmp, T val) { \ + __device__ int roc_shmem_##TNAME##_test(T *ivars, int cmp, T val) { \ return roc_shmem_test(ivars, cmp, val); \ } diff --git a/src/templates.hpp b/src/templates.hpp index 1b5e82d191..bba411151b 100644 --- a/src/templates.hpp +++ b/src/templates.hpp @@ -388,7 +388,7 @@ __device__ void roc_shmem_atomic_set(T *dest, T value, int pe); * */ template -__device__ void roc_shmem_wait_until(T *ivars, roc_shmem_cmps cmp, T val); +__device__ void roc_shmem_wait_until(T *ivars, int cmp, T val); /** * @brief test if the condition (* \p ptr \p cmps \p val) is @@ -407,7 +407,7 @@ __device__ void roc_shmem_wait_until(T *ivars, roc_shmem_cmps cmp, T val); * */ template -__device__ int roc_shmem_test(T *ivars, roc_shmem_cmps cmp, T val); +__device__ int roc_shmem_test(T *ivars, int cmp, T val); /** * @brief Perform a broadcast between PEs in the active set. The caller diff --git a/src/templates_host.hpp b/src/templates_host.hpp index ea4818a25a..97c92e87da 100644 --- a/src/templates_host.hpp +++ b/src/templates_host.hpp @@ -136,35 +136,35 @@ __host__ void roc_shmem_to_all(roc_shmem_ctx_t ctx, T *dest, const T *source, int PE_size, T *pWrk, long *pSync); template -__host__ void roc_shmem_wait_until(T *ivars, roc_shmem_cmps cmp, T val); +__host__ void roc_shmem_wait_until(T *ivars, int cmp, T val); template __host__ void wait_until_all(T* ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T val); + int cmp, T val); template __host__ size_t wait_until_any(T* ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T val); + int cmp, T val); template __host__ size_t wait_until_some(T* ivars, size_t nelems, size_t* indices, - const int *status, roc_shmem_cmps cmp, T val); + const int *status, int cmp, T val); template __host__ void wait_until_all_vector(T* ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals); + int cmp, T* vals); template __host__ size_t wait_until_any_vector(T* ivars, size_t nelems, const int *status, - roc_shmem_cmps cmp, T* vals); + int cmp, T* vals); template __host__ size_t wait_until_some_vector(T* ivars, size_t nelems, size_t* indices, const int *status, - roc_shmem_cmps cmp, T* vals); + int cmp, T* vals); template -__host__ int roc_shmem_test(T *ivars, roc_shmem_cmps cmp, T val); +__host__ int roc_shmem_test(T *ivars, int cmp, T val); } // namespace rocshmem