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,
155 const int64_t bucket_size);
157 #define EMPTY_KEY_64 std::numeric_limits<int64_t>::max()
158 #define EMPTY_KEY_32 std::numeric_limits<int32_t>::max()
159 #define EMPTY_KEY_16 std::numeric_limits<int16_t>::max()
160 #define EMPTY_KEY_8 std::numeric_limits<int8_t>::max()
163 const uint32_t key_qw_count,
164 const uint32_t key_byte_width);
167 int64_t* groups_buffer,
168 const uint32_t groups_buffer_entry_count,
170 const uint32_t key_count,
171 const uint32_t key_width,
172 const uint32_t row_size_quad);
181 int64_t* groups_buffer,
182 const uint32_t groups_buffer_entry_count,
184 const uint32_t key_count,
185 const uint32_t key_width,
186 const uint32_t row_size_quad);
189 int64_t* groups_buffer,
190 const uint32_t groups_buffer_entry_count,
192 const uint32_t key_qw_count);
195 int64_t* groups_buffer,
196 const uint32_t groups_buffer_entry_count,
198 const uint32_t key_qw_count);
202 const int64_t min_key,
203 const int64_t bucket,
204 const uint32_t row_size_quad);
207 int64_t* groups_buffer,
209 const int64_t orig_key,
210 const int64_t min_key,
211 const int64_t bucket,
212 const uint32_t row_size_quad);
216 const int64_t min_key,
217 const int64_t bucket);
220 int64_t* groups_buffer,
223 const uint32_t key_qw_count,
224 const uint32_t row_size_quad);
227 int64_t* groups_buffer,
228 const uint32_t hashed_index,
229 const uint32_t row_size_quad);
234 const int64_t min_key,
235 const int64_t translated_null_val,
236 const int64_t bucket_normalization = 1);
241 const int64_t min_key,
242 const int64_t translated_null_val);
246 const int64_t min_key);
251 const int64_t min_key,
252 const uint32_t entry_count_per_shard,
253 const uint32_t num_shards,
254 const uint32_t device_count);
259 const int64_t min_key,
260 const int64_t translated_null_val,
261 const uint32_t entry_count_per_shard,
262 const uint32_t num_shards,
263 const uint32_t device_count,
264 const int64_t bucket_normalization);
269 const int64_t min_key,
270 const uint32_t entry_count_per_shard,
271 const uint32_t shard,
272 const uint32_t num_shards,
273 const uint32_t device_count);
278 const int64_t min_key,
279 const int64_t translated_null_val,
280 const uint32_t entry_count_per_shard,
281 const uint32_t shard,
282 const uint32_t num_shards,
283 const uint32_t device_count,
284 const int64_t bucket_normalization);
288 const int32_t invalid_slot_val);
293 const int32_t invalid_slot_val);
296 const uint32_t bitmap_bytes,
297 const uint8_t* key_bytes,
298 const uint32_t key_len);
305 const int32_t byte_width,
310 const int32_t byte_width,
314 const int8_t* byte_stream,
318 const int8_t* byte_stream,
323 const int32_t byte_width,
324 const int32_t null_val,
325 const int64_t ret_null_val,
330 const int32_t null_val,
331 const int64_t ret_null_val);
333 template <
typename T =
int64_t>
335 static_assert(std::is_same<T, int64_t>::value,
336 "Unsupported template parameter other than int64_t for now");
345 #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 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)
RUNTIME_EXPORT ALWAYS_INLINE void agg_count_distinct_bitmap(int64_t *agg, const int64_t val, const int64_t min_val, const int64_t bucket_size)
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)
ALWAYS_INLINE DEVICE int32_t *SUFFIX() get_hash_slot_bitwise_eq(int32_t *buff, const int64_t key, const int64_t min_key, const int64_t translated_null_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)