GroupByRuntime.cpp File Reference

#include "JoinHashTable/Runtime/JoinHashImpl.h"
#include "MurmurHash.h"

Include dependency graph for GroupByRuntime.cpp:

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

Go to the source code of this file.

Macros
#define	DEF_TRANSLATE_NULL_KEY(key_type)

Functions
RUNTIME_EXPORT ALWAYS_INLINE DEVICE uint32_t	key_hash (const int64_t *key, const uint32_t key_count, const uint32_t key_byte_width)
RUNTIME_EXPORT NEVER_INLINE DEVICE int64_t *	get_group_value (int64_t *groups_buffer, const uint32_t groups_buffer_entry_count, const int64_t *key, const uint32_t key_count, const uint32_t key_width, const uint32_t row_size_quad)
RUNTIME_EXPORT NEVER_INLINE DEVICE bool	dynamic_watchdog ()
RUNTIME_EXPORT NEVER_INLINE DEVICE int64_t *	get_group_value_with_watchdog (int64_t *groups_buffer, const uint32_t groups_buffer_entry_count, const int64_t *key, const uint32_t key_count, const uint32_t key_width, const uint32_t row_size_quad)
RUNTIME_EXPORT NEVER_INLINE DEVICE int32_t	get_group_value_columnar_slot (int64_t *groups_buffer, const uint32_t groups_buffer_entry_count, const int64_t *key, const uint32_t key_count, const uint32_t key_width)
RUNTIME_EXPORT NEVER_INLINE DEVICE int32_t	get_group_value_columnar_slot_with_watchdog (int64_t *groups_buffer, const uint32_t groups_buffer_entry_count, const int64_t *key, const uint32_t key_count, const uint32_t key_width)
RUNTIME_EXPORT NEVER_INLINE DEVICE int64_t *	get_group_value_columnar (int64_t *groups_buffer, const uint32_t groups_buffer_entry_count, const int64_t *key, const uint32_t key_qw_count)
RUNTIME_EXPORT NEVER_INLINE DEVICE int64_t *	get_group_value_columnar_with_watchdog (int64_t *groups_buffer, const uint32_t groups_buffer_entry_count, const int64_t *key, const uint32_t key_qw_count)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t *	get_group_value_fast (int64_t *groups_buffer, const int64_t key, const int64_t min_key, const int64_t bucket, const uint32_t row_size_quad)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t *	get_group_value_fast_with_original_key (int64_t *groups_buffer, const int64_t key, const int64_t orig_key, const int64_t min_key, const int64_t bucket, const uint32_t row_size_quad)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE uint32_t	get_columnar_group_bin_offset (int64_t *key_base_ptr, const int64_t key, const int64_t min_key, const int64_t bucket)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t *	get_scan_output_slot (int64_t *output_buffer, const uint32_t output_buffer_entry_count, const uint32_t pos, const int64_t offset_in_fragment, const uint32_t row_size_quad)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE int32_t	get_columnar_scan_output_offset (int64_t *output_buffer, const uint32_t output_buffer_entry_count, const uint32_t pos, const int64_t offset_in_fragment)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t	bucketized_hash_join_idx (int64_t hash_buff, int64_t const key, int64_t const min_key, int64_t const max_key, const int64_t translated_null_val, int64_t bucket_normalization)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t	hash_join_idx (int64_t hash_buff, const int64_t key, const int64_t min_key, const int64_t max_key)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t	bucketized_hash_join_idx_nullable (int64_t hash_buff, const int64_t key, const int64_t min_key, const int64_t max_key, const int64_t null_val, const int64_t bucket_normalization)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t	hash_join_idx_nullable (int64_t hash_buff, const int64_t key, const int64_t min_key, const int64_t max_key, const int64_t null_val)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t	bucketized_hash_join_idx_bitwise (int64_t hash_buff, const int64_t key, const int64_t min_key, const int64_t max_key, const int64_t null_val, const int64_t translated_val, const int64_t bucket_normalization)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t	hash_join_idx_bitwise (int64_t hash_buff, const int64_t key, const int64_t min_key, const int64_t max_key, const int64_t null_val, const int64_t translated_val)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t	hash_join_idx_sharded (int64_t hash_buff, const int64_t key, const int64_t min_key, const int64_t max_key, const uint32_t entry_count_per_shard, const uint32_t num_shards, const uint32_t device_count)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t	hash_join_idx_sharded_nullable (int64_t hash_buff, const int64_t key, const int64_t min_key, const int64_t max_key, const uint32_t entry_count_per_shard, const uint32_t num_shards, const uint32_t device_count, const int64_t null_val)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t	hash_join_idx_bitwise_sharded (int64_t hash_buff, const int64_t key, const int64_t min_key, const int64_t max_key, const uint32_t entry_count_per_shard, const uint32_t num_shards, const uint32_t device_count, const int64_t null_val, const int64_t translated_val)

Macro Definition Documentation

#define DEF_TRANSLATE_NULL_KEY

(

key_type

)

Value:

extern "C" RUNTIME_EXPORT NEVER_INLINE DEVICE int64_t translate_null_key_##key_type( \
      const key_type key, const key_type null_val, const int64_t translated_val) {     \
    if (key == null_val) {                                                             \
      return translated_val;                                                           \
    }                                                                                  \
    return key;                                                                        \
  }

Definition at line 413 of file GroupByRuntime.cpp.

Function Documentation

RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t bucketized_hash_join_idx	(	int64_t	hash_buff,
		int64_t const	key,
		int64_t const	min_key,
		int64_t const	max_key,
		const int64_t	translated_null_val,
		int64_t	bucket_normalization
	)

Definition at line 269 of file GroupByRuntime.cpp.

References get_bucketized_hash_slot(), and SUFFIX.

Referenced by bucketized_hash_join_idx_bitwise(), and bucketized_hash_join_idx_nullable().

                                                       {
  if (key >= min_key && key <= max_key) {
    return *SUFFIX(get_bucketized_hash_slot)(reinterpret_cast<int32_t*>(hash_buff),
                                             key,
                                             min_key / bucket_normalization,
                                             translated_null_val,
                                             bucket_normalization);
  }
  return -1;
}

Here is the call graph for this function:

Here is the caller graph for this function:

RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t bucketized_hash_join_idx_bitwise	(	int64_t	hash_buff,
		const int64_t	key,
		const int64_t	min_key,
		const int64_t	max_key,
		const int64_t	null_val,
		const int64_t	translated_val,
		const int64_t	bucket_normalization
	)

Definition at line 319 of file GroupByRuntime.cpp.

References bucketized_hash_join_idx().

                                                                     {
  return key != null_val
             ? bucketized_hash_join_idx(
                   hash_buff, key, min_key, max_key, translated_val, bucket_normalization)
             : bucketized_hash_join_idx(hash_buff,
                                        translated_val,
                                        min_key,
                                        max_key,
                                        translated_val,
                                        bucket_normalization);
}

Here is the call graph for this function:

RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t bucketized_hash_join_idx_nullable	(	int64_t	hash_buff,
		const int64_t	key,
		const int64_t	min_key,
		const int64_t	max_key,
		const int64_t	null_val,
		const int64_t	bucket_normalization
	)

Definition at line 297 of file GroupByRuntime.cpp.

References bucketized_hash_join_idx().

                                                                      {
  return key != null_val
             ? bucketized_hash_join_idx(
                   hash_buff, key, min_key, max_key, null_val, bucket_normalization)
             : -1;
}

Here is the call graph for this function:

RUNTIME_EXPORT NEVER_INLINE DEVICE bool dynamic_watchdog

(

)

Definition at line 115 of file cuda_mapd_rt.cu.

References dw_abort, dw_cycle_budget, DW_DEADLINE, dw_sm_cycle_start, dynamic_watchdog_init(), get_smid(), logger::INFO, LOG, and read_cycle_counter().

Referenced by anonymous_namespace{ResultSetReduction.cpp}::check_watchdog(), check_watchdog_rt(), anonymous_namespace{ResultSetReduction.cpp}::check_watchdog_with_seed(), get_group_value_columnar_slot_with_watchdog(), get_group_value_columnar_with_watchdog(), and get_group_value_with_watchdog().

                                              {
  // 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;
}

Here is the call graph for this function:

Here is the caller graph for this function:

RUNTIME_EXPORT ALWAYS_INLINE DEVICE uint32_t get_columnar_group_bin_offset	(	int64_t *	key_base_ptr,
		const int64_t	key,
		const int64_t	min_key,
		const int64_t	bucket
	)

Definition at line 228 of file GroupByRuntime.cpp.

References EMPTY_KEY_64.

                                                    {
  int64_t off = key - min_key;
  if (bucket) {
    off /= bucket;
  }
  if (key_base_ptr[off] == EMPTY_KEY_64) {
    key_base_ptr[off] = key;
  }
  return off;
}

RUNTIME_EXPORT ALWAYS_INLINE DEVICE int32_t get_columnar_scan_output_offset	(	int64_t *	output_buffer,
		const uint32_t	output_buffer_entry_count,
		const uint32_t	pos,
		const int64_t	offset_in_fragment
	)

Definition at line 257 of file GroupByRuntime.cpp.

                                                                  {
  if (pos < output_buffer_entry_count) {
    output_buffer[pos] = offset_in_fragment;
    return pos;
  }
  return -1;
}

RUNTIME_EXPORT NEVER_INLINE DEVICE int64_t* get_group_value	(	int64_t *	groups_buffer,
		const uint32_t	groups_buffer_entry_count,
		const int64_t *	key,
		const uint32_t	key_count,
		const uint32_t	key_width,
		const uint32_t	row_size_quad
	)

Definition at line 25 of file GroupByRuntime.cpp.

References get_matching_group_value(), and key_hash().

Referenced by ResultSetStorage::moveOneEntryToBuffer().

                                  {
  uint32_t h = key_hash(key, key_count, key_width) % groups_buffer_entry_count;
  int64_t* matching_group = get_matching_group_value(
      groups_buffer, h, key, key_count, key_width, row_size_quad);
  if (matching_group) {
    return matching_group;
  }
  uint32_t h_probe = (h + 1) % groups_buffer_entry_count;
  while (h_probe != h) {
    matching_group = get_matching_group_value(
        groups_buffer, h_probe, key, key_count, key_width, row_size_quad);
    if (matching_group) {
      return matching_group;
    }
    h_probe = (h_probe + 1) % groups_buffer_entry_count;
  }
  return NULL;
}

Here is the call graph for this function:

Here is the caller graph for this function:

RUNTIME_EXPORT NEVER_INLINE DEVICE int64_t* get_group_value_columnar	(	int64_t *	groups_buffer,
		const uint32_t	groups_buffer_entry_count,
		const int64_t *	key,
		const uint32_t	key_qw_count
	)

Definition at line 139 of file GroupByRuntime.cpp.

References get_matching_group_value_columnar(), and key_hash().

Referenced by ResultSetStorage::moveOneEntryToBuffer().

                                 {
  uint32_t h = key_hash(key, key_qw_count, sizeof(int64_t)) % groups_buffer_entry_count;
  int64_t* matching_group = get_matching_group_value_columnar(
      groups_buffer, h, key, key_qw_count, groups_buffer_entry_count);
  if (matching_group) {
    return matching_group;
  }
  uint32_t h_probe = (h + 1) % groups_buffer_entry_count;
  while (h_probe != h) {
    matching_group = get_matching_group_value_columnar(
        groups_buffer, h_probe, key, key_qw_count, groups_buffer_entry_count);
    if (matching_group) {
      return matching_group;
    }
    h_probe = (h_probe + 1) % groups_buffer_entry_count;
  }
  return NULL;
}

Here is the call graph for this function:

Here is the caller graph for this function:

RUNTIME_EXPORT NEVER_INLINE DEVICE int32_t get_group_value_columnar_slot	(	int64_t *	groups_buffer,
		const uint32_t	groups_buffer_entry_count,
		const int64_t *	key,
		const uint32_t	key_count,
		const uint32_t	key_width
	)

Definition at line 85 of file GroupByRuntime.cpp.

References get_matching_group_value_columnar_slot(), and key_hash().

                                                        {
  uint32_t h = key_hash(key, key_count, key_width) % groups_buffer_entry_count;
  int32_t matching_slot = get_matching_group_value_columnar_slot(
      groups_buffer, groups_buffer_entry_count, h, key, key_count, key_width);
  if (matching_slot != -1) {
    return h;
  }
  uint32_t h_probe = (h + 1) % groups_buffer_entry_count;
  while (h_probe != h) {
    matching_slot = get_matching_group_value_columnar_slot(
        groups_buffer, groups_buffer_entry_count, h_probe, key, key_count, key_width);
    if (matching_slot != -1) {
      return h_probe;
    }
    h_probe = (h_probe + 1) % groups_buffer_entry_count;
  }
  return -1;
}

Here is the call graph for this function:

RUNTIME_EXPORT NEVER_INLINE DEVICE int32_t get_group_value_columnar_slot_with_watchdog	(	int64_t *	groups_buffer,
		const uint32_t	groups_buffer_entry_count,
		const int64_t *	key,
		const uint32_t	key_count,
		const uint32_t	key_width
	)

Definition at line 109 of file GroupByRuntime.cpp.

References dynamic_watchdog(), get_matching_group_value_columnar_slot(), and key_hash().

                                                                      {
  uint32_t h = key_hash(key, key_count, key_width) % groups_buffer_entry_count;
  int32_t matching_slot = get_matching_group_value_columnar_slot(
      groups_buffer, groups_buffer_entry_count, h, key, key_count, key_width);
  if (matching_slot != -1) {
    return h;
  }
  uint32_t watchdog_countdown = 100;
  uint32_t h_probe = (h + 1) % groups_buffer_entry_count;
  while (h_probe != h) {
    matching_slot = get_matching_group_value_columnar_slot(
        groups_buffer, groups_buffer_entry_count, h_probe, key, key_count, key_width);
    if (matching_slot != -1) {
      return h_probe;
    }
    h_probe = (h_probe + 1) % groups_buffer_entry_count;
    if (--watchdog_countdown == 0) {
      if (dynamic_watchdog()) {
        return -1;
      }
      watchdog_countdown = 100;
    }
  }
  return -1;
}

Here is the call graph for this function:

RUNTIME_EXPORT NEVER_INLINE DEVICE int64_t* get_group_value_columnar_with_watchdog	(	int64_t *	groups_buffer,
		const uint32_t	groups_buffer_entry_count,
		const int64_t *	key,
		const uint32_t	key_qw_count
	)

Definition at line 163 of file GroupByRuntime.cpp.

References dynamic_watchdog(), get_matching_group_value_columnar(), and key_hash().

                                                                    {
  uint32_t h = key_hash(key, key_qw_count, sizeof(int64_t)) % groups_buffer_entry_count;
  int64_t* matching_group = get_matching_group_value_columnar(
      groups_buffer, h, key, key_qw_count, groups_buffer_entry_count);
  if (matching_group) {
    return matching_group;
  }
  uint32_t watchdog_countdown = 100;
  uint32_t h_probe = (h + 1) % groups_buffer_entry_count;
  while (h_probe != h) {
    matching_group = get_matching_group_value_columnar(
        groups_buffer, h_probe, key, key_qw_count, groups_buffer_entry_count);
    if (matching_group) {
      return matching_group;
    }
    h_probe = (h_probe + 1) % groups_buffer_entry_count;
    if (--watchdog_countdown == 0) {
      if (dynamic_watchdog()) {
        return NULL;
      }
      watchdog_countdown = 100;
    }
  }
  return NULL;
}

Here is the call graph for this function:

RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t* get_group_value_fast	(	int64_t *	groups_buffer,
		const int64_t	key,
		const int64_t	min_key,
		const int64_t	bucket,
		const uint32_t	row_size_quad
	)

Definition at line 192 of file GroupByRuntime.cpp.

References EMPTY_KEY_64.

                                  {
  int64_t key_diff = key - min_key;
  if (bucket) {
    key_diff /= bucket;
  }
  int64_t off = key_diff * row_size_quad;
  if (groups_buffer[off] == EMPTY_KEY_64) {
    groups_buffer[off] = key;
  }
  return groups_buffer + off + 1;
}

RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t* get_group_value_fast_with_original_key	(	int64_t *	groups_buffer,
		const int64_t	key,
		const int64_t	orig_key,
		const int64_t	min_key,
		const int64_t	bucket,
		const uint32_t	row_size_quad
	)

Definition at line 210 of file GroupByRuntime.cpp.

References EMPTY_KEY_64.

                                                                     {
  int64_t key_diff = key - min_key;
  if (bucket) {
    key_diff /= bucket;
  }
  int64_t off = key_diff * row_size_quad;
  if (groups_buffer[off] == EMPTY_KEY_64) {
    groups_buffer[off] = orig_key;
  }
  return groups_buffer + off + 1;
}

RUNTIME_EXPORT NEVER_INLINE DEVICE int64_t* get_group_value_with_watchdog	(	int64_t *	groups_buffer,
		const uint32_t	groups_buffer_entry_count,
		const int64_t *	key,
		const uint32_t	key_count,
		const uint32_t	key_width,
		const uint32_t	row_size_quad
	)

Definition at line 52 of file GroupByRuntime.cpp.

References dynamic_watchdog(), get_matching_group_value(), and key_hash().

                                  {
  uint32_t h = key_hash(key, key_count, key_width) % groups_buffer_entry_count;
  int64_t* matching_group = get_matching_group_value(
      groups_buffer, h, key, key_count, key_width, row_size_quad);
  if (matching_group) {
    return matching_group;
  }
  uint32_t watchdog_countdown = 100;
  uint32_t h_probe = (h + 1) % groups_buffer_entry_count;
  while (h_probe != h) {
    matching_group = get_matching_group_value(
        groups_buffer, h_probe, key, key_count, key_width, row_size_quad);
    if (matching_group) {
      return matching_group;
    }
    h_probe = (h_probe + 1) % groups_buffer_entry_count;
    if (--watchdog_countdown == 0) {
      if (dynamic_watchdog()) {
        return NULL;
      }
      watchdog_countdown = 100;
    }
  }
  return NULL;
}

Here is the call graph for this function:

RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t* get_scan_output_slot	(	int64_t *	output_buffer,
		const uint32_t	output_buffer_entry_count,
		const uint32_t	pos,
		const int64_t	offset_in_fragment,
		const uint32_t	row_size_quad
	)

Definition at line 242 of file GroupByRuntime.cpp.

Referenced by SqliteMemDatabase::runSelect().

                                  {
  uint64_t off = static_cast<uint64_t>(pos) * static_cast<uint64_t>(row_size_quad);
  if (pos < output_buffer_entry_count) {
    output_buffer[off] = offset_in_fragment;
    return output_buffer + off + 1;
  }
  return NULL;
}

Here is the caller graph for this function:

RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t hash_join_idx	(	int64_t	hash_buff,
		const int64_t	key,
		const int64_t	min_key,
		const int64_t	max_key
	)

Definition at line 286 of file GroupByRuntime.cpp.

References get_hash_slot(), and SUFFIX.

Referenced by hash_join_idx_bitwise(), and hash_join_idx_nullable().

                                     {
  if (key >= min_key && key <= max_key) {
    return *SUFFIX(get_hash_slot)(reinterpret_cast<int32_t*>(hash_buff), key, min_key);
  }
  return -1;
}

Here is the call graph for this function:

Here is the caller graph for this function:

RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t hash_join_idx_bitwise	(	int64_t	hash_buff,
		const int64_t	key,
		const int64_t	min_key,
		const int64_t	max_key,
		const int64_t	null_val,
		const int64_t	translated_val
	)

Definition at line 338 of file GroupByRuntime.cpp.

References hash_join_idx().

                                                    {
  return key != null_val
             ? hash_join_idx(hash_buff, key, min_key, max_key)
             : hash_join_idx(hash_buff, translated_val, min_key, translated_val);
}

Here is the call graph for this function:

RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t hash_join_idx_bitwise_sharded	(	int64_t	hash_buff,
		const int64_t	key,
		const int64_t	min_key,
		const int64_t	max_key,
		const uint32_t	entry_count_per_shard,
		const uint32_t	num_shards,
		const uint32_t	device_count,
		const int64_t	null_val,
		const int64_t	translated_val
	)

Definition at line 388 of file GroupByRuntime.cpp.

References hash_join_idx_sharded().

                                                            {
  return key != null_val ? hash_join_idx_sharded(hash_buff,
                                                 key,
                                                 min_key,
                                                 max_key,
                                                 entry_count_per_shard,
                                                 num_shards,
                                                 device_count)
                         : hash_join_idx_sharded(hash_buff,
                                                 translated_val,
                                                 min_key,
                                                 translated_val,
                                                 entry_count_per_shard,
                                                 num_shards,
                                                 device_count);
}

Here is the call graph for this function:

RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t hash_join_idx_nullable	(	int64_t	hash_buff,
		const int64_t	key,
		const int64_t	min_key,
		const int64_t	max_key,
		const int64_t	null_val
	)

Definition at line 310 of file GroupByRuntime.cpp.

References hash_join_idx().

                                               {
  return key != null_val ? hash_join_idx(hash_buff, key, min_key, max_key) : -1;
}

Here is the call graph for this function:

RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t hash_join_idx_sharded	(	int64_t	hash_buff,
		const int64_t	key,
		const int64_t	min_key,
		const int64_t	max_key,
		const uint32_t	entry_count_per_shard,
		const uint32_t	num_shards,
		const uint32_t	device_count
	)

Definition at line 350 of file GroupByRuntime.cpp.

References get_hash_slot_sharded(), and SUFFIX.

Referenced by hash_join_idx_bitwise_sharded(), and hash_join_idx_sharded_nullable().

                                                   {
  if (key >= min_key && key <= max_key) {
    return *SUFFIX(get_hash_slot_sharded)(reinterpret_cast<int32_t*>(hash_buff),
                                          key,
                                          min_key,
                                          entry_count_per_shard,
                                          num_shards,
                                          device_count);
  }
  return -1;
}

Here is the call graph for this function:

Here is the caller graph for this function:

RUNTIME_EXPORT ALWAYS_INLINE DEVICE int64_t hash_join_idx_sharded_nullable	(	int64_t	hash_buff,
		const int64_t	key,
		const int64_t	min_key,
		const int64_t	max_key,
		const uint32_t	entry_count_per_shard,
		const uint32_t	num_shards,
		const uint32_t	device_count,
		const int64_t	null_val
	)

Definition at line 369 of file GroupByRuntime.cpp.

References hash_join_idx_sharded().

                                                       {
  return key != null_val ? hash_join_idx_sharded(hash_buff,
                                                 key,
                                                 min_key,
                                                 max_key,
                                                 entry_count_per_shard,
                                                 num_shards,
                                                 device_count)
                         : -1;
}

Here is the call graph for this function:

RUNTIME_EXPORT ALWAYS_INLINE DEVICE uint32_t key_hash	(	const int64_t *	key,
		const uint32_t	key_count,
		const uint32_t	key_byte_width
	)

Definition at line 21 of file GroupByRuntime.cpp.

References MurmurHash3().

Referenced by get_group_value(), get_group_value_columnar(), anonymous_namespace{ResultSetReduction.cpp}::get_group_value_columnar_reduction(), get_group_value_columnar_slot(), get_group_value_columnar_slot_with_watchdog(), get_group_value_columnar_with_watchdog(), result_set::get_group_value_reduction(), and get_group_value_with_watchdog().

                                                                                      {
  return MurmurHash3(key, key_byte_width * key_count, 0);
}

Here is the call graph for this function:

Here is the caller graph for this function: