From c3be453fe6ac6ca753cf883fad0278a34b26ed42 Mon Sep 17 00:00:00 2001
From: Sourabh Betigeri <sourabh.betigeri@amd.com>
Date: Mon, 12 Jul 2021 23:38:17 -0700
Subject: [PATCH] SWDEV-290685 - Dispatch a barrier packet with acquire system
 scope to ensure ordering before a hipStreamWrite()

Change-Id: I8853ad86a6634d55a98173ca3f79d93b85c08f85


[ROCm/clr commit: e4ddb5e08ae8e56f5d4dac762ff555eeaf85d406]
---
 projects/clr/rocclr/device/rocm/rocvirtual.cpp | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp
index 0433435a5c..624a807d4e 100644
--- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp
+++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp
@@ -2276,6 +2276,9 @@ void VirtualGPU::submitStreamOperation(amd::StreamOperationCommand& cmd) {
     amd::Coord3D size(sizeBytes);
     bool entire = amdMemory->isEntirelyCovered(origin, size);
 
+    // Ensure memory ordering preceding the write
+    dispatchBarrierPacket(kBarrierPacketAcquireHeader);
+
     // Use GPU Blit to write
     bool result = blitMgr().fillBuffer(*memory, &value, sizeBytes, origin, size, entire, true);
     ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Writting value: 0x%lx", value);