From e2e30b53395cdc605d4d4d99dc983f15cf8fcebf Mon Sep 17 00:00:00 2001 From: Edgar Gabriel Date: Tue, 27 Aug 2024 06:25:58 -0700 Subject: [PATCH] remove device wait_until functions adding the device versions of the wait_until* and test functions in the ipc folder leads to linking errors of the functional tests. Remove them and use for now the upper level versions of the functions, similarly to the RO conduit. Might have to revisit this later again. --- src/ipc/context_ipc_device.hpp | 38 --------------------- src/ipc/context_ipc_tmpl_device.hpp | 52 ----------------------------- 2 files changed, 90 deletions(-) diff --git a/src/ipc/context_ipc_device.hpp b/src/ipc/context_ipc_device.hpp index afadb9828c..a8cefb446a 100644 --- a/src/ipc/context_ipc_device.hpp +++ b/src/ipc/context_ipc_device.hpp @@ -194,44 +194,6 @@ class IPCContext : public Context { template __device__ void get_nbi_wave(T *dest, const T *source, size_t nelems, int pe); - // Wait / Test functions - template - __device__ void wait_until(T* ptr, roc_shmem_cmps cmp, T val); - - template - __device__ void wait_until_all(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T val); - - template - __device__ size_t wait_until_any(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T val); - - template - __device__ size_t wait_until_some(T* ptr, size_t nelems, - size_t* indices, - const int *status, - roc_shmem_cmps cmp, T val); - - template - __device__ void wait_until_all_vector(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T* vals); - - template - __device__ size_t wait_until_any_vector(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T* vals); - template - __device__ size_t wait_until_some_vector(T* ptr, size_t nelems, - size_t* indices, - const int *status, - roc_shmem_cmps cmp, T* vals); - - template - __device__ int test(T* ptr, roc_shmem_cmps cmp, T val); - private: //context class has IpcImpl object (ipcImpl_) diff --git a/src/ipc/context_ipc_tmpl_device.hpp b/src/ipc/context_ipc_tmpl_device.hpp index 28cd718f0a..4bf652781e 100644 --- a/src/ipc/context_ipc_tmpl_device.hpp +++ b/src/ipc/context_ipc_tmpl_device.hpp @@ -223,58 +223,6 @@ __device__ void IPCContext::get_nbi_wave(T *dest, const T *source, getmem_nbi_wave(dest, source, nelems * sizeof(T), pe); } - -//Wait/test functions -template -__device__ void wait_until(T* ptr, roc_shmem_cmps cmp, T val) { -} - -template -__device__ void wait_until_all(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T val) { -} - -template -__device__ size_t wait_until_any(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T val) { - return 0; -} - -template -__device__ size_t wait_until_some(T* ptr, size_t nelems, - size_t* indices, - const int *status, - roc_shmem_cmps cmp, T val){ - return 0; -} - -template -__device__ void wait_until_all_vector(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T* vals) { -} - -template -__device__ size_t wait_until_any_vector(T* ptr, size_t nelems, - const int *status, - roc_shmem_cmps cmp, T* vals){ - return 0; -} - -template -__device__ size_t wait_until_some_vector(T* ptr, size_t nelems, - size_t* indices, - const int *status, - roc_shmem_cmps cmp, T* vals) { -} - -template -__device__ int test(T* ptr, roc_shmem_cmps cmp, T val) { - return 0; -} - } // namespace rocshmem #endif // LIBRARY_SRC_IPC_CONTEXT_TMPL_DEVICE_HPP_