SWDEV-301667 - Cleanup code and better log

Change-Id: Ie2345264e84026156a9f81b421eed3cf4aeeeffc


[ROCm/clr commit: 81b8598af9]

projects/clr/hipamd/src/hip_graph_internal.cpp

@@ -66,38 +66,21 @@ hipError_t GraphMemcpyNode1D::ValidateParams(void* dst, const void* src, size_t
   if (status != hipSuccess) {
     return status;
   }
-  size_t sOffsetOrig = 0;
-  amd::Memory* origSrcMemory = getMemoryObject(src, sOffsetOrig);
-  size_t dOffsetOrig = 0;
-  amd::Memory* origDstMemory = getMemoryObject(dst, dOffsetOrig);
-
   size_t sOffset = 0;
   amd::Memory* srcMemory = getMemoryObject(src, sOffset);
   size_t dOffset = 0;
   amd::Memory* dstMemory = getMemoryObject(dst, dOffset);
 
   if ((srcMemory == nullptr) && (dstMemory != nullptr)) {  // host to device
-    if (origDstMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) {
-      return hipErrorInvalidValue;
-    }
     if ((kind != hipMemcpyHostToDevice) && (kind != hipMemcpyDefault)) {
       return hipErrorInvalidValue;
     }
   } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) {  // device to host
-    if (origSrcMemory->getContext().devices()[0] != srcMemory->getContext().devices()[0]) {
-      return hipErrorInvalidValue;
-    }
     if ((kind != hipMemcpyDeviceToHost) && (kind != hipMemcpyDefault)) {
       return hipErrorInvalidValue;
     }
-  } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) {
-    if (origDstMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) {
-      return hipErrorInvalidValue;
-    }
-    if (origSrcMemory->getContext().devices()[0] != srcMemory->getContext().devices()[0]) {
-      return hipErrorInvalidValue;
-    }
   }
+
   return hipSuccess;
 }
 

projects/clr/rocclr/device/rocm/rocvirtual.cpp

@@ -900,7 +900,8 @@ bool VirtualGPU::dispatchGenericAqlPacket(
             "HWq=0x%zx, Dispatch Header = "
             "0x%x (type=%d, barrier=%d, acquire=%d, release=%d), "
             "setup=%d, grid=[%zu, %zu, %zu], workgroup=[%zu, %zu, %zu], private_seg_size=%zu, "
-            "group_seg_size=%zu, kernel_obj=0x%zx, kernarg_address=0x%zx, completion_signal=0x%zx",
+            "group_seg_size=%zu, kernel_obj=0x%zx, kernarg_address=0x%zx, completion_signal=0x%zx "
+            "rptr=%u, wptr=%u",
             gpu_queue_->base_address, header,
             extractAqlBits(header, HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_WIDTH_TYPE),
             extractAqlBits(header, HSA_PACKET_HEADER_BARRIER,
@@ -919,7 +920,7 @@ bool VirtualGPU::dispatchGenericAqlPacket(
             reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->group_segment_size,
             reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->kernel_object,
             reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->kernarg_address,
-            reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->completion_signal);
+            reinterpret_cast<hsa_kernel_dispatch_packet_t*>(packet)->completion_signal, read, index);
   }
 
   hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index - 1);