Skip to content

Commit

Permalink
Fix memory alignment issues in hash computation (#338)
Browse files Browse the repository at this point in the history
When hashing large keys, e.g., strings, we traverse the input key
iteratively in chunks of 4/8 bytes. The current implementation of the
`load_chunk` function falsely assumes that the start of the key is
always aligned to the chunk size, which is not always the case (see
[discussion](rapidsai/cudf#13612 (comment))).

Additionally, this PR fixes some uncaught `[-Wmaybe-uninitialized]`
warnings when compiling the unit tests.
  • Loading branch information
sleeepyjack authored Jul 29, 2023
1 parent a2833db commit ff16201
Showing 1 changed file with 4 additions and 2 deletions.
6 changes: 4 additions & 2 deletions include/cuco/detail/hash_functions/utils.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -21,8 +21,10 @@ namespace cuco::detail {
template <typename T, typename U, typename Extent>
constexpr __host__ __device__ T load_chunk(U const* const data, Extent index) noexcept
{
auto const chunks = reinterpret_cast<T const*>(data);
return chunks[index];
auto const bytes = reinterpret_cast<std::byte const*>(data);
T chunk;
memcpy(&chunk, bytes + index * sizeof(T), sizeof(T));
return chunk;
}

}; // namespace cuco::detail

0 comments on commit ff16201

Please sign in to comment.