From fb0a7c8236276879d9da56b184bc21f4b653d93c Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Wed, 22 Jul 2026 00:23:47 +0000 Subject: [PATCH 01/10] Encode non-negative nanos for negative ORC timestamps --- cpp/src/io/orc/stripe_enc.cu | 23 ++++++++++--- cpp/tests/io/orc_test.cpp | 67 ++++++++++++++++++++++++++++++++++++ 2 files changed, 86 insertions(+), 4 deletions(-) diff --git a/cpp/src/io/orc/stripe_enc.cu b/cpp/src/io/orc/stripe_enc.cu index 61a96c84f08f..d5cecb6191c6 100644 --- a/cpp/src/io/orc/stripe_enc.cu +++ b/cpp/src/io/orc/stripe_enc.cu @@ -800,10 +800,25 @@ CUDF_KERNEL void __launch_bounds__(block_size) case BOOLEAN: case BYTE: s->vals.u8[nz_idx] = column.element(row); break; case TIMESTAMP: { - int64_t ts = column.element(row); - int32_t ts_scale = cudf::detail::powers_of_ten[9 - min(s->chunk.scale, 9)]; - int64_t seconds = ts / ts_scale; - int64_t nanos = (ts - seconds * ts_scale); + auto const ts = column.element(row); + auto const ts_scale = cudf::detail::powers_of_ten[9 - min(s->chunk.scale, 9)]; + auto seconds = ts / ts_scale; + auto nanos = (ts - seconds * ts_scale); + // Integer division truncates toward zero, so a negative timestamp with a fractional + // part produces a negative remainder here. The ORC SECONDARY stream is unsigned and the + // Apache ORC reference implementation expects a non-negative nanos value, so normalize + // the remainder to be non-negative. See https://github.com/rapidsai/cudf/issues/19350. + if (nanos < 0) { + nanos += ts_scale; + // On read, Apache ORC (and the libcudf reader) subtracts one second for a negative + // timestamp only when the stored nanos are >= 1 ms; see the borrow logic in the + // reader and Apache ORC-306/ORC-763. For a sub-millisecond remainder the reader does + // not borrow, so the writer must carry that borrow itself to keep the value (and + // interop with Apache readers) correct. + if (nanos * cudf::detail::powers_of_ten[min(s->chunk.scale, 9)] < 1'000'000) { + seconds -= 1; + } + } s->vals.i64[nz_idx] = seconds - orc_utc_epoch; if (nanos != 0) { // Trailing zeroes are encoded in the lower 3-bits diff --git a/cpp/tests/io/orc_test.cpp b/cpp/tests/io/orc_test.cpp index 6159facde017..b65e47fcff4d 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -619,6 +619,73 @@ TEST_F(OrcWriterTest, negTimestampsNano) CUDF_TEST_EXPECT_TABLES_EQUAL(expected, result.tbl->view()); } +// Regression test for https://github.com/rapidsai/cudf/issues/19350. +// A negative timestamp with a sub-second fractional part produces a negative remainder when the +// writer splits it into (seconds, nanos). ORC's SECONDARY stream is unsigned and the Apache ORC +// reference implementation expects a non-negative nanos value; the writer must normalize the +// remainder to be non-negative. Because the reader (and Apache ORC, per ORC-306/ORC-763) borrows +// one second on read only when the stored nanos are >= 1 ms, the writer must additionally carry the +// borrow itself for sub-millisecond remainders. This test exercises both the >= 1 ms and the +// sub-millisecond negative cases across resolutions and verifies the libcudf round-trip is +// lossless. Note: the definitive interop check (reading the generated file with Apache ORC / Spark +// and confirming it no longer throws "nanos > 999999999 or < 0") must be performed with an external +// reader. +template +void test_negative_fractional_timestamp_roundtrip(std::vector const& values) +{ + cudf::test::fixed_width_column_wrapper const ts(values.begin(), values.end()); + cudf::table_view const expected({ts}); + + std::vector out_buffer; + cudf::io::orc_writer_options const out_opts = + cudf::io::orc_writer_options::builder(cudf::io::sink_info{&out_buffer}, expected); + cudf::io::write_orc(out_opts); + + cudf::io::orc_reader_options const in_opts = + cudf::io::orc_reader_options::builder(cudf::io::source_info{cudf::host_span{ + reinterpret_cast(out_buffer.data()), + out_buffer.size()}}) + .use_index(false) + .timestamp_type(cudf::data_type{cudf::type_to_id()}); + auto const result = cudf::io::read_orc(in_opts); + + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + expected.column(0), result.tbl->view().column(0), cudf::test::debug_output_level::ALL_ERRORS); +} + +TEST_F(OrcWriterTest, NegativeFractionalTimestampsInterop) +{ + // Microseconds: mix of positive, negative with >= 1 ms fraction, and negative sub-millisecond + // fractions (e.g. -5'999'500 us == -5.9995 s, fraction -0.5 ms). The sub-ms cases are misread by + // one second unless the writer carries the borrow. + test_negative_fractional_timestamp_roundtrip({ + 45'045'557'685'074'778L, // positive, unaffected by the fix + 116'614'807'755'579'786L, // positive, unaffected by the fix + 942'496L, // small positive, sub-second + -7'713'116'127L, // negative, fraction >= 1 ms + -33'426'545'118'057'504L, // negative, fraction >= 1 ms + -54'218'791'351'223'251L, // negative, fraction >= 1 ms + -5'999'500L, // negative, sub-millisecond fraction + -5'999'999L, // negative, sub-millisecond fraction + -100'000'500L, // negative, sub-millisecond fraction + }); + + // Nanoseconds: values < -1 s with sub-millisecond fractions. + test_negative_fractional_timestamp_roundtrip({ + -5'999'500'000L, + -5'999'999'999L, + -9'000'000'500L, + -131'968'727'238'000'000L, + }); + + // Milliseconds: negative fractional values (always >= 1 ms granularity). + test_negative_fractional_timestamp_roundtrip({ + -5'999L, + -123'456L, + -1'000L, + }); +} + TEST_F(OrcWriterTest, Slice) { int32_col col{{1, 2, 3, 4, 5}, cudf::test::iterators::null_at(3)}; From 908f7dfac23afd3511eb36a6cdab827684d865d4 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Wed, 12 Aug 2026 02:20:01 +0000 Subject: [PATCH 02/10] Match the Apache ORC timestamp encoding exactly Express the writer as a direct transcription of the Apache ORC rule and pin the near-epoch behavior that the format cannot represent in a test. --- cpp/src/io/orc/stripe_enc.cu | 40 +++++++------ cpp/tests/io/orc_test.cpp | 108 ++++++++++++++++++++++------------- 2 files changed, 89 insertions(+), 59 deletions(-) diff --git a/cpp/src/io/orc/stripe_enc.cu b/cpp/src/io/orc/stripe_enc.cu index d5cecb6191c6..a9dc5b9afeb7 100644 --- a/cpp/src/io/orc/stripe_enc.cu +++ b/cpp/src/io/orc/stripe_enc.cu @@ -800,30 +800,34 @@ CUDF_KERNEL void __launch_bounds__(block_size) case BOOLEAN: case BYTE: s->vals.u8[nz_idx] = column.element(row); break; case TIMESTAMP: { - auto const ts = column.element(row); - auto const ts_scale = cudf::detail::powers_of_ten[9 - min(s->chunk.scale, 9)]; - auto seconds = ts / ts_scale; - auto nanos = (ts - seconds * ts_scale); - // Integer division truncates toward zero, so a negative timestamp with a fractional - // part produces a negative remainder here. The ORC SECONDARY stream is unsigned and the - // Apache ORC reference implementation expects a non-negative nanos value, so normalize - // the remainder to be non-negative. See https://github.com/rapidsai/cudf/issues/19350. + auto const ts = column.element(row); + auto const ticks_per_sec = cudf::detail::powers_of_ten[9 - min(s->chunk.scale, 9)]; + auto const nanos_per_tick = cudf::detail::powers_of_ten[min(s->chunk.scale, 9)]; + + // ORC splits a timestamp into whole seconds plus a nanosecond remainder that must be in + // [0, 1e9); the unsigned SECONDARY stream cannot represent a negative remainder, and + // Apache ORC readers reject one. Integer division truncates toward zero, so use the + // floor of the timestamp instead to keep the remainder non-negative. + // See https://github.com/rapidsai/cudf/issues/19350. + auto seconds = ts / ticks_per_sec; + auto nanos = (ts - seconds * ticks_per_sec) * nanos_per_tick; if (nanos < 0) { - nanos += ts_scale; - // On read, Apache ORC (and the libcudf reader) subtracts one second for a negative - // timestamp only when the stored nanos are >= 1 ms; see the borrow logic in the - // reader and Apache ORC-306/ORC-763. For a sub-millisecond remainder the reader does - // not borrow, so the writer must carry that borrow itself to keep the value (and - // interop with Apache readers) correct. - if (nanos * cudf::detail::powers_of_ten[min(s->chunk.scale, 9)] < 1'000'000) { - seconds -= 1; - } + seconds -= 1; + nanos += 1'000'000'000; } + // Apache ORC readers borrow a second from the seconds stream when the stored seconds + // are negative and the stored nanos are at least 1 ms (see ORC-306/ORC-763 and the + // matching logic in the reader), so the writer has to give that second back here. This + // makes the encoding identical to the Apache ORC writer, which also means that + // timestamps within 999 ms before the epoch end up stored with zero seconds; readers + // cannot tell those apart from the same nanos one second later. Apache ORC has the same + // limitation and asserts the resulting one second shift in its own tests (ORC-771). + if (seconds < 0 and nanos > 999'999) { seconds += 1; } + s->vals.i64[nz_idx] = seconds - orc_utc_epoch; if (nanos != 0) { // Trailing zeroes are encoded in the lower 3-bits uint32_t zeroes = 0; - nanos *= cudf::detail::powers_of_ten[min(s->chunk.scale, 9)]; if (!(nanos % 100)) { nanos /= 100; zeroes = 1; diff --git a/cpp/tests/io/orc_test.cpp b/cpp/tests/io/orc_test.cpp index b65e47fcff4d..7e8d73a32d1a 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -619,71 +619,97 @@ TEST_F(OrcWriterTest, negTimestampsNano) CUDF_TEST_EXPECT_TABLES_EQUAL(expected, result.tbl->view()); } -// Regression test for https://github.com/rapidsai/cudf/issues/19350. -// A negative timestamp with a sub-second fractional part produces a negative remainder when the -// writer splits it into (seconds, nanos). ORC's SECONDARY stream is unsigned and the Apache ORC -// reference implementation expects a non-negative nanos value; the writer must normalize the -// remainder to be non-negative. Because the reader (and Apache ORC, per ORC-306/ORC-763) borrows -// one second on read only when the stored nanos are >= 1 ms, the writer must additionally carry the -// borrow itself for sub-millisecond remainders. This test exercises both the >= 1 ms and the -// sub-millisecond negative cases across resolutions and verifies the libcudf round-trip is -// lossless. Note: the definitive interop check (reading the generated file with Apache ORC / Spark -// and confirming it no longer throws "nanos > 999999999 or < 0") must be performed with an external -// reader. +// Writes `values` as a timestamp column of type `T`, reads it back at the same resolution and +// compares the result against `expected_values`. template -void test_negative_fractional_timestamp_roundtrip(std::vector const& values) +void test_timestamp_roundtrip(std::vector const& values, + std::vector const& expected_values) { - cudf::test::fixed_width_column_wrapper const ts(values.begin(), values.end()); - cudf::table_view const expected({ts}); + cudf::test::fixed_width_column_wrapper const input(values.begin(), + values.end()); + cudf::test::fixed_width_column_wrapper const expected(expected_values.begin(), + expected_values.end()); + cudf::table_view const input_table({input}); std::vector out_buffer; cudf::io::orc_writer_options const out_opts = - cudf::io::orc_writer_options::builder(cudf::io::sink_info{&out_buffer}, expected); + cudf::io::orc_writer_options::builder(cudf::io::sink_info{&out_buffer}, input_table); cudf::io::write_orc(out_opts); cudf::io::orc_reader_options const in_opts = - cudf::io::orc_reader_options::builder(cudf::io::source_info{cudf::host_span{ - reinterpret_cast(out_buffer.data()), - out_buffer.size()}}) + cudf::io::orc_reader_options::builder( + cudf::io::source_info{cudf::host_span{ + reinterpret_cast(out_buffer.data()), out_buffer.size()}}) .use_index(false) .timestamp_type(cudf::data_type{cudf::type_to_id()}); auto const result = cudf::io::read_orc(in_opts); CUDF_TEST_EXPECT_COLUMNS_EQUAL( - expected.column(0), result.tbl->view().column(0), cudf::test::debug_output_level::ALL_ERRORS); + expected, result.tbl->view().column(0), cudf::test::debug_output_level::ALL_ERRORS); } -TEST_F(OrcWriterTest, NegativeFractionalTimestampsInterop) -{ - // Microseconds: mix of positive, negative with >= 1 ms fraction, and negative sub-millisecond - // fractions (e.g. -5'999'500 us == -5.9995 s, fraction -0.5 ms). The sub-ms cases are misread by - // one second unless the writer carries the borrow. - test_negative_fractional_timestamp_roundtrip({ - 45'045'557'685'074'778L, // positive, unaffected by the fix - 116'614'807'755'579'786L, // positive, unaffected by the fix - 942'496L, // small positive, sub-second - -7'713'116'127L, // negative, fraction >= 1 ms - -33'426'545'118'057'504L, // negative, fraction >= 1 ms - -54'218'791'351'223'251L, // negative, fraction >= 1 ms - -5'999'500L, // negative, sub-millisecond fraction - -5'999'999L, // negative, sub-millisecond fraction - -100'000'500L, // negative, sub-millisecond fraction - }); +// Regression test for https://github.com/rapidsai/cudf/issues/19350. +// A negative timestamp with a sub-second fractional part used to be split into a negative nanos +// remainder, which the unsigned SECONDARY stream stored as a huge value. Apache ORC readers reject +// such files with "nanos > 999999999 or < 0". The writer now emits the same (seconds, nanos) pair +// as the Apache ORC writer, which round-trips losslessly for every timestamp outside of one second +// before the epoch. Note that the definitive interop check (reading the generated file with Apache +// ORC / Spark) has to be performed with an external reader. +TEST_F(OrcWriterTest, NegativeFractionalTimestamps) +{ + auto const timestamps_us = std::vector{ + 45'045'557'685'074'778L, // positive + 116'614'807'755'579'786L, // positive + 942'496L, // positive, sub-second + -7'713'116'127L, // fraction >= 1 ms + -33'426'545'118'057'504L, // fraction >= 1 ms + -54'218'791'351'223'251L, // fraction >= 1 ms + -5'999'500L, // fraction < 1 ms + -5'999'999L, // fraction < 1 ms + -100'000'500L, // fraction < 1 ms + }; + test_timestamp_roundtrip(timestamps_us, timestamps_us); - // Nanoseconds: values < -1 s with sub-millisecond fractions. - test_negative_fractional_timestamp_roundtrip({ + auto const timestamps_ns = std::vector{ -5'999'500'000L, -5'999'999'999L, -9'000'000'500L, -131'968'727'238'000'000L, - }); + }; + test_timestamp_roundtrip(timestamps_ns, timestamps_ns); - // Milliseconds: negative fractional values (always >= 1 ms granularity). - test_negative_fractional_timestamp_roundtrip({ + auto const timestamps_ms = std::vector{ -5'999L, -123'456L, -1'000L, - }); + }; + test_timestamp_roundtrip(timestamps_ms, timestamps_ms); +} + +// ORC cannot represent the timestamps in [-999 ms, -1 ns]: they are stored with zero seconds and a +// nanos value that readers cannot tell apart from the same nanos one second later, so they are read +// back one second too late. Apache ORC has the same limitation and asserts this exact behavior in +// its own tests (ORC-763, ORC-771), so the values below are what any ORC implementation produces +// for this input. Timestamps of -999'000'001 ns and earlier are stored with negative seconds and +// are not affected. +TEST_F(OrcWriterTest, NegativeTimestampsNearEpoch) +{ + test_timestamp_roundtrip( + {-1L, + -500L, + -500'000L, + -999'000L, // last value shifted by one second + -999'001L, // first value stored losslessly + -999'999L, + -1'000'001L, + -5'999'500L}, + {999'999L, 999'500L, 500'000L, 1'000L, -999'001L, -999'999L, -1'000'001L, -5'999'500L}); + + test_timestamp_roundtrip( + {-1L, -999'000'000L, -999'000'001L, -1'000'000'000L}, + {999'999'999L, 1'000'000L, -999'000'001L, -1'000'000'000L}); + + test_timestamp_roundtrip({-1L, -999L, -1'000L}, {999L, 1L, -1'000L}); } TEST_F(OrcWriterTest, Slice) From e9a8bba2da2d1090fa7efe5c7221d10bc0fdf019 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Thu, 13 Aug 2026 18:17:07 +0000 Subject: [PATCH 03/10] Document the near-epoch timestamp limitation Add the limitation to the write_orc and to_orc docs, clarify which test guards the encoding, and cover the null path in the near-epoch test. --- cpp/include/cudf/io/orc.hpp | 9 ++++++ cpp/src/io/orc/stripe_enc.cu | 6 ++-- cpp/tests/io/orc_test.cpp | 53 +++++++++++++++++-------------- python/cudf/cudf/utils/ioutils.py | 10 ++++++ 4 files changed, 52 insertions(+), 26 deletions(-) diff --git a/cpp/include/cudf/io/orc.hpp b/cpp/include/cudf/io/orc.hpp index 53d35e6c33b4..327170a015b6 100644 --- a/cpp/include/cudf/io/orc.hpp +++ b/cpp/include/cudf/io/orc.hpp @@ -1070,6 +1070,12 @@ class orc_writer_options_builder { * * @note If an exception is thrown during encoding or compression, no data is written to the sink. * + * @note Timestamps in the last 999 milliseconds before the UNIX epoch cannot be represented in the + * ORC format. Such a timestamp is stored with zero seconds and a nanosecond remainder that no + * reader can distinguish from the same remainder one second later, so it is read back one second + * later than it was written, and the column statistics describe the value as it was written rather + * than as it is read. The Apache ORC writer has the same behavior; see Apache ORC-763 and ORC-771. + * * @param options Settings for controlling reading behavior * @param stream CUDA stream used for device memory operations and kernel launches */ @@ -1487,6 +1493,9 @@ class chunked_orc_writer_options_builder { * ... * writer.close(); * @endcode + * + * @note Timestamps in the last 999 milliseconds before the UNIX epoch cannot be represented in the + * ORC format; see `write_orc()` for details. */ class orc_chunked_writer { public: diff --git a/cpp/src/io/orc/stripe_enc.cu b/cpp/src/io/orc/stripe_enc.cu index a9dc5b9afeb7..6c27d0461267 100644 --- a/cpp/src/io/orc/stripe_enc.cu +++ b/cpp/src/io/orc/stripe_enc.cu @@ -820,8 +820,10 @@ CUDF_KERNEL void __launch_bounds__(block_size) // matching logic in the reader), so the writer has to give that second back here. This // makes the encoding identical to the Apache ORC writer, which also means that // timestamps within 999 ms before the epoch end up stored with zero seconds; readers - // cannot tell those apart from the same nanos one second later. Apache ORC has the same - // limitation and asserts the resulting one second shift in its own tests (ORC-771). + // cannot tell those apart from the same nanos one second later, and the statistics + // (which are computed from the input values) describe them as written rather than as + // read. Apache ORC has the same limitation and asserts the resulting one second shift + // in its own tests (ORC-771). if (seconds < 0 and nanos > 999'999) { seconds += 1; } s->vals.i64[nz_idx] = seconds - orc_utc_epoch; diff --git a/cpp/tests/io/orc_test.cpp b/cpp/tests/io/orc_test.cpp index 7e8d73a32d1a..bfa7b890af5d 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -620,15 +620,19 @@ TEST_F(OrcWriterTest, negTimestampsNano) } // Writes `values` as a timestamp column of type `T`, reads it back at the same resolution and -// compares the result against `expected_values`. +// compares the result against `expected_values`. Makes every other value null when `with_nulls` is +// set. template void test_timestamp_roundtrip(std::vector const& values, - std::vector const& expected_values) -{ - cudf::test::fixed_width_column_wrapper const input(values.begin(), - values.end()); - cudf::test::fixed_width_column_wrapper const expected(expected_values.begin(), - expected_values.end()); + std::vector const& expected_values, + bool with_nulls = false) +{ + auto const validity = cudf::detail::make_counting_transform_iterator( + 0, [with_nulls](auto i) { return not with_nulls or i % 2 == 0; }); + cudf::test::fixed_width_column_wrapper const input( + values.begin(), values.end(), validity); + cudf::test::fixed_width_column_wrapper const expected( + expected_values.begin(), expected_values.end(), validity); cudf::table_view const input_table({input}); std::vector out_buffer; @@ -648,13 +652,16 @@ void test_timestamp_roundtrip(std::vector const& values, expected, result.tbl->view().column(0), cudf::test::debug_output_level::ALL_ERRORS); } -// Regression test for https://github.com/rapidsai/cudf/issues/19350. -// A negative timestamp with a sub-second fractional part used to be split into a negative nanos -// remainder, which the unsigned SECONDARY stream stored as a huge value. Apache ORC readers reject -// such files with "nanos > 999999999 or < 0". The writer now emits the same (seconds, nanos) pair -// as the Apache ORC writer, which round-trips losslessly for every timestamp outside of one second -// before the epoch. Note that the definitive interop check (reading the generated file with Apache -// ORC / Spark) has to be performed with an external reader. +// Verifies that the timestamps that ORC can represent - everything except the last 999 ms before +// the epoch, see `NegativeTimestampsNearEpoch` - round-trip losslessly. +// +// This does not detect a return of the encoding that https://github.com/rapidsai/cudf/issues/19350 +// is about, where a negative fractional timestamp was split into a negative nanos remainder that +// the unsigned SECONDARY stream stored as a huge value: the libcudf reader sign-extends the stored +// value and inverts that encoding exactly, so these values round-trip either way. Only an Apache +// ORC reader rejects such a file ("nanos > 999999999 or < 0"), so the encoding itself has to be +// verified externally; `NegativeTimestampsNearEpoch` below is the in-repo guard, because its +// expected values only hold for the Apache encoding. TEST_F(OrcWriterTest, NegativeFractionalTimestamps) { auto const timestamps_us = std::vector{ @@ -694,16 +701,14 @@ TEST_F(OrcWriterTest, NegativeFractionalTimestamps) // are not affected. TEST_F(OrcWriterTest, NegativeTimestampsNearEpoch) { - test_timestamp_roundtrip( - {-1L, - -500L, - -500'000L, - -999'000L, // last value shifted by one second - -999'001L, // first value stored losslessly - -999'999L, - -1'000'001L, - -5'999'500L}, - {999'999L, 999'500L, 500'000L, 1'000L, -999'001L, -999'999L, -1'000'001L, -5'999'500L}); + // The two values around the -999 ms boundary are the smallest shifted and the largest lossless + // timestamp at this resolution. + auto const timestamps_us = std::vector{ + -1L, -500L, -500'000L, -999'000L, -999'001L, -999'999L, -1'000'001L, -5'999'500L}; + auto const read_back_us = std::vector{ + 999'999L, 999'500L, 500'000L, 1'000L, -999'001L, -999'999L, -1'000'001L, -5'999'500L}; + test_timestamp_roundtrip(timestamps_us, read_back_us); + test_timestamp_roundtrip(timestamps_us, read_back_us, /* with_nulls */ true); test_timestamp_roundtrip( {-1L, -999'000'000L, -999'000'001L, -1'000'000'000L}, diff --git a/python/cudf/cudf/utils/ioutils.py b/python/cudf/cudf/utils/ioutils.py index 0db154c3c1a5..3a61c462ebc9 100644 --- a/python/cudf/cudf/utils/ioutils.py +++ b/python/cudf/cudf/utils/ioutils.py @@ -579,6 +579,16 @@ doesn't require much space and is faster. Other indexes will be included as columns in the file output. +Notes +----- +Timestamps in the last 999 milliseconds before the UNIX epoch cannot be +represented in the ORC format. Such a timestamp is stored with zero seconds and +a nanosecond remainder that no reader can distinguish from the same remainder +one second later, so it is read back one second later than it was written, and +the column statistics describe the value as it was written rather than as it is +read. The Apache ORC writer has the same behavior; see Apache ORC-763 and +ORC-771. + See Also -------- cudf.read_orc From d2807dd0c92640cee0b5c452f676644a5bff66b1 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Thu, 13 Aug 2026 20:40:33 +0000 Subject: [PATCH 04/10] Shorten the near-epoch timestamp notes --- cpp/include/cudf/io/orc.hpp | 12 +++++------- cpp/src/io/orc/stripe_enc.cu | 5 +---- python/cudf/cudf/utils/ioutils.py | 11 ++++------- 3 files changed, 10 insertions(+), 18 deletions(-) diff --git a/cpp/include/cudf/io/orc.hpp b/cpp/include/cudf/io/orc.hpp index 327170a015b6..3b37a070647d 100644 --- a/cpp/include/cudf/io/orc.hpp +++ b/cpp/include/cudf/io/orc.hpp @@ -1070,11 +1070,9 @@ class orc_writer_options_builder { * * @note If an exception is thrown during encoding or compression, no data is written to the sink. * - * @note Timestamps in the last 999 milliseconds before the UNIX epoch cannot be represented in the - * ORC format. Such a timestamp is stored with zero seconds and a nanosecond remainder that no - * reader can distinguish from the same remainder one second later, so it is read back one second - * later than it was written, and the column statistics describe the value as it was written rather - * than as it is read. The Apache ORC writer has the same behavior; see Apache ORC-763 and ORC-771. + * @note Timestamps in the last 999 milliseconds before the UNIX epoch are read back one second + * later than they were written, because the ORC encoding cannot tell them apart from that later + * value. The Apache ORC writer behaves the same way; see Apache ORC-763 and ORC-771. * * @param options Settings for controlling reading behavior * @param stream CUDA stream used for device memory operations and kernel launches @@ -1494,8 +1492,8 @@ class chunked_orc_writer_options_builder { * writer.close(); * @endcode * - * @note Timestamps in the last 999 milliseconds before the UNIX epoch cannot be represented in the - * ORC format; see `write_orc()` for details. + * @note Timestamps in the last 999 milliseconds before the UNIX epoch are read back one second + * later than they were written; see `write_orc()` for details. */ class orc_chunked_writer { public: diff --git a/cpp/src/io/orc/stripe_enc.cu b/cpp/src/io/orc/stripe_enc.cu index 6c27d0461267..fbea3bb2986f 100644 --- a/cpp/src/io/orc/stripe_enc.cu +++ b/cpp/src/io/orc/stripe_enc.cu @@ -804,10 +804,7 @@ CUDF_KERNEL void __launch_bounds__(block_size) auto const ticks_per_sec = cudf::detail::powers_of_ten[9 - min(s->chunk.scale, 9)]; auto const nanos_per_tick = cudf::detail::powers_of_ten[min(s->chunk.scale, 9)]; - // ORC splits a timestamp into whole seconds plus a nanosecond remainder that must be in - // [0, 1e9); the unsigned SECONDARY stream cannot represent a negative remainder, and - // Apache ORC readers reject one. Integer division truncates toward zero, so use the - // floor of the timestamp instead to keep the remainder non-negative. + // Use the floor of the timestamp to keep the nanosecond remainder non-negative. // See https://github.com/rapidsai/cudf/issues/19350. auto seconds = ts / ticks_per_sec; auto nanos = (ts - seconds * ticks_per_sec) * nanos_per_tick; diff --git a/python/cudf/cudf/utils/ioutils.py b/python/cudf/cudf/utils/ioutils.py index 3a61c462ebc9..d5d992183921 100644 --- a/python/cudf/cudf/utils/ioutils.py +++ b/python/cudf/cudf/utils/ioutils.py @@ -581,13 +581,10 @@ Notes ----- -Timestamps in the last 999 milliseconds before the UNIX epoch cannot be -represented in the ORC format. Such a timestamp is stored with zero seconds and -a nanosecond remainder that no reader can distinguish from the same remainder -one second later, so it is read back one second later than it was written, and -the column statistics describe the value as it was written rather than as it is -read. The Apache ORC writer has the same behavior; see Apache ORC-763 and -ORC-771. +Timestamps in the last 999 milliseconds before the UNIX epoch are read back one +second later than they were written, because the ORC encoding cannot tell them +apart from that later value. The Apache ORC writer behaves the same way; see +Apache ORC-763 and ORC-771. See Also -------- From 8d660297ec670605e9d757d18bbcae26f871f686 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Thu, 13 Aug 2026 20:42:09 +0000 Subject: [PATCH 05/10] Shorten the writer's borrow comment --- cpp/src/io/orc/stripe_enc.cu | 13 ++++--------- 1 file changed, 4 insertions(+), 9 deletions(-) diff --git a/cpp/src/io/orc/stripe_enc.cu b/cpp/src/io/orc/stripe_enc.cu index fbea3bb2986f..a18ef1d7e3b0 100644 --- a/cpp/src/io/orc/stripe_enc.cu +++ b/cpp/src/io/orc/stripe_enc.cu @@ -812,15 +812,10 @@ CUDF_KERNEL void __launch_bounds__(block_size) seconds -= 1; nanos += 1'000'000'000; } - // Apache ORC readers borrow a second from the seconds stream when the stored seconds - // are negative and the stored nanos are at least 1 ms (see ORC-306/ORC-763 and the - // matching logic in the reader), so the writer has to give that second back here. This - // makes the encoding identical to the Apache ORC writer, which also means that - // timestamps within 999 ms before the epoch end up stored with zero seconds; readers - // cannot tell those apart from the same nanos one second later, and the statistics - // (which are computed from the input values) describe them as written rather than as - // read. Apache ORC has the same limitation and asserts the resulting one second shift - // in its own tests (ORC-771). + // Apache ORC readers borrow a second when the stored seconds are negative and the + // stored nanos are at least 1 ms (ORC-306/ORC-763, mirrored in the reader), so give + // that second back here. Timestamps in the last 999 ms before the epoch then store zero + // seconds and are read back a second later, as documented on `write_orc` (ORC-771). if (seconds < 0 and nanos > 999'999) { seconds += 1; } s->vals.i64[nz_idx] = seconds - orc_utc_epoch; From 3f35c40afcdf203b4716f8fd693d17c390749d82 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Thu, 13 Aug 2026 21:11:32 +0000 Subject: [PATCH 06/10] Shorten the near-epoch timestamp notes further --- cpp/include/cudf/io/orc.hpp | 9 ++++----- cpp/src/io/orc/stripe_enc.cu | 13 +++++++------ python/cudf/cudf/utils/ioutils.py | 7 +++---- 3 files changed, 14 insertions(+), 15 deletions(-) diff --git a/cpp/include/cudf/io/orc.hpp b/cpp/include/cudf/io/orc.hpp index 3b37a070647d..47eac12966a3 100644 --- a/cpp/include/cudf/io/orc.hpp +++ b/cpp/include/cudf/io/orc.hpp @@ -1070,9 +1070,8 @@ class orc_writer_options_builder { * * @note If an exception is thrown during encoding or compression, no data is written to the sink. * - * @note Timestamps in the last 999 milliseconds before the UNIX epoch are read back one second - * later than they were written, because the ORC encoding cannot tell them apart from that later - * value. The Apache ORC writer behaves the same way; see Apache ORC-763 and ORC-771. + * @note Timestamps in the last 999 milliseconds before the UNIX epoch are not representable in ORC; + * they are read back one second later, as with the Apache ORC writer (ORC-763, ORC-771). * * @param options Settings for controlling reading behavior * @param stream CUDA stream used for device memory operations and kernel launches @@ -1492,8 +1491,8 @@ class chunked_orc_writer_options_builder { * writer.close(); * @endcode * - * @note Timestamps in the last 999 milliseconds before the UNIX epoch are read back one second - * later than they were written; see `write_orc()` for details. + * @note Timestamps in the last 999 milliseconds before the UNIX epoch are not representable in ORC; + * see `write_orc()` for details. */ class orc_chunked_writer { public: diff --git a/cpp/src/io/orc/stripe_enc.cu b/cpp/src/io/orc/stripe_enc.cu index a18ef1d7e3b0..7b60d197b1ba 100644 --- a/cpp/src/io/orc/stripe_enc.cu +++ b/cpp/src/io/orc/stripe_enc.cu @@ -804,18 +804,19 @@ CUDF_KERNEL void __launch_bounds__(block_size) auto const ticks_per_sec = cudf::detail::powers_of_ten[9 - min(s->chunk.scale, 9)]; auto const nanos_per_tick = cudf::detail::powers_of_ten[min(s->chunk.scale, 9)]; - // Use the floor of the timestamp to keep the nanosecond remainder non-negative. - // See https://github.com/rapidsai/cudf/issues/19350. auto seconds = ts / ticks_per_sec; auto nanos = (ts - seconds * ticks_per_sec) * nanos_per_tick; + + // Adjust to keep the nanosecond remainder non-negative. + // See https://github.com/rapidsai/cudf/issues/19350. if (nanos < 0) { seconds -= 1; nanos += 1'000'000'000; } - // Apache ORC readers borrow a second when the stored seconds are negative and the - // stored nanos are at least 1 ms (ORC-306/ORC-763, mirrored in the reader), so give - // that second back here. Timestamps in the last 999 ms before the epoch then store zero - // seconds and are read back a second later, as documented on `write_orc` (ORC-771). + // ORC readers borrow a second when the stored seconds are negative and the stored nanos + // are at least 1 ms (ORC-306/ORC-763, mirrored in the reader), so give that second back + // here. Timestamps in the last 999 ms before the epoch then store zero seconds and are + // read back a second later, as documented on `write_orc` (ORC-771). if (seconds < 0 and nanos > 999'999) { seconds += 1; } s->vals.i64[nz_idx] = seconds - orc_utc_epoch; diff --git a/python/cudf/cudf/utils/ioutils.py b/python/cudf/cudf/utils/ioutils.py index d5d992183921..c6d8f0df0264 100644 --- a/python/cudf/cudf/utils/ioutils.py +++ b/python/cudf/cudf/utils/ioutils.py @@ -581,10 +581,9 @@ Notes ----- -Timestamps in the last 999 milliseconds before the UNIX epoch are read back one -second later than they were written, because the ORC encoding cannot tell them -apart from that later value. The Apache ORC writer behaves the same way; see -Apache ORC-763 and ORC-771. +Timestamps in the last 999 milliseconds before the UNIX epoch are not +representable in ORC; they are read back one second later, as with the Apache +ORC writer (ORC-763, ORC-771). See Also -------- From cc264206b00b8a8dc0addeab046e8cdc241b4e8e Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Thu, 13 Aug 2026 23:32:58 +0000 Subject: [PATCH 07/10] trim tests --- cpp/tests/io/orc_test.cpp | 41 +++++++++++---------------------------- 1 file changed, 11 insertions(+), 30 deletions(-) diff --git a/cpp/tests/io/orc_test.cpp b/cpp/tests/io/orc_test.cpp index 8ab29fe2e429..621e2f6659fd 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -653,9 +653,6 @@ TEST_F(OrcWriterTest, negTimestampsNano) CUDF_TEST_EXPECT_TABLES_EQUAL(expected, result.tbl->view()); } -// Writes `values` as a timestamp column of type `T`, reads it back at the same resolution and -// compares the result against `expected_values`. Makes every other value null when `with_nulls` is -// set. template void test_timestamp_roundtrip(std::vector const& values, std::vector const& expected_values, @@ -686,53 +683,37 @@ void test_timestamp_roundtrip(std::vector const& values, expected, result.tbl->view().column(0), cudf::test::debug_output_level::ALL_ERRORS); } -// Verifies that the timestamps that ORC can represent - everything except the last 999 ms before -// the epoch, see `NegativeTimestampsNearEpoch` - round-trip losslessly. -// -// This does not detect a return of the encoding that https://github.com/rapidsai/cudf/issues/19350 -// is about, where a negative fractional timestamp was split into a negative nanos remainder that -// the unsigned SECONDARY stream stored as a huge value: the libcudf reader sign-extends the stored -// value and inverts that encoding exactly, so these values round-trip either way. Only an Apache -// ORC reader rejects such a file ("nanos > 999999999 or < 0"), so the encoding itself has to be -// verified externally; `NegativeTimestampsNearEpoch` below is the in-repo guard, because its -// expected values only hold for the Apache encoding. +// Verifies that the timestamps ORC can represent - everything except the last 999 ms before the +// epoch, see `NegativeTimestampsNearEpoch` - round-trip losslessly +// (https://github.com/rapidsai/cudf/issues/19350). TEST_F(OrcWriterTest, NegativeFractionalTimestamps) { auto const timestamps_us = std::vector{ - 45'045'557'685'074'778L, // positive 116'614'807'755'579'786L, // positive 942'496L, // positive, sub-second - -7'713'116'127L, // fraction >= 1 ms - -33'426'545'118'057'504L, // fraction >= 1 ms -54'218'791'351'223'251L, // fraction >= 1 ms - -5'999'500L, // fraction < 1 ms + -7'713'116'127L, // fraction >= 1 ms, sub-second -5'999'999L, // fraction < 1 ms - -100'000'500L, // fraction < 1 ms }; test_timestamp_roundtrip(timestamps_us, timestamps_us); auto const timestamps_ns = std::vector{ - -5'999'500'000L, - -5'999'999'999L, - -9'000'000'500L, - -131'968'727'238'000'000L, + -131'968'727'238'000'000L, // fraction >= 1 ms + -5'999'999'999L, // fraction < 1 ms }; test_timestamp_roundtrip(timestamps_ns, timestamps_ns); + // Milliseconds have no fraction below 1 ms, so only the borrowing path applies. auto const timestamps_ms = std::vector{ - -5'999L, -123'456L, - -1'000L, + -5'999L, }; test_timestamp_roundtrip(timestamps_ms, timestamps_ms); } -// ORC cannot represent the timestamps in [-999 ms, -1 ns]: they are stored with zero seconds and a -// nanos value that readers cannot tell apart from the same nanos one second later, so they are read -// back one second too late. Apache ORC has the same limitation and asserts this exact behavior in -// its own tests (ORC-763, ORC-771), so the values below are what any ORC implementation produces -// for this input. Timestamps of -999'000'001 ns and earlier are stored with negative seconds and -// are not affected. +// Timestamps in the last 999 ms before the epoch are not representable in ORC; they are read back +// one second later, as with the Apache ORC writer, whose own tests assert the same values (ORC-763, +// ORC-771). TEST_F(OrcWriterTest, NegativeTimestampsNearEpoch) { // The two values around the -999 ms boundary are the smallest shifted and the largest lossless From 4f937e3a78ac2310dc37ae8ee136fd2136d3ca9c Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Fri, 14 Aug 2026 00:23:25 +0000 Subject: [PATCH 08/10] trim tests some more --- cpp/tests/io/orc_test.cpp | 21 ++++++++------------- 1 file changed, 8 insertions(+), 13 deletions(-) diff --git a/cpp/tests/io/orc_test.cpp b/cpp/tests/io/orc_test.cpp index 621e2f6659fd..3bc23a636ca2 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -683,30 +683,25 @@ void test_timestamp_roundtrip(std::vector const& values, expected, result.tbl->view().column(0), cudf::test::debug_output_level::ALL_ERRORS); } -// Verifies that the timestamps ORC can represent - everything except the last 999 ms before the -// epoch, see `NegativeTimestampsNearEpoch` - round-trip losslessly -// (https://github.com/rapidsai/cudf/issues/19350). TEST_F(OrcWriterTest, NegativeFractionalTimestamps) { + // ORC readers handle remainders of >= 1 ms above the lower second differently, so cover both auto const timestamps_us = std::vector{ - 116'614'807'755'579'786L, // positive - 942'496L, // positive, sub-second - -54'218'791'351'223'251L, // fraction >= 1 ms - -7'713'116'127L, // fraction >= 1 ms, sub-second - -5'999'999L, // fraction < 1 ms + -54'218'791'351'223'251L, // 776.749 ms above the lower second, so >= 1 ms + -5'999'999L, // 1 us above the lower second, so < 1 ms }; test_timestamp_roundtrip(timestamps_us, timestamps_us); auto const timestamps_ns = std::vector{ - -131'968'727'238'000'000L, // fraction >= 1 ms - -5'999'999'999L, // fraction < 1 ms + -131'968'727'238'000'000L, // 762 ms above the lower second, so >= 1 ms + -5'999'999'999L, // 1 ns above the lower second, so < 1 ms }; test_timestamp_roundtrip(timestamps_ns, timestamps_ns); - // Milliseconds have no fraction below 1 ms, so only the borrowing path applies. + // Millisecond timestamps cannot have a < 1 ms remainder auto const timestamps_ms = std::vector{ - -123'456L, - -5'999L, + -123'456L, // 544 ms above the lower second + -5'999L, // 1 ms above the lower second, the smallest remainder at this resolution }; test_timestamp_roundtrip(timestamps_ms, timestamps_ms); } From f7b57a6fc87e4c7952d1185298b4fe1f91fa6c05 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Fri, 14 Aug 2026 00:47:59 +0000 Subject: [PATCH 09/10] and a bit more --- cpp/tests/io/orc_test.cpp | 28 ++++++++++------------------ 1 file changed, 10 insertions(+), 18 deletions(-) diff --git a/cpp/tests/io/orc_test.cpp b/cpp/tests/io/orc_test.cpp index 3bc23a636ca2..98c9eae171da 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -655,15 +655,12 @@ TEST_F(OrcWriterTest, negTimestampsNano) template void test_timestamp_roundtrip(std::vector const& values, - std::vector const& expected_values, - bool with_nulls = false) + std::vector const& expected_values) { - auto const validity = cudf::detail::make_counting_transform_iterator( - 0, [with_nulls](auto i) { return not with_nulls or i % 2 == 0; }); - cudf::test::fixed_width_column_wrapper const input( - values.begin(), values.end(), validity); + cudf::test::fixed_width_column_wrapper const input(values.begin(), + values.end()); cudf::test::fixed_width_column_wrapper const expected( - expected_values.begin(), expected_values.end(), validity); + expected_values.begin(), expected_values.end()); cudf::table_view const input_table({input}); std::vector out_buffer; @@ -711,20 +708,15 @@ TEST_F(OrcWriterTest, NegativeFractionalTimestamps) // ORC-771). TEST_F(OrcWriterTest, NegativeTimestampsNearEpoch) { - // The two values around the -999 ms boundary are the smallest shifted and the largest lossless - // timestamp at this resolution. - auto const timestamps_us = std::vector{ - -1L, -500L, -500'000L, -999'000L, -999'001L, -999'999L, -1'000'001L, -5'999'500L}; - auto const read_back_us = std::vector{ - 999'999L, 999'500L, 500'000L, 1'000L, -999'001L, -999'999L, -1'000'001L, -5'999'500L}; + auto const timestamps_us = + std::vector{-1L, -500L, -500'000L, -999'000L}; + auto const read_back_us = + std::vector{999'999L, 999'500L, 500'000L, 1'000L}; test_timestamp_roundtrip(timestamps_us, read_back_us); - test_timestamp_roundtrip(timestamps_us, read_back_us, /* with_nulls */ true); - test_timestamp_roundtrip( - {-1L, -999'000'000L, -999'000'001L, -1'000'000'000L}, - {999'999'999L, 1'000'000L, -999'000'001L, -1'000'000'000L}); + test_timestamp_roundtrip({-1L, -999'000'000L}, {999'999'999L, 1'000'000L}); - test_timestamp_roundtrip({-1L, -999L, -1'000L}, {999L, 1L, -1'000L}); + test_timestamp_roundtrip({-1L, -999L}, {999L, 1L}); } TEST_F(OrcWriterTest, Slice) From aa31766148738a601f2cb2ec8b13d9d0e1bd7514 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Fri, 14 Aug 2026 02:00:44 +0000 Subject: [PATCH 10/10] style --- cpp/tests/io/orc_test.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/cpp/tests/io/orc_test.cpp b/cpp/tests/io/orc_test.cpp index 98c9eae171da..7beaa48a9a4d 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -658,9 +658,9 @@ void test_timestamp_roundtrip(std::vector const& values, std::vector const& expected_values) { cudf::test::fixed_width_column_wrapper const input(values.begin(), - values.end()); - cudf::test::fixed_width_column_wrapper const expected( - expected_values.begin(), expected_values.end()); + values.end()); + cudf::test::fixed_width_column_wrapper const expected(expected_values.begin(), + expected_values.end()); cudf::table_view const input_table({input}); std::vector out_buffer; @@ -708,8 +708,7 @@ TEST_F(OrcWriterTest, NegativeFractionalTimestamps) // ORC-771). TEST_F(OrcWriterTest, NegativeTimestampsNearEpoch) { - auto const timestamps_us = - std::vector{-1L, -500L, -500'000L, -999'000L}; + auto const timestamps_us = std::vector{-1L, -500L, -500'000L, -999'000L}; auto const read_back_us = std::vector{999'999L, 999'500L, 500'000L, 1'000L}; test_timestamp_roundtrip(timestamps_us, read_back_us);