diff --git a/cpp/include/cudf/io/orc.hpp b/cpp/include/cudf/io/orc.hpp index 53d35e6c33b4..47eac12966a3 100644 --- a/cpp/include/cudf/io/orc.hpp +++ b/cpp/include/cudf/io/orc.hpp @@ -1070,6 +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 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 */ @@ -1487,6 +1490,9 @@ class chunked_orc_writer_options_builder { * ... * writer.close(); * @endcode + * + * @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 d76c0dcc57ee..a19530f824a0 100644 --- a/cpp/src/io/orc/stripe_enc.cu +++ b/cpp/src/io/orc/stripe_enc.cu @@ -799,15 +799,29 @@ 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 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)]; + + 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; + } + // 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; 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 41c341b64021..7beaa48a9a4d 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -653,6 +653,71 @@ TEST_F(OrcWriterTest, negTimestampsNano) CUDF_TEST_EXPECT_TABLES_EQUAL(expected, result.tbl->view()); } +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()); + 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}, 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()}}) + .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, result.tbl->view().column(0), cudf::test::debug_output_level::ALL_ERRORS); +} + +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{ + -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, // 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); + + // Millisecond timestamps cannot have a < 1 ms remainder + auto const timestamps_ms = std::vector{ + -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); +} + +// 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) +{ + 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({-1L, -999'000'000L}, {999'999'999L, 1'000'000L}); + + test_timestamp_roundtrip({-1L, -999L}, {999L, 1L}); +} + TEST_F(OrcWriterTest, Slice) { int32_col col{{1, 2, 3, 4, 5}, cudf::test::iterators::null_at(3)}; diff --git a/python/cudf/cudf/utils/ioutils.py b/python/cudf/cudf/utils/ioutils.py index 0db154c3c1a5..c6d8f0df0264 100644 --- a/python/cudf/cudf/utils/ioutils.py +++ b/python/cudf/cudf/utils/ioutils.py @@ -579,6 +579,12 @@ 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 are not +representable in ORC; they are read back one second later, as with the Apache +ORC writer (ORC-763, ORC-771). + See Also -------- cudf.read_orc