Substitute pow2bin allocator with a dlmalloc based allocator (#71)

* Add dlmalloc_strat allocator strategy
 - Use mspace variant to ease encapsulation
 - Make pow2bins and dlmalloc cmake selectable
* Add unit tester for dlmalloc, rework single_heap, pow2bins unit testers
accordingly
 - add dlmalloc get_used/get_avail, and have all strats allocators also have a get_used
 - Rework memallocator unit tests: bin size is per strat, alignment is verified in singleheap
* bugfix: dlmalloc exposed that the pingpong test would write past end of
allocation with -w 32
* iostream leakage/mixed usage of cerr and fprintf(stderr

---------

Signed-off-by: Aurelien Bouteiller <aurelien.bouteiller@amd.com>

CMakeLists.txt

@@ -49,6 +49,8 @@ option(USE_COHERENT_HEAP "Enable support for coherent systems" OFF)
 option(USE_MANAGED_HEAP "Enable managed memory" OFF)
 option(USE_HOST_HEAP "Enable host memory using malloc/free" OFF)
 option(USE_HIP_HOST_HEAP "Enable host memory using hip api" OFF)
+option(USE_ALLOC_DLMALLOC "Enable dlmalloc device memory allocator" ON)
+option(USE_ALLOC_POW2BINS "Enable legacy Pow2Bins device memory allocator" OFF)
 option(USE_FUNC_CALL "Force compiler to use function calls on library API" OFF)
 option(USE_SHARED_CTX "Request support for shared ctx between WG" OFF)
 option(USE_SINGLE_NODE "Enable single node support only." OFF)

cmake/rocshmem_config.h.in

@@ -33,6 +33,8 @@
 #cmakedefine USE_MANAGED_HEAP
 #cmakedefine USE_HOST_HEAP
 #cmakedefine USE_HIP_HOST_HEAP
+#cmakedefine USE_ALLOC_DLMALLOC
+#cmakedefine USE_ALLOC_POW2BINS
 #cmakedefine USE_FUNC_CALL
 #cmakedefine USE_SINGLE_NODE
 #cmakedefine USE_HOST_SIDE_HDP_FLUSH

scripts/functional_tests/driver.sh

@@ -429,6 +429,10 @@ TestOther() {
   ExecTest  "pingpong"         2       8            1
   ExecTest  "pingpong"         2       32           1
 
+  ExecTest  "pingall"          2       1            1
+  ExecTest  "pingall"          2       8            1
+  ExecTest  "pingall"          2       32           1
+
   # This test requires more contexts than workgroups
   export ROCSHMEM_MAX_NUM_CONTEXTS=1024
   ExecTest  "teamctxinfra"     2       1            1

src/backend_bc.cpp

@@ -93,7 +93,7 @@ void Backend::init_mpi_once(MPI_Comm comm) {
   if (!init_done) {
     NET_CHECK(MPI_Init_thread(0, 0, MPI_THREAD_MULTIPLE, &provided));
     if (provided != MPI_THREAD_MULTIPLE) {
-      std::cerr << "MPI_THREAD_MULTIPLE support disabled.\n";
+      fprintf(stderr, "MPI_THREAD_MULTIPLE support disabled.\n");
     }
   }
   if (comm == MPI_COMM_NULL) comm = MPI_COMM_WORLD;

src/memory/CMakeLists.txt

@@ -30,4 +30,5 @@ target_sources(
   PRIVATE
     single_heap.cpp
     memory_allocator.cpp
+    dlmalloc.cpp
 )

src/memory/dlmalloc.hpp

@@ -0,0 +1,186 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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_DLMALLOC_HPP_
+#define LIBRARY_SRC_MEMORY_DLMALLOC_HPP_
+
+#include <cassert>
+#include <map>
+
+#include "../constants.hpp"
+#include "shmem_allocator_strategy.hpp"
+
+/**
+ * @file dlmalloc.hpp
+ *
+ * @brief Contains an allocator strategy for the heap using dlmalloc.
+ *
+ * This strategy uses DLMalloc to allocate/free from the symmetric heap
+ *
+ */
+
+namespace rocshmem {
+
+
+/*
+ * @brief an encapsulation class for the C-static functions inherited from dlmalloc
+ *
+ * @note only a subset of dlmalloc is exposed, not intended for external use
+ *
+ * @rationale static functions cannot be defined in the templated class DLAllocatorStategy
+ */
+class DLMalloc {
+ public:
+  typedef void* mspace;
+  static size_t destroy_mspace(mspace msp);
+  static mspace create_mspace_with_base(void* base, size_t capacity, int locked);
+  static void* mspace_malloc(mspace msp, size_t bytes);
+  static void mspace_free(mspace msp, void* mem);
+  static void* mspace_memalign(mspace msp, size_t alignment, size_t bytes);
+  static size_t mspace_footprint(mspace msp);
+  static size_t mspace_max_footprint(mspace msp);
+  static size_t mspace_avail(mspace msp);
+  static size_t mspace_used(mspace msp);
+};
+
+template <typename HM_T>
+class DLAllocatorStrategy : public ShmemAllocatorStrategy {
+
+public:
+  /**
+   * @brief Required for default construction of other objects
+   *
+   * @note Not intended for direct usage.
+   */
+  DLAllocatorStrategy() = default;
+
+  /**
+   * @brief Primary constructor type
+   *
+   * Construct the dlmalloc mspace
+   *
+   * @param[in] Raw pointer to heap memory type
+   */
+  explicit DLAllocatorStrategy(HM_T* heap_mem) {
+    mspace_ = DLMalloc::create_mspace_with_base(heap_mem->get_ptr(), heap_mem->get_size(), false);
+  }
+
+  /**
+   * @brief Primary destructor
+   *
+   * Destroy the dlmalloc mspace
+   */
+  ~DLAllocatorStrategy() {
+    if(mspace_) {
+      DLMalloc::destroy_mspace(mspace_);
+    }
+  }
+
+  /**
+   * @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;
+    }
+    *ptr = static_cast<char*>(DLMalloc::mspace_malloc(mspace_, 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
+   *
+   * @note Not implemented
+   */
+  __device__ void alloc([[maybe_unused]] char** ptr,
+                        [[maybe_unused]] size_t request_size) override {}
+
+  /**
+   * @brief Frees memory from the heap
+   *
+   * Released memory is tracked by bookkeeping structures within this class.
+   *
+   * @param[in] Raw pointer to heap memory
+   *
+   */
+  void free(char* ptr) override {
+    DLMalloc::mspace_free(mspace_, ptr);
+  }
+
+  /**
+   * @brief Frees memory from the heap
+   *
+   * Released memory is tracked by bookkeeping structures within this class.
+   *
+   * @param[in] Raw pointer to heap memory
+   *
+   * @note Not implemented
+   */
+  __device__ void free([[maybe_unused]] char* ptr) override {}
+
+  /**
+   * @brief Used heap memory
+   *
+   * @return memory size
+   *
+   * @note The used size may be larger than the sum of the user allocation sizes
+   * (due to chunk tracking overhead and alignment).
+   *
+   */
+  size_t get_used() override {
+    size_t size{0};
+    size = DLMalloc::mspace_used(mspace_);
+    return size;
+  }
+
+  /**
+   * @brief Available heap memory
+   *
+   * @return memory size
+   *
+   * @note The available size may be smaller than the total heap size minus the sum
+   * of user allocation sizes (due to chunk tracking overhead and alignment).
+   */
+  size_t get_avail() {
+    size_t size{0};
+    size = DLMalloc::mspace_avail(mspace_);
+    return size;
+  }
+
+ private:
+  DLMalloc::mspace mspace_{nullptr};
+};
+
+
+}  // namespace rocshmem
+
+#endif  // LIBRARY_SRC_MEMORY_DLMALLOC_HPP_

src/memory/pow2_bins.hpp

@@ -164,6 +164,17 @@ class Pow2Bins : public ShmemAllocatorStrategy {
    */
   __device__ void free([[maybe_unused]] char* ptr) override {}
 
+  /**
+   * @brief Used memory from the heap
+   *
+   * Sum of all proffered_ memory sizes
+   *
+   * @return memory size
+   */
+  size_t get_used() override {
+    return amount_proffered();
+  }
+
   /**
    * @brief Sum of all proffered_ memory sizes
    *

src/memory/shmem_allocator_strategy.hpp

@@ -85,6 +85,17 @@ class ShmemAllocatorStrategy {
    * @param[in] Raw pointer to symmetric heap memory
    */
   __device__ virtual void free(char* ptr) = 0;
+
+  /**
+   * @brief Used heap memory
+   *
+   * @return memory size
+   *
+   * @note The used size may be larger than the sum of the user allocation sizes
+   * (due to chunk tracking overhead and alignment).
+   *
+   */
+  virtual size_t get_used() = 0;
 };
 
 }  // namespace rocshmem

src/memory/single_heap.cpp

@@ -61,7 +61,7 @@ char* SingleHeap::get_base_ptr() { return heap_mem_.get_ptr(); }
 
 size_t SingleHeap::get_size() { return heap_mem_.get_size(); }
 
-size_t SingleHeap::get_used() { return strat_.amount_proffered(); }
+size_t SingleHeap::get_used() { return strat_.get_used(); }
 
 size_t SingleHeap::get_avail() { return get_size() - get_used(); }
 

src/memory/single_heap.hpp

@@ -25,10 +25,16 @@
 #ifndef LIBRARY_SRC_MEMORY_SINGLE_HEAP_HPP_
 #define LIBRARY_SRC_MEMORY_SINGLE_HEAP_HPP_
 
-#include "address_record.hpp"
 #include "heap_memory.hpp"
 #include "heap_type.hpp"
+#if defined USE_ALLOC_DLMALLOC
+#include "dlmalloc.hpp"
+#elif defined USE_ALLOC_POW2BINS
+#include "address_record.hpp"
 #include "pow2_bins.hpp"
+#else
+#error "You need to have one of USE_ALLOC_DLMALLOC, USE_ALLOC_POW2BINS set to ON"
+#endif
 
 /**
  * @file single_heap.hpp
@@ -42,15 +48,21 @@
 namespace rocshmem {
 
 class SingleHeap {
+#if defined USE_ALLOC_DLMALLOC
+  /**
+   * @brief Helper type for allocation strategy
+   */
+  using STRAT_T = DLAllocatorStrategy<HEAP_T>;
+#elif defined USE_ALLOC_POW2BINS
   /**
    * @brief Helper type for address records
    */
   using AR_T = AddressRecord;
-
   /**
    * @brief Helper type for allocation strategy
    */
   using STRAT_T = Pow2Bins<AR_T, HEAP_T>;
+#endif // defined USE_ALLOC_POW2BINS
 
  public:
   /**

src/reverse_offload/backend_ro.cpp

@@ -30,7 +30,6 @@
 
 #include <cstdio>
 #include <cstdlib>
-#include <iostream>
 #include <memory>
 #include <thread>  // NOLINT
 

src/reverse_offload/mpi_transport.cpp

@@ -53,7 +53,7 @@ MPITransport::MPITransport(MPI_Comm comm, Queue* q)
   if (!init_done) {
     NET_CHECK(MPI_Init_thread(0, 0, MPI_THREAD_MULTIPLE, &provided));
     if (provided != MPI_THREAD_MULTIPLE) {
-      std::cerr << "MPI_THREAD_MULTIPLE support disabled.\n";
+      fprintf(stderr, "MPI_THREAD_MULTIPLE support disabled.\n");
     }
   }
   assert(comm != MPI_COMM_NULL);

tests/functional_tests/ping_all_tester.cpp

@@ -47,8 +47,8 @@ __global__ void PingAllTest(int loop, int skip, long long int *start_time,
     status[j] = 0;
   }
 
-  if (hipThreadIdx_x == 0) {
-    auto blk_pe_off {hipBlockIdx_x * num_pe};
+  if (is_thread_zero_in_block()) {
+    auto blk_pe_off {wg_id * num_pe};
 
     for (int i = 0; i < loop + skip; i++) {
       if (i == skip) {
@@ -61,6 +61,7 @@ __global__ void PingAllTest(int loop, int skip, long long int *start_time,
       rocshmem_int_wait_until_all(&r_buf[blk_pe_off], num_pe, status, ROCSHMEM_CMP_EQ, 1);
     }
     end_time[wg_id] = wall_clock64();
+    rocshmem_ctx_quiet(ctx);
   }
   rocshmem_wg_ctx_destroy(&ctx);
   rocshmem_wg_finalize();
@@ -71,14 +72,14 @@ __global__ void PingAllTest(int loop, int skip, long long int *start_time,
  *****************************************************************************/
 PingAllTester::PingAllTester(TesterArguments args) : Tester(args) {
   int num_pes {rocshmem_n_pes()};
-  r_buf = (int *)rocshmem_malloc(sizeof(int) * args.wg_size * num_pes);
+  r_buf = (int *)rocshmem_malloc(sizeof(int) * args.num_wgs * num_pes);
 }
 
 PingAllTester::~PingAllTester() { rocshmem_free(r_buf); }
 
 void PingAllTester::resetBuffers(uint64_t size) {
   int num_pes {rocshmem_n_pes()};
-  memset(r_buf, 0, sizeof(int) * args.wg_size * num_pes);
+  memset(r_buf, 0, sizeof(int) * args.num_wgs * num_pes);
 }
 
 void PingAllTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,

tests/functional_tests/ping_pong_tester.cpp

@@ -42,7 +42,7 @@ __global__ void PingPongTest(int loop, int skip, long long int *start_time,
 
   int pe = rocshmem_ctx_my_pe(ctx);
 
-  if (hipThreadIdx_x == 0) {
+  if (is_thread_zero_in_block()) {
 
     for (int i = 0; i < loop + skip; i++) {
       if (i == skip) {
@@ -60,7 +60,10 @@ __global__ void PingPongTest(int loop, int skip, long long int *start_time,
       }
     }
     end_time[wg_id] = wall_clock64();
+
+    rocshmem_ctx_quiet(ctx);
   }
+
   rocshmem_wg_ctx_destroy(&ctx);
   rocshmem_wg_finalize();
 }
@@ -69,13 +72,13 @@ __global__ void PingPongTest(int loop, int skip, long long int *start_time,
  * HOST TESTER CLASS METHODS
  *****************************************************************************/
 PingPongTester::PingPongTester(TesterArguments args) : Tester(args) {
-  r_buf = (int *)rocshmem_malloc(sizeof(int) * args.wg_size);
+  r_buf = (int *)rocshmem_malloc(sizeof(int) * args.num_wgs);
 }
 
 PingPongTester::~PingPongTester() { rocshmem_free(r_buf); }
 
 void PingPongTester::resetBuffers(uint64_t size) {
-  memset(r_buf, 0, sizeof(int) * args.wg_size);
+  memset(r_buf, 0, sizeof(int) * args.num_wgs);
 }
 
 void PingPongTester::launchKernel(dim3 gridSize, dim3 blockSize, int loop,

tests/unit_tests/CMakeLists.txt

@@ -84,6 +84,7 @@ target_sources(
     single_heap_gtest.cpp
     symmetric_heap_gtest.cpp
     pow2_bins_gtest.cpp
+    dlmalloc_gtest.cpp
     remote_heap_info_gtest.cpp
     mpi_init_singleton_gtest.cpp
     abql_block_mutex_gtest.cpp

tests/unit_tests/dlmalloc_gtest.cpp

@@ -0,0 +1,130 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 "dlmalloc_gtest.hpp"
+#include <cstdint>
+
+using namespace rocshmem;
+
+// a small portion of the heap is not available due to cost of dlmalloc bookkeeping
+#define DLMALLOC_BOOKKEEPING static_cast<size_t>(128 * ALIGNMENT)
+
+TEST_F(DLMallocTestFixture, used_0_bytes) {
+  size_t heap_size{1 << 30};
+  ASSERT_LE(strat_.get_used(), DLMALLOC_BOOKKEEPING);
+  ASSERT_EQ(strat_.get_used() + strat_.get_avail(), heap_size);
+}
+
+TEST_F(DLMallocTestFixture, alloc_0_bytes) {
+  size_t initial_used{strat_.get_used()};
+  char* c_ptr{nullptr};
+  size_t size{0};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_EQ(c_ptr, nullptr);
+  ASSERT_EQ(strat_.get_used(), initial_used);
+}
+
+TEST_F(DLMallocTestFixture, alloc_1_byte) {
+  char* c_ptr{nullptr};
+  size_t size{1};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_128_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{128};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_256_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{256};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_512_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{512};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_513_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{513};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_4KB) {
+  char* c_ptr{nullptr};
+  size_t size{4096};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_128KB) {
+  char* c_ptr{nullptr};
+  size_t size{131072};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+}
+
+TEST_F(DLMallocTestFixture, alloc_1GB) {
+  char* c_ptr{nullptr};
+  size_t heap_size{1 << 30};
+  size_t size{heap_size - DLMALLOC_BOOKKEEPING};
+
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+  ASSERT_EQ(strat_.get_used() + strat_.get_avail(), heap_size);
+}
+
+TEST_F(DLMallocTestFixture, alloc_256_bytes_X2_free_256_bytes_X2) {
+  char* c_ptr_1{nullptr};
+  char* c_ptr_2{nullptr};
+  size_t size{256};
+  strat_.alloc(&c_ptr_1, size);
+  ASSERT_NE(c_ptr_1, nullptr);
+  strat_.alloc(&c_ptr_2, size);
+  ASSERT_NE(c_ptr_2, nullptr);
+
+  strat_.free(c_ptr_1);
+  strat_.free(c_ptr_2);
+}
+
+TEST_F(DLMallocTestFixture, alloc_1GB_free_1GB) {
+  char* c_ptr{nullptr};
+  size_t heap_size{1 << 30};
+  size_t size{heap_size - DLMALLOC_BOOKKEEPING};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+  ASSERT_EQ(strat_.get_used() + strat_.get_avail(), heap_size);
+
+  strat_.free(c_ptr);
+}
+

tests/unit_tests/dlmalloc_gtest.hpp

@@ -0,0 +1,62 @@
+/******************************************************************************
+ * Copyright (c) Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * 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 ROCSHMEM_DLMALLOC_GTEST_HPP
+#define ROCSHMEM_DLMALLOC_GTEST_HPP
+
+#include "gtest/gtest.h"
+
+#include "../src/memory/heap_memory.hpp"
+#include "../src/memory/hip_allocator.hpp"
+#include "../src/memory/dlmalloc.hpp"
+
+namespace rocshmem {
+
+class DLMallocTestFixture : public ::testing::Test
+{
+    /**
+     * @brief Helper type for heap memory
+     */
+    using HEAP_T = HeapMemory<HIPAllocator>;
+
+    /**
+     * @brief Helper type for allocation strategy
+     */
+    using STRAT_T = DLAllocatorStrategy<HEAP_T>;
+
+  protected:
+    /**
+     * @brief Heap memory object
+     */
+    HEAP_T heap_mem_ {};
+
+    /**
+     * @brief Allocation strategy object
+     */
+    STRAT_T strat_ {&heap_mem_};
+};
+
+} // namespace rocshmem
+
+#endif // ROCSHMEM_DLMALLOC_GTEST_HPP

tests/unit_tests/pow2_bins_gtest.cpp

@@ -26,35 +26,47 @@
 
 using namespace rocshmem;
 
+TEST_F(Pow2BinsTestFixture, used_0_bytes) {
+  ASSERT_EQ(strat_.get_used(), 0);
+}
+
 TEST_F(Pow2BinsTestFixture, alloc_0_bytes) {
   char* c_ptr{nullptr};
   size_t size{0};
+  size_t expected_used{0};
   strat_.alloc(&c_ptr, size);
   ASSERT_EQ(c_ptr, nullptr);
+  ASSERT_EQ(strat_.get_used(), expected_used);
 }
 
 TEST_F(Pow2BinsTestFixture, alloc_1_byte) {
   char* c_ptr{nullptr};
   size_t size{1};
+  size_t align_size{ALIGNMENT * (1 + (size - 1) / ALIGNMENT)};
   strat_.alloc(&c_ptr, size);
   ASSERT_NE(c_ptr, nullptr);
 
   size_t min_size{256};
+  ASSERT_LE(align_size, 256); // test fixture won't work for larger values
   auto bins{strat_.get_bins()};
   auto bin{(*bins)[min_size]};
   ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), align_size);
 }
 
 TEST_F(Pow2BinsTestFixture, alloc_128_bytes) {
   char* c_ptr{nullptr};
   size_t size{128};
+  size_t align_size{ALIGNMENT * (1 + (size - 1) / ALIGNMENT)};
   strat_.alloc(&c_ptr, size);
   ASSERT_NE(c_ptr, nullptr);
 
   size_t min_size{256};
+  ASSERT_LE(align_size, 256);
   auto bins{strat_.get_bins()};
   auto bin{(*bins)[min_size]};
   ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), align_size);
 }
 
 TEST_F(Pow2BinsTestFixture, alloc_256_bytes) {
@@ -66,6 +78,7 @@ TEST_F(Pow2BinsTestFixture, alloc_256_bytes) {
   auto bins{strat_.get_bins()};
   auto bin{(*bins)[size]};
   ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), size);
 }
 
 TEST_F(Pow2BinsTestFixture, alloc_512_bytes) {
@@ -77,6 +90,7 @@ TEST_F(Pow2BinsTestFixture, alloc_512_bytes) {
   auto bins{strat_.get_bins()};
   auto bin{(*bins)[size]};
   ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), size);
 }
 
 TEST_F(Pow2BinsTestFixture, alloc_513_bytes) {
@@ -89,6 +103,20 @@ TEST_F(Pow2BinsTestFixture, alloc_513_bytes) {
   auto bins{strat_.get_bins()};
   auto bin{(*bins)[min_size]};
   ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), min_size);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_4095_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{4095};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+
+  size_t min_size{4096};
+  auto bins{strat_.get_bins()};
+  auto bin{(*bins)[min_size]};
+  ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), min_size);
 }
 
 TEST_F(Pow2BinsTestFixture, alloc_4KB) {
@@ -100,6 +128,20 @@ TEST_F(Pow2BinsTestFixture, alloc_4KB) {
   auto bins{strat_.get_bins()};
   auto bin{(*bins)[size]};
   ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), size);
+}
+
+TEST_F(Pow2BinsTestFixture, alloc_4097_bytes) {
+  char* c_ptr{nullptr};
+  size_t size{4097};
+  strat_.alloc(&c_ptr, size);
+  ASSERT_NE(c_ptr, nullptr);
+
+  size_t min_size{8192};
+  auto bins{strat_.get_bins()};
+  auto bin{(*bins)[min_size]};
+  ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), min_size);
 }
 
 TEST_F(Pow2BinsTestFixture, alloc_128KB) {
@@ -111,6 +153,7 @@ TEST_F(Pow2BinsTestFixture, alloc_128KB) {
   auto bins{strat_.get_bins()};
   auto bin{(*bins)[size]};
   ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), size);
 }
 
 TEST_F(Pow2BinsTestFixture, alloc_1GB) {
@@ -122,6 +165,7 @@ TEST_F(Pow2BinsTestFixture, alloc_1GB) {
   auto bins{strat_.get_bins()};
   auto bin{(*bins)[size]};
   ASSERT_EQ(bin.size(), 0);
+  ASSERT_EQ(strat_.get_used(), size);
 }
 
 TEST_F(Pow2BinsTestFixture, alloc_256_bytes_X2_free_256_bytes_X2) {
@@ -136,11 +180,13 @@ TEST_F(Pow2BinsTestFixture, alloc_256_bytes_X2_free_256_bytes_X2) {
   auto bins{strat_.get_bins()};
   auto& bin{(*bins)[size]};
   ASSERT_EQ(bin.size(), 0);
+  ASSERT_EQ(strat_.get_used(), 2 * size);
 
   strat_.free(c_ptr_1);
   strat_.free(c_ptr_2);
 
   ASSERT_EQ(bin.size(), 2);
+  ASSERT_EQ(strat_.get_used(), 0);
 }
 
 TEST_F(Pow2BinsTestFixture, alloc_1GB_free_1GB) {
@@ -152,8 +198,10 @@ TEST_F(Pow2BinsTestFixture, alloc_1GB_free_1GB) {
   auto bins{strat_.get_bins()};
   auto& bin{(*bins)[size]};
   ASSERT_EQ(bin.size(), 0);
+  ASSERT_EQ(strat_.get_used(), size);
 
   strat_.free(c_ptr);
 
   ASSERT_EQ(bin.size(), 1);
+  ASSERT_EQ(strat_.get_used(), 0);
 }

tests/unit_tests/single_heap_gtest.cpp

@@ -30,115 +30,157 @@ TEST_F(SingleHeapTestFixture, unallocated_size_check) {
   ASSERT_EQ(single_heap_.get_size(), 1 << 30);
 }
 
-TEST_F(SingleHeapTestFixture, unallocated_avail_check) {
-  ASSERT_EQ(single_heap_.get_avail(), 1 << 30);
-}
-
-TEST_F(SingleHeapTestFixture, unallocated_used_check) {
-  ASSERT_EQ(single_heap_.get_used(), 0);
-}
-
 TEST_F(SingleHeapTestFixture, free_null) {
   void* ptr{nullptr};
   single_heap_.free(ptr);
 }
 
 TEST_F(SingleHeapTestFixture, alloc_0) {
+  // some allocators (e.g. dlmalloc) use memory for internal bookkeeping
+  size_t initial_used{single_heap_.get_used()};
   size_t request_size{0};
   void* ptr{nullptr};
 
   single_heap_.malloc(&ptr, request_size);
   ASSERT_EQ(ptr, nullptr);
-
-  size_t expected_used{request_size};
-  ASSERT_EQ(single_heap_.get_used(), expected_used);
-  size_t expected_avail{single_heap_.get_size() - expected_used};
-  ASSERT_EQ(single_heap_.get_avail(), expected_avail);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
 
   single_heap_.free(ptr);
-
-  expected_used = 0;
-  ASSERT_EQ(single_heap_.get_used(), expected_used);
-  expected_avail = single_heap_.get_size();
-  ASSERT_EQ(single_heap_.get_avail(), expected_avail);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
 }
 
 TEST_F(SingleHeapTestFixture, alloc_1) {
+  size_t initial_used{single_heap_.get_used()};
   size_t request_size{1};
   void* ptr{nullptr};
 
   single_heap_.malloc(&ptr, request_size);
   ASSERT_NE(ptr, nullptr);
-
-  size_t expected_used{128};
-  ASSERT_EQ(single_heap_.get_used(), expected_used);
-  size_t expected_avail{single_heap_.get_size() - expected_used};
-  ASSERT_EQ(single_heap_.get_avail(), expected_avail);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr) & (ALIGNMENT-1), 0);
 
   single_heap_.free(ptr);
-
-  expected_used = 0;
-  ASSERT_EQ(single_heap_.get_used(), expected_used);
-  expected_avail = single_heap_.get_size();
-  ASSERT_EQ(single_heap_.get_avail(), expected_avail);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
 }
 
 TEST_F(SingleHeapTestFixture, alloc_256) {
+  size_t initial_used{single_heap_.get_used()};
   size_t request_size{256};
   void* ptr{nullptr};
 
   single_heap_.malloc(&ptr, request_size);
   ASSERT_NE(ptr, nullptr);
-
-  size_t expected_used{request_size};
-  ASSERT_EQ(single_heap_.get_used(), expected_used);
-  size_t expected_avail{single_heap_.get_size() - expected_used};
-  ASSERT_EQ(single_heap_.get_avail(), expected_avail);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr) & (ALIGNMENT-1), 0);
 
   single_heap_.free(ptr);
-
-  expected_used = 0;
-  ASSERT_EQ(single_heap_.get_used(), expected_used);
-  expected_avail = single_heap_.get_size();
-  ASSERT_EQ(single_heap_.get_avail(), expected_avail);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
 }
 
 TEST_F(SingleHeapTestFixture, alloc_1024) {
+  size_t initial_used{single_heap_.get_used()};
   size_t request_size{1024};
   void* ptr{nullptr};
 
   single_heap_.malloc(&ptr, request_size);
   ASSERT_NE(ptr, nullptr);
-
-  size_t expected_used{request_size};
-  ASSERT_EQ(single_heap_.get_used(), expected_used);
-  size_t expected_avail{single_heap_.get_size() - expected_used};
-  ASSERT_EQ(single_heap_.get_avail(), expected_avail);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr) & (ALIGNMENT-1), 0);
 
   single_heap_.free(ptr);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+}
 
-  expected_used = 0;
-  ASSERT_EQ(single_heap_.get_used(), expected_used);
-  expected_avail = single_heap_.get_size();
-  ASSERT_EQ(single_heap_.get_avail(), expected_avail);
+TEST_F(SingleHeapTestFixture, alloc_1MB) {
+  size_t initial_used{single_heap_.get_used()};
+  size_t request_size{1 << 20};
+  void* ptr{nullptr};
+
+  single_heap_.malloc(&ptr, request_size);
+  ASSERT_NE(ptr, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
 }
 
 TEST_F(SingleHeapTestFixture, alloc_4097) {
+  size_t initial_used{single_heap_.get_used()};
   size_t request_size{4097};
   void* ptr{nullptr};
 
   single_heap_.malloc(&ptr, request_size);
   ASSERT_NE(ptr, nullptr);
-
-  size_t expected_used{8192};
-  ASSERT_EQ(single_heap_.get_used(), expected_used);
-  size_t expected_avail{single_heap_.get_size() - expected_used};
-  ASSERT_EQ(single_heap_.get_avail(), expected_avail);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr) & (ALIGNMENT-1), 0);
 
   single_heap_.free(ptr);
-
-  expected_used = 0;
-  ASSERT_EQ(single_heap_.get_used(), expected_used);
-  expected_avail = single_heap_.get_size();
-  ASSERT_EQ(single_heap_.get_avail(), expected_avail);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+}
+
+TEST_F(SingleHeapTestFixture, alloc_X2_8191) {
+  size_t initial_used{single_heap_.get_used()};
+  size_t request_size{8191};
+  void* ptr_1{nullptr};
+  void* ptr_2{nullptr};
+
+  single_heap_.malloc(&ptr_1, request_size);
+  ASSERT_NE(ptr_1, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_1) & (ALIGNMENT-1), 0);
+
+  single_heap_.malloc(&ptr_2, request_size);
+  ASSERT_NE(ptr_2, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_2) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr_1);
+  single_heap_.free(ptr_2);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+}
+
+TEST_F(SingleHeapTestFixture, alloc_X2_free_alloc_free_X2_1MB) {
+  size_t initial_used{single_heap_.get_used()};
+  void* ptr_1{nullptr};
+  void* ptr_2{nullptr};
+  void* ptr_3{nullptr};
+  size_t request_size{1 << 20};
+
+  single_heap_.malloc(&ptr_1, request_size);
+  ASSERT_NE(ptr_1, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_1) & (ALIGNMENT-1), 0);
+
+  single_heap_.malloc(&ptr_2, request_size);
+  ASSERT_NE(ptr_1, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_2) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr_1);
+
+  single_heap_.malloc(&ptr_3, request_size);
+  ASSERT_NE(ptr_3, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_3) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr_3);
+  single_heap_.free(ptr_2);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
+}
+
+TEST_F(SingleHeapTestFixture, alloc_X2_free_alloc_free_X2_63) {
+  size_t initial_used{single_heap_.get_used()};
+  void* ptr_1{nullptr};
+  void* ptr_2{nullptr};
+  void* ptr_3{nullptr};
+  size_t request_size{63};
+
+  single_heap_.malloc(&ptr_1, request_size);
+  ASSERT_NE(ptr_1, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_1) & (ALIGNMENT-1), 0);
+
+  single_heap_.malloc(&ptr_2, request_size);
+  ASSERT_NE(ptr_1, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_2) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr_1);
+
+  single_heap_.malloc(&ptr_3, request_size);
+  ASSERT_NE(ptr_3, nullptr);
+  ASSERT_EQ(reinterpret_cast<uintptr_t>(ptr_3) & (ALIGNMENT-1), 0);
+
+  single_heap_.free(ptr_3);
+  single_heap_.free(ptr_2);
+  ASSERT_EQ(single_heap_.get_used(), initial_used);
 }