SWDEV-413997 - Handling cases where multiple phys_mem is mapped into a single virtual address.

Change-Id: Ie678e607a64f2e5c35a10b9083185f041c5527ac


[ROCm/clr commit: 49b4aef8c9]

projects/clr/hipamd/src/hip_memory.cpp

@@ -367,11 +367,25 @@ hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes,
   amd::Memory* srcMemory = getMemoryObject(src, sOffset);
   size_t dOffset = 0;
   amd::Memory* dstMemory = getMemoryObject(dst, dOffset);
+
+  // If the mem object is a VMM sub buffer (subbuffer has parent set),
+  // then use parent's size for validation.
+  if (srcMemory && srcMemory->parent() && (srcMemory->getMemFlags() & CL_MEM_VA_RANGE_AMD)) {
+      srcMemory = srcMemory->parent();
+  }
+
+  // If the mem object is a VMM sub buffer (subbuffer has parent set),
+  // then use parent's size for validation.
+  if (dstMemory && dstMemory->parent() && (dstMemory->getMemFlags() & CL_MEM_VA_RANGE_AMD)) {
+      dstMemory = dstMemory->parent();
+  }
+
   // Return error if sizeBytes passed to memcpy is more than the actual size allocated
   if ((dstMemory && sizeBytes > (dstMemory->getSize() - dOffset)) ||
       (srcMemory && sizeBytes > (srcMemory->getSize() - sOffset))) {
     return hipErrorInvalidValue;
   }
+
   //If src and dst ptr are null then kind must be either h2h or def.
   if (!IsHtoHMemcpyValid(dst, src, kind)) {
     return hipErrorInvalidValue;

projects/clr/rocclr/device/device.cpp

@@ -405,6 +405,8 @@ amd::Memory* Device::CreateVirtualBuffer(amd::Context& device_context, void* vpt
 
   amd::Memory* vaddr_base_obj = nullptr;
   amd::Memory* vaddr_sub_obj = nullptr;
+  constexpr bool kSysMemAlloc = false;
+  constexpr bool kSkipAlloc = false;
 
   if (parent) {
     vaddr_base_obj = new (device_context) amd::Buffer(device_context, CL_MEM_VA_RANGE_AMD, size,
@@ -415,8 +417,6 @@ amd::Memory* Device::CreateVirtualBuffer(amd::Context& device_context, void* vpt
     }
     // This curr_mem_obj->create() does not create an actual memory but stores the memory info
     // with given vptr on ROCr backend.
-    constexpr bool kSysMemAlloc = false;
-    constexpr bool kSkipAlloc = false;
     if (!vaddr_base_obj->create(nullptr, kSysMemAlloc, kSkipAlloc, kForceAlloc)) {
       LogError("failed to create a va range mem object");
       vaddr_base_obj->release();
@@ -431,6 +431,7 @@ amd::Memory* Device::CreateVirtualBuffer(amd::Context& device_context, void* vpt
       LogPrintfError("Cannot find entry in VirtualMemObjMap: 0x%x \n", vptr);
       return nullptr;
     }
+    assert(vaddr_base_obj->getMemFlags() & CL_MEM_VA_RANGE_AMD);
 
     size_t offset = (reinterpret_cast<address>(vptr)
                      - reinterpret_cast<address>(vaddr_base_obj->getSvmPtr()));
@@ -439,8 +440,6 @@ amd::Memory* Device::CreateVirtualBuffer(amd::Context& device_context, void* vpt
 
     // This curr_mem_obj->create() does not create an actual memory but stores the memory info
     // with given vptr on ROCr backend.
-    constexpr bool kSysMemAlloc = false;
-    constexpr bool kSkipAlloc = false;
     if (!vaddr_sub_obj->create(nullptr, kSysMemAlloc, kSkipAlloc, kForceAlloc)) {
       LogError("failed to create a va range mem object");
       vaddr_sub_obj->release();