From 4842e57b4190cd6b34ee1be44499b5be1b4d04e3 Mon Sep 17 00:00:00 2001
From: foreman <dl.constructicon@amd.com>
Date: Wed, 24 Oct 2018 18:00:42 -0400
Subject: [PATCH] P4 to Git Change 1623601 by
 skudchad@skudchad_test2_win_opencl on 2018/10/24 17:52:57

	SWDEV-145570 - [HIP] Refactor hipStreamAddCallback

	ReviewBoardURL = http://ocltc.amd.com/reviews/r/16045/diff/

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#15 edit


[ROCm/clr commit: 5a126bf67857abd9c79b1364073699eec457c764]
---
 projects/clr/hipamd/api/hip/hip_stream.cpp | 45 +++++++++++++++++-----
 1 file changed, 35 insertions(+), 10 deletions(-)

diff --git a/projects/clr/hipamd/api/hip/hip_stream.cpp b/projects/clr/hipamd/api/hip/hip_stream.cpp
index af00b0f104..7acd3d90dd 100644
--- a/projects/clr/hipamd/api/hip/hip_stream.cpp
+++ b/projects/clr/hipamd/api/hip/hip_stream.cpp
@@ -28,6 +28,20 @@ THE SOFTWARE.
 static amd::Monitor streamSetLock("Guards global stream set");
 static std::unordered_set<amd::HostQueue*> streamSet;
 
+// Internal structure for stream callback handler
+class StreamCallback {
+   public:
+    StreamCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,
+                  amd::Command* command)
+        : stream_(stream), callBack_(callback),
+          userData_(userData), command_(command) {
+        };
+    hipStream_t stream_;
+    hipStreamCallback_t callBack_;
+    void* userData_;
+    amd::Command* command_;
+};
+
 namespace hip {
 
 void syncStreams() {
@@ -40,6 +54,15 @@ void syncStreams() {
 
 };
 
+void ihipStreamCallback(cl_event event, cl_int command_exec_status, void* user_data) {
+
+  hipError_t status = hipSuccess;
+  StreamCallback* cbo = reinterpret_cast<StreamCallback*>(user_data);
+  cbo->callBack_(cbo->stream_, status, cbo->userData_);
+  cbo->command_->release();
+  delete cbo;
+}
+
 static hipError_t ihipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) {
   amd::Device* device = hip::getCurrentContext()->devices()[0];
 
@@ -66,15 +89,6 @@ static hipError_t ihipStreamCreateWithFlags(hipStream_t *stream, unsigned int fl
   return hipSuccess;
 }
 
-
-void ihipStreamCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData) {
-    //Stream synchronize
-    hipError_t status = hipStreamSynchronize(stream);
-
-    // Call the callback function
-    callback(stream, status, userData);
-}
-
 hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) {
   HIP_INIT_API(stream, flags);
 
@@ -189,7 +203,18 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback
                                 unsigned int flags) {
   HIP_INIT_API(stream, callback, userData, flags);
 
-  std::thread (ihipStreamCallback, stream, callback, userData).detach();
+  amd::HostQueue* hostQueue = as_amd(reinterpret_cast<cl_command_queue>
+                              (stream))->asHostQueue();
+  amd::Command* command = hostQueue->getLastQueuedCommand(true);
+  amd::Event& event = command->event();
+  StreamCallback* cbo = new StreamCallback(stream, callback, userData, command);
+
+  if(!event.setCallback(CL_COMPLETE, ihipStreamCallback, reinterpret_cast<void*>(cbo))) {
+    command->release();
+    return hipErrorInvalidResourceHandle;
+  }
+
+  event.notifyCmdQueue();
 
   HIP_RETURN(hipSuccess);
 }