Merge pull request #54 from Yiltan/sig-ops
Added SHMEM Signalling Operations
[ROCm/rocshmem commit: 0c5524d7df]
This commit is contained in:
Yiltan
GitHub
@@ -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(
|
||||
|
||||
@@ -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
|
||||
|
||||
|
||||
@@ -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);
|
||||
|
||||
|
||||
@@ -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);
|
||||
|
||||
|
||||
@@ -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
|
||||
|
||||
|
||||
@@ -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 <iostream>
|
||||
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <roc_shmem/roc_shmem.hpp>
|
||||
|
||||
#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<nelem; i++) {
|
||||
message[i] = rank;
|
||||
}
|
||||
|
||||
CHECK_HIP(hipMemset(data, 0, (nelem * sizeof(uint64_t))));
|
||||
CHECK_HIP(hipDeviceSynchronize());
|
||||
|
||||
int threadsPerBlock=256;
|
||||
simple_put_signal_test<<<dim3(1), dim3(threadsPerBlock), 0, 0>>>(data, message, nelem, sig_addr, rank, dst_pe);
|
||||
roc_shmem_barrier_all();
|
||||
CHECK_HIP(hipDeviceSynchronize());
|
||||
|
||||
bool pass = true;
|
||||
for (int i=0; i<nelem; i++) {
|
||||
if (data[i] != 0) {
|
||||
pass = false;
|
||||
#if VERBOSE
|
||||
printf("[%d] Error in element %d expected 0 got %d\n", rank, i, dst[i]);
|
||||
#endif
|
||||
}
|
||||
}
|
||||
printf("[%d] Test %s \t %s\n", rank, argv[0], pass ? "[PASS]" : "[FAIL]");
|
||||
|
||||
roc_shmem_free(data);
|
||||
roc_shmem_free(message);
|
||||
roc_shmem_finalize();
|
||||
return 0;
|
||||
}
|
||||
@@ -66,6 +66,11 @@ enum ROC_SHMEM_OP {
|
||||
ROC_SHMEM_REPLACE
|
||||
};
|
||||
|
||||
enum ROC_SHMEM_SIGNAL_OPS {
|
||||
ROC_SHMEM_SIGNAL_SET,
|
||||
ROC_SHMEM_SIGNAL_ADD,
|
||||
};
|
||||
|
||||
/**
|
||||
* @brief Types defined for roc_shmem_wait() operations.
|
||||
*/
|
||||
@@ -2627,6 +2632,70 @@ GET_NBI_API_EXT_GEN(wg, unsigned long, ulong) // NOLINT(runtime/int)
|
||||
GET_NBI_API_EXT_GEN(wg, unsigned long long, ulonglong) // NOLINT(runtime/int)
|
||||
///@}
|
||||
|
||||
|
||||
/*
|
||||
* ROC_SHMEM Signalling Operations
|
||||
*/
|
||||
#define PUTMEM_SIGNAL_DEC(SUFFIX) \
|
||||
__device__ ATTR_NO_INLINE 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); \
|
||||
__device__ ATTR_NO_INLINE 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);
|
||||
|
||||
#define PUT_SIGNAL_TYPED_DEC(T, TNAME, SUFFIX) \
|
||||
__device__ ATTR_NO_INLINE 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); \
|
||||
__device__ ATTR_NO_INLINE 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);
|
||||
|
||||
#define PUT_SIGNAL_DEC(SUFFIX) \
|
||||
PUT_SIGNAL_TYPED_DEC(float, float, SUFFIX) \
|
||||
PUT_SIGNAL_TYPED_DEC(double, double, SUFFIX) \
|
||||
PUT_SIGNAL_TYPED_DEC(char, char, SUFFIX) \
|
||||
PUT_SIGNAL_TYPED_DEC(signed char, schar, SUFFIX) \
|
||||
PUT_SIGNAL_TYPED_DEC(short, short, SUFFIX) \
|
||||
PUT_SIGNAL_TYPED_DEC(int, int, SUFFIX) \
|
||||
PUT_SIGNAL_TYPED_DEC(long, long, SUFFIX) \
|
||||
PUT_SIGNAL_TYPED_DEC(long long, longlong, SUFFIX) \
|
||||
PUT_SIGNAL_TYPED_DEC(unsigned char, uchar, SUFFIX) \
|
||||
PUT_SIGNAL_TYPED_DEC(unsigned short, ushort, SUFFIX) \
|
||||
PUT_SIGNAL_TYPED_DEC(unsigned int, uint, SUFFIX) \
|
||||
PUT_SIGNAL_TYPED_DEC(unsigned long, ulong, SUFFIX) \
|
||||
PUT_SIGNAL_TYPED_DEC(unsigned long long, ulonglong, SUFFIX)
|
||||
|
||||
#define SIGNALING_API_DEC(SUFFIX) \
|
||||
PUTMEM_SIGNAL_DEC(SUFFIX) \
|
||||
PUT_SIGNAL_DEC(SUFFIX)
|
||||
|
||||
SIGNALING_API_DEC()
|
||||
SIGNALING_API_DEC(_wg)
|
||||
SIGNALING_API_DEC(_wave)
|
||||
SIGNALING_API_DEC(_nbi)
|
||||
SIGNALING_API_DEC(_nbi_wg)
|
||||
SIGNALING_API_DEC(_nbi_wave)
|
||||
|
||||
__device__ ATTR_NO_INLINE uint64_t roc_shmem_signal_fetch(const uint64_t *sig_addr);
|
||||
__device__ ATTR_NO_INLINE uint64_t roc_shmem_signal_fetch_wg(const uint64_t *sig_addr);
|
||||
__device__ ATTR_NO_INLINE uint64_t roc_shmem_signal_fetch_wave(const uint64_t *sig_addr);
|
||||
|
||||
} // namespace rocshmem
|
||||
|
||||
#endif // LIBRARY_INCLUDE_ROC_SHMEM_HPP
|
||||
|
||||
@@ -220,6 +220,33 @@ case $2 in
|
||||
echo "amoset_n2_w8_z1"
|
||||
ROC_SHMEM_MAX_NUM_CONTEXTS=8 mpirun -np 2 $1 -w 8 -z 1 -a 44 > $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
|
||||
|
||||
@@ -271,6 +271,33 @@ class Context {
|
||||
template <typename T>
|
||||
__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 <typename T> \
|
||||
__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 ******************************
|
||||
*************************************************************************/
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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 <typename T> \
|
||||
__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_
|
||||
|
||||
@@ -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<uint64_t>(static_cast<void*>(sig_addr), signal, pe);
|
||||
break;
|
||||
case ROC_SHMEM_SIGNAL_ADD:
|
||||
amo_add<uint64_t>(static_cast<void*>(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<uint64_t>(static_cast<void*>(sig_addr), signal, pe);
|
||||
break;
|
||||
case ROC_SHMEM_SIGNAL_ADD:
|
||||
amo_add<uint64_t>(static_cast<void*>(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<uint64_t>(static_cast<void*>(sig_addr), signal, pe);
|
||||
break;
|
||||
case ROC_SHMEM_SIGNAL_ADD:
|
||||
amo_add<uint64_t>(static_cast<void*>(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<uint64_t*>(sig_addr);
|
||||
return amo_fetch_add<uint64_t>(static_cast<void*>(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<uint64_t*>(sig_addr);
|
||||
value = amo_fetch_add<uint64_t>(static_cast<void*>(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<uint64_t*>(sig_addr);
|
||||
value = amo_fetch_add<uint64_t>(static_cast<void*>(dst), 0, my_pe);
|
||||
}
|
||||
__threadfence_block();
|
||||
value = __shfl(value, 0);
|
||||
return value;
|
||||
}
|
||||
|
||||
} // namespace rocshmem
|
||||
|
||||
@@ -186,6 +186,27 @@ class IPCContext : public Context {
|
||||
template <typename T>
|
||||
__device__ void get_nbi_wave(T *dest, const T *source, size_t nelems, int pe);
|
||||
|
||||
#define IPC_CONTEXT_PUT_SIGNAL_DEC(SUFFIX) \
|
||||
template <typename T> \
|
||||
__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_)
|
||||
|
||||
@@ -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 <typename T> \
|
||||
__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 <typename T> \
|
||||
__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_
|
||||
|
||||
@@ -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 <typename T> \
|
||||
__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 <typename T> \
|
||||
__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<T>(ivars, cmp, val); \
|
||||
}
|
||||
|
||||
#define RMA_SIGNAL_SUFFIX_DEC(SUFFIX) \
|
||||
template <typename T> \
|
||||
__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 <typename T> \
|
||||
__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<T>(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 **************************
|
||||
*****************************************************************************/
|
||||
|
||||
@@ -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
|
||||
};
|
||||
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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 <roc_shmem/roc_shmem.hpp>
|
||||
|
||||
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);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -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
|
||||
@@ -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*> 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));
|
||||
|
||||
@@ -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 };
|
||||
|
||||
Reference in New Issue
Block a user