From 04b696abee9211f56825d9c48981a6eb94cffb94 Mon Sep 17 00:00:00 2001 From: German Date: Thu, 4 May 2023 14:54:19 -0400 Subject: [PATCH] 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 | 10 +- hipamd/src/hip_graph.cpp | 25 ++-- hipamd/src/hip_graph_internal.cpp | 2 + hipamd/src/hip_graph_internal.hpp | 211 ++++++++++++++++++++++++++++-- hipamd/src/hip_mempool.cpp | 53 ++++++-- hipamd/src/hip_mempool_impl.cpp | 38 +++++- hipamd/src/hip_mempool_impl.hpp | 7 +- hipamd/src/hip_vm.hpp | 11 +- rocclr/device/pal/palvirtual.cpp | 13 +- rocclr/device/pal/palvirtual.hpp | 2 +- rocclr/platform/command.hpp | 8 +- rocclr/platform/memory.cpp | 7 +- rocclr/utils/debug.hpp | 1 + rocclr/utils/flags.hpp | 2 + 14 files changed, 341 insertions(+), 49 deletions(-) diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp index d545a5ba47..2908dd3610 100644 --- a/hipamd/src/hip_device.cpp +++ b/hipamd/src/hip_device.cpp @@ -53,10 +53,12 @@ bool Device::Create() { return false; } - uint64_t max_size = std::numeric_limits::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::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_; diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index 2725d59f37..42f9039a50 100644 --- a/hipamd/src/hip_graph.cpp +++ b/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); diff --git a/hipamd/src/hip_graph_internal.cpp b/hipamd/src/hip_graph_internal.cpp index 7573b8b652..c4c99f4936 100644 --- a/hipamd/src/hip_graph_internal.cpp +++ b/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: diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index 02c425397a..04085249cd 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/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>& 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(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(*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(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(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(*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; } diff --git a/hipamd/src/hip_mempool.cpp b/hipamd/src/hip_mempool.cpp index cdbe1929b3..36dd674e55 100644 --- a/hipamd/src/hip_mempool.cpp +++ b/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(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(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(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(stream); + auto cmd = new FreeAsyncCommand(*hip_stream, dev_ptr); + if (cmd == nullptr) { + HIP_RETURN(hipErrorUnknown); } + cmd->enqueue(); + cmd->release(); HIP_RETURN(hipSuccess); } diff --git a/hipamd/src/hip_mempool_impl.cpp b/hipamd/src/hip_mempool_impl.cpp index 2606688ff0..45b7bdc680 100644 --- a/hipamd/src/hip_mempool_impl.cpp +++ b/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(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 diff --git a/hipamd/src/hip_mempool_impl.hpp b/hipamd/src/hip_mempool_impl.hpp index 5e18cb3599..75674742c0 100644 --- a/hipamd/src/hip_mempool_impl.hpp +++ b/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; } diff --git a/hipamd/src/hip_vm.hpp b/hipamd/src/hip_vm.hpp index a38acf63e8..1b8db6a1cb 100644 --- a/hipamd/src/hip_vm.hpp +++ b/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_; diff --git a/rocclr/device/pal/palvirtual.cpp b/rocclr/device/pal/palvirtual.cpp index e80e22298f..10a896bc60 100644 --- a/rocclr/device/pal/palvirtual.cpp +++ b/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 diff --git a/rocclr/device/pal/palvirtual.hpp b/rocclr/device/pal/palvirtual.hpp index 74fb039fae..bdf8615b50 100644 --- a/rocclr/device/pal/palvirtual.hpp +++ b/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); } } diff --git a/rocclr/platform/command.hpp b/rocclr/platform/command.hpp index 981cef5c94..6d133b18c8 100644 --- a/rocclr/platform/command.hpp +++ b/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) diff --git a/rocclr/platform/memory.cpp b/rocclr/platform/memory.cpp index b95d731484..b239d51b8c 100644 --- a/rocclr/platform/memory.cpp +++ b/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()); + } } } diff --git a/rocclr/utils/debug.hpp b/rocclr/utils/debug.hpp index 07cb3b61e8..2a5dd77acc 100644 --- a/rocclr/utils/debug.hpp +++ b/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 }; diff --git a/rocclr/utils/flags.hpp b/rocclr/utils/flags.hpp index 186d059b41..6e4d9d34e9 100644 --- a/rocclr/utils/flags.hpp +++ b/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, \