Omit Parquet min/max statistics for float/double columns containing NaN

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
  */
 

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
  */
 

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
  */
 

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); }
   }
 

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 {

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 <cudf/wrappers/durations.hpp>
 #include <cudf/wrappers/timestamps.hpp>
 
+#include <cuda/std/cmath>
+#include <cuda/std/type_traits>
+
 #include <tuple>
 
 namespace cudf {
@@ -206,8 +209,8 @@ class aggregation_type {
       return val.size_bytes();
     } else if constexpr (std::is_integral_v<T>) {
       return val;
-    } else if constexpr (std::is_floating_point_v<T>) {
-      return isnan(val) ? 0 : val;
+    } else if constexpr (cuda::std::is_floating_point_v<T>) {
+      return cuda::std::isnan(val) ? 0 : val;
     } else if constexpr (cudf::is_fixed_point<T>()) {
       return val.value();
     } else if constexpr (cudf::is_duration<T>()) {

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 <cudf/wrappers/timestamps.hpp>
 
 #include <cuda/functional>
+#include <cuda/std/cmath>
 #include <cuda/std/functional>
 #include <cuda/std/limits>
+#include <cuda/std/type_traits>
 #include <math_constants.h>
 
 namespace cudf {
@@ -112,6 +114,7 @@ struct typed_statistics_chunk<T, true> {
 
   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<E>()),
@@ -123,6 +126,7 @@ struct typed_statistics_chunk<T, true> {
   __device__ void reduce(T const& elem)
   {
     non_nulls++;
+    if constexpr (cuda::std::is_floating_point_v<T>) { has_nan |= cuda::std::isnan(elem); }
     minimum_value = cuda::std::min<E>(minimum_value, detail::extrema_type<T>::convert(elem));
     maximum_value = cuda::std::max<E>(maximum_value, detail::extrema_type<T>::convert(elem));
     aggregate += detail::aggregation_type<T>::convert(elem);
@@ -138,6 +142,7 @@ struct typed_statistics_chunk<T, true> {
     if (chunk.has_sum) { aggregate += union_member::get<A>(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<T, false> {
 
   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<E>()), maximum_value(detail::maximum_identity<E>())
@@ -162,6 +168,7 @@ struct typed_statistics_chunk<T, false> {
   __device__ void reduce(T const& elem)
   {
     non_nulls++;
+    if constexpr (cuda::std::is_floating_point_v<T>) { has_nan |= cuda::std::isnan(elem); }
     minimum_value = cuda::std::min<E>(minimum_value, detail::extrema_type<T>::convert(elem));
     maximum_value = cuda::std::max<E>(maximum_value, detail::extrema_type<T>::convert(elem));
     has_minmax    = true;
@@ -175,6 +182,7 @@ struct typed_statistics_chunk<T, false> {
     }
     non_nulls += chunk.non_nulls;
     null_count += chunk.null_count;
+    has_nan |= chunk.has_nan;
   }
 };
 
@@ -209,6 +217,7 @@ __inline__ __device__ typed_statistics_chunk<T, include_aggregate> block_reduce(
     count_reduce(storage.template get<uint32_t>()).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<T, include_aggregate> 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<E> or std::is_integral_v<E>) {

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 <cudf/utilities/pinned_memory.hpp>
 
 #include <rmm/cuda_device.hpp>
+#include <rmm/detail/aligned.hpp>
 #include <rmm/mr/pinned_host_memory_resource.hpp>
 #include <rmm/mr/pool_memory_resource.hpp>
 #include <rmm/resource_ref.hpp>

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
  */
 

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
  */
 

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<float>::quiet_NaN();
+  auto constexpr nand = std::numeric_limits<double>::quiet_NaN();
+
+  column_wrapper<float> col_f_nan{{1.0f, nanf, 3.0f, 2.0f}};     // NaN mixed with non-NaN
+  column_wrapper<double> col_d_nan{{1.0, 2.0, nand, 4.0}};       // double variant
+  column_wrapper<float> col_f_allnan{{nanf, nanf, nanf, nanf}};  // all NaN
+  column_wrapper<float> col_f_nonan{{1.0f, 2.0f, 3.0f, 4.0f}};   // control: no NaN
+  column_wrapper<float> 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<float>::quiet_NaN();
+  auto constexpr num_rows = 20000;  // > default 5000-row page fragment -> multiple fragments merged
+  std::vector<float> data(num_rows);
+  for (int i = 0; i < num_rows; ++i) {
+    data[i] = static_cast<float>(i);
+  }
+  data[num_rows / 2] = nanf;  // a single NaN in a middle fragment
+  column_wrapper<float> 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<float>::quiet_NaN();
+  cudf::test::lists_column_wrapper<float> 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"