From 48cebf7ac498439ebae3d5a50581a552f8fed88d Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 18 Mar 2022 16:27:23 -0500 Subject: [PATCH] Add CUDF_UNREACHABLE macro. (#9727) Resolves #7753. I replaced all instances of `cudf_assert(false && "message");` with `CUDF_UNREACHABLE("message");`. There are a few instances where the condition of the assertion is not always `false`, and thus the code following it may still be reachable. I did not change those cases. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Jake Hemstad (https://github.com/jrhemstad) - https://github.com/nvdbaranec URL: https://github.com/rapidsai/cudf/pull/9727 --- .../type_dispatcher/type_dispatcher.cu | 2 +- .../cudf/ast/detail/expression_evaluator.cuh | 22 +++++------- cpp/include/cudf/ast/detail/operators.hpp | 14 ++++---- .../cudf/column/column_device_view.cuh | 3 +- .../cudf/detail/aggregation/aggregation.cuh | 2 +- .../cudf/detail/aggregation/aggregation.hpp | 12 +------ cpp/include/cudf/detail/indexalator.cuh | 7 ++-- cpp/include/cudf/detail/utilities/assert.cuh | 28 ++++++++++++++- .../cudf/detail/utilities/hash_functions.cuh | 17 ++++----- cpp/include/cudf/table/row_operators.cuh | 12 +++---- .../cudf/utilities/type_dispatcher.hpp | 13 ++----- cpp/src/datetime/datetime_ops.cu | 6 ++-- cpp/src/groupby/hash/multi_pass_kernels.cuh | 4 +-- cpp/src/hash/md5_hash.cu | 6 ++-- cpp/src/io/orc/reader_impl.cu | 4 +-- cpp/src/io/orc/writer_impl.cu | 2 +- cpp/src/io/parquet/chunk_dict.cu | 12 +++---- cpp/src/io/parquet/reader_impl.cu | 2 +- cpp/src/io/parquet/writer_impl.cu | 6 ++-- cpp/src/quantiles/quantiles_util.hpp | 36 +++++++++---------- cpp/tests/reductions/scan_tests.cpp | 10 ++---- cpp/tests/unary/cast_tests.cpp | 4 +-- 22 files changed, 102 insertions(+), 122 deletions(-) diff --git a/cpp/benchmarks/type_dispatcher/type_dispatcher.cu b/cpp/benchmarks/type_dispatcher/type_dispatcher.cu index 3be599e8c41..6ab6f9a2095 100644 --- a/cpp/benchmarks/type_dispatcher/type_dispatcher.cu +++ b/cpp/benchmarks/type_dispatcher/type_dispatcher.cu @@ -120,7 +120,7 @@ struct RowHandle { template ())> __device__ void operator()(cudf::mutable_column_device_view source, cudf::size_type index) { - cudf_assert(false && "Unsupported type."); + CUDF_UNREACHABLE("Unsupported type."); } }; diff --git a/cpp/include/cudf/ast/detail/expression_evaluator.cuh b/cpp/include/cudf/ast/detail/expression_evaluator.cuh index 2bfe1b03dd3..a179fe10774 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, 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) { _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 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* thread_intermediate_storage, possibly_null_value_t 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* 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* 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."); } }; diff --git a/cpp/include/cudf/ast/detail/operators.hpp b/cpp/include/cudf/ast/detail/operators.hpp index 63048a993c7..3c6c9de055d 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()(std::forward(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 } }; diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh index ea6552a1d61..3ee73282438 100644 --- a/cpp/include/cudf/column/column_device_view.cuh +++ b/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() and std::is_unsigned_v))> __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"); } }; diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh index 60ea10cef88..02121957184 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.cuh +++ b/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."); } }; diff --git a/cpp/include/cudf/detail/aggregation/aggregation.hpp b/cpp/include/cudf/detail/aggregation/aggregation.hpp index eba24dd2d13..dfdc653b2a3 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.hpp +++ b/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()(std::forward(args)...)); - return return_type(); + CUDF_UNREACHABLE("Unsupported aggregation."); #endif } } diff --git a/cpp/include/cudf/detail/indexalator.cuh b/cpp/include/cudf/detail/indexalator.cuh index 11c82da8097..a63faa40e1d 100644 --- a/cpp/include/cudf/detail/indexalator.cuh +++ b/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 { template ()>* = 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 { template ()>* = 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"); } }; diff --git a/cpp/include/cudf/detail/utilities/assert.cuh b/cpp/include/cudf/detail/utilities/assert.cuh index 69f9e2d3791..fd2a0e44200 100644 --- a/cpp/include/cudf/detail/utilities/assert.cuh +++ b/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(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 327d3acea6c..264c80f223c 100644 --- a/cpp/include/cudf/detail/utilities/hash_functions.cuh +++ b/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() && !is_chrono()) { return thrust::make_pair(reinterpret_cast(&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::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::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 @@ -551,16 +549,14 @@ template <> hash_value_type __device__ inline SparkMurmurHash3_32::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::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, 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 diff --git a/cpp/include/cudf/table/row_operators.cuh b/cpp/include/cudf/table/row_operators.cuh index b5d46935fe8..20845818b0f 100644 --- a/cpp/include/cudf/table/row_operators.cuh +++ b/cpp/include/cudf/table/row_operators.cuh @@ -217,8 +217,7 @@ class element_equality_comparator { std::enable_if_t()>* = 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()>* = 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 ())> __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 ())> __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: diff --git a/cpp/include/cudf/utilities/type_dispatcher.hpp b/cpp/include/cudf/utilities/type_dispatcher.hpp index 7b111b11eb4..5c3a6b128b5 100644 --- a/cpp/include/cudf/utilities/type_dispatcher.hpp +++ b/cpp/include/cudf/utilities/type_dispatcher.hpp @@ -511,18 +511,9 @@ CUDF_HOST_DEVICE __forceinline__ constexpr decltype(auto) type_dispatcher(cudf:: std::forward(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()(std::forward(args)...)); - return return_type(); + CUDF_UNREACHABLE("Invalid type_id."); #endif } } diff --git a/cpp/src/datetime/datetime_ops.cu b/cpp/src/datetime/datetime_ops.cu index 4dbe9faaa47..1bac2df8d2b 100644 --- a/cpp/src/datetime/datetime_ops.cu +++ b/cpp/src/datetime/datetime_ops.cu @@ -105,9 +105,8 @@ struct RoundFunctor { case rounding_function::CEIL: return cuda::std::chrono::ceil(dt); case rounding_function::FLOOR: return cuda::std::chrono::floor(dt); case rounding_function::ROUND: return cuda::std::chrono::round(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( RoundFunctor{}(round_kind, ts)); - default: cudf_assert(false && "Unsupported datetime rounding resolution."); + default: CUDF_UNREACHABLE("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 02977bb4ece..15a38029bc4 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, 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. @@ -67,7 +67,7 @@ struct var_hash_functor { size_type source_index, size_type target_index) noexcept { - cudf_assert(false and "Invalid source type for std, var aggregation combination."); + CUDF_UNREACHABLE("Invalid source type for std, var aggregation combination."); } template diff --git a/cpp/src/hash/md5_hash.cu b/cpp/src/hash/md5_hash.cu index 589e443d21d..a1531a7b094 100644 --- a/cpp/src/hash/md5_hash.cu +++ b/cpp/src/hash/md5_hash.cu @@ -175,7 +175,7 @@ struct HasherDispatcher { hasher->process(input_col.element(row_index)); } else { (void)row_index; - cudf_assert(false && "Unsupported type for hash function."); + CUDF_UNREACHABLE("Unsupported type for hash function."); } } }; @@ -201,7 +201,7 @@ struct ListHasherDispatcher { } else { (void)offset_begin; (void)offset_end; - cudf_assert(false && "Unsupported type for hash function."); + CUDF_UNREACHABLE("Unsupported type for hash function."); } } }; @@ -262,7 +262,7 @@ std::unique_ptr 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_assert(false && "Nested list unsupported"); + CUDF_UNREACHABLE("Nested list unsupported"); } auto const offset_begin = offsets.element(row_index); auto const offset_end = offsets.element(row_index + 1); diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index f133b79a27e..89488d75735 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2021, 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. @@ -408,7 +408,7 @@ rmm::device_buffer reader::impl::decompress_stripe_data( gpu_unsnap(inflate_in.data(), inflate_out.data(), num_compressed_blocks, stream)); } break; - default: CUDF_EXPECTS(false, "Unexpected decompression dispatch"); break; + default: CUDF_FAIL("Unexpected decompression dispatch"); break; } } if (num_uncompressed_blocks > 0) { diff --git a/cpp/src/io/orc/writer_impl.cu b/cpp/src/io/orc/writer_impl.cu index a917dbf93a5..c2cf873e5bf 100644 --- a/cpp/src/io/orc/writer_impl.cu +++ b/cpp/src/io/orc/writer_impl.cu @@ -75,7 +75,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_EXPECTS(false, "Unsupported compression type"); return orc::CompressionKind::NONE; + default: CUDF_FAIL("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 5589f87e57e..f61cfa83579 100644 --- a/cpp/src/io/parquet/chunk_dict.cu +++ b/cpp/src/io/parquet/chunk_dict.cu @@ -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. @@ -69,9 +69,8 @@ struct map_insert_fn { auto equality_fn = equality_functor{col}; return map.insert(std::make_pair(i, i), hash_fn, equality_fn); } else { - cudf_assert(false && "Unsupported type to insert in map"); + CUDF_UNREACHABLE("Unsupported type to insert in map"); } - return false; } }; @@ -79,16 +78,15 @@ struct map_find_fn { map_type::device_view& map; template - __device__ auto operator()(column_device_view const& col, size_type i) + __device__ map_type::device_view::iterator operator()(column_device_view const& col, size_type i) { if constexpr (column_device_view::has_element_accessor()) { auto hash_fn = hash_functor{col}; auto equality_fn = equality_functor{col}; return map.find(i, hash_fn, equality_fn); } else { - cudf_assert(false && "Unsupported type to insert in map"); + CUDF_UNREACHABLE("Unsupported type to find in map"); } - return map.end(); } }; @@ -161,7 +159,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_assert(false && "Unsupported type for dictionary encoding"); return 0; + default: CUDF_UNREACHABLE("Unsupported type for dictionary encoding"); } }(); } diff --git a/cpp/src/io/parquet/reader_impl.cu b/cpp/src/io/parquet/reader_impl.cu index 885f36aeca4..4b619133d6a 100644 --- a/cpp/src/io/parquet/reader_impl.cu +++ b/cpp/src/io/parquet/reader_impl.cu @@ -1222,7 +1222,7 @@ rmm::device_buffer reader::impl::decompress_page_data( argc - start_pos, stream)); break; - default: CUDF_EXPECTS(false, "Unexpected decompression dispatch"); break; + default: CUDF_FAIL("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 a9306275b26..c259b4383e7 100644 --- a/cpp/src/io/parquet/writer_impl.cu +++ b/cpp/src/io/parquet/writer_impl.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2021, 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. @@ -71,9 +71,7 @@ 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_EXPECTS(false, "Unsupported compression type"); - return parquet::Compression::UNCOMPRESSED; + default: CUDF_FAIL("Unsupported compression type"); } } diff --git a/cpp/src/quantiles/quantiles_util.hpp b/cpp/src/quantiles/quantiles_util.hpp index a5dc643a688..171b81152ff 100644 --- a/cpp/src/quantiles/quantiles_util.hpp +++ b/cpp/src/quantiles/quantiles_util.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019, 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. @@ -144,13 +144,13 @@ CUDF_HOST_DEVICE inline Result select_quantile(ValueAccessor get_value, case interpolation::NEAREST: return static_cast(get_value(idx.nearest)); - default: -#if defined(__CUDA_ARCH__) - cudf_assert(false && "Invalid interpolation operation for quantiles"); - return Result(); -#else + default: { +#ifndef __CUDA_ARCH__ CUDF_FAIL("Invalid interpolation operation for quantiles."); +#else + CUDF_UNREACHABLE("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(*(begin + idx.lower), *(begin + idx.higher)); - } - -#if defined(__CUDA_ARCH__) - cudf_assert(false && "Invalid interpolation operation for quantiles"); - return Result(); + default: { +#ifndef __CUDA_ARCH__ + CUDF_FAIL("Invalid interpolation operation for quantiles."); #else - CUDF_FAIL("Invalid interpolation operation for quantiles."); + CUDF_UNREACHABLE("Invalid interpolation operation for quantiles"); #endif + } + } } template @@ -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); - } - -#if defined(__CUDA_ARCH__) - cudf_assert(false && "Invalid interpolation operation for quantiles"); - return false; + default: { +#ifndef __CUDA_ARCH__ + CUDF_FAIL("Invalid interpolation operation for quantiles."); #else - CUDF_FAIL("Invalid interpolation operation for quantiles."); + CUDF_UNREACHABLE("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 94eb36eee8e..75ea5900ec8 100644 --- a/cpp/tests/reductions/scan_tests.cpp +++ b/cpp/tests/reductions/scan_tests.cpp @@ -126,10 +126,7 @@ struct ScanTest : public BaseScanTest { switch (agg->kind) { case aggregation::MIN: return std::string{"\xF7\xBF\xBF\xBF"}; case aggregation::MAX: return std::string{}; - default: { - CUDF_FAIL("Unsupported aggregation"); - return HostType{}; - } + default: CUDF_FAIL("Unsupported aggregation"); } } else { switch (agg->kind) { @@ -137,10 +134,7 @@ struct ScanTest : public BaseScanTest { case aggregation::PRODUCT: return HostType{1}; case aggregation::MIN: return std::numeric_limits::max(); case aggregation::MAX: return std::numeric_limits::lowest(); - default: { - CUDF_FAIL("Unsupported aggregation"); - return HostType{}; - } + default: CUDF_FAIL("Unsupported aggregation"); } } } diff --git a/cpp/tests/unary/cast_tests.cpp b/cpp/tests/unary/cast_tests.cpp index 906380f5e87..ceaff4b7c58 100644 --- a/cpp/tests/unary/cast_tests.cpp +++ b/cpp/tests/unary/cast_tests.cpp @@ -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. @@ -150,7 +150,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(""); + default: CUDF_FAIL("Unsupported type_id"); } };