Add CUDF_UNREACHABLE macro.

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/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/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)

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

@@ -139,7 +139,7 @@ auto __device__ inline 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.");
+    CUDF_UNREACHABLE("Unsupported type.");
   }
 }
 
@@ -334,16 +334,14 @@ template <>
 hash_value_type __device__ inline MurmurHash3_32<cudf::list_view>::operator()(
   cudf::list_view const& key) const
 {
-  cudf_assert(false && "List column hashing is not supported");
-  return 0;
+  CUDF_UNREACHABLE("List column hashing is not supported");
 }
 
 template <>
 hash_value_type __device__ inline MurmurHash3_32<cudf::struct_view>::operator()(
   cudf::struct_view const& key) const
 {
-  cudf_assert(false && "Direct hashing of struct_view is not supported");
-  return 0;
+  CUDF_UNREACHABLE("Direct hashing of struct_view is not supported");
 }
 
 template <typename Key>
@@ -551,16 +549,14 @@ template <>
 hash_value_type __device__ inline SparkMurmurHash3_32<cudf::list_view>::operator()(
   cudf::list_view const& key) const
 {
-  cudf_assert(false && "List column hashing is not supported");
-  return 0;
+  CUDF_UNREACHABLE("List column hashing is not supported");
 }
 
 template <>
 hash_value_type __device__ inline SparkMurmurHash3_32<cudf::struct_view>::operator()(
   cudf::struct_view const& key) const
 {
-  cudf_assert(false && "Direct hashing of struct_view is not supported");
-  return 0;
+  CUDF_UNREACHABLE("Direct hashing of struct_view is not supported");
 }
 
 /**
@@ -577,8 +573,7 @@ struct IdentityHash {
   constexpr std::enable_if_t<!std::is_arithmetic_v<Key>, return_type> operator()(
     Key const& key) const
   {
-    cudf_assert(false && "IdentityHash does not support this data type");
-    return 0;
+    CUDF_UNREACHABLE("IdentityHash does not support this data type");
   }
 
   template <typename return_type = result_type>

cpp/include/cudf/table/row_operators.cuh

@@ -217,8 +217,7 @@ class element_equality_comparator {
             std::enable_if_t<not cudf::is_equality_comparable<Element, Element>()>* = nullptr>
   __device__ bool operator()(size_type lhs_element_index, size_type rhs_element_index)
   {
-    cudf_assert(false && "Attempted to compare elements of uncomparable types.");
-    return false;
+    CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types.");
   }
 
  private:
@@ -323,8 +322,7 @@ class element_relational_comparator {
             std::enable_if_t<not cudf::is_relationally_comparable<Element, Element>()>* = nullptr>
   __device__ weak_ordering operator()(size_type lhs_element_index, size_type rhs_element_index)
   {
-    cudf_assert(false && "Attempted to compare elements of uncomparable types.");
-    return weak_ordering::LESS;
+    CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types.");
   }
 
  private:
@@ -442,8 +440,7 @@ class element_hasher {
   template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
   __device__ hash_value_type operator()(column_device_view col, size_type row_index) const
   {
-    cudf_assert(false && "Unsupported type in hash.");
-    return {};
+    CUDF_UNREACHABLE("Unsupported type in hash.");
   }
 
   Nullate has_nulls;
@@ -472,8 +469,7 @@ class element_hasher_with_seed {
   template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
   __device__ hash_value_type operator()(column_device_view col, size_type row_index) const
   {
-    cudf_assert(false && "Unsupported type in hash.");
-    return {};
+    CUDF_UNREACHABLE("Unsupported type in hash.");
   }
 
  private:

cpp/include/cudf/utilities/type_dispatcher.hpp

@@ -511,18 +511,9 @@ CUDF_HOST_DEVICE __forceinline__ constexpr decltype(auto) type_dispatcher(cudf::
         std::forward<Ts>(args)...);
     default: {
 #ifndef __CUDA_ARCH__
-      CUDF_FAIL("Unsupported type_id.");
+      CUDF_FAIL("Invalid type_id.");
 #else
-      cudf_assert(false && "Unsupported type_id.");
-
-      // 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()<int8_t>(std::forward<Ts>(args)...));
-      return return_type();
+      CUDF_UNREACHABLE("Invalid type_id.");
 #endif
     }
   }

cpp/src/datetime/datetime_ops.cu

@@ -105,9 +105,8 @@ struct RoundFunctor {
       case rounding_function::CEIL: return cuda::std::chrono::ceil<DurationType>(dt);
       case rounding_function::FLOOR: return cuda::std::chrono::floor<DurationType>(dt);
       case rounding_function::ROUND: return cuda::std::chrono::round<DurationType>(dt);
-      default: cudf_assert(false && "Unsupported rounding kind.");
+      default: CUDF_UNREACHABLE("Unsupported rounding kind.");
     }
-    __builtin_unreachable();
   }
 };
 
@@ -145,9 +144,8 @@ struct RoundingDispatcher {
       case rounding_frequency::NANOSECOND:
         return time_point_cast<typename Timestamp::duration>(
           RoundFunctor<duration_ns>{}(round_kind, ts));
-      default: cudf_assert(false && "Unsupported datetime rounding resolution.");
+      default: CUDF_UNREACHABLE("Unsupported datetime rounding resolution.");
     }
-    __builtin_unreachable();
   }
 };