From 3468260c7f12f2ef7440fd8570d78b40189ccae5 Mon Sep 17 00:00:00 2001 From: Allen Xu Date: Wed, 21 Oct 2026 12:41:34 +0800 Subject: [PATCH 1/5] Persist proleptic Gregorian calendar in ORC footers Record the standard calendar field for one-shot and chunked writers, expose it through ORC metadata, and add matching C++ and Java/JNI options. Reject unsupported legacy output rather than mislabeling proleptic data. Encode DATE statistics using DateStatistics instead of IntegerStatistics. Cover historical and nested dates, nulls, empty input, calendar defaults, and Apache ORC interoperability through both JNI sink paths. Contributes to NVIDIA/cudf#23794. Signed-off-by: Allen Xu --- cpp/include/cudf/io/orc.hpp | 62 +++++++ cpp/include/cudf/io/orc_metadata.hpp | 22 ++- cpp/include/cudf/io/orc_types.hpp | 9 + cpp/src/io/functions.cpp | 3 +- cpp/src/io/orc/orc.cpp | 4 +- cpp/src/io/orc/orc.hpp | 1 + cpp/src/io/orc/writer_impl.cu | 4 +- cpp/tests/io/orc_test.cpp | 108 ++++++++++++ java/pom.xml | 26 +++ .../java/ai/rapids/cudf/ORCWriterOptions.java | 35 ++++ java/src/main/java/ai/rapids/cudf/Table.java | 6 + java/src/main/native/src/TableJni.cpp | 46 ++--- .../java/ai/rapids/cudf/ORCCalendarTest.java | 160 ++++++++++++++++++ 13 files changed, 461 insertions(+), 25 deletions(-) create mode 100644 java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java diff --git a/cpp/include/cudf/io/orc.hpp b/cpp/include/cudf/io/orc.hpp index 2c0d86d7551b..a00ea177e7a6 100644 --- a/cpp/include/cudf/io/orc.hpp +++ b/cpp/include/cudf/io/orc.hpp @@ -6,6 +6,7 @@ #pragma once #include +#include #include #include #include @@ -664,6 +665,8 @@ class orc_writer_options { bool _enable_dictionary_sort = true; // Timezone that the written timestamps are relative to, recorded in the stripe footers std::string _writer_timezone = "UTC"; + // Calendar recorded in the file footer; libcudf values use the proleptic Gregorian calendar + orc::CalendarKind _calendar_kind = orc::CalendarKind::PROLEPTIC_GREGORIAN; friend orc_writer_options_builder; @@ -800,6 +803,12 @@ class orc_writer_options { */ [[nodiscard]] std::string const& get_writer_timezone() const { return _writer_timezone; } + /** + * @brief Returns the calendar recorded in the file footer. + * @return Calendar kind + */ + [[nodiscard]] orc::CalendarKind get_calendar_kind() const { return _calendar_kind; } + // Setters /** @@ -931,6 +940,22 @@ class orc_writer_options { * @param timezone Timezone name, for example "America/Los_Angeles" */ void set_writer_timezone(std::string timezone) { _writer_timezone = std::move(timezone); } + + /** + * @brief Sets the calendar recorded in the ORC file footer. + * + * Only `PROLEPTIC_GREGORIAN` is supported, matching libcudf's DATE and TIMESTAMP values. + * This option does not rebase values. Legacy Julian/Gregorian output is not supported. + * + * @param calendar_kind Calendar kind + * @throw cudf::logic_error if the calendar is not `PROLEPTIC_GREGORIAN` + */ + void set_calendar_kind(orc::CalendarKind calendar_kind) + { + CUDF_EXPECTS(calendar_kind == orc::CalendarKind::PROLEPTIC_GREGORIAN, + "Only the proleptic Gregorian calendar is supported for ORC writes"); + _calendar_kind = calendar_kind; + } }; /** @@ -1093,6 +1118,16 @@ class orc_writer_options_builder { return *this; } + /** + * @copydoc orc_writer_options::set_calendar_kind + * @return this for chaining + */ + orc_writer_options_builder& calendar_kind(orc::CalendarKind calendar_kind) + { + options.set_calendar_kind(calendar_kind); + return *this; + } + /** * @brief move orc_writer_options member once it's built. */ @@ -1162,6 +1197,8 @@ class chunked_orc_writer_options { bool _enable_dictionary_sort = true; // Timezone that the written timestamps are relative to, recorded in the stripe footers std::string _writer_timezone = "UTC"; + // Calendar recorded in the file footer + orc::CalendarKind _calendar_kind = orc::CalendarKind::PROLEPTIC_GREGORIAN; friend chunked_orc_writer_options_builder; @@ -1276,6 +1313,11 @@ class chunked_orc_writer_options { */ [[nodiscard]] std::string const& get_writer_timezone() const { return _writer_timezone; } + /** + * @copydoc orc_writer_options::get_calendar_kind + */ + [[nodiscard]] orc::CalendarKind get_calendar_kind() const { return _calendar_kind; } + // Setters /** @@ -1383,6 +1425,16 @@ class chunked_orc_writer_options { * @copydoc orc_writer_options::set_writer_timezone */ void set_writer_timezone(std::string timezone) { _writer_timezone = std::move(timezone); } + + /** + * @copydoc orc_writer_options::set_calendar_kind + */ + void set_calendar_kind(orc::CalendarKind calendar_kind) + { + CUDF_EXPECTS(calendar_kind == orc::CalendarKind::PROLEPTIC_GREGORIAN, + "Only the proleptic Gregorian calendar is supported for ORC writes"); + _calendar_kind = calendar_kind; + } }; /** @@ -1531,6 +1583,16 @@ class chunked_orc_writer_options_builder { return *this; } + /** + * @copydoc chunked_orc_writer_options::set_calendar_kind + * @return this for chaining + */ + chunked_orc_writer_options_builder& calendar_kind(orc::CalendarKind calendar_kind) + { + options.set_calendar_kind(calendar_kind); + return *this; + } + /** * @brief move chunked_orc_writer_options member once it's built. */ diff --git a/cpp/include/cudf/io/orc_metadata.hpp b/cpp/include/cudf/io/orc_metadata.hpp index 33613590f8a7..3b6f35e6d674 100644 --- a/cpp/include/cudf/io/orc_metadata.hpp +++ b/cpp/include/cudf/io/orc_metadata.hpp @@ -326,9 +326,16 @@ class orc_metadata { * @param schema ORC schema * @param num_rows number of rows * @param num_stripes number of stripes + * @param calendar_kind calendar recorded in the footer, or `std::nullopt` if absent */ - orc_metadata(orc_schema schema, uint64_t num_rows, size_type num_stripes) - : _schema{std::move(schema)}, _num_rows{num_rows}, _num_stripes{num_stripes} + orc_metadata(orc_schema schema, + uint64_t num_rows, + size_type num_stripes, + std::optional calendar_kind = std::nullopt) + : _schema{std::move(schema)}, + _num_rows{num_rows}, + _num_stripes{num_stripes}, + _calendar_kind{calendar_kind} { } @@ -356,10 +363,21 @@ class orc_metadata { */ [[nodiscard]] auto num_stripes() const { return _num_stripes; } + /** + * @brief Returns the calendar recorded in the ORC footer. + * + * A missing value means the writer did not specify a calendar. This accessor does not + * convert the values returned by `read_orc()` from legacy calendars. + * + * @return Calendar kind, or `std::nullopt` if the footer has no calendar field + */ + [[nodiscard]] auto calendar_kind() const { return _calendar_kind; } + private: orc_schema _schema; uint64_t _num_rows; size_type _num_stripes; + std::optional _calendar_kind; }; /** diff --git a/cpp/include/cudf/io/orc_types.hpp b/cpp/include/cudf/io/orc_types.hpp index 85a020706df4..f21f312165b0 100644 --- a/cpp/include/cudf/io/orc_types.hpp +++ b/cpp/include/cudf/io/orc_types.hpp @@ -33,6 +33,15 @@ enum CompressionKind : uint8_t { ZSTD = 5, }; +/** + * @brief Identifies the calendar used to encode dates and timestamps in an ORC file. + */ +enum class CalendarKind : uint8_t { + UNKNOWN_CALENDAR = 0, ///< Calendar is unspecified + JULIAN_GREGORIAN = 1, ///< Hybrid calendar with the October 1582 cutover + PROLEPTIC_GREGORIAN = 2, ///< Gregorian calendar extended to all dates +}; + /** * @brief Identifies a data type in an orc file */ diff --git a/cpp/src/io/functions.cpp b/cpp/src/io/functions.cpp index 8fa8218e979d..14094f08ca5c 100644 --- a/cpp/src/io/functions.cpp +++ b/cpp/src/io/functions.cpp @@ -474,7 +474,8 @@ orc_metadata read_orc_metadata(source_info const& src_info, cuda::stream_ref str return {{make_orc_column_schema(footer.types, 0, "")}, footer.numberOfRows, - static_cast(footer.stripes.size())}; + static_cast(footer.stripes.size()), + footer.calendar}; } /** diff --git a/cpp/src/io/orc/orc.cpp b/cpp/src/io/orc/orc.cpp index 7fb690bb31bd..0087de022fa0 100644 --- a/cpp/src/io/orc/orc.cpp +++ b/cpp/src/io/orc/orc.cpp @@ -80,7 +80,8 @@ void protobuf_reader::read(Footer& s, size_t maxlen) field_reader(6, s.numberOfRows), raw_field_reader(7, s.statistics), field_reader(8, s.rowIndexStride), - field_reader(9, s.writer)); + field_reader(9, s.writer), + field_reader(11, s.calendar)); function_builder(s, maxlen, op); } @@ -320,6 +321,7 @@ size_t protobuf_writer::write(Footer const& s) w.field_repeated_struct_blob(7, s.statistics); w.field_uint(8, s.rowIndexStride); if (s.writer) w.field_uint(9, *s.writer); + if (s.calendar) w.field_uint(11, static_cast(*s.calendar)); return w.value(); } diff --git a/cpp/src/io/orc/orc.hpp b/cpp/src/io/orc/orc.hpp index 33731493ff4f..9770ab1b509b 100644 --- a/cpp/src/io/orc/orc.hpp +++ b/cpp/src/io/orc/orc.hpp @@ -120,6 +120,7 @@ struct Footer { std::vector statistics; // Column statistics blobs uint32_t rowIndexStride = 0; // the maximum number of rows in each index entry std::optional writer; // Writer code + std::optional calendar; // Calendar used to encode dates and timestamps }; struct Stream { diff --git a/cpp/src/io/orc/writer_impl.cu b/cpp/src/io/orc/writer_impl.cu index dfb3837c5ece..e26a58648425 100644 --- a/cpp/src/io/orc/writer_impl.cu +++ b/cpp/src/io/orc/writer_impl.cu @@ -1363,7 +1363,7 @@ cudf::detail::hostdevice_vector allocate_and_encode_blobs( case TypeKind::FLOAT: return dtype_float32; case TypeKind::DOUBLE: return dtype_float64; case TypeKind::STRING: return dtype_string; - case TypeKind::DATE: return dtype_int32; + case TypeKind::DATE: return dtype_date32; case TypeKind::TIMESTAMP: return dtype_timestamp64; case TypeKind::DECIMAL: return dtype_decimal64; default: return dtype_none; @@ -2738,6 +2738,7 @@ writer::impl::impl(std::unique_ptr sink, _timezone(options.get_writer_timezone()), _out_sink(std::move(sink)) { + _footer.calendar = options.get_calendar_kind(); if (options.get_metadata()) { _table_meta = std::make_unique(*options.get_metadata()); } @@ -2762,6 +2763,7 @@ writer::impl::impl(std::unique_ptr sink, _timezone(options.get_writer_timezone()), _out_sink(std::move(sink)) { + _footer.calendar = options.get_calendar_kind(); if (options.get_metadata()) { _table_meta = std::make_unique(*options.get_metadata()); } diff --git a/cpp/tests/io/orc_test.cpp b/cpp/tests/io/orc_test.cpp index b00c870dd81e..eba01b03d404 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -2091,6 +2091,114 @@ TEST_F(OrcReaderTest, EmptyColumnsParam) EXPECT_EQ(result.tbl->num_rows(), 0); } +TEST_F(OrcWriterTest, ProlepticCalendar) +{ + using calendar = cudf::io::orc::CalendarKind; + // 1001-01-01, 1200-01-01, both sides of the cutover, a date inside the hybrid + // calendar's gap, the Unix epoch, and a modern date. The last row is null. + column_wrapper dates( + {-353920, -281237, -141438, -141432, -141427, 0, 19849, 0}, + {true, true, true, true, true, true, true, false}); + struct_col nested{dates}; + cudf::test::lists_column_wrapper lists{ + {-353920}, {-281237}, {-141438}, {-141432}, {-141427}, {}, {19849}, {0}}; + table_view expected({dates, nested, lists}); + + for (auto compression : {cudf::io::compression_type::NONE, cudf::io::compression_type::SNAPPY}) { + for (bool chunked : {false, true}) { + SCOPED_TRACE(static_cast(compression)); + SCOPED_TRACE(chunked); + std::vector buffer; + if (chunked) { + auto opts = cudf::io::chunked_orc_writer_options::builder(cudf::io::sink_info{&buffer}) + .compression(compression) + .calendar_kind(calendar::PROLEPTIC_GREGORIAN) + .build(); + cudf::io::orc_chunked_writer writer(opts); + auto const chunks = cudf::slice(expected, {0, 3, 3, 8}); + writer.write(chunks[0]).write(chunks[1]); + writer.close(); + } else { + // Exercise the default, which must describe the values the writer actually encodes. + auto opts = cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, expected) + .compression(compression) + .build(); + EXPECT_EQ(opts.get_calendar_kind(), calendar::PROLEPTIC_GREGORIAN); + cudf::io::write_orc(opts); + } + cudf::io::source_info source{cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}}; + EXPECT_EQ(cudf::io::read_orc_metadata(source).calendar_kind(), calendar::PROLEPTIC_GREGORIAN); + auto const result = cudf::io::read_orc( + cudf::io::orc_reader_options::builder(source).use_np_dtypes(false).build()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected, result.tbl->view()); + auto const stats = cudf::io::read_parsed_orc_statistics(source); + auto const date_stats = + std::get(stats.file_stats[1].type_specific_stats); + EXPECT_EQ(date_stats.minimum, -353920); + EXPECT_EQ(date_stats.maximum, 19849); + } + } +} + +TEST_F(OrcWriterTest, CalendarRejectsUnsupportedModes) +{ + using calendar = cudf::io::orc::CalendarKind; + cudf::io::orc_writer_options opts; + cudf::io::chunked_orc_writer_options chunked_opts; + for (auto kind : + {calendar::UNKNOWN_CALENDAR, calendar::JULIAN_GREGORIAN, static_cast(3)}) { + EXPECT_THROW(opts.set_calendar_kind(kind), cudf::logic_error); + EXPECT_THROW(chunked_opts.set_calendar_kind(kind), cudf::logic_error); + } +} + +TEST_F(OrcMetadataReaderTest, CalendarPresenceAndKinds) +{ + // A minimal uncompressed ORC file, expressed as wire bytes independently of libcudf's + // serializer. The footer contains header/content lengths, an empty root STRUCT, and no rows. + using calendar = cudf::io::orc::CalendarKind; + for (auto kind : std::vector>{std::nullopt, + calendar::UNKNOWN_CALENDAR, + calendar::JULIAN_GREGORIAN, + calendar::PROLEPTIC_GREGORIAN}) { + std::vector buffer{ + 'O', 'R', 'C', 0x08, 0x03, 0x10, 0x03, 0x22, 0x02, 0x08, 0x0c, 0x30, 0x00}; + if (kind) { + buffer.push_back(0x58); // Footer.calendar, protobuf field 11, varint + buffer.push_back(static_cast(*kind)); + } + auto const footer_size = static_cast(buffer.size() - 3); + buffer.insert(buffer.end(), {0x08, footer_size, 0x10, 0x00, 0x04}); + cudf::io::source_info source{cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}}; + EXPECT_EQ(cudf::io::read_orc_metadata(source).calendar_kind(), kind); + } +} + +TEST_F(OrcWriterTest, CalendarOnEmptyInput) +{ + column_wrapper dates{}; + for (bool chunked : {false, true}) { + std::vector buffer; + table_view empty({dates}); + if (chunked) { + cudf::io::orc_chunked_writer writer( + cudf::io::chunked_orc_writer_options::builder(cudf::io::sink_info{&buffer})); + writer.write(empty); + writer.close(); + } else { + cudf::io::write_orc( + cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, empty)); + } + cudf::io::source_info source{cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}}; + auto const meta = cudf::io::read_orc_metadata(source); + EXPECT_EQ(meta.num_rows(), 0); + EXPECT_EQ(meta.calendar_kind(), cudf::io::orc::CalendarKind::PROLEPTIC_GREGORIAN); + } +} + TEST_F(OrcMetadataReaderTest, TestBasic) { auto const num_rows = 1'200'000; diff --git a/java/pom.xml b/java/pom.xml index ae5d47016506..b168ddc2ca35 100644 --- a/java/pom.xml +++ b/java/pom.xml @@ -136,6 +136,32 @@ + + org.apache.orc + orc-core + 1.9.4 + shaded-protobuf + test + + + + org.apache.hadoop + hadoop-client-api + + + + + org.apache.hadoop + hadoop-hdfs-client + 3.4.0 + test + + + org.apache.hive + hive-storage-api + 2.8.1 + test + org.slf4j slf4j-api diff --git a/java/src/main/java/ai/rapids/cudf/ORCWriterOptions.java b/java/src/main/java/ai/rapids/cudf/ORCWriterOptions.java index 21f186ffc5d6..0e1d2efa7906 100644 --- a/java/src/main/java/ai/rapids/cudf/ORCWriterOptions.java +++ b/java/src/main/java/ai/rapids/cudf/ORCWriterOptions.java @@ -12,13 +12,27 @@ * that will be used by the ORC writer to write the file. */ public class ORCWriterOptions extends CompressionMetadataWriterOptions { + /** Calendar recorded in the standard ORC file footer. */ + public enum CalendarKind { + JULIAN_GREGORIAN(1), + PROLEPTIC_GREGORIAN(2); + + final int nativeId; + + CalendarKind(int nativeId) { + this.nativeId = nativeId; + } + } + private int stripeSizeRows; private String writerTimezone; + private CalendarKind calendarKind; private ORCWriterOptions(Builder builder) { super(builder); this.stripeSizeRows = builder.stripeSizeRows; this.writerTimezone = builder.writerTimezone; + this.calendarKind = builder.calendarKind; } public static Builder builder() { @@ -33,11 +47,16 @@ public String getWriterTimezone() { return writerTimezone; } + public CalendarKind getCalendarKind() { + return calendarKind; + } + public static class Builder extends CompressionMetadataWriterOptions.Builder { // < 1M rows default orc stripe rows, defined in cudf/cpp/include/cudf/io/orc.hpp private int stripeSizeRows = 1000000; private String writerTimezone = "UTC"; + private CalendarKind calendarKind = CalendarKind.PROLEPTIC_GREGORIAN; public Builder withStripeSizeRows(int stripeSizeRows) { // maximum stripe size cannot be smaller than 512 @@ -69,6 +88,22 @@ public Builder withWriterTimezone(String writerTimezone) { return this; } + /** + * Sets the calendar recorded in the file footer. Only {@code PROLEPTIC_GREGORIAN} is + * supported, matching cuDF's date and timestamp values. This does not rebase values. + * + * @param calendarKind calendar kind + * @throws IllegalArgumentException if the calendar is null or is not proleptic Gregorian + */ + public Builder withCalendarKind(CalendarKind calendarKind) { + if (calendarKind != CalendarKind.PROLEPTIC_GREGORIAN) { + throw new IllegalArgumentException( + "Only the proleptic Gregorian calendar is supported for ORC writes"); + } + this.calendarKind = calendarKind; + return this; + } + public ORCWriterOptions build() { return new ORCWriterOptions(this); } diff --git a/java/src/main/java/ai/rapids/cudf/Table.java b/java/src/main/java/ai/rapids/cudf/Table.java index 78061d5710b6..659e7593bfd6 100644 --- a/java/src/main/java/ai/rapids/cudf/Table.java +++ b/java/src/main/java/ai/rapids/cudf/Table.java @@ -363,6 +363,7 @@ private static native long[] readORCFromDataSource(String[] filterColumnNames, * the columns * @param isMapValues true if a column is a map * @param writerTimezone timezone that the written timestamps are relative to + * @param calendarKind native calendar kind recorded in the ORC footer * @param filename local output path * @return a handle that is used in later calls to writeORCChunk and writeORCEnd. */ @@ -377,6 +378,7 @@ private static native long writeORCFileBegin(String[] columnNames, boolean[] isMapValues, int stripeSizeRows, String writerTimezone, + int calendarKind, String filename) throws CudfException; /** @@ -392,6 +394,7 @@ private static native long writeORCFileBegin(String[] columnNames, * the columns * @param isMapValues true if a column is a map * @param writerTimezone timezone that the written timestamps are relative to + * @param calendarKind native calendar kind recorded in the ORC footer * @param consumer consumer of host buffers produced. * @return a handle that is used in later calls to writeORCChunk and writeORCEnd. */ @@ -406,6 +409,7 @@ private static native long writeORCBufferBegin(String[] columnNames, boolean[] isMapValues, int stripeSizeRows, String writerTimezone, + int calendarKind, HostBufferConsumer consumer, HostMemoryAllocator hostMemoryAllocator ) throws CudfException; @@ -1698,6 +1702,7 @@ private ORCTableWriter(ORCWriterOptions options, File outputFile) { options.getFlatIsMap(), options.getStripeSizeRows(), options.getWriterTimezone(), + options.getCalendarKind().nativeId, outputFile.getAbsolutePath())); this.consumer = null; } @@ -1715,6 +1720,7 @@ private ORCTableWriter(ORCWriterOptions options, HostBufferConsumer consumer, options.getFlatIsMap(), options.getStripeSizeRows(), options.getWriterTimezone(), + options.getCalendarKind().nativeId, consumer, hostMemoryAllocator)); this.consumer = consumer; } diff --git a/java/src/main/native/src/TableJni.cpp b/java/src/main/native/src/TableJni.cpp index 405af0b85f8e..942e66c2e763 100644 --- a/java/src/main/native/src/TableJni.cpp +++ b/java/src/main/native/src/TableJni.cpp @@ -2599,6 +2599,7 @@ Java_ai_rapids_cudf_Table_writeORCBufferBegin(JNIEnv* env, jbooleanArray j_is_map, jint j_stripe_size_rows, jstring j_writer_timezone, + jint j_calendar_kind, jobject consumer, jobject host_memory_allocator) { @@ -2650,16 +2651,18 @@ Java_ai_rapids_cudf_Table_writeORCBufferBegin(JNIEnv* env, new cudf::jni::jni_writer_data_sink(env, consumer, host_memory_allocator)); sink_info sink{data_sink.get()}; - auto stats = std::make_shared(); - chunked_orc_writer_options opts = chunked_orc_writer_options::builder(sink) - .metadata(std::move(metadata)) - .compression(static_cast(j_compression)) - .enable_statistics(ORC_STATISTICS_ROW_GROUP) - .key_value_metadata(kv_metadata) - .compression_statistics(stats) - .stripe_size_rows(j_stripe_size_rows) - .writer_timezone(writer_timezone.get()) - .build(); + auto stats = std::make_shared(); + chunked_orc_writer_options opts = + chunked_orc_writer_options::builder(sink) + .metadata(std::move(metadata)) + .compression(static_cast(j_compression)) + .enable_statistics(ORC_STATISTICS_ROW_GROUP) + .key_value_metadata(kv_metadata) + .compression_statistics(stats) + .stripe_size_rows(j_stripe_size_rows) + .writer_timezone(writer_timezone.get()) + .calendar_kind(static_cast(j_calendar_kind)) + .build(); auto writer_ptr = std::make_unique(opts); cudf::jni::native_orc_writer_handle* ret = new cudf::jni::native_orc_writer_handle( std::move(writer_ptr), std::move(data_sink), std::move(stats)); @@ -2681,6 +2684,7 @@ JNIEXPORT long JNICALL Java_ai_rapids_cudf_Table_writeORCFileBegin(JNIEnv* env, jbooleanArray j_is_map, jint j_stripe_size_rows, jstring j_writer_timezone, + jint j_calendar_kind, jstring j_output_path) { JNI_NULL_CHECK(env, j_col_names, "null columns", 0); @@ -2728,16 +2732,18 @@ JNIEXPORT long JNICALL Java_ai_rapids_cudf_Table_writeORCFileBegin(JNIEnv* env, [](std::string const& k, std::string const& v) { return std::make_pair(k, v); }); sink_info sink{output_path}; - auto stats = std::make_shared(); - chunked_orc_writer_options opts = chunked_orc_writer_options::builder(sink) - .metadata(std::move(metadata)) - .compression(static_cast(j_compression)) - .enable_statistics(ORC_STATISTICS_ROW_GROUP) - .key_value_metadata(kv_metadata) - .compression_statistics(stats) - .stripe_size_rows(j_stripe_size_rows) - .writer_timezone(writer_timezone.get()) - .build(); + auto stats = std::make_shared(); + chunked_orc_writer_options opts = + chunked_orc_writer_options::builder(sink) + .metadata(std::move(metadata)) + .compression(static_cast(j_compression)) + .enable_statistics(ORC_STATISTICS_ROW_GROUP) + .key_value_metadata(kv_metadata) + .compression_statistics(stats) + .stripe_size_rows(j_stripe_size_rows) + .writer_timezone(writer_timezone.get()) + .calendar_kind(static_cast(j_calendar_kind)) + .build(); auto writer_ptr = std::make_unique(opts); cudf::jni::native_orc_writer_handle* ret = new cudf::jni::native_orc_writer_handle(std::move(writer_ptr), nullptr, std::move(stats)); diff --git a/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java b/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java new file mode 100644 index 000000000000..6f2a31393305 --- /dev/null +++ b/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java @@ -0,0 +1,160 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +package ai.rapids.cudf; + +import org.apache.hadoop.conf.Configuration; +import org.apache.hadoop.fs.Path; +import org.apache.hadoop.hive.ql.exec.vector.ListColumnVector; +import org.apache.hadoop.hive.ql.exec.vector.LongColumnVector; +import org.apache.hadoop.hive.ql.exec.vector.StructColumnVector; +import org.apache.hadoop.hive.ql.exec.vector.VectorizedRowBatch; +import org.apache.orc.DateColumnStatistics; +import org.apache.orc.OrcFile; +import org.apache.orc.Reader; +import org.apache.orc.RecordReader; +import org.junit.jupiter.api.Test; + +import java.io.ByteArrayOutputStream; +import java.io.IOException; +import java.nio.file.Files; +import java.time.LocalDate; +import java.util.Arrays; +import java.util.Collections; + +import static ai.rapids.cudf.AssertUtils.assertTablesAreEqual; +import static org.junit.jupiter.api.Assertions.assertEquals; +import static org.junit.jupiter.api.Assertions.assertFalse; +import static org.junit.jupiter.api.Assertions.assertThrows; +import static org.junit.jupiter.api.Assertions.assertTrue; + +class ORCCalendarTest { + private static final Integer[] DAYS = Arrays.stream(new String[] { + "1001-01-01", "1200-01-01", "1582-10-04", "1582-10-10", "1582-10-15", + "1970-01-01", "2024-05-06", null + }).map(value -> value == null ? null : Math.toIntExact(LocalDate.parse(value).toEpochDay())) + .toArray(Integer[]::new); + + @Test + void rejectsUnsupportedCalendars() { + assertEquals(ORCWriterOptions.CalendarKind.PROLEPTIC_GREGORIAN, + ORCWriterOptions.builder().build().getCalendarKind()); + assertThrows(IllegalArgumentException.class, + () -> ORCWriterOptions.builder().withCalendarKind(null)); + assertThrows(IllegalArgumentException.class, () -> ORCWriterOptions.builder() + .withCalendarKind(ORCWriterOptions.CalendarKind.JULIAN_GREGORIAN)); + } + + @Test + @SuppressWarnings("unchecked") + void datesInteroperateWithApacheOrc() throws IOException { + HostColumnVector.ListType listType = new HostColumnVector.ListType(true, + new HostColumnVector.BasicType(true, DType.TIMESTAMP_DAYS)); + try (ColumnVector dates = ColumnVector.timestampDaysFromBoxedInts(DAYS); + ColumnVector structs = ColumnVector.makeStruct(dates); + ColumnVector lists = ColumnVector.fromLists(listType, + Arrays.stream(DAYS).map(Collections::singletonList).toArray(java.util.List[]::new)); + Table input = new Table(dates, structs, lists)) { + for (CompressionType compression : new CompressionType[] { + CompressionType.NONE, CompressionType.SNAPPY}) { + for (boolean bufferSink : new boolean[] {false, true}) { + for (boolean explicitCalendar : new boolean[] {false, true}) { + java.nio.file.Path file = Files.createTempFile("orc-calendar", ".orc"); + try { + ORCWriterOptions.Builder builder = ORCWriterOptions.builder() + .withColumns(true, "d") + .withStructColumn(ColumnWriterOptions.structBuilder("s") + .withColumns(true, "d").build()) + .withListColumn(ColumnWriterOptions.listBuilder("l") + .withColumns(true, "d").build()) + .withCompressionType(compression); + if (explicitCalendar) { + builder.withCalendarKind(ORCWriterOptions.CalendarKind.PROLEPTIC_GREGORIAN); + } + ORCWriterOptions options = builder.build(); + if (bufferSink) { + ByteArrayOutputStream output = new ByteArrayOutputStream(); + HostBufferConsumer consumer = (buffer, length) -> { + try { + byte[] bytes = new byte[Math.toIntExact(length)]; + buffer.getBytes(bytes, 0, 0, bytes.length); + output.write(bytes, 0, bytes.length); + } finally { + buffer.close(); + } + }; + try (TableWriter writer = Table.writeORCChunked(options, consumer)) { + writer.write(input); + writer.write(input); + } + Files.write(file, output.toByteArray()); + } else { + try (TableWriter writer = Table.writeORCChunked(options, file.toFile())) { + writer.write(input); + writer.write(input); + } + } + try (Table expected = Table.concatenate(input, input); + Table actual = Table.readORC( + ORCOptions.builder().withNumPyTypes(false).build(), file.toFile())) { + assertTablesAreEqual(expected, actual); + } + // Both missing-field defaults must be overridden by the file's explicit calendar. + for (boolean defaultProleptic : new boolean[] {false, true}) { + Configuration conf = new Configuration(); + conf.setBoolean("orc.proleptic.gregorian.default", defaultProleptic); + try (Reader reader = OrcFile.createReader(new Path(file.toUri()), + OrcFile.readerOptions(conf).convertToProlepticGregorian(true))) { + assertTrue(reader.writerUsedProlepticGregorian()); + assertEquals(DAYS.length * 2, reader.getNumberOfRows()); + // DATE leaves in the top-level column, STRUCT, and LIST. + for (int columnId : new int[] {1, 3, 5}) { + DateColumnStatistics stats = + (DateColumnStatistics) reader.getStatistics()[columnId]; + assertEquals(DAYS[0].longValue(), stats.getMinimumDayOfEpoch()); + assertEquals(DAYS[6].longValue(), stats.getMaximumDayOfEpoch()); + assertEquals((DAYS.length - 1) * 2, stats.getNumberOfValues()); + } + VectorizedRowBatch batch = reader.getSchema().createRowBatch(); + int rowCount = 0; + try (RecordReader rows = reader.rows()) { + while (rows.nextBatch(batch)) { + LongColumnVector top = (LongColumnVector) batch.cols[0]; + LongColumnVector nested = (LongColumnVector) + ((StructColumnVector) batch.cols[1]).fields[0]; + ListColumnVector list = (ListColumnVector) batch.cols[2]; + LongColumnVector child = (LongColumnVector) list.child; + for (int row = 0; row < batch.size; row++, rowCount++) { + Integer expected = DAYS[rowCount % DAYS.length]; + assertDate(top, row, expected); + assertDate(nested, row, expected); + assertEquals(1, list.lengths[row]); + assertDate(child, Math.toIntExact(list.offsets[row]), expected); + } + } + } + assertEquals(DAYS.length * 2, rowCount); + } + } + } finally { + Files.deleteIfExists(file); + } + } + } + } + } + } + + private static void assertDate(LongColumnVector column, int row, Integer expected) { + int index = column.isRepeating ? 0 : row; + if (expected == null) { + assertFalse(column.noNulls); + assertTrue(column.isNull[index]); + } else { + assertTrue(column.noNulls || !column.isNull[index]); + assertEquals(expected.longValue(), column.vector[index]); + } + } +} From 70eeb18bb7470868d8e57db9537ff231c2a564a0 Mon Sep 17 00:00:00 2001 From: Allen Xu Date: Wed, 21 Oct 2026 13:45:39 +0800 Subject: [PATCH 2/5] Support legacy ORC DATE calendar rebasing Rebase DATE leaves before encoding and statistics, and add opt-in conversion when reading explicitly marked hybrid files. Keep raw reader behavior for existing callers, and reject lossy dates and legacy TIMESTAMP schemas. Cover nested slices, hidden nulls, mixed calendars, retries, wide dates, and Apache ORC interoperability. Signed-off-by: Allen Xu --- cpp/include/cudf/io/orc.hpp | 53 +++- cpp/include/cudf/io/orc_metadata.hpp | 2 +- cpp/src/io/orc/calendar.cuh | 60 +++++ cpp/src/io/orc/orc_gpu.hpp | 9 +- cpp/src/io/orc/reader_impl.cu | 1 + cpp/src/io/orc/reader_impl.hpp | 7 +- cpp/src/io/orc/reader_impl_decode.cu | 4 + cpp/src/io/orc/stripe_data.cu | 15 +- cpp/src/io/orc/writer_impl.cu | 80 ++++++ cpp/tests/io/orc_test.cpp | 243 ++++++++++++++++-- .../main/java/ai/rapids/cudf/ORCOptions.java | 20 +- .../java/ai/rapids/cudf/ORCWriterOptions.java | 13 +- java/src/main/java/ai/rapids/cudf/Table.java | 12 +- java/src/main/native/src/TableJni.cpp | 6 +- .../java/ai/rapids/cudf/ORCCalendarTest.java | 173 ++++++++++++- 15 files changed, 635 insertions(+), 63 deletions(-) create mode 100644 cpp/src/io/orc/calendar.cuh diff --git a/cpp/include/cudf/io/orc.hpp b/cpp/include/cudf/io/orc.hpp index a00ea177e7a6..951110a99f8d 100644 --- a/cpp/include/cudf/io/orc.hpp +++ b/cpp/include/cudf/io/orc.hpp @@ -83,7 +83,8 @@ class orc_reader_options { bool _use_index = true; // Whether to use numpy-compatible dtypes - bool _use_np_dtypes = true; + bool _use_np_dtypes = true; + bool _convert_dates_to_proleptic = false; // Cast timestamp columns to a specific type data_type _timestamp_type{type_id::EMPTY}; @@ -168,6 +169,12 @@ class orc_reader_options { */ [[nodiscard]] bool is_enabled_use_np_dtypes() const { return _use_np_dtypes; } + /** + * @brief Whether explicitly marked hybrid DATE values are converted to proleptic Gregorian. + * @return Whether DATE conversion is enabled + */ + [[nodiscard]] bool get_convert_dates_to_proleptic() const { return _convert_dates_to_proleptic; } + /** * @brief Returns timestamp type to which timestamp column will be cast. * @@ -275,6 +282,19 @@ class orc_reader_options { */ void enable_use_np_dtypes(bool use) { _use_np_dtypes = use; } + /** + * @brief Convert DATE values from explicitly marked JULIAN_GREGORIAN files to proleptic + * Gregorian. + * + * Disabled by default to preserve raw epoch-day reads used by callers that rebase themselves. + * Missing and UNKNOWN calendar fields are left unchanged. TIMESTAMP columns are unchanged. + * Conversion precedes numpy dtype promotion and applies recursively to nested DATE columns. + * Julian-only leap days and dates outside the int32 proleptic domain cause a read error. + * + * @param convert Whether to convert explicitly marked hybrid DATE values + */ + void set_convert_dates_to_proleptic(bool convert) { _convert_dates_to_proleptic = convert; } + /** * @brief Sets timestamp type to which timestamp column will be cast. * @@ -401,6 +421,16 @@ class orc_reader_options_builder { return *this; } + /** + * @copydoc orc_reader_options::set_convert_dates_to_proleptic + * @return this for chaining + */ + orc_reader_options_builder& convert_dates_to_proleptic(bool convert) + { + options.set_convert_dates_to_proleptic(convert); + return *this; + } + /** * @brief Sets timestamp type to which timestamp column will be cast. * @@ -665,7 +695,7 @@ class orc_writer_options { bool _enable_dictionary_sort = true; // Timezone that the written timestamps are relative to, recorded in the stripe footers std::string _writer_timezone = "UTC"; - // Calendar recorded in the file footer; libcudf values use the proleptic Gregorian calendar + // Calendar used to encode DATE values and recorded in the file footer orc::CalendarKind _calendar_kind = orc::CalendarKind::PROLEPTIC_GREGORIAN; friend orc_writer_options_builder; @@ -944,16 +974,20 @@ class orc_writer_options { /** * @brief Sets the calendar recorded in the ORC file footer. * - * Only `PROLEPTIC_GREGORIAN` is supported, matching libcudf's DATE and TIMESTAMP values. - * This option does not rebase values. Legacy Julian/Gregorian output is not supported. + * `PROLEPTIC_GREGORIAN` is the default and retains libcudf's DATE and TIMESTAMP values. + * Legacy mode rebases DATE values, including nested values, before encoding and computing + * statistics. TIMESTAMP columns, dates in 1582-10-05 through 1582-10-14, and dates outside + * the int32 hybrid domain are rejected. Use convert_dates_to_proleptic when reading legacy + * output back into libcudf. * * @param calendar_kind Calendar kind - * @throw cudf::logic_error if the calendar is not `PROLEPTIC_GREGORIAN` + * @throw cudf::logic_error if the calendar kind is unknown or invalid */ void set_calendar_kind(orc::CalendarKind calendar_kind) { - CUDF_EXPECTS(calendar_kind == orc::CalendarKind::PROLEPTIC_GREGORIAN, - "Only the proleptic Gregorian calendar is supported for ORC writes"); + CUDF_EXPECTS(calendar_kind == orc::CalendarKind::PROLEPTIC_GREGORIAN or + calendar_kind == orc::CalendarKind::JULIAN_GREGORIAN, + "Unsupported ORC calendar kind"); _calendar_kind = calendar_kind; } }; @@ -1431,8 +1465,9 @@ class chunked_orc_writer_options { */ void set_calendar_kind(orc::CalendarKind calendar_kind) { - CUDF_EXPECTS(calendar_kind == orc::CalendarKind::PROLEPTIC_GREGORIAN, - "Only the proleptic Gregorian calendar is supported for ORC writes"); + CUDF_EXPECTS(calendar_kind == orc::CalendarKind::PROLEPTIC_GREGORIAN or + calendar_kind == orc::CalendarKind::JULIAN_GREGORIAN, + "Unsupported ORC calendar kind"); _calendar_kind = calendar_kind; } }; diff --git a/cpp/include/cudf/io/orc_metadata.hpp b/cpp/include/cudf/io/orc_metadata.hpp index 3b6f35e6d674..485c3a497908 100644 --- a/cpp/include/cudf/io/orc_metadata.hpp +++ b/cpp/include/cudf/io/orc_metadata.hpp @@ -367,7 +367,7 @@ class orc_metadata { * @brief Returns the calendar recorded in the ORC footer. * * A missing value means the writer did not specify a calendar. This accessor does not - * convert the values returned by `read_orc()` from legacy calendars. + * convert values. See `orc_reader_options::set_convert_dates_to_proleptic` for DATE conversion. * * @return Calendar kind, or `std::nullopt` if the footer has no calendar field */ diff --git a/cpp/src/io/orc/calendar.cuh b/cpp/src/io/orc/calendar.cuh new file mode 100644 index 000000000000..284c7048e254 --- /dev/null +++ b/cpp/src/io/orc/calendar.cuh @@ -0,0 +1,60 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include + +#include + +namespace cudf::io::orc::detail { + +// Epoch days of 1582-10-15. Earlier hybrid dates use the Julian calendar. +constexpr int32_t calendar_cutover = -141427; + +CUDF_HOST_DEVICE constexpr int64_t floor_div(int64_t value, int64_t divisor) +{ + return value / divisor - (value % divisor < 0); +} + +/** + * @brief Convert epoch days while preserving the civil year, month, and day. + * + * March-based 400-year Gregorian and 4-year Julian cycles keep intermediate values positive + * within each cycle, including for astronomical years zero and below. All arithmetic is 64 bit + * so that the complete int32 DATE domain can be checked before narrowing the result. + */ +template +CUDF_HOST_DEVICE constexpr int64_t rebase_days(int64_t days) +{ + if (days >= calendar_cutover) { return days; } + + auto const z = days + (to_hybrid ? 719468 : 719470); + auto const era = floor_div(z, to_hybrid ? 146097 : 1461); + auto const doe = z - era * (to_hybrid ? 146097 : 1461); + auto const yoe = to_hybrid ? (doe - doe / 1460 + doe / 36524 - doe / 146096) / 365 + : (doe == 1460 ? 3 : doe / 365); + auto const year = yoe + era * (to_hybrid ? 400 : 4); + auto const doy = doe - (365 * yoe + (to_hybrid ? yoe / 4 - yoe / 100 : 0)); + auto const mp = (5 * doy + 2) / 153; + auto const day = doy - (153 * mp + 2) / 5 + 1; + + if constexpr (to_hybrid) { + return 365 * year + floor_div(year, 4) + (153 * mp + 2) / 5 + day - 1 - 719470; + } else { + return 365 * year + floor_div(year, 4) - floor_div(year, 100) + floor_div(year, 400) + + (153 * mp + 2) / 5 + day - 1 - 719468; + } +} + +template +CUDF_HOST_DEVICE constexpr bool is_representable_date(int64_t original, int64_t rebased) +{ + // The hybrid cutover gap and Julian-only leap days have no lossless counterpart. + return rebased >= INT32_MIN and rebased <= INT32_MAX and + rebase_days(rebased) == original; +} + +} // namespace cudf::io::orc::detail diff --git a/cpp/src/io/orc/orc_gpu.hpp b/cpp/src/io/orc/orc_gpu.hpp index 22b92e67af69..703575fc9767 100644 --- a/cpp/src/io/orc/orc_gpu.hpp +++ b/cpp/src/io/orc/orc_gpu.hpp @@ -116,10 +116,11 @@ struct column_desc { uint32_t rowgroup_id; // row group position ColumnEncodingKind encoding_kind; // column encoding kind TypeKind type_kind; // column data type - uint8_t dtype_len; // data type length (for types that can be mapped to different sizes) - type_id dtype_id; // TODO - int32_t decimal_scale; // number of fractional decimal digits for decimal type - type_id timestamp_type_id; // output timestamp type id (type_id::EMPTY by default) + uint8_t dtype_len; // data type length (for types that can be mapped to different sizes) + bool convert_dates_to_proleptic; // explicitly marked hybrid DATE values + type_id dtype_id; // TODO + int32_t decimal_scale; // number of fractional decimal digits for decimal type + type_id timestamp_type_id; // output timestamp type id (type_id::EMPTY by default) column_validity_info parent_validity_info; // consists of parent column valid_map and null count uint32_t* parent_null_count_prefix_sums; // per-stripe prefix sums of parent column's null count }; diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index e10f8508009f..80312f42948d 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -156,6 +156,7 @@ reader_impl::reader_impl(std::size_t chunk_read_limit, _mr(mr), _options{options.get_timestamp_type(), options.is_enabled_use_index(), + options.get_convert_dates_to_proleptic(), options.is_enabled_use_np_dtypes(), options.get_decimal128_columns(), options.get_ignore_timezone_in_stripe_footer(), diff --git a/cpp/src/io/orc/reader_impl.hpp b/cpp/src/io/orc/reader_impl.hpp index 12ee62e83d68..8adcddc7941a 100644 --- a/cpp/src/io/orc/reader_impl.hpp +++ b/cpp/src/io/orc/reader_impl.hpp @@ -163,9 +163,10 @@ class reader_impl { // Reader configs. struct { - data_type timestamp_type; // override output timestamp resolution - bool use_index; // enable or disable attempt to use row index for parsing - bool use_np_dtypes; // enable or disable the conversion to numpy-compatible dtypes + data_type timestamp_type; // override output timestamp resolution + bool use_index; // enable or disable attempt to use row index for parsing + bool convert_dates_to_proleptic; // rebase explicitly marked hybrid DATE values + bool use_np_dtypes; // enable or disable the conversion to numpy-compatible dtypes std::vector decimal128_columns; // control decimals conversion bool ignore_timezone_in_stripe_footer; // ignore writer timezone in strip footer diff --git a/cpp/src/io/orc/reader_impl_decode.cu b/cpp/src/io/orc/reader_impl_decode.cu index 4a8dd4a90e79..24d5118308f3 100644 --- a/cpp/src/io/orc/reader_impl_decode.cu +++ b/cpp/src/io/orc/reader_impl_decode.cu @@ -869,6 +869,10 @@ void reader_impl::decompress_and_decode_stripes(read_mode mode) // num_child_rows for a struct column will be same, for other nested types it will be // calculated. + chunk.convert_dates_to_proleptic = + _options.convert_dates_to_proleptic and + _metadata.per_file_metadata[stripe.source_idx].ff.calendar == + CalendarKind::JULIAN_GREGORIAN; chunk.num_child_rows = (chunk.type_kind != STRUCT) ? 0 : chunk.num_rows; chunk.dtype_id = column_types[col_idx].id(); chunk.decimal_scale = _metadata.per_file_metadata[stripe.source_idx] diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index 37d0e6066f3e..a59a3c85b24c 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -3,6 +3,7 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include "calendar.cuh" #include "io/utilities/block_utils.cuh" #include "io/utilities/column_buffer.hpp" #include "orc_gpu.hpp" @@ -1999,16 +2000,24 @@ CUDF_KERNEL void __launch_bounds__(block_size) static_cast(data_out)[row] = (s->vals.u8[(t + vals_skipped) >> 3] >> ((~(t + vals_skipped)) & 7)) & 1; break; - case DATE: + case DATE: { + auto const original = s->vals.i32[t + vals_skipped]; + auto const rebased = + s->chunk.convert_dates_to_proleptic ? rebase_days(original) : original; + if (s->chunk.convert_dates_to_proleptic and + not is_representable_date(original, rebased)) { + atomicAdd(error_count, 1); + } if (s->chunk.dtype_len == 8) { - cudf::duration_D days{s->vals.i32[t + vals_skipped]}; + cudf::duration_D days{static_cast(rebased)}; // Convert from days to milliseconds static_cast(data_out)[row] = cuda::std::chrono::duration_cast(days).count(); } else { - static_cast(data_out)[row] = s->vals.u32[t + vals_skipped]; + static_cast(data_out)[row] = static_cast(rebased); } break; + } case STRING: case BINARY: case VARCHAR: diff --git a/cpp/src/io/orc/writer_impl.cu b/cpp/src/io/orc/writer_impl.cu index e26a58648425..eed92b5f654c 100644 --- a/cpp/src/io/orc/writer_impl.cu +++ b/cpp/src/io/orc/writer_impl.cu @@ -8,6 +8,7 @@ * @brief cuDF-IO ORC writer class implementation */ +#include "calendar.cuh" #include "io/comp/compression.hpp" #include "io/orc/orc_gpu.hpp" #include "io/statistics/column_statistics.cuh" @@ -56,6 +57,7 @@ #include #include #include +#include #include #include @@ -218,6 +220,7 @@ class orc_column_view { void add_child(uint32_t child_idx) { children.emplace_back(child_idx); } + auto const& view() const noexcept { return cudf_column; } auto type() const noexcept { return cudf_column.type(); } auto is_string() const noexcept { return cudf_column.type().id() == type_id::STRING; } @@ -325,6 +328,8 @@ struct orc_table_view { cuda::device_buffer d_columns; std::vector string_column_indices; cuda::device_buffer d_string_column_indices; + std::vector> rebased_dates; + std::vector> rebased_date_masks; [[nodiscard]] auto num_columns() const noexcept { return columns.size(); } [[nodiscard]] size_type num_rows() const noexcept @@ -2144,6 +2149,76 @@ hostdevice_2dvector calculate_rowgroup_bounds(orc_table_view cons return rowgroup_bounds; } +/** + * @brief Rebase DATE leaves before encoding and statistics, using the same device column views. + * + * Rowgroup bounds restrict validation to rows actually written from sliced nested columns. + * Pushdown masks also exclude values hidden by null ancestors. Owned buffers live in orc_table. + */ +void rebase_date_columns(orc_table_view& orc_table, + pushdown_null_masks const& pd_masks, + host_2dspan bounds, + cuda::stream_ref stream) +{ + for (auto const& col : orc_table.columns) { + CUDF_EXPECTS(col.orc_kind() != TIMESTAMP, + "Legacy ORC calendar does not support TIMESTAMP columns"); + } + if (orc_table.num_rows() == 0) { return; } + for (auto const& col : orc_table.columns) { + if (col.orc_kind() != DATE) { continue; } + auto const& view = col.view(); + auto const begin = bounds[0][col.index()].begin; + auto const end = bounds[bounds.size().first - 1][col.index()].end; + auto& data = orc_table.rebased_dates.emplace_back( + static_cast(view.offset()) + view.size(), stream); + auto const columns = device_span{orc_table.d_columns}; + auto const parent_mask = col.is_child() ? pd_masks.masks[col.parent_index()] : nullptr; + auto mask = view.null_mask(); + auto null_count = view.null_count(); + if (parent_mask != nullptr) { + // Statistics use the base column view, so they need the ancestor validity as well. + std::vector masks{parent_mask}; + if (mask != nullptr) { masks.push_back(mask); } + std::vector offsets(masks.size(), 0); + auto [combined, count] = cudf::detail::bitmask_and(masks, + offsets, + view.offset() + view.size(), + stream, + cudf::get_current_device_resource_ref()); + mask = orc_table.rebased_date_masks.emplace_back(std::move(combined)).data(); + null_count = view.offset() == 0 ? count + : cudf::detail::null_count( + mask, view.offset(), view.offset() + view.size(), stream); + } + auto const invalid = thrust::transform_reduce( + rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator{begin}, + cuda::counting_iterator{end}, + [columns, idx = col.index(), mask, out = data.data()] __device__(auto row) -> bool { + auto const& column = columns[idx]; + if (not bit_value_or(mask, row + column.offset(), true)) { + out[row + column.offset()] = 0; + return false; + } + auto const original = column.element(row); + auto const rebased = rebase_days(original); + out[row + column.offset()] = static_cast(rebased); + return not is_representable_date(original, rebased); + }, + false, + cuda::std::logical_or{}); + CUDF_EXPECTS(not invalid, "DATE cannot be represented losslessly in the legacy ORC calendar"); + auto const rebased_view = column_device_view::create( + column_view{view.type(), view.size(), data.data(), mask, null_count, view.offset()}, stream); + cudf::detail::device_single_thread( + [columns, idx = col.index(), rebased = *rebased_view] __device__() { + static_cast(columns[idx]) = rebased; + }, + stream); + } +} + // returns host vector of per-rowgroup sizes encoder_decimal_info decimal_chunk_sizes(orc_table_view& orc_table, file_segmentation const& segmentation, @@ -2537,6 +2612,7 @@ auto convert_table_to_orc_data(table_view const& input, stripe_size_limits max_stripe_size, size_type row_index_stride, writer_timezone const& timezone, + CalendarKind calendar, bool enable_dictionary, bool sort_dictionaries, compression_type compression, @@ -2555,6 +2631,9 @@ auto convert_table_to_orc_data(table_view const& input, [[maybe_unused]] auto const pd_masks = init_pushdown_null_masks(orc_table, stream); auto rowgroup_bounds = calculate_rowgroup_bounds(orc_table, row_index_stride, stream); + if (calendar == CalendarKind::JULIAN_GREGORIAN) { + rebase_date_columns(orc_table, pd_masks, rowgroup_bounds.host_view(), stream); + } [[maybe_unused]] auto const rg_char_counts_data = set_rowgroup_char_counts(orc_table, rowgroup_bounds, stream); @@ -2812,6 +2891,7 @@ void writer::impl::write(table_view const& input) _max_stripe_size, _row_index_stride, _timezone, + *_footer.calendar, _enable_dictionary, _sort_dictionaries, _compression, diff --git a/cpp/tests/io/orc_test.cpp b/cpp/tests/io/orc_test.cpp index eba01b03d404..206bf586b97f 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -2141,13 +2141,215 @@ TEST_F(OrcWriterTest, ProlepticCalendar) } } +TEST_F(OrcWriterTest, LegacyCalendar) +{ + using calendar = cudf::io::orc::CalendarKind; + column_wrapper dates( + {-719528, -719162, -353920, -281237, -141438, -141427, 0, 19849, -141432}, + {true, true, true, true, true, true, true, true, false}); + struct_col nested{dates}; + cudf::test::lists_column_wrapper lists{ + {-719528}, {}, {-353920}, {-281237}, {-141438}, {-141427}, {0}, {19849}, {}}; + table_view input({dates, nested, lists}); + for (auto compression : {cudf::io::compression_type::NONE, cudf::io::compression_type::SNAPPY}) { + for (bool chunked : {false, true}) { + std::vector buffer; + if (chunked) { + cudf::io::orc_chunked_writer writer( + cudf::io::chunked_orc_writer_options::builder(cudf::io::sink_info{&buffer}) + .compression(compression) + .calendar_kind(calendar::JULIAN_GREGORIAN)); + auto chunks = cudf::slice(input, {0, 3, 3, 9}); + writer.write(chunks[0]).write(chunks[1]); + writer.close(); + } else { + cudf::io::write_orc( + cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, input) + .compression(compression) + .calendar_kind(calendar::JULIAN_GREGORIAN)); + } + cudf::io::source_info source{cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}}; + EXPECT_EQ(cudf::io::read_orc_metadata(source).calendar_kind(), calendar::JULIAN_GREGORIAN); + auto result = cudf::io::read_orc(cudf::io::orc_reader_options::builder(source) + .use_np_dtypes(false) + .convert_dates_to_proleptic(true)); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(input, result.tbl->view()); + // Independent Julian epoch-day values, including astronomical year zero. + column_wrapper raw( + {-719530, -719164, -353914, -281230, -141428, -141427, 0, 19849, 0}, + {true, true, true, true, true, true, true, true, false}); + auto raw_result = cudf::io::read_orc( + cudf::io::orc_reader_options::builder(source).columns({"_col0"}).use_np_dtypes(false)); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(raw, raw_result.tbl->view().column(0)); + auto stats = cudf::io::read_parsed_orc_statistics(source); + auto date_stats = + std::get(stats.file_stats[1].type_specific_stats); + EXPECT_EQ(date_stats.minimum, -719530); + EXPECT_EQ(date_stats.maximum, 19849); + } + } +} + +TEST_F(OrcWriterTest, LegacyCalendarRejectsLossyValues) +{ + using calendar = cudf::io::orc::CalendarKind; + std::vector buffer; + auto options = cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, table_view{}) + .calendar_kind(calendar::JULIAN_GREGORIAN) + .build(); + for (int32_t day : std::vector{INT32_MIN, + -141437, + -141436, + -141435, + -141434, + -141433, + -141432, + -141431, + -141430, + -141429, + -141428}) { + column_wrapper dates{day}; + options.set_table(table_view{{dates}}); + EXPECT_THROW(cudf::io::write_orc(options), cudf::logic_error); + EXPECT_TRUE(buffer.empty()); + } + column_wrapper timestamps{0}; + struct_col nested{timestamps}; + options.set_table(table_view{{nested}}); + EXPECT_THROW(cudf::io::write_orc(options), cudf::logic_error); + EXPECT_TRUE(buffer.empty()); +} + +TEST_F(OrcWriterTest, LegacyCalendarSliceIgnoresUnwrittenGap) +{ + column_wrapper dates{-141432, -353920, -141432}; + struct_col nested{dates}; + cudf::test::lists_column_wrapper lists{ + {-141432}, {-353920}, {-141432}}; + auto input = cudf::slice(table_view{{nested, lists}}, {1, 2})[0]; + std::vector buffer; + cudf::io::write_orc(cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, input) + .calendar_kind(cudf::io::orc::CalendarKind::JULIAN_GREGORIAN)); + cudf::io::source_info source{cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}}; + auto result = cudf::io::read_orc( + cudf::io::orc_reader_options::builder(source).use_np_dtypes(false).convert_dates_to_proleptic( + true)); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(input, result.tbl->view()); +} + +TEST_F(OrcWriterTest, LegacyMixedCalendarsAndChunkedRead) +{ + using calendar = cudf::io::orc::CalendarKind; + column_wrapper dates{-353920, -281237, -141438, -141427, 19849}; + table_view input{{dates}}; + std::vector paths; + for (auto kind : {calendar::JULIAN_GREGORIAN, calendar::PROLEPTIC_GREGORIAN}) { + paths.push_back( + temp_env->get_temp_filepath("calendar-" + std::to_string(static_cast(kind)))); + cudf::io::write_orc( + cudf::io::orc_writer_options::builder(cudf::io::sink_info{paths.back()}, input) + .calendar_kind(kind)); + } + auto opts = cudf::io::orc_reader_options::builder(cudf::io::source_info{paths}) + .use_np_dtypes(false) + .convert_dates_to_proleptic(true) + .build(); + auto expected = cudf::concatenate(std::vector{input, input}); + cudf::io::chunked_orc_reader reader(16, 0, 1, opts); + std::vector> chunks; + while (reader.has_next()) { + chunks.push_back(std::move(reader.read_chunk().tbl)); + } + std::vector views; + for (auto const& chunk : chunks) { + views.push_back(chunk->view()); + } + auto result = cudf::concatenate(views); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected->view(), result->view()); + for (bool index : {false, true}) { + opts.enable_use_index(index); + opts.set_skip_rows(1); + opts.set_num_rows(8); + auto subset = cudf::io::read_orc(opts); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(cudf::slice(expected->view(), {1, 9})[0], + subset.tbl->view()); + } +} + +TEST_F(OrcWriterTest, LegacyCalendarRetryAfterInvalidChunk) +{ + column_wrapper good{-353920}; + column_wrapper gap{-141432}; + std::vector buffer; + cudf::io::orc_chunked_writer writer( + cudf::io::chunked_orc_writer_options::builder(cudf::io::sink_info{&buffer}) + .calendar_kind(cudf::io::orc::CalendarKind::JULIAN_GREGORIAN)); + writer.write(table_view{{good}}); + auto const size = buffer.size(); + EXPECT_THROW(writer.write(table_view{{gap}}), cudf::logic_error); + EXPECT_EQ(buffer.size(), size); + writer.write(table_view{{good}}); + writer.close(); + cudf::io::source_info source{cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}}; + auto result = cudf::io::read_orc( + cudf::io::orc_reader_options::builder(source).use_np_dtypes(false).convert_dates_to_proleptic( + true)); + column_wrapper expected{-353920, -353920}; + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected, result.tbl->view().column(0)); +} + +TEST_F(OrcWriterTest, LegacyCalendarIgnoresHiddenListValues) +{ + column_wrapper dates{-141432, -353920, -281237}; + int32_col offsets{0, 1, 2, 3}; + auto validity = std::vector{false, true, true}; + auto [mask, null_count] = cudf::test::detail::make_null_mask(validity.begin(), validity.end()); + auto lists = + cudf::make_lists_column(3, offsets.release(), dates.release(), null_count, std::move(mask)); + std::vector buffer; + cudf::io::write_orc( + cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, table_view{{lists->view()}}) + .calendar_kind(cudf::io::orc::CalendarKind::JULIAN_GREGORIAN)); + cudf::io::source_info source{cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}}; + auto result = cudf::io::read_orc( + cudf::io::orc_reader_options::builder(source).use_np_dtypes(false).convert_dates_to_proleptic( + true)); + cudf::test::lists_column_wrapper expected({{}, {-353920}, {-281237}}, + validity.begin()); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected, result.tbl->view().column(0)); + auto stats = cudf::io::read_parsed_orc_statistics(source); + auto date_stats = std::get(stats.file_stats[2].type_specific_stats); + EXPECT_EQ(date_stats.minimum, -353914); + EXPECT_EQ(date_stats.maximum, -281230); + EXPECT_EQ(stats.file_stats[2].number_of_values, 2); +} + +TEST_F(OrcWriterTest, LegacyCalendarWideDateDomain) +{ + column_wrapper dates{ + INT32_MIN + 100000, -1800000000, -1000000000, -365000000, -10000000, INT32_MAX}; + std::vector buffer; + table_view input{{dates}}; + cudf::io::write_orc(cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, input) + .calendar_kind(cudf::io::orc::CalendarKind::JULIAN_GREGORIAN)); + cudf::io::source_info source{cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}}; + auto result = cudf::io::read_orc( + cudf::io::orc_reader_options::builder(source).use_np_dtypes(false).convert_dates_to_proleptic( + true)); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(input, result.tbl->view()); +} + TEST_F(OrcWriterTest, CalendarRejectsUnsupportedModes) { using calendar = cudf::io::orc::CalendarKind; cudf::io::orc_writer_options opts; cudf::io::chunked_orc_writer_options chunked_opts; - for (auto kind : - {calendar::UNKNOWN_CALENDAR, calendar::JULIAN_GREGORIAN, static_cast(3)}) { + for (auto kind : {calendar::UNKNOWN_CALENDAR, static_cast(3)}) { EXPECT_THROW(opts.set_calendar_kind(kind), cudf::logic_error); EXPECT_THROW(chunked_opts.set_calendar_kind(kind), cudf::logic_error); } @@ -2179,23 +2381,28 @@ TEST_F(OrcMetadataReaderTest, CalendarPresenceAndKinds) TEST_F(OrcWriterTest, CalendarOnEmptyInput) { column_wrapper dates{}; - for (bool chunked : {false, true}) { - std::vector buffer; - table_view empty({dates}); - if (chunked) { - cudf::io::orc_chunked_writer writer( - cudf::io::chunked_orc_writer_options::builder(cudf::io::sink_info{&buffer})); - writer.write(empty); - writer.close(); - } else { - cudf::io::write_orc( - cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, empty)); + for (auto kind : {cudf::io::orc::CalendarKind::PROLEPTIC_GREGORIAN, + cudf::io::orc::CalendarKind::JULIAN_GREGORIAN}) { + for (bool chunked : {false, true}) { + std::vector buffer; + table_view empty({dates}); + if (chunked) { + cudf::io::orc_chunked_writer writer( + cudf::io::chunked_orc_writer_options::builder(cudf::io::sink_info{&buffer}) + .calendar_kind(kind)); + writer.write(empty); + writer.close(); + } else { + cudf::io::write_orc( + cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, empty) + .calendar_kind(kind)); + } + cudf::io::source_info source{cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}}; + auto const meta = cudf::io::read_orc_metadata(source); + EXPECT_EQ(meta.num_rows(), 0); + EXPECT_EQ(meta.calendar_kind(), kind); } - cudf::io::source_info source{cudf::host_span{ - reinterpret_cast(buffer.data()), buffer.size()}}; - auto const meta = cudf::io::read_orc_metadata(source); - EXPECT_EQ(meta.num_rows(), 0); - EXPECT_EQ(meta.calendar_kind(), cudf::io::orc::CalendarKind::PROLEPTIC_GREGORIAN); } } diff --git a/java/src/main/java/ai/rapids/cudf/ORCOptions.java b/java/src/main/java/ai/rapids/cudf/ORCOptions.java index 9fa8d89a23cf..d07b64a8faf6 100644 --- a/java/src/main/java/ai/rapids/cudf/ORCOptions.java +++ b/java/src/main/java/ai/rapids/cudf/ORCOptions.java @@ -1,6 +1,6 @@ /* * - * 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 * */ @@ -19,6 +19,7 @@ public class ORCOptions extends ColumnFilterOptions { public static ORCOptions DEFAULT = new ORCOptions(new Builder()); private final boolean useNumPyTypes; + private final boolean convertDatesToProleptic; private final DType unit; private final String[] decimal128Columns; private final boolean ignoreTimezoneInStripeFooter; @@ -27,6 +28,7 @@ private ORCOptions(Builder builder) { super(builder); decimal128Columns = builder.decimal128Columns.toArray(new String[0]); useNumPyTypes = builder.useNumPyTypes; + convertDatesToProleptic = builder.convertDatesToProleptic; unit = builder.unit; ignoreTimezoneInStripeFooter = builder.ignoreTimezoneInStripeFooter; } @@ -35,6 +37,8 @@ boolean usingNumPyTypes() { return useNumPyTypes; } + boolean convertDatesToProleptic() { return convertDatesToProleptic; } + DType timeUnit() { return unit; } @@ -53,6 +57,7 @@ public static Builder builder() { public static class Builder extends ColumnFilterOptions.Builder { private boolean useNumPyTypes = true; + private boolean convertDatesToProleptic = false; private DType unit = DType.EMPTY; final List decimal128Columns = new ArrayList<>(); @@ -99,6 +104,19 @@ public Builder decimal128Column(String... names) { return this; } + /** + * Convert DATE columns explicitly marked JULIAN_GREGORIAN to cuDF proleptic dates. + * Disabled by default for callers that perform their own rebasing. Missing and UNKNOWN + * calendar fields and TIMESTAMP columns are unchanged. Conversion applies to nested DATE + * columns and precedes NumPy promotion. Unrepresentable dates cause a read error. + * @param convert whether DATE conversion is enabled + * @return builder for chaining + */ + public Builder convertDatesToProleptic(boolean convert) { + this.convertDatesToProleptic = convert; + return this; + } + public Builder ignoreTimezoneInStripeFooter() { this.ignoreTimezoneInStripeFooter = true; return this; diff --git a/java/src/main/java/ai/rapids/cudf/ORCWriterOptions.java b/java/src/main/java/ai/rapids/cudf/ORCWriterOptions.java index 0e1d2efa7906..98d7c47190c0 100644 --- a/java/src/main/java/ai/rapids/cudf/ORCWriterOptions.java +++ b/java/src/main/java/ai/rapids/cudf/ORCWriterOptions.java @@ -89,16 +89,19 @@ public Builder withWriterTimezone(String writerTimezone) { } /** - * Sets the calendar recorded in the file footer. Only {@code PROLEPTIC_GREGORIAN} is - * supported, matching cuDF's date and timestamp values. This does not rebase values. + * Sets the calendar used for DATE encoding and recorded in the file footer. + * Legacy mode rebases nested and top-level DATE values. It rejects TIMESTAMP columns, + * dates in October 5-14, 1582, and dates outside the int32 hybrid domain. + * Use {@link ORCOptions.Builder#convertDatesToProleptic(boolean)} to read legacy output + * back as cuDF proleptic Gregorian dates. * * @param calendarKind calendar kind - * @throws IllegalArgumentException if the calendar is null or is not proleptic Gregorian + * @throws IllegalArgumentException if the calendar is null */ public Builder withCalendarKind(CalendarKind calendarKind) { - if (calendarKind != CalendarKind.PROLEPTIC_GREGORIAN) { + if (calendarKind == null) { throw new IllegalArgumentException( - "Only the proleptic Gregorian calendar is supported for ORC writes"); + "ORC calendar must not be null"); } this.calendarKind = calendarKind; return this; diff --git a/java/src/main/java/ai/rapids/cudf/Table.java b/java/src/main/java/ai/rapids/cudf/Table.java index 659e7593bfd6..0d853314822f 100644 --- a/java/src/main/java/ai/rapids/cudf/Table.java +++ b/java/src/main/java/ai/rapids/cudf/Table.java @@ -342,12 +342,14 @@ private static native long[] readORC(String[] filterColumnNames, String filePath, long address, long length, boolean usingNumPyTypes, int timeUnit, String[] decimal128Columns, - boolean ignoreTimezoneInStripeFooter) throws CudfException; + boolean ignoreTimezoneInStripeFooter, + boolean convertDatesToProleptic) throws CudfException; private static native long[] readORCFromDataSource(String[] filterColumnNames, boolean usingNumPyTypes, int timeUnit, String[] decimal128Columns, boolean ignoreTimezoneInStripeFooter, + boolean convertDatesToProleptic, long dataSourceHandle) throws CudfException; /** @@ -1523,7 +1525,8 @@ public static Table readORC(ORCOptions opts, File path) { return new Table(readORC(opts.getIncludeColumnNames(), path.getAbsolutePath(), 0, 0, opts.usingNumPyTypes(), opts.timeUnit().typeId.getNativeId(), - opts.getDecimal128Columns(), opts.ignoreTimezoneInStripeFooter())); + opts.getDecimal128Columns(), opts.ignoreTimezoneInStripeFooter(), + opts.convertDatesToProleptic())); } /** @@ -1592,7 +1595,8 @@ public static Table readORC(ORCOptions opts, HostMemoryBuffer buffer, return new Table(readORC(opts.getIncludeColumnNames(), null, buffer.getAddress() + offset, len, opts.usingNumPyTypes(), opts.timeUnit().typeId.getNativeId(), - opts.getDecimal128Columns(), opts.ignoreTimezoneInStripeFooter())); + opts.getDecimal128Columns(), opts.ignoreTimezoneInStripeFooter(), + opts.convertDatesToProleptic())); } public static Table readORC(ORCOptions opts, DataSource ds) { @@ -1601,7 +1605,7 @@ public static Table readORC(ORCOptions opts, DataSource ds) { return new Table(readORCFromDataSource(opts.getIncludeColumnNames(), opts.usingNumPyTypes(), opts.timeUnit().typeId.getNativeId(), opts.getDecimal128Columns(), opts.ignoreTimezoneInStripeFooter(), - dataSourceHandle)); + opts.convertDatesToProleptic(), dataSourceHandle)); } finally { DataSourceHelper.destroyWrapperDataSource(dataSourceHandle); } diff --git a/java/src/main/native/src/TableJni.cpp b/java/src/main/native/src/TableJni.cpp index 942e66c2e763..f0c202473640 100644 --- a/java/src/main/native/src/TableJni.cpp +++ b/java/src/main/native/src/TableJni.cpp @@ -2493,6 +2493,7 @@ Java_ai_rapids_cudf_Table_readORCFromDataSource(JNIEnv* env, jint unit, jobjectArray dec128_col_names, jboolean ignore_timezone_in_stripe_footer, + jboolean convert_dates_to_proleptic, jlong ds_handle) { JNI_NULL_CHECK(env, ds_handle, "no data source handle given", 0); @@ -2519,6 +2520,7 @@ Java_ai_rapids_cudf_Table_readORCFromDataSource(JNIEnv* env, .timestamp_type(cudf::data_type(static_cast(unit))) .decimal128_columns(n_dec128_col_names.as_cpp_vector()) .ignore_timezone_in_stripe_footer(ignore_timezone_in_stripe_footer) + .convert_dates_to_proleptic(convert_dates_to_proleptic) .build(); return convert_table_for_return(env, cudf::io::read_orc(opts).tbl); } @@ -2535,7 +2537,8 @@ Java_ai_rapids_cudf_Table_readORC(JNIEnv* env, jboolean usingNumPyTypes, jint unit, jobjectArray dec128_col_names, - jboolean ignore_timezone_in_stripe_footer) + jboolean ignore_timezone_in_stripe_footer, + jboolean convert_dates_to_proleptic) { bool read_buffer = true; if (buffer == 0) { @@ -2579,6 +2582,7 @@ Java_ai_rapids_cudf_Table_readORC(JNIEnv* env, .timestamp_type(cudf::data_type(static_cast(unit))) .decimal128_columns(n_dec128_col_names.as_cpp_vector()) .ignore_timezone_in_stripe_footer(ignore_timezone_in_stripe_footer) + .convert_dates_to_proleptic(convert_dates_to_proleptic) .build(); return convert_table_for_return(env, cudf::io::read_orc(opts).tbl); } diff --git a/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java b/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java index 6f2a31393305..92cf961a4991 100644 --- a/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java +++ b/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java @@ -43,24 +43,146 @@ void rejectsUnsupportedCalendars() { ORCWriterOptions.builder().build().getCalendarKind()); assertThrows(IllegalArgumentException.class, () -> ORCWriterOptions.builder().withCalendarKind(null)); - assertThrows(IllegalArgumentException.class, () -> ORCWriterOptions.builder() - .withCalendarKind(ORCWriterOptions.CalendarKind.JULIAN_GREGORIAN)); + assertEquals(ORCWriterOptions.CalendarKind.JULIAN_GREGORIAN, ORCWriterOptions.builder() + .withCalendarKind(ORCWriterOptions.CalendarKind.JULIAN_GREGORIAN).build().getCalendarKind()); } @Test @SuppressWarnings("unchecked") void datesInteroperateWithApacheOrc() throws IOException { + checkDates(DAYS, false); + } + + @Test + void legacyDatesInteroperateWithApacheOrc() throws IOException { + Integer[] days = Arrays.stream(new String[] { + "0001-01-01", "0100-03-01", "0400-02-29", "1001-01-01", "1200-01-01", "1500-02-28", + "1500-03-01", "1582-10-04", "1582-10-15", "1970-01-01", "2024-05-06", null + }).map(value -> value == null ? null : Math.toIntExact(LocalDate.parse(value).toEpochDay())) + .toArray(Integer[]::new); + checkDates(days, true); + } + + @Test + void apacheLegacyDatesReadAsProleptic() throws IOException { + // Independent Apache ORC conversion provides the encoded day oracle, including BCE dates + // and a deterministic wide sample within Apache HybridChronology's negative-year range. + java.util.Random random = new java.util.Random(23794); + Integer[] proleptic = new Integer[4096]; + Integer[] hybrid = new Integer[proleptic.length]; + // Hive 2.8.1's vector calendar conversion uses SimpleDateFormat without an era field. + // Check BCE encoding against Apache ORC's DateUtils directly, which preserves the era. + int[] anchors = Arrays.stream(new String[] { + "-10000-02-28", "-0400-02-29", "-0001-03-01", "0000-02-29", "0001-01-01" + }).mapToInt(value -> Math.toIntExact(LocalDate.parse(value).toEpochDay())).toArray(); + for (int i = 0; i < proleptic.length; ++i) { + int day = i < anchors.length ? anchors[i] : random.nextInt(); + if (day < -365000000) { day %= 365000000; } + if (day < Integer.MIN_VALUE + 100000 || (day >= -141437 && day < -141427)) { + day = 0; + } + proleptic[i] = day; + hybrid[i] = org.apache.orc.impl.DateUtils.convertDateToHybrid(day); + } + java.nio.file.Path file = Files.createTempFile("apache-legacy", ".orc"); + try { + org.apache.orc.TypeDescription schema = org.apache.orc.TypeDescription.fromString("struct"); + try (org.apache.orc.Writer writer = OrcFile.createWriter(new Path(file.toUri()), + OrcFile.writerOptions(new Configuration()).overwrite(true).setSchema(schema).setProlepticGregorian(false) + .compress(org.apache.orc.CompressionKind.NONE))) { + VectorizedRowBatch batch = schema.createRowBatch(); + for (int i = 0; i < hybrid.length; ++i) { + ((LongColumnVector) batch.cols[0]).vector[batch.size++] = hybrid[i]; + if (batch.size == batch.getMaxSize()) { + writer.addRowBatch(batch); + batch.reset(); + } + } + if (batch.size > 0) { writer.addRowBatch(batch); } + } + try (ColumnVector dates = ColumnVector.timestampDaysFromBoxedInts(proleptic); + ColumnVector encoded = ColumnVector.timestampDaysFromBoxedInts(hybrid); + Table expected = new Table(dates); + Table rawExpected = new Table(encoded); + Table actual = Table.readORC(ORCOptions.builder().withNumPyTypes(false) + .convertDatesToProleptic(true).build(), file.toFile()); + Table raw = Table.readORC(ORCOptions.builder().withNumPyTypes(false).build(), file.toFile())) { + assertTablesAreEqual(expected, actual); + assertTablesAreEqual(rawExpected, raw); + java.nio.file.Path gpuFile = Files.createTempFile("gpu-legacy", ".orc"); + try { + try (TableWriter writer = Table.writeORCChunked(ORCWriterOptions.builder() + .withColumns(true, "d").withCalendarKind(ORCWriterOptions.CalendarKind.JULIAN_GREGORIAN) + .build(), gpuFile.toFile())) { + writer.write(expected); + } + try (Table gpuRaw = Table.readORC(ORCOptions.builder().withNumPyTypes(false).build(), + gpuFile.toFile())) { + assertTablesAreEqual(rawExpected, gpuRaw); + } + } finally { Files.deleteIfExists(gpuFile); } + } + byte[] originalBytes = Files.readAllBytes(file); + for (org.apache.orc.OrcProto.CalendarKind calendar : new org.apache.orc.OrcProto.CalendarKind[] { + null, org.apache.orc.OrcProto.CalendarKind.UNKNOWN_CALENDAR, + org.apache.orc.OrcProto.CalendarKind.PROLEPTIC_GREGORIAN, + org.apache.orc.OrcProto.CalendarKind.JULIAN_GREGORIAN}) { + byte[] variant = withCalendar(originalBytes, calendar); + try (ColumnVector dates = ColumnVector.timestampDaysFromBoxedInts( + calendar == org.apache.orc.OrcProto.CalendarKind.JULIAN_GREGORIAN ? proleptic : hybrid); + Table expected = new Table(dates); + Table actual = Table.readORC(ORCOptions.builder().withNumPyTypes(false) + .convertDatesToProleptic(true).build(), variant)) { + assertTablesAreEqual(expected, actual); + } + } + Long[] millis = Arrays.stream(proleptic).map(day -> day.longValue() * 86400000L) + .toArray(Long[]::new); + try (ColumnVector dates = ColumnVector.timestampMilliSecondsFromBoxedLongs(millis); + Table expected = new Table(dates); + Table actual = Table.readORC(ORCOptions.builder().convertDatesToProleptic(true).build(), + Files.readAllBytes(file))) { + assertTablesAreEqual(expected, actual); + } + } finally { Files.deleteIfExists(file); } + } + + @Test + void rejectsJulianOnlyLeapDayOnConversion() throws IOException { + int day = org.apache.orc.impl.DateUtils.parseDate("1500-02-29", false); + java.nio.file.Path file = Files.createTempFile("apache-julian-leap", ".orc"); + try { + org.apache.orc.TypeDescription schema = org.apache.orc.TypeDescription.fromString("struct"); + try (org.apache.orc.Writer writer = OrcFile.createWriter(new Path(file.toUri()), + OrcFile.writerOptions(new Configuration()).overwrite(true).setSchema(schema) + .setProlepticGregorian(false))) { + VectorizedRowBatch batch = schema.createRowBatch(); + ((LongColumnVector) batch.cols[0]).vector[0] = day; + batch.size = 1; + writer.addRowBatch(batch); + } + try (Table raw = Table.readORC(ORCOptions.builder().withNumPyTypes(false).build(), file.toFile())) { + assertEquals(1, raw.getRowCount()); + } + assertThrows(CudfException.class, () -> Table.readORC(ORCOptions.builder() + .convertDatesToProleptic(true).build(), file.toFile())); + } finally { Files.deleteIfExists(file); } + } + + @SuppressWarnings("unchecked") + private void checkDates(Integer[] days, boolean legacy) throws IOException { HostColumnVector.ListType listType = new HostColumnVector.ListType(true, new HostColumnVector.BasicType(true, DType.TIMESTAMP_DAYS)); - try (ColumnVector dates = ColumnVector.timestampDaysFromBoxedInts(DAYS); + try (ColumnVector dates = ColumnVector.timestampDaysFromBoxedInts(days); ColumnVector structs = ColumnVector.makeStruct(dates); ColumnVector lists = ColumnVector.fromLists(listType, - Arrays.stream(DAYS).map(Collections::singletonList).toArray(java.util.List[]::new)); + Arrays.stream(days).map(Collections::singletonList).toArray(java.util.List[]::new)); Table input = new Table(dates, structs, lists)) { for (CompressionType compression : new CompressionType[] { CompressionType.NONE, CompressionType.SNAPPY}) { for (boolean bufferSink : new boolean[] {false, true}) { for (boolean explicitCalendar : new boolean[] {false, true}) { + if (legacy && !explicitCalendar) { continue; } java.nio.file.Path file = Files.createTempFile("orc-calendar", ".orc"); try { ORCWriterOptions.Builder builder = ORCWriterOptions.builder() @@ -70,8 +192,9 @@ void datesInteroperateWithApacheOrc() throws IOException { .withListColumn(ColumnWriterOptions.listBuilder("l") .withColumns(true, "d").build()) .withCompressionType(compression); - if (explicitCalendar) { - builder.withCalendarKind(ORCWriterOptions.CalendarKind.PROLEPTIC_GREGORIAN); + if (legacy || explicitCalendar) { + builder.withCalendarKind(legacy ? ORCWriterOptions.CalendarKind.JULIAN_GREGORIAN + : ORCWriterOptions.CalendarKind.PROLEPTIC_GREGORIAN); } ORCWriterOptions options = builder.build(); if (bufferSink) { @@ -98,7 +221,7 @@ void datesInteroperateWithApacheOrc() throws IOException { } try (Table expected = Table.concatenate(input, input); Table actual = Table.readORC( - ORCOptions.builder().withNumPyTypes(false).build(), file.toFile())) { + ORCOptions.builder().withNumPyTypes(false).convertDatesToProleptic(true).build(), file.toFile())) { assertTablesAreEqual(expected, actual); } // Both missing-field defaults must be overridden by the file's explicit calendar. @@ -107,15 +230,16 @@ void datesInteroperateWithApacheOrc() throws IOException { conf.setBoolean("orc.proleptic.gregorian.default", defaultProleptic); try (Reader reader = OrcFile.createReader(new Path(file.toUri()), OrcFile.readerOptions(conf).convertToProlepticGregorian(true))) { - assertTrue(reader.writerUsedProlepticGregorian()); - assertEquals(DAYS.length * 2, reader.getNumberOfRows()); + assertEquals(!legacy, reader.writerUsedProlepticGregorian()); + assertEquals(days.length * 2, reader.getNumberOfRows()); // DATE leaves in the top-level column, STRUCT, and LIST. for (int columnId : new int[] {1, 3, 5}) { DateColumnStatistics stats = (DateColumnStatistics) reader.getStatistics()[columnId]; - assertEquals(DAYS[0].longValue(), stats.getMinimumDayOfEpoch()); - assertEquals(DAYS[6].longValue(), stats.getMaximumDayOfEpoch()); - assertEquals((DAYS.length - 1) * 2, stats.getNumberOfValues()); + assertEquals(days[0].longValue(), stats.getMinimumDayOfEpoch()); + assertEquals(Arrays.stream(days).filter(java.util.Objects::nonNull).mapToLong(Integer::longValue) + .max().getAsLong(), stats.getMaximumDayOfEpoch()); + assertEquals((days.length - 1) * 2, stats.getNumberOfValues()); } VectorizedRowBatch batch = reader.getSchema().createRowBatch(); int rowCount = 0; @@ -127,7 +251,7 @@ void datesInteroperateWithApacheOrc() throws IOException { ListColumnVector list = (ListColumnVector) batch.cols[2]; LongColumnVector child = (LongColumnVector) list.child; for (int row = 0; row < batch.size; row++, rowCount++) { - Integer expected = DAYS[rowCount % DAYS.length]; + Integer expected = days[rowCount % days.length]; assertDate(top, row, expected); assertDate(nested, row, expected); assertEquals(1, list.lengths[row]); @@ -135,7 +259,7 @@ void datesInteroperateWithApacheOrc() throws IOException { } } } - assertEquals(DAYS.length * 2, rowCount); + assertEquals(days.length * 2, rowCount); } } } finally { @@ -147,6 +271,27 @@ void datesInteroperateWithApacheOrc() throws IOException { } } + // Rewrite only the uncompressed footer using Apache's independently generated protobuf classes. + private static byte[] withCalendar(byte[] file, org.apache.orc.OrcProto.CalendarKind calendar) + throws IOException { + int postscriptLength = Byte.toUnsignedInt(file[file.length - 1]); + int postscriptStart = file.length - 1 - postscriptLength; + org.apache.orc.OrcProto.PostScript ps = org.apache.orc.OrcProto.PostScript.parseFrom( + Arrays.copyOfRange(file, postscriptStart, file.length - 1)); + int footerStart = postscriptStart - Math.toIntExact(ps.getFooterLength()); + org.apache.orc.OrcProto.Footer.Builder footer = org.apache.orc.OrcProto.Footer.parseFrom( + Arrays.copyOfRange(file, footerStart, postscriptStart)).toBuilder(); + if (calendar == null) { footer.clearCalendar(); } else { footer.setCalendar(calendar); } + byte[] footerBytes = footer.build().toByteArray(); + byte[] psBytes = ps.toBuilder().setFooterLength(footerBytes.length).build().toByteArray(); + ByteArrayOutputStream output = new ByteArrayOutputStream(); + output.write(file, 0, footerStart); + output.write(footerBytes); + output.write(psBytes); + output.write(psBytes.length); + return output.toByteArray(); + } + private static void assertDate(LongColumnVector column, int row, Integer expected) { int index = column.isRepeating ? 0 : row; if (expected == null) { From db4aedbfc29abcde09f87c60d99f9bd41f8390fe Mon Sep 17 00:00:00 2001 From: Allen Xu Date: Wed, 21 Oct 2026 15:49:53 +0800 Subject: [PATCH 3/5] Address ORC calendar review findings Preserve unknown protobuf calendar values and make shared decode errors sticky across columns. Bound legacy DATE scratch memory to written slices, skip empty nested ranges, and retain compacted column ownership through statistics persistence. Compact top-level DATE slices before device-view construction to avoid duplicate setup in chunked writes. Cover raw enum values, decoder overflow state, bounded slice allocations, nested DATE statistics, and indexed reading. Clarify calendar constants and parameter documentation. Contributes to NVIDIA/cudf#23794. Signed-off-by: Allen Xu --- cpp/include/cudf/io/orc_types.hpp | 2 +- cpp/src/io/orc/calendar.cuh | 23 ++-- cpp/src/io/orc/decode_error.cuh | 26 +++++ cpp/src/io/orc/stripe_data.cu | 7 +- cpp/src/io/orc/writer_impl.cu | 61 ++++++++++- cpp/tests/io/orc_chunked_reader_test.cu | 24 +++++ cpp/tests/io/orc_test.cpp | 100 +++++++++++++++++- .../java/ai/rapids/cudf/ORCCalendarTest.java | 29 ++++- 8 files changed, 250 insertions(+), 22 deletions(-) create mode 100644 cpp/src/io/orc/decode_error.cuh diff --git a/cpp/include/cudf/io/orc_types.hpp b/cpp/include/cudf/io/orc_types.hpp index f21f312165b0..89ccc7c4c834 100644 --- a/cpp/include/cudf/io/orc_types.hpp +++ b/cpp/include/cudf/io/orc_types.hpp @@ -36,7 +36,7 @@ enum CompressionKind : uint8_t { /** * @brief Identifies the calendar used to encode dates and timestamps in an ORC file. */ -enum class CalendarKind : uint8_t { +enum class CalendarKind : uint32_t { UNKNOWN_CALENDAR = 0, ///< Calendar is unspecified JULIAN_GREGORIAN = 1, ///< Hybrid calendar with the October 1582 cutover PROLEPTIC_GREGORIAN = 2, ///< Gregorian calendar extended to all dates diff --git a/cpp/src/io/orc/calendar.cuh b/cpp/src/io/orc/calendar.cuh index 284c7048e254..85c1a10a6754 100644 --- a/cpp/src/io/orc/calendar.cuh +++ b/cpp/src/io/orc/calendar.cuh @@ -5,14 +5,14 @@ #pragma once -#include +#include #include namespace cudf::io::orc::detail { // Epoch days of 1582-10-15. Earlier hybrid dates use the Julian calendar. -constexpr int32_t calendar_cutover = -141427; +constexpr int32_t calendar_cutover = -141'427; CUDF_HOST_DEVICE constexpr int64_t floor_div(int64_t value, int64_t divisor) { @@ -31,21 +31,26 @@ CUDF_HOST_DEVICE constexpr int64_t rebase_days(int64_t days) { if (days >= calendar_cutover) { return days; } - auto const z = days + (to_hybrid ? 719468 : 719470); - auto const era = floor_div(z, to_hybrid ? 146097 : 1461); - auto const doe = z - era * (to_hybrid ? 146097 : 1461); - auto const yoe = to_hybrid ? (doe - doe / 1460 + doe / 36524 - doe / 146096) / 365 - : (doe == 1460 ? 3 : doe / 365); + constexpr int64_t gregorian_epoch_offset = 719'468; + constexpr int64_t julian_epoch_offset = 719'470; + constexpr int64_t gregorian_era_days = 146'097; + constexpr int64_t julian_era_days = 1'461; + constexpr auto source_era_days = to_hybrid ? gregorian_era_days : julian_era_days; + auto const z = days + (to_hybrid ? gregorian_epoch_offset : julian_epoch_offset); + auto const era = floor_div(z, source_era_days); + auto const doe = z - era * source_era_days; + auto const yoe = to_hybrid ? (doe - doe / 1'460 + doe / 36'524 - doe / 146'096) / 365 + : (doe == 1'460 ? 3 : doe / 365); auto const year = yoe + era * (to_hybrid ? 400 : 4); auto const doy = doe - (365 * yoe + (to_hybrid ? yoe / 4 - yoe / 100 : 0)); auto const mp = (5 * doy + 2) / 153; auto const day = doy - (153 * mp + 2) / 5 + 1; if constexpr (to_hybrid) { - return 365 * year + floor_div(year, 4) + (153 * mp + 2) / 5 + day - 1 - 719470; + return 365 * year + floor_div(year, 4) + (153 * mp + 2) / 5 + day - 1 - julian_epoch_offset; } else { return 365 * year + floor_div(year, 4) - floor_div(year, 100) + floor_div(year, 400) + - (153 * mp + 2) / 5 + day - 1 - 719468; + (153 * mp + 2) / 5 + day - 1 - gregorian_epoch_offset; } } diff --git a/cpp/src/io/orc/decode_error.cuh b/cpp/src/io/orc/decode_error.cuh new file mode 100644 index 000000000000..d5415d1e764b --- /dev/null +++ b/cpp/src/io/orc/decode_error.cuh @@ -0,0 +1,26 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include + +#include + +namespace cudf::io::orc::detail { + +/** + * @brief Record a decoding failure without a shared error count wrapping back to zero. + * + * The reader checks for any error after synchronizing the decoding stream. Columns and + * stripes share the same status, so their combined number of failures is not row-count bounded. + */ +inline __device__ void set_decode_error(size_type* status) +{ + cuda::atomic_ref error{*status}; + error.fetch_or(1, cuda::std::memory_order_relaxed); +} + +} // namespace cudf::io::orc::detail diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index a59a3c85b24c..a859a0ea304a 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -4,6 +4,7 @@ */ #include "calendar.cuh" +#include "decode_error.cuh" #include "io/utilities/block_utils.cuh" #include "io/utilities/column_buffer.hpp" #include "orc_gpu.hpp" @@ -1668,10 +1669,10 @@ CUDF_KERNEL void __launch_bounds__(block_size) // If we have an index, seek to the initial run and update row positions if (num_rowgroups > 0) { if (s->top.data.index.strm_offset[0] > s->chunk.strm_len[CI_DATA]) { - atomicAdd(error_count, 1); + set_decode_error(error_count); } if (s->top.data.index.strm_offset[1] > s->chunk.strm_len[CI_DATA2]) { - atomicAdd(error_count, 1); + set_decode_error(error_count); } auto const ofs0 = min(s->top.data.index.strm_offset[0], s->chunk.strm_len[CI_DATA]); auto const ofs1 = min(s->top.data.index.strm_offset[1], s->chunk.strm_len[CI_DATA2]); @@ -2006,7 +2007,7 @@ CUDF_KERNEL void __launch_bounds__(block_size) s->chunk.convert_dates_to_proleptic ? rebase_days(original) : original; if (s->chunk.convert_dates_to_proleptic and not is_representable_date(original, rebased)) { - atomicAdd(error_count, 1); + set_decode_error(error_count); } if (s->chunk.dtype_len == 8) { cudf::duration_D days{static_cast(rebased)}; diff --git a/cpp/src/io/orc/writer_impl.cu b/cpp/src/io/orc/writer_impl.cu index eed92b5f654c..30c700a9a796 100644 --- a/cpp/src/io/orc/writer_impl.cu +++ b/cpp/src/io/orc/writer_impl.cu @@ -14,6 +14,7 @@ #include "io/statistics/column_statistics.cuh" #include "writer_impl.hpp" +#include #include #include #include @@ -330,6 +331,8 @@ struct orc_table_view { cuda::device_buffer d_string_column_indices; std::vector> rebased_dates; std::vector> rebased_date_masks; + // Keep selectively compacted input alive through statistics persistence and sink writes. + std::vector> compacted_columns; [[nodiscard]] auto num_columns() const noexcept { return columns.size(); } [[nodiscard]] size_type num_rows() const noexcept @@ -2170,7 +2173,8 @@ void rebase_date_columns(orc_table_view& orc_table, auto const& view = col.view(); auto const begin = bounds[0][col.index()].begin; auto const end = bounds[bounds.size().first - 1][col.index()].end; - auto& data = orc_table.rebased_dates.emplace_back( + if (begin == end) { continue; } + auto& data = orc_table.rebased_dates.emplace_back( static_cast(view.offset()) + view.size(), stream); auto const columns = device_span{orc_table.d_columns}; auto const parent_mask = col.is_child() ? pd_masks.masks[col.parent_index()] : nullptr; @@ -2596,6 +2600,7 @@ struct stripe_stream_size_less { * @param max_stripe_size Maximum size of stripes in the output file * @param row_index_stride The row index stride * @param timezone Timezone that the written timestamps are relative to + * @param calendar Calendar used for DATE encoding and statistics * @param enable_dictionary Whether dictionary is enabled * @param sort_dictionaries Whether to sort the dictionaries * @param compression The compression format @@ -2623,14 +2628,62 @@ auto convert_table_to_orc_data(table_view const& input, data_sink const& out_sink, cuda::stream_ref stream) { - auto const input_tview = table_device_view::create(input, stream); + std::vector> compacted_columns; + auto input_view = input; + if (calendar == CalendarKind::JULIAN_GREGORIAN and input.num_rows() != 0) { + // Top-level DATE offsets are already known. Compact them before constructing device + // views so the common chunked-write path only builds its views and rowgroup bounds once. + std::vector views(input.begin(), input.end()); + for (size_type index = 0; index < input.num_columns(); ++index) { + auto const& col = input.column(index); + if (col.type().id() != type_id::TIMESTAMP_DAYS or col.offset() == 0) { continue; } + auto& compacted = compacted_columns.emplace_back( + std::make_unique(col, stream, cudf::get_current_device_resource_ref())); + views[index] = compacted->view(); + } + if (not compacted_columns.empty()) { input_view = table_view{views}; } + } + auto input_tview = table_device_view::create(input_view, stream); + + auto orc_table = make_orc_table_view(input_view, *input_tview, table_meta, stream); - auto orc_table = make_orc_table_view(input, *input_tview, table_meta, stream); + auto rowgroup_bounds = calculate_rowgroup_bounds(orc_table, row_index_stride, stream); + if (calendar == CalendarKind::JULIAN_GREGORIAN and orc_table.num_rows() != 0) { + // A DATE view retains its backing offset/extent even when only a small nested slice is + // written. Compact just the affected top-level columns before allocating rebased buffers. + std::vector compact(input.num_columns(), false); + size_type root_index = -1; + for (auto const& col : orc_table.columns) { + if (not col.is_child()) { ++root_index; } + if (col.orc_kind() != DATE) { continue; } + auto const begin = rowgroup_bounds[0][col.index()].begin; + auto const end = rowgroup_bounds[rowgroup_bounds.size().first - 1][col.index()].end; + if (begin != end and (col.view().offset() != 0 or begin != 0 or end != col.size())) { + compact[root_index] = true; + } + } + if (std::any_of(compact.begin(), compact.end(), [](bool value) { return value; })) { + std::vector views; + for (size_type index = 0; index < input.num_columns(); ++index) { + if (compact[index]) { + auto& compacted = compacted_columns.emplace_back(std::make_unique( + input_view.column(index), stream, cudf::get_current_device_resource_ref())); + views.push_back(compacted->view()); + } else { + views.push_back(input_view.column(index)); + } + } + table_view compacted_input{views}; + input_tview = table_device_view::create(compacted_input, stream); + orc_table = make_orc_table_view(compacted_input, *input_tview, table_meta, stream); + rowgroup_bounds = calculate_rowgroup_bounds(orc_table, row_index_stride, stream); + } + } + orc_table.compacted_columns = std::move(compacted_columns); // This is unused but it holds memory buffers for later access thus needs to be kept alive. [[maybe_unused]] auto const pd_masks = init_pushdown_null_masks(orc_table, stream); - auto rowgroup_bounds = calculate_rowgroup_bounds(orc_table, row_index_stride, stream); if (calendar == CalendarKind::JULIAN_GREGORIAN) { rebase_date_columns(orc_table, pd_masks, rowgroup_bounds.host_view(), stream); } diff --git a/cpp/tests/io/orc_chunked_reader_test.cu b/cpp/tests/io/orc_chunked_reader_test.cu index e46bd1cbaa92..a313c4055896 100644 --- a/cpp/tests/io/orc_chunked_reader_test.cu +++ b/cpp/tests/io/orc_chunked_reader_test.cu @@ -4,6 +4,7 @@ */ #include "compression_common.hpp" +#include "io/orc/decode_error.cuh" #include #include @@ -20,6 +21,7 @@ #include #include #include +#include #include #include #include @@ -31,6 +33,7 @@ #include #include +#include #include #include @@ -192,10 +195,31 @@ void expect_selected_nested_empty_struct_chunked_round_trip( expect_selected_nested_empty_struct_table(expected->view(), chunk.tbl->view()); } +void record_two_decode_errors(cudf::size_type* status, cuda::stream_ref stream) +{ + cudf::detail::device_single_thread( + [status] __device__ { + cudf::io::orc::detail::set_decode_error(status); + cudf::io::orc::detail::set_decode_error(status); + }, + stream); +} + } // namespace struct OrcChunkedReaderTest : public cudf::test::BaseFixture {}; +TEST_F(OrcChunkedReaderTest, CalendarDecodeErrorDoesNotWrap) +{ + auto const stream = cudf::get_default_stream(); + for (auto initial : {cudf::size_type{0}, cudf::size_type{-2}}) { + rmm::device_scalar status{initial, stream}; + // -2 represents the state after 2^32 - 2 failures. Two more increments used to clear it. + record_two_decode_errors(status.data(), stream); + EXPECT_NE(status.value(stream), 0); + } +} + using OrcChunkedDecompressionTest = DecompressionTest; TEST_F(OrcChunkedReaderTest, TestChunkedReadNoData) diff --git a/cpp/tests/io/orc_test.cpp b/cpp/tests/io/orc_test.cpp index 206bf586b97f..20f2a414ed4c 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -2237,6 +2237,87 @@ TEST_F(OrcWriterTest, LegacyCalendarSliceIgnoresUnwrittenGap) cudf::io::orc_reader_options::builder(source).use_np_dtypes(false).convert_dates_to_proleptic( true)); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(input, result.tbl->view()); + auto const stats = cudf::io::read_parsed_orc_statistics(source); + for (auto column_id : {2, 4}) { + auto const date_stats = + std::get(stats.file_stats[column_id].type_specific_stats); + EXPECT_EQ(date_stats.minimum, -353'914); + EXPECT_EQ(date_stats.maximum, -353'914); + EXPECT_EQ(stats.file_stats[column_id].number_of_values, 1); + } +} + +TEST_F(OrcWriterTest, LegacyCalendarSliceAllocation) +{ + // All backing columns are allocated before installing the write-only allocation budget. + // An offset-sized DATE scratch buffer would request 4 MB for a one-row slice or empty list. + constexpr cudf::size_type rows = 1'000'000; + std::vector values(rows, -353'920); + std::vector text(rows, "x"); + column_wrapper dates(values.begin(), values.end()); + str_col strings(text.begin(), text.end()); + struct_col nested{dates, strings}; + auto const write_with_budget = [](table_view const& input) { + rmm::mr::cuda_memory_resource upstream; + rmm::mr::callback_memory_resource mr{ + [&](cuda::stream_ref stream, std::size_t bytes, std::size_t alignment, void*) -> void* { + if (bytes > 1'048'576) { throw rmm::bad_alloc{"ORC slice allocation exceeds budget"}; } + return upstream.allocate(stream, bytes, alignment); + }, + [&](cuda::stream_ref stream, void* ptr, std::size_t bytes, std::size_t alignment, void*) { + upstream.deallocate(stream, ptr, bytes, alignment); + }}; + std::vector buffer; + { + cudf::test::scoped_current_device_resource scope{mr}; + cudf::io::write_orc(cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, input) + .compression(cudf::io::compression_type::NONE) + .calendar_kind(cudf::io::orc::CalendarKind::JULIAN_GREGORIAN)); + } + return buffer; + }; + auto const check = [&](table_view const& input) { + auto buffer = write_with_budget(input); + cudf::io::source_info source{cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}}; + auto result = cudf::io::read_orc( + cudf::io::orc_reader_options::builder(source).use_np_dtypes(false).convert_dates_to_proleptic( + true)); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(input, result.tbl->view()); + }; + // The compacted STRUCT also owns a STRING child whose statistics outlive encoding. + check(cudf::slice(table_view{{dates, nested, strings}}, {rows - 1, rows})[0]); + int32_col offsets{0, 0, rows - 1, rows}; + auto lists = cudf::make_lists_column(3, + offsets.release(), + dates.release(), + 0, + cudf::create_null_mask(0, cudf::mask_state::UNALLOCATED)); + check(cudf::slice(table_view{{lists->view()}}, {0, 1})[0]); + check(cudf::slice(table_view{{lists->view()}}, {2, 3})[0]); +} + +TEST_F(OrcWriterTest, LegacyCalendarIndexedRead) +{ + column_wrapper dates{-353'920, -281'237, -141'438, -141'427, 19'849}; + table_view seed{{dates}}; + auto input = cudf::concatenate(std::vector(129, seed)); + std::vector buffer; + cudf::io::write_orc( + cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, input->view()) + .row_index_stride(512) + .calendar_kind(cudf::io::orc::CalendarKind::JULIAN_GREGORIAN)); + cudf::io::source_info source{cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}}; + auto opts = cudf::io::orc_reader_options::builder(source) + .use_np_dtypes(false) + .convert_dates_to_proleptic(true) + .build(); + for (bool index : {false, true}) { + opts.enable_use_index(index); + auto result = cudf::io::read_orc(opts); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(input->view(), result.tbl->view()); + } } TEST_F(OrcWriterTest, LegacyMixedCalendarsAndChunkedRead) @@ -2349,7 +2430,11 @@ TEST_F(OrcWriterTest, CalendarRejectsUnsupportedModes) using calendar = cudf::io::orc::CalendarKind; cudf::io::orc_writer_options opts; cudf::io::chunked_orc_writer_options chunked_opts; - for (auto kind : {calendar::UNKNOWN_CALENDAR, static_cast(3)}) { + for (auto kind : {calendar::UNKNOWN_CALENDAR, + static_cast(3), + static_cast(257), + static_cast(258), + static_cast(UINT32_MAX)}) { EXPECT_THROW(opts.set_calendar_kind(kind), cudf::logic_error); EXPECT_THROW(chunked_opts.set_calendar_kind(kind), cudf::logic_error); } @@ -2363,12 +2448,21 @@ TEST_F(OrcMetadataReaderTest, CalendarPresenceAndKinds) for (auto kind : std::vector>{std::nullopt, calendar::UNKNOWN_CALENDAR, calendar::JULIAN_GREGORIAN, - calendar::PROLEPTIC_GREGORIAN}) { + calendar::PROLEPTIC_GREGORIAN, + static_cast(3), + static_cast(257), + static_cast(258), + static_cast(UINT32_MAX)}) { std::vector buffer{ 'O', 'R', 'C', 0x08, 0x03, 0x10, 0x03, 0x22, 0x02, 0x08, 0x0c, 0x30, 0x00}; if (kind) { buffer.push_back(0x58); // Footer.calendar, protobuf field 11, varint - buffer.push_back(static_cast(*kind)); + auto value = static_cast(*kind); + while (value > 0x7f) { + buffer.push_back(static_cast((value & 0x7f) | 0x80)); + value >>= 7; + } + buffer.push_back(static_cast(value)); } auto const footer_size = static_cast(buffer.size() - 3); buffer.insert(buffer.end(), {0x08, footer_size, 0x10, 0x00, 0x04}); diff --git a/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java b/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java index 92cf961a4991..88961ef3cd74 100644 --- a/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java +++ b/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java @@ -136,6 +136,16 @@ void apacheLegacyDatesReadAsProleptic() throws IOException { assertTablesAreEqual(expected, actual); } } + // Unknown protobuf enum numbers must not alias recognized calendars after narrowing. + for (int calendar : new int[] {3, 257, 258, -1}) { + byte[] variant = withCalendarValue(originalBytes, calendar); + try (ColumnVector dates = ColumnVector.timestampDaysFromBoxedInts(hybrid); + Table expected = new Table(dates); + Table actual = Table.readORC(ORCOptions.builder().withNumPyTypes(false) + .convertDatesToProleptic(true).build(), variant)) { + assertTablesAreEqual(expected, actual); + } + } Long[] millis = Arrays.stream(proleptic).map(day -> day.longValue() * 86400000L) .toArray(Long[]::new); try (ColumnVector dates = ColumnVector.timestampMilliSecondsFromBoxedLongs(millis); @@ -274,6 +284,10 @@ private void checkDates(Integer[] days, boolean legacy) throws IOException { // Rewrite only the uncompressed footer using Apache's independently generated protobuf classes. private static byte[] withCalendar(byte[] file, org.apache.orc.OrcProto.CalendarKind calendar) throws IOException { + return withCalendarValue(file, calendar == null ? null : calendar.getNumber()); + } + + private static byte[] withCalendarValue(byte[] file, Integer calendar) throws IOException { int postscriptLength = Byte.toUnsignedInt(file[file.length - 1]); int postscriptStart = file.length - 1 - postscriptLength; org.apache.orc.OrcProto.PostScript ps = org.apache.orc.OrcProto.PostScript.parseFrom( @@ -281,8 +295,19 @@ private static byte[] withCalendar(byte[] file, org.apache.orc.OrcProto.Calendar int footerStart = postscriptStart - Math.toIntExact(ps.getFooterLength()); org.apache.orc.OrcProto.Footer.Builder footer = org.apache.orc.OrcProto.Footer.parseFrom( Arrays.copyOfRange(file, footerStart, postscriptStart)).toBuilder(); - if (calendar == null) { footer.clearCalendar(); } else { footer.setCalendar(calendar); } - byte[] footerBytes = footer.build().toByteArray(); + footer.clearCalendar(); + ByteArrayOutputStream rewrittenFooter = new ByteArrayOutputStream(); + rewrittenFooter.write(footer.build().toByteArray()); + if (calendar != null) { + rewrittenFooter.write(0x58); // Footer.calendar, field 11, varint. + long value = Integer.toUnsignedLong(calendar); + while ((value & ~0x7fL) != 0) { + rewrittenFooter.write((int) (value & 0x7f) | 0x80); + value >>>= 7; + } + rewrittenFooter.write((int) value); + } + byte[] footerBytes = rewrittenFooter.toByteArray(); byte[] psBytes = ps.toBuilder().setFooterLength(footerBytes.length).build().toByteArray(); ByteArrayOutputStream output = new ByteArrayOutputStream(); output.write(file, 0, footerStart); From c941fc6c67f6a19252ace8ebc9724135e33fd995 Mon Sep 17 00:00:00 2001 From: Allen Xu Date: Wed, 28 Oct 2026 16:59:06 +0800 Subject: [PATCH 4/5] Address ORC calendar review and share kernel error reporting Signed-off-by: Allen Xu --- cpp/include/cudf/io/orc.hpp | 17 +- cpp/include/cudf/io/orc_metadata.hpp | 3 +- cpp/src/io/orc/calendar.cuh | 26 ++- cpp/src/io/orc/decode_error.cuh | 14 +- cpp/src/io/orc/orc_gpu.hpp | 5 +- cpp/src/io/orc/reader_impl.cu | 2 +- cpp/src/io/orc/reader_impl_decode.cu | 15 +- cpp/src/io/orc/stripe_data.cu | 14 +- cpp/src/io/orc/writer_impl.cu | 159 +++++++--------- .../experimental/hybrid_scan_preprocess.cu | 41 ++-- cpp/src/io/parquet/page_decode.cuh | 2 +- cpp/src/io/parquet/page_hdr.cu | 2 +- cpp/src/io/parquet/page_string_decode.cu | 2 +- cpp/src/io/parquet/parquet_gpu.hpp | 2 +- cpp/src/io/parquet/reader_impl.cpp | 2 +- cpp/src/io/parquet/reader_impl_preprocess.cu | 2 +- .../parquet/reader_impl_preprocess_utils.cu | 2 +- .../error.hpp => utilities/kernel_error.hpp} | 4 +- cpp/tests/io/orc_chunked_reader_test.cu | 14 +- cpp/tests/io/orc_test.cpp | 180 +++++++----------- cpp/tests/io/parquet_reader_test.cpp | 10 +- 21 files changed, 238 insertions(+), 280 deletions(-) rename cpp/src/io/{parquet/error.hpp => utilities/kernel_error.hpp} (96%) diff --git a/cpp/include/cudf/io/orc.hpp b/cpp/include/cudf/io/orc.hpp index 951110a99f8d..164da828708f 100644 --- a/cpp/include/cudf/io/orc.hpp +++ b/cpp/include/cudf/io/orc.hpp @@ -173,7 +173,10 @@ class orc_reader_options { * @brief Whether explicitly marked hybrid DATE values are converted to proleptic Gregorian. * @return Whether DATE conversion is enabled */ - [[nodiscard]] bool get_convert_dates_to_proleptic() const { return _convert_dates_to_proleptic; } + [[nodiscard]] bool is_enabled_convert_dates_to_proleptic() const + { + return _convert_dates_to_proleptic; + } /** * @brief Returns timestamp type to which timestamp column will be cast. @@ -286,14 +289,12 @@ class orc_reader_options { * @brief Convert DATE values from explicitly marked JULIAN_GREGORIAN files to proleptic * Gregorian. * - * Disabled by default to preserve raw epoch-day reads used by callers that rebase themselves. - * Missing and UNKNOWN calendar fields are left unchanged. TIMESTAMP columns are unchanged. - * Conversion precedes numpy dtype promotion and applies recursively to nested DATE columns. - * Julian-only leap days and dates outside the int32 proleptic domain cause a read error. + * Disabled by default. Only DATE columns, including nested ones, are converted. Dates with no + * proleptic Gregorian equivalent cause a read error. Conversion precedes numpy dtype promotion. * * @param convert Whether to convert explicitly marked hybrid DATE values */ - void set_convert_dates_to_proleptic(bool convert) { _convert_dates_to_proleptic = convert; } + void enable_convert_dates_to_proleptic(bool convert) { _convert_dates_to_proleptic = convert; } /** * @brief Sets timestamp type to which timestamp column will be cast. @@ -422,12 +423,12 @@ class orc_reader_options_builder { } /** - * @copydoc orc_reader_options::set_convert_dates_to_proleptic + * @copydoc orc_reader_options::enable_convert_dates_to_proleptic * @return this for chaining */ orc_reader_options_builder& convert_dates_to_proleptic(bool convert) { - options.set_convert_dates_to_proleptic(convert); + options.enable_convert_dates_to_proleptic(convert); return *this; } diff --git a/cpp/include/cudf/io/orc_metadata.hpp b/cpp/include/cudf/io/orc_metadata.hpp index 485c3a497908..84cf1eff437c 100644 --- a/cpp/include/cudf/io/orc_metadata.hpp +++ b/cpp/include/cudf/io/orc_metadata.hpp @@ -367,7 +367,8 @@ class orc_metadata { * @brief Returns the calendar recorded in the ORC footer. * * A missing value means the writer did not specify a calendar. This accessor does not - * convert values. See `orc_reader_options::set_convert_dates_to_proleptic` for DATE conversion. + * convert values. See `orc_reader_options::enable_convert_dates_to_proleptic` for DATE + * conversion. * * @return Calendar kind, or `std::nullopt` if the footer has no calendar field */ diff --git a/cpp/src/io/orc/calendar.cuh b/cpp/src/io/orc/calendar.cuh index 85c1a10a6754..7095ea9e4d42 100644 --- a/cpp/src/io/orc/calendar.cuh +++ b/cpp/src/io/orc/calendar.cuh @@ -35,22 +35,30 @@ CUDF_HOST_DEVICE constexpr int64_t rebase_days(int64_t days) constexpr int64_t julian_epoch_offset = 719'470; constexpr int64_t gregorian_era_days = 146'097; constexpr int64_t julian_era_days = 1'461; + constexpr int64_t gregorian_century_days = gregorian_era_days / 4; + constexpr int64_t five_month_days = 153; constexpr auto source_era_days = to_hybrid ? gregorian_era_days : julian_era_days; - auto const z = days + (to_hybrid ? gregorian_epoch_offset : julian_epoch_offset); - auto const era = floor_div(z, source_era_days); - auto const doe = z - era * source_era_days; - auto const yoe = to_hybrid ? (doe - doe / 1'460 + doe / 36'524 - doe / 146'096) / 365 - : (doe == 1'460 ? 3 : doe / 365); + auto const z = days + (to_hybrid ? gregorian_epoch_offset : julian_epoch_offset); + auto const era = floor_div(z, source_era_days); + auto const doe = z - era * source_era_days; + // Discount leap days to recover the year within the source calendar's era. + auto const yoe = to_hybrid ? (doe - doe / (julian_era_days - 1) + doe / gregorian_century_days - + doe / (gregorian_era_days - 1)) / + 365 + : (doe == julian_era_days - 1 ? 3 : doe / 365); auto const year = yoe + era * (to_hybrid ? 400 : 4); auto const doy = doe - (365 * yoe + (to_hybrid ? yoe / 4 - yoe / 100 : 0)); - auto const mp = (5 * doy + 2) / 153; - auto const day = doy - (153 * mp + 2) / 5 + 1; + // March starts the year so the leap day is last. Five-month blocks contain 153 days; + // mp maps March through February to 0 through 11, with doy measured from March 1. + auto const mp = (5 * doy + 2) / five_month_days; + auto const day = doy - (five_month_days * mp + 2) / 5 + 1; if constexpr (to_hybrid) { - return 365 * year + floor_div(year, 4) + (153 * mp + 2) / 5 + day - 1 - julian_epoch_offset; + return 365 * year + floor_div(year, 4) + (five_month_days * mp + 2) / 5 + day - 1 - + julian_epoch_offset; } else { return 365 * year + floor_div(year, 4) - floor_div(year, 100) + floor_div(year, 400) + - (153 * mp + 2) / 5 + day - 1 - gregorian_epoch_offset; + (five_month_days * mp + 2) / 5 + day - 1 - gregorian_epoch_offset; } } diff --git a/cpp/src/io/orc/decode_error.cuh b/cpp/src/io/orc/decode_error.cuh index d5415d1e764b..e8631ef7f1c0 100644 --- a/cpp/src/io/orc/decode_error.cuh +++ b/cpp/src/io/orc/decode_error.cuh @@ -5,22 +5,28 @@ #pragma once -#include +#include "io/utilities/kernel_error.hpp" #include namespace cudf::io::orc::detail { +enum class decode_error : kernel_error::value_type { + INVALID_DATA_STREAM = 1 << 0, + UNREPRESENTABLE_DATE = 1 << 1, +}; + /** * @brief Record a decoding failure without a shared error count wrapping back to zero. * * The reader checks for any error after synchronizing the decoding stream. Columns and * stripes share the same status, so their combined number of failures is not row-count bounded. */ -inline __device__ void set_decode_error(size_type* status) +inline __device__ void set_decode_error(kernel_error::pointer status, + decode_error kind = decode_error::INVALID_DATA_STREAM) { - cuda::atomic_ref error{*status}; - error.fetch_or(1, cuda::std::memory_order_relaxed); + cuda::atomic_ref error{*status}; + error.fetch_or(static_cast(kind), cuda::std::memory_order_relaxed); } } // namespace cudf::io::orc::detail diff --git a/cpp/src/io/orc/orc_gpu.hpp b/cpp/src/io/orc/orc_gpu.hpp index 703575fc9767..ab24f5c3ac1b 100644 --- a/cpp/src/io/orc/orc_gpu.hpp +++ b/cpp/src/io/orc/orc_gpu.hpp @@ -6,6 +6,7 @@ #pragma once #include "io/statistics/statistics.cuh" +#include "io/utilities/kernel_error.hpp" #include "orc.hpp" #include @@ -331,7 +332,7 @@ void decode_nulls_and_string_dictionaries(column_desc* chunks, * @param[in] num_rowgroups Number of row groups in row index data * @param[in] rowidx_stride Row index stride * @param[in] level Current nesting level being processed - * @param[out] error_count Number of errors during decode + * @param[out] error_code Sticky bitmask of errors during decode * @param[in] stream CUDA stream used for device memory operations and kernel launches */ void decode_column_data(column_desc* chunks, @@ -345,7 +346,7 @@ void decode_column_data(column_desc* chunks, int64_t num_rowgroups, size_type rowidx_stride, size_t level, - size_type* error_count, + kernel_error::pointer error_code, cuda::stream_ref stream); /** diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 80312f42948d..53f32b66810d 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -156,7 +156,7 @@ reader_impl::reader_impl(std::size_t chunk_read_limit, _mr(mr), _options{options.get_timestamp_type(), options.is_enabled_use_index(), - options.get_convert_dates_to_proleptic(), + options.is_enabled_convert_dates_to_proleptic(), options.is_enabled_use_np_dtypes(), options.get_decimal128_columns(), options.get_ignore_timezone_in_stripe_footer(), diff --git a/cpp/src/io/orc/reader_impl_decode.cu b/cpp/src/io/orc/reader_impl_decode.cu index 24d5118308f3..c7ad044aa893 100644 --- a/cpp/src/io/orc/reader_impl_decode.cu +++ b/cpp/src/io/orc/reader_impl_decode.cu @@ -3,6 +3,7 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include "decode_error.cuh" #include "io/comp/decompression.hpp" #include "io/comp/gpuinflate.hpp" #include "io/orc/reader_impl.hpp" @@ -424,8 +425,7 @@ void decode_stream_data(int64_t num_dicts, update_null_mask(chunks, out_buffers, stream, mr); } - cudf::detail::device_scalar error_count( - 0, stream, cudf::get_current_device_resource_ref()); + kernel_error error_code(stream); decode_column_data(chunks.base_device_ptr(), global_dict.data(), row_groups, @@ -437,12 +437,15 @@ void decode_stream_data(int64_t num_dicts, row_groups.size().first, row_index_stride, level, - error_count.data(), + error_code.data(), stream); chunks.device_to_host_async(stream); - // `value` synchronizes - auto const num_errors = error_count.value(stream); - CUDF_EXPECTS(num_errors == 0, "ORC data decode failed"); + auto const errors = error_code.value_sync(stream); + CUDF_EXPECTS( + (errors & static_cast(decode_error::UNREPRESENTABLE_DATE)) == 0, + "ORC DATE cannot be represented in the proleptic Gregorian calendar"); + CUDF_EXPECTS(errors == 0, + "ORC data decode failed with code(s) " + kernel_error::to_string(errors)); std::for_each(col_idx_it + 0, col_idx_it + num_columns, [&](auto col_idx) { out_buffers[col_idx].null_count() = diff --git a/cpp/src/io/orc/stripe_data.cu b/cpp/src/io/orc/stripe_data.cu index a859a0ea304a..006905e12b7d 100644 --- a/cpp/src/io/orc/stripe_data.cu +++ b/cpp/src/io/orc/stripe_data.cu @@ -1628,7 +1628,7 @@ CUDF_KERNEL void __launch_bounds__(block_size) int64_t first_row, size_type rowidx_stride, size_t level, - size_type* error_count) + kernel_error::pointer error_code) { __shared__ __align__(16) orcdec_state_s state_g; using block_reduce = cub::BlockReduce; @@ -1669,10 +1669,10 @@ CUDF_KERNEL void __launch_bounds__(block_size) // If we have an index, seek to the initial run and update row positions if (num_rowgroups > 0) { if (s->top.data.index.strm_offset[0] > s->chunk.strm_len[CI_DATA]) { - set_decode_error(error_count); + set_decode_error(error_code); } if (s->top.data.index.strm_offset[1] > s->chunk.strm_len[CI_DATA2]) { - set_decode_error(error_count); + set_decode_error(error_code); } auto const ofs0 = min(s->top.data.index.strm_offset[0], s->chunk.strm_len[CI_DATA]); auto const ofs1 = min(s->top.data.index.strm_offset[1], s->chunk.strm_len[CI_DATA2]); @@ -2007,7 +2007,7 @@ CUDF_KERNEL void __launch_bounds__(block_size) s->chunk.convert_dates_to_proleptic ? rebase_days(original) : original; if (s->chunk.convert_dates_to_proleptic and not is_representable_date(original, rebased)) { - set_decode_error(error_count); + set_decode_error(error_code, decode_error::UNREPRESENTABLE_DATE); } if (s->chunk.dtype_len == 8) { cudf::duration_D days{static_cast(rebased)}; @@ -2192,7 +2192,7 @@ void __host__ decode_nulls_and_string_dictionaries(column_desc* chunks, * @param[in] num_rowgroups Number of row groups in row index data * @param[in] rowidx_stride Row index stride * @param[in] level nesting level being processed - * @param[in] error_count Pointer to error count output + * @param[in] error_code Pointer to the shared decoding error bitmask * @param[in] stream CUDA stream used for device memory operations and kernel launches */ void __host__ decode_column_data(column_desc* chunks, @@ -2206,7 +2206,7 @@ void __host__ decode_column_data(column_desc* chunks, int64_t num_rowgroups, size_type rowidx_stride, size_t level, - size_type* error_count, + kernel_error::pointer error_code, cuda::stream_ref stream) { auto const num_blocks = num_columns * (num_rowgroups > 0 ? num_rowgroups : num_stripes); @@ -2219,7 +2219,7 @@ void __host__ decode_column_data(column_desc* chunks, first_row, rowidx_stride, level, - error_count); + error_code); CUDF_CUDA_TRY(cudaGetLastError()); } diff --git a/cpp/src/io/orc/writer_impl.cu b/cpp/src/io/orc/writer_impl.cu index 30c700a9a796..e7192bf08214 100644 --- a/cpp/src/io/orc/writer_impl.cu +++ b/cpp/src/io/orc/writer_impl.cu @@ -9,6 +9,7 @@ */ #include "calendar.cuh" +#include "decode_error.cuh" #include "io/comp/compression.hpp" #include "io/orc/orc_gpu.hpp" #include "io/statistics/column_statistics.cuh" @@ -221,7 +222,7 @@ class orc_column_view { void add_child(uint32_t child_idx) { children.emplace_back(child_idx); } - auto const& view() const noexcept { return cudf_column; } + [[nodiscard]] auto const& view() const noexcept { return cudf_column; } auto type() const noexcept { return cudf_column.type(); } auto is_string() const noexcept { return cudf_column.type().id() == type_id::STRING; } @@ -330,7 +331,6 @@ struct orc_table_view { std::vector string_column_indices; cuda::device_buffer d_string_column_indices; std::vector> rebased_dates; - std::vector> rebased_date_masks; // Keep selectively compacted input alive through statistics persistence and sink writes. std::vector> compacted_columns; @@ -2153,13 +2153,44 @@ hostdevice_2dvector calculate_rowgroup_bounds(orc_table_view cons } /** - * @brief Rebase DATE leaves before encoding and statistics, using the same device column views. + * @brief Compact sliced DATE inputs before allocating rebased data. + * + * Return both the owning columns and their table view so ownership remains explicit. + */ +[[nodiscard]] auto compact_legacy_date_inputs(table_view const& input, + orc_table_view const& orc_table, + host_2dspan bounds, + cuda::stream_ref stream) +{ + std::vector compact(input.num_columns(), false); + size_type root_index = -1; + for (auto const& col : orc_table.columns) { + if (not col.is_child()) { ++root_index; } + if (col.orc_kind() != DATE) { continue; } + auto const begin = bounds[0][col.index()].begin; + auto const end = bounds[bounds.size().first - 1][col.index()].end; + if (begin != end and (col.view().offset() != 0 or begin != 0 or end != col.size())) { + compact[root_index] = true; + } + } + std::vector> owned_columns; + std::vector views(input.begin(), input.end()); + for (size_type index = 0; index < input.num_columns(); ++index) { + if (not compact[index]) { continue; } + auto& compacted = owned_columns.emplace_back(std::make_unique( + input.column(index), stream, cudf::get_current_device_resource_ref())); + views[index] = compacted->view(); + } + return std::pair{std::move(owned_columns), table_view{views}}; +} + +/** + * @brief Rebase DATE leaves before encoding and statistics, preserving their null masks. * * Rowgroup bounds restrict validation to rows actually written from sliced nested columns. - * Pushdown masks also exclude values hidden by null ancestors. Owned buffers live in orc_table. + * Invalid dates share one device error flag, checked after processing all DATE leaves. */ void rebase_date_columns(orc_table_view& orc_table, - pushdown_null_masks const& pd_masks, host_2dspan bounds, cuda::stream_ref stream) { @@ -2168,6 +2199,7 @@ void rebase_date_columns(orc_table_view& orc_table, "Legacy ORC calendar does not support TIMESTAMP columns"); } if (orc_table.num_rows() == 0) { return; } + kernel_error error_code(stream); for (auto const& col : orc_table.columns) { if (col.orc_kind() != DATE) { continue; } auto const& view = col.view(); @@ -2176,51 +2208,39 @@ void rebase_date_columns(orc_table_view& orc_table, if (begin == end) { continue; } auto& data = orc_table.rebased_dates.emplace_back( static_cast(view.offset()) + view.size(), stream); - auto const columns = device_span{orc_table.d_columns}; - auto const parent_mask = col.is_child() ? pd_masks.masks[col.parent_index()] : nullptr; - auto mask = view.null_mask(); - auto null_count = view.null_count(); - if (parent_mask != nullptr) { - // Statistics use the base column view, so they need the ancestor validity as well. - std::vector masks{parent_mask}; - if (mask != nullptr) { masks.push_back(mask); } - std::vector offsets(masks.size(), 0); - auto [combined, count] = cudf::detail::bitmask_and(masks, - offsets, - view.offset() + view.size(), - stream, - cudf::get_current_device_resource_ref()); - mask = orc_table.rebased_date_masks.emplace_back(std::move(combined)).data(); - null_count = view.offset() == 0 ? count - : cudf::detail::null_count( - mask, view.offset(), view.offset() + view.size(), stream); - } - auto const invalid = thrust::transform_reduce( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::counting_iterator{begin}, - cuda::counting_iterator{end}, - [columns, idx = col.index(), mask, out = data.data()] __device__(auto row) -> bool { - auto const& column = columns[idx]; - if (not bit_value_or(mask, row + column.offset(), true)) { - out[row + column.offset()] = 0; - return false; - } - auto const original = column.element(row); - auto const rebased = rebase_days(original); - out[row + column.offset()] = static_cast(rebased); - return not is_representable_date(original, rebased); - }, - false, - cuda::std::logical_or{}); - CUDF_EXPECTS(not invalid, "DATE cannot be represented losslessly in the legacy ORC calendar"); + auto const columns = device_span{orc_table.d_columns}; + auto const mask = view.null_mask(); + thrust::for_each_n(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator{begin}, + end - begin, + [columns, + idx = col.index(), + mask, + out = data.data(), + error_status = error_code.data()] __device__(auto row) { + auto const& column = columns[idx]; + if (not bit_value_or(mask, row + column.offset(), true)) { + out[row + column.offset()] = 0; + return; + } + auto const original = column.element(row); + auto const rebased = rebase_days(original); + out[row + column.offset()] = static_cast(rebased); + if (not is_representable_date(original, rebased)) { + set_decode_error(error_status, decode_error::UNREPRESENTABLE_DATE); + } + }); auto const rebased_view = column_device_view::create( - column_view{view.type(), view.size(), data.data(), mask, null_count, view.offset()}, stream); + column_view{view.type(), view.size(), data.data(), mask, view.null_count(), view.offset()}, + stream); cudf::detail::device_single_thread( [columns, idx = col.index(), rebased = *rebased_view] __device__() { static_cast(columns[idx]) = rebased; }, stream); } + CUDF_EXPECTS(error_code.value_sync(stream) == 0, + "DATE cannot be represented losslessly in the legacy ORC calendar"); } // returns host vector of per-rowgroup sizes @@ -2628,64 +2648,25 @@ auto convert_table_to_orc_data(table_view const& input, data_sink const& out_sink, cuda::stream_ref stream) { - std::vector> compacted_columns; - auto input_view = input; - if (calendar == CalendarKind::JULIAN_GREGORIAN and input.num_rows() != 0) { - // Top-level DATE offsets are already known. Compact them before constructing device - // views so the common chunked-write path only builds its views and rowgroup bounds once. - std::vector views(input.begin(), input.end()); - for (size_type index = 0; index < input.num_columns(); ++index) { - auto const& col = input.column(index); - if (col.type().id() != type_id::TIMESTAMP_DAYS or col.offset() == 0) { continue; } - auto& compacted = compacted_columns.emplace_back( - std::make_unique(col, stream, cudf::get_current_device_resource_ref())); - views[index] = compacted->view(); - } - if (not compacted_columns.empty()) { input_view = table_view{views}; } - } - auto input_tview = table_device_view::create(input_view, stream); - - auto orc_table = make_orc_table_view(input_view, *input_tview, table_meta, stream); - + auto input_tview = table_device_view::create(input, stream); + auto orc_table = make_orc_table_view(input, *input_tview, table_meta, stream); auto rowgroup_bounds = calculate_rowgroup_bounds(orc_table, row_index_stride, stream); if (calendar == CalendarKind::JULIAN_GREGORIAN and orc_table.num_rows() != 0) { - // A DATE view retains its backing offset/extent even when only a small nested slice is - // written. Compact just the affected top-level columns before allocating rebased buffers. - std::vector compact(input.num_columns(), false); - size_type root_index = -1; - for (auto const& col : orc_table.columns) { - if (not col.is_child()) { ++root_index; } - if (col.orc_kind() != DATE) { continue; } - auto const begin = rowgroup_bounds[0][col.index()].begin; - auto const end = rowgroup_bounds[rowgroup_bounds.size().first - 1][col.index()].end; - if (begin != end and (col.view().offset() != 0 or begin != 0 or end != col.size())) { - compact[root_index] = true; - } - } - if (std::any_of(compact.begin(), compact.end(), [](bool value) { return value; })) { - std::vector views; - for (size_type index = 0; index < input.num_columns(); ++index) { - if (compact[index]) { - auto& compacted = compacted_columns.emplace_back(std::make_unique( - input_view.column(index), stream, cudf::get_current_device_resource_ref())); - views.push_back(compacted->view()); - } else { - views.push_back(input_view.column(index)); - } - } - table_view compacted_input{views}; + auto [compacted_columns, compacted_input] = + compact_legacy_date_inputs(input, orc_table, rowgroup_bounds.host_view(), stream); + if (not compacted_columns.empty()) { input_tview = table_device_view::create(compacted_input, stream); orc_table = make_orc_table_view(compacted_input, *input_tview, table_meta, stream); rowgroup_bounds = calculate_rowgroup_bounds(orc_table, row_index_stride, stream); } + orc_table.compacted_columns = std::move(compacted_columns); } - orc_table.compacted_columns = std::move(compacted_columns); // This is unused but it holds memory buffers for later access thus needs to be kept alive. [[maybe_unused]] auto const pd_masks = init_pushdown_null_masks(orc_table, stream); if (calendar == CalendarKind::JULIAN_GREGORIAN) { - rebase_date_columns(orc_table, pd_masks, rowgroup_bounds.host_view(), stream); + rebase_date_columns(orc_table, rowgroup_bounds.host_view(), stream); } [[maybe_unused]] auto const rg_char_counts_data = diff --git a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu index a1b2848002aa..286c639e8784 100644 --- a/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu +++ b/cpp/src/io/parquet/experimental/hybrid_scan_preprocess.cu @@ -65,7 +65,7 @@ void decode_dictionary_page_headers( chunk_page_offsets.end(), 0); - parquet::kernel_error error_code(stream); + kernel_error error_code(stream); // Decode dictionary page headers, one thread per page parquet::detail::decode_page_headers_from_page_data( @@ -78,33 +78,32 @@ void decode_dictionary_page_headers( stream); if (auto const error = error_code.value_sync(stream); error != 0) { - CUDF_FAIL("Parquet header parsing failed with code(s) " + - parquet::kernel_error::to_string(error)); + CUDF_FAIL("Parquet header parsing failed with code(s) " + kernel_error::to_string(error)); } // Point each chunk at its dictionary page. A chunk whose span is empty keeps a zeroed page and is // not pruned with. A non-empty span must hold a dictionary page. - thrust::for_each(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::counting_iterator(0), - cuda::counting_iterator(chunks.size()), - [chunks = chunks.device_begin(), - pages = pages.device_begin(), - error_code = error_code.data()] __device__(auto chunk_idx) { - auto const& page = pages[chunk_idx]; - auto& chunk = chunks[chunk_idx]; - if (page.flags & parquet::detail::PAGEINFO_FLAGS_DICTIONARY) { - chunk.dict_page = &page; - } else if (chunk.compressed_size > 0) { - parquet::detail::set_error( - static_cast( - parquet::detail::decode_error::INVALID_PAGE_TYPE), - error_code); - } - }); + thrust::for_each( + rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator(0), + cuda::counting_iterator(chunks.size()), + [chunks = chunks.device_begin(), + pages = pages.device_begin(), + error_code = error_code.data()] __device__(auto chunk_idx) { + auto const& page = pages[chunk_idx]; + auto& chunk = chunks[chunk_idx]; + if (page.flags & parquet::detail::PAGEINFO_FLAGS_DICTIONARY) { + chunk.dict_page = &page; + } else if (chunk.compressed_size > 0) { + parquet::detail::set_error( + static_cast(parquet::detail::decode_error::INVALID_PAGE_TYPE), + error_code); + } + }); if (auto const error = error_code.value_sync(stream); error != 0) { CUDF_FAIL("Dictionary page data holds a non-dictionary page, code(s) " + - parquet::kernel_error::to_string(error)); + kernel_error::to_string(error)); } pages.device_to_host_async(stream); diff --git a/cpp/src/io/parquet/page_decode.cuh b/cpp/src/io/parquet/page_decode.cuh index 44e1f520a621..777a380a17cf 100644 --- a/cpp/src/io/parquet/page_decode.cuh +++ b/cpp/src/io/parquet/page_decode.cuh @@ -5,9 +5,9 @@ #pragma once -#include "error.hpp" #include "io/parquet/timestamp_utils.cuh" #include "io/utilities/block_utils.cuh" +#include "io/utilities/kernel_error.hpp" #include "parquet_gpu.hpp" #include "rle_stream.cuh" diff --git a/cpp/src/io/parquet/page_hdr.cu b/cpp/src/io/parquet/page_hdr.cu index edb979ee0ab3..c2fdebd21847 100644 --- a/cpp/src/io/parquet/page_hdr.cu +++ b/cpp/src/io/parquet/page_hdr.cu @@ -4,8 +4,8 @@ */ #include "cuda/std/__utility/cmp.h" -#include "error.hpp" #include "io/utilities/block_utils.cuh" +#include "io/utilities/kernel_error.hpp" #include "parquet_gpu.hpp" #include diff --git a/cpp/src/io/parquet/page_string_decode.cu b/cpp/src/io/parquet/page_string_decode.cu index b16bab32bbff..45022bcf5aad 100644 --- a/cpp/src/io/parquet/page_string_decode.cu +++ b/cpp/src/io/parquet/page_string_decode.cu @@ -4,7 +4,7 @@ */ #include "delta_binary.cuh" -#include "error.hpp" +#include "io/utilities/kernel_error.hpp" #include "page_decode.cuh" #include "page_state_composed.cuh" #include "page_string_utils.cuh" diff --git a/cpp/src/io/parquet/parquet_gpu.hpp b/cpp/src/io/parquet/parquet_gpu.hpp index f5268131679d..f117cc924d4f 100644 --- a/cpp/src/io/parquet/parquet_gpu.hpp +++ b/cpp/src/io/parquet/parquet_gpu.hpp @@ -5,11 +5,11 @@ #pragma once -#include "error.hpp" #include "io/comp/compression.hpp" #include "io/parquet/parquet_common.hpp" #include "io/statistics/statistics.cuh" #include "io/utilities/column_buffer.hpp" +#include "io/utilities/kernel_error.hpp" #include #include diff --git a/cpp/src/io/parquet/reader_impl.cpp b/cpp/src/io/parquet/reader_impl.cpp index dd25ff56c231..758b20ec0ac4 100644 --- a/cpp/src/io/parquet/reader_impl.cpp +++ b/cpp/src/io/parquet/reader_impl.cpp @@ -6,7 +6,7 @@ #include "reader_impl.hpp" #include "column_path_helpers.hpp" -#include "error.hpp" +#include "io/utilities/kernel_error.hpp" #include "runtime/context.hpp" #include "synthetic_column_helpers.hpp" diff --git a/cpp/src/io/parquet/reader_impl_preprocess.cu b/cpp/src/io/parquet/reader_impl_preprocess.cu index d221e76425c3..1f05b65aa101 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess.cu @@ -3,8 +3,8 @@ * SPDX-License-Identifier: Apache-2.0 */ -#include "error.hpp" #include "io/comp/common.hpp" +#include "io/utilities/kernel_error.hpp" #include "page_decode.cuh" #include "reader_impl.hpp" #include "reader_impl_chunking_utils.cuh" diff --git a/cpp/src/io/parquet/reader_impl_preprocess_utils.cu b/cpp/src/io/parquet/reader_impl_preprocess_utils.cu index a4ea8f842073..c8d39bd27248 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess_utils.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess_utils.cu @@ -3,7 +3,7 @@ * SPDX-License-Identifier: Apache-2.0 */ -#include "error.hpp" +#include "io/utilities/kernel_error.hpp" #include "reader_impl_preprocess_utils.cuh" #include diff --git a/cpp/src/io/parquet/error.hpp b/cpp/src/io/utilities/kernel_error.hpp similarity index 96% rename from cpp/src/io/parquet/error.hpp rename to cpp/src/io/utilities/kernel_error.hpp index 23abeefe0ddd..bc80de0acf15 100644 --- a/cpp/src/io/parquet/error.hpp +++ b/cpp/src/io/utilities/kernel_error.hpp @@ -12,7 +12,7 @@ #include #include -namespace cudf::io::parquet { +namespace cudf::io { /** * @brief Specialized device scalar for use in reporting errors that occur in @@ -77,4 +77,4 @@ class kernel_error { } }; -} // namespace cudf::io::parquet +} // namespace cudf::io diff --git a/cpp/tests/io/orc_chunked_reader_test.cu b/cpp/tests/io/orc_chunked_reader_test.cu index a313c4055896..11208f7c3436 100644 --- a/cpp/tests/io/orc_chunked_reader_test.cu +++ b/cpp/tests/io/orc_chunked_reader_test.cu @@ -195,11 +195,12 @@ void expect_selected_nested_empty_struct_chunked_round_trip( expect_selected_nested_empty_struct_table(expected->view(), chunk.tbl->view()); } -void record_two_decode_errors(cudf::size_type* status, cuda::stream_ref stream) +void record_two_decode_errors(cudf::io::kernel_error::pointer status, cuda::stream_ref stream) { cudf::detail::device_single_thread( [status] __device__ { - cudf::io::orc::detail::set_decode_error(status); + cudf::io::orc::detail::set_decode_error( + status, cudf::io::orc::detail::decode_error::UNREPRESENTABLE_DATE); cudf::io::orc::detail::set_decode_error(status); }, stream); @@ -212,11 +213,12 @@ struct OrcChunkedReaderTest : public cudf::test::BaseFixture {}; TEST_F(OrcChunkedReaderTest, CalendarDecodeErrorDoesNotWrap) { auto const stream = cudf::get_default_stream(); - for (auto initial : {cudf::size_type{0}, cudf::size_type{-2}}) { - rmm::device_scalar status{initial, stream}; - // -2 represents the state after 2^32 - 2 failures. Two more increments used to clear it. + using value_type = cudf::io::kernel_error::value_type; + for (auto initial : {value_type{0}, value_type{UINT32_MAX - 1}}) { + rmm::device_scalar status{initial, stream}; + // 2^32 - 2 followed by two more increments used to wrap back to zero. record_two_decode_errors(status.data(), stream); - EXPECT_NE(status.value(stream), 0); + EXPECT_EQ(status.value(stream), initial | value_type{3}); } } diff --git a/cpp/tests/io/orc_test.cpp b/cpp/tests/io/orc_test.cpp index 20f2a414ed4c..f7383dea0b46 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -2104,40 +2104,39 @@ TEST_F(OrcWriterTest, ProlepticCalendar) {-353920}, {-281237}, {-141438}, {-141432}, {-141427}, {}, {19849}, {0}}; table_view expected({dates, nested, lists}); - for (auto compression : {cudf::io::compression_type::NONE, cudf::io::compression_type::SNAPPY}) { - for (bool chunked : {false, true}) { - SCOPED_TRACE(static_cast(compression)); - SCOPED_TRACE(chunked); - std::vector buffer; - if (chunked) { - auto opts = cudf::io::chunked_orc_writer_options::builder(cudf::io::sink_info{&buffer}) - .compression(compression) - .calendar_kind(calendar::PROLEPTIC_GREGORIAN) - .build(); - cudf::io::orc_chunked_writer writer(opts); - auto const chunks = cudf::slice(expected, {0, 3, 3, 8}); - writer.write(chunks[0]).write(chunks[1]); - writer.close(); - } else { - // Exercise the default, which must describe the values the writer actually encodes. - auto opts = cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, expected) - .compression(compression) - .build(); - EXPECT_EQ(opts.get_calendar_kind(), calendar::PROLEPTIC_GREGORIAN); - cudf::io::write_orc(opts); - } - cudf::io::source_info source{cudf::host_span{ - reinterpret_cast(buffer.data()), buffer.size()}}; - EXPECT_EQ(cudf::io::read_orc_metadata(source).calendar_kind(), calendar::PROLEPTIC_GREGORIAN); - auto const result = cudf::io::read_orc( - cudf::io::orc_reader_options::builder(source).use_np_dtypes(false).build()); - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected, result.tbl->view()); - auto const stats = cudf::io::read_parsed_orc_statistics(source); - auto const date_stats = - std::get(stats.file_stats[1].type_specific_stats); - EXPECT_EQ(date_stats.minimum, -353920); - EXPECT_EQ(date_stats.maximum, 19849); + auto const compression = cudf::io::compression_type::SNAPPY; + + for (bool chunked : {false, true}) { + SCOPED_TRACE(chunked); + std::vector buffer; + if (chunked) { + auto opts = cudf::io::chunked_orc_writer_options::builder(cudf::io::sink_info{&buffer}) + .compression(compression) + .calendar_kind(calendar::PROLEPTIC_GREGORIAN) + .build(); + cudf::io::orc_chunked_writer writer(opts); + auto const chunks = cudf::slice(expected, {0, 3, 3, 8}); + writer.write(chunks[0]).write(chunks[1]); + writer.close(); + } else { + // Exercise the default, which must describe the values the writer actually encodes. + auto opts = cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, expected) + .compression(compression) + .build(); + EXPECT_EQ(opts.get_calendar_kind(), calendar::PROLEPTIC_GREGORIAN); + cudf::io::write_orc(opts); } + cudf::io::source_info source{cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}}; + EXPECT_EQ(cudf::io::read_orc_metadata(source).calendar_kind(), calendar::PROLEPTIC_GREGORIAN); + auto const result = cudf::io::read_orc( + cudf::io::orc_reader_options::builder(source).use_np_dtypes(false).build()); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(expected, result.tbl->view()); + auto const stats = cudf::io::read_parsed_orc_statistics(source); + auto const date_stats = + std::get(stats.file_stats[1].type_specific_stats); + EXPECT_EQ(date_stats.minimum, -353920); + EXPECT_EQ(date_stats.maximum, 19849); } } @@ -2151,43 +2150,41 @@ TEST_F(OrcWriterTest, LegacyCalendar) cudf::test::lists_column_wrapper lists{ {-719528}, {}, {-353920}, {-281237}, {-141438}, {-141427}, {0}, {19849}, {}}; table_view input({dates, nested, lists}); - for (auto compression : {cudf::io::compression_type::NONE, cudf::io::compression_type::SNAPPY}) { - for (bool chunked : {false, true}) { - std::vector buffer; - if (chunked) { - cudf::io::orc_chunked_writer writer( - cudf::io::chunked_orc_writer_options::builder(cudf::io::sink_info{&buffer}) - .compression(compression) - .calendar_kind(calendar::JULIAN_GREGORIAN)); - auto chunks = cudf::slice(input, {0, 3, 3, 9}); - writer.write(chunks[0]).write(chunks[1]); - writer.close(); - } else { - cudf::io::write_orc( - cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, input) - .compression(compression) - .calendar_kind(calendar::JULIAN_GREGORIAN)); - } - cudf::io::source_info source{cudf::host_span{ - reinterpret_cast(buffer.data()), buffer.size()}}; - EXPECT_EQ(cudf::io::read_orc_metadata(source).calendar_kind(), calendar::JULIAN_GREGORIAN); - auto result = cudf::io::read_orc(cudf::io::orc_reader_options::builder(source) - .use_np_dtypes(false) - .convert_dates_to_proleptic(true)); - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(input, result.tbl->view()); - // Independent Julian epoch-day values, including astronomical year zero. - column_wrapper raw( - {-719530, -719164, -353914, -281230, -141428, -141427, 0, 19849, 0}, - {true, true, true, true, true, true, true, true, false}); - auto raw_result = cudf::io::read_orc( - cudf::io::orc_reader_options::builder(source).columns({"_col0"}).use_np_dtypes(false)); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(raw, raw_result.tbl->view().column(0)); - auto stats = cudf::io::read_parsed_orc_statistics(source); - auto date_stats = - std::get(stats.file_stats[1].type_specific_stats); - EXPECT_EQ(date_stats.minimum, -719530); - EXPECT_EQ(date_stats.maximum, 19849); + auto const compression = cudf::io::compression_type::SNAPPY; + + for (bool chunked : {false, true}) { + std::vector buffer; + if (chunked) { + cudf::io::orc_chunked_writer writer( + cudf::io::chunked_orc_writer_options::builder(cudf::io::sink_info{&buffer}) + .compression(compression) + .calendar_kind(calendar::JULIAN_GREGORIAN)); + auto chunks = cudf::slice(input, {0, 3, 3, 9}); + writer.write(chunks[0]).write(chunks[1]); + writer.close(); + } else { + cudf::io::write_orc(cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, input) + .compression(compression) + .calendar_kind(calendar::JULIAN_GREGORIAN)); } + cudf::io::source_info source{cudf::host_span{ + reinterpret_cast(buffer.data()), buffer.size()}}; + EXPECT_EQ(cudf::io::read_orc_metadata(source).calendar_kind(), calendar::JULIAN_GREGORIAN); + auto result = cudf::io::read_orc( + cudf::io::orc_reader_options::builder(source).use_np_dtypes(false).convert_dates_to_proleptic( + true)); + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(input, result.tbl->view()); + // Independent Julian epoch-day values, including astronomical year zero. + column_wrapper raw( + {-719530, -719164, -353914, -281230, -141428, -141427, 0, 19849, 0}, + {true, true, true, true, true, true, true, true, false}); + auto raw_result = cudf::io::read_orc( + cudf::io::orc_reader_options::builder(source).columns({"_col0"}).use_np_dtypes(false)); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(raw, raw_result.tbl->view().column(0)); + auto stats = cudf::io::read_parsed_orc_statistics(source); + auto date_stats = std::get(stats.file_stats[1].type_specific_stats); + EXPECT_EQ(date_stats.minimum, -719530); + EXPECT_EQ(date_stats.maximum, 19849); } } @@ -2198,17 +2195,7 @@ TEST_F(OrcWriterTest, LegacyCalendarRejectsLossyValues) auto options = cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, table_view{}) .calendar_kind(calendar::JULIAN_GREGORIAN) .build(); - for (int32_t day : std::vector{INT32_MIN, - -141437, - -141436, - -141435, - -141434, - -141433, - -141432, - -141431, - -141430, - -141429, - -141428}) { + for (int32_t day : {INT32_MIN, -141437, -141428}) { column_wrapper dates{day}; options.set_table(table_view{{dates}}); EXPECT_THROW(cudf::io::write_orc(options), cudf::logic_error); @@ -2382,33 +2369,6 @@ TEST_F(OrcWriterTest, LegacyCalendarRetryAfterInvalidChunk) CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected, result.tbl->view().column(0)); } -TEST_F(OrcWriterTest, LegacyCalendarIgnoresHiddenListValues) -{ - column_wrapper dates{-141432, -353920, -281237}; - int32_col offsets{0, 1, 2, 3}; - auto validity = std::vector{false, true, true}; - auto [mask, null_count] = cudf::test::detail::make_null_mask(validity.begin(), validity.end()); - auto lists = - cudf::make_lists_column(3, offsets.release(), dates.release(), null_count, std::move(mask)); - std::vector buffer; - cudf::io::write_orc( - cudf::io::orc_writer_options::builder(cudf::io::sink_info{&buffer}, table_view{{lists->view()}}) - .calendar_kind(cudf::io::orc::CalendarKind::JULIAN_GREGORIAN)); - cudf::io::source_info source{cudf::host_span{ - reinterpret_cast(buffer.data()), buffer.size()}}; - auto result = cudf::io::read_orc( - cudf::io::orc_reader_options::builder(source).use_np_dtypes(false).convert_dates_to_proleptic( - true)); - cudf::test::lists_column_wrapper expected({{}, {-353920}, {-281237}}, - validity.begin()); - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected, result.tbl->view().column(0)); - auto stats = cudf::io::read_parsed_orc_statistics(source); - auto date_stats = std::get(stats.file_stats[2].type_specific_stats); - EXPECT_EQ(date_stats.minimum, -353914); - EXPECT_EQ(date_stats.maximum, -281230); - EXPECT_EQ(stats.file_stats[2].number_of_values, 2); -} - TEST_F(OrcWriterTest, LegacyCalendarWideDateDomain) { column_wrapper dates{ @@ -2430,11 +2390,7 @@ TEST_F(OrcWriterTest, CalendarRejectsUnsupportedModes) using calendar = cudf::io::orc::CalendarKind; cudf::io::orc_writer_options opts; cudf::io::chunked_orc_writer_options chunked_opts; - for (auto kind : {calendar::UNKNOWN_CALENDAR, - static_cast(3), - static_cast(257), - static_cast(258), - static_cast(UINT32_MAX)}) { + for (auto kind : {calendar::UNKNOWN_CALENDAR, static_cast(3)}) { EXPECT_THROW(opts.set_calendar_kind(kind), cudf::logic_error); EXPECT_THROW(chunked_opts.set_calendar_kind(kind), cudf::logic_error); } diff --git a/cpp/tests/io/parquet_reader_test.cpp b/cpp/tests/io/parquet_reader_test.cpp index 586f97de8ec8..57300843e5b2 100644 --- a/cpp/tests/io/parquet_reader_test.cpp +++ b/cpp/tests/io/parquet_reader_test.cpp @@ -399,7 +399,7 @@ TEST_F(ParquetReaderTest, RequiredStructUserBoundsAcrossPages) .max_page_fragment_size(32); cudf::io::write_parquet(out_opts); - for (auto const [skip_rows, num_rows_to_read] : page_boundary_slices(num_rows)) { + for (auto const& [skip_rows, num_rows_to_read] : page_boundary_slices(num_rows)) { auto read_opts = cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}) .skip_rows(skip_rows) .num_rows(num_rows_to_read); @@ -475,7 +475,7 @@ TEST_F(ParquetReaderTest, NullableStructBoolAndStringUserBounds) .max_page_fragment_size(32) .write_v2_headers(true)); - for (auto const [skip_rows, num_rows_to_read] : page_boundary_slices(row_count)) { + for (auto const& [skip_rows, num_rows_to_read] : page_boundary_slices(row_count)) { auto const result = cudf::io::read_parquet( cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}) .skip_rows(skip_rows) @@ -558,7 +558,7 @@ TEST_F(ParquetReaderTest, NullableStructDictionaryAndByteStreamSplit) bss_encodings.end()); for (auto const& filepath : {dictionary_filepath, bss_filepath}) { - for (auto const [skip_rows, num_rows_to_read] : page_boundary_slices(row_count)) { + for (auto const& [skip_rows, num_rows_to_read] : page_boundary_slices(row_count)) { auto const result = cudf::io::read_parquet( cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}) .skip_rows(skip_rows) @@ -636,7 +636,7 @@ TEST_F(ParquetReaderTest, ByteStreamSplitNestedListFloat) std::find(encodings.begin(), encodings.end(), cudf::io::parquet::Encoding::BYTE_STREAM_SPLIT), encodings.end()); - for (auto const [skip_rows, num_rows_to_read] : page_boundary_slices(num_rows)) { + for (auto const& [skip_rows, num_rows_to_read] : page_boundary_slices(num_rows)) { auto const result = cudf::io::read_parquet( cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}) .skip_rows(skip_rows) @@ -842,7 +842,7 @@ TEST_F(ParquetReaderTest, NullableListAllNullPages) .max_page_size_rows(32) .max_page_fragment_size(32)); - for (auto const [skip_rows, num_rows_to_read] : page_boundary_slices(num_rows)) { + for (auto const& [skip_rows, num_rows_to_read] : page_boundary_slices(num_rows)) { auto const result = cudf::io::read_parquet( cudf::io::parquet_reader_options::builder(cudf::io::source_info{filepath}) .skip_rows(skip_rows) From a507a9958740d23f807542319e847ae851a69cb3 Mon Sep 17 00:00:00 2001 From: Allen Xu Date: Fri, 6 Nov 2026 16:47:43 +0800 Subject: [PATCH 5/5] Expose ORC calendar controls through Java and JNI Signed-off-by: Allen Xu --- java/pom.xml | 48 +++ .../main/java/ai/rapids/cudf/ORCOptions.java | 20 +- .../java/ai/rapids/cudf/ORCWriterOptions.java | 38 +++ java/src/main/java/ai/rapids/cudf/Table.java | 18 +- java/src/main/native/src/TableJni.cpp | 52 +-- .../rapids/cudf/ORCCalendarInteropTest.java | 320 ++++++++++++++++++ .../java/ai/rapids/cudf/ORCCalendarTest.java | 100 ++++++ 7 files changed, 570 insertions(+), 26 deletions(-) create mode 100644 java/src/test/interop/ai/rapids/cudf/ORCCalendarInteropTest.java create mode 100644 java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java diff --git a/java/pom.xml b/java/pom.xml index ae5d47016506..db088169a4ef 100644 --- a/java/pom.xml +++ b/java/pom.xml @@ -136,6 +136,32 @@ + + org.apache.orc + orc-core + 1.9.4 + shaded-protobuf + test + + + + org.apache.hadoop + hadoop-client-api + + + + + org.apache.hadoop + hadoop-hdfs-client + 3.4.0 + test + + + org.apache.hive + hive-storage-api + 2.8.1 + test + org.slf4j slf4j-api @@ -664,6 +690,28 @@ + + + org.codehaus.mojo + build-helper-maven-plugin + 3.6.0 + + + add-interop-tests + generate-test-sources + + add-test-source + + + + ${project.basedir}/src/test/interop + + + + + maven-antrun-plugin 1.8 diff --git a/java/src/main/java/ai/rapids/cudf/ORCOptions.java b/java/src/main/java/ai/rapids/cudf/ORCOptions.java index 9fa8d89a23cf..d07b64a8faf6 100644 --- a/java/src/main/java/ai/rapids/cudf/ORCOptions.java +++ b/java/src/main/java/ai/rapids/cudf/ORCOptions.java @@ -1,6 +1,6 @@ /* * - * 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 * */ @@ -19,6 +19,7 @@ public class ORCOptions extends ColumnFilterOptions { public static ORCOptions DEFAULT = new ORCOptions(new Builder()); private final boolean useNumPyTypes; + private final boolean convertDatesToProleptic; private final DType unit; private final String[] decimal128Columns; private final boolean ignoreTimezoneInStripeFooter; @@ -27,6 +28,7 @@ private ORCOptions(Builder builder) { super(builder); decimal128Columns = builder.decimal128Columns.toArray(new String[0]); useNumPyTypes = builder.useNumPyTypes; + convertDatesToProleptic = builder.convertDatesToProleptic; unit = builder.unit; ignoreTimezoneInStripeFooter = builder.ignoreTimezoneInStripeFooter; } @@ -35,6 +37,8 @@ boolean usingNumPyTypes() { return useNumPyTypes; } + boolean convertDatesToProleptic() { return convertDatesToProleptic; } + DType timeUnit() { return unit; } @@ -53,6 +57,7 @@ public static Builder builder() { public static class Builder extends ColumnFilterOptions.Builder { private boolean useNumPyTypes = true; + private boolean convertDatesToProleptic = false; private DType unit = DType.EMPTY; final List decimal128Columns = new ArrayList<>(); @@ -99,6 +104,19 @@ public Builder decimal128Column(String... names) { return this; } + /** + * Convert DATE columns explicitly marked JULIAN_GREGORIAN to cuDF proleptic dates. + * Disabled by default for callers that perform their own rebasing. Missing and UNKNOWN + * calendar fields and TIMESTAMP columns are unchanged. Conversion applies to nested DATE + * columns and precedes NumPy promotion. Unrepresentable dates cause a read error. + * @param convert whether DATE conversion is enabled + * @return builder for chaining + */ + public Builder convertDatesToProleptic(boolean convert) { + this.convertDatesToProleptic = convert; + return this; + } + public Builder ignoreTimezoneInStripeFooter() { this.ignoreTimezoneInStripeFooter = true; return this; diff --git a/java/src/main/java/ai/rapids/cudf/ORCWriterOptions.java b/java/src/main/java/ai/rapids/cudf/ORCWriterOptions.java index 21f186ffc5d6..98d7c47190c0 100644 --- a/java/src/main/java/ai/rapids/cudf/ORCWriterOptions.java +++ b/java/src/main/java/ai/rapids/cudf/ORCWriterOptions.java @@ -12,13 +12,27 @@ * that will be used by the ORC writer to write the file. */ public class ORCWriterOptions extends CompressionMetadataWriterOptions { + /** Calendar recorded in the standard ORC file footer. */ + public enum CalendarKind { + JULIAN_GREGORIAN(1), + PROLEPTIC_GREGORIAN(2); + + final int nativeId; + + CalendarKind(int nativeId) { + this.nativeId = nativeId; + } + } + private int stripeSizeRows; private String writerTimezone; + private CalendarKind calendarKind; private ORCWriterOptions(Builder builder) { super(builder); this.stripeSizeRows = builder.stripeSizeRows; this.writerTimezone = builder.writerTimezone; + this.calendarKind = builder.calendarKind; } public static Builder builder() { @@ -33,11 +47,16 @@ public String getWriterTimezone() { return writerTimezone; } + public CalendarKind getCalendarKind() { + return calendarKind; + } + public static class Builder extends CompressionMetadataWriterOptions.Builder { // < 1M rows default orc stripe rows, defined in cudf/cpp/include/cudf/io/orc.hpp private int stripeSizeRows = 1000000; private String writerTimezone = "UTC"; + private CalendarKind calendarKind = CalendarKind.PROLEPTIC_GREGORIAN; public Builder withStripeSizeRows(int stripeSizeRows) { // maximum stripe size cannot be smaller than 512 @@ -69,6 +88,25 @@ public Builder withWriterTimezone(String writerTimezone) { return this; } + /** + * Sets the calendar used for DATE encoding and recorded in the file footer. + * Legacy mode rebases nested and top-level DATE values. It rejects TIMESTAMP columns, + * dates in October 5-14, 1582, and dates outside the int32 hybrid domain. + * Use {@link ORCOptions.Builder#convertDatesToProleptic(boolean)} to read legacy output + * back as cuDF proleptic Gregorian dates. + * + * @param calendarKind calendar kind + * @throws IllegalArgumentException if the calendar is null + */ + public Builder withCalendarKind(CalendarKind calendarKind) { + if (calendarKind == null) { + throw new IllegalArgumentException( + "ORC calendar must not be null"); + } + this.calendarKind = calendarKind; + return this; + } + public ORCWriterOptions build() { return new ORCWriterOptions(this); } diff --git a/java/src/main/java/ai/rapids/cudf/Table.java b/java/src/main/java/ai/rapids/cudf/Table.java index 78061d5710b6..0d853314822f 100644 --- a/java/src/main/java/ai/rapids/cudf/Table.java +++ b/java/src/main/java/ai/rapids/cudf/Table.java @@ -342,12 +342,14 @@ private static native long[] readORC(String[] filterColumnNames, String filePath, long address, long length, boolean usingNumPyTypes, int timeUnit, String[] decimal128Columns, - boolean ignoreTimezoneInStripeFooter) throws CudfException; + boolean ignoreTimezoneInStripeFooter, + boolean convertDatesToProleptic) throws CudfException; private static native long[] readORCFromDataSource(String[] filterColumnNames, boolean usingNumPyTypes, int timeUnit, String[] decimal128Columns, boolean ignoreTimezoneInStripeFooter, + boolean convertDatesToProleptic, long dataSourceHandle) throws CudfException; /** @@ -363,6 +365,7 @@ private static native long[] readORCFromDataSource(String[] filterColumnNames, * the columns * @param isMapValues true if a column is a map * @param writerTimezone timezone that the written timestamps are relative to + * @param calendarKind native calendar kind recorded in the ORC footer * @param filename local output path * @return a handle that is used in later calls to writeORCChunk and writeORCEnd. */ @@ -377,6 +380,7 @@ private static native long writeORCFileBegin(String[] columnNames, boolean[] isMapValues, int stripeSizeRows, String writerTimezone, + int calendarKind, String filename) throws CudfException; /** @@ -392,6 +396,7 @@ private static native long writeORCFileBegin(String[] columnNames, * the columns * @param isMapValues true if a column is a map * @param writerTimezone timezone that the written timestamps are relative to + * @param calendarKind native calendar kind recorded in the ORC footer * @param consumer consumer of host buffers produced. * @return a handle that is used in later calls to writeORCChunk and writeORCEnd. */ @@ -406,6 +411,7 @@ private static native long writeORCBufferBegin(String[] columnNames, boolean[] isMapValues, int stripeSizeRows, String writerTimezone, + int calendarKind, HostBufferConsumer consumer, HostMemoryAllocator hostMemoryAllocator ) throws CudfException; @@ -1519,7 +1525,8 @@ public static Table readORC(ORCOptions opts, File path) { return new Table(readORC(opts.getIncludeColumnNames(), path.getAbsolutePath(), 0, 0, opts.usingNumPyTypes(), opts.timeUnit().typeId.getNativeId(), - opts.getDecimal128Columns(), opts.ignoreTimezoneInStripeFooter())); + opts.getDecimal128Columns(), opts.ignoreTimezoneInStripeFooter(), + opts.convertDatesToProleptic())); } /** @@ -1588,7 +1595,8 @@ public static Table readORC(ORCOptions opts, HostMemoryBuffer buffer, return new Table(readORC(opts.getIncludeColumnNames(), null, buffer.getAddress() + offset, len, opts.usingNumPyTypes(), opts.timeUnit().typeId.getNativeId(), - opts.getDecimal128Columns(), opts.ignoreTimezoneInStripeFooter())); + opts.getDecimal128Columns(), opts.ignoreTimezoneInStripeFooter(), + opts.convertDatesToProleptic())); } public static Table readORC(ORCOptions opts, DataSource ds) { @@ -1597,7 +1605,7 @@ public static Table readORC(ORCOptions opts, DataSource ds) { return new Table(readORCFromDataSource(opts.getIncludeColumnNames(), opts.usingNumPyTypes(), opts.timeUnit().typeId.getNativeId(), opts.getDecimal128Columns(), opts.ignoreTimezoneInStripeFooter(), - dataSourceHandle)); + opts.convertDatesToProleptic(), dataSourceHandle)); } finally { DataSourceHelper.destroyWrapperDataSource(dataSourceHandle); } @@ -1698,6 +1706,7 @@ private ORCTableWriter(ORCWriterOptions options, File outputFile) { options.getFlatIsMap(), options.getStripeSizeRows(), options.getWriterTimezone(), + options.getCalendarKind().nativeId, outputFile.getAbsolutePath())); this.consumer = null; } @@ -1715,6 +1724,7 @@ private ORCTableWriter(ORCWriterOptions options, HostBufferConsumer consumer, options.getFlatIsMap(), options.getStripeSizeRows(), options.getWriterTimezone(), + options.getCalendarKind().nativeId, consumer, hostMemoryAllocator)); this.consumer = consumer; } diff --git a/java/src/main/native/src/TableJni.cpp b/java/src/main/native/src/TableJni.cpp index 405af0b85f8e..f0c202473640 100644 --- a/java/src/main/native/src/TableJni.cpp +++ b/java/src/main/native/src/TableJni.cpp @@ -2493,6 +2493,7 @@ Java_ai_rapids_cudf_Table_readORCFromDataSource(JNIEnv* env, jint unit, jobjectArray dec128_col_names, jboolean ignore_timezone_in_stripe_footer, + jboolean convert_dates_to_proleptic, jlong ds_handle) { JNI_NULL_CHECK(env, ds_handle, "no data source handle given", 0); @@ -2519,6 +2520,7 @@ Java_ai_rapids_cudf_Table_readORCFromDataSource(JNIEnv* env, .timestamp_type(cudf::data_type(static_cast(unit))) .decimal128_columns(n_dec128_col_names.as_cpp_vector()) .ignore_timezone_in_stripe_footer(ignore_timezone_in_stripe_footer) + .convert_dates_to_proleptic(convert_dates_to_proleptic) .build(); return convert_table_for_return(env, cudf::io::read_orc(opts).tbl); } @@ -2535,7 +2537,8 @@ Java_ai_rapids_cudf_Table_readORC(JNIEnv* env, jboolean usingNumPyTypes, jint unit, jobjectArray dec128_col_names, - jboolean ignore_timezone_in_stripe_footer) + jboolean ignore_timezone_in_stripe_footer, + jboolean convert_dates_to_proleptic) { bool read_buffer = true; if (buffer == 0) { @@ -2579,6 +2582,7 @@ Java_ai_rapids_cudf_Table_readORC(JNIEnv* env, .timestamp_type(cudf::data_type(static_cast(unit))) .decimal128_columns(n_dec128_col_names.as_cpp_vector()) .ignore_timezone_in_stripe_footer(ignore_timezone_in_stripe_footer) + .convert_dates_to_proleptic(convert_dates_to_proleptic) .build(); return convert_table_for_return(env, cudf::io::read_orc(opts).tbl); } @@ -2599,6 +2603,7 @@ Java_ai_rapids_cudf_Table_writeORCBufferBegin(JNIEnv* env, jbooleanArray j_is_map, jint j_stripe_size_rows, jstring j_writer_timezone, + jint j_calendar_kind, jobject consumer, jobject host_memory_allocator) { @@ -2650,16 +2655,18 @@ Java_ai_rapids_cudf_Table_writeORCBufferBegin(JNIEnv* env, new cudf::jni::jni_writer_data_sink(env, consumer, host_memory_allocator)); sink_info sink{data_sink.get()}; - auto stats = std::make_shared(); - chunked_orc_writer_options opts = chunked_orc_writer_options::builder(sink) - .metadata(std::move(metadata)) - .compression(static_cast(j_compression)) - .enable_statistics(ORC_STATISTICS_ROW_GROUP) - .key_value_metadata(kv_metadata) - .compression_statistics(stats) - .stripe_size_rows(j_stripe_size_rows) - .writer_timezone(writer_timezone.get()) - .build(); + auto stats = std::make_shared(); + chunked_orc_writer_options opts = + chunked_orc_writer_options::builder(sink) + .metadata(std::move(metadata)) + .compression(static_cast(j_compression)) + .enable_statistics(ORC_STATISTICS_ROW_GROUP) + .key_value_metadata(kv_metadata) + .compression_statistics(stats) + .stripe_size_rows(j_stripe_size_rows) + .writer_timezone(writer_timezone.get()) + .calendar_kind(static_cast(j_calendar_kind)) + .build(); auto writer_ptr = std::make_unique(opts); cudf::jni::native_orc_writer_handle* ret = new cudf::jni::native_orc_writer_handle( std::move(writer_ptr), std::move(data_sink), std::move(stats)); @@ -2681,6 +2688,7 @@ JNIEXPORT long JNICALL Java_ai_rapids_cudf_Table_writeORCFileBegin(JNIEnv* env, jbooleanArray j_is_map, jint j_stripe_size_rows, jstring j_writer_timezone, + jint j_calendar_kind, jstring j_output_path) { JNI_NULL_CHECK(env, j_col_names, "null columns", 0); @@ -2728,16 +2736,18 @@ JNIEXPORT long JNICALL Java_ai_rapids_cudf_Table_writeORCFileBegin(JNIEnv* env, [](std::string const& k, std::string const& v) { return std::make_pair(k, v); }); sink_info sink{output_path}; - auto stats = std::make_shared(); - chunked_orc_writer_options opts = chunked_orc_writer_options::builder(sink) - .metadata(std::move(metadata)) - .compression(static_cast(j_compression)) - .enable_statistics(ORC_STATISTICS_ROW_GROUP) - .key_value_metadata(kv_metadata) - .compression_statistics(stats) - .stripe_size_rows(j_stripe_size_rows) - .writer_timezone(writer_timezone.get()) - .build(); + auto stats = std::make_shared(); + chunked_orc_writer_options opts = + chunked_orc_writer_options::builder(sink) + .metadata(std::move(metadata)) + .compression(static_cast(j_compression)) + .enable_statistics(ORC_STATISTICS_ROW_GROUP) + .key_value_metadata(kv_metadata) + .compression_statistics(stats) + .stripe_size_rows(j_stripe_size_rows) + .writer_timezone(writer_timezone.get()) + .calendar_kind(static_cast(j_calendar_kind)) + .build(); auto writer_ptr = std::make_unique(opts); cudf::jni::native_orc_writer_handle* ret = new cudf::jni::native_orc_writer_handle(std::move(writer_ptr), nullptr, std::move(stats)); diff --git a/java/src/test/interop/ai/rapids/cudf/ORCCalendarInteropTest.java b/java/src/test/interop/ai/rapids/cudf/ORCCalendarInteropTest.java new file mode 100644 index 000000000000..a74728a6b611 --- /dev/null +++ b/java/src/test/interop/ai/rapids/cudf/ORCCalendarInteropTest.java @@ -0,0 +1,320 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +package ai.rapids.cudf; + +import org.apache.hadoop.conf.Configuration; +import org.apache.hadoop.fs.Path; +import org.apache.hadoop.hive.ql.exec.vector.ListColumnVector; +import org.apache.hadoop.hive.ql.exec.vector.LongColumnVector; +import org.apache.hadoop.hive.ql.exec.vector.StructColumnVector; +import org.apache.hadoop.hive.ql.exec.vector.VectorizedRowBatch; +import org.apache.orc.DateColumnStatistics; +import org.apache.orc.OrcFile; +import org.apache.orc.Reader; +import org.apache.orc.RecordReader; +import org.junit.jupiter.api.Test; + +import java.io.ByteArrayOutputStream; +import java.io.IOException; +import java.nio.file.Files; +import java.time.LocalDate; +import java.util.Arrays; +import java.util.Collections; + +import static ai.rapids.cudf.AssertUtils.assertTablesAreEqual; +import static org.junit.jupiter.api.Assertions.assertEquals; +import static org.junit.jupiter.api.Assertions.assertFalse; +import static org.junit.jupiter.api.Assertions.assertThrows; +import static org.junit.jupiter.api.Assertions.assertTrue; + +class ORCCalendarInteropTest { + private static final Integer[] DAYS = Arrays.stream(new String[] { + "1001-01-01", "1200-01-01", "1582-10-04", "1582-10-10", "1582-10-15", + "1970-01-01", "2024-05-06", null + }).map(value -> value == null ? null : Math.toIntExact(LocalDate.parse(value).toEpochDay())) + .toArray(Integer[]::new); + + @Test + @SuppressWarnings("unchecked") + void datesInteroperateWithApacheOrc() throws IOException { + checkDates(DAYS, false); + } + + @Test + void legacyDatesInteroperateWithApacheOrc() throws IOException { + Integer[] days = Arrays.stream(new String[] { + "0001-01-01", "0100-03-01", "0400-02-29", "1001-01-01", "1200-01-01", "1500-02-28", + "1500-03-01", "1582-10-04", "1582-10-15", "1970-01-01", "2024-05-06", null + }).map(value -> value == null ? null : Math.toIntExact(LocalDate.parse(value).toEpochDay())) + .toArray(Integer[]::new); + checkDates(days, true); + } + + @Test + void apacheLegacyDatesReadAsProleptic() throws IOException { + // Independent Apache ORC conversion provides the encoded day oracle, including BCE dates + // and a deterministic wide sample within Apache HybridChronology's negative-year range. + java.util.Random random = new java.util.Random(23794); + Integer[] proleptic = new Integer[4096]; + Integer[] hybrid = new Integer[proleptic.length]; + // Hive 2.8.1's vector calendar conversion uses SimpleDateFormat without an era field. + // Check BCE encoding against Apache ORC's DateUtils directly, which preserves the era. + int[] anchors = Arrays.stream(new String[] { + "-10000-02-28", "-0400-02-29", "-0001-03-01", "0000-02-29", "0001-01-01" + }).mapToInt(value -> Math.toIntExact(LocalDate.parse(value).toEpochDay())).toArray(); + for (int i = 0; i < proleptic.length; ++i) { + int day = i < anchors.length ? anchors[i] : random.nextInt(); + if (day < -365000000) { day %= 365000000; } + if (day < Integer.MIN_VALUE + 100000 || (day >= -141437 && day < -141427)) { + day = 0; + } + proleptic[i] = day; + hybrid[i] = org.apache.orc.impl.DateUtils.convertDateToHybrid(day); + } + java.nio.file.Path file = Files.createTempFile("apache-legacy", ".orc"); + try { + org.apache.orc.TypeDescription schema = org.apache.orc.TypeDescription.fromString("struct"); + try (org.apache.orc.Writer writer = OrcFile.createWriter(new Path(file.toUri()), + OrcFile.writerOptions(new Configuration()).overwrite(true).setSchema(schema).setProlepticGregorian(false) + .compress(org.apache.orc.CompressionKind.NONE))) { + VectorizedRowBatch batch = schema.createRowBatch(); + for (int i = 0; i < hybrid.length; ++i) { + ((LongColumnVector) batch.cols[0]).vector[batch.size++] = hybrid[i]; + if (batch.size == batch.getMaxSize()) { + writer.addRowBatch(batch); + batch.reset(); + } + } + if (batch.size > 0) { writer.addRowBatch(batch); } + } + try (ColumnVector dates = ColumnVector.timestampDaysFromBoxedInts(proleptic); + ColumnVector encoded = ColumnVector.timestampDaysFromBoxedInts(hybrid); + Table expected = new Table(dates); + Table rawExpected = new Table(encoded); + Table actual = Table.readORC(ORCOptions.builder().withNumPyTypes(false) + .convertDatesToProleptic(true).build(), file.toFile()); + Table raw = Table.readORC(ORCOptions.builder().withNumPyTypes(false).build(), file.toFile())) { + assertTablesAreEqual(expected, actual); + assertTablesAreEqual(rawExpected, raw); + java.nio.file.Path gpuFile = Files.createTempFile("gpu-legacy", ".orc"); + try { + try (TableWriter writer = Table.writeORCChunked(ORCWriterOptions.builder() + .withColumns(true, "d").withCalendarKind(ORCWriterOptions.CalendarKind.JULIAN_GREGORIAN) + .build(), gpuFile.toFile())) { + writer.write(expected); + } + try (Table gpuRaw = Table.readORC(ORCOptions.builder().withNumPyTypes(false).build(), + gpuFile.toFile())) { + assertTablesAreEqual(rawExpected, gpuRaw); + } + } finally { Files.deleteIfExists(gpuFile); } + } + byte[] originalBytes = Files.readAllBytes(file); + for (org.apache.orc.OrcProto.CalendarKind calendar : new org.apache.orc.OrcProto.CalendarKind[] { + null, org.apache.orc.OrcProto.CalendarKind.UNKNOWN_CALENDAR, + org.apache.orc.OrcProto.CalendarKind.PROLEPTIC_GREGORIAN, + org.apache.orc.OrcProto.CalendarKind.JULIAN_GREGORIAN}) { + byte[] variant = withCalendar(originalBytes, calendar); + try (ColumnVector dates = ColumnVector.timestampDaysFromBoxedInts( + calendar == org.apache.orc.OrcProto.CalendarKind.JULIAN_GREGORIAN ? proleptic : hybrid); + Table expected = new Table(dates); + Table actual = Table.readORC(ORCOptions.builder().withNumPyTypes(false) + .convertDatesToProleptic(true).build(), variant)) { + assertTablesAreEqual(expected, actual); + } + } + // Unknown protobuf enum numbers must not alias recognized calendars after narrowing. + for (int calendar : new int[] {3, 257, 258, -1}) { + byte[] variant = withCalendarValue(originalBytes, calendar); + try (ColumnVector dates = ColumnVector.timestampDaysFromBoxedInts(hybrid); + Table expected = new Table(dates); + Table actual = Table.readORC(ORCOptions.builder().withNumPyTypes(false) + .convertDatesToProleptic(true).build(), variant)) { + assertTablesAreEqual(expected, actual); + } + } + Long[] millis = Arrays.stream(proleptic).map(day -> day.longValue() * 86400000L) + .toArray(Long[]::new); + try (ColumnVector dates = ColumnVector.timestampMilliSecondsFromBoxedLongs(millis); + Table expected = new Table(dates); + Table actual = Table.readORC(ORCOptions.builder().convertDatesToProleptic(true).build(), + Files.readAllBytes(file))) { + assertTablesAreEqual(expected, actual); + } + } finally { Files.deleteIfExists(file); } + } + + @Test + void rejectsJulianOnlyLeapDayOnConversion() throws IOException { + int day = org.apache.orc.impl.DateUtils.parseDate("1500-02-29", false); + java.nio.file.Path file = Files.createTempFile("apache-julian-leap", ".orc"); + try { + org.apache.orc.TypeDescription schema = org.apache.orc.TypeDescription.fromString("struct"); + try (org.apache.orc.Writer writer = OrcFile.createWriter(new Path(file.toUri()), + OrcFile.writerOptions(new Configuration()).overwrite(true).setSchema(schema) + .setProlepticGregorian(false))) { + VectorizedRowBatch batch = schema.createRowBatch(); + ((LongColumnVector) batch.cols[0]).vector[0] = day; + batch.size = 1; + writer.addRowBatch(batch); + } + try (Table raw = Table.readORC(ORCOptions.builder().withNumPyTypes(false).build(), file.toFile())) { + assertEquals(1, raw.getRowCount()); + } + assertThrows(CudfException.class, () -> Table.readORC(ORCOptions.builder() + .convertDatesToProleptic(true).build(), file.toFile())); + } finally { Files.deleteIfExists(file); } + } + + @SuppressWarnings("unchecked") + private void checkDates(Integer[] days, boolean legacy) throws IOException { + HostColumnVector.ListType listType = new HostColumnVector.ListType(true, + new HostColumnVector.BasicType(true, DType.TIMESTAMP_DAYS)); + try (ColumnVector dates = ColumnVector.timestampDaysFromBoxedInts(days); + ColumnVector structs = ColumnVector.makeStruct(dates); + ColumnVector lists = ColumnVector.fromLists(listType, + Arrays.stream(days).map(Collections::singletonList).toArray(java.util.List[]::new)); + Table input = new Table(dates, structs, lists)) { + for (CompressionType compression : new CompressionType[] { + CompressionType.NONE, CompressionType.SNAPPY}) { + for (boolean bufferSink : new boolean[] {false, true}) { + for (boolean explicitCalendar : new boolean[] {false, true}) { + if (legacy && !explicitCalendar) { continue; } + java.nio.file.Path file = Files.createTempFile("orc-calendar", ".orc"); + try { + ORCWriterOptions.Builder builder = ORCWriterOptions.builder() + .withColumns(true, "d") + .withStructColumn(ColumnWriterOptions.structBuilder("s") + .withColumns(true, "d").build()) + .withListColumn(ColumnWriterOptions.listBuilder("l") + .withColumns(true, "d").build()) + .withCompressionType(compression); + if (legacy || explicitCalendar) { + builder.withCalendarKind(legacy ? ORCWriterOptions.CalendarKind.JULIAN_GREGORIAN + : ORCWriterOptions.CalendarKind.PROLEPTIC_GREGORIAN); + } + ORCWriterOptions options = builder.build(); + if (bufferSink) { + ByteArrayOutputStream output = new ByteArrayOutputStream(); + HostBufferConsumer consumer = (buffer, length) -> { + try { + byte[] bytes = new byte[Math.toIntExact(length)]; + buffer.getBytes(bytes, 0, 0, bytes.length); + output.write(bytes, 0, bytes.length); + } finally { + buffer.close(); + } + }; + try (TableWriter writer = Table.writeORCChunked(options, consumer)) { + writer.write(input); + writer.write(input); + } + Files.write(file, output.toByteArray()); + } else { + try (TableWriter writer = Table.writeORCChunked(options, file.toFile())) { + writer.write(input); + writer.write(input); + } + } + try (Table expected = Table.concatenate(input, input); + Table actual = Table.readORC( + ORCOptions.builder().withNumPyTypes(false).convertDatesToProleptic(true).build(), file.toFile())) { + assertTablesAreEqual(expected, actual); + } + // Both missing-field defaults must be overridden by the file's explicit calendar. + for (boolean defaultProleptic : new boolean[] {false, true}) { + Configuration conf = new Configuration(); + conf.setBoolean("orc.proleptic.gregorian.default", defaultProleptic); + try (Reader reader = OrcFile.createReader(new Path(file.toUri()), + OrcFile.readerOptions(conf).convertToProlepticGregorian(true))) { + assertEquals(!legacy, reader.writerUsedProlepticGregorian()); + assertEquals(days.length * 2, reader.getNumberOfRows()); + // DATE leaves in the top-level column, STRUCT, and LIST. + for (int columnId : new int[] {1, 3, 5}) { + DateColumnStatistics stats = + (DateColumnStatistics) reader.getStatistics()[columnId]; + assertEquals(days[0].longValue(), stats.getMinimumDayOfEpoch()); + assertEquals(Arrays.stream(days).filter(java.util.Objects::nonNull).mapToLong(Integer::longValue) + .max().getAsLong(), stats.getMaximumDayOfEpoch()); + assertEquals((days.length - 1) * 2, stats.getNumberOfValues()); + } + VectorizedRowBatch batch = reader.getSchema().createRowBatch(); + int rowCount = 0; + try (RecordReader rows = reader.rows()) { + while (rows.nextBatch(batch)) { + LongColumnVector top = (LongColumnVector) batch.cols[0]; + LongColumnVector nested = (LongColumnVector) + ((StructColumnVector) batch.cols[1]).fields[0]; + ListColumnVector list = (ListColumnVector) batch.cols[2]; + LongColumnVector child = (LongColumnVector) list.child; + for (int row = 0; row < batch.size; row++, rowCount++) { + Integer expected = days[rowCount % days.length]; + assertDate(top, row, expected); + assertDate(nested, row, expected); + assertEquals(1, list.lengths[row]); + assertDate(child, Math.toIntExact(list.offsets[row]), expected); + } + } + } + assertEquals(days.length * 2, rowCount); + } + } + } finally { + Files.deleteIfExists(file); + } + } + } + } + } + } + + // Rewrite only the uncompressed footer using Apache's independently generated protobuf classes. + private static byte[] withCalendar(byte[] file, org.apache.orc.OrcProto.CalendarKind calendar) + throws IOException { + return withCalendarValue(file, calendar == null ? null : calendar.getNumber()); + } + + private static byte[] withCalendarValue(byte[] file, Integer calendar) throws IOException { + int postscriptLength = Byte.toUnsignedInt(file[file.length - 1]); + int postscriptStart = file.length - 1 - postscriptLength; + org.apache.orc.OrcProto.PostScript ps = org.apache.orc.OrcProto.PostScript.parseFrom( + Arrays.copyOfRange(file, postscriptStart, file.length - 1)); + int footerStart = postscriptStart - Math.toIntExact(ps.getFooterLength()); + org.apache.orc.OrcProto.Footer.Builder footer = org.apache.orc.OrcProto.Footer.parseFrom( + Arrays.copyOfRange(file, footerStart, postscriptStart)).toBuilder(); + footer.clearCalendar(); + ByteArrayOutputStream rewrittenFooter = new ByteArrayOutputStream(); + rewrittenFooter.write(footer.build().toByteArray()); + if (calendar != null) { + rewrittenFooter.write(0x58); // Footer.calendar, field 11, varint. + long value = Integer.toUnsignedLong(calendar); + while ((value & ~0x7fL) != 0) { + rewrittenFooter.write((int) (value & 0x7f) | 0x80); + value >>>= 7; + } + rewrittenFooter.write((int) value); + } + byte[] footerBytes = rewrittenFooter.toByteArray(); + byte[] psBytes = ps.toBuilder().setFooterLength(footerBytes.length).build().toByteArray(); + ByteArrayOutputStream output = new ByteArrayOutputStream(); + output.write(file, 0, footerStart); + output.write(footerBytes); + output.write(psBytes); + output.write(psBytes.length); + return output.toByteArray(); + } + + private static void assertDate(LongColumnVector column, int row, Integer expected) { + int index = column.isRepeating ? 0 : row; + if (expected == null) { + assertFalse(column.noNulls); + assertTrue(column.isNull[index]); + } else { + assertTrue(column.noNulls || !column.isNull[index]); + assertEquals(expected.longValue(), column.vector[index]); + } + } +} diff --git a/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java b/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java new file mode 100644 index 000000000000..7eda5ee6c88f --- /dev/null +++ b/java/src/test/java/ai/rapids/cudf/ORCCalendarTest.java @@ -0,0 +1,100 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +package ai.rapids.cudf; + +import org.junit.jupiter.api.Test; + +import java.io.IOException; +import java.nio.file.Files; +import java.time.LocalDate; +import java.util.Arrays; +import java.util.Collections; + +import static ai.rapids.cudf.AssertUtils.assertTablesAreEqual; +import static org.junit.jupiter.api.Assertions.assertEquals; +import static org.junit.jupiter.api.Assertions.assertFalse; +import static org.junit.jupiter.api.Assertions.assertThrows; +import static org.junit.jupiter.api.Assertions.assertTrue; + +class ORCCalendarTest { + @Test + void rejectsUnsupportedCalendars() { + assertEquals(ORCWriterOptions.CalendarKind.PROLEPTIC_GREGORIAN, + ORCWriterOptions.builder().build().getCalendarKind()); + assertThrows(IllegalArgumentException.class, + () -> ORCWriterOptions.builder().withCalendarKind(null)); + assertEquals(ORCWriterOptions.CalendarKind.JULIAN_GREGORIAN, ORCWriterOptions.builder() + .withCalendarKind(ORCWriterOptions.CalendarKind.JULIAN_GREGORIAN).build().getCalendarKind()); + } + + @Test + void readerConversionIsOptIn() { + assertFalse(ORCOptions.builder().build().convertDatesToProleptic()); + assertTrue(ORCOptions.builder().convertDatesToProleptic(true).build() + .convertDatesToProleptic()); + assertFalse(ORCOptions.builder().convertDatesToProleptic(true) + .convertDatesToProleptic(false).build().convertDatesToProleptic()); + } + + @Test + void prolepticDatesRoundTrip() throws IOException { + Integer[] days = epochDays("1001-01-01", "1200-01-01", "1582-10-04", "1582-10-10", + "1582-10-15", "1970-01-01", "2024-05-06", null); + checkRoundTrip(ORCWriterOptions.CalendarKind.PROLEPTIC_GREGORIAN, days, days); + } + + @Test + void legacyDatesRoundTrip() throws IOException { + Integer[] days = epochDays("1001-01-01", "1200-01-01", "1582-10-04", "1582-10-15", + "1970-01-01", "2024-05-06", null); + // Independent hybrid epoch-day anchors: raw legacy reads must not silently rebase. + Integer[] encoded = {-353914, -281230, -141428, -141427, 0, 19849, null}; + checkRoundTrip(ORCWriterOptions.CalendarKind.JULIAN_GREGORIAN, days, encoded); + } + + private static Integer[] epochDays(String... values) { + return Arrays.stream(values) + .map(value -> value == null ? null : Math.toIntExact(LocalDate.parse(value).toEpochDay())) + .toArray(Integer[]::new); + } + + @SuppressWarnings("unchecked") + private static void checkRoundTrip(ORCWriterOptions.CalendarKind calendar, Integer[] days, + Integer[] encoded) throws IOException { + HostColumnVector.ListType listType = new HostColumnVector.ListType(true, + new HostColumnVector.BasicType(true, DType.TIMESTAMP_DAYS)); + java.nio.file.Path file = Files.createTempFile("orc-calendar-roundtrip", ".orc"); + try (ColumnVector dates = ColumnVector.timestampDaysFromBoxedInts(days); + ColumnVector structs = ColumnVector.makeStruct(dates); + ColumnVector lists = ColumnVector.fromLists(listType, + Arrays.stream(days).map(Collections::singletonList).toArray(java.util.List[]::new)); + Table input = new Table(dates, structs, lists)) { + ORCWriterOptions options = ORCWriterOptions.builder() + .withColumns(true, "d") + .withStructColumn(ColumnWriterOptions.structBuilder("s") + .withColumns(true, "d").build()) + .withListColumn(ColumnWriterOptions.listBuilder("l") + .withColumns(true, "d").build()) + .withCalendarKind(calendar) + .build(); + try (TableWriter writer = Table.writeORCChunked(options, file.toFile())) { + writer.write(input); + } + try (Table actual = Table.readORC(ORCOptions.builder().withNumPyTypes(false) + .convertDatesToProleptic(true).build(), file.toFile())) { + assertTablesAreEqual(input, actual); + } + try (ColumnVector rawDates = ColumnVector.timestampDaysFromBoxedInts(encoded); + Table expected = new Table(rawDates); + Table raw = Table.readORC(ORCOptions.builder().withNumPyTypes(false) + .includeColumn("d").build(), file.toFile())) { + assertTablesAreEqual(expected, raw); + } + } finally { + Files.deleteIfExists(file); + } + } +}