From 290b5f17382b6fd192aed06de6095621c4c97fef Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 19 Dec 2017 16:06:14 +0530 Subject: [PATCH 1/3] Implement hipStreamAddCallback Change-Id: Ib851e4d86ba9c8406ca37b88162ea483ccbc9d36 [ROCm/clr commit: cd9ba0d1e18ca6bc8242357e8d10346b324e27a3] --- projects/clr/hipamd/src/hip_hcc.cpp | 30 +++++++++ projects/clr/hipamd/src/hip_hcc_internal.h | 20 ++++++ projects/clr/hipamd/src/hip_stream.cpp | 36 +++++----- .../stream/hipStreamAddCallback.cpp | 66 +++++++++++++------ 4 files changed, 114 insertions(+), 38 deletions(-) diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index f15a0eb1d8..0ca170152b 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -37,6 +37,7 @@ THE SOFTWARE. #include #include #include +#include #include #include @@ -1409,9 +1410,38 @@ void ihipInit() tprintf(DB_SYNC, "pid=%u %-30s g_numLogicalThreads=%u\n", getpid(), "", g_numLogicalThreads); } +hipError_t ihipStreamSynchronize(hipStream_t stream) +{ + hipError_t e = hipSuccess; + if (stream == hipStreamNull) { + ihipCtx_t *ctx = ihipGetTlsDefaultCtx(); + ctx->locked_syncDefaultStream(true/*waitOnSelf*/, true/*syncToHost*/); + } else { + // note this does not synchornize with the NULL stream: + stream->locked_wait(); + e = hipSuccess; + } + return e; +} +void ihipStreamCallbackHandler(ihipStreamCallback_t *cb) +{ + hipError_t e = hipSuccess; + + // Notify hipStreamAddCallback that callback handler thread is active + std::lock_guard guard(cb->_mtx); + cb->_ready = true; + + // Synchronize stream + tprintf(DB_SYNC, "ihipStreamCallbackHandler wait on stream %s\n", ToString(cb->_stream).c_str()); + e = ihipStreamSynchronize(cb->_stream); + + // Call registered callback function + cb->_callback(cb->_stream, e, cb->_userData); + delete cb; +} //--- // Get the stream to use for a command submission. diff --git a/projects/clr/hipamd/src/hip_hcc_internal.h b/projects/clr/hipamd/src/hip_hcc_internal.h index 4891f54fee..601b66f343 100644 --- a/projects/clr/hipamd/src/hip_hcc_internal.h +++ b/projects/clr/hipamd/src/hip_hcc_internal.h @@ -622,6 +622,24 @@ private: // Data }; +//---- +// Internal structure for stream callback handler +class ihipStreamCallback_t { +public: + ihipStreamCallback_t(hipStream_t stream, hipStreamCallback_t callback, void *userData) : + _stream(stream), + _callback(callback), + _userData(userData) + { + _ready = false; + }; + hipStream_t _stream; + hipStreamCallback_t _callback; + void* _userData; + std::mutex _mtx; + bool _ready; +}; + //---- // Internal event structure: @@ -931,6 +949,8 @@ ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex); hipStream_t ihipSyncAndResolveStream(hipStream_t); +hipError_t ihipStreamSynchronize(hipStream_t stream); +void ihipStreamCallbackHandler(ihipStreamCallback_t *cb); // Stream printf functions: inline std::ostream& operator<<(std::ostream& os, const ihipStream_t& s) diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index dab31dad62..94fc436b75 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -20,6 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +#include +#include #include "hip/hip_runtime.h" #include "hip_hcc_internal.h" #include "trace_helper.h" @@ -147,20 +149,8 @@ hipError_t hipStreamSynchronize(hipStream_t stream) { HIP_INIT_SPECIAL_API(TRACE_SYNC, stream); - hipError_t e = hipSuccess; - - if (stream == hipStreamNull) { - ihipCtx_t *ctx = ihipGetTlsDefaultCtx(); - ctx->locked_syncDefaultStream(true/*waitOnSelf*/, true/*syncToHost*/); - } else { - // note this does not synchornize with the NULL stream: - stream->locked_wait(); - e = hipSuccess; - } - - - return ihipLogStatus(e); -}; + return ihipLogStatus(ihipStreamSynchronize(stream)); +} //--- @@ -216,8 +206,20 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback { HIP_INIT_API(stream, callback, userData, flags); hipError_t e = hipSuccess; - //--- explicitly synchronize stream to add callback routines - hipStreamSynchronize(stream); - callback(stream, e, userData); + + // Create a thread in detached mode to handle callback + ihipStreamCallback_t *cb = new ihipStreamCallback_t(stream, callback, userData); + std::thread (ihipStreamCallbackHandler, cb).detach(); + + // Wait for thread to be ready + cb->_mtx.lock(); + while(cb->_ready != true) + { + cb->_mtx.unlock(); + std::this_thread::sleep_for(std::chrono::milliseconds(10)); + cb->_mtx.lock(); + } + cb->_mtx.unlock(); + return ihipLogStatus(e); } diff --git a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp index 32a2793479..692d090509 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp @@ -23,8 +23,7 @@ THE SOFTWARE. * HIT_END */ -// Test under-development. Call hipStreamAddCallback function and see if it works as expected. - +#include #include "hip/hip_runtime.h" #include "test_common.h" @@ -32,32 +31,57 @@ THE SOFTWARE. #define HIPRT_CB #endif -class CallbackClass +__global__ void vector_square(float *C_d, float *A_d, size_t N) { -public: - static void HIPRT_CB Callback(hipStream_t stream, hipError_t status, void *userData); + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x ; -private: - void callbackFunc(hipError_t status); -}; - -void HIPRT_CB CallbackClass::Callback(hipStream_t stream, hipError_t status, void *userData) -{ - CallbackClass* obj = (CallbackClass*) userData; - obj->callbackFunc(status); + for (size_t i=offset; i Date: Tue, 30 Jan 2018 12:32:46 +0530 Subject: [PATCH 2/3] Fix hipStreamAddCallback testcase for nvcc Change-Id: Ieec4b8d7933d8d68394d21d27132da206111efc8 [ROCm/clr commit: 6a37c60eaa11832d2c99b63465645fa090e6ce3c] --- .../tests/src/runtimeApi/stream/hipStreamAddCallback.cpp | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp index 692d090509..0f9f32a253 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp @@ -24,6 +24,7 @@ THE SOFTWARE. */ #include +#include #include "hip/hip_runtime.h" #include "test_common.h" @@ -42,16 +43,17 @@ __global__ void vector_square(float *C_d, float *A_d, size_t N) } float *A_h, *C_h; -size_t N = 1000000; +bool cbDone = false; static void HIPRT_CB Callback(hipStream_t stream, hipError_t status, void *userData) { for (size_t i=0; i Date: Wed, 31 Jan 2018 10:48:17 +0530 Subject: [PATCH 3/3] hipStreamAddCallback.cpp: Replace unistd sleep with sleep_for Change-Id: I7a5d40a1acd8be76a0f175bcfa731ad89fb88d81 [ROCm/clr commit: 3aab52d7ccc37146a82642b03d0395c387bc2aa8] --- .../tests/src/runtimeApi/stream/hipStreamAddCallback.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp index 0f9f32a253..02912e14c6 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp @@ -18,13 +18,14 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../../test_common.cpp + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 * RUN: %t * HIT_END */ #include -#include +#include +#include #include "hip/hip_runtime.h" #include "test_common.h" @@ -87,5 +88,6 @@ int main(int argc, char *argv[]) HIPCHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream)); HIPCHECK(hipStreamAddCallback(mystream, Callback, NULL, 0)); - while(!cbDone) sleep(1); + while(!cbDone) + std::this_thread::sleep_for(std::chrono::milliseconds(10)); }