diff --git a/projects/clr/hipamd/src/hip_graph_internal.cpp b/projects/clr/hipamd/src/hip_graph_internal.cpp index eb627b90f3..7e33caa2b1 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.cpp +++ b/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; } diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index 5a504feb09..5ccd2701ff 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/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(packet)->group_segment_size, reinterpret_cast(packet)->kernel_object, reinterpret_cast(packet)->kernarg_address, - reinterpret_cast(packet)->completion_signal); + reinterpret_cast(packet)->completion_signal, read, index); } hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index - 1);