OmniSciDB  6686921089
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
GpuInterrupt.cpp
Go to the documentation of this file.
1 /*
2  * Copyright 2017 MapD Technologies, Inc.
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16 
17 #include "DynamicWatchdog.h"
18 #include "Execute.h"
19 
20 void Executor::registerActiveModule(void* module, const int device_id) const {
21 #ifdef HAVE_CUDA
22  std::lock_guard<std::mutex> lock(gpu_active_modules_mutex_);
23  CHECK_LT(device_id, max_gpu_count);
24  gpu_active_modules_device_mask_ |= (1 << device_id);
25  gpu_active_modules_[device_id] = module;
26  VLOG(1) << "Executor " << executor_id_ << ", mask 0x" << std::hex
27  << gpu_active_modules_device_mask_ << ": Registered module " << module
28  << " on device " << std::to_string(device_id);
29 #endif
30 }
31 
32 void Executor::unregisterActiveModule(void* module, const int device_id) const {
33 #ifdef HAVE_CUDA
34  std::lock_guard<std::mutex> lock(gpu_active_modules_mutex_);
35  CHECK_LT(device_id, max_gpu_count);
36  if ((gpu_active_modules_device_mask_ & (1 << device_id)) == 0) {
37  return;
38  }
39  CHECK_EQ(gpu_active_modules_[device_id], module);
40  gpu_active_modules_device_mask_ ^= (1 << device_id);
41  VLOG(1) << "Executor " << executor_id_ << ", mask 0x" << std::hex
42  << gpu_active_modules_device_mask_ << ": Unregistered module " << module
43  << " on device " << std::to_string(device_id);
44 #endif
45 }
46 
47 void Executor::interrupt(const std::string& query_session,
48  const std::string& interrupt_session) {
49  const auto allow_interrupt =
51  if (allow_interrupt) {
52  bool is_running_query = false;
53  {
54  // here we validate the requested query session is valid (is already enrolled)
55  // if not, we skip the interrupt request
56  mapd_shared_lock<mapd_shared_mutex> session_read_lock(executor_session_mutex_);
57  if (!checkIsQuerySessionEnrolled(query_session, session_read_lock)) {
58  VLOG(1) << "Skip the interrupt request (no query has been submitted from the "
59  "given query session)";
60  return;
61  }
62  if (checkIsQuerySessionInterrupted(query_session, session_read_lock)) {
63  VLOG(1) << "Skip the interrupt request (already interrupted query session)";
64  return;
65  }
66  // if a query is pending query, we just need to turn interrupt flag for the session
67  // on (not sending interrupt signal to "RUNNING" kernel, see the below code)
68  is_running_query = checkCurrentQuerySession(query_session, session_read_lock);
69  }
70  {
71  // We have to cover interrupt request from *any* session because we don't know
72  // whether the request is for the running query or pending query
73  // or for non-kernel time interrupt
74  // (or just false alarm that indicates unregistered session in a queue).
75  // So we try to set a session has been interrupted once we confirm
76  // the session has been enrolled and is not interrupted at this moment
77  mapd_unique_lock<mapd_shared_mutex> session_write_lock(executor_session_mutex_);
78  setQuerySessionAsInterrupted(query_session, session_write_lock);
79  }
80  if (!is_running_query) {
81  return;
82  }
83  // mark the interrupted status of this executor
84  interrupted_.store(true);
85  }
86 
87  // for both GPU and CPU kernel execution, interrupt flag that running kernel accesses
88  // is a global variable from a view of Executors
89  // but it's okay for now since we hold a kernel_lock when starting the query execution
90  // this indicates we should revisit this logic when starting to use multi-query
91  // execution for supporting per-kernel interrupt
92  bool CPU_execution_mode = true;
93 
94 #ifdef HAVE_CUDA
95  // The below code is basically for runtime query interrupt for GPU.
96  // It is also possible that user forces to use CPU-mode even if the user has GPU(s).
97  // In this case, we should not execute the code in below to avoid runtime failure
99  auto cuda_mgr = data_mgr_->getCudaMgr();
100  if (cuda_mgr && (g_enable_dynamic_watchdog || allow_interrupt)) {
101  // we additionally allow sending interrupt signal for
102  // `g_enable_non_kernel_time_query_interrupt` especially for CTAS/ITAS queries: data
103  // population happens on CPU but select_query can be processed via GPU
104  CHECK_GE(cuda_mgr->getDeviceCount(), 1);
105  std::lock_guard<std::mutex> lock(gpu_active_modules_mutex_);
106  VLOG(1) << "Executor " << executor_id_ << ": Interrupting Active Modules: mask 0x"
107  << std::hex << gpu_active_modules_device_mask_;
108  CUcontext old_cu_context;
109  checkCudaErrors(cuCtxGetCurrent(&old_cu_context));
110  for (int device_id = 0; device_id < max_gpu_count; device_id++) {
111  if (gpu_active_modules_device_mask_ & (1 << device_id)) {
112  void* module = gpu_active_modules_[device_id];
113  auto cu_module = static_cast<CUmodule>(module);
114  if (!cu_module) {
115  continue;
116  } else {
117  VLOG(1) << "Try to interrupt the running query on GPU assigned to Executor "
118  << executor_id_;
119  CPU_execution_mode = false;
120  }
121  VLOG(1) << "Executor " << executor_id_ << ": Interrupting Active Modules: mask 0x"
122  << std::hex << gpu_active_modules_device_mask_ << " on device "
123  << std::to_string(device_id);
124 
125  cuda_mgr->setContext(device_id);
126 
127  // Create high priority non-blocking communication stream
128  CUstream cu_stream1;
130  cuStreamCreateWithPriority(&cu_stream1, CU_STREAM_NON_BLOCKING, 1));
131 
132  CUevent start, stop;
133  cuEventCreate(&start, 0);
134  cuEventCreate(&stop, 0);
135  cuEventRecord(start, cu_stream1);
136 
139  size_t dw_abort_size;
140  if (cuModuleGetGlobal(&dw_abort, &dw_abort_size, cu_module, "dw_abort") ==
141  CUDA_SUCCESS) {
142  CHECK_EQ(dw_abort_size, sizeof(uint32_t));
143  int32_t abort_val = 1;
144  checkCudaErrors(cuMemcpyHtoDAsync(dw_abort,
145  reinterpret_cast<void*>(&abort_val),
146  sizeof(int32_t),
147  cu_stream1));
148 
149  if (device_id == 0) {
150  VLOG(1) << "GPU: Async Abort submitted to Device "
151  << std::to_string(device_id);
152  }
153  }
154  }
155 
156  if (allow_interrupt) {
158  size_t runtime_interrupt_flag_size;
159  auto status = cuModuleGetGlobal(&runtime_interrupt_flag,
160  &runtime_interrupt_flag_size,
161  cu_module,
162  "runtime_interrupt_flag");
163  if (status == CUDA_SUCCESS) {
164  VLOG(1) << "Executor " << executor_id_
165  << " retrieves interrupt status from GPU " << device_id;
166  CHECK_EQ(runtime_interrupt_flag_size, sizeof(uint32_t));
167  int32_t abort_val = 1;
168  checkCudaErrors(cuMemcpyHtoDAsync(runtime_interrupt_flag,
169  reinterpret_cast<void*>(&abort_val),
170  sizeof(int32_t),
171  cu_stream1));
172  if (device_id == 0) {
173  VLOG(1) << "GPU: send interrupt signal from Executor " << executor_id_
174  << " to Device " << std::to_string(device_id);
175  }
176  } else if (status == CUDA_ERROR_NOT_FOUND) {
177  std::runtime_error(
178  "Runtime query interrupt on Executor " + std::to_string(executor_id_) +
179  " has failed: an interrupt flag on the GPU could "
180  "not be initialized (CUDA_ERROR_CODE: CUDA_ERROR_NOT_FOUND)");
181  } else {
182  // if we reach here, query runtime interrupt is failed due to
183  // one of the following error: CUDA_ERROR_NOT_INITIALIZED,
184  // CUDA_ERROR_DEINITIALIZED. CUDA_ERROR_INVALID_CONTEXT, and
185  // CUDA_ERROR_INVALID_VALUE. All those error codes are due to device failure.
186  const char* error_ret_str = nullptr;
187  cuGetErrorName(status, &error_ret_str);
188  if (!error_ret_str) {
189  error_ret_str = "UNKNOWN";
190  }
191  std::string error_str(error_ret_str);
192  std::runtime_error(
193  "Runtime interrupt on Executor " + std::to_string(executor_id_) +
194  " has failed due to a device " + std::to_string(device_id) +
195  "'s issue "
196  "(CUDA_ERROR_CODE: " +
197  error_str + ")");
198  }
199 
200  cuEventRecord(stop, cu_stream1);
201  cuEventSynchronize(stop);
202  float milliseconds = 0;
203  cuEventElapsedTime(&milliseconds, start, stop);
204  VLOG(1) << "Device " << std::to_string(device_id)
205  << ": submitted async interrupt request from Executor " << executor_id_
206  << " : SUCCESS: " << std::to_string(milliseconds) << " ms";
207  checkCudaErrors(cuStreamDestroy(cu_stream1));
208  }
209  }
210  checkCudaErrors(cuCtxSetCurrent(old_cu_context));
211  }
212  }
213 #endif
215  dynamic_watchdog_init(static_cast<unsigned>(DW_ABORT));
216  }
217 
218  if (allow_interrupt && CPU_execution_mode) {
219  // turn interrupt flag on for CPU mode
220  VLOG(1) << "Try to interrupt the running query on CPU from Executor " << executor_id_;
221  check_interrupt_init(static_cast<unsigned>(INT_ABORT));
222  }
223 }
224 
226 #ifdef HAVE_CUDA
227  std::lock_guard<std::mutex> lock(gpu_active_modules_mutex_);
228 #endif
229  const auto allow_interrupt =
232  dynamic_watchdog_init(static_cast<unsigned>(DW_RESET));
233  } else if (allow_interrupt) {
234  VLOG(1) << "Reset interrupt flag for CPU execution kernel on Executor "
235  << executor_id_;
236  check_interrupt_init(static_cast<unsigned>(INT_RESET));
237  }
238 
239  if (interrupted_.load()) {
240  VLOG(1) << "RESET Executor " << executor_id_
241  << " that had previously been interrupted";
242  interrupted_.store(false);
243  }
244 }
#define CHECK_EQ(x, y)
Definition: Logger.h:217
int CUcontext
Definition: nocuda.h:22
void checkCudaErrors(CUresult err)
Definition: sample.cpp:38
void registerActiveModule(void *module, const int device_id) const
unsigned long long CUdeviceptr
Definition: nocuda.h:27
#define CHECK_GE(x, y)
Definition: Logger.h:222
bool g_enable_dynamic_watchdog
Definition: Execute.cpp:77
bool g_enable_non_kernel_time_query_interrupt
Definition: Execute.cpp:119
std::string to_string(char const *&&v)
__device__ int32_t runtime_interrupt_flag
Definition: cuda_mapd_rt.cu:96
RUNTIME_EXPORT uint64_t dynamic_watchdog_init(unsigned ms_budget)
executor_id_(executor_id)
#define CHECK_LT(x, y)
Definition: Logger.h:219
data_mgr_(data_mgr)
void unregisterActiveModule(void *module, const int device_id) const
void resetInterrupt()
bool check_interrupt_init(unsigned command)
#define CHECK(condition)
Definition: Logger.h:209
void interrupt(const QuerySessionId &query_session="", const QuerySessionId &interrupt_session="")
__device__ int32_t dw_abort
Definition: cuda_mapd_rt.cu:95
bool g_enable_runtime_query_interrupt
Definition: Execute.cpp:118
#define VLOG(n)
Definition: Logger.h:303
void * CUmodule
Definition: nocuda.h:23