github.com/rohankumardubey/aresdb@v0.0.2-0.20190517170215-e54e3ca06b9c/query/functor.hpp (about) 1 // Copyright (c) 2017-2018 Uber Technologies, Inc. 2 // 3 // Licensed under the Apache License, Version 2.0 (the "License"); 4 // you may not use this file except in compliance with the License. 5 // You may obtain a copy of the License at 6 // 7 // http://www.apache.org/licenses/LICENSE-2.0 8 // 9 // Unless required by applicable law or agreed to in writing, software 10 // distributed under the License is distributed on an "AS IS" BASIS, 11 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. 12 // See the License for the specific language governing permissions and 13 // limitations under the License. 14 15 16 #ifndef QUERY_FUNCTOR_HPP_ 17 #define QUERY_FUNCTOR_HPP_ 18 #include <cuda_runtime.h> 19 #include <thrust/tuple.h> 20 #include <iostream> 21 #include <tuple> 22 #include "query/iterator.hpp" 23 #include "query/time_series_aggregate.h" 24 #include "query/utils.hpp" 25 26 namespace ares { 27 28 // logical operators. 29 30 struct AndFunctor { 31 __host__ __device__ 32 thrust::tuple<bool, bool> operator()( 33 const thrust::tuple<bool, bool> t1, 34 const thrust::tuple<bool, bool> t2) const { 35 // if one of them is null, the result is null. 36 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 37 return thrust::make_tuple(false, false); 38 } 39 return thrust::make_tuple(thrust::get<0>(t1) && thrust::get<0>(t2) != 0, 40 true); 41 } 42 }; 43 44 struct OrFunctor { 45 __host__ __device__ 46 thrust::tuple<bool, bool> operator()( 47 const thrust::tuple<bool, bool> t1, 48 const thrust::tuple<bool, bool> t2) const { 49 bool value1 = thrust::get<0>(t1); 50 bool valid1 = thrust::get<1>(t1); 51 bool value2 = thrust::get<0>(t2); 52 bool valid2 = thrust::get<1>(t2); 53 54 // If one of them is true, the result is true. 55 if ((value1 && valid1) || (value2 && valid2)) { 56 return thrust::make_tuple(true, true); 57 } 58 59 // Otherwise if one of them is null, the result is null. 60 if (!valid1 || !valid2) { 61 return thrust::make_tuple(false, false); 62 } 63 64 return thrust::make_tuple(false, true); 65 } 66 }; 67 68 struct NotFunctor { 69 __host__ __device__ 70 thrust::tuple<bool, bool> operator()( 71 const thrust::tuple<bool, bool> t) const { 72 // Not null is null. 73 if (!thrust::get<1>(t)) { 74 return thrust::make_tuple(false, false); 75 } 76 return thrust::make_tuple(!thrust::get<0>(t), true); 77 } 78 }; 79 80 // comparison operators 81 82 template<typename T> 83 struct EqualFunctor { 84 __host__ __device__ 85 thrust::tuple<bool, bool> operator()(const thrust::tuple<T, bool> t1, 86 const thrust::tuple<T, bool> t2) const { 87 // if one of them is null, the result is null. 88 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 89 return thrust::make_tuple(false, false); 90 } 91 92 return thrust::make_tuple(thrust::get<0>(t1) == thrust::get<0>(t2), true); 93 } 94 }; 95 96 // Equal functor for GeoPointT. 97 template<> 98 struct EqualFunctor<GeoPointT> { 99 __host__ __device__ 100 thrust::tuple<bool, bool> operator()( 101 const thrust::tuple<const GeoPointT, bool> t1, 102 const thrust::tuple<const GeoPointT, bool> t2) const { 103 // if one of them is null, the result is null. 104 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 105 return thrust::make_tuple(false, false); 106 } 107 return thrust::make_tuple( 108 thrust::get<0>(t1).Lat == thrust::get<0>(t2).Lat && 109 thrust::get<0>(t1).Long == thrust::get<0>(t2).Long, 110 true); 111 } 112 }; 113 114 template<typename T> 115 struct NotEqualFunctor { 116 __host__ __device__ 117 thrust::tuple<bool, bool> operator()(const thrust::tuple<T, bool> t1, 118 const thrust::tuple<T, bool> t2) const { 119 // if one of them is null, the result is null. 120 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 121 return thrust::make_tuple(false, false); 122 } 123 124 return thrust::make_tuple(thrust::get<0>(t1) != thrust::get<0>(t2), true); 125 } 126 }; 127 128 template<typename T> 129 struct LessThanFunctor { 130 __host__ __device__ 131 thrust::tuple<bool, bool> operator()(const thrust::tuple<T, bool> t1, 132 const thrust::tuple<T, bool> t2) const { 133 // if one of them is null, the result is null. 134 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 135 return thrust::make_tuple(false, false); 136 } 137 138 return thrust::make_tuple(thrust::get<0>(t1) < thrust::get<0>(t2), true); 139 } 140 }; 141 142 template<typename T> 143 struct LessThanOrEqualFunctor { 144 __host__ __device__ 145 thrust::tuple<bool, bool> operator()(const thrust::tuple<T, bool> t1, 146 const thrust::tuple<T, bool> t2) const { 147 // if one of them is null, the result is null. 148 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 149 return thrust::make_tuple(false, false); 150 } 151 152 return thrust::make_tuple(thrust::get<0>(t1) <= thrust::get<0>(t2), true); 153 } 154 }; 155 156 template<typename T> 157 struct GreaterThanFunctor { 158 __host__ __device__ 159 thrust::tuple<bool, bool> operator()(const thrust::tuple<T, bool> t1, 160 const thrust::tuple<T, bool> t2) const { 161 // if one of them is null, the result is null. 162 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 163 return thrust::make_tuple(false, false); 164 } 165 166 return thrust::make_tuple(thrust::get<0>(t1) > thrust::get<0>(t2), true); 167 } 168 }; 169 170 template<typename T> 171 struct GreaterThanOrEqualFunctor { 172 __host__ __device__ 173 thrust::tuple<bool, bool> operator()(const thrust::tuple<T, bool> t1, 174 const thrust::tuple<T, bool> t2) const { 175 // if one of them is null, the result is null. 176 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 177 return thrust::make_tuple(false, false); 178 } 179 180 return thrust::make_tuple(thrust::get<0>(t1) >= thrust::get<0>(t2), true); 181 } 182 }; 183 184 // arithmetic operators 185 186 template<typename T> 187 struct PlusFunctor { 188 __host__ __device__ 189 thrust::tuple<T, bool> operator()(const thrust::tuple<T, bool> t1, 190 const thrust::tuple<T, bool> t2) const { 191 // if one of them is null, the result is null. 192 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 193 return thrust::make_tuple(0, false); 194 } 195 return thrust::make_tuple(thrust::get<0>(t1) + thrust::get<0>(t2), true); 196 } 197 }; 198 199 template<typename T> 200 struct MinusFunctor { 201 __host__ __device__ 202 thrust::tuple<T, bool> operator()(const thrust::tuple<T, bool> t1, 203 const thrust::tuple<T, bool> t2) const { 204 // if one of them is null, the result is null. 205 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 206 return thrust::make_tuple(0, false); 207 } 208 209 return thrust::make_tuple(thrust::get<0>(t1) - thrust::get<0>(t2), true); 210 } 211 }; 212 213 template<typename T> 214 struct MultiplyFunctor { 215 __host__ __device__ 216 thrust::tuple<T, bool> operator()(const thrust::tuple<T, bool> t1, 217 const thrust::tuple<T, bool> t2) const { 218 // if one of them is null, the result is null. 219 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 220 return thrust::make_tuple(0, false); 221 } 222 223 return thrust::make_tuple(thrust::get<0>(t1) * thrust::get<0>(t2), true); 224 } 225 }; 226 227 template<typename T> 228 struct DivideFunctor { 229 __host__ __device__ 230 thrust::tuple<T, bool> operator()(const thrust::tuple<T, bool> t1, 231 const thrust::tuple<T, bool> t2) const { 232 // if one of them is null, the result is null. 233 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 234 return thrust::make_tuple(0, false); 235 } 236 237 return thrust::make_tuple(thrust::get<0>(t1) / thrust::get<0>(t2), true); 238 } 239 }; 240 241 template<typename T> 242 struct ModFunctor { 243 __host__ __device__ 244 thrust::tuple<T, bool> operator()(const thrust::tuple<T, bool> t1, 245 const thrust::tuple<T, bool> t2) const { 246 // if one of them is null, the result is null. 247 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 248 return thrust::make_tuple(0, false); 249 } 250 251 return thrust::make_tuple(thrust::get<0>(t1) % thrust::get<0>(t2), true); 252 } 253 }; 254 255 template<typename T> 256 struct NegateFunctor { 257 __host__ __device__ 258 thrust::tuple<T, bool> operator()(const thrust::tuple<T, bool> t) const { 259 bool valid = thrust::get<1>(t); 260 if (!valid) { 261 return thrust::make_tuple(0, valid); 262 } 263 return thrust::make_tuple(-thrust::get<0>(t), valid); 264 } 265 }; 266 267 // bitwise operators 268 269 template<typename T> 270 struct BitwiseAndFunctor { 271 __host__ __device__ 272 thrust::tuple<T, bool> operator()(const thrust::tuple<T, bool> t1, 273 const thrust::tuple<T, bool> t2) const { 274 // if one of them is null, the result is null. 275 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 276 return thrust::make_tuple(0, false); 277 } 278 return thrust::make_tuple(thrust::get<0>(t1) & thrust::get<0>(t2), true); 279 } 280 }; 281 282 template<typename T> 283 struct BitwiseOrFunctor { 284 __host__ __device__ 285 thrust::tuple<T, bool> operator()(const thrust::tuple<T, bool> t1, 286 const thrust::tuple<T, bool> t2) const { 287 // if one of them is null, the result is null. 288 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 289 return thrust::make_tuple(0, false); 290 } 291 return thrust::make_tuple(thrust::get<0>(t1) | thrust::get<0>(t2), true); 292 } 293 }; 294 295 template<typename T> 296 struct BitwiseXorFunctor { 297 __host__ __device__ 298 thrust::tuple<T, bool> operator()(const thrust::tuple<T, bool> t1, 299 const thrust::tuple<T, bool> t2) const { 300 // if one of them is null, the result is null. 301 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 302 return thrust::make_tuple(0, false); 303 } 304 return thrust::make_tuple(thrust::get<0>(t1) ^ thrust::get<0>(t2), true); 305 } 306 }; 307 308 template<typename T> 309 struct BitwiseNotFunctor { 310 __host__ __device__ 311 thrust::tuple<T, bool> operator()(const thrust::tuple<T, bool> t) const { 312 if (!thrust::get<1>(t)) { 313 return thrust::make_tuple(0, false); 314 } 315 return thrust::make_tuple(~thrust::get<0>(t), true); 316 } 317 }; 318 319 template<typename T> 320 struct FloorFunctor { 321 __host__ __device__ 322 thrust::tuple<T, bool> operator()(const thrust::tuple<T, bool> t1, 323 const thrust::tuple<T, bool> t2) const { 324 // if one of them is null, the result is null. 325 if (!thrust::get<1>(t1) || !thrust::get<1>(t2)) { 326 return thrust::make_tuple(0, false); 327 } 328 329 return thrust::make_tuple( 330 thrust::get<0>(t1) - thrust::get<0>(t1) % thrust::get<0>(t2), 331 true); 332 } 333 }; 334 335 336 // misc operators 337 338 struct IsNullFunctor { 339 __host__ __device__ 340 thrust::tuple<bool, bool> operator()( 341 const thrust::tuple<bool, bool> t) const { 342 return thrust::make_tuple(!thrust::get<1>(t), true); 343 } 344 }; 345 346 struct IsNotNullFunctor { 347 __host__ __device__ 348 thrust::tuple<bool, bool> operator()( 349 const thrust::tuple<bool, bool> t) const { 350 return thrust::make_tuple(thrust::get<1>(t), true); 351 } 352 }; 353 354 // Functor used for filtering based on a single column directly. 355 // It just return the argument as it is. 356 template<typename T> 357 struct NoopFunctor { 358 __host__ __device__ 359 thrust::tuple<T, bool> operator()(const thrust::tuple<T, bool> t) const { 360 return t; 361 } 362 }; 363 364 // date operators. 365 struct GetWeekStartFunctor { 366 __host__ __device__ 367 thrust::tuple<uint32_t, bool> operator()( 368 const thrust::tuple<uint32_t, bool> t) const; 369 }; 370 371 struct GetMonthStartFunctor { 372 __host__ __device__ 373 thrust::tuple<uint32_t, bool> operator()( 374 const thrust::tuple<uint32_t, bool> t) const; 375 }; 376 377 struct GetQuarterStartFunctor { 378 __host__ __device__ 379 thrust::tuple<uint32_t, bool> operator()( 380 const thrust::tuple<uint32_t, bool> t) const; 381 }; 382 383 struct GetYearStartFunctor { 384 __host__ __device__ 385 thrust::tuple<uint32_t, bool> operator()( 386 const thrust::tuple<uint32_t, bool> t) const; 387 }; 388 389 struct GetDayOfMonthFunctor { 390 __host__ __device__ 391 thrust::tuple<uint32_t, bool> operator()( 392 const thrust::tuple<uint32_t, bool> t) const; 393 }; 394 395 struct GetDayOfYearFunctor { 396 __host__ __device__ 397 thrust::tuple<uint32_t, bool> operator()( 398 const thrust::tuple<uint32_t, bool> t) const; 399 }; 400 401 struct GetMonthOfYearFunctor { 402 __host__ __device__ 403 thrust::tuple<uint32_t, bool> operator()( 404 const thrust::tuple<uint32_t, bool> t) const; 405 }; 406 407 struct GetQuarterOfYearFunctor { 408 __host__ __device__ 409 thrust::tuple<uint32_t, bool> operator()( 410 const thrust::tuple<uint32_t, bool> t) const; 411 }; 412 413 template <typename I> 414 inline __host__ __device__ uint64_t hll_hash(I value) { 415 uint64_t hashedOutput[2]; 416 murmur3sum128(reinterpret_cast<uint8_t *>(&value), sizeof(value), 0, 417 hashedOutput); 418 return hashedOutput[0]; 419 } 420 421 template <> 422 inline __host__ __device__ uint64_t hll_hash(UUIDT uuid) { 423 return uuid.p1 ^ uuid.p2; 424 } 425 426 // GetHLLValueFunctor calcuates the register 427 template <typename I> 428 struct GetHLLValueFunctor { 429 __host__ __device__ thrust::tuple<uint32_t, bool> operator()( 430 thrust::tuple<I, bool> input) const { 431 if (!thrust::get<1>(input)) { 432 return thrust::make_tuple<uint32_t>(0, false); 433 } 434 I value = thrust::get<0>(input); 435 uint64_t hashed = hll_hash(value); 436 uint32_t group = static_cast<uint32_t>(hashed & ((1 << HLL_BITS) - 1)); 437 uint32_t rho = 0; 438 while (true) { 439 uint32_t h = hashed & (1 << (rho + HLL_BITS)); 440 if (rho + HLL_BITS < 64 && h == 0) { 441 rho++; 442 } else { 443 break; 444 } 445 } 446 return thrust::make_tuple(rho << 16 | group, true); 447 } 448 }; 449 450 // We combine all unary functors into a single functor class and use the 451 // UnaryFunctorType enum to do RTTI function call. Thereby we reduce number 452 // of class bindings for thrust template functions. 453 template<typename O, typename I, typename Enable = void> 454 struct UnaryFunctor { 455 typedef thrust::tuple<I, bool> argument_type; 456 typedef thrust::tuple<O, bool> result_type; 457 458 explicit UnaryFunctor(UnaryFunctorType functorType) 459 : functorType(functorType) { 460 } 461 462 UnaryFunctorType functorType; 463 464 __host__ __device__ 465 result_type operator()(const argument_type t) const { 466 switch (functorType) { 467 case Not:return NotFunctor()(t); 468 case IsNull:return IsNullFunctor()(t); 469 case IsNotNull:return IsNotNullFunctor()(t); 470 case Negate:return NegateFunctor<I>()(t); 471 case BitwiseNot:return BitwiseNotFunctor<I>()(t); 472 case Noop:return NoopFunctor<I>()(t); 473 case GetWeekStart: return GetWeekStartFunctor()(t); 474 case GetMonthStart: return GetMonthStartFunctor()(t); 475 case GetQuarterStart: return GetQuarterStartFunctor()(t); 476 case GetYearStart: return GetYearStartFunctor()(t); 477 case GetDayOfMonth: return GetDayOfMonthFunctor()(t); 478 case GetDayOfYear: return GetDayOfYearFunctor()(t); 479 case GetMonthOfYear: return GetMonthOfYearFunctor()(t); 480 case GetQuarterOfYear: return GetQuarterOfYearFunctor()(t); 481 case GetHLLValue: return GetHLLValueFunctor<I>()(t); 482 default: 483 // We will not handle uncaught enum here since the AQL compiler 484 // should ensure that. 485 return t; 486 } 487 } 488 }; 489 490 // disable unary transformation from any type to UUIDT 491 template <typename I> 492 struct UnaryFunctor<UUIDT, I> { 493 typedef thrust::tuple<I, bool> argument_type; 494 typedef thrust::tuple<UUIDT, bool> result_type; 495 496 explicit UnaryFunctor(UnaryFunctorType functorType) 497 : functorType(functorType) {} 498 499 UnaryFunctorType functorType; 500 501 __host__ __device__ result_type operator()(const argument_type t) const { 502 UUIDT uuid = {0, 0}; 503 return thrust::make_tuple<UUIDT, bool>(uuid, false); 504 } 505 }; 506 507 // Specialization with float type to avoid illegal functor type template 508 // generation. 509 template <typename O> 510 struct UnaryFunctor< 511 O, float_t, typename std::enable_if<!std::is_same<O, UUIDT>::value>::type> { 512 typedef thrust::tuple<float_t, bool> argument_type; 513 typedef thrust::tuple<O, bool> result_type; 514 515 explicit UnaryFunctor(UnaryFunctorType functorType) 516 : functorType(functorType) {} 517 518 UnaryFunctorType functorType; 519 520 __host__ __device__ result_type operator()(const argument_type t) const { 521 switch (functorType) { 522 case Not: 523 return NotFunctor()(t); 524 case IsNull: 525 return IsNullFunctor()(t); 526 case IsNotNull: 527 return IsNotNullFunctor()(t); 528 case Negate: 529 return NegateFunctor<float_t>()(t); 530 case Noop: 531 return NoopFunctor<float_t>()(t); 532 default: 533 // We will not handle uncaught enum here since the AQL compiler 534 // should ensure that. 535 return t; 536 } 537 } 538 }; 539 540 // specialize UnaryFunctor for UUIDT input input type to types other than UUIDT 541 template <typename O> 542 struct UnaryFunctor< 543 O, 544 UUIDT, 545 typename std::enable_if<!std::is_same<O, GeoPointT>::value>::type 546 > { 547 typedef thrust::tuple<UUIDT, bool> argument_type; 548 typedef thrust::tuple<O, bool> result_type; 549 550 explicit UnaryFunctor(UnaryFunctorType functorType) 551 : functorType(functorType) {} 552 553 UnaryFunctorType functorType; 554 555 __host__ __device__ result_type operator()(const argument_type t) const { 556 switch (functorType) { 557 case GetHLLValue: 558 return GetHLLValueFunctor<UUIDT>()(t); 559 default: 560 O o; 561 return thrust::make_tuple<O, bool>(o, false); 562 } 563 } 564 }; 565 566 // Specialize unary transformation from UUIDT to UUIDT 567 template <> 568 struct UnaryFunctor<UUIDT, UUIDT> { 569 typedef thrust::tuple<UUIDT, bool> argument_type; 570 typedef thrust::tuple<UUIDT, bool> result_type; 571 572 explicit UnaryFunctor(UnaryFunctorType functorType) 573 : functorType(functorType) {} 574 575 UnaryFunctorType functorType; 576 577 __host__ __device__ result_type operator()(const argument_type t) const { 578 return NoopFunctor<UUIDT>()(t); 579 } 580 }; 581 582 // Specialize UnaryFunctor for GeoPointT input type to types other than 583 // GeoPointT 584 template <typename O> 585 struct UnaryFunctor< 586 O, 587 GeoPointT, 588 typename std::enable_if< 589 !std::is_same<O, GeoPointT>::value && !std::is_same<O, UUIDT>::value 590 >::type 591 >{ 592 typedef thrust::tuple<GeoPointT, bool> argument_type; 593 typedef thrust::tuple<O, bool> result_type; 594 595 explicit UnaryFunctor(UnaryFunctorType functorType) 596 : functorType(functorType) {} 597 598 UnaryFunctorType functorType; 599 600 __host__ __device__ result_type operator()(const argument_type t) const { 601 O o; 602 return thrust::make_tuple<O, bool>(o, false); 603 } 604 }; 605 606 // Specialize UnaryFunctor for input type other than GeoPointT to 607 // GeoPointT output type 608 template <typename I> 609 struct UnaryFunctor< 610 GeoPointT, 611 I, 612 typename std::enable_if<!std::is_same<I, GeoPointT>::value>::type 613 > { 614 typedef thrust::tuple<I, bool> argument_type; 615 typedef thrust::tuple<GeoPointT, bool> result_type; 616 617 explicit UnaryFunctor(UnaryFunctorType functorType) 618 : functorType(functorType) {} 619 620 UnaryFunctorType functorType; 621 622 __host__ __device__ result_type operator()(const argument_type t) const { 623 GeoPointT o; 624 return thrust::make_tuple<GeoPointT, bool>(o, false); 625 } 626 }; 627 628 // Specialize from GeoPointT to GeoPointT 629 template <> 630 struct UnaryFunctor<GeoPointT, GeoPointT> { 631 typedef thrust::tuple<GeoPointT, bool> argument_type; 632 typedef thrust::tuple<GeoPointT, bool> result_type; 633 634 explicit UnaryFunctor(UnaryFunctorType functorType) 635 : functorType(functorType) {} 636 637 UnaryFunctorType functorType; 638 639 __host__ __device__ result_type operator()(const argument_type t) const { 640 return NoopFunctor<GeoPointT>()(t); 641 } 642 }; 643 644 645 // Specialize from GeoPointT to float_t(to resolve partial specialization tie) 646 template <> 647 struct UnaryFunctor<GeoPointT, float_t> { 648 typedef thrust::tuple<float_t, bool> argument_type; 649 typedef thrust::tuple<GeoPointT, bool> result_type; 650 651 explicit UnaryFunctor(UnaryFunctorType functorType) 652 : functorType(functorType) {} 653 654 UnaryFunctorType functorType; 655 656 __host__ __device__ result_type operator()(const argument_type t) const { 657 GeoPointT g; 658 return thrust::make_tuple<GeoPointT, bool>(g, false); 659 } 660 }; 661 662 // UnaryPredicateFunctor simply applies the UnaryFunctor f on the argument 663 // and extract the 1st element of the result tuple which should usually 664 // be a boolean value. 665 template<typename O, typename I> 666 struct UnaryPredicateFunctor { 667 explicit UnaryPredicateFunctor(UnaryFunctorType functorType) 668 : f(UnaryFunctor<O, I>(functorType)) { 669 } 670 671 typedef typename UnaryFunctor<O, I>::argument_type argument_type; 672 673 UnaryFunctor<O, I> f; 674 675 __host__ __device__ 676 bool operator()(const argument_type t) { 677 return thrust::get<0>(f(t)); 678 } 679 }; 680 681 // Same as the single UnaryFunctor class to avoid generating too many class 682 // bindings for thrust template functions. 683 template <typename O, typename I, typename Enable = void> 684 struct BinaryFunctor { 685 typedef thrust::tuple<I, bool> argument_type; 686 typedef thrust::tuple<O, bool> result_type; 687 688 explicit BinaryFunctor(BinaryFunctorType functorType) 689 : functorType(functorType) { 690 } 691 692 BinaryFunctorType functorType; 693 694 __host__ __device__ 695 result_type operator()(const argument_type t1, const argument_type t2) const { 696 switch (functorType) { 697 case And:return AndFunctor()(t1, t2); 698 case Or:return OrFunctor()(t1, t2); 699 case Equal:return EqualFunctor<I>()(t1, t2); 700 case NotEqual:return NotEqualFunctor<I>()(t1, t2); 701 case LessThan:return LessThanFunctor<I>()(t1, t2); 702 case LessThanOrEqual:return LessThanOrEqualFunctor<I>()(t1, t2); 703 case GreaterThan:return GreaterThanFunctor<I>()(t1, t2); 704 case GreaterThanOrEqual:return GreaterThanOrEqualFunctor<I>()(t1, t2); 705 case Plus:return PlusFunctor<I>()(t1, t2); 706 case Minus:return MinusFunctor<I>()(t1, t2); 707 case Multiply:return MultiplyFunctor<I>()(t1, t2); 708 case Divide:return DivideFunctor<I>()(t1, t2); 709 case Mod:return ModFunctor<I>()(t1, t2); 710 case BitwiseAnd:return BitwiseAndFunctor<I>()(t1, t2); 711 case BitwiseOr:return BitwiseOrFunctor<I>()(t1, t2); 712 case BitwiseXor:return BitwiseXorFunctor<I>()(t1, t2); 713 case Floor:return FloorFunctor<I>()(t1, t2); 714 default: 715 // We will not handle uncaught enum here since the AQL compiler 716 // should ensure that. 717 return t1; 718 } 719 } 720 }; 721 722 template <typename I> 723 struct BinaryFunctor<UUIDT, I> { 724 typedef thrust::tuple<I, bool> argument_type; 725 typedef thrust::tuple<UUIDT, bool> result_type; 726 727 explicit BinaryFunctor(BinaryFunctorType functorType) 728 : functorType(functorType) {} 729 730 BinaryFunctorType functorType; 731 732 __host__ __device__ result_type operator()(const argument_type t1, 733 const argument_type t2) const { 734 UUIDT uuid = {0, 0}; 735 return thrust::make_tuple(uuid, false); 736 } 737 }; 738 739 // Specialization with float type to avoid illegal functor type template 740 // generation. 741 template <typename O> 742 struct BinaryFunctor< 743 O, float_t, typename std::enable_if<!std::is_same<O, UUIDT>::value>::type> { 744 typedef thrust::tuple<float_t, bool> argument_type; 745 typedef thrust::tuple<O, bool> result_type; 746 747 explicit BinaryFunctor(BinaryFunctorType functorType) 748 : functorType(functorType) { 749 } 750 751 BinaryFunctorType functorType; 752 753 __host__ __device__ 754 result_type operator()(const argument_type t1, const argument_type t2) const { 755 switch (functorType) { 756 case And:return AndFunctor()(t1, t2); 757 case Or:return OrFunctor()(t1, t2); 758 case Equal:return EqualFunctor<float_t>()(t1, t2); 759 case NotEqual:return NotEqualFunctor<float_t>()(t1, t2); 760 case LessThan:return LessThanFunctor<float_t>()(t1, t2); 761 case LessThanOrEqual:return LessThanOrEqualFunctor<float_t>()(t1, t2); 762 case GreaterThan:return GreaterThanFunctor<float_t>()(t1, t2); 763 case GreaterThanOrEqual: 764 return GreaterThanOrEqualFunctor<float_t>()(t1, 765 t2); 766 case Plus:return PlusFunctor<float_t>()(t1, t2); 767 case Minus:return MinusFunctor<float_t>()(t1, t2); 768 case Multiply:return MultiplyFunctor<float_t>()(t1, t2); 769 case Divide:return DivideFunctor<float_t>()(t1, t2); 770 default: 771 // We will not handle uncaught enum here since the AQL compiler 772 // should ensure that. 773 return t1; 774 } 775 } 776 }; 777 778 // Specialization with GeoPointT type to avoid illegal functor type template 779 // generation. 780 template <typename O> 781 struct BinaryFunctor< 782 O, GeoPointT, 783 typename std::enable_if<!std::is_same<O, UUIDT>::value>::type> { 784 typedef thrust::tuple<GeoPointT, bool> argument_type; 785 typedef thrust::tuple<O, bool> result_type; 786 787 explicit BinaryFunctor(BinaryFunctorType functorType) 788 : functorType(functorType) {} 789 790 BinaryFunctorType functorType; 791 792 __host__ __device__ result_type operator()(const argument_type t1, 793 const argument_type t2) const { 794 switch (functorType) { 795 case Equal: 796 return EqualFunctor<GeoPointT>()(t1, t2); 797 default: 798 // should not came here, GeoPoint only support equal function 799 return false; 800 } 801 } 802 }; 803 804 // BinaryPredicateFunctor simply applies the BinaryFunctor f on <lhs, rhs> 805 // and extract the 1st element of the result tuple which should usually 806 // be a boolean value. 807 template<typename O, typename I> 808 struct BinaryPredicateFunctor { 809 explicit BinaryPredicateFunctor(BinaryFunctorType functorType) 810 : f(BinaryFunctor<O, I>(functorType)) { 811 } 812 813 typedef typename BinaryFunctor<O, I>::argument_type argument_type; 814 815 BinaryFunctor<O, I> f; 816 817 __host__ __device__ 818 bool operator()(const argument_type t1, const argument_type t2) { 819 return thrust::get<0>(f(t1, t2)); 820 } 821 }; 822 823 // RemoveFilter is a functor to tell whether we need to remove an index 824 // from index vector given a pre-computed predicate vector. Note the 825 // predicate vector tells us whether we need to "keep" the row. So we 826 // need to negate the predicate value.The argumenttype is a tuple of 827 // <index seq, index value> where we will only use the 1st element of 828 // the tuple to indexing the predicate vector. 829 template<typename Value, typename Predicate> 830 struct RemoveFilter { 831 explicit RemoveFilter(Predicate *predicates) : predicates(predicates) {} 832 833 Predicate *predicates; 834 835 __host__ __device__ 836 bool operator()(const Value &index) { 837 return !predicates[thrust::get<0>(index)]; 838 } 839 }; 840 841 // This functor lookup a value in a hash table 842 template<typename I> 843 struct HashLookupFunctor { 844 uint8_t *buckets; 845 uint8_t *stash; 846 uint32_t seeds[4] = {0}; 847 848 int keyBytes; 849 int bucketBytes; 850 // numHashes might be less then 4, but never more than 4 851 int numHashes; 852 int numBuckets; 853 854 int offsetToSignature; 855 int offsetToKey; 856 857 typedef thrust::tuple<I, bool> argument_type; 858 859 explicit HashLookupFunctor(uint8_t *_buckets, uint32_t *_seeds, int _keyBytes, 860 int _numHashes, int _numBuckets) { 861 buckets = _buckets; 862 keyBytes = _keyBytes; 863 numHashes = _numHashes; 864 numBuckets = _numBuckets; 865 866 // recordIDBytes + keyBytes + signatureByte 867 // No event time here for dimension table join 868 int cellBytes = 8 + keyBytes + 1; 869 bucketBytes = HASH_BUCKET_SIZE * cellBytes; 870 int totalBucketBytes = bucketBytes * numBuckets; 871 stash = buckets + totalBucketBytes; 872 for (int i = 0; i < _numHashes; i++) { 873 seeds[i] = _seeds[i]; 874 } 875 876 offsetToSignature = HASH_BUCKET_SIZE * 8; 877 offsetToKey = offsetToSignature + HASH_BUCKET_SIZE * 1; 878 } 879 880 __host__ __device__ 881 uint8_t getSignature(uint8_t *bucket, int index) const { 882 return bucket[offsetToSignature + index]; 883 } 884 885 __host__ __device__ 886 uint8_t *getKey(uint8_t *bucket, int index) const { 887 return bucket + (offsetToKey + index * keyBytes); 888 } 889 890 __host__ __device__ 891 RecordID getRecordID(uint8_t *bucket, 892 int index) const { 893 return reinterpret_cast<RecordID *>(bucket)[index]; 894 } 895 896 // Note: RecordID{0,0} is used to represent unfound record, 897 // for all live batch ids are larger than 0 898 // all archive batch ids (epoch date) are larger than 0 899 __host__ __device__ 900 RecordID operator()(const argument_type t) const { 901 if (!thrust::get<1>(t)) { 902 RecordID recordID = {0, 0}; 903 return recordID; 904 } 905 906 I v = thrust::get<0>(t); 907 uint8_t *key = reinterpret_cast<uint8_t *>(&v); 908 for (int i = 0; i < numHashes; i++) { 909 uint32_t hashValue = murmur3sum32(key, keyBytes, seeds[i]); 910 int bucketIndex = hashValue % numBuckets; 911 uint8_t *bucket = buckets + bucketIndex * bucketBytes; 912 uint8_t signature = (uint8_t)(hashValue >> 24); 913 if (signature < 1) { 914 signature = 1; 915 } 916 917 for (int j = 0; j < HASH_BUCKET_SIZE; j++) { 918 if (signature == getSignature(bucket, j) && 919 memequal(getKey(bucket, j), key, keyBytes)) { 920 return getRecordID(bucket, j); 921 } 922 } 923 } 924 925 for (int j = 0; j < HASH_STASH_SIZE; j++) { 926 if (getSignature(stash, j) != 0 && 927 memequal(getKey(stash, j), key, keyBytes)) { 928 return getRecordID(stash, j); 929 } 930 } 931 932 RecordID recordID = {0, 0}; 933 return recordID; 934 } 935 }; 936 937 // ReduceByHashFunctor is the binaryOp for reduce dimIndexVector and values 938 // together, according to hash vector as key for reduction AssociativeOperator 939 // is the binary operator for reducing values eg. we have: 940 // hashVector: [1,1,2,2] 941 // dimIndexVector: [3,1,2,0] 942 // valueVector: [1.0,2.0,3.0,4.0] 943 // if we have AssociativeOperator of thrust::plus<uint32_t> 944 // then we have value_type of uint32_t 945 // the result of dimIndexVector is [3,2] and 946 // the result of valueVector is [3.0,7.0] 947 template<typename AssociativeOperator> 948 struct ReduceByHashFunctor { 949 typedef typename AssociativeOperator::first_argument_type value_type; 950 AssociativeOperator binaryOp; 951 952 __host__ __device__ 953 explicit ReduceByHashFunctor(AssociativeOperator binaryOp) 954 : binaryOp(binaryOp) {} 955 956 __host__ __device__ 957 thrust::tuple<uint32_t, value_type> operator()( 958 const thrust::tuple<uint32_t, value_type> t1, 959 const thrust::tuple<uint32_t, value_type> t2) const { 960 value_type ret = binaryOp(thrust::get<1>(t1), thrust::get<1>(t2)); 961 return thrust::make_tuple(thrust::get<0>(t1), ret); 962 } 963 }; 964 965 // HLLHashFunctor calculate hash value combining dimension 64bit hash 966 // with reg_id of hll value The final 64bit will use higher 48bit of the 967 // dimension hash and lower 16bit of hll value 968 struct HLLHashFunctor { 969 __host__ __device__ 970 uint64_t operator()(uint64_t hash, 971 uint32_t hllValue) const { 972 return (hash & 0xFFFFFFFFFFFF0000) | (hllValue & 0x3FFF); 973 } 974 }; 975 976 // HLLDimNotEqualFunctor only compares the first 48bit 977 // of the hash to determine whether two hashes represent the same dimension 978 struct HLLDimNotEqualFunctor { 979 __host__ __device__ 980 bool operator()(uint64_t v1, uint64_t v2) const { 981 return (v1 >> 16) ^ (v2 >> 16); 982 } 983 }; 984 985 // HLLMergeComparator 986 // first sort by hash value, then break tie by hll value 987 struct HLLMergeComparator { 988 __host__ __device__ 989 bool operator()( 990 thrust::tuple<uint64_t, uint32_t> t1, 991 thrust::tuple<uint64_t, uint32_t> t2) const { 992 if (thrust::get<0>(t1) == thrust::get<0>(t2)) { 993 return thrust::get<1>(t1) > thrust::get<1>(t2); 994 } 995 return thrust::get<0>(t1) < thrust::get<0>(t2); 996 } 997 }; 998 999 // HLLDimByteCountFunctor 1000 struct HLLDimByteCountFunctor { 1001 __host__ __device__ 1002 uint64_t operator()(uint16_t regCount) const { 1003 if (regCount < HLL_DENSE_THRESHOLD) { 1004 return (uint64_t) regCount * 4; 1005 } else { 1006 return HLL_DENSE_SIZE; 1007 } 1008 } 1009 }; 1010 1011 template<typename I, typename O> 1012 struct CastFunctor { 1013 __host__ __device__ 1014 O operator()(I in) const { 1015 return static_cast<O>(in); 1016 } 1017 }; 1018 1019 // CopyHLLFunctor 1020 struct CopyHLLFunctor { 1021 uint8_t *hllVector; 1022 __host__ __device__ 1023 explicit CopyHLLFunctor(uint8_t *hllVector) 1024 : hllVector(hllVector) {} 1025 1026 __host__ __device__ 1027 int operator()( 1028 thrust::tuple<uint64_t, uint32_t, int> t) const { 1029 uint64_t offset = thrust::get<0>(t); 1030 uint32_t value = thrust::get<1>(t); 1031 uint16_t regID = static_cast<uint16_t>(value & 0x3FFF); 1032 // rho must plus 1 1033 uint8_t rho = static_cast<uint8_t>((value >> 16) & 0xFF) + 1; 1034 int numBytes = thrust::get<2>(t); 1035 if (numBytes == 4) { 1036 *reinterpret_cast<uint32_t *>(hllVector + offset) = 1037 static_cast<uint32_t>(rho << 16 | regID); 1038 } else { 1039 *reinterpret_cast<uint8_t *>(hllVector + offset) = rho; 1040 } 1041 return 0; 1042 } 1043 }; 1044 1045 // dateutils 1046 1047 enum TimeBucketizer { 1048 YEAR, 1049 QUATER, 1050 MONTH, 1051 DAY_OF_MONTH, 1052 DAY_OF_YEAR, 1053 MONTH_OF_YEAR, 1054 QUARTER_OF_YEAR, 1055 }; 1056 1057 // resolveTimeBucketizer returns the start timestamp of a time that ts 1058 // represents if the timeBucketizer is a time series bucketizer ( 1059 // YEAR/QUARTER/MONTH). If the time bucketizer is a recurring time bucketizer( 1060 // DAY_OF_MONTH/DAY_OF_YEAR/MONTH_OF_YEAR/QUARTER_OF_YEAR), it will return the 1061 // number of units of the bucketizer. Note daysBeforeMonth need to be passed by 1062 // caller so that the device/host logic is determined outside of this function. 1063 __host__ __device__ 1064 uint32_t 1065 resolveTimeBucketizer(int64_t ts, enum TimeBucketizer timeBucketizer, 1066 const uint16_t *daysBeforeMonth); 1067 1068 __host__ __device__ 1069 uint32_t 1070 getWeekStartTimestamp(uint32_t ts); 1071 1072 // Forward declaration of the struct. 1073 struct EmptyStruct; 1074 1075 // VoidFunctor is the functor to take empty struct as argument 1076 // and produce no result. Notice it does not occupy any memory space. 1077 struct VoidFunctor { 1078 __host__ __device__ 1079 void operator()(EmptyStruct) const { 1080 } 1081 }; 1082 1083 } // namespace ares 1084 #endif // QUERY_FUNCTOR_HPP_