hsa_signal_value_t DefaultSignal::WaitRelaxed(hsa_signal_condition_t condition,
hsa_signal_value_t compare_value,
uint64_t timeout,
hsa_wait_state_t wait_hint) {
atomic::Increment(&waiting_);
MAKE_SCOPE_GUARD([&]() { atomic::Decrement(&waiting_); });
bool condition_met = false;
int64_t value;
assert(!g_use_interrupt_wait && "Use of non-host signal in host signal wait API.");
timer::fast_clock::time_point start_time, time;
start_time = timer::fast_clock::now();
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));
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);
time = timer::fast_clock::now();
if (time - start_time > fast_timeout) {
value = atomic::Load(&signal_.value, std::memory_order_relaxed);
return hsa_signal_value_t(value);
}
}
}
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);
}
}
}
}
