Improve debug print messages.

- Remove "call-to-call" for hipStreamCreate and hipEventCreate.
  These now call an internal functions rather than calling through
  hipStreamCreateWithFalgs and hipEventCreateWithFlags.

- Add HIP_INIT_API for more functions so they trace correctly.

- Use stream#DEVICE.STREAMID in debug messages via new specialization in
  tace_helper.


[ROCm/clr commit: 82f57ca610]

projects/clr/hipamd/include/hcc_detail/hip_hcc.h

@@ -341,7 +341,7 @@ typedef uint64_t SeqNum_t ;
 
     //-- Non-racy accessors:
     // These functions access fields set at initialization time and are non-racy (so do not acquire mutex)
-    ihipDevice_t *       getDevice() const;
+    ihipDevice_t *              getDevice() const;
     StreamMutex &               mutex() {return _mutex;};
 
     //---
@@ -375,9 +375,20 @@ private:
     std::deque<ihipSignal_t>    _signalPool;   // Pool of signals for use by this stream.
 
     StreamMutex                  _mutex;
+
+    friend std::ostream& operator<<(std::ostream& os, const ihipStream_t& s);
 };
 
 
+inline std::ostream& operator<<(std::ostream& os, const ihipStream_t& s)
+{
+    os << "stream#";
+    os << s._device_index;
+    os << '.';
+    os << s._id;
+    return os;
+}
+
 
 //----
 // Internal event structure:

projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h

@@ -443,10 +443,7 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags);
  * @see hipStreamDestroy
  *
  */
-static inline hipError_t hipStreamCreate(hipStream_t *stream)
-{
-    return hipStreamCreateWithFlags(stream, hipStreamDefault);
-}
+hipError_t hipStreamCreate(hipStream_t *stream);
 
 
 /**
@@ -544,13 +541,10 @@ hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags);
 /**
  *  Create an event
  *
- *  @param[in] event Creates an event
+ * @param[in,out] event Returns the newly created event.
  *
  */
-static inline hipError_t hipEventCreate(hipEvent_t* event)
-{
-    return hipEventCreateWithFlags(event, 0);
-}
+hipError_t hipEventCreate(hipEvent_t* event);
 
 
 /**

projects/clr/hipamd/include/hcc_detail/trace_helper.h

@@ -44,16 +44,38 @@ std::string ToHexString(T v)
 
 
 //---
-// Template overloads for ToString to handle various types:
-// Note these use C++11 variadic templates
+// Template overloads for ToString to handle specific types
+
+// This is the default which works for most types:
 template <typename T>
-std::string ToString(T v) {
+std::string ToString(T v) 
+{
     std::ostringstream ss;
     ss << v;
     return ss.str();
 };
 
 
+//  hipEvent_t specialization. TODO - maybe add an event ID for debug?
+template <>
+std::string ToString(hipEvent_t v) 
+{
+    return ToString(&v);
+};
+
+
+
+//  hipStream_t
+template <>
+std::string ToString(hipStream_t v) 
+{
+    std::ostringstream ss;
+    ss << *v;
+
+    return ss.str();
+};
+
+//  hipMemcpyKind specialization
 template <>
 std::string ToString(hipMemcpyKind v) {
     switch(v) {

projects/clr/hipamd/src/hip_event.cpp

@@ -25,16 +25,13 @@ THE SOFTWARE.
 //-------------------------------------------------------------------------------------------------
 // Events
 //---
-/**
- * @warning : flags must be 0.
- */
-hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags)
-{
-    // TODO - support hipEventDefault, hipEventBlockingSync, hipEventDisableTiming
-    std::call_once(hip_initialized, ihipInit);
 
+
+hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags)
+{
     hipError_t e = hipSuccess;
 
+    // TODO - support hipEventDefault, hipEventBlockingSync, hipEventDisableTiming
     if (flags == 0) {
         ihipEvent_t *eh = event->_handle = new ihipEvent_t();
 
@@ -47,8 +44,25 @@ hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags)
         e = hipErrorInvalidValue;
     }
 
+    return e;
+}
 
-    return ihipLogStatus(e);
+/**
+ * @warning : flags must be 0.
+ */
+hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags)
+{
+    HIP_INIT_API(event, flags);
+
+    return ihipLogStatus(ihipEventCreate(event, flags));
+}
+
+
+hipError_t hipEventCreate(hipEvent_t* event)
+{
+    HIP_INIT_API(event);
+
+    return ihipLogStatus(ihipEventCreate(event, 0));
 }
 
 

projects/clr/hipamd/src/hip_stream.cpp

@@ -28,10 +28,8 @@ THE SOFTWARE.
 //
 
 //---
-hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
+hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags)
 {
-    std::call_once(hip_initialized, ihipInit);
-
     ihipDevice_t *device = ihipGetTlsDefaultDevice();
     hc::accelerator acc = device->_acc;
 
@@ -48,7 +46,25 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
     *stream = istream;
     tprintf(DB_SYNC, "hipStreamCreate, stream=%p\n", *stream);
 
-    return ihipLogStatus(hipSuccess);
+    return hipSuccess;
+}
+
+
+//---
+hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
+{
+    HIP_INIT_API(stream, flags);
+
+    return ihipLogStatus(ihipStreamCreate(stream, flags));
+
+}
+
+//---
+hipError_t hipStreamCreate(hipStream_t *stream) 
+{
+    HIP_INIT_API(stream);
+
+    return ihipLogStatus(ihipStreamCreate(stream, hipStreamDefault));
 }
 
 
@@ -58,8 +74,7 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
  */
 hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags)
 {
-
-    std::call_once(hip_initialized, ihipInit);
+    HIP_INIT_API(stream, event, flags);
 
     hipError_t e = hipSuccess;
 
@@ -78,7 +93,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
 //---
 hipError_t hipStreamSynchronize(hipStream_t stream)
 {
-    std::call_once(hip_initialized, ihipInit);
+    HIP_INIT_API(stream);
 
     hipError_t e = hipSuccess;
 
@@ -101,7 +116,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream)
  */
 hipError_t hipStreamDestroy(hipStream_t stream)
 {
-    std::call_once(hip_initialized, ihipInit);
+    HIP_INIT_API(stream);
 
     hipError_t e = hipSuccess;
 
@@ -130,7 +145,7 @@ hipError_t hipStreamDestroy(hipStream_t stream)
 //---
 hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags)
 {
-    std::call_once(hip_initialized, ihipInit);
+    HIP_INIT_API(stream, flags);
 
     if (flags == NULL) {
         return ihipLogStatus(hipErrorInvalidValue);

projects/clr/hipamd/util/vim/hip.vim

@@ -150,6 +150,7 @@ syn keyword hipFunctionName hipMemset2D
 syn keyword hipFunctionName hipMemset3D 
 syn keyword hipFunctionName hipSetDevice 
 syn keyword hipFunctionName hipSetupArgument 
+syn keyword hipFunctionName hipStreamCreateWithFlags
 syn keyword hipFunctionName hipStreamCreate 
 syn keyword hipFunctionName hipStreamDestroy 
 syn keyword hipFunctionName hipStreamQuery