11 #if CUDA_VERSION < 10000
12 static_assert(
false,
"CUDA v10.0 or later is required.");
15 #if (defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 350)
16 static_assert(
false,
"CUDA Compute Capability of 3.5 or greater is required.");
27 extern "C" __device__ int32_t
pos_start_impl(
const int32_t* row_index_resume) {
28 return blockIdx.x * blockDim.x + threadIdx.x;
36 return blockDim.x * gridDim.x;
40 return threadIdx.x % warp_sz;
44 const int64_t* groups_buffer,
45 const int32_t groups_buffer_size) {
57 extern __shared__ int64_t shared_mem_buffer[];
58 return shared_mem_buffer;
66 extern "C" __device__
const int64_t*
init_shared_mem(
const int64_t* global_groups_buffer,
67 const int32_t groups_buffer_size) {
69 extern __shared__ int64_t shared_groups_buffer[];
73 const int32_t buffer_units = groups_buffer_size >> 3;
75 for (int32_t pos = threadIdx.x; pos < buffer_units; pos += blockDim.x) {
76 shared_groups_buffer[pos] = global_groups_buffer[pos];
79 return shared_groups_buffer;
82 #define init_group_by_buffer_gpu_impl init_group_by_buffer_gpu
86 #undef init_group_by_buffer_gpu_impl
97 __inline__ __device__ uint32_t
get_smid(
void) {
99 asm(
"mov.u32 %0, %%smid;" :
"=r"(ret));
127 __shared__
volatile int64_t dw_block_cycle_start;
128 __shared__
volatile bool
135 if (threadIdx.x == 0) {
136 dw_block_cycle_start = 0LL;
137 int64_t cycle_count =
static_cast<int64_t
>(clock64());
140 dw_block_cycle_start =
static_cast<int64_t
>(
143 static_cast<unsigned long long>(cycle_count)));
146 int64_t cycles = cycle_count - dw_block_cycle_start;
147 if ((smid ==
get_smid()) && (dw_block_cycle_start > 0LL) &&
150 dw_should_terminate =
true;
152 dw_should_terminate =
false;
156 return dw_should_terminate;
163 template <
typename T =
unsigned long long>
173 template <
typename T>
177 const uint32_t key_count,
178 const uint32_t row_size_quad) {
179 const T empty_key = get_empty_key<T>();
180 uint32_t off = h * row_size_quad;
181 auto row_ptr =
reinterpret_cast<T*
>(groups_buffer + off);
183 const T old = atomicCAS(row_ptr, empty_key, *key);
184 if (empty_key == old && key_count > 1) {
185 for (
size_t i = 1; i <= key_count - 1; ++i) {
186 atomicExch(row_ptr + i, key[i]);
191 while (atomicAdd(row_ptr + key_count - 1, 0) == empty_key) {
197 for (uint32_t i = 0; i < key_count; ++i) {
198 if (row_ptr[i] != key[i]) {
205 auto row_ptr_i8 =
reinterpret_cast<int8_t*
>(row_ptr + key_count);
214 const uint32_t key_count,
215 const uint32_t key_width,
216 const uint32_t row_size_quad) {
221 reinterpret_cast<const unsigned int*>(key),
227 reinterpret_cast<const unsigned long long*>(key),
235 template <
typename T>
237 const uint32_t entry_count,
240 const uint32_t key_count) {
241 const T empty_key = get_empty_key<T>();
243 atomicCAS(reinterpret_cast<T*>(groups_buffer + h), empty_key, *key);
245 if (old == empty_key) {
246 uint32_t offset = h + entry_count;
247 for (
size_t i = 1; i < key_count; ++i) {
248 *
reinterpret_cast<T*
>(groups_buffer + offset) = key[i];
249 offset += entry_count;
257 if (old != empty_key) {
259 for (uint32_t i = 0; i < key_count; ++i) {
260 if (*reinterpret_cast<T*>(groups_buffer + offset) != key[i]) {
263 offset += entry_count;
269 extern "C" __device__ int32_t
271 const uint32_t entry_count,
274 const uint32_t key_count,
275 const uint32_t key_width) {
282 reinterpret_cast<const unsigned int*>(key),
289 reinterpret_cast<const unsigned long long*>(key),
297 int64_t* groups_buffer,
300 const uint32_t key_qw_count,
301 const size_t entry_count) {
304 const uint64_t old = atomicCAS(
305 reinterpret_cast<unsigned long long*>(groups_buffer + off),
EMPTY_KEY_64, *key);
307 for (
size_t i = 0; i < key_qw_count; ++i) {
308 groups_buffer[off] = key[i];
311 return &groups_buffer[off];
316 for (
size_t i = 0; i < key_qw_count; ++i) {
317 if (groups_buffer[off] != key[i]) {
322 return &groups_buffer[off];
331 unsigned long long int* address_as_ull = (
unsigned long long int*)address;
332 unsigned long long int old = *address_as_ull, assumed;
336 old = atomicCAS(address_as_ull, assumed, max((
long long)val, (
long long)assumed));
337 }
while (assumed != old);
343 unsigned long long int* address_as_ull = (
unsigned long long int*)address;
344 unsigned long long int old = *address_as_ull, assumed;
348 old = atomicCAS(address_as_ull, assumed, min((
long long)val, (
long long)assumed));
349 }
while (assumed != old);
354 #if (defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 600)
355 __device__
double atomicAdd(
double* address,
double val) {
356 unsigned long long int* address_as_ull = (
unsigned long long int*)address;
357 unsigned long long int old = *address_as_ull, assumed;
361 old = atomicCAS(address_as_ull,
363 __double_as_longlong(val + __longlong_as_double(assumed)));
366 }
while (assumed != old);
368 return __longlong_as_double(old);
372 __device__
double atomicMax(
double* address,
double val) {
373 unsigned long long int* address_as_ull = (
unsigned long long int*)address;
374 unsigned long long int old = *address_as_ull, assumed;
378 old = atomicCAS(address_as_ull,
380 __double_as_longlong(max(val, __longlong_as_double(assumed))));
383 }
while (assumed != old);
385 return __longlong_as_double(old);
389 int* address_as_int = (
int*)address;
390 int old = *address_as_int, assumed;
395 address_as_int, assumed, __float_as_int(max(val, __int_as_float(assumed))));
398 }
while (assumed != old);
400 return __int_as_float(old);
403 __device__
double atomicMin(
double* address,
double val) {
404 unsigned long long int* address_as_ull = (
unsigned long long int*)address;
405 unsigned long long int old = *address_as_ull, assumed;
409 old = atomicCAS(address_as_ull,
411 __double_as_longlong(min(val, __longlong_as_double(assumed))));
412 }
while (assumed != old);
414 return __longlong_as_double(old);
417 __device__
double atomicMin(
float* address,
float val) {
418 int* address_as_ull = (
int*)address;
419 int old = *address_as_ull, assumed;
424 address_as_ull, assumed, __float_as_int(min(val, __int_as_float(assumed))));
425 }
while (assumed != old);
427 return __int_as_float(old);
431 return static_cast<uint64_t
>(atomicAdd(reinterpret_cast<uint32_t*>(agg), 1UL));
435 return atomicAdd(agg, 1UL);
447 return atomicAdd(reinterpret_cast<unsigned long long*>(agg), val);
451 return atomicAdd(agg, val);
455 atomicAdd(reinterpret_cast<float*>(agg), val);
459 atomicAdd(reinterpret_cast<double*>(agg), val);
471 atomicMax(reinterpret_cast<double*>(agg), val);
475 atomicMax(reinterpret_cast<float*>(agg), val);
486 #if CUDA_VERSION > 10000 && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
487 __device__
void atomicMax16(int16_t* agg,
const int16_t val) {
488 unsigned short int* address_as_us =
reinterpret_cast<unsigned short int*
>(agg);
489 unsigned short int old = *address_as_us, assumed;
493 old = atomicCAS(address_as_us,
495 static_cast<unsigned short>(max(static_cast<short int>(val),
496 static_cast<short int>(assumed))));
497 }
while (assumed != old);
502 unsigned int* base_address_u32 =
503 reinterpret_cast<unsigned int*
>(
reinterpret_cast<size_t>(agg) & ~0x3);
505 unsigned int old_value = *base_address_u32;
506 unsigned int swap_value, compare_value;
508 compare_value = old_value;
510 (
reinterpret_cast<size_t>(agg) & 0x2)
511 ?
static_cast<unsigned int>(max(static_cast<int16_t>(old_value >> 16), val))
514 : (old_value & 0xFFFF0000) |
515 static_cast<unsigned int>(
516 max(static_cast<int16_t>(old_value & 0xFFFF), val));
517 old_value = atomicCAS(base_address_u32, compare_value, swap_value);
518 }
while (old_value != compare_value);
524 unsigned int* base_address_u32 =
525 reinterpret_cast<unsigned int*
>(
reinterpret_cast<size_t>(agg) & ~0x3);
532 constexpr
unsigned int byte_permutations[] = {0x3214, 0x3240, 0x3410, 0x4210};
533 unsigned int old_value = *base_address_u32;
534 unsigned int swap_value, compare_value;
536 compare_value = old_value;
537 auto max_value =
static_cast<unsigned int>(
540 static_cast<int8_t>(__byte_perm(
541 compare_value, 0, (reinterpret_cast<size_t>(agg) & 0x3) | 0x4440))));
542 swap_value = __byte_perm(
543 compare_value, max_value, byte_permutations[reinterpret_cast<size_t>(agg) & 0x3]);
544 old_value = atomicCAS(base_address_u32, compare_value, swap_value);
545 }
while (compare_value != old_value);
548 #if CUDA_VERSION > 10000 && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
549 __device__
void atomicMin16(int16_t* agg,
const int16_t val) {
550 unsigned short int* address_as_us =
reinterpret_cast<unsigned short int*
>(agg);
551 unsigned short int old = *address_as_us, assumed;
555 old = atomicCAS(address_as_us,
557 static_cast<unsigned short>(min(static_cast<short int>(val),
558 static_cast<short int>(assumed))));
559 }
while (assumed != old);
564 unsigned int* base_address_u32 =
565 reinterpret_cast<unsigned int*
>(
reinterpret_cast<size_t>(agg) & ~0x3);
567 unsigned int old_value = *base_address_u32;
568 unsigned int swap_value, compare_value;
570 compare_value = old_value;
572 (
reinterpret_cast<size_t>(agg) & 0x2)
573 ?
static_cast<unsigned int>(min(static_cast<int16_t>(old_value >> 16), val))
576 : (old_value & 0xFFFF0000) |
577 static_cast<unsigned int>(
578 min(static_cast<int16_t>(old_value & 0xFFFF), val));
579 old_value = atomicCAS(base_address_u32, compare_value, swap_value);
580 }
while (old_value != compare_value);
586 const int16_t skip_val) {
588 unsigned int* base_address_u32 =
589 reinterpret_cast<unsigned int*
>(
reinterpret_cast<size_t>(agg) & ~0x3);
591 unsigned int old_value = *base_address_u32;
592 unsigned int swap_value, compare_value;
594 compare_value = old_value;
595 int16_t selected_old_val = (
reinterpret_cast<size_t>(agg) & 0x2)
596 ?
static_cast<int16_t
>(old_value >> 16)
597 : static_cast<int16_t>(old_value & 0xFFFF);
600 (
reinterpret_cast<size_t>(agg) & 0x2)
601 ?
static_cast<unsigned int>(
602 selected_old_val == skip_val ? val : min(selected_old_val, val))
605 : (old_value & 0xFFFF0000) |
606 static_cast<unsigned int>(
607 selected_old_val == skip_val ? val : min(selected_old_val, val));
608 old_value = atomicCAS(base_address_u32, compare_value, swap_value);
609 }
while (old_value != compare_value);
614 unsigned int* base_address_u32 =
615 reinterpret_cast<unsigned int*
>(
reinterpret_cast<size_t>(agg) & ~0x3);
617 constexpr
unsigned int byte_permutations[] = {0x3214, 0x3240, 0x3410, 0x4210};
618 unsigned int old_value = *base_address_u32;
619 unsigned int swap_value, compare_value;
621 compare_value = old_value;
622 auto min_value =
static_cast<unsigned int>(
624 static_cast<int8_t>(__byte_perm(
625 compare_value, 0, (reinterpret_cast<size_t>(agg) & 0x3) | 0x4440))));
626 swap_value = __byte_perm(
627 compare_value, min_value, byte_permutations[reinterpret_cast<size_t>(agg) & 0x3]);
628 old_value = atomicCAS(base_address_u32, compare_value, swap_value);
629 }
while (compare_value != old_value);
634 unsigned int* base_address_u32 =
635 reinterpret_cast<unsigned int*
>(
reinterpret_cast<size_t>(agg) & ~0x3);
637 constexpr
unsigned int byte_permutations[] = {0x3214, 0x3240, 0x3410, 0x4210};
638 unsigned int old_value = *base_address_u32;
639 unsigned int swap_value, compare_value;
641 compare_value = old_value;
642 int8_t selected_old_val =
static_cast<int8_t
>(
643 __byte_perm(compare_value, 0, (reinterpret_cast<size_t>(agg) & 0x3) | 0x4440));
644 auto min_value =
static_cast<unsigned int>(
645 selected_old_val == skip_val ? val : min(val, selected_old_val));
646 swap_value = __byte_perm(
647 compare_value, min_value, byte_permutations[reinterpret_cast<size_t>(agg) & 0x3]);
648 old_value = atomicCAS(base_address_u32, compare_value, swap_value);
649 }
while (compare_value != old_value);
669 atomicMin(reinterpret_cast<double*>(agg), val);
673 atomicMin(reinterpret_cast<float*>(agg), val);
681 const int64_t offset,
683 const int64_t size_bytes) {
684 for (
auto i = 0; i < size_bytes; i++) {
685 varlen_buffer[offset + i] = value[i];
687 return &varlen_buffer[offset];
692 const int64_t null_val) {
693 unsigned long long int* address_as_ull =
reinterpret_cast<unsigned long long int*
>(agg);
694 unsigned long long int old = *address_as_ull, assumed;
696 if (val == null_val) {
701 if (static_cast<int64_t>(old) != null_val) {
702 if (static_cast<int64_t>(old) != val) {
711 old = atomicCAS(address_as_ull, assumed, val);
712 }
while (assumed != old);
717 #define DEF_AGG_ID_INT_SHARED(n) \
718 extern "C" __device__ void agg_id_int##n##_shared(int##n##_t* agg, \
719 const int##n##_t val) { \
727 #undef DEF_AGG_ID_INT_SHARED
730 *agg = *(
reinterpret_cast<const int64_t*
>(&val));
735 const double null_val) {
736 unsigned long long int* address_as_ull =
reinterpret_cast<unsigned long long int*
>(agg);
737 unsigned long long int old = *address_as_ull, assumed;
739 if (val == null_val) {
744 if (static_cast<int64_t>(old) != __double_as_longlong(null_val)) {
745 if (static_cast<int64_t>(old) != __double_as_longlong(val)) {
754 old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val));
755 }
while (assumed != old);
761 *agg = *(
reinterpret_cast<const int64_t*
>(val));
764 extern "C" __device__ int32_t
767 const double null_val) {
768 unsigned long long int* address_as_ull =
reinterpret_cast<unsigned long long int*
>(agg);
769 unsigned long long int old = *address_as_ull, assumed;
772 if (val == null_val) {
777 if (static_cast<int64_t>(old) != __double_as_longlong(null_val)) {
778 if (static_cast<int64_t>(old) != __double_as_longlong(val)) {
787 old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val));
788 }
while (assumed != old);
794 *agg = __float_as_int(val);
799 const float null_val) {
800 int* address_as_ull =
reinterpret_cast<int*
>(agg);
801 int old = *address_as_ull, assumed;
803 if (val == null_val) {
808 if (old != __float_as_int(null_val)) {
809 if (old != __float_as_int(val)) {
818 old = atomicCAS(address_as_ull, assumed, __float_as_int(val));
819 }
while (assumed != old);
824 #define DEF_SKIP_AGG(base_agg_func) \
825 extern "C" __device__ ADDR_T base_agg_func##_skip_val_shared( \
826 ADDR_T* agg, const DATA_T val, const DATA_T skip_val) { \
827 if (val != skip_val) { \
828 return base_agg_func##_shared(agg, val); \
833 #define DATA_T int64_t
834 #define ADDR_T uint64_t
839 #define DATA_T int32_t
840 #define ADDR_T uint32_t
848 const int32_t skip_val) {
849 if (val != skip_val) {
856 const int16_t skip_val) {
857 if (val != skip_val) {
864 const int16_t skip_val) {
865 if (val != skip_val) {
872 const int8_t skip_val) {
873 if (val != skip_val) {
880 const int8_t skip_val) {
881 if (val != skip_val) {
888 const int32_t skip_val) {
889 int32_t old = atomicExch(address, INT_MAX);
890 return atomicMin(address, old == skip_val ? val : min(old, val));
895 const int32_t skip_val) {
896 if (val != skip_val) {
903 const int32_t skip_val) {
904 unsigned int* address_as_int = (
unsigned int*)address;
905 int32_t old = atomicExch(address_as_int, 0);
906 int32_t old2 = atomicAdd(address_as_int, old == skip_val ? val : (val + old));
907 return old == skip_val ? old2 : (old2 + old);
912 const int32_t skip_val) {
913 if (val != skip_val) {
922 const int64_t skip_val) {
923 unsigned long long int* address_as_ull = (
unsigned long long int*)address;
924 int64_t old = atomicExch(address_as_ull, 0);
925 int64_t old2 = atomicAdd(address_as_ull, old == skip_val ? val : (val + old));
926 return old == skip_val ? old2 : (old2 + old);
931 const int64_t skip_val) {
932 if (val != skip_val) {
940 const int64_t skip_val) {
941 unsigned long long int* address_as_ull =
942 reinterpret_cast<unsigned long long int*
>(address);
943 unsigned long long int old = *address_as_ull, assumed;
947 old = atomicCAS(address_as_ull,
949 assumed == skip_val ? val : min((
long long)val, (
long long)assumed));
950 }
while (assumed != old);
957 const int64_t skip_val) {
958 if (val != skip_val) {
965 const int64_t skip_val) {
966 unsigned long long int* address_as_ull =
967 reinterpret_cast<unsigned long long int*
>(address);
968 unsigned long long int old = *address_as_ull, assumed;
972 old = atomicCAS(address_as_ull,
974 assumed == skip_val ? val : max((
long long)val, (
long long)assumed));
975 }
while (assumed != old);
982 const int64_t skip_val) {
983 if (val != skip_val) {
989 #define DEF_SKIP_AGG(base_agg_func) \
990 extern "C" __device__ ADDR_T base_agg_func##_skip_val_shared( \
991 ADDR_T* agg, const DATA_T val, const DATA_T skip_val) { \
992 if (val != skip_val) { \
993 return base_agg_func##_shared(agg, val); \
998 #define DATA_T double
999 #define ADDR_T uint64_t
1004 #define DATA_T float
1005 #define ADDR_T uint32_t
1013 const float skip_val) {
1014 if (__float_as_int(val) != __float_as_int(skip_val)) {
1015 float old = atomicExch(reinterpret_cast<float*>(agg), -FLT_MAX);
1016 atomicMax(reinterpret_cast<float*>(agg),
1017 __float_as_int(old) == __float_as_int(skip_val) ? val : fmaxf(old, val));
1022 float old = atomicExch(reinterpret_cast<float*>(address), FLT_MAX);
1024 reinterpret_cast<float*>(address),
1025 __float_as_int(old) == __float_as_int(skip_val) ? val : fminf(old, val));
1030 const float skip_val) {
1031 if (__float_as_int(val) != __float_as_int(skip_val)) {
1038 const float skip_val) {
1039 float old = atomicExch(address, 0.
f);
1040 atomicAdd(address, __float_as_int(old) == __float_as_int(skip_val) ? val : (val + old));
1045 const float skip_val) {
1046 if (__float_as_int(val) != __float_as_int(skip_val)) {
1053 const double skip_val) {
1054 unsigned long long int* address_as_ull = (
unsigned long long int*)address;
1055 double old = __longlong_as_double(atomicExch(address_as_ull, __double_as_longlong(0.)));
1058 __double_as_longlong(old) == __double_as_longlong(skip_val) ? val : (val + old));
1063 const double skip_val) {
1064 if (__double_as_longlong(val) != __double_as_longlong(skip_val)) {
1071 const double skip_val) {
1072 unsigned long long int* address_as_ull =
1073 reinterpret_cast<unsigned long long int*
>(address);
1074 unsigned long long int old = *address_as_ull;
1075 unsigned long long int skip_val_as_ull =
1076 *
reinterpret_cast<const unsigned long long*
>(&skip_val);
1077 unsigned long long int assumed;
1081 old = atomicCAS(address_as_ull,
1083 assumed == skip_val_as_ull
1084 ? *reinterpret_cast<unsigned long long*>(&val)
1085 : __double_as_longlong(min(val, __longlong_as_double(assumed))));
1086 }
while (assumed != old);
1088 return __longlong_as_double(old);
1093 const double skip_val) {
1094 if (val != skip_val) {
1101 const double skip_val) {
1102 if (__double_as_longlong(val) != __double_as_longlong(skip_val)) {
1103 double old = __longlong_as_double(atomicExch(
1104 reinterpret_cast<unsigned long long int*>(agg), __double_as_longlong(-DBL_MAX)));
1105 atomicMax(reinterpret_cast<double*>(agg),
1106 __double_as_longlong(old) == __double_as_longlong(skip_val)
1117 auto slot_address =
reinterpret_cast<unsigned long long int*
>(slot);
1118 const auto empty_key =
1119 static_cast<unsigned long long int*
>(
static_cast<void*
>(&init_val));
1120 const auto new_val_cast =
1121 static_cast<unsigned long long int*
>(
static_cast<void*
>(&new_val));
1123 const auto old_val = atomicCAS(slot_address, *empty_key, *new_val_cast);
1124 if (old_val == *empty_key) {
1134 unsigned int* slot_address =
reinterpret_cast<unsigned int*
>(slot);
1135 unsigned int compare_value =
static_cast<unsigned int>(init_val);
1136 unsigned int swap_value =
static_cast<unsigned int>(new_val);
1138 const unsigned int old_value = atomicCAS(slot_address, compare_value, swap_value);
1139 return old_value == compare_value;
1145 unsigned int* base_slot_address =
1146 reinterpret_cast<unsigned int*
>(
reinterpret_cast<size_t>(slot) & ~0x3);
1147 unsigned int old_value = *base_slot_address;
1148 unsigned int swap_value, compare_value;
1150 compare_value = old_value;
1153 if (static_cast<unsigned int>(init_val) !=
1155 compare_value, 0, (reinterpret_cast<size_t>(slot) & 0x2 ? 0x3244 : 0x4410))) {
1158 swap_value = __byte_perm(compare_value,
1159 static_cast<unsigned int>(new_val),
1160 (reinterpret_cast<size_t>(slot) & 0x2) ? 0x5410 : 0x3254);
1161 old_value = atomicCAS(base_slot_address, compare_value, swap_value);
1162 }
while (compare_value != old_value);
1170 unsigned int* base_slot_address =
1171 reinterpret_cast<unsigned int*
>(
reinterpret_cast<size_t>(slot) & ~0x3);
1172 constexpr
unsigned int byte_permutations[] = {0x3214, 0x3240, 0x3410, 0x4210};
1173 unsigned int old_value = *base_slot_address;
1174 unsigned int swap_value, compare_value;
1176 compare_value = old_value;
1179 if (static_cast<unsigned int>(init_val) !=
1180 __byte_perm(compare_value, 0, (reinterpret_cast<size_t>(slot) & 0x3) | 0x4440)) {
1183 swap_value = __byte_perm(compare_value,
1184 static_cast<unsigned int>(new_val),
1185 byte_permutations[reinterpret_cast<size_t>(slot) & 0x3]);
1186 old_value = atomicCAS(base_slot_address, compare_value, swap_value);
1187 }
while (compare_value != old_value);
1191 #include "../Utils/ChunkIter.cpp"
1194 #define EXECUTE_INCLUDE
1199 #undef EXECUTE_INCLUDE
1200 #include "../Utils/Regexp.cpp"
1201 #include "../Utils/StringLike.cpp"
1203 extern "C" __device__ uint64_t
string_decode(int8_t* chunk_iter_, int64_t pos) {
1210 : (
reinterpret_cast<uint64_t
>(vd.
pointer) & 0xffffffffffff) |
1211 (
static_cast<uint64_t
>(vd.
length) << 48);
1215 const uint32_t bitmap_bytes,
1216 const uint8_t* key_bytes,
1217 const uint32_t key_len) {
1218 const uint32_t bit_pos =
MurmurHash3(key_bytes, key_len, 0) % (bitmap_bytes * 8);
1219 const uint32_t word_idx = bit_pos / 32;
1220 const uint32_t bit_idx = bit_pos % 32;
1221 atomicOr(((uint32_t*)bitmap) + word_idx, 1 << bit_idx);
1226 const int64_t min_val,
1227 const int64_t base_dev_addr,
1228 const int64_t base_host_addr,
1229 const uint64_t sub_bitmap_count,
1230 const uint64_t bitmap_bytes) {
1231 const uint64_t bitmap_idx = val - min_val;
1232 const uint32_t byte_idx = bitmap_idx >> 3;
1233 const uint32_t word_idx = byte_idx >> 2;
1234 const uint32_t byte_word_idx = byte_idx & 3;
1235 const int64_t host_addr = *agg;
1236 uint32_t* bitmap = (uint32_t*)(base_dev_addr + host_addr - base_host_addr +
1237 (threadIdx.x & (sub_bitmap_count - 1)) * bitmap_bytes);
1238 switch (byte_word_idx) {
1240 atomicOr(&bitmap[word_idx], 1 << (bitmap_idx & 7));
1243 atomicOr(&bitmap[word_idx], 1 << ((bitmap_idx & 7) + 8));
1246 atomicOr(&bitmap[word_idx], 1 << ((bitmap_idx & 7) + 16));
1249 atomicOr(&bitmap[word_idx], 1 << ((bitmap_idx & 7) + 24));
1259 const int64_t min_val,
1260 const int64_t skip_val,
1261 const int64_t base_dev_addr,
1262 const int64_t base_host_addr,
1263 const uint64_t sub_bitmap_count,
1264 const uint64_t bitmap_bytes) {
1265 if (val != skip_val) {
1267 agg, val, min_val, base_dev_addr, base_host_addr, sub_bitmap_count, bitmap_bytes);
1275 const int64_t base_dev_addr,
1276 const int64_t base_host_addr) {
1278 const uint32_t index = hash >> (64 - b);
1279 const int32_t rank =
get_rank(hash << b, 64 - b);
1280 const int64_t host_addr = *agg;
1281 int32_t* M = (int32_t*)(base_dev_addr + host_addr - base_host_addr);
1286 __threadfence_block();
1303 if ((((row_count - 1) | 0x1F) - thread_pos) >= 32) {
1321 int64_t* output_buffer,
1322 const int32_t agg_idx) {
1323 if (threadIdx.x == agg_idx) {
__device__ void sync_warp_protected(int64_t thread_pos, int64_t row_count)
__device__ int32_t checked_single_agg_id_double_shared_slow(int64_t *agg, const double *valp, const double null_val)
__device__ void agg_max_float_shared(int32_t *agg, const float val)
__device__ uint32_t agg_count_float_shared(uint32_t *agg, const float val)
__device__ void agg_count_distinct_bitmap_skip_val_gpu(int64_t *agg, const int64_t val, const int64_t min_val, const int64_t skip_val, const int64_t base_dev_addr, const int64_t base_host_addr, const uint64_t sub_bitmap_count, const uint64_t bitmap_bytes)
__device__ bool dynamic_watchdog()
__device__ int64_t * get_matching_group_value_columnar(int64_t *groups_buffer, const uint32_t h, const int64_t *key, const uint32_t key_qw_count, const size_t entry_count)
__device__ void agg_max_shared(int64_t *agg, const int64_t val)
__device__ void write_back_nop(int64_t *dest, int64_t *src, const int32_t sz)
__device__ void agg_sum_float_skip_val_shared(int32_t *agg, const float val, const float skip_val)
FORCE_INLINE uint8_t get_rank(uint64_t x, uint32_t b)
__device__ void agg_min_int32_shared(int32_t *agg, const int32_t val)
__device__ int8_t thread_warp_idx(const int8_t warp_sz)
__device__ int64_t dw_sm_cycle_start[128]
__device__ void agg_id_float_shared(int32_t *agg, const float val)
__device__ void agg_min_double_shared(int64_t *agg, const double val)
__device__ int64_t get_thread_index()
RUNTIME_EXPORT NEVER_INLINE DEVICE uint64_t MurmurHash64A(const void *key, int len, uint64_t seed)
__device__ int32_t atomicMin32SkipVal(int32_t *address, int32_t val, const int32_t skip_val)
__device__ int32_t pos_step_impl()
__device__ void write_back_non_grouped_agg(int64_t *input_buffer, int64_t *output_buffer, const int32_t agg_idx)
Structures and runtime functions of streaming top-k heap.
__device__ void agg_min_int8_shared(int8_t *agg, const int8_t val)
__device__ int32_t checked_single_agg_id_double_shared(int64_t *agg, const double val, const double null_val)
__device__ float atomicMinFltSkipVal(int32_t *address, float val, const float skip_val)
__device__ const int64_t * init_shared_mem_nop(const int64_t *groups_buffer, const int32_t groups_buffer_size)
__device__ double atomicMin(double *address, double val)
__device__ void agg_max_int8_shared(int8_t *agg, const int8_t val)
__device__ int32_t checked_single_agg_id_float_shared(int32_t *agg, const float val, const float null_val)
__device__ void atomicMin8SkipVal(int8_t *agg, const int8_t val, const int8_t skip_val)
Functions to support geospatial operations used by the executor.
Macros and functions for groupby buffer compaction.
__device__ int64_t * get_matching_group_value(int64_t *groups_buffer, const uint32_t h, const T *key, const uint32_t key_count, const uint32_t row_size_quad)
__device__ uint32_t agg_count_int32_shared(uint32_t *agg, const int32_t val)
__device__ int64_t dw_cycle_budget
__device__ int64_t agg_sum_shared(int64_t *agg, const int64_t val)
__device__ void agg_id_double_shared_slow(int64_t *agg, const double *val)
DEVICE void ChunkIter_get_nth(ChunkIter *it, int n, bool uncompress, VarlenDatum *result, bool *is_end)
__device__ void agg_min_float_shared(int32_t *agg, const float val)
__device__ int8_t * agg_id_varlen_shared(int8_t *varlen_buffer, const int64_t offset, const int8_t *value, const int64_t size_bytes)
__device__ int64_t atomicMin64(int64_t *address, int64_t val)
__device__ int64_t * declare_dynamic_shared_memory()
__device__ void agg_max_double_shared(int64_t *agg, const double val)
__device__ void atomicSumDblSkipVal(double *address, const double val, const double skip_val)
__device__ int32_t agg_sum_int32_shared(int32_t *agg, const int32_t val)
__device__ int64_t agg_sum_skip_val_shared(int64_t *agg, const int64_t val, const int64_t skip_val)
__device__ void agg_sum_float_shared(int32_t *agg, const float val)
__device__ void agg_id_double_shared(int64_t *agg, const double val)
__device__ void agg_max_skip_val_shared(int64_t *agg, const int64_t val, const int64_t skip_val)
__device__ void atomicMax16(int16_t *agg, const int16_t val)
#define DEF_SKIP_AGG(base_agg_func)
__device__ int64_t get_block_index()
__device__ void agg_min_float_skip_val_shared(int32_t *agg, const float val, const float skip_val)
__device__ bool check_interrupt()
__device__ bool slotEmptyKeyCAS_int32(int32_t *slot, int32_t new_val, int32_t init_val)
__device__ int64_t atomicSum64SkipVal(int64_t *address, const int64_t val, const int64_t skip_val)
__device__ int32_t agg_sum_int32_skip_val_shared(int32_t *agg, const int32_t val, const int32_t skip_val)
__device__ void agg_min_int32_skip_val_shared(int32_t *agg, const int32_t val, const int32_t skip_val)
__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 uint64_t agg_count_double(uint64_t *agg, const double val)
__device__ void agg_count_distinct_bitmap_gpu(int64_t *agg, const int64_t val, const int64_t min_val, const int64_t base_dev_addr, const int64_t base_host_addr, const uint64_t sub_bitmap_count, const uint64_t bitmap_bytes)
__device__ void atomicSumFltSkipVal(float *address, const float val, const float skip_val)
__device__ void agg_sum_double_shared(int64_t *agg, const double val)
__inline__ __device__ uint32_t get_smid(void)
__device__ void agg_min_skip_val_shared(int64_t *agg, const int64_t val, const int64_t skip_val)
__device__ uint64_t agg_count_shared(uint64_t *agg, const int64_t val)
__device__ int64_t atomicMax64(int64_t *address, int64_t val)
__device__ bool slotEmptyKeyCAS(int64_t *slot, int64_t new_val, int64_t init_val)
__device__ int32_t pos_start_impl(const int32_t *row_index_resume)
__device__ int64_t atomicMax64SkipVal(int64_t *address, int64_t val, const int64_t skip_val)
__device__ void atomicMin16(int16_t *agg, const int16_t val)
__device__ void agg_max_float_skip_val_shared(int32_t *agg, const float val, const float skip_val)
__device__ int32_t runtime_interrupt_flag
__device__ void agg_approximate_count_distinct_gpu(int64_t *agg, const int64_t key, const uint32_t b, const int64_t base_dev_addr, const int64_t base_host_addr)
__device__ void sync_warp()
__device__ void atomicMin16SkipVal(int16_t *agg, const int16_t val, const int16_t skip_val)
__device__ void agg_sum_double_skip_val_shared(int64_t *agg, const double val, const double skip_val)
__device__ void agg_max_int8_skip_val_shared(int8_t *agg, const int8_t val, const int8_t skip_val)
__device__ void agg_max_int16_skip_val_shared(int16_t *agg, const int16_t val, const int16_t skip_val)
__device__ void atomicMin8(int8_t *agg, const int8_t val)
RUNTIME_EXPORT NEVER_INLINE DEVICE uint32_t MurmurHash3(const void *key, int len, const uint32_t seed)
__device__ void agg_min_int16_shared(int16_t *agg, const int16_t val)
__device__ void agg_max_int16_shared(int16_t *agg, const int16_t val)
__device__ const int64_t * init_shared_mem(const int64_t *global_groups_buffer, const int32_t groups_buffer_size)
RUNTIME_EXPORT ALWAYS_INLINE uint32_t agg_count_int32(uint32_t *agg, const int32_t)
__device__ void agg_min_double_skip_val_shared(int64_t *agg, const double val, const double skip_val)
#define DEF_AGG_ID_INT_SHARED(n)
__device__ uint64_t agg_count_double_shared(uint64_t *agg, const double val)
__device__ T get_empty_key()
__device__ void agg_min_int16_skip_val_shared(int16_t *agg, const int16_t val, const int16_t skip_val)
__device__ void sync_threadblock()
__device__ void agg_min_int8_skip_val_shared(int8_t *agg, const int8_t val, const int8_t skip_val)
RUNTIME_EXPORT ALWAYS_INLINE uint64_t agg_count(uint64_t *agg, const int64_t)
__device__ void atomicMax8(int8_t *agg, const int8_t val)
__device__ void agg_id_shared(int64_t *agg, const int64_t val)
__device__ double atomicMax(double *address, double val)
__device__ uint64_t string_decode(int8_t *chunk_iter_, int64_t pos)
__device__ int32_t atomicSum32SkipVal(int32_t *address, const int32_t val, const int32_t skip_val)
__device__ double atomicMinDblSkipVal(double *address, double val, const double skip_val)
__device__ int32_t get_matching_group_value_columnar_slot(int64_t *groups_buffer, const uint32_t entry_count, const uint32_t h, const T *key, const uint32_t key_count)
__device__ void agg_max_int32_shared(int32_t *agg, const int32_t val)
__device__ int32_t checked_single_agg_id_shared(int64_t *agg, const int64_t val, const int64_t null_val)
__device__ void agg_max_int32_skip_val_shared(int32_t *agg, const int32_t val, const int32_t skip_val)
__device__ int32_t dw_abort
__device__ bool slotEmptyKeyCAS_int16(int16_t *slot, int16_t new_val, int16_t init_val)
__device__ void agg_max_double_skip_val_shared(int64_t *agg, const double val, const double skip_val)
FORCE_INLINE HOST DEVICE T align_to_int64(T addr)
__device__ int64_t atomicMin64SkipVal(int64_t *address, int64_t val, const int64_t skip_val)
Functions to support array operations used by the executor.
__device__ void force_sync()
RUNTIME_EXPORT ALWAYS_INLINE uint32_t agg_count_float(uint32_t *agg, const float val)
__device__ void agg_min_shared(int64_t *agg, const int64_t val)
__device__ bool slotEmptyKeyCAS_int8(int8_t *slot, int8_t new_val, int8_t init_val)
__device__ int32_t group_buff_idx_impl()