diff --git a/cpp/benchmarks/type_dispatcher/type_dispatcher.cu b/cpp/benchmarks/type_dispatcher/type_dispatcher.cu
index 6ab6f9a2095..3be599e8c41 100644
--- a/cpp/benchmarks/type_dispatcher/type_dispatcher.cu
+++ b/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.");
   }
 };
 
diff --git a/cpp/include/cudf/ast/detail/expression_evaluator.cuh b/cpp/include/cudf/ast/detail/expression_evaluator.cuh
index a179fe10774..2bfe1b03dd3 100644
--- a/cpp/include/cudf/ast/detail/expression_evaluator.cuh
+++ b/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.");
     }
   };
 
diff --git a/cpp/include/cudf/ast/detail/operators.hpp b/cpp/include/cudf/ast/detail/operators.hpp
index 3c6c9de055d..63048a993c7 100644
--- a/cpp/include/cudf/ast/detail/operators.hpp
+++ b/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
   }
 };
diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh
index 070ca80858b..5b31181fd6c 100644
--- a/cpp/include/cudf/column/column_device_view.cuh
+++ b/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;
     }
   };
 
diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh
index 818e8cd7cc6..4184d9bcbf2 100644
--- a/cpp/include/cudf/detail/aggregation/aggregation.cuh
+++ b/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.");
   }
 };
 
diff --git a/cpp/include/cudf/detail/aggregation/aggregation.hpp b/cpp/include/cudf/detail/aggregation/aggregation.hpp
index 886151fb9d6..0f59eeb6bb6 100644
--- a/cpp/include/cudf/detail/aggregation/aggregation.hpp
+++ b/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
     }
   }
diff --git a/cpp/include/cudf/detail/indexalator.cuh b/cpp/include/cudf/detail/indexalator.cuh
index 3657d700397..407001ecb71 100644
--- a/cpp/include/cudf/detail/indexalator.cuh
+++ b/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");
     }
   };
 
diff --git a/cpp/include/cudf/detail/utilities/assert.cuh b/cpp/include/cudf/detail/utilities/assert.cuh
index fd2a0e44200..69f9e2d3791 100644
--- a/cpp/include/cudf/detail/utilities/assert.cuh
+++ b/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)
diff --git a/cpp/include/cudf/detail/utilities/hash_functions.cuh b/cpp/include/cudf/detail/utilities/hash_functions.cuh
index 09d94d10e79..3375fbc39fc 100644
--- a/cpp/include/cudf/detail/utilities/hash_functions.cuh
+++ b/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>
diff --git a/cpp/include/cudf/table/row_operators.cuh b/cpp/include/cudf/table/row_operators.cuh
index 4eca03a800c..4d1611c7a6e 100644
--- a/cpp/include/cudf/table/row_operators.cuh
+++ b/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:
diff --git a/cpp/include/cudf/utilities/type_dispatcher.hpp b/cpp/include/cudf/utilities/type_dispatcher.hpp
index 5c3a6b128b5..7b111b11eb4 100644
--- a/cpp/include/cudf/utilities/type_dispatcher.hpp
+++ b/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
     }
   }
diff --git a/cpp/src/datetime/datetime_ops.cu b/cpp/src/datetime/datetime_ops.cu
index 866dae46327..304b577afb0 100644
--- a/cpp/src/datetime/datetime_ops.cu
+++ b/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();
   }
 };
 
diff --git a/cpp/src/groupby/hash/multi_pass_kernels.cuh b/cpp/src/groupby/hash/multi_pass_kernels.cuh
index 15a38029bc4..02977bb4ece 100644
--- a/cpp/src/groupby/hash/multi_pass_kernels.cuh
+++ b/cpp/src/groupby/hash/multi_pass_kernels.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 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.
@@ -67,7 +67,7 @@ struct var_hash_functor {
                                                                   size_type source_index,
                                                                   size_type target_index) noexcept
   {
-    CUDF_UNREACHABLE("Invalid source type for std, var aggregation combination.");
+    cudf_assert(false and "Invalid source type for std, var aggregation combination.");
   }
 
   template <typename Source>
diff --git a/cpp/src/hash/md5_hash.cu b/cpp/src/hash/md5_hash.cu
index 0b04cd86029..41a57503230 100644
--- a/cpp/src/hash/md5_hash.cu
+++ b/cpp/src/hash/md5_hash.cu
@@ -176,7 +176,7 @@ struct HasherDispatcher {
       hasher->process(input_col.element<Element>(row_index));
     } else {
       (void)row_index;
-      CUDF_UNREACHABLE("Unsupported type for hash function.");
+      cudf_assert(false && "Unsupported type for hash function.");
     }
   }
 };
@@ -202,7 +202,7 @@ struct ListHasherDispatcher {
     } else {
       (void)offset_begin;
       (void)offset_end;
-      CUDF_UNREACHABLE("Unsupported type for hash function.");
+      cudf_assert(false && "Unsupported type for hash function.");
     }
   }
 };
@@ -263,7 +263,7 @@ std::unique_ptr<column> md5_hash(table_view const& input,
             auto const data_col = col.child(lists_column_view::child_column_index);
             auto const offsets  = col.child(lists_column_view::offsets_column_index);
             if (data_col.type().id() == type_id::LIST) {
-              CUDF_UNREACHABLE("Nested list unsupported");
+              cudf_assert(false && "Nested list unsupported");
             }
             auto const offset_begin = offsets.element<size_type>(row_index);
             auto const offset_end   = offsets.element<size_type>(row_index + 1);
diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu
index 059df283c94..6a71ae38ad7 100644
--- a/cpp/src/io/orc/reader_impl.cu
+++ b/cpp/src/io/orc/reader_impl.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2019-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.
@@ -444,7 +444,7 @@ rmm::device_buffer reader::impl::decompress_stripe_data(
             gpu_unsnap(inflate_in.data(), inflate_out.data(), num_compressed_blocks, stream));
         }
         break;
-      default: CUDF_FAIL("Unexpected decompression dispatch"); break;
+      default: CUDF_EXPECTS(false, "Unexpected decompression dispatch"); break;
     }
     decompress_check(inflate_out_view, any_block_failure.device_ptr(), stream);
   }
diff --git a/cpp/src/io/orc/writer_impl.cu b/cpp/src/io/orc/writer_impl.cu
index 30385d395f1..6e1f9bcfcb9 100644
--- a/cpp/src/io/orc/writer_impl.cu
+++ b/cpp/src/io/orc/writer_impl.cu
@@ -86,7 +86,7 @@ orc::CompressionKind to_orc_compression(compression_type compression)
     case compression_type::AUTO:
     case compression_type::SNAPPY: return orc::CompressionKind::SNAPPY;
     case compression_type::NONE: return orc::CompressionKind::NONE;
-    default: CUDF_FAIL("Unsupported compression type"); return orc::CompressionKind::NONE;
+    default: CUDF_EXPECTS(false, "Unsupported compression type"); return orc::CompressionKind::NONE;
   }
 }
 
diff --git a/cpp/src/io/parquet/chunk_dict.cu b/cpp/src/io/parquet/chunk_dict.cu
index f61cfa83579..5589f87e57e 100644
--- a/cpp/src/io/parquet/chunk_dict.cu
+++ b/cpp/src/io/parquet/chunk_dict.cu
@@ -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.
@@ -69,8 +69,9 @@ struct map_insert_fn {
       auto equality_fn = equality_functor<T>{col};
       return map.insert(std::make_pair(i, i), hash_fn, equality_fn);
     } else {
-      CUDF_UNREACHABLE("Unsupported type to insert in map");
+      cudf_assert(false && "Unsupported type to insert in map");
     }
+    return false;
   }
 };
 
@@ -78,15 +79,16 @@ struct map_find_fn {
   map_type::device_view& map;
 
   template <typename T>
-  __device__ map_type::device_view::iterator operator()(column_device_view const& col, size_type i)
+  __device__ auto operator()(column_device_view const& col, size_type i)
   {
     if constexpr (column_device_view::has_element_accessor<T>()) {
       auto hash_fn     = hash_functor<T>{col};
       auto equality_fn = equality_functor<T>{col};
       return map.find(i, hash_fn, equality_fn);
     } else {
-      CUDF_UNREACHABLE("Unsupported type to find in map");
+      cudf_assert(false && "Unsupported type to insert in map");
     }
+    return map.end();
   }
 };
 
@@ -159,7 +161,7 @@ __global__ void __launch_bounds__(block_size, 1)
               }
             case Type::FIXED_LEN_BYTE_ARRAY:
               if (data_col.type().id() == type_id::DECIMAL128) { return sizeof(__int128_t); }
-            default: CUDF_UNREACHABLE("Unsupported type for dictionary encoding");
+            default: cudf_assert(false && "Unsupported type for dictionary encoding"); return 0;
           }
         }();
       }
diff --git a/cpp/src/io/parquet/reader_impl.cu b/cpp/src/io/parquet/reader_impl.cu
index df940bb15d3..6899884603f 100644
--- a/cpp/src/io/parquet/reader_impl.cu
+++ b/cpp/src/io/parquet/reader_impl.cu
@@ -1279,7 +1279,7 @@ rmm::device_buffer reader::impl::decompress_page_data(
                                 argc - start_pos,
                                 stream));
           break;
-        default: CUDF_FAIL("Unexpected decompression dispatch"); break;
+        default: CUDF_EXPECTS(false, "Unexpected decompression dispatch"); break;
       }
       CUDA_TRY(cudaMemcpyAsync(inflate_out.host_ptr(start_pos),
                                inflate_out.device_ptr(start_pos),
diff --git a/cpp/src/io/parquet/writer_impl.cu b/cpp/src/io/parquet/writer_impl.cu
index a29164ba051..f8097c8f262 100644
--- a/cpp/src/io/parquet/writer_impl.cu
+++ b/cpp/src/io/parquet/writer_impl.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2019-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.
@@ -78,7 +78,9 @@ parquet::Compression to_parquet_compression(compression_type compression)
     case compression_type::AUTO:
     case compression_type::SNAPPY: return parquet::Compression::SNAPPY;
     case compression_type::NONE: return parquet::Compression::UNCOMPRESSED;
-    default: CUDF_FAIL("Unsupported compression type");
+    default:
+      CUDF_EXPECTS(false, "Unsupported compression type");
+      return parquet::Compression::UNCOMPRESSED;
   }
 }
 
diff --git a/cpp/src/quantiles/quantiles_util.hpp b/cpp/src/quantiles/quantiles_util.hpp
index 171b81152ff..a5dc643a688 100644
--- a/cpp/src/quantiles/quantiles_util.hpp
+++ b/cpp/src/quantiles/quantiles_util.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2019, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -144,13 +144,13 @@ CUDF_HOST_DEVICE inline Result select_quantile(ValueAccessor get_value,
 
     case interpolation::NEAREST: return static_cast<Result>(get_value(idx.nearest));
 
-    default: {
-#ifndef __CUDA_ARCH__
-      CUDF_FAIL("Invalid interpolation operation for quantiles.");
+    default:
+#if defined(__CUDA_ARCH__)
+      cudf_assert(false && "Invalid interpolation operation for quantiles");
+      return Result();
 #else
-      CUDF_UNREACHABLE("Invalid interpolation operation for quantiles");
+      CUDF_FAIL("Invalid interpolation operation for quantiles.");
 #endif
-    }
   }
 }
 
@@ -176,14 +176,14 @@ CUDF_HOST_DEVICE inline Result select_quantile_data(Iterator begin,
 
     case interpolation::MIDPOINT:
       return interpolate::midpoint<Result>(*(begin + idx.lower), *(begin + idx.higher));
-    default: {
-#ifndef __CUDA_ARCH__
-      CUDF_FAIL("Invalid interpolation operation for quantiles.");
+  }
+
+#if defined(__CUDA_ARCH__)
+  cudf_assert(false && "Invalid interpolation operation for quantiles");
+  return Result();
 #else
-      CUDF_UNREACHABLE("Invalid interpolation operation for quantiles");
+  CUDF_FAIL("Invalid interpolation operation for quantiles.");
 #endif
-    }
-  }
 }
 
 template <typename Iterator>
@@ -203,14 +203,14 @@ CUDF_HOST_DEVICE inline bool select_quantile_validity(Iterator begin,
 
     case interpolation::LINEAR:
     case interpolation::MIDPOINT: return *(begin + idx.lower) and *(begin + idx.higher);
-    default: {
-#ifndef __CUDA_ARCH__
-      CUDF_FAIL("Invalid interpolation operation for quantiles.");
+  }
+
+#if defined(__CUDA_ARCH__)
+  cudf_assert(false && "Invalid interpolation operation for quantiles");
+  return false;
 #else
-      CUDF_UNREACHABLE("Invalid interpolation operation for quantiles");
+  CUDF_FAIL("Invalid interpolation operation for quantiles.");
 #endif
-    }
-  }
 }
 
 }  // namespace detail
diff --git a/cpp/tests/reductions/scan_tests.cpp b/cpp/tests/reductions/scan_tests.cpp
index d533a91f4d0..efd56ad45ee 100644
--- a/cpp/tests/reductions/scan_tests.cpp
+++ b/cpp/tests/reductions/scan_tests.cpp
@@ -128,7 +128,10 @@ struct ScanTest : public BaseScanTest<T> {
       switch (agg->kind) {
         case aggregation::MIN: return std::string{"\xF7\xBF\xBF\xBF"};
         case aggregation::MAX: return std::string{};
-        default: CUDF_FAIL("Unsupported aggregation");
+        default: {
+          CUDF_FAIL("Unsupported aggregation");
+          return HostType{};
+        }
       }
     } else {
       switch (agg->kind) {
@@ -136,7 +139,10 @@ struct ScanTest : public BaseScanTest<T> {
         case aggregation::PRODUCT: return HostType{1};
         case aggregation::MIN: return std::numeric_limits<HostType>::max();
         case aggregation::MAX: return std::numeric_limits<HostType>::lowest();
-        default: CUDF_FAIL("Unsupported aggregation");
+        default: {
+          CUDF_FAIL("Unsupported aggregation");
+          return HostType{};
+        }
       }
     }
   }
diff --git a/cpp/tests/unary/cast_tests.cpp b/cpp/tests/unary/cast_tests.cpp
index f53498bccec..aaee898141a 100644
--- a/cpp/tests/unary/cast_tests.cpp
+++ b/cpp/tests/unary/cast_tests.cpp
@@ -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.
@@ -153,7 +153,7 @@ inline cudf::column make_exp_chrono_column(cudf::type_id type_id)
         rmm::device_buffer{test_durations_ns.data(),
                            test_durations_ns.size() * sizeof(test_durations_ns.front()),
                            rmm::cuda_stream_default});
-    default: CUDF_FAIL("Unsupported type_id");
+    default: CUDF_FAIL("");
   }
 };