Merge branch 'branch-24.12' into migrate-orc-dict-encoding

rapidsai · Oct 11, 2024 · 7993cf7 · 7993cf7
2 parents 5dacd83 + b8f3e21
commit 7993cf7
Show file tree

Hide file tree

Showing 50 changed files with 2,907 additions and 2,368 deletions.
diff --git a/ci/build_docs.sh b/ci/build_docs.sh
@@ -3,8 +3,8 @@
 
 set -euo pipefail
 
-RAPIDS_VERSION_MAJOR_MINOR="$(rapids-version-major-minor)"
-export RAPIDS_VERSION_NUMBER="$RAPIDS_VERSION_MAJOR_MINOR"
+export RAPIDS_VERSION="$(rapids-version)"
+export RAPIDS_VERSION_MAJOR_MINOR="$(rapids-version-major-minor)"
 
 rapids-logger "Create test conda environment"
 . /opt/conda/etc/profile.d/conda.sh
@@ -28,16 +28,16 @@ PYTHON_CHANNEL=$(rapids-download-conda-from-s3 python)
 rapids-mamba-retry install \
   --channel "${CPP_CHANNEL}" \
   --channel "${PYTHON_CHANNEL}" \
-  "libcudf=${RAPIDS_VERSION_MAJOR_MINOR}" \
-  "pylibcudf=${RAPIDS_VERSION_MAJOR_MINOR}" \
-  "cudf=${RAPIDS_VERSION_MAJOR_MINOR}" \
-  "dask-cudf=${RAPIDS_VERSION_MAJOR_MINOR}"
+  "libcudf=${RAPIDS_VERSION}" \
+  "pylibcudf=${RAPIDS_VERSION}" \
+  "cudf=${RAPIDS_VERSION}" \
+  "dask-cudf=${RAPIDS_VERSION}"
 
 export RAPIDS_DOCS_DIR="$(mktemp -d)"
 
 rapids-logger "Build CPP docs"
 pushd cpp/doxygen
-aws s3 cp s3://rapidsai-docs/librmm/html/${RAPIDS_VERSION_NUMBER}/rmm.tag . || echo "Failed to download rmm Doxygen tag"
+aws s3 cp s3://rapidsai-docs/librmm/html/${RAPIDS_VERSION_MAJOR_MINOR}/rmm.tag . || echo "Failed to download rmm Doxygen tag"
 doxygen Doxyfile
 mkdir -p "${RAPIDS_DOCS_DIR}/libcudf/html"
 mv html/* "${RAPIDS_DOCS_DIR}/libcudf/html"
@@ -57,4 +57,4 @@ mkdir -p "${RAPIDS_DOCS_DIR}/dask-cudf/html"
 mv build/dirhtml/* "${RAPIDS_DOCS_DIR}/dask-cudf/html"
 popd
 
-rapids-upload-docs
+RAPIDS_VERSION_NUMBER="${RAPIDS_VERSION_MAJOR_MINOR}" rapids-upload-docs
diff --git a/ci/test_cpp_common.sh b/ci/test_cpp_common.sh
@@ -5,7 +5,7 @@ set -euo pipefail
 
 . /opt/conda/etc/profile.d/conda.sh
 
-RAPIDS_VERSION_MAJOR_MINOR="$(rapids-version-major-minor)"
+RAPIDS_VERSION="$(rapids-version)"
 
 rapids-logger "Generate C++ testing dependencies"
 
@@ -33,10 +33,10 @@ rapids-print-env
 
 rapids-mamba-retry install \
   --channel "${CPP_CHANNEL}" \
-  "libcudf=${RAPIDS_VERSION_MAJOR_MINOR}" \
-  "libcudf_kafka=${RAPIDS_VERSION_MAJOR_MINOR}" \
-  "libcudf-tests=${RAPIDS_VERSION_MAJOR_MINOR}" \
-  "libcudf-example=${RAPIDS_VERSION_MAJOR_MINOR}"
+  "libcudf=${RAPIDS_VERSION}" \
+  "libcudf_kafka=${RAPIDS_VERSION}" \
+  "libcudf-tests=${RAPIDS_VERSION}" \
+  "libcudf-example=${RAPIDS_VERSION}"
 
 rapids-logger "Check GPU usage"
 nvidia-smi
diff --git a/ci/test_java.sh b/ci/test_java.sh
@@ -5,7 +5,7 @@ set -euo pipefail
 
 . /opt/conda/etc/profile.d/conda.sh
 
-RAPIDS_VERSION_MAJOR_MINOR="$(rapids-version-major-minor)"
+RAPIDS_VERSION="$(rapids-version)"
 
 rapids-logger "Generate Java testing dependencies"
 
@@ -32,7 +32,7 @@ CPP_CHANNEL=$(rapids-download-conda-from-s3 cpp)
 
 rapids-mamba-retry install \
   --channel "${CPP_CHANNEL}" \
-  "libcudf=${RAPIDS_VERSION_MAJOR_MINOR}"
+  "libcudf=${RAPIDS_VERSION}"
 
 rapids-logger "Check GPU usage"
 nvidia-smi

diff --git a/ci/test_notebooks.sh b/ci/test_notebooks.sh
@@ -5,7 +5,7 @@ set -euo pipefail
 
 . /opt/conda/etc/profile.d/conda.sh
 
-RAPIDS_VERSION_MAJOR_MINOR="$(rapids-version-major-minor)"
+RAPIDS_VERSION="$(rapids-version)"
 
 rapids-logger "Generate notebook testing dependencies"
 
@@ -32,8 +32,8 @@ PYTHON_CHANNEL=$(rapids-download-conda-from-s3 python)
 rapids-mamba-retry install \
   --channel "${CPP_CHANNEL}" \
   --channel "${PYTHON_CHANNEL}" \
-  "cudf=${RAPIDS_VERSION_MAJOR_MINOR}" \
-  "libcudf=${RAPIDS_VERSION_MAJOR_MINOR}"
+  "cudf=${RAPIDS_VERSION}" \
+  "libcudf=${RAPIDS_VERSION}"
 
 NBTEST="$(realpath "$(dirname "$0")/utils/nbtest.sh")"
 pushd notebooks

diff --git a/ci/test_python_common.sh b/ci/test_python_common.sh
@@ -7,7 +7,7 @@ set -euo pipefail
 
 . /opt/conda/etc/profile.d/conda.sh
 
-RAPIDS_VERSION_MAJOR_MINOR="$(rapids-version-major-minor)"
+RAPIDS_VERSION="$(rapids-version)"
 
 rapids-logger "Generate Python testing dependencies"
 
@@ -40,5 +40,5 @@ rapids-print-env
 rapids-mamba-retry install \
   --channel "${CPP_CHANNEL}" \
   --channel "${PYTHON_CHANNEL}" \
-  "cudf=${RAPIDS_VERSION_MAJOR_MINOR}" \
-  "libcudf=${RAPIDS_VERSION_MAJOR_MINOR}"
+  "cudf=${RAPIDS_VERSION}" \
+  "libcudf=${RAPIDS_VERSION}"
diff --git a/ci/test_python_other.sh b/ci/test_python_other.sh
@@ -7,14 +7,14 @@ cd "$(dirname "$(realpath "${BASH_SOURCE[0]}")")"/../
 # Common setup steps shared by Python test jobs
 source ./ci/test_python_common.sh test_python_other
 
-RAPIDS_VERSION_MAJOR_MINOR="$(rapids-version-major-minor)"
+RAPIDS_VERSION="$(rapids-version)"
 
 rapids-mamba-retry install \
   --channel "${CPP_CHANNEL}" \
   --channel "${PYTHON_CHANNEL}" \
-  "dask-cudf=${RAPIDS_VERSION_MAJOR_MINOR}" \
-  "cudf_kafka=${RAPIDS_VERSION_MAJOR_MINOR}" \
-  "custreamz=${RAPIDS_VERSION_MAJOR_MINOR}"
+  "dask-cudf=${RAPIDS_VERSION}" \
+  "cudf_kafka=${RAPIDS_VERSION}" \
+  "custreamz=${RAPIDS_VERSION}"
 
 rapids-logger "Check GPU usage"
 nvidia-smi

diff --git a/cpp/include/cudf/hashing/detail/helper_functions.cuh b/cpp/include/cudf/hashing/detail/helper_functions.cuh
@@ -47,197 +47,3 @@ inline size_t compute_hash_table_size(cudf::size_type num_keys_to_insert,
 
   return hash_table_size;
 }
-
-template <typename pair_type>
-__forceinline__ __device__ pair_type load_pair_vectorized(pair_type const* __restrict__ const ptr)
-{
-  if (sizeof(uint4) == sizeof(pair_type)) {
-    union pair_type2vec_type {
-      uint4 vec_val;
-      pair_type pair_val;
-    };
-    pair_type2vec_type converter = {0, 0, 0, 0};
-    converter.vec_val            = *reinterpret_cast<uint4 const*>(ptr);
-    return converter.pair_val;
-  } else if (sizeof(uint2) == sizeof(pair_type)) {
-    union pair_type2vec_type {
-      uint2 vec_val;
-      pair_type pair_val;
-    };
-    pair_type2vec_type converter = {0, 0};
-    converter.vec_val            = *reinterpret_cast<uint2 const*>(ptr);
-    return converter.pair_val;
-  } else if (sizeof(int) == sizeof(pair_type)) {
-    union pair_type2vec_type {
-      int vec_val;
-      pair_type pair_val;
-    };
-    pair_type2vec_type converter = {0};
-    converter.vec_val            = *reinterpret_cast<int const*>(ptr);
-    return converter.pair_val;
-  } else if (sizeof(short) == sizeof(pair_type)) {
-    union pair_type2vec_type {
-      short vec_val;
-      pair_type pair_val;
-    };
-    pair_type2vec_type converter = {0};
-    converter.vec_val            = *reinterpret_cast<short const*>(ptr);
-    return converter.pair_val;
-  } else {
-    return *ptr;
-  }
-}
-
-template <typename pair_type>
-__forceinline__ __device__ void store_pair_vectorized(pair_type* __restrict__ const ptr,
-                                                      pair_type const val)
-{
-  if (sizeof(uint4) == sizeof(pair_type)) {
-    union pair_type2vec_type {
-      uint4 vec_val;
-      pair_type pair_val;
-    };
-    pair_type2vec_type converter   = {0, 0, 0, 0};
-    converter.pair_val             = val;
-    *reinterpret_cast<uint4*>(ptr) = converter.vec_val;
-  } else if (sizeof(uint2) == sizeof(pair_type)) {
-    union pair_type2vec_type {
-      uint2 vec_val;
-      pair_type pair_val;
-    };
-    pair_type2vec_type converter   = {0, 0};
-    converter.pair_val             = val;
-    *reinterpret_cast<uint2*>(ptr) = converter.vec_val;
-  } else if (sizeof(int) == sizeof(pair_type)) {
-    union pair_type2vec_type {
-      int vec_val;
-      pair_type pair_val;
-    };
-    pair_type2vec_type converter = {0};
-    converter.pair_val           = val;
-    *reinterpret_cast<int*>(ptr) = converter.vec_val;
-  } else if (sizeof(short) == sizeof(pair_type)) {
-    union pair_type2vec_type {
-      short vec_val;
-      pair_type pair_val;
-    };
-    pair_type2vec_type converter   = {0};
-    converter.pair_val             = val;
-    *reinterpret_cast<short*>(ptr) = converter.vec_val;
-  } else {
-    *ptr = val;
-  }
-}
-
-template <typename value_type, typename size_type, typename key_type, typename elem_type>
-CUDF_KERNEL void init_hashtbl(value_type* __restrict__ const hashtbl_values,
-                              size_type const n,
-                              key_type const key_val,
-                              elem_type const elem_val)
-{
-  size_type const idx = blockIdx.x * blockDim.x + threadIdx.x;
-  if (idx < n) {
-    store_pair_vectorized(hashtbl_values + idx, thrust::make_pair(key_val, elem_val));
-  }
-}
-
-template <typename T>
-struct equal_to {
-  using result_type          = bool;
-  using first_argument_type  = T;
-  using second_argument_type = T;
-  __forceinline__ __host__ __device__ constexpr bool operator()(
-    first_argument_type const& lhs, second_argument_type const& rhs) const
-  {
-    return lhs == rhs;
-  }
-};
-
-template <typename Iterator>
-class cycle_iterator_adapter {
- public:
-  using value_type      = typename std::iterator_traits<Iterator>::value_type;
-  using difference_type = typename std::iterator_traits<Iterator>::difference_type;
-  using pointer         = typename std::iterator_traits<Iterator>::pointer;
-  using reference       = typename std::iterator_traits<Iterator>::reference;
-  using iterator_type   = Iterator;
-
-  cycle_iterator_adapter() = delete;
-
-  __host__ __device__ explicit cycle_iterator_adapter(iterator_type const& begin,
-                                                      iterator_type const& end,
-                                                      iterator_type const& current)
-    : m_begin(begin), m_end(end), m_current(current)
-  {
-  }
-
-  __host__ __device__ cycle_iterator_adapter& operator++()
-  {
-    if (m_end == (m_current + 1))
-      m_current = m_begin;
-    else
-      ++m_current;
-    return *this;
-  }
-
-  __host__ __device__ cycle_iterator_adapter const& operator++() const
-  {
-    if (m_end == (m_current + 1))
-      m_current = m_begin;
-    else
-      ++m_current;
-    return *this;
-  }
-
-  __host__ __device__ cycle_iterator_adapter& operator++(int)
-  {
-    cycle_iterator_adapter<iterator_type> old(m_begin, m_end, m_current);
-    if (m_end == (m_current + 1))
-      m_current = m_begin;
-    else
-      ++m_current;
-    return old;
-  }
-
-  __host__ __device__ cycle_iterator_adapter const& operator++(int) const
-  {
-    cycle_iterator_adapter<iterator_type> old(m_begin, m_end, m_current);
-    if (m_end == (m_current + 1))
-      m_current = m_begin;
-    else
-      ++m_current;
-    return old;
-  }
-
-  __host__ __device__ bool equal(cycle_iterator_adapter<iterator_type> const& other) const
-  {
-    return m_current == other.m_current && m_begin == other.m_begin && m_end == other.m_end;
-  }
-
-  __host__ __device__ reference& operator*() { return *m_current; }
-
-  __host__ __device__ reference const& operator*() const { return *m_current; }
-
-  __host__ __device__ const pointer operator->() const { return m_current.operator->(); }
-
-  __host__ __device__ pointer operator->() { return m_current; }
-
- private:
-  iterator_type m_current;
-  iterator_type m_begin;
-  iterator_type m_end;
-};
-
-template <class T>
-__host__ __device__ bool operator==(cycle_iterator_adapter<T> const& lhs,
-                                    cycle_iterator_adapter<T> const& rhs)
-{
-  return lhs.equal(rhs);
-}
-
-template <class T>
-__host__ __device__ bool operator!=(cycle_iterator_adapter<T> const& lhs,
-                                    cycle_iterator_adapter<T> const& rhs)
-{
-  return !lhs.equal(rhs);
-}
diff --git a/cpp/src/io/parquet/chunk_dict.cu b/cpp/src/io/parquet/chunk_dict.cu
@@ -84,7 +84,7 @@ struct map_insert_fn {
                                                storage_ref};
 
       // Create a map ref with `cuco::insert` operator
-      auto map_insert_ref = hash_map_ref.with_operators(cuco::insert);
+      auto map_insert_ref = hash_map_ref.rebind_operators(cuco::insert);
       auto const t        = threadIdx.x;
 
       // Create atomic refs to the current chunk's num_dict_entries and uniq_data_size
@@ -186,7 +186,7 @@ struct map_find_fn {
                                                storage_ref};
 
       // Create a map ref with `cuco::find` operator
-      auto const map_find_ref = hash_map_ref.with_operators(cuco::find);
+      auto const map_find_ref = hash_map_ref.rebind_operators(cuco::find);
       auto const t            = threadIdx.x;
 
       // Note: Adjust the following loop to use `cg::tiles<map_cg_size>` if needed in the future.