DynamicWatchdog.h File Reference

#include "Shared/funcannotations.h"
#include <cstdint>

Include dependency graph for DynamicWatchdog.h:

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Enumerations
enum	DynamicWatchdogFlags { DW_DEADLINE = 0, DW_ABORT = -1, DW_RESET = -2 }

Functions
RUNTIME_EXPORT uint64_t	dynamic_watchdog_init (unsigned ms_budget)
RUNTIME_EXPORT bool	dynamic_watchdog ()

Enumeration Type Documentation

enum DynamicWatchdogFlags

Enumerator
DW_DEADLINE
DW_ABORT
DW_RESET

Definition at line 24 of file DynamicWatchdog.h.

{ DW_DEADLINE = 0, DW_ABORT = -1, DW_RESET = -2 };

Function Documentation

RUNTIME_EXPORT bool dynamic_watchdog

(

)

Definition at line 115 of file cuda_mapd_rt.cu.

                                              {
  // check for dynamic watchdog, if triggered all threads return true
  if (dw_cycle_budget == 0LL) {
    return false;  // Uninitialized watchdog can't check time
  }
  if (dw_abort == 1) {
    return true;  // Received host request to abort
  }
  uint32_t smid = get_smid();
  if (smid >= 128) {
    return false;
  }
  __shared__ volatile int64_t dw_block_cycle_start;  // Thread block shared cycle start
  __shared__ volatile bool
      dw_should_terminate;  // all threads within a block should return together if
                            // watchdog criteria is met

  // thread 0 either initializes or read the initial clock cycle, the result is stored
  // into shared memory. Since all threads wihtin a block shares the same SM, there's no
  // point in using more threads here.
  if (threadIdx.x == 0) {
    dw_block_cycle_start = 0LL;
    int64_t cycle_count = static_cast<int64_t>(clock64());
    // Make sure the block hasn't switched SMs
    if (smid == get_smid()) {
      dw_block_cycle_start = static_cast<int64_t>(
          atomicCAS(reinterpret_cast<unsigned long long*>(&dw_sm_cycle_start[smid]),
                    0ULL,
                    static_cast<unsigned long long>(cycle_count)));
    }

    int64_t cycles = cycle_count - dw_block_cycle_start;
    if ((smid == get_smid()) && (dw_block_cycle_start > 0LL) &&
        (cycles > dw_cycle_budget)) {
      // Check if we're out of time on this particular SM
      dw_should_terminate = true;
    } else {
      dw_should_terminate = false;
    }
  }
  __syncthreads();
  return dw_should_terminate;
}

RUNTIME_EXPORT uint64_t dynamic_watchdog_init

(

unsigned

ms_budget

)

Definition at line 36 of file DynamicWatchdog.cpp.

References DW_ABORT, dw_abort, dw_cycle_budget, DW_DEADLINE, DW_RESET, read_cycle_counter(), and VLOG.

Referenced by dynamic_watchdog(), Executor::interrupt(), Executor::resetInterrupt(), and ExecutionKernel::runImpl().

                                                                             {
  static uint64_t dw_cycle_start = 0ULL;
  static uint64_t dw_cycle_budget = 0ULL;
  static std::atomic_bool dw_abort{false};

  if (ms_budget == static_cast<unsigned>(DW_DEADLINE)) {
    if (dw_abort.load()) {
      { return 0LL; }
    }
    return dw_cycle_start + dw_cycle_budget;
  }
  if (ms_budget == static_cast<unsigned>(DW_ABORT)) {
    dw_abort = true;
    return 0LL;
  }
  if (ms_budget == static_cast<unsigned>(DW_RESET)) {
    dw_abort = false;
    return 0LL;
  }

  // Init cycle start, measure freq, set and return cycle budget
  dw_cycle_start = read_cycle_counter();
  std::this_thread::sleep_for(std::chrono::milliseconds(1));
  auto freq_kHz = read_cycle_counter() - dw_cycle_start;
  dw_cycle_budget = freq_kHz * static_cast<uint64_t>(ms_budget);
  VLOG(1) << "INIT: thread " << std::this_thread::get_id() << ": ms_budget " << ms_budget
          << ", cycle_start " << dw_cycle_start << ", cycle_budget " << dw_cycle_budget
          << ", dw_deadline " << dw_cycle_start + dw_cycle_budget;
  return dw_cycle_budget;
}

Here is the call graph for this function:

Here is the caller graph for this function: