From 9a524046fe1c6fa875ee0a083c2fa3fb6e1bf139 Mon Sep 17 00:00:00 2001 From: avinashkethineedi Date: Tue, 29 Oct 2024 16:18:46 +0000 Subject: [PATCH] Remove active-set-based broadcast test from the functional tests suite --- tests/functional_tests/broadcast_tester.cpp | 159 -------------------- tests/functional_tests/broadcast_tester.hpp | 61 -------- tests/functional_tests/tester.cpp | 39 ----- tests/functional_tests/tester.hpp | 2 +- tests/functional_tests/tester_arguments.cpp | 7 +- 5 files changed, 4 insertions(+), 264 deletions(-) delete mode 100644 tests/functional_tests/broadcast_tester.cpp delete mode 100644 tests/functional_tests/broadcast_tester.hpp diff --git a/tests/functional_tests/broadcast_tester.cpp b/tests/functional_tests/broadcast_tester.cpp deleted file mode 100644 index f5c936c2ed..0000000000 --- a/tests/functional_tests/broadcast_tester.cpp +++ /dev/null @@ -1,159 +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_broadcast(roc_shmem_ctx_t ctx, T *dest, const T *source, - int nelem, int pe_root, int pe_start, - int log_pe_stride, int pe_size, long *p_sync) { - return; -} - -/* Define templates to call ROC_SHMEM */ -#define BROADCAST_DEF_GEN(T, TNAME) \ - template <> \ - __device__ void wg_broadcast( \ - roc_shmem_ctx_t ctx, T * dest, const T *source, int nelem, int pe_root, \ - int pe_start, int log_pe_stride, int pe_size, long *p_sync) { \ - roc_shmem_ctx_##TNAME##_wg_broadcast(ctx, dest, source, nelem, pe_root, \ - pe_start, log_pe_stride, pe_size, \ - p_sync); \ - } - -BROADCAST_DEF_GEN(float, float) -BROADCAST_DEF_GEN(double, double) -BROADCAST_DEF_GEN(char, char) -// BROADCAST_DEF_GEN(long double, longdouble) -BROADCAST_DEF_GEN(signed char, schar) -BROADCAST_DEF_GEN(short, short) -BROADCAST_DEF_GEN(int, int) -BROADCAST_DEF_GEN(long, long) -BROADCAST_DEF_GEN(long long, longlong) -BROADCAST_DEF_GEN(unsigned char, uchar) -BROADCAST_DEF_GEN(unsigned short, ushort) -BROADCAST_DEF_GEN(unsigned int, uint) -BROADCAST_DEF_GEN(unsigned long, ulong) -BROADCAST_DEF_GEN(unsigned long long, ulonglong) - -/****************************************************************************** - * DEVICE TEST KERNEL - *****************************************************************************/ -template -__global__ void BroadcastTest(int loop, int skip, uint64_t *timer, - T1 *source_buf, T1 *dest_buf, long *pSync, - int size, 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; i++) { - if (i == skip && hipThreadIdx_x == 0) { - start = roc_shmem_timer(); - } - - wg_broadcast(ctx, - dest_buf, // T* dest - source_buf, // const T* source - size, // int nelement - 0, // int PE_root - 0, // int PE_start - 0, // int logPE_stride - n_pes, // int PE_size - pSync); // long *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 -BroadcastTester::BroadcastTester( - TesterArguments args, std::function f1, - std::function(const T1 &)> f2) - : Tester(args), init_buf{f1}, verify_buf{f2} { - source_buf = (T1 *)roc_shmem_malloc(args.max_msg_size * sizeof(T1)); - dest_buf = (T1 *)roc_shmem_malloc(args.max_msg_size * sizeof(T1)); - - size_t p_sync_size = ROC_SHMEM_BCAST_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 -BroadcastTester::~BroadcastTester() { - roc_shmem_free(source_buf); - roc_shmem_free(dest_buf); - roc_shmem_free(pSync); -} - -template -void BroadcastTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop, - uint64_t size) { - size_t shared_bytes = 0; - - hipLaunchKernelGGL(BroadcastTest, gridSize, blockSize, shared_bytes, - stream, loop, args.skip, timer, source_buf, dest_buf, - pSync, size, _shmem_context); - - num_msgs = loop + args.skip; - num_timed_msgs = loop; -} - -template -void BroadcastTester::resetBuffers(uint64_t size) { - for (int i = 0; i < args.max_msg_size; i++) { - init_buf(source_buf[i], dest_buf[i]); - } -} - -template -void BroadcastTester::verifyResults(uint64_t size) { - for (int i = 0; i < size; i++) { - auto r = verify_buf(dest_buf[i]); - 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/broadcast_tester.hpp b/tests/functional_tests/broadcast_tester.hpp deleted file mode 100644 index b436756aa4..0000000000 --- a/tests/functional_tests/broadcast_tester.hpp +++ /dev/null @@ -1,61 +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 _BROADCAST_TESTER_HPP_ -#define _BROADCAST_TESTER_HPP_ - -#include -#include - -#include "tester.hpp" - -/****************************************************************************** - * HOST TESTER CLASS - *****************************************************************************/ -template -class BroadcastTester : public Tester { - public: - explicit BroadcastTester( - TesterArguments args, std::function f1, - std::function(const T1 &)> f2); - virtual ~BroadcastTester(); - - 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 *source_buf; - T1 *dest_buf; - long *pSync; - - private: - std::function init_buf; - std::function(const T1 &)> verify_buf; -}; - -#include "broadcast_tester.cpp" - -#endif diff --git a/tests/functional_tests/tester.cpp b/tests/functional_tests/tester.cpp index 5458c69316..b73ae67b6c 100644 --- a/tests/functional_tests/tester.cpp +++ b/tests/functional_tests/tester.cpp @@ -35,7 +35,6 @@ #include "amo_extended_tester.hpp" #include "amo_standard_tester.hpp" #include "barrier_all_tester.hpp" -#include "broadcast_tester.hpp" #include "empty_tester.hpp" #include "extended_primitives.hpp" #include "fcollect_tester.hpp" @@ -155,43 +154,6 @@ std::vector Tester::create(TesterArguments args) { std::to_string(n_pes)); })); return testers; - case BroadcastTestType: - if (rank == 0) { - std::cout << "Broadcast Test ###" << std::endl; - } - testers.push_back(new BroadcastTester( - args, - [](long& f1, long& f2) { - f1 = 1; - f2 = 2; - }, - [rank](long v) { - long expected_val; - /** - * The verification routine here requires that the - * PE_root value is 0 which denotes that the - * sending processing element is rank 0. - * - * The difference in expected values arises from - * the specification for broadcast where the - * PE_root processing element does not copy the - * contents from its own source to dest during - * the broadcast. - */ - if (rank == 0) { - expected_val = 2; - } else { - expected_val = 1; - } - - return (v == expected_val) - ? std::make_pair(true, "") - : std::make_pair( - false, "Rank " + std::to_string(rank) + ", Got " + - std::to_string(v) + ", Expect " + - std::to_string(expected_val)); - })); - return testers; case TeamBroadcastTestType: if (rank == 0) { std::cout << "Team Broadcast Test ###" << std::endl; @@ -571,7 +533,6 @@ bool Tester::peLaunchesKernel() { * Some test types are active on both sides. */ is_launcher = is_launcher || (_type == TeamReductionTestType) || - (_type == BroadcastTestType) || (_type == TeamBroadcastTestType) || (_type == AllToAllTestType) || (_type == FCollectTestType) || (_type == PingPongTestType) || (_type == BarrierAllTestType) || diff --git a/tests/functional_tests/tester.hpp b/tests/functional_tests/tester.hpp index 58a76f895d..d5da3a6bc7 100644 --- a/tests/functional_tests/tester.hpp +++ b/tests/functional_tests/tester.hpp @@ -52,7 +52,7 @@ enum TestType { BarrierAllTestType = 17, SyncAllTestType = 18, SyncTestType = 19, - BroadcastTestType = 20, + // BroadcastTestType = 20, CollectTestType = 21, FCollectTestType = 22, AllToAllTestType = 23, diff --git a/tests/functional_tests/tester_arguments.cpp b/tests/functional_tests/tester_arguments.cpp index 87f73dc21b..a8c5067492 100644 --- a/tests/functional_tests/tester_arguments.cpp +++ b/tests/functional_tests/tester_arguments.cpp @@ -137,10 +137,9 @@ void TesterArguments::get_rocshmem_arguments() { TestType type = (TestType)algorithm; if ((type != BarrierAllTestType) && (type != SyncAllTestType) && - (type != SyncTestType) && (type != BroadcastTestType) && - (type != AllToAllTestType) && (type != FCollectTestType) && - (type != TeamReductionTestType) && (type != TeamBroadcastTestType) && - (type != PingAllTestType)) { + (type != SyncTestType) && (type != AllToAllTestType) && + (type != FCollectTestType) && (type != TeamReductionTestType) && + (type != TeamBroadcastTestType) && (type != PingAllTestType)) { if (numprocs != 2) { if (myid == 0) { std::cerr << "This test requires exactly two processes, we have "