VITIS-12792 HIP Binding: Memory pool APIs (#8268)

* add hipMemsetD32Async(). Signed-off-by: Chiming Zhang <[email protected]> * fix typo. Signed-off-by: Chiming Zhang <[email protected]> * Replace hard coded name with __func__. Signed-off-by: Chiming Zhang <[email protected]> * fix staled pointer error. Signed-off-by: Chiming Zhang <[email protected]> * fix double allocation issue. Signed-off-by: Chiming Zhang <[email protected]> * fix incorrect share_ptr use. Signed-off-by: Chiming Zhang <[email protected]> * add hipMemsetD32Async(). Signed-off-by: Chiming Zhang <[email protected]> * fix typo. Signed-off-by: Chiming Zhang <[email protected]> * Replace hard coded name with __func__. Signed-off-by: Chiming Zhang <[email protected]> * fix staled pointer error. Signed-off-by: Chiming Zhang <[email protected]> * fix double allocation issue. Signed-off-by: Chiming Zhang <[email protected]> * 1) add lock in memory_database::get_hip_mem_from_addr() 2) avoid crash from getting hip mem object from null pointer 3) allow setting 0 for kernel arguments Signed-off-by: Chiming Zhang <[email protected]> * fix incorrect share_ptr use. Signed-off-by: Chiming Zhang <[email protected]> * Fix rebase error. Signed-off-by: Chiming Zhang <[email protected]> * move to_hex() into core/common/utils.h Signed-off-by: Chiming Zhang <[email protected]> * remove extra member in class copy_buffer. Signed-off-by: Chiming Zhang <[email protected]> * fix a typo. Signed-off-by: Chiming Zhang <[email protected]> * fix a typo. Signed-off-by: Chiming Zhang <[email protected]> * fix some typo and remove template use from copy_buffer class implementation. Signed-off-by: Chiming Zhang <[email protected]> * change type of err_msg argument of helper function throw_if() to const char*. Signed-off-by: Chiming <[email protected]> * fix incorrect usage of shared_ptr in copy_buffer constructor. Signed-off-by: Chiming Zhang <[email protected]> * change the place where host_vec is moved. Signed-off-by: Chiming Zhang <[email protected]> * add back missing std::move in copy_from_host_buffer_commad constructor. Signed-off-by: Chiming Zhang <[email protected]> * Add initial implementation of hip stream ordered memory allocator. Signed-off-by: Chiming Zhang <[email protected]> * Fix the size alignment in memory pool allocator. Signed-off-by: Chiming Zhang <[email protected]> * fix rebase error. Signed-off-by: Chiming Zhang <[email protected]> * fix rebase error. Signed-off-by: Chiming Zhang <[email protected]> * fix rebase error. Signed-off-by: Chiming Zhang <[email protected]> * Add comment for the choice of shared_ptr vs unique_ptr in in enqueing the async memcpy commands. Signed-off-by: Chiming Zhang <[email protected]> * Add comment for using shared_ptr for storage of pointer to memory_pool pointers. Signed-off-by: Chiming Zhang <[email protected]> * use unique_ptr for device_cache. Signed-off-by: Chiming <[email protected]> * Fix issues raised in code review. Signed-off-by: Chiming <[email protected]> * fix the error in memory::write(). Signed-off-by: Chiming <[email protected]> * remove curly braces. Signed-off-by: Chiming <[email protected]> * Fix issues in hipMallocAsync() and hipFreeAsync(). Signed-off-by: Chiming <[email protected]> * fix error found in unit testing. Signed-off-by: Chiming <[email protected]> * use sub class of hip::memory for async allocation from hip memory pool. Signed-off-by: Chiming <[email protected]> * fix the error in sub mem lookup from memory_database. Signed-off-by: Chiming <[email protected]> * remove sub_mem address map. Signed-off-by: Chiming <[email protected]> * code clean up. Signed-off-by: Chiming <[email protected]> * add code for hipDeviceGetDefaultMemPool(), hipDeviceGetMemPool() and hipDeviceSetMemPool(). Signed-off-by: Chiming <[email protected]> * fix nullptr error. Signed-off-by: Chiming <[email protected]> * Fix compile error on Linux. Signed-off-by: Chiming Zhang <[email protected]> * Fix compile warning caused by using "int" type. Signed-off-by: Chiming Zhang <[email protected]> * fix compile error in release builds. Signed-off-by: Chiming Zhang <[email protected]> * Fix compile error. Signed-off-by: Chiming Zhang <[email protected]> --------- Signed-off-by: Chiming Zhang <[email protected]> Signed-off-by: Chiming <[email protected]> Co-authored-by: Chiming <[email protected]>
Xilinx · Oct 6, 2024 · 1733fa0 · 1733fa0
1 parent b7c2c09
commit 1733fa0
Show file tree

Hide file tree

Showing 17 changed files with 1,172 additions and 97 deletions.
diff --git a/src/runtime_src/core/common/api/handle.h b/src/runtime_src/core/common/api/handle.h
@@ -40,6 +40,12 @@ class handle_map<HandleType, std::shared_ptr<ImplType>>
   }
 
 public:
+  const std::map<HandleType, std::shared_ptr<ImplType>>&
+  get_map() const
+  {
+    return handles;
+  }
+
   const std::shared_ptr<ImplType>&
   get_or_error(HandleType handle) const
   {
@@ -105,6 +111,12 @@ class handle_map<HandleType, std::unique_ptr<ImplType>>
   std::map<HandleType, std::unique_ptr<ImplType>> handles;
 
 public:
+  const std::map<HandleType, std::unique_ptr<ImplType>>&
+  get_map() const
+  {
+    return handles;
+  }
+
   ImplType*
   get_or_error(HandleType handle) const
   {
@@ -153,6 +165,13 @@ class handle_map<HandleType, std::unique_ptr<ImplType>>
     std::lock_guard<std::mutex> lk(mutex);
     return handles.count(handle);
   }
+
+  size_t
+  size() const
+  {
+    std::lock_guard<std::mutex> lk(mutex);
+    return handles.size();
+  }
 };
 
 } // xrt_core
diff --git a/src/runtime_src/hip/api/hip_device.cpp b/src/runtime_src/hip/api/hip_device.cpp
@@ -5,6 +5,7 @@
 
 #include "hip/core/common.h"
 #include "hip/core/device.h"
+#include "hip/core/memory_pool.h"
 
 #include <cstring>
 #include <mutex>
@@ -46,8 +47,12 @@ device_init()
   for (uint32_t i = 0; i < dev_count; i++) {
     if (device_cache.count(static_cast<device_handle>(i)) > 0)
       continue;
-    auto dev = std::make_shared<xrt::core::hip::device>(i);
+    auto dev = std::make_unique<xrt::core::hip::device>(i);
     device_cache.add(i, std::move(dev));
+    auto default_mem_pool = std::make_shared<xrt::core::hip::memory_pool>(device_cache.get_or_error(i), MAX_MEMORY_POOL_SIZE_NPU, MEMORY_POOL_BLOCK_SIZE_NPU);
+    memory_pool_db[i].push_front(default_mem_pool);
+    current_memory_pool_db[i] = default_mem_pool;
+    insert_in_map(mem_pool_cache, default_mem_pool);
   }
   // make first device as default device
   if (dev_count > 0)