Merge remote-tracking branch 'origin/branch-22.04' into refactor/remo…

…ve_frame_index

codecov.yml

@@ -3,6 +3,9 @@ coverage:
   status:
     project: off
     patch: on
+      default:
+        target: auto
+        threshold: 0%
 
 github_checks:
     annotations: true

cpp/CMakeLists.txt

@@ -300,12 +300,12 @@ add_library(
   src/io/orc/stripe_init.cu
   src/io/orc/timezone.cpp
   src/io/orc/writer_impl.cu
+  src/io/parquet/compact_protocol_reader.cpp
   src/io/parquet/compact_protocol_writer.cpp
   src/io/parquet/page_data.cu
   src/io/parquet/chunk_dict.cu
   src/io/parquet/page_enc.cu
   src/io/parquet/page_hdr.cu
-  src/io/parquet/parquet.cpp
   src/io/parquet/reader_impl.cu
   src/io/parquet/writer_impl.cu
   src/io/statistics/orc_column_statistics.cu
@@ -480,8 +480,10 @@ add_library(
   src/text/normalize.cu
   src/text/replace.cu
   src/text/stemmer.cu
+  src/text/subword/bpe_tokenizer.cu
   src/text/subword/data_normalizer.cu
   src/text/subword/load_hash_file.cu
+  src/text/subword/load_merges_file.cu
   src/text/subword/subword_tokenize.cu
   src/text/subword/wordpiece_tokenizer.cu
   src/text/tokenize.cu

cpp/benchmarks/hashing/hash.cpp

@@ -54,11 +54,11 @@ static void BM_hash(benchmark::State& state, cudf::hash_id hid, contains_nulls h
 #define HASH_BENCHMARK_DEFINE(hid, n) H_BENCHMARK_DEFINE(concat(hid, _, n), hid, n)
 
 HASH_BENCHMARK_DEFINE(HASH_MURMUR3, nulls)
-HASH_BENCHMARK_DEFINE(HASH_MD5, nulls)
 HASH_BENCHMARK_DEFINE(HASH_SERIAL_MURMUR3, nulls)
 HASH_BENCHMARK_DEFINE(HASH_SPARK_MURMUR3, nulls)
+HASH_BENCHMARK_DEFINE(HASH_MD5, nulls)
 
 HASH_BENCHMARK_DEFINE(HASH_MURMUR3, no_nulls)
-HASH_BENCHMARK_DEFINE(HASH_MD5, no_nulls)
 HASH_BENCHMARK_DEFINE(HASH_SERIAL_MURMUR3, no_nulls)
 HASH_BENCHMARK_DEFINE(HASH_SPARK_MURMUR3, no_nulls)
+HASH_BENCHMARK_DEFINE(HASH_MD5, no_nulls)

cpp/benchmarks/type_dispatcher/type_dispatcher.cu

@@ -120,7 +120,7 @@ struct RowHandle {
   template <typename T, CUDF_ENABLE_IF(not cudf::is_rep_layout_compatible<T>())>
   __device__ void operator()(cudf::mutable_column_device_view source, cudf::size_type index)
   {
-    cudf_assert(false && "Unsupported type.");
+    CUDF_UNREACHABLE("Unsupported type.");
   }
 };
 

cpp/include/cudf/ast/detail/expression_evaluator.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -103,7 +103,7 @@ struct value_expression_result
     if constexpr (std::is_same_v<Element, T>) {
       _obj = result;
     } else {
-      cudf_assert(false && "Output type does not match container type.");
+      CUDF_UNREACHABLE("Output type does not match container type.");
     }
   }
 
@@ -178,9 +178,7 @@ struct mutable_column_expression_result
   {
     // Not implemented since it would require modifying the API in the parent class to accept an
     // index.
-    cudf_assert(false && "This method is not implemented.");
-    // Unreachable return used to silence compiler warnings.
-    return {};
+    CUDF_UNREACHABLE("This method is not implemented.");
   }
 
   /**
@@ -190,7 +188,7 @@ struct mutable_column_expression_result
   {
     // Not implemented since it would require modifying the API in the parent class to accept an
     // index.
-    cudf_assert(false && "This method is not implemented.");
+    CUDF_UNREACHABLE("This method is not implemented.");
   }
 
   mutable_column_device_view& _obj;  ///< The column to which the data is written.
@@ -334,9 +332,7 @@ struct expression_evaluator {
     cudf::size_type left_row_index,
     thrust::optional<cudf::size_type> right_row_index = {}) const
   {
-    cudf_assert(false && "Unsupported type in resolve_input.");
-    // Unreachable return used to silence compiler warnings.
-    return {};
+    CUDF_UNREACHABLE("Unsupported type in resolve_input.");
   }
 
   /**
@@ -498,7 +494,7 @@ struct expression_evaluator {
                         op,
                         thread_intermediate_storage);
       } else {
-        cudf_assert(false && "Invalid operator arity.");
+        CUDF_UNREACHABLE("Invalid operator arity.");
       }
     }
   }
@@ -567,7 +563,7 @@ struct expression_evaluator {
       IntermediateDataType<has_nulls>* thread_intermediate_storage,
       possibly_null_value_t<Element, has_nulls> const& result) const
     {
-      cudf_assert(false && "Invalid type in resolve_output.");
+      CUDF_UNREACHABLE("Invalid type in resolve_output.");
     }
   };
 
@@ -630,7 +626,7 @@ struct expression_evaluator {
       detail::device_data_reference const& output,
       IntermediateDataType<has_nulls>* thread_intermediate_storage) const
     {
-      cudf_assert(false && "Invalid unary dispatch operator for the provided input.");
+      CUDF_UNREACHABLE("Invalid unary dispatch operator for the provided input.");
     }
   };
 
@@ -698,7 +694,7 @@ struct expression_evaluator {
       detail::device_data_reference const& output,
       IntermediateDataType<has_nulls>* thread_intermediate_storage) const
     {
-      cudf_assert(false && "Invalid binary dispatch operator for the provided input.");
+      CUDF_UNREACHABLE("Invalid binary dispatch operator for the provided input.");
     }
   };
 

cpp/include/cudf/ast/detail/operators.hpp

@@ -199,13 +199,13 @@ CUDF_HOST_DEVICE inline constexpr void ast_operator_dispatcher(ast_operator op,
     case ast_operator::CAST_TO_FLOAT64:
       f.template operator()<ast_operator::CAST_TO_FLOAT64>(std::forward<Ts>(args)...);
       break;
-    default:
+    default: {
 #ifndef __CUDA_ARCH__
       CUDF_FAIL("Invalid operator.");
 #else
-      cudf_assert(false && "Invalid operator.");
+      CUDF_UNREACHABLE("Invalid operator.");
 #endif
-      break;
+    }
   }
 }
 
@@ -934,7 +934,7 @@ struct single_dispatch_binary_operator_types {
 #ifndef __CUDA_ARCH__
     CUDF_FAIL("Invalid binary operation.");
 #else
-    cudf_assert(false && "Invalid binary operation.");
+    CUDF_UNREACHABLE("Invalid binary operation.");
 #endif
   }
 };
@@ -1023,7 +1023,7 @@ struct dispatch_unary_operator_types {
 #ifndef __CUDA_ARCH__
     CUDF_FAIL("Invalid unary operation.");
 #else
-    cudf_assert(false && "Invalid unary operation.");
+    CUDF_UNREACHABLE("Invalid unary operation.");
 #endif
   }
 };
@@ -1097,7 +1097,7 @@ struct return_type_functor {
 #ifndef __CUDA_ARCH__
     CUDF_FAIL("Invalid binary operation. Return type cannot be determined.");
 #else
-    cudf_assert(false && "Invalid binary operation. Return type cannot be determined.");
+    CUDF_UNREACHABLE("Invalid binary operation. Return type cannot be determined.");
 #endif
   }
 
@@ -1125,7 +1125,7 @@ struct return_type_functor {
 #ifndef __CUDA_ARCH__
     CUDF_FAIL("Invalid unary operation. Return type cannot be determined.");
 #else
-    cudf_assert(false && "Invalid unary operation. Return type cannot be determined.");
+    CUDF_UNREACHABLE("Invalid unary operation. Return type cannot be determined.");
 #endif
   }
 };

cpp/include/cudf/column/column_device_view.cuh

@@ -391,8 +391,7 @@ class alignas(16) column_device_view : public detail::column_device_view_base {
               CUDF_ENABLE_IF(not(is_index_type<IndexType>() and std::is_unsigned_v<IndexType>))>
     __device__ size_type operator()(Args&&... args)
     {
-      cudf_assert(false and "dictionary indices must be an unsigned integral type");
-      return 0;
+      CUDF_UNREACHABLE("dictionary indices must be an unsigned integral type");
     }
   };
 

cpp/include/cudf/copying.hpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2018-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -371,7 +371,7 @@ std::unique_ptr<column> copy_range(
  * @param fill_value Fill value for indeterminable outputs.
  * @param mr         Device memory resource used to allocate the returned result's device memory
  *
- * @throw cudf::logic_error if @p input dtype is not fixed-with.
+ * @throw cudf::logic_error if @p input dtype is neither fixed-width nor string type
  * @throw cudf::logic_error if @p fill_value dtype does not match @p input dtype.
  */
 std::unique_ptr<column> shift(

cpp/include/cudf/detail/aggregation/aggregation.cuh

@@ -122,7 +122,7 @@ struct update_target_element {
                              column_device_view source,
                              size_type source_index) const noexcept
   {
-    cudf_assert(false and "Invalid source type and aggregation combination.");
+    CUDF_UNREACHABLE("Invalid source type and aggregation combination.");
   }
 };
 

cpp/include/cudf/detail/aggregation/aggregation.hpp

@@ -1483,17 +1483,7 @@ CUDF_HOST_DEVICE inline decltype(auto) aggregation_dispatcher(aggregation::Kind
 #ifndef __CUDA_ARCH__
       CUDF_FAIL("Unsupported aggregation.");
 #else
-      cudf_assert(false && "Unsupported aggregation.");
-
-      // The following code will never be reached, but the compiler generates a
-      // warning if there isn't a return value.
-
-      // Need to find out what the return type is in order to have a default
-      // return value and solve the compiler warning for lack of a default
-      // return
-      using return_type =
-        decltype(f.template operator()<aggregation::SUM>(std::forward<Ts>(args)...));
-      return return_type();
+      CUDF_UNREACHABLE("Unsupported aggregation.");
 #endif
     }
   }

cpp/include/cudf/detail/hashing.hpp

@@ -42,15 +42,15 @@ std::unique_ptr<column> 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> md5_hash(
+template <template <typename> class hash_function>
+std::unique_ptr<column> serial_murmur_hash3_32(
   table_view const& input,
+  uint32_t seed                       = 0,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
-template <template <typename> class hash_function>
-std::unique_ptr<column> serial_murmur_hash3_32(
+std::unique_ptr<column> md5_hash(
   table_view const& input,
-  uint32_t seed                       = 0,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 

cpp/include/cudf/detail/indexalator.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -267,8 +267,7 @@ struct input_indexalator : base_indexalator<input_indexalator> {
     template <typename T, std::enable_if_t<not is_index_type<T>()>* = nullptr>
     __device__ size_type operator()(void const* tp)
     {
-      cudf_assert(false and "only index types are supported");
-      return 0;
+      CUDF_UNREACHABLE("only index types are supported");
     }
   };
   /**
@@ -365,7 +364,7 @@ struct output_indexalator : base_indexalator<output_indexalator> {
     template <typename T, std::enable_if_t<not is_index_type<T>()>* = nullptr>
     __device__ void operator()(void* tp, size_type const value)
     {
-      cudf_assert(false and "only index types are supported");
+      CUDF_UNREACHABLE("only index types are supported");
     }
   };
 

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

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -35,3 +35,29 @@
 #else
 #define cudf_assert(e) (static_cast<void>(0))
 #endif
+
+/**
+ * @brief Macro indicating that a location in the code is unreachable.
+ *
+ * The CUDF_UNREACHABLE macro should only be used where CUDF_FAIL cannot be used
+ * due to performance or due to being used in device code. In the majority of
+ * host code situations, an exception should be thrown in "unreachable" code
+ * paths as those usually aren't tight inner loops like they are in device code.
+ *
+ * One example where this macro may be used is in conjunction with dispatchers
+ * to indicate that a function does not need to return a default value because
+ * it has already exhausted all possible cases in a `switch` statement.
+ *
+ * The assert in this macro can be used when compiling in debug mode to help
+ * debug functions that may reach the supposedly unreachable code.
+ *
+ * Example usage:
+ * ```
+ * CUDF_UNREACHABLE("Invalid type_id.");
+ * ```
+ */
+#define CUDF_UNREACHABLE(msg)             \
+  do {                                    \
+    assert(false && "Unreachable: " msg); \
+    __builtin_unreachable();              \
+  } while (0)