From 788bd5483c82219df8b37b16e4f040c8e6359ffe Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 10 Jul 2023 15:39:41 -0400 Subject: [PATCH 1/8] Separate MurmurHash32 from hash_functions.cuh --- conda/recipes/libcudf/meta.yaml | 2 +- .../cudf/detail/aggregation/result_cache.hpp | 2 +- cpp/include/cudf/detail/join.hpp | 3 +- .../cudf/hashing/detail/hash_functions.cuh | 212 ++++++++++++++++++ .../cudf/{ => hashing}/detail/hashing.hpp | 11 - .../cudf/hashing/detail/murmur_hash32.cuh | 196 ++++++++++++++++ .../cudf/table/experimental/row_operators.cuh | 4 +- cpp/include/cudf/table/row_operators.cuh | 4 +- cpp/src/column/column_view.cpp | 2 +- cpp/src/groupby/hash/groupby.cu | 2 +- cpp/src/hash/concurrent_unordered_map.cuh | 2 +- cpp/src/hash/hashing.cu | 40 +--- cpp/src/hash/md5_hash.cu | 14 +- cpp/src/hash/murmur_hash.cu | 16 +- cpp/src/hash/spark_murmur_hash.cu | 15 +- cpp/src/hash/unordered_multiset.cuh | 3 +- cpp/src/io/json/json_gpu.cu | 2 +- cpp/src/io/json/json_tree.cu | 4 +- cpp/src/io/parquet/page_data.cu | 3 +- cpp/src/join/join_common_utils.hpp | 2 +- cpp/src/partitioning/partitioning.cu | 2 +- .../stream_compaction_common.hpp | 2 +- cpp/src/text/minhash.cu | 4 +- cpp/src/text/subword/bpe_tokenizer.cu | 2 +- cpp/src/text/subword/bpe_tokenizer.cuh | 4 +- cpp/src/text/subword/load_merges_file.cu | 2 +- cpp/tests/io/json_tree.cpp | 2 +- 27 files changed, 476 insertions(+), 81 deletions(-) create mode 100644 cpp/include/cudf/hashing/detail/hash_functions.cuh rename cpp/include/cudf/{ => hashing}/detail/hashing.hpp (89%) create mode 100644 cpp/include/cudf/hashing/detail/murmur_hash32.cuh diff --git a/conda/recipes/libcudf/meta.yaml b/conda/recipes/libcudf/meta.yaml index 4e9b5e2fdc1..f4b9945de0f 100644 --- a/conda/recipes/libcudf/meta.yaml +++ b/conda/recipes/libcudf/meta.yaml @@ -139,7 +139,6 @@ outputs: - test -f $PREFIX/include/cudf/detail/groupby.hpp - test -f $PREFIX/include/cudf/detail/groupby/group_replace_nulls.hpp - test -f $PREFIX/include/cudf/detail/groupby/sort_helper.hpp - - test -f $PREFIX/include/cudf/detail/hashing.hpp - test -f $PREFIX/include/cudf/detail/interop.hpp - test -f $PREFIX/include/cudf/detail/is_element_valid.hpp - test -f $PREFIX/include/cudf/detail/join.hpp @@ -192,6 +191,7 @@ outputs: - test -f $PREFIX/include/cudf/fixed_point/temporary.hpp - test -f $PREFIX/include/cudf/groupby.hpp - test -f $PREFIX/include/cudf/hashing.hpp + - test -f $PREFIX/include/cudf/hashing/detail/hashing.hpp - test -f $PREFIX/include/cudf/interop.hpp - test -f $PREFIX/include/cudf/io/avro.hpp - test -f $PREFIX/include/cudf/io/csv.hpp diff --git a/cpp/include/cudf/detail/aggregation/result_cache.hpp b/cpp/include/cudf/detail/aggregation/result_cache.hpp index b1a2a369d22..41eec156c47 100644 --- a/cpp/include/cudf/detail/aggregation/result_cache.hpp +++ b/cpp/include/cudf/detail/aggregation/result_cache.hpp @@ -18,7 +18,7 @@ #include #include -#include +#include #include #include diff --git a/cpp/include/cudf/detail/join.hpp b/cpp/include/cudf/detail/join.hpp index 4a34eb6b328..4cbb59c035c 100644 --- a/cpp/include/cudf/detail/join.hpp +++ b/cpp/include/cudf/detail/join.hpp @@ -16,7 +16,8 @@ #pragma once #include -#include +#include +#include #include #include #include diff --git a/cpp/include/cudf/hashing/detail/hash_functions.cuh b/cpp/include/cudf/hashing/detail/hash_functions.cuh new file mode 100644 index 00000000000..96f5b5b067d --- /dev/null +++ b/cpp/include/cudf/hashing/detail/hash_functions.cuh @@ -0,0 +1,212 @@ +/* + * Copyright (c) 2017-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. + */ + +#pragma once + +#include +#include + +#include + +#include + +namespace cudf { +namespace detail { + +template +struct MurmurHash3_32; + +template +using default_hash = MurmurHash3_32; + +/** + * Normalization of floating point NaNs, passthrough for all other values. + */ +template +T __device__ inline normalize_nans(T const& key) +{ + if constexpr (cudf::is_floating_point()) { + if (std::isnan(key)) { return std::numeric_limits::quiet_NaN(); } + } + return key; +} + +/** + * Normalization of floating point NaNs and zeros, passthrough for all other values. + */ +template +T __device__ inline normalize_nans_and_zeros(T const& key) +{ + if constexpr (cudf::is_floating_point()) { + if (key == T{0.0}) { return T{0.0}; } + } + return normalize_nans(key); +} + +__device__ inline uint32_t rotate_bits_left(uint32_t x, uint32_t r) +{ + // This function is equivalent to (x << r) | (x >> (32 - r)) + return __funnelshift_l(x, x, r); +} + +__device__ inline uint32_t rotate_bits_right(uint32_t x, uint32_t r) +{ + // This function is equivalent to (x >> r) | (x << (32 - r)) + return __funnelshift_r(x, x, r); +} + +__device__ inline uint64_t rotate_bits_right(uint64_t x, uint32_t r) +{ + return (x >> r) | (x << (64 - r)); +} + +// Swap the endianness of a 32 bit value +__device__ inline uint32_t swap_endian(uint32_t x) +{ + // The selector 0x0123 reverses the byte order + return __byte_perm(x, 0, 0x0123); +} + +// Swap the endianness of a 64 bit value +// There is no CUDA intrinsic for permuting bytes in 64 bit integers +__device__ inline uint64_t swap_endian(uint64_t x) +{ + // Reverse the endianness of each 32 bit section + uint32_t low_bits = swap_endian(static_cast(x)); + uint32_t high_bits = swap_endian(static_cast(x >> 32)); + // Reassemble a 64 bit result, swapping the low bits and high bits + return (static_cast(low_bits) << 32) | (static_cast(high_bits)); +}; + +template +struct hash_circular_buffer { + uint8_t storage[capacity]; + uint8_t* cur; + int available_space{capacity}; + hash_step_callable hash_step; + + __device__ inline hash_circular_buffer(hash_step_callable hash_step) + : cur{storage}, hash_step{hash_step} + { + } + + __device__ inline void put(uint8_t const* in, int size) + { + int copy_start = 0; + while (size >= available_space) { + // The buffer will be filled by this chunk of data. Copy a chunk of the + // data to fill the buffer and trigger a hash step. + memcpy(cur, in + copy_start, available_space); + hash_step(storage); + size -= available_space; + copy_start += available_space; + cur = storage; + available_space = capacity; + } + // The buffer will not be filled by the remaining data. That is, `size >= 0 + // && size < capacity`. We copy the remaining data into the buffer but do + // not trigger a hash step. + memcpy(cur, in + copy_start, size); + cur += size; + available_space -= size; + } + + __device__ inline void pad(int const space_to_leave) + { + if (space_to_leave > available_space) { + memset(cur, 0x00, available_space); + hash_step(storage); + cur = storage; + available_space = capacity; + } + memset(cur, 0x00, available_space - space_to_leave); + cur += available_space - space_to_leave; + available_space = space_to_leave; + } + + __device__ inline uint8_t const& operator[](int idx) const { return storage[idx]; } +}; + +// Get a uint8_t pointer to a column element and its size as a pair. +template +auto __device__ inline get_element_pointer_and_size(Element const& element) +{ + if constexpr (is_fixed_width() && !is_chrono()) { + return thrust::make_pair(reinterpret_cast(&element), sizeof(Element)); + } else { + CUDF_UNREACHABLE("Unsupported type."); + } +} + +template <> +auto __device__ inline get_element_pointer_and_size(string_view const& element) +{ + return thrust::make_pair(reinterpret_cast(element.data()), element.size_bytes()); +} + +/** + * Modified GPU implementation of + * https://johnnylee-sde.github.io/Fast-unsigned-integer-to-hex-string/ + * Copyright (c) 2015 Barry Clark + * Licensed under the MIT license. + * See file LICENSE for detail or copy at https://opensource.org/licenses/MIT + */ +void __device__ inline uint32ToLowercaseHexString(uint32_t num, char* destination) +{ + // Transform 0xABCD'1234 => 0x0000'ABCD'0000'1234 => 0x0B0A'0D0C'0201'0403 + uint64_t x = num; + x = ((x & 0xFFFF'0000u) << 16) | ((x & 0xFFFF)); + x = ((x & 0x000F'0000'000Fu) << 8) | ((x & 0x00F0'0000'00F0u) >> 4) | + ((x & 0x0F00'0000'0F00u) << 16) | ((x & 0xF000'0000'F000) << 4); + + // Calculate a mask of ascii value offsets for bytes that contain alphabetical hex digits + uint64_t offsets = (((x + 0x0606'0606'0606'0606) >> 4) & 0x0101'0101'0101'0101) * 0x27; + + x |= 0x3030'3030'3030'3030; + x += offsets; + std::memcpy(destination, reinterpret_cast(&x), 8); +} + +/** + * @brief This hash function simply returns the value that is asked to be hash + * reinterpreted as the result_type of the functor. + */ +template +struct IdentityHash { + using result_type = uint32_t; + IdentityHash() = default; + constexpr IdentityHash(uint32_t seed) : m_seed(seed) {} + + template + constexpr std::enable_if_t, return_type> operator()( + Key const& key) const + { + CUDF_UNREACHABLE("IdentityHash does not support this data type"); + } + + template + constexpr std::enable_if_t, return_type> operator()( + Key const& key) const + { + return static_cast(key); + } + + private: + uint32_t m_seed{0}; +}; + +} // namespace detail +} // namespace cudf diff --git a/cpp/include/cudf/detail/hashing.hpp b/cpp/include/cudf/hashing/detail/hashing.hpp similarity index 89% rename from cpp/include/cudf/detail/hashing.hpp rename to cpp/include/cudf/hashing/detail/hashing.hpp index 0447384ffdc..08e1a1d03c0 100644 --- a/cpp/include/cudf/detail/hashing.hpp +++ b/cpp/include/cudf/hashing/detail/hashing.hpp @@ -84,17 +84,6 @@ constexpr std::size_t hash_combine(std::size_t lhs, std::size_t rhs) } } // namespace detail - -/** - * @copydoc cudf::hash - * - * @param stream CUDA stream used for device memory operations and kernel launches. - */ -std::unique_ptr hash(table_view const& input, - hash_id hash_function, - uint32_t seed, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr); } // namespace hashing } // namespace cudf diff --git a/cpp/include/cudf/hashing/detail/murmur_hash32.cuh b/cpp/include/cudf/hashing/detail/murmur_hash32.cuh new file mode 100644 index 00000000000..26425875068 --- /dev/null +++ b/cpp/include/cudf/hashing/detail/murmur_hash32.cuh @@ -0,0 +1,196 @@ +/* + * Copyright (c) 2017-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. + */ + +#pragma once + +#include +#include +#include +#include +#include +#include +#include + +#include + +namespace cudf { +namespace detail { + +// MurmurHash3_32 implementation from +// https://github.com/aappleby/smhasher/blob/master/src/MurmurHash3.cpp +//----------------------------------------------------------------------------- +// MurmurHash3 was written by Austin Appleby, and is placed in the public +// domain. The author hereby disclaims copyright to this source code. +// Note - The x86 and x64 versions do _not_ produce the same results, as the +// algorithms are optimized for their respective platforms. You can still +// compile and run any of them on any platform, but your performance with the +// non-native version will be less than optimal. +template +struct MurmurHash3_32 { + using result_type = hash_value_type; + + constexpr MurmurHash3_32() = default; + constexpr MurmurHash3_32(uint32_t seed) : m_seed(seed) {} + + [[nodiscard]] __device__ inline uint32_t fmix32(uint32_t h) const + { + h ^= h >> 16; + h *= 0x85ebca6b; + h ^= h >> 13; + h *= 0xc2b2ae35; + h ^= h >> 16; + return h; + } + + [[nodiscard]] __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 const block = reinterpret_cast(data + offset); + return block[0] | (block[1] << 8) | (block[2] << 16) | (block[3] << 24); + } + + [[nodiscard]] result_type __device__ inline operator()(Key const& key) const + { + return compute(detail::normalize_nans_and_zeros(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 len, + cudf::size_type tail_offset, + result_type h) const + { + // Process remaining bytes that do not fill a four-byte chunk. + uint32_t k1 = 0; + switch (len % 4) { + case 3: k1 ^= std::to_integer(data[tail_offset + 2]) << 16; [[fallthrough]]; + case 2: k1 ^= std::to_integer(data[tail_offset + 1]) << 8; [[fallthrough]]; + case 1: + k1 ^= std::to_integer(data[tail_offset]); + k1 *= c1; + k1 = cudf::detail::rotate_bits_left(k1, rot_c1); + k1 *= c2; + h ^= k1; + }; + return h; + } + + result_type __device__ compute_bytes(std::byte const* data, cudf::size_type const len) const + { + constexpr cudf::size_type BLOCK_SIZE = 4; + cudf::size_type const nblocks = len / BLOCK_SIZE; + cudf::size_type const tail_offset = nblocks * BLOCK_SIZE; + result_type h = m_seed; + + // Process all four-byte chunks. + for (cudf::size_type i = 0; i < nblocks; i++) { + uint32_t k1 = getblock32(data, i * BLOCK_SIZE); + k1 *= c1; + k1 = cudf::detail::rotate_bits_left(k1, rot_c1); + k1 *= c2; + h ^= k1; + h = cudf::detail::rotate_bits_left(h, rot_c2); + h = h * 5 + c3; + } + + h = compute_remaining_bytes(data, len, tail_offset, h); + + // Finalize hash. + h ^= len; + h = fmix32(h); + return h; + } + + private: + uint32_t m_seed{cudf::DEFAULT_HASH_SEED}; + static constexpr uint32_t c1 = 0xcc9e2d51; + static constexpr uint32_t c2 = 0x1b873593; + static constexpr uint32_t c3 = 0xe6546b64; + static constexpr uint32_t rot_c1 = 15; + static constexpr uint32_t rot_c2 = 13; +}; + +template <> +hash_value_type __device__ inline MurmurHash3_32::operator()(bool const& key) const +{ + return compute(static_cast(key)); +} + +template <> +hash_value_type __device__ inline MurmurHash3_32::operator()(float const& key) const +{ + return compute(detail::normalize_nans_and_zeros(key)); +} + +template <> +hash_value_type __device__ inline MurmurHash3_32::operator()(double const& key) const +{ + return compute(detail::normalize_nans_and_zeros(key)); +} + +template <> +hash_value_type __device__ inline MurmurHash3_32::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 MurmurHash3_32::operator()( + numeric::decimal32 const& key) const +{ + return compute(key.value()); +} + +template <> +hash_value_type __device__ inline MurmurHash3_32::operator()( + numeric::decimal64 const& key) const +{ + return compute(key.value()); +} + +template <> +hash_value_type __device__ inline MurmurHash3_32::operator()( + numeric::decimal128 const& key) const +{ + return compute(key.value()); +} + +template <> +hash_value_type __device__ inline MurmurHash3_32::operator()( + cudf::list_view const& key) const +{ + CUDF_UNREACHABLE("List column hashing is not supported"); +} + +template <> +hash_value_type __device__ inline MurmurHash3_32::operator()( + cudf::struct_view const& key) const +{ + CUDF_UNREACHABLE("Direct hashing of struct_view is not supported"); +} + +} // namespace detail +} // namespace cudf diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index ce6dd024622..7e876d6cbad 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -17,11 +17,11 @@ #pragma once #include -#include #include #include #include -#include +#include +#include #include #include #include diff --git a/cpp/include/cudf/table/row_operators.cuh b/cpp/include/cudf/table/row_operators.cuh index f90c78b9ba6..599a85c8a54 100644 --- a/cpp/include/cudf/table/row_operators.cuh +++ b/cpp/include/cudf/table/row_operators.cuh @@ -17,9 +17,9 @@ #pragma once #include -#include #include -#include +#include +#include #include #include #include diff --git a/cpp/src/column/column_view.cpp b/cpp/src/column/column_view.cpp index bf7c84b2735..75722ede9d2 100644 --- a/cpp/src/column/column_view.cpp +++ b/cpp/src/column/column_view.cpp @@ -15,8 +15,8 @@ */ #include -#include #include +#include #include #include #include diff --git a/cpp/src/groupby/hash/groupby.cu b/cpp/src/groupby/hash/groupby.cu index 62d83e8f6ae..f7ecc40c20b 100644 --- a/cpp/src/groupby/hash/groupby.cu +++ b/cpp/src/groupby/hash/groupby.cu @@ -33,10 +33,10 @@ #include #include #include -#include #include #include #include +#include #include #include #include diff --git a/cpp/src/hash/concurrent_unordered_map.cuh b/cpp/src/hash/concurrent_unordered_map.cuh index 5acfba0a8bf..89fa8442f21 100644 --- a/cpp/src/hash/concurrent_unordered_map.cuh +++ b/cpp/src/hash/concurrent_unordered_map.cuh @@ -21,7 +21,7 @@ #include #include -#include +#include #include #include diff --git a/cpp/src/hash/hashing.cu b/cpp/src/hash/hashing.cu index ca32bd14aff..875d17587c6 100644 --- a/cpp/src/hash/hashing.cu +++ b/cpp/src/hash/hashing.cu @@ -13,22 +13,11 @@ * 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 - -#include namespace cudf { namespace hashing { @@ -49,33 +38,6 @@ std::unique_ptr hash(table_view const& input, } } // namespace detail - -std::unique_ptr murmur_hash3_32(table_view const& input, - uint32_t seed, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - CUDF_FUNC_RANGE(); - return detail::murmur_hash3_32(input, seed, stream, mr); -} - -std::unique_ptr spark_murmur_hash3_32(table_view const& input, - uint32_t seed, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - CUDF_FUNC_RANGE(); - return detail::spark_murmur_hash3_32(input, seed, stream, mr); -} - -std::unique_ptr md5(table_view const& input, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - CUDF_FUNC_RANGE(); - return detail::md5(input, stream, mr); -} - } // namespace hashing std::unique_ptr hash(table_view const& input, diff --git a/cpp/src/hash/md5_hash.cu b/cpp/src/hash/md5_hash.cu index 62946902960..7a3c3526e75 100644 --- a/cpp/src/hash/md5_hash.cu +++ b/cpp/src/hash/md5_hash.cu @@ -15,10 +15,11 @@ */ #include #include -#include #include #include -#include +#include +#include +#include #include #include #include @@ -281,5 +282,14 @@ std::unique_ptr md5(table_view const& input, } } // namespace detail + +std::unique_ptr md5(table_view const& input, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::md5(input, stream, mr); +} + } // namespace hashing } // namespace cudf diff --git a/cpp/src/hash/murmur_hash.cu b/cpp/src/hash/murmur_hash.cu index 3683a45246f..83bef54d32f 100644 --- a/cpp/src/hash/murmur_hash.cu +++ b/cpp/src/hash/murmur_hash.cu @@ -14,9 +14,11 @@ * limitations under the License. */ #include -#include -#include +#include #include +#include +#include +#include #include #include @@ -57,5 +59,15 @@ std::unique_ptr murmur_hash3_32(table_view const& input, } } // namespace detail + +std::unique_ptr murmur_hash3_32(table_view const& input, + uint32_t seed, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::murmur_hash3_32(input, seed, stream, mr); +} + } // namespace hashing } // namespace cudf diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu index a201f1133ea..239ded002d3 100644 --- a/cpp/src/hash/spark_murmur_hash.cu +++ b/cpp/src/hash/spark_murmur_hash.cu @@ -14,9 +14,10 @@ * limitations under the License. */ #include -#include -#include +#include #include +#include +#include #include #include @@ -425,5 +426,15 @@ std::unique_ptr spark_murmur_hash3_32(table_view const& input, } } // namespace detail + +std::unique_ptr spark_murmur_hash3_32(table_view const& input, + uint32_t seed, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::spark_murmur_hash3_32(input, seed, stream, mr); +} + } // namespace hashing } // namespace cudf diff --git a/cpp/src/hash/unordered_multiset.cuh b/cpp/src/hash/unordered_multiset.cuh index 96e6728df81..e51e603dfb5 100644 --- a/cpp/src/hash/unordered_multiset.cuh +++ b/cpp/src/hash/unordered_multiset.cuh @@ -18,8 +18,9 @@ #include -#include +#include #include +#include #include #include diff --git a/cpp/src/io/json/json_gpu.cu b/cpp/src/io/json/json_gpu.cu index 16aba0e70dc..0cc3b79a679 100644 --- a/cpp/src/io/json/json_gpu.cu +++ b/cpp/src/io/json/json_gpu.cu @@ -19,8 +19,8 @@ #include #include -#include #include +#include #include #include #include diff --git a/cpp/src/io/json/json_tree.cu b/cpp/src/io/json/json_tree.cu index 1a8ddeefdf5..8f94fbcd9c6 100644 --- a/cpp/src/io/json/json_tree.cu +++ b/cpp/src/io/json/json_tree.cu @@ -19,12 +19,12 @@ #include #include -#include #include #include #include -#include #include +#include +#include #include #include diff --git a/cpp/src/io/parquet/page_data.cu b/cpp/src/io/parquet/page_data.cu index e49378485fc..b54d487d5a3 100644 --- a/cpp/src/io/parquet/page_data.cu +++ b/cpp/src/io/parquet/page_data.cu @@ -18,7 +18,8 @@ #include -#include +#include +#include namespace cudf { namespace io { diff --git a/cpp/src/join/join_common_utils.hpp b/cpp/src/join/join_common_utils.hpp index cbccd78049a..64686c7a763 100644 --- a/cpp/src/join/join_common_utils.hpp +++ b/cpp/src/join/join_common_utils.hpp @@ -16,7 +16,7 @@ #pragma once #include -#include +#include #include #include #include diff --git a/cpp/src/partitioning/partitioning.cu b/cpp/src/partitioning/partitioning.cu index 12e1a00b8ba..3e389fcfd42 100644 --- a/cpp/src/partitioning/partitioning.cu +++ b/cpp/src/partitioning/partitioning.cu @@ -21,8 +21,8 @@ #include #include #include -#include #include +#include #include #include #include diff --git a/cpp/src/stream_compaction/stream_compaction_common.hpp b/cpp/src/stream_compaction/stream_compaction_common.hpp index eb57a62fd71..4e887b1199b 100644 --- a/cpp/src/stream_compaction/stream_compaction_common.hpp +++ b/cpp/src/stream_compaction/stream_compaction_common.hpp @@ -15,7 +15,7 @@ */ #pragma once -#include +#include #include #include diff --git a/cpp/src/text/minhash.cu b/cpp/src/text/minhash.cu index 6658d574dcc..e73654fdf4c 100644 --- a/cpp/src/text/minhash.cu +++ b/cpp/src/text/minhash.cu @@ -20,12 +20,12 @@ #include #include #include -#include #include #include #include #include -#include +#include +#include #include #include #include diff --git a/cpp/src/text/subword/bpe_tokenizer.cu b/cpp/src/text/subword/bpe_tokenizer.cu index 2ce29ec8d5c..1cd83cf4c48 100644 --- a/cpp/src/text/subword/bpe_tokenizer.cu +++ b/cpp/src/text/subword/bpe_tokenizer.cu @@ -22,7 +22,7 @@ #include #include #include -#include +#include #include #include #include diff --git a/cpp/src/text/subword/bpe_tokenizer.cuh b/cpp/src/text/subword/bpe_tokenizer.cuh index 24b10fc4a36..df367f49a18 100644 --- a/cpp/src/text/subword/bpe_tokenizer.cuh +++ b/cpp/src/text/subword/bpe_tokenizer.cuh @@ -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. @@ -21,7 +21,7 @@ #include #include -#include +#include #include diff --git a/cpp/src/text/subword/load_merges_file.cu b/cpp/src/text/subword/load_merges_file.cu index 68294ac882b..5d465be9cbe 100644 --- a/cpp/src/text/subword/load_merges_file.cu +++ b/cpp/src/text/subword/load_merges_file.cu @@ -21,8 +21,8 @@ #include #include #include -#include #include +#include #include #include diff --git a/cpp/tests/io/json_tree.cpp b/cpp/tests/io/json_tree.cpp index a81348872cf..2e95fe6cdd9 100644 --- a/cpp/tests/io/json_tree.cpp +++ b/cpp/tests/io/json_tree.cpp @@ -17,8 +17,8 @@ #include #include -#include #include +#include #include #include #include From 8c4187b258677f2279c5f3cf2f0079a3317a7934 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 11 Jul 2023 15:23:30 -0400 Subject: [PATCH 2/8] put hash utils in hashing namespace --- .../cudf/detail/utilities/hash_functions.cuh | 381 ------------------ .../cudf/hashing/detail/hash_functions.cuh | 113 +----- .../cudf/hashing/detail/murmur_hash32.cuh | 18 +- cpp/include/cudf/join.hpp | 7 +- .../cudf/table/experimental/row_operators.cuh | 2 +- cpp/src/groupby/hash/groupby.cu | 12 +- cpp/src/hash/concurrent_unordered_map.cuh | 2 +- cpp/src/hash/md5_hash.cu | 99 ++++- cpp/src/hash/murmur_hash.cu | 2 +- cpp/src/hash/spark_murmur_hash.cu | 12 +- cpp/src/hash/unordered_multiset.cuh | 4 +- cpp/src/io/json/json_gpu.cu | 4 +- cpp/src/io/json/json_tree.cu | 10 +- cpp/src/io/parquet/chunk_dict.cu | 2 +- cpp/src/io/parquet/page_data.cu | 3 +- cpp/src/join/join_common_utils.hpp | 3 +- cpp/src/join/mixed_join_common_utils.cuh | 3 +- cpp/src/partitioning/partitioning.cu | 4 +- cpp/src/text/minhash.cu | 3 +- cpp/src/text/subword/bpe_tokenizer.cuh | 2 +- 20 files changed, 144 insertions(+), 542 deletions(-) delete mode 100644 cpp/include/cudf/detail/utilities/hash_functions.cuh diff --git a/cpp/include/cudf/detail/utilities/hash_functions.cuh b/cpp/include/cudf/detail/utilities/hash_functions.cuh deleted file mode 100644 index e57822f3fdb..00000000000 --- a/cpp/include/cudf/detail/utilities/hash_functions.cuh +++ /dev/null @@ -1,381 +0,0 @@ -/* - * Copyright (c) 2017-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. - */ - -#pragma once - -#include - -#include -#include -#include -#include -#include -#include - -#include -#include -#include -#include -#include -#include - -namespace cudf { -namespace detail { - -/** - * Normalization of floating point NaNs, passthrough for all other values. - */ -template -T __device__ inline normalize_nans(T const& key) -{ - if constexpr (cudf::is_floating_point()) { - if (std::isnan(key)) { return std::numeric_limits::quiet_NaN(); } - } - return key; -} - -/** - * Normalization of floating point NaNs and zeros, passthrough for all other values. - */ -template -T __device__ inline normalize_nans_and_zeros(T const& key) -{ - if constexpr (cudf::is_floating_point()) { - if (key == T{0.0}) { return T{0.0}; } - } - return normalize_nans(key); -} - -__device__ inline uint32_t rotate_bits_left(uint32_t x, uint32_t r) -{ - // This function is equivalent to (x << r) | (x >> (32 - r)) - return __funnelshift_l(x, x, r); -} - -__device__ inline uint32_t rotate_bits_right(uint32_t x, uint32_t r) -{ - // This function is equivalent to (x >> r) | (x << (32 - r)) - return __funnelshift_r(x, x, r); -} - -__device__ inline uint64_t rotate_bits_right(uint64_t x, uint32_t r) -{ - return (x >> r) | (x << (64 - r)); -} - -// Swap the endianness of a 32 bit value -__device__ inline uint32_t swap_endian(uint32_t x) -{ - // The selector 0x0123 reverses the byte order - return __byte_perm(x, 0, 0x0123); -} - -// Swap the endianness of a 64 bit value -// There is no CUDA intrinsic for permuting bytes in 64 bit integers -__device__ inline uint64_t swap_endian(uint64_t x) -{ - // Reverse the endianness of each 32 bit section - uint32_t low_bits = swap_endian(static_cast(x)); - uint32_t high_bits = swap_endian(static_cast(x >> 32)); - // Reassemble a 64 bit result, swapping the low bits and high bits - return (static_cast(low_bits) << 32) | (static_cast(high_bits)); -}; - -template -struct hash_circular_buffer { - uint8_t storage[capacity]; - uint8_t* cur; - int available_space{capacity}; - hash_step_callable hash_step; - - __device__ inline hash_circular_buffer(hash_step_callable hash_step) - : cur{storage}, hash_step{hash_step} - { - } - - __device__ inline void put(uint8_t const* in, int size) - { - int copy_start = 0; - while (size >= available_space) { - // The buffer will be filled by this chunk of data. Copy a chunk of the - // data to fill the buffer and trigger a hash step. - memcpy(cur, in + copy_start, available_space); - hash_step(storage); - size -= available_space; - copy_start += available_space; - cur = storage; - available_space = capacity; - } - // The buffer will not be filled by the remaining data. That is, `size >= 0 - // && size < capacity`. We copy the remaining data into the buffer but do - // not trigger a hash step. - memcpy(cur, in + copy_start, size); - cur += size; - available_space -= size; - } - - __device__ inline void pad(int const space_to_leave) - { - if (space_to_leave > available_space) { - memset(cur, 0x00, available_space); - hash_step(storage); - cur = storage; - available_space = capacity; - } - memset(cur, 0x00, available_space - space_to_leave); - cur += available_space - space_to_leave; - available_space = space_to_leave; - } - - __device__ inline uint8_t const& operator[](int idx) const { return storage[idx]; } -}; - -// Get a uint8_t pointer to a column element and its size as a pair. -template -auto __device__ inline get_element_pointer_and_size(Element const& element) -{ - if constexpr (is_fixed_width() && !is_chrono()) { - return thrust::make_pair(reinterpret_cast(&element), sizeof(Element)); - } else { - CUDF_UNREACHABLE("Unsupported type."); - } -} - -template <> -auto __device__ inline get_element_pointer_and_size(string_view const& element) -{ - return thrust::make_pair(reinterpret_cast(element.data()), element.size_bytes()); -} - -/** - * Modified GPU implementation of - * https://johnnylee-sde.github.io/Fast-unsigned-integer-to-hex-string/ - * Copyright (c) 2015 Barry Clark - * Licensed under the MIT license. - * See file LICENSE for detail or copy at https://opensource.org/licenses/MIT - */ -void __device__ inline uint32ToLowercaseHexString(uint32_t num, char* destination) -{ - // Transform 0xABCD'1234 => 0x0000'ABCD'0000'1234 => 0x0B0A'0D0C'0201'0403 - uint64_t x = num; - x = ((x & 0xFFFF'0000u) << 16) | ((x & 0xFFFF)); - x = ((x & 0x000F'0000'000Fu) << 8) | ((x & 0x00F0'0000'00F0u) >> 4) | - ((x & 0x0F00'0000'0F00u) << 16) | ((x & 0xF000'0000'F000) << 4); - - // Calculate a mask of ascii value offsets for bytes that contain alphabetical hex digits - uint64_t offsets = (((x + 0x0606'0606'0606'0606) >> 4) & 0x0101'0101'0101'0101) * 0x27; - - x |= 0x3030'3030'3030'3030; - x += offsets; - std::memcpy(destination, reinterpret_cast(&x), 8); -} - -// MurmurHash3_32 implementation from -// https://github.com/aappleby/smhasher/blob/master/src/MurmurHash3.cpp -//----------------------------------------------------------------------------- -// MurmurHash3 was written by Austin Appleby, and is placed in the public -// domain. The author hereby disclaims copyright to this source code. -// Note - The x86 and x64 versions do _not_ produce the same results, as the -// algorithms are optimized for their respective platforms. You can still -// compile and run any of them on any platform, but your performance with the -// non-native version will be less than optimal. -template -struct MurmurHash3_32 { - using result_type = hash_value_type; - - constexpr MurmurHash3_32() = default; - constexpr MurmurHash3_32(uint32_t seed) : m_seed(seed) {} - - [[nodiscard]] __device__ inline uint32_t fmix32(uint32_t h) const - { - h ^= h >> 16; - h *= 0x85ebca6b; - h ^= h >> 13; - h *= 0xc2b2ae35; - h ^= h >> 16; - return h; - } - - [[nodiscard]] __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 const block = reinterpret_cast(data + offset); - return block[0] | (block[1] << 8) | (block[2] << 16) | (block[3] << 24); - } - - [[nodiscard]] result_type __device__ inline operator()(Key const& key) const - { - return compute(detail::normalize_nans_and_zeros(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 len, - cudf::size_type tail_offset, - result_type h) const - { - // Process remaining bytes that do not fill a four-byte chunk. - uint32_t k1 = 0; - switch (len % 4) { - case 3: k1 ^= std::to_integer(data[tail_offset + 2]) << 16; [[fallthrough]]; - case 2: k1 ^= std::to_integer(data[tail_offset + 1]) << 8; [[fallthrough]]; - case 1: - k1 ^= std::to_integer(data[tail_offset]); - k1 *= c1; - k1 = cudf::detail::rotate_bits_left(k1, rot_c1); - k1 *= c2; - h ^= k1; - }; - return h; - } - - result_type __device__ compute_bytes(std::byte const* data, cudf::size_type const len) const - { - constexpr cudf::size_type BLOCK_SIZE = 4; - cudf::size_type const nblocks = len / BLOCK_SIZE; - cudf::size_type const tail_offset = nblocks * BLOCK_SIZE; - result_type h = m_seed; - - // Process all four-byte chunks. - for (cudf::size_type i = 0; i < nblocks; i++) { - uint32_t k1 = getblock32(data, i * BLOCK_SIZE); - k1 *= c1; - k1 = cudf::detail::rotate_bits_left(k1, rot_c1); - k1 *= c2; - h ^= k1; - h = cudf::detail::rotate_bits_left(h, rot_c2); - h = h * 5 + c3; - } - - h = compute_remaining_bytes(data, len, tail_offset, h); - - // Finalize hash. - h ^= len; - h = fmix32(h); - return h; - } - - private: - uint32_t m_seed{cudf::DEFAULT_HASH_SEED}; - static constexpr uint32_t c1 = 0xcc9e2d51; - static constexpr uint32_t c2 = 0x1b873593; - static constexpr uint32_t c3 = 0xe6546b64; - static constexpr uint32_t rot_c1 = 15; - static constexpr uint32_t rot_c2 = 13; -}; - -template <> -hash_value_type __device__ inline MurmurHash3_32::operator()(bool const& key) const -{ - return compute(static_cast(key)); -} - -template <> -hash_value_type __device__ inline MurmurHash3_32::operator()(float const& key) const -{ - return compute(detail::normalize_nans_and_zeros(key)); -} - -template <> -hash_value_type __device__ inline MurmurHash3_32::operator()(double const& key) const -{ - return compute(detail::normalize_nans_and_zeros(key)); -} - -template <> -hash_value_type __device__ inline MurmurHash3_32::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 MurmurHash3_32::operator()( - numeric::decimal32 const& key) const -{ - return compute(key.value()); -} - -template <> -hash_value_type __device__ inline MurmurHash3_32::operator()( - numeric::decimal64 const& key) const -{ - return compute(key.value()); -} - -template <> -hash_value_type __device__ inline MurmurHash3_32::operator()( - numeric::decimal128 const& key) const -{ - return compute(key.value()); -} - -template <> -hash_value_type __device__ inline MurmurHash3_32::operator()( - cudf::list_view const& key) const -{ - CUDF_UNREACHABLE("List column hashing is not supported"); -} - -template <> -hash_value_type __device__ inline MurmurHash3_32::operator()( - cudf::struct_view const& key) const -{ - CUDF_UNREACHABLE("Direct hashing of struct_view is not supported"); -} - -/** - * @brief This hash function simply returns the value that is asked to be hash - * reinterpreted as the result_type of the functor. - */ -template -struct IdentityHash { - using result_type = hash_value_type; - IdentityHash() = default; - constexpr IdentityHash(uint32_t seed) : m_seed(seed) {} - - template - constexpr std::enable_if_t, return_type> operator()( - Key const& key) const - { - CUDF_UNREACHABLE("IdentityHash does not support this data type"); - } - - template - constexpr std::enable_if_t, return_type> operator()( - Key const& key) const - { - return static_cast(key); - } - - private: - uint32_t m_seed{cudf::DEFAULT_HASH_SEED}; -}; - -template -using default_hash = MurmurHash3_32; - -} // namespace detail -} // namespace cudf diff --git a/cpp/include/cudf/hashing/detail/hash_functions.cuh b/cpp/include/cudf/hashing/detail/hash_functions.cuh index 96f5b5b067d..769d1ad552e 100644 --- a/cpp/include/cudf/hashing/detail/hash_functions.cuh +++ b/cpp/include/cudf/hashing/detail/hash_functions.cuh @@ -23,8 +23,7 @@ #include -namespace cudf { -namespace detail { +namespace cudf::hashing::detail { template struct MurmurHash3_32; @@ -73,113 +72,6 @@ __device__ inline uint64_t rotate_bits_right(uint64_t x, uint32_t r) return (x >> r) | (x << (64 - r)); } -// Swap the endianness of a 32 bit value -__device__ inline uint32_t swap_endian(uint32_t x) -{ - // The selector 0x0123 reverses the byte order - return __byte_perm(x, 0, 0x0123); -} - -// Swap the endianness of a 64 bit value -// There is no CUDA intrinsic for permuting bytes in 64 bit integers -__device__ inline uint64_t swap_endian(uint64_t x) -{ - // Reverse the endianness of each 32 bit section - uint32_t low_bits = swap_endian(static_cast(x)); - uint32_t high_bits = swap_endian(static_cast(x >> 32)); - // Reassemble a 64 bit result, swapping the low bits and high bits - return (static_cast(low_bits) << 32) | (static_cast(high_bits)); -}; - -template -struct hash_circular_buffer { - uint8_t storage[capacity]; - uint8_t* cur; - int available_space{capacity}; - hash_step_callable hash_step; - - __device__ inline hash_circular_buffer(hash_step_callable hash_step) - : cur{storage}, hash_step{hash_step} - { - } - - __device__ inline void put(uint8_t const* in, int size) - { - int copy_start = 0; - while (size >= available_space) { - // The buffer will be filled by this chunk of data. Copy a chunk of the - // data to fill the buffer and trigger a hash step. - memcpy(cur, in + copy_start, available_space); - hash_step(storage); - size -= available_space; - copy_start += available_space; - cur = storage; - available_space = capacity; - } - // The buffer will not be filled by the remaining data. That is, `size >= 0 - // && size < capacity`. We copy the remaining data into the buffer but do - // not trigger a hash step. - memcpy(cur, in + copy_start, size); - cur += size; - available_space -= size; - } - - __device__ inline void pad(int const space_to_leave) - { - if (space_to_leave > available_space) { - memset(cur, 0x00, available_space); - hash_step(storage); - cur = storage; - available_space = capacity; - } - memset(cur, 0x00, available_space - space_to_leave); - cur += available_space - space_to_leave; - available_space = space_to_leave; - } - - __device__ inline uint8_t const& operator[](int idx) const { return storage[idx]; } -}; - -// Get a uint8_t pointer to a column element and its size as a pair. -template -auto __device__ inline get_element_pointer_and_size(Element const& element) -{ - if constexpr (is_fixed_width() && !is_chrono()) { - return thrust::make_pair(reinterpret_cast(&element), sizeof(Element)); - } else { - CUDF_UNREACHABLE("Unsupported type."); - } -} - -template <> -auto __device__ inline get_element_pointer_and_size(string_view const& element) -{ - return thrust::make_pair(reinterpret_cast(element.data()), element.size_bytes()); -} - -/** - * Modified GPU implementation of - * https://johnnylee-sde.github.io/Fast-unsigned-integer-to-hex-string/ - * Copyright (c) 2015 Barry Clark - * Licensed under the MIT license. - * See file LICENSE for detail or copy at https://opensource.org/licenses/MIT - */ -void __device__ inline uint32ToLowercaseHexString(uint32_t num, char* destination) -{ - // Transform 0xABCD'1234 => 0x0000'ABCD'0000'1234 => 0x0B0A'0D0C'0201'0403 - uint64_t x = num; - x = ((x & 0xFFFF'0000u) << 16) | ((x & 0xFFFF)); - x = ((x & 0x000F'0000'000Fu) << 8) | ((x & 0x00F0'0000'00F0u) >> 4) | - ((x & 0x0F00'0000'0F00u) << 16) | ((x & 0xF000'0000'F000) << 4); - - // Calculate a mask of ascii value offsets for bytes that contain alphabetical hex digits - uint64_t offsets = (((x + 0x0606'0606'0606'0606) >> 4) & 0x0101'0101'0101'0101) * 0x27; - - x |= 0x3030'3030'3030'3030; - x += offsets; - std::memcpy(destination, reinterpret_cast(&x), 8); -} - /** * @brief This hash function simply returns the value that is asked to be hash * reinterpreted as the result_type of the functor. @@ -208,5 +100,4 @@ struct IdentityHash { uint32_t m_seed{0}; }; -} // namespace detail -} // namespace cudf +} // namespace cudf::hashing::detail diff --git a/cpp/include/cudf/hashing/detail/murmur_hash32.cuh b/cpp/include/cudf/hashing/detail/murmur_hash32.cuh index 26425875068..1a63f8c45ca 100644 --- a/cpp/include/cudf/hashing/detail/murmur_hash32.cuh +++ b/cpp/include/cudf/hashing/detail/murmur_hash32.cuh @@ -26,8 +26,7 @@ #include -namespace cudf { -namespace detail { +namespace cudf::hashing::detail { // MurmurHash3_32 implementation from // https://github.com/aappleby/smhasher/blob/master/src/MurmurHash3.cpp @@ -66,7 +65,7 @@ struct MurmurHash3_32 { [[nodiscard]] result_type __device__ inline operator()(Key const& key) const { - return compute(detail::normalize_nans_and_zeros(key)); + return compute(normalize_nans_and_zeros(key)); } template @@ -88,7 +87,7 @@ struct MurmurHash3_32 { case 1: k1 ^= std::to_integer(data[tail_offset]); k1 *= c1; - k1 = cudf::detail::rotate_bits_left(k1, rot_c1); + k1 = rotate_bits_left(k1, rot_c1); k1 *= c2; h ^= k1; }; @@ -106,10 +105,10 @@ struct MurmurHash3_32 { for (cudf::size_type i = 0; i < nblocks; i++) { uint32_t k1 = getblock32(data, i * BLOCK_SIZE); k1 *= c1; - k1 = cudf::detail::rotate_bits_left(k1, rot_c1); + k1 = rotate_bits_left(k1, rot_c1); k1 *= c2; h ^= k1; - h = cudf::detail::rotate_bits_left(h, rot_c2); + h = rotate_bits_left(h, rot_c2); h = h * 5 + c3; } @@ -139,13 +138,13 @@ hash_value_type __device__ inline MurmurHash3_32::operator()(bool const& k template <> hash_value_type __device__ inline MurmurHash3_32::operator()(float const& key) const { - return compute(detail::normalize_nans_and_zeros(key)); + return compute(normalize_nans_and_zeros(key)); } template <> hash_value_type __device__ inline MurmurHash3_32::operator()(double const& key) const { - return compute(detail::normalize_nans_and_zeros(key)); + return compute(normalize_nans_and_zeros(key)); } template <> @@ -192,5 +191,4 @@ hash_value_type __device__ inline MurmurHash3_32::operator()( CUDF_UNREACHABLE("Direct hashing of struct_view is not supported"); } -} // namespace detail -} // namespace cudf +} // namespace cudf::hashing::detail diff --git a/cpp/include/cudf/join.hpp b/cpp/include/cudf/join.hpp index 61f8c13bb77..e1455bd325c 100644 --- a/cpp/include/cudf/join.hpp +++ b/cpp/include/cudf/join.hpp @@ -34,10 +34,11 @@ namespace cudf { // forward declaration -namespace detail { +namespace hashing::detail { template class MurmurHash3_32; - +} // namespace hashing::detail +namespace detail { template class hash_join; } // namespace detail @@ -272,7 +273,7 @@ enum class nullable_join : bool { YES, NO }; class hash_join { public: using impl_type = typename cudf::detail::hash_join< - cudf::detail::MurmurHash3_32>; ///< Implementation type + cudf::hashing::detail::MurmurHash3_32>; ///< Implementation type hash_join() = delete; ~hash_join(); diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index 7e876d6cbad..1aa0f21fca2 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -1942,7 +1942,7 @@ class row_hasher { * @param seed The seed to use for the hash function * @return A hash operator to use on the device */ - template