diff --git a/docs/markdown/hip_terms2.md b/docs/markdown/hip_terms2.md index 82174405cd..9603807925 100644 --- a/docs/markdown/hip_terms2.md +++ b/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. diff --git a/include/hcc_detail/hip_hcc.h b/include/hcc_detail/hip_hcc.h index 1a4e9780aa..939d57c062 100644 --- a/include/hcc_detail/hip_hcc.h +++ b/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); }; diff --git a/include/hcc_detail/hip_runtime.h b/include/hcc_detail/hip_runtime.h index 4e67cb7292..547df405a2 100644 --- a/include/hcc_detail/hip_runtime.h +++ b/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); diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h index fc390bbd23..1f4a1fb8a9 100644 --- a/include/hcc_detail/hip_runtime_api.h +++ b/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); diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index f4d20021a9..9d9b667f8b 100644 --- a/src/hip_hcc.cpp +++ b/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 << "<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 << "<grid_dim + << " groupDim:" << lp->group_dim + << " sharedMem:+" << lp->dynamic_group_mem_bytes + << " " << *stream; - std::cerr << os.str(); - //fprintf(stderr, KGRN "<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); } diff --git a/src/hip_stream.cpp b/src/hip_stream.cpp index b5d7d8cf5b..7b3dc31f07 100644 --- a/src/hip_stream.cpp +++ b/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) { diff --git a/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp b/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp index 7148f50628..4ad093a16e 100644 --- a/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp +++ b/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::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 +void Streamer::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; irunAsync(); + if (p_tests & 0x1) { + printf ("==> Test 0x1 runAsnc\n"); + for (int i=0; irunAsync(); + } + HIPCHECK(hipDeviceSynchronize()); + } + + if (p_tests & 0x2) { + printf ("==> Test 0x2 queryUntilComplete\n"); + for (int i=0; irunAsync(); + streamers[i]->queryUntilComplete(); + } + HIPCHECK(hipDeviceSynchronize()); + } + + if (p_tests & 0x4) { + hipStreamQuery(0/* try null stream*/); + } - HIPCHECK(hipDeviceSynchronize()); passed(); }