diff --git a/cpp/include/cudf/io/experimental/variant.hpp b/cpp/include/cudf/io/experimental/variant.hpp index 7c53a89e4e7..43da7a27638 100644 --- a/cpp/include/cudf/io/experimental/variant.hpp +++ b/cpp/include/cudf/io/experimental/variant.hpp @@ -7,6 +7,7 @@ #include #include +#include #include #include #include @@ -14,6 +15,7 @@ #include #include +#include #include /** @@ -48,6 +50,9 @@ namespace io::parquet::experimental { * @param variant_column Struct column (VARIANT materialization) with `list` children * (`metadata`, `value`), plus optional shredded siblings * @param path JSONPath-like path string identifying the target field + * @param status Optional. When provided, filled with `variant_operation_status` values, one per + * row. Must be non-nullable, `UINT8`, and have the same row count as + * `variant_column` * @param stream CUDA stream * @param mr Device memory resource * @return `list` column with the extracted value's encoded bytes. A row is null when the @@ -55,13 +60,16 @@ namespace io::parquet::experimental { * the current value. * * @throws std::invalid_argument on empty path or malformed syntax (`[*]` wildcards, negative - * indices, out-of-range indices, and quoted names inside `[...]` are not supported) + * indices, out-of-range indices, and quoted names inside `[...]` are not supported); or if + * `status` is provided but is nullable, not `UINT8`, or has a different row count than + * `variant_column` */ [[nodiscard]] std::unique_ptr get_variant_field( column_view const& variant_column, std::string_view path, - rmm::cuda_stream_view stream = cudf::get_default_stream(), - rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + std::optional status = std::nullopt, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** * @brief Decode a VARIANT value column's blobs into a typed cuDF column. @@ -72,19 +80,27 @@ namespace io::parquet::experimental { * @param values `list` column of VARIANT-encoded value bytes * @param desired_type Target cuDF type (`STRING`, `INT8`/`INT16`/`INT32`/`INT64`, * `FLOAT32`/`FLOAT64`, or `BOOL8`) + * @param incoming_status Optional status column from a prior `get_variant_field` call. When + * provided, non-success rows are propagated directly to the output without decoding. + * Must be non-nullable, `UINT8`, and have the same row count as `values`. + * @param status Optional. When provided, filled with `variant_operation_status` values, one per + * row. Must be non-nullable, `UINT8`, and have the same row count as `values` * @param stream CUDA stream * @param mr Device memory resource * @return Typed column decoded from the VARIANT value blobs * - * @throws std::invalid_argument if `values` is not a `list` column, or if `desired_type` + * @throws std::invalid_argument if `values` is not a `list` column; if `desired_type` * is not one of the supported types (`STRING`, `INT8`/`INT16`/`INT32`/`INT64`, - * `FLOAT32`/`FLOAT64`, or `BOOL8`) + * `FLOAT32`/`FLOAT64`, or `BOOL8`); or if `incoming_status` or `status` is provided but is + * nullable, not `UINT8`, or has a different row count than `values` */ [[nodiscard]] std::unique_ptr cast_variant( column_view const& values, data_type desired_type, - rmm::cuda_stream_view stream = cudf::get_default_stream(), - rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + std::optional incoming_status = std::nullopt, + std::optional status = std::nullopt, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** * @brief Convenience wrapper: extract a nested object value by path and decode into a typed column. @@ -96,18 +112,23 @@ namespace io::parquet::experimental { * @param path JSONPath-like path string (see `get_variant_field` for syntax) * @param desired_type Target type: `STRING`, `INT8`/`INT16`/`INT32`/`INT64`, * `FLOAT32`/`FLOAT64`, or `BOOL8` + * @param status Optional. When provided, filled with `variant_operation_status` values, one per + * row. Must be non-nullable, `UINT8`, and have the same row count as + * `variant_column` * @param stream CUDA stream * @param mr Device memory resource * @return Column of `desired_type` * - * @throws std::invalid_argument on empty path or malformed syntax + * @throws std::invalid_argument on empty path or malformed syntax; or if `status` is provided but + * is nullable, not `UINT8`, or has a different row count than `variant_column` */ [[nodiscard]] std::unique_ptr extract_variant_field( column_view const& variant_column, std::string_view path, data_type desired_type, - rmm::cuda_stream_view stream = cudf::get_default_stream(), - rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + std::optional status = std::nullopt, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @} */ } // namespace io::parquet::experimental diff --git a/cpp/include/cudf/io/experimental/variant_spec.hpp b/cpp/include/cudf/io/experimental/variant_spec.hpp index 6b71dc57385..aad48f7f1a1 100644 --- a/cpp/include/cudf/io/experimental/variant_spec.hpp +++ b/cpp/include/cudf/io/experimental/variant_spec.hpp @@ -46,4 +46,18 @@ enum class variant_primitive_type : uint8_t { UUID = 20, }; +/** + * @brief VARIANT status types. + */ +enum class variant_operation_status : uint8_t { + SUCCESS = 0, ///< operation completed successfully + ROW_NULL = 1, ///< the SQL row itself was null (no VARIANT data to decode) + MISSING_PATH = 2, ///< the requested path does not exist in the VARIANT + VARIANT_NULL = 3, ///< the value at the path is a VARIANT null + TYPE_MISMATCH = 4, ///< the value's type does not match the requested type + MALFORMED_VARIANT = 5, ///< the VARIANT binary encoding is invalid + OVERFLOW = 6, ///< the value overflows the target numeric type + INVALID_CONVERSION = 7, ///< the value cannot be converted to the requested type +}; + } // namespace cudf::io::parquet::experimental diff --git a/cpp/src/io/parquet/experimental/variant_extract.cu b/cpp/src/io/parquet/experimental/variant_extract.cu index 8327e6c5cde..b4d227abefd 100644 --- a/cpp/src/io/parquet/experimental/variant_extract.cu +++ b/cpp/src/io/parquet/experimental/variant_extract.cu @@ -44,6 +44,7 @@ #include #include +#include #include #include @@ -62,6 +63,9 @@ using basic_type = variant_basic_type; // For a primitive value, the value_header is the physical type id of the payload. using primitive_type = variant_primitive_type; +// The status of a VARIANT operation. +using op_status = variant_operation_status; + __device__ cuda::std::optional read_uint64(device_span data, size_type pos, int width) @@ -72,7 +76,6 @@ __device__ cuda::std::optional read_uint64(device_span return v; } -// Safely narrow a decoded value to size_type __device__ cuda::std::optional narrow_cast(cuda::std::optional value) { if (!value.has_value() || @@ -150,7 +153,6 @@ __device__ cuda::std::optional variant_value_length(device_span(value_header)) { case primitive_type::NULLVAL: @@ -229,46 +231,57 @@ __device__ cuda::std::optional variant_value_length(device_span find_key_in_metadata(device_span meta, - cudf::string_view key) +__device__ cuda::std::pair, op_status> find_key_in_metadata( + device_span meta, cudf::string_view key) { auto const meta_len = static_cast(meta.size()); - if (meta_len < 1) { return cuda::std::nullopt; } + if (meta_len < 1) { return {cuda::std::nullopt, op_status::MALFORMED_VARIANT}; } auto const header = meta[0]; int const version = header & 0x0F; - if (version != variant_version_v1) { return cuda::std::nullopt; } + if (version != variant_version_v1) { return {cuda::std::nullopt, op_status::MALFORMED_VARIANT}; } int const offset_size = ((header >> 6) & 0x03) + 1; size_type pos = 1; auto const num_entries = narrow_cast(read_uint64(meta, pos, offset_size)); - if (!num_entries.has_value()) { return cuda::std::nullopt; } + if (!num_entries.has_value()) { return {cuda::std::nullopt, op_status::MALFORMED_VARIANT}; } pos += offset_size; auto const offsets_start = pos; auto const offsets_bytes = (static_cast(num_entries.value()) + 1) * offset_size; if (cuda::std::cmp_greater(offsets_bytes, meta_len - offsets_start)) { - return cuda::std::nullopt; + return {cuda::std::nullopt, op_status::MALFORMED_VARIANT}; } auto start_off = read_uint64(meta, offsets_start, offset_size); - if (!start_off.has_value()) { return cuda::std::nullopt; } - auto const strings_base = offsets_start + static_cast(offsets_bytes); - // Bytes available for dictionary string payloads + // Parquet VARIANT spec requires offsets[0] == 0; any other value is malformed. + if (!start_off.has_value() || start_off.value() != 0) { + return {cuda::std::nullopt, op_status::MALFORMED_VARIANT}; + } + auto const strings_base = offsets_start + static_cast(offsets_bytes); auto const strings_extent = meta_len - strings_base; + // Read the terminal offset offsets[num_entries] before scanning entries. An early key match + // must be bounded by this declared extent, not the physical buffer, so validate it upfront. + auto const terminal_off_pos = + offsets_start + static_cast(num_entries.value()) * offset_size; + auto const terminal_off = read_uint64(meta, terminal_off_pos, offset_size); + if (!terminal_off.has_value() || cuda::std::cmp_greater(terminal_off.value(), strings_extent)) { + return {cuda::std::nullopt, op_status::MALFORMED_VARIANT}; + } + auto const strings_declared = static_cast(terminal_off.value()); for (size_type i = 0; i < num_entries.value(); ++i) { auto const end_off = read_uint64(meta, offsets_start + (i + 1) * offset_size, offset_size); - if (!end_off.has_value()) { return cuda::std::nullopt; } - if (end_off.value() < start_off.value() || end_off.value() > strings_extent) { - return cuda::std::nullopt; + if (!end_off.has_value()) { return {cuda::std::nullopt, op_status::MALFORMED_VARIANT}; } + if (end_off.value() < start_off.value() || end_off.value() > strings_declared) { + return {cuda::std::nullopt, op_status::MALFORMED_VARIANT}; } cudf::string_view const entry{ reinterpret_cast(meta.data() + strings_base + start_off.value()), static_cast(end_off.value() - start_off.value())}; - if (entry == key) { return i; } + if (entry == key) { return {i, op_status::SUCCESS}; } start_off = end_off; } - return cuda::std::nullopt; + return {cuda::std::nullopt, op_status::MISSING_PATH}; } /** @@ -297,57 +310,71 @@ __device__ cuda::std::optional find_key_in_metadata(device_span locate_object_field(device_span val, int id) +__device__ cuda::std::pair, op_status> locate_object_field( + device_span val, int id) { auto const val_len = static_cast(val.size()); - if (val_len < 1) { return {}; } + if (val_len < 1) { return {{}, op_status::MALFORMED_VARIANT}; } auto const value_metadata = val[0]; - if (decode_basic_type(value_metadata) != basic_type::OBJECT) { return {}; } + if (decode_basic_type(value_metadata) != basic_type::OBJECT) { + return {{}, op_status::MISSING_PATH}; + } auto const [offset_size, id_size, num_elements_size] = decode_object_array_header(variant_value_header(value_metadata), true); size_type pos = 1; auto const num_fields = narrow_cast(read_uint64(val, pos, num_elements_size)); - if (!num_fields.has_value()) { return {}; } + if (!num_fields.has_value()) { return {{}, op_status::MALFORMED_VARIANT}; } pos += num_elements_size; auto const ids_start = pos; auto const ids_bytes = static_cast(num_fields.value()) * id_size; - if (ids_bytes > val_len - ids_start) { return {}; } + if (ids_bytes > val_len - ids_start) { return {{}, op_status::MALFORMED_VARIANT}; } auto const offsets_start = ids_start + static_cast(ids_bytes); auto const offsets_bytes = (static_cast(num_fields.value()) + 1) * offset_size; - if (offsets_bytes > val_len - offsets_start) { return {}; } + if (offsets_bytes > val_len - offsets_start) { return {{}, op_status::MALFORMED_VARIANT}; } - auto const values_base = offsets_start + static_cast(offsets_bytes); - // Maximum legitimate field-offset value: bytes available after values_base + auto const values_base = offsets_start + static_cast(offsets_bytes); auto const values_extent = val_len - values_base; - // Find the matching field ID and its start offset + // Read the sentinel (terminal offset at index num_fields) to get the authoritative end of the + // values region. Using the physical remainder (values_extent) would allow a malformed object + // to reference bytes beyond the sentinel, passing validation despite corrupt data. + auto const sentinel_raw = + read_uint64(val, offsets_start + num_fields.value() * offset_size, offset_size); + if (!sentinel_raw.has_value() || sentinel_raw.value() > static_cast(values_extent)) { + return {{}, op_status::MALFORMED_VARIANT}; + } + auto const values_region = static_cast(sentinel_raw.value()); + bool found = false; uint64_t match_start = 0; for (size_type i = 0; i < num_fields.value(); ++i) { auto const current_id = read_uint64(val, ids_start + i * id_size, id_size); - if (!current_id.has_value()) { return {}; } + if (!current_id.has_value()) { return {{}, op_status::MALFORMED_VARIANT}; } if (cuda::std::cmp_not_equal(current_id.value(), id)) { continue; } auto const match_offset = read_uint64(val, offsets_start + i * offset_size, offset_size); - if (!match_offset.has_value()) { return {}; } - if (match_offset.value() > values_extent) { return {}; } + if (!match_offset.has_value()) { return {{}, op_status::MALFORMED_VARIANT}; } + if (match_offset.value() > static_cast(values_region)) { + return {{}, op_status::MALFORMED_VARIANT}; + } match_start = match_offset.value(); found = true; break; } - if (!found) { return {}; } + if (!found) { return {{}, op_status::MISSING_PATH}; } - // Derive field's value length from its header auto const value = val.subspan(values_base + match_start); auto const value_len = variant_value_length(value); - if (!value_len.has_value()) { return {}; } + if (!value_len.has_value()) { return {{}, op_status::MALFORMED_VARIANT}; } auto const match_end = match_start + value_len.value(); - if (match_end > values_extent) { return {}; } - return val.subspan(values_base + match_start, value_len.value()); + if (match_end > static_cast(values_region)) { + return {{}, op_status::MALFORMED_VARIANT}; + } + return {val.subspan(values_base + match_start, value_len.value()), op_status::SUCCESS}; } // Parse an array value header and return the sub-span of the element at `index` (0-based) within @@ -364,15 +391,17 @@ __device__ device_span locate_object_field(device_span locate_array_element(device_span value, - size_type index) +__device__ cuda::std::pair, op_status> locate_array_element( + device_span value, size_type index) { - if (index < 0) { return {}; } + if (index < 0) { return {{}, op_status::MISSING_PATH}; } auto const value_size = static_cast(value.size()); - if (value_size < 1) { return {}; } + if (value_size < 1) { return {{}, op_status::MALFORMED_VARIANT}; } uint8_t const value_metadata = value[0]; - if (decode_basic_type(value_metadata) != basic_type::ARRAY) { return {}; } + if (decode_basic_type(value_metadata) != basic_type::ARRAY) { + return {{}, op_status::MISSING_PATH}; + } int const value_header = variant_value_header(value_metadata); [[maybe_unused]] auto const [offset_size, _, num_elements_size] = @@ -380,32 +409,60 @@ __device__ device_span locate_array_element(device_span= num_elements) { return {}; } + if (index >= num_elements) { return {{}, op_status::MISSING_PATH}; } position += num_elements_size; size_type const offsets_start = position; + // Computed in 64-bit because (num_elements + 1) * offset_size can exceed the signed `size_type` // range (which would be UB); the check below then rejects any array that overruns the value blob. auto const offsets_bytes = (static_cast(num_elements) + 1) * offset_size; - if (cuda::std::cmp_greater(offsets_bytes, value_size - offsets_start)) { return {}; } + if (cuda::std::cmp_greater(offsets_bytes, value_size - offsets_start)) { + return {{}, op_status::MALFORMED_VARIANT}; + } size_type const values_base = offsets_start + static_cast(offsets_bytes); auto const values_extent = value_size - values_base; + // Read the terminal offset offsets[num_elements]; it is the spec-declared bound on the + // values region and must be used instead of the physical extent so that an element whose + // offset escapes the declared boundary is caught as malformed even when physical bytes + // are present beyond it. + auto const terminal_off_pos = offsets_start + static_cast(num_elements) * offset_size; + auto const terminal_off = read_uint64(value, terminal_off_pos, offset_size); + if (!terminal_off.has_value() || cuda::std::cmp_greater(*terminal_off, values_extent)) { + return {{}, op_status::MALFORMED_VARIANT}; + } + // The spec requires offsets[0] == 0; a nonzero first offset silently skips leading + // value bytes and can return a plausible result from a malformed array. + auto const first_off = read_uint64(value, offsets_start, offset_size); + if (!first_off.has_value() || *first_off != 0) { return {{}, op_status::MALFORMED_VARIANT}; } auto const start_offset_pos = offsets_start + static_cast(index) * offset_size; auto const end_offset_pos = offsets_start + (static_cast(index) + 1) * offset_size; - if (cuda::std::cmp_greater(end_offset_pos + offset_size, value_size)) { return {}; } - + if (cuda::std::cmp_greater(end_offset_pos + offset_size, value_size)) { + return {{}, op_status::MALFORMED_VARIANT}; + } auto const start_offset = read_uint64(value, start_offset_pos, offset_size); auto const end_offset = read_uint64(value, end_offset_pos, offset_size); - if (!start_offset.has_value() || !end_offset.has_value()) { return {}; } + if (!start_offset.has_value() || !end_offset.has_value()) { + return {{}, op_status::MALFORMED_VARIANT}; + } auto const element_start = *start_offset; auto const element_end = *end_offset; - if (element_end < element_start || cuda::std::cmp_greater(element_end, values_extent)) { - return {}; + if (element_end < element_start || cuda::std::cmp_greater(element_end, *terminal_off)) { + return {{}, op_status::MALFORMED_VARIANT}; } - return value.subspan(values_base + element_start, element_end - element_start); + return {value.subspan(values_base + element_start, element_end - element_start), + op_status::SUCCESS}; +} + +__device__ bool is_variant_null(device_span enc) +{ + if (enc.empty()) { return false; } + auto const vm = enc[0]; + return decode_basic_type(vm) == basic_type::PRIMITIVE && + variant_value_header(vm) == static_cast(primitive_type::NULLVAL); } // The fixed-width signed integers a VARIANT value can be cast to: INT{8,16,32,64}. Matches the @@ -515,26 +572,36 @@ __device__ cuda::std::optional parse_index_step(cudf::string_view ste // - "" -> descend into an object by dictionary key, or // - "[]" -> descend into an array by zero-based integer index. // The step kind is inferred from the first byte (`'['` means index). -__device__ device_span resolve_path(device_span meta, - device_span val, - column_device_view path) +__device__ cuda::std::pair, op_status> resolve_path( + device_span meta, device_span val, column_device_view path) { device_span sub_val = val; for (size_type i = 0; i < path.size(); ++i) { auto const step = path.element(i); - if (step.size_bytes() >= 1 && step.data()[0] == '[') { auto const index = parse_index_step(step); - if (!index.has_value()) { return {}; } - sub_val = locate_array_element(sub_val, index.value()); + if (!index.has_value()) { return {{}, op_status::MISSING_PATH}; } + auto const [span, st] = locate_array_element(sub_val, index.value()); + if (st != op_status::SUCCESS) { return {{}, st}; } + sub_val = span; } else { - auto const field_id = find_key_in_metadata(meta, step); - if (!field_id.has_value()) { return {}; } - sub_val = locate_object_field(sub_val, field_id.value()); + auto const [field_id, meta_st] = find_key_in_metadata(meta, step); + if (meta_st == op_status::MALFORMED_VARIANT) { return {{}, op_status::MALFORMED_VARIANT}; } + if (!field_id.has_value()) { return {{}, op_status::MISSING_PATH}; } + auto const [span, st] = locate_object_field(sub_val, field_id.value()); + if (st != op_status::SUCCESS) { return {{}, st}; } + sub_val = span; } - if (sub_val.empty()) { return {}; } + + // VARIANT null before the end of the path is missing_path per spec. + if (i + 1 < path.size() && is_variant_null(sub_val)) { return {{}, op_status::MISSING_PATH}; } + // A zero-length resolved value is not decodable; the value-only path drops the row. + if (sub_val.empty()) { return {{}, op_status::MALFORMED_VARIANT}; } } - return sub_val; + + // Terminal VARIANT null: return the bytes with variant_null status. + if (is_variant_null(sub_val)) { return {sub_val, op_status::VARIANT_NULL}; } + return {sub_val, op_status::SUCCESS}; } __device__ cuda::std::optional> decode_string( @@ -573,7 +640,6 @@ __device__ device_span list_row_span(cudf::lists_column_device_vi return {col.child().data() + begin, static_cast(end - begin)}; } -// Returns the metadata and value list bytes for a given row from device views __device__ cuda::std::pair, device_span> metadata_and_value_at(cudf::lists_column_device_view const& metadata, cudf::lists_column_device_view const& values, @@ -591,13 +657,15 @@ constexpr int block_size = 256; * `d_sizes[row]` and its offset within the row's value blob to `d_src_offsets[row]`. Rows that are * null, or whose path does not resolve, are marked null in `d_null_mask` with a size of 0. */ +template CUDF_KERNEL __launch_bounds__(block_size) void locate_variant_fields_kernel( cudf::lists_column_device_view metadata, cudf::lists_column_device_view values, column_device_view path, device_span d_sizes, device_span d_src_offsets, - bitmask_type* d_null_mask) + bitmask_type* d_null_mask, + device_span d_status) // only read/written when HasStatus { auto const num_rows = static_cast(d_sizes.size()); auto const tid = cudf::detail::grid_1d::global_thread_id(); @@ -607,24 +675,74 @@ CUDF_KERNEL __launch_bounds__(block_size) void locate_variant_fields_kernel( if (!cudf::bit_is_set(d_null_mask, row)) { d_sizes[row] = 0; d_src_offsets[row] = 0; + if constexpr (HasStatus) { d_status[row] = op_status::ROW_NULL; } continue; } auto const [meta, val] = metadata_and_value_at(metadata, values, row); + auto const [field, st] = resolve_path(meta, val, path); - auto const field = resolve_path(meta, val, path); if (field.empty()) { d_sizes[row] = 0; d_src_offsets[row] = 0; cudf::clear_bit(d_null_mask, row); - continue; + if constexpr (HasStatus) { d_status[row] = st; } + } else { + d_sizes[row] = static_cast(field.size()); + d_src_offsets[row] = static_cast(field.data() - val.data()); + if constexpr (HasStatus) { d_status[row] = st; } } + } +} - d_sizes[row] = static_cast(field.size()); - d_src_offsets[row] = static_cast(field.data() - val.data()); +// Returns true for every primitive_type ID that variant_value_length maps to a known payload +// size in its `basic_type::PRIMITIVE` switch, i.e. every ID other than its `default` case. +__device__ bool is_recognized_primitive_type(primitive_type ptype) +{ + switch (ptype) { + case primitive_type::NULLVAL: + case primitive_type::BOOLEAN_TRUE: + case primitive_type::BOOLEAN_FALSE: + case primitive_type::INT8: + case primitive_type::INT16: + case primitive_type::INT32: + case primitive_type::INT64: + case primitive_type::FLOAT64: + case primitive_type::DECIMAL4: + case primitive_type::DECIMAL8: + case primitive_type::DECIMAL16: + case primitive_type::DATE: + case primitive_type::TIMESTAMP_MICROS: + case primitive_type::TIMESTAMP_NTZ_MICROS: + case primitive_type::FLOAT32: + case primitive_type::BINARY: + case primitive_type::LONG_STRING: + case primitive_type::TIME_NTZ_MICROS: + case primitive_type::TIMESTAMP_NANOS: + case primitive_type::TIMESTAMP_NTZ_NANOS: + case primitive_type::UUID: return true; + default: return false; } } +/** + * @brief Status helper for fixed-width primitive targets: classifies why `decode_primitive` + * failed to decode `val`, per `variant_operation_status` semantics. + */ +template + requires(is_variant_numerical) +__device__ op_status cast_status_for_primitive(device_span val) +{ + if (val.empty()) { return op_status::MALFORMED_VARIANT; } + if (is_variant_null(val)) { return op_status::VARIANT_NULL; } + if (decode_primitive(val).has_value()) { return op_status::SUCCESS; } + if (decode_basic_type(val[0]) != basic_type::PRIMITIVE) { return op_status::TYPE_MISMATCH; } + auto const ptype = static_cast(variant_value_header(val[0])); + if (ptype == primitive_type_for()) { return op_status::MALFORMED_VARIANT; } + return is_recognized_primitive_type(ptype) ? op_status::TYPE_MISMATCH + : op_status::MALFORMED_VARIANT; +} + /** * @brief Per-row kernel: decode each VARIANT value blob into a fixed-width primitive of type `T`. * @@ -634,32 +752,97 @@ CUDF_KERNEL __launch_bounds__(block_size) void locate_variant_fields_kernel( * that are null, or whose value is not an exact-width match for `T`, are marked null in * `d_null_mask` with an output of 0. */ -template +// `d_status` may be an empty span when HasStatus=true but no `status` output was requested (only +// incoming_status propagation to the null mask is desired). All writes to d_status are guarded +// by d_status.data() so the empty-span case is safe. +template CUDF_KERNEL __launch_bounds__(block_size) void cast_variant_primitive_kernel( - cudf::lists_column_device_view values, device_span d_output, bitmask_type* d_null_mask) + cudf::lists_column_device_view values, + device_span d_output, + bitmask_type* d_null_mask, + column_device_view incoming_status, // only used when HasStatus + bool has_incoming, // only meaningful when HasStatus + device_span d_status) // only used when HasStatus; may be empty { auto const num_rows = static_cast(d_output.size()); auto const tid = cudf::detail::grid_1d::global_thread_id(); auto const stride = cudf::detail::grid_1d::grid_stride(); for (auto row = tid; row < num_rows; row += stride) { - if (!cudf::bit_is_set(d_null_mask, row)) { - d_output[row] = 0; - continue; + if constexpr (HasStatus) { + if (has_incoming) { + // Status column is always non-nullable; row_null replaces the null bit. + auto const s = static_cast(incoming_status.element(row)); + if (s != op_status::SUCCESS) { + d_output[row] = T{}; + if (cudf::bit_is_set(d_null_mask, row)) { cudf::clear_bit(d_null_mask, row); } + if (d_status.data()) { d_status[row] = s; } + continue; + } + // incoming success → fall through to decode (value null bit is set) + } else { + if (!cudf::bit_is_set(d_null_mask, row)) { + d_output[row] = T{}; + if (d_status.data()) { d_status[row] = op_status::ROW_NULL; } + continue; + } + } + } else { + if (!cudf::bit_is_set(d_null_mask, row)) { + d_output[row] = T{}; + continue; + } } - auto const val = list_row_span(values, row); - + auto const val = list_row_span(values, row); auto const decoded = decode_primitive(val); if (decoded.has_value()) { d_output[row] = *decoded; + if constexpr (HasStatus) { + if (d_status.data()) { d_status[row] = op_status::SUCCESS; } + } } else { - d_output[row] = 0; + d_output[row] = T{}; cudf::clear_bit(d_null_mask, row); + if constexpr (HasStatus) { + if (d_status.data()) { d_status[row] = cast_status_for_primitive(val); } + } } } } +__device__ op_status cast_status_for_bool(device_span val) +{ + if (val.empty()) { return op_status::MALFORMED_VARIANT; } + if (is_variant_null(val)) { return op_status::VARIANT_NULL; } + if (decode_bool(val).has_value()) { return op_status::SUCCESS; } + if (decode_basic_type(val[0]) != basic_type::PRIMITIVE) { return op_status::TYPE_MISMATCH; } + // Boolean values carry no payload, so a BOOLEAN_TRUE/FALSE header can never be truncated; + // decode_bool would have succeeded above. Any remaining primitive ID is a type mismatch when + // recognised, or malformed when not. + auto const ptype = static_cast(variant_value_header(val[0])); + return is_recognized_primitive_type(ptype) ? op_status::TYPE_MISMATCH + : op_status::MALFORMED_VARIANT; +} + +__device__ op_status cast_status_for_string(device_span val) +{ + if (val.empty()) { return op_status::MALFORMED_VARIANT; } + if (is_variant_null(val)) { return op_status::VARIANT_NULL; } + if (decode_string(val).has_value()) { return op_status::SUCCESS; } + auto const btype = decode_basic_type(val[0]); + if (btype == basic_type::SHORT_STRING) { return op_status::MALFORMED_VARIANT; } + if (btype == basic_type::PRIMITIVE) { + auto const ptype = static_cast(variant_value_header(val[0])); + // LONG_STRING is a recognized string type whose payload was truncated. + if (ptype == primitive_type::LONG_STRING) { return op_status::MALFORMED_VARIANT; } + return is_recognized_primitive_type(ptype) ? op_status::TYPE_MISMATCH + : op_status::MALFORMED_VARIANT; + } + // OBJECT, ARRAY, or other non-primitive basic types: well-formed, just not a string. + return op_status::TYPE_MISMATCH; +} + /** * @brief Strings-children functor: decode each VARIANT value blob into a string. * @@ -674,28 +857,52 @@ struct cast_variant_string_fn { size_type* d_sizes; char* d_chars; cudf::detail::input_offsetalator d_offsets; + // Status tracking (optional: d_status non-null to enable; status is always non-nullable) + op_status* d_status{nullptr}; + column_device_view incoming_status; + bool has_incoming{false}; __device__ void operator()(size_type row) { - if (!cudf::bit_is_set(d_null_mask, row)) { - if (!d_chars) { d_sizes[row] = 0; } - return; + // Status is only written on the sizing pass (d_chars == nullptr). On the writing pass the + // null mask may already be cleared from the sizing pass, so we must not re-inspect it to + // write status (that would misidentify a decode-failed row as a SQL-null row). + bool const sizing = (d_chars == nullptr); + + if (has_incoming) { + // Status column is always non-nullable; row_null replaces the null bit. + auto const s = static_cast(incoming_status.element(row)); + if (s != op_status::SUCCESS) { + if (sizing) { d_sizes[row] = 0; } + if (cudf::bit_is_set(d_null_mask, row)) { cudf::clear_bit(d_null_mask, row); } + if (sizing && d_status) { d_status[row] = s; } + return; + } + // incoming success: fall through to decode + } else { + if (!cudf::bit_is_set(d_null_mask, row)) { + if (sizing) { d_sizes[row] = 0; } + if (sizing && d_status) { d_status[row] = op_status::ROW_NULL; } + return; + } } auto const val = list_row_span(d_values, row); auto const str = decode_string(val); if (!str) { - if (!d_chars) { d_sizes[row] = 0; } + if (sizing) { d_sizes[row] = 0; } cudf::clear_bit(d_null_mask, row); + if (sizing && d_status) { d_status[row] = cast_status_for_string(val); } return; } - if (!d_chars) { + if (sizing) { d_sizes[row] = str->size(); } else { cuda::std::memcpy(d_chars + d_offsets[row], str->data(), str->size()); } + if (sizing && d_status) { d_status[row] = op_status::SUCCESS; } } }; @@ -717,16 +924,40 @@ struct cast_variant_fn { rmm::device_buffer null_mask; rmm::cuda_stream_view stream; rmm::device_async_resource_ref mr; + // Optional status tracking + column_device_view incoming_status_view; + bool has_incoming{false}; + op_status* d_status{nullptr}; template std::unique_ptr operator()() requires(is_variant_numerical) { rmm::device_buffer data{num_rows * sizeof(T), stream, mr}; - auto grid = cudf::detail::grid_1d{num_rows, block_size}; - cast_variant_primitive_kernel<<>>( - values, {static_cast(data.data()), static_cast(num_rows)}, d_null_mask); - CUDF_CUDA_TRY(cudaGetLastError()); + auto const grid = cudf::detail::grid_1d{num_rows, block_size}; + auto const d_out = + device_span{static_cast(data.data()), static_cast(num_rows)}; + if (d_status != nullptr) { + cast_variant_primitive_kernel<<>>( + values, + d_out, + d_null_mask, + incoming_status_view, + has_incoming, + {d_status, static_cast(num_rows)}); + CUDF_CUDA_TRY(cudaGetLastError()); + } else if (has_incoming) { + // No status output requested, but incoming_status still needs to be applied to the null + // mask. Use HasStatus=true with an empty d_status span so writes are guarded (no-op). + cast_variant_primitive_kernel<<>>( + values, d_out, d_null_mask, incoming_status_view, true, {}); + CUDF_CUDA_TRY(cudaGetLastError()); + } else { + cast_variant_primitive_kernel<<>>( + values, d_out, d_null_mask, incoming_status_view, false, {}); + CUDF_CUDA_TRY(cudaGetLastError()); + } + auto const null_count = num_rows - cudf::detail::count_set_bits(d_null_mask, 0, num_rows, stream); return std::make_unique(desired_type, @@ -742,19 +973,46 @@ struct cast_variant_fn { { rmm::device_buffer data{num_rows * sizeof(bool), stream, mr}; - thrust::transform( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::counting_iterator(0), - cuda::counting_iterator(num_rows), - static_cast(data.data()), - [values = this->values, d_null_mask = this->d_null_mask] __device__(size_type row) -> bool { - if (!cudf::bit_is_set(d_null_mask, row)) { return false; } - auto const val = list_row_span(values, row); - auto const decoded = decode_bool(val); - if (decoded.has_value()) { return *decoded; } - cudf::clear_bit(d_null_mask, row); - return false; - }); + auto* dp_s = d_status; + + auto const inc_view = incoming_status_view; + auto const hi = has_incoming; + thrust::for_each(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator(0), + cuda::counting_iterator(num_rows), + [vals = this->values, + d_out = static_cast(data.data()), + dnm = this->d_null_mask, + dp_s, + inc_view, + hi] __device__(size_type row) { + // Status column is always non-nullable; row_null replaces the null bit. + if (hi) { + auto const s = static_cast(inc_view.element(row)); + if (s != op_status::SUCCESS) { + d_out[row] = false; + if (cudf::bit_is_set(dnm, row)) { cudf::clear_bit(dnm, row); } + if (dp_s) { dp_s[row] = s; } + return; + } + } else { + if (!cudf::bit_is_set(dnm, row)) { + d_out[row] = false; + if (dp_s) { dp_s[row] = op_status::ROW_NULL; } + return; + } + } + auto const val = list_row_span(vals, row); + auto const decoded = decode_bool(val); + if (decoded.has_value()) { + d_out[row] = *decoded; + if (dp_s) { dp_s[row] = op_status::SUCCESS; } + } else { + d_out[row] = false; + cudf::clear_bit(dnm, row); + if (dp_s) { dp_s[row] = cast_status_for_bool(val); } + } + }); auto const null_count = num_rows - cudf::detail::count_set_bits(d_null_mask, 0, num_rows, stream); @@ -769,9 +1027,11 @@ struct cast_variant_fn { std::unique_ptr operator()() requires(cuda::std::is_same_v) { - cast_variant_string_fn fn{values, d_null_mask, nullptr, nullptr, {}}; + cast_variant_string_fn fn{ + values, d_null_mask, nullptr, nullptr, {}, d_status, incoming_status_view, has_incoming}; auto [offsets_column, chars] = cudf::strings::detail::make_strings_children(fn, num_rows, stream, mr); + auto const null_count = num_rows - cudf::detail::count_set_bits(d_null_mask, 0, num_rows, stream); return make_strings_column(num_rows, @@ -822,10 +1082,10 @@ namespace detail { std::unique_ptr get_variant_field(column_view const& variant_column, std::string_view path, + std::optional status, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - // Validate the variant column CUDF_EXPECTS(variant_column.type().id() == type_id::STRUCT, "VARIANT column must be struct type", std::invalid_argument); @@ -839,6 +1099,18 @@ std::unique_ptr get_variant_field(column_view const& variant_column, auto const steps = parse_variant_path(path); auto const num_rows = variant_column.size(); + + if (status.has_value()) { + CUDF_EXPECTS(!status->nullable(), + "status column must not be nullable; use row_null for SQL-null rows", + std::invalid_argument); + CUDF_EXPECTS( + status->type().id() == type_id::UINT8, "status column must be UINT8", std::invalid_argument); + CUDF_EXPECTS(status->size() == num_rows, + "status column must have the same number of rows as variant_column", + std::invalid_argument); + } + if (num_rows == 0) { return cudf::make_lists_column( 0, make_empty_column(type_id::INT32), make_empty_column(type_id::UINT8), 0, {}); @@ -868,18 +1140,30 @@ std::unique_ptr get_variant_field(column_view const& variant_column, : cudf::create_null_mask(variant_column.size(), mask_state::ALL_VALID, stream, mr); auto* d_null_mask = static_cast(null_mask.data()); - // Parse the path per row and compute the output sizes auto grid = cudf::detail::grid_1d{num_rows, block_size}; - locate_variant_fields_kernel<<>>( - meta_lists_device_view, - val_lists_device_view, - *path_device_view, - d_sizes, - d_src_offsets, - d_null_mask); - CUDF_CUDA_TRY(cudaGetLastError()); - - // Convert sizes to offsets + + if (status.has_value()) { + locate_variant_fields_kernel<<>>( + meta_lists_device_view, + val_lists_device_view, + *path_device_view, + d_sizes, + d_src_offsets, + d_null_mask, + {reinterpret_cast(status->data()), static_cast(num_rows)}); + CUDF_CUDA_TRY(cudaGetLastError()); + } else { + locate_variant_fields_kernel + <<>>(meta_lists_device_view, + val_lists_device_view, + *path_device_view, + d_sizes, + d_src_offsets, + d_null_mask, + {}); + CUDF_CUDA_TRY(cudaGetLastError()); + } + auto [offsets_column, total_bytes] = cudf::strings::detail::make_offsets_child_column(d_sizes.begin(), d_sizes.end(), stream, mr); CUDF_EXPECTS(total_bytes <= std::numeric_limits::max(), @@ -888,7 +1172,6 @@ std::unique_ptr get_variant_field(column_view const& variant_column, device_span d_offsets{offsets_column->view().data(), static_cast(num_rows + 1)}; - // Copy values into the output buffer auto val_child = make_numeric_column( data_type{type_id::UINT8}, total_bytes, mask_state::UNALLOCATED, stream, mr); if (total_bytes > 0) { @@ -919,6 +1202,8 @@ std::unique_ptr get_variant_field(column_view const& variant_column, std::unique_ptr cast_variant(column_view const& values, data_type desired_type, + std::optional incoming_status, + std::optional status, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { @@ -937,57 +1222,109 @@ std::unique_ptr cast_variant(column_view const& values, } size_type const num_rows = values.size(); + + // Validate incoming_status before the empty-values fast path so callers always get + // std::invalid_argument for a malformed status column, even when values is empty. + if (incoming_status.has_value()) { + CUDF_EXPECTS(!incoming_status->nullable(), + "incoming status column must not be nullable; use row_null for SQL-null rows", + std::invalid_argument); + CUDF_EXPECTS(incoming_status->type().id() == type_id::UINT8, + "incoming status column must be UINT8", + std::invalid_argument); + CUDF_EXPECTS(incoming_status->size() == num_rows, + "incoming status column must have the same number of rows as the values column", + std::invalid_argument); + } + + if (status.has_value()) { + CUDF_EXPECTS(!status->nullable(), + "status column must not be nullable; use row_null for SQL-null rows", + std::invalid_argument); + CUDF_EXPECTS( + status->type().id() == type_id::UINT8, "status column must be UINT8", std::invalid_argument); + CUDF_EXPECTS(status->size() == num_rows, + "status column must have the same number of rows as the values column", + std::invalid_argument); + } + if (num_rows == 0) { return make_empty_column(desired_type); } auto val_device_view = column_device_view::create(values, stream); cudf::lists_column_device_view val_lists_device_view(*val_device_view); - // Initialize the null mask from the values column (or all-valid) auto null_mask = values.nullable() ? cudf::detail::copy_bitmask(values, stream, mr) : cudf::create_null_mask(num_rows, mask_state::ALL_VALID, stream, mr); auto* d_null_mask = static_cast(null_mask.data()); - return cudf::type_dispatcher(desired_type, - cast_variant_fn{val_lists_device_view, - num_rows, - desired_type, - d_null_mask, - std::move(null_mask), - stream, - mr}); + // Build device view for incoming status if provided; keep a placeholder when absent so that + // cast_variant_fn always holds a valid column_device_view (kernel ignores it when !has_incoming). + auto placeholder_col = make_empty_column(data_type{type_id::UINT8}); + auto incoming_dev_view = incoming_status.has_value() + ? column_device_view::create(*incoming_status, stream) + : column_device_view::create(*placeholder_col, stream); + bool const has_incoming = incoming_status.has_value(); + + return cudf::type_dispatcher( + desired_type, + cast_variant_fn{ + val_lists_device_view, + num_rows, + desired_type, + d_null_mask, + std::move(null_mask), + stream, + mr, + *incoming_dev_view, + has_incoming, + status.has_value() ? reinterpret_cast(status->data()) : nullptr}); } } // namespace detail std::unique_ptr get_variant_field(column_view const& variant_column, std::string_view path, + std::optional status, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); - return detail::get_variant_field(variant_column, path, stream, mr); + return detail::get_variant_field(variant_column, path, status, stream, mr); } std::unique_ptr cast_variant(column_view const& values, data_type desired_type, + std::optional incoming_status, + std::optional status, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); - return detail::cast_variant(values, desired_type, stream, mr); + return detail::cast_variant(values, desired_type, incoming_status, status, stream, mr); } std::unique_ptr extract_variant_field(column_view const& variant_column, std::string_view path, data_type desired_type, + std::optional status, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); - auto value = detail::get_variant_field( - variant_column, path, stream, cudf::get_current_device_resource_ref()); - return detail::cast_variant(value->view(), desired_type, stream, mr); + auto const temp_mr = cudf::get_current_device_resource_ref(); + + if (status.has_value()) { + auto extract_status = make_numeric_column( + data_type{type_id::UINT8}, variant_column.size(), mask_state::UNALLOCATED, stream, temp_mr); + auto value = detail::get_variant_field( + variant_column, path, extract_status->mutable_view(), stream, temp_mr); + return detail::cast_variant( + value->view(), desired_type, extract_status->view(), status, stream, mr); + } + + auto value = detail::get_variant_field(variant_column, path, std::nullopt, stream, temp_mr); + return detail::cast_variant(value->view(), desired_type, std::nullopt, std::nullopt, stream, mr); } } // namespace io::parquet::experimental diff --git a/cpp/tests/io/experimental/variant_extract_test.cpp b/cpp/tests/io/experimental/variant_extract_test.cpp index 03c334f250e..c9739f322f1 100644 --- a/cpp/tests/io/experimental/variant_extract_test.cpp +++ b/cpp/tests/io/experimental/variant_extract_test.cpp @@ -53,34 +53,35 @@ namespace { using cudf::io::parquet::experimental::variant_basic_type; using cudf::io::parquet::experimental::variant_primitive_type; -// Compose a value-metadata header byte from a basic type and its 6-bit value_header. constexpr uint8_t make_variant_header(variant_basic_type basic, uint8_t value_header) { CUDF_EXPECTS(value_header <= 0x3F, "VARIANT value_header must fit in 6 bits"); return static_cast(static_cast(basic) | (value_header << 2)); } -// Header byte for a primitive value of the given physical type. constexpr uint8_t make_variant_primitive(variant_primitive_type type) { return make_variant_header(variant_basic_type::PRIMITIVE, static_cast(type)); } -// Header byte for a short string of the given length (must fit in 6 bits: 0..63). +/** + * @brief Header byte for a short string of the given length (must fit in 6 bits: 0..63). + */ constexpr uint8_t make_variant_short_string_header(std::size_t length) { CUDF_EXPECTS(length <= 0x3F, "VARIANT short string length must fit in 6 bits"); return make_variant_header(variant_basic_type::SHORT_STRING, static_cast(length)); } -// Header byte for an object value with 1-byte field ids and 1-byte offsets -// (is_large=false), i.e. value_header == 0. +/** + * @brief Header byte for an object value with 1-byte field ids and 1-byte offsets + * (is_large=false), i.e. value_header == 0. + */ constexpr uint8_t make_variant_object_header() { return make_variant_header(variant_basic_type::OBJECT, 0); } -// Build a struct `column_view` over (metadata, value) without copying. inline cudf::column_view wrap_variant_view(cudf::column_view const& metadata, cudf::column_view const& value) { @@ -95,7 +96,6 @@ inline cudf::column_view wrap_variant_view(cudf::column_view const& metadata, {metadata, value}}; } -// Wrap a single-row (metadata, value) pair as a VARIANT struct column. inline cudf::test::structs_column_wrapper wrap_single_variant(std::vector const& meta, std::vector const& val) { @@ -104,7 +104,6 @@ inline cudf::test::structs_column_wrapper wrap_single_variant(std::vector cudf::test::structs_column_wrapper make_apache_variant(avf::fixture const& f) { @@ -113,10 +112,12 @@ cudf::test::structs_column_wrapper make_apache_variant(avf::fixture const& return cudf::test::structs_column_wrapper{{m, v}}; } -// Three-row VARIANT fixture reused by multiple multi-row tests below. -// Row 0: dict {x,y}, value { x: INT32(7), y: "hi" } -// Row 1: dict {x,z}, value { x: INT32(42), z: INT32(99) } -// Row 2: dict {y}, value { y: "zzz" } +/** + * @brief Three-row VARIANT fixture reused by multiple multi-row tests below. + * Row 0: dict {x,y}, value { x: INT32(7), y: "hi" } + * Row 1: dict {x,z}, value { x: INT32(42), z: INT32(99) } + * Row 2: dict {y}, value { y: "zzz" } + */ inline cudf::test::structs_column_wrapper make_xyz_three_row_variant() { std::vector const m1 = {0x01, 0x02, 0x00, 0x01, 0x02, 'x', 'y'}; @@ -153,8 +154,12 @@ TEST_F(ExtractVariantFieldTest, NullStructRow) // Use the validity vector to mask the second row null. cudf::test::structs_column_wrapper col{{meta, val}, std::vector{true, false}}; - auto got = cudf::io::parquet::experimental::extract_variant_field( - col, "x", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, + "x", + cudf::data_type{cudf::type_id::INT32}, + std::nullopt, + cudf::test::get_default_stream()); cudf::test::fixed_width_column_wrapper expected({7, 0}, {true, false}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); @@ -167,8 +172,12 @@ TEST_F(ExtractVariantFieldTest, NonObjectValueYieldsNull) std::vector const valb = {0x14, 0x07, 0x00, 0x00, 0x00}; auto col = wrap_single_variant(metab, valb); - auto got = cudf::io::parquet::experimental::extract_variant_field( - col, "x", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, + "x", + cudf::data_type{cudf::type_id::INT32}, + std::nullopt, + cudf::test::get_default_stream()); cudf::test::fixed_width_column_wrapper expected({0}, {false}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); @@ -181,8 +190,12 @@ TEST_F(ExtractVariantFieldTest, InvalidMetadataYieldsNull) std::vector const valb = {0x02, 0x01, 0x00, 0x00, 0x05, 0x14, 0x07, 0x00, 0x00, 0x00}; auto col = wrap_single_variant(metab, valb); - auto got = cudf::io::parquet::experimental::extract_variant_field( - col, "x", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, + "x", + cudf::data_type{cudf::type_id::INT32}, + std::nullopt, + cudf::test::get_default_stream()); cudf::test::fixed_width_column_wrapper expected({0}, {false}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); @@ -195,8 +208,12 @@ TEST_F(ExtractVariantFieldTest, UnsupportedMetadataVersionYieldsNull) std::vector const valb = {0x02, 0x01, 0x00, 0x00, 0x05, 0x14, 0x07, 0x00, 0x00, 0x00}; auto col = wrap_single_variant(metab, valb); - auto got = cudf::io::parquet::experimental::extract_variant_field( - col, "x", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, + "x", + cudf::data_type{cudf::type_id::INT32}, + std::nullopt, + cudf::test::get_default_stream()); cudf::test::fixed_width_column_wrapper expected({0}, {false}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); @@ -209,8 +226,12 @@ TEST_F(ExtractVariantFieldTest, TruncatedObjectValueYieldsNull) std::vector const valb = {0x02}; auto col = wrap_single_variant(metab, valb); - auto got = cudf::io::parquet::experimental::extract_variant_field( - col, "x", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, + "x", + cudf::data_type{cudf::type_id::INT32}, + std::nullopt, + cudf::test::get_default_stream()); cudf::test::fixed_width_column_wrapper expected({0}, {false}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); @@ -221,17 +242,17 @@ TEST_F(ExtractVariantFieldTest, MultiRow) auto col = make_xyz_three_row_variant(); auto stream = cudf::test::get_default_stream(); auto x = cudf::io::parquet::experimental::extract_variant_field( - col, "x", cudf::data_type{cudf::type_id::INT32}, stream); + col, "x", cudf::data_type{cudf::type_id::INT32}, std::nullopt, stream); cudf::test::fixed_width_column_wrapper x_exp({7, 42, 0}, {true, true, false}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*x, x_exp); auto y = cudf::io::parquet::experimental::extract_variant_field( - col, "y", cudf::data_type{cudf::type_id::STRING}, stream); + col, "y", cudf::data_type{cudf::type_id::STRING}, std::nullopt, stream); cudf::test::strings_column_wrapper y_exp({"hi", "", "zzz"}, {true, false, true}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*y, y_exp); auto z = cudf::io::parquet::experimental::extract_variant_field( - col, "z", cudf::data_type{cudf::type_id::INT32}, stream); + col, "z", cudf::data_type{cudf::type_id::INT32}, std::nullopt, stream); cudf::test::fixed_width_column_wrapper z_exp({0, 99, 0}, {false, true, false}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*z, z_exp); } @@ -242,8 +263,12 @@ TEST_F(ExtractVariantFieldTest, SlicedInput) auto const col = make_xyz_three_row_variant(); auto const sliced = cudf::slice(col, {1, 3}).front(); - auto got = cudf::io::parquet::experimental::extract_variant_field( - sliced, "x", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + auto got = + cudf::io::parquet::experimental::extract_variant_field(sliced, + "x", + cudf::data_type{cudf::type_id::INT32}, + std::nullopt, + cudf::test::get_default_stream()); cudf::test::fixed_width_column_wrapper expected({42, 0}, {true, false}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); @@ -259,7 +284,8 @@ TEST_F(ExtractVariantFieldTest, ApacheObjectPrimitiveStringFields) {std::pair{"string_field", "Apache Parquet"}, std::pair{"timestamp_field", "2025-04-16T12:34:56.78"}}) { SCOPED_TRACE(std::string{"field: "} + field); - auto got = cudf::io::parquet::experimental::extract_variant_field(col, field, s, stream); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, field, s, std::nullopt, stream); cudf::test::strings_column_wrapper expected({expected_str}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -273,7 +299,8 @@ TEST_F(ExtractVariantFieldTest, ApacheObjectPrimitiveNullCases) for (auto const& field : {"no_such_field", "null_field"}) { SCOPED_TRACE(std::string{"field: "} + field); - auto got = cudf::io::parquet::experimental::extract_variant_field(col, field, s, stream); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, field, s, std::nullopt, stream); ASSERT_EQ(got->size(), 1); EXPECT_EQ(got->null_count(), 1); } @@ -282,8 +309,12 @@ TEST_F(ExtractVariantFieldTest, ApacheObjectPrimitiveNullCases) TEST_F(ExtractVariantFieldTest, ApacheObjectPrimitiveIntField) { auto col = make_apache_variant(avf::object_primitive); - auto got = cudf::io::parquet::experimental::extract_variant_field( - col, "int_field", cudf::data_type{cudf::type_id::INT8}, cudf::test::get_default_stream()); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, + "int_field", + cudf::data_type{cudf::type_id::INT8}, + std::nullopt, + cudf::test::get_default_stream()); cudf::test::fixed_width_column_wrapper expected{int8_t{1}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -297,12 +328,12 @@ TEST_F(ExtractVariantFieldTest, ApacheObjectNested) SCOPED_TRACE(std::string{"path: "} + path); if constexpr (std::is_same_v) { auto got = cudf::io::parquet::experimental::extract_variant_field( - col, path, cudf::data_type{cudf::type_id::STRING}, stream); + col, path, cudf::data_type{cudf::type_id::STRING}, std::nullopt, stream); cudf::test::strings_column_wrapper expected({expected_val}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } else { auto got = cudf::io::parquet::experimental::extract_variant_field( - col, path, cudf::data_type{cudf::type_to_id()}, stream); + col, path, cudf::data_type{cudf::type_to_id()}, std::nullopt, stream); cudf::test::fixed_width_column_wrapper expected{expected_val}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -319,8 +350,12 @@ TEST_F(ExtractVariantFieldTest, ApacheObjectNested) TEST_F(ExtractVariantFieldTest, ApacheObjectEmpty) { auto col = make_apache_variant(avf::object_empty); - auto got = cudf::io::parquet::experimental::extract_variant_field( - col, "foo", cudf::data_type{cudf::type_id::STRING}, cudf::test::get_default_stream()); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, + "foo", + cudf::data_type{cudf::type_id::STRING}, + std::nullopt, + cudf::test::get_default_stream()); ASSERT_EQ(got->size(), 1); EXPECT_EQ(got->null_count(), 1); } @@ -331,14 +366,15 @@ TEST_F(ExtractVariantFieldTest, ApacheObjectNestedChainedCalls) auto stream = cudf::test::get_default_stream(); auto single = cudf::io::parquet::experimental::get_variant_field( - col, "$.observation.value.temperature", stream); + col, "$.observation.value.temperature", std::nullopt, stream); auto const meta_v = cudf::structs_column_view{col}.get_sliced_child(0, stream); - auto obs = cudf::io::parquet::experimental::get_variant_field(col, "observation", stream); + auto obs = + cudf::io::parquet::experimental::get_variant_field(col, "observation", std::nullopt, stream); auto vobj = cudf::io::parquet::experimental::get_variant_field( - wrap_variant_view(meta_v, obs->view()), "value", stream); + wrap_variant_view(meta_v, obs->view()), "value", std::nullopt, stream); auto chained = cudf::io::parquet::experimental::get_variant_field( - wrap_variant_view(meta_v, vobj->view()), "temperature", stream); + wrap_variant_view(meta_v, vobj->view()), "temperature", std::nullopt, stream); EXPECT_EQ(single->type().id(), cudf::type_id::LIST); EXPECT_EQ(chained->type().id(), cudf::type_id::LIST); @@ -351,7 +387,7 @@ TEST_F(ExtractVariantFieldTest, ApacheObjectNestedMissingIntermediate) auto stream = cudf::test::get_default_stream(); auto got = cudf::io::parquet::experimental::extract_variant_field( - col, "$.species.nope", cudf::data_type{cudf::type_id::STRING}, stream); + col, "$.species.nope", cudf::data_type{cudf::type_id::STRING}, std::nullopt, stream); cudf::test::strings_column_wrapper expected({"donotread"}, {false}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); @@ -366,8 +402,12 @@ TEST_F(ExtractVariantFieldTest, NestedPathNonObjectIntermediate) auto col = wrap_single_variant(metab, valb); // Descending into "a" fails because it is a primitive, not an object. - auto got = cudf::io::parquet::experimental::extract_variant_field( - col, "$.a.b", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, + "$.a.b", + cudf::data_type{cudf::type_id::INT32}, + std::nullopt, + cudf::test::get_default_stream()); cudf::test::fixed_width_column_wrapper expected({0}, {false}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); @@ -378,15 +418,18 @@ TEST_F(ExtractVariantFieldTest, BareNameEqualsDollarPath) auto col = make_xyz_three_row_variant(); auto stream = cudf::test::get_default_stream(); - auto bare = cudf::io::parquet::experimental::get_variant_field(col, "x", stream); - auto dollar = cudf::io::parquet::experimental::get_variant_field(col, "$.x", stream); + auto bare = cudf::io::parquet::experimental::get_variant_field(col, "x", std::nullopt, stream); + auto dollar = + cudf::io::parquet::experimental::get_variant_field(col, "$.x", std::nullopt, stream); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*bare, *dollar); } namespace { -// INT32 primitive blob: primitive int32 header + little-endian 4-byte payload. +/** + * @brief INT32 primitive blob: primitive int32 header + little-endian 4-byte payload. + */ inline std::vector enc_int32(int32_t v) { auto const u = static_cast(v); @@ -397,7 +440,6 @@ inline std::vector enc_int32(int32_t v) static_cast((u >> 24) & 0xff)}; } -// Short-string primitive blob (single-byte header). inline std::vector enc_short_string(std::string_view s) { CUDF_EXPECTS(s.size() < 64, "short-string length must fit in 6 bits of the single-byte header"); @@ -406,7 +448,9 @@ inline std::vector enc_short_string(std::string_view s) return out; } -// Append `width` little-endian bytes of `bits` to `out`. +/** + * @brief Append `width` little-endian bytes of `bits` to `out`. + */ inline void append_le(std::vector& out, uint64_t bits, int width) { for (int i = 0; i < width; ++i) { @@ -414,7 +458,10 @@ inline void append_le(std::vector& out, uint64_t bits, int width) } } -// Primitive value blobs (header + fixed payload) for every physical type the cast matrix exercises. +/** + * @brief Primitive value blobs (header + fixed payload) for every physical type the cast matrix + * exercises. + */ inline std::vector enc_null() { return {make_variant_primitive(variant_primitive_type::NULLVAL)}; @@ -452,7 +499,9 @@ inline std::vector enc_float64(double v) return out; } -// Long-string primitive blob: header + 4-byte LE length + payload. +/** + * @brief Long-string primitive blob: header + 4-byte LE length + payload. + */ inline std::vector enc_long_string(std::string_view s) { std::vector out{make_variant_primitive(variant_primitive_type::LONG_STRING)}; @@ -463,8 +512,10 @@ inline std::vector enc_long_string(std::string_view s) return out; } -// Build a single-field object value wrapping `inner` under field id `fid`. -// field_off_size=1, field_id_size=1, is_large=false. +/** + * @brief Build a single-field object value wrapping `inner` under field id `fid`. + * field_off_size=1, field_id_size=1, is_large=false. + */ inline std::vector build_single_field_object(uint8_t fid, std::vector const& inner) { @@ -476,10 +527,12 @@ inline std::vector build_single_field_object(uint8_t fid, return out; } -// Build a VARIANT object blob with `n_fields` fields. Field ids are 0..n_fields-1 -// (in ascending order, matching the dictionary positions) and each field holds a bare INT32 equal -// to its field id. Uses 1-byte field_id_size and 1-byte field_off_size; n_fields must be -// <= 51 so the total value bytes (5 * n_fields) still fit in 1-byte offsets. +/** + * @brief Build a VARIANT object blob with `n_fields` fields. Field ids are 0..n_fields-1 + * (in ascending order, matching the dictionary positions) and each field holds a bare INT32 equal + * to its field id. Uses 1-byte field_id_size and 1-byte field_off_size; n_fields must be + * <= 51 so the total value bytes (5 * n_fields) still fit in 1-byte offsets. + */ inline std::vector build_sequential_int32_object(int n_fields) { std::vector out{make_variant_object_header(), static_cast(n_fields)}; @@ -496,7 +549,9 @@ inline std::vector build_sequential_int32_object(int n_fields) return out; } -// Lexicographically ordered dictionary of N zero-padded two-digit keys "k". +/** + * @brief Lexicographically ordered dictionary of N zero-padded two-digit keys "k". + */ inline std::vector make_numeric_keys(int n) { std::vector out; @@ -509,9 +564,11 @@ inline std::vector make_numeric_keys(int n) return out; } -// Wrap per-row (metadata, value) byte vectors into a VARIANT struct column. Built -// with make_lists_column + structs_column_wrapper directly so the helper stays -// self-contained within this test file for dynamic row counts. +/** + * @brief Wrap per-row (metadata, value) byte vectors into a VARIANT struct column. Built + * with make_lists_column + structs_column_wrapper directly so the helper stays + * self-contained within this test file for dynamic row counts. + */ inline cudf::test::structs_column_wrapper wrap_multi_row_variant( std::vector> const& meta_rows, std::vector> const& val_rows) @@ -535,7 +592,9 @@ inline cudf::test::structs_column_wrapper wrap_multi_row_variant( return cudf::test::structs_column_wrapper{std::move(children)}; } -// Build a metadata blob (version 1, offset_size=1) for the given ordered string dictionary. +/** + * @brief Build a metadata blob (version 1, offset_size=1) for the given ordered string dictionary. + */ inline std::vector build_metadata(std::vector const& keys) { std::vector out{0x01, static_cast(keys.size())}; @@ -576,8 +635,12 @@ TEST_F(ExtractVariantFieldTest, NestedPathMultiRowMixedNulls) {v0.begin(), v0.end()}, {v1.begin(), v1.end()}, {v2.begin(), v2.end()}}; cudf::test::structs_column_wrapper col{{meta, val}}; - auto got = cudf::io::parquet::experimental::extract_variant_field( - col, "$.1st.foo-bar", cudf::data_type{cudf::type_id::INT32}, cudf::test::get_default_stream()); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, + "$.1st.foo-bar", + cudf::data_type{cudf::type_id::INT32}, + std::nullopt, + cudf::test::get_default_stream()); cudf::test::fixed_width_column_wrapper expected({1, 0, 0}, {true, false, false}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); @@ -587,14 +650,14 @@ TEST_F(ExtractVariantFieldTest, EmptyPathRejected) { auto col = wrap_single_variant(build_metadata({}), enc_int32(1)); auto stream = cudf::test::get_default_stream(); - EXPECT_THROW( - static_cast(cudf::io::parquet::experimental::get_variant_field(col, "", stream)), - std::invalid_argument); - EXPECT_THROW( - static_cast(cudf::io::parquet::experimental::get_variant_field(col, "$", stream)), - std::invalid_argument); + EXPECT_THROW(static_cast( + cudf::io::parquet::experimental::get_variant_field(col, "", std::nullopt, stream)), + std::invalid_argument); + EXPECT_THROW(static_cast(cudf::io::parquet::experimental::get_variant_field( + col, "$", std::nullopt, stream)), + std::invalid_argument); EXPECT_THROW(static_cast(cudf::io::parquet::experimental::extract_variant_field( - col, "", cudf::data_type{cudf::type_id::INT32}, stream)), + col, "", cudf::data_type{cudf::type_id::INT32}, std::nullopt, stream)), std::invalid_argument); } @@ -616,9 +679,9 @@ TEST_F(ExtractVariantFieldTest, SyntaxErrors) "$.a[01x]", "$.a[1", "$.a[99999999999999999999]"}) { - EXPECT_THROW( - static_cast(cudf::io::parquet::experimental::get_variant_field(col, bad, stream)), - std::invalid_argument) + EXPECT_THROW(static_cast(cudf::io::parquet::experimental::get_variant_field( + col, bad, std::nullopt, stream)), + std::invalid_argument) << "path that should have thrown: " << bad; } } @@ -630,7 +693,8 @@ TEST_F(ExtractVariantFieldTest, ApacheArrayPrimitiveIndexing) auto stream = cudf::test::get_default_stream(); auto const i8 = cudf::data_type{cudf::type_id::INT8}; auto const get = [&](char const* path) { - return cudf::io::parquet::experimental::extract_variant_field(col, path, i8, stream); + return cudf::io::parquet::experimental::extract_variant_field( + col, path, i8, std::nullopt, stream); }; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*get("$[0]"), @@ -660,7 +724,8 @@ TEST_F(ExtractVariantFieldTest, ApacheArrayPrimitiveIndexing) value.insert(value.end(), {0x0c, 42}); // INT8(42) auto wide_col = wrap_single_variant(build_metadata({}), value); - auto got = cudf::io::parquet::experimental::extract_variant_field(wide_col, "$[0]", i8, stream); + auto got = cudf::io::parquet::experimental::extract_variant_field( + wide_col, "$[0]", i8, std::nullopt, stream); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, cudf::test::fixed_width_column_wrapper{int8_t{42}}); } @@ -677,12 +742,12 @@ TEST_F(ExtractVariantFieldTest, ArrayIndexingTypeMismatchAndBounds) // Object-key descent into an array value: no such key -> null. auto key_on_array = - cudf::io::parquet::experimental::extract_variant_field(col, "$.foo", i8, stream); + cudf::io::parquet::experimental::extract_variant_field(col, "$.foo", i8, std::nullopt, stream); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*key_on_array, null_expected); // Index step against a primitive element (after first descending into it): non-array -> null. - auto index_on_primitive = - cudf::io::parquet::experimental::extract_variant_field(col, "$[0][0]", i8, stream); + auto index_on_primitive = cudf::io::parquet::experimental::extract_variant_field( + col, "$[0][0]", i8, std::nullopt, stream); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*index_on_primitive, null_expected); } @@ -695,18 +760,31 @@ TEST_F(ExtractVariantFieldTest, EmptyArrayIndexing) for (auto const* path : {"$[0]", "$[1]"}) { SCOPED_TRACE(std::string{"path: "} + path); - auto got = cudf::io::parquet::experimental::extract_variant_field(col, path, i8, stream); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, path, i8, std::nullopt, stream); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, null_expected); } - // Truncated counts/tables, decreasing offsets, and offsets beyond the values region yield null. - for (auto const& value : std::vector>{{0x13}, - {0x03, 0x01, 0x00}, - {0x03, 0x01, 0x02, 0x01, 0x0c, 42}, - {0x03, 0x01, 0x00, 0x03, 0x0c, 42}}) { + // Truncated counts/tables, decreasing offsets, offsets beyond the values region, and an + // element end that escapes the terminal offset all yield null. + for (auto const& value : std::vector>{ + {0x13}, + {0x03, 0x01, 0x00}, + {0x03, 0x01, 0x02, 0x01, 0x0c, 42}, + {0x03, 0x01, 0x00, 0x03, 0x0c, 42}, + // 2-element array: offsets[0]=0, offsets[1]=5, terminal offsets[2]=1. + // 5 physical value bytes are present (passes the old physical-extent check), + // but the terminal offset declares only 1 value byte, so element 0's end (5) + // escapes the declared boundary → malformed. + {0x03, 0x02, 0x00, 0x05, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00}, + // 1-element array where offsets[0]=1 (must be 0 per spec). + // offsets: [1, 2], values: [0x0c, 0x2a] (an int8 variant for 42). + // terminal_off=2 <= values_extent=2 so the terminal check passes; the + // nonzero first offset is caught by the new offsets[0]==0 guard. + {0x03, 0x01, 0x01, 0x02, 0x0c, 0x2a}}) { auto malformed_col = wrap_single_variant(build_metadata({}), value); - auto got = - cudf::io::parquet::experimental::extract_variant_field(malformed_col, "$[0]", i8, stream); + auto got = cudf::io::parquet::experimental::extract_variant_field( + malformed_col, "$[0]", i8, std::nullopt, stream); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, null_expected); } } @@ -723,14 +801,14 @@ TEST_F(ExtractVariantFieldTest, MixedObjectArrayTraversal) auto const check_str = [&](char const* path, char const* expected) { SCOPED_TRACE(std::string{"path: "} + path); auto got = cudf::io::parquet::experimental::extract_variant_field( - col, path, cudf::data_type{cudf::type_id::STRING}, stream); + col, path, cudf::data_type{cudf::type_id::STRING}, std::nullopt, stream); cudf::test::strings_column_wrapper const expected_col({expected}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected_col); }; auto const check_null = [&](char const* path) { SCOPED_TRACE(std::string{"path: "} + path); auto got = cudf::io::parquet::experimental::extract_variant_field( - col, path, cudf::data_type{cudf::type_id::STRING}, stream); + col, path, cudf::data_type{cudf::type_id::STRING}, std::nullopt, stream); cudf::test::strings_column_wrapper const null_col({""}, {false}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, null_col); }; @@ -757,12 +835,12 @@ TEST_F(ExtractVariantFieldTest, LargeDictionaryAndObjectScan) auto const int32_dtype = cudf::data_type{cudf::type_id::INT32}; // First, middle, and last keys each decode to their own field id. - auto first = - cudf::io::parquet::experimental::extract_variant_field(col, "k00", int32_dtype, stream); - auto mid = - cudf::io::parquet::experimental::extract_variant_field(col, "k24", int32_dtype, stream); - auto last = - cudf::io::parquet::experimental::extract_variant_field(col, "k49", int32_dtype, stream); + auto first = cudf::io::parquet::experimental::extract_variant_field( + col, "k00", int32_dtype, std::nullopt, stream); + auto mid = cudf::io::parquet::experimental::extract_variant_field( + col, "k24", int32_dtype, std::nullopt, stream); + auto last = cudf::io::parquet::experimental::extract_variant_field( + col, "k49", int32_dtype, std::nullopt, stream); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*first, cudf::test::fixed_width_column_wrapper{0}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*mid, cudf::test::fixed_width_column_wrapper{24}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*last, cudf::test::fixed_width_column_wrapper{49}); @@ -791,13 +869,24 @@ TEST_F(ExtractVariantFieldTest, MalformedVariantDataYieldsNull) {"object declares more fields than the value buffer holds", build_metadata({"x"}), {make_variant_object_header(), 0xFF}}, + // Two-key dict: offsets[0]=0, offsets[1]=1 ("x" is key 0), terminal offsets[2]=0. + // 2 physical string bytes are present so the old per-entry check passes (1 <= 2), + // but the terminal offset declares the string region as 0 bytes. The key "x" + // matches at i=0 before the terminal is consulted → must be malformed. + {"metadata terminal offset below first key's declared end", + {0x01, 0x02, 0x00, 0x01, 0x00, 'x', 'y'}, + build_single_field_object(0, enc_int32(42))}, + // Single-key dict where offsets[0] != 0. The Parquet VARIANT spec requires offsets[0] == 0; + // a non-zero first offset makes the string region ill-defined. + // Layout: num_entries=1, offsets[0]=1 (invalid), offsets[1]=2, string bytes "x". + {"metadata first offset non-zero", {0x01, 0x01, 0x01, 0x02, 'x'}, valid_object}, }; for (auto const& c : cases) { SCOPED_TRACE(c.label); auto col = wrap_single_variant(c.meta, c.val); - auto got = - cudf::io::parquet::experimental::extract_variant_field(col, "x", int32_dtype, stream); + auto got = cudf::io::parquet::experimental::extract_variant_field( + col, "x", int32_dtype, std::nullopt, stream); ASSERT_EQ(got->size(), 1); EXPECT_EQ(got->null_count(), 1); } @@ -846,8 +935,12 @@ TEST_F(ExtractVariantFieldTest, NullsAtDifferentDepths) auto col = wrap_multi_row_variant(meta_rows, val_rows); - auto got = cudf::io::parquet::experimental::extract_variant_field( - col, "$.a.b.c.d", cudf::data_type{cudf::type_id::STRING}, cudf::test::get_default_stream()); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, + "$.a.b.c.d", + cudf::data_type{cudf::type_id::STRING}, + std::nullopt, + cudf::test::get_default_stream()); cudf::test::strings_column_wrapper expected(exp_strs.begin(), exp_strs.end(), exp_valid.begin()); CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*got, expected); @@ -859,7 +952,7 @@ TEST_F(ExtractVariantFieldTest, EmptyInput) auto const variant = cudf::empty_like(make_xyz_three_row_variant()); auto got = cudf::io::parquet::experimental::extract_variant_field( - *variant, "x", cudf::data_type{cudf::type_id::INT32}, stream); + *variant, "x", cudf::data_type{cudf::type_id::INT32}, std::nullopt, stream); EXPECT_EQ(got->type().id(), cudf::type_id::INT32); EXPECT_EQ(got->size(), 0); EXPECT_EQ(got->null_count(), 0); @@ -872,14 +965,15 @@ TEST_F(GetVariantFieldTest, ApacheObjectPrimitive) auto col = make_apache_variant(avf::object_primitive); auto stream = cudf::test::get_default_stream(); - auto got = cudf::io::parquet::experimental::get_variant_field(col, "int_field", stream); + auto got = + cudf::io::parquet::experimental::get_variant_field(col, "int_field", std::nullopt, stream); EXPECT_EQ(got->type().id(), cudf::type_id::LIST); EXPECT_EQ(got->size(), 1); EXPECT_EQ(cudf::lists_column_view{got->view()}.child().type().id(), cudf::type_id::UINT8); auto casted = cudf::io::parquet::experimental::cast_variant( - got->view(), cudf::data_type{cudf::type_id::INT8}, stream); + got->view(), cudf::data_type{cudf::type_id::INT8}, {}, std::nullopt, stream); cudf::test::fixed_width_column_wrapper expected{int8_t{1}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*casted, expected); } @@ -888,7 +982,7 @@ TEST_F(GetVariantFieldTest, ApacheObjectPrimitiveMissingKeyAllNull) { auto col = make_apache_variant(avf::object_primitive); auto got = cudf::io::parquet::experimental::get_variant_field( - col, "no_such_field", cudf::test::get_default_stream()); + col, "no_such_field", std::nullopt, cudf::test::get_default_stream()); EXPECT_EQ(got->type().id(), cudf::type_id::LIST); EXPECT_EQ(got->size(), 1); @@ -901,11 +995,12 @@ TEST_F(GetVariantFieldTest, GetAndCastMatchesExtract) auto stream = cudf::test::get_default_stream(); auto extract_x = cudf::io::parquet::experimental::extract_variant_field( - col, "x", cudf::data_type{cudf::type_id::INT32}, stream); + col, "x", cudf::data_type{cudf::type_id::INT32}, std::nullopt, stream); - auto intermediate = cudf::io::parquet::experimental::get_variant_field(col, "x", stream); - auto two_step_x = cudf::io::parquet::experimental::cast_variant( - intermediate->view(), cudf::data_type{cudf::type_id::INT32}, stream); + auto intermediate = + cudf::io::parquet::experimental::get_variant_field(col, "x", std::nullopt, stream); + auto two_step_x = cudf::io::parquet::experimental::cast_variant( + intermediate->view(), cudf::data_type{cudf::type_id::INT32}, {}, std::nullopt, stream); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*extract_x, *two_step_x); } @@ -915,7 +1010,8 @@ TEST_F(GetVariantFieldTest, EmptyInput) auto const stream = cudf::test::get_default_stream(); auto const variant = cudf::empty_like(make_xyz_three_row_variant()); - auto got = cudf::io::parquet::experimental::get_variant_field(*variant, "x", stream); + auto got = + cudf::io::parquet::experimental::get_variant_field(*variant, "x", std::nullopt, stream); EXPECT_EQ(got->type().id(), cudf::type_id::LIST); EXPECT_EQ(got->size(), 0); EXPECT_EQ(got->null_count(), 0); @@ -929,7 +1025,7 @@ std::unique_ptr cast_apache_primitive(avf::fixture const& fi auto col = make_apache_variant(fixture); auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); return cudf::io::parquet::experimental::cast_variant( - value, cudf::data_type{cudf::type_to_id()}, stream); + value, cudf::data_type{cudf::type_to_id()}, {}, std::nullopt, stream); } struct CastVariantTest : public cudf::test::BaseFixture {}; @@ -966,7 +1062,7 @@ TEST_F(CastVariantTest, ApachePrimitiveFloats) auto col = make_apache_variant(fixture); auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); auto got = cudf::io::parquet::experimental::cast_variant( - value, cudf::data_type{cudf::type_to_id()}, stream); + value, cudf::data_type{cudf::type_to_id()}, {}, std::nullopt, stream); cudf::test::fixed_width_column_wrapper expected{expected_val}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); }; @@ -982,7 +1078,7 @@ TEST_F(CastVariantTest, ApachePrimitiveBooleans) auto col = make_apache_variant(fixture); auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); auto got = cudf::io::parquet::experimental::cast_variant( - value, cudf::data_type{cudf::type_id::BOOL8}, stream); + value, cudf::data_type{cudf::type_id::BOOL8}, {}, std::nullopt, stream); cudf::test::fixed_width_column_wrapper expected{expected_val}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); }; @@ -995,7 +1091,7 @@ TEST_F(CastVariantTest, ApachePrimitiveBooleans) auto col = make_apache_variant(avf::primitive_null); auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); auto got = cudf::io::parquet::experimental::cast_variant( - value, cudf::data_type{cudf::type_id::BOOL8}, stream); + value, cudf::data_type{cudf::type_id::BOOL8}, {}, std::nullopt, stream); cudf::test::fixed_width_column_wrapper expected({false}, {false}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1043,7 +1139,7 @@ TEST_F(CastVariantTest, ApachePrimitiveBooleans) auto const sliced = cudf::slice(col, {slice_beg, slice_end}).front(); auto const value = cudf::structs_column_view{sliced}.get_sliced_child(1, stream); auto got = cudf::io::parquet::experimental::cast_variant( - value, cudf::data_type{cudf::type_id::BOOL8}, stream); + value, cudf::data_type{cudf::type_id::BOOL8}, {}, std::nullopt, stream); cudf::test::fixed_width_column_wrapper expected( exp_vals.begin() + slice_beg, exp_vals.begin() + slice_end, exp_valid.begin() + slice_beg); @@ -1058,7 +1154,7 @@ TEST_F(CastVariantTest, ApacheShortString) auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); auto got = cudf::io::parquet::experimental::cast_variant( - value, cudf::data_type{cudf::type_id::STRING}, stream); + value, cudf::data_type{cudf::type_id::STRING}, {}, std::nullopt, stream); // Decoded from short_string.value: skip the 1-byte header, take the rest. std::string const expected_str(reinterpret_cast(avf::short_string.value.data() + 1), @@ -1074,7 +1170,7 @@ TEST_F(CastVariantTest, ApachePrimitiveString) auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); auto got = cudf::io::parquet::experimental::cast_variant( - value, cudf::data_type{cudf::type_id::STRING}, stream); + value, cudf::data_type{cudf::type_id::STRING}, {}, std::nullopt, stream); // Long-string layout: 1 header byte + 4-byte LE length + payload. std::string const expected_str( @@ -1091,7 +1187,7 @@ TEST_F(CastVariantTest, MismatchedTypeYieldsNull) auto col = make_apache_variant(avf::object_primitive); auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); auto got = cudf::io::parquet::experimental::cast_variant( - value, cudf::data_type{cudf::type_id::INT32}, stream); + value, cudf::data_type{cudf::type_id::INT32}, {}, std::nullopt, stream); ASSERT_EQ(got->size(), 1); EXPECT_EQ(got->null_count(), 1); } @@ -1107,7 +1203,8 @@ TEST_F(CastVariantTest, EmptyInput) cudf::type_id::FLOAT32, cudf::type_id::FLOAT64, cudf::type_id::BOOL8}) { - auto got = cudf::io::parquet::experimental::cast_variant(*values, cudf::data_type{id}, stream); + auto got = cudf::io::parquet::experimental::cast_variant( + *values, cudf::data_type{id}, {}, std::nullopt, stream); EXPECT_EQ(got->type().id(), id); EXPECT_EQ(got->size(), 0); EXPECT_EQ(got->null_count(), 0); @@ -1136,7 +1233,7 @@ TEST_F(CastVariantTest, UnsupportedTypeThrows) cudf::empty_like(cudf::structs_column_view{make_xyz_three_row_variant()}.child(1)); for (auto const id : ids) { EXPECT_THROW(static_cast(cudf::io::parquet::experimental::cast_variant( - *empty_values, cudf::data_type{id}, stream)), + *empty_values, cudf::data_type{id}, {}, std::nullopt, stream)), std::invalid_argument) << std::format("expected throw for type_id {} on empty input", static_cast(id)); } @@ -1146,7 +1243,7 @@ TEST_F(CastVariantTest, UnsupportedTypeThrows) auto const value = cudf::structs_column_view{col}.get_sliced_child(1, stream); for (auto const id : ids) { EXPECT_THROW(static_cast(cudf::io::parquet::experimental::cast_variant( - value, cudf::data_type{id}, stream)), + value, cudf::data_type{id}, {}, std::nullopt, stream)), std::invalid_argument) << std::format("expected throw for type_id {} on non-empty input", static_cast(id)); } @@ -1189,7 +1286,8 @@ TEST_F(CastVariantTest, CastSourceTargetMatrix) for (auto const& src : sources) { SCOPED_TRACE(std::string{"int target "} + match_label + ", source " + src.label); auto values = values_of(src.bytes); - auto got = cudf::io::parquet::experimental::cast_variant(values, target, stream); + auto got = + cudf::io::parquet::experimental::cast_variant(values, target, {}, std::nullopt, stream); if (std::string_view{src.label} == match_label) { cudf::test::fixed_width_column_wrapper const expected{match_value}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); @@ -1209,7 +1307,8 @@ TEST_F(CastVariantTest, CastSourceTargetMatrix) for (auto const& src : sources) { SCOPED_TRACE(std::string{"string target, source "} + src.label); auto values = values_of(src.bytes); - auto got = cudf::io::parquet::experimental::cast_variant(values, string_type, stream); + auto got = + cudf::io::parquet::experimental::cast_variant(values, string_type, {}, std::nullopt, stream); std::string_view const label{src.label}; if (label == "short_string" || label == "long_string") { std::string const expected_str = (label == "short_string") ? "hi" : std::string(70, 'a'); @@ -1229,7 +1328,7 @@ TEST_F(CastVariantTest, ShortStringLengthZero) std::vector const val{make_variant_short_string_header(0)}; cudf::test::lists_column_wrapper values(val.begin(), val.end()); auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::STRING}, stream); + values, cudf::data_type{cudf::type_id::STRING}, {}, std::nullopt, stream); cudf::test::strings_column_wrapper expected({""}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1244,7 +1343,7 @@ TEST_F(CastVariantTest, ShortStringMaxLength) val.insert(val.end(), payload.begin(), payload.end()); cudf::test::lists_column_wrapper values(val.begin(), val.end()); auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::STRING}, stream); + values, cudf::data_type{cudf::type_id::STRING}, {}, std::nullopt, stream); cudf::test::strings_column_wrapper expected({payload}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1257,7 +1356,7 @@ TEST_F(CastVariantTest, LongStringLengthZero) make_variant_primitive(variant_primitive_type::LONG_STRING), 0x00, 0x00, 0x00, 0x00}; cudf::test::lists_column_wrapper values(val.begin(), val.end()); auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::STRING}, stream); + values, cudf::data_type{cudf::type_id::STRING}, {}, std::nullopt, stream); cudf::test::strings_column_wrapper expected({""}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1277,7 +1376,7 @@ TEST_F(CastVariantTest, LongStringDeclaredLengthExceedsPayloadYieldsNull) SCOPED_TRACE(std::string{"payload bytes present: "} + std::to_string(val.size() - 5)); cudf::test::lists_column_wrapper values(val.begin(), val.end()); auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::STRING}, stream); + values, cudf::data_type{cudf::type_id::STRING}, {}, std::nullopt, stream); ASSERT_EQ(got->size(), 1); EXPECT_EQ(got->null_count(), 1); } @@ -1294,7 +1393,7 @@ TEST_F(CastVariantTest, LongStringPayloadExceedsDeclaredLength) hdr, 0x03, 0x00, 0x00, 0x00, 'a', 'b', 'c', 'x', 'x', 'x', 'x', 'x'}; cudf::test::lists_column_wrapper values(val.begin(), val.end()); auto got = cudf::io::parquet::experimental::cast_variant( - values, cudf::data_type{cudf::type_id::STRING}, stream); + values, cudf::data_type{cudf::type_id::STRING}, {}, std::nullopt, stream); cudf::test::strings_column_wrapper expected({"abc"}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); } @@ -1303,25 +1402,27 @@ struct InvalidInputShapeTest : public cudf::test::BaseFixture {}; namespace { -// A well-formed VARIANT child: a single-row list holding `bytes`. inline std::unique_ptr list_u8(std::vector const& bytes) { return cudf::test::lists_column_wrapper(bytes.begin(), bytes.end()).release(); } -// A single-row list (wrong element type for a VARIANT child). +/** + * @brief A single-row list (wrong element type for a VARIANT child). + */ inline std::unique_ptr list_i32(std::vector const& values) { return cudf::test::lists_column_wrapper(values.begin(), values.end()).release(); } -// A single-row fixed-width int32 column (a non-list child). +/** + * @brief A single-row fixed-width int32 column (a non-list child). + */ inline std::unique_ptr scalar_i32() { return cudf::test::fixed_width_column_wrapper{42}.release(); } -// A single-row STRUCT column adopting `children`. inline std::unique_ptr struct_of(std::vector> children) { return cudf::make_structs_column(1, std::move(children), 0, rmm::device_buffer{}); @@ -1336,7 +1437,6 @@ inline std::vector> two_children(std::unique_ptr column; @@ -1344,11 +1444,11 @@ struct broken_shape { } // namespace -// A VARIANT column must be a STRUCT whose first two children are each a list. Enumerate the -// distinct ways that column-shape contract can be broken; get_variant_field must reject every one -// with std::invalid_argument. TEST_F(InvalidInputShapeTest, GetVariantFieldRejectsMalformedInput) { + // A VARIANT column must be a STRUCT whose first two children are each a list. Enumerate + // the distinct ways that column-shape contract can be broken; get_variant_field must reject every + // one with std::invalid_argument. auto stream = cudf::test::get_default_stream(); std::vector cases; @@ -1370,15 +1470,15 @@ TEST_F(InvalidInputShapeTest, GetVariantFieldRejectsMalformedInput) for (auto const& c : cases) { SCOPED_TRACE(c.label); EXPECT_THROW(static_cast(cudf::io::parquet::experimental::get_variant_field( - c.column->view(), "x", stream)), + c.column->view(), "x", std::nullopt, stream)), std::invalid_argument); } } -// cast_variant requires a list input; every other shape must be rejected with -// std::invalid_argument. TEST_F(InvalidInputShapeTest, CastVariantRejectsMalformedInput) { + // cast_variant requires a list input; every other shape must be rejected with + // std::invalid_argument. auto stream = cudf::test::get_default_stream(); std::vector cases; @@ -1387,8 +1487,617 @@ TEST_F(InvalidInputShapeTest, CastVariantRejectsMalformedInput) for (auto const& c : cases) { SCOPED_TRACE(c.label); - EXPECT_THROW(static_cast(cudf::io::parquet::experimental::cast_variant( - c.column->view(), cudf::data_type{cudf::type_id::INT32}, stream)), - std::invalid_argument); + EXPECT_THROW( + static_cast(cudf::io::parquet::experimental::cast_variant( + c.column->view(), cudf::data_type{cudf::type_id::INT32}, {}, std::nullopt, stream)), + std::invalid_argument); + } +} + +TEST_F(InvalidInputShapeTest, CastVariantRejectsNullableIncomingStatus) +{ + // cast_variant must reject a nullable incoming_status column (SQL-null rows must be represented + // by the row_null enum value, not by null bits). + auto stream = cudf::test::get_default_stream(); + // One-row valid values column. + auto values = + list_u8({make_variant_primitive(variant_primitive_type::INT32), 0x01, 0x00, 0x00, 0x00}); + // Incoming status with a null entry (row 0 is null) — must be rejected. + // Use uint8_t{0} (== op_status::SUCCESS) directly; ST_SUCCESS is not in scope here. + std::vector const sv{uint8_t{0}}; + std::vector const sv_valid{false}; + cudf::test::fixed_width_column_wrapper nullable_status( + sv.begin(), sv.end(), sv_valid.begin()); + auto const status_col = nullable_status.release(); + auto const status_view = status_col->view(); + EXPECT_THROW( + static_cast(cudf::io::parquet::experimental::cast_variant( + values->view(), cudf::data_type{cudf::type_id::INT32}, status_view, std::nullopt, stream)), + std::invalid_argument); +} + +TEST_F(InvalidInputShapeTest, CastVariantRejectsInvalidIncomingStatusOnEmptyValues) +{ + // Regression: incoming_status validation must fire even when values is empty (zero rows). + // Prior to the fix, the empty-values fast path returned before the validation block, so a + // nullable, non-UINT8, or row-count-mismatched status column was silently accepted. + auto stream = cudf::test::get_default_stream(); + // Build a zero-row list values column. + auto const empty_values = + cudf::empty_like(cudf::structs_column_view{make_xyz_three_row_variant()}.child(1)); + + // Case 1: nullable incoming_status (one row, but values has zero rows — catch nullable first). + { + std::vector const sv{uint8_t{0}}; + std::vector const sv_valid{false}; + cudf::test::fixed_width_column_wrapper nullable_status( + sv.begin(), sv.end(), sv_valid.begin()); + auto const status_col = nullable_status.release(); + auto const status_view = status_col->view(); + EXPECT_THROW( + static_cast(cudf::io::parquet::experimental::cast_variant( + *empty_values, cudf::data_type{cudf::type_id::INT32}, status_view, std::nullopt, stream)), + std::invalid_argument) + << "nullable incoming_status must be rejected even when values is empty"; + } + + // Case 2: non-UINT8 incoming_status (zero-row INT32 column, non-nullable). + { + cudf::test::fixed_width_column_wrapper wrong_type_status{}; + auto const status_col = wrong_type_status.release(); + auto const status_view = status_col->view(); + EXPECT_THROW( + static_cast(cudf::io::parquet::experimental::cast_variant( + *empty_values, cudf::data_type{cudf::type_id::INT32}, status_view, std::nullopt, stream)), + std::invalid_argument) + << "non-UINT8 incoming_status must be rejected even when values is empty"; } + + // Case 3: row-count mismatch (one-row status vs zero-row values). + { + cudf::test::fixed_width_column_wrapper mismatched_status({uint8_t{0}}); + auto const status_col = mismatched_status.release(); + auto const status_view = status_col->view(); + EXPECT_THROW( + static_cast(cudf::io::parquet::experimental::cast_variant( + *empty_values, cudf::data_type{cudf::type_id::INT32}, status_view, std::nullopt, stream)), + std::invalid_argument) + << "row-count-mismatched incoming_status must be rejected even when values is empty"; + } +} + +// --------------------------------------------------------------------------- +// Status column tests +// --------------------------------------------------------------------------- +using op_status = cudf::io::parquet::experimental::variant_operation_status; +namespace expns = cudf::io::parquet::experimental; +auto const& cmr = cudf::get_current_device_resource_ref; + +/** + * @brief Helper using fixed_width_column_wrapper comparison for the common case where the status + * column has no nulls. + */ +static void expect_status_values(cudf::column_view const& status, + std::vector const& expected) +{ + cudf::test::fixed_width_column_wrapper exp(expected.begin(), expected.end()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(status, exp); +} + +/** + * @brief Allocates a non-nullable UINT8 column of `num_rows` rows for callers to pass as the + * `status` output parameter of `get_variant_field`/`cast_variant`/`extract_variant_field`. + */ +static std::unique_ptr make_status_buffer(cudf::size_type num_rows) +{ + return cudf::make_numeric_column( + cudf::data_type{cudf::type_id::UINT8}, num_rows, cudf::mask_state::UNALLOCATED); +} + +constexpr uint8_t ST_SUCCESS = static_cast(op_status::SUCCESS); +constexpr uint8_t ST_ROW_NULL = static_cast(op_status::ROW_NULL); +constexpr uint8_t ST_MISSING = static_cast(op_status::MISSING_PATH); +constexpr uint8_t ST_VNULL = static_cast(op_status::VARIANT_NULL); +constexpr uint8_t ST_MISMATCH = static_cast(op_status::TYPE_MISMATCH); +constexpr uint8_t ST_MALFORMED = static_cast(op_status::MALFORMED_VARIANT); + +// --------------------------------------------------------------------------- +// GetVariantField status tests +// --------------------------------------------------------------------------- + +struct GetVariantFieldStatusTest : public cudf::test::BaseFixture {}; + +TEST_F(GetVariantFieldStatusTest, SqlNullInputProducesRowNullStatus) +{ + // SQL-null input row → null output + row_null status (status column is always non-nullable) + cudf::test::lists_column_wrapper meta{{0x01, 0x01, 0x00, 0x01, 'x'}}; + cudf::test::lists_column_wrapper val{{0x14, 0x07, 0x00, 0x00, 0x00}}; + cudf::test::structs_column_wrapper col{{meta, val}, std::vector{false}}; + + auto stream = cudf::test::get_default_stream(); + auto status = make_status_buffer(cudf::column_view{col}.size()); + auto got = cudf::io::parquet::experimental::get_variant_field( + col, "x", status->mutable_view(), stream, cmr()); + + ASSERT_EQ(status->null_count(), 0); + expect_status_values(*status, {ST_ROW_NULL}); + ASSERT_EQ(got->null_count(), 1); +} + +TEST_F(GetVariantFieldStatusTest, SuccessStatus) +{ + // Successful extraction → success status + auto col = make_xyz_three_row_variant(); + auto stream = cudf::test::get_default_stream(); + + auto status = make_status_buffer(cudf::column_view{col}.size()); + auto got = cudf::io::parquet::experimental::get_variant_field( + col, "x", status->mutable_view(), stream, cmr()); + + // Row 0: x=INT32(7) → success; Row 1: x=INT32(42) → success; Row 2: no x → missing_path + expect_status_values(*status, {ST_SUCCESS, ST_SUCCESS, ST_MISSING}); + // Output rows 0,1 valid; row 2 null + EXPECT_EQ(got->null_count(), 1); +} + +TEST_F(GetVariantFieldStatusTest, MissingKeyProducesMissingPathStatus) +{ + // Missing key → missing_path status + auto col = make_apache_variant(avf::object_primitive); + auto stream = cudf::test::get_default_stream(); + + auto status = make_status_buffer(cudf::column_view{col}.size()); + auto got = cudf::io::parquet::experimental::get_variant_field( + col, "no_such_field", status->mutable_view(), stream, cmr()); + + expect_status_values(*status, {ST_MISSING}); + EXPECT_EQ(got->null_count(), 1); +} + +TEST_F(GetVariantFieldStatusTest, VariantNullPreservedWithStatus) +{ + // VARIANT null terminal value → variant_null status, preserved bytes (non-null output) + // Build a single-row VARIANT: object {null_field: VARIANT_NULL} + // metadata: {null_field}, value: object wrapping NULLVAL primitive + auto const m = build_metadata({"null_field"}); + auto const v = build_single_field_object(/*fid=*/0, enc_null()); + auto col = wrap_single_variant(m, v); + auto stream = cudf::test::get_default_stream(); + + auto status = make_status_buffer(cudf::column_view{col}.size()); + auto got = cudf::io::parquet::experimental::get_variant_field( + col, "null_field", status->mutable_view(), stream, cmr()); + + expect_status_values(*status, {ST_VNULL}); + // With status requested, the VARIANT null bytes are preserved (output is NOT SQL null) + EXPECT_EQ(got->null_count(), 0); + auto const null_bytes = enc_null(); + cudf::test::lists_column_wrapper expected_bytes{{null_bytes.begin(), null_bytes.end()}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected_bytes); +} + +TEST_F(GetVariantFieldStatusTest, VariantNullReturnedAsBytesWithoutStatus) +{ + // Without status_out, VARIANT null is returned as bytes (non-null list row), same as with status. + // Only cast_variant turns a VARIANT null blob into a SQL null. + auto const m = build_metadata({"null_field"}); + auto const v = build_single_field_object(/*fid=*/0, enc_null()); + auto col = wrap_single_variant(m, v); + auto stream = cudf::test::get_default_stream(); + + // No status_out: get_variant_field returns the VARIANT null bytes as a non-null list row. + auto got = + cudf::io::parquet::experimental::get_variant_field(col, "null_field", std::nullopt, stream); + EXPECT_EQ(got->null_count(), 0); + EXPECT_EQ(got->size(), 1); + auto const null_bytes = enc_null(); + cudf::test::lists_column_wrapper expected_bytes{{null_bytes.begin(), null_bytes.end()}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected_bytes); +} + +TEST_F(GetVariantFieldStatusTest, MalformedMetadataProducesMalformedStatus) +{ + // Malformed metadata → malformed_variant status + std::vector const bad_meta = {0x02}; // too short / version ≠ 1 + std::vector const val = {0x14, 0x07, 0x00, 0x00, 0x00}; + auto col = wrap_single_variant(bad_meta, val); + auto stream = cudf::test::get_default_stream(); + + auto status = make_status_buffer(cudf::column_view{col}.size()); + auto got = cudf::io::parquet::experimental::get_variant_field( + col, "x", status->mutable_view(), stream, cmr()); + + expect_status_values(*status, {ST_MALFORMED}); + EXPECT_EQ(got->null_count(), 1); +} + +TEST_F(GetVariantFieldStatusTest, VariantNullBeforeEndIsMissingPath) +{ + // VARIANT null before end of a nested path → missing_path + // Object {a: VARIANT_NULL}; path "$.a.b" should be missing_path (null intermediate) + auto const m = build_metadata({"a"}); + auto const v = build_single_field_object(/*fid=*/0, enc_null()); + auto col = wrap_single_variant(m, v); + auto stream = cudf::test::get_default_stream(); + + auto status = make_status_buffer(cudf::column_view{col}.size()); + auto got = cudf::io::parquet::experimental::get_variant_field( + col, "$.a.b", status->mutable_view(), stream, cmr()); + + expect_status_values(*status, {ST_MISSING}); + EXPECT_EQ(got->null_count(), 1); +} + +TEST_F(GetVariantFieldStatusTest, MixedRows) +{ + // Mixed rows: success / missing / variant_null / malformed / SQL null + auto stream = cudf::test::get_default_stream(); + + auto const dict = build_metadata({"x"}); + + // Row 0: {x: INT32(5)} → success + auto const v0 = build_single_field_object(/*fid=*/0, enc_int32(5)); + // Row 1: {x: NULLVAL} → variant_null + auto const v1 = build_single_field_object(/*fid=*/0, enc_null()); + // Row 2: {} (no x key) → missing_path + auto const m2 = build_metadata({}); + auto const v2 = build_single_field_object(/*fid=*/0, enc_int32(0)); // fid 0 but dict empty + // Row 3: SQL null → row_null status (status column is always non-nullable) + auto const v3 = enc_int32(0); + + cudf::test::lists_column_wrapper meta{{dict.begin(), dict.end()}, + {dict.begin(), dict.end()}, + {m2.begin(), m2.end()}, + {dict.begin(), dict.end()}}; + cudf::test::lists_column_wrapper val{ + {v0.begin(), v0.end()}, {v1.begin(), v1.end()}, {v2.begin(), v2.end()}, {v3.begin(), v3.end()}}; + // Row 3 is SQL null + cudf::test::structs_column_wrapper col{{meta, val}, std::vector{true, true, true, false}}; + + auto status = make_status_buffer(cudf::column_view{col}.size()); + auto got = cudf::io::parquet::experimental::get_variant_field( + col, "x", status->mutable_view(), stream, cmr()); + + ASSERT_EQ(status->null_count(), 0); + expect_status_values(*status, {ST_SUCCESS, ST_VNULL, ST_MISSING, ST_ROW_NULL}); + + // Row 0: valid (INT32 bytes), Row 1: valid (VARIANT null bytes preserved), Row 2+3: null + EXPECT_EQ(got->null_count(), 2); +} + +TEST_F(GetVariantFieldStatusTest, EmptyInput) +{ + // Empty input → empty status column + auto const stream = cudf::test::get_default_stream(); + auto const variant = cudf::empty_like(make_xyz_three_row_variant()); + + auto status = make_status_buffer(variant->size()); + auto got = cudf::io::parquet::experimental::get_variant_field( + *variant, "x", status->mutable_view(), stream, cmr()); + + EXPECT_EQ(status->size(), 0); + EXPECT_EQ(got->size(), 0); +} + +// --------------------------------------------------------------------------- +// CastVariant status tests +// --------------------------------------------------------------------------- + +struct CastVariantStatusTest : public cudf::test::BaseFixture {}; + +namespace { + +inline cudf::test::lists_column_wrapper make_value_col(std::vector const& bytes) +{ + return cudf::test::lists_column_wrapper(bytes.begin(), bytes.end()); +} + +} // namespace + +TEST_F(CastVariantStatusTest, SuccessProducesSuccessStatus) +{ + // Success → success status + auto stream = cudf::test::get_default_stream(); + auto values = make_value_col(enc_int32(42)); + auto status = make_status_buffer(cudf::column_view{values}.size()); + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::INT32}, {}, status->mutable_view(), stream, cmr()); + + expect_status_values(*status, {ST_SUCCESS}); + cudf::test::fixed_width_column_wrapper expected{42}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(CastVariantStatusTest, VariantNullProducesVariantNullStatus) +{ + // VARIANT null → variant_null status + auto stream = cudf::test::get_default_stream(); + auto values = make_value_col(enc_null()); + auto status = make_status_buffer(cudf::column_view{values}.size()); + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::INT32}, {}, status->mutable_view(), stream, cmr()); + + expect_status_values(*status, {ST_VNULL}); + EXPECT_EQ(got->null_count(), 1); +} + +TEST_F(CastVariantStatusTest, TypeMismatchStatus) +{ + // Type mismatch → type_mismatch status + auto stream = cudf::test::get_default_stream(); + auto values = make_value_col(enc_int8(5)); // INT8 cast to INT32 target → mismatch + auto status = make_status_buffer(cudf::column_view{values}.size()); + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::INT32}, {}, status->mutable_view(), stream, cmr()); + + expect_status_values(*status, {ST_MISMATCH}); + EXPECT_EQ(got->null_count(), 1); +} + +TEST_F(CastVariantStatusTest, SqlNullInputProducesRowNullStatus) +{ + // SQL-null input (null list row) → row_null status (status column is always non-nullable) + auto stream = cudf::test::get_default_stream(); + + // Build the values list column directly (two rows), then mask row 1 null. + auto b0 = enc_int32(42); + auto b1 = enc_int32(0); + // offsets: 0, b0.size(), b0.size()+b1.size() + std::vector offsets{ + 0, static_cast(b0.size()), static_cast(b0.size() + b1.size())}; + std::vector flat; + flat.insert(flat.end(), b0.begin(), b0.end()); + flat.insert(flat.end(), b1.begin(), b1.end()); + auto offs_col = + cudf::test::fixed_width_column_wrapper(offsets.begin(), offsets.end()).release(); + auto data_col = + cudf::test::fixed_width_column_wrapper(flat.begin(), flat.end()).release(); + auto values_col = cudf::make_lists_column(2, std::move(offs_col), std::move(data_col), 0, {}); + // Mask row 1 SQL null + auto null_mask = cudf::create_null_mask(2, cudf::mask_state::ALL_VALID, stream, cmr()); + cudf::set_null_mask(static_cast(null_mask.data()), 1, 2, false); + stream.synchronize(); + values_col->set_null_mask(std::move(null_mask), 1); + + auto status = make_status_buffer(values_col->size()); + auto got = cudf::io::parquet::experimental::cast_variant(values_col->view(), + cudf::data_type{cudf::type_id::INT32}, + {}, + status->mutable_view(), + stream, + cmr()); + + // Row 0: success; row 1: row_null (status column is always non-nullable) + ASSERT_EQ(status->null_count(), 0); + expect_status_values(*status, {ST_SUCCESS, ST_ROW_NULL}); + EXPECT_EQ(got->null_count(), 1); +} + +TEST_F(CastVariantStatusTest, IncomingStatusPropagation) +{ + // Incoming status propagation: non-success upstream → propagated status + auto stream = cudf::test::get_default_stream(); + + // 3 rows: success, missing_path, variant_null (from a prior get_variant_field) + // The values column: row 0 = INT32(7), rows 1+2 = anything (won't be decoded for non-success) + std::vector> const val_rows{enc_int32(7), enc_int32(0), enc_null()}; + auto col = + wrap_multi_row_variant(std::vector>(3, build_metadata({})), val_rows); + auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + + // Build incoming_status column: {success, missing_path, variant_null} + cudf::test::fixed_width_column_wrapper incoming_status_w( + {ST_SUCCESS, ST_MISSING, ST_VNULL}); + auto incoming_status_col = incoming_status_w.release(); + auto const incoming_view1 = incoming_status_col->view(); + + auto status = make_status_buffer(values.size()); + auto got = cudf::io::parquet::experimental::cast_variant(values, + cudf::data_type{cudf::type_id::INT32}, + incoming_view1, + status->mutable_view(), + stream, + cmr()); + + // Row 0: success (decoded), Row 1: missing_path (propagated), Row 2: variant_null (propagated) + expect_status_values(*status, {ST_SUCCESS, ST_MISSING, ST_VNULL}); + cudf::test::fixed_width_column_wrapper expected({7, 0, 0}, {true, false, false}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(CastVariantStatusTest, IncomingRowNullStatusProducesRowNullStatus) +{ + // Incoming row_null status → null output and row_null status for that row. + // The status column produced by get_variant_field is non-nullable; SQL-null rows carry row_null. + auto stream = cudf::test::get_default_stream(); + + std::vector> const val_rows{enc_int32(7), enc_int32(1)}; + auto col = + wrap_multi_row_variant(std::vector>(2, build_metadata({})), val_rows); + auto const values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + + // Row 0: success, Row 1: row_null (non-nullable incoming status, as produced by + // get_variant_field) + cudf::test::fixed_width_column_wrapper incoming_status_w2({ST_SUCCESS, ST_ROW_NULL}); + auto incoming_status_col2 = incoming_status_w2.release(); + auto const incoming_view2 = incoming_status_col2->view(); + + auto status = make_status_buffer(values.size()); + auto got = cudf::io::parquet::experimental::cast_variant(values, + cudf::data_type{cudf::type_id::INT32}, + incoming_view2, + status->mutable_view(), + stream, + cmr()); + + ASSERT_EQ(status->null_count(), 0); + expect_status_values(*status, {ST_SUCCESS, ST_ROW_NULL}); + + cudf::test::fixed_width_column_wrapper expected({7, 0}, {true, false}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(CastVariantStatusTest, BoolStatusTracking) +{ + // Status for bool target + auto stream = cudf::test::get_default_stream(); + + // 3 rows: bool_true (success), null (variant_null), int32 (type_mismatch) + std::vector> const val_rows{enc_bool(true), enc_null(), enc_int32(1)}; + auto col = + wrap_multi_row_variant(std::vector>(3, build_metadata({})), val_rows); + auto values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + + auto status = make_status_buffer(values.size()); + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::BOOL8}, {}, status->mutable_view(), stream, cmr()); + + expect_status_values(*status, {ST_SUCCESS, ST_VNULL, ST_MISMATCH}); + cudf::test::fixed_width_column_wrapper expected({true, false, false}, {true, false, false}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(CastVariantStatusTest, StringStatusTracking) +{ + // Status for string target: short_string, variant_null, type_mismatch, malformed long_string, + // truncated short_string, and unrecognized primitive id. + auto stream = cudf::test::get_default_stream(); + + // A SHORT_STRING header that claims 5 bytes of content but provides none. + std::vector const truncated_short_string{make_variant_short_string_header(5)}; + + // An unrecognized primitive type id (0x3F maps to the value_header field of a PRIMITIVE byte + // and does not correspond to any defined primitive_type enum value). + std::vector const unknown_primitive_id{ + make_variant_header(variant_basic_type::PRIMITIVE, 0x3F)}; + + std::vector> const val_rows{ + enc_short_string("hi"), // success + enc_null(), // variant_null + enc_int32(5), // type_mismatch (recognized non-string primitive) + // malformed long_string: header + declares 10 bytes but only 2 present + {make_variant_primitive(cudf::io::parquet::experimental::variant_primitive_type::LONG_STRING), + 0x0A, + 0x00, + 0x00, + 0x00, + 'a', + 'b'}, + truncated_short_string, // malformed: SHORT_STRING with truncated payload + unknown_primitive_id, // malformed: unrecognized primitive id + }; + auto col = + wrap_multi_row_variant(std::vector>(6, build_metadata({})), val_rows); + auto values = cudf::structs_column_view{col}.get_sliced_child(1, stream); + + auto status = make_status_buffer(values.size()); + auto got = cudf::io::parquet::experimental::cast_variant( + values, cudf::data_type{cudf::type_id::STRING}, {}, status->mutable_view(), stream, cmr()); + + expect_status_values( + *status, {ST_SUCCESS, ST_VNULL, ST_MISMATCH, ST_MALFORMED, ST_MALFORMED, ST_MALFORMED}); + EXPECT_EQ(got->null_count(), 5); // all but row 0 are null +} + +TEST_F(CastVariantStatusTest, EmptyInput) +{ + // Empty input → empty status column + auto const stream = cudf::test::get_default_stream(); + auto const values = + cudf::empty_like(cudf::structs_column_view{make_xyz_three_row_variant()}.child(1)); + auto status = make_status_buffer(values->size()); + auto got = cudf::io::parquet::experimental::cast_variant( + *values, cudf::data_type{cudf::type_id::INT32}, {}, status->mutable_view(), stream, cmr()); + + EXPECT_EQ(status->size(), 0); + EXPECT_EQ(got->size(), 0); +} + +// --------------------------------------------------------------------------- +// ExtractVariantField status tests (end-to-end: extraction + decode) +// --------------------------------------------------------------------------- + +struct ExtractVariantFieldStatusTest : public cudf::test::BaseFixture {}; + +TEST_F(ExtractVariantFieldStatusTest, SuccessStatus) +{ + // Success path: object {x: INT32(7)} extracted as INT32 + auto col = make_xyz_three_row_variant(); + auto stream = cudf::test::get_default_stream(); + + auto status = make_status_buffer(cudf::column_view{col}.size()); + auto got = cudf::io::parquet::experimental::extract_variant_field( + col, "x", cudf::data_type{cudf::type_id::INT32}, status->mutable_view(), stream, cmr()); + + // Rows 0,1 have x as INT32 → success; row 2 has no x → missing_path + expect_status_values(*status, {ST_SUCCESS, ST_SUCCESS, ST_MISSING}); + cudf::test::fixed_width_column_wrapper expected({7, 42, 0}, {true, true, false}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*got, expected); +} + +TEST_F(ExtractVariantFieldStatusTest, SqlNullInputProducesRowNullStatus) +{ + // SQL null input → row_null status (status column is always non-nullable) + cudf::test::lists_column_wrapper meta{{0x01, 0x01, 0x00, 0x01, 'x'}}; + cudf::test::lists_column_wrapper val{{0x14, 0x07, 0x00, 0x00, 0x00}}; + cudf::test::structs_column_wrapper col{{meta, val}, std::vector{false}}; + + auto stream = cudf::test::get_default_stream(); + auto status = make_status_buffer(cudf::column_view{col}.size()); + auto got = cudf::io::parquet::experimental::extract_variant_field( + col, "x", cudf::data_type{cudf::type_id::INT32}, status->mutable_view(), stream, cmr()); + + expect_status_values(*status, {ST_ROW_NULL}); + EXPECT_EQ(got->null_count(), 1); +} + +TEST_F(ExtractVariantFieldStatusTest, VariantNullStatus) +{ + // VARIANT null → variant_null status (from extraction phase) + auto const m = build_metadata({"f"}); + auto const v = build_single_field_object(/*fid=*/0, enc_null()); + auto col = wrap_single_variant(m, v); + auto stream = cudf::test::get_default_stream(); + + auto status = make_status_buffer(cudf::column_view{col}.size()); + auto got = cudf::io::parquet::experimental::extract_variant_field( + col, "f", cudf::data_type{cudf::type_id::INT32}, status->mutable_view(), stream, cmr()); + + expect_status_values(*status, {ST_VNULL}); + EXPECT_EQ(got->null_count(), 1); +} + +TEST_F(ExtractVariantFieldStatusTest, TypeMismatchStatus) +{ + // Type mismatch: field exists but is a string, requested as INT32 + auto const m = build_metadata({"s"}); + auto const v = build_single_field_object(/*fid=*/0, enc_short_string("hello")); + auto col = wrap_single_variant(m, v); + auto stream = cudf::test::get_default_stream(); + + auto status = make_status_buffer(cudf::column_view{col}.size()); + auto got = cudf::io::parquet::experimental::extract_variant_field( + col, "s", cudf::data_type{cudf::type_id::INT32}, status->mutable_view(), stream, cmr()); + + expect_status_values(*status, {ST_MISMATCH}); + EXPECT_EQ(got->null_count(), 1); +} + +TEST_F(ExtractVariantFieldStatusTest, MissingNestedPathStatus) +{ + // Missing path for a multi-step path + auto col = make_apache_variant(avf::object_nested); + auto stream = cudf::test::get_default_stream(); + + auto status = make_status_buffer(cudf::column_view{col}.size()); + auto got = + cudf::io::parquet::experimental::extract_variant_field(col, + "$.species.nope", + cudf::data_type{cudf::type_id::STRING}, + status->mutable_view(), + stream, + cmr()); + + expect_status_values(*status, {ST_MISSING}); + EXPECT_EQ(got->null_count(), 1); } diff --git a/java/src/main/native/src/VariantUtilsJni.cpp b/java/src/main/native/src/VariantUtilsJni.cpp index d427da7b6bd..a9c0d97ca9b 100644 --- a/java/src/main/native/src/VariantUtilsJni.cpp +++ b/java/src/main/native/src/VariantUtilsJni.cpp @@ -10,6 +10,8 @@ #include #include +#include + extern "C" { JNIEXPORT jlong JNICALL Java_ai_rapids_cudf_VariantUtils_getVariantFieldValue( @@ -25,6 +27,7 @@ JNIEXPORT jlong JNICALL Java_ai_rapids_cudf_VariantUtils_getVariantFieldValue( return cudf::jni::release_as_jlong( cudf::io::parquet::experimental::get_variant_field(variant_struct, path.get(), + std::nullopt, cudf::get_default_stream(), cudf::get_current_device_resource_ref())); } @@ -44,6 +47,8 @@ JNIEXPORT jlong JNICALL Java_ai_rapids_cudf_VariantUtils_castVariantValue(JNIEnv return cudf::jni::release_as_jlong(cudf::io::parquet::experimental::cast_variant( value_bytes, cudf::data_type{static_cast(cudf_type_id)}, + std::nullopt, + std::nullopt, cudf::get_default_stream(), cudf::get_current_device_resource_ref())); } @@ -64,6 +69,7 @@ JNIEXPORT jlong JNICALL Java_ai_rapids_cudf_VariantUtils_extractVariantField( variant_struct, path.get(), cudf::data_type{static_cast(cudf_type_id)}, + std::nullopt, cudf::get_default_stream(), cudf::get_current_device_resource_ref())); }