SWDEV-290685 - Dispatch a barrier packet with acquire system scope to ensure ordering before a hipStreamWrite()

Change-Id: I8853ad86a6634d55a98173ca3f79d93b85c08f85


[ROCm/clr commit: e4ddb5e08a]

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);