Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Clean up xxhash_64 implementations #17455

Merged
merged 5 commits into from
Nov 28, 2024
Merged

Clean up xxhash_64 implementations #17455

Changes from 4 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
70e2497
Updates
PointKernel Nov 26, 2024
95fe1f6
Cleanups
PointKernel Nov 26, 2024
6b721da
Fix ctor
PointKernel Nov 26, 2024
2b41edb
Merge branch 'branch-25.02' into improve-xxhash64
mhaseeb123 Nov 27, 2024
757d517
Merge remote-tracking branch 'upstream/branch-25.02' into improve-xxh…
PointKernel Nov 27, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
47 changes: 25 additions & 22 deletions cpp/include/cudf/hashing/detail/xxhash_64.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,9 +16,9 @@

#pragma once

#include "hash_functions.cuh"

#include <cudf/fixed_point/fixed_point.hpp>
#include <cudf/hashing.hpp>
#include <cudf/hashing/detail/hash_functions.cuh>
#include <cudf/strings/string_view.cuh>
#include <cudf/types.hpp>

Expand All @@ -28,72 +28,75 @@
namespace cudf::hashing::detail {

template <typename Key>
struct XXHash_64 : public cuco::xxhash_64<Key> {
using result_type = typename cuco::xxhash_64<Key>::result_type;
struct XXHash_64 {
using result_type = std::uint64_t;

__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<Key>::operator()(key);
return this->_impl.compute_hash(bytes, size);
}

template <typename Extent>
__device__ result_type compute_hash(cuda::std::byte const* bytes, Extent size) const
private:
template <typename T>
__device__ constexpr result_type compute(T const& key) const
{
return cuco::xxhash_64<Key>::compute_hash(bytes, size);
return this->compute_bytes(reinterpret_cast<cuda::std::byte const*>(&key), sizeof(T));
}

cuco::xxhash_64<Key> _impl;
};

template <>
XXHash_64<bool>::result_type __device__ inline XXHash_64<bool>::operator()(bool const& key) const
{
return this->compute_hash(reinterpret_cast<cuda::std::byte const*>(&key), sizeof(key));
return this->compute(static_cast<uint8_t>(key));
}

template <>
XXHash_64<float>::result_type __device__ inline XXHash_64<float>::operator()(float const& key) const
{
return cuco::xxhash_64<float>::operator()(normalize_nans(key));
return this->compute(normalize_nans(key));
}

template <>
XXHash_64<double>::result_type __device__ inline XXHash_64<double>::operator()(
double const& key) const
{
return cuco::xxhash_64<double>::operator()(normalize_nans(key));
return this->compute(normalize_nans(key));
}

template <>
XXHash_64<cudf::string_view>::result_type
__device__ inline XXHash_64<cudf::string_view>::operator()(cudf::string_view const& key) const
{
return this->compute_hash(reinterpret_cast<cuda::std::byte const*>(key.data()), key.size_bytes());
return this->compute_bytes(reinterpret_cast<cuda::std::byte const*>(key.data()),
key.size_bytes());
}

template <>
XXHash_64<numeric::decimal32>::result_type
__device__ inline XXHash_64<numeric::decimal32>::operator()(numeric::decimal32 const& key) const
{
auto const val = key.value();
auto const len = sizeof(val);
return this->compute_hash(reinterpret_cast<cuda::std::byte const*>(&val), len);
return this->compute(key.value());
}

template <>
XXHash_64<numeric::decimal64>::result_type
__device__ inline XXHash_64<numeric::decimal64>::operator()(numeric::decimal64 const& key) const
{
auto const val = key.value();
auto const len = sizeof(val);
return this->compute_hash(reinterpret_cast<cuda::std::byte const*>(&val), len);
return this->compute(key.value());
}

template <>
XXHash_64<numeric::decimal128>::result_type
__device__ inline XXHash_64<numeric::decimal128>::operator()(numeric::decimal128 const& key) const
{
auto const val = key.value();
auto const len = sizeof(val);
return this->compute_hash(reinterpret_cast<cuda::std::byte const*>(&val), len);
return this->compute(key.value());
}

} // namespace cudf::hashing::detail
Loading