From b7613a38c1adff29d04b5c1311dca949173b3a3e Mon Sep 17 00:00:00 2001 From: Aurelien Bouteiller Date: Tue, 15 Apr 2025 09:57:59 -0400 Subject: [PATCH] Remove dev-mono-linear (#81) * Remove dev_mono_linear (followup to removal of slab_heap) * cleanup: use CHECK_HIP rather than ad-hoc error checking --------- Signed-off-by: Aurelien Bouteiller [ROCm/rocshmem commit: a1a0560ca3489a77803b35b18f7f0050ac48a9cf] --- projects/rocshmem/src/backend_bc.cpp | 9 +- .../rocshmem/src/memory/dev_mono_linear.hpp | 144 ------------------ projects/rocshmem/src/rocshmem.cpp | 4 +- 3 files changed, 4 insertions(+), 153 deletions(-) delete mode 100644 projects/rocshmem/src/memory/dev_mono_linear.hpp diff --git a/projects/rocshmem/src/backend_bc.cpp b/projects/rocshmem/src/backend_bc.cpp index 56e57bcd52..9df442474a 100644 --- a/projects/rocshmem/src/backend_bc.cpp +++ b/projects/rocshmem/src/backend_bc.cpp @@ -42,14 +42,11 @@ namespace rocshmem { } Backend::Backend(MPI_Comm comm) : heap{comm} { - int num_cus{}; - if (hipDeviceGetAttribute(&num_cus, hipDeviceAttributeMultiprocessorCount, - 0)) { - abort(); - } - CHECK_HIP(hipGetDevice(&hip_dev_id)); + int num_cus{}; + CHECK_HIP(hipDeviceGetAttribute(&num_cus, hipDeviceAttributeMultiprocessorCount, hip_dev_id)); + /* * Initialize 'print_lock' global and copy to the device memory space. */ diff --git a/projects/rocshmem/src/memory/dev_mono_linear.hpp b/projects/rocshmem/src/memory/dev_mono_linear.hpp deleted file mode 100644 index 477bd02600..0000000000 --- a/projects/rocshmem/src/memory/dev_mono_linear.hpp +++ /dev/null @@ -1,144 +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 LIBRARY_SRC_MEMORY_DEV_MONO_LINEAR_HPP_ -#define LIBRARY_SRC_MEMORY_DEV_MONO_LINEAR_HPP_ - -#include - -#include "shmem_allocator_strategy.hpp" -#include "../util.hpp" - -/** - * @file dev_mono_linear.hpp - * - * @brief Contains an allocator strategy for the heap. - * - * This strategy returns memory chunks by monotonically increasing a pointer. - */ - -namespace rocshmem { - -template -class DevMonoLinear : public ShmemAllocatorStrategy { - public: - /** - * @brief Required for default construction of other objects - * - * @note Not intended for direct usage. - */ - DevMonoLinear() = default; - - /** - * @brief Primary constructor type - * - * @param[in] Raw pointer to heap memory type - */ - explicit DevMonoLinear(HM_T* heap_mem) - : heap_mem_{heap_mem}, current_ptr_{heap_mem_->get_ptr()} {} - - /** - * @brief Allocates memory from the heap - * - * @param[in, out] Address of raw pointer (&pointer_to_char) - * @param[in] Size in bytes of memory allocation - */ - void alloc(char** ptr, size_t request_size) override { - assert(ptr); - *ptr = nullptr; - - if (!request_size) { - return; - } - - char* heap_end{heap_mem_->get_ptr() + heap_mem_->get_size()}; - - if (current_ptr_ + request_size < heap_end) { - *ptr = current_ptr_; - current_ptr_ += request_size; - } - } - - /** - * @brief Allocates memory from the heap - * - * @param[in, out] Address of raw pointer (&pointer_to_char) - * @param[in] Size in bytes of memory allocation - */ - __device__ void alloc(char** ptr, size_t request_size) override { - if (is_thread_zero_in_block()) { - assert(ptr); - *ptr = nullptr; - - if (!request_size) { - return; - } - - char* heap_end{heap_mem_->get_ptr() + heap_mem_->get_size()}; - - if (current_ptr_ + request_size < heap_end) { - *ptr = current_ptr_; - current_ptr_ += request_size; - } - } - } - - /** - * @brief Frees memory from the heap - * - * Released memory ignored. - * - * @param[in] Raw pointer to heap memory - */ - __host__ void free([[maybe_unused]] char* ptr) override {} - - /** - * @brief Frees memory from the heap - * - * Released memory ignored. - * - * @param[in] Raw pointer to heap memory - */ - __device__ void free([[maybe_unused]] char* ptr) override {} - - /** - * @brief Return pointer to monotonic linear ptr - * - * @return Raw pointer to heap memory. - */ - __host__ __device__ char* current() { return current_ptr_; } - - private: - /** - * The pointer is used to keep reference to heap memory. - */ - HM_T* heap_mem_{nullptr}; - - /** - * The pointer is used to track the next allocation point. - */ - char* current_ptr_{nullptr}; -}; - -} // namespace rocshmem - -#endif // LIBRARY_SRC_MEMORY_DEV_MONO_LINEAR_HPP_ diff --git a/projects/rocshmem/src/rocshmem.cpp b/projects/rocshmem/src/rocshmem.cpp index 49be967068..e491646884 100644 --- a/projects/rocshmem/src/rocshmem.cpp +++ b/projects/rocshmem/src/rocshmem.cpp @@ -75,9 +75,7 @@ rocshmem_ctx_t ROCSHMEM_HOST_CTX_DEFAULT; [[maybe_unused]] __host__ void inline library_init(MPI_Comm comm) { assert(!backend); int count = 0; - if (hipGetDeviceCount(&count) != hipSuccess) { - abort(); - } + CHECK_HIP(hipGetDeviceCount(&count)); if (count == 0) { printf("No GPU found! \n");