OmniSciDB  d2f719934e
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
TableFunctionExecutionContext Class Reference

#include <TableFunctionExecutionContext.h>

Public Member Functions

 TableFunctionExecutionContext (std::shared_ptr< RowSetMemoryOwner > row_set_mem_owner)
 
 TableFunctionExecutionContext (const TableFunctionExecutionContext &)=delete
 
TableFunctionExecutionContextoperator= (const TableFunctionExecutionContext &)=delete
 
ResultSetPtr execute (const TableFunctionExecutionUnit &exe_unit, const std::vector< InputTableInfo > &table_infos, const std::shared_ptr< CompilationContext > &compilation_context, const ColumnFetcher &column_fetcher, const ExecutorDeviceType device_type, Executor *executor, bool is_pre_launch_udtf)
 

Private Member Functions

void launchPreCodeOnCpu (const TableFunctionExecutionUnit &exe_unit, const std::shared_ptr< CpuCompilationContext > &compilation_context, std::vector< const int8_t * > &col_buf_ptrs, std::vector< int64_t > &col_sizes, const size_t elem_count, Executor *executor)
 
ResultSetPtr launchCpuCode (const TableFunctionExecutionUnit &exe_unit, const std::shared_ptr< CpuCompilationContext > &compilation_context, std::vector< const int8_t * > &col_buf_ptrs, std::vector< int64_t > &col_sizes, const size_t elem_count, Executor *executor)
 
ResultSetPtr launchGpuCode (const TableFunctionExecutionUnit &exe_unit, const std::shared_ptr< GpuCompilationContext > &compilation_context, std::vector< const int8_t * > &col_buf_ptrs, std::vector< int64_t > &col_sizes, const size_t elem_count, const int device_id, Executor *executor)
 

Private Attributes

std::shared_ptr
< RowSetMemoryOwner
row_set_mem_owner_
 

Detailed Description

Definition at line 28 of file TableFunctionExecutionContext.h.

Constructor & Destructor Documentation

TableFunctionExecutionContext::TableFunctionExecutionContext ( std::shared_ptr< RowSetMemoryOwner row_set_mem_owner)
inline

Definition at line 30 of file TableFunctionExecutionContext.h.

31  : row_set_mem_owner_(row_set_mem_owner) {}
std::shared_ptr< RowSetMemoryOwner > row_set_mem_owner_
TableFunctionExecutionContext::TableFunctionExecutionContext ( const TableFunctionExecutionContext )
delete

Member Function Documentation

ResultSetPtr TableFunctionExecutionContext::execute ( const TableFunctionExecutionUnit exe_unit,
const std::vector< InputTableInfo > &  table_infos,
const std::shared_ptr< CompilationContext > &  compilation_context,
const ColumnFetcher column_fetcher,
const ExecutorDeviceType  device_type,
Executor executor,
bool  is_pre_launch_udtf 
)

Definition at line 117 of file TableFunctionExecutionContext.cpp.

References CHECK, CHECK_EQ, ColumnFetcher::columnarized_table_cache_, table_functions::TableFunction::containsRequireFnCheck(), CPU, Data_Namespace::CPU_LEVEL, anonymous_namespace{TableFunctionExecutionContext.cpp}::create_literal_buffer(), get_bit_width(), ColumnFetcher::getOneColumnFragment(), GPU, Data_Namespace::GPU_LEVEL, table_functions::TableFunction::hasOutputSizeIndependentOfInputSize(), TableFunctionExecutionUnit::input_exprs, launchCpuCode(), launchGpuCode(), launchPreCodeOnCpu(), TableFunctionExecutionUnit::table_func, and UNREACHABLE.

Referenced by Executor::executeTableFunction().

124  {
125  CHECK(compilation_context);
126  std::vector<std::shared_ptr<Chunk_NS::Chunk>> chunks_owner;
127  std::vector<std::unique_ptr<char[]>> literals_owner;
128 
129  const int device_id = 0; // TODO(adb): support multi-gpu table functions
130  std::unique_ptr<CudaAllocator> device_allocator;
131  if (device_type == ExecutorDeviceType::GPU) {
132  auto data_mgr = executor->getDataMgr();
133  device_allocator.reset(new CudaAllocator(data_mgr, device_id));
134  }
135  std::vector<const int8_t*> col_buf_ptrs;
136  std::vector<int64_t> col_sizes;
137  std::optional<size_t> input_num_rows;
138 
139  int col_index = -1;
140  // TODO: col_list_bufs are allocated on CPU memory, so UDTFs with column_list
141  // arguments are not supported on GPU atm.
142  std::vector<std::vector<const int8_t*>> col_list_bufs;
143  for (const auto& input_expr : exe_unit.input_exprs) {
144  auto ti = input_expr->get_type_info();
145  if (!ti.is_column_list()) {
146  CHECK_EQ(col_index, -1);
147  }
148  if (auto col_var = dynamic_cast<Analyzer::ColumnVar*>(input_expr)) {
149  auto table_id = col_var->get_table_id();
150  auto table_info_it = std::find_if(
151  table_infos.begin(), table_infos.end(), [&table_id](const auto& table_info) {
152  return table_info.table_id == table_id;
153  });
154  CHECK(table_info_it != table_infos.end());
155  auto [col_buf, buf_elem_count] = ColumnFetcher::getOneColumnFragment(
156  executor,
157  *col_var,
158  table_info_it->info.fragments.front(),
161  device_id,
162  device_allocator.get(),
163  /*thread_idx=*/0,
164  chunks_owner,
165  column_fetcher.columnarized_table_cache_);
166  // We use the number of entries in the first column to be the number of rows to base
167  // the output off of (optionally depending on the sizing parameter)
168  if (!input_num_rows) {
169  input_num_rows = (buf_elem_count ? buf_elem_count : 1);
170  }
171  if (ti.is_column_list()) {
172  if (col_index == -1) {
173  col_list_bufs.push_back({});
174  col_list_bufs.back().reserve(ti.get_dimension());
175  } else {
176  CHECK_EQ(col_sizes.back(), buf_elem_count);
177  }
178  col_index++;
179  col_list_bufs.back().push_back(col_buf);
180  // append col_buf to column_list col_buf
181  if (col_index + 1 == ti.get_dimension()) {
182  col_index = -1;
183  }
184  // columns in the same column_list point to column_list data
185  col_buf_ptrs.push_back((const int8_t*)col_list_bufs.back().data());
186  } else {
187  col_buf_ptrs.push_back(col_buf);
188  }
189  col_sizes.push_back(buf_elem_count);
190  } else if (const auto& constant_val = dynamic_cast<Analyzer::Constant*>(input_expr)) {
191  // TODO(adb): Unify literal handling with rest of system, either in Codegen or as a
192  // separate serialization component
193  col_sizes.push_back(0);
194  const auto const_val_datum = constant_val->get_constval();
195  const auto& ti = constant_val->get_type_info();
196  if (ti.is_fp()) {
197  switch (get_bit_width(ti)) {
198  case 32:
199  col_buf_ptrs.push_back(create_literal_buffer(const_val_datum.floatval,
200  device_type,
201  literals_owner,
202  device_allocator.get()));
203  break;
204  case 64:
205  col_buf_ptrs.push_back(create_literal_buffer(const_val_datum.doubleval,
206  device_type,
207  literals_owner,
208  device_allocator.get()));
209  break;
210  default:
211  UNREACHABLE();
212  }
213  } else if (ti.is_integer()) {
214  switch (get_bit_width(ti)) {
215  case 8:
216  col_buf_ptrs.push_back(create_literal_buffer(const_val_datum.tinyintval,
217  device_type,
218  literals_owner,
219  device_allocator.get()));
220  break;
221  case 16:
222  col_buf_ptrs.push_back(create_literal_buffer(const_val_datum.smallintval,
223  device_type,
224  literals_owner,
225  device_allocator.get()));
226  break;
227  case 32:
228  col_buf_ptrs.push_back(create_literal_buffer(const_val_datum.intval,
229  device_type,
230  literals_owner,
231  device_allocator.get()));
232  break;
233  case 64:
234  col_buf_ptrs.push_back(create_literal_buffer(const_val_datum.bigintval,
235  device_type,
236  literals_owner,
237  device_allocator.get()));
238  break;
239  default:
240  UNREACHABLE();
241  }
242  } else if (ti.is_boolean()) {
243  col_buf_ptrs.push_back(create_literal_buffer(const_val_datum.boolval,
244  device_type,
245  literals_owner,
246  device_allocator.get()));
247  } else if (ti.is_bytes()) { // text encoding none string
248  col_buf_ptrs.push_back(create_literal_buffer(const_val_datum.stringval,
249  device_type,
250  literals_owner,
251  device_allocator.get()));
252  } else {
253  throw TableFunctionError("Literal value " + constant_val->toString() +
254  " is not yet supported.");
255  }
256  }
257  }
258  CHECK_EQ(col_buf_ptrs.size(), exe_unit.input_exprs.size());
259  CHECK_EQ(col_sizes.size(), exe_unit.input_exprs.size());
260  if (!exe_unit.table_func
261  .hasOutputSizeIndependentOfInputSize()) { // includes compile-time constants,
262  // user-specified constants,
263  // and runtime table funtion
264  // specified sizing, only
265  // user-specified row-multipliers
266  // currently take into account input
267  // row size
268  CHECK(input_num_rows);
269  }
270  if (is_pre_launch_udtf) {
273  exe_unit,
274  std::dynamic_pointer_cast<CpuCompilationContext>(compilation_context),
275  col_buf_ptrs,
276  col_sizes,
277  *input_num_rows,
278  executor);
279  return nullptr;
280  } else {
281  switch (device_type) {
283  return launchCpuCode(
284  exe_unit,
285  std::dynamic_pointer_cast<CpuCompilationContext>(compilation_context),
286  col_buf_ptrs,
287  col_sizes,
288  *input_num_rows,
289  executor);
291  return launchGpuCode(
292  exe_unit,
293  std::dynamic_pointer_cast<GpuCompilationContext>(compilation_context),
294  col_buf_ptrs,
295  col_sizes,
296  *input_num_rows,
297  /*device_id=*/0,
298  executor);
299  }
300  }
301  UNREACHABLE();
302  return nullptr;
303 }
ResultSetPtr launchGpuCode(const TableFunctionExecutionUnit &exe_unit, const std::shared_ptr< GpuCompilationContext > &compilation_context, std::vector< const int8_t * > &col_buf_ptrs, std::vector< int64_t > &col_sizes, const size_t elem_count, const int device_id, Executor *executor)
#define CHECK_EQ(x, y)
Definition: Logger.h:219
void launchPreCodeOnCpu(const TableFunctionExecutionUnit &exe_unit, const std::shared_ptr< CpuCompilationContext > &compilation_context, std::vector< const int8_t * > &col_buf_ptrs, std::vector< int64_t > &col_sizes, const size_t elem_count, Executor *executor)
std::vector< Analyzer::Expr * > input_exprs
const table_functions::TableFunction table_func
#define UNREACHABLE()
Definition: Logger.h:255
ColumnCacheMap columnarized_table_cache_
size_t get_bit_width(const SQLTypeInfo &ti)
ResultSetPtr launchCpuCode(const TableFunctionExecutionUnit &exe_unit, const std::shared_ptr< CpuCompilationContext > &compilation_context, std::vector< const int8_t * > &col_buf_ptrs, std::vector< int64_t > &col_sizes, const size_t elem_count, Executor *executor)
static std::pair< const int8_t *, size_t > getOneColumnFragment(Executor *executor, const Analyzer::ColumnVar &hash_col, const Fragmenter_Namespace::FragmentInfo &fragment, const Data_Namespace::MemoryLevel effective_mem_lvl, const int device_id, DeviceAllocator *device_allocator, const size_t thread_idx, std::vector< std::shared_ptr< Chunk_NS::Chunk >> &chunks_owner, ColumnCacheMap &column_cache)
Gets one chunk&#39;s pointer and element count on either CPU or GPU.
const int8_t * create_literal_buffer(const T literal, const ExecutorDeviceType device_type, std::vector< std::unique_ptr< char[]>> &literals_owner, CudaAllocator *gpu_allocator)
#define CHECK(condition)
Definition: Logger.h:211

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

ResultSetPtr TableFunctionExecutionContext::launchCpuCode ( const TableFunctionExecutionUnit exe_unit,
const std::shared_ptr< CpuCompilationContext > &  compilation_context,
std::vector< const int8_t * > &  col_buf_ptrs,
std::vector< int64_t > &  col_sizes,
const size_t  elem_count,
Executor executor 
)
private

Definition at line 365 of file TableFunctionExecutionContext.cpp.

References CHECK, CHECK_EQ, DEBUG_TIMER, GenericError, anonymous_namespace{TableFunctionExecutionContext.cpp}::get_output_row_count(), i, row_set_mem_owner_, t, and to_string().

Referenced by execute().

371  {
372  int64_t output_row_count = 0;
373 
374  // If TableFunctionManager must be a singleton but it has been
375  // initialized from another thread, TableFunctionManager constructor
376  // blocks via TableFunctionManager_singleton_mutex until the
377  // existing singleton is deconstructed.
378  auto mgr = std::make_unique<TableFunctionManager>(
379  exe_unit,
380  executor,
381  col_buf_ptrs,
383  /*is_singleton=*/!exe_unit.table_func.usesManager());
384 
385  if (exe_unit.table_func.hasOutputSizeKnownPreLaunch()) {
386  // allocate output buffers because the size is known up front, from
387  // user specified parameters (and table size in the case of a user
388  // specified row multiplier)
389  output_row_count = get_output_row_count(exe_unit, elem_count);
390  }
391 
392  // setup the inputs
393  // We can have an empty col_buf_ptrs vector if there are no arguments to the function
394  const auto byte_stream_ptr = !col_buf_ptrs.empty()
395  ? reinterpret_cast<const int8_t**>(col_buf_ptrs.data())
396  : nullptr;
397  if (!col_buf_ptrs.empty()) {
398  CHECK(byte_stream_ptr);
399  }
400  const auto col_sizes_ptr = !col_sizes.empty() ? col_sizes.data() : nullptr;
401  if (!col_sizes.empty()) {
402  CHECK(col_sizes_ptr);
403  }
404 
405  // execute
406  auto timer = DEBUG_TIMER(__func__);
407  const auto err = compilation_context->table_function_entry_point()(
408  reinterpret_cast<const int8_t*>(mgr.get()),
409  byte_stream_ptr, // input columns buffer
410  col_sizes_ptr, // input column sizes
411  nullptr,
412  &output_row_count);
413 
415  throw UserTableFunctionError("Error executing table function: " +
416  std::string(mgr->get_error_message()));
417  }
418 
419  else if (err) {
420  throw UserTableFunctionError("Error executing table function: " +
421  std::to_string(err));
422  }
423 
424  if (exe_unit.table_func.hasCompileTimeOutputSizeConstant()) {
425  if (static_cast<size_t>(output_row_count) != mgr->get_nrows()) {
426  throw TableFunctionError(
427  "Table function with constant sizing parameter must return " +
428  std::to_string(mgr->get_nrows()) + " (got " + std::to_string(output_row_count) +
429  ")");
430  }
431  } else {
432  if (output_row_count < 0 || (size_t)output_row_count > mgr->get_nrows()) {
433  output_row_count = mgr->get_nrows();
434  }
435  }
436  // Update entry count, it may differ from allocated mem size
437  if (exe_unit.table_func.hasTableFunctionSpecifiedParameter() && !mgr->query_buffers) {
438  // set_output_row_size has not been called
439  if (output_row_count == 0) {
440  // allocate for empty output columns
441  mgr->allocate_output_buffers(0);
442  } else {
443  throw TableFunctionError("Table function must call set_output_row_size");
444  }
445  }
446 
447  mgr->query_buffers->getResultSet(0)->updateStorageEntryCount(output_row_count);
448 
449  const size_t column_size = output_row_count * sizeof(int64_t);
450  const size_t allocated_column_size = mgr->get_nrows() * sizeof(int64_t);
451  auto group_by_buffers_ptr = mgr->query_buffers->getGroupByBuffersPtr();
452  CHECK(group_by_buffers_ptr);
453  auto output_buffers_ptr = reinterpret_cast<int64_t*>(group_by_buffers_ptr[0]);
454 
455  auto num_out_columns = exe_unit.target_exprs.size();
456  int8_t* src = reinterpret_cast<int8_t*>(output_buffers_ptr);
457  int8_t* dst = reinterpret_cast<int8_t*>(output_buffers_ptr);
458  for (size_t i = 0; i < num_out_columns; i++) {
459  if (src != dst) {
460  auto t = memmove(dst, src, column_size);
461  CHECK_EQ(dst, t);
462  }
463  src += allocated_column_size;
464  dst += column_size;
465  }
466  return mgr->query_buffers->getResultSetOwned(0);
467 }
#define CHECK_EQ(x, y)
Definition: Logger.h:219
size_t get_output_row_count(const TableFunctionExecutionUnit &exe_unit, size_t input_element_count)
std::string to_string(char const *&&v)
std::shared_ptr< RowSetMemoryOwner > row_set_mem_owner_
#define CHECK(condition)
Definition: Logger.h:211
#define DEBUG_TIMER(name)
Definition: Logger.h:358
char * t

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

ResultSetPtr TableFunctionExecutionContext::launchGpuCode ( const TableFunctionExecutionUnit exe_unit,
const std::shared_ptr< GpuCompilationContext > &  compilation_context,
std::vector< const int8_t * > &  col_buf_ptrs,
std::vector< int64_t > &  col_sizes,
const size_t  elem_count,
const int  device_id,
Executor executor 
)
private

Definition at line 481 of file TableFunctionExecutionContext.cpp.

References QueryMemoryDescriptor::addColSlotInfo(), CHECK, CHECK_EQ, checkCudaErrors(), anonymous_namespace{TableFunctionExecutionContext.cpp}::COL_BUFFERS, anonymous_namespace{TableFunctionExecutionContext.cpp}::COL_SIZES, anonymous_namespace{TableFunctionExecutionContext.cpp}::ERROR_BUFFER, anonymous_namespace{TableFunctionExecutionContext.cpp}::get_output_row_count(), GPU, table_functions::TableFunction::hasTableFunctionSpecifiedParameter(), i, anonymous_namespace{TableFunctionExecutionContext.cpp}::KERNEL_PARAM_COUNT, anonymous_namespace{TableFunctionExecutionContext.cpp}::MANAGER, anonymous_namespace{TableFunctionExecutionContext.cpp}::OUTPUT_BUFFERS, anonymous_namespace{TableFunctionExecutionContext.cpp}::OUTPUT_ROW_COUNT, Projection, query_mem_desc, row_set_mem_owner_, QueryMemoryDescriptor::setOutputColumnar(), TableFunctionExecutionUnit::table_func, TableFunctionExecutionUnit::target_exprs, to_string(), and UNREACHABLE.

Referenced by execute().

488  {
489 #ifdef HAVE_CUDA
491  throw QueryMustRunOnCpu();
492  }
493 
494  auto num_out_columns = exe_unit.target_exprs.size();
495  auto data_mgr = executor->getDataMgr();
496  auto gpu_allocator = std::make_unique<CudaAllocator>(data_mgr, device_id);
497  CHECK(gpu_allocator);
498  std::vector<CUdeviceptr> kernel_params(KERNEL_PARAM_COUNT, 0);
499 
500  // TODO: implement table function manager for CUDA
501  // kernels. kernel_params[MANAGER] ought to contain a device pointer
502  // to a struct that a table function kernel with a
503  // TableFunctionManager argument can access from the device.
504  kernel_params[MANAGER] =
505  reinterpret_cast<CUdeviceptr>(gpu_allocator->alloc(sizeof(int8_t*)));
506 
507  // setup the inputs
508  auto byte_stream_ptr = !(col_buf_ptrs.empty())
509  ? gpu_allocator->alloc(col_buf_ptrs.size() * sizeof(int64_t))
510  : nullptr;
511  if (byte_stream_ptr) {
512  gpu_allocator->copyToDevice(byte_stream_ptr,
513  reinterpret_cast<int8_t*>(col_buf_ptrs.data()),
514  col_buf_ptrs.size() * sizeof(int64_t));
515  }
516  kernel_params[COL_BUFFERS] = reinterpret_cast<CUdeviceptr>(byte_stream_ptr);
517 
518  auto col_sizes_ptr = !(col_sizes.empty())
519  ? gpu_allocator->alloc(col_sizes.size() * sizeof(int64_t))
520  : nullptr;
521  if (col_sizes_ptr) {
522  gpu_allocator->copyToDevice(col_sizes_ptr,
523  reinterpret_cast<int8_t*>(col_sizes.data()),
524  col_sizes.size() * sizeof(int64_t));
525  }
526  kernel_params[COL_SIZES] = reinterpret_cast<CUdeviceptr>(col_sizes_ptr);
527 
528  kernel_params[ERROR_BUFFER] =
529  reinterpret_cast<CUdeviceptr>(gpu_allocator->alloc(sizeof(int32_t)));
530  // initialize output memory
532  executor, elem_count, QueryDescriptionType::Projection, /*is_table_function=*/true);
533  query_mem_desc.setOutputColumnar(true);
534 
535  for (size_t i = 0; i < num_out_columns; i++) {
536  // All outputs padded to 8 bytes
537  query_mem_desc.addColSlotInfo({std::make_tuple(8, 8)});
538  }
539  const auto allocated_output_row_count = get_output_row_count(exe_unit, elem_count);
540  auto query_buffers = std::make_unique<QueryMemoryInitializer>(
541  exe_unit,
543  device_id,
545  (allocated_output_row_count == 0 ? 1 : allocated_output_row_count),
546  std::vector<std::vector<const int8_t*>>{col_buf_ptrs},
547  std::vector<std::vector<uint64_t>>{{0}}, // frag offsets
549  gpu_allocator.get(),
550  executor);
551 
552  // setup the output
553  int64_t output_row_count = allocated_output_row_count;
554 
555  kernel_params[OUTPUT_ROW_COUNT] =
556  reinterpret_cast<CUdeviceptr>(gpu_allocator->alloc(sizeof(int64_t*)));
557  gpu_allocator->copyToDevice(reinterpret_cast<int8_t*>(kernel_params[OUTPUT_ROW_COUNT]),
558  reinterpret_cast<int8_t*>(&output_row_count),
559  sizeof(output_row_count));
560 
561  const unsigned block_size_x = executor->blockSize();
562  const unsigned block_size_y = 1;
563  const unsigned block_size_z = 1;
564  const unsigned grid_size_x = executor->gridSize();
565  const unsigned grid_size_y = 1;
566  const unsigned grid_size_z = 1;
567 
568  auto gpu_output_buffers = query_buffers->setupTableFunctionGpuBuffers(
569  query_mem_desc, device_id, block_size_x, grid_size_x);
570 
571  kernel_params[OUTPUT_BUFFERS] = reinterpret_cast<CUdeviceptr>(gpu_output_buffers.ptrs);
572 
573  // execute
574  CHECK_EQ(static_cast<size_t>(KERNEL_PARAM_COUNT), kernel_params.size());
575 
576  std::vector<void*> param_ptrs;
577  for (auto& param : kernel_params) {
578  param_ptrs.push_back(&param);
579  }
580 
581  // Get cu func
582 
583  CHECK(compilation_context);
584  const auto native_code = compilation_context->getNativeCode(device_id);
585  auto cu_func = static_cast<CUfunction>(native_code.first);
586  checkCudaErrors(cuLaunchKernel(cu_func,
587  grid_size_x,
588  grid_size_y,
589  grid_size_z,
590  block_size_x,
591  block_size_y,
592  block_size_z,
593  0, // shared mem bytes
594  nullptr,
595  &param_ptrs[0],
596  nullptr));
597  // TODO(adb): read errors
598 
599  // read output row count from GPU
600  gpu_allocator->copyFromDevice(
601  reinterpret_cast<int8_t*>(&output_row_count),
602  reinterpret_cast<int8_t*>(kernel_params[OUTPUT_ROW_COUNT]),
603  sizeof(int64_t));
604  if (exe_unit.table_func.hasNonUserSpecifiedOutputSize()) {
605  if (static_cast<size_t>(output_row_count) != allocated_output_row_count) {
606  throw TableFunctionError(
607  "Table function with constant sizing parameter must return " +
608  std::to_string(allocated_output_row_count) + " (got " +
609  std::to_string(output_row_count) + ")");
610  }
611  } else {
612  if (output_row_count < 0 || (size_t)output_row_count > allocated_output_row_count) {
613  output_row_count = allocated_output_row_count;
614  }
615  }
616 
617  // Update entry count, it may differ from allocated mem size
618  query_buffers->getResultSet(0)->updateStorageEntryCount(output_row_count);
619 
620  // Copy back to CPU storage
621  query_buffers->copyFromTableFunctionGpuBuffers(data_mgr,
622  query_mem_desc,
623  output_row_count,
624  gpu_output_buffers,
625  device_id,
626  block_size_x,
627  grid_size_x);
628 
629  return query_buffers->getResultSetOwned(0);
630 #else
631  UNREACHABLE();
632  return nullptr;
633 #endif
634 }
#define CHECK_EQ(x, y)
Definition: Logger.h:219
size_t get_output_row_count(const TableFunctionExecutionUnit &exe_unit, size_t input_element_count)
const table_functions::TableFunction table_func
void checkCudaErrors(CUresult err)
Definition: sample.cpp:38
unsigned long long CUdeviceptr
Definition: nocuda.h:27
#define UNREACHABLE()
Definition: Logger.h:255
std::string to_string(char const *&&v)
void * CUfunction
Definition: nocuda.h:24
std::shared_ptr< RowSetMemoryOwner > row_set_mem_owner_
#define CHECK(condition)
Definition: Logger.h:211
std::vector< Analyzer::Expr * > target_exprs

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

void TableFunctionExecutionContext::launchPreCodeOnCpu ( const TableFunctionExecutionUnit exe_unit,
const std::shared_ptr< CpuCompilationContext > &  compilation_context,
std::vector< const int8_t * > &  col_buf_ptrs,
std::vector< int64_t > &  col_sizes,
const size_t  elem_count,
Executor executor 
)
private

Definition at line 307 of file TableFunctionExecutionContext.cpp.

References CHECK, DEBUG_TIMER, GenericError, anonymous_namespace{TableFunctionExecutionContext.cpp}::get_output_row_count(), row_set_mem_owner_, and to_string().

Referenced by execute().

313  {
314  int64_t output_row_count = 0;
315 
316  // If TableFunctionManager must be a singleton but it has been
317  // initialized from another thread, TableFunctionManager constructor
318  // blocks via TableFunctionManager_singleton_mutex until the
319  // existing singleton is deconstructed.
320  auto mgr = std::make_unique<TableFunctionManager>(
321  exe_unit,
322  executor,
323  col_buf_ptrs,
325  /*is_singleton=*/!exe_unit.table_func.usesManager());
326 
327  if (exe_unit.table_func.hasOutputSizeKnownPreLaunch()) {
328  // allocate output buffers because the size is known up front, from
329  // user specified parameters (and table size in the case of a user
330  // specified row multiplier)
331  output_row_count = get_output_row_count(exe_unit, elem_count);
332  }
333 
334  // setup the inputs
335  // We can have an empty col_buf_ptrs vector if there are no arguments to the function
336  const auto byte_stream_ptr = !col_buf_ptrs.empty()
337  ? reinterpret_cast<const int8_t**>(col_buf_ptrs.data())
338  : nullptr;
339  if (!col_buf_ptrs.empty()) {
340  CHECK(byte_stream_ptr);
341  }
342  const auto col_sizes_ptr = !col_sizes.empty() ? col_sizes.data() : nullptr;
343  if (!col_sizes.empty()) {
344  CHECK(col_sizes_ptr);
345  }
346 
347  // execute
348  auto timer = DEBUG_TIMER(__func__);
349  const auto err = compilation_context->table_function_entry_point()(
350  reinterpret_cast<const int8_t*>(mgr.get()),
351  byte_stream_ptr, // input columns buffer
352  col_sizes_ptr, // input column sizes
353  nullptr,
354  &output_row_count);
355 
357  throw UserTableFunctionError("Error executing table function require check: " +
358  std::string(mgr->get_error_message()));
359  } else if (err) {
360  throw UserTableFunctionError("Error executing table function require check: " +
361  std::to_string(err));
362  }
363 }
size_t get_output_row_count(const TableFunctionExecutionUnit &exe_unit, size_t input_element_count)
std::string to_string(char const *&&v)
std::shared_ptr< RowSetMemoryOwner > row_set_mem_owner_
#define CHECK(condition)
Definition: Logger.h:211
#define DEBUG_TIMER(name)
Definition: Logger.h:358

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

TableFunctionExecutionContext& TableFunctionExecutionContext::operator= ( const TableFunctionExecutionContext )
delete

Member Data Documentation

std::shared_ptr<RowSetMemoryOwner> TableFunctionExecutionContext::row_set_mem_owner_
private

The documentation for this class was generated from the following files: