From 9504ce2bc4036f7e85bb33120e54cd8b0577c714 Mon Sep 17 00:00:00 2001
From: foreman <dl.constructicon@amd.com>
Date: Fri, 30 Mar 2018 01:06:00 -0400
Subject: [PATCH] P4 to Git Change 1534798 by cpaquot@cpaquot-ocl-lc-lnx on
 2018/03/30 00:56:35

	SWDEV-145570 - [HIP] Implemented hipStream create/destroy
	Use the provided stream in hipModuleLaunchKernel

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#4 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#2 edit
---
 hipamd/api/hip/hip_module.cpp |  7 ++----
 hipamd/api/hip/hip_stream.cpp | 40 ++++++++++++++++++++++++++---------
 2 files changed, 32 insertions(+), 15 deletions(-)

diff --git a/hipamd/api/hip/hip_module.cpp b/hipamd/api/hip/hip_module.cpp
index 6fd0cc9ac9..0bcae3551b 100644
--- a/hipamd/api/hip/hip_module.cpp
+++ b/hipamd/api/hip/hip_module.cpp
@@ -144,9 +144,8 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
   amd::Kernel* kernel = as_amd(reinterpret_cast<cl_kernel>(f));
   amd::Device* device = g_context->devices()[0];
 
-  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
-                                             amd::CommandQueue::RealTimeDisabled,
-                                             amd::CommandQueue::Priority::Normal);
+  amd::HostQueue* queue = as_amd(reinterpret_cast<cl_command_queue>(hStream))->asHostQueue();
+
   if (!queue) {
     return hipErrorOutOfMemory;
   }
@@ -184,8 +183,6 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
   command->awaitCompletion();
   command->release();
 
-  queue->release();
-
   return hipSuccess;
 }
 
diff --git a/hipamd/api/hip/hip_stream.cpp b/hipamd/api/hip/hip_stream.cpp
index efecb5174d..46014ebfef 100644
--- a/hipamd/api/hip/hip_stream.cpp
+++ b/hipamd/api/hip/hip_stream.cpp
@@ -24,13 +24,30 @@ THE SOFTWARE.
 
 #include "hip_internal.hpp"
 
+static hipError_t ihipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
+{
+  assert(flags == 0); // we don't handle flags yet
+
+  amd::Device* device = g_context->devices()[0];
+
+  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
+                                             amd::CommandQueue::RealTimeDisabled,
+                                             amd::CommandQueue::Priority::Normal);
+
+  if (queue == nullptr) {
+    return hipErrorOutOfMemory;
+  }
+
+  *stream = reinterpret_cast<hipStream_t>(as_cl(queue));
+
+  return hipSuccess;
+}
+
 hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
 {
   HIP_INIT_API(stream, flags);
 
-  assert(0 && "Unimplemented");
-
-  return hipErrorUnknown;
+  return ihipStreamCreateWithFlags(stream, flags);
 }
 
 
@@ -38,9 +55,7 @@ hipError_t hipStreamCreate(hipStream_t *stream)
 {
   HIP_INIT_API(stream);
 
-  assert(0 && "Unimplemented");
-
-  return hipErrorUnknown;
+  return ihipStreamCreateWithFlags(stream, hipStreamDefault);
 }
 
 
@@ -58,9 +73,14 @@ hipError_t hipStreamSynchronize(hipStream_t stream)
 {
   HIP_INIT_API(stream);
 
-  assert(0 && "Unimplemented");
+  amd::HostQueue* hostQueue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
+  if (hostQueue == nullptr) {
+    return hipErrorUnknown;
+  }
 
-  return hipErrorUnknown;
+  hostQueue->finish();
+
+  return hipSuccess;
 }
 
 
@@ -68,9 +88,9 @@ hipError_t hipStreamDestroy(hipStream_t stream)
 {
   HIP_INIT_API(stream);
 
-  assert(0 && "Unimplemented");
+  as_amd(reinterpret_cast<cl_command_queue>(stream))->release();
 
-  return hipErrorUnknown;
+  return hipSuccess;
 }