Add CUDA error checks after every kernel launch - #22755
Merged
Merged
Conversation
…rement Based on offline discussions, this PR updates the C++ review guideline to remove the kernel error checking requirement after each kernel launch.
PointKernel
added
doc
Member
Author
|
related Slack discussions: https://nvidia.slack.com/archives/C01CW5L51QC/p1780351009811349 |
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (2)
💤 Files with no reviewable changes (1)
🚧 Files skipped from review as they are similar to previous changes (1)
📝 WalkthroughSummary by CodeRabbit
WalkthroughThis PR inserts immediate post-kernel CUDA error checks (CUDF_CUDA_TRY(cudaGetLastError())) after many GPU kernel launches across bitmask, copying, I/O (CSV/Parquet/ORC/Avro), join, string, text, transform, and utility modules. No public APIs or kernel logic were modified. ChangesSystematic CUDA Kernel Launch Error Checking
🎯 3 (Moderate) | ⏱️ ~20 minutes Suggested labels: Suggested reviewers:
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing Touches📝 Generate docstrings
🧪 Generate unit tests (beta)
Comment |
vuule
approved these changes
Jun 2, 2026
bdice
approved these changes
Jun 2, 2026
PointKernel
marked this pull request as draft
June 2, 2026 20:46
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
Member
Author
|
/ok to test 4da31c2 |
PointKernel
changed the title
PointKernel
marked this pull request as ready for review
June 3, 2026 16:12
Member
Author
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 2
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (2)
cpp/src/io/text/multibyte_split.cu (1)
405-412:⚠️ Potential issue | 🔴 Critical | ⚡ Quick winMissing CUDA error check after kernel launch.
The
multibyte_split_init_kernellaunch at lines 405–412 is missing the post-launch error check. For consistency with the PR objective and to match the error checks added after the other kernel launches in this file, addCUDF_CUDA_TRY(cudaGetLastError());immediately after line 412.🔧 Proposed fix
tile_multistates, tile_offsets); + CUDF_CUDA_TRY(cudaGetLastError()); CUDF_CUDA_TRY(cudaStreamWaitEvent(scan_stream.value(), last_launch_event));As per coding guidelines: Check for unchecked CUDA errors in kernel launches, memory operations, and synchronization calls.
🤖 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/text/multibyte_split.cu` around lines 405 - 412, The kernel launch of multibyte_split_init_kernel using tiles_in_launch, THREADS_PER_TILE and scan_stream is missing a post-launch CUDA error check; add a CUDF_CUDA_TRY(cudaGetLastError()); immediately after the multibyte_split_init_kernel<<<...>>>(...) call (the same pattern used after other kernel launches in this file) to catch launch errors and follow the project's error-checking convention.cpp/src/text/wordpiece_tokenize.cu (1)
827-831:⚠️ Potential issue | 🟠 Major | ⚡ Quick winGuard the limited-tokenization launch when no words were found.
total_wordscan be 0 for valid inputs such as non-null empty strings. In that case this constructs a zero-block launch, and the newCUDF_CUDA_TRY(cudaGetLastError())will now fail withcudaErrorInvalidConfiguration.💡 Suggested fix
- cudf::detail::grid_1d grid{total_words, 512}; - tokenize_kernel<decltype(map_ref), decltype(sub_map_ref)> - <<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>( - start_words, word_sizes, d_input_chars, map_ref, sub_map_ref, unk_id, d_tokens.data()); - CUDF_CUDA_TRY(cudaGetLastError()); + if (total_words > 0) { + cudf::detail::grid_1d grid{total_words, 512}; + tokenize_kernel<decltype(map_ref), decltype(sub_map_ref)> + <<<grid.num_blocks, grid.num_threads_per_block, 0, stream.value()>>>( + start_words, word_sizes, d_input_chars, map_ref, sub_map_ref, unk_id, d_tokens.data()); + CUDF_CUDA_TRY(cudaGetLastError()); + }Based on learnings: Verify kernel launches have valid grid/block dimensions (non-zero blocks/threads).
🤖 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/text/wordpiece_tokenize.cu` around lines 827 - 831, total_words can be zero which creates a zero-block kernel launch causing cudaErrorInvalidConfiguration when CUDF_CUDA_TRY(cudaGetLastError()) is called; guard the tokenize_kernel launch by checking total_words > 0 (or grid.num_blocks > 0) before invoking tokenize_kernel<<<...>>>(start_words, word_sizes, d_input_chars, map_ref, sub_map_ref, unk_id, d_tokens.data()) and only call CUDF_CUDA_TRY(cudaGetLastError()) when the kernel was actually launched so no zero-block launch occurs.
🤖 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/comp/unsnap.cu`:
- Around line 719-721: The post-launch cudaGetLastError() checks can throw when
kernels were launched with zero blocks for valid empty inputs; update the launch
sites (e.g., unsnap_kernel_no_racecheck<<<dim_grid, dim_block,...>>>(...) and
the other kernel at lines ~758-760) to first guard the launch and the subsequent
cudaGetLastError() by checking that the grid and block dimensions are non-zero
(or that inputs.size() > 0) and only launch + call cudaGetLastError() when those
conditions hold; mirror the pattern used by gpuinflate/gpu_snap to skip launches
for empty inputs.
In `@cpp/src/io/parquet/decode_preprocess.cu`:
- Around line 485-492: The kernel launch in compute_page_sizes() can be invoked
with dim_grid.x == 0 when pages.size() == 0, causing an invalid configuration
error due to the new cudaGetLastError() check; add an early no-op guard in
compute_page_sizes() (the caller is preprocess_levels()) to return immediately
when there are no pages/chunks to process (e.g., pages.size()==0 or chunks==0 or
page_mask indicates empty) before computing dim_grid/dim_block and before
launching compute_page_sizes_kernel<uint8_t> /
compute_page_sizes_kernel<uint16_t>, so no kernel is launched with zero grid
dimensions.
---
Outside diff comments:
In `@cpp/src/io/text/multibyte_split.cu`:
- Around line 405-412: The kernel launch of multibyte_split_init_kernel using
tiles_in_launch, THREADS_PER_TILE and scan_stream is missing a post-launch CUDA
error check; add a CUDF_CUDA_TRY(cudaGetLastError()); immediately after the
multibyte_split_init_kernel<<<...>>>(...) call (the same pattern used after
other kernel launches in this file) to catch launch errors and follow the
project's error-checking convention.
In `@cpp/src/text/wordpiece_tokenize.cu`:
- Around line 827-831: total_words can be zero which creates a zero-block kernel
launch causing cudaErrorInvalidConfiguration when
CUDF_CUDA_TRY(cudaGetLastError()) is called; guard the tokenize_kernel launch by
checking total_words > 0 (or grid.num_blocks > 0) before invoking
tokenize_kernel<<<...>>>(start_words, word_sizes, d_input_chars, map_ref,
sub_map_ref, unk_id, d_tokens.data()) and only call
CUDF_CUDA_TRY(cudaGetLastError()) when the kernel was actually launched so no
zero-block launch occurs.
🪄 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: a1955893-e695-4bb4-8e95-3f821e6455ec
📒 Files selected for processing (75)
cpp/src/bitmask/null_mask.cucpp/src/copying/concatenate.cucpp/src/copying/contiguous_split.cucpp/src/copying/scatter.cucpp/src/groupby/hash/compute_mapping_indices.cuhcpp/src/groupby/hash/compute_shared_memory_aggs.cucpp/src/io/avro/avro_gpu.cucpp/src/io/comp/debrotli.cucpp/src/io/comp/gpuinflate.cucpp/src/io/comp/snap.cucpp/src/io/comp/unsnap.cucpp/src/io/csv/csv_gpu.cucpp/src/io/fst/dispatch_dfa.cuhcpp/src/io/orc/dict_enc.cucpp/src/io/orc/stats_enc.cucpp/src/io/orc/stripe_data.cucpp/src/io/orc/stripe_enc.cucpp/src/io/orc/stripe_init.cucpp/src/io/orc/writer_impl.cucpp/src/io/parquet/chunk_dict.cucpp/src/io/parquet/decode_fixed.cucpp/src/io/parquet/decode_preprocess.cucpp/src/io/parquet/experimental/dictionary_page_filter.cucpp/src/io/parquet/page_data.cucpp/src/io/parquet/page_delta_decode.cucpp/src/io/parquet/page_enc.cucpp/src/io/parquet/page_hdr.cucpp/src/io/parquet/page_string_decode.cucpp/src/io/statistics/column_statistics.cuhcpp/src/io/text/multibyte_split.cucpp/src/io/utilities/data_casting.cucpp/src/io/utilities/type_inference.cucpp/src/join/conditional_join.cucpp/src/join/filter_join_indices_kernel.cuhcpp/src/join/filtered_join.cucpp/src/join/hash_join/partitioned_count_kernels.cuhcpp/src/join/hash_join/partitioned_retrieve_kernels.cuhcpp/src/join/key_remapping.cucpp/src/join/mark_join.cucpp/src/join/mixed_join_kernel.cuhcpp/src/join/mixed_join_kernels_semi.cucpp/src/join/mixed_join_size_kernel.cuhcpp/src/json/json_path.cucpp/src/merge/merge.cucpp/src/partitioning/partitioning.cucpp/src/quantiles/tdigest/tdigest.cucpp/src/quantiles/tdigest/tdigest_aggregation.cucpp/src/replace/nulls.cucpp/src/replace/replace.cucpp/src/rolling/detail/rolling.cuhcpp/src/sort/segmented_top_k.cucpp/src/strings/attributes.cucpp/src/strings/case.cucpp/src/strings/convert/convert_urls.cucpp/src/strings/copying/concatenate.cucpp/src/strings/like.cucpp/src/strings/regex/utilities.cuhcpp/src/strings/replace/multi.cucpp/src/strings/replace/replace.cucpp/src/strings/search/contains_multiple.cucpp/src/strings/search/find.cucpp/src/strings/search/find_instance.cucpp/src/strings/slice.cucpp/src/strings/split/split.cuhcpp/src/strings/strings_column_factories.cucpp/src/text/bpe/byte_pair_encoding.cucpp/src/text/edit_distance.cucpp/src/text/generate_ngrams.cucpp/src/text/jaccard.cucpp/src/text/minhash.cucpp/src/text/normalize.cucpp/src/text/vocabulary_tokenize.cucpp/src/text/wordpiece_tokenize.cucpp/src/transform/compute_column_kernel.cuhcpp/src/transform/row_bit_count.cu
✅ Files skipped from review due to trivial changes (6)
- cpp/src/io/statistics/column_statistics.cuh
- cpp/src/bitmask/null_mask.cu
- cpp/src/strings/split/split.cuh
- cpp/src/join/mark_join.cu
- cpp/src/io/parquet/decode_fixed.cu
- cpp/src/io/parquet/page_enc.cu
bdice
approved these changes
Jun 3, 2026
PointKernel
added
improvement
mhaseeb123
approved these changes
Jun 4, 2026
# Conflicts: # cpp/src/join/mark_join.cu
Member
Author
|
/merge |
abigalekim
pushed a commit
to abigalekim/cudf
that referenced
this pull request
Jun 12, 2026
This PR updates the cuDF codebase to follow CUDA best practices by consistently checking for errors after kernel launches. It also fixes a place where raw `__global__` was used instead of the dedicated `CUDF_KERNEL` macro. Authors: - Yunsong Wang (https://github.com/PointKernel) Approvers: - Vukasin Milovanovic (https://github.com/vuule) - Bradley Dice (https://github.com/bdice) - Muhammad Haseeb (https://github.com/mhaseeb123) URL: NVIDIA#22755
3 tasks
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
This PR updates the cuDF codebase to follow CUDA best practices by consistently checking for errors after kernel launches. It also fixes a place where raw
__global__was used instead of the dedicatedCUDF_KERNELmacro.Checklist