OmniSciDB  0fdbebe030
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
HashJoinRuntime.h File Reference
#include <cstddef>
#include <cstdint>
#include <vector>
#include "../Shared/SqlTypesLayout.h"
#include "../Shared/sqltypes.h"
#include "RuntimeFunctions.h"
#include "../Shared/funcannotations.h"
+ Include dependency graph for HashJoinRuntime.h:
+ This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Classes

struct  HashEntryInfo
 
struct  JoinChunk
 
struct  JoinColumn
 
struct  JoinColumnTypeInfo
 
struct  JoinBucketInfo
 
struct  ShardInfo
 

Enumerations

enum  ColumnType { SmallDate = 0, Signed = 1, Unsigned = 2, Double = 3 }
 

Functions

void init_hash_join_buff (int32_t *buff, const int32_t entry_count, const int32_t invalid_slot_val, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
 
void init_hash_join_buff_on_device (int32_t *buff, const int32_t entry_count, const int32_t invalid_slot_val, const size_t block_size_x, const size_t grid_size_x)
 
void init_baseline_hash_join_buff_32 (int8_t *hash_join_buff, const int32_t entry_count, const size_t key_component_count, const bool with_val_slot, const int32_t invalid_slot_val, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
 
void init_baseline_hash_join_buff_64 (int8_t *hash_join_buff, const int32_t entry_count, const size_t key_component_count, const bool with_val_slot, const int32_t invalid_slot_val, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
 
void init_baseline_hash_join_buff_on_device_32 (int8_t *hash_join_buff, const int32_t entry_count, const size_t key_component_count, const bool with_val_slot, const int32_t invalid_slot_val, const size_t block_size_x, const size_t grid_size_x)
 
void init_baseline_hash_join_buff_on_device_64 (int8_t *hash_join_buff, const int32_t entry_count, const size_t key_component_count, const bool with_val_slot, const int32_t invalid_slot_val, const size_t block_size_x, const size_t grid_size_x)
 
ColumnType get_join_column_type_kind (const SQLTypeInfo &ti)
 
int fill_hash_join_buff_bucketized (int32_t *buff, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const void *sd_inner, const void *sd_outer, const int32_t cpu_thread_idx, const int32_t cpu_thread_count, const int64_t bucket_normalization)
 
int fill_hash_join_buff (int32_t *buff, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const void *sd_inner, const void *sd_outer, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
 
void fill_hash_join_buff_on_device (int32_t *buff, const int32_t invalid_slot_val, int *dev_err_buff, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const size_t block_size_x, const size_t grid_size_x)
 
void fill_hash_join_buff_on_device_bucketized (int32_t *buff, const int32_t invalid_slot_val, int *dev_err_buff, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const size_t block_size_x, const size_t grid_size_x, const int64_t bucket_normalization)
 
void fill_hash_join_buff_on_device_sharded (int32_t *buff, const int32_t invalid_slot_val, int *dev_err_buff, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const ShardInfo shard_info, const size_t block_size_x, const size_t grid_size_x)
 
void fill_hash_join_buff_on_device_sharded_bucketized (int32_t *buff, const int32_t invalid_slot_val, int *dev_err_buff, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const ShardInfo shard_info, const size_t block_size_x, const size_t grid_size_x, const int64_t bucket_normalization)
 
void fill_one_to_many_hash_table (int32_t *buff, const HashEntryInfo hash_entry_info, const int32_t invalid_slot_val, const JoinColumn &join_column, const JoinColumnTypeInfo &type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const unsigned cpu_thread_count)
 
void fill_one_to_many_hash_table_bucketized (int32_t *buff, const HashEntryInfo hash_entry_info, const int32_t invalid_slot_val, const JoinColumn &join_column, const JoinColumnTypeInfo &type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const unsigned cpu_thread_count)
 
void fill_one_to_many_hash_table_sharded_bucketized (int32_t *buff, const HashEntryInfo hash_entry_info, const int32_t invalid_slot_val, const JoinColumn &join_column, const JoinColumnTypeInfo &type_info, const ShardInfo &shard_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const unsigned cpu_thread_count)
 
void fill_one_to_many_hash_table_on_device (int32_t *buff, const HashEntryInfo hash_entry_info, const int32_t invalid_slot_val, const JoinColumn &join_column, const JoinColumnTypeInfo &type_info, const size_t block_size_x, const size_t grid_size_x)
 
void fill_one_to_many_hash_table_on_device_bucketized (int32_t *buff, const HashEntryInfo hash_entry_info, const int32_t invalid_slot_val, const JoinColumn &join_column, const JoinColumnTypeInfo &type_info, const size_t block_size_x, const size_t grid_size_x)
 
void fill_one_to_many_hash_table_on_device_sharded (int32_t *buff, const HashEntryInfo hash_entry_info, const int32_t invalid_slot_val, const JoinColumn &join_column, const JoinColumnTypeInfo &type_info, const ShardInfo &shard_info, const size_t block_size_x, const size_t grid_size_x)
 
int fill_baseline_hash_join_buff_32 (int8_t *hash_buff, const size_t entry_count, const int32_t invalid_slot_val, const size_t key_component_count, const bool with_val_slot, const GenericKeyHandler *key_handler, const size_t num_elems, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
 
int overlaps_fill_baseline_hash_join_buff_32 (int8_t *hash_buff, const size_t entry_count, const int32_t invalid_slot_val, const size_t key_component_count, const bool with_val_slot, const OverlapsKeyHandler *key_handler, const size_t num_elems, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
 
int fill_baseline_hash_join_buff_64 (int8_t *hash_buff, const size_t entry_count, const int32_t invalid_slot_val, const size_t key_component_count, const bool with_val_slot, const GenericKeyHandler *key_handler, const size_t num_elems, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
 
int overlaps_fill_baseline_hash_join_buff_64 (int8_t *hash_buff, const size_t entry_count, const int32_t invalid_slot_val, const size_t key_component_count, const bool with_val_slot, const OverlapsKeyHandler *key_handler, const size_t num_elems, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
 
void fill_baseline_hash_join_buff_on_device_32 (int8_t *hash_buff, const size_t entry_count, const int32_t invalid_slot_val, const size_t key_component_count, const bool with_val_slot, int *dev_err_buff, const GenericKeyHandler *key_handler, const size_t num_elems, const size_t block_size_x, const size_t grid_size_x)
 
void fill_baseline_hash_join_buff_on_device_64 (int8_t *hash_buff, const size_t entry_count, const int32_t invalid_slot_val, const size_t key_component_count, const bool with_val_slot, int *dev_err_buff, const GenericKeyHandler *key_handler, const size_t num_elems, const size_t block_size_x, const size_t grid_size_x)
 
void overlaps_fill_baseline_hash_join_buff_on_device_64 (int8_t *hash_buff, const size_t entry_count, const int32_t invalid_slot_val, const size_t key_component_count, const bool with_val_slot, int *dev_err_buff, const OverlapsKeyHandler *key_handler, const size_t num_elems, const size_t block_size_x, const size_t grid_size_x)
 
void fill_one_to_many_baseline_hash_table_32 (int32_t *buff, const int32_t *composite_key_dict, const size_t hash_entry_count, const int32_t invalid_slot_val, const size_t key_component_count, const std::vector< JoinColumn > &join_column_per_key, const std::vector< JoinColumnTypeInfo > &type_info_per_key, const std::vector< JoinBucketInfo > &join_bucket_info, const std::vector< const void * > &sd_inner_proxy_per_key, const std::vector< const void * > &sd_outer_proxy_per_key, const int32_t cpu_thread_count)
 
void fill_one_to_many_baseline_hash_table_64 (int32_t *buff, const int64_t *composite_key_dict, const size_t hash_entry_count, const int32_t invalid_slot_val, const size_t key_component_count, const std::vector< JoinColumn > &join_column_per_key, const std::vector< JoinColumnTypeInfo > &type_info_per_key, const std::vector< JoinBucketInfo > &join_bucket_info, const std::vector< const void * > &sd_inner_proxy_per_key, const std::vector< const void * > &sd_outer_proxy_per_key, const int32_t cpu_thread_count)
 
void fill_one_to_many_baseline_hash_table_on_device_32 (int32_t *buff, const int32_t *composite_key_dict, const size_t hash_entry_count, const int32_t invalid_slot_val, const size_t key_component_count, const GenericKeyHandler *key_handler, const size_t num_elems, const size_t block_size_x, const size_t grid_size_x)
 
void fill_one_to_many_baseline_hash_table_on_device_64 (int32_t *buff, const int64_t *composite_key_dict, const size_t hash_entry_count, const int32_t invalid_slot_val, const GenericKeyHandler *key_handler, const size_t num_elems, const size_t block_size_x, const size_t grid_size_x)
 
void overlaps_fill_one_to_many_baseline_hash_table_on_device_64 (int32_t *buff, const int64_t *composite_key_dict, const size_t hash_entry_count, const int32_t invalid_slot_val, const OverlapsKeyHandler *key_handler, const size_t num_elems, const size_t block_size_x, const size_t grid_size_x)
 
void approximate_distinct_tuples (uint8_t *hll_buffer_all_cpus, const uint32_t b, const size_t padded_size_bytes, const std::vector< JoinColumn > &join_column_per_key, const std::vector< JoinColumnTypeInfo > &type_info_per_key, const int thread_count)
 
void approximate_distinct_tuples_overlaps (uint8_t *hll_buffer_all_cpus, std::vector< int32_t > &row_counts, const uint32_t b, const size_t padded_size_bytes, const std::vector< JoinColumn > &join_column_per_key, const std::vector< JoinColumnTypeInfo > &type_info_per_key, const std::vector< JoinBucketInfo > &join_buckets_per_key, const int thread_count)
 
void approximate_distinct_tuples_on_device (uint8_t *hll_buffer, const uint32_t b, const GenericKeyHandler *key_handler, const size_t num_elems, const size_t block_size_x, const size_t grid_size_x)
 
void approximate_distinct_tuples_on_device_overlaps (uint8_t *hll_buffer, const uint32_t b, int32_t *row_counts_buffer, const OverlapsKeyHandler *key_handler, const size_t num_elems, const size_t block_size_x, const size_t grid_size_x)
 
void compute_bucket_sizes (std::vector< double > &bucket_sizes_for_dimension, const JoinColumn &join_column, const JoinColumnTypeInfo &type_info, const double bucket_size_threshold, const int thread_count)
 
void compute_bucket_sizes_on_device (double *bucket_sizes_buffer, const JoinColumn *join_column, const JoinColumnTypeInfo *type_info, const double bucket_sz_threshold, const size_t block_size_x, const size_t grid_size_x)
 

Variables

const size_t g_maximum_conditions_to_coalesce
 

Enumeration Type Documentation

enum ColumnType
Enumerator
SmallDate 
Signed 
Unsigned 
Double 

Definition at line 106 of file HashJoinRuntime.h.

Function Documentation

void approximate_distinct_tuples ( uint8_t *  hll_buffer_all_cpus,
const uint32_t  b,
const size_t  padded_size_bytes,
const std::vector< JoinColumn > &  join_column_per_key,
const std::vector< JoinColumnTypeInfo > &  type_info_per_key,
const int  thread_count 
)

Definition at line 1922 of file HashJoinRuntime.cpp.

References approximate_distinct_tuples_impl(), CHECK(), and CHECK_EQ.

Referenced by BaselineJoinHashTable::approximateTupleCount().

1927  {
1928  CHECK_EQ(join_column_per_key.size(), type_info_per_key.size());
1929  CHECK(!join_column_per_key.empty());
1930 
1931  std::vector<std::future<void>> approx_distinct_threads;
1932  for (int thread_idx = 0; thread_idx < thread_count; ++thread_idx) {
1933  approx_distinct_threads.push_back(std::async(
1934  std::launch::async,
1935  [&join_column_per_key,
1936  &type_info_per_key,
1937  b,
1938  hll_buffer_all_cpus,
1939  padded_size_bytes,
1940  thread_idx,
1941  thread_count] {
1942  auto hll_buffer = hll_buffer_all_cpus + thread_idx * padded_size_bytes;
1943 
1944  const auto key_handler = GenericKeyHandler(join_column_per_key.size(),
1945  false,
1946  &join_column_per_key[0],
1947  &type_info_per_key[0],
1948  nullptr,
1949  nullptr);
1951  nullptr,
1952  b,
1953  join_column_per_key[0].num_elems,
1954  &key_handler,
1955  thread_idx,
1956  thread_count);
1957  }));
1958  }
1959  for (auto& child : approx_distinct_threads) {
1960  child.get();
1961  }
1962 }
#define CHECK_EQ(x, y)
Definition: Logger.h:205
GLOBAL void SUFFIX() approximate_distinct_tuples_impl(uint8_t *hll_buffer, int32_t *row_count_buffer, const uint32_t b, const size_t num_elems, const KEY_HANDLER *f, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
CHECK(cgen_state)

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

void approximate_distinct_tuples_on_device ( uint8_t *  hll_buffer,
const uint32_t  b,
const GenericKeyHandler key_handler,
const size_t  num_elems,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 546 of file HashJoinRuntimeGpu.cu.

Referenced by BaselineJoinHashTable::approximateTupleCount().

551  {
552  approximate_distinct_tuples_impl_gpu<<<grid_size_x, block_size_x>>>(
553  hll_buffer, nullptr, b, num_elems, key_handler);
554 }

+ Here is the caller graph for this function:

void approximate_distinct_tuples_on_device_overlaps ( uint8_t *  hll_buffer,
const uint32_t  b,
int32_t *  row_counts_buffer,
const OverlapsKeyHandler key_handler,
const size_t  num_elems,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 531 of file HashJoinRuntimeGpu.cu.

References inclusive_scan().

Referenced by OverlapsJoinHashTable::approximateTupleCount().

537  {
538  approximate_distinct_tuples_impl_gpu<<<grid_size_x, block_size_x>>>(
539  hll_buffer, row_counts_buffer, b, num_elems, key_handler);
540 
541  auto row_counts_buffer_ptr = thrust::device_pointer_cast(row_counts_buffer);
543  row_counts_buffer_ptr, row_counts_buffer_ptr + num_elems, row_counts_buffer_ptr);
544 }
void inclusive_scan(InputIterator first, InputIterator last, OutputIterator out, const size_t thread_count)

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

void approximate_distinct_tuples_overlaps ( uint8_t *  hll_buffer_all_cpus,
std::vector< int32_t > &  row_counts,
const uint32_t  b,
const size_t  padded_size_bytes,
const std::vector< JoinColumn > &  join_column_per_key,
const std::vector< JoinColumnTypeInfo > &  type_info_per_key,
const std::vector< JoinBucketInfo > &  join_buckets_per_key,
const int  thread_count 
)

Definition at line 1964 of file HashJoinRuntime.cpp.

References approximate_distinct_tuples_impl(), CHECK(), CHECK_EQ, and inclusive_scan().

Referenced by OverlapsJoinHashTable::approximateTupleCount().

1972  {
1973  CHECK_EQ(join_column_per_key.size(), join_buckets_per_key.size());
1974  CHECK_EQ(join_column_per_key.size(), type_info_per_key.size());
1975  CHECK(!join_column_per_key.empty());
1976 
1977  std::vector<std::future<void>> approx_distinct_threads;
1978  for (int thread_idx = 0; thread_idx < thread_count; ++thread_idx) {
1979  approx_distinct_threads.push_back(std::async(
1980  std::launch::async,
1981  [&join_column_per_key,
1982  &join_buckets_per_key,
1983  &row_counts,
1984  b,
1985  hll_buffer_all_cpus,
1986  padded_size_bytes,
1987  thread_idx,
1988  thread_count] {
1989  auto hll_buffer = hll_buffer_all_cpus + thread_idx * padded_size_bytes;
1990 
1991  const auto key_handler = OverlapsKeyHandler(
1992  join_buckets_per_key[0].bucket_sizes_for_dimension.size(),
1993  &join_column_per_key[0],
1994  join_buckets_per_key[0].bucket_sizes_for_dimension.data());
1996  row_counts.data(),
1997  b,
1998  join_column_per_key[0].num_elems,
1999  &key_handler,
2000  thread_idx,
2001  thread_count);
2002  }));
2003  }
2004  for (auto& child : approx_distinct_threads) {
2005  child.get();
2006  }
2007 
2009  row_counts.begin(), row_counts.end(), row_counts.begin(), thread_count);
2010 }
#define CHECK_EQ(x, y)
Definition: Logger.h:205
GLOBAL void SUFFIX() approximate_distinct_tuples_impl(uint8_t *hll_buffer, int32_t *row_count_buffer, const uint32_t b, const size_t num_elems, const KEY_HANDLER *f, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
void inclusive_scan(InputIterator first, InputIterator last, OutputIterator out, const size_t thread_count)
CHECK(cgen_state)

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

void compute_bucket_sizes ( std::vector< double > &  bucket_sizes_for_dimension,
const JoinColumn join_column,
const JoinColumnTypeInfo type_info,
const double  bucket_size_threshold,
const int  thread_count 
)

Definition at line 2012 of file HashJoinRuntime.cpp.

Referenced by OverlapsJoinHashTable::computeBucketSizes().

2016  {
2017  std::vector<std::vector<double>> bucket_sizes_for_threads;
2018  for (int thread_idx = 0; thread_idx < thread_count; ++thread_idx) {
2019  bucket_sizes_for_threads.emplace_back(bucket_sizes_for_dimension.size(),
2020  std::numeric_limits<double>::max());
2021  }
2022  std::vector<std::future<void>> threads;
2023  for (int thread_idx = 0; thread_idx < thread_count; ++thread_idx) {
2024  threads.push_back(std::async(std::launch::async,
2025  compute_bucket_sizes_impl<2>,
2026  bucket_sizes_for_threads[thread_idx].data(),
2027  &join_column,
2028  &type_info,
2029  bucket_size_threshold,
2030  thread_idx,
2031  thread_count));
2032  }
2033  for (auto& child : threads) {
2034  child.get();
2035  }
2036 
2037  for (int thread_idx = 0; thread_idx < thread_count; ++thread_idx) {
2038  for (size_t i = 0; i < bucket_sizes_for_dimension.size(); i++) {
2039  if (bucket_sizes_for_threads[thread_idx][i] < bucket_sizes_for_dimension[i]) {
2040  bucket_sizes_for_dimension[i] = bucket_sizes_for_threads[thread_idx][i];
2041  }
2042  }
2043  }
2044 }

+ Here is the caller graph for this function:

void compute_bucket_sizes_on_device ( double *  bucket_sizes_buffer,
const JoinColumn join_column,
const JoinColumnTypeInfo type_info,
const double  bucket_sz_threshold,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 556 of file HashJoinRuntimeGpu.cu.

Referenced by OverlapsJoinHashTable::computeBucketSizes().

561  {
562  compute_bucket_sizes_impl_gpu<2><<<grid_size_x, block_size_x>>>(bucket_sizes_buffer,
563  join_column,
564  type_info,
565  bucket_sz_threshold,
566  block_size_x,
567  grid_size_x);
568 }

+ Here is the caller graph for this function:

int fill_baseline_hash_join_buff_32 ( int8_t *  hash_buff,
const size_t  entry_count,
const int32_t  invalid_slot_val,
const size_t  key_component_count,
const bool  with_val_slot,
const GenericKeyHandler key_handler,
const size_t  num_elems,
const int32_t  cpu_thread_idx,
const int32_t  cpu_thread_count 
)

Definition at line 1630 of file HashJoinRuntime.cpp.

Referenced by BaselineJoinHashTable::initHashTableOnCpu().

1638  {
1639  return fill_baseline_hash_join_buff<int32_t>(hash_buff,
1640  entry_count,
1641  invalid_slot_val,
1642  key_component_count,
1643  with_val_slot,
1644  key_handler,
1645  num_elems,
1646  cpu_thread_idx,
1647  cpu_thread_count);
1648 }

+ Here is the caller graph for this function:

int fill_baseline_hash_join_buff_64 ( int8_t *  hash_buff,
const size_t  entry_count,
const int32_t  invalid_slot_val,
const size_t  key_component_count,
const bool  with_val_slot,
const GenericKeyHandler key_handler,
const size_t  num_elems,
const int32_t  cpu_thread_idx,
const int32_t  cpu_thread_count 
)

Definition at line 1670 of file HashJoinRuntime.cpp.

Referenced by BaselineJoinHashTable::initHashTableOnCpu().

1678  {
1679  return fill_baseline_hash_join_buff<int64_t>(hash_buff,
1680  entry_count,
1681  invalid_slot_val,
1682  key_component_count,
1683  with_val_slot,
1684  key_handler,
1685  num_elems,
1686  cpu_thread_idx,
1687  cpu_thread_count);
1688 }

+ Here is the caller graph for this function:

void fill_baseline_hash_join_buff_on_device_32 ( int8_t *  hash_buff,
const size_t  entry_count,
const int32_t  invalid_slot_val,
const size_t  key_component_count,
const bool  with_val_slot,
int *  dev_err_buff,
const GenericKeyHandler key_handler,
const size_t  num_elems,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 409 of file HashJoinRuntimeGpu.cu.

Referenced by BaselineJoinHashTable::initHashTableOnGpu().

418  {
419  fill_baseline_hash_join_buff_wrapper<int32_t>
420  <<<grid_size_x, block_size_x>>>(hash_buff,
421  entry_count,
422  invalid_slot_val,
423  key_component_count,
424  with_val_slot,
425  dev_err_buff,
426  key_handler,
427  num_elems);
428 }

+ Here is the caller graph for this function:

void fill_baseline_hash_join_buff_on_device_64 ( int8_t *  hash_buff,
const size_t  entry_count,
const int32_t  invalid_slot_val,
const size_t  key_component_count,
const bool  with_val_slot,
int *  dev_err_buff,
const GenericKeyHandler key_handler,
const size_t  num_elems,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 430 of file HashJoinRuntimeGpu.cu.

Referenced by BaselineJoinHashTable::initHashTableOnGpu().

439  {
440  fill_baseline_hash_join_buff_wrapper<unsigned long long>
441  <<<grid_size_x, block_size_x>>>(hash_buff,
442  entry_count,
443  invalid_slot_val,
444  key_component_count,
445  with_val_slot,
446  dev_err_buff,
447  key_handler,
448  num_elems);
449 }

+ Here is the caller graph for this function:

int fill_hash_join_buff ( int32_t *  buff,
const int32_t  invalid_slot_val,
const JoinColumn  join_column,
const JoinColumnTypeInfo  type_info,
const void *  sd_inner,
const void *  sd_outer,
const int32_t  cpu_thread_idx,
const int32_t  cpu_thread_count 
)

Definition at line 187 of file HashJoinRuntime.cpp.

References fill_hash_join_buff_impl(), get_hash_slot(), and SUFFIX.

Referenced by fill_hash_join_buff_wrapper().

194  {
195  auto slot_selector = [&](auto elem) {
196  return SUFFIX(get_hash_slot)(buff, elem, type_info.min_val);
197  };
198  return fill_hash_join_buff_impl(buff,
199  invalid_slot_val,
200  join_column,
201  type_info,
202  sd_inner_proxy,
203  sd_outer_proxy,
204  cpu_thread_idx,
205  cpu_thread_count,
206  slot_selector);
207 }
#define SUFFIX(name)
ALWAYS_INLINE DEVICE int32_t *SUFFIX() get_hash_slot(int32_t *buff, const int64_t key, const int64_t min_key)
Definition: JoinHashImpl.h:39
DEVICE auto fill_hash_join_buff_impl(int32_t *buff, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const int32_t cpu_thread_idx, const int32_t cpu_thread_count, SLOT_SELECTOR slot_sel)
const int64_t min_val

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

int fill_hash_join_buff_bucketized ( int32_t *  buff,
const int32_t  invalid_slot_val,
const JoinColumn  join_column,
const JoinColumnTypeInfo  type_info,
const void *  sd_inner,
const void *  sd_outer,
const int32_t  cpu_thread_idx,
const int32_t  cpu_thread_count,
const int64_t  bucket_normalization 
)

Definition at line 163 of file HashJoinRuntime.cpp.

References fill_hash_join_buff_impl(), get_bucketized_hash_slot(), and SUFFIX.

Referenced by fill_hash_join_buff_bucketized_wrapper(), and JoinHashTable::initOneToOneHashTableOnCpu().

171  {
172  auto slot_selector = [&](auto elem) {
174  buff, elem, type_info.min_val, bucket_normalization);
175  };
176  return fill_hash_join_buff_impl(buff,
177  invalid_slot_val,
178  join_column,
179  type_info,
180  sd_inner_proxy,
181  sd_outer_proxy,
182  cpu_thread_idx,
183  cpu_thread_count,
184  slot_selector);
185 }
ALWAYS_INLINE DEVICE int32_t *SUFFIX() get_bucketized_hash_slot(int32_t *buff, const int64_t key, const int64_t min_key, const int64_t bucket_normalization)
Definition: JoinHashImpl.h:31
#define SUFFIX(name)
DEVICE auto fill_hash_join_buff_impl(int32_t *buff, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const int32_t cpu_thread_idx, const int32_t cpu_thread_count, SLOT_SELECTOR slot_sel)
const int64_t min_val

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

void fill_hash_join_buff_on_device ( int32_t *  buff,
const int32_t  invalid_slot_val,
int *  dev_err_buff,
const JoinColumn  join_column,
const JoinColumnTypeInfo  type_info,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 62 of file HashJoinRuntimeGpu.cu.

68  {
69  fill_hash_join_buff_wrapper<<<grid_size_x, block_size_x>>>(
70  buff, invalid_slot_val, join_column, type_info, dev_err_buff);
71 }
void fill_hash_join_buff_on_device_bucketized ( int32_t *  buff,
const int32_t  invalid_slot_val,
int *  dev_err_buff,
const JoinColumn  join_column,
const JoinColumnTypeInfo  type_info,
const size_t  block_size_x,
const size_t  grid_size_x,
const int64_t  bucket_normalization 
)

Definition at line 50 of file HashJoinRuntimeGpu.cu.

Referenced by JoinHashTable::initOneToOneHashTable().

57  {
58  fill_hash_join_buff_bucketized_wrapper<<<grid_size_x, block_size_x>>>(
59  buff, invalid_slot_val, join_column, type_info, dev_err_buff, bucket_normalization);
60 }

+ Here is the caller graph for this function:

void fill_hash_join_buff_on_device_sharded ( int32_t *  buff,
const int32_t  invalid_slot_val,
int *  dev_err_buff,
const JoinColumn  join_column,
const JoinColumnTypeInfo  type_info,
const ShardInfo  shard_info,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 125 of file HashJoinRuntimeGpu.cu.

132  {
133  fill_hash_join_buff_wrapper_sharded<<<grid_size_x, block_size_x>>>(
134  buff, invalid_slot_val, join_column, type_info, shard_info, dev_err_buff);
135 }
void fill_hash_join_buff_on_device_sharded_bucketized ( int32_t *  buff,
const int32_t  invalid_slot_val,
int *  dev_err_buff,
const JoinColumn  join_column,
const JoinColumnTypeInfo  type_info,
const ShardInfo  shard_info,
const size_t  block_size_x,
const size_t  grid_size_x,
const int64_t  bucket_normalization 
)

Definition at line 105 of file HashJoinRuntimeGpu.cu.

Referenced by JoinHashTable::initOneToOneHashTable().

114  {
115  fill_hash_join_buff_wrapper_sharded_bucketized<<<grid_size_x, block_size_x>>>(
116  buff,
117  invalid_slot_val,
118  join_column,
119  type_info,
120  shard_info,
121  dev_err_buff,
122  bucket_normalization);
123 }

+ Here is the caller graph for this function:

void fill_one_to_many_baseline_hash_table_32 ( int32_t *  buff,
const int32_t *  composite_key_dict,
const size_t  hash_entry_count,
const int32_t  invalid_slot_val,
const size_t  key_component_count,
const std::vector< JoinColumn > &  join_column_per_key,
const std::vector< JoinColumnTypeInfo > &  type_info_per_key,
const std::vector< JoinBucketInfo > &  join_bucket_info,
const std::vector< const void * > &  sd_inner_proxy_per_key,
const std::vector< const void * > &  sd_outer_proxy_per_key,
const int32_t  cpu_thread_count 
)

Definition at line 1872 of file HashJoinRuntime.cpp.

Referenced by OverlapsJoinHashTable::initHashTableOnCpu(), and BaselineJoinHashTable::initHashTableOnCpu().

1883  {
1884  fill_one_to_many_baseline_hash_table<int32_t>(buff,
1885  composite_key_dict,
1886  hash_entry_count,
1887  invalid_slot_val,
1888  key_component_count,
1889  join_column_per_key,
1890  type_info_per_key,
1891  join_bucket_info,
1892  sd_inner_proxy_per_key,
1893  sd_outer_proxy_per_key,
1894  cpu_thread_count);
1895 }

+ Here is the caller graph for this function:

void fill_one_to_many_baseline_hash_table_64 ( int32_t *  buff,
const int64_t *  composite_key_dict,
const size_t  hash_entry_count,
const int32_t  invalid_slot_val,
const size_t  key_component_count,
const std::vector< JoinColumn > &  join_column_per_key,
const std::vector< JoinColumnTypeInfo > &  type_info_per_key,
const std::vector< JoinBucketInfo > &  join_bucket_info,
const std::vector< const void * > &  sd_inner_proxy_per_key,
const std::vector< const void * > &  sd_outer_proxy_per_key,
const int32_t  cpu_thread_count 
)

Definition at line 1897 of file HashJoinRuntime.cpp.

Referenced by OverlapsJoinHashTable::initHashTableOnCpu(), and BaselineJoinHashTable::initHashTableOnCpu().

1908  {
1909  fill_one_to_many_baseline_hash_table<int64_t>(buff,
1910  composite_key_dict,
1911  hash_entry_count,
1912  invalid_slot_val,
1913  key_component_count,
1914  join_column_per_key,
1915  type_info_per_key,
1916  join_bucket_info,
1917  sd_inner_proxy_per_key,
1918  sd_outer_proxy_per_key,
1919  cpu_thread_count);
1920 }

+ Here is the caller graph for this function:

void fill_one_to_many_baseline_hash_table_on_device_32 ( int32_t *  buff,
const int32_t *  composite_key_dict,
const size_t  hash_entry_count,
const int32_t  invalid_slot_val,
const size_t  key_component_count,
const GenericKeyHandler key_handler,
const size_t  num_elems,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 473 of file HashJoinRuntimeGpu.cu.

Referenced by BaselineJoinHashTable::initHashTableOnGpu().

482  {
483  fill_one_to_many_baseline_hash_table_on_device<int32_t>(buff,
484  composite_key_dict,
485  hash_entry_count,
486  invalid_slot_val,
487  key_handler,
488  num_elems,
489  block_size_x,
490  grid_size_x);
491 }

+ Here is the caller graph for this function:

void fill_one_to_many_baseline_hash_table_on_device_64 ( int32_t *  buff,
const int64_t *  composite_key_dict,
const size_t  hash_entry_count,
const int32_t  invalid_slot_val,
const GenericKeyHandler key_handler,
const size_t  num_elems,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 493 of file HashJoinRuntimeGpu.cu.

Referenced by BaselineJoinHashTable::initHashTableOnGpu().

501  {
502  fill_one_to_many_baseline_hash_table_on_device<int64_t>(buff,
503  composite_key_dict,
504  hash_entry_count,
505  invalid_slot_val,
506  key_handler,
507  num_elems,
508  block_size_x,
509  grid_size_x);
510 }

+ Here is the caller graph for this function:

void fill_one_to_many_hash_table ( int32_t *  buff,
const HashEntryInfo  hash_entry_info,
const int32_t  invalid_slot_val,
const JoinColumn join_column,
const JoinColumnTypeInfo type_info,
const void *  sd_inner_proxy,
const void *  sd_outer_proxy,
const unsigned  cpu_thread_count 
)

Definition at line 1337 of file HashJoinRuntime.cpp.

References count_matches(), fill_one_to_many_hash_table_impl(), fill_row_ids(), HashEntryInfo::hash_entry_count, and SUFFIX.

Referenced by JoinHashTable::initOneToManyHashTableOnCpu().

1344  {
1345  auto launch_count_matches = [count_buff = buff + hash_entry_info.hash_entry_count,
1346  invalid_slot_val,
1347  &join_column,
1348  &type_info,
1349  sd_inner_proxy,
1350  sd_outer_proxy](auto cpu_thread_idx,
1351  auto cpu_thread_count) {
1353  (count_buff,
1354  invalid_slot_val,
1355  join_column,
1356  type_info,
1357  sd_inner_proxy,
1358  sd_outer_proxy,
1359  cpu_thread_idx,
1360  cpu_thread_count);
1361  };
1362  auto launch_fill_row_ids = [hash_entry_count = hash_entry_info.hash_entry_count,
1363  buff,
1364  invalid_slot_val,
1365  &join_column,
1366  &type_info,
1367  sd_inner_proxy,
1368  sd_outer_proxy](auto cpu_thread_idx,
1369  auto cpu_thread_count) {
1371  (buff,
1372  hash_entry_count,
1373  invalid_slot_val,
1374  join_column,
1375  type_info,
1376  sd_inner_proxy,
1377  sd_outer_proxy,
1378  cpu_thread_idx,
1379  cpu_thread_count);
1380  };
1381 
1383  hash_entry_info.hash_entry_count,
1384  invalid_slot_val,
1385  join_column,
1386  type_info,
1387  sd_inner_proxy,
1388  sd_outer_proxy,
1389  cpu_thread_count,
1390  launch_count_matches,
1391  launch_fill_row_ids);
1392 }
#define SUFFIX(name)
GLOBAL void SUFFIX() fill_row_ids(int32_t *buff, const int32_t hash_entry_count, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
size_t hash_entry_count
GLOBAL void SUFFIX() count_matches(int32_t *count_buff, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
void fill_one_to_many_hash_table_impl(int32_t *buff, const int32_t hash_entry_count, const int32_t invalid_slot_val, const JoinColumn &join_column, const JoinColumnTypeInfo &type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const unsigned cpu_thread_count, COUNT_MATCHES_LAUNCH_FUNCTOR count_matches_func, FILL_ROW_IDS_LAUNCH_FUNCTOR fill_row_ids_func)

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

void fill_one_to_many_hash_table_bucketized ( int32_t *  buff,
const HashEntryInfo  hash_entry_info,
const int32_t  invalid_slot_val,
const JoinColumn join_column,
const JoinColumnTypeInfo type_info,
const void *  sd_inner_proxy,
const void *  sd_outer_proxy,
const unsigned  cpu_thread_count 
)

Definition at line 1394 of file HashJoinRuntime.cpp.

References HashEntryInfo::bucket_normalization, count_matches_bucketized(), fill_one_to_many_hash_table_impl(), fill_row_ids_bucketized(), HashEntryInfo::getNormalizedHashEntryCount(), and SUFFIX.

Referenced by JoinHashTable::initOneToManyHashTableOnCpu().

1401  {
1402  auto bucket_normalization = hash_entry_info.bucket_normalization;
1403  auto hash_entry_count = hash_entry_info.getNormalizedHashEntryCount();
1404  auto launch_count_matches = [bucket_normalization,
1405  count_buff = buff + hash_entry_count,
1406  invalid_slot_val,
1407  &join_column,
1408  &type_info,
1409  sd_inner_proxy,
1410  sd_outer_proxy](auto cpu_thread_idx,
1411  auto cpu_thread_count) {
1413  (count_buff,
1414  invalid_slot_val,
1415  join_column,
1416  type_info,
1417  sd_inner_proxy,
1418  sd_outer_proxy,
1419  cpu_thread_idx,
1420  cpu_thread_count,
1421  bucket_normalization);
1422  };
1423  auto launch_fill_row_ids = [bucket_normalization,
1424  hash_entry_count,
1425  buff,
1426  invalid_slot_val,
1427  &join_column,
1428  &type_info,
1429  sd_inner_proxy,
1430  sd_outer_proxy](auto cpu_thread_idx,
1431  auto cpu_thread_count) {
1433  (buff,
1434  hash_entry_count,
1435  invalid_slot_val,
1436  join_column,
1437  type_info,
1438  sd_inner_proxy,
1439  sd_outer_proxy,
1440  cpu_thread_idx,
1441  cpu_thread_count,
1442  bucket_normalization);
1443  };
1444 
1446  hash_entry_count,
1447  invalid_slot_val,
1448  join_column,
1449  type_info,
1450  sd_inner_proxy,
1451  sd_outer_proxy,
1452  cpu_thread_count,
1453  launch_count_matches,
1454  launch_fill_row_ids);
1455 }
GLOBAL void SUFFIX() count_matches_bucketized(int32_t *count_buff, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const int32_t cpu_thread_idx, const int32_t cpu_thread_count, const int64_t bucket_normalization)
GLOBAL void SUFFIX() fill_row_ids_bucketized(int32_t *buff, const int32_t hash_entry_count, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const int32_t cpu_thread_idx, const int32_t cpu_thread_count, const int64_t bucket_normalization)
#define SUFFIX(name)
int64_t bucket_normalization
size_t getNormalizedHashEntryCount() const
void fill_one_to_many_hash_table_impl(int32_t *buff, const int32_t hash_entry_count, const int32_t invalid_slot_val, const JoinColumn &join_column, const JoinColumnTypeInfo &type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const unsigned cpu_thread_count, COUNT_MATCHES_LAUNCH_FUNCTOR count_matches_func, FILL_ROW_IDS_LAUNCH_FUNCTOR fill_row_ids_func)

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

void fill_one_to_many_hash_table_on_device ( int32_t *  buff,
const HashEntryInfo  hash_entry_info,
const int32_t  invalid_slot_val,
const JoinColumn join_column,
const JoinColumnTypeInfo type_info,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 204 of file HashJoinRuntimeGpu.cu.

References count_matches(), fill_one_to_many_hash_table_on_device_impl(), fill_row_ids(), HashEntryInfo::hash_entry_count, and SUFFIX.

Referenced by JoinHashTable::initOneToManyHashTable().

210  {
211  auto hash_entry_count = hash_entry_info.hash_entry_count;
212  auto count_matches_func = [hash_entry_count,
213  grid_size_x,
214  block_size_x,
215  count_buff = buff + hash_entry_count,
216  invalid_slot_val,
217  join_column,
218  type_info] {
219  SUFFIX(count_matches)<<<grid_size_x, block_size_x>>>(
220  count_buff, invalid_slot_val, join_column, type_info);
221  };
222 
223  auto fill_row_ids_func = [grid_size_x,
224  block_size_x,
225  buff,
226  hash_entry_count,
227  invalid_slot_val,
228  join_column,
229  type_info] {
230  SUFFIX(fill_row_ids)<<<grid_size_x, block_size_x>>>(
231  buff, hash_entry_count, invalid_slot_val, join_column, type_info);
232  };
233 
235  hash_entry_count,
236  invalid_slot_val,
237  join_column,
238  type_info,
239  block_size_x,
240  grid_size_x,
241  count_matches_func,
242  fill_row_ids_func);
243 }
#define SUFFIX(name)
void fill_one_to_many_hash_table_on_device_impl(int32_t *buff, const int32_t hash_entry_count, const int32_t invalid_slot_val, const JoinColumn &join_column, const JoinColumnTypeInfo &type_info, const size_t block_size_x, const size_t grid_size_x, COUNT_MATCHES_FUNCTOR count_matches_func, FILL_ROW_IDS_FUNCTOR fill_row_ids_func)
GLOBAL void SUFFIX() fill_row_ids(int32_t *buff, const int32_t hash_entry_count, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
size_t hash_entry_count
GLOBAL void SUFFIX() count_matches(int32_t *count_buff, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

void fill_one_to_many_hash_table_on_device_bucketized ( int32_t *  buff,
const HashEntryInfo  hash_entry_info,
const int32_t  invalid_slot_val,
const JoinColumn join_column,
const JoinColumnTypeInfo type_info,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 245 of file HashJoinRuntimeGpu.cu.

References HashEntryInfo::bucket_normalization, count_matches_bucketized(), fill_one_to_many_hash_table_on_device_impl(), fill_row_ids_bucketized(), HashEntryInfo::getNormalizedHashEntryCount(), and SUFFIX.

Referenced by JoinHashTable::initOneToManyHashTable().

251  {
252  auto hash_entry_count = hash_entry_info.getNormalizedHashEntryCount();
253  auto count_matches_func = [grid_size_x,
254  block_size_x,
255  count_buff = buff + hash_entry_count,
256  invalid_slot_val,
257  join_column,
258  type_info,
259  bucket_normalization =
260  hash_entry_info.bucket_normalization] {
261  SUFFIX(count_matches_bucketized)<<<grid_size_x, block_size_x>>>(
262  count_buff, invalid_slot_val, join_column, type_info, bucket_normalization);
263  };
264 
265  auto fill_row_ids_func = [grid_size_x,
266  block_size_x,
267  buff,
268  hash_entry_count =
269  hash_entry_info.getNormalizedHashEntryCount(),
270  invalid_slot_val,
271  join_column,
272  type_info,
273  bucket_normalization = hash_entry_info.bucket_normalization] {
274  SUFFIX(fill_row_ids_bucketized)<<<grid_size_x, block_size_x>>>(buff,
275  hash_entry_count,
276  invalid_slot_val,
277  join_column,
278  type_info,
279  bucket_normalization);
280  };
281 
283  hash_entry_count,
284  invalid_slot_val,
285  join_column,
286  type_info,
287  block_size_x,
288  grid_size_x,
289  count_matches_func,
290  fill_row_ids_func);
291 }
GLOBAL void SUFFIX() count_matches_bucketized(int32_t *count_buff, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const int32_t cpu_thread_idx, const int32_t cpu_thread_count, const int64_t bucket_normalization)
GLOBAL void SUFFIX() fill_row_ids_bucketized(int32_t *buff, const int32_t hash_entry_count, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const int32_t cpu_thread_idx, const int32_t cpu_thread_count, const int64_t bucket_normalization)
#define SUFFIX(name)
void fill_one_to_many_hash_table_on_device_impl(int32_t *buff, const int32_t hash_entry_count, const int32_t invalid_slot_val, const JoinColumn &join_column, const JoinColumnTypeInfo &type_info, const size_t block_size_x, const size_t grid_size_x, COUNT_MATCHES_FUNCTOR count_matches_func, FILL_ROW_IDS_FUNCTOR fill_row_ids_func)
int64_t bucket_normalization
size_t getNormalizedHashEntryCount() const

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

void fill_one_to_many_hash_table_on_device_sharded ( int32_t *  buff,
const HashEntryInfo  hash_entry_info,
const int32_t  invalid_slot_val,
const JoinColumn join_column,
const JoinColumnTypeInfo type_info,
const ShardInfo shard_info,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 293 of file HashJoinRuntimeGpu.cu.

References count_matches_sharded(), fill_row_ids_sharded(), HashEntryInfo::hash_entry_count, inclusive_scan(), and SUFFIX.

Referenced by JoinHashTable::initOneToManyHashTable().

300  {
301  auto hash_entry_count = hash_entry_info.hash_entry_count;
302  int32_t* pos_buff = buff;
303  int32_t* count_buff = buff + hash_entry_count;
304  cudaMemset(count_buff, 0, hash_entry_count * sizeof(int32_t));
305  SUFFIX(count_matches_sharded)<<<grid_size_x, block_size_x>>>(
306  count_buff, invalid_slot_val, join_column, type_info, shard_info);
307 
308  set_valid_pos_flag<<<grid_size_x, block_size_x>>>(
309  pos_buff, count_buff, hash_entry_count);
310 
311  auto count_buff_dev_ptr = thrust::device_pointer_cast(count_buff);
313  count_buff_dev_ptr, count_buff_dev_ptr + hash_entry_count, count_buff_dev_ptr);
314  set_valid_pos<<<grid_size_x, block_size_x>>>(pos_buff, count_buff, hash_entry_count);
315  cudaMemset(count_buff, 0, hash_entry_count * sizeof(int32_t));
316  SUFFIX(fill_row_ids_sharded)<<<grid_size_x, block_size_x>>>(
317  buff, hash_entry_count, invalid_slot_val, join_column, type_info, shard_info);
318 }
GLOBAL void SUFFIX() fill_row_ids_sharded(int32_t *buff, const int32_t hash_entry_count, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const ShardInfo shard_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
#define SUFFIX(name)
GLOBAL void SUFFIX() count_matches_sharded(int32_t *count_buff, const int32_t invalid_slot_val, const JoinColumn join_column, const JoinColumnTypeInfo type_info, const ShardInfo shard_info, const void *sd_inner_proxy, const void *sd_outer_proxy, const int32_t cpu_thread_idx, const int32_t cpu_thread_count)
void inclusive_scan(InputIterator first, InputIterator last, OutputIterator out, const size_t thread_count)
size_t hash_entry_count

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

void fill_one_to_many_hash_table_sharded_bucketized ( int32_t *  buff,
const HashEntryInfo  hash_entry_info,
const int32_t  invalid_slot_val,
const JoinColumn join_column,
const JoinColumnTypeInfo type_info,
const ShardInfo shard_info,
const void *  sd_inner_proxy,
const void *  sd_outer_proxy,
const unsigned  cpu_thread_count 
)
ColumnType get_join_column_type_kind ( const SQLTypeInfo ti)
inline

Definition at line 133 of file HashJoinRuntime.h.

References SQLTypeInfo::is_date_in_days(), is_unsigned_type(), Signed, SmallDate, and Unsigned.

Referenced by OverlapsJoinHashTable::fetchColumnsForDevice(), BaselineJoinHashTable::fetchColumnsForDevice(), JoinHashTable::initOneToManyHashTable(), JoinHashTable::initOneToManyHashTableOnCpu(), JoinHashTable::initOneToOneHashTable(), and JoinHashTable::initOneToOneHashTableOnCpu().

133  {
134  if (ti.is_date_in_days()) {
135  return SmallDate;
136  } else {
137  return is_unsigned_type(ti) ? Unsigned : Signed;
138  }
139 }
bool is_date_in_days() const
Definition: sqltypes.h:611
bool is_unsigned_type(const SQLTypeInfo &ti)

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

void init_baseline_hash_join_buff_32 ( int8_t *  hash_join_buff,
const int32_t  entry_count,
const size_t  key_component_count,
const bool  with_val_slot,
const int32_t  invalid_slot_val,
const int32_t  cpu_thread_idx,
const int32_t  cpu_thread_count 
)

Definition at line 1598 of file HashJoinRuntime.cpp.

Referenced by OverlapsJoinHashTable::initHashTableOnCpu(), and BaselineJoinHashTable::initHashTableOnCpu().

1604  {
1605  init_baseline_hash_join_buff<int32_t>(hash_join_buff,
1606  entry_count,
1607  key_component_count,
1608  with_val_slot,
1609  invalid_slot_val,
1610  cpu_thread_idx,
1611  cpu_thread_count);
1612 }

+ Here is the caller graph for this function:

void init_baseline_hash_join_buff_64 ( int8_t *  hash_join_buff,
const int32_t  entry_count,
const size_t  key_component_count,
const bool  with_val_slot,
const int32_t  invalid_slot_val,
const int32_t  cpu_thread_idx,
const int32_t  cpu_thread_count 
)

Definition at line 1614 of file HashJoinRuntime.cpp.

Referenced by OverlapsJoinHashTable::initHashTableOnCpu(), and BaselineJoinHashTable::initHashTableOnCpu().

1620  {
1621  init_baseline_hash_join_buff<int64_t>(hash_join_buff,
1622  entry_count,
1623  key_component_count,
1624  with_val_slot,
1625  invalid_slot_val,
1626  cpu_thread_idx,
1627  cpu_thread_count);
1628 }

+ Here is the caller graph for this function:

void init_baseline_hash_join_buff_on_device_32 ( int8_t *  hash_join_buff,
const int32_t  entry_count,
const size_t  key_component_count,
const bool  with_val_slot,
const int32_t  invalid_slot_val,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 366 of file HashJoinRuntimeGpu.cu.

Referenced by OverlapsJoinHashTable::initHashTableOnGpu(), and BaselineJoinHashTable::initHashTableOnGpu().

372  {
373  init_baseline_hash_join_buff_wrapper<int32_t><<<grid_size_x, block_size_x>>>(
374  hash_join_buff, entry_count, key_component_count, with_val_slot, invalid_slot_val);
375 }

+ Here is the caller graph for this function:

void init_baseline_hash_join_buff_on_device_64 ( int8_t *  hash_join_buff,
const int32_t  entry_count,
const size_t  key_component_count,
const bool  with_val_slot,
const int32_t  invalid_slot_val,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 377 of file HashJoinRuntimeGpu.cu.

Referenced by OverlapsJoinHashTable::initHashTableOnGpu(), and BaselineJoinHashTable::initHashTableOnGpu().

383  {
384  init_baseline_hash_join_buff_wrapper<int64_t><<<grid_size_x, block_size_x>>>(
385  hash_join_buff, entry_count, key_component_count, with_val_slot, invalid_slot_val);
386 }

+ Here is the caller graph for this function:

void init_hash_join_buff ( int32_t *  buff,
const int32_t  entry_count,
const int32_t  invalid_slot_val,
const int32_t  cpu_thread_idx,
const int32_t  cpu_thread_count 
)

Definition at line 91 of file HashJoinRuntime.cpp.

References groups_buffer, and Asio::start().

Referenced by init_hash_join_buff_wrapper(), OverlapsJoinHashTable::initHashTableOnCpu(), BaselineJoinHashTable::initHashTableOnCpu(), JoinHashTable::initOneToManyHashTableOnCpu(), and JoinHashTable::initOneToOneHashTableOnCpu().

95  {
96 #ifdef __CUDACC__
97  int32_t start = threadIdx.x + blockDim.x * blockIdx.x;
98  int32_t step = blockDim.x * gridDim.x;
99 #else
100  int32_t start = cpu_thread_idx;
101  int32_t step = cpu_thread_count;
102 #endif
103  for (int32_t i = start; i < hash_entry_count; i += step) {
104  groups_buffer[i] = invalid_slot_val;
105  }
106 }
const int32_t groups_buffer_size return groups_buffer
void start()
Definition: Asio.cpp:33

+ Here is the call graph for this function:

+ Here is the caller graph for this function:

void init_hash_join_buff_on_device ( int32_t *  buff,
const int32_t  entry_count,
const int32_t  invalid_slot_val,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 143 of file HashJoinRuntimeGpu.cu.

Referenced by OverlapsJoinHashTable::initHashTableOnGpu(), BaselineJoinHashTable::initHashTableOnGpu(), JoinHashTable::initOneToManyHashTable(), and JoinHashTable::initOneToOneHashTable().

147  {
148  init_hash_join_buff_wrapper<<<grid_size_x, block_size_x>>>(
149  buff, hash_entry_count, invalid_slot_val);
150 }

+ Here is the caller graph for this function:

int overlaps_fill_baseline_hash_join_buff_32 ( int8_t *  hash_buff,
const size_t  entry_count,
const int32_t  invalid_slot_val,
const size_t  key_component_count,
const bool  with_val_slot,
const OverlapsKeyHandler key_handler,
const size_t  num_elems,
const int32_t  cpu_thread_idx,
const int32_t  cpu_thread_count 
)

Definition at line 1650 of file HashJoinRuntime.cpp.

Referenced by OverlapsJoinHashTable::initHashTableOnCpu().

1658  {
1659  return fill_baseline_hash_join_buff<int32_t>(hash_buff,
1660  entry_count,
1661  invalid_slot_val,
1662  key_component_count,
1663  with_val_slot,
1664  key_handler,
1665  num_elems,
1666  cpu_thread_idx,
1667  cpu_thread_count);
1668 }

+ Here is the caller graph for this function:

int overlaps_fill_baseline_hash_join_buff_64 ( int8_t *  hash_buff,
const size_t  entry_count,
const int32_t  invalid_slot_val,
const size_t  key_component_count,
const bool  with_val_slot,
const OverlapsKeyHandler key_handler,
const size_t  num_elems,
const int32_t  cpu_thread_idx,
const int32_t  cpu_thread_count 
)

Definition at line 1690 of file HashJoinRuntime.cpp.

Referenced by OverlapsJoinHashTable::initHashTableOnCpu().

1698  {
1699  return fill_baseline_hash_join_buff<int64_t>(hash_buff,
1700  entry_count,
1701  invalid_slot_val,
1702  key_component_count,
1703  with_val_slot,
1704  key_handler,
1705  num_elems,
1706  cpu_thread_idx,
1707  cpu_thread_count);
1708 }

+ Here is the caller graph for this function:

void overlaps_fill_baseline_hash_join_buff_on_device_64 ( int8_t *  hash_buff,
const size_t  entry_count,
const int32_t  invalid_slot_val,
const size_t  key_component_count,
const bool  with_val_slot,
int *  dev_err_buff,
const OverlapsKeyHandler key_handler,
const size_t  num_elems,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 451 of file HashJoinRuntimeGpu.cu.

Referenced by OverlapsJoinHashTable::initHashTableOnGpu().

461  {
462  fill_baseline_hash_join_buff_wrapper<unsigned long long>
463  <<<grid_size_x, block_size_x>>>(hash_buff,
464  entry_count,
465  invalid_slot_val,
466  key_component_count,
467  with_val_slot,
468  dev_err_buff,
469  key_handler,
470  num_elems);
471 }

+ Here is the caller graph for this function:

void overlaps_fill_one_to_many_baseline_hash_table_on_device_64 ( int32_t *  buff,
const int64_t *  composite_key_dict,
const size_t  hash_entry_count,
const int32_t  invalid_slot_val,
const OverlapsKeyHandler key_handler,
const size_t  num_elems,
const size_t  block_size_x,
const size_t  grid_size_x 
)

Definition at line 512 of file HashJoinRuntimeGpu.cu.

Referenced by OverlapsJoinHashTable::initHashTableOnGpu().

520  {
521  fill_one_to_many_baseline_hash_table_on_device<int64_t>(buff,
522  composite_key_dict,
523  hash_entry_count,
524  invalid_slot_val,
525  key_handler,
526  num_elems,
527  block_size_x,
528  grid_size_x);
529 }

+ Here is the caller graph for this function:

Variable Documentation