diff --git a/cpp/include/cudf/hashing/detail/xxhash_64.cuh b/cpp/include/cudf/hashing/detail/xxhash_64.cuh index 49432ffd0a4..7d72349e340 100644 --- a/cpp/include/cudf/hashing/detail/xxhash_64.cuh +++ b/cpp/include/cudf/hashing/detail/xxhash_64.cuh @@ -28,77 +28,75 @@ namespace cudf::hashing::detail { template -struct XXHash_64 : public cuco::xxhash_64 { - using result_type = typename cuco::xxhash_64::result_type; +struct XXHash_64 { + using result_type = std::uint64_t; - __host__ __device__ constexpr XXHash_64(uint64_t seed = cudf::DEFAULT_HASH_SEED) - : cuco::xxhash_64{seed} - { - } + __host__ __device__ constexpr XXHash_64(uint64_t seed = cudf::DEFAULT_HASH_SEED) : _impl{seed} {} - __device__ result_type operator()(Key const& key) const + __device__ constexpr result_type operator()(Key const& key) const { return this->_impl(key); } + + __device__ constexpr result_type compute_bytes(cuda::std::byte const* bytes, + std::uint64_t size) const { - return cuco::xxhash_64::operator()(key); + return this->_impl.compute_hash(bytes, size); } - template - __device__ result_type compute_hash(cuda::std::byte const* bytes, Extent size) const + private: + template + __device__ constexpr result_type compute(T const& key) const { - return cuco::xxhash_64::compute_hash(bytes, size); + return this->compute_bytes(reinterpret_cast(&key), sizeof(T)); } + + cuco::xxhash_64 _impl; }; template <> XXHash_64::result_type __device__ inline XXHash_64::operator()(bool const& key) const { - return this->compute_hash(reinterpret_cast(&key), sizeof(key)); + return this->compute(static_cast(key)); } template <> XXHash_64::result_type __device__ inline XXHash_64::operator()(float const& key) const { - return cuco::xxhash_64::operator()(normalize_nans(key)); + return this->compute(normalize_nans(key)); } template <> XXHash_64::result_type __device__ inline XXHash_64::operator()( double const& key) const { - return cuco::xxhash_64::operator()(normalize_nans(key)); + return this->compute(normalize_nans(key)); } template <> XXHash_64::result_type __device__ inline XXHash_64::operator()(cudf::string_view const& key) const { - return this->compute_hash(reinterpret_cast(key.data()), key.size_bytes()); + return this->compute_bytes(reinterpret_cast(key.data()), + key.size_bytes()); } template <> XXHash_64::result_type __device__ inline XXHash_64::operator()(numeric::decimal32 const& key) const { - auto const val = key.value(); - auto const len = sizeof(val); - return this->compute_hash(reinterpret_cast(&val), len); + return this->compute(key.value()); } template <> XXHash_64::result_type __device__ inline XXHash_64::operator()(numeric::decimal64 const& key) const { - auto const val = key.value(); - auto const len = sizeof(val); - return this->compute_hash(reinterpret_cast(&val), len); + return this->compute(key.value()); } template <> XXHash_64::result_type __device__ inline XXHash_64::operator()(numeric::decimal128 const& key) const { - auto const val = key.value(); - auto const len = sizeof(val); - return this->compute_hash(reinterpret_cast(&val), len); + return this->compute(key.value()); } } // namespace cudf::hashing::detail