From e9484bbb8665eb57b299d75f2cf60c4d30c6d109 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Mon, 28 Oct 2024 21:22:46 +0000 Subject: [PATCH] Remove active-set-based reduction test from the functional tests suite --- tests/functional_tests/reduction_tester.cpp | 168 -------------------- tests/functional_tests/reduction_tester.hpp | 63 -------- tests/functional_tests/tester.cpp | 91 +---------- tests/functional_tests/tester.hpp | 2 +- 4 files changed, 2 insertions(+), 322 deletions(-) delete mode 100644 tests/functional_tests/reduction_tester.cpp delete mode 100644 tests/functional_tests/reduction_tester.hpp diff --git a/tests/functional_tests/reduction_tester.cpp b/tests/functional_tests/reduction_tester.cpp deleted file mode 100644 index 4c1cbf66ba..0000000000 --- a/tests/functional_tests/reduction_tester.cpp +++ /dev/null @@ -1,168 +0,0 @@ -/****************************************************************************** - * 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. - *****************************************************************************/ - -using namespace rocshmem; - -/* Declare the template with a generic implementation */ -template -__device__ void wg_to_all(roc_shmem_ctx_t ctx, T *dest, const T *source, - int nreduce, int PE_start, int logPE_stride, - int PE_size, T *pWrk, long *pSync) { - return; -} - -/* Define templates to call ROC_SHMEM */ -#define REDUCTION_DEF_GEN(T, TNAME, Op_API, Op) \ - template <> \ - __device__ void wg_to_all( \ - roc_shmem_ctx_t ctx, T * dest, const T *source, int nreduce, \ - int PE_start, int logPE_stride, int PE_size, T *pWrk, long *pSync) { \ - roc_shmem_ctx_##TNAME##_##Op_API##_wg_to_all(ctx, dest, source, nreduce, \ - PE_start, logPE_stride, \ - PE_size, pWrk, pSync); \ - } - -#define ARITH_REDUCTION_DEF_GEN(T, TNAME) \ - REDUCTION_DEF_GEN(T, TNAME, sum, ROC_SHMEM_SUM) \ - REDUCTION_DEF_GEN(T, TNAME, min, ROC_SHMEM_MIN) \ - REDUCTION_DEF_GEN(T, TNAME, max, ROC_SHMEM_MAX) \ - REDUCTION_DEF_GEN(T, TNAME, prod, ROC_SHMEM_PROD) - -#define BITWISE_REDUCTION_DEF_GEN(T, TNAME) \ - REDUCTION_DEF_GEN(T, TNAME, or, ROC_SHMEM_OR) \ - REDUCTION_DEF_GEN(T, TNAME, and, ROC_SHMEM_AND) \ - REDUCTION_DEF_GEN(T, TNAME, xor, ROC_SHMEM_XOR) - -#define INT_REDUCTION_DEF_GEN(T, TNAME) \ - ARITH_REDUCTION_DEF_GEN(T, TNAME) \ - BITWISE_REDUCTION_DEF_GEN(T, TNAME) - -#define FLOAT_REDUCTION_DEF_GEN(T, TNAME) ARITH_REDUCTION_DEF_GEN(T, TNAME) - -INT_REDUCTION_DEF_GEN(int, int) -INT_REDUCTION_DEF_GEN(short, short) -INT_REDUCTION_DEF_GEN(long, long) -INT_REDUCTION_DEF_GEN(long long, longlong) -FLOAT_REDUCTION_DEF_GEN(float, float) -FLOAT_REDUCTION_DEF_GEN(double, double) -// long double reduction fails. hipcc/device may not support long double. -// so disable it for now. -// FLOAT_REDUCTION_DEF_GEN(long double, longdouble) - -/****************************************************************************** - * DEVICE TEST KERNEL - *****************************************************************************/ -template -__global__ void ReductionTest(int loop, int skip, uint64_t *timer, T1 *s_buf, - T1 *r_buf, T1 *pWrk, long *pSync, int size, - TestType type, ShmemContextType ctx_type) { - __shared__ roc_shmem_ctx_t ctx; - - roc_shmem_wg_init(); - roc_shmem_wg_ctx_create(ctx_type, &ctx); - - int n_pes = roc_shmem_ctx_n_pes(ctx); - - __syncthreads(); - - uint64_t start; - for (int i = 0; i < loop + skip; i++) { - if (i == skip && hipThreadIdx_x == 0) { - start = roc_shmem_timer(); - } - wg_to_all(ctx, r_buf, s_buf, size, 0, 0, n_pes, pWrk, pSync); - roc_shmem_ctx_wg_barrier_all(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 - *****************************************************************************/ -template -ReductionTester::ReductionTester( - TesterArguments args, std::function f1, - std::function(const T1 &, const T1 &)> f2) - : Tester(args), init_buf{f1}, verify_buf{f2} { - s_buf = (T1 *)roc_shmem_malloc(args.max_msg_size * sizeof(T1)); - r_buf = (T1 *)roc_shmem_malloc(args.max_msg_size * sizeof(T1)); - - size_t p_wrk_size = - std::max(args.max_msg_size / 2 + 1, ROC_SHMEM_REDUCE_MIN_WRKDATA_SIZE); - pWrk = (T1 *)roc_shmem_malloc(p_wrk_size * sizeof(T1)); - - size_t p_sync_size = ROC_SHMEM_REDUCE_SYNC_SIZE; - pSync = (long *)roc_shmem_malloc(p_sync_size * sizeof(long)); - - for (int i = 0; i < p_sync_size; i++) { - pSync[i] = ROC_SHMEM_SYNC_VALUE; - } -} - -template -ReductionTester::~ReductionTester() { - roc_shmem_free(s_buf); - roc_shmem_free(r_buf); - roc_shmem_free(pWrk); - roc_shmem_free(pSync); -} - -template -void ReductionTester::launchKernel(dim3 gridSize, dim3 blockSize, - int loop, uint64_t size) { - size_t shared_bytes = 0; - - hipLaunchKernelGGL(HIP_KERNEL_NAME(ReductionTest), gridSize, - blockSize, shared_bytes, stream, loop, args.skip, timer, - s_buf, r_buf, pWrk, pSync, size, _type, _shmem_context); - - num_msgs = loop + args.skip; - num_timed_msgs = loop; -} - -template -void ReductionTester::resetBuffers(uint64_t size) { - for (int i = 0; i < args.max_msg_size; i++) { - init_buf(s_buf[i], r_buf[i]); - } -} - -template -void ReductionTester::verifyResults(uint64_t size) { - int n_pes = roc_shmem_n_pes(); - for (int i = 0; i < size; i++) { - auto r = verify_buf(r_buf[i], (T1)n_pes); - if (r.first == false) { - fprintf(stderr, "Data validation error at idx %d\n", i); - fprintf(stderr, "%s.\n", r.second.c_str()); - exit(-1); - } - } -} diff --git a/tests/functional_tests/reduction_tester.hpp b/tests/functional_tests/reduction_tester.hpp deleted file mode 100644 index 7342fec9ff..0000000000 --- a/tests/functional_tests/reduction_tester.hpp +++ /dev/null @@ -1,63 +0,0 @@ -/****************************************************************************** - * 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 _REDUCTION_TESTER_HPP_ -#define _REDUCTION_TESTER_HPP_ - -#include -#include - -#include "tester.hpp" - -/****************************************************************************** - * HOST TESTER CLASS - *****************************************************************************/ -template -class ReductionTester : public Tester { - public: - explicit ReductionTester( - TesterArguments args, std::function f1, - std::function(const T1 &, const T1 &)> f2); - virtual ~ReductionTester(); - - 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; - - T1 *s_buf; - T1 *r_buf; - T1 *pWrk; - long *pSync; - - private: - std::function init_buf; - std::function(const T1 &, const T1 &)> - verify_buf; -}; - -#include "reduction_tester.cpp" - -#endif diff --git a/tests/functional_tests/tester.cpp b/tests/functional_tests/tester.cpp index 8d081d976b..5458c69316 100644 --- a/tests/functional_tests/tester.cpp +++ b/tests/functional_tests/tester.cpp @@ -44,7 +44,6 @@ #include "primitive_mr_tester.hpp" #include "primitive_tester.hpp" #include "random_access_tester.hpp" -#include "reduction_tester.hpp" #include "shmem_ptr_tester.hpp" #include "swarm_tester.hpp" #include "sync_tester.hpp" @@ -156,93 +155,6 @@ std::vector Tester::create(TesterArguments args) { std::to_string(n_pes)); })); return testers; - case ReductionTestType: - if (rank == 0) std::cout << "All-to-All Reduction ###" << std::endl; - - testers.push_back(new ReductionTester( - args, - [](float& f1, float& f2) { - f1 = 1; - f2 = 1; - }, - [](float v, float n_pes) { - return (v == n_pes) - ? std::make_pair(true, "") - : std::make_pair(false, "Got " + std::to_string(v) + - ", Expect " + - std::to_string(n_pes)); - })); - -#if 0 - testers.push_back( - new ReductionTester( - args, - [](double& f1, double& f2) - { - f1=1; - f2=1; - }, - [](double v, double n_pes) - { - return (v == n_pes) ? - std::make_pair(true, "") : - std::make_pair(false, - "Got "+ std::to_string(v) + ", Expect " + std::to_string(n_pes)); - } - ) - ); - - testers.push_back( new ReductionTester(args, - [](long double& f1,long double& f2){f1=1; f2=1;}, - [](long double v){ return (v==2.0) ? std::make_pair(true, "") : - std::make_pair(false, "Got "+ std::to_string(v) +", Expect 2.0. [r3]."); })); - testers.push_back( new ReductionTester(args, - [](short& f1, short& f2){f1=1; f2=2;}, - [](short v){ return (v==3) ? std::make_pair(true, "") : - std::make_pair(false, "Got "+ std::to_string(v) +", Expect 3. [r4]."); })); - testers.push_back( new ReductionTester(args, - [](int& f1, int& f2){f1=1; f2=2;}, - [](int v){ return (v==3) ? std::make_pair(true, "") : - std::make_pair(false, "Got "+ std::to_string(v) +", Expect 3. [r5]."); })); - testers.push_back( new ReductionTester(args, - [](long long& f1, long long& f2){f1=1; f2=2;}, - [](long long v){ return (v==3) ? std::make_pair(true, "") : - std::make_pair(false, "Got "+ std::to_string(v) +", Expect 3. [r6]."); })); - testers.push_back( new ReductionTester(args, - [](long long& f1, long long& f2){f1=1; f2=2;}, - [](long long v){ return (v==3) ? std::make_pair(true, "") : - std::make_pair(false, "Got "+ std::to_string(v) +", Expect 3. [r7]."); })); - // seems like deadlock or soemthing, this test hang forever - testers.push_back( new ReductionTester(args, - [](short& f1, short& f2){f1=1; f2=2;}, - [](short v){ return (v==1) ? std::make_pair(true, "") : - std::make_pair(false, "Got "+ std::to_string(v) +", Expect 1. [r8]."); })); - testers.push_back( new ReductionTester(args, - [](int& f1, int& f2){f1=1; f2=2;}, - [](int v){ return (v==1) ? std::make_pair(true, "") : - std::make_pair(false, "Got "+ std::to_string(v) +", Expect 1. [r9]."); })); - testers.push_back( new ReductionTester(args, - [](long long& f1, long long& f2){f1=1; f2=2;}, - [](long long v){ return (v==1) ? std::make_pair(true, "") : - std::make_pair(false, "Got "+ std::to_string(v) +", Expect 1. [r10]."); })); - testers.push_back( new ReductionTester(args, - [](long long& f1, long long& f2){f1=1; f2=2;}, - [](long long v){ return (v==1) ? std::make_pair(true, "") : - std::make_pair(false, "Got "+ std::to_string(v) +", Expect 1. [r11]."); })); - testers.push_back( new ReductionTester(args, - [](int& f1, int& f2){f1=1; f2=2;}, - [](int v){ return (v==2) ? std::make_pair(true, "") : - std::make_pair(false, "Got "+ std::to_string(v) +", Expect 2. [r12]."); })); - testers.push_back( new ReductionTester(args, - [](long long& f1, long long& f2){f1=1; f2=2;}, - [](long long v){ return (v==2) ? std::make_pair(true, "") : - std::make_pair(false, "Got "+ std::to_string(v) +", Expect 2. [r13]."); })); - testers.push_back( new ReductionTester(args, - [](long long& f1, long long& f2){f1=1; f2=2;}, - [](long long v){ return (v==2) ? std::make_pair(true, "") : - std::make_pair(false, "Got "+ std::to_string(v) +", Expect 2. [r14]."); })); -#endif - return testers; case BroadcastTestType: if (rank == 0) { std::cout << "Broadcast Test ###" << std::endl; @@ -658,8 +570,7 @@ bool Tester::peLaunchesKernel() { /** * Some test types are active on both sides. */ - is_launcher = is_launcher || (_type == ReductionTestType) || - (_type == TeamReductionTestType) || + is_launcher = is_launcher || (_type == TeamReductionTestType) || (_type == BroadcastTestType) || (_type == TeamBroadcastTestType) || (_type == AllToAllTestType) || (_type == FCollectTestType) || diff --git a/tests/functional_tests/tester.hpp b/tests/functional_tests/tester.hpp index ad639a63ae..58a76f895d 100644 --- a/tests/functional_tests/tester.hpp +++ b/tests/functional_tests/tester.hpp @@ -37,7 +37,7 @@ enum TestType { PutTestType = 2, PutNBITestType = 3, GetSwarmTestType = 4, - ReductionTestType = 5, + // ReductionTestType = 5, AMO_FAddTestType = 6, AMO_FIncTestType = 7, AMO_FetchTestType = 8,