Add SHA-1 and SHA-2 hash functions.

cpp/CMakeLists.txt

@@ -249,6 +249,7 @@ add_library(cudf
     src/hash/hashing.cu
     src/hash/md5_hash.cu
     src/hash/murmur_hash.cu
+    src/hash/sha_hash.cu
     src/interop/dlpack.cpp
     src/interop/from_arrow.cu
     src/interop/to_arrow.cu
@@ -468,6 +469,8 @@ set_target_properties(cudf
                INTERFACE_POSITION_INDEPENDENT_CODE ON
 )
 
+set_source_files_properties(src/hash/sha_hash.cu PROPERTIES COMPILE_OPTIONS "-G")
+
 target_compile_options(cudf
             PRIVATE "$<$<COMPILE_LANGUAGE:CXX>:${CUDF_CXX_FLAGS}>"
                     "$<$<COMPILE_LANGUAGE:CUDA>:${CUDF_CUDA_FLAGS}>"

cpp/benchmarks/hashing/hash_benchmark.cpp

@@ -41,11 +41,16 @@ static void BM_hash(benchmark::State& state, cudf::hash_id hid)
   (::benchmark::State & st) { BM_hash(st, cudf::hash_id::name); } \
   BENCHMARK_REGISTER_F(HashBenchmark, name)                       \
     ->RangeMultiplier(4)                                          \
-    ->Ranges({{1 << 14, 1 << 24}})                                \
+    ->Ranges({{1 << 14, 1 << 23}})                                \
     ->UseManualTime()                                             \
     ->Unit(benchmark::kMillisecond);
 
 HASH_BENCHMARK_DEFINE(HASH_MURMUR3)
 HASH_BENCHMARK_DEFINE(HASH_MD5)
 HASH_BENCHMARK_DEFINE(HASH_SERIAL_MURMUR3)
 HASH_BENCHMARK_DEFINE(HASH_SPARK_MURMUR3)
+HASH_BENCHMARK_DEFINE(HASH_SHA1)
+HASH_BENCHMARK_DEFINE(HASH_SHA224)
+HASH_BENCHMARK_DEFINE(HASH_SHA256)
+HASH_BENCHMARK_DEFINE(HASH_SHA384)
+HASH_BENCHMARK_DEFINE(HASH_SHA512)

cpp/include/cudf/detail/hashing.hpp

@@ -56,6 +56,31 @@ std::unique_ptr<column> serial_murmur_hash3_32(
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
+std::unique_ptr<column> sha1_hash(
+  table_view const& input,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+std::unique_ptr<column> sha224_hash(
+  table_view const& input,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+std::unique_ptr<column> sha256_hash(
+  table_view const& input,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+std::unique_ptr<column> sha384_hash(
+  table_view const& input,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+std::unique_ptr<column> sha512_hash(
+  table_view const& input,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
 /* Copyright 2005-2014 Daniel James.
  *
  * Use, modification and distribution is subject to the Boost Software

cpp/include/cudf/detail/utilities/hash_functions.cuh

@@ -43,6 +43,107 @@ T CUDA_DEVICE_CALLABLE normalize_nans_and_zeros(T const& key)
   return key;
 }
 
+CUDA_DEVICE_CALLABLE uint32_t rotate_bits_left(uint32_t x, int8_t r)
+{
+  // This function is equivalent to (x << r) | (x >> (32 - r))
+  return __funnelshift_l(x, x, r);
+}
+
+CUDA_DEVICE_CALLABLE uint32_t rotate_bits_right(uint32_t x, int8_t r)
+{
+  // This function is equivalent to (x >> r) | (x << (32 - r))
+  return __funnelshift_r(x, x, r);
+}
+
+CUDA_DEVICE_CALLABLE uint64_t rotate_bits_right(uint64_t x, int8_t r)
+{
+  return (x >> r) | (x << (64 - r));
+}
+
+// Swap the endianness of a 32 bit value
+CUDA_DEVICE_CALLABLE 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
+CUDA_DEVICE_CALLABLE 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));
+};
+
+template <int capacity, typename hash_step_callable>
+struct hash_circular_buffer {
+  uint8_t storage[capacity];
+  uint8_t* cur;
+  int available_space{capacity};
+  hash_step_callable hash_step;
+
+  CUDA_DEVICE_CALLABLE hash_circular_buffer(hash_step_callable hash_step)
+    : cur{storage}, hash_step{hash_step}
+  {
+  }
+
+  CUDA_DEVICE_CALLABLE 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;
+  }
+
+  CUDA_DEVICE_CALLABLE 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;
+  }
+
+  CUDA_DEVICE_CALLABLE const uint8_t& operator[](int idx) const { return storage[idx]; }
+};
+
+// Get a uint8_t pointer to a column element and its size as a pair.
+template <typename Element>
+auto CUDA_DEVICE_CALLABLE get_element_pointer_and_size(Element const& element)
+{
+  if constexpr (is_fixed_width<Element>() && !is_chrono<Element>()) {
+    return thrust::make_pair(reinterpret_cast<uint8_t const*>(&element), sizeof(Element));
+  } else {
+    cudf_assert(false && "Unsupported type.");
+  }
+}
+
+template <>
+auto CUDA_DEVICE_CALLABLE 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/

cpp/include/cudf/types.hpp

@@ -335,7 +335,12 @@ enum class hash_id {
   HASH_MURMUR3,         ///< Murmur3 hash function
   HASH_MD5,             ///< MD5 hash function
   HASH_SERIAL_MURMUR3,  ///< Serial Murmur3 hash function
-  HASH_SPARK_MURMUR3    ///< Spark Murmur3 hash function
+  HASH_SPARK_MURMUR3,   ///< Spark Murmur3 hash function
+  HASH_SHA1,            ///< SHA-1 hash function
+  HASH_SHA224,          ///< SHA-224 hash function
+  HASH_SHA256,          ///< SHA-256 hash function
+  HASH_SHA384,          ///< SHA-384 hash function
+  HASH_SHA512           ///< SHA-512 hash function
 };
 
 /**

cpp/src/hash/hashing.cu

@@ -119,7 +119,12 @@ std::unique_ptr<column> hash(table_view const& input,
       return serial_murmur_hash3_32<MurmurHash3_32>(input, seed, stream, mr);
     case (hash_id::HASH_SPARK_MURMUR3):
       return serial_murmur_hash3_32<SparkMurmurHash3_32>(input, seed, stream, mr);
-    default: return nullptr;
+    case (hash_id::HASH_SHA1): return sha1_hash(input, stream, mr);
+    case (hash_id::HASH_SHA224): return sha224_hash(input, stream, mr);
+    case (hash_id::HASH_SHA256): return sha256_hash(input, stream, mr);
+    case (hash_id::HASH_SHA384): return sha384_hash(input, stream, mr);
+    case (hash_id::HASH_SHA512): return sha512_hash(input, stream, mr);
+    default: CUDF_FAIL("Unsupported hash function.");
   }
 }
 

cpp/src/hash/md5_hash.cu

@@ -13,9 +13,11 @@
  * See the License for the specific language governing permissions and
  * limitations under the License.
  */
+#include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/hashing.hpp>
 #include <cudf/detail/iterator.cuh>
+#include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/utilities/hash_functions.cuh>
 #include <cudf/lists/lists_column_view.hpp>
 #include <cudf/scalar/scalar.hpp>
@@ -55,72 +57,6 @@ const __constant__ uint32_t md5_hash_constants[64] = {
   0x6fa87e4f, 0xfe2ce6e0, 0xa3014314, 0x4e0811a1, 0xf7537e82, 0xbd3af235, 0x2ad7d2bb, 0xeb86d391,
 };
 
-template <int capacity, typename hash_step_callable>
-struct hash_circular_buffer {
-  uint8_t storage[capacity];
-  uint8_t* cur;
-  int available_space{capacity};
-  hash_step_callable hash_step;
-
-  CUDA_DEVICE_CALLABLE hash_circular_buffer(hash_step_callable hash_step)
-    : cur{storage}, hash_step{hash_step}
-  {
-  }
-
-  CUDA_DEVICE_CALLABLE 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;
-  }
-
-  CUDA_DEVICE_CALLABLE 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;
-  }
-
-  CUDA_DEVICE_CALLABLE const uint8_t& operator[](int idx) const { return storage[idx]; }
-};
-
-// Get a uint8_t pointer to a column element and its size as a pair.
-template <typename Element>
-auto CUDA_DEVICE_CALLABLE get_element_pointer_and_size(Element const& element)
-{
-  if constexpr (is_fixed_width<Element>() && !is_chrono<Element>()) {
-    return thrust::make_pair(reinterpret_cast<uint8_t const*>(&element), sizeof(Element));
-  } else {
-    cudf_assert(false && "Unsupported type.");
-  }
-}
-
-template <>
-auto CUDA_DEVICE_CALLABLE get_element_pointer_and_size(string_view const& element)
-{
-  return thrust::make_pair(reinterpret_cast<uint8_t const*>(element.data()), element.size_bytes());
-}
-
 struct MD5Hasher {
   static constexpr int message_chunk_size = 64;
 
@@ -205,7 +141,7 @@ struct MD5Hasher {
         A = D;
         D = C;
         C = B;
-        B = B + __funnelshift_l(F, F, md5_shift_constants[((j / 16) * 4) + (j % 4)]);
+        B = B + rotate_bits_left(F, md5_shift_constants[((j / 16) * 4) + (j % 4)]);
       }
 
       hash_values[0] += A;
@@ -309,7 +245,8 @@ std::unique_ptr<column> md5_hash(table_view const& input,
   auto chars_view = chars_column->mutable_view();
   auto d_chars    = chars_view.data<char>();
 
-  rmm::device_buffer null_mask{0, stream, mr};
+  // Build an output null mask from the logical AND of all input columns' null masks.
+  rmm::device_buffer null_mask{cudf::detail::bitmask_and(input, stream)};
 
   auto const device_input = table_device_view::create(input, stream);