diff --git a/cpp/src/io/json/host_tree_algorithms.cu b/cpp/src/io/json/host_tree_algorithms.cu index 74c7e5e31e6e..5b54d58f6ebf 100644 --- a/cpp/src/io/json/host_tree_algorithms.cu +++ b/cpp/src/io/json/host_tree_algorithms.cu @@ -898,12 +898,19 @@ void scatter_offsets(tree_meta_t const& tree, if (d_ignore_vals[col_ids[i]]) return; auto const node_category = column_categories[col_ids[i]]; switch (node_category) { - case NC_STRUCT: set_bit(d_columns_data[col_ids[i]].validity, row_offsets[i]); break; - case NC_LIST: set_bit(d_columns_data[col_ids[i]].validity, row_offsets[i]); break; + case NC_STRUCT: + if (d_columns_data[col_ids[i]].validity) + set_bit(d_columns_data[col_ids[i]].validity, row_offsets[i]); + break; + case NC_LIST: + if (d_columns_data[col_ids[i]].validity) + set_bit(d_columns_data[col_ids[i]].validity, row_offsets[i]); + break; case NC_STR: [[fallthrough]]; case NC_VAL: if (d_ignore_vals[col_ids[i]]) break; - set_bit(d_columns_data[col_ids[i]].validity, row_offsets[i]); + if (d_columns_data[col_ids[i]].validity) + set_bit(d_columns_data[col_ids[i]].validity, row_offsets[i]); d_columns_data[col_ids[i]].string_offsets[row_offsets[i]] = range_begin[i]; d_columns_data[col_ids[i]].string_lengths[row_offsets[i]] = range_end[i] - range_begin[i]; break; diff --git a/cpp/src/io/json/json_tree.cu b/cpp/src/io/json/json_tree.cu index 4c46b0617676..f38c4f0c55cc 100644 --- a/cpp/src/io/json/json_tree.cu +++ b/cpp/src/io/json/json_tree.cu @@ -26,8 +26,10 @@ #include #include #include +#include #include #include +#include #include #include #include @@ -44,8 +46,6 @@ #include #include -#include - namespace cudf::io::json { namespace detail { @@ -129,6 +129,23 @@ struct is_nested_end { } }; +struct checked_token_level_output { + bool* depth_out_of_range; + + __device__ TreeDepthT operator()(size_type level) const + { + static_assert(sizeof(TreeDepthT) < sizeof(size_type)); + if (level < static_cast(cuda::std::numeric_limits::min()) || + level > static_cast(cuda::std::numeric_limits::max())) { + cuda::atomic_ref flag{*depth_out_of_range}; + if (!flag.load(cuda::std::memory_order_relaxed)) { + flag.store(true, cuda::std::memory_order_relaxed); + } + } + return static_cast(level); + } +}; + /** * @brief Returns stable sorted keys and its sorted order * @@ -284,10 +301,15 @@ tree_meta_t get_tree_representation(device_span tokens, [does_push, does_pop] __device__(PdaTokenT const token) -> size_type { return does_push(token) - does_pop(token); })); + auto depth_out_of_range = + cudf::detail::device_scalar(false, stream, cudf::get_current_device_resource_ref()); + auto const token_level_output_it = thrust::make_transform_output_iterator( + token_levels.begin(), checked_token_level_output{depth_out_of_range.data()}); thrust::exclusive_scan(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), push_pop_it, push_pop_it + num_tokens, - token_levels.begin()); + token_level_output_it, + size_type{0}); auto const node_levels_end = cudf::detail::copy_if(token_levels.begin(), token_levels.end(), @@ -295,6 +317,12 @@ tree_meta_t get_tree_representation(device_span tokens, node_levels.begin(), is_node, stream); + CUDF_EXPECTS( + !depth_out_of_range.value(stream), + "JSON token nesting depth is outside the supported range for TreeDepthT [" + + std::to_string(static_cast(cuda::std::numeric_limits::min())) + + ", " + + std::to_string(static_cast(cuda::std::numeric_limits::max())) + "]"); CUDF_EXPECTS(cuda::std::distance(node_levels.begin(), node_levels_end) == static_cast(num_nodes), "node level count mismatch"); diff --git a/cpp/src/io/json/nested_json.hpp b/cpp/src/io/json/nested_json.hpp index 9ad36b56cda0..2b2962d4e2cb 100644 --- a/cpp/src/io/json/nested_json.hpp +++ b/cpp/src/io/json/nested_json.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -276,6 +276,7 @@ std::pair, rmm::device_uvector> pr * @param options Parsing options specifying the parsing behaviour * @param stream The cuda stream to dispatch GPU kernels to */ +CUDF_EXPORT void validate_token_stream(device_span d_input, device_span tokens, device_span token_indices, diff --git a/cpp/src/io/json/nested_json_gpu.cu b/cpp/src/io/json/nested_json_gpu.cu index a11d2aefbf79..63a46710f495 100644 --- a/cpp/src/io/json/nested_json_gpu.cu +++ b/cpp/src/io/json/nested_json_gpu.cu @@ -609,10 +609,13 @@ struct PdaSymbolToSymbolGroupId { // escape, comma, colon or whitespace characters. auto constexpr newline = '\n'; auto constexpr whitespace = ' '; + // Cast to unsigned char first so high-bit bytes (>= 0x80) are not sign-extended to negative + // int32_t values, which would underflow the min() clamp used as the lookup index. auto const symbol_position = symbol == delimiter ? static_cast(newline) - : (symbol == newline ? static_cast(whitespace) : static_cast(symbol)); + : (symbol == newline ? static_cast(whitespace) + : static_cast(static_cast(symbol))); PdaSymbolGroupIdT symbol_gid = tos_sg_to_pda_sgid[min(symbol_position, pda_sgid_lookup_size - 1)]; return stack_idx * static_cast(symbol_group_id::NUM_PDA_INPUT_SGS) + diff --git a/cpp/src/io/json/process_tokens.cu b/cpp/src/io/json/process_tokens.cu index cd57ebc1a3e2..3a249189e3ff 100644 --- a/cpp/src/io/json/process_tokens.cu +++ b/cpp/src/io/json/process_tokens.cu @@ -113,7 +113,7 @@ void validate_token_stream(device_span d_input, if (is_nonnumeric) { return true; } } auto c = data[start]; - if ('-' == c || c <= '9' && 'c' >= '0') { + if ('-' == c || (c <= '9' && c >= '0')) { // number auto num_state = number_state::START; for (auto at = start; at < end; at++) { diff --git a/cpp/tests/io/json/json_test.cpp b/cpp/tests/io/json/json_test.cpp index 8095d945994e..edde943c9392 100644 --- a/cpp/tests/io/json/json_test.cpp +++ b/cpp/tests/io/json/json_test.cpp @@ -22,16 +22,11 @@ #include #include #include -#include #include #include #include -#include -#include - #include -#include #include #include @@ -1607,15 +1602,15 @@ TEST_F(JsonReaderTest, TestColumnOrder) // Read in data using nested JSON reader cudf::io::table_with_metadata new_reader_table = cudf::io::read_json(json_lines_options); - // Verify root column order (assert to avoid OOB access) + // Verify root column order before accessing schema entries. ASSERT_EQ(new_reader_table.metadata.schema_info.size(), root_col_names.size()); - for (std::size_t i = 0; i < a_child_col_names.size(); i++) { + for (std::size_t i = 0; i < root_col_names.size(); i++) { auto const& root_col_name = root_col_names[i]; EXPECT_EQ(new_reader_table.metadata.schema_info[i].name, root_col_name); } - // Verify nested child column order (assert to avoid OOB access) + // Verify nested child column order before accessing schema entries. ASSERT_EQ(new_reader_table.metadata.schema_info[2].children.size(), a_child_col_names.size()); for (std::size_t i = 0; i < a_child_col_names.size(); i++) { auto const& a_child_col_name = a_child_col_names[i]; @@ -3635,4 +3630,56 @@ TEST_F(JsonReaderTest, DeviceWriteAsyncThrows) } } +TEST_F(JsonReaderTest, MalformedFieldNameWithBrace) +{ + // Garbled field name containing '{' creates structural ambiguity in the + // token tree. Combined with an invalid byte in a sibling row, the parser + // may produce column-tree nodes that are never materialised as columns. + // Recovery mode must handle this without accessing uninitialised memory. + std::string json_string = + R"({"name":"Alice","address":{"city":"NYC","zip":"1"},"score{":[95,2]})" + "\n" + R"({"name":"Bob","address":{"city":"LA","zip":"1"},"scores":[)" + "\xbf" + R"(8,82]})" + "\n" + R"({"name":"C","address":{"city":"Chi","zip":"60601"},"scores":[1,97]})"; + + cudf::io::json_reader_options options = + cudf::io::json_reader_options::builder( + cudf::io::source_info{cudf::host_span{ + reinterpret_cast(json_string.data()), json_string.size()}}) + .lines(true) + .recovery_mode(cudf::io::json_recovery_mode_t::RECOVER_WITH_NULL); + CUDF_EXPECT_NO_THROW(cudf::io::read_json(options)); +} + +TEST_F(JsonReaderTest, MalformedStructuralCharsInValues) +{ + // A malformed middle line must not affect surrounding well-formed records. + std::string json_string = R"({"a":1,"b":[10,20]})" + "\n" + R"({"phantom":{"nested":[30,{"c":40)" + "\n" + R"({"a":3,"b":[50,60]})"; + + cudf::io::json_reader_options options = + cudf::io::json_reader_options::builder( + cudf::io::source_info{cudf::host_span{ + reinterpret_cast(json_string.data()), json_string.size()}}) + .lines(true) + .recovery_mode(cudf::io::json_recovery_mode_t::RECOVER_WITH_NULL); + cudf::io::table_with_metadata tbl; + ASSERT_NO_THROW(tbl = cudf::io::read_json(options)); + ASSERT_EQ(tbl.tbl->num_rows(), 3); + ASSERT_EQ(tbl.tbl->num_columns(), 2); + cudf::test::fixed_width_column_wrapper expected_a{{1, 0, 3}, {true, false, true}}; + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(tbl.tbl->view().column(0), expected_a); + EXPECT_EQ(tbl.tbl->get_column(0).null_count(), 1); + // Column "b": list [[10,20], null, [50,60]]. + using LCWI = cudf::test::lists_column_wrapper; + LCWI expected_b{{LCWI{10, 20}, LCWI{}, LCWI{50, 60}}, std::vector{1, 0, 1}.begin()}; + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(tbl.tbl->view().column(1), expected_b); +} + CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/io/json/nested_json_test.cpp b/cpp/tests/io/json/nested_json_test.cpp index 9c61796aaf9e..c5aa6f028723 100644 --- a/cpp/tests/io/json/nested_json_test.cpp +++ b/cpp/tests/io/json/nested_json_test.cpp @@ -14,7 +14,6 @@ #include #include -#include #include #include #include @@ -22,6 +21,8 @@ #include #include +#include +#include #include namespace cuio_json = cudf::io::json; @@ -1358,4 +1359,91 @@ TEST_P(JsonDelimiterParamTest, RecoveringTokenStreamNewlineAsWSAndDelimiter) } } +TEST_F(JsonTest, StringContainingNonAsciiBytes) +{ + for (int b : {0x80, 0xA0, 0xC3, 0xE2, 0xF0, 0xFF}) { + std::string const expected_cell{static_cast(b)}; + std::string s{R"({"k":")"}; + s += expected_cell; + s += R"("})"; + s += '\n'; + auto const opts = cudf::io::json_reader_options::builder( + cudf::io::source_info{cudf::host_span{ + reinterpret_cast(s.data()), s.size()}}) + .lines(true) + .build(); + cudf::io::table_with_metadata tbl; + ASSERT_NO_THROW(tbl = cudf::io::read_json(opts)) << "byte 0x" << std::hex << b; + ASSERT_EQ(tbl.tbl->num_columns(), 1) << "byte 0x" << std::hex << b; + ASSERT_EQ(tbl.tbl->num_rows(), 1) << "byte 0x" << std::hex << b; + cudf::test::strings_column_wrapper expected({expected_cell}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(tbl.tbl->view().column(0), expected); + } +} + +// Rejects unquoted values whose first byte is neither '-' nor a digit. +TEST_F(JsonTest, RejectsUnquotedValuesWithInvalidLeadingChar) +{ + for (char const* bad : {"!", "+", " ", ".", "x", "/"}) { + std::string s = std::string{R"({"k":)"} + bad + "}\n"; + auto const opts = cudf::io::json_reader_options::builder(cudf::io::source_info{}) + .lines(true) + .recovery_mode(cudf::io::json_recovery_mode_t::RECOVER_WITH_NULL) + .strict_validation(true) + .build(); + cudf::string_scalar const d_scalar(s, true); + auto const d_input = cudf::device_span{ + d_scalar.data(), static_cast(d_scalar.size())}; + auto const stream = cudf::get_default_stream(); + using token_t = cuio_json::token_t; + std::vector tokens{token_t::StructBegin, + token_t::StructMemberBegin, + token_t::FieldNameBegin, + token_t::FieldNameEnd, + token_t::ValueBegin, + token_t::ValueEnd, + token_t::StructMemberEnd, + token_t::StructEnd, + token_t::LineEnd}; + std::vector token_indices{0, 0, 1, 3, 5, 6, 6, 6, 7}; + auto d_tokens = cudf::detail::make_device_uvector_async( + tokens, stream, cudf::get_current_device_resource_ref()); + auto d_token_indices = cudf::detail::make_device_uvector_async( + token_indices, stream, cudf::get_current_device_resource_ref()); + + cuio_json::detail::validate_token_stream(d_input, d_tokens, d_token_indices, opts, stream); + auto const validated_tokens = cudf::detail::make_std_vector_async(d_tokens, stream); + stream.synchronize(); + EXPECT_NE(std::find(validated_tokens.begin(), validated_tokens.end(), token_t::ErrorBegin), + validated_tokens.end()) + << "value " << bad << " was unexpectedly accepted as a number"; + } +} + +// Rejects JSON inputs whose nesting depth exceeds the supported maximum. +TEST_F(JsonTest, NestedInputAboveDepthLimit) +{ + std::size_t const depth = std::numeric_limits::max() + std::size_t{1}; + std::string s(depth, '['); + s.append(depth, ']'); + auto const opts = cudf::io::json_reader_options::builder( + cudf::io::source_info{cudf::host_span{ + reinterpret_cast(s.data()), s.size()}}) + .build(); + EXPECT_THROW(cudf::io::read_json(opts), cudf::logic_error); +} + +// Accepts JSON inputs at the maximum supported nesting depth. +TEST_F(JsonTest, NestedInputAtDepthLimit) +{ + std::size_t const depth = std::numeric_limits::max(); + std::string s(depth, '['); + s.append(depth, ']'); + auto const opts = cudf::io::json_reader_options::builder( + cudf::io::source_info{cudf::host_span{ + reinterpret_cast(s.data()), s.size()}}) + .build(); + EXPECT_NO_THROW(cudf::io::read_json(opts)); +} + CUDF_TEST_PROGRAM_MAIN()