SWDEV-353281 - VM support in mempool for graphs

The change enables VM support in graphs on Windows. That allows
to avoid caching of all allocations at the cost of map/unmap
overhead during memory create/destroy.

Change-Id: I792be00fba099e5e5d3cd44a963e1dfd6976a86d

hipamd/src/hip_device.cpp

@@ -53,10 +53,12 @@ bool Device::Create() {
     return false;
   }
 
-  uint64_t max_size = std::numeric_limits<uint64_t>::max();
-  // Use maximum value to hold memory, because current implementation doesn't support VM
-  // Note: the call for the threshold is always successful
-  auto error = graph_mem_pool_->SetAttribute(hipMemPoolAttrReleaseThreshold, &max_size);
+  if (!HIP_MEM_POOL_USE_VM) {
+    uint64_t max_size = std::numeric_limits<uint64_t>::max();
+    // Use maximum value to hold memory, because current implementation doesn't support VM
+    // Note: the call for the threshold is always successful
+    auto error = graph_mem_pool_->SetAttribute(hipMemPoolAttrReleaseThreshold, &max_size);
+  }
 
   // Current is default pool after device creation
   current_mem_pool_ = default_mem_pool_;

hipamd/src/hip_graph.cpp

@@ -884,8 +884,8 @@ hipError_t capturehipMallocAsync(hipStream_t stream, hipMemPool_t mem_pool,
   if (status != hipSuccess) {
     return status;
   }
-  // Execute the node during capture, so runtime can return a valid device pointer
-  *dev_ptr = mem_alloc_node->Execute(s);
+  // Without VM runtime executes the node during capture, so it can return a valid device pointer
+  *dev_ptr = (HIP_MEM_POOL_USE_VM) ? mem_alloc_node->ReserveAddress() : mem_alloc_node->Execute(s);
   s->SetLastCapturedNode(mem_alloc_node);
 
   return hipSuccess;
@@ -900,8 +900,10 @@ hipError_t capturehipFreeAsync(hipStream_t stream, void* dev_ptr) {
   if (status != hipSuccess) {
     return status;
   }
-  // Execute the node during capture, so runtime can release memory into cache
-  mem_free_node->Execute(s);
+  // Execute the node without VM support, so runtime can release memory into cache
+  if (!HIP_MEM_POOL_USE_VM) {
+    mem_free_node->Execute(s);
+  }
   s->SetLastCapturedNode(mem_free_node);
   return hipSuccess;
 }
@@ -2202,7 +2204,8 @@ hipError_t hipGraphAddMemAllocNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
   *pGraphNode = mem_alloc_node;
   auto status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies);
   // The address must be provided during the node creation time
-  pNodeParams->dptr = mem_alloc_node->Execute();
+  pNodeParams->dptr =
+      (HIP_MEM_POOL_USE_VM) ? mem_alloc_node->ReserveAddress() : mem_alloc_node->Execute();
   HIP_RETURN(status);
 }
 
@@ -2231,9 +2234,15 @@ hipError_t hipGraphAddMemFreeNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
 
   // Is memory passed to be free'd valid
   size_t offset = 0;
-  amd::Memory* memory_object = getMemoryObject(dev_ptr, offset);
-  if (memory_object == nullptr) {
-    HIP_RETURN(hipErrorInvalidValue);
+  auto memory = getMemoryObject(dev_ptr, offset);
+  if (memory == nullptr) {
+    if (HIP_MEM_POOL_USE_VM) {
+      // When VM is on the address must be valid and may point to a VA object
+      memory = amd::MemObjMap::FindVirtualMemObj(dev_ptr);
+    }
+    if (memory == nullptr) {
+      HIP_RETURN(hipErrorInvalidValue);
+    }
   }
 
   auto mem_free_node = new hipGraphMemFreeNode(dev_ptr);

hipamd/src/hip_graph_internal.cpp

@@ -38,6 +38,8 @@ const char* GetGraphNodeTypeString(uint32_t op) {
     CASE_STRING(hipGraphNodeTypeEventRecord, EventRecordNode)
     CASE_STRING(hipGraphNodeTypeExtSemaphoreSignal, ExtSemaphoreSignalNode)
     CASE_STRING(hipGraphNodeTypeExtSemaphoreWait, ExtSemaphoreWaitNode)
+    CASE_STRING(hipGraphNodeTypeMemAlloc, MemAllocNode)
+    CASE_STRING(hipGraphNodeTypeMemFree, MemFreeNode)
     CASE_STRING(hipGraphNodeTypeMemcpyFromSymbol, MemcpyFromSymbolNode)
     CASE_STRING(hipGraphNodeTypeMemcpyToSymbol, MemcpyToSymbolNode)
     default:

hipamd/src/hip_graph_internal.hpp

@@ -1,4 +1,4 @@
-/* Copyright (c) 2021 - 2021 Advanced Micro Devices, Inc.
+/* Copyright (c) 2021 - 2023 Advanced Micro Devices, Inc.
 
  Permission is hereby granted, free of charge, to any person obtaining a copy
  of this software and associated documentation files (the "Software"), to deal
@@ -33,6 +33,7 @@
 #include "hip_event.hpp"
 #include "hip_platform.hpp"
 #include "hip_mempool_impl.hpp"
+#include "hip_vm.hpp"
 
 typedef hipGraphNode* Node;
 hipError_t FillCommands(std::vector<std::vector<Node>>& parallelLists,
@@ -508,6 +509,37 @@ struct ihipGraph {
     return ptr;
   }
 
+  void* ReserveAddress(size_t size) const {
+    void* startAddress = nullptr;
+    void* ptr;
+    for (auto& dev : g_devices) {
+      const auto& dev_info = dev->devices()[0]->info();
+      ptr = dev->devices()[0]->virtualAlloc(startAddress, size,
+          dev_info.virtualMemAllocGranularity_);
+
+      // if addr==0 then runtime will use the first VA on other devices
+      if (startAddress == nullptr) {
+        startAddress = ptr;
+      } else if (ptr != startAddress) {
+        // if runtime cannot reserve the same VA on other devices, just fail
+        for (auto& d : g_devices) {
+          if (d == dev) {
+            d->devices()[0]->virtualFree(ptr);
+            return nullptr;
+          }
+          d->devices()[0]->virtualFree(startAddress);
+        }
+      }
+    }
+    return ptr;
+  }
+
+  void FreeAddress(void* ptr) const {
+    for (auto& dev : g_devices) {
+      dev->devices()[0]->virtualFree(ptr);
+    }
+  }
+
   void FreeMemory(void* dev_ptr, hip::Stream* stream) const {
     size_t offset = 0;
     auto memory = getMemoryObject(dev_ptr, offset);
@@ -766,7 +798,7 @@ class hipGraphKernelNode : public hipGraphNode {
       sprintf(buffer,
               "{\n%s\n| {ID | %d | %s\\<\\<\\<(%u,%u,%u),(%u,%u,%u),%u\\>\\>\\>}\n| {{node "
               "handle | func handle} | {%p | %p}}\n| {accessPolicyWindow | {base_ptr | num_bytes | "
-              "hitRatio | hitProp | missProp} | {%p | %ld | %f | %d | %d}}\n| {cooperative | "
+              "hitRatio | hitProp | missProp} | {%p | %zu | %f | %d | %d}}\n| {cooperative | "
               "%u}\n| {priority | 0}\n}",
               label_.c_str(), GetID(), function->name().c_str(), pKernelParams_->gridDim.x,
               pKernelParams_->gridDim.y, pKernelParams_->gridDim.z, pKernelParams_->blockDim.x,
@@ -781,7 +813,7 @@ class hipGraphKernelNode : public hipGraphNode {
       sprintf(buffer,
               "{\n%s\n| {ID | %d | %s}\n"
               "| {accessPolicyWindow | {base_ptr | num_bytes | "
-              "hitRatio | hitProp | missProp} |\n| {%p | %ld | %f | %d | %d}}\n| {cooperative | "
+              "hitRatio | hitProp | missProp} |\n| {%p | %zu | %f | %d | %d}}\n| {cooperative | "
               "%u}\n| {priority | 0}\n}",
               label_.c_str(), GetID(), function->name().c_str(),
               kernelAttr_.accessPolicyWindow.base_ptr, kernelAttr_.accessPolicyWindow.num_bytes,
@@ -1927,26 +1959,125 @@ class hipGraphEmptyNode : public hipGraphNode {
   }
 };
 
+// ================================================================================================
 class hipGraphMemAllocNode : public hipGraphNode {
   hipMemAllocNodeParams node_params_; // Node parameters for memory allocation
+  amd::Memory* va_ = nullptr;         // Memory object, which holds a virtual address
+
+  // Derive the new class for VirtualMapCommand,
+  // so runtime can allocate memory during the execution of command
+  class VirtualMemAllocNode : public amd::VirtualMapCommand {
+   public:
+    VirtualMemAllocNode(amd::HostQueue& queue, const amd::Event::EventWaitList& eventWaitList,
+                        amd::Memory* va, size_t size, amd::Memory* memory, ihipGraph* graph)
+        : VirtualMapCommand(queue, eventWaitList, va->getSvmPtr(), size, memory),
+          va_(va), graph_(graph) {}
+
+    virtual void submit(device::VirtualDevice& device) final {
+      // Remove VA reference from the global mapping. Runtime has to keep a dummy reference for
+      // validation logic during the capture or creation of the nodes
+      amd::MemObjMap::RemoveMemObj(va_->getSvmPtr());
+      // Allocate real memory for mapping
+      const auto& dev_info = queue()->device().info();
+      auto aligned_size = amd::alignUp(size_, dev_info.virtualMemAllocGranularity_);
+      auto dptr = graph_->AllocateMemory(aligned_size, static_cast<hip::Stream*>(queue()), nullptr);
+      if (dptr == nullptr) {
+        setStatus(CL_INVALID_OPERATION);
+        return;
+      }
+      size_t offset = 0;
+      // Get memory object associated with the real allocation
+      memory_ = getMemoryObject(dptr, offset);
+      // Retain memory object because command release will release it
+      memory_->retain();
+      size_ = aligned_size;
+      // Save geenric allocation info to match VM interfaces
+      memory_->getUserData().data = new hip::MemMapAllocUserData(dptr, aligned_size, va_);
+      // Execute the original mapping command
+      VirtualMapCommand::submit(device);
+      // Update the internal svm address to ptr
+      memory()->setSvmPtr(va_->getSvmPtr());
+      // Can't destroy VA, because it's used in mapping even if the node will be destroyed
+      va_->retain();
+      ClPrint(amd::LOG_INFO, amd::LOG_MEM_POOL, "Graph MemAlloc execute: %p, %p",
+          va_->getSvmPtr(), memory());
+    }
+
+   private:
+    amd::Memory* va_;   // Memory object with the new virtual address for mapping
+    ihipGraph* graph_;  // Graph which allocates/maps memory
+  };
 
  public:
   hipGraphMemAllocNode(const hipMemAllocNodeParams* node_params)
       : hipGraphNode(hipGraphNodeTypeMemAlloc, "solid", "rectangle", "MEM_ALLOC") {
-        node_params_ = *node_params;
-      }
-  ~hipGraphMemAllocNode() {}
+    node_params_ = *node_params;
+  }
 
-  hipGraphNode* clone() const {
+  hipGraphMemAllocNode(const hipGraphMemAllocNode& rhs)
+      : hipGraphNode(rhs) {
+    node_params_ = rhs.node_params_;
+    if (HIP_MEM_POOL_USE_VM) {
+      assert(rhs.va_ != nullptr && "Graph MemAlloc runtime can't clone an invalid node!");
+      va_ = rhs.va_;
+      va_->retain();
+    }
+  }
+
+  virtual ~hipGraphMemAllocNode() final {
+    if (va_ != nullptr) {
+      va_->release();
+    }
+  }
+
+  virtual hipGraphNode* clone() const final {
     return new hipGraphMemAllocNode(static_cast<hipGraphMemAllocNode const&>(*this));
   }
 
-  virtual hipError_t CreateCommand(hip::Stream* stream) {
+  virtual hipError_t CreateCommand(hip::Stream* stream) final {
     auto error = hipGraphNode::CreateCommand(stream);
-    auto ptr = Execute(stream_);
+    if (!HIP_MEM_POOL_USE_VM) {
+      auto ptr = Execute(stream_);
+    } else {
+      auto graph = GetParentGraph();
+      if (graph != nullptr) {
+        assert(va_ != nullptr && "Runtime can't create a command for an invalid node!");
+        // Create command for memory mapping
+        auto cmd = new VirtualMemAllocNode(*stream, amd::Event::EventWaitList{},
+            va_, node_params_.bytesize, nullptr, graph);
+        commands_.push_back(cmd);
+        size_t offset = 0;
+        // Check if memory was already added after first reserve
+        if (getMemoryObject(node_params_.dptr, offset) == nullptr) {
+          // Map VA in the accessible space because the graph execution still has
+          // pointers validation and must find a valid object
+          // @note: Memory can be released outside of the graph and
+          // runtime can't keep a valid mapping since it doesn't know if the graph will
+          // be executed again
+          amd::MemObjMap::AddMemObj(node_params_.dptr, va_);
+        }
+        ClPrint(amd::LOG_INFO, amd::LOG_MEM_POOL, "Graph MemAlloc create: %p",
+            node_params_.dptr);
+      }
+    }
     return error;
   }
 
+  void* ReserveAddress() {
+    auto graph = GetParentGraph();
+    if (graph != nullptr) {
+      node_params_.dptr = graph->ReserveAddress(node_params_.bytesize);
+      if (node_params_.dptr != nullptr) {
+        // Find VA and map in the accessible space so capture can find a valid object
+        va_ = amd::MemObjMap::FindVirtualMemObj(node_params_.dptr);
+        amd::MemObjMap::AddMemObj(node_params_.dptr, va_);
+      }
+      ClPrint(amd::LOG_INFO, amd::LOG_MEM_POOL, "Graph MemAlloc reserve VA: %p",
+          node_params_.dptr);
+    }
+    return node_params_.dptr;
+  }
+
   void* Execute(hip::Stream* stream = nullptr) {
     auto graph = GetParentGraph();
     if (graph != nullptr) {
@@ -1969,28 +2100,80 @@ class hipGraphMemAllocNode : public hipGraphNode {
     return graph->ProbeMemory(node_params_.dptr);
   }
 
-
   void GetParams(hipMemAllocNodeParams* params) const {
     std::memcpy(params, &node_params_, sizeof(hipMemAllocNodeParams));
   }
 };
 
+// ================================================================================================
 class hipGraphMemFreeNode : public hipGraphNode {
   void* device_ptr_;    // Device pointer of the freed memory
 
+  // Derive the new class for VirtualMap command, since runtime has to free
+  // real allocation after unmap is complete
+  class VirtualMemFreeNode : public amd::VirtualMapCommand {
+   public:
+    VirtualMemFreeNode(ihipGraph* graph, int device_id, amd::HostQueue& queue,
+        const amd::Event::EventWaitList& eventWaitList, void* ptr, size_t size,
+        amd::Memory* memory) : VirtualMapCommand(queue, eventWaitList, ptr, size, memory)
+        , graph_(graph), device_id_(device_id) {}
+
+    virtual void submit(device::VirtualDevice& device) final {
+      // Find memory object before unmap logic
+      auto alloc = amd::MemObjMap::FindMemObj(ptr());
+      VirtualMapCommand::submit(device);
+      // Restore the original address of the generic allocation
+      auto ga = reinterpret_cast<hip::MemMapAllocUserData*>(alloc->getUserData().data);
+      alloc->setSvmPtr(ga->ptr_);
+      if (!AMD_DIRECT_DISPATCH) {
+        // Update the current device, since hip event, used in mem pools, requires device
+        hip::setCurrentDevice(device_id_);
+      }
+      // Free virtual address
+      ga->va_->release();
+      alloc->getUserData().data = nullptr;
+      // Release the allocation back to graph's pool
+      graph_->FreeMemory(ga->ptr_, static_cast<hip::Stream*>(queue()));
+      amd::MemObjMap::AddMemObj(ptr(), ga->va_);
+      delete ga;
+      ClPrint(amd::LOG_INFO, amd::LOG_MEM_POOL, "Graph MemFree execute: %p, %p",
+          ptr(), alloc);
+    }
+
+   private:
+    ihipGraph* graph_;  // Graph, which has the execution of this command
+    int device_id_;     // Device ID where this command is executed
+  };
+
  public:
   hipGraphMemFreeNode(void* dptr)
     : hipGraphNode(hipGraphNodeTypeMemFree, "solid", "rectangle", "MEM_FREE")
     , device_ptr_(dptr) {}
-  ~hipGraphMemFreeNode() {}
+  hipGraphMemFreeNode(const hipGraphMemFreeNode& rhs) : hipGraphNode(rhs) {
+    device_ptr_ = rhs.device_ptr_;
+  }
 
-  hipGraphNode* clone() const {
+  virtual hipGraphNode* clone() const final {
     return new hipGraphMemFreeNode(static_cast<hipGraphMemFreeNode const&>(*this));
   }
 
-  virtual hipError_t CreateCommand(hip::Stream* stream) {
+  virtual hipError_t CreateCommand(hip::Stream* stream) final {
     auto error = hipGraphNode::CreateCommand(stream);
-    Execute(stream_);
+    if (!HIP_MEM_POOL_USE_VM) {
+      Execute(stream_);
+    } else {
+      auto graph = GetParentGraph();
+      if (graph != nullptr) {
+        const auto& dev_info = stream->device().info();
+        auto va = amd::MemObjMap::FindVirtualMemObj(device_ptr_);
+        // Unmap virtual address from memory
+        amd::Command* cmd = new VirtualMemFreeNode(graph, stream->DeviceId(), *stream,
+            amd::Command::EventWaitList{}, device_ptr_,
+            amd::alignUp(va->getSize(), dev_info.virtualMemAllocGranularity_), nullptr);
+        commands_.push_back(cmd);
+        ClPrint(amd::LOG_INFO, amd::LOG_MEM_POOL, "Graph FreeMem create: %p", device_ptr_);
+      }
+    }
     return error;
   }
 

hipamd/src/hip_mempool.cpp

@@ -82,9 +82,44 @@ hipError_t hipMallocAsync(void** dev_ptr, size_t size, hipStream_t stream) {
   STREAM_CAPTURE(hipMallocAsync, stream, reinterpret_cast<hipMemPool_t>(mem_pool), size, dev_ptr);
 
   *dev_ptr = mem_pool->AllocateMemory(size, hip_stream);
+  if (*dev_ptr == nullptr) {
+    HIP_RETURN(hipErrorOutOfMemory);
+  }
   HIP_RETURN(hipSuccess);
 }
 
+// ================================================================================================
+// @note: Runtime needs the new command for MT path, since the app can execute hipFreeAsync()
+// before the graph execution is done. Hence there could be a race condition between
+// memory allocatiom in graph, which occurs in a worker thread, and host execution of hipFreeAsync
+class FreeAsyncCommand : public amd::Command {
+ private:
+  void* ptr_;     //!< Virtual address for asynchronious free
+
+ public:
+  FreeAsyncCommand(amd::HostQueue& queue, void* ptr)
+      : amd::Command(queue, 1, amd::Event::nullWaitList), ptr_(ptr) {}
+
+  virtual void submit(device::VirtualDevice& device) final {
+    size_t offset = 0;
+    auto memory = getMemoryObject(ptr_, offset);
+    if (memory != nullptr) {
+      auto id = memory->getUserData().deviceId;
+      if (!AMD_DIRECT_DISPATCH) {
+        // Required for HIP events
+        hip::setCurrentDevice(id);
+      }
+      if (!g_devices[id]->FreeMemory(memory, static_cast<hip::Stream*>(queue()))) {
+        // @note It's not the most optimal logic.
+        // The current implementation has unconditional waits
+        if (ihipFree(ptr_) != hipSuccess) {
+          setStatus(CL_INVALID_OPERATION);
+        }
+      }
+    }
+  }
+};
+
 // ================================================================================================
 hipError_t hipFreeAsync(void* dev_ptr, hipStream_t stream) {
   HIP_INIT_API(hipFreeAsync, dev_ptr, stream);
@@ -92,17 +127,15 @@ hipError_t hipFreeAsync(void* dev_ptr, hipStream_t stream) {
     HIP_RETURN(hipErrorInvalidValue);
   }
   STREAM_CAPTURE(hipFreeAsync, stream, dev_ptr);
-  size_t offset = 0;
-  auto memory = getMemoryObject(dev_ptr, offset);
-  if (memory != nullptr) {
-    auto id = memory->getUserData().deviceId;
-    auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream() :
-      reinterpret_cast<hip::Stream*>(stream);
-    if (!g_devices[id]->FreeMemory(memory, hip_stream)) {
-      //! @todo It's not the most optimal logic. The current implementation has unconditional waits
-      HIP_RETURN(ihipFree(dev_ptr));
-    }
+
+  auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream()
+                    : reinterpret_cast<hip::Stream*>(stream);
+  auto cmd = new FreeAsyncCommand(*hip_stream, dev_ptr);
+  if (cmd == nullptr) {
+    HIP_RETURN(hipErrorUnknown);
   }
+  cmd->enqueue();
+  cmd->release();
 
   HIP_RETURN(hipSuccess);
 }

hipamd/src/hip_mempool_impl.cpp

@@ -1,4 +1,4 @@
-/* Copyright (c) 2022 Advanced Micro Devices, Inc.
+/* Copyright (c) 2022-2023 Advanced Micro Devices, Inc.
 
  Permission is hereby granted, free of charge, to any person obtaining a copy
  of this software and associated documentation files (the "Software"), to deal
@@ -19,6 +19,8 @@
  THE SOFTWARE. */
 
 #include "hip_mempool_impl.hpp"
+#include "hip_vm.hpp"
+#include "platform/command.hpp"
 
 namespace hip {
 
@@ -40,10 +42,16 @@ void Heap::AddMemory(amd::Memory* memory, const MemoryTimestamp& ts) {
 amd::Memory* Heap::FindMemory(size_t size, hip::Stream* stream, bool opportunistic, void* dptr) {
   amd::Memory* memory = nullptr;
   for (auto it = allocations_.begin(); it != allocations_.end();) {
-    bool check_address = (dptr == nullptr) || (it->first->getSvmPtr() == dptr);
-    // Check if size can match and it's safe to use this resource
-    if ((it->first->getSize() >= size) && check_address &&
-       (it->second.IsSafeFind(stream, opportunistic))) {
+    bool check_address = (dptr == nullptr);
+    if (it->first->getSvmPtr() == dptr) {
+      // If the search is done for the specified address then runtime must wait
+      it->second.Wait();
+      check_address = true;
+    }
+    // Check if size can match and it's safe to use this resource.
+    // Runtime can accept an allocation with 12.5% on the size threshold
+    if ((it->first->getSize() >= size) && (it->first->getSize() <= (size / 8) * 9) &&
+        check_address && (it->second.IsSafeFind(stream, opportunistic))) {
       memory = it->first;
       total_size_ -= memory->getSize();
       // Remove found allocation from the map
@@ -197,6 +205,8 @@ void* MemoryPool::AllocateMemory(size_t size, hip::Stream* stream, void* dptr) {
   // Increment the reference counter on the pool
   retain();
 
+  ClPrint(amd::LOG_INFO, amd::LOG_MEM_POOL, "Pool AllocMem: %p, %p", memory->getSvmPtr(), memory);
+
   return dev_ptr;
 }
 
@@ -210,6 +220,24 @@ bool MemoryPool::FreeMemory(amd::Memory* memory, hip::Stream* stream) {
     // This pool doesn't contain memory
     return false;
   }
+  ClPrint(amd::LOG_INFO, amd::LOG_MEM_POOL, "Pool FreeMem: %p, %p", memory->getSvmPtr(), memory);
+
+  auto ga = reinterpret_cast<hip::MemMapAllocUserData*>(memory->getUserData().data);
+  if (ga != nullptr) {
+    if (stream == nullptr) {
+      stream = g_devices[memory->getUserData().deviceId]->NullStream();
+    }
+    // Unmap virtual address from memory
+    auto cmd = new amd::VirtualMapCommand(*stream, amd::Command::EventWaitList{},
+                                          memory->getSvmPtr(), ga->size_, nullptr);
+    cmd->enqueue();
+    cmd->release();
+    memory->setSvmPtr(ga->ptr_);
+    // Free virtual address and destroy generic allocation object
+    ga->va_->release();
+    delete ga;
+    memory->getUserData().data = nullptr;
+  }
 
   if (stream != nullptr) {
     // The stream of destruction is a safe stream, because the app must handle sync

hipamd/src/hip_mempool_impl.hpp

@@ -33,7 +33,9 @@ class Stream;
 
 struct MemoryTimestamp {
   MemoryTimestamp(hip::Stream* stream, hip::Event* event = nullptr): event_(event) {
-    safe_streams_.insert(stream);
+    if (stream != nullptr) {
+      safe_streams_.insert(stream);
+    }
   }
   MemoryTimestamp(): event_(nullptr) {}
 
@@ -63,6 +65,9 @@ struct MemoryTimestamp {
     } else if (opportunistic && (event_ != nullptr)) {
       // Check HIP event for a retired status
       result = (event_->query() == hipSuccess) ? true : false;
+    } else if (event_ == nullptr) {
+      // Event doesn't exist. It was a safe release with explicit wait
+      return true;
     }
     return result;
   }

hipamd/src/hip_vm.hpp

@@ -1,4 +1,4 @@
-/* Copyright (c) 2015 - 2022 Advanced Micro Devices, Inc.
+/* Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc.
 
  Permission is hereby granted, free of charge, to any person obtaining a copy
  of this software and associated documentation files (the "Software"), to deal
@@ -27,6 +27,15 @@
 hipError_t ihipFree(void* ptr);
 
 namespace hip {
+
+struct MemMapAllocUserData {
+  void* ptr_;       // Original pointer of the allocation
+  size_t size_;     // Aligned size of the allocation
+  amd::Memory* va_; // Memory object for the virtual address
+
+  MemMapAllocUserData(void* ptr, size_t size, amd::Memory* va) : ptr_(ptr), size_(size), va_(va) {}
+};
+
 class GenericAllocation {
   void* ptr_;
   size_t size_;

rocclr/device/pal/palvirtual.cpp

@@ -865,6 +865,7 @@ bool VirtualGPU::createVirtualQueue(uint deviceQueueSize) {
   return true;
 }
 
+// ================================================================================================
 VirtualGPU::VirtualGPU(Device& device)
     : device::VirtualDevice(device),
       engineID_(MainEngine),
@@ -897,6 +898,7 @@ VirtualGPU::VirtualGPU(Device& device)
   hostcallBuffer_ = nullptr;
 }
 
+// ================================================================================================
 bool VirtualGPU::create(bool profiling, uint deviceQueueSize, uint rtCUs,
                         amd::CommandQueue::Priority priority) {
   device::BlitManager::Setup blitSetup;
@@ -1046,10 +1048,11 @@ bool VirtualGPU::create(bool profiling, uint deviceQueueSize, uint rtCUs,
     dev().rgpCaptureMgr()->RegisterTimedQueue(2 * index() + 1, queue(SdmaEngine).iQueue_,
                                               &dbg_vmid);
   }
-
+  
   return true;
 }
 
+// ================================================================================================
 bool VirtualGPU::allocHsaQueueMem() {
   // Allocate a dummy HSA queue
   hsaQueueMem_ = new Memory(dev(), sizeof(amd_queue_t));
@@ -2212,6 +2215,14 @@ void VirtualGPU::submitVirtualMap(amd::VirtualMapCommand& vcmd) {
     vcmd.size(),
     Pal::VirtualGpuMemAccessMode::NoAccess
   };
+
+  // Wait for previous operations before unmap
+  if (vcmd.memory() == nullptr) {
+    // @note: Need to verify if compute requires a wait or IB flush is enough
+    WaitForIdleCompute();
+    WaitForIdleSdma();
+  }
+
   eventBegin(MainEngine);
   auto result = queue(MainEngine).iQueue_->RemapVirtualMemoryPages(1, &range, false, nullptr);
   // Capture GPU event for the paging operation

rocclr/device/pal/palvirtual.hpp

@@ -707,7 +707,7 @@ inline void VirtualGPU::logVmMemory(const std::string name, const Memory* memory
     if (PAL_EMBED_KERNEL_MD) {
       iCmd()->CmdCommentString(buf);
     }
-    LogPrintfInfo("%s", buf);
+    ClPrint(amd::LOG_INFO, amd::LOG_KERN, "%s", buf);
   }
 }
 

rocclr/platform/command.hpp

@@ -1696,10 +1696,12 @@ class SvmPrefetchAsyncCommand : public Command {
 class VirtualMapCommand : public Command {
  private:
   const void* ptr_;     //!< Virtual address to map to the memory
-  size_t size_;         //!< Size of the mapping in bytes
-  Memory* memory_;      //!< Memory to map, nullptr means unmap
 
- public:
+protected:
+  Memory* memory_;  //!< Memory to map, nullptr means unmap
+  size_t size_;     //!< Size of the mapping in bytes
+
+public:
   //! Construct a new VirtualMapCommand
   VirtualMapCommand(HostQueue& queue, const EventWaitList& eventWaitList,
                    void* ptr, size_t size, Memory* memory)

rocclr/platform/memory.cpp

@@ -449,9 +449,14 @@ Memory::~Memory() {
     parent_->release();
   }
   hostMemRef_.deallocateMemory(context_());
-
   if (getMemFlags() & CL_MEM_VA_RANGE_AMD) {
     amd::MemObjMap::RemoveVirtualMemObj(getSvmPtr());
+    // If runtime executes graph mempool with VM, then VA can be mapped in space
+    // for graph validation logic during execution. And the reason it's not unmaped
+    // in graph itself because the app can have a graph without a free node
+    if (amd::MemObjMap::FindMemObj(getSvmPtr())) {
+      amd::MemObjMap::RemoveMemObj(getSvmPtr());
+    }
   }
 }
 

rocclr/utils/debug.hpp

@@ -56,6 +56,7 @@ enum LogMask {
   LOG_CMD2      = 0x00008000, //!< More detailed command info, including barrier commands
   LOG_LOCATION  = 0x00010000, //!< Log message location
   LOG_MEM       = 0x00020000, //!< Memory allocation
+  LOG_MEM_POOL  = 0x00040000, //!< Memory pool allocation, including memory in graphs
   LOG_ALWAYS    = 0xFFFFFFFF, //!< Log always even mask flag is zero
 };
 

rocclr/utils/flags.hpp

@@ -251,6 +251,8 @@ release(bool, GPU_FORCE_QUEUE_PROFILING, false,                               \
         "Force command queue profiling by default")                           \
 release(bool, HIP_MEM_POOL_SUPPORT, false,                                    \
         "Enables memory pool support in HIP")                                 \
+release(bool, HIP_MEM_POOL_USE_VM, IS_WINDOWS,                                \
+        "Enables memory pool support in HIP")                                 \
 release(uint, PAL_FORCE_ASIC_REVISION, 0,                                     \
         "Force a specific asic revision for all devices")                     \
 release(bool, PAL_EMBED_KERNEL_MD, false,                                     \