P4 to Git Change 1555193 by cpaquot@cpaquot-ocl-lc-lnx on 2018/05/15 16:19:50

SWDEV-145570 - [HIP] Implemented events

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.cpp#4 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.hpp#2 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#25 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#9 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#9 edit

hipamd/api/hip/hip_event.cpp

@@ -24,37 +24,68 @@ THE SOFTWARE.
 
 #include "hip_event.hpp"
 
+hipError_t ihipEventCreateWithFlags(hipEvent_t* event, unsigned flags) {
+  if (event == nullptr) {
+    return hipErrorInvalidValue;
+  }
+
+  unsigned supportedFlags = hipEventDefault | hipEventBlockingSync | hipEventDisableTiming |
+                            hipEventReleaseToDevice | hipEventReleaseToSystem;
+  const unsigned releaseFlags = (hipEventReleaseToDevice | hipEventReleaseToSystem);
+
+  const bool illegalFlags =
+      (flags & ~supportedFlags) ||             // can't set any unsupported flags.
+      (flags & releaseFlags) == releaseFlags;  // can't set both release flags
+
+  if (!illegalFlags) {
+    hip::Event* e = new hip::Event(flags);
+
+    if (e == nullptr) {
+      return hipErrorOutOfMemory;
+    }
+
+    *event = reinterpret_cast<hipEvent_t>(e);
+  } else {
+    return hipErrorInvalidValue;
+  }
+  return hipSuccess;
+}
+
+hipError_t ihipEventQuery(hipEvent_t event) {
+  if (event == nullptr) {
+    return hipErrorInvalidResourceHandle;
+  }
+
+  hip::Event* e = reinterpret_cast<hip::Event*>(event);
+
+  if (e->event_ == nullptr) {
+    return hipErrorInvalidResourceHandle;
+  }
+
+  e->event_->notifyCmdQueue();
+
+  return (e->event_->status() == CL_COMPLETE)? hipSuccess : hipErrorNotReady;
+}
+
 hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) {
   HIP_INIT_API(event, flags);
 
-  hip::Event* e = new hip::Event(flags);
-
-  if (e == nullptr) {
-    return hipErrorOutOfMemory;
-  }
-
-  *event = reinterpret_cast<hipEvent_t>(e);
-
-  return hipSuccess;
-}
+  return ihipEventCreateWithFlags(event, flags);
+}  
 
 hipError_t hipEventCreate(hipEvent_t* event) {
   HIP_INIT_API(event);
 
-  hip::Event* e = new hip::Event(0);
-
-  if (e == nullptr) {
-    return hipErrorOutOfMemory;
-  }
-
-  *event = reinterpret_cast<hipEvent_t>(e);
-
-  return hipSuccess;
+  return ihipEventCreateWithFlags(event, 0);
 }
 
 hipError_t hipEventDestroy(hipEvent_t event) {
   HIP_INIT_API(event);
 
+  if (event == nullptr) {
+    return hipErrorInvalidResourceHandle;
+  }
+
   delete reinterpret_cast<hip::Event*>(event);
 
   return hipSuccess;
@@ -63,29 +94,83 @@ hipError_t hipEventDestroy(hipEvent_t event) {
 hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) {
   HIP_INIT_API(ms, start, stop);
 
-  return hipErrorUnknown;
+  if (start == nullptr || stop == nullptr) {
+    return hipErrorInvalidResourceHandle;
+  }
+
+  hip::Event* eStart = reinterpret_cast<hip::Event*>(start);
+  hip::Event* eStop  = reinterpret_cast<hip::Event*>(stop);
+
+  if (eStart->event_ == nullptr ||
+      eStop->event_  == nullptr) {
+    return hipErrorInvalidResourceHandle;
+  }
+
+  if ((eStart->flags | eStop->flags) & hipEventDisableTiming) {
+    return hipErrorInvalidResourceHandle;
+  }
+
+  if (ihipEventQuery(start) == hipErrorNotReady ||
+      ihipEventQuery(stop) == hipErrorNotReady) {
+    return hipErrorNotReady;
+  }
+
+  if (ms == nullptr) {
+    return hipErrorInvalidValue;
+  }
+
+  *ms = static_cast<float>(eStop->event_->profilingInfo().submitted_ - eStart->event_->profilingInfo().submitted_)/1000000.f;
+
+  return hipSuccess;
 }
 
 hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
   HIP_INIT_API(event, stream);
 
-  assert(0 && "Unimplemented");
+  if (event == nullptr) {
+    return hipErrorInvalidResourceHandle;
+  }
 
-  return hipErrorUnknown;
+  hip::Event* e = reinterpret_cast<hip::Event*>(event);
+
+  if (stream == nullptr) {
+    e->stream_ = hip::getNullStream();
+  } else {
+    e->stream_ = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
+  }
+  amd::Command* command = (e->flags & hipEventDisableTiming)? new amd::Marker(*e->stream_, true) :
+    new hip::TimerMarker(*e->stream_);
+  command->enqueue();
+
+  if (e->event_ != nullptr) {
+    e->event_->release();
+  }
+
+  e->event_ = &command->event();
+
+  return hipSuccess;
 }
 
 hipError_t hipEventSynchronize(hipEvent_t event) {
   HIP_INIT_API(event);
 
-  assert(0 && "Unimplemented");
+  if (event == nullptr) {
+    return hipErrorInvalidResourceHandle;
+  }
 
-  return hipErrorUnknown;
+  hip::Event* e = reinterpret_cast<hip::Event*>(event);
+
+  if (e->event_ == nullptr) {
+    return hipErrorInvalidResourceHandle;
+  }
+
+  e->event_->awaitCompletion();
+
+  return hipSuccess;
 }
 
 hipError_t hipEventQuery(hipEvent_t event) {
   HIP_INIT_API(event);
 
-  assert(0 && "Unimplemented");
-
-  return hipErrorUnknown;
+  return ihipEventQuery(event);
 }

hipamd/api/hip/hip_event.hpp

@@ -27,12 +27,27 @@ THE SOFTWARE.
 
 namespace hip {
 
+class TimerMarker: public amd::Marker {
+public:
+  TimerMarker(amd::HostQueue& queue) : amd::Marker(queue, true) {
+    profilingInfo_.enabled_ = true;
+    profilingInfo_.callback_ = nullptr;
+    profilingInfo_.start_ = profilingInfo_.end_ = 0;
+  }
+};
+
 class Event {
 public:
-  Event(unsigned int flags) : flags(flags) {}
-  ~Event() {}
+  Event(unsigned int flags) : flags(flags), stream_(nullptr), event_(nullptr) {}
+  ~Event() {
+    if (event_ != nullptr) {
+      event_->release();
+    }
+  }
   unsigned int flags;
-private:
+
+  amd::HostQueue* stream_;
+  amd::Event* event_;
 };
 
 };

hipamd/api/hip/hip_memory.cpp

@@ -68,7 +68,6 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
 
 hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
                       amd::HostQueue& queue, bool isAsync = false) {
-
   amd::Command* command = nullptr;
   amd::Command::EventWaitList waitList;
   size_t sOffset = 0;
@@ -116,15 +115,6 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
     return hipErrorOutOfMemory;
   }
 
-// FIXME: virtualize MemoryCommand::validateMemory()
-#if 0
-  // Make sure we have memory for the command execution
-  if (CL_SUCCESS != command->validateMemory()) {
-    delete command;
-    return hipErrorMemoryAllocation;
-  }
-#endif
-
   command->enqueue();
   if (!isAsync) {
     command->awaitCompletion();
@@ -193,6 +183,7 @@ hipError_t hipFree(void* ptr) {
 hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
   HIP_INIT_API(dst, src, sizeBytes, kind);
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
   return ihipMemcpy(dst, src, sizeBytes, kind, *queue);
 }
@@ -203,8 +194,10 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st
   amd::HostQueue* queue;
 
   if (stream == nullptr) {
+    hip::syncStreams();
     queue = hip::getNullStream();
   } else {
+    hip::getNullStream()->finish();
     queue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
   }
 
@@ -214,6 +207,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st
 hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
   HIP_INIT_API(dst, value, sizeBytes);
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
 
   return ihipMemset(dst, value, sizeBytes, *queue);
@@ -533,6 +527,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co
 hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
   HIP_INIT_API(dst, src, sizeBytes);
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
 
   return ihipMemcpy(reinterpret_cast<void*>(dst), (const void*) src, sizeBytes, hipMemcpyHostToDevice, *queue);
@@ -541,6 +536,7 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
 hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) {
   HIP_INIT_API(dst, src, sizeBytes);
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
 
   return ihipMemcpy(reinterpret_cast<void*>(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToHost, *queue);
@@ -549,6 +545,7 @@ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) {
 hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) {
   HIP_INIT_API(dst, src, sizeBytes);
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
 
   return ihipMemcpy(reinterpret_cast<void*>(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToDevice, *queue);
@@ -557,6 +554,7 @@ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeByte
 hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
   HIP_INIT_API(dst, src, sizeBytes);
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
 
   return ihipMemcpy(reinterpret_cast<void*>(dst), (const void*) src, sizeBytes, hipMemcpyHostToHost, *queue);
@@ -569,8 +567,10 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes,
   amd::HostQueue* queue = nullptr;
 
   if (stream == nullptr) {
+    hip::syncStreams();
     queue = hip::getNullStream();
   } else {
+    hip::getNullStream()->finish();
     queue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
   }
 
@@ -585,8 +585,10 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes,
   amd::HostQueue* queue = nullptr;
 
   if (stream == nullptr) {
+    hip::syncStreams();
     queue = hip::getNullStream();
   } else {
+    hip::getNullStream()->finish();
     queue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
   }
 
@@ -601,8 +603,10 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz
   amd::HostQueue* queue = nullptr;
 
   if (stream == nullptr) {
+    hip::syncStreams();
     queue = hip::getNullStream();
   } else {
+    hip::getNullStream()->finish();
     queue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
   }
 
@@ -617,8 +621,10 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes,
   amd::HostQueue* queue = nullptr;
 
   if (stream == nullptr) {
+    hip::syncStreams();
     queue = hip::getNullStream();
   } else {
+    hip::getNullStream()->finish();
     queue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
   }
 
@@ -723,6 +729,7 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
                        size_t height, hipMemcpyKind kind) {
   HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind);
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
 
   return ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, *queue);
@@ -736,8 +743,10 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
   amd::HostQueue* queue;
 
   if (stream == nullptr) {
+    hip::syncStreams();
     queue = hip::getNullStream();
   } else {
+    hip::getNullStream()->finish();
     queue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
   }
 
@@ -757,6 +766,7 @@ hipError_t hipMemcpyToArray(hipArray* dstArray, size_t wOffset, size_t hOffset,
                             size_t count, hipMemcpyKind kind) {
   HIP_INIT_API(dstArray, wOffset, hOffset, src, count, kind);
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
 
   amd::Command* command = nullptr;
@@ -795,6 +805,7 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffs
                               size_t count, hipMemcpyKind kind) {
   HIP_INIT_API(dst, srcArray, wOffset, hOffset, count, kind);
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
 
   amd::Command* command = nullptr;
@@ -833,6 +844,7 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffs
 hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count) {
   HIP_INIT_API(dstArray, dstOffset, srcHost, count);
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
 
   amd::Command::EventWaitList waitList;
@@ -856,6 +868,7 @@ hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHo
 hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count) {
   HIP_INIT_API(dst, srcArray, srcOffset, count);
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
 
   amd::Command::EventWaitList waitList;
@@ -879,6 +892,7 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t
 hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) {
   HIP_INIT_API(p);
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
 
   size_t byteSize;
@@ -1035,6 +1049,7 @@ hipError_t ihipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t
 hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) {
   HIP_INIT_API(dst, pitch, value, width, height);
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
   return ihipMemset2D(dst, pitch, value, width, height, *queue);
 }
@@ -1045,8 +1060,10 @@ hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value,
 
   amd::HostQueue* queue = nullptr;
   if (stream == nullptr) {
+    hip::syncStreams();
     queue = hip::getNullStream();
   } else {
+    hip::getNullStream()->finish();
     queue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
   }
 
@@ -1060,6 +1077,7 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes
     return hipErrorInvalidValue;
   }
 
+  hip::syncStreams();
   amd::HostQueue* queue = hip::getNullStream();
   size_t offset = 0;
   amd::Command::EventWaitList waitList;

hipamd/api/hip/hip_module.cpp

@@ -146,8 +146,10 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
 
   amd::HostQueue* queue;
   if (hStream == nullptr) {
+    hip::syncStreams();
     queue = hip::getNullStream();
   } else {
+    hip::getNullStream()->finish();
     queue = as_amd(reinterpret_cast<cl_command_queue>(hStream))->asHostQueue();
   }
   if (!queue) {
@@ -195,7 +197,6 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
   }
 
   command->enqueue();
-  command->awaitCompletion();
   command->release();
 
   return hipSuccess;

hipamd/api/hip/hip_stream.cpp

@@ -22,6 +22,7 @@ THE SOFTWARE.
 
 #include <hip/hip_runtime.h>
 #include "hip_internal.hpp"
+#include "hip_event.hpp"
 #include "thread/monitor.hpp"
 
 static amd::Monitor streamSetLock("Guards global stream set");
@@ -97,8 +98,12 @@ hipError_t hipStreamSynchronize(hipStream_t stream) {
   amd::HostQueue* hostQueue;
 
   if (stream == nullptr) {
+    hip::syncStreams();
+
     hostQueue = hip::getNullStream();
   } else {
+    hip::getNullStream()->finish();
+
     hostQueue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
   }
 
@@ -130,9 +135,29 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
 hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) {
   HIP_INIT_API(stream, event, flags);
 
-  assert(0 && "Unimplemented");
+  if (stream == nullptr || event == nullptr) {
+    return hipErrorInvalidResourceHandle;
+  }
 
-  return hipErrorUnknown;
+  amd::HostQueue* hostQueue = as_amd(reinterpret_cast<cl_command_queue>(stream))->asHostQueue();
+  hip::Event* e = reinterpret_cast<hip::Event*>(event);
+  cl_event clEvent = as_cl(e->event_);
+
+  amd::Command::EventWaitList eventWaitList;
+  cl_int err = amd::clSetEventWaitList(eventWaitList, hostQueue->context(), 1,
+                                       &clEvent);
+  if (err != CL_SUCCESS) {
+    return hipErrorUnknown;
+  }
+
+  amd::Command* command = new amd::Marker(*hostQueue, true, eventWaitList);
+  if (command == NULL) {
+    return hipErrorOutOfMemory;
+  }
+  command->enqueue();
+  command->release();
+
+  return hipSuccess;
 }
 
 hipError_t hipStreamQuery(hipStream_t stream) {