OmniSciDB  72c90bc290
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
RuntimeFunctions.h
Go to the documentation of this file.
1 /*
2  * Copyright 2022 HEAVY.AI, Inc.
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16 
17 #ifndef QUERYENGINE_RUNTIMEFUNCTIONS_H
18 #define QUERYENGINE_RUNTIMEFUNCTIONS_H
19 
20 #include "Shared/funcannotations.h"
21 
22 #include <cassert>
23 #include <cstdint>
24 #include <ctime>
25 #include <limits>
26 #include <type_traits>
27 
28 extern "C" RUNTIME_EXPORT int64_t agg_sum(int64_t* agg, const int64_t val);
29 
30 extern "C" RUNTIME_EXPORT int64_t agg_sum_if(int64_t* agg,
31  const int64_t val,
32  const int8_t cond);
33 
34 extern "C" RUNTIME_EXPORT void agg_max(int64_t* agg, const int64_t val);
35 
36 extern "C" RUNTIME_EXPORT void agg_min(int64_t* agg, const int64_t val);
37 
38 extern "C" RUNTIME_EXPORT void agg_sum_double(int64_t* agg, const double val);
39 
40 extern "C" RUNTIME_EXPORT void agg_sum_if_double(int64_t* agg,
41  const double val,
42  const int8_t cond);
43 
44 extern "C" RUNTIME_EXPORT void agg_max_double(int64_t* agg, const double val);
45 
46 extern "C" RUNTIME_EXPORT void agg_min_double(int64_t* agg, const double val);
47 
48 extern "C" RUNTIME_EXPORT int32_t agg_sum_int32_skip_val(int32_t* agg,
49  const int32_t val,
50  const int32_t skip_val);
51 
52 extern "C" RUNTIME_EXPORT int64_t agg_sum_skip_val(int64_t* agg,
53  const int64_t val,
54  const int64_t skip_val);
55 
56 extern "C" RUNTIME_EXPORT int32_t agg_sum_if_int32_skip_val(int32_t* agg,
57  const int32_t val,
58  const int32_t skip_val,
59  const int8_t cond);
60 
61 extern "C" RUNTIME_EXPORT int64_t agg_sum_if_skip_val(int64_t* agg,
62  const int64_t val,
63  const int64_t skip_val,
64  const int8_t cond);
65 
66 extern "C" RUNTIME_EXPORT void agg_max_skip_val(int64_t* agg,
67  const int64_t val,
68  const int64_t skip_val);
69 
70 extern "C" RUNTIME_EXPORT void agg_min_skip_val(int64_t* agg,
71  const int64_t val,
72  const int64_t skip_val);
73 
74 extern "C" RUNTIME_EXPORT void agg_sum_float_skip_val(int32_t* agg,
75  const float val,
76  const float skip_val);
77 
78 extern "C" RUNTIME_EXPORT void agg_sum_double_skip_val(int64_t* agg,
79  const double val,
80  const double skip_val);
81 
82 extern "C" RUNTIME_EXPORT void agg_sum_if_float_skip_val(int32_t* agg,
83  const float val,
84  const float skip_val,
85  const int8_t cond);
86 
87 extern "C" RUNTIME_EXPORT void agg_sum_if_double_skip_val(int64_t* agg,
88  const double val,
89  const double skip_val,
90  const int8_t cond);
91 
92 extern "C" RUNTIME_EXPORT void agg_max_double_skip_val(int64_t* agg,
93  const double val,
94  const double skip_val);
95 
96 extern "C" RUNTIME_EXPORT void agg_min_double_skip_val(int64_t* agg,
97  const double val,
98  const double skip_val);
99 
100 extern "C" RUNTIME_EXPORT int32_t agg_sum_int32(int32_t* agg, const int32_t val);
101 
102 extern "C" RUNTIME_EXPORT int32_t agg_sum_if_int32(int32_t* agg,
103  const int32_t val,
104  const int8_t cond);
105 
106 extern "C" RUNTIME_EXPORT void agg_max_int32(int32_t* agg, const int32_t val);
107 extern "C" RUNTIME_EXPORT void agg_max_int16(int16_t* agg, const int16_t val);
108 extern "C" RUNTIME_EXPORT void agg_max_int8(int8_t* agg, const int8_t val);
109 
110 extern "C" RUNTIME_EXPORT void agg_min_int32(int32_t* agg, const int32_t val);
111 extern "C" RUNTIME_EXPORT void agg_min_int16(int16_t* agg, const int16_t val);
112 extern "C" RUNTIME_EXPORT void agg_min_int8(int8_t* agg, const int8_t val);
113 
114 extern "C" RUNTIME_EXPORT void agg_sum_float(int32_t* agg, const float val);
115 
116 extern "C" RUNTIME_EXPORT void agg_sum_if_float(int32_t* agg,
117  const float val,
118  const int8_t cond);
119 
120 extern "C" RUNTIME_EXPORT void agg_max_float(int32_t* agg, const float val);
121 
122 extern "C" RUNTIME_EXPORT void agg_min_float(int32_t* agg, const float val);
123 
124 extern "C" RUNTIME_EXPORT void agg_max_int32_skip_val(int32_t* agg,
125  const int32_t val,
126  const int32_t skip_val);
127 extern "C" RUNTIME_EXPORT void agg_max_int16_skip_val(int16_t* agg,
128  const int16_t val,
129  const int16_t skip_val);
130 extern "C" RUNTIME_EXPORT void agg_max_int8_skip_val(int8_t* agg,
131  const int8_t val,
132  const int8_t skip_val);
133 
134 extern "C" RUNTIME_EXPORT void agg_min_int32_skip_val(int32_t* agg,
135  const int32_t val,
136  const int32_t skip_val);
137 extern "C" RUNTIME_EXPORT void agg_min_int16_skip_val(int16_t* agg,
138  const int16_t val,
139  const int16_t skip_val);
140 extern "C" RUNTIME_EXPORT void agg_min_int8_skip_val(int8_t* agg,
141  const int8_t val,
142  const int8_t skip_val);
143 
144 extern "C" RUNTIME_EXPORT void agg_max_float_skip_val(int32_t* agg,
145  const float val,
146  const float skip_val);
147 
148 extern "C" RUNTIME_EXPORT void agg_min_float_skip_val(int32_t* agg,
149  const float val,
150  const float skip_val);
151 
152 extern "C" RUNTIME_EXPORT void agg_count_distinct_bitmap(int64_t* agg,
153  const int64_t val,
154  const int64_t min_val,
155  const int64_t bucket_size);
156 
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()
161 
162 extern "C" RUNTIME_EXPORT uint32_t key_hash(const int64_t* key,
163  const uint32_t key_qw_count,
164  const uint32_t key_byte_width);
165 
166 extern "C" RUNTIME_EXPORT int64_t* get_group_value(
167  int64_t* groups_buffer,
168  const uint32_t groups_buffer_entry_count,
169  const int64_t* key,
170  const uint32_t key_count,
171  const uint32_t key_width,
172  const uint32_t row_size_quad);
173 
175 
176 extern "C" bool RUNTIME_EXPORT check_interrupt();
177 
178 extern "C" bool RUNTIME_EXPORT check_interrupt_init(unsigned command);
179 
181  int64_t* groups_buffer,
182  const uint32_t groups_buffer_entry_count,
183  const int64_t* key,
184  const uint32_t key_count,
185  const uint32_t key_width,
186  const uint32_t row_size_quad);
187 
188 extern "C" RUNTIME_EXPORT int64_t* get_group_value_columnar(
189  int64_t* groups_buffer,
190  const uint32_t groups_buffer_entry_count,
191  const int64_t* key,
192  const uint32_t key_qw_count);
193 
195  int64_t* groups_buffer,
196  const uint32_t groups_buffer_entry_count,
197  const int64_t* key,
198  const uint32_t key_qw_count);
199 
200 extern "C" RUNTIME_EXPORT int64_t* get_group_value_fast(int64_t* groups_buffer,
201  const int64_t key,
202  const int64_t min_key,
203  const int64_t bucket,
204  const uint32_t row_size_quad);
205 
207  int64_t* groups_buffer,
208  const int64_t key,
209  const int64_t orig_key,
210  const int64_t min_key,
211  const int64_t bucket,
212  const uint32_t row_size_quad);
213 
214 extern "C" RUNTIME_EXPORT uint32_t get_columnar_group_bin_offset(int64_t* key_base_ptr,
215  const int64_t key,
216  const int64_t min_key,
217  const int64_t bucket);
218 
220  int64_t* groups_buffer,
221  const uint32_t h,
222  const int64_t* key,
223  const uint32_t key_qw_count,
224  const uint32_t row_size_quad);
225 
227  int64_t* groups_buffer,
228  const uint32_t hashed_index,
229  const uint32_t row_size_quad);
230 
231 extern "C" RUNTIME_EXPORT int32_t* get_bucketized_hash_slot(
232  int32_t* buff,
233  const int64_t key,
234  const int64_t min_key,
235  const int64_t translated_null_val,
236  const int64_t bucket_normalization = 1);
237 
238 extern "C" RUNTIME_EXPORT int32_t* get_hash_slot_bitwise_eq(
239  int32_t* buff,
240  const int64_t key,
241  const int64_t min_key,
242  const int64_t translated_null_val);
243 
244 extern "C" RUNTIME_EXPORT int32_t* get_hash_slot(int32_t* buff,
245  const int64_t key,
246  const int64_t min_key);
247 
248 extern "C" RUNTIME_EXPORT int32_t* get_hash_slot_sharded(
249  int32_t* buff,
250  const int64_t 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);
255 
257  int32_t* buff,
258  const int64_t key,
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);
265 
266 extern "C" RUNTIME_EXPORT int32_t* get_hash_slot_sharded_opt(
267  int32_t* buff,
268  const int64_t key,
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);
274 
276  int32_t* buff,
277  const int64_t key,
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);
285 
286 extern "C" RUNTIME_EXPORT int fill_one_to_one_hashtable(size_t idx,
287  int32_t* entry_ptr,
288  const int32_t invalid_slot_val);
289 
291  size_t idx,
292  int32_t* entry_ptr,
293  const int32_t invalid_slot_val);
294 
295 extern "C" RUNTIME_EXPORT void linear_probabilistic_count(uint8_t* bitmap,
296  const uint32_t bitmap_bytes,
297  const uint8_t* key_bytes,
298  const uint32_t key_len);
299 
300 // Regular fixed_width_*_decode are only available from the JIT,
301 // we need to call them for lazy fetch columns -- create wrappers.
302 
303 extern "C" RUNTIME_EXPORT int64_t
304 fixed_width_int_decode_noinline(const int8_t* byte_stream,
305  const int32_t byte_width,
306  const int64_t pos);
307 
308 extern "C" RUNTIME_EXPORT int64_t
309 fixed_width_unsigned_decode_noinline(const int8_t* byte_stream,
310  const int32_t byte_width,
311  const int64_t pos);
312 
314  const int8_t* byte_stream,
315  const int64_t pos);
316 
318  const int8_t* byte_stream,
319  const int64_t pos);
320 
321 extern "C" RUNTIME_EXPORT int64_t
322 fixed_width_small_date_decode_noinline(const int8_t* byte_stream,
323  const int32_t byte_width,
324  const int32_t null_val,
325  const int64_t ret_null_val,
326  const int64_t pos);
327 
328 extern "C" DEVICE NEVER_INLINE int64_t
329  SUFFIX(fixed_width_date_encode_noinline)(const int64_t cur_col_val,
330  const int32_t null_val,
331  const int64_t ret_null_val);
332 
333 template <typename T = int64_t>
334 inline T get_empty_key() {
335  static_assert(std::is_same<T, int64_t>::value,
336  "Unsupported template parameter other than int64_t for now");
337  return EMPTY_KEY_64;
338 }
339 
340 template <>
341 inline int32_t get_empty_key() {
342  return EMPTY_KEY_32;
343 }
344 
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)
Definition: DecodersImpl.h:91
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)
#define EMPTY_KEY_64
RUNTIME_EXPORT ALWAYS_INLINE void agg_min_float(int32_t *agg, const float val)
#define SUFFIX(name)
RuntimeInterruptFlags
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)
Definition: JoinHashImpl.h:54
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)
#define DEVICE
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)
Definition: JoinHashImpl.h:82
__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)
Definition: JoinHashImpl.h:76
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)
Definition: JoinHashImpl.h:66
__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)
Definition: DecodersImpl.h:98
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)
#define RUNTIME_EXPORT
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)
Definition: DecodersImpl.h:149
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)
Definition: JoinHashImpl.h:140
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)
Definition: JoinHashImpl.h:108
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)
#define NEVER_INLINE
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)
Definition: JoinHashImpl.h:44
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)
Definition: JoinHashImpl.h:122
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)
Definition: DecodersImpl.h:121
DEVICE NEVER_INLINE double SUFFIX() fixed_width_double_decode_noinline(const int8_t *byte_stream, const int64_t pos)
Definition: DecodersImpl.h:134
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)
Definition: JoinHashImpl.h:90
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)
#define EMPTY_KEY_32
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)
Definition: DecodersImpl.h:173