Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Add CUDF_UNREACHABLE macro. #9727

Merged
merged 19 commits into from
Mar 18, 2022
Merged
Changes from 1 commit
Commits
Show all changes
19 commits
Select commit Hold shift + click to select a range
937a512
Add CUDF_UNREACHABLE macro.
bdice Nov 18, 2021
106b708
Add unreachable macro to type dispatcher.
bdice Nov 18, 2021
f9894d1
Add unreachable macro to AST operator dispatcher.
bdice Nov 18, 2021
8f41a74
Merge remote-tracking branch 'upstream/branch-22.02' into add-cudf-un…
bdice Nov 19, 2021
ab1cdb9
Update CUDF_UNREACHABLE docs.
bdice Nov 19, 2021
77c97ae
Throw error in host-side type dispatcher if invalid type is passed.
bdice Nov 19, 2021
2c39210
Merge remote-tracking branch 'upstream/branch-22.04' into add-cudf-un…
bdice Jan 24, 2022
2d61ea9
Use CUDF_UNREACHABLE only on device in AST operator dispatch.
bdice Jan 24, 2022
a0d9baa
Merge remote-tracking branch 'upstream/branch-22.04' into add-cudf-un…
bdice Jan 26, 2022
8c7d978
Merge remote-tracking branch 'upstream/branch-22.04' into add-cudf-un…
bdice Jan 28, 2022
52ad3a0
Use CUDF_UNREACHABLE (may be too aggressive, some cases may need to f…
bdice Jan 31, 2022
7de2aa0
Merge remote-tracking branch 'upstream/branch-22.04' into add-cudf-un…
bdice Mar 2, 2022
3bcfd60
Replace cudf_assert(false...) in device functions with CUDF_UNREACHABLE.
bdice Mar 2, 2022
c320c1f
Merge remote-tracking branch 'upstream/branch-22.04' into add-cudf-un…
bdice Mar 15, 2022
6d8d498
Merge remote-tracking branch 'upstream/branch-22.04' into add-cudf-un…
bdice Mar 17, 2022
8c2b1dc
Add comment about assert in debug mode.
bdice Mar 17, 2022
1407315
Remove unnecessary returns.
bdice Mar 17, 2022
b66c4c4
Add explicit type since auto deduction fails with CUDF_UNREACHABLE.
bdice Mar 17, 2022
3d3000c
Merge remote-tracking branch 'upstream/branch-22.04' into add-cudf-un…
bdice Mar 18, 2022
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
Use CUDF_UNREACHABLE (may be too aggressive, some cases may need to f…
…ail rather than be unreachable).
bdice committed Jan 31, 2022
commit 52ad3a05c760c26b988a1dfc941f14aab935c955
2 changes: 1 addition & 1 deletion cpp/benchmarks/type_dispatcher/type_dispatcher.cu
Original file line number Diff line number Diff line change
@@ -119,7 +119,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.");
}
};

16 changes: 7 additions & 9 deletions cpp/include/cudf/ast/detail/expression_evaluator.cuh
Original file line number Diff line number Diff line change
@@ -178,7 +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.");
CUDF_UNREACHABLE("This method is not implemented.");
vuule marked this conversation as resolved.
Show resolved Hide resolved
// Unreachable return used to silence compiler warnings.
return {};
}
@@ -190,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_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 +334,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 +496,7 @@ struct expression_evaluator {
op,
thread_intermediate_storage);
} else {
cudf_assert(false && "Invalid operator arity.");
CUDF_UNREACHABLE("Invalid operator arity.");
}
}
}
@@ -567,7 +565,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 +628,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 +696,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.");
}
};

12 changes: 6 additions & 6 deletions cpp/include/cudf/ast/detail/operators.hpp
Original file line number Diff line number Diff line change
@@ -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.");
#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
}
};
12 changes: 1 addition & 11 deletions cpp/include/cudf/detail/aggregation/aggregation.hpp
Original file line number Diff line number Diff line change
@@ -1431,17 +1431,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
}
}
6 changes: 2 additions & 4 deletions cpp/src/datetime/datetime_ops.cu
Original file line number Diff line number Diff line change
@@ -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();
}
};

2 changes: 1 addition & 1 deletion cpp/src/hash/hashing.cu
Original file line number Diff line number Diff line change
@@ -99,7 +99,7 @@ std::unique_ptr<column> hash(table_view const& input,
return serial_murmur_hash3_32<MurmurHash3_32>(input, seed, stream, mr);
case (hash_id::HASH_SPARK_MURMUR3):
return serial_murmur_hash3_32<SparkMurmurHash3_32>(input, seed, stream, mr);
default: return nullptr;
default: CUDF_FAIL("Unsupported hash function.");
bdice marked this conversation as resolved.
Show resolved Hide resolved
}
}

2 changes: 1 addition & 1 deletion cpp/src/io/orc/reader_impl.cu
Original file line number Diff line number Diff line change
@@ -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) {
2 changes: 1 addition & 1 deletion cpp/src/io/orc/writer_impl.cu
Original file line number Diff line number Diff line change
@@ -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;
}
}

2 changes: 1 addition & 1 deletion cpp/src/io/parquet/chunk_dict.cu
Original file line number Diff line number Diff line change
@@ -161,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_assert(false && "Unsupported type for dictionary encoding"); return 0;
default: CUDF_UNREACHABLE("Unsupported type for dictionary encoding");
}
}();
}
2 changes: 1 addition & 1 deletion cpp/src/io/parquet/reader_impl.cu
Original file line number Diff line number Diff line change
@@ -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),
4 changes: 1 addition & 3 deletions cpp/src/io/parquet/writer_impl.cu
Original file line number Diff line number Diff line change
@@ -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"); return parquet::Compression::UNCOMPRESSED;
bdice marked this conversation as resolved.
Show resolved Hide resolved
}
}

34 changes: 17 additions & 17 deletions cpp/src/quantiles/quantiles_util.hpp
Original file line number Diff line number Diff line change
@@ -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:
#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<Result>(*(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 <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);
}

#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
10 changes: 2 additions & 8 deletions cpp/tests/reductions/scan_tests.cpp
Original file line number Diff line number Diff line change
@@ -125,21 +125,15 @@ 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");
return HostType{};
}
default: CUDF_FAIL("Unsupported aggregation");
}
} else {
switch (agg->kind) {
case aggregation::SUM: return HostType{0};
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");
return HostType{};
}
default: CUDF_FAIL("Unsupported aggregation");
}
}
}
2 changes: 1 addition & 1 deletion cpp/tests/unary/cast_tests.cpp
Original file line number Diff line number Diff line change
@@ -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");
}
};