diff --git a/projects/clr/rocclr/device/rocm/rocblit.cpp b/projects/clr/rocclr/device/rocm/rocblit.cpp index 3bf6720f42..031f9652a1 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/device/rocm/rocblit.cpp @@ -661,7 +661,7 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory, LogError("Async copy failed"); status = HSA_STATUS_ERROR; } else { - gpu().hasPendingDispatch(); + gpu().addSystemScope(); } } else { LogPrintfError("Hsa copy from host to device failed with code %d", status); @@ -749,7 +749,7 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_ offset += size; } - gpu().hasPendingDispatch(); + gpu().addSystemScope(); return true; } diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index 7a7e3efc1b..dafdc21f0d 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp @@ -436,6 +436,12 @@ bool VirtualGPU::dispatchGenericAqlPacket( uint64_t read = hsa_queue_load_read_index_relaxed(gpu_queue_); hsa_signal_t signal; + if (addSystemScope_) { + header &= ~(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE); + header |= (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE); + addSystemScope_ = false; + } + // TODO: placeholder to setup the kernel to populate start and end timestamp. if (timestamp_ != nullptr) { // Find signal slot diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.hpp b/projects/clr/rocclr/device/rocm/rocvirtual.hpp index 23ac9ca075..90f7a11420 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.hpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.hpp @@ -261,7 +261,8 @@ class VirtualGPU : public device::VirtualDevice { void enableSyncBlit() const; - void hasPendingDispatch() { hasPendingDispatch_ = true;} + void hasPendingDispatch() { hasPendingDispatch_ = true; } + void addSystemScope() { addSystemScope_ = true; } // } roc OpenCL integration private: @@ -319,6 +320,7 @@ class VirtualGPU : public device::VirtualDevice { uint32_t imageBufferWrtBack_ : 1; //!< Image buffer write back is required uint32_t profiling_ : 1; //!< Profiling is enabled uint32_t cooperative_ : 1; //!< Cooperative launch is enabled + uint32_t addSystemScope_ : 1; //!< Insert a system scope to the next aql }; uint32_t state_; };