CCCL Review Guidelines#
Guidelines for reviewing CCCL changes, distilled from past regressions. Each guideline states what a human or AI reviewer should flag in a diff.
Format#
Each guideline is a section of the form:
## <area>.<slug> (<severity>, <scope>)
<!-- provenance:
#<introducing PR>→#<fixing PR> <short note>;
review feedback on #<PR> (<link>) <short note>;
...
-->
<3–8 lines of imperative guidance: what to flag in a diff, and under which conditions it is acceptable>
<area>.<slug>— stable rule id;areais one ofbuild,correctness,perf,api,abi,test,infra,docs. New areas and slugs can be added as needed.<severity>—critical(must be addressed),important(not addressing requires a justification), orsuggestion(worth considering, no action required).<scope>— which files/diffs the rule applies to.The provenance comment lists what the rule was distilled from: historical regressions (introducing PR → fixing PR) or review feedback that prevented a defect from shipping (link to the review comment); it is metadata for maintainers, not part of the rule.
Rules are grouped by area, in the order
build,correctness,api,abi,perf,test,infra,docs.
build.compiler-matrix (important, all C++ code)#
Flag compiler-specific constructs unless guarded by a feature macro: __int128, GNU inline
asm, __attribute__, VLAs, visibility attributes, one-compiler warning suppressions, etc.
CCCL must build with GCC, Clang, MSVC, NVHPC/NVC++ (incl. -stdpar), and NVRTC; PR CI does
not cover all of these, so green CI is not sufficient. Compiler-detection refactors must
stay equivalent for every supported compiler — beware masquerading (Clang and NVHPC define
__GNUC__). Benchmarks and tests count too.
build.fp16-implicit-ops (important, C++ code using __half/__nv_bfloat16 or similar extended FP types)#
Flag any use of the vendor-provided conversions or operators on CUDA extended floating-point types
(__half, __half2, __nv_bfloat16, __nv_bfloat162): static_cast<__half>(f),
static_cast<float>(h), arithmetic/comparison operators (h1 + h2, h1 < h2), and compound
assignments. The vendor headers compile these out when a user defines
__CUDA_NO_HALF_CONVERSIONS__/__CUDA_NO_HALF_OPERATORS__ (or the HALF2/BFLOAT16/BFLOAT162
equivalents), so such code silently fails to build for those configurations. Use the explicit
intrinsics instead (__float2half/__half2float, __hadd/__hgt, …). __nv_bfloat16
arithmetic/comparison intrinsics require SM80 while CCCL supports sm75+: gate them with
NV_IF_ELSE_TARGET(NV_PROVIDES_SM_80, …) and a float round-trip fallback via the explicit conversion
intrinsics (__bfloat162float/__float2bfloat16_rn; for __nv_bfloat162,
__bfloat1622float2/__float22bfloat162_rn). Candidate for a pre-commit grep.
build.narrow-arithmetic-then-widen (important, C++/CUDA code computing a size/count/capacity/offset, including test files)#
When a diff computes a size, count, capacity, or offset by adding or multiplying two operands and widening the RESULT afterward — explicitly with a cast or implicitly through a wider destination type, a wider function parameter, or return type — flag it as a code smell: the operation executes in the narrower type and can silently overflow before the result is widened. The author should either widen an operand before the operation or, if the narrow result is intended, narrow the destination type so no widening occurs.
correctness.pdl-restrict-aliasing (critical, CUDA kernels that call _CCCL_PDL_GRID_DEPENDENCY_SYNC() / cudaGridDependencySynchronize())#
Flag a kernel parameter marked both const and _CCCL_RESTRICT (or raw __restrict/__restrict__)
whose pointee is read after a _CCCL_PDL_GRID_DEPENDENCY_SYNC()/cudaGridDependencySynchronize()
call, when that memory is written by the preceding kernel. const + restrict together assert the
memory is read-only and uniquely accessed through this pointer, licensing the compiler to reorder
the load ahead of the sync. Acceptable only if a comment at the parameter explains why the memory can
never be concurrently aliased (e.g. no producer kernel writes it).
correctness.pdl-cc-gate (important, CUB dispatch code in cub/device/dispatch/*.cuh and launcher factories launching kernels with programmatic dependent launch)#
When a kernel is launched with PDL enabled (either directly or via a launcher factory like
TripleChevronFactory or CudaDriverLauncherFactory), flag any argument for
dependent_launch/use_pdl that does not consider the current device’s compute capability queried
via cub::detail::ptx_compute_cap. PDL may only be enabled if
cc >= ::cuda::compute_capability{9, 0} (see dispatch_find.cuh).
correctness.pdl-sync (critical, kernels launched with programmatic dependent launch)#
When a diff enables programmatic dependent launch for a kernel by setting dependent_launch to true
at the kernel launcher, flag any global memory access in the kernel’s body (i.e., a load from or a
store to a pointer passed at the kernel’s interface) that happens before any call to
_CCCL_PDL_GRID_DEPENDENCY_SYNC — the previous kernel may still be writing that memory — unless the
access has a comment explaining why a PDL sync can come later.
correctness.trivially-copyable-trait (important, generic code constraining or branching on trivial copyability)#
Flag cuda::std::is_trivially_copyable(_v) or std::is_trivially_copyable(_v) applied to a generic
value-type parameter;
use cuda::is_trivially_copyable(_v) instead, which supports more cases. The vendor headers give
__half/__nv_bfloat16 non-trivial special members, so the standard trait reports false for them
(and aggregates of them) even though they are functionally copyable. Candidate for a pre-commit grep.
api.internal-symbol-exposure (important, new implementation-detail types/functions)#
Flag any entity added to a public namespace which can be recognized as intended to be internal by its
spelling (e.g. snake_case in CUB, or prefixed with __ in libcu++/cudax) or usage pattern (e.g. used
as utility for other functions, not documented, etc.). The entity should be marked internal as
appropriate.
abi.missing-hide-from-abi (critical, inline functions in public headers whose behavior depends on the build configuration)#
Flag an inline function in a public header whose result can differ between two copies of CCCL linked
into one binary (e.g., built against different CUDA toolkits or CCCL versions) — especially functions
returning kernel or function pointers — unless it is marked _CCCL_HIDE_FROM_ABI (or an attribute macro
that includes it, like _CCCL_HOST_API). With default visibility the linker keeps ONE definition
across all copies, so the losing copy’s callers get the other build’s result (e.g. a wrong kernel
pointer) — no build error, just wrong behavior at run time.
correctness.cuda-driver-symbol-version-guard (critical, code calling CUDA Driver API symbols)#
When a diff gates a call to a CUDA Driver API symbol introduced in a specific CUDA version behind a
build-time-only check (_CCCL_CTK_AT_LEAST(...)), flag it: libcuda.so/nvcuda.dll comes from the
installed display driver, which is independent of — and often older than — the CTK the binary was
built against, so the symbol can be absent at run time regardless of any build-time guard. Resolve
driver entry points through the versioned cuGetProcAddress bootstrap in
cuda/__driver/driver_api.h (which reports availability), or verify cudaDriverGetVersion before
the call. PR CI builds and runs with matched driver/CTK, so this only reproduces in the field.
infra.pin-deps (important, CMake/CI/submodules)#
Flag dependencies fetched by branch name (CPMAddPackage("gh:org/repo#main"), GIT_TAG main); pin a commit or tag. Candidate for a pre-commit grep.
correctness.stale-refs-after-rename (important, renames, moves, or splits of files, symbols, or modules anywhere in the repo)#
When a diff renames, moves, or splits a file, macro, symbol, or module, git grep for the old name:
each remaining hit must be updated, or be classified as an unrelated entity that merely shares
the name. Doc directives (automodule::/toctree::) can go stale without containing the old name and
still build cleanly (autodoc renders emptied packages as blank pages) — verify the rendered docs, not
just the grep.
correctness.workaround-breaks-constexpr (important, constexpr-marked)#
When a diff introduces any workaround inside a constexpr function, verify the function is still
usable during constant evaluation on every supported compiler, not merely that it compiles as a
runtime call. The break only surfaces when a caller uses the function during constant evaluation,
which the unit tests may not exercise. Adding a test that evaluates the function at compile time is
recommended.
perf.partial-pdl (important, kernels launched with programmatic dependent launch)#
When a diff enables programmatic dependent launch for a kernel by setting dependent_launch to true
at the kernel launcher, open the full dispatch function (or equivalent) and enumerate EVERY kernel
launch it makes. All kernels should be launched using PDL, otherwise the performance gain is marginal.
Replacing calls to cudaMemset by kernels launched using PDL should be strongly considered and
pointed out as suggestions.
api.type-replacement (critical, public types in thrust/libcudacxx/cub)#
When a diff reimplements, re-derives, or aliases any public type (thrust::pair/tuple/complex,
iterators, …), verify every observable property of the old type is preserved: trivial copyability and
layout (downstream code memcpys them), size/alignment, implicit conversions and promotions, overload
resolution, and numerical behavior.
build.windows-min-max-macro (important, C++ code calling .max()/.min() or naming a new member/trait max/min)#
Flag an unparenthesized call to a function literally named max/min (e.g.
std::numeric_limits<T>::max()), and any new member or trait named max/min. On Windows,
<windows.h> defines max/min as function-like macros, breaking such code. Headers sandwiched
between <cuda/std/__cccl/prologue.h>/epilogue.h (libcudacxx, cudax) are safe; everywhere else
(CUB, Thrust, tests, examples), require the macro-safe spelling (std::numeric_limits<T>::max)()
and prefer other member names. Candidate for a pre-commit grep.
perf.benchmark-exec-tag-sync-without-sync-call (important, nvbench benchmark harness state.exec(...) calls)#
When a diff makes an nvbench state.exec(...) call use nvbench::exec_tag::sync (which tells
nvbench that the benchmark region will perform CUDA synchronization itself), or changes the lambda
body of a call already using such a tag, verify the lambda actually performs any explicit CUDA
synchronization (like launch.get_stream().sync(), cudaStreamSynchronize). Parallel algorithms in
Thrust and cuda::std:: synchronize internally, except under thrust::cuda::par_nosync. Without a
sync, the measured time silently excludes some or all of the kernel’s execution. If the lambda does
not sync, exec_tag::no_batch or exec_tag::timer is likely what was intended.
perf.tuning-refactor-verification (important, CUB tuning-policy selectors in cub/device/dispatch/tuning/*.cuh and perf-critical type/arch dispatch)#
A policy selector is a constexpr function from a compute capability and type information to a tuning
policy value. When a diff refactors one (or anything it calls, or the *Policy structs it returns)
without claiming a perf change, the function must return the same policy values for every input as
before — check equivalence per (architecture, type, size, operation) combination, not per code branch.
Typical silent breaks: a combination that used to hit tuned values now falls through to a fallback or
empty optional (no compile error, just an untuned policy); a field added/removed/reordered in a
policy struct shifting every positional brace-init (ScanLookaheadPolicy{6, 104 - 1, 8, 2, 2}); a
dispatch wrapper no longer forwarding its policy selector to an inner dispatch call. A claim of “no
SASS changes” verified on one test type is not sufficient; demand a SASS diff or benchmark sweep over
non-default/non-primitive value types and every affected architecture.
perf.vector-init-before-overwrite (suggestion, thrust::host_vector/thrust::device_vector and c2h’s host_vector/device_vector construction anywhere in the diff)#
Flag a sized thrust/c2h host_vector/device_vector construction (vector(n), vector(n, value))
or resize(n) that doesn’t use the thrust::no_init sentinel when the vector’s content is never read
before being fully overwritten — as an algorithm’s output, a kernel/copy destination, or an explicit
full assignment right after. The default overload value-initializes every element first, work the
overwrite immediately discards. Temporary-storage allocations are almost always a no_init candidate,
since callers assume that memory starts uninitialized. Does not apply when the vector is read before
being overwritten, or when the initial values are semantically required (e.g. an accumulator seed, or a
vector only partially written). When the element type is not trivially default-constructible,
or generic (e.g., in a template context), use thrust::default_init instead.
