AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

SWDEV-350025 - Added hipLaunchHostFunc functionality

Telusuri Sumber 
Change-Id: I6b9d547f7ddf84e617e35ff5d5a490ff458272ef


[ROCm/clr commit: 52a22a1263]
This commit is contained in:
Anusha Godavarthy Surya
2022-08-29 12:12:26 +00:00
orang tua 5053dbe40e
melakukan 1cb6a65b16
9 mengubah file dengan 89 tambahan dan 29 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h
+4 -1
Melihat File
@@ -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
#endif //HIP_INCLUDE_HIP_HIP_RUNTIME_PT_API_H
projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h
+3
Melihat File
@@ -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) { \
projects/clr/hipamd/src/amdhip.def
+2 -1
Melihat File
@@ -432,4 +432,5 @@ hipUserObjectRelease
hipUserObjectRetain
hipGraphRetainUserObject
hipGraphReleaseUserObject
hipLaunchHostFunc
hipLaunchHostFunc_spt
projects/clr/hipamd/src/hip_event.hpp
+33 -8
Melihat File
@@ -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);
projects/clr/hipamd/src/hip_event_ipc.cpp
+3 -2
Melihat File
@@ -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<hipStreamCallback_t>(WaitThenDecrementSignal), t, command);
StreamCallback* cbo = new StreamAddCallback(
stream, reinterpret_cast<hipStreamCallback_t>(WaitThenDecrementSignal), t);
if (!command->setCallback(CL_COMPLETE, ihipStreamCallback, cbo)) {
command->release();
return hipErrorInvalidHandle;
}
command->enqueue();
command->release();
command->awaitCompletion();
return hipSuccess;
}
projects/clr/hipamd/src/hip_graph_capture.hpp
+2
Melihat File
@@ -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);
projects/clr/hipamd/src/hip_hcc.def.in
+2
Melihat File
@@ -446,3 +446,5 @@ hipUserObjectRelease
hipUserObjectRetain
hipGraphRetainUserObject
hipGraphReleaseUserObject
hipLaunchHostFunc
hipLaunchHostFunc_spt
projects/clr/hipamd/src/hip_hcc.map.in
+2
Melihat File
@@ -500,6 +500,8 @@ global:
hipUserObjectRetain;
hipGraphRetainUserObject;
hipGraphReleaseUserObject;
hipLaunchHostFunc;
hipLaunchHostFunc_spt;
local:
*;
} hip_5.2;
projects/clr/hipamd/src/hip_stream.cpp
+38 -17
Melihat File
@@ -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<StreamCallback*>(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) {