diff --git a/cpp/include/cudf/detail/aggregation/device_aggregators.cuh b/cpp/include/cudf/detail/aggregation/device_aggregators.cuh index 4397af433621..ef0410f958cc 100644 --- a/cpp/include/cudf/detail/aggregation/device_aggregators.cuh +++ b/cpp/include/cudf/detail/aggregation/device_aggregators.cuh @@ -16,7 +16,7 @@ #include #include -#include +#include #include namespace cudf::detail { @@ -134,9 +134,7 @@ template (cudf::is_fixed_point() && cudf::has_atomic_support>()) || cuda::std::is_same_v) struct update_target_element { - using DeviceType = device_storage_type_t; - static constexpr auto type_max = cuda::std::numeric_limits::max(); - static constexpr auto type_min = cuda::std::numeric_limits::min(); + using DeviceType = device_storage_type_t; __device__ void operator()(mutable_column_device_view target, size_type target_index, @@ -155,12 +153,9 @@ struct update_target_element { *(overflow_column.data() + target_index)}; if (bool_ref.load(cuda::memory_order_relaxed)) { return; } - // TODO: to be replaced by CCCL equivalents once https://github.com/NVIDIA/cccl/pull/3755 is - // ready - auto const overflow = - source_value > 0 ? old_sum > type_max - source_value : old_sum < type_min - source_value; - - if (overflow) { cudf::detail::atomic_max(&overflow_column.element(target_index), true); } + if (cuda::add_overflow(old_sum, source_value).overflow) { + cudf::detail::atomic_max(&overflow_column.element(target_index), true); + } } }; diff --git a/cpp/include/cudf/fixed_point/fixed_point.hpp b/cpp/include/cudf/fixed_point/fixed_point.hpp index 6dcc4aed20a1..0b7c4428a563 100644 --- a/cpp/include/cudf/fixed_point/fixed_point.hpp +++ b/cpp/include/cudf/fixed_point/fixed_point.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -9,6 +9,7 @@ #include #include +#include #include #include #include @@ -603,8 +604,7 @@ class fixed_point { template CUDF_HOST_DEVICE inline auto addition_overflow(T lhs, T rhs) { - return rhs > 0 ? lhs > cuda::std::numeric_limits::max() - rhs - : lhs < cuda::std::numeric_limits::min() - rhs; + return cuda::add_overflow(lhs, rhs).overflow; } /** @brief Function for identifying integer overflow when subtracting @@ -618,8 +618,7 @@ CUDF_HOST_DEVICE inline auto addition_overflow(T lhs, T rhs) template CUDF_HOST_DEVICE inline auto subtraction_overflow(T lhs, T rhs) { - return rhs > 0 ? lhs < cuda::std::numeric_limits::min() + rhs - : lhs > cuda::std::numeric_limits::max() + rhs; + return cuda::sub_overflow(lhs, rhs).overflow; } /** @brief Function for identifying integer overflow when dividing @@ -633,7 +632,7 @@ CUDF_HOST_DEVICE inline auto subtraction_overflow(T lhs, T rhs) template CUDF_HOST_DEVICE inline auto division_overflow(T lhs, T rhs) { - return lhs == cuda::std::numeric_limits::min() && rhs == -1; + return cuda::div_overflow(lhs, rhs).overflow; } /** @brief Function for identifying integer overflow when multiplying @@ -647,11 +646,7 @@ CUDF_HOST_DEVICE inline auto division_overflow(T lhs, T rhs) template CUDF_HOST_DEVICE inline auto multiplication_overflow(T lhs, T rhs) { - auto const min = cuda::std::numeric_limits::min(); - auto const max = cuda::std::numeric_limits::max(); - if (rhs > 0) { return lhs > max / rhs || lhs < min / rhs; } - if (rhs < -1) { return lhs > min / rhs || lhs < max / rhs; } - return rhs == -1 && lhs == min; + return cuda::mul_overflow(lhs, rhs).overflow; } // PLUS Operation diff --git a/cpp/src/bitmask/null_mask.cu b/cpp/src/bitmask/null_mask.cu index 4bb7d6d2dc94..cdbbb64660cf 100644 --- a/cpp/src/bitmask/null_mask.cu +++ b/cpp/src/bitmask/null_mask.cu @@ -25,11 +25,11 @@ #include #include #include +#include #include #include #include -#include #include namespace cudf { @@ -211,7 +211,7 @@ void set_null_masks(cudf::host_span bitmasks, auto const num_words = num_bitmask_words(end_bits[i]) - begin_bits[i] / detail::size_in_bits(); // Handle overflow if any - if (num_words >= std::numeric_limits::max() - cumulative_null_mask_words) { + if (cuda::add_overflow(cumulative_null_mask_words, num_words).overflow) { average_nullmask_words += cudf::util::div_rounding_up_safe(cumulative_null_mask_words, num_bitmasks); cumulative_null_mask_words = 0; diff --git a/cpp/src/io/orc/orc.cpp b/cpp/src/io/orc/orc.cpp index be87d51c82ac..03ebffef76c1 100644 --- a/cpp/src/io/orc/orc.cpp +++ b/cpp/src/io/orc/orc.cpp @@ -12,9 +12,9 @@ #include #include +#include #include -#include #include #include @@ -432,9 +432,9 @@ host_span orc_decompressor::decompress_blocks(host_span::max() - m_blockSize, - "ORC decompression: compression block size overflow"); - max_dst_length += m_blockSize; + auto const next = cuda::add_overflow(max_dst_length, m_blockSize); + CUDF_EXPECTS(!next.overflow, "ORC decompression: compression block size overflow"); + max_dst_length = next.value; } i += block_len; CUDF_EXPECTS(i <= src.size() and block_len <= m_blockSize, "Error in decompression"); diff --git a/cpp/src/io/parquet/reader_impl_helpers.cpp b/cpp/src/io/parquet/reader_impl_helpers.cpp index 7c4462c3463d..31da9b5bebab 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.cpp +++ b/cpp/src/io/parquet/reader_impl_helpers.cpp @@ -20,6 +20,7 @@ #include #include +#include #include #include @@ -1368,36 +1369,38 @@ std::vector> aggregate_reader_metadata::apply_byte_bounds auto filtered_row_group_indices = std::vector>(input_row_group_indices.size()); - std::for_each( - input_row_group_indices.front().begin(), - input_row_group_indices.front().end(), - [&](auto const& rg_idx) { - // Get the file offset of this row group - auto const row_group_file_offset = [&]() { - auto const& rg = per_file_metadata.front().row_groups[rg_idx]; - if (rg.file_offset.has_value()) { - return rg.file_offset.value(); - } else if (rg.columns.front().file_offset != 0) { - return rg.columns.front().file_offset; - } else { - auto const& col_meta = rg.columns.front().meta_data; - return col_meta.dictionary_page_offset != 0 - ? std::min(col_meta.dictionary_page_offset, col_meta.data_page_offset) - : col_meta.data_page_offset; - } - }(); - - // Check if the row group starts within the byte range: row group file offset is >= - // bytes_to_skip AND (bytes_to_read is not specified OR the max byte offset overflows - // size_t OR row group file offset is < bytes_to_skip + bytes_to_read) - auto const is_within_byte_range = - std::cmp_greater_equal(row_group_file_offset, bytes_to_skip) and - (not bytes_to_read.has_value() or - (std::numeric_limits::max() - bytes_to_read.value() <= bytes_to_skip) or - std::cmp_less(row_group_file_offset, bytes_to_skip + bytes_to_read.value())); - - if (is_within_byte_range) { filtered_row_group_indices.front().emplace_back(rg_idx); } - }); + std::for_each(input_row_group_indices.front().begin(), + input_row_group_indices.front().end(), + [&](auto const& rg_idx) { + // Get the file offset of this row group + auto const row_group_file_offset = [&]() { + auto const& rg = per_file_metadata.front().row_groups[rg_idx]; + if (rg.file_offset.has_value()) { + return rg.file_offset.value(); + } else if (rg.columns.front().file_offset != 0) { + return rg.columns.front().file_offset; + } else { + auto const& col_meta = rg.columns.front().meta_data; + return col_meta.dictionary_page_offset != 0 + ? std::min(col_meta.dictionary_page_offset, + col_meta.data_page_offset) + : col_meta.data_page_offset; + } + }(); + + // Check if the row group starts within the byte range: row group file offset is + // >= bytes_to_skip AND (bytes_to_read is not specified OR the max byte offset + // overflows size_t OR row group file offset is < bytes_to_skip + bytes_to_read) + auto const is_within_byte_range = + std::cmp_greater_equal(row_group_file_offset, bytes_to_skip) and + (not bytes_to_read.has_value() or + cuda::add_overflow(bytes_to_skip, bytes_to_read.value()).overflow or + std::cmp_less(row_group_file_offset, bytes_to_skip + bytes_to_read.value())); + + if (is_within_byte_range) { + filtered_row_group_indices.front().emplace_back(rg_idx); + } + }); return filtered_row_group_indices; }