From 76e04c2cf58545dcfc2e36a9bbd3625134b360e1 Mon Sep 17 00:00:00 2001
From: Sourabh Betigeri <sourabh.betigeri@amd.com>
Date: Wed, 14 Jul 2021 13:46:18 -0700
Subject: [PATCH] SWDEV-290685 - Relacing release fence instead of acquire for
 hipStreamWrite()

Change-Id: Ic2946b68c427d3e058948c0813863a27c21b903d


[ROCm/clr commit: 6408b9b906b82397a73629d57e70ab8952025ee4]
---
 projects/clr/rocclr/device/rocm/rocvirtual.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp
index 624a807d4e..07147600fb 100644
--- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp
+++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp
@@ -2277,7 +2277,7 @@ void VirtualGPU::submitStreamOperation(amd::StreamOperationCommand& cmd) {
     bool entire = amdMemory->isEntirelyCovered(origin, size);
 
     // Ensure memory ordering preceding the write
-    dispatchBarrierPacket(kBarrierPacketAcquireHeader);
+    dispatchBarrierPacket(kBarrierPacketReleaseHeader);
 
     // Use GPU Blit to write
     bool result = blitMgr().fillBuffer(*memory, &value, sizeBytes, origin, size, entire, true);