AI/rocm-systems
2
0
Fork 0
程式碼 問題 合併請求 Actions 套件 專案 版本發布 Wiki 動態

updated *_wait* APIs to use int rather than roc_shmem_cmps

瀏覽代碼
此提交包含在：
Yiltan Hassan Temucin
2024-10-10 14:13:44 -07:00
父節點 bcf3fdff10
當前提交 722a5f0731
共有 12 個檔案被更改，包括 152 行新增 和 152 行删除
下載 Patch 檔 下載差異檔 展開所有檔案 折疊所有檔案
include/roc_shmem/roc_shmem.hpp
+16 -16
查看文件
@@ -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
src/context.hpp
+16 -16
查看文件
@@ -74,42 +74,42 @@ class Context {
***************************** DEVICE METHODS *****************************
*************************************************************************/
template <typename T>
__device__ void wait_until(T *ivars, roc_shmem_cmps cmp, T val);
__device__ void wait_until(T *ivars, int cmp, T val);
template <typename T>
__device__ void wait_until_all(T *ivars, size_t nelems,
const int *status,
roc_shmem_cmps cmp, T val);
int cmp, T val);
template <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__host__ void wait_until(T *ivars, roc_shmem_cmps cmp, T val);
__host__ void wait_until(T *ivars, int cmp, T val);
template <typename T>
__host__ void wait_until_all(T *ivars, size_t nelems,
const int *status,
roc_shmem_cmps cmp, T val);
int cmp, T val);
template <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__host__ int test(T *ivars, roc_shmem_cmps cmp, T val);
__host__ int test(T *ivars, int cmp, T val);
public:
/**
src/context_tmpl_device.hpp
+8 -8
查看文件
@@ -199,7 +199,7 @@ __device__ void Context::broadcast(T *dest, const T *source, int nelems,
}
template <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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<T *>(ivars);
src/context_tmpl_host.hpp
+8 -8
查看文件
@@ -234,7 +234,7 @@ __host__ void Context::to_all(roc_shmem_team_t team, T *dest, const T *source,
}
template <typename T>
__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<T>(ivars, cmp, val));
@@ -243,7 +243,7 @@ __host__ void Context::wait_until(T *ivars, roc_shmem_cmps cmp, T val) {
template <typename T>
__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<T>(ivars, nelems, status, cmp, val));
@@ -252,7 +252,7 @@ __host__ size_t Context::wait_until_any(T *ivars, size_t nelems,
template <typename T>
__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<T>(ivars, nelems, status, cmp, val));
@@ -262,7 +262,7 @@ template <typename T>
__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<T>(ivars, nelems, indices, status, cmp, val));
@@ -271,7 +271,7 @@ __host__ size_t Context::wait_until_some(T *ivars, size_t nelems,
template <typename T>
__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<T>(ivars, nelems, status, cmp, vals));
@@ -280,7 +280,7 @@ __host__ void Context::wait_until_all_vector(T *ivars, size_t nelems,
template <typename T>
__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<T>(ivars, nelems, status, cmp, vals));
@@ -290,14 +290,14 @@ template <typename T>
__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<T>(ivars, nelems, indices, status, cmp, vals));
}
template <typename T>
__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<T>(ivars, cmp, val));
src/host/host.hpp
+10 -10
查看文件
@@ -211,41 +211,41 @@ class HostInterface {
int nreduce);
template <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T, ROC_SHMEM_OP Op>
__host__ void to_all_internal(MPI_Comm mpi_comm, T* dest, const T* source,
src/host/host_templates.hpp
+10 -10
查看文件
@@ -392,7 +392,7 @@ __host__ void HostInterface::to_all(roc_shmem_team_t team, T* dest,
}
template <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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");
src/ipc/context_ipc_host.hpp
+8 -8
查看文件
@@ -99,42 +99,42 @@ class IPCHostContext : public Context {
int nreduce);
template <typename T>
__host__ void wait_until(T *ivars, roc_shmem_cmps cmp, T val);
__host__ void wait_until(T *ivars, int cmp, T val);
template <typename T>
__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 <typename T>
__host__ void wait_until_all(T *ivars, size_t nelems,
const int *status,
roc_shmem_cmps cmp, T val);
int cmp, T val);
template <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 */
src/ipc/context_ipc_tmpl_host.hpp
+8 -8
查看文件
@@ -115,21 +115,21 @@ __host__ void IPCHostContext::to_all(roc_shmem_team_t team, T *dest,
}
template <typename T>
__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<T>(ivars, cmp, val, context_window_info);
}
template <typename T>
__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<T>(ivars, nelems, status, cmp, val, context_window_info);
}
template <typename T>
__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<T>(ivars, nelems, status, cmp, val, context_window_info);
}
@@ -137,21 +137,21 @@ template <typename T>
__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<T>(ivars, nelems, indices, status, cmp, val, context_window_info);
}
template <typename T>
__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<T>(ivars, nelems, status, cmp, vals, context_window_info);
}
template <typename T>
__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<T>(ivars, nelems, status, cmp, vals, context_window_info);
}
@@ -159,12 +159,12 @@ template <typename T>
__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<T>(ivars, nelems, indices, status, cmp, vals, context_window_info);
}
template <typename T>
__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<T>(ivars, cmp, val, context_window_info);
}
src/roc_shmem.cpp
+26 -26
查看文件
@@ -744,7 +744,7 @@ __host__ void roc_shmem_to_all([[maybe_unused]] roc_shmem_ctx_t ctx,
}
template <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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>(T *ivars, roc_shmem_cmps cmp, \
template __host__ void roc_shmem_wait_until<T>(T *ivars, int cmp, \
T val); \
template __host__ int roc_shmem_test<T>(T *ivars, roc_shmem_cmps cmp, T val);\
template __host__ void Context::wait_until<T>(T *ivars, roc_shmem_cmps cmp, \
template __host__ int roc_shmem_test<T>(T *ivars, int cmp, T val); \
template __host__ void Context::wait_until<T>(T *ivars, int cmp, \
T val); \
template __host__ size_t roc_shmem_wait_until_any<T>(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>(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>(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>(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>(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>(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>(T *ivars, roc_shmem_cmps cmp, T val);
template __host__ int Context::test<T>(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<T>(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<T>(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<T>(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<T>(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<T>(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<T>(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<T>(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<T>(ivars, cmp, val); \
}
src/roc_shmem_gpu.cpp
+32 -32
查看文件
@@ -488,7 +488,7 @@ __device__ void roc_shmem_wg_fcollect(roc_shmem_ctx_t ctx,
}
template <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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>(T *ivars, \
roc_shmem_cmps cmp, T val); \
int cmp, T val); \
template __device__ size_t roc_shmem_wait_until_any<T>(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>(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>(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>(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>(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>(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>(T *ivars, roc_shmem_cmps cmp, \
template __device__ int roc_shmem_test<T>(T *ivars, int cmp, \
T val); \
template __device__ void Context::wait_until<T>(T *ivars, roc_shmem_cmps cmp,\
template __device__ void Context::wait_until<T>(T *ivars, int cmp, \
T val); \
template __device__ size_t Context::wait_until_any<T>(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>(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>(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>(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>(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>(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>(T *ivars, roc_shmem_cmps cmp, T val);
template __device__ int Context::test<T>(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<T>(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<T>(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<T>(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<T>(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<T>(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<T>(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<T>(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<T>(ivars, cmp, val); \
}
src/templates.hpp
+2 -2
查看文件
@@ -388,7 +388,7 @@ __device__ void roc_shmem_atomic_set(T *dest, T value, int pe);
*
*/
template <typename T>
__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 <typename T>
__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
src/templates_host.hpp
+8 -8
查看文件
@@ -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 <typename T>
__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 <typename T>
__host__ void wait_until_all(T* ivars, size_t nelems, const int *status,
roc_shmem_cmps cmp, T val);
int cmp, T val);
template <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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 <typename T>
__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