diff --git a/projects/rocshmem/scripts/functional_tests/driver.sh b/projects/rocshmem/scripts/functional_tests/driver.sh index 5fa93051ef..c2a76559a1 100755 --- a/projects/rocshmem/scripts/functional_tests/driver.sh +++ b/projects/rocshmem/scripts/functional_tests/driver.sh @@ -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 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 };