Beispiel #1
0
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) {
    HIP_INIT_SPECIAL_API(hipStreamWaitEvent, TRACE_SYNC, stream, event, flags);

    hipError_t e = hipSuccess;

    auto ecd = event->locked_copyCrit();

    if (event == nullptr) {
        e = hipErrorInvalidResourceHandle;

    } else if ((ecd._state != hipEventStatusUnitialized) && (ecd._state != hipEventStatusCreated)) {
        if (HIP_SYNC_STREAM_WAIT || (HIP_SYNC_NULL_STREAM && (stream == 0))) {
            // conservative wait on host for the specified event to complete:
            // return _stream->locked_eventWaitComplete(this, waitMode);
            //
            ecd._stream->locked_eventWaitComplete(
                ecd.marker(), (event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked
                                                                     : hc::hcWaitModeActive);
        } else {
            stream = ihipSyncAndResolveStream(stream);
            // This will use create_blocking_marker to wait on the specified queue.
            stream->locked_streamWaitEvent(ecd);
        }

    }  // else event not recorded, return immediately and don't create marker.

    return ihipLogStatus(e);
};