Revert "Add CUDF_UNREACHABLE macro. (rapidsai#9727)"

This reverts commit 48cebf7.

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_UNREACHABLE("Unsupported type.");
+    cudf_assert(false && "Unsupported type.");
   }
 };
 

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

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, 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_UNREACHABLE("Output type does not match container type.");
+      cudf_assert(false && "Output type does not match container type.");
     }
   }
 
@@ -178,7 +178,9 @@ struct mutable_column_expression_result
   {
     // Not implemented since it would require modifying the API in the parent class to accept an
     // index.
-    CUDF_UNREACHABLE("This method is not implemented.");
+    cudf_assert(false && "This method is not implemented.");
+    // Unreachable return used to silence compiler warnings.
+    return {};
   }
 
   /**
@@ -188,7 +190,7 @@ struct mutable_column_expression_result
   {
     // Not implemented since it would require modifying the API in the parent class to accept an
     // index.
-    CUDF_UNREACHABLE("This method is not implemented.");
+    cudf_assert(false && "This method is not implemented.");
   }
 
   mutable_column_device_view& _obj;  ///< The column to which the data is written.
@@ -332,7 +334,9 @@ struct expression_evaluator {
     cudf::size_type left_row_index,
     thrust::optional<cudf::size_type> right_row_index = {}) const
   {
-    CUDF_UNREACHABLE("Unsupported type in resolve_input.");
+    cudf_assert(false && "Unsupported type in resolve_input.");
+    // Unreachable return used to silence compiler warnings.
+    return {};
   }
 
   /**
@@ -494,7 +498,7 @@ struct expression_evaluator {
                         op,
                         thread_intermediate_storage);
       } else {
-        CUDF_UNREACHABLE("Invalid operator arity.");
+        cudf_assert(false && "Invalid operator arity.");
       }
     }
   }
@@ -563,7 +567,7 @@ struct expression_evaluator {
       IntermediateDataType<has_nulls>* thread_intermediate_storage,
       possibly_null_value_t<Element, has_nulls> const& result) const
     {
-      CUDF_UNREACHABLE("Invalid type in resolve_output.");
+      cudf_assert(false && "Invalid type in resolve_output.");
     }
   };
 
@@ -626,7 +630,7 @@ struct expression_evaluator {
       detail::device_data_reference const& output,
       IntermediateDataType<has_nulls>* thread_intermediate_storage) const
     {
-      CUDF_UNREACHABLE("Invalid unary dispatch operator for the provided input.");
+      cudf_assert(false && "Invalid unary dispatch operator for the provided input.");
     }
   };
 
@@ -694,7 +698,7 @@ struct expression_evaluator {
       detail::device_data_reference const& output,
       IntermediateDataType<has_nulls>* thread_intermediate_storage) const
     {
-      CUDF_UNREACHABLE("Invalid binary dispatch operator for the provided input.");
+      cudf_assert(false && "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_UNREACHABLE("Invalid operator.");
+      cudf_assert(false && "Invalid operator.");
 #endif
-    }
+      break;
   }
 }
 
@@ -934,7 +934,7 @@ struct single_dispatch_binary_operator_types {
 #ifndef __CUDA_ARCH__
     CUDF_FAIL("Invalid binary operation.");
 #else
-    CUDF_UNREACHABLE("Invalid binary operation.");
+    cudf_assert(false && "Invalid binary operation.");
 #endif
   }
 };
@@ -1023,7 +1023,7 @@ struct dispatch_unary_operator_types {
 #ifndef __CUDA_ARCH__
     CUDF_FAIL("Invalid unary operation.");
 #else
-    CUDF_UNREACHABLE("Invalid unary operation.");
+    cudf_assert(false && "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_UNREACHABLE("Invalid binary operation. Return type cannot be determined.");
+    cudf_assert(false && "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_UNREACHABLE("Invalid unary operation. Return type cannot be determined.");
+    cudf_assert(false && "Invalid unary operation. Return type cannot be determined.");
 #endif
   }
 };

cpp/include/cudf/column/column_device_view.cuh

@@ -418,7 +418,8 @@ 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_UNREACHABLE("dictionary indices must be an unsigned integral type");
+      cudf_assert(false and "dictionary indices must be an unsigned integral type");
+      return 0;
     }
   };
 

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

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

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

@@ -1483,7 +1483,17 @@ CUDF_HOST_DEVICE inline decltype(auto) aggregation_dispatcher(aggregation::Kind
 #ifndef __CUDA_ARCH__
       CUDF_FAIL("Unsupported aggregation.");
 #else
-      CUDF_UNREACHABLE("Unsupported aggregation.");
+      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();
 #endif
     }
   }

cpp/include/cudf/detail/indexalator.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -272,7 +272,8 @@ 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_UNREACHABLE("only index types are supported");
+      cudf_assert(false and "only index types are supported");
+      return 0;
     }
   };
   /**
@@ -369,7 +370,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_UNREACHABLE("only index types are supported");
+      cudf_assert(false and "only index types are supported");
     }
   };
 

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

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2020, 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,29 +35,3 @@
 #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

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

cpp/include/cudf/table/row_operators.cuh

@@ -219,7 +219,8 @@ 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_UNREACHABLE("Attempted to compare elements of uncomparable types.");
+    cudf_assert(false && "Attempted to compare elements of uncomparable types.");
+    return false;
   }
 
  private:
@@ -324,7 +325,8 @@ 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_UNREACHABLE("Attempted to compare elements of uncomparable types.");
+    cudf_assert(false && "Attempted to compare elements of uncomparable types.");
+    return weak_ordering::LESS;
   }
 
  private:
@@ -442,7 +444,8 @@ 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_UNREACHABLE("Unsupported type in hash.");
+    cudf_assert(false && "Unsupported type in hash.");
+    return {};
   }
 
   Nullate has_nulls;
@@ -471,7 +474,8 @@ 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_UNREACHABLE("Unsupported type in hash.");
+    cudf_assert(false && "Unsupported type in hash.");
+    return {};
   }
 
  private:

cpp/include/cudf/utilities/type_dispatcher.hpp

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

cpp/src/datetime/datetime_ops.cu

@@ -107,8 +107,9 @@ 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_UNREACHABLE("Unsupported rounding kind.");
+      default: cudf_assert(false && "Unsupported rounding kind.");
     }
+    __builtin_unreachable();
   }
 };
 
@@ -146,8 +147,9 @@ struct RoundingDispatcher {
       case rounding_frequency::NANOSECOND:
         return time_point_cast<typename Timestamp::duration>(
           RoundFunctor<duration_ns>{}(round_kind, ts));
-      default: CUDF_UNREACHABLE("Unsupported datetime rounding resolution.");
+      default: cudf_assert(false && "Unsupported datetime rounding resolution.");
     }
+    __builtin_unreachable();
   }
 };