Fix parquet preprocess strings page bounds check for pages with zero rows - #23241
Conversation
pmattione-nvidia
added
! - Hotfix
|
Related to NVIDIA/cudf-spark#15235. |
|
Running the failing job with the fix. |
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (4)
🚧 Files skipped from review as they are similar to previous changes (3)
📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesList string offset preprocessing
Estimated code review effort: 3 (Moderate) | ~25 minutes Suggested labels: Suggested reviewers: 🚥 Pre-merge checks | ✅ 5✅ Passed checks (5 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
🧹 Nitpick comments (1)
cpp/src/io/parquet/reader_impl_preprocess.cu (1)
380-401: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winClarify that
subpass_start_row/subpass_end_roware actually pass-level bounds.At the call site (
compute_page_string_offset_indices(pass.skip_rows, pass.num_rows)), these values are the pass's bounds, not the subpass's —subpass.skip_rows/subpass.num_rowsaren't computed until later inpreprocess_subpass_pages. Naming themsubpass_start_row/subpass_end_rowhere is misleading given how subtle this exact pass-vs-subpass boundary reasoning is (it's the root of the bug being fixed). A rename (e.g.read_start_row/read_end_row) or a short comment noting they're the full pass range (deliberately over-inclusive for later trimming) would reduce the odds of a future regression here.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cpp/src/io/parquet/reader_impl_preprocess.cu` around lines 380 - 401, Rename subpass_start_row and subpass_end_row in the surrounding page-processing logic to pass-level names such as read_start_row and read_end_row, or add a concise comment identifying them as the full pass bounds used deliberately before subpass trimming. Update all references in process_page while preserving the existing boundary checks.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Nitpick comments:
In `@cpp/src/io/parquet/reader_impl_preprocess.cu`:
- Around line 380-401: Rename subpass_start_row and subpass_end_row in the
surrounding page-processing logic to pass-level names such as read_start_row and
read_end_row, or add a concise comment identifying them as the full pass bounds
used deliberately before subpass trimming. Update all references in process_page
while preserving the existing boundary checks.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 1422c434-52ca-4dab-a366-06141c8f4e26
📒 Files selected for processing (2)
cpp/src/io/parquet/reader_impl_preprocess.cucpp/tests/io/parquet_chunked_reader_test.cu
| bool const has_rows = (page.num_rows > 0) && (page_start_row < subpass_end_row) && | ||
| (page_end_row > subpass_start_row); | ||
| if (has_rows || !is_list_col) { return has_rows; } | ||
|
|
||
| // A single list row can span pages, so a list page can carry values (and offsets) with 0 | ||
| // rows; such a trailing page's page_start_row is already incremented past its spanning last | ||
| // row, so match the decode path's inclusive bounds/contained tests to keep it. | ||
| bool const is_bounds_page = | ||
| (page_start_row <= subpass_start_row && page_end_row >= subpass_start_row) || | ||
| (page_start_row <= subpass_end_row && page_end_row >= subpass_end_row); | ||
| bool const is_contained = | ||
| (page_start_row >= subpass_start_row && page_end_row <= subpass_end_row); | ||
| return is_bounds_page || is_contained; |
There was a problem hiding this comment.
This piece seems to have been extracted from setup_local_page_info(). Is it possible to move it to a common function and use it at both places so they don't go out of sync?
There was a problem hiding this comment.
refactored
There was a problem hiding this comment.
Actionable comments posted: 1
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@cpp/src/io/parquet/decode_preprocess.cu`:
- Around line 446-480: Initialize the shared copy_barrier exactly once before
the first staged-copy decoder pass, rather than calling init() again for the
definition pass. Update the repetition/definition flow around
decoders[level_type::REPETITION].init and decoders[level_type::DEFINITION].init
so both reuse the initialized barrier through arrive_and_wait(), preserving
synchronization between phases.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 0ab7872e-f737-4933-aae9-d3fdb1131958
📒 Files selected for processing (4)
cpp/src/io/parquet/decode_preprocess.cucpp/src/io/parquet/page_string_decode.cucpp/src/io/parquet/reader_impl_preprocess.cucpp/tests/io/parquet_chunked_reader_test.cu
🚧 Files skipped from review as they are similar to previous changes (3)
- cpp/tests/io/parquet_chunked_reader_test.cu
- cpp/src/io/parquet/reader_impl_preprocess.cu
- cpp/src/io/parquet/page_string_decode.cu
There was a problem hiding this comment.
Caution
Inline review comments failed to post. This is likely due to GitHub's internal server error or limits when posting large numbers of comments. If you are seeing this consistently it is likely a permissions issue. Please check "Moderation" -> "Code review limits" under your organization settings.
Actionable comments posted: 1
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@cpp/src/io/parquet/decode_preprocess.cu`:
- Around line 446-480: Initialize the shared copy_barrier exactly once before
the first staged-copy decoder pass, rather than calling init() again for the
definition pass. Update the repetition/definition flow around
decoders[level_type::REPETITION].init and decoders[level_type::DEFINITION].init
so both reuse the initialized barrier through arrive_and_wait(), preserving
synchronization between phases.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 0ab7872e-f737-4933-aae9-d3fdb1131958
📒 Files selected for processing (4)
cpp/src/io/parquet/decode_preprocess.cucpp/src/io/parquet/page_string_decode.cucpp/src/io/parquet/reader_impl_preprocess.cucpp/tests/io/parquet_chunked_reader_test.cu
🚧 Files skipped from review as they are similar to previous changes (3)
- cpp/tests/io/parquet_chunked_reader_test.cu
- cpp/src/io/parquet/reader_impl_preprocess.cu
- cpp/src/io/parquet/page_string_decode.cu
🛑 Comments failed to post (1)
cpp/src/io/parquet/decode_preprocess.cu (1)
446-480: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/usr/bin/env bash set -euo pipefail rg -n -C 8 'cuda::barrier|copy_barrier|arrive_and_wait|init\(' cpp/src/io/parquet rg -n -C 8 'class barrier|struct barrier|void init' . --glob '*barrier*'Repository: rapidsai/cudf
Length of output: 28378
🏁 Script executed:
#!/usr/bin/env bash set -euo pipefail sed -n '200,260p' cpp/src/io/parquet/rle_stream.cuh printf '\n----\n' rg -n -C 3 'init\(©_barrier|arrive_and_wait\(\)' cpp/src/io/parquet printf '\n----\n' rg -n -C 3 'cuda::barrier<cuda::thread_scope_block>|init\(&.*barrier|arrive_and_wait\(\)' cpp/src -g '*.{cu,cuh,cpp,h,hpp}'Repository: rapidsai/cudf
Length of output: 9988
🌐 Web query:
CUDA cuda::barrier reinitialize init arrive_and_wait semantics shared barrier phase reuse💡 Result:
In the CUDA programming model, the
cuda::barrierobject is a reusable synchronization mechanism that manages thread coordination across a sequence of phases [1][2][3]. Initialization and Reinitialization Before any thread participates in a barrier, it must be explicitly initialized using thecuda::barrier::init()friend function [4][3]. Forsharedbarriers, this initialization must happen before any participating thread arrives at the barrier [3][5]. Because initialization itself is a synchronization-dependent task, threads typically use a collective operation likeblock.sync()orsyncthreads()to bootstrap the barrier before its first use [3][5]. There is no distinctreinitializefunction forcuda::barrier. Instead, the barrier is designed to be self-resetting [3]. When the expected arrival count (the number ofarrive()calls) reaches zero, the barrier automatically and atomically resets its internal counter to the expected arrival count (adjusted by anyarrive_and_dropcalls) and increments the phase [2][3]. Semantics and Phase Reuse Thecuda::barrierlifetime consists of a sequence of phases [2]. This phase-based design is what enables reusability [2][6]. - arrive: This function decrements the barrier's expected arrival count [2][6]. It returns anarrival_token, which represents the current phase and must be passed towait()[2][7]. - arrive_and_wait: This is a convenience function that performs the arrival and immediately waits on the same phase [8]. - Phase Completion: Once all expected threads have calledarrive(orarrive_and_drop), the barrier completes the current phase, executes any associated completion function, and transitions to the next phase [2][3][6]. - Thread Wait: A thread that callsarrivecan defer waiting on the phase synchronization point [2][6]. This allows threads to perform independent work between signaling their arrival and blocking for the phase to complete, which is a key mechanism for hiding synchronization latency [2][9]. Crucially, the barrier is only ready to be reused for the next phase once the current phase's completion step has finished [3]. If a thread'sarrive()call triggers the countdown to zero, it must ensure it does not attempt to reuse the barrier for anotherarrive()call until the phase has transitioned [3]. In a correctly synchronized program, the automatic phase increment ensures that subsequentarrive()calls are correctly associated with the next phase [2][3].Citations:
- 1: https://nvidia.github.io/cccl/unstable/libcudacxx/extended_api/synchronization_primitives/barrier.html
- 2: https://gevtushenko.github.io/cccl/libcudacxx/extended_api/synchronization_primitives/barrier.html
- 3: https://docs.nvidia.com/cuda/archive/13.2.1/cuda-programming-guide/04-special-topics/async-barriers.html
- 4: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/async-barriers.html
- 5: https://docs.nvidia.com/cuda/archive/13.2.0/cuda-programming-guide/04-special-topics/async-barriers.html
- 6: https://cppreference.dev/w/cpp/thread/barrier
- 7: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/cooperative-groups.html
- 8: https://github.com/NVIDIA/cccl/blob/main/libcudacxx/include/cuda/std/__barrier/barrier.h
- 9: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/device-callable-apis.html
🌐 Web query:
NVIDIA libcudacxx cuda::barrier init shared barrier may be called again same object after phase completion💡 Result:
In the libcudacxx cuda::barrier API, you should not call the init function multiple times on the same object after phase completion [1][2][3]. According to official NVIDIA documentation, the barrier must be initialized using the cuda::barrier::init friend function before any thread participates [1][2]. Once the barrier is in use, it is designed to manage phases automatically [2]. When the last call to arrive occurs (completing the expected count), the barrier automatically and atomically resets its countdown for the next phase [2]. Re-initializing a barrier that is already in use or has completed a phase can lead to undefined behavior, as the internal state—including the phase bit and arrival count—is managed by the arrival and wait operations to ensure continuous, multi-phase synchronization [4][2]. If you need to reset the barrier or change its expected arrival count, you should not re-initialize the object itself; instead, you should manage the barrier lifecycle according to the specific scope requirements (such as thread_scope_block or thread_scope_device) [2][5]. For example, if a barrier is intended to be used across different, non-concurrent thread groups, ensure that any necessary memory synchronization or reuse logic is handled at the application level without attempting to re-run the init constructor on the same memory location while other threads might still be observing it [5].
Citations:
- 1: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/async-barriers.html
- 2: https://docs.nvidia.com/cuda/archive/13.2.1/cuda-programming-guide/04-special-topics/async-barriers.html
- 3: https://docs.nvidia.com/cuda/archive/13.1.0/cuda-programming-guide/04-special-topics/async-barriers.html
- 4: https://github.com/NVIDIA/cccl/blob/main/libcudacxx/include/cuda/std/__barrier/barrier.h
- 5: NVIDIA/libcudacxx#325
🌐 Web query:
CUDA barrier init function semantics reused barrier phases docs💡 Result:
In CUDA, a
cuda::barrierobject requires explicit initialization before any thread can participate in it [1][2]. The initialization must occur using thecuda::barrier::init()friend function, which defines the expected arrival count (the number of threads or operations that must arrive before the barrier completes its current phase) [1][2][3]. The lifecycle of acuda::barrierconsists of a sequence of barrier phases [4][5]. The semantics regarding initialization, phase advancement, and reuse are as follows: 1. Initialization: Theinit()function must be called before any thread arrives at the barrier [1][2]. Because this is a collective requirement, threads typically use a prior synchronization mechanism (such assyncthreads()or a cooperative group sync) to bootstrap the initialization [2]. 2. Phase Advancement: An asynchronous barrier maintains a countdown starting from the expected arrival count. As participating threads callarrive(), the countdown decrements [3][6]. When the countdown reaches zero, the barrier automatically and atomically resets the countdown to the original expected arrival count and advances to the next phase [3][6]. 3. Reuse and Sequencing: - Tokens: When a thread callsarrive(), it receives anarrival_tokenassociated with the current phase [3][6]. - Blocking/Waiting: A call towait(std::move(token))blocks the thread if the barrier is still in the phase associated with that token [3][6]. If the phase has already advanced, the thread does not block [3][6]. - Reuse Constraint: If a thread's call toarrive()completes the countdown (reaching zero), the thread must callwait()before the barrier can be reused for a subsequentarrive()call [3][6]. - Token Validity:wait()must only be called using a token from the current phase or the immediately preceding phase; using tokens from other phases results in undefined behavior [3][6]. Forcuda::barrier<cuda::thread_scope_block>objects residing in shared memory, a specific implementation constraint exists: they should only be used by the threads in the thread block (CTA) that constructed them [7]. Other barrier scopes or memory locations may allow for reuse across different thread groups, provided no concurrent accesses occur between the different groups [7].Citations:
- 1: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/async-barriers.html
- 2: https://docs.nvidia.com/cuda/archive/13.1.0/cuda-programming-guide/04-special-topics/async-barriers.html
- 3: https://docs.nvidia.com/cuda/archive/13.2.1/cuda-programming-guide/04-special-topics/async-barriers.html
- 4: https://nvidia.github.io/cccl/unstable/libcudacxx/extended_api/synchronization_primitives/barrier.html
- 5: https://gevtushenko.github.io/cccl/libcudacxx/extended_api/synchronization_primitives/barrier.html
- 6: https://docs.nvidia.com/cuda/archive/13.1.1/cuda-programming-guide/04-special-topics/async-barriers.html
- 7: NVIDIA/libcudacxx#325
Initialize
copy_barrieronce
cuda::barrieris reusable across phases; callinginit()again on the same shared object is not the intended lifecycle and can trigger undefined behavior. Bootstrap it once before the first decoder that uses staged copies, then reuse the same barrier witharrive_and_wait()for the repetition and definition passes.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cpp/src/io/parquet/decode_preprocess.cu` around lines 446 - 480, Initialize the shared copy_barrier exactly once before the first staged-copy decoder pass, rather than calling init() again for the definition pass. Update the repetition/definition flow around decoders[level_type::REPETITION].init and decoders[level_type::DEFINITION].init so both reuse the initialized barrier through arrive_and_wait(), preserving synchronization between phases.
This has nothing to do with the current PR. This can be evaluated/handled elsewhere at an elsetime. |
jihoonson
left a comment
jihoonson
left a comment
There was a problem hiding this comment.
LGTM. The Spark job previously failing succeeds with this patch.
|
/merge |
This adds a patch to release 26.06 for applying [this fix](NVIDIA/cudf#23241) for a parquet decoding bug in cuDF. --------- Signed-off-by: Paul Mattione <pmattione@nvidia.com>
6 participants
This fixes a bug in the string offset preprocessing for pages with zero rows. This happens when you have a very large list row that spans multiple pages. The bounds checking did not handle this case directly, it does now.
Checklist