From 05868c97a1668c7940f99dfea402b9303b7f28c7 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Fri, 23 Jun 2023 16:17:59 -0400 Subject: [PATCH 01/19] Add XXHash_64 hash function to cudf --- cpp/CMakeLists.txt | 1 + cpp/include/cudf/detail/hashing.hpp | 5 + .../cudf/detail/utilities/hash_functions.cuh | 5 + cpp/include/cudf/hashing.hpp | 30 +- cpp/src/hash/hashing.cu | 52 ++- cpp/src/hash/xxhash64.cu | 323 ++++++++++++++++++ cpp/tests/hashing/hash_test.cpp | 21 ++ 7 files changed, 406 insertions(+), 31 deletions(-) create mode 100644 cpp/src/hash/xxhash64.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index d9f3824f706..6722744d44f 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -346,6 +346,7 @@ add_library( src/hash/md5_hash.cu src/hash/murmur_hash.cu src/hash/spark_murmur_hash.cu + src/hash/xxhash64.cu src/interop/dlpack.cpp src/interop/from_arrow.cu src/interop/to_arrow.cu diff --git a/cpp/include/cudf/detail/hashing.hpp b/cpp/include/cudf/detail/hashing.hpp index 771b3e150ec..3cd2e4fbb14 100644 --- a/cpp/include/cudf/detail/hashing.hpp +++ b/cpp/include/cudf/detail/hashing.hpp @@ -51,6 +51,11 @@ std::unique_ptr md5_hash(table_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr); +std::unique_ptr xxhash64(table_view const& input, + uint64_t seed, + rmm::cuda_stream_view, + rmm::mr::device_memory_resource* mr); + /* Copyright 2005-2014 Daniel James. * * Use, modification and distribution is subject to the Boost Software diff --git a/cpp/include/cudf/detail/utilities/hash_functions.cuh b/cpp/include/cudf/detail/utilities/hash_functions.cuh index e57822f3fdb..102c735b579 100644 --- a/cpp/include/cudf/detail/utilities/hash_functions.cuh +++ b/cpp/include/cudf/detail/utilities/hash_functions.cuh @@ -65,6 +65,11 @@ __device__ inline uint32_t rotate_bits_left(uint32_t x, uint32_t r) return __funnelshift_l(x, x, r); } +__device__ inline uint64_t rotate_bits_left(uint64_t h, int8_t r) +{ + return ((h << r) | (h >> (64 - r))); +} + __device__ inline uint32_t rotate_bits_right(uint32_t x, uint32_t r) { // This function is equivalent to (x >> r) | (x << (32 - r)) diff --git a/cpp/include/cudf/hashing.hpp b/cpp/include/cudf/hashing.hpp index d8b31776cc8..c23ab2a5e0e 100644 --- a/cpp/include/cudf/hashing.hpp +++ b/cpp/include/cudf/hashing.hpp @@ -22,7 +22,8 @@ namespace cudf { -using hash_value_type = uint32_t; ///< Type of hash value +using hash_value_type = uint32_t; ///< Type of hash 32-bit value +using hash64_value_type = uint64_t; ///< Type of hash 64-bit value /** * @addtogroup column_hash @@ -40,6 +41,13 @@ enum class hash_id { HASH_MD5 ///< MD5 hash function }; +/** + * @brief Identifies the 64-bit hash function to be used + */ +enum class hash64_id { + XXHASH_64 ///< XXHash_64 hash function +}; + /** * @brief The default seed value for hash functions */ @@ -63,5 +71,25 @@ std::unique_ptr hash( rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +/** + * @brief Computes the hash value of each row in the input set of columns + * + * This function takes a 64-bit seed value and returns a column of type UINT64. + * + * @param input The table of columns to hash + * @param hash_function The hash function enum to use + * @param seed Optional seed value to use for the hash function + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned column's device memory + * + * @returns A column where each row is the hash of a column from the input + */ +std::unique_ptr hash64( + table_view const& input, + hash64_id hash_function = hash64_id::XXHASH_64, + uint64_t seed = DEFAULT_HASH_SEED, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + /** @} */ // end of group } // namespace cudf diff --git a/cpp/src/hash/hashing.cu b/cpp/src/hash/hashing.cu index 66c9a9a70bb..0b75ad05fa7 100644 --- a/cpp/src/hash/hashing.cu +++ b/cpp/src/hash/hashing.cu @@ -13,44 +13,14 @@ * See the License for the specific language governing permissions and * limitations under the License. */ -#include #include #include -#include -#include -#include -#include #include #include -#include - -#include -#include - -#include namespace cudf { namespace detail { -namespace { - -template -std::vector to_leaf_columns(IterType iter_begin, IterType iter_end) -{ - std::vector leaf_columns; - std::for_each(iter_begin, iter_end, [&leaf_columns](column_view const& col) { - if (is_nested(col.type())) { - CUDF_EXPECTS(col.type().id() == type_id::STRUCT, "unsupported nested type"); - auto child_columns = to_leaf_columns(col.child_begin(), col.child_end()); - leaf_columns.insert(leaf_columns.end(), child_columns.begin(), child_columns.end()); - } else { - leaf_columns.emplace_back(col); - } - }); - return leaf_columns; -} - -} // namespace std::unique_ptr hash(table_view const& input, hash_id hash_function, @@ -66,6 +36,18 @@ std::unique_ptr hash(table_view const& input, } } +std::unique_ptr hash64(table_view const& input, + hash64_id hash_function, + uint64_t seed, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + switch (hash_function) { + case (hash64_id::XXHASH_64): return xxhash64(input, seed, stream, mr); + default: CUDF_FAIL("Unsupported hash function."); + } +} + } // namespace detail std::unique_ptr hash(table_view const& input, @@ -78,4 +60,14 @@ std::unique_ptr hash(table_view const& input, return detail::hash(input, hash_function, seed, stream, mr); } +std::unique_ptr hash64(table_view const& input, + hash64_id hash_function, + uint64_t seed, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::hash64(input, hash_function, seed, stream, mr); +} + } // namespace cudf diff --git a/cpp/src/hash/xxhash64.cu b/cpp/src/hash/xxhash64.cu new file mode 100644 index 00000000000..92e7022cf6f --- /dev/null +++ b/cpp/src/hash/xxhash64.cu @@ -0,0 +1,323 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include +#include +#include +#include +#include + +#include +#include + +#include + +namespace cudf { +namespace detail { + +namespace { + +using hash_value_type = uint64_t; + +template +struct XXHash_64 { + using result_type = hash_value_type; + + constexpr XXHash_64() = default; + constexpr XXHash_64(hash_value_type seed) : m_seed(seed) {} + + __device__ inline uint32_t getblock32(std::byte const* data, cudf::size_type offset) const + { + // Read a 4-byte value from the data pointer as individual bytes for safe + // unaligned access (very likely for string types). + auto block = reinterpret_cast(data + offset); + return block[0] | (block[1] << 8) | (block[2] << 16) | (block[3] << 24); + } + + __device__ inline uint64_t getblock64(std::byte const* data, cudf::size_type offset) const + { + uint64_t result = getblock32(data, offset); + result = result << 32; + return result | getblock32(data, offset + 4); + } + + result_type __device__ inline operator()(Key const& key) const { return compute(key); } + + template + result_type __device__ inline compute(T const& key) const + { + return compute_bytes(reinterpret_cast(&key), sizeof(T)); + } + + result_type __device__ inline compute_remaining_bytes(std::byte const* data, + cudf::size_type nbytes, + cudf::size_type offset, + result_type h64) const + { + // remaining data can be processed in 8-byte chunks + if ((nbytes % 32) >= 8) { + for (; offset <= nbytes - 8; offset += 8) { + uint64_t k1 = getblock64(data, offset) * prime2; + + k1 = rotate_bits_left(k1, 31) * prime1; + h64 ^= k1; + h64 = rotate_bits_left(h64, 27) * prime1 + prime4; + } + } + + // remaining data can be processed in 4-byte chunks + if (((nbytes % 32) % 8) >= 4) { + for (; offset <= nbytes - 4; offset += 4) { + h64 ^= (getblock32(data, offset) & 0xffffffffull) * prime1; + h64 = rotate_bits_left(h64, 23) * prime2 + prime3; + } + } + + // and the rest + if (nbytes % 4) { + while (offset < nbytes) { + h64 += (static_cast(data[offset]) & 0xff) * prime5; + h64 = rotate_bits_left(h64, 11) * prime1; + ++offset; + } + } + return h64; + } + + result_type __device__ compute_bytes(std::byte const* data, cudf::size_type const nbytes) const + { + uint64_t offset = 0; + uint64_t h64; + // data can be processed in 32-byte chunks + if (nbytes >= 32) { + auto limit = nbytes - 32; + uint64_t v1 = m_seed + prime1 + prime2; + uint64_t v2 = m_seed + prime2; + uint64_t v3 = m_seed; + uint64_t v4 = m_seed - prime1; + + do { + // pipeline 4*8byte computations + v1 += getblock64(data, offset) * prime2; + v1 = rotate_bits_left(v1, 31); + v1 *= prime1; + offset += 8; + v2 += getblock64(data, offset) * prime2; + v2 = rotate_bits_left(v2, 31); + v2 *= prime1; + offset += 8; + v3 += getblock64(data, offset) * prime2; + v3 = rotate_bits_left(v3, 31); + v3 *= prime1; + offset += 8; + v4 += getblock64(data, offset) * prime2; + v4 = rotate_bits_left(v4, 31); + v4 *= prime1; + offset += 8; + } while (offset <= limit); + + h64 = rotate_bits_left(v1, 1) + rotate_bits_left(v2, 7) + rotate_bits_left(v3, 12) + + rotate_bits_left(v4, 18); + + v1 *= prime2; + v1 = rotate_bits_left(v1, 31); + v1 *= prime1; + h64 ^= v1; + h64 = h64 * prime1 + prime4; + + v2 *= prime2; + v2 = rotate_bits_left(v2, 31); + v2 *= prime1; + h64 ^= v2; + h64 = h64 * prime1 + prime4; + + v3 *= prime2; + v3 = rotate_bits_left(v3, 31); + v3 *= prime1; + h64 ^= v3; + h64 = h64 * prime1 + prime4; + + v4 *= prime2; + v4 = rotate_bits_left(v4, 31); + v4 *= prime1; + h64 ^= v4; + h64 = h64 * prime1 + prime4; + } else { + h64 = m_seed + prime5; + } + + h64 += nbytes; + + h64 = compute_remaining_bytes(data, nbytes, offset, h64); + + return finalize(h64); + } + + constexpr __host__ __device__ std::uint64_t finalize(std::uint64_t h) const noexcept + { + h ^= h >> 33; + h *= prime2; + h ^= h >> 29; + h *= prime3; + h ^= h >> 32; + return h; + } + + private: + hash_value_type m_seed{}; + static constexpr uint64_t prime1 = 11400714785074694791ul; + static constexpr uint64_t prime2 = 14029467366897019727ul; + static constexpr uint64_t prime3 = 1609587929392839161ul; + static constexpr uint64_t prime4 = 9650029242287828579ul; + static constexpr uint64_t prime5 = 2870177450012600261ul; +}; + +template <> +hash_value_type __device__ inline XXHash_64::operator()(bool const& key) const +{ + return compute(static_cast(key)); +} + +template <> +hash_value_type __device__ inline XXHash_64::operator()(float const& key) const +{ + return compute(detail::normalize_nans_and_zeros(key)); +} + +template <> +hash_value_type __device__ inline XXHash_64::operator()(double const& key) const +{ + return compute(detail::normalize_nans_and_zeros(key)); +} + +template <> +hash_value_type __device__ inline XXHash_64::operator()( + cudf::string_view const& key) const +{ + auto const data = reinterpret_cast(key.data()); + auto const len = key.size_bytes(); + return compute_bytes(data, len); +} + +template <> +hash_value_type __device__ inline XXHash_64::operator()( + numeric::decimal32 const& key) const +{ + return compute(key.value()); +} + +template <> +hash_value_type __device__ inline XXHash_64::operator()( + numeric::decimal64 const& key) const +{ + return compute(key.value()); +} + +template <> +hash_value_type __device__ inline XXHash_64::operator()( + numeric::decimal128 const& key) const +{ + return compute(key.value()); +} + +/** + * @brief Computes the hash value of a row in the given table. + * + * @tparam Nullate A cudf::nullate type describing whether to check for nulls. + */ +template +class device_row_hasher { + public: + device_row_hasher(Nullate nulls, table_device_view const& t, hash_value_type seed) + : _check_nulls(nulls), _table(t), _seed(seed) + { + } + + __device__ auto operator()(size_type row_index) const noexcept + { + return detail::accumulate( + _table.begin(), + _table.end(), + _seed, + [row_index, nulls = _check_nulls] __device__(auto hash, auto column) { + return cudf::type_dispatcher( + column.type(), element_hasher_adapter{}, column, row_index, nulls, hash); + }); + } + + /** + * @brief Computes the hash value of an element in the given column. + */ + class element_hasher_adapter { + public: + template ())> + __device__ hash_value_type operator()(column_device_view const& col, + size_type row_index, + Nullate const _check_nulls, + hash_value_type const _seed) const noexcept + { + if (_check_nulls && col.is_null(row_index)) { + return std::numeric_limits::max(); + } + auto const hasher = XXHash_64{_seed}; + return hasher(col.element(row_index)); + } + + template ())> + __device__ hash_value_type operator()(column_device_view const&, + size_type, + Nullate const, + hash_value_type const) const noexcept + { + CUDF_UNREACHABLE("Unsupported type for MurmurHash64"); + } + }; + + Nullate const _check_nulls; + table_device_view const _table; + hash_value_type const _seed; +}; + +} // namespace + +std::unique_ptr xxhash64(table_view const& input, + uint64_t seed, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + auto output = make_numeric_column(data_type(type_to_id()), + input.num_rows(), + mask_state::UNALLOCATED, + stream, + mr); + + // Return early if there's nothing to hash + if (input.num_columns() == 0 || input.num_rows() == 0) { return output; } + + bool const nullable = has_nulls(input); + auto const input_view = table_device_view::create(input, stream); + auto output_view = output->mutable_view(); + + // Compute the hash value for each row + thrust::tabulate(rmm::exec_policy(stream), + output_view.begin(), + output_view.end(), + device_row_hasher(nullable, *input_view, seed)); + + return output; +} + +} // namespace detail +} // namespace cudf diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp index 45dcb3f80a1..2edafe7428f 100644 --- a/cpp/tests/hashing/hash_test.cpp +++ b/cpp/tests/hashing/hash_test.cpp @@ -1106,4 +1106,25 @@ TYPED_TEST(MD5HashTestFloatTyped, TestListExtremes) CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view(), verbosity); } +template +class HashXX64TestTyped : public cudf::test::BaseFixture {}; + +TYPED_TEST_SUITE(HashXX64TestTyped, NumericTypesNoBools); + +TYPED_TEST(HashXX64TestTyped, TestNumeric) +{ + using T = TypeParam; + auto col1 = cudf::test::fixed_width_column_wrapper{ + {-1, -1, 0, 2, 22, 0, 11, 12, 116, 32, 0, 42, 7, 62, 1, -22, 0, 0}, + {1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0}}; + auto col2 = cudf::test::fixed_width_column_wrapper{ + {-1, -1, 0, 2, 22, 1, 11, 12, 116, 32, 0, 42, 7, 62, 1, -22, 1, -22}, + {1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0}}; + + auto const output1 = cudf::hash64(cudf::table_view({col1})); + auto const output2 = cudf::hash64(cudf::table_view({col2})); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view()); +} + CUDF_TEST_PROGRAM_MAIN() From f39bba0096fab8c3b8604a587c7e31a6df452fd7 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 26 Jun 2023 15:51:34 -0400 Subject: [PATCH 02/19] fix typo in constant --- cpp/src/hash/xxhash64.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/hash/xxhash64.cu b/cpp/src/hash/xxhash64.cu index 92e7022cf6f..b7ff5807f32 100644 --- a/cpp/src/hash/xxhash64.cu +++ b/cpp/src/hash/xxhash64.cu @@ -80,7 +80,7 @@ struct XXHash_64 { // remaining data can be processed in 4-byte chunks if (((nbytes % 32) % 8) >= 4) { for (; offset <= nbytes - 4; offset += 4) { - h64 ^= (getblock32(data, offset) & 0xffffffffull) * prime1; + h64 ^= (getblock32(data, offset) & 0xfffffffful) * prime1; h64 = rotate_bits_left(h64, 23) * prime2 + prime3; } } From f9436d30e628d60c61f71023367df626f3374ea5 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 27 Jun 2023 09:59:54 -0400 Subject: [PATCH 03/19] fix rotate function --- cpp/include/cudf/detail/utilities/hash_functions.cuh | 2 +- cpp/src/hash/spark_murmur_hash.cu | 6 +++--- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/include/cudf/detail/utilities/hash_functions.cuh b/cpp/include/cudf/detail/utilities/hash_functions.cuh index 102c735b579..74dcf45c0c0 100644 --- a/cpp/include/cudf/detail/utilities/hash_functions.cuh +++ b/cpp/include/cudf/detail/utilities/hash_functions.cuh @@ -65,7 +65,7 @@ __device__ inline uint32_t rotate_bits_left(uint32_t x, uint32_t r) return __funnelshift_l(x, x, r); } -__device__ inline uint64_t rotate_bits_left(uint64_t h, int8_t r) +__device__ inline uint64_t rotate_bits_left(uint64_t h, uint32_t r) { return ((h << r) | (h >> (64 - r))); } diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu index ed3bc2a3605..b94059f68a3 100644 --- a/cpp/src/hash/spark_murmur_hash.cu +++ b/cpp/src/hash/spark_murmur_hash.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -87,7 +87,7 @@ struct SparkMurmurHash3_32 { k1 = cudf::detail::rotate_bits_left(k1, rot_c1); k1 *= c2; h ^= k1; - h = cudf::detail::rotate_bits_left(h, rot_c2); + h = cudf::detail::rotate_bits_left(static_cast(h), rot_c2); h = h * 5 + c3; } return h; @@ -107,7 +107,7 @@ struct SparkMurmurHash3_32 { k1 = cudf::detail::rotate_bits_left(k1, rot_c1); k1 *= c2; h ^= k1; - h = cudf::detail::rotate_bits_left(h, rot_c2); + h = cudf::detail::rotate_bits_left(static_cast(h), rot_c2); h = h * 5 + c3; } From de00dc40dd4e595cb39c56d5bed0b479aa87f410 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 27 Jun 2023 14:52:17 -0400 Subject: [PATCH 04/19] add string test; convert primes to hex; fix getblock logic --- cpp/src/hash/xxhash64.cu | 14 ++++++------ cpp/tests/hashing/hash_test.cpp | 40 +++++++++++++++++++++++++++++++++ 2 files changed, 47 insertions(+), 7 deletions(-) diff --git a/cpp/src/hash/xxhash64.cu b/cpp/src/hash/xxhash64.cu index b7ff5807f32..3c4330b4e52 100644 --- a/cpp/src/hash/xxhash64.cu +++ b/cpp/src/hash/xxhash64.cu @@ -48,9 +48,9 @@ struct XXHash_64 { __device__ inline uint64_t getblock64(std::byte const* data, cudf::size_type offset) const { - uint64_t result = getblock32(data, offset); + uint64_t result = getblock32(data, offset + 4); result = result << 32; - return result | getblock32(data, offset + 4); + return result | getblock32(data, offset); } result_type __device__ inline operator()(Key const& key) const { return compute(key); } @@ -177,11 +177,11 @@ struct XXHash_64 { private: hash_value_type m_seed{}; - static constexpr uint64_t prime1 = 11400714785074694791ul; - static constexpr uint64_t prime2 = 14029467366897019727ul; - static constexpr uint64_t prime3 = 1609587929392839161ul; - static constexpr uint64_t prime4 = 9650029242287828579ul; - static constexpr uint64_t prime5 = 2870177450012600261ul; + static constexpr uint64_t prime1 = 0x9e3779b185ebca87ul; + static constexpr uint64_t prime2 = 0xc2b2ae3d27d4eb4ful; + static constexpr uint64_t prime3 = 0x165667b19e3779f9ul; + static constexpr uint64_t prime4 = 0x85ebca77c2b2ae63ul; + static constexpr uint64_t prime5 = 0x27d4eb2f165667c5ul; }; template <> diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp index 2edafe7428f..5f910c6cbf7 100644 --- a/cpp/tests/hashing/hash_test.cpp +++ b/cpp/tests/hashing/hash_test.cpp @@ -1127,4 +1127,44 @@ TYPED_TEST(HashXX64TestTyped, TestNumeric) CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view()); } +class HashXX64Test : public cudf::test::BaseFixture {}; + +TEST_F(HashXX64Test, StringType) +{ + auto col1 = cudf::test::strings_column_wrapper( + {"The", + "quick", + "brown fox", + "jumps over the lazy dog.", + "I am Jack's complete lack of null value", + "A very long (greater than 128 bytes/char string) to test a a very long string", + "Some multi-byte characters here: ééé", + "ééé", + "ééé ééé", + "ééé ééé ééé ééé", + "", + "!@#$%^&*(())", + "0123456789", + "{}|:<>?,./;[]=-"}); + + auto output = cudf::hash64(cudf::table_view({col1})); + + // these were generated using the CPU compiled version of the cuco xxhash64 source + auto expected = cudf::test::fixed_width_column_wrapper({11648823711624848724ul, + 10848020664967373619ul, + 14871996948511285677ul, + 17291005374665645904ul, + 12382650615907311857ul, + 11838841145970261371ul, + 3765709498017562316ul, + 8794998527445624677ul, + 17430199718277149821ul, + 16032571691393850730ul, + 17241709254077376921ul, + 7379359170906687646ul, + 71989731308565429ul, + 11645213193975881208ul}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected); +} + CUDF_TEST_PROGRAM_MAIN() From 660357eeafe8ef9701518e88a2fd997480b6f881 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 27 Jun 2023 22:39:22 -0400 Subject: [PATCH 05/19] fully-qualify calls to detail functions --- cpp/src/hash/xxhash64.cu | 34 ++++++++++++++++----------------- cpp/tests/hashing/hash_test.cpp | 1 + 2 files changed, 18 insertions(+), 17 deletions(-) diff --git a/cpp/src/hash/xxhash64.cu b/cpp/src/hash/xxhash64.cu index 3c4330b4e52..fd5d95f0530 100644 --- a/cpp/src/hash/xxhash64.cu +++ b/cpp/src/hash/xxhash64.cu @@ -71,9 +71,9 @@ struct XXHash_64 { for (; offset <= nbytes - 8; offset += 8) { uint64_t k1 = getblock64(data, offset) * prime2; - k1 = rotate_bits_left(k1, 31) * prime1; + k1 = cudf::detail::rotate_bits_left(k1, 31) * prime1; h64 ^= k1; - h64 = rotate_bits_left(h64, 27) * prime1 + prime4; + h64 = cudf::detail::rotate_bits_left(h64, 27) * prime1 + prime4; } } @@ -81,7 +81,7 @@ struct XXHash_64 { if (((nbytes % 32) % 8) >= 4) { for (; offset <= nbytes - 4; offset += 4) { h64 ^= (getblock32(data, offset) & 0xfffffffful) * prime1; - h64 = rotate_bits_left(h64, 23) * prime2 + prime3; + h64 = cudf::detail::rotate_bits_left(h64, 23) * prime2 + prime3; } } @@ -89,7 +89,7 @@ struct XXHash_64 { if (nbytes % 4) { while (offset < nbytes) { h64 += (static_cast(data[offset]) & 0xff) * prime5; - h64 = rotate_bits_left(h64, 11) * prime1; + h64 = cudf::detail::rotate_bits_left(h64, 11) * prime1; ++offset; } } @@ -111,46 +111,46 @@ struct XXHash_64 { do { // pipeline 4*8byte computations v1 += getblock64(data, offset) * prime2; - v1 = rotate_bits_left(v1, 31); + v1 = cudf::detail::rotate_bits_left(v1, 31); v1 *= prime1; offset += 8; v2 += getblock64(data, offset) * prime2; - v2 = rotate_bits_left(v2, 31); + v2 = cudf::detail::rotate_bits_left(v2, 31); v2 *= prime1; offset += 8; v3 += getblock64(data, offset) * prime2; - v3 = rotate_bits_left(v3, 31); + v3 = cudf::detail::rotate_bits_left(v3, 31); v3 *= prime1; offset += 8; v4 += getblock64(data, offset) * prime2; - v4 = rotate_bits_left(v4, 31); + v4 = cudf::detail::rotate_bits_left(v4, 31); v4 *= prime1; offset += 8; } while (offset <= limit); - h64 = rotate_bits_left(v1, 1) + rotate_bits_left(v2, 7) + rotate_bits_left(v3, 12) + - rotate_bits_left(v4, 18); + h64 = cudf::detail::rotate_bits_left(v1, 1) + cudf::detail::rotate_bits_left(v2, 7) + + cudf::detail::rotate_bits_left(v3, 12) + cudf::detail::rotate_bits_left(v4, 18); v1 *= prime2; - v1 = rotate_bits_left(v1, 31); + v1 = cudf::detail::rotate_bits_left(v1, 31); v1 *= prime1; h64 ^= v1; h64 = h64 * prime1 + prime4; v2 *= prime2; - v2 = rotate_bits_left(v2, 31); + v2 = cudf::detail::rotate_bits_left(v2, 31); v2 *= prime1; h64 ^= v2; h64 = h64 * prime1 + prime4; v3 *= prime2; - v3 = rotate_bits_left(v3, 31); + v3 = cudf::detail::rotate_bits_left(v3, 31); v3 *= prime1; h64 ^= v3; h64 = h64 * prime1 + prime4; v4 *= prime2; - v4 = rotate_bits_left(v4, 31); + v4 = cudf::detail::rotate_bits_left(v4, 31); v4 *= prime1; h64 ^= v4; h64 = h64 * prime1 + prime4; @@ -193,13 +193,13 @@ hash_value_type __device__ inline XXHash_64::operator()(bool const& key) c template <> hash_value_type __device__ inline XXHash_64::operator()(float const& key) const { - return compute(detail::normalize_nans_and_zeros(key)); + return compute(cudf::detail::normalize_nans_and_zeros(key)); } template <> hash_value_type __device__ inline XXHash_64::operator()(double const& key) const { - return compute(detail::normalize_nans_and_zeros(key)); + return compute(cudf::detail::normalize_nans_and_zeros(key)); } template <> @@ -247,7 +247,7 @@ class device_row_hasher { __device__ auto operator()(size_type row_index) const noexcept { - return detail::accumulate( + return cudf::detail::accumulate( _table.begin(), _table.end(), _seed, diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp index 5f910c6cbf7..6bc5ec57eb2 100644 --- a/cpp/tests/hashing/hash_test.cpp +++ b/cpp/tests/hashing/hash_test.cpp @@ -1150,6 +1150,7 @@ TEST_F(HashXX64Test, StringType) auto output = cudf::hash64(cudf::table_view({col1})); // these were generated using the CPU compiled version of the cuco xxhash64 source + // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh auto expected = cudf::test::fixed_width_column_wrapper({11648823711624848724ul, 10848020664967373619ul, 14871996948511285677ul, From a5a0d4bf49f82936737238276c9669b0e1155f53 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Thu, 29 Jun 2023 17:44:16 -0400 Subject: [PATCH 06/19] fix bug in xxhash finalize step --- cpp/src/hash/xxhash64.cu | 2 +- cpp/tests/hashing/hash_test.cpp | 22 +++++++++++----------- 2 files changed, 12 insertions(+), 12 deletions(-) diff --git a/cpp/src/hash/xxhash64.cu b/cpp/src/hash/xxhash64.cu index fd5d95f0530..a6e27e5a69f 100644 --- a/cpp/src/hash/xxhash64.cu +++ b/cpp/src/hash/xxhash64.cu @@ -88,7 +88,7 @@ struct XXHash_64 { // and the rest if (nbytes % 4) { while (offset < nbytes) { - h64 += (static_cast(data[offset]) & 0xff) * prime5; + h64 ^= (static_cast(data[offset]) & 0xff) * prime5; h64 = cudf::detail::rotate_bits_left(h64, 11) * prime1; ++offset; } diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp index 6bc5ec57eb2..c90cd974443 100644 --- a/cpp/tests/hashing/hash_test.cpp +++ b/cpp/tests/hashing/hash_test.cpp @@ -1151,20 +1151,20 @@ TEST_F(HashXX64Test, StringType) // these were generated using the CPU compiled version of the cuco xxhash64 source // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh - auto expected = cudf::test::fixed_width_column_wrapper({11648823711624848724ul, - 10848020664967373619ul, - 14871996948511285677ul, + auto expected = cudf::test::fixed_width_column_wrapper({4686269239494003989ul, + 6715983472207430822ul, + 8148134898123095730ul, 17291005374665645904ul, - 12382650615907311857ul, - 11838841145970261371ul, - 3765709498017562316ul, - 8794998527445624677ul, - 17430199718277149821ul, - 16032571691393850730ul, + 2631835514925512071ul, + 17518648592268952189ul, + 8749004388517322364ul, + 17701789113925815768ul, + 8612485687958712810ul, + 5148645515269989956ul, 17241709254077376921ul, 7379359170906687646ul, - 71989731308565429ul, - 11645213193975881208ul}); + 4566581271137380327ul, + 17962149534752128981ul}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected); } From aaafd8e2dbea66111549936e38e79a2da8fd457f Mon Sep 17 00:00:00 2001 From: David Wendt Date: Fri, 30 Jun 2023 20:22:00 -0400 Subject: [PATCH 07/19] add gtests for integer, double, fixed-point --- cpp/src/hash/xxhash64.cu | 4 +- cpp/tests/CMakeLists.txt | 2 +- cpp/tests/hashing/hash_test.cpp | 62 ----------- cpp/tests/hashing/xxhash64_test.cpp | 167 ++++++++++++++++++++++++++++ 4 files changed, 170 insertions(+), 65 deletions(-) create mode 100644 cpp/tests/hashing/xxhash64_test.cpp diff --git a/cpp/src/hash/xxhash64.cu b/cpp/src/hash/xxhash64.cu index a6e27e5a69f..c367e933a4a 100644 --- a/cpp/src/hash/xxhash64.cu +++ b/cpp/src/hash/xxhash64.cu @@ -193,13 +193,13 @@ hash_value_type __device__ inline XXHash_64::operator()(bool const& key) c template <> hash_value_type __device__ inline XXHash_64::operator()(float const& key) const { - return compute(cudf::detail::normalize_nans_and_zeros(key)); + return compute(cudf::detail::normalize_nans(key)); } template <> hash_value_type __device__ inline XXHash_64::operator()(double const& key) const { - return compute(cudf::detail::normalize_nans_and_zeros(key)); + return compute(cudf::detail::normalize_nans(key)); } template <> diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index f51a1292a48..ddcc9fb4885 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -163,7 +163,7 @@ ConfigureTest(DATETIME_OPS_TEST datetime/datetime_ops_test.cpp) # ################################################################################################## # * hashing tests --------------------------------------------------------------------------------- -ConfigureTest(HASHING_TEST hashing/hash_test.cpp) +ConfigureTest(HASHING_TEST hashing/hash_test.cpp hashing/xxhash64_test.cpp) # ################################################################################################## # * partitioning tests ---------------------------------------------------------------------------- diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp index c90cd974443..45dcb3f80a1 100644 --- a/cpp/tests/hashing/hash_test.cpp +++ b/cpp/tests/hashing/hash_test.cpp @@ -1106,66 +1106,4 @@ TYPED_TEST(MD5HashTestFloatTyped, TestListExtremes) CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view(), verbosity); } -template -class HashXX64TestTyped : public cudf::test::BaseFixture {}; - -TYPED_TEST_SUITE(HashXX64TestTyped, NumericTypesNoBools); - -TYPED_TEST(HashXX64TestTyped, TestNumeric) -{ - using T = TypeParam; - auto col1 = cudf::test::fixed_width_column_wrapper{ - {-1, -1, 0, 2, 22, 0, 11, 12, 116, 32, 0, 42, 7, 62, 1, -22, 0, 0}, - {1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0}}; - auto col2 = cudf::test::fixed_width_column_wrapper{ - {-1, -1, 0, 2, 22, 1, 11, 12, 116, 32, 0, 42, 7, 62, 1, -22, 1, -22}, - {1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0}}; - - auto const output1 = cudf::hash64(cudf::table_view({col1})); - auto const output2 = cudf::hash64(cudf::table_view({col2})); - - CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view()); -} - -class HashXX64Test : public cudf::test::BaseFixture {}; - -TEST_F(HashXX64Test, StringType) -{ - auto col1 = cudf::test::strings_column_wrapper( - {"The", - "quick", - "brown fox", - "jumps over the lazy dog.", - "I am Jack's complete lack of null value", - "A very long (greater than 128 bytes/char string) to test a a very long string", - "Some multi-byte characters here: ééé", - "ééé", - "ééé ééé", - "ééé ééé ééé ééé", - "", - "!@#$%^&*(())", - "0123456789", - "{}|:<>?,./;[]=-"}); - - auto output = cudf::hash64(cudf::table_view({col1})); - - // these were generated using the CPU compiled version of the cuco xxhash64 source - // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh - auto expected = cudf::test::fixed_width_column_wrapper({4686269239494003989ul, - 6715983472207430822ul, - 8148134898123095730ul, - 17291005374665645904ul, - 2631835514925512071ul, - 17518648592268952189ul, - 8749004388517322364ul, - 17701789113925815768ul, - 8612485687958712810ul, - 5148645515269989956ul, - 17241709254077376921ul, - 7379359170906687646ul, - 4566581271137380327ul, - 17962149534752128981ul}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected); -} - CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/hashing/xxhash64_test.cpp b/cpp/tests/hashing/xxhash64_test.cpp new file mode 100644 index 00000000000..29a6b4b0866 --- /dev/null +++ b/cpp/tests/hashing/xxhash64_test.cpp @@ -0,0 +1,167 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include + +#include +#include +#include +#include +#include + +using NumericTypesNoBools = + cudf::test::Concat; + +template +class HashXX64TestTyped : public cudf::test::BaseFixture {}; + +TYPED_TEST_SUITE(HashXX64TestTyped, NumericTypesNoBools); + +TYPED_TEST(HashXX64TestTyped, TestAllNumeric) +{ + using T = TypeParam; + auto col1 = cudf::test::fixed_width_column_wrapper{ + {-1, -1, 0, 2, 22, 0, 11, 12, 116, 32, 0, 42, 7, 62, 1, -22, 0, 0}, + {1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0}}; + auto col2 = cudf::test::fixed_width_column_wrapper{ + {-1, -1, 0, 2, 22, 1, 11, 12, 116, 32, 0, 42, 7, 62, 1, -22, 1, -22}, + {1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0}}; + + auto const output1 = cudf::hash64(cudf::table_view({col1})); + auto const output2 = cudf::hash64(cudf::table_view({col2})); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view()); +} + +class HashXX64Test : public cudf::test::BaseFixture {}; + +TEST_F(HashXX64Test, TestInteger) +{ + auto col1 = + cudf::test::fixed_width_column_wrapper{{-127, + -70000, + 0, + 200000, + 128, + std::numeric_limits::max(), + std::numeric_limits::min(), + std::numeric_limits::lowest()}}; + + auto const output = cudf::hash64(cudf::table_view({col1})); + + // these were generated using the CPU compiled version of the cuco xxhash64 source + // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh + auto expected = cudf::test::fixed_width_column_wrapper({4827426872506142937ul, + 13867166853951622683ul, + 4246796580750024372ul, + 17339819992360460003ul, + 7292178400482025765ul, + 2971168436322821236ul, + 9380524276503839603ul, + 9380524276503839603ul}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected); +} + +TEST_F(HashXX64Test, TestDouble) +{ + auto col1 = + cudf::test::fixed_width_column_wrapper{{-127., + -70000.125, + 0.0, + 200000.5, + 128.5, + -0.0, + std::numeric_limits::infinity(), + std::numeric_limits::quiet_NaN(), + std::numeric_limits::max(), + std::numeric_limits::min(), + std::numeric_limits::lowest()}}; + + auto const output = cudf::hash64(cudf::table_view({col1})); + + // these were generated using the CPU compiled version of the cuco xxhash64 source + // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh + auto expected = cudf::test::fixed_width_column_wrapper({16892115221677838993ul, + 1686446903308179321ul, + 3803688792395291579ul, + 18250447068822614389ul, + 3511911086082166358ul, + 4558309869707674848ul, + 18031741628920313605ul, + 16838308782748609196ul, + 3127544388062992779ul, + 1692401401506680154ul, + 13770442912356326755ul}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected); +} + +TEST_F(HashXX64Test, StringType) +{ + auto col1 = cudf::test::strings_column_wrapper( + {"The", + "quick", + "brown fox", + "jumps over the lazy dog.", + "I am Jack's complete lack of null value", + "A very long (greater than 128 bytes/char string) to test a a very long string", + "Some multi-byte characters here: ééé", + "ééé", + "ééé ééé", + "ééé ééé ééé ééé", + "", + "!@#$%^&*(())", + "0123456789", + "{}|:<>?,./;[]=-"}); + + auto output = cudf::hash64(cudf::table_view({col1})); + + // these were generated using the CPU compiled version of the cuco xxhash64 source + // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh + auto expected = cudf::test::fixed_width_column_wrapper({4686269239494003989ul, + 6715983472207430822ul, + 8148134898123095730ul, + 17291005374665645904ul, + 2631835514925512071ul, + 17518648592268952189ul, + 8749004388517322364ul, + 17701789113925815768ul, + 8612485687958712810ul, + 5148645515269989956ul, + 17241709254077376921ul, + 7379359170906687646ul, + 4566581271137380327ul, + 17962149534752128981ul}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected); +} + +TEST_F(HashXX64Test, TestFixedPoint) +{ + auto const col1 = cudf::test::fixed_point_column_wrapper( + {0, 100, -100, -999999999, 999999999}, numeric::scale_type{-3}); + auto const output = cudf::hash64(cudf::table_view({col1})); + + // these were generated using the CPU compiled version of the cuco xxhash64 source + // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh + // and passing the 'value' of each input (without the scale) as the decimal-type + auto expected = cudf::test::fixed_width_column_wrapper({4246796580750024372ul, + 5959467639951725378ul, + 4122185689695768261ul, + 3249245648192442585ul, + 8009575895491381648ul}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected); +} From 3b4dbb8aa668118c5af5b765b0a7a145bde5f358 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 10 Jul 2023 12:53:30 -0400 Subject: [PATCH 08/19] rename hash64 to xxhash64 --- cpp/include/cudf/hashing.hpp | 43 ++++++++++++----------------- cpp/src/hash/hashing.cu | 23 --------------- cpp/src/hash/xxhash64.cu | 12 +++++++- cpp/tests/hashing/xxhash64_test.cpp | 15 ++++++---- 4 files changed, 37 insertions(+), 56 deletions(-) diff --git a/cpp/include/cudf/hashing.hpp b/cpp/include/cudf/hashing.hpp index 9f8ed1b37bf..0d5d05efe34 100644 --- a/cpp/include/cudf/hashing.hpp +++ b/cpp/include/cudf/hashing.hpp @@ -28,8 +28,7 @@ namespace cudf { * @file */ -using hash_value_type = uint32_t; ///< Type of hash 32-bit value -using hash64_value_type = uint64_t; ///< Type of hash 64-bit value +using hash_value_type = uint32_t; ///< Type of hash 32-bit value /** * @brief Identifies the hash function to be used @@ -42,13 +41,6 @@ enum class hash_id { HASH_MD5 ///< MD5 hash function }; -/** - * @brief Identifies the 64-bit hash function to be used - */ -enum class hash64_id { - XXHASH_64 ///< XXHash_64 hash function -}; - /** * @brief The default seed value for hash functions */ @@ -97,54 +89,53 @@ std::unique_ptr murmur_hash3_32( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Computes the hash value of each row in the input set of columns + * @brief Computes the MurmurHash3 32-bit of each row in the given table * - * This function takes a 64-bit seed value and returns a column of type UINT64. + * This function computes the hash similar to MurmurHash3_32 with special processing + * to match Spark's implementation results. * * @param input The table of columns to hash - * @param hash_function The hash function enum to use + * @param seed Optional seed value to use for the hash function * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource used to allocate the returned column's device memory * - * @returns A column where each row is the hash of a row from the input + * @returns A column where each row is the hash of a column from the input */ -std::unique_ptr hash64( +std::unique_ptr spark_murmur_hash3_32( table_view const& input, - hash64_id hash_function = hash64_id::XXHASH_64, - uint64_t seed = DEFAULT_HASH_SEED, + uint32_t seed = DEFAULT_HASH_SEED, rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Computes the MurmurHash3 32-bit of each row in the given table - * - * This function computes the hash similar to MurmurHash3_32 with special processing - * to match Spark's implementation results. + * @brief Computes the MD5 hash of each row in the given table * * @param input The table of columns to hash - * @param seed Optional seed value to use for the hash function * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource used to allocate the returned column's device memory * - * @returns A column where each row is the hash of a column from the input + * @returns A column where each row is the hash of a row from the input */ -std::unique_ptr spark_murmur_hash3_32( +std::unique_ptr md5( table_view const& input, - uint32_t seed = DEFAULT_HASH_SEED, rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Computes the MD5 hash of each row in the given table + * @brief Computes the hash value of each row in the input set of columns + * + * This function takes a 64-bit seed value and returns a column of type UINT64. * * @param input The table of columns to hash + * @param seed Optional seed value to use for the hash function * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource used to allocate the returned column's device memory * * @returns A column where each row is the hash of a row from the input */ -std::unique_ptr md5( +std::unique_ptr xxhash64( table_view const& input, + uint64_t seed = DEFAULT_HASH_SEED, rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); diff --git a/cpp/src/hash/hashing.cu b/cpp/src/hash/hashing.cu index 5d7c42d3574..c5a2d479e29 100644 --- a/cpp/src/hash/hashing.cu +++ b/cpp/src/hash/hashing.cu @@ -37,18 +37,6 @@ std::unique_ptr hash(table_view const& input, } } -std::unique_ptr hash64(table_view const& input, - hash64_id hash_function, - uint64_t seed, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - switch (hash_function) { - case (hash64_id::XXHASH_64): return xxhash64(input, seed, stream, mr); - default: CUDF_FAIL("Unsupported hash function."); - } -} - } // namespace detail std::unique_ptr murmur_hash3_32(table_view const& input, @@ -77,16 +65,6 @@ std::unique_ptr md5(table_view const& input, return detail::md5(input, stream, mr); } -std::unique_ptr hash64(table_view const& input, - hash64_id hash_function, - uint64_t seed, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - CUDF_FUNC_RANGE(); - return hashing::detail::hash64(input, hash_function, seed, stream, mr); -} - } // namespace hashing std::unique_ptr hash(table_view const& input, @@ -100,4 +78,3 @@ std::unique_ptr hash(table_view const& input, } } // namespace cudf - diff --git a/cpp/src/hash/xxhash64.cu b/cpp/src/hash/xxhash64.cu index 40041bbd334..c27baf1a85e 100644 --- a/cpp/src/hash/xxhash64.cu +++ b/cpp/src/hash/xxhash64.cu @@ -15,6 +15,7 @@ */ #include #include +#include #include #include #include @@ -321,6 +322,15 @@ std::unique_ptr xxhash64(table_view const& input, } } // namespace detail + +std::unique_ptr xxhash64(table_view const& input, + uint64_t seed, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::xxhash64(input, seed, stream, mr); +} + } // namespace hashing } // namespace cudf - diff --git a/cpp/tests/hashing/xxhash64_test.cpp b/cpp/tests/hashing/xxhash64_test.cpp index fa390f08454..ea2c23b5009 100644 --- a/cpp/tests/hashing/xxhash64_test.cpp +++ b/cpp/tests/hashing/xxhash64_test.cpp @@ -42,9 +42,12 @@ TYPED_TEST(HashXX64TestTyped, TestAllNumeric) {-1, -1, 0, 2, 22, 1, 11, 12, 116, 32, 0, 42, 7, 62, 1, -22, 1, -22}, {1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0}}; - auto const output1 = cudf::hashing::hash64(cudf::table_view({col1})); - auto const output2 = cudf::hashing::hash64(cudf::table_view({col2})); + auto output1 = cudf::hashing::xxhash64(cudf::table_view({col1})); + auto output2 = cudf::hashing::xxhash64(cudf::table_view({col2})); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view()); + output1 = cudf::hashing::xxhash64(cudf::table_view({col1}), 7); + output2 = cudf::hashing::xxhash64(cudf::table_view({col2}), 7); CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view()); } @@ -62,7 +65,7 @@ TEST_F(HashXX64Test, TestInteger) std::numeric_limits::min(), std::numeric_limits::lowest()}}; - auto const output = cudf::hashing::hash64(cudf::table_view({col1})); + auto const output = cudf::hashing::xxhash64(cudf::table_view({col1})); // these were generated using the CPU compiled version of the cuco xxhash64 source // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh @@ -92,7 +95,7 @@ TEST_F(HashXX64Test, TestDouble) std::numeric_limits::min(), std::numeric_limits::lowest()}}; - auto const output = cudf::hashing::hash64(cudf::table_view({col1})); + auto const output = cudf::hashing::xxhash64(cudf::table_view({col1})); // these were generated using the CPU compiled version of the cuco xxhash64 source // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh @@ -128,7 +131,7 @@ TEST_F(HashXX64Test, StringType) "0123456789", "{}|:<>?,./;[]=-"}); - auto output = cudf::hashing::hash64(cudf::table_view({col1})); + auto output = cudf::hashing::xxhash64(cudf::table_view({col1})); // these were generated using the CPU compiled version of the cuco xxhash64 source // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh @@ -153,7 +156,7 @@ TEST_F(HashXX64Test, TestFixedPoint) { auto const col1 = cudf::test::fixed_point_column_wrapper( {0, 100, -100, -999999999, 999999999}, numeric::scale_type{-3}); - auto const output = cudf::hashing::hash64(cudf::table_view({col1})); + auto const output = cudf::hashing::xxhash64(cudf::table_view({col1})); // these were generated using the CPU compiled version of the cuco xxhash64 source // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh From 1ce895d8263c8f2b9feeef3a65e39ee1d594c0ab Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 11 Jul 2023 16:48:06 -0400 Subject: [PATCH 09/19] fix cmake format style violation --- cpp/tests/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 324e184c1c4..080e73b2e6c 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -164,7 +164,7 @@ ConfigureTest(DATETIME_OPS_TEST datetime/datetime_ops_test.cpp) # ################################################################################################## # * hashing tests --------------------------------------------------------------------------------- ConfigureTest( - HASHING_TEST hashing/md5_test.cpp hashing/murmur3_test.cpp hashing/spark_murmur3_test.cpp hashing/xxhash64_test.cpp + HASHING_TEST hashing/md5_test.cpp hashing/murmur3_test.cpp hashing/spark_murmur3_test.cpp hashing/xxhash64_test.cpp ) # ################################################################################################## From 65b6bad599f8d00e04e57631ff13fb539eb6b71f Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 11 Jul 2023 17:02:40 -0400 Subject: [PATCH 10/19] forgot to include the style fix --- cpp/tests/CMakeLists.txt | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 080e73b2e6c..0dea2593376 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -164,7 +164,8 @@ ConfigureTest(DATETIME_OPS_TEST datetime/datetime_ops_test.cpp) # ################################################################################################## # * hashing tests --------------------------------------------------------------------------------- ConfigureTest( - HASHING_TEST hashing/md5_test.cpp hashing/murmur3_test.cpp hashing/spark_murmur3_test.cpp hashing/xxhash64_test.cpp + HASHING_TEST hashing/md5_test.cpp hashing/murmur3_test.cpp hashing/spark_murmur3_test.cpp + hashing/xxhash64_test.cpp ) # ################################################################################################## From 9e2ab2c9c1218418692c3913593797898f1d34e2 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Wed, 12 Jul 2023 13:04:53 -0400 Subject: [PATCH 11/19] undo unintentional unchanges --- cpp/include/cudf/hashing.hpp | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/cpp/include/cudf/hashing.hpp b/cpp/include/cudf/hashing.hpp index 0d5d05efe34..1ad0bbe6a86 100644 --- a/cpp/include/cudf/hashing.hpp +++ b/cpp/include/cudf/hashing.hpp @@ -28,7 +28,11 @@ namespace cudf { * @file */ -using hash_value_type = uint32_t; ///< Type of hash 32-bit value +/** + * @brief Type of hash value + * + */ +using hash_value_type = uint32_t; /** * @brief Identifies the hash function to be used @@ -66,6 +70,7 @@ std::unique_ptr hash( rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +//! Hash APIs namespace hashing { /** @@ -99,7 +104,7 @@ std::unique_ptr murmur_hash3_32( * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource used to allocate the returned column's device memory * - * @returns A column where each row is the hash of a column from the input + * @returns A column where each row is the hash of a row from the input */ std::unique_ptr spark_murmur_hash3_32( table_view const& input, From a586b5b9ef09fea5ac5b56a414f16ae890a2c010 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 17 Jul 2023 11:27:59 -0400 Subject: [PATCH 12/19] change names of the gtests --- .../cudf/hashing/detail/hash_functions.cuh | 4 +-- cpp/src/hash/xxhash64.cu | 4 +-- cpp/tests/hashing/xxhash64_test.cpp | 25 +++++++++++-------- 3 files changed, 19 insertions(+), 14 deletions(-) diff --git a/cpp/include/cudf/hashing/detail/hash_functions.cuh b/cpp/include/cudf/hashing/detail/hash_functions.cuh index 437ce3aade4..4089abe7e30 100644 --- a/cpp/include/cudf/hashing/detail/hash_functions.cuh +++ b/cpp/include/cudf/hashing/detail/hash_functions.cuh @@ -52,9 +52,9 @@ __device__ inline uint32_t rotate_bits_left(uint32_t x, uint32_t r) return __funnelshift_l(x, x, r); } -__device__ inline uint64_t rotate_bits_left(uint64_t h, uint32_t r) +__device__ inline uint64_t rotate_bits_left(uint64_t x, uint32_t r) { - return ((h << r) | (h >> (64 - r))); + return ((x << r) | (x >> (64 - r))); } __device__ inline uint32_t rotate_bits_right(uint32_t x, uint32_t r) diff --git a/cpp/src/hash/xxhash64.cu b/cpp/src/hash/xxhash64.cu index c27baf1a85e..990a95c0a5e 100644 --- a/cpp/src/hash/xxhash64.cu +++ b/cpp/src/hash/xxhash64.cu @@ -14,10 +14,10 @@ * limitations under the License. */ #include -#include #include #include #include +#include #include #include @@ -283,7 +283,7 @@ class device_row_hasher { Nullate const, hash_value_type const) const noexcept { - CUDF_UNREACHABLE("Unsupported type for MurmurHash64"); + CUDF_UNREACHABLE("Unsupported type for XXHash_64"); } }; diff --git a/cpp/tests/hashing/xxhash64_test.cpp b/cpp/tests/hashing/xxhash64_test.cpp index ea2c23b5009..5747b40220b 100644 --- a/cpp/tests/hashing/xxhash64_test.cpp +++ b/cpp/tests/hashing/xxhash64_test.cpp @@ -28,11 +28,11 @@ using NumericTypesNoBools = cudf::test::Concat; template -class HashXX64TestTyped : public cudf::test::BaseFixture {}; +class XXHash_64_TestTyped : public cudf::test::BaseFixture {}; -TYPED_TEST_SUITE(HashXX64TestTyped, NumericTypesNoBools); +TYPED_TEST_SUITE(XXHash_64_TestTyped, NumericTypesNoBools); -TYPED_TEST(HashXX64TestTyped, TestAllNumeric) +TYPED_TEST(XXHash_64_TestTyped, TestAllNumeric) { using T = TypeParam; auto col1 = cudf::test::fixed_width_column_wrapper{ @@ -51,9 +51,9 @@ TYPED_TEST(HashXX64TestTyped, TestAllNumeric) CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view()); } -class HashXX64Test : public cudf::test::BaseFixture {}; +class XXHash_64_Test : public cudf::test::BaseFixture {}; -TEST_F(HashXX64Test, TestInteger) +TEST_F(XXHash_64_Test, TestInteger) { auto col1 = cudf::test::fixed_width_column_wrapper{{-127, @@ -80,7 +80,7 @@ TEST_F(HashXX64Test, TestInteger) CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected); } -TEST_F(HashXX64Test, TestDouble) +TEST_F(XXHash_64_Test, TestDouble) { auto col1 = cudf::test::fixed_width_column_wrapper{{-127., @@ -113,15 +113,17 @@ TEST_F(HashXX64Test, TestDouble) CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected); } -TEST_F(HashXX64Test, StringType) +TEST_F(XXHash_64_Test, StringType) { + // clang-format off auto col1 = cudf::test::strings_column_wrapper( {"The", "quick", "brown fox", "jumps over the lazy dog.", "I am Jack's complete lack of null value", - "A very long (greater than 128 bytes/char string) to test a a very long string", + "A very long (greater than 128 bytes/characters) to test a very long string. " + "2nd half of the very long string to verify the long string hashing happening.", "Some multi-byte characters here: ééé", "ééé", "ééé ééé", @@ -130,17 +132,20 @@ TEST_F(HashXX64Test, StringType) "!@#$%^&*(())", "0123456789", "{}|:<>?,./;[]=-"}); + // clang-format on auto output = cudf::hashing::xxhash64(cudf::table_view({col1})); // these were generated using the CPU compiled version of the cuco xxhash64 source // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh + // Also verified these with https://pypi.org/project/xxhash/ + // using xxhash.xxh64(bytes(s,'utf-8')).intdigest() auto expected = cudf::test::fixed_width_column_wrapper({4686269239494003989ul, 6715983472207430822ul, 8148134898123095730ul, 17291005374665645904ul, 2631835514925512071ul, - 17518648592268952189ul, + 4181420602165187991ul, 8749004388517322364ul, 17701789113925815768ul, 8612485687958712810ul, @@ -152,7 +157,7 @@ TEST_F(HashXX64Test, StringType) CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected); } -TEST_F(HashXX64Test, TestFixedPoint) +TEST_F(XXHash_64_Test, TestFixedPoint) { auto const col1 = cudf::test::fixed_point_column_wrapper( {0, 100, -100, -999999999, 999999999}, numeric::scale_type{-3}); From 25856575372781cd449102b9c3529d83bd2ea369 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 17 Jul 2023 11:37:12 -0400 Subject: [PATCH 13/19] fix hash_functions.cuh references --- cpp/src/hash/xxhash64.cu | 34 ++++++++++++++--------------- cpp/tests/hashing/xxhash64_test.cpp | 2 +- 2 files changed, 18 insertions(+), 18 deletions(-) diff --git a/cpp/src/hash/xxhash64.cu b/cpp/src/hash/xxhash64.cu index 990a95c0a5e..139b5b02db4 100644 --- a/cpp/src/hash/xxhash64.cu +++ b/cpp/src/hash/xxhash64.cu @@ -16,7 +16,7 @@ #include #include #include -#include +#include #include #include @@ -73,9 +73,9 @@ struct XXHash_64 { for (; offset <= nbytes - 8; offset += 8) { uint64_t k1 = getblock64(data, offset) * prime2; - k1 = cudf::detail::rotate_bits_left(k1, 31) * prime1; + k1 = rotate_bits_left(k1, 31) * prime1; h64 ^= k1; - h64 = cudf::detail::rotate_bits_left(h64, 27) * prime1 + prime4; + h64 = rotate_bits_left(h64, 27) * prime1 + prime4; } } @@ -83,7 +83,7 @@ struct XXHash_64 { if (((nbytes % 32) % 8) >= 4) { for (; offset <= nbytes - 4; offset += 4) { h64 ^= (getblock32(data, offset) & 0xfffffffful) * prime1; - h64 = cudf::detail::rotate_bits_left(h64, 23) * prime2 + prime3; + h64 = rotate_bits_left(h64, 23) * prime2 + prime3; } } @@ -91,7 +91,7 @@ struct XXHash_64 { if (nbytes % 4) { while (offset < nbytes) { h64 ^= (static_cast(data[offset]) & 0xff) * prime5; - h64 = cudf::detail::rotate_bits_left(h64, 11) * prime1; + h64 = rotate_bits_left(h64, 11) * prime1; ++offset; } } @@ -113,46 +113,46 @@ struct XXHash_64 { do { // pipeline 4*8byte computations v1 += getblock64(data, offset) * prime2; - v1 = cudf::detail::rotate_bits_left(v1, 31); + v1 = rotate_bits_left(v1, 31); v1 *= prime1; offset += 8; v2 += getblock64(data, offset) * prime2; - v2 = cudf::detail::rotate_bits_left(v2, 31); + v2 = rotate_bits_left(v2, 31); v2 *= prime1; offset += 8; v3 += getblock64(data, offset) * prime2; - v3 = cudf::detail::rotate_bits_left(v3, 31); + v3 = rotate_bits_left(v3, 31); v3 *= prime1; offset += 8; v4 += getblock64(data, offset) * prime2; - v4 = cudf::detail::rotate_bits_left(v4, 31); + v4 = rotate_bits_left(v4, 31); v4 *= prime1; offset += 8; } while (offset <= limit); - h64 = cudf::detail::rotate_bits_left(v1, 1) + cudf::detail::rotate_bits_left(v2, 7) + - cudf::detail::rotate_bits_left(v3, 12) + cudf::detail::rotate_bits_left(v4, 18); + h64 = rotate_bits_left(v1, 1) + rotate_bits_left(v2, 7) + rotate_bits_left(v3, 12) + + rotate_bits_left(v4, 18); v1 *= prime2; - v1 = cudf::detail::rotate_bits_left(v1, 31); + v1 = rotate_bits_left(v1, 31); v1 *= prime1; h64 ^= v1; h64 = h64 * prime1 + prime4; v2 *= prime2; - v2 = cudf::detail::rotate_bits_left(v2, 31); + v2 = rotate_bits_left(v2, 31); v2 *= prime1; h64 ^= v2; h64 = h64 * prime1 + prime4; v3 *= prime2; - v3 = cudf::detail::rotate_bits_left(v3, 31); + v3 = rotate_bits_left(v3, 31); v3 *= prime1; h64 ^= v3; h64 = h64 * prime1 + prime4; v4 *= prime2; - v4 = cudf::detail::rotate_bits_left(v4, 31); + v4 = rotate_bits_left(v4, 31); v4 *= prime1; h64 ^= v4; h64 = h64 * prime1 + prime4; @@ -195,13 +195,13 @@ hash_value_type __device__ inline XXHash_64::operator()(bool const& key) c template <> hash_value_type __device__ inline XXHash_64::operator()(float const& key) const { - return compute(cudf::detail::normalize_nans(key)); + return compute(normalize_nans(key)); } template <> hash_value_type __device__ inline XXHash_64::operator()(double const& key) const { - return compute(cudf::detail::normalize_nans(key)); + return compute(normalize_nans(key)); } template <> diff --git a/cpp/tests/hashing/xxhash64_test.cpp b/cpp/tests/hashing/xxhash64_test.cpp index 5747b40220b..cf63db5a382 100644 --- a/cpp/tests/hashing/xxhash64_test.cpp +++ b/cpp/tests/hashing/xxhash64_test.cpp @@ -132,7 +132,7 @@ TEST_F(XXHash_64_Test, StringType) "!@#$%^&*(())", "0123456789", "{}|:<>?,./;[]=-"}); - // clang-format on + // clang-format on auto output = cudf::hashing::xxhash64(cudf::table_view({col1})); From e1e7b8d8d9a6c1341fd7b9bfc6902a4c1bb61b79 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 17 Jul 2023 13:58:49 -0400 Subject: [PATCH 14/19] rename test source file --- cpp/src/hash/xxhash64.cu | 2 +- cpp/tests/CMakeLists.txt | 2 +- cpp/tests/hashing/{xxhash64_test.cpp => xxhash_64_test.cpp} | 0 3 files changed, 2 insertions(+), 2 deletions(-) rename cpp/tests/hashing/{xxhash64_test.cpp => xxhash_64_test.cpp} (100%) diff --git a/cpp/src/hash/xxhash64.cu b/cpp/src/hash/xxhash64.cu index 139b5b02db4..c84c1f7f4df 100644 --- a/cpp/src/hash/xxhash64.cu +++ b/cpp/src/hash/xxhash64.cu @@ -80,7 +80,7 @@ struct XXHash_64 { } // remaining data can be processed in 4-byte chunks - if (((nbytes % 32) % 8) >= 4) { + if ((nbytes % 8) >= 4) { for (; offset <= nbytes - 4; offset += 4) { h64 ^= (getblock32(data, offset) & 0xfffffffful) * prime1; h64 = rotate_bits_left(h64, 23) * prime2 + prime3; diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index e85451fb6ec..68203846281 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -165,7 +165,7 @@ ConfigureTest(DATETIME_OPS_TEST datetime/datetime_ops_test.cpp) # * hashing tests --------------------------------------------------------------------------------- ConfigureTest( HASHING_TEST hashing/md5_test.cpp hashing/murmurhash3_x86_32_test.cpp - hashing/spark_murmurhash3_x86_32_test.cpp hashing/xxhash64_test.cpp + hashing/spark_murmurhash3_x86_32_test.cpp hashing/xxhash_64_test.cpp ) # ################################################################################################## diff --git a/cpp/tests/hashing/xxhash64_test.cpp b/cpp/tests/hashing/xxhash_64_test.cpp similarity index 100% rename from cpp/tests/hashing/xxhash64_test.cpp rename to cpp/tests/hashing/xxhash_64_test.cpp From cdec016db3dabd5f6dae8a010e7b96c4f21574bd Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 17 Jul 2023 19:01:45 -0400 Subject: [PATCH 15/19] rename xxhash64 to xxhash_64 --- cpp/CMakeLists.txt | 2 +- cpp/include/cudf/hashing.hpp | 2 +- cpp/include/cudf/hashing/detail/hashing.hpp | 8 +++---- cpp/src/hash/{xxhash64.cu => xxhash_64.cu} | 18 ++++++++-------- cpp/tests/hashing/xxhash_64_test.cpp | 24 ++++++++++----------- 5 files changed, 27 insertions(+), 27 deletions(-) rename cpp/src/hash/{xxhash64.cu => xxhash_64.cu} (94%) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 6f71e5d7391..fdc42dcb3d1 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -346,7 +346,7 @@ add_library( src/hash/md5_hash.cu src/hash/murmurhash3_x86_32.cu src/hash/spark_murmurhash3_x86_32.cu - src/hash/xxhash64.cu + src/hash/xxhash_64.cu src/interop/dlpack.cpp src/interop/from_arrow.cu src/interop/to_arrow.cu diff --git a/cpp/include/cudf/hashing.hpp b/cpp/include/cudf/hashing.hpp index f2c89ecabc3..f9e53967cf8 100644 --- a/cpp/include/cudf/hashing.hpp +++ b/cpp/include/cudf/hashing.hpp @@ -138,7 +138,7 @@ std::unique_ptr md5( * * @returns A column where each row is the hash of a row from the input */ -std::unique_ptr xxhash64( +std::unique_ptr xxhash_64( table_view const& input, uint64_t seed = DEFAULT_HASH_SEED, rmm::cuda_stream_view stream = cudf::get_default_stream(), diff --git a/cpp/include/cudf/hashing/detail/hashing.hpp b/cpp/include/cudf/hashing/detail/hashing.hpp index 3a6e0870f02..ac9d09649df 100644 --- a/cpp/include/cudf/hashing/detail/hashing.hpp +++ b/cpp/include/cudf/hashing/detail/hashing.hpp @@ -41,10 +41,10 @@ std::unique_ptr md5(table_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr); -std::unique_ptr xxhash64(table_view const& input, - uint64_t seed, - rmm::cuda_stream_view, - rmm::mr::device_memory_resource* mr); +std::unique_ptr xxhash_64(table_view const& input, + uint64_t seed, + rmm::cuda_stream_view, + rmm::mr::device_memory_resource* mr); /* Copyright 2005-2014 Daniel James. * diff --git a/cpp/src/hash/xxhash64.cu b/cpp/src/hash/xxhash_64.cu similarity index 94% rename from cpp/src/hash/xxhash64.cu rename to cpp/src/hash/xxhash_64.cu index c84c1f7f4df..b89a398c778 100644 --- a/cpp/src/hash/xxhash64.cu +++ b/cpp/src/hash/xxhash_64.cu @@ -294,10 +294,10 @@ class device_row_hasher { } // namespace -std::unique_ptr xxhash64(table_view const& input, - uint64_t seed, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::unique_ptr xxhash_64(table_view const& input, + uint64_t seed, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { auto output = make_numeric_column(data_type(type_to_id()), input.num_rows(), @@ -323,13 +323,13 @@ std::unique_ptr xxhash64(table_view const& input, } // namespace detail -std::unique_ptr xxhash64(table_view const& input, - uint64_t seed, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::unique_ptr xxhash_64(table_view const& input, + uint64_t seed, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::xxhash64(input, seed, stream, mr); + return detail::xxhash_64(input, seed, stream, mr); } } // namespace hashing diff --git a/cpp/tests/hashing/xxhash_64_test.cpp b/cpp/tests/hashing/xxhash_64_test.cpp index cf63db5a382..9deee632832 100644 --- a/cpp/tests/hashing/xxhash_64_test.cpp +++ b/cpp/tests/hashing/xxhash_64_test.cpp @@ -42,12 +42,12 @@ TYPED_TEST(XXHash_64_TestTyped, TestAllNumeric) {-1, -1, 0, 2, 22, 1, 11, 12, 116, 32, 0, 42, 7, 62, 1, -22, 1, -22}, {1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0}}; - auto output1 = cudf::hashing::xxhash64(cudf::table_view({col1})); - auto output2 = cudf::hashing::xxhash64(cudf::table_view({col2})); + auto output1 = cudf::hashing::xxhash_64(cudf::table_view({col1})); + auto output2 = cudf::hashing::xxhash_64(cudf::table_view({col2})); CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view()); - output1 = cudf::hashing::xxhash64(cudf::table_view({col1}), 7); - output2 = cudf::hashing::xxhash64(cudf::table_view({col2}), 7); + output1 = cudf::hashing::xxhash_64(cudf::table_view({col1}), 7); + output2 = cudf::hashing::xxhash_64(cudf::table_view({col2}), 7); CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view()); } @@ -65,9 +65,9 @@ TEST_F(XXHash_64_Test, TestInteger) std::numeric_limits::min(), std::numeric_limits::lowest()}}; - auto const output = cudf::hashing::xxhash64(cudf::table_view({col1})); + auto const output = cudf::hashing::xxhash_64(cudf::table_view({col1})); - // these were generated using the CPU compiled version of the cuco xxhash64 source + // these were generated using the CPU compiled version of the cuco xxhash_64 source // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh auto expected = cudf::test::fixed_width_column_wrapper({4827426872506142937ul, 13867166853951622683ul, @@ -95,9 +95,9 @@ TEST_F(XXHash_64_Test, TestDouble) std::numeric_limits::min(), std::numeric_limits::lowest()}}; - auto const output = cudf::hashing::xxhash64(cudf::table_view({col1})); + auto const output = cudf::hashing::xxhash_64(cudf::table_view({col1})); - // these were generated using the CPU compiled version of the cuco xxhash64 source + // these were generated using the CPU compiled version of the cuco xxhash_64 source // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh auto expected = cudf::test::fixed_width_column_wrapper({16892115221677838993ul, 1686446903308179321ul, @@ -134,9 +134,9 @@ TEST_F(XXHash_64_Test, StringType) "{}|:<>?,./;[]=-"}); // clang-format on - auto output = cudf::hashing::xxhash64(cudf::table_view({col1})); + auto output = cudf::hashing::xxhash_64(cudf::table_view({col1})); - // these were generated using the CPU compiled version of the cuco xxhash64 source + // these were generated using the CPU compiled version of the cuco xxhash_64 source // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh // Also verified these with https://pypi.org/project/xxhash/ // using xxhash.xxh64(bytes(s,'utf-8')).intdigest() @@ -161,9 +161,9 @@ TEST_F(XXHash_64_Test, TestFixedPoint) { auto const col1 = cudf::test::fixed_point_column_wrapper( {0, 100, -100, -999999999, 999999999}, numeric::scale_type{-3}); - auto const output = cudf::hashing::xxhash64(cudf::table_view({col1})); + auto const output = cudf::hashing::xxhash_64(cudf::table_view({col1})); - // these were generated using the CPU compiled version of the cuco xxhash64 source + // these were generated using the CPU compiled version of the cuco xxhash_64 source // https://github.com/NVIDIA/cuCollections/blob/dev/include/cuco/detail/hash_functions/xxhash.cuh // and passing the 'value' of each input (without the scale) as the decimal-type auto expected = cudf::test::fixed_width_column_wrapper({4246796580750024372ul, From 74cda04891e1c54fddcde8ad6fe5d0b84b6a5c5e Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 18 Jul 2023 07:08:19 -0400 Subject: [PATCH 16/19] fix cmake style violation --- cpp/tests/CMakeLists.txt | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index f3fbcf75b9f..8a0942eec0d 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -164,8 +164,11 @@ ConfigureTest(DATETIME_OPS_TEST datetime/datetime_ops_test.cpp) # ################################################################################################## # * hashing tests --------------------------------------------------------------------------------- ConfigureTest( - HASHING_TEST hashing/md5_test.cpp hashing/murmurhash3_x86_32_test.cpp - hashing/murmurhash3_x64_128_test.cpp hashing/spark_murmurhash3_x86_32_test.cpp + HASHING_TEST + hashing/md5_test.cpp + hashing/murmurhash3_x86_32_test.cpp + hashing/murmurhash3_x64_128_test.cpp + hashing/spark_murmurhash3_x86_32_test.cpp hashing/xxhash_64_test.cpp ) From d47d62474eadb27307da617ce4c0bd94905db23b Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 18 Jul 2023 15:53:56 -0400 Subject: [PATCH 17/19] add some const decls --- cpp/src/hash/xxhash_64.cu | 4 ++-- cpp/tests/hashing/xxhash_64_test.cpp | 6 ++++-- 2 files changed, 6 insertions(+), 4 deletions(-) diff --git a/cpp/src/hash/xxhash_64.cu b/cpp/src/hash/xxhash_64.cu index b89a398c778..559c791b496 100644 --- a/cpp/src/hash/xxhash_64.cu +++ b/cpp/src/hash/xxhash_64.cu @@ -266,7 +266,7 @@ class device_row_hasher { public: template ())> __device__ hash_value_type operator()(column_device_view const& col, - size_type row_index, + size_type const row_index, Nullate const _check_nulls, hash_value_type const _seed) const noexcept { @@ -279,7 +279,7 @@ class device_row_hasher { template ())> __device__ hash_value_type operator()(column_device_view const&, - size_type, + size_type const, Nullate const, hash_value_type const) const noexcept { diff --git a/cpp/tests/hashing/xxhash_64_test.cpp b/cpp/tests/hashing/xxhash_64_test.cpp index 9deee632832..5916c4c2fb9 100644 --- a/cpp/tests/hashing/xxhash_64_test.cpp +++ b/cpp/tests/hashing/xxhash_64_test.cpp @@ -46,8 +46,10 @@ TYPED_TEST(XXHash_64_TestTyped, TestAllNumeric) auto output2 = cudf::hashing::xxhash_64(cudf::table_view({col2})); CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view()); - output1 = cudf::hashing::xxhash_64(cudf::table_view({col1}), 7); - output2 = cudf::hashing::xxhash_64(cudf::table_view({col2}), 7); + constexpr uint64_t seed = 7; + + output1 = cudf::hashing::xxhash_64(cudf::table_view({col1}), seed); + output2 = cudf::hashing::xxhash_64(cudf::table_view({col2}), seed); CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view()); } From 606b736d51b638d326872f395a3c085d7407797a Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 18 Jul 2023 16:39:34 -0400 Subject: [PATCH 18/19] fix doxygen wording for the hash APIs --- cpp/include/cudf/hashing.hpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/cpp/include/cudf/hashing.hpp b/cpp/include/cudf/hashing.hpp index 5daead79952..72e32715ed4 100644 --- a/cpp/include/cudf/hashing.hpp +++ b/cpp/include/cudf/hashing.hpp @@ -74,7 +74,7 @@ std::unique_ptr hash( namespace hashing { /** - * @brief Computes the MurmurHash3 32-bit of each row in the given table + * @brief Computes the MurmurHash3 32-bit hash value of each row in the given table * * This function computes the hash of each column using the `seed` for the first column * and the resulting hash as a seed for the next column and so on. @@ -94,7 +94,7 @@ std::unique_ptr murmurhash3_x86_32( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Computes the hash values of each row in the input set of columns + * @brief Computes the MurmurHash3 64-bit hash value of each row in the given table * * This function takes a 64-bit seed value and returns hash values using the * MurmurHash3_x64_128 algorithm. The hash produces in two uint64 values per row. @@ -113,7 +113,7 @@ std::unique_ptr murmurhash3_x64_128( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Computes the MurmurHash3 32-bit of each row in the given table + * @brief Computes the MurmurHash3 32-bit hash value of each row in the given table * * This function computes the hash similar to MurmurHash3_x86_32 with special processing * to match Spark's implementation results. @@ -132,7 +132,7 @@ std::unique_ptr spark_murmurhash3_x86_32( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Computes the MD5 hash of each row in the given table + * @brief Computes the MD5 hash value of each row in the given table * * @param input The table of columns to hash * @param stream CUDA stream used for device memory operations and kernel launches @@ -146,7 +146,7 @@ std::unique_ptr md5( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Computes the hash value of each row in the input set of columns + * @brief Computes the XXHash_64 hash value of each row in the given table * * This function takes a 64-bit seed value and returns a column of type UINT64. * From a5e38388ca5f34d2fcd377038807eaff7c7c7a6d Mon Sep 17 00:00:00 2001 From: David Wendt Date: Wed, 19 Jul 2023 09:08:33 -0400 Subject: [PATCH 19/19] use device-span --- cpp/src/hash/xxhash_64.cu | 55 ++++++++++++++++++++------------------- 1 file changed, 28 insertions(+), 27 deletions(-) diff --git a/cpp/src/hash/xxhash_64.cu b/cpp/src/hash/xxhash_64.cu index 559c791b496..e17bc134420 100644 --- a/cpp/src/hash/xxhash_64.cu +++ b/cpp/src/hash/xxhash_64.cu @@ -19,6 +19,7 @@ #include #include #include +#include #include #include @@ -40,7 +41,7 @@ struct XXHash_64 { constexpr XXHash_64() = default; constexpr XXHash_64(hash_value_type seed) : m_seed(seed) {} - __device__ inline uint32_t getblock32(std::byte const* data, cudf::size_type offset) const + __device__ inline uint32_t getblock32(std::byte const* data, std::size_t offset) const { // Read a 4-byte value from the data pointer as individual bytes for safe // unaligned access (very likely for string types). @@ -48,7 +49,7 @@ struct XXHash_64 { return block[0] | (block[1] << 8) | (block[2] << 16) | (block[3] << 24); } - __device__ inline uint64_t getblock64(std::byte const* data, cudf::size_type offset) const + __device__ inline uint64_t getblock64(std::byte const* data, std::size_t offset) const { uint64_t result = getblock32(data, offset + 4); result = result << 32; @@ -60,18 +61,18 @@ struct XXHash_64 { template result_type __device__ inline compute(T const& key) const { - return compute_bytes(reinterpret_cast(&key), sizeof(T)); + auto data = device_span(reinterpret_cast(&key), sizeof(T)); + return compute_bytes(data); } - result_type __device__ inline compute_remaining_bytes(std::byte const* data, - cudf::size_type nbytes, - cudf::size_type offset, + result_type __device__ inline compute_remaining_bytes(device_span& in, + std::size_t offset, result_type h64) const { // remaining data can be processed in 8-byte chunks - if ((nbytes % 32) >= 8) { - for (; offset <= nbytes - 8; offset += 8) { - uint64_t k1 = getblock64(data, offset) * prime2; + if ((in.size() % 32) >= 8) { + for (; offset <= in.size() - 8; offset += 8) { + uint64_t k1 = getblock64(in.data(), offset) * prime2; k1 = rotate_bits_left(k1, 31) * prime1; h64 ^= k1; @@ -80,17 +81,17 @@ struct XXHash_64 { } // remaining data can be processed in 4-byte chunks - if ((nbytes % 8) >= 4) { - for (; offset <= nbytes - 4; offset += 4) { - h64 ^= (getblock32(data, offset) & 0xfffffffful) * prime1; + if ((in.size() % 8) >= 4) { + for (; offset <= in.size() - 4; offset += 4) { + h64 ^= (getblock32(in.data(), offset) & 0xfffffffful) * prime1; h64 = rotate_bits_left(h64, 23) * prime2 + prime3; } } // and the rest - if (nbytes % 4) { - while (offset < nbytes) { - h64 ^= (static_cast(data[offset]) & 0xff) * prime5; + if (in.size() % 4) { + while (offset < in.size()) { + h64 ^= (std::to_integer(in[offset]) & 0xff) * prime5; h64 = rotate_bits_left(h64, 11) * prime1; ++offset; } @@ -98,13 +99,13 @@ struct XXHash_64 { return h64; } - result_type __device__ compute_bytes(std::byte const* data, cudf::size_type const nbytes) const + result_type __device__ compute_bytes(device_span& in) const { uint64_t offset = 0; uint64_t h64; // data can be processed in 32-byte chunks - if (nbytes >= 32) { - auto limit = nbytes - 32; + if (in.size() >= 32) { + auto limit = in.size() - 32; uint64_t v1 = m_seed + prime1 + prime2; uint64_t v2 = m_seed + prime2; uint64_t v3 = m_seed; @@ -112,19 +113,19 @@ struct XXHash_64 { do { // pipeline 4*8byte computations - v1 += getblock64(data, offset) * prime2; + v1 += getblock64(in.data(), offset) * prime2; v1 = rotate_bits_left(v1, 31); v1 *= prime1; offset += 8; - v2 += getblock64(data, offset) * prime2; + v2 += getblock64(in.data(), offset) * prime2; v2 = rotate_bits_left(v2, 31); v2 *= prime1; offset += 8; - v3 += getblock64(data, offset) * prime2; + v3 += getblock64(in.data(), offset) * prime2; v3 = rotate_bits_left(v3, 31); v3 *= prime1; offset += 8; - v4 += getblock64(data, offset) * prime2; + v4 += getblock64(in.data(), offset) * prime2; v4 = rotate_bits_left(v4, 31); v4 *= prime1; offset += 8; @@ -160,9 +161,9 @@ struct XXHash_64 { h64 = m_seed + prime5; } - h64 += nbytes; + h64 += in.size(); - h64 = compute_remaining_bytes(data, nbytes, offset, h64); + h64 = compute_remaining_bytes(in, offset, h64); return finalize(h64); } @@ -208,9 +209,9 @@ template <> hash_value_type __device__ inline XXHash_64::operator()( cudf::string_view const& key) const { - auto const data = reinterpret_cast(key.data()); - auto const len = key.size_bytes(); - return compute_bytes(data, len); + auto const len = key.size_bytes(); + auto data = device_span(reinterpret_cast(key.data()), len); + return compute_bytes(data); } template <>