17 #ifndef QUERYENGINE_RUNTIMEFUNCTIONS_H
18 #define QUERYENGINE_RUNTIMEFUNCTIONS_H
26 #include <type_traits>
50 const int32_t skip_val);
54 const int64_t skip_val);
58 const int32_t skip_val,
63 const int64_t skip_val,
68 const int64_t skip_val);
72 const int64_t skip_val);
76 const float skip_val);
80 const double skip_val);
89 const double skip_val,
94 const double skip_val);
98 const double skip_val);
126 const int32_t skip_val);
129 const int16_t skip_val);
132 const int8_t skip_val);
136 const int32_t skip_val);
139 const int16_t skip_val);
142 const int8_t skip_val);
146 const float skip_val);
150 const float skip_val);
154 const int64_t min_val);
156 #define EMPTY_KEY_64 std::numeric_limits<int64_t>::max()
157 #define EMPTY_KEY_32 std::numeric_limits<int32_t>::max()
158 #define EMPTY_KEY_16 std::numeric_limits<int16_t>::max()
159 #define EMPTY_KEY_8 std::numeric_limits<int8_t>::max()
162 const uint32_t key_qw_count,
163 const uint32_t key_byte_width);
166 int64_t* groups_buffer,
167 const uint32_t groups_buffer_entry_count,
169 const uint32_t key_count,
170 const uint32_t key_width,
171 const uint32_t row_size_quad);
180 int64_t* groups_buffer,
181 const uint32_t groups_buffer_entry_count,
183 const uint32_t key_count,
184 const uint32_t key_width,
185 const uint32_t row_size_quad);
188 int64_t* groups_buffer,
189 const uint32_t groups_buffer_entry_count,
191 const uint32_t key_qw_count);
194 int64_t* groups_buffer,
195 const uint32_t groups_buffer_entry_count,
197 const uint32_t key_qw_count);
201 const int64_t min_key,
202 const int64_t bucket,
203 const uint32_t row_size_quad);
206 int64_t* groups_buffer,
208 const int64_t orig_key,
209 const int64_t min_key,
210 const int64_t bucket,
211 const uint32_t row_size_quad);
215 const int64_t min_key,
216 const int64_t bucket);
219 int64_t* groups_buffer,
222 const uint32_t key_qw_count,
223 const uint32_t row_size_quad);
226 int64_t* groups_buffer,
227 const uint32_t hashed_index,
228 const uint32_t row_size_quad);
233 const int64_t min_key,
234 const int64_t translated_null_val,
235 const int64_t bucket_normalization = 1);
239 const int64_t min_key);
244 const int64_t min_key,
245 const uint32_t entry_count_per_shard,
246 const uint32_t num_shards,
247 const uint32_t device_count);
252 const int64_t min_key,
253 const int64_t translated_null_val,
254 const uint32_t entry_count_per_shard,
255 const uint32_t num_shards,
256 const uint32_t device_count,
257 const int64_t bucket_normalization);
262 const int64_t min_key,
263 const uint32_t entry_count_per_shard,
264 const uint32_t shard,
265 const uint32_t num_shards,
266 const uint32_t device_count);
271 const int64_t min_key,
272 const int64_t translated_null_val,
273 const uint32_t entry_count_per_shard,
274 const uint32_t shard,
275 const uint32_t num_shards,
276 const uint32_t device_count,
277 const int64_t bucket_normalization);
281 const int32_t invalid_slot_val);
286 const int32_t invalid_slot_val);
289 const uint32_t bitmap_bytes,
290 const uint8_t* key_bytes,
291 const uint32_t key_len);
298 const int32_t byte_width,
303 const int32_t byte_width,
307 const int8_t* byte_stream,
311 const int8_t* byte_stream,
316 const int32_t byte_width,
317 const int32_t null_val,
318 const int64_t ret_null_val,
323 const int32_t null_val,
324 const int64_t ret_null_val);
326 template <
typename T =
int64_t>
328 static_assert(std::is_same<T, int64_t>::value,
329 "Unsupported template parameter other than int64_t for now");
338 #endif // QUERYENGINE_RUNTIMEFUNCTIONS_H
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 void agg_min_int8(int8_t *agg, const int8_t val)
RUNTIME_EXPORT ALWAYS_INLINE int32_t agg_sum_if_int32_skip_val(int32_t *agg, const int32_t val, const int32_t skip_val, const int8_t cond)
RUNTIME_EXPORT void agg_max_int32(int32_t *agg, const int32_t val)
RUNTIME_EXPORT ALWAYS_INLINE int64_t * get_matching_group_value_perfect_hash(int64_t *groups_buffer, const uint32_t hashed_index, const int64_t *key, const uint32_t key_count, const uint32_t row_size_quad)
RUNTIME_EXPORT void agg_min_int16(int16_t *agg, const int16_t val)
RUNTIME_EXPORT void agg_max_int16(int16_t *agg, const int16_t val)
RUNTIME_EXPORT void agg_min_int32_skip_val(int32_t *agg, const int32_t val, const int32_t skip_val)
DEVICE NEVER_INLINE int64_t SUFFIX() fixed_width_int_decode_noinline(const int8_t *byte_stream, const int32_t byte_width, const int64_t pos)
RUNTIME_EXPORT ALWAYS_INLINE void agg_max_double(int64_t *agg, const double val)
RUNTIME_EXPORT ALWAYS_INLINE int64_t agg_sum_if(int64_t *agg, const int64_t val, const int8_t cond)
RUNTIME_EXPORT ALWAYS_INLINE void agg_max(int64_t *agg, const int64_t val)
RUNTIME_EXPORT ALWAYS_INLINE void agg_min_float(int32_t *agg, const float val)
RUNTIME_EXPORT ALWAYS_INLINE void agg_count_distinct_bitmap(int64_t *agg, const int64_t val, const int64_t min_val)
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 void agg_min_int8_skip_val(int8_t *agg, const int8_t val, const int8_t skip_val)
ALWAYS_INLINE DEVICE int SUFFIX() fill_hashtable_for_semi_join(size_t idx, int32_t *entry_ptr, const int32_t invalid_slot_val)
RUNTIME_EXPORT void agg_max_int16_skip_val(int16_t *agg, const int16_t val, const int16_t skip_val)
RUNTIME_EXPORT void agg_sum_if_float(int32_t *agg, const float val, const int8_t cond)
RUNTIME_EXPORT ALWAYS_INLINE int64_t * get_matching_group_value_perfect_hash_keyless(int64_t *groups_buffer, const uint32_t hashed_index, const uint32_t row_size_quad)
RUNTIME_EXPORT void agg_sum_float_skip_val(int32_t *agg, const float val, const float skip_val)
__device__ bool check_interrupt()
RUNTIME_EXPORT void agg_sum_double_skip_val(int64_t *agg, const double val, const double skip_val)
ALWAYS_INLINE DEVICE int32_t *SUFFIX() get_hash_slot(int32_t *buff, const int64_t key, const int64_t min_key)
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 translated_null_val, const int64_t bucket_normalization)
__device__ void linear_probabilistic_count(uint8_t *bitmap, const uint32_t bitmap_bytes, const uint8_t *key_bytes, const uint32_t key_len)
RUNTIME_EXPORT ALWAYS_INLINE void agg_min_double(int64_t *agg, const double val)
RUNTIME_EXPORT void agg_min_int16_skip_val(int16_t *agg, const int16_t val, const int16_t skip_val)
RUNTIME_EXPORT void agg_max_double_skip_val(int64_t *agg, const double val, const double skip_val)
DEVICE NEVER_INLINE int64_t SUFFIX() fixed_width_unsigned_decode_noinline(const int8_t *byte_stream, const int32_t byte_width, const int64_t pos)
RUNTIME_EXPORT ALWAYS_INLINE void agg_min(int64_t *agg, const int64_t val)
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 void agg_sum_double(int64_t *agg, const double val)
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 void agg_min_skip_val(int64_t *agg, const int64_t val, const int64_t skip_val)
RUNTIME_EXPORT ALWAYS_INLINE int64_t agg_sum_if_skip_val(int64_t *agg, const int64_t val, const int64_t skip_val, const int8_t cond)
RUNTIME_EXPORT ALWAYS_INLINE int32_t agg_sum_int32(int32_t *agg, const int32_t val)
DEVICE NEVER_INLINE int64_t SUFFIX() fixed_width_small_date_decode_noinline(const int8_t *byte_stream, const int32_t byte_width, const int32_t null_val, const int64_t ret_null_val, const int64_t pos)
ALWAYS_INLINE DEVICE int32_t *SUFFIX() get_hash_slot_sharded_opt(int32_t *buff, const int64_t key, const int64_t min_key, const uint32_t entry_count_per_shard, const uint32_t shard, const uint32_t num_shards, const uint32_t device_count)
RUNTIME_EXPORT void agg_sum_if_float_skip_val(int32_t *agg, const float val, const float skip_val, const int8_t cond)
RUNTIME_EXPORT void agg_max_int8(int8_t *agg, const int8_t val)
RUNTIME_EXPORT ALWAYS_INLINE void agg_max_float(int32_t *agg, const float val)
RUNTIME_EXPORT void agg_max_int8_skip_val(int8_t *agg, const int8_t val, const int8_t skip_val)
ALWAYS_INLINE DEVICE int32_t *SUFFIX() get_hash_slot_sharded(int32_t *buff, const int64_t key, const int64_t min_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 * 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 void agg_max_int32_skip_val(int32_t *agg, const int32_t val, const int32_t skip_val)
RUNTIME_EXPORT void agg_min_double_skip_val(int64_t *agg, const double val, const double skip_val)
RUNTIME_EXPORT ALWAYS_INLINE DEVICE uint32_t key_hash(const int64_t *key, const uint32_t key_count, const uint32_t key_byte_width)
ALWAYS_INLINE DEVICE int SUFFIX() fill_one_to_one_hashtable(size_t idx, int32_t *entry_ptr, const int32_t invalid_slot_val)
ALWAYS_INLINE DEVICE int32_t *SUFFIX() get_bucketized_hash_slot_sharded_opt(int32_t *buff, const int64_t key, const int64_t min_key, const int64_t translated_null_val, const uint32_t entry_count_per_shard, const uint32_t shard, const uint32_t num_shards, const uint32_t device_count, const int64_t bucket_normalization)
RUNTIME_EXPORT void agg_max_float_skip_val(int32_t *agg, const float val, const float skip_val)
RUNTIME_EXPORT ALWAYS_INLINE int32_t agg_sum_if_int32(int32_t *agg, const int32_t val, const int8_t cond)
__device__ T get_empty_key()
RUNTIME_EXPORT ALWAYS_INLINE int64_t agg_sum_skip_val(int64_t *agg, const int64_t val, const int64_t skip_val)
DEVICE NEVER_INLINE float SUFFIX() fixed_width_float_decode_noinline(const int8_t *byte_stream, const int64_t pos)
DEVICE NEVER_INLINE double SUFFIX() fixed_width_double_decode_noinline(const int8_t *byte_stream, const int64_t pos)
RUNTIME_EXPORT ALWAYS_INLINE int32_t agg_sum_int32_skip_val(int32_t *agg, const int32_t val, const int32_t skip_val)
RUNTIME_EXPORT ALWAYS_INLINE void agg_sum_float(int32_t *agg, const float val)
RUNTIME_EXPORT bool check_interrupt_init(unsigned command)
RUNTIME_EXPORT void agg_min_float_skip_val(int32_t *agg, const float val, const float skip_val)
RUNTIME_EXPORT ALWAYS_INLINE int64_t agg_sum(int64_t *agg, const int64_t val)
ALWAYS_INLINE DEVICE int32_t *SUFFIX() get_bucketized_hash_slot_sharded(int32_t *buff, const int64_t key, const int64_t min_key, const int64_t translated_null_val, const uint32_t entry_count_per_shard, const uint32_t num_shards, const uint32_t device_count, const int64_t bucket_normalization)
RUNTIME_EXPORT void agg_max_skip_val(int64_t *agg, const int64_t val, const int64_t skip_val)
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_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 void agg_sum_if_double_skip_val(int64_t *agg, const double val, const double skip_val, const int8_t cond)
RUNTIME_EXPORT void agg_min_int32(int32_t *agg, const int32_t val)
RUNTIME_EXPORT ALWAYS_INLINE void agg_sum_if_double(int64_t *agg, const double val, const int8_t cond)
DEVICE NEVER_INLINE int64_t SUFFIX() fixed_width_date_encode_noinline(const int64_t cur_col_val, const int32_t ret_null_val, const int64_t null_val)