diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h index d2a0f7247b..544cd3c8a7 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h @@ -66,6 +66,7 @@ THE SOFTWARE. #define hipStreamGetPriority __HIP_API_SPT(hipStreamGetPriority) #define hipStreamWaitEvent __HIP_API_SPT(hipStreamWaitEvent) #define hipStreamAddCallback __HIP_API_SPT(hipStreamAddCallback) + #define hipLaunchHostFunc __HIP_API_SPT(hipLaunchHostFunc) // Event APIs #define hipEventRecord __HIP_API_SPT(hipEventRecord) @@ -179,10 +180,12 @@ hipError_t hipStreamGetCaptureInfo_v2_spt(hipStream_t stream, hipStreamCaptureSt unsigned long long* id_out, hipGraph_t* graph_out, const hipGraphNode_t** dependencies_out, size_t* numDependencies_out); +hipError_t hipLaunchHostFunc_spt(hipStream_t stream, hipHostFn_t fn, void* userData); + #ifdef __cplusplus } #endif // extern "C" #endif //(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__)) -#endif //HIP_INCLUDE_HIP_HIP_RUNTIME_PT_API_H \ No newline at end of file +#endif //HIP_INCLUDE_HIP_HIP_RUNTIME_PT_API_H diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h index 02a92e32bb..73c2fd124d 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -4177,6 +4177,9 @@ typedef struct hip_api_data_s { }; // hipLaunchHostFunc[('hipStream_t', 'stream'), ('hipHostFn_t', 'fn'), ('void*', 'userData')] #define INIT_hipLaunchHostFunc_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipLaunchHostFunc.stream = (hipStream_t)stream; \ + cb_data.args.hipLaunchHostFunc.fn = (hipHostFn_t)fn; \ + cb_data.args.hipLaunchHostFunc.userData = (void*)userData; \ }; // hipLaunchKernel[('const void*', 'function_address'), ('dim3', 'numBlocks'), ('dim3', 'dimBlocks'), ('void**', 'args'), ('size_t', 'sharedMemBytes'), ('hipStream_t', 'stream')] #define INIT_hipLaunchKernel_CB_ARGS_DATA(cb_data) { \ diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index a6890da096..9f06a70a36 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -432,4 +432,5 @@ hipUserObjectRelease hipUserObjectRetain hipGraphRetainUserObject hipGraphReleaseUserObject - +hipLaunchHostFunc +hipLaunchHostFunc_spt diff --git a/projects/clr/hipamd/src/hip_event.hpp b/projects/clr/hipamd/src/hip_event.hpp index 4f50f1d523..d6418b1607 100644 --- a/projects/clr/hipamd/src/hip_event.hpp +++ b/projects/clr/hipamd/src/hip_event.hpp @@ -24,17 +24,42 @@ #include "hip_internal.hpp" #include "thread/monitor.hpp" - // Internal structure for stream callback handler class StreamCallback { - public: - StreamCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData, - amd::Command* command) - : stream_(stream), callBack_(callback), userData_(userData), command_(command){}; - hipStream_t stream_; - hipStreamCallback_t callBack_; +protected: void* userData_; - amd::Command* command_; + public: + StreamCallback(void* userData) + : userData_(userData) {} + + virtual void CL_CALLBACK callback() = 0; +}; + +class StreamAddCallback : public StreamCallback { + hipStreamCallback_t callBack_; + hipStream_t stream_; +public: + StreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData) + : StreamCallback(userData) { + stream_ = stream; + callBack_ = callback; + } + + void CL_CALLBACK callback() { + hipError_t status = hipSuccess; + callBack_(stream_, status, userData_); + } +}; + +class LaunchHostFuncCallback : public StreamCallback { + hipHostFn_t callBack_; + public: + LaunchHostFuncCallback(hipHostFn_t callback, void* userData) + : StreamCallback(userData) { + callBack_ = callback; + } + + void CL_CALLBACK callback() { callBack_(userData_); } }; void CL_CALLBACK ihipStreamCallback(cl_event event, cl_int command_exec_status, void* user_data); diff --git a/projects/clr/hipamd/src/hip_event_ipc.cpp b/projects/clr/hipamd/src/hip_event_ipc.cpp index 5f0ec0c6e5..1aa658b185 100644 --- a/projects/clr/hipamd/src/hip_event_ipc.cpp +++ b/projects/clr/hipamd/src/hip_event_ipc.cpp @@ -112,13 +112,14 @@ hipError_t IPCEvent::streamWaitCommand(amd::Command*& command, amd::HostQueue* q hipError_t IPCEvent::enqueueStreamWaitCommand(hipStream_t stream, amd::Command* command) { auto t{new CallbackData{ipc_evt_.ipc_shmem_->read_index, ipc_evt_.ipc_shmem_}}; - StreamCallback* cbo = new StreamCallback( - stream, reinterpret_cast(WaitThenDecrementSignal), t, command); + StreamCallback* cbo = new StreamAddCallback( + stream, reinterpret_cast(WaitThenDecrementSignal), t); if (!command->setCallback(CL_COMPLETE, ihipStreamCallback, cbo)) { command->release(); return hipErrorInvalidHandle; } command->enqueue(); + command->release(); command->awaitCompletion(); return hipSuccess; } diff --git a/projects/clr/hipamd/src/hip_graph_capture.hpp b/projects/clr/hipamd/src/hip_graph_capture.hpp index 9fbcd01224..39ddab15d0 100644 --- a/projects/clr/hipamd/src/hip_graph_capture.hpp +++ b/projects/clr/hipamd/src/hip_graph_capture.hpp @@ -103,3 +103,5 @@ hipError_t capturehipMemset3DAsync(hipStream_t& stream, hipPitchedPtr& pitchedDe hipError_t capturehipEventRecord(hipStream_t& stream, hipEvent_t& event); hipError_t capturehipStreamWaitEvent(hipEvent_t& event, hipStream_t& stream, unsigned int& flags); + +hipError_t capturehipLaunchHostFunc(hipStream_t& stream, hipHostFn_t& fn, void*& userData); diff --git a/projects/clr/hipamd/src/hip_hcc.def.in b/projects/clr/hipamd/src/hip_hcc.def.in index 73be353b75..469d99d1aa 100644 --- a/projects/clr/hipamd/src/hip_hcc.def.in +++ b/projects/clr/hipamd/src/hip_hcc.def.in @@ -446,3 +446,5 @@ hipUserObjectRelease hipUserObjectRetain hipGraphRetainUserObject hipGraphReleaseUserObject +hipLaunchHostFunc +hipLaunchHostFunc_spt diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index cd991dff21..965321884b 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -500,6 +500,8 @@ global: hipUserObjectRetain; hipGraphRetainUserObject; hipGraphReleaseUserObject; + hipLaunchHostFunc; + hipLaunchHostFunc_spt; local: *; } hip_5.2; diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 2772969178..5e51a1ebd2 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -265,10 +265,8 @@ void iHipWaitActiveStreams(amd::HostQueue* blocking_queue, bool wait_null_stream // ================================================================================================ void CL_CALLBACK ihipStreamCallback(cl_event event, cl_int command_exec_status, void* user_data) { - hipError_t status = hipSuccess; StreamCallback* cbo = reinterpret_cast(user_data); - cbo->callBack_(cbo->stream_, status, cbo->userData_); - cbo->command_->release(); + cbo->callback(); delete cbo; } @@ -573,14 +571,7 @@ hipError_t hipStreamQuery_spt(hipStream_t stream) { HIP_RETURN(hipStreamQuery_common(stream)); } -// ================================================================================================ -hipError_t hipStreamAddCallback_common(hipStream_t stream, hipStreamCallback_t callback, void* userData, - unsigned int flags) { - //flags - Reserved for future use, must be 0 - if (callback == nullptr || flags != 0) { - return hipErrorInvalidValue; - } - +hipError_t streamCallback_common(hipStream_t stream, StreamCallback* cbo, void* userData) { if (!hip::isValid(stream)) { return hipErrorContextIsDestroyed; } @@ -595,8 +586,6 @@ hipError_t hipStreamAddCallback_common(hipStream_t stream, hipStreamCallback_t c if (command == nullptr) { return hipErrorInvalidValue; } - StreamCallback* cbo = new StreamCallback(stream, callback, userData, command); - if ((cbo == nullptr) || !command->setCallback(CL_COMPLETE, ihipStreamCallback, cbo)) { command->release(); if (last_command != nullptr) { @@ -604,8 +593,6 @@ hipError_t hipStreamAddCallback_common(hipStream_t stream, hipStreamCallback_t c } return hipErrorInvalidHandle; } - // Retain callback command for the blocking marker - command->retain(); command->enqueue(); // @note: don't release the command here, because it will be released after HIP callback if (last_command != nullptr) { @@ -626,6 +613,17 @@ hipError_t hipStreamAddCallback_common(hipStream_t stream, hipStreamCallback_t c return hipSuccess; } +// ================================================================================================ +hipError_t hipStreamAddCallback_common(hipStream_t stream, hipStreamCallback_t callback, + void* userData, unsigned int flags) { + // flags - Reserved for future use, must be 0 + if (callback == nullptr || flags != 0) { + return hipErrorInvalidValue; + } + StreamCallback* cbo = new StreamAddCallback(stream, callback, userData); + return streamCallback_common(stream, cbo, userData); +} + // ================================================================================================ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData, unsigned int flags) { @@ -634,13 +632,36 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback } // ================================================================================================ -hipError_t hipStreamAddCallback_spt(hipStream_t stream, hipStreamCallback_t callback, void* userData, - unsigned int flags) { +hipError_t hipStreamAddCallback_spt(hipStream_t stream, hipStreamCallback_t callback, + void* userData, unsigned int flags) { HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags); PER_THREAD_DEFAULT_STREAM(stream); HIP_RETURN(hipStreamAddCallback_common(stream, callback, userData, flags)); } +// ================================================================================================ +hipError_t hipLaunchHostFunc_common(hipStream_t stream, hipHostFn_t fn, void* userData) { + STREAM_CAPTURE(hipLaunchHostFunc, stream, fn, userData); + if (fn == nullptr) { + return hipErrorInvalidValue; + } + StreamCallback* cbo = new LaunchHostFuncCallback(fn, userData); + return streamCallback_common(stream, cbo, userData); +} + +// ================================================================================================ +hipError_t hipLaunchHostFunc_spt(hipStream_t stream, hipHostFn_t fn, void* userData) { + HIP_INIT_API(hipLaunchHostFunc, stream, fn, userData); + PER_THREAD_DEFAULT_STREAM(stream); + HIP_RETURN(hipLaunchHostFunc_common(stream, fn, userData)); +} + +// ================================================================================================ +hipError_t hipLaunchHostFunc(hipStream_t stream, hipHostFn_t fn, void* userData) { + HIP_INIT_API(hipLaunchHostFunc, stream, fn, userData); + HIP_RETURN(hipLaunchHostFunc_common(stream, fn, userData)); +} + // ================================================================================================ hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize, const uint32_t* cuMask) {