From da0a525982e10c323935f1d309c8ea1bf11a82d9 Mon Sep 17 00:00:00 2001 From: Jason Tang Date: Mon, 17 Aug 2020 19:49:25 -0400 Subject: [PATCH] SWDEV-239502 - fix interop regression When header==0, the legitimate packet->header is wiped out, so also add an assert. Change-Id: I6b3037d4618719262b0d7c1792bd54f768a63660 [ROCm/clr commit: 19d1497fa25fe7a4f8954606f7fd99010593b5ef] --- projects/clr/rocclr/device/rocm/rocvirtual.cpp | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index bc6fb0bb31..891397eed9 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp @@ -451,6 +451,7 @@ bool VirtualGPU::dispatchGenericAqlPacket( uint64_t read = hsa_queue_load_read_index_relaxed(gpu_queue_); hsa_signal_t signal; + assert(header != 0); if (addSystemScope_) { header &= ~(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE); header |= (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE); @@ -2454,8 +2455,6 @@ void VirtualGPU::submitAcquireExtObjects(amd::AcquireExtObjectsCommand& vcmd) { amd::ScopedLock lock(execution()); profilingBegin(vcmd); - auto fence = kBarrierAcquirePacket; - dispatchAqlPacket(&fence, 0, 0, false); profilingEnd(vcmd); } @@ -2463,8 +2462,6 @@ void VirtualGPU::submitReleaseExtObjects(amd::ReleaseExtObjectsCommand& vcmd) { // Make sure VirtualGPU has an exclusive access to the resources amd::ScopedLock lock(execution()); profilingBegin(vcmd); - auto fence = kBarrierReleasePacket; - dispatchAqlPacket(&fence, 0, 0, false); profilingEnd(vcmd); }