P4 to Git Change 1552198 by skudchad@skudchad_rocm on 2018/05/08 18:57:32

SWDEV-145570 - [HIP] - Synchronize the legacy default stream with other blocking streams as per the spec.

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

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_context.cpp#11 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#10 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#5 edit


[ROCm/hip commit: 7176a5fe81]

projects/hip/api/hip/hip_context.cpp

@@ -24,8 +24,6 @@ THE SOFTWARE.
 #include "hip_internal.hpp"
 #include "platform/runtime.hpp"
 #include "utils/versions.hpp"
-#include <stack>
-#include <thread>
 
 std::vector<amd::Context*> g_devices;
 
@@ -35,7 +33,7 @@ thread_local amd::Context* g_context = nullptr;
 thread_local std::stack<amd::Context*> g_ctxtStack;
 std::once_flag g_ihipInitialized;
 
-std::map<amd::Context*,amd::HostQueue*> g_nullStreams;
+std::map<amd::Context*, amd::HostQueue*> g_nullStreams;
 
 void init() {
   if (!amd::Runtime::initialized()) {
@@ -66,6 +64,12 @@ void setCurrentContext(unsigned int index) {
   g_context = g_devices[index];
 }
 
+void syncStreams() {
+  for (const auto& it : streamSet) {
+    it->finish();
+  }
+}
+
 amd::HostQueue* getNullStream() {
   auto stream = g_nullStreams.find(getCurrentContext());
   if (stream == g_nullStreams.end()) {
@@ -76,6 +80,7 @@ amd::HostQueue* getNullStream() {
     g_nullStreams[getCurrentContext()] = queue;
     return queue;
   }
+  syncStreams();
   return stream->second;
 }
 

projects/hip/api/hip/hip_internal.hpp

@@ -24,8 +24,9 @@ THE SOFTWARE.
 #define HIP_SRC_HIP_INTERNAL_H
 
 #include "cl_common.hpp"
-
+#include <unordered_set>
 #include <thread>
+#include <stack>
 
 #define HIP_INIT() \
   std::call_once(hip::g_ihipInitialized, hip::init); \
@@ -57,9 +58,10 @@ namespace hip {
   extern void setCurrentContext(unsigned int index);
 
   extern amd::HostQueue* getNullStream();
+  extern void syncStreams();
 };
 extern std::vector<amd::Context*> g_devices;
-
+extern thread_local std::unordered_set<amd::HostQueue*> streamSet;
 extern hipError_t ihipDeviceGetCount(int* count);
 
 #endif // HIP_SRC_HIP_INTERNAL_H

projects/hip/api/hip/hip_stream.cpp

@@ -21,18 +21,22 @@ THE SOFTWARE.
 */
 
 #include <hip/hip_runtime.h>
-
 #include "hip_internal.hpp"
 
-static hipError_t ihipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) {
-  assert(flags == 0); // we don't handle flags yet
+thread_local std::unordered_set<amd::HostQueue*> streamSet;
 
+static hipError_t ihipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) {
   amd::Device* device = hip::getCurrentContext()->devices()[0];
 
   amd::HostQueue* queue = new amd::HostQueue(*hip::getCurrentContext(), *device, 0,
                                              amd::CommandQueue::RealTimeDisabled,
                                              amd::CommandQueue::Priority::Normal);
 
+  if (!(flags & hipStreamNonBlocking)) {
+    hip::syncStreams();
+    streamSet.insert(queue);
+  }
+
   if (queue == nullptr) {
     return hipErrorOutOfMemory;
   }
@@ -48,23 +52,27 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) {
   return ihipStreamCreateWithFlags(stream, flags);
 }
 
-
 hipError_t hipStreamCreate(hipStream_t *stream) {
   HIP_INIT_API(stream);
 
   return ihipStreamCreateWithFlags(stream, hipStreamDefault);
 }
 
-
 hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) {
   HIP_INIT_API(stream, flags);
 
-  assert(0 && "Unimplemented");
+  amd::HostQueue* hostQueue = reinterpret_cast<amd::HostQueue*>(stream);
+  auto it = streamSet.find(hostQueue);
 
-  return hipErrorUnknown;
+  if(flags != nullptr) {
+    *flags = (it != streamSet.end()) ? hipStreamNonBlocking : hipStreamDefault;
+  } else {
+    return hipErrorInvalidValue;
+  }
+
+  return hipSuccess;
 }
 
-
 hipError_t hipStreamSynchronize(hipStream_t stream) {
   HIP_INIT_API(stream);
 
@@ -85,7 +93,6 @@ hipError_t hipStreamSynchronize(hipStream_t stream) {
   return hipSuccess;
 }
 
-
 hipError_t hipStreamDestroy(hipStream_t stream) {
   HIP_INIT_API(stream);
 
@@ -93,12 +100,14 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
     return hipErrorInvalidResourceHandle;
   }
 
+  amd::HostQueue* hostQueue = reinterpret_cast<amd::HostQueue*>(stream);
+  streamSet.erase(hostQueue);
+
   as_amd(reinterpret_cast<cl_command_queue>(stream))->release();
 
   return hipSuccess;
 }
 
-
 hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) {
   HIP_INIT_API(stream, event, flags);