From b3f3af2b1b2bee93b2fdc0530a930c0216b11c63 Mon Sep 17 00:00:00 2001
From: German Andryeyev <German.Andryeyev@amd.com>
Date: Thu, 4 Mar 2021 17:19:47 -0500
Subject: [PATCH] SWDEV-272496 - enqueue a marker for callbacks always

The current implementation in ROCclr for callback is
based on OCL specification. If in HIP the same command
could get multiple callbacks, then ROCclr will process them in
a reverse order. Unique Markers for each callback will make
sure it won't happen.
Add a dependency wait for callbacks, since HSA signal callback
doesn't guarantee the order.

Change-Id: I9d514734e258312fe9a74d48132361eb17c52d67


[ROCm/hip commit: 84f878528803de6898f43a8cce1c2e2b6f32aaaf]
---
 projects/hip/rocclr/hip_stream.cpp | 31 ++++++++++++++++--------------
 1 file changed, 17 insertions(+), 14 deletions(-)

diff --git a/projects/hip/rocclr/hip_stream.cpp b/projects/hip/rocclr/hip_stream.cpp
index 91311d3232..339b61acb2 100755
--- a/projects/hip/rocclr/hip_stream.cpp
+++ b/projects/hip/rocclr/hip_stream.cpp
@@ -367,28 +367,31 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
     HIP_RETURN(hipErrorInvalidValue);
   }
   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
-      // command can be processed
-      (command->type() == 0)) {
-    amd::Command::EventWaitList eventWaitList;
-    command = new amd::Marker(*hostQueue, kMarkerDisableFlush, eventWaitList);
-    cmd_created = true;
+  amd::Command* last_command = hostQueue->getLastQueuedCommand(true);
+  amd::Command::EventWaitList eventWaitList;
+  if (last_command != nullptr) {
+    eventWaitList.push_back(last_command);
+  }
+  amd::Command* command = new amd::Marker(*hostQueue, !kMarkerDisableFlush, eventWaitList);
+  if (command == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
   }
   amd::Event& event = command->event();
   StreamCallback* cbo = new StreamCallback(stream, callback, userData, command);
 
-  if(!event.setCallback(CL_COMPLETE, ihipStreamCallback, reinterpret_cast<void*>(cbo))) {
+  if ((cbo == nullptr) || !event.setCallback(CL_COMPLETE, ihipStreamCallback, cbo)) {
     command->release();
+    if (last_command != nullptr) {
+      last_command->release();
+    }
     return hipErrorInvalidHandle;
   }
-  if (cmd_created) {
-    command->enqueue();
+  command->enqueue();
+  // @note: don't release the command here, because it will be released after HIP callback
+
+  if (last_command != nullptr) {
+    last_command->release();
   }
-  event.notifyCmdQueue();
 
   HIP_RETURN(hipSuccess);
 }