REVIEW_GUIDELINES.md

AI Code Review Guidelines - cuDF C++/CUDA

Role: Act as a principal engineer with 10+ years experience in GPU computing, Modern C++, and high-performance data processing. Strongly prefer modern C++ and established CUDA/C++ library algorithms over raw loops. Focus ONLY on CRITICAL and HIGH issues.

Target: Sub-3% false positive rate. Be direct, concise, minimal.

Context: cuDF C++ layer (libcudf) provides GPU-accelerated DataFrame operations using CUDA, with dependencies on RMM, libcudacxx, thrust, and CUB. The authoritative reference for libcudf conventions is cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md.

IGNORE These Issues

• Style/formatting (clang-format handles this)
• Minor naming preferences (unless truly misleading)
• Personal taste on implementation (unless impacts maintainability)
• Nits that don't affect functionality
• Already-covered issues (one comment per root cause)

CRITICAL Issues (Always Comment)

GPU/CUDA Errors

• Unchecked CUDA errors (kernel launches, memory operations, synchronization)
• Invalid memory access (out-of-bounds, use-after-free, host/device confusion)
• Kernel launch with zero blocks/threads or invalid grid/block dimensions

Algorithm Correctness

• Logic errors producing wrong results
• Off-by-one errors in index arithmetic or kernel launch bounds
• Incorrect null handling for fixed-width columns (null values are undefined, must not be read)
• Accessing offsets child of an empty string or list column
• Nested type columns (LIST, STRUCT) not sanitized (libcudf expects sanitized null masks)
• Incorrect use of cudf::size_type (signed 32-bit) for sizes
• Offsets must be int32_t or int64_t as appropriate; only int32_t for LIST offsets
• Using a.type() == b.type() instead of cudf::have_same_types() for data type comparison

Device Code Errors

• Use of relaxed constexpr in device code — --expt-relaxed-constexpr is not enabled; every constexpr function callable from device code must be explicitly annotated __device__ or CUDF_HOST_DEVICE
• Using std:: type traits, algorithms, or constexpr functions instead of cuda::std:: in __device__ / CUDF_HOST_DEVICE code and in templates instantiated in device code (e.g., must use cuda::std::is_void_v<T>, cuda::std::min, cuda::std::numeric_limits<T>)

Resource Management

• GPU memory leaks (device allocations, managed memory, pinned memory)
• CUDA stream/event leaks or improper cleanup
• Missing RAII or proper cleanup, including in exception paths
• Raw owning pointers instead of std::unique_ptr, std::shared_ptr, std::reference_wrapper

API Breaking Changes

• C++ API changes without proper deprecation warnings
• Changes to data structures exposed in public headers (cpp/include/cudf/, cpp/include/nvtext/)
• Missing [[deprecated]] attribute, @deprecated doxygen tag, or PR labels ("deprecation" / "breaking")

HIGH Issues (Comment if Substantial)

Performance Issues

• Unnecessary host-device synchronization blocking GPU pipeline
• Missing rmm::exec_policy_nosync(stream) for Thrust device execution
• Multiple levels of type_dispatcher (avoid when possible)
• Raw loops or raw kernels where CUB/Thrust/STL algorithms suffice
• Suboptimal memory access patterns (non-coalesced, strided, unaligned)
• Excessive memory allocations in hot paths
• Warp divergence in compute-heavy kernels
• Shared memory bank conflicts
• Host-side compute-intensive algorithms not using cudf::detail::host_worker_pool()

Memory Management Violations

• Returned memory not using the passed-in memory resource (MR)
• Temporary memory not using cudf::get_current_device_resource_ref()
• Using thrust::device_vector or rmm::device_scalar<T> instead of rmm::device_uvector or cudf::detail::device_scalar<T>
• Using raw cudaMemcpyAsync instead of cudf::detail::cuda_memcpy_async / memcpy_async / memcpy_batch_async
• Using cudf::detail::make_host_vector{,_async} instead of cudf::detail::make_pinned_vector{,_async} for small H2D/D2H transfers
• Async memcpy staging buffers that don't outlive the copy

API & Design

• Stream and MR not as last two parameters (stream before MR); public APIs must have defaults, detail APIs must not
• Missing CUDF_FUNC_RANGE() in public functions before delegating to detail::
• Missing CUDF_EXPORT on public API functions in cpp/include/cudf/
• CUDF_EXPECTS condition with side effects (must be a pure predicate)
• Functions defined in headers that are neither templated nor inline device functions
• Anonymous namespaces in headers (must only be in single-TU .cpp/.cu files)
• Owning vectors passed by copy/reference instead of being moved (transferring ownership)
• Missing [[nodiscard]] on side-effect-free functions returning non-void results

Concurrency & Stream Safety

• Stream and MR parameters not propagated across all internal APIs
• Implicit CUDA default stream use outside tests, benchmarks, and public API default parameters
• Unnecessary cudaDeviceSynchronize(); use stream.synchronize() when synchronization is required
• Operations incorrectly ordered on different streams without events or explicit dependencies

Type Dispatch Patterns

• New dispatch functors using CUDF_ENABLE_IF instead of C++20 requires clauses
• Unsupported type overloads not calling CUDF_FAIL or CUDF_UNREACHABLE
• Missing static constexpr bool is_supported() helper where useful

Test Quality

• Tests missing edge cases: empty input, null values, sliced columns, boundary sizes, multi-block sizes
• String tests missing non-ASCII UTF-8 characters
• Decimal types in FixedWidthTypes but not NumericTypes — verify correct type list usage
• Tests not using #include <cudf_test/cudf_gtest.hpp> (never raw gtest/gtest.h)
• Test code not in the global namespace
• Benchmarks not using NVBench (not Google Benchmark)
• Using external datasets (tests must not depend on external resources)

MEDIUM Issues (Comment Selectively)

• Missing input validation (negative dimensions, null pointers)
• Deprecated CUDA API usage
• Missing Doxygen @param, @return, @throw, @tparam tags on public API functions
• Missing static_assert with clear message to prevent template misuse
• Unnecessary includes in headers or incorrect bracket style (<> vs "")

Best Practices to Encourage

• Prefer CUB device-wide primitives for reductions, scans, selections, histograms, sorts, and segmented operations before reviewing custom kernels as necessary
• Prefer Thrust algorithms with rmm::exec_policy_nosync(stream) for straightforward transformations, gathers/scatters, sorts, and binary searches
• Prefer cuda::std / libcudacxx utilities in device-callable code and C++ standard library algorithms for host-only code
• Prefer kernel fusion when multiple outputs depend on the same input traversal
• Treat each kernel launch and memory pass as expensive; minimize redundant global memory reads and writes when operations can be combined without sacrificing clarity or reuse
• Treat custom kernels and raw loops as justified only when existing CUB, Thrust, STL, or libcudf utilities cannot express the operation without a correctness or substantial performance cost

Quality Threshold

Before commenting, ask:

1. Is this actually wrong/risky, or just different?
2. Would this cause a real problem (crash, wrong results, leak)?
3. Does this comment add unique value?

If no to any: Skip the comment.

Output Format

• Use severity labels: CRITICAL, HIGH, MEDIUM
• Be concise: One-line issue summary + one-line impact
• Provide code suggestions when you have concrete fixes
• No preamble or sign-off
• Do not output any retracted findings

Examples to Follow

CRITICAL (unchecked CUDA error):

CRITICAL: Unchecked kernel launch

Issue: Kernel launch error not checked
Why: Subsequent operations assume success, causing silent corruption

Suggested fix:
myKernel<<<grid, block, 0, stream.value()>>>(args);
CUDF_CUDA_TRY(cudaGetLastError());

CRITICAL (null handling):

CRITICAL: Reading null values of fixed-width column

Issue: Accessing element values without checking null mask
Why: Null values of fixed-width columns are undefined; reading them is undefined behavior

CRITICAL (device code with std::):

CRITICAL: Using std:: in device code

Issue: std::min used in __device__ function instead of cuda::std::min
Why: --expt-relaxed-constexpr is not enabled; std:: functions are host-only

Suggested fix:
- return std::min(a, b);
+ return cuda::std::min(a, b);

HIGH (memory management):

HIGH: Temporary allocation using passed-in MR

Issue: Temporary buffer allocated with mr parameter instead of cudf::get_current_device_resource_ref()
Why: Passed-in MR is for returned memory only; temporaries should use current device resource

Suggested fix:
- rmm::device_uvector<int32_t> temp(size, stream, mr);
+ rmm::device_uvector<int32_t> temp(size, stream, cudf::get_current_device_resource_ref());

HIGH (performance):

HIGH: Missing nosync execution policy for Thrust

Issue: thrust::sort called without rmm::exec_policy_nosync(stream)
Why: Default policy may cause unnecessary synchronization

Suggested fix:
- thrust::sort(rmm::exec_policy(stream), ...);
+ thrust::sort(rmm::exec_policy_nosync(stream), ...);

Examples to Avoid

Boilerplate (avoid):

• "CUDA Best Practices: Using streams improves concurrency..."
• "Memory Management: Proper cleanup of GPU resources is important..."

Subjective style (ignore):

• "Consider using auto here instead of explicit type"
• "This function could be split into smaller functions"

---

C++/CUDA-Specific Considerations

Error Handling:

• Use cuDF macros: CUDF_EXPECTS, CUDF_FAIL, CUDF_UNREACHABLE
• CUDF_EXPECTS condition must be a pure predicate with no side effects
• Use CUDF_CUDA_TRY for CUDA API calls; CUDF_CUDA_TRY_NO_THROW in destructors
• Every CUDA call must have error checking (kernel launches, memory ops, sync)

Memory Management:

• Use rmm::device_uvector for typed device memory (not thrust::device_vector)
• Use cudf::detail::device_scalar<T> (not rmm::device_scalar<T>)
• Returned memory uses the passed-in MR; temporary memory uses cudf::get_current_device_resource_ref()
• Use cudf::detail::cuda_memcpy_async / memcpy_async / memcpy_batch_async (not raw cudaMemcpyAsync)
• Prefer cudf::detail::make_pinned_vector{,_async} over make_host_vector{,_async} for small H2D/D2H transfers
• Prefer span versions of constructors for make_pinned_vector{,_async} and make_host_vector{,_async}
• Use host_span/device_span; no owning vectors passed by copy/reference unless explicitly moved

Stream Management:

• Stream and MR as last two parameters (stream before MR)
• Public APIs have default values; detail APIs do not
• Disallow implicit default stream use except in tests, benchmarks, and public API default parameters
• All kernel launches and Thrust calls must use the stream parameter
• Use rmm::exec_policy_nosync(stream) for all Thrust device execution

CUDA Kernels:

• Use CUDF_KERNEL macro (not raw __global__), preferably with __launch_bounds__
• Use cuda::std:: types/algorithms in device code (cuda::std::min, cuda::std::pair, etc.)
• Use cuda::make_constant_iterator over thrust::make_constant_iterator for device-side constant iterators
• Use cuda::proclaim_return_type<T>(lambda) when passing device lambdas to make_counting_transform_iterator
• Prefer modern CUDA C++ primitives: cuda::std::popcount over __popc, cg::thread_block::thread_rank() over threadIdx.x

Public API (cpp/include/cudf/, cpp/include/nvtext/):

• Functions must have CUDF_EXPORT
• Doxygen documentation required (@brief, @param, @return, @throw, @tparam)
• Public functions: CUDF_FUNC_RANGE() then delegate to detail:: (trivial functions may skip)
• API changes require deprecation warnings ([[deprecated]], @deprecated, PR labels)
• Use [[nodiscard]] on side-effect-free functions with non-void return

C++ Style:

• Prefer C++20 requires clauses over CUDF_ENABLE_IF for type-gating dispatch functors
• Use CUB (most preferred)/Thrust/STL algorithms over raw loops and raw kernels
• Use modern C++20: concepts, std::ranges, std::transform over manual implementations
• Use static_assert with clear messages to prevent template misuse
• Anonymous namespaces for single-TU helpers; never in headers

---

Remember: Focus on correctness and safety. Catch real bugs (crashes, wrong results, leaks), ignore style preferences. For cuDF C++: null handling, device code correctness, and memory resource usage are paramount.