HSAKMT_STATUS HSAKMTAPI hsaKmtWaitOnMultipleEvents( HsaEvent* Events[], //IN HSAuint32 NumEvents, //IN bool WaitOnAll, //IN HSAuint32 Milliseconds //IN ) { CHECK_KFD_OPEN(); return hsaKmtWaitOnEvent(NULL, Milliseconds); }
hsa_signal_value_t InterruptSignal::WaitRelaxed( hsa_signal_condition_t condition, hsa_signal_value_t compare_value, uint64_t timeout, hsa_wait_state_t wait_hint) { uint32_t prior = atomic::Increment(&waiting_); // assert(prior == 0 && "Multiple waiters on interrupt signal!"); // Allow only the first waiter to sleep (temporary, known to be bad). if (prior != 0) wait_hint = HSA_WAIT_STATE_ACTIVE; MAKE_SCOPE_GUARD([&]() { atomic::Decrement(&waiting_); }); int64_t value; timer::fast_clock::time_point start_time = timer::fast_clock::now(); // Set a polling timeout value // Exact time is not hugely important, it should just be a short while which // is smaller than the thread scheduling quantum (usually around 16ms) const timer::fast_clock::duration kMaxElapsed = std::chrono::milliseconds(5); uint64_t hsa_freq; HSA::hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &hsa_freq); const timer::fast_clock::duration fast_timeout = timer::duration_from_seconds<timer::fast_clock::duration>( double(timeout) / double(hsa_freq)); bool condition_met = false; while (true) { if (invalid_) return 0; value = atomic::Load(&signal_.value, std::memory_order_relaxed); switch (condition) { case HSA_SIGNAL_CONDITION_EQ: { condition_met = (value == compare_value); break; } case HSA_SIGNAL_CONDITION_NE: { condition_met = (value != compare_value); break; } case HSA_SIGNAL_CONDITION_GTE: { condition_met = (value >= compare_value); break; } case HSA_SIGNAL_CONDITION_LT: { condition_met = (value < compare_value); break; } default: return 0; } if (condition_met) return hsa_signal_value_t(value); timer::fast_clock::time_point time = timer::fast_clock::now(); if (time - start_time > kMaxElapsed) { if (time - start_time > fast_timeout) { value = atomic::Load(&signal_.value, std::memory_order_relaxed); return hsa_signal_value_t(value); } if (wait_on_event_ && wait_hint != HSA_WAIT_STATE_ACTIVE) { uint32_t wait_ms; auto time_remaining = fast_timeout - (time - start_time); if ((timeout == -1) || (time_remaining > std::chrono::milliseconds(uint32_t(-1)))) wait_ms = uint32_t(-1); else wait_ms = timer::duration_cast<std::chrono::milliseconds>( time_remaining).count(); hsaKmtWaitOnEvent(event_, wait_ms); } } } }