Add SHA-1 and SHA-2 hash functions. (#14391)

This PR adds support for SHA-1 and SHA-2 (SHA-256, SHA-512, and truncated digests SHA-224, SHA-384) hash functions. Resolves #8641. Replaces #9215. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Robert Maynard (https://github.com/robertmaynard) - Matthew Roeschke (https://github.com/mroeschke) - David Wendt (https://github.com/davidwendt) - https://github.com/nvdbaranec URL: #14391
rapidsai · Jan 22, 2024 · 42d8d78 · 42d8d78
1 parent b1468a5
commit 42d8d78
Show file tree

Hide file tree

Showing 23 changed files with 2,348 additions and 97 deletions.
diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
@@ -344,6 +344,11 @@ add_library(
   src/hash/md5_hash.cu
   src/hash/murmurhash3_x86_32.cu
   src/hash/murmurhash3_x64_128.cu
+  src/hash/sha1_hash.cu
+  src/hash/sha224_hash.cu
+  src/hash/sha256_hash.cu
+  src/hash/sha384_hash.cu
+  src/hash/sha512_hash.cu
   src/hash/spark_murmurhash3_x86_32.cu
   src/hash/xxhash_64.cu
   src/interop/dlpack.cpp

diff --git a/cpp/benchmarks/hashing/hash.cpp b/cpp/benchmarks/hashing/hash.cpp
@@ -67,6 +67,36 @@ static void bench_hash(nvbench::state& state)
 
     state.exec(nvbench::exec_tag::sync,
                [&](nvbench::launch& launch) { auto result = cudf::hashing::md5(data->view()); });
+  } else if (hash_name == "sha1") {
+    // sha1 creates a 40-byte string
+    state.add_global_memory_writes<nvbench::int8_t>(40 * num_rows);
+
+    state.exec(nvbench::exec_tag::sync,
+               [&](nvbench::launch& launch) { auto result = cudf::hashing::sha1(data->view()); });
+  } else if (hash_name == "sha224") {
+    // sha224 creates a 56-byte string
+    state.add_global_memory_writes<nvbench::int8_t>(56 * num_rows);
+
+    state.exec(nvbench::exec_tag::sync,
+               [&](nvbench::launch& launch) { auto result = cudf::hashing::sha224(data->view()); });
+  } else if (hash_name == "sha256") {
+    // sha256 creates a 64-byte string
+    state.add_global_memory_writes<nvbench::int8_t>(64 * num_rows);
+
+    state.exec(nvbench::exec_tag::sync,
+               [&](nvbench::launch& launch) { auto result = cudf::hashing::sha256(data->view()); });
+  } else if (hash_name == "sha384") {
+    // sha384 creates a 96-byte string
+    state.add_global_memory_writes<nvbench::int8_t>(96 * num_rows);
+
+    state.exec(nvbench::exec_tag::sync,
+               [&](nvbench::launch& launch) { auto result = cudf::hashing::sha384(data->view()); });
+  } else if (hash_name == "sha512") {
+    // sha512 creates a 128-byte string
+    state.add_global_memory_writes<nvbench::int8_t>(128 * num_rows);
+
+    state.exec(nvbench::exec_tag::sync,
+               [&](nvbench::launch& launch) { auto result = cudf::hashing::sha512(data->view()); });
   } else if (hash_name == "spark_murmurhash3_x86_32") {
     state.add_global_memory_writes<nvbench::int32_t>(num_rows);
 
@@ -82,4 +112,12 @@ NVBENCH_BENCH(bench_hash)
   .set_name("hashing")
   .add_int64_axis("num_rows", {65536, 16777216})
   .add_float64_axis("nulls", {0.0, 0.1})
-  .add_string_axis("hash_name", {"murmurhash3_x86_32", "md5", "spark_murmurhash3_x86_32"});
+  .add_string_axis("hash_name",
+                   {"murmurhash3_x86_32",
+                    "md5",
+                    "sha1",
+                    "sha224",
+                    "sha256",
+                    "sha384",
+                    "sha512",
+                    "spark_murmurhash3_x86_32"});
diff --git a/cpp/include/cudf/hashing.hpp b/cpp/include/cudf/hashing.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -145,6 +145,76 @@ std::unique_ptr<column> md5(
   rmm::cuda_stream_view stream        = cudf::get_default_stream(),
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
+/**
+ * @brief Computes the SHA-1 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
+ * @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<column> sha1(
+  table_view const& input,
+  rmm::cuda_stream_view stream        = cudf::get_default_stream(),
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+/**
+ * @brief Computes the SHA-224 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
+ * @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<column> sha224(
+  table_view const& input,
+  rmm::cuda_stream_view stream        = cudf::get_default_stream(),
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+/**
+ * @brief Computes the SHA-256 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
+ * @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<column> sha256(
+  table_view const& input,
+  rmm::cuda_stream_view stream        = cudf::get_default_stream(),
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+/**
+ * @brief Computes the SHA-384 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
+ * @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<column> sha384(
+  table_view const& input,
+  rmm::cuda_stream_view stream        = cudf::get_default_stream(),
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+/**
+ * @brief Computes the SHA-512 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
+ * @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<column> sha512(
+  table_view const& input,
+  rmm::cuda_stream_view stream        = cudf::get_default_stream(),
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
 /**
  * @brief Computes the XXHash_64 hash value of each row in the given table
  *

diff --git a/cpp/include/cudf/hashing/detail/hash_functions.cuh b/cpp/include/cudf/hashing/detail/hash_functions.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2017-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -68,4 +68,45 @@ __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<uint32_t>(x));
+  uint32_t high_bits = swap_endian(static_cast<uint32_t>(x >> 32));
+  // Reassemble a 64 bit result, swapping the low bits and high bits
+  return (static_cast<uint64_t>(low_bits) << 32) | (static_cast<uint64_t>(high_bits));
+};
+
+/**
+ * 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
+ */
+__device__ inline void 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<uint8_t*>(&x), 8);
+}
+
 }  // namespace cudf::hashing::detail
diff --git a/cpp/include/cudf/hashing/detail/hashing.hpp b/cpp/include/cudf/hashing/detail/hashing.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -46,6 +46,26 @@ std::unique_ptr<column> md5(table_view const& input,
                             rmm::cuda_stream_view stream,
                             rmm::mr::device_memory_resource* mr);
 
+std::unique_ptr<column> sha1(table_view const& input,
+                             rmm::cuda_stream_view stream,
+                             rmm::mr::device_memory_resource* mr);
+
+std::unique_ptr<column> sha224(table_view const& input,
+                               rmm::cuda_stream_view stream,
+                               rmm::mr::device_memory_resource* mr);
+
+std::unique_ptr<column> sha256(table_view const& input,
+                               rmm::cuda_stream_view stream,
+                               rmm::mr::device_memory_resource* mr);
+
+std::unique_ptr<column> sha384(table_view const& input,
+                               rmm::cuda_stream_view stream,
+                               rmm::mr::device_memory_resource* mr);
+
+std::unique_ptr<column> sha512(table_view const& input,
+                               rmm::cuda_stream_view stream,
+                               rmm::mr::device_memory_resource* mr);
+
 std::unique_ptr<column> xxhash_64(table_view const& input,
                                   uint64_t seed,
                                   rmm::cuda_stream_view,

diff --git a/cpp/src/hash/md5_hash.cu b/cpp/src/hash/md5_hash.cu
@@ -108,29 +108,6 @@ auto __device__ inline get_element_pointer_and_size(string_view const& element)
   return thrust::make_pair(reinterpret_cast<uint8_t const*>(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<uint8_t*>(&x), 8);
-}
-
 // The MD5 algorithm and its hash/shift constants are officially specified in
 // RFC 1321. For convenience, these values can also be found on Wikipedia:
 // https://en.wikipedia.org/wiki/MD5

diff --git a/cpp/src/hash/sha1_hash.cu b/cpp/src/hash/sha1_hash.cu
@@ -0,0 +1,81 @@
+/*
+ * Copyright (c) 2024, 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 "sha_hash.cuh"
+
+#include <cudf/column/column.hpp>
+#include <cudf/scalar/scalar.hpp>
+#include <cudf/table/table_view.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
+#include <rmm/mr/device/device_memory_resource.hpp>
+
+#include <memory>
+
+namespace cudf {
+namespace hashing {
+namespace detail {
+
+namespace {
+
+struct sha1_hash_state {
+  uint64_t message_length = 0;
+  uint32_t buffer_length  = 0;
+  uint32_t hash_value[5]  = {0x67452301, 0xefcdab89, 0x98badcfe, 0x10325476, 0xc3d2e1f0};
+  uint8_t buffer[64];
+};
+
+struct SHA1Hash : HashBase<SHA1Hash> {
+  __device__ inline SHA1Hash(char* result_location) : HashBase<SHA1Hash>(result_location) {}
+
+  // Intermediate data type storing the hash state
+  using hash_state = sha1_hash_state;
+  // The word type used by this hash function
+  using sha_word_type = uint32_t;
+  // Number of bytes processed in each hash step
+  static constexpr uint32_t message_chunk_size = 64;
+  // Digest size in bytes
+  static constexpr uint32_t digest_size = 40;
+  // Number of bytes used for the message length
+  static constexpr uint32_t message_length_size = 8;
+
+  __device__ inline void hash_step(hash_state& state) { sha1_hash_step(state); }
+
+  hash_state state;
+};
+
+}  // namespace
+
+std::unique_ptr<column> sha1(table_view const& input,
+                             rmm::cuda_stream_view stream,
+                             rmm::mr::device_memory_resource* mr)
+{
+  return sha_hash<SHA1Hash>(input, stream, mr);
+}
+
+}  // namespace detail
+
+std::unique_ptr<column> sha1(table_view const& input,
+                             rmm::cuda_stream_view stream,
+                             rmm::mr::device_memory_resource* mr)
+{
+  CUDF_FUNC_RANGE();
+  return detail::sha1(input, stream, mr);
+}
+
+}  // namespace hashing
+}  // namespace cudf