diff --git a/cpp/src/io/functions.cpp b/cpp/src/io/functions.cpp index 2db331184395..55d1a750bc06 100644 --- a/cpp/src/io/functions.cpp +++ b/cpp/src/io/functions.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/io/parquet/experimental/deletion_vectors.cu b/cpp/src/io/parquet/experimental/deletion_vectors.cu index a04d1767a854..1d5ada49c69a 100644 --- a/cpp/src/io/parquet/experimental/deletion_vectors.cu +++ b/cpp/src/io/parquet/experimental/deletion_vectors.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/io/parquet/reader_impl_helpers.cpp b/cpp/src/io/parquet/reader_impl_helpers.cpp index a571d89ce4fc..667449fdb9ae 100644 --- a/cpp/src/io/parquet/reader_impl_helpers.cpp +++ b/cpp/src/io/parquet/reader_impl_helpers.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/io/statistics/column_statistics.cuh b/cpp/src/io/statistics/column_statistics.cuh index 81964cf2eabe..27d886f832d4 100644 --- a/cpp/src/io/statistics/column_statistics.cuh +++ b/cpp/src/io/statistics/column_statistics.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -202,6 +202,12 @@ struct merge_group_statistics_functor { chunk = block_reduce(chunk, storage); + // PARQUET-1246: if a float/double column contains any NaN, min/max must be omitted, + // else a reader doing NaN predicate pushdown skips the row group. spark-rapids#15004. + if constexpr (IO == detail::io_file_format::PARQUET) { + if (chunk.has_nan) { chunk.has_minmax = false; } + } + if (t == 0) { s.ck = get_untyped_chunk(chunk); } } diff --git a/cpp/src/io/statistics/statistics.cuh b/cpp/src/io/statistics/statistics.cuh index 57854d10f87b..1897cf272a2c 100644 --- a/cpp/src/io/statistics/statistics.cuh +++ b/cpp/src/io/statistics/statistics.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -88,6 +88,7 @@ struct statistics_chunk { statistics_val sum{}; //!< sum of chunk uint8_t has_minmax{}; //!< Nonzero if min_value and max_values are valid uint8_t has_sum{}; //!< Nonzero if sum is valid + uint8_t has_nan{}; //!< Nonzero if a NaN was seen (floating point only) }; struct statistics_group { diff --git a/cpp/src/io/statistics/statistics_type_identification.cuh b/cpp/src/io/statistics/statistics_type_identification.cuh index ed7812419afe..4324e9eb4ec8 100644 --- a/cpp/src/io/statistics/statistics_type_identification.cuh +++ b/cpp/src/io/statistics/statistics_type_identification.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -19,6 +19,9 @@ #include #include +#include +#include + #include namespace cudf { @@ -206,8 +209,8 @@ class aggregation_type { return val.size_bytes(); } else if constexpr (std::is_integral_v) { return val; - } else if constexpr (std::is_floating_point_v) { - return isnan(val) ? 0 : val; + } else if constexpr (cuda::std::is_floating_point_v) { + return cuda::std::isnan(val) ? 0 : val; } else if constexpr (cudf::is_fixed_point()) { return val.value(); } else if constexpr (cudf::is_duration()) { diff --git a/cpp/src/io/statistics/typed_statistics_chunk.cuh b/cpp/src/io/statistics/typed_statistics_chunk.cuh index a6125c42a78b..dd0748486ffe 100644 --- a/cpp/src/io/statistics/typed_statistics_chunk.cuh +++ b/cpp/src/io/statistics/typed_statistics_chunk.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -20,8 +20,10 @@ #include #include +#include #include #include +#include #include namespace cudf { @@ -112,6 +114,7 @@ struct typed_statistics_chunk { uint8_t has_minmax{false}; //!< Nonzero if min_value and max_values are valid uint8_t has_sum{false}; //!< Nonzero if sum is valid + uint8_t has_nan{false}; //!< Nonzero if a NaN was seen (floating point only) __device__ typed_statistics_chunk() : minimum_value(detail::minimum_identity()), @@ -123,6 +126,7 @@ struct typed_statistics_chunk { __device__ void reduce(T const& elem) { non_nulls++; + if constexpr (cuda::std::is_floating_point_v) { has_nan |= cuda::std::isnan(elem); } minimum_value = cuda::std::min(minimum_value, detail::extrema_type::convert(elem)); maximum_value = cuda::std::max(maximum_value, detail::extrema_type::convert(elem)); aggregate += detail::aggregation_type::convert(elem); @@ -138,6 +142,7 @@ struct typed_statistics_chunk { if (chunk.has_sum) { aggregate += union_member::get(chunk.sum); } non_nulls += chunk.non_nulls; null_count += chunk.null_count; + has_nan |= chunk.has_nan; } }; @@ -153,6 +158,7 @@ struct typed_statistics_chunk { uint8_t has_minmax{false}; //!< Nonzero if min_value and max_values are valid uint8_t has_sum{false}; //!< Nonzero if sum is valid + uint8_t has_nan{false}; //!< Nonzero if a NaN was seen (floating point only) __device__ typed_statistics_chunk() : minimum_value(detail::minimum_identity()), maximum_value(detail::maximum_identity()) @@ -162,6 +168,7 @@ struct typed_statistics_chunk { __device__ void reduce(T const& elem) { non_nulls++; + if constexpr (cuda::std::is_floating_point_v) { has_nan |= cuda::std::isnan(elem); } minimum_value = cuda::std::min(minimum_value, detail::extrema_type::convert(elem)); maximum_value = cuda::std::max(maximum_value, detail::extrema_type::convert(elem)); has_minmax = true; @@ -175,6 +182,7 @@ struct typed_statistics_chunk { } non_nulls += chunk.non_nulls; null_count += chunk.null_count; + has_nan |= chunk.has_nan; } }; @@ -209,6 +217,7 @@ __inline__ __device__ typed_statistics_chunk block_reduce( count_reduce(storage.template get()).Sum(output_chunk.null_count); __syncthreads(); output_chunk.has_minmax = __syncthreads_or(output_chunk.has_minmax); + output_chunk.has_nan = __syncthreads_or(output_chunk.has_nan); // FIXME : Is another syncthreads needed here? if constexpr (include_aggregate) { @@ -237,6 +246,7 @@ get_untyped_chunk(typed_statistics_chunk const& chunk) stat.non_nulls = chunk.non_nulls; stat.null_count = chunk.null_count; stat.has_minmax = chunk.has_minmax; + stat.has_nan = chunk.has_nan; stat.has_sum = [&]() { // invalidate the sum if overflow or underflow is possible if constexpr (std::is_floating_point_v or std::is_integral_v) { diff --git a/cpp/src/utilities/host_memory.cpp b/cpp/src/utilities/host_memory.cpp index 76bfca2ebcad..3d6e56471c50 100644 --- a/cpp/src/utilities/host_memory.cpp +++ b/cpp/src/utilities/host_memory.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -11,6 +11,7 @@ #include #include +#include #include #include #include diff --git a/cpp/tests/io/io_test_utils.hpp b/cpp/tests/io/io_test_utils.hpp index c51246353e5a..9fb6c64bb108 100644 --- a/cpp/tests/io/io_test_utils.hpp +++ b/cpp/tests/io/io_test_utils.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/tests/io/parquet_chunked_reader_test.cu b/cpp/tests/io/parquet_chunked_reader_test.cu index 3f8b8708d4c4..e36d14fd531c 100644 --- a/cpp/tests/io/parquet_chunked_reader_test.cu +++ b/cpp/tests/io/parquet_chunked_reader_test.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/tests/io/parquet_writer_test.cpp b/cpp/tests/io/parquet_writer_test.cpp index fd285daabe60..bcd79595049c 100644 --- a/cpp/tests/io/parquet_writer_test.cpp +++ b/cpp/tests/io/parquet_writer_test.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -853,6 +853,107 @@ TEST_F(ParquetWriterTest, Decimal128Stats) EXPECT_EQ(expected_max, stats.max_value); } +TEST_F(ParquetWriterTest, FloatingPointWithNaNStatsOmitted) +{ + // PARQUET-1246: a float/double column containing a NaN must not expose min/max, or a + // reader doing `= NaN` predicate pushdown skips the row group. NVIDIA/spark-rapids#15004. + auto constexpr nanf = std::numeric_limits::quiet_NaN(); + auto constexpr nand = std::numeric_limits::quiet_NaN(); + + column_wrapper col_f_nan{{1.0f, nanf, 3.0f, 2.0f}}; // NaN mixed with non-NaN + column_wrapper col_d_nan{{1.0, 2.0, nand, 4.0}}; // double variant + column_wrapper col_f_allnan{{nanf, nanf, nanf, nanf}}; // all NaN + column_wrapper col_f_nonan{{1.0f, 2.0f, 3.0f, 4.0f}}; // control: no NaN + column_wrapper col_f_nan_null{{1.0f, nanf, 3.0f, 5.0f}, + {true, true, true, false}}; // NaN alongside a null + + auto const expected = + table_view{{col_f_nan, col_d_nan, col_f_allnan, col_f_nonan, col_f_nan_null}}; + + auto const filepath = temp_env->get_temp_filepath("FloatingPointWithNaNStats.parquet"); + cudf::io::parquet_writer_options const out_opts = + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, expected); + cudf::io::write_parquet(out_opts); + + auto const source = cudf::io::datasource::create(filepath); + cudf::io::parquet::FileMetaData fmd; + read_footer(source, &fmd); + + auto const stats_f_nan = get_statistics(fmd.row_groups[0].columns[0]); + auto const stats_d_nan = get_statistics(fmd.row_groups[0].columns[1]); + auto const stats_f_allnan = get_statistics(fmd.row_groups[0].columns[2]); + auto const stats_f_nonan = get_statistics(fmd.row_groups[0].columns[3]); + auto const stats_f_nan_null = get_statistics(fmd.row_groups[0].columns[4]); + + // any column containing a NaN must not expose min/max + EXPECT_FALSE(stats_f_nan.min_value.has_value()); + EXPECT_FALSE(stats_f_nan.max_value.has_value()); + EXPECT_FALSE(stats_d_nan.min_value.has_value()); + EXPECT_FALSE(stats_d_nan.max_value.has_value()); + EXPECT_FALSE(stats_f_allnan.min_value.has_value()); + EXPECT_FALSE(stats_f_allnan.max_value.has_value()); + + // a column with no NaN is unaffected and still carries min/max + EXPECT_TRUE(stats_f_nonan.min_value.has_value()); + EXPECT_TRUE(stats_f_nonan.max_value.has_value()); + + // a null alongside the NaN does not interfere with NaN detection + EXPECT_FALSE(stats_f_nan_null.min_value.has_value()); + EXPECT_FALSE(stats_f_nan_null.max_value.has_value()); +} + +TEST_F(ParquetWriterTest, FloatingPointWithNaNStatsOmittedAcrossFragments) +{ + // A NaN in any page fragment must propagate through the fragment -> column-chunk statistics + // merge, so a multi-fragment column chunk with a single NaN still omits min/max. + // NVIDIA/spark-rapids#15004. + auto constexpr nanf = std::numeric_limits::quiet_NaN(); + auto constexpr num_rows = 20000; // > default 5000-row page fragment -> multiple fragments merged + std::vector data(num_rows); + for (int i = 0; i < num_rows; ++i) { + data[i] = static_cast(i); + } + data[num_rows / 2] = nanf; // a single NaN in a middle fragment + column_wrapper col(data.begin(), data.end()); + auto const expected = table_view{{col}}; + + auto const filepath = temp_env->get_temp_filepath("FloatingPointNaNStatsFragments.parquet"); + cudf::io::parquet_writer_options const out_opts = + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, expected); + cudf::io::write_parquet(out_opts); + + auto const source = cudf::io::datasource::create(filepath); + cudf::io::parquet::FileMetaData fmd; + read_footer(source, &fmd); + + ASSERT_EQ(fmd.row_groups.size(), 1); + auto const stats = get_statistics(fmd.row_groups[0].columns[0]); + EXPECT_FALSE(stats.min_value.has_value()); + EXPECT_FALSE(stats.max_value.has_value()); +} + +TEST_F(ParquetWriterTest, FloatingPointWithNaNStatsOmittedNested) +{ + // NaN detection must reach a float leaf nested in a LIST column (rapidsai/cudf#22817). + auto constexpr nanf = std::numeric_limits::quiet_NaN(); + cudf::test::lists_column_wrapper list_col{{1.0f, nanf, 3.0f}, {4.0f, 5.0f}}; + auto const expected = table_view{{list_col}}; + + auto const filepath = temp_env->get_temp_filepath("FloatingPointNaNStatsNested.parquet"); + cudf::io::parquet_writer_options const out_opts = + cudf::io::parquet_writer_options::builder(cudf::io::sink_info{filepath}, expected); + cudf::io::write_parquet(out_opts); + + auto const source = cudf::io::datasource::create(filepath); + cudf::io::parquet::FileMetaData fmd; + read_footer(source, &fmd); + + // the leaf float column (list element) contains a NaN -> min/max omitted + auto const stats = get_statistics(fmd.row_groups[0].columns[0]); + EXPECT_FALSE(stats.min_value.has_value()); + EXPECT_FALSE(stats.max_value.has_value()); +} + TEST_F(ParquetWriterTest, CheckColumnIndexTruncation) { std::array coldata{// in-range 7 bit. should truncate to "yyyyyyyz"