SWDEV-413997 - Changes to use GlobalContext in views.

Change-Id: I1f8411eae9ed49632667e244a25f223fed92c720


[ROCm/clr commit: 0e1a0572e6]

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);

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;

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<DeviceMemory*>(reinterpret_cast<char*>(this) + sizeof(Buffer));
   memset(deviceMemories_, 0, NumDevicesWithP2P() * sizeof(DeviceMemory));

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,