Add a protocol buffer decode kernel

[thirtiseven]

This PR adds a protocol buffer decoder with a large subset of Proto2 features, to support spark expression from_protobuf.

The code is ready to me but too large to review. I'm splitting it into small parts:

Part 0: #4373

PR Split Plan

This PR is being split into a chain of focused PRs to make review manageable. Each PR is independently buildable and testable, and each subsequent PR is largely additive.

Please check NVIDIA/cudf-spark#14069 for detailed pr split plan.

Summary

This PR adds a GPU-accelerated protobuf decoder that converts LIST<INT8/UINT8> columns (one serialized protobuf message per row) into nested cuDF STRUCT columns via JNI. This is the native kernel layer that powers from_protobuf() GPU acceleration in the spark-rapids plugin.

The implementation spans ~6,500 lines of new CUDA/C++/Java code and ~3,900 lines of tests, organized into a clean four-file C++ architecture plus a validated Java schema API.

Key capabilities

• All scalar protobuf types: int32, int64, uint32, uint64, sint32/sint64 (zigzag), fixed32/sfixed32/fixed64/sfixed64, float, double, bool, string, bytes
• Nested messages: up to 10 levels deep, recursive decode
• Repeated fields: both packed and unpacked encoding, auto-detected per-row
• Repeated messages: ArrayType(StructType) — repeated nested messages with arbitrary child fields
• Repeated-in-nested: repeated fields inside nested messages, repeated fields inside repeated messages
• Enum-as-string: varint → validated enum → UTF-8 string name conversion, with lookup tables for GPU-parallel name resolution
• Default values: per-field defaults for all scalar types and strings
• Required field validation: proto2-style required field checks
• PERMISSIVE / FAILFAST modes: configurable error handling — permissive mode nullifies malformed rows instead of throwing; invalid enum values nullify the entire struct row and propagate nulls to all descendants
• Schema projection ready: the flattened schema representation supports decoding arbitrary subsets of fields

Performance characteristics

• Multi-pass algorithm optimized for GPU occupancy: many simple kernels > one complex kernel
• O(1) field-number lookup tables for scan/count kernels (up to FIELD_LOOKUP_TABLE_MAX = 4096)
• Batched scalar extraction groups fields by type to minimize kernel launches
• Two-phase string construction (compute lengths → prefix sum → copy) avoids pre-allocation guessing

Architecture

File structure

src/main/cpp/src/
├── protobuf.hpp            (279 lines)  Public API: types, context, validation
├── protobuf_common.cuh    (1823 lines)  Shared types, device helpers, template kernels
├── protobuf_kernels.cu    (1307 lines)  Non-template CUDA kernels
├── protobuf_builders.cu   (1719 lines)  Column builder functions
├── protobuf.cu            (1196 lines)  Entry point: decode_protobuf_to_struct
└── ProtobufJni.cpp         (278 lines)  JNI bridge

src/main/java/.../jni/
├── Protobuf.java            (116 lines)  Java public API
└── ProtobufSchemaDescriptor.java (319 lines)  Immutable schema with validation

src/test/java/.../jni/
├── ProtobufTest.java       (3565 lines)  107 decode tests
└── ProtobufSchemaDescriptorTest.java (338 lines)  13 schema validation tests

src/main/cpp/benchmarks/
└── protobuf_decode.cu      (1322 lines)  8 NVBench benchmarks

Dependency graph

Protobuf.java ──► ProtobufSchemaDescriptor.java
     │
     │ JNI
     ▼
ProtobufJni.cpp ──► protobuf.hpp
                         │
                         ▼
                    protobuf.cu (entry point, orchestration)
                         │
                         │ #include
                         ▼
                    protobuf_common.cuh (shared foundation)
                    ▲            ▲
                    │            │
          protobuf_kernels.cu  protobuf_builders.cu

Multi-pass decode algorithm

The decoder processes each batch of messages through multiple GPU passes:

1. Count pass (count_repeated_fields_kernel): One thread per row. Scans message bytes to count repeated field occurrences and record nested message locations. Handles both packed and unpacked repeated encoding.
2. Offset pass (thrust::exclusive_scan): Prefix sum on repeated counts to compute output array offsets.
3. Scan pass (scan_all_fields_kernel + scan_all_repeated_occurrences_kernel): Records exact byte locations (offset + length) for every target field. Last-one-wins semantics for duplicate scalar fields.
4. Extract pass (type-specific kernels): Parallel data extraction using pre-computed locations. Batched 2D kernel launches group fields by type (varint, fixed32, fixed64, zigzag, etc.) to minimize launch overhead.
5. Build pass (recursive column builders): Assembles cuDF columns bottom-up. Nested structs and repeated messages are processed recursively up to MAX_NESTED_STRUCT_DECODE_DEPTH = 10.

Flattened schema representation

The protobuf schema is represented as parallel arrays passed through JNI. Fields are ordered in pre-order traversal (parent before children):

Array	Description
fieldNumbers[]	Protobuf field numbers
parentIndices[]	Parent index in flat array (-1 for top-level)
depthLevels[]	Nesting depth (0 for top-level)
wireTypes[]	Expected protobuf wire type (0=varint, 1=64bit, 2=len, 5=32bit)
outputTypeIds[]	cuDF type IDs for output columns
encodings[]	Encoding (0=default, 1=fixed, 2=zigzag, 3=enum_string)
isRepeated[]	Whether field is repeated (output becomes LIST)
isRequired[]	Whether field is required (proto2)
hasDefaultValue[]	Whether a default value exists
defaultInts/Floats/Bools/Strings[]	Default values per field
enumValidValues[][]	Sorted valid enum values per field (for binary search)
enumNames[][][]	Enum name UTF-8 bytes per field (for enum-as-string)

Example for message Outer { int32 a = 1; Inner b = 2; } message Inner { int32 x = 1; string y = 2; }:

Index 0: a  (parentIdx=-1, depth=0, wireType=VARINT, type=INT32)
Index 1: b  (parentIdx=-1, depth=0, wireType=LEN,    type=STRUCT)
Index 2: x  (parentIdx=1,  depth=1, wireType=VARINT, type=INT32)
Index 3: y  (parentIdx=1,  depth=1, wireType=LEN,    type=STRING)

Test coverage

107 JUnit tests in ProtobufTest.java + 13 tests in ProtobufSchemaDescriptorTest.java, organized by feature:

Category	Tests	What is covered
Basic scalar types	3	INT32/64, FLOAT32/64, BOOL, STRING end-to-end
Varint & zigzag	9	Max values, zero, over-encoded zero, 10th-byte validation, zigzag min/max/negative
Wire format errors	14	Malformed varint, truncated fields (varint/string/fixed32/fixed64), partial data, wrong wire type
Unknown field skip	4	Skip varint, fixed32, fixed64, length-delimited unknowns
Last-one-wins	2	Duplicate field handling for scalars and strings
Float/double specials	2	NaN, +Inf, -Inf
Schema projection	2	Partial field decode, decode-none
Required fields	12	Present/missing in permissive/failfast, multi-row, nested required, absent parent skip
Default values	13	All scalar types, strings, empty strings, mixed defaults, multi-row
Repeated fields	10	Unpacked int32, packed double, uint32/64, packed-in-nested, packed-in-repeated-message
Nested messages	6	1-level, 3-level deep, repeated-inside-nested, repeated-in-repeated
Enum (as INT32)	8	Valid, zero, unknown, negative, multiple fields, missing
Enum-as-string	18	Valid/unknown/mixed, permissive null propagation, repeated enum, nested repeated enum, sibling field visibility
FAILFAST mode	13	All error types throw, valid data does not throw
Packed edge cases	4	Misaligned packed fixed32/64, large repeated, mixed packed/unpacked
Deep nesting	6	9-level, 10-level, zero-length nested, empty packed, large field numbers
Schema descriptor	13	Repeated+default reject, struct/list default reject, enum metadata, duplicate fields, encoding compat, depth limit, serialization roundtrip
Performance	1	Multi-field batched extraction correctness

Benchmarks

8 NVBench benchmarks in protobuf_decode.cu:

Benchmark	What it stresses
BM_protobuf_flat_scalars	Top-level scalar extraction throughput
BM_protobuf_nested	Nested message recursive decode
BM_protobuf_repeated	Top-level repeated field count/scan/extract
BM_protobuf_wide_repeated_message	Wide repeated struct (many children)
BM_protobuf_repeated_child_lists	Repeated-in-repeated (nested LIST)
BM_protobuf_repeated_child_string_count_scan	Nested repeated string count+scan isolation
BM_protobuf_repeated_child_string_build	Nested repeated string build isolation
BM_protobuf_many_repeated	Many independent repeated fields

---

Review Guide

This PR is large (~12,000 lines total) but has a well-defined layered architecture. This guide provides a recommended reading order, key areas to focus on per file, and a mental model for understanding the code.

Recommended reading order

Read bottom-up from the API surface to the kernel internals:

Order	File	Focus	Time estimate
1	protobuf.hpp	Data structures, API contract, validation logic	15 min
2	ProtobufSchemaDescriptor.java	Java-side schema validation (mirrors C++ validation)	15 min
3	Protobuf.java	Public Java API, PERMISSIVE mode semantics	5 min
4	ProtobufJni.cpp	JNI bridge: array conversion, local ref management	15 min
5	protobuf_common.cuh §1: types	Lines 54–165: field_location, device_nested_field_descriptor, etc.	10 min
6	protobuf_common.cuh §2: device helpers	Lines 167–400: read_varint, skip_field, get_field_data_location, decode_tag, lookup_field	20 min
7	protobuf_common.cuh §3: LocationProviders	Lines ~400–650: TopLevelLocationProvider, NestedLocationProvider, etc. — these abstract how extraction kernels compute byte offsets	15 min
8	protobuf_common.cuh §4: template kernels	extract_varint_kernel, extract_fixed_kernel, extract_lengths_kernel, copy_varlen_data_kernel, batched variants	20 min
9	protobuf_common.cuh §5: template host functions	extract_typed_column, build_repeated_scalar_column, extract_and_build_string_or_bytes_column, validate_enum_and_propagate_rows	20 min
10	protobuf_kernels.cu §1: scan	scan_all_fields_kernel — the core single-pass field scanner	20 min
11	protobuf_kernels.cu §2: count/scan repeated	count_repeated_fields_kernel, scan_all_repeated_occurrences_kernel, shared __device__ helpers	20 min
12	protobuf_kernels.cu §3: nested	scan_nested_message_fields_kernel, scan_repeated_message_children_kernel, compute kernels	20 min
13	protobuf_kernels.cu §4: validation	check_required_fields_kernel, validate_enum_values_kernel, enum-string kernels	10 min
14	protobuf_builders.cu §1: utilities	make_null_column, make_empty_column_safe, make_null_list_column_with_child	10 min
15	protobuf_builders.cu §2: enum-string	make_enum_string_lookup_tables, build_enum_string_column, build_repeated_enum_string_column	15 min
16	protobuf_builders.cu §3: nested struct	build_nested_struct_column — most complex builder, recursive depth handling	25 min
17	protobuf_builders.cu §4: repeated struct	build_repeated_struct_column, build_repeated_child_list_column — repeated-in-repeated	20 min
18	protobuf.cu	Entry point decode_protobuf_to_struct: orchestration, batched scalar extraction, PERMISSIVE null propagation	30 min
19	ProtobufTest.java	Tests — skim by category, deep-read the tricky ones (enum-as-string, nested repeated)	30 min
20	protobuf_decode.cu	Benchmarks — helper encoders and benchmark configurations	15 min

Total estimated review time: ~5-6 hours for a thorough review.

Key review areas by priority

P0: Correctness-critical

1. scan_all_fields_kernel (protobuf_kernels.cu): This is the most important kernel. One thread per row, scans all bytes of each message, records field locations. Key things to verify:

   • Last-one-wins semantics for duplicate fields
   • Correct wire type dispatch and field skipping
   • Bounds checking at every byte read (cur < msg_end)
   • Error flag setting uses atomicCAS (no races)
   • Permissive mode: row_has_invalid_data is set on parse errors so the row can be nullified

2. count_repeated_fields_kernel (protobuf_kernels.cu): Counts repeated field occurrences. Must correctly distinguish packed vs unpacked encoding. Packed detection: wire_type == WT_LEN but expected_wire_type != WT_LEN.

3. build_nested_struct_column (protobuf_builders.cu): Recursive builder for nested messages. Verify:

   • Depth limit enforcement (MAX_NESTED_STRUCT_DECODE_DEPTH)
   • Correct parent-child location derivation
   • is_repeated children inside nested messages get proper LIST wrapping
   • 0-row / empty-child edge cases create correctly typed columns

4. PERMISSIVE mode null propagation (protobuf.cu, lines ~1163-1190): Invalid enum values must nullify the entire struct row AND propagate to all descendants via apply_parent_mask_to_row_aligned_column + propagate_nulls_to_descendants.

5. JNI memory safety (ProtobufJni.cpp): Every GetObjectArrayElement / GetByteArrayElements / GetIntArrayElements must have a matching DeleteLocalRef / ReleaseXxxArrayElements. Verify no leaks in the enum_names triple-nested loop.

P1: Robustness

6. validate_decode_context (protobuf.hpp): Validates all schema invariants before any GPU work. Check completeness — duplicate field numbers under same parent, encoding compatibility, enum metadata non-empty for ENC_ENUM_STRING.

7. ProtobufSchemaDescriptor.validate (Java): Mirrors C++ validation. Defensive copies in constructor, re-validation on deserialization.

8. Varint parsing (protobuf_common.cuh read_varint): 10th byte must only use lowest bit. Truncated/malformed varints must return false.

9. Wire type handling (skip_field, get_wire_type_size): Verify WT_SGROUP uses iterative handling with depth cap of 32 (not recursive). WT_EGROUP is rejected as standalone.

P2: Performance

10. Batched scalar extraction (protobuf.cu, lines ~400-470): Fields are grouped into 12 categories by type and extracted with 2D kernel launches. Verify grouping logic covers all type+encoding combinations.

11. Field lookup tables (build_field_lookup_table, build_index_lookup_table in protobuf_common.cuh): O(1) field_number → index mapping when max field number ≤ FIELD_LOOKUP_TABLE_MAX. Falls back to linear scan otherwise.

12. String two-phase construction: extract_lengths_kernel → make_offsets_child_column → copy_varlen_data_kernel. Verify no off-by-one in offset calculations.

Things to watch for

• Evaluation order: Never call uvector.size() after uvector.release() in the same expression. The code caches sizes before releasing.
• Offsets column size: LIST columns require offsets with exactly num_rows + 1 elements.
• Packed repeated in nested: count_repeated_in_nested_kernel and scan_repeated_in_nested_kernel handle both packed and unpacked within nested message boundaries.
• Empty struct children: When building 0-row struct columns, repeated children must still be wrapped in empty LIST columns to maintain correct schema.
• cudf::logic_error for data errors: The code uses cudf::logic_error for wire-format errors in strict mode. This is semantically imprecise (it conventionally signals API misuse), but functionally correct.

Mapping review to test coverage

If you're reviewing...	Verify these test categories pass
scan_all_fields_kernel	Basic scalar, varint/zigzag, wire format errors, unknown field skip, last-one-wins
count/scan_repeated	Repeated fields, packed edge cases
build_nested_struct_column	Nested messages, deep nesting
validate_enum + enum-string	Enum (INT32), enum-as-string, PERMISSIVE mode
check_required_fields_kernel	Required fields
Default value handling	Default values
JNI bridge	All tests (they all go through JNI)
Schema validation	ProtobufSchemaDescriptorTest

[Copilot]

Pull request overview

This PR adds a GPU-accelerated protocol buffer decoder with intentionally limited features, focusing on simple scalar field types. The implementation provides a JNI interface for decoding binary protobuf messages into cuDF STRUCT columns.

Key changes:

• Implements GPU kernels for decoding protobuf varint, fixed32/64, and length-delimited (string) fields
• Adds JNI bindings between Java and CUDA implementation
• Provides basic test coverage for INT64 and STRING field types

Reviewed changes

Copilot reviewed 6 out of 6 changed files in this pull request and generated 14 comments.

Show a summary per file

File	Description
src/main/java/com/nvidia/spark/rapids/jni/ProtobufSimple.java	Java API providing decodeToStruct() method with parameter validation
src/test/java/com/nvidia/spark/rapids/jni/ProtobufSimpleTest.java	Basic test case covering varint (INT64) and string decoding with missing fields and null messages
src/main/cpp/src/protobuf_simple.hpp	C++ API declaration with documentation of supported types
src/main/cpp/src/protobuf_simple.cu	CUDA implementation with three specialized kernels for varint, fixed-width, and string extraction
src/main/cpp/src/ProtobufSimpleJni.cpp	JNI bridge translating Java arrays to C++ vectors and invoking decode logic
src/main/cpp/CMakeLists.txt	Build configuration adding new source files to compilation targets

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[thirtiseven]

@greptile full review

[greptile-apps]

Greptile Summary

This PR adds a GPU-accelerated protocol buffer decode kernel using CUDA, implementing a multi-pass algorithm (count → scan → extract → build) that supports nested STRUCT, LIST, STRING, scalar types, proto2 closed enum semantics with deferred error reporting, and singular nested message merging via cub::DeviceMemcpy::Batched.

• Multi-pass decode: a unified field_scan_view covers both the count and scan passes; enum recognition is integrated directly into the scan via is_recognized_enum_value, and a new validate_message_fragments_kernel handles singular nested-message fragment validation.
• Enum handling: proto2 last-one-wins semantics for repeated occurrences; an enum_error_scope controls whether enum errors propagate to the root or stay local; a deferred_enum_error buffer separates enum validation from required-field validation.
• Schema validation (Java): ProtobufSchemaDescriptor.validateEnumMetadata now enforces that enum metadata only appears on compatible field types and that the declared default value is among the valid enum values.

Confidence Score: 3/5

Largely safe but two P1 logic issues need author confirmation before merge: premature deferred enum error under last-wins semantics, and a potentially dead drop_unknown_repeated_enum_values function.

The PR is a large, well-structured addition with a clear design. Two P1 findings were identified: (1) the deferred enum error is set for every invalid occurrence during the scan pass, meaning an intermediate invalid occurrence that is later overridden by a valid last-wins value still leaves the deferred error flag set, which may cause spurious errors in strict mode; (2) drop_unknown_repeated_enum_values is defined but has no call site in the PR, suggesting it is dead code or a missing call. Two P2 findings cover a stack-size/depth-guard decoupling in skip_group and incomplete fragment validation coverage. The overall architecture is sound and the Java-side schema validation improvements are correct.

Files Needing Attention: src/main/cpp/src/protobuf/protobuf_kernels.cu (deferred enum error timing), src/main/cpp/src/protobuf/protobuf_builders.cu (drop_unknown_repeated_enum_values dead code)

Important Files Changed

Filename	Overview
src/main/cpp/src/protobuf/protobuf_kernels.cu	Core CUDA kernels refactored significantly: scan now uses message_scan_context with lookup_view template, enum recognition integrated via is_recognized_enum_value, new validate_message_fragments_kernel added; two P1 issues found: deferred enum error set prematurely for intermediate invalid occurrences in last-wins semantics, and validate_message_fragments_kernel fragment coverage worth verifying.
src/main/cpp/src/protobuf/protobuf_builders.cu	Builder refactored to delegate to build_protobuf_field_values_column_shared; new build_merged_singular_struct_column uses cub::DeviceMemcpy::Batched; drop_unknown_repeated_enum_values is newly defined but has no apparent call site in the PR, making it dead code.
src/main/cpp/src/protobuf/protobuf_types.cuh	Major type-system additions: field_descriptor gains valid_enum_values/num_valid_enum_values/output_index; new batched_scalar_desc, scalar_decode_options, enum_error_scope, protobuf_value_domain_view, required_field_input_view; device_nested_field_descriptor removed. Changes are well-structured and consistent with new multi-pass design.
src/main/cpp/src/protobuf/protobuf_kernels.cuh	Host-side template helpers updated: scalar_kinds constexpr array and dispatch_scalar_decoder added; make_null_mask_from_valid now clears mask when null_count==0; new launch helpers for validate_message_fragments, compute_virtual_parents_for_nested_repeated, compute_msg_locations_from_occurrences. Design is clean.
src/main/cpp/src/protobuf/protobuf_device_helpers.cuh	skip_group refactored to standalone device noinline function with iterative group tracking; skip_field now takes field_number and max_group_depth; decode_tag returns proto_wire_type enum; set_row_invalid helper added. Stack array sized to compile-time constant while depth guard uses runtime parameter — P2 mismatch.
src/main/cpp/src/protobuf/protobuf_host_helpers.hpp	protobuf_decode_runtime_context extended with deferred_enum_error and fail_on_errors; new singular_message_merge_work, repeated_field_work_bundle structs; make_offsets_column, make_top_row_indices, validate_nonempty_repeated_field_work helpers added. Design is consistent; no issues found.
src/main/cpp/src/protobuf/protobuf.cu	Entry point refactored: d_deferred_enum_error buffer added, field_scan_view unified for count/scan passes, repeated STRUCT support added (previously rejected), batched scalar extraction via scalar_kinds fold expression, singular_message_merge_work added for top-level.
src/main/java/com/nvidia/spark/rapids/jni/ProtobufSchemaDescriptor.java	validateEnumMetadata extended to check enum metadata only appears on INT32/DEFAULT or STRING/ENUM_STRING fields and that the default value is in enumValidValues. Validation logic is correct and thorough.
src/main/java/com/nvidia/spark/rapids/jni/Protobuf.java	Minor javadoc updates only — clarifies permissive mode behavior. No functional changes.

Sequence Diagram

sequenceDiagram
    participant Java as Protobuf.java (JNI)
    participant Entry as protobuf.cu
    participant Scan as scan_all_field_occurrences_kernel
    participant Val as validate_message_fragments_kernel
    participant Scalar as extract_scalar_batched_kernel
    participant Build as protobuf_builders.cu

    Java->>Entry: decode(data, schema, options)
    Entry->>Entry: allocate deferred_enum_error, field_scan_view
    Entry->>Scan: count pass - count occurrences per field per row
    Scan-->>Entry: h_occurrence_counts
    Entry->>Entry: prefix-sum to offsets, allocate occurrence buffers
    Entry->>Scan: scan pass - fill field_scan_view with locations
    Scan-->>Entry: field locations, wire type errors, deferred enum errors
    Entry->>Entry: check h_error and h_deferred_enum_error
    Entry->>Scalar: extract_scalar_batched_kernel for all scalar_kinds
    Scalar-->>Entry: typed scalar buffers
    Entry->>Val: validate_message_fragments_kernel for singular nested messages
    Val-->>Entry: merged fragment validation results
    Entry->>Build: build_protobuf_field_values_column_shared
    Build->>Build: build_nested_struct_column recursive for STRUCT fields
    Build->>Build: build_merged_singular_struct_column via cub DeviceMemcpy Batched
    Build-->>Entry: cudf column result
    Entry-->>Java: column handle

Loading

Comments Outside Diff (1)

1. src/main/cpp/src/protobuf/protobuf_kernels.cu, line 227-270 (link)

   Deferred enum error set for any invalid occurrence even if a later valid one wins

   In record_singular, when a duplicate scalar field appears multiple times and the first (or any earlier) occurrence has an unrecognized enum value, set_error_once(fields.deferred_enum_error, INVALID_ENUM) is called and the flag remains set even if a subsequent recognized occurrence overwrites field_locations[f]. In strict mode this causes an INVALID_ENUM exception to be thrown despite the final field value being valid.

   Concretely: a message with field 1 encoded twice as [unknown_val, known_val] should decode to known_val (proto2 last-one-wins), but in strict mode the deferred flag from the first occurrence causes h_error = INVALID_ENUM and an exception is thrown. The fix is to only set the deferred error if the location is ultimately unrecognized (i.e., after scanning completes), rather than eagerly flagging every intermediate unrecognized occurrence.

_{Reviews (8): Last reviewed commit: "fix" | Re-trigger Greptile}

[greptile-apps]

Additional Comments (6)

1. src/main/cpp/src/protobuf_simple.cu, line 89-91 (link)

   logic: potential overflow: len64 can be up to 2^64-1, but casting to int on line 90 can overflow if len64 > INT_MAX

2. src/main/cpp/src/protobuf_simple.cu, line 323-324 (link)

   logic: potential overflow: len64 can be larger than INT_MAX, but casting to int on line 324 will overflow

3. src/main/cpp/src/protobuf_simple.cu, line 375-376 (link)

   logic: race condition: multiple threads write to *error_flag without atomics, causing undefined behavior when multiple threads encounter errors simultaneously

   Then in kernels, use atomicOr(error_flag, 1) instead of *error_flag = 1

4. src/main/cpp/src/protobuf_simple.cu, line 398-407 (link)

   logic: protobuf uses zigzag encoding for signed integers (sint32/sint64), but varint decoding here treats them as unsigned - decoding negative values will produce incorrect results. Are you only supporting unsigned int32/int64, or should zigzag decoding be implemented for signed types?

5. src/main/cpp/src/protobuf_simple.cu, line 240 (link)

   syntax: type punning through reinterpret_cast of incompatible pointer types is undefined behavior in C++

6. src/main/cpp/src/protobuf_simple.cu, line 248 (link)

   syntax: type punning through reinterpret_cast of incompatible pointer types is undefined behavior in C++

_{6 files reviewed, 6 comments}

_{Edit Code Review Agent Settings | Greptile}

[thirtiseven]

@greptile full review

[greptile-apps]

Additional Comments (2)

1. src/main/cpp/src/protobuf_simple.cu, line 186 (link)

   logic: zigzag decode uses signed right shift which is correct for signed types, but applied to unsigned v

2. src/main/cpp/src/ProtobufSimpleJni.cpp, line 55 (link)

   logic: encodings is constructed from n_type_scales but then out_types is also constructed using n_type_scales[i] as the scale parameter, which would be wrong for non-decimal types where this represents encoding

_{6 files reviewed, 2 comments}

_{Edit Code Review Agent Settings | Greptile}

[Copilot]

Pull request overview

Copilot reviewed 6 out of 6 changed files in this pull request and generated 13 comments.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[Copilot]

Pull request overview

Copilot reviewed 6 out of 6 changed files in this pull request and generated 10 comments.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[thirtiseven]

@greptileai full review

[greptile-apps]

_{4 files reviewed, 4 comments}

_{Edit Code Review Agent Settings | Greptile}

[nvauto]

NOTE: release/26.02 has been created from main. Please retarget your PR to release/26.02 if it should be included in the release.

[nvauto]

NOTE: release/26.08 has been created from main. Please retarget your PR to release/26.08 if it should be included in the release.