From e4870581d383fc8884de266386cd968abab6dd82 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Tue, 5 May 2026 18:38:55 +0000 Subject: [PATCH 1/8] fix validity nullptr check in json reader scatter offsets --- cpp/src/io/json/host_tree_algorithms.cu | 13 ++++-- cpp/tests/io/json/json_test.cpp | 61 +++++++++++++++++++++++++ 2 files changed, 71 insertions(+), 3 deletions(-) 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/tests/io/json/json_test.cpp b/cpp/tests/io/json/json_test.cpp index 8095d945994e..30ac8363ecd7 100644 --- a/cpp/tests/io/json/json_test.cpp +++ b/cpp/tests/io/json/json_test.cpp @@ -3635,4 +3635,65 @@ 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) +{ + // Stray structural characters inside values cause the tokenizer to create + // phantom column-tree entries. The reader must not dereference uninitialised + // column data for those entries during offset scattering. + { + std::string json_string = R"({"a":1,"b":{"c":"ok"},"d":[10,20]})" + "\n" + R"({"a":2,"b":{"c":"x}y"},"d{":[30]})" + "\n" + R"({"a":3,"b":{"c":"ok"},"d":[40,50]})"; + + 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)); + } + { + std::string json_string = R"({"x":[1,2],"y":{"k":1}})" + "\n" + R"({"x":["a]b"],"y":{"k":2}})" + "\n" + R"({"x":[3,4],"y{":{})"; + + 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)); + } +} + CUDF_TEST_PROGRAM_MAIN() From 2ba7c7dfaaf05455dbc9b0a7555dcfa2c7a07513 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Mon, 11 May 2026 11:09:45 -0500 Subject: [PATCH 2/8] add unit tests for few corner cases --- cpp/tests/io/json/json_test.cpp | 64 +++++++----------- cpp/tests/io/json/nested_json_test.cpp | 91 +++++++++++++++++++++++++- 2 files changed, 115 insertions(+), 40 deletions(-) diff --git a/cpp/tests/io/json/json_test.cpp b/cpp/tests/io/json/json_test.cpp index 30ac8363ecd7..b2f6e791e9c1 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,7 +1602,7 @@ 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++) { @@ -1615,7 +1610,7 @@ TEST_F(JsonReaderTest, TestColumnOrder) 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]; @@ -3661,39 +3656,30 @@ TEST_F(JsonReaderTest, MalformedFieldNameWithBrace) TEST_F(JsonReaderTest, MalformedStructuralCharsInValues) { - // Stray structural characters inside values cause the tokenizer to create - // phantom column-tree entries. The reader must not dereference uninitialised - // column data for those entries during offset scattering. - { - std::string json_string = R"({"a":1,"b":{"c":"ok"},"d":[10,20]})" - "\n" - R"({"a":2,"b":{"c":"x}y"},"d{":[30]})" - "\n" - R"({"a":3,"b":{"c":"ok"},"d":[40,50]})"; + // 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_EXPECT_NO_THROW(cudf::io::read_json(options)); - } - { - std::string json_string = R"({"x":[1,2],"y":{"k":1}})" - "\n" - R"({"x":["a]b"],"y":{"k":2}})" - "\n" - R"({"x":[3,4],"y{":{})"; - - 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)); - } + 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..d6db2697e5c4 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,92 @@ TEST_P(JsonDelimiterParamTest, RecoveringTokenStreamNewlineAsWSAndDelimiter) } } +// Parses string values that contain non-ASCII bytes (>= 0x80). +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, DeeplyNestedInput) +{ + std::size_t const depth = std::numeric_limits::max() + 73; + 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 minus one. +TEST_F(JsonTest, NestedInputBelowDepthLimit) +{ + std::size_t const depth = std::numeric_limits::max() - 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_NO_THROW(cudf::io::read_json(opts)); +} + CUDF_TEST_PROGRAM_MAIN() From 9557471acf3c984c6f1e9e2bd9bee57464c0620c Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Mon, 11 May 2026 11:11:05 -0500 Subject: [PATCH 3/8] fix type in json validation code --- cpp/src/io/json/process_tokens.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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++) { From 861314ba36cb6090b8c8ab33557a3cdf4833fc72 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Mon, 11 May 2026 11:12:46 -0500 Subject: [PATCH 4/8] export to enable test code --- cpp/src/io/json/nested_json.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) 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, From 71623d0b84fd173f48983b0e8c6d3f78d3bc3e7b Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Mon, 11 May 2026 11:13:38 -0500 Subject: [PATCH 5/8] Fix casting issue --- cpp/src/io/json/nested_json_gpu.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/src/io/json/nested_json_gpu.cu b/cpp/src/io/json/nested_json_gpu.cu index a11d2aefbf79..6d34dbb42601 100644 --- a/cpp/src/io/json/nested_json_gpu.cu +++ b/cpp/src/io/json/nested_json_gpu.cu @@ -612,7 +612,8 @@ struct PdaSymbolToSymbolGroupId { 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) + From c8fa5a9a8008aa9a62f4e649b28b05249958fedb Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Mon, 11 May 2026 11:17:51 -0500 Subject: [PATCH 6/8] add error if json max depth exceeds limit --- cpp/src/io/json/json_tree.cu | 36 +++++++++++++++++++++++++++++++----- 1 file changed, 31 insertions(+), 5 deletions(-) diff --git a/cpp/src/io/json/json_tree.cu b/cpp/src/io/json/json_tree.cu index 4c46b0617676..07fcd9b10519 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,22 @@ struct is_nested_end { } }; +struct checked_token_level_output { + bool* depth_out_of_range; + + __device__ TreeDepthT operator()(int32_t level) const + { + 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 * @@ -280,14 +296,24 @@ tree_meta_t get_tree_representation(device_span tokens, rmm::device_uvector token_levels(num_tokens, stream); auto const push_pop_it = thrust::make_transform_iterator( tokens.begin(), - cuda::proclaim_return_type( - [does_push, does_pop] __device__(PdaTokenT const token) -> size_type { + cuda::proclaim_return_type( + [does_push, does_pop] __device__(PdaTokenT const token) -> int32_t { 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, + int32_t{0}); + 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())) + "]"); auto const node_levels_end = cudf::detail::copy_if(token_levels.begin(), token_levels.end(), From 2a8e21d09998db35fe9464e9dfd77efadd41a3d3 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Tue, 12 May 2026 13:04:35 -0500 Subject: [PATCH 7/8] address review comments --- cpp/src/io/json/json_tree.cu | 11 ++++++----- cpp/tests/io/json/nested_json_test.cpp | 4 ++-- 2 files changed, 8 insertions(+), 7 deletions(-) diff --git a/cpp/src/io/json/json_tree.cu b/cpp/src/io/json/json_tree.cu index 07fcd9b10519..9775cc781cad 100644 --- a/cpp/src/io/json/json_tree.cu +++ b/cpp/src/io/json/json_tree.cu @@ -134,6 +134,7 @@ struct checked_token_level_output { __device__ TreeDepthT operator()(int32_t level) const { + static_assert(sizeof(TreeDepthT) < sizeof(int32_t)); 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}; @@ -309,11 +310,6 @@ tree_meta_t get_tree_representation(device_span tokens, push_pop_it + num_tokens, token_level_output_it, int32_t{0}); - 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())) + "]"); auto const node_levels_end = cudf::detail::copy_if(token_levels.begin(), token_levels.end(), @@ -321,6 +317,11 @@ 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/tests/io/json/nested_json_test.cpp b/cpp/tests/io/json/nested_json_test.cpp index d6db2697e5c4..888f575c0e85 100644 --- a/cpp/tests/io/json/nested_json_test.cpp +++ b/cpp/tests/io/json/nested_json_test.cpp @@ -1424,7 +1424,7 @@ TEST_F(JsonTest, RejectsUnquotedValuesWithInvalidLeadingChar) // Rejects JSON inputs whose nesting depth exceeds the supported maximum. TEST_F(JsonTest, DeeplyNestedInput) { - std::size_t const depth = std::numeric_limits::max() + 73; + 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( @@ -1437,7 +1437,7 @@ TEST_F(JsonTest, DeeplyNestedInput) // Accepts JSON inputs at the maximum supported nesting depth minus one. TEST_F(JsonTest, NestedInputBelowDepthLimit) { - std::size_t const depth = std::numeric_limits::max() - 1; + 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( From 758714ac624489617fbadd1aba1116da7f400855 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Wed, 13 May 2026 22:35:54 -0500 Subject: [PATCH 8/8] address review comments --- cpp/src/io/json/json_tree.cu | 19 ++++++++++--------- cpp/src/io/json/nested_json_gpu.cu | 2 ++ cpp/tests/io/json/json_test.cpp | 2 +- cpp/tests/io/json/nested_json_test.cpp | 7 +++---- 4 files changed, 16 insertions(+), 14 deletions(-) diff --git a/cpp/src/io/json/json_tree.cu b/cpp/src/io/json/json_tree.cu index 9775cc781cad..f38c4f0c55cc 100644 --- a/cpp/src/io/json/json_tree.cu +++ b/cpp/src/io/json/json_tree.cu @@ -132,11 +132,11 @@ struct is_nested_end { struct checked_token_level_output { bool* depth_out_of_range; - __device__ TreeDepthT operator()(int32_t level) const + __device__ TreeDepthT operator()(size_type level) const { - static_assert(sizeof(TreeDepthT) < sizeof(int32_t)); - if (level < static_cast(cuda::std::numeric_limits::min()) || - level > static_cast(cuda::std::numeric_limits::max())) { + 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); @@ -297,8 +297,8 @@ tree_meta_t get_tree_representation(device_span tokens, rmm::device_uvector token_levels(num_tokens, stream); auto const push_pop_it = thrust::make_transform_iterator( tokens.begin(), - cuda::proclaim_return_type( - [does_push, does_pop] __device__(PdaTokenT const token) -> int32_t { + cuda::proclaim_return_type( + [does_push, does_pop] __device__(PdaTokenT const token) -> size_type { return does_push(token) - does_pop(token); })); auto depth_out_of_range = @@ -309,7 +309,7 @@ tree_meta_t get_tree_representation(device_span tokens, push_pop_it, push_pop_it + num_tokens, token_level_output_it, - int32_t{0}); + size_type{0}); auto const node_levels_end = cudf::detail::copy_if(token_levels.begin(), token_levels.end(), @@ -320,8 +320,9 @@ tree_meta_t get_tree_representation(device_span tokens, 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())) + "]"); + 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_gpu.cu b/cpp/src/io/json/nested_json_gpu.cu index 6d34dbb42601..63a46710f495 100644 --- a/cpp/src/io/json/nested_json_gpu.cu +++ b/cpp/src/io/json/nested_json_gpu.cu @@ -609,6 +609,8 @@ 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) diff --git a/cpp/tests/io/json/json_test.cpp b/cpp/tests/io/json/json_test.cpp index b2f6e791e9c1..edde943c9392 100644 --- a/cpp/tests/io/json/json_test.cpp +++ b/cpp/tests/io/json/json_test.cpp @@ -1605,7 +1605,7 @@ TEST_F(JsonReaderTest, TestColumnOrder) // 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); } diff --git a/cpp/tests/io/json/nested_json_test.cpp b/cpp/tests/io/json/nested_json_test.cpp index 888f575c0e85..c5aa6f028723 100644 --- a/cpp/tests/io/json/nested_json_test.cpp +++ b/cpp/tests/io/json/nested_json_test.cpp @@ -1359,7 +1359,6 @@ TEST_P(JsonDelimiterParamTest, RecoveringTokenStreamNewlineAsWSAndDelimiter) } } -// Parses string values that contain non-ASCII bytes (>= 0x80). TEST_F(JsonTest, StringContainingNonAsciiBytes) { for (int b : {0x80, 0xA0, 0xC3, 0xE2, 0xF0, 0xFF}) { @@ -1422,7 +1421,7 @@ TEST_F(JsonTest, RejectsUnquotedValuesWithInvalidLeadingChar) } // Rejects JSON inputs whose nesting depth exceeds the supported maximum. -TEST_F(JsonTest, DeeplyNestedInput) +TEST_F(JsonTest, NestedInputAboveDepthLimit) { std::size_t const depth = std::numeric_limits::max() + std::size_t{1}; std::string s(depth, '['); @@ -1434,8 +1433,8 @@ TEST_F(JsonTest, DeeplyNestedInput) EXPECT_THROW(cudf::io::read_json(opts), cudf::logic_error); } -// Accepts JSON inputs at the maximum supported nesting depth minus one. -TEST_F(JsonTest, NestedInputBelowDepthLimit) +// 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, '[');