AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter

SWDEV-24806 - Added support for event wait and event record graph APIs

Bläddra i källkod 
hip graph implementation for APIs AddEventRecordNode, AddEventWaitNode, EventRecordNodeGetEvent, EventWaitNodeGetEvent, EventRecordNodeSetEvent,EventWaitNodeSetEvent, ExecEventRecordNodeSetEvent, ExecEventWaitNodeSetEvent

Change-Id: Ifde51f4591067795d6cbf5595372cbdbd8a7bc26
This commit is contained in:
anusha GodavarthySurya
2021-10-05 10:06:40 -07:00
förälder c4eda85835
incheckning f6456c5e27
11 ändrade filer med 353 tillägg och 157 borttagningar
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/src/hip_stream.cpp
+8 -29
Visa fil
@@ -415,18 +415,13 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
HIP_RETURN(hipSuccess);
}
struct CallbackData {
int previous_read_index;
hip::ihipIpcEventShmem_t *shmem;
};
void WaitThenDecrementSignal(hipStream_t stream, hipError_t status, void* user_data){
CallbackData *data = (CallbackData *)user_data;
int offset = data->previous_read_index % IPC_SIGNALS_PER_EVENT;
while (data->shmem->read_index < data->previous_read_index+IPC_SIGNALS_PER_EVENT
&& data->shmem->signal[offset] != 0) {
}
delete data;
void WaitThenDecrementSignal(hipStream_t stream, hipError_t status, void* user_data) {
CallbackData* data = (CallbackData*)user_data;
int offset = data->previous_read_index % IPC_SIGNALS_PER_EVENT;
while (data->shmem->read_index < data->previous_read_index + IPC_SIGNALS_PER_EVENT &&
data->shmem->signal[offset] != 0) {
}
delete data;
}
// ================================================================================================
@@ -443,24 +438,8 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
HIP_RETURN(hipErrorContextIsDestroyed);
}
amd::HostQueue* queue = hip::getQueue(stream);
hip::Event* e = reinterpret_cast<hip::Event*>(event);
if (e->flags & hipEventInterprocess) {
amd::Command* command = new amd::Marker(*queue, false);
auto t{new CallbackData{e->ipc_evt_.ipc_shmem_->read_index, e->ipc_evt_.ipc_shmem_}};
StreamCallback* cbo = new StreamCallback(stream,
reinterpret_cast<hipStreamCallback_t> (WaitThenDecrementSignal), t, command);
if (!command->setCallback(CL_COMPLETE, ihipStreamCallback,cbo)) {
command->release();
return hipErrorInvalidHandle;
}
command->enqueue();
command->awaitCompletion();
HIP_RETURN(hipSuccess);
} else {
HIP_RETURN(e->streamWait(queue, flags));
}
HIP_RETURN(e->streamWait(stream, flags));
}
// ================================================================================================