Add hipStreamQuery

Change-Id: Ib0813b1065feba4fe9ae861d24cfc6f9c5f580be

docs/markdown/hip_terms2.md

@@ -15,3 +15,5 @@ The default device can be set with hipSetDevice.
 - hipify - tool to convert CUDA(R) code to portable C++ code.
 - hipconfig - tool to report various confoguration properties of the target platform.
 
+- nvcc = nvcc compiler, do not capitalize.
+- hcc  = heterogeneous compute compiler, do not capitalize.

include/hcc_detail/hip_hcc.h

@@ -515,6 +515,7 @@ private: // Data
 
     // Friends:
     friend std::ostream& operator<<(std::ostream& os, const ihipStream_t& s);
+    friend hipError_t hipStreamQuery(hipStream_t);
 };
 
 

include/hcc_detail/hip_runtime.h

@@ -622,9 +622,9 @@ __device__ static inline void* memset(void* ptr, uint8_t val, size_t size)
 
 #ifdef __HCC_CPP__
 extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm *lp, const char *kernelNameStr);
-extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp);
-extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp);
-extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp);
+extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr);
+extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp, const char *kernelNameStr);
+extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr);
 extern void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp);
 
 

include/hcc_detail/hip_runtime_api.h

@@ -474,6 +474,20 @@ hipError_t hipStreamCreate(hipStream_t *stream);
 hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags);
 
 
+/**
+ * @brief Return #hipSuccess if all of the operations in the specified @p stream have completed, or #hipErrorNotReady if not.
+ *
+ * @param[in] stream stream to query
+ *
+ * @return #hipSuccess, #hipErrorNotReady
+ *
+ * This is thread-safe and returns a snapshot of the current state of the queue.  However, if other host threads are sending work to the stream,
+ * the status may change immediately after the function is called.  It is typically used for debug.
+ */
+hipError_t hipStreamQuery(hipStream_t stream);
+
+
+
 /**
  * @brief Wait for all commands in stream to complete.
  *
@@ -726,7 +740,7 @@ hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags) __attribute
 hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, unsigned int flags) ;
 
 /**
- *  @brief Return flags associated with host pointer 
+ *  @brief Return flags associated with host pointer
  *
  *  @param[out]  flagsPtr Memory location to store flags
  *  @param[in] hostPtr Host Pointer allocated through hipHostMalloc
@@ -1299,10 +1313,10 @@ hipError_t hipDriverGetVersion(int *driverVersion) ;
  *
  * @param [in] fname
  * @param [out] module
- * 
+ *
  * @returns hipSuccess, hipErrorInvalidValue, hipErrorInvalidContext, hipErrorFileNotFound, hipErrorOutOfMemory, hipErrorSharedObjectInitFailed, hipErrorNotInitialized
  *
- * 
+ *
  */
 hipError_t hipModuleLoad(hipModule_t *module, const char *fname);
 
@@ -1313,7 +1327,7 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname);
  *
  * @returns hipSuccess, hipInvalidValue
  * module is freed and the code objects associated with it are destroyed
- * 
+ *
  */
 
 hipError_t hipModuleUnload(hipModule_t module);
@@ -1325,7 +1339,7 @@ hipError_t hipModuleUnload(hipModule_t module);
  * @param [in] kname
  * @param [out] function
  *
- * @returns hipSuccess, hipErrorInvalidValue, hipErrorInvalidContext, hipErrorNotInitialized, hipErrorNotFound, 
+ * @returns hipSuccess, hipErrorInvalidValue, hipErrorInvalidContext, hipErrorNotInitialized, hipErrorNotFound,
  */
 hipError_t hipModuleGetFunction(hipFunction_t *function, hipModule_t module, const char *kname);
 

src/hip_hcc.cpp

@@ -1404,19 +1404,20 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
 
 void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, const hipStream_t stream)
 {
-    std::stringstream os;
-    os << API_COLOR << "<<hip-api: hipLaunchKernel '" << kernelName << "'"
-        << " gridDim:"  << lp->grid_dim
-        << " groupDim:" << lp->group_dim
-        << " sharedMem:+" << lp->dynamic_group_mem_bytes
-        << " " << *stream
-        << API_COLOR_END << std::endl;
+    if (HIP_ATP_MARKER || (COMPILE_HIP_DB && HIP_TRACE_API)) {
+        std::stringstream os;
+        os  << "<<hip-api: hipLaunchKernel '" << kernelName << "'"
+            << " gridDim:"  << lp->grid_dim
+            << " groupDim:" << lp->group_dim
+            << " sharedMem:+" << lp->dynamic_group_mem_bytes
+            << " " << *stream;
 
-    std::cerr << os.str();
 
-    //fprintf(stderr, KGRN "<<hip-api: hipLaunchKernel '%s' gridDim:(%d,%d,%d) groupDim:(%d,%d,%d) groupMem:+%d %s\n" KNRM, \
-    //        kernelName, lp->grid_dim.x, lp->grid_dim.y, lp->grid_dim.z, lp->group_dim.x, lp->group_dim.y, lp->group_dim.z,
-    //        lp->dynamic_group_mem_bytes, streamString.c_str());
+        if (COMPILE_HIP_DB && HIP_TRACE_API) {
+            std::cerr << API_COLOR << os.str() << API_COLOR_END << std::endl;
+        }
+        SCOPED_MARKER(os.str().c_str(), "HIP", NULL);
+    }
 }
 
 // TODO - data-up to data-down:
@@ -1439,15 +1440,13 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_
     lp->av = &(crit->_av);
     lp->cf = new hc::completion_future;
 
-    if (HIP_TRACE_API) {
-        ihipPrintKernelLaunch(kernelNameStr, lp, stream);
-    }
+    ihipPrintKernelLaunch(kernelNameStr, lp, stream);
 
     return (stream);
 }
 
 
-hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp)
+hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp, const char *kernelNameStr)
 {
     HIP_INIT();
     stream = ihipSyncAndResolveStream(stream);
@@ -1463,11 +1462,13 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, gri
     auto crit = stream->lockopen_preKernelCommand();
     lp->av = &(crit->_av);
     lp->cf = new hc::completion_future;
+
+    ihipPrintKernelLaunch(kernelNameStr, lp, stream);
     return (stream);
 }
 
 
-hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp)
+hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr)
 {
     HIP_INIT();
     stream = ihipSyncAndResolveStream(stream);
@@ -1483,11 +1484,12 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, gri
     auto crit = stream->lockopen_preKernelCommand();
     lp->av = &(crit->_av);
     lp->cf = new hc::completion_future;
+    ihipPrintKernelLaunch(kernelNameStr, lp, stream);
     return (stream);
 }
 
 
-hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp)
+hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr)
 {
     HIP_INIT();
     stream = ihipSyncAndResolveStream(stream);
@@ -1503,6 +1505,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g
     auto crit = stream->lockopen_preKernelCommand();
     lp->av = &(crit->_av);
     lp->cf = new hc::completion_future; // TODO, is this necessary?
+    ihipPrintKernelLaunch(kernelNameStr, lp, stream);
     return (stream);
 }
 

src/hip_stream.cpp

@@ -115,6 +115,27 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
 };
 
 
+//---
+hipError_t hipStreamQuery(hipStream_t stream)
+{
+    HIP_INIT_API(stream);
+
+    // Use default stream if 0 specified:
+    if (stream == hipStreamNull) {
+        ihipCtx_t *device = ihipGetTlsDefaultCtx();
+        stream =  device->_defaultStream;
+    }
+
+    LockedAccessor_StreamCrit_t crit(stream->_criticalData);
+    int pendingOps = crit->_av.get_pending_async_ops();
+
+
+    hipError_t e = (pendingOps > 0) ? hipErrorNotReady : hipSuccess;
+
+    return ihipLogStatus(e);
+}
+
+
 //---
 hipError_t hipStreamSynchronize(hipStream_t stream)
 {

tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp

@@ -33,7 +33,8 @@ public:
     Streamer(size_t numElements);
     ~Streamer();
     void runAsync();
-    void waitComplete();
+    void queryUntilComplete();
+
 
 private:
     T *_A_h;
@@ -66,11 +67,24 @@ void Streamer<T>::runAsync()
     printf ("testing: %s  numElements=%zu size=%6.2fMB\n", __func__, _numElements, _numElements * sizeof(T) / 1024.0/1024.0);
     unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, _numElements);
     hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, _stream, _A_d, _B_d, _C_d, _numElements);
+
+    // Test case where hipStreamWaitEvent waits on same event we just placed into the queue.
     HIPCHECK(hipEventRecord(_event, _stream));
     HIPCHECK(hipStreamWaitEvent(_stream, _event, 0));
-
 }
 
+template <typename T>
+void Streamer<T>::queryUntilComplete()
+{
+    int numQueries = 0;
+    hipError_t e = hipSuccess;
+    do {
+        numQueries++;
+        e = hipStreamQuery(_stream);
+    } while (e != hipSuccess) ;
+
+    printf ("completed after %d queries\n", numQueries);
+};
 
 
 
@@ -113,11 +127,28 @@ int main(int argc, char *argv[])
         streamers.push_back(s);
     }
 
-    for (int i=0; i<p_streams; i++) {
-        streamers[i]->runAsync();
+    if (p_tests & 0x1) {
+        printf ("==> Test 0x1 runAsnc\n");
+        for (int i=0; i<p_streams; i++) {
+            streamers[i]->runAsync();
+        }
+        HIPCHECK(hipDeviceSynchronize());
+    }
+
+    if (p_tests & 0x2) {
+        printf ("==> Test 0x2 queryUntilComplete\n");
+        for (int i=0; i<p_streams; i++) {
+            streamers[i]->runAsync();
+            streamers[i]->queryUntilComplete();
+        }
+        HIPCHECK(hipDeviceSynchronize());
+    }
+
+    if (p_tests & 0x4) {
+        hipStreamQuery(0/* try null stream*/);
+
     }
 
-    HIPCHECK(hipDeviceSynchronize());
 
     passed();
 }