From e6bdfcccf7d6de43fd660669ae7454574da390c0 Mon Sep 17 00:00:00 2001
From: Anusha GodavarthySurya <anusha.godavarthysurya@amd.com>
Date: Thu, 20 Jul 2023 05:55:26 +0000
Subject: [PATCH] SWDEV-392732 - CP processes until doorbell ring

Change-Id: I3a0974435d3f95a40cea6ed326f744bcab5b4c0f


[ROCm/clr commit: 0711320cb4e451497f7446399dd09ad23c846765]
---
 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 97c38d2437..32c184c664 100644
--- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp
+++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp
@@ -871,7 +871,7 @@ bool VirtualGPU::dispatchAqlBuffer() {
 
     packet_store_release(reinterpret_cast<uint32_t*>(aql_loc), firstPacketHeader,
                          aqlBuffer_.front().setup);
-    hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index);
+    hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index + size - 1);
 
     aqlBuffer_.clear();
   }