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 <aurelien.bouteiller@amd.com>

[ROCm/rocshmem commit: a1a0560ca3]

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.
    */

projects/rocshmem/src/memory/dev_mono_linear.hpp

@@ -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 <cassert>
-
-#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 <typename HM_T>
-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_

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");