diff --git a/projects/rocshmem/examples/CMakeLists.txt b/projects/rocshmem/examples/CMakeLists.txt index b12a468883..41d0fcbcac 100644 --- a/projects/rocshmem/examples/CMakeLists.txt +++ b/projects/rocshmem/examples/CMakeLists.txt @@ -28,17 +28,18 @@ set(EXAMPLE_SOURCES rocshmem_alltoall_test.cc rocshmem_broadcast_test.cc rocshmem_getmem_test.cc + rocshmem_put_signal_test.cc ) foreach(SOURCE_FILE IN LISTS EXAMPLE_SOURCES) get_filename_component(EXECUTABLE_NAME ${SOURCE_FILE} NAME_WE) - + add_executable(${EXECUTABLE_NAME} ${SOURCE_FILE}) target_include_directories( ${EXECUTABLE_NAME} PRIVATE - rocshmem::rocshmem + rocshmem::rocshmem ) target_link_libraries( diff --git a/projects/rocshmem/examples/rocshmem_allreduce_test.cc b/projects/rocshmem/examples/rocshmem_allreduce_test.cc index 671130d1f7..0ecefbd807 100644 --- a/projects/rocshmem/examples/rocshmem_allreduce_test.cc +++ b/projects/rocshmem/examples/rocshmem_allreduce_test.cc @@ -2,11 +2,11 @@ hipcc -c -fgpu-rdc -x hip rocshmem_allreduce_test.cc \ -I/opt/rocm/include \ -I$ROCSHMEM_SRC_DIR/include \ - -I$ROCHSMEM_INSTALL_DIR/include \ + -I$ROCSHMEM_INSTALL_DIR/include \ -I$OPENMPI_UCX_INSTALL_DIR/include/ hipcc -fgpu-rdc --hip-link rocshmem_allreduce_test.o -o rocshmem_allreduce_test \ - $ROCHSMEM_INSTALL_DIR/lib/librocshmem.a \ + $ROCSHMEM_INSTALL_DIR/lib/librocshmem.a \ $OPENMPI_UCX_INSTALL_DIR/lib/libmpi.so \ -L/opt/rocm/lib -lamdhip64 -lhsa-runtime64 diff --git a/projects/rocshmem/examples/rocshmem_alltoall_test.cc b/projects/rocshmem/examples/rocshmem_alltoall_test.cc index 6f48f07653..ba4c09680a 100644 --- a/projects/rocshmem/examples/rocshmem_alltoall_test.cc +++ b/projects/rocshmem/examples/rocshmem_alltoall_test.cc @@ -2,11 +2,11 @@ hipcc -c -fgpu-rdc -x hip rocshmem_alltoall_test.cc \ -I/opt/rocm/include \ -I$ROCSHMEM_SRC_DIR/include \ - -I$ROCHSMEM_INSTALL_DIR/include \ + -I$ROCSHMEM_INSTALL_DIR/include \ -I$OPENMPI_UCX_INSTALL_DIR/include/ hipcc -fgpu-rdc --hip-link rocshmem_alltoall_test.o -o rocshmem_alltoall_test \ - $ROCHSMEM_INSTALL_DIR/lib/librocshmem.a \ + $ROCSHMEM_INSTALL_DIR/lib/librocshmem.a \ $OPENMPI_UCX_INSTALL_DIR/lib/libmpi.so \ -L/opt/rocm/lib -lamdhip64 -lhsa-runtime64 @@ -126,7 +126,7 @@ int main (int argc, char **argv) bool pass = check_recvbuf(dest, nelem, my_pe, npes); printf("Test %s \t nelem %d %s\n", argv[0], nelem, pass ? "[PASS]" : "[FAIL]"); - + roc_shmem_free(source); roc_shmem_free(dest); diff --git a/projects/rocshmem/examples/rocshmem_broadcast_test.cc b/projects/rocshmem/examples/rocshmem_broadcast_test.cc index c66eca48f7..b1eba15989 100644 --- a/projects/rocshmem/examples/rocshmem_broadcast_test.cc +++ b/projects/rocshmem/examples/rocshmem_broadcast_test.cc @@ -2,11 +2,11 @@ hipcc -c -fgpu-rdc -x hip rocshmem_broadcast_test.cc \ -I/opt/rocm/include \ -I$ROCSHMEM_SRC_DIR/include \ - -I$ROCHSMEM_INSTALL_DIR/include \ + -I$ROCSHMEM_INSTALL_DIR/include \ -I$OPENMPI_UCX_INSTALL_DIR/include/ hipcc -fgpu-rdc --hip-link rocshmem_broadcast_test.o -o rocshmem_broadcast_test \ - $ROCHSMEM_INSTALL_DIR/lib/librocshmem.a \ + $ROCSHMEM_INSTALL_DIR/lib/librocshmem.a \ $OPENMPI_UCX_INSTALL_DIR/lib/libmpi.so \ -L/opt/rocm/lib -lamdhip64 -lhsa-runtime64 @@ -121,7 +121,7 @@ int main(int argc, char **argv) bool pass = check_recvbuf(dest, nelem, my_pe, npes); printf("Test %s \t nelem %d %s\n", argv[0], nelem, pass ? "[PASS]" : "[FAIL]"); } - + roc_shmem_free(source); roc_shmem_free(dest); diff --git a/projects/rocshmem/examples/rocshmem_getmem_test.cc b/projects/rocshmem/examples/rocshmem_getmem_test.cc index e041520c01..470dbbc037 100644 --- a/projects/rocshmem/examples/rocshmem_getmem_test.cc +++ b/projects/rocshmem/examples/rocshmem_getmem_test.cc @@ -2,11 +2,11 @@ hipcc -c -fgpu-rdc -x hip rocshmem_getmem_test.cc \ -I/opt/rocm/include \ -I$ROCSHMEM_SRC_DIR/include \ - -I$ROCHSMEM_INSTALL_DIR/include \ + -I$ROCSHMEM_INSTALL_DIR/include \ -I$OPENMPI_UCX_INSTALL_DIR/include/ hipcc -fgpu-rdc --hip-link rocshmem_getmem_test.o -o rocshmem_getmem_test \ - $ROCHSMEM_INSTALL_DIR/lib/librocshmem.a \ + $ROCSHMEM_INSTALL_DIR/lib/librocshmem.a \ $OPENMPI_UCX_INSTALL_DIR/lib/libmpi.so \ -L/opt/rocm/lib -lamdhip64 -lhsa-runtime64 diff --git a/projects/rocshmem/examples/rocshmem_put_signal_test.cc b/projects/rocshmem/examples/rocshmem_put_signal_test.cc new file mode 100644 index 0000000000..6ed047b13f --- /dev/null +++ b/projects/rocshmem/examples/rocshmem_put_signal_test.cc @@ -0,0 +1,111 @@ +/* +hipcc -c -fgpu-rdc -x hip rocshmem_put_signal_test.cc \ + -I/opt/rocm/include \ + -I$ROCSHMEM_SRC_DIR/include \ + -I$ROCSHMEM_INSTALL_DIR/include \ + -I$OPENMPI_UCX_INSTALL_DIR/include/ + +hipcc -fgpu-rdc --hip-link rocshmem_put_signal_test.o -o rocshmem_getmem_test \ + $ROCSHMEM_INSTALL_DIR/lib/librocshmem.a \ + $OPENMPI_UCX_INSTALL_DIR/lib/libmpi.so \ + -L/opt/rocm/lib -lamdhip64 -lhsa-runtime64 + +ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 ./rocshmem_put_signal_test +*/ + +#include + +#include +#include +#include + +#define CHECK_HIP(condition) { \ + hipError_t error = condition; \ + if(error != hipSuccess){ \ + fprintf(stderr,"HIP error: %d line: %d\n", error, __LINE__); \ + MPI_Abort(MPI_COMM_WORLD, error); \ + } \ + } + +using namespace rocshmem; + +__global__ void simple_put_signal_test(uint64_t *data, uint64_t *message, size_t nelem, + uint64_t *sig_addr, int my_pe, int dst_pe) +{ + roc_shmem_wg_init(); + + int threadId = blockIdx.x * blockDim.x + threadIdx.x; + + if (threadId == 0) { + if (my_pe == 0) { + roc_shmem_ulong_put_signal(data, message, nelem, sig_addr, 1, ROC_SHMEM_SIGNAL_SET, dst_pe); + } + else { + roc_shmem_ulong_wait_until(sig_addr, ROC_SHMEM_CMP_EQ, 1); + roc_shmem_ulong_put_signal(data, data, nelem, sig_addr, 1, ROC_SHMEM_SIGNAL_SET, dst_pe); + } + } + + __syncthreads(); + roc_shmem_wg_finalize(); +} + +#define MAX_ELEM 256 + +int main (int argc, char **argv) +{ + int rank = roc_shmem_my_pe(); + int ndevices, my_device = 0; + CHECK_HIP(hipGetDeviceCount(&ndevices)); + my_device = rank % ndevices; + CHECK_HIP(hipSetDevice(my_device)); + int nelem = MAX_ELEM; + + if (argc > 1) { + nelem = atoi(argv[1]); + } + + roc_shmem_init(); + int npes = roc_shmem_n_pes(); + int dst_pe = (rank + 1) % npes; + uint64_t *message = (uint64_t*)roc_shmem_malloc(nelem * sizeof(uint64_t)); + uint64_t *data = (uint64_t*)roc_shmem_malloc(nelem * sizeof(uint64_t)); + uint64_t *sig_addr = (uint64_t*)roc_shmem_malloc(sizeof(uint64_t)); + if (NULL == data || NULL == message || NULL == sig_addr) { + std::cout << "Error allocating memory from symmetric heap" << std::endl; + std::cout << "data: " << data + << ", message: " << message + << ", size: " << sizeof(uint64_t) * nelem + << ", sig_addr: " << sig_addr + << std::endl; + roc_shmem_global_exit(1); + } + + for (int i=0; i>>(data, message, nelem, sig_addr, rank, dst_pe); + roc_shmem_barrier_all(); + CHECK_HIP(hipDeviceSynchronize()); + + bool pass = true; + for (int i=0; i $3/amoset_n2_w8_z1.log check amoset_n2_w8_z1 + echo "putsignal_n2_w1_z1" + mpirun -np 2 $1 -w 1 -z 1 -s 1048576 -a 53 > $3/putsignal_n2_w1_z1.log + check putsignal_n2_w1_z1 + echo "putsignalwg_n2_w2_z32" + ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 32 -s 1048576 -a 55 > $3/putsignalwg_n2_w2_z32.log + check putsignalwg_n2_w2_z32 + echo "putsignalwave_n2_w1_z32" + mpirun -np 2 $1 -w 1 -z 32 -s 1048576 -a 55 > $3/putsignalwave_n2_w1_z32.log + check putsignalwave_n2_w1_z32 + echo "putsignalnbi" + mpirun -np 2 $1 -w 1 -z 1 -s 1048576 -a 56 > $3/putsignalnbi_n2_w1_z1.log + check putsignalnbi_n2_w1_z1 + echo "putsignalnbiwg" + ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 32 -s 1048576 -a 57 > $3/putsignalnbiwg_n2_w2_z32.log + check putsignalnbiwg_n2_w2_z32 + echo "putsignalnbiwave" + mpirun -np 2 $1 -w 1 -z 32 -s 1048576 -a 58 > $3/putsignalnbiwave_n2_w1_z32.log + check putsignalnbiwave_n2_w1_z32 + echo "signalfetch" + mpirun -np 2 $1 -w 1 -z 1 -a 59 > $3/signalfetch_n2_w1_z1.log + check signalfetch_n2_w1_z1 + echo "signalfetchwg" + ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 32 -a 60 > $3/signalfetchwg_n2_w2_z32.log + check signalfetchwg_n2_w2_z32 + echo "signalfetchwave" + mpirun -np 2 $1 -w 1 -z 32 -a 60 > $3/signalfetchwave_n2_w2_z32.log + check signalfetchwave_n2_w2_z32 ;; ########################################################################### @@ -576,6 +603,33 @@ case $2 in *"amoxor") mpirun -np 2 $1 -w 1 -z 1 -a 51 ;; + *"putsignal") + mpirun -np 2 $1 -w 1 -z 1 -a 53 + ;; + *"putsignalwg") + ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 32 -a 55 + ;; + *"putsignalwave") + mpirun -np 2 $1 -w 1 -z 32 -a 55 + ;; + *"putsignalnbi") + mpirun -np 2 $1 -w 1 -z 1 -a 56 + ;; + *"putsignalnbiwg") + ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 32 -a 57 + ;; + *"putsignalnbiwave") + mpirun -np 2 $1 -w 1 -z 32 -a 58 + ;; + *"signalfetch") + mpirun -np 2 $1 -w 1 -z 1 -a 59 + ;; + *"signalfetchwg") + ROC_SHMEM_MAX_NUM_CONTEXTS=2 mpirun -np 2 $1 -w 2 -z 32 -a 60 + ;; + *"signalfetchwave") + mpirun -np 2 $1 -w 1 -z 32 -a 60 + ;; *) echo "UNKNOWN TEST TYPE: $2" exit -1 diff --git a/projects/rocshmem/src/context.hpp b/projects/rocshmem/src/context.hpp index 4516a5d0a1..50f9713c4b 100644 --- a/projects/rocshmem/src/context.hpp +++ b/projects/rocshmem/src/context.hpp @@ -271,6 +271,33 @@ class Context { template __device__ void get_nbi_wave(T* dest, const T* source, size_t nelems, int pe); +#define CONTEXT_PUTMEM_SIGNAL_DEC(SUFFIX) \ + __device__ void putmem_signal##SUFFIX(void *dest, const void *source, size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, int sig_op, int pe); + + CONTEXT_PUTMEM_SIGNAL_DEC() + CONTEXT_PUTMEM_SIGNAL_DEC(_wg) + CONTEXT_PUTMEM_SIGNAL_DEC(_wave) + CONTEXT_PUTMEM_SIGNAL_DEC(_nbi) + CONTEXT_PUTMEM_SIGNAL_DEC(_nbi_wg) + CONTEXT_PUTMEM_SIGNAL_DEC(_nbi_wave) + +#define CONTEXT_PUT_SIGNAL_DEC(SUFFIX) \ + template \ + __device__ void put_signal##SUFFIX(T *dest, const T *source, size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, int sig_op, int pe); + + CONTEXT_PUT_SIGNAL_DEC() + CONTEXT_PUT_SIGNAL_DEC(_wg) + CONTEXT_PUT_SIGNAL_DEC(_wave) + CONTEXT_PUT_SIGNAL_DEC(_nbi) + CONTEXT_PUT_SIGNAL_DEC(_nbi_wg) + CONTEXT_PUT_SIGNAL_DEC(_nbi_wave) + + __device__ uint64_t signal_fetch(const uint64_t *sig_addr); + __device__ uint64_t signal_fetch_wg(const uint64_t *sig_addr); + __device__ uint64_t signal_fetch_wave(const uint64_t *sig_addr); + /************************************************************************** ****************************** HOST METHODS ****************************** *************************************************************************/ diff --git a/projects/rocshmem/src/context_device.cpp b/projects/rocshmem/src/context_device.cpp index 33f468d72d..518ad06a5e 100644 --- a/projects/rocshmem/src/context_device.cpp +++ b/projects/rocshmem/src/context_device.cpp @@ -248,4 +248,34 @@ __device__ void Context::getmem_nbi_wave(void* dest, const void* source, DISPATCH(getmem_nbi_wave(dest, source, size, pe)); } +#define CONTEXT_PUTMEM_SIGNAL_DEF(SUFFIX, STATS_SUFFIX) \ + __device__ void Context::putmem_signal##SUFFIX(void *dest, const void *source, size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, int sig_op, \ + int pe) { \ + if (nelems == 0) { \ + return; \ + } \ + \ + ctxStats.incStat(NUM_PUT_SIGNAL##STATS_SUFFIX); \ + \ + DISPATCH(putmem_signal##SUFFIX(dest, source, nelems, sig_addr, signal, sig_op, pe)); \ + } + +CONTEXT_PUTMEM_SIGNAL_DEF(,) +CONTEXT_PUTMEM_SIGNAL_DEF(_wg, _WG) +CONTEXT_PUTMEM_SIGNAL_DEF(_wave, _WAVE) +CONTEXT_PUTMEM_SIGNAL_DEF(_nbi, _NBI) +CONTEXT_PUTMEM_SIGNAL_DEF(_nbi_wg, _NBI_WG) +CONTEXT_PUTMEM_SIGNAL_DEF(_nbi_wave, _NBI_WAVE) + +#define CONTEXT_SIGNAL_FETCH_DEF(SUFFIX) \ +__device__ uint64_t Context::signal_fetch##SUFFIX(const uint64_t *sig_addr) \ +{ \ + DISPATCH_RET(signal_fetch##SUFFIX(sig_addr)); \ +} + +CONTEXT_SIGNAL_FETCH_DEF() +CONTEXT_SIGNAL_FETCH_DEF(_wg) +CONTEXT_SIGNAL_FETCH_DEF(_wave) + } // namespace rocshmem diff --git a/projects/rocshmem/src/context_tmpl_device.hpp b/projects/rocshmem/src/context_tmpl_device.hpp index b504fce229..532c64813a 100644 --- a/projects/rocshmem/src/context_tmpl_device.hpp +++ b/projects/rocshmem/src/context_tmpl_device.hpp @@ -556,6 +556,27 @@ __device__ void Context::amo_cas(void *dst, T value, T cond, int pe) { DISPATCH(amo_cas(dst, value, cond, pe)); } +#define CONTEXT_PUT_SIGNAL_DEF(SUFFIX, STATS_SUFFIX) \ + template \ + __device__ void Context::put_signal##SUFFIX(T *dest, const T *source, size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, int sig_op, \ + int pe) { \ + if (nelems == 0) { \ + return; \ + } \ + \ + ctxStats.incStat(NUM_PUT_SIGNAL##STATS_SUFFIX); \ + \ + DISPATCH(put_signal##SUFFIX(dest, source, nelems, sig_addr, signal, sig_op, pe)); \ + } + +CONTEXT_PUT_SIGNAL_DEF(,) +CONTEXT_PUT_SIGNAL_DEF(_wg, _WG) +CONTEXT_PUT_SIGNAL_DEF(_wave, _WAVE) +CONTEXT_PUT_SIGNAL_DEF(_nbi, _NBI) +CONTEXT_PUT_SIGNAL_DEF(_nbi_wg, _NBI_WG) +CONTEXT_PUT_SIGNAL_DEF(_nbi_wave, _NBI_WAVE) + } // namespace rocshmem #endif // LIBRARY_SRC_CONTEXT_TMPL_DEVICE_HPP_ diff --git a/projects/rocshmem/src/ipc/context_ipc_device.cpp b/projects/rocshmem/src/ipc/context_ipc_device.cpp index 47d45565ef..81d1475b9e 100644 --- a/projects/rocshmem/src/ipc/context_ipc_device.cpp +++ b/projects/rocshmem/src/ipc/context_ipc_device.cpp @@ -222,4 +222,109 @@ __device__ void IPCContext::internal_getmem_wave(void *dest, ipcImpl_.ipcFence(); } +__device__ void IPCContext::putmem_signal(void *dest, const void *source, size_t nelems, + uint64_t *sig_addr, uint64_t signal, int sig_op, + int pe) { + putmem(dest, source, nelems, pe); + fence(); + + switch (sig_op) { + case ROC_SHMEM_SIGNAL_SET: + amo_set(static_cast(sig_addr), signal, pe); + break; + case ROC_SHMEM_SIGNAL_ADD: + amo_add(static_cast(dest), signal, pe); + break; + default: + DPRINTF("[%s] Invalid sig_op value (%d)\n", __func__, sig_op); + break; + } +} + +__device__ void IPCContext::putmem_signal_wg(void *dest, const void *source, size_t nelems, + uint64_t *sig_addr, uint64_t signal, int sig_op, + int pe) { + putmem_wg(dest, source, nelems, pe); + fence(); + + if (is_thread_zero_in_block()) { + switch (sig_op) { + case ROC_SHMEM_SIGNAL_SET: + amo_set(static_cast(sig_addr), signal, pe); + break; + case ROC_SHMEM_SIGNAL_ADD: + amo_add(static_cast(dest), signal, pe); + break; + default: + DPRINTF("[%s] Invalid sig_op value (%d)\n", __func__, sig_op); + break; + } + } +} + +__device__ void IPCContext::putmem_signal_wave(void *dest, const void *source, size_t nelems, + uint64_t *sig_addr, uint64_t signal, int sig_op, + int pe) { + putmem_wave(dest, source, nelems, pe); + fence(); + + if (is_thread_zero_in_wave()) { + switch (sig_op) { + case ROC_SHMEM_SIGNAL_SET: + amo_set(static_cast(sig_addr), signal, pe); + break; + case ROC_SHMEM_SIGNAL_ADD: + amo_add(static_cast(dest), signal, pe); + break; + default: + DPRINTF("[%s] Invalid sig_op value (%d)\n", __func__, sig_op); + break; + } + } +} + +__device__ void IPCContext::putmem_signal_nbi(void *dest, const void *source, size_t nelems, + uint64_t *sig_addr, uint64_t signal, int sig_op, + int pe) { + putmem_signal(dest, source, nelems, sig_addr, signal, sig_op, pe); +} + +__device__ void IPCContext::putmem_signal_nbi_wg(void *dest, const void *source, size_t nelems, + uint64_t *sig_addr, uint64_t signal, int sig_op, + int pe) { + putmem_signal_wg(dest, source, nelems, sig_addr, signal, sig_op, pe); +} + +__device__ void IPCContext::putmem_signal_nbi_wave(void *dest, const void *source, size_t nelems, + uint64_t *sig_addr, uint64_t signal, int sig_op, + int pe) { + putmem_signal_wave(dest, source, nelems, sig_addr, signal, sig_op, pe); +} + +__device__ uint64_t IPCContext::signal_fetch(const uint64_t *sig_addr) { + uint64_t *dst = const_cast(sig_addr); + return amo_fetch_add(static_cast(dst), 0, my_pe); +} + +__device__ uint64_t IPCContext::signal_fetch_wg(const uint64_t *sig_addr) { + __shared__ uint64_t value; + if (is_thread_zero_in_block()) { + uint64_t *dst = const_cast(sig_addr); + value = amo_fetch_add(static_cast(dst), 0, my_pe); + } + __threadfence_block(); + return value; +} + +__device__ uint64_t IPCContext::signal_fetch_wave(const uint64_t *sig_addr) { + uint64_t value; + if (is_thread_zero_in_wave()) { + uint64_t *dst = const_cast(sig_addr); + value = amo_fetch_add(static_cast(dst), 0, my_pe); + } + __threadfence_block(); + value = __shfl(value, 0); + return value; +} + } // namespace rocshmem diff --git a/projects/rocshmem/src/ipc/context_ipc_device.hpp b/projects/rocshmem/src/ipc/context_ipc_device.hpp index 03b72efe89..2f3e899546 100644 --- a/projects/rocshmem/src/ipc/context_ipc_device.hpp +++ b/projects/rocshmem/src/ipc/context_ipc_device.hpp @@ -186,6 +186,27 @@ class IPCContext : public Context { template __device__ void get_nbi_wave(T *dest, const T *source, size_t nelems, int pe); +#define IPC_CONTEXT_PUT_SIGNAL_DEC(SUFFIX) \ + template \ + __device__ void put_signal##SUFFIX(T *dest, const T *source, size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, int sig_op, \ + int pe); \ + \ + __device__ void putmem_signal##SUFFIX(void *dest, const void *source, size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, int sig_op, \ + int pe); + + IPC_CONTEXT_PUT_SIGNAL_DEC() + IPC_CONTEXT_PUT_SIGNAL_DEC(_wg) + IPC_CONTEXT_PUT_SIGNAL_DEC(_wave) + IPC_CONTEXT_PUT_SIGNAL_DEC(_nbi) + IPC_CONTEXT_PUT_SIGNAL_DEC(_nbi_wg) + IPC_CONTEXT_PUT_SIGNAL_DEC(_nbi_wave) + + __device__ uint64_t signal_fetch(const uint64_t *sig_addr); + __device__ uint64_t signal_fetch_wg(const uint64_t *sig_addr); + __device__ uint64_t signal_fetch_wave(const uint64_t *sig_addr); + private: //context class has IpcImpl object (ipcImpl_) diff --git a/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp b/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp index 4da49b6f7b..f239672243 100644 --- a/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp +++ b/projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp @@ -578,6 +578,25 @@ __device__ void IPCContext::get_nbi_wave(T *dest, const T *source, getmem_nbi_wave(dest, source, nelems * sizeof(T), pe); } +#define IPC_CONTEXT_PUT_SIGNAL_DEF(SUFFIX) \ + template \ + __device__ void IPCContext::put_signal##SUFFIX(T *dest, const T *source, size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, int sig_op, \ + int pe) { \ + putmem_signal##SUFFIX(dest, source, nelems * sizeof(T), sig_addr, signal, sig_op, pe); \ + } \ + \ + template \ + __device__ void IPCContext::put_signal_nbi##SUFFIX(T *dest, const T *source, size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, int sig_op, \ + int pe) { \ + putmem_signal##SUFFIX(dest, source, nelems * sizeof(T), sig_addr, signal, sig_op, pe); \ + } + +IPC_CONTEXT_PUT_SIGNAL_DEF() +IPC_CONTEXT_PUT_SIGNAL_DEF(_wg) +IPC_CONTEXT_PUT_SIGNAL_DEF(_wave) + } // namespace rocshmem #endif // LIBRARY_SRC_IPC_CONTEXT_TMPL_DEVICE_HPP_ diff --git a/projects/rocshmem/src/roc_shmem_gpu.cpp b/projects/rocshmem/src/roc_shmem_gpu.cpp index 2deff5d0b7..949f6619b1 100644 --- a/projects/rocshmem/src/roc_shmem_gpu.cpp +++ b/projects/rocshmem/src/roc_shmem_gpu.cpp @@ -235,6 +235,31 @@ __device__ void roc_shmem_atomic_xor(T *dest, T value, int pe) { roc_shmem_atomic_xor(ROC_SHMEM_CTX_DEFAULT, dest, value, pe); } +#define ROC_SHMEM_PUTMEM_SIGNAL_DEF(SUFFIX) \ + __device__ void roc_shmem_putmem_signal##SUFFIX(void *dest, const void *source, size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, \ + int sig_op, int pe) { \ + roc_shmem_ctx_putmem_signal##SUFFIX(ROC_SHMEM_CTX_DEFAULT, \ + dest, source, nelems, \ + sig_addr, signal, sig_op, pe); \ + } \ + \ + template \ + __device__ void roc_shmem_put_signal##SUFFIX(T *dest, const T *source, size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, \ + int sig_op, int pe) { \ + roc_shmem_ctx_put_signal##SUFFIX(ROC_SHMEM_CTX_DEFAULT, \ + dest, source, nelems, \ + sig_addr, signal, sig_op, pe); \ + } + +ROC_SHMEM_PUTMEM_SIGNAL_DEF() +ROC_SHMEM_PUTMEM_SIGNAL_DEF(_wg) +ROC_SHMEM_PUTMEM_SIGNAL_DEF(_wave) +ROC_SHMEM_PUTMEM_SIGNAL_DEF(_nbi) +ROC_SHMEM_PUTMEM_SIGNAL_DEF(_nbi_wg) +ROC_SHMEM_PUTMEM_SIGNAL_DEF(_nbi_wave) + /****************************************************************************** ************************* Private Context Interfaces ************************* *****************************************************************************/ @@ -846,6 +871,47 @@ __device__ void roc_shmem_get_nbi_wave(roc_shmem_ctx_t ctx, T *dest, get_internal_ctx(ctx)->get_nbi_wave(dest, source, nelems, pe); } +#define ROC_SHMEM_CTX_PUTMEM_SIGNAL_DEF(SUFFIX) \ + __device__ void roc_shmem_ctx_putmem_signal##SUFFIX(roc_shmem_ctx_t ctx, \ + void *dest, const void *source, \ + size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, \ + int sig_op, \ + int pe) { \ + GPU_DPRINTF("Function: roc_shmem_ctx_putmem_signal##SUFFIX\n"); \ + \ + get_internal_ctx(ctx)->putmem_signal##SUFFIX(dest, source, nelems, \ + sig_addr, signal, sig_op, pe); \ + } \ + \ + template \ + __device__ void roc_shmem_ctx_put_signal##SUFFIX(roc_shmem_ctx_t ctx, \ + T *dest, const T *source, \ + size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, \ + int sig_op, int pe) { \ + GPU_DPRINTF("Function: roc_shmem_ctx_put_signal##SUFFIX\n"); \ + \ + get_internal_ctx(ctx)->put_signal##SUFFIX(dest, source, nelems, \ + sig_addr, signal, sig_op, pe); \ + } + +ROC_SHMEM_CTX_PUTMEM_SIGNAL_DEF() +ROC_SHMEM_CTX_PUTMEM_SIGNAL_DEF(_wg) +ROC_SHMEM_CTX_PUTMEM_SIGNAL_DEF(_wave) +ROC_SHMEM_CTX_PUTMEM_SIGNAL_DEF(_nbi) +ROC_SHMEM_CTX_PUTMEM_SIGNAL_DEF(_nbi_wg) +ROC_SHMEM_CTX_PUTMEM_SIGNAL_DEF(_nbi_wave) + +#define ROC_SHMEM_SIGNAL_FETCH_DEF(SUFFIX) \ + __device__ uint64_t roc_shmem_signal_fetch##SUFFIX(const uint64_t *sig_addr) { \ + return get_internal_ctx(ROC_SHMEM_CTX_DEFAULT)->signal_fetch##SUFFIX(sig_addr); \ + } + +ROC_SHMEM_SIGNAL_FETCH_DEF() +ROC_SHMEM_SIGNAL_FETCH_DEF(_wg) +ROC_SHMEM_SIGNAL_FETCH_DEF(_wave) + /****************************************************************************** ****************************** Teams Interface ******************************* *****************************************************************************/ @@ -1384,6 +1450,67 @@ __device__ int roc_shmem_team_translate_pe(roc_shmem_team_t src_team, return roc_shmem_test(ivars, cmp, val); \ } +#define RMA_SIGNAL_SUFFIX_DEC(SUFFIX) \ + template \ + __device__ void roc_shmem__ctx_put_signal##SUFFIX(roc_shmem_ctx_t ctx, \ + T *dest, const T *source, \ + size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, \ + int sig_op, int pe); \ + \ + template \ + __device__ void roc_shmem__put_signal##SUFFIX(T *dest, const T *source, size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, \ + int sig_op, int pe); \ + +#define RMA_SIGNAL_SUFFIX_DEF(T, TNAME, SUFFIX) \ + __device__ void roc_shmem_ctx_##TNAME##_put_signal##SUFFIX(roc_shmem_ctx_t ctx, \ + T *dest, const T *source, \ + size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, \ + int sig_op, int pe) { \ + roc_shmem_ctx_put_signal##SUFFIX(ctx, dest, source, nelems, sig_addr, signal, sig_op, pe); \ + } \ + \ + __device__ void roc_shmem_##TNAME##_put_signal##SUFFIX(T *dest, const T *source, size_t nelems, \ + uint64_t *sig_addr, uint64_t signal, \ + int sig_op, int pe) { \ + roc_shmem_put_signal##SUFFIX(dest, source, nelems, sig_addr, signal, sig_op, pe); \ + } + +#define RMA_SIGNAL_GEN(SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEC(SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(float, float, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(double, double, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(char, char, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(signed char, schar, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(short, short, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(int, int, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(long, long, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(long long, longlong, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(unsigned char, uchar, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(unsigned short, ushort, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(unsigned int, uint, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(unsigned long, ulong, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(unsigned long long, ulonglong, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(int8_t, int8, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(int16_t, int16, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(int32_t, int32, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(int64_t, int64, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(uint8_t, uint8, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(uint16_t, uint16, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(uint32_t, uint32, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(uint64_t, uint64, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(size_t, size, SUFFIX) \ + RMA_SIGNAL_SUFFIX_DEF(ptrdiff_t, ptrdiff, SUFFIX) + +RMA_SIGNAL_GEN(_wg) +RMA_SIGNAL_GEN() +RMA_SIGNAL_GEN(_wave) +RMA_SIGNAL_GEN(_nbi) +RMA_SIGNAL_GEN(_nbi_wg) +RMA_SIGNAL_GEN(_nbi_wave) + /****************************************************************************** ************************* Macro Invocation Per Type ************************** *****************************************************************************/ diff --git a/projects/rocshmem/src/stats.hpp b/projects/rocshmem/src/stats.hpp index 704313f102..15cb16779c 100644 --- a/projects/rocshmem/src/stats.hpp +++ b/projects/rocshmem/src/stats.hpp @@ -82,6 +82,12 @@ enum roc_shmem_stats { NUM_CREATE, NUM_ALLTOALL, NUM_FCOLLECT, + NUM_PUT_SIGNAL, + NUM_PUT_SIGNAL_WG, + NUM_PUT_SIGNAL_WAVE, + NUM_PUT_SIGNAL_NBI, + NUM_PUT_SIGNAL_NBI_WG, + NUM_PUT_SIGNAL_NBI_WAVE, NUM_STATS }; diff --git a/projects/rocshmem/tests/functional_tests/CMakeLists.txt b/projects/rocshmem/tests/functional_tests/CMakeLists.txt index 542462261d..1309f58a90 100644 --- a/projects/rocshmem/tests/functional_tests/CMakeLists.txt +++ b/projects/rocshmem/tests/functional_tests/CMakeLists.txt @@ -53,6 +53,8 @@ target_sources( swarm_tester.cpp random_access_tester.cpp shmem_ptr_tester.cpp + signaling_operations_tester.cpp + signaling_operations_tester.hpp extended_primitives.cpp empty_tester.cpp wave_level_primitives.cpp diff --git a/projects/rocshmem/tests/functional_tests/signaling_operations_tester.cpp b/projects/rocshmem/tests/functional_tests/signaling_operations_tester.cpp new file mode 100644 index 0000000000..bb65a7b679 --- /dev/null +++ b/projects/rocshmem/tests/functional_tests/signaling_operations_tester.cpp @@ -0,0 +1,163 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#include "signaling_operations_tester.hpp" + +#include + +using namespace rocshmem; + +/****************************************************************************** + * DEVICE TEST KERNEL + *****************************************************************************/ +__global__ void SignalingOperationsTest(int loop, int skip, uint64_t *timer, char *s_buf, + char *r_buf, int size, uint64_t *sig_addr, + uint64_t *fetched_value, + TestType type, ShmemContextType ctx_type) { + __shared__ roc_shmem_ctx_t ctx; + roc_shmem_wg_init(); + roc_shmem_wg_ctx_create(ctx_type, &ctx); + + uint64_t start; + uint64_t signal = 0; + int sig_op = ROC_SHMEM_SIGNAL_SET; + + for (int i = 0; i < loop + skip; i++) { + if (i == skip) { + __syncthreads(); + start = roc_shmem_timer(); + } + + switch (type) { + case PutSignalTestType: + roc_shmem_ctx_putmem_signal(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1); + break; + case WGPutSignalTestType: + roc_shmem_ctx_putmem_signal_wg(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1); + break; + case WAVEPutSignalTestType: + roc_shmem_ctx_putmem_signal_wave(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1); + break; + case PutSignalNBITestType: + roc_shmem_ctx_putmem_signal_nbi(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1); + break; + case WGPutSignalNBITestType: + roc_shmem_ctx_putmem_signal_nbi_wg(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1); + break; + case WAVEPutSignalNBITestType: + roc_shmem_ctx_putmem_signal_nbi_wave(ctx, r_buf, s_buf, size, sig_addr, signal, sig_op, 1); + break; + case SignalFetchTestType: + *fetched_value = roc_shmem_signal_fetch(sig_addr); + break; + case WGSignalFetchTestType: + *fetched_value = roc_shmem_signal_fetch_wg(sig_addr); + break; + case WAVESignalFetchTestType: + *fetched_value = roc_shmem_signal_fetch_wave(sig_addr); + break; + default: + break; + } + } + + roc_shmem_ctx_quiet(ctx); + + __syncthreads(); + + if (hipThreadIdx_x == 0) { + timer[hipBlockIdx_x] = roc_shmem_timer() - start; + } + + roc_shmem_wg_ctx_destroy(&ctx); + roc_shmem_wg_finalize(); +} + +/****************************************************************************** + * HOST TESTER CLASS METHODS + *****************************************************************************/ +SignalingOperationsTester::SignalingOperationsTester(TesterArguments args) : Tester(args) { + s_buf = (char *)roc_shmem_malloc(args.max_msg_size * args.wg_size); + r_buf = (char *)roc_shmem_malloc(args.max_msg_size * args.wg_size); + sig_addr = (uint64_t *)roc_shmem_malloc(sizeof(uint64_t)); + CHECK_HIP(hipMallocManaged(&fetched_value, sizeof(uint64_t), hipMemAttachHost)); +} + +SignalingOperationsTester::~SignalingOperationsTester() { + roc_shmem_free(s_buf); + roc_shmem_free(r_buf); + roc_shmem_free(sig_addr); + CHECK_HIP(hipFree(fetched_value)); +} + +void SignalingOperationsTester::resetBuffers(uint64_t size) { + memset(s_buf, '0', args.max_msg_size * args.wg_size); + memset(r_buf, '1', args.max_msg_size * args.wg_size); + *fetched_value = -1; + *sig_addr = args.myid + 123; +} + +void SignalingOperationsTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, + uint64_t size) { + size_t shared_bytes = 0; + + hipLaunchKernelGGL(SignalingOperationsTest, gridSize, blockSize, shared_bytes, stream, + loop, args.skip, timer, s_buf, r_buf, size, sig_addr, fetched_value, + _type, _shmem_context); + + num_msgs = (loop + args.skip) * gridSize.x; + num_timed_msgs = loop; +} + +void SignalingOperationsTester::verifyResults(uint64_t size) { + int check_data_id = (_type == PutSignalTestType || + _type == PutSignalNBITestType || + _type == WAVEPutSignalTestType || + _type == WAVEPutSignalNBITestType || + _type == WGPutSignalTestType || + _type == WGPutSignalNBITestType) + ? 1 : -1; // do not check if it doesn't match a test + + int check_fetched_value_id = (_type == SignalFetchTestType || + _type == WAVESignalFetchTestType || + _type == WGSignalFetchTestType) + ? 0 : -1; // do not check if it doesn't match a test + + if (args.myid == check_data_id) { + for (int i = 0; i < size; i++) { + if (r_buf[i] != '0') { + fprintf(stderr, "Data validation error at idx %d\n", i); + fprintf(stderr, "Got %c, Expected %c\n", r_buf[i], '0'); + exit(-1); + } + } + } + + if (args.myid == check_fetched_value_id) { + uint64_t value = *fetched_value; + uint64_t expected_value = (args.myid + 123); + if (value != expected_value) { + fprintf(stderr, "Fetched Value %lu, Expected %lu\n", value, expected_value); + exit(-1); + } + } +} diff --git a/projects/rocshmem/tests/functional_tests/signaling_operations_tester.hpp b/projects/rocshmem/tests/functional_tests/signaling_operations_tester.hpp new file mode 100644 index 0000000000..acff9ea057 --- /dev/null +++ b/projects/rocshmem/tests/functional_tests/signaling_operations_tester.hpp @@ -0,0 +1,50 @@ +/****************************************************************************** + * Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#ifndef _SIGNALING_OPERATIONS_ +#define _SIGNALING_OPERATIONS_ + +#include "tester.hpp" + +/****************************************************************************** + * HOST TESTER CLASS + *****************************************************************************/ +class SignalingOperationsTester : public Tester { + public: + explicit SignalingOperationsTester(TesterArguments args); + virtual ~SignalingOperationsTester(); + + protected: + virtual void resetBuffers(uint64_t size) override; + + virtual void launchKernel(dim3 gridSize, dim3 blockSize, int loop, + uint64_t size) override; + + virtual void verifyResults(uint64_t size) override; + + char *s_buf = nullptr; + char *r_buf = nullptr; + uint64_t *sig_addr; + uint64_t *fetched_value; +}; + +#endif diff --git a/projects/rocshmem/tests/functional_tests/tester.cpp b/projects/rocshmem/tests/functional_tests/tester.cpp index b73ae67b6c..a6d6b979db 100644 --- a/projects/rocshmem/tests/functional_tests/tester.cpp +++ b/projects/rocshmem/tests/functional_tests/tester.cpp @@ -44,6 +44,7 @@ #include "primitive_tester.hpp" #include "random_access_tester.hpp" #include "shmem_ptr_tester.hpp" +#include "signaling_operations_tester.hpp" #include "swarm_tester.hpp" #include "sync_tester.hpp" #include "team_broadcast_tester.hpp" @@ -419,6 +420,42 @@ std::vector Tester::create(TesterArguments args) { } testers.push_back(new WaveLevelPrimitiveTester(args)); return testers; + case PutSignalTestType: + if (rank == 0) std::cout << "Putmem Signal ###" << std::endl; + testers.push_back(new SignalingOperationsTester(args)); + return testers; + case WGPutSignalTestType: + if (rank == 0) std::cout << "WG Putmem Signal ###" << std::endl; + testers.push_back(new SignalingOperationsTester(args)); + return testers; + case WAVEPutSignalTestType: + if (rank == 0) std::cout << "Wave Putmem Signal ###" << std::endl; + testers.push_back(new SignalingOperationsTester(args)); + return testers; + case PutSignalNBITestType: + if (rank == 0) std::cout << "Non-Blocking Putmem Signal ###" << std::endl; + testers.push_back(new SignalingOperationsTester(args)); + return testers; + case WGPutSignalNBITestType: + if (rank == 0) std::cout << "Non-Blocking WG Putmem Signal ###" << std::endl; + testers.push_back(new SignalingOperationsTester(args)); + return testers; + case WAVEPutSignalNBITestType: + if (rank == 0) std::cout << "Non-Blocking Wave Putmem Signal ###" << std::endl; + testers.push_back(new SignalingOperationsTester(args)); + return testers; + case SignalFetchTestType: + if (rank == 0) std::cout << "Signal Fetch ###" << std::endl; + testers.push_back(new SignalingOperationsTester(args)); + return testers; + case WGSignalFetchTestType: + if (rank == 0) std::cout << "WG Signal Fetch ###" << std::endl; + testers.push_back(new SignalingOperationsTester(args)); + return testers; + case WAVESignalFetchTestType: + if (rank == 0) std::cout << "Wave Signal Fetch ###" << std::endl; + testers.push_back(new SignalingOperationsTester(args)); + return testers; default: if (rank == 0) std::cout << "Unknown ###" << std::endl; testers.push_back(new PrimitiveTester(args)); diff --git a/projects/rocshmem/tests/functional_tests/tester.hpp b/projects/rocshmem/tests/functional_tests/tester.hpp index d5da3a6bc7..f15f0ea7c1 100644 --- a/projects/rocshmem/tests/functional_tests/tester.hpp +++ b/projects/rocshmem/tests/functional_tests/tester.hpp @@ -84,7 +84,16 @@ enum TestType { AMO_AndTestType = 49, AMO_OrTestType = 50, AMO_XorTestType = 51, - PingAllTestType = 52 + PingAllTestType = 52, + PutSignalTestType = 53, + WGPutSignalTestType = 54, + WAVEPutSignalTestType = 55, + PutSignalNBITestType = 56, + WGPutSignalNBITestType = 57, + WAVEPutSignalNBITestType = 58, + SignalFetchTestType = 59, + WGSignalFetchTestType = 60, + WAVESignalFetchTestType = 61, }; enum OpType { PutType = 0, GetType = 1 };