Migrate hashing operations to pylibcudf

cpp/src/hash/sha_hash.cuh

@@ -512,7 +512,8 @@ std::unique_ptr<column> sha_hash(table_view const& input,
   CUDF_EXPECTS(
     std::all_of(
       input.begin(), input.end(), [](auto const& col) { return sha_leaf_type_check(col.type()); }),
-    "Unsupported column type for hash function.");
+    "Unsupported column type for hash function.",
+    cudf::data_type_error);
 
   // Result column allocation and creation
   auto begin = thrust::make_constant_iterator(Hasher::digest_size);

docs/cudf/source/user_guide/api_docs/pylibcudf/hashing.rst

@@ -0,0 +1,6 @@
+========
+hashing
+========
+
+.. automodule:: cudf._lib.pylibcudf.hashing
+    :members:

docs/cudf/source/user_guide/api_docs/pylibcudf/index.rst

@@ -16,6 +16,7 @@ This page provides API documentation for pylibcudf.
     filling
     gpumemoryview
     groupby
+    hashing
     join
     lists
     merge

python/cudf/cudf/_lib/cpp/hash.pxd

@@ -7,40 +7,52 @@ from libcpp.vector cimport vector
 from cudf._lib.cpp.column.column cimport column
 from cudf._lib.cpp.table.table cimport table
 from cudf._lib.cpp.table.table_view cimport table_view
+from cudf._lib.exception_handler cimport cudf_exception_handler
 
 
 cdef extern from "cudf/hashing.hpp" namespace "cudf::hashing" nogil:
 
     cdef unique_ptr[column] murmurhash3_x86_32 "cudf::hashing::murmurhash3_x86_32" (
         const table_view& input,
         const uint32_t seed
-    ) except +
+    ) except +cudf_exception_handler
+
+    cdef unique_ptr[table] murmurhash3_x64_128 "cudf::hashing::murmurhash3_x64_128" (
+        const table_view& input,
+        const uint64_t seed
+    ) except +cudf_exception_handler
 
     cdef unique_ptr[column] md5 "cudf::hashing::md5" (
         const table_view& input
-    ) except +
+    ) except +cudf_exception_handler
 
     cdef unique_ptr[column] sha1 "cudf::hashing::sha1" (
         const table_view& input
-    ) except +
+    ) except +cudf_exception_handler
 
     cdef unique_ptr[column] sha224 "cudf::hashing::sha224" (
         const table_view& input
-    ) except +
+    ) except +cudf_exception_handler
 
     cdef unique_ptr[column] sha256 "cudf::hashing::sha256" (
         const table_view& input
-    ) except +
+    ) except +cudf_exception_handler
 
     cdef unique_ptr[column] sha384 "cudf::hashing::sha384" (
         const table_view& input
-    ) except +
+    ) except +cudf_exception_handler
 
     cdef unique_ptr[column] sha512 "cudf::hashing::sha512" (
         const table_view& input
-    ) except +
+    ) except +cudf_exception_handler
 
     cdef unique_ptr[column] xxhash_64 "cudf::hashing::xxhash_64" (
         const table_view& input,
         const uint64_t seed
-    ) except +
+    ) except +cudf_exception_handler
+
+    cpdef enum class hash_id(int):
+        HASH_IDENTITY
+        HASH_MURMUR3
+        HASH_SPARK_MURMUR3
+        HASH_MD5

python/cudf/cudf/_lib/hash.pyx

@@ -7,10 +7,14 @@ from libcpp.pair cimport pair
 from libcpp.utility cimport move
 from libcpp.vector cimport vector
 
+from cudf._lib import pylibcudf
+
 cimport cudf._lib.cpp.types as libcudf_types
 from cudf._lib.column cimport Column
-from cudf._lib.cpp.column.column cimport column
-from cudf._lib.cpp.hash cimport (
+from cudf._lib.cpp.partitioning cimport hash_partition as cpp_hash_partition
+from cudf._lib.cpp.table.table cimport table
+from cudf._lib.cpp.table.table_view cimport table_view
+from cudf._lib.pylibcudf.hashing cimport (
     md5,
     murmurhash3_x86_32,
     sha1,
@@ -20,9 +24,6 @@ from cudf._lib.cpp.hash cimport (
     sha512,
     xxhash_64,
 )
-from cudf._lib.cpp.partitioning cimport hash_partition as cpp_hash_partition
-from cudf._lib.cpp.table.table cimport table
-from cudf._lib.cpp.table.table_view cimport table_view
 from cudf._lib.utils cimport columns_from_unique_ptr, table_view_from_columns
 
 
@@ -51,32 +52,24 @@ def hash_partition(list source_columns, object columns_to_hash,
 
 @acquire_spill_lock()
 def hash(list source_columns, str method, int seed=0):
-    cdef table_view c_source_view = table_view_from_columns(source_columns)
-    cdef unique_ptr[column] c_result
+    ctbl = pylibcudf.Table([c.to_pylibcudf(mode="read") for c in source_columns])
     if method == "murmur3":
-        with nogil:
-            c_result = move(murmurhash3_x86_32(c_source_view, seed))
+        return Column.from_pylibcudf(murmurhash3_x86_32(ctbl, seed))
+    elif method == "xxhash64":
+        return Column.from_pylibcudf(xxhash_64(ctbl, seed))
     elif method == "md5":
-        with nogil:
-            c_result = move(md5(c_source_view))
+        return Column.from_pylibcudf(md5(ctbl))
     elif method == "sha1":
-        with nogil:
-            c_result = move(sha1(c_source_view))
+        return Column.from_pylibcudf(sha1(ctbl))
     elif method == "sha224":
-        with nogil:
-            c_result = move(sha224(c_source_view))
+        return Column.from_pylibcudf(sha224(ctbl))
     elif method == "sha256":
-        with nogil:
-            c_result = move(sha256(c_source_view))
+        return Column.from_pylibcudf(sha256(ctbl))
     elif method == "sha384":
-        with nogil:
-            c_result = move(sha384(c_source_view))
+        return Column.from_pylibcudf(sha384(ctbl))
     elif method == "sha512":
-        with nogil:
-            c_result = move(sha512(c_source_view))
-    elif method == "xxhash64":
-        with nogil:
-            c_result = move(xxhash_64(c_source_view, seed))
+        return Column.from_pylibcudf(sha512(ctbl))
     else:
-        raise ValueError(f"Unsupported hash function: {method}")
-    return Column.from_unique_ptr(move(c_result))
+        raise ValueError(
+            f"Unsupported hashing algorithm {method}."
+        )

python/cudf/cudf/_lib/pylibcudf/CMakeLists.txt

@@ -21,6 +21,7 @@ set(cython_sources
     filling.pyx
     gpumemoryview.pyx
     groupby.pyx
+    hashing.pyx
     interop.pyx
     join.pyx
     lists.pyx

python/cudf/cudf/_lib/pylibcudf/__init__.pxd

@@ -8,6 +8,7 @@ from . cimport (
     copying,
     filling,
     groupby,
+    hashing,
     join,
     lists,
     merge,
@@ -40,6 +41,7 @@ __all__ = [
     "filling",
     "gpumemoryview",
     "groupby",
+    "hashing",
     "join",
     "lists",
     "merge",

python/cudf/cudf/_lib/pylibcudf/__init__.py

@@ -7,6 +7,7 @@
     copying,
     filling,
     groupby,
+    hashing,
     interop,
     join,
     lists,
@@ -39,6 +40,7 @@
     "filling",
     "gpumemoryview",
     "groupby",
+    "hashing",
     "interop",
     "join",
     "lists",

python/cudf/cudf/_lib/pylibcudf/hashing.pxd

@@ -0,0 +1,34 @@
+# Copyright (c) 2024, NVIDIA CORPORATION.
+
+from libc.stdint cimport uint32_t, uint64_t
+from cudf._lib.cpp.types cimport size_type
+from cudf._lib.cpp.table cimport table
+from libcpp.vector cimport vector
+
+#from cudf._lib.cpp.hash cimport hash_id
+
+from .column cimport Column
+from .table cimport Table
+
+
+cpdef Column murmurhash3_x86_32(
+    Table input, 
+    uint32_t seed
+)
+
+cpdef Table murmurhash3_x64_128(
+    Table input,
+    uint64_t seed
+)
+
+cpdef Column md5(Table input)
+cpdef Column sha1(Table input)
+cpdef Column sha224(Table input)
+cpdef Column sha256(Table input)
+cpdef Column sha384(Table input)
+cpdef Column sha512(Table input)
+
+cpdef Column xxhash_64(
+    Table input,
+    uint64_t seed
+)

python/cudf/cudf/_lib/pylibcudf/hashing.pyx

@@ -0,0 +1,105 @@
+# Copyright (c) 2024, NVIDIA CORPORATION.
+from libc.stdint cimport uint32_t, uint64_t
+from libcpp.memory cimport unique_ptr
+from libcpp.utility cimport move
+
+from cudf._lib.cpp.column.column cimport column
+from cudf._lib.cpp.hash cimport (
+    md5 as cpp_md5,
+    murmurhash3_x64_128 as cpp_murmurhash3_x64_128,
+    murmurhash3_x86_32 as cpp_murmurhash3_x86_32,
+    sha1 as cpp_sha1,
+    sha224 as cpp_sha224,
+    sha256 as cpp_sha256,
+    sha384 as cpp_sha384,
+    sha512 as cpp_sha512,
+    xxhash_64 as cpp_xxhash_64,
+)
+from cudf._lib.cpp.table.table cimport table
+
+from .column cimport Column
+from .table cimport Table
+
+
+cpdef Column murmurhash3_x86_32(
+    Table input,
+    uint32_t seed
+):
+    cdef unique_ptr[column] c_result
+    with nogil:
+        c_result = move(
+            cpp_murmurhash3_x86_32(
+                input.view(),
+                seed
+            )
+        )
+
+    return Column.from_libcudf(move(c_result))
+
+cpdef Table murmurhash3_x64_128(
+    Table input,
+    uint64_t seed
+):
+    cdef unique_ptr[table] c_result
+    with nogil:
+        c_result = move(
+            cpp_murmurhash3_x64_128(
+                input.view(),
+                seed
+            )
+        )
+
+    return Table.from_libcudf(move(c_result))
+
+
+cpdef Column xxhash_64(
+    Table input,
+    uint64_t seed
+):
+    cdef unique_ptr[column] c_result
+    with  nogil:
+        c_result = move(
+            cpp_xxhash_64(
+                input.view(),
+                seed
+            )
+        )
+
+    return Column.from_libcudf(move(c_result))
+
+
+cpdef Column md5(Table input):
+    cdef unique_ptr[column] c_result
+    with nogil:
+        c_result = move(cpp_md5(input.view()))
+    return Column.from_libcudf(move(c_result))
+
+cpdef Column sha1(Table input):
+    cdef unique_ptr[column] c_result
+    with nogil:
+        c_result = move(cpp_sha1(input.view()))
+    return Column.from_libcudf(move(c_result))
+
+cpdef Column sha224(Table input):
+    cdef unique_ptr[column] c_result
+    with nogil:
+        c_result = move(cpp_sha224(input.view()))
+    return Column.from_libcudf(move(c_result))
+
+cpdef Column sha256(Table input):
+    cdef unique_ptr[column] c_result
+    with nogil:
+        c_result = move(cpp_sha256(input.view()))
+    return Column.from_libcudf(move(c_result))
+
+cpdef Column sha384(Table input):
+    cdef unique_ptr[column] c_result
+    with nogil:
+        c_result = move(cpp_sha384(input.view()))
+    return Column.from_libcudf(move(c_result))
+
+cpdef Column sha512(Table input):
+    cdef unique_ptr[column] c_result
+    with nogil:
+        c_result = move(cpp_sha512(input.view()))
+    return Column.from_libcudf(move(c_result))

python/cudf/cudf/pylibcudf_tests/common/utils.py

@@ -35,7 +35,6 @@ def assert_column_eq(plc_column: plc.Column, pa_array: pa.Array) -> None:
         plc_pa = plc_pa.combine_chunks()
     if isinstance(pa_array, pa.ChunkedArray):
         pa_array = pa_array.combine_chunks()
-
     assert plc_pa.equals(pa_array)
 
 

python/cudf/cudf/pylibcudf_tests/conftest.py

@@ -11,6 +11,8 @@
 
 from utils import DEFAULT_STRUCT_TESTING_TYPE
 
+import cudf._lib.pylibcudf as plc
+
 
 # This fixture defines the standard set of types that all tests should default to
 # running on. If there is a need for some tests to run on a different set of types, that
@@ -29,3 +31,25 @@
 )
 def pa_type(request):
     return request.param
+
+
+# TODO: Test nullable data
+@pytest.fixture(scope="session")
+def pa_input_column(pa_type):
+    if pa.types.is_integer(pa_type) or pa.types.is_floating(pa_type):
+        return pa.array([1, 2, 3], type=pa_type)
+    elif pa.types.is_string(pa_type):
+        return pa.array(["a", "b", "c"], type=pa_type)
+    elif pa.types.is_boolean(pa_type):
+        return pa.array([True, True, False], type=pa_type)
+    elif pa.types.is_list(pa_type):
+        # TODO: Add heterogenous sizes
+        return pa.array([[1], [2], [3]], type=pa_type)
+    elif pa.types.is_struct(pa_type):
+        return pa.array([{"v": 1}, {"v": 2}, {"v": 3}], type=pa_type)
+    raise ValueError("Unsupported type")
+
+
+@pytest.fixture(scope="session")
+def input_column(pa_input_column):
+    return plc.interop.from_arrow(pa_input_column)