Fix JSON reader guards for scatter validity, validation, and max nesting depth

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;

cpp/src/io/json/json_tree.cu

@@ -26,8 +26,10 @@
 #include <cub/device/device_radix_sort.cuh>
 #include <cuco/static_map.cuh>
 #include <cuco/static_set.cuh>
+#include <cuda/atomic>
 #include <cuda/functional>
 #include <cuda/iterator>
+#include <cuda/std/limits>
 #include <cuda/std/tuple>
 #include <thrust/binary_search.h>
 #include <thrust/count.h>
@@ -44,8 +46,6 @@
 #include <thrust/tabulate.h>
 #include <thrust/transform.h>
 
-#include <limits>
-
 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<size_type>(cuda::std::numeric_limits<TreeDepthT>::min()) ||
+        level > static_cast<size_type>(cuda::std::numeric_limits<TreeDepthT>::max())) {
+      cuda::atomic_ref<bool, cuda::thread_scope_device> flag{*depth_out_of_range};
+      if (!flag.load(cuda::std::memory_order_relaxed)) {
+        flag.store(true, cuda::std::memory_order_relaxed);
+      }
+    }
+    return static_cast<TreeDepthT>(level);
+  }
+};
+
 /**
  * @brief Returns stable sorted keys and its sorted order
  *
@@ -284,17 +301,28 @@ tree_meta_t get_tree_representation(device_span<PdaTokenT const> 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<bool>(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(),
                                                        tokens.begin(),
                                                        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<size_type>(cuda::std::numeric_limits<TreeDepthT>::min())) +
+        ", " +
+        std::to_string(static_cast<size_type>(cuda::std::numeric_limits<TreeDepthT>::max())) + "]");
     CUDF_EXPECTS(cuda::std::distance(node_levels.begin(), node_levels_end) ==
                    static_cast<std::ptrdiff_t>(num_nodes),
                  "node level count mismatch");

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<PdaTokenT>, rmm::device_uvector<SymbolOffsetT>> 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<char const> d_input,
                            device_span<PdaTokenT> tokens,
                            device_span<SymbolOffsetT> token_indices,

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<int32_t>(newline)
-        : (symbol == newline ? static_cast<int32_t>(whitespace) : static_cast<int32_t>(symbol));
+        : (symbol == newline ? static_cast<int32_t>(whitespace)
+                             : static_cast<int32_t>(static_cast<unsigned char>(symbol)));
     PdaSymbolGroupIdT symbol_gid =
       tos_sg_to_pda_sgid[min(symbol_position, pda_sgid_lookup_size - 1)];
     return stack_idx * static_cast<PdaSymbolGroupIdT>(symbol_group_id::NUM_PDA_INPUT_SGS) +

cpp/src/io/json/process_tokens.cu

@@ -113,7 +113,7 @@ void validate_token_stream(device_span<char const> 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++) {

cpp/tests/io/json/json_test.cpp

@@ -22,16 +22,11 @@
 #include <cudf/strings/convert/convert_fixed_point.hpp>
 #include <cudf/strings/repeat_strings.hpp>
 #include <cudf/strings/strings_column_view.hpp>
-#include <cudf/table/table.hpp>
 #include <cudf/table/table_view.hpp>
 #include <cudf/types.hpp>
 #include <cudf/utilities/pinned_memory.hpp>
 
-#include <rmm/mr/pinned_host_memory_resource.hpp>
-#include <rmm/mr/pool_memory_resource.hpp>
-
 #include <cuda/iterator>
-#include <cuda/memory_resource>
 
 #include <fstream>
 #include <limits>
@@ -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<std::byte const>{
+        reinterpret_cast<std::byte const*>(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<std::byte const>{
+        reinterpret_cast<std::byte const*>(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<int64_t> 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<int64> [[10,20], null, [50,60]].
+  using LCWI = cudf::test::lists_column_wrapper<int64_t>;
+  LCWI expected_b{{LCWI{10, 20}, LCWI{}, LCWI{50, 60}}, std::vector<bool>{1, 0, 1}.begin()};
+  CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(tbl.tbl->view().column(1), expected_b);
+}
+
 CUDF_TEST_PROGRAM_MAIN()

cpp/tests/io/json/nested_json_test.cpp

@@ -14,14 +14,15 @@
 #include <cudf_test/testing_main.hpp>
 
 #include <cudf/io/json.hpp>
-#include <cudf/io/types.hpp>
 #include <cudf/scalar/scalar.hpp>
 #include <cudf/utilities/default_stream.hpp>
 #include <cudf/utilities/span.hpp>
 
 #include <cuda/std/tuple>
 #include <thrust/iterator/zip_iterator.h>
 
+#include <algorithm>
+#include <limits>
 #include <string>
 
 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<char>(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<std::byte const>{
+                          reinterpret_cast<std::byte const*>(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<cuio_json::SymbolT const>{
+      d_scalar.data(), static_cast<size_t>(d_scalar.size())};
+    auto const stream = cudf::get_default_stream();
+    using token_t     = cuio_json::token_t;
+    std::vector<cuio_json::PdaTokenT> 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<cuio_json::SymbolOffsetT> 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<int8_t>::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<std::byte const>{
+                        reinterpret_cast<std::byte const*>(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<int8_t>::max();
+  std::string s(depth, '[');
+  s.append(depth, ']');
+  auto const opts = cudf::io::json_reader_options::builder(
+                      cudf::io::source_info{cudf::host_span<std::byte const>{
+                        reinterpret_cast<std::byte const*>(s.data()), s.size()}})
+                      .build();
+  EXPECT_NO_THROW(cudf::io::read_json(opts));
+}
+
 CUDF_TEST_PROGRAM_MAIN()