26 #if (defined(__x86_64__) || defined(__x86_64))
28 __asm__ __volatile__(
"rdtsc" :
"=a"(lo),
"=d"(hi));
29 return (static_cast<uint64_t>(hi) << 32) |
static_cast<uint64_t
>(lo);
37 static uint64_t dw_cycle_start = 0ULL;
39 static std::atomic_bool
dw_abort{
false};
41 if (ms_budget == static_cast<unsigned>(
DW_DEADLINE)) {
47 if (ms_budget == static_cast<unsigned>(
DW_ABORT)) {
51 if (ms_budget == static_cast<unsigned>(
DW_RESET)) {
58 std::this_thread::sleep_for(std::chrono::milliseconds(1));
60 dw_cycle_budget = freq_kHz *
static_cast<uint64_t
>(ms_budget);
61 VLOG(1) <<
"INIT: thread " << std::this_thread::get_id() <<
": ms_budget " << ms_budget
62 <<
", cycle_start " << dw_cycle_start <<
", cycle_budget " << dw_cycle_budget
71 if (clock > dw_deadline) {
72 LOG(
INFO) <<
"TIMEOUT: thread " << std::this_thread::get_id() <<
": clock " << clock
73 <<
", deadline " << dw_deadline;
__device__ bool dynamic_watchdog()
__device__ int64_t dw_cycle_budget
RUNTIME_EXPORT uint64_t dynamic_watchdog_init(unsigned ms_budget)
static FORCE_INLINE uint64_t read_cycle_counter(void)
__device__ int32_t dw_abort