add IPC event support (#1996)

src/hip_hcc.cpp

@@ -332,12 +332,55 @@ void ihipStream_t::locked_wait() {
     locked_wait(waited);
 };
 
+typedef struct {
+    int previous_read_index;
+    ihipIpcEventShmem_t *shmem;
+    hsa_signal_t signal;
+} callback_data_t;
+
+static void WaitThenDecrementSignal(callback_data_t *data) {
+    int offset = data->previous_read_index % IPC_SIGNALS_PER_EVENT;
+    // While event valid and locked, spin.
+    while (data->shmem->read_index < data->previous_read_index+IPC_SIGNALS_PER_EVENT && data->shmem->signal[offset] != 0) {
+    }
+    hsa_signal_store_relaxed(data->signal, 0);
+    delete data;
+}
+
 // Causes current stream to wait for specified event to complete:
 // Note this does not provide any kind of host serialization.
 void ihipStream_t::locked_streamWaitEvent(ihipEventData_t& ecd) {
     LockedAccessor_StreamCrit_t crit(_criticalData);
 
-    crit->_av.create_blocking_marker(ecd.marker(), hc::accelerator_scope);
+    // if event is an IPC event, it doesn't have a marker
+    // we use a host callback to block stream with a signal wait
+    if (ecd._ipc_shmem) {
+        // create first marker
+        auto cf = crit->_av.create_marker(hc::no_scope);
+        // get its signal
+        auto signal = *reinterpret_cast<hsa_signal_t*>(cf.get_native_handle());
+        // increment its signal value
+        hsa_signal_add_relaxed(signal, 1);
+
+        // create callback that can be passed to hsa_amd_signal_async_handler
+        // this function will host wait on IPC signal, then sets first packet's signal to 0 to indicate completion
+        auto t{new callback_data_t{ecd._ipc_shmem->read_index, ecd._ipc_shmem, signal}};
+
+        // register above callback with HSA runtime to be called when first packet's signal
+        // is decremented from 2 to 1 by CP (or it is already at 1)
+        // the HSA async handler is single threaded, so we can't block, therefore use a detached thread
+        hsa_amd_signal_async_handler(signal, HSA_SIGNAL_CONDITION_EQ, 1,
+            [](hsa_signal_value_t x, void* p) {
+                std::thread(WaitThenDecrementSignal, static_cast<decltype(t)>(p)).detach();
+                return false;
+            }, t);
+
+        // create additional marker that blocks on the first one
+        crit->_av.create_blocking_marker(cf, hc::accelerator_scope);
+    }
+    else {
+        crit->_av.create_blocking_marker(ecd.marker(), hc::accelerator_scope);
+    }
 }