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

cpp/src/io/statistics/column_statistics.cuh

@@ -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.
  * 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/typed_statistics_chunk.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -112,6 +112,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 +124,9 @@ struct typed_statistics_chunk<T, true> {
   __device__ void reduce(T const& elem)
   {
     non_nulls++;
+    if constexpr (std::is_floating_point_v<T>) {
+      if (isnan(elem)) { has_nan = true; }
+    }
     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,9 @@ struct typed_statistics_chunk<T, false> {
   __device__ void reduce(T const& elem)
   {
     non_nulls++;
+    if constexpr (std::is_floating_point_v<T>) {
+      if (isnan(elem)) { has_nan = true; }
+    }
     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 +184,7 @@ struct typed_statistics_chunk<T, false> {
     }
     non_nulls += chunk.non_nulls;
     null_count += chunk.null_count;
+    has_nan |= chunk.has_nan;
   }
 };
 
@@ -209,6 +219,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 +248,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/tests/io/parquet_writer_test.cpp

@@ -853,6 +853,47 @@ 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 const nanf = std::numeric_limits<float>::quiet_NaN();
+  auto const 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
+
+  auto const expected = table_view{{col_f_nan, col_d_nan, col_f_allnan, col_f_nonan}};
+
+  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]);
+
+  // 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());
+}
+
 TEST_F(ParquetWriterTest, CheckColumnIndexTruncation)
 {
   std::array coldata{// in-range 7 bit.  should truncate to "yyyyyyyz"