From 94afcd643ca7bb1433af8eeb8abac852cbb699ff Mon Sep 17 00:00:00 2001
From: German Andryeyev <German.Andryeyev@amd.com>
Date: Wed, 3 Mar 2021 13:14:28 -0500
Subject: [PATCH] SWDEV-272496 - Delay enqueue until callback setup

With direct dispatch enqueue occurs before callback update and
it can't be tracked in the device backend

Change-Id: Ie8793e3ddb68cc5bb36348f7a8dcdbdc87a2487c


[ROCm/hip commit: ffc89dff8aa2d653b8cd4a5ef26b13b26ac4687e]
---
 projects/hip/rocclr/hip_stream.cpp | 7 +++++--
 1 file changed, 5 insertions(+), 2 deletions(-)

diff --git a/projects/hip/rocclr/hip_stream.cpp b/projects/hip/rocclr/hip_stream.cpp
index 01820476cb..91311d3232 100755
--- a/projects/hip/rocclr/hip_stream.cpp
+++ b/projects/hip/rocclr/hip_stream.cpp
@@ -368,6 +368,7 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
   }
   amd::HostQueue* hostQueue = hip::getQueue(stream);
   amd::Command* command = hostQueue->getLastQueuedCommand(true);
+  bool  cmd_created = false;
   if (command == nullptr ||
       // Add a marker if the last command is a notification,
       // since with multithreading AddCallback may occur before the previous
@@ -375,7 +376,7 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
       (command->type() == 0)) {
     amd::Command::EventWaitList eventWaitList;
     command = new amd::Marker(*hostQueue, kMarkerDisableFlush, eventWaitList);
-    command->enqueue();
+    cmd_created = true;
   }
   amd::Event& event = command->event();
   StreamCallback* cbo = new StreamCallback(stream, callback, userData, command);
@@ -384,7 +385,9 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
     command->release();
     return hipErrorInvalidHandle;
   }
-
+  if (cmd_created) {
+    command->enqueue();
+  }
   event.notifyCmdQueue();
 
   HIP_RETURN(hipSuccess);