From 8028d327e99d6e524d287e7323795b60500ce0ce Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Tue, 17 Jan 2023 16:27:11 -0800 Subject: [PATCH] SWDEV-345213 - Use the right accessor - Use correct accessor to fetch memory objects. This checks the svm map and arena maps Change-Id: I84515330bb530cfe2b39abf30e1e659938f06806 --- hipamd/src/hip_graph_internal.cpp | 18 +++++++++++------ hipamd/src/hip_graph_internal.hpp | 13 +++++++++---- hipamd/src/hip_memory.cpp | 32 ++++++++++++++++++++++--------- hipamd/src/hip_vm.hpp | 5 ++++- 4 files changed, 48 insertions(+), 20 deletions(-) diff --git a/hipamd/src/hip_graph_internal.cpp b/hipamd/src/hip_graph_internal.cpp index f4060a0fe9..328cf4ca3b 100644 --- a/hipamd/src/hip_graph_internal.cpp +++ b/hipamd/src/hip_graph_internal.cpp @@ -103,6 +103,7 @@ hipError_t hipGraphMemcpyNode::ValidateParams(const hipMemcpy3DParms* pNodeParam if (status != hipSuccess) { return status; } + size_t offset = 0; const HIP_MEMCPY3D pCopy = hip::getDrvMemcpy3DDesc(*pNodeParams); // If {src/dst}MemoryType is hipMemoryTypeUnified, {src/dst}Device and {src/dst}Pitch specify the // (unified virtual address space) base address of the source data and the bytes per row to apply. @@ -110,36 +111,42 @@ hipError_t hipGraphMemcpyNode::ValidateParams(const hipMemcpy3DParms* pNodeParam hipMemoryType srcMemoryType = pCopy.srcMemoryType; if (srcMemoryType == hipMemoryTypeUnified) { srcMemoryType = - amd::MemObjMap::FindMemObj(pCopy.srcDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost; + getMemoryObject(pCopy.srcDevice, offset) ? hipMemoryTypeDevice : hipMemoryTypeHost; if (srcMemoryType == hipMemoryTypeHost) { // {src/dst}Host may be unitialized. Copy over {src/dst}Device into it if we detect system // memory. const_cast(&pCopy)->srcHost = pCopy.srcDevice; + const_cast(&pCopy)->srcXInBytes += offset; } } + offset = 0; hipMemoryType dstMemoryType = pCopy.dstMemoryType; if (dstMemoryType == hipMemoryTypeUnified) { dstMemoryType = - amd::MemObjMap::FindMemObj(pCopy.dstDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost; + getMemoryObject(pCopy.dstDevice, offset) ? hipMemoryTypeDevice : hipMemoryTypeHost; if (srcMemoryType == hipMemoryTypeHost) { const_cast(&pCopy)->dstHost = pCopy.dstDevice; + const_cast(&pCopy)->dstXInBytes += offset; } } - + offset = 0; // If {src/dst}MemoryType is hipMemoryTypeHost, check if the memory was prepinned. // In that case upgrade the copy type to hipMemoryTypeDevice to avoid extra pinning. if (srcMemoryType == hipMemoryTypeHost) { - amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy.srcHost); + amd::Memory* mem = getMemoryObject(pCopy.srcHost, offset); srcMemoryType = mem ? hipMemoryTypeDevice : hipMemoryTypeHost; if (srcMemoryType == hipMemoryTypeDevice) { const_cast(&pCopy)->srcDevice = const_cast(pCopy.srcHost); + const_cast(&pCopy)->srcXInBytes += offset; } } + offset = 0; if (dstMemoryType == hipMemoryTypeHost) { - amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy.dstHost); + amd::Memory* mem = getMemoryObject(pCopy.dstHost, offset); dstMemoryType = mem ? hipMemoryTypeDevice : hipMemoryTypeHost; if (dstMemoryType == hipMemoryTypeDevice) { const_cast(&pCopy)->dstDevice = const_cast(pCopy.dstDevice); + const_cast(&pCopy)->dstXInBytes += offset; } } @@ -255,7 +262,6 @@ bool ihipGraph::isGraphValid(ihipGraph* pGraph) { return true; } - void ihipGraph::AddNode(const Node& node) { vertices_.emplace_back(node); ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] Add %s(%p)\n", diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index 819125c1f3..a5799081ec 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/hipamd/src/hip_graph_internal.hpp @@ -1126,28 +1126,33 @@ class hipGraphMemcpyNode : public hipGraphNode { const hipGraphMemcpyNode* memcpyNode = static_cast(node); return SetParams(memcpyNode->pCopyParams_); } + // ToDo: use this when commands are cloned and command params are to be updated hipError_t ValidateParams(const hipMemcpy3DParms* pNodeParams); + std::string GetLabel(hipGraphDebugDotFlags flag) { + size_t offset = 0; const HIP_MEMCPY3D pCopy = hip::getDrvMemcpy3DDesc(*pCopyParams_); hipMemoryType srcMemoryType = pCopy.srcMemoryType; if (srcMemoryType == hipMemoryTypeUnified) { srcMemoryType = - amd::MemObjMap::FindMemObj(pCopy.srcDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost; + getMemoryObject(pCopy.srcDevice, offset) ? hipMemoryTypeDevice : hipMemoryTypeHost; } + offset = 0; hipMemoryType dstMemoryType = pCopy.dstMemoryType; if (dstMemoryType == hipMemoryTypeUnified) { dstMemoryType = - amd::MemObjMap::FindMemObj(pCopy.dstDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost; + getMemoryObject(pCopy.dstDevice, offset) ? hipMemoryTypeDevice : hipMemoryTypeHost; } // If {src/dst}MemoryType is hipMemoryTypeHost, check if the memory was prepinned. // In that case upgrade the copy type to hipMemoryTypeDevice to avoid extra pinning. + offset = 0; if (srcMemoryType == hipMemoryTypeHost) { - amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy.srcHost); + amd::Memory* mem = getMemoryObject(pCopy.srcHost, offset); srcMemoryType = mem ? hipMemoryTypeDevice : hipMemoryTypeHost; } if (dstMemoryType == hipMemoryTypeHost) { - amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy.dstHost); + amd::Memory* mem = getMemoryObject(pCopy.dstHost, offset); dstMemoryType = mem ? hipMemoryTypeDevice : hipMemoryTypeHost; } std::string memcpyDirection; diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 732ae9c490..fc29d0c1ec 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -2041,42 +2041,50 @@ hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, hipArray* srcArray, voi hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3D* pCopy, hip::Stream* stream) { + size_t offset = 0; // If {src/dst}MemoryType is hipMemoryTypeUnified, {src/dst}Device and {src/dst}Pitch specify the // (unified virtual address space) base address of the source data and the bytes per row to apply. // {src/dst}Array is ignored. hipMemoryType srcMemoryType = pCopy->srcMemoryType; if (srcMemoryType == hipMemoryTypeUnified) { srcMemoryType = - amd::MemObjMap::FindMemObj(pCopy->srcDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost; + getMemoryObject(pCopy->srcDevice, offset) ? hipMemoryTypeDevice : hipMemoryTypeHost; if (srcMemoryType == hipMemoryTypeHost) { // {src/dst}Host may be unitialized. Copy over {src/dst}Device into it if we detect system // memory. const_cast(pCopy)->srcHost = pCopy->srcDevice; + const_cast(pCopy)->srcXInBytes += offset; } } + offset = 0; hipMemoryType dstMemoryType = pCopy->dstMemoryType; if (dstMemoryType == hipMemoryTypeUnified) { dstMemoryType = - amd::MemObjMap::FindMemObj(pCopy->dstDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost; + getMemoryObject(pCopy->dstDevice, offset) ? hipMemoryTypeDevice : hipMemoryTypeHost; if (srcMemoryType == hipMemoryTypeHost) { const_cast(pCopy)->dstHost = pCopy->dstDevice; + const_cast(pCopy)->dstXInBytes += offset; } } // If {src/dst}MemoryType is hipMemoryTypeHost, check if the memory was prepinned. // In that case upgrade the copy type to hipMemoryTypeDevice to avoid extra pinning. + offset = 0; if (srcMemoryType == hipMemoryTypeHost) { - amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy->srcHost); + amd::Memory* mem = getMemoryObject(pCopy->srcHost, offset); srcMemoryType = mem ? hipMemoryTypeDevice : hipMemoryTypeHost; if (srcMemoryType == hipMemoryTypeDevice) { const_cast(pCopy)->srcDevice = const_cast(pCopy->srcHost); + const_cast(pCopy)->srcXInBytes += offset; } } + offset = 0; if (dstMemoryType == hipMemoryTypeHost) { - amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy->dstHost); + amd::Memory* mem = getMemoryObject(pCopy->dstHost, offset); dstMemoryType = mem ? hipMemoryTypeDevice : hipMemoryTypeHost; if (dstMemoryType == hipMemoryTypeDevice) { const_cast(pCopy)->dstDevice = const_cast(pCopy->dstHost); + const_cast(pCopy)->dstXInBytes += offset; } } @@ -2147,6 +2155,7 @@ inline hipError_t ihipMemcpyCmdEnqueue(amd::Command* command, bool isAsync = fal hipError_t ihipMemcpyParam3D(const HIP_MEMCPY3D* pCopy, hipStream_t stream, bool isAsync = false) { hipError_t status; + size_t offset = 0; if (pCopy == nullptr) { return hipErrorInvalidValue; } @@ -2162,27 +2171,31 @@ hipError_t ihipMemcpyParam3D(const HIP_MEMCPY3D* pCopy, hipStream_t stream, bool // base address of the source data and the bytes per row to apply. {src/dst}Array is ignored. hipMemoryType srcMemoryType = pCopy->srcMemoryType; if (srcMemoryType == hipMemoryTypeUnified) { - srcMemoryType = amd::MemObjMap::FindMemObj(pCopy->srcDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost; + srcMemoryType = getMemoryObject(pCopy->srcDevice, offset) ? hipMemoryTypeDevice : hipMemoryTypeHost; if (srcMemoryType == hipMemoryTypeHost) { // {src/dst}Host may be unitialized. Copy over {src/dst}Device into it if we detect system memory. const_cast(pCopy)->srcHost = pCopy->srcDevice; + const_cast(pCopy)->srcXInBytes += offset; } } + offset = 0; hipMemoryType dstMemoryType = pCopy->dstMemoryType; if (dstMemoryType == hipMemoryTypeUnified) { - dstMemoryType = amd::MemObjMap::FindMemObj(pCopy->dstDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost; + dstMemoryType = getMemoryObject(pCopy->dstDevice, offset) ? hipMemoryTypeDevice : hipMemoryTypeHost; if (srcMemoryType == hipMemoryTypeHost) { const_cast(pCopy)->dstHost = pCopy->dstDevice; + const_cast(pCopy)->dstXInBytes += offset; } } // If {src/dst}MemoryType is hipMemoryTypeHost, check if the memory was prepinned. // In that case upgrade the copy type to hipMemoryTypeDevice to avoid extra pinning. + offset = 0; if (srcMemoryType == hipMemoryTypeHost) { - amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy->srcHost); + amd::Memory* mem = getMemoryObject(pCopy->srcHost, offset); srcMemoryType = mem ? hipMemoryTypeDevice : hipMemoryTypeHost; } if (dstMemoryType == hipMemoryTypeHost) { - amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy->dstHost); + amd::Memory* mem = getMemoryObject(pCopy->dstHost, offset); dstMemoryType = mem ? hipMemoryTypeDevice : hipMemoryTypeHost; } if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeHost)) { @@ -2781,7 +2794,8 @@ hipError_t ihipGraphMemsetParams_validate(const hipMemsetParams* pNodeParams) { return hipErrorInvalidValue; } - amd::Memory *memObj = amd::MemObjMap::FindMemObj(pNodeParams->dst); + size_t discardOffset = 0; + amd::Memory *memObj = getMemoryObject(pNodeParams->dst, discardOffset); if (memObj != nullptr) { if ((pNodeParams->pitch * pNodeParams->height) > memObj->getSize()) { return hipErrorInvalidValue; diff --git a/hipamd/src/hip_vm.hpp b/hipamd/src/hip_vm.hpp index d87768ac46..a38acf63e8 100644 --- a/hipamd/src/hip_vm.hpp +++ b/hipamd/src/hip_vm.hpp @@ -38,7 +38,10 @@ public: const hipMemAllocationProp& GetProperties() const { return properties_; } hipMemGenericAllocationHandle_t asMemGenericAllocationHandle() { return reinterpret_cast(this); } - amd::Memory& asAmdMemory() { return *amd::MemObjMap::FindMemObj(genericAddress()); } + amd::Memory& asAmdMemory() { + size_t discardOffset; + return *getMemoryObject(genericAddress(), discardOffset); + } void* genericAddress() const { return ptr_; } }; };