From 154a113a882a41cd07a16cf4412b94f6bb44c54b Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Wed, 24 Apr 2024 17:15:03 -0400 Subject: [PATCH] SWDEV-413997 - Changes to use GlobalContext in views. Change-Id: I1f8411eae9ed49632667e244a25f223fed92c720 [ROCm/clr commit: 0e1a0572e62fc23e0912811d1e148e1ad98caf43] --- projects/clr/hipamd/src/hip_memory.cpp | 26 ++++++++++++------------- projects/clr/rocclr/device/device.cpp | 3 +-- projects/clr/rocclr/platform/memory.cpp | 5 +++++ projects/clr/rocclr/platform/memory.hpp | 3 +++ 4 files changed, 22 insertions(+), 15 deletions(-) diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 3fa6eb078c..8ff94a6df7 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -405,8 +405,8 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, amd::CopyMetadata copyMetadata(isAsync, amd::CopyMetadata::CopyEnginePreference::NONE); if ((srcMemory == nullptr) && (dstMemory != nullptr)) { hip::Stream* pStream = &stream; - if (queueDevice != dstMemory->getContext().devices()[0]) { - pStream = hip::getNullStream(dstMemory->getContext()); + if (queueDevice != dstMemory->GetDeviceById()) { + pStream = hip::getNullStream(dstMemory->GetDeviceById()->context()); amd::Command* cmd = stream.getLastQueuedCommand(true); if (cmd != nullptr) { waitList.push_back(cmd); @@ -416,8 +416,8 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, *dstMemory->asBuffer(), dOffset, sizeBytes, src, 0, 0, copyMetadata); } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) { hip::Stream* pStream = &stream; - if (queueDevice != srcMemory->getContext().devices()[0]) { - pStream = hip::getNullStream(srcMemory->getContext()); + if (queueDevice != srcMemory->GetDeviceById()) { + pStream = hip::getNullStream(srcMemory->GetDeviceById()->context()); amd::Command* cmd = stream.getLastQueuedCommand(true); if (cmd != nullptr) { waitList.push_back(cmd); @@ -429,7 +429,7 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, // Check if the queue device doesn't match the device on any memory object. // And any of them are not host allocation. // Hence it's a P2P transfer, because the app has requested access to another GPU - if ((srcMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) && + if ((srcMemory->GetDeviceById() != dstMemory->GetDeviceById()) && ((srcMemory->getContext().devices().size() == 1) && (dstMemory->getContext().devices().size() == 1))) { command = new amd::CopyMemoryP2PCommand(stream, CL_COMMAND_COPY_BUFFER, waitList, @@ -445,26 +445,26 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, } } else { hip::Stream* pStream = &stream; - if ((srcMemory->getContext().devices()[0] == dstMemory->getContext().devices()[0]) && + if ((srcMemory->GetDeviceById() == dstMemory->GetDeviceById()) && (queueDevice != srcMemory->getContext().devices()[0])) { - pStream = hip::getNullStream(srcMemory->getContext()); + pStream = hip::getNullStream(srcMemory->GetDeviceById()->context()); amd::Command* cmd = stream.getLastQueuedCommand(true); if (cmd != nullptr) { waitList.push_back(cmd); } - } else if (srcMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) { + } else if (srcMemory->GetDeviceById() != dstMemory->GetDeviceById()) { // Scenarios such as DtoH where dst is pinned memory if ((queueDevice != srcMemory->getContext().devices()[0]) && (dstMemory->getContext().devices().size() != 1)) { - pStream = hip::getNullStream(srcMemory->getContext()); + pStream = hip::getNullStream(srcMemory->GetDeviceById()->context()); amd::Command* cmd = stream.getLastQueuedCommand(true); if (cmd != nullptr) { waitList.push_back(cmd); } // Scenarios such as HtoD where src is pinned memory - } else if ((queueDevice != dstMemory->getContext().devices()[0]) && + } else if ((queueDevice != dstMemory->GetDeviceById()) && (srcMemory->getContext().devices().size() != 1)) { - pStream = hip::getNullStream(dstMemory->getContext()); + pStream = hip::getNullStream(dstMemory->GetDeviceById()->context()); amd::Command* cmd = stream.getLastQueuedCommand(true); if (cmd != nullptr) { waitList.push_back(cmd); @@ -530,7 +530,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin } else if (((srcMemory == nullptr) && (dstMemory != nullptr)) || ((srcMemory != nullptr) && (dstMemory == nullptr))) { isHostAsync = false; - } else if (srcMemory->getContext().devices()[0] == dstMemory->getContext().devices()[0]) { + } else if (srcMemory->GetDeviceById() == dstMemory->GetDeviceById()) { hipMemoryType srcMemoryType = ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) & srcMemory->getMemFlags())? hipMemoryTypeHost : hipMemoryTypeDevice; hipMemoryType dstMemoryType = ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) & @@ -550,7 +550,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin if (!isHostAsync) { command->awaitCompletion(); } else if (!isGPUAsync) { - hip::Stream* pStream = hip::getNullStream(dstMemory->getContext()); + hip::Stream* pStream = hip::getNullStream(dstMemory->GetDeviceById()->context()); amd::Command::EventWaitList waitList; waitList.push_back(command); amd::Command* depdentMarker = new amd::Marker(*pStream, false, waitList); diff --git a/projects/clr/rocclr/device/device.cpp b/projects/clr/rocclr/device/device.cpp index 902016cd72..08ce8eb6e9 100644 --- a/projects/clr/rocclr/device/device.cpp +++ b/projects/clr/rocclr/device/device.cpp @@ -409,8 +409,7 @@ amd::Memory* Device::CreateVirtualBuffer(amd::Context& device_context, void* vpt constexpr bool kSkipAlloc = false; if (parent) { - vaddr_base_obj = new (device_context) amd::Buffer(device_context, CL_MEM_VA_RANGE_AMD, size, - vptr); + vaddr_base_obj = new (GlbCtx()) amd::Buffer(GlbCtx(), CL_MEM_VA_RANGE_AMD, size, vptr); if (vaddr_base_obj == nullptr) { LogError("failed to new a va range curr_mem_obj object!"); return nullptr; diff --git a/projects/clr/rocclr/platform/memory.cpp b/projects/clr/rocclr/platform/memory.cpp index d2629d6377..1242ccd15d 100644 --- a/projects/clr/rocclr/platform/memory.cpp +++ b/projects/clr/rocclr/platform/memory.cpp @@ -540,6 +540,11 @@ void Memory::uncommitSvmMemory() { } } +Device* Memory::GetDeviceById() { + size_t device_idx = (userData_.deviceId < getContext().devices().size()) ? userData_.deviceId : 0; + return getContext().devices()[device_idx]; +} + void Buffer::initDeviceMemory() { deviceMemories_ = reinterpret_cast(reinterpret_cast(this) + sizeof(Buffer)); memset(deviceMemories_, 0, NumDevicesWithP2P() * sizeof(DeviceMemory)); diff --git a/projects/clr/rocclr/platform/memory.hpp b/projects/clr/rocclr/platform/memory.hpp index 26da54ba0b..87f67cd927 100644 --- a/projects/clr/rocclr/platform/memory.hpp +++ b/projects/clr/rocclr/platform/memory.hpp @@ -413,6 +413,9 @@ class Memory : public amd::RuntimeObject { //!find if memory object is Arena memory virtual bool isArena() { return false; } + + //! get device by id when glb ctx is used. + Device* GetDeviceById(); }; //! Buffers are a specialization of memory. Just a wrapper, really,