| 516 | |
| 517 | template <typename Fn> |
| 518 | XGBOOST_HOST_DEV_INLINE decltype(auto) DispatchCall(Fn func) const { |
| 519 | using T = ArrayInterfaceHandler::Type; |
| 520 | switch (type) { |
| 521 | case T::kF2: { |
| 522 | #if defined(XGBOOST_USE_CUDA) |
| 523 | return func(reinterpret_cast<__half const *>(data)); |
| 524 | #endif // defined(XGBOOST_USE_CUDA) |
| 525 | } |
| 526 | case T::kF4: |
| 527 | return func(reinterpret_cast<float const *>(data)); |
| 528 | case T::kF8: |
| 529 | return func(reinterpret_cast<double const *>(data)); |
| 530 | #ifdef __CUDA_ARCH__ |
| 531 | case T::kF16: { |
| 532 | // CUDA device code doesn't support long double. |
| 533 | SPAN_CHECK(false); |
| 534 | return func(reinterpret_cast<double const *>(data)); |
| 535 | } |
| 536 | #else |
| 537 | case T::kF16: |
| 538 | return func(reinterpret_cast<long double const *>(data)); |
| 539 | #endif |
| 540 | case T::kI1: |
| 541 | return func(reinterpret_cast<int8_t const *>(data)); |
| 542 | case T::kI2: |
| 543 | return func(reinterpret_cast<int16_t const *>(data)); |
| 544 | case T::kI4: |
| 545 | return func(reinterpret_cast<int32_t const *>(data)); |
| 546 | case T::kI8: |
| 547 | return func(reinterpret_cast<int64_t const *>(data)); |
| 548 | case T::kU1: |
| 549 | return func(reinterpret_cast<uint8_t const *>(data)); |
| 550 | case T::kU2: |
| 551 | return func(reinterpret_cast<uint16_t const *>(data)); |
| 552 | case T::kU4: |
| 553 | return func(reinterpret_cast<uint32_t const *>(data)); |
| 554 | case T::kU8: |
| 555 | return func(reinterpret_cast<uint64_t const *>(data)); |
| 556 | } |
| 557 | SPAN_CHECK(false); |
| 558 | return func(reinterpret_cast<uint64_t const *>(data)); |
| 559 | } |
| 560 | |
| 561 | [[nodiscard]] XGBOOST_DEVICE std::size_t ElementSize() const { |
| 562 | return this->DispatchCall([](auto *typed_data_ptr) { |
no outgoing calls
no test coverage detected