[STF] cudax sharded: concepts for sharded structures + concept-generic algorithms (pilot) - #11118
Open
caugonnet wants to merge 87 commits into
Open
[STF] cudax sharded: concepts for sharded structures + concept-generic algorithms (pilot)#11118caugonnet wants to merge 87 commits into
caugonnet wants to merge 87 commits into
Conversation
…sources) A grid of places names where things can run; a place_group owns what it takes to execute there: lazily created per-place stream pools and per-place memory resources, with explicit ownership and teardown. Standalone groups own their exec_place_resources registry; a group can also borrow the registry of an STF async_resources_handle so exactly one pool owner exists when both layers coexist. Includes: - place_group with grid/vector constructors and by_devices / by_locality_domains factories - place_memory_resource: a cuda::mr resource over data_place, usable in CUB single-call environments (place_group::env) - free helpers: all_device_ids, places_from_devices, places_from_grid, places_from_locality_domains, make_stream_wait_for - tests: construction/factories, stream pools (lazy creation, colors, isolation between groups), memory resources, STF borrowing, moves - docs: place-group section in places.rst Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…cit link request Doxygen resolves `::std::vector` in a doc comment as an explicit link request and errors when the target is not in the tag files; the plain spelling documents the same thing without asking Doxygen to link it. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ze stream indexing, doc fixes - place_memory_resource::deallocate/deallocate_sync stay noexcept (the cuda::mr resource convention) but the bodies no longer throw through the noexcept boundary: deallocation failures are reported to stderr. - place_group::get_stream indexes modulo the place's ACTUAL cached stream-pool size (custom pool sizes are safe); documented that colors wrap per place, while num_stream_colors() reports the group default. - make_stream_wait_for destroys its event on every path and surfaces the first failure after cleanup. - Move construction documented as requiring exclusive access to the source; direct includes for the __stf utility headers used here; places.rst examples use the public umbrella header and define their inputs. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…eader The __places idiom (review suggestion): the five test functions move from a standalone executable into UNITTEST blocks at the bottom of place_group.cuh, run by the already-registered unittested-header target — one full test TU + link less to compile. The STF async_resources_handle include rides under the UNITTESTED_FILE guard, so the header stays STF-free for normal consumers; the stream-execution check uses cudaMemsetAsync so the header defines no __global__ symbols even in the unittest TU. std::move becomes mv() per review. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…der; drop the duplicate stream helper; polish - place_memory_resource moves to __places/place_memory_resource.cuh: it is a data_place -> cuda::mr adapter with no dependency on place_group (the group only builds its memory_resource(i)/env(i) conveniences on top). Registered as a unittested header; the future delegation to the core per-locality-domain default pools now touches one header. - make_stream_wait_for deleted: cuda::stream_ref::wait(stream_ref) is the same operation, with an RAII event and the driver-API wait that avoids pushing device 0 on an empty context stack. - make_stream_wait_for's interim error handling ported into... (n/a — helper gone); place_group::sync() now synchronizes only streams that exist (lazy pools stay lazy; documented). - Direct std includes; const/constexpr at the flagged sites; [[nodiscard]]/noexcept on the simple getters (the layer-wide _CCCL_HOST_API sweep is deliberately NOT adopted here: the __places layer is STF-origin style throughout — a layer-wide annotation pass is separate work). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
… copy/move doc The __places layer is STF-origin, and the STF convention for internal namespaces is 'reserved' (35 files) rather than 'detail'. place_group is non-copyable and not move-assignable; move-CONSTRUCTIBLE only, so the factories and ownership transfer work — now stated verbatim on the members. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…leave the public surface
Review point: the free helpers (all_device_ids, places_from_devices,
places_from_grid, places_from_locality_domains) both duplicated the
existing place vocabulary (exec_place::device, grids, place_partition)
and introduced a second dialect next to the by_* factories. They move
into the reserved namespace as the factories' implementation detail.
Public construction is now exactly: the ctors (vector of places / grid
/ scalar exec_place — e.g. place_group{exec_place::device(0)}) plus the
two machine-level factories by_devices()/by_locality_domains(). Tests
and docs use the existing vocabulary.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…n failures, deadlock-free sync - place_memory_resource::allocate/allocate_sync validate the requested alignment (at most, and dividing, cuda::mr::default_cuda_malloc_alignment — what every data_place path guarantees) and refuse the rest with std::invalid_argument instead of silently under-aligning. - Deallocation failures are reported to stderr AND debug-assert (the legacy_pinned_memory_resource semantics) rather than reading as success. - place_group::sync() snapshots the cached streams under the mutex and synchronizes after unlocking: a host function enqueued on a cached stream may itself call get_stream(), which would deadlock against cudaStreamSynchronize under the lock. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
is_valid_alignment(0) would divide by zero; zero is not a valid alignment and is now refused up front. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
- get_or_create_streams refuses a place that is not a member of the group with std::invalid_argument (the assert compiled out with NDEBUG and the public get_stream/sync_stream paths then indexed the cache out of range). - place_memory_resource rejects allocation sizes above PTRDIFF_MAX before the size_t -> ptrdiff_t conversion. - place_memory_resource gets the cuda::mr-world annotations (_CCCL_HOST_API, [[nodiscard]]/noexcept getters, const locals) — this header models the annotated libcu++ resource family; place_group.cuh stays STF-layer style. - UNITTEST arms for every refusal: zero/unsupported alignment, oversize, foreign place. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ization
sync_stream(place, color) had no callers and duplicates the existing
spelling: cuda::stream_ref{group.get_stream(place, color)}.sync().
place_group keeps sync() (the whole-group operation, which has no
one-line equivalent).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
A new sibling capability cudax/__sharded (public umbrella
<cuda/experimental/sharded.cuh>, namespace cuda::experimental::sharded),
peer to __multi_gpu: placement-owning containers for the in-process
places rung.
- shard<T>: one placed piece (data/size/capacity/global_offset,
data_place/exec_place/reference stream, global<->local index math).
- sharded_array<T>: allocation from explicit specs or a place_group;
adoption of existing buffers under the two-word rule (adopt() =
zero-copy view, from_* = owned copies); allocate_like; host transfer;
slice with place correspondence; each_shard visitation; copy_between.
- Contiguous backing contract (allocate_contiguous): shards become views
into ONE contiguous VA range whose physical pages are owned per place
(VMM via places::localized_array); exact logical boundaries,
fixed sizes, contiguous_data() for single-pointer consumers.
Tests: containers/{sharded_array,contiguous}.cu + include_only + header
tests.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
… combine) Every algorithm takes (place_group&, containers...); per-place temporaries come from each shard's own place via the group's stream + memory-resource environments. - elementwise (no cross-place stage): fill, sequence, iota, tabulate, generate, for_each, transform (in-place/unary/binary). - reduce/sum/min/max: per-place cub::DeviceReduce + combine of the P partials; inclusive/exclusive scan: per-place cub::DeviceScan, host prefix of shard totals, in-place fold over the shared address space. - adjacent_difference; count/count_if; histogram_even (read-only, so available on every array including contiguous ones). - copy_if/filter/remove_if (per-place in-place cub::DeviceSelect::If) and unique (std::unique semantics across shard boundaries, O(1) trim against the previous non-empty shard). Size-mutators enforce the contiguous-backing contract with std::invalid_argument. Tests vs host references over locality domains + the sharded_reduce example; docs/cudax/sharded.rst opens with the cooperation-scope ladder. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…arded_array Ordering declarations, not synchronizations (the host returns immediately): fork_from(stream) makes every shard stream depend on the caller stream's enqueued work; join_into(stream) is the mirror. Events come from a small container-owned lazy pool (disableTiming; join events created under the shard's exec scope; fork events keyed by the caller stream's device), so adopted arrays over foreign streams are supported identically. Both members are capture-safe: under an active CUDA graph capture the record/wait pairs become graph dependencies. Tests: eager producer->fork->consumers->join->reader chain with one host sync total, an adopted-array variant over foreign streams, a capture instantiate+relaunch variant, and degenerate no-ops. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
What captures: the non-blocking elementwise family — pure per-shard kernel launches — with the fork_from/join_into idiom. Pinned by tests: per-place (green-context) SM confinement survives inside the instantiated graph; replays recompute from current shard contents; cross-stream dependencies inside capture work; contiguous arrays are transparent to capture and compose with whole-array kernels in one graph. What refuses: everything that allocates containers, transfers host data or synchronizes throws std::runtime_error under an active capture, detected via the new safe query places::stream_in_capture(stream) BEFORE any CUDA call — so the capture stays valid and keeps accepting supported work (under global-mode capture the first illegal call would otherwise both error and invalidate the capture, and several of these paths run under cuda_safe_call, which aborts). Benign by construction (tested): adoption and slice during capture; place_group construction / lazy stream materialization records nothing. Graph-owned memory: place_memory_resource allocation is deliberately left capturable (a graph mem-alloc node from the place's pool); both the balanced-pair and unfreed-allocation shapes are pinned by tests. docs/cudax/sharded.rst gains the contract as a user-facing section. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The round-3 review-response commit (scan init/identity split) reflowed the file's header section with a line-wrap that broke RST syntax throughout: the `.. _cudax-sharded:` label was split across three lines, so docs/cudax/index.rst's `:ref:`sharded containers ... <cudax-sharded>`` never resolved -- undefined label, warnings-as-errors, "Build documentation" CI failure. Restores correct RST for the header/sharded_array section (the round-3 content addition -- the self-contained `const size_t n` line in the usage example -- is preserved); the rest of the file, added by later commits, was already correctly formatted. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ound, per-shard and per-call environments, stream_scope Adds <cuda/experimental/__sharded/concepts.cuh>: concepts describing what a sharded structure is (an indexed collection of shard descriptors — element range, global region, equality-comparable place identity — with documented semantic guarantees and a debug validate()), what per-shard environments supply (get_stream mandatory; get_memory_resource for scratch-bearing algorithms), the optional self-binding capability (default_envs, in the spirit of std::execution's get_env), and the per-call environment machinery (sync_policy query with a throw-before-blocking guard). Adds stream_scope: device currency derived from a stream, sufficient for generic per-shard work — kernels launched into a place's stream execute in the stream's context with its SM confinement regardless of the calling thread's current context (see the accompanying test). Adds default_envs(sharded_array): per-shard environments derived from the binding the container recorded at construction, reusing place_group::env. The concepts are descriptive: the shipped shard/sharded_array model them as-is, with zero container changes. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
concepts/models.cu: static_asserts that the shipped types model the concepts as-is (shard<T> is a shard_descriptor; sharded_array is a sharded_view, owning_sharded and self_bound; place_group::env results are alloc envs), that hand-rolled foreign structures model sharded_view via the descriptor shape (basic_shard_view) and become self_bound by providing default_envs through ADL, and negative checks; runtime checks for validate() semantics (gaps/overlaps rejected, empty shards permitted), default_envs stream identity, and the sync_policy::forbid guard. stream_scope.cu: launches into a domain's pool stream carry the stream's context and SM confinement under four currency regimes (activated / stream_scope / adversarial other-context-current / thrust), with per-domain SM sets asserted identical across regimes and results checked; events created without activation compose across pool streams. This is the evidence for the concepts' no-activation-required design. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…w + environments Adds the generic tier's first algorithm: in-place unary transform over any sharded_view with explicit per-shard environments or, for self-bound structures, environments derived via default_envs. The per-call environment selects the contract: a stream present = asynchronous (fork on entry, join on exit against that stream, no host synchronization; capture-composable), no stream = synchronous convenience (refused under sync_policy::forbid). Generic bodies use stream_scope instead of execution-place activation. The place_group-parameterized signatures are unchanged; algorithms/ generic_tier.cu checks the two tiers agree and exercises explicit envs, the async form against a caller stream, and the forbid refusal. docs/cudax/sharded.rst gains the 'Concepts and the generic tier' section. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…-as-locations form reduce (generic, synchronous): per-shard cub::DeviceReduce on each shard's environment (stream + memory resource), pinned staging, deterministic host fold in shard order; refuses before any work under capture and under sync_policy::forbid. reduce_into (asynchronous): writes the aggregate through a device-writable output iterator on the call environment's stream — per-shard partials reduced directly into stream-ordered scratch, then one deterministic fold kernel whose order (including the empty-shard init contribution) reproduces the synchronous fold bit for bit. No host synchronization: legal under sync_policy::forbid and under CUDA graph capture — the reduce_into test captures transform + reduce_into into one graph, replays it twice, and reads the evolving aggregate from a pinned slot after each replay (the iterative- solver residual shape, previously outside the capturable surface). Output locations: device memory, pinned host memory, or any device-writable output iterator (storage optional). Transient-event fork/join helper added to stream_scope.cuh; scratch and CUB temporaries stay stream-ordered and enclosed, so captured graphs own their memory nodes. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…dependently-written structures, zero adapters Two models exercised end to end through the generic transform / reduce / reduce_into: a fully foreign structure (hand-rolled descriptors over raw cudaMallocAsync buffers, caller-created streams, a hand-rolled cuda::mr resource, its own ADL default_envs — no container, no place_group, no places types in the model), and the adoption model (sharded_array::adopt over caller-owned memory and caller streams). Both agree with references; the foreign model also runs the asynchronous reduce_into against a caller stream. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
… (measured) Two measured fixes: 1. Pinned staging arena (pinned_staging.cuh): a per-call cudaMallocHost/ cudaFreeHost pair costs ~0.9 ms on this system (OS page pinning) and was 86-89% of a small reduce's total time. The combine-bearing algorithms now default to a cached thread-local pinned block (grown on demand, reused); a memory resource carried on the call environment overrides it. Applied to reduce (both tiers) and scan; measured: generic reduce at 200K elements 1050 -> 38 us (27x). Remaining per-call staging sites (count/histogram/copy_if/unique/adjacent_difference) follow in the sweep. 2. Fork/join structure in the asynchronous forms: fork + enqueue all shards first, join all second. The previous per-shard fork/work/join sequence routed each shard's start through the previous shard's completion via the caller's timeline, serializing shards (measured at 100M doubles: reduce_into 487 us vs 267 us for the concurrent synchronous form). After the fix: reduce_into+sync 269.5 us vs sync 265.7 us (parity within 1.5%), and the pipelined form (no per-iteration sync) wins 2.3x at small sizes (17.5 vs 41 us) and 11% at 100M (238.7 us) from cross-iteration overlap. All affected tests re-verified green (generic_tier, reduce_into, foreign_models, reduce_scan). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…e, adjacent_difference Same measured fix as reduce/scan: each of these staged through a per-call pinned allocation (~0.9 ms per call on this system); all five now use the cached thread-local arena. Every algorithm stages through exactly one host block per call (unique already packed its buffers), so the single-block arena serves them all. count_histogram and compaction suites re-verified. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…artitioned views One cub::DeviceTransform call per shard over a tuple of shard pointers: out[i] = op(in1[i], in2[i], ...), any arity, in-place into an input supported, co-partitioning (same shard count + identical per-shard global regions) checked up front. Closes the measured gap where a 3-input update (the iterative-solver reflected/current/initial shape) previously needed two binary passes through a temporary (-6..-9% at GiB scale, worse at small sizes). Same call-environment contract as transform (async when a stream is present, synchronous convenience otherwise). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Contributor
Contributor
Author
|
/ok to test 87ca6d2 |
Contributor
|
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:
📝 SummarySummary by CodeRabbit
WalkthroughThe change adds CUDAX sharded containers, execution contracts, algorithms, examples, documentation, and test infrastructure. It also adds place-group resources, locality-domain grids, stream-capture handling, and dedicated sharded-header tests. ChangesSharded CUDA support
Suggested reviewers: Merge Risk: 🟠 High · up to The new sharded CUDA APIs can fail for valid odd shard sizes and retain unresolved memory-lifetime, synchronization, and correctness defects. These issues can cause exceptions, stale or dangling data, leaks, and misleading test or example results, so the change is not ready to merge. Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Actionable comments posted: 17
🧹 Nitpick comments (7)
cudax/test/sharded/algorithms/generic_tier.cu (1)
246-246: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion:
threwshadows the outerthrewdeclared at Line 211. This PR adds strict-warning test targets, so rename the inner variable to keep-Wshadowbuilds clean.cudax/test/sharded/containers/sharded_array.cu (1)
306-308: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion:
test_empty_shard_allocationsits outside the anonymous namespace that encloses every other test helper, so it gets external linkage. Move the closing}of the namespace after this function.cudax/test/sharded/graph_capture/graph_owned_memory.cu (1)
177-179: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Assert
cudaErrorInvalidValuefor the relaunch. The graph has an allocation node without a matching free node, so relaunching while the allocation remains live can returncudaErrorInvalidValue. The current assertion also accepts unrelated errors.Source: Coding guidelines
cudax/test/sharded/containers/fork_join.cu (1)
226-226: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion: Use the three-argument
cudaGraphInstantiate(&exec, graph, 0)form. CUDA 12 and later removed the legacy five-argument form.cudax/include/cuda/experimental/__places/machine.cuh (1)
84-85: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: both new release-threshold hunks use fixed-width integer symbols from
<cstdint>without including it, and rely on a transitive include from the CUDA headers.
cudax/include/cuda/experimental/__places/machine.cuh#L84-L85: add#include <cstdint>foruint64_tandUINT64_MAX.cudax/include/cuda/experimental/__places/exec/locality_domain.cuh#L240-L241: add#include <cstdint>forUINT64_MAX.As per coding guidelines: "Include all headers needed by the symbols being used; do not rely on transitive includes."
Source: Coding guidelines
cudax/include/cuda/experimental/__sharded/concepts.cuh (1)
435-435: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion:
get_sync_policy_t::operator()carries no CCCL API annotation, while the siblingget_composition_t::operator()on Line 397 uses_CCCL_API. Annotate it for consistent visibility and host/device availability.- [[nodiscard]] constexpr sync_policy operator()(const _Env& __env) const noexcept + [[nodiscard]] _CCCL_API constexpr sync_policy operator()(const _Env& __env) const noexceptSource: Coding guidelines
cudax/include/cuda/experimental/__sharded/default_envs.cuh (1)
71-75: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: shards allocated without a reference stream have
s.stream == nullptr, so every manufactured environment reports the legacy default stream.self_boundthen holds for such an array, but the generic algorithms enqueue all shards onto one stream and get no per-lane concurrency. Document that outcome here, or reject null streams so callers learn that they must supply streams to get lanes.As per path instructions for
cudax/**/*, this review focuses on "stream ordering, host/device annotations, experimental API clarity".Source: Path instructions
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 6a7dc9b1-12a9-4609-bc8c-fc3267e3c611
📒 Files selected for processing (64)
cudax/cmake/cudaxHeaderTesting.cmakecudax/examples/places/CMakeLists.txtcudax/examples/places/sharded_graph.cucudax/examples/places/sharded_multi_field_pipeline.cucudax/examples/places/sharded_pipeline.cucudax/examples/places/sharded_reduce.cucudax/include/cuda/experimental/__places/exec/locality_domain.cuhcudax/include/cuda/experimental/__places/machine.cuhcudax/include/cuda/experimental/__places/place_group.cuhcudax/include/cuda/experimental/__places/place_memory_resource.cuhcudax/include/cuda/experimental/__places/place_partition.cuhcudax/include/cuda/experimental/__places/places.cuhcudax/include/cuda/experimental/__sharded/adjacent_difference.cuhcudax/include/cuda/experimental/__sharded/composition.cuhcudax/include/cuda/experimental/__sharded/concepts.cuhcudax/include/cuda/experimental/__sharded/copy_if.cuhcudax/include/cuda/experimental/__sharded/count.cuhcudax/include/cuda/experimental/__sharded/cuda_safe_call.cuhcudax/include/cuda/experimental/__sharded/default_envs.cuhcudax/include/cuda/experimental/__sharded/fill.cuhcudax/include/cuda/experimental/__sharded/fork_join.cuhcudax/include/cuda/experimental/__sharded/histogram.cuhcudax/include/cuda/experimental/__sharded/pinned_staging.cuhcudax/include/cuda/experimental/__sharded/reduce.cuhcudax/include/cuda/experimental/__sharded/scan.cuhcudax/include/cuda/experimental/__sharded/segmented_reduce.cuhcudax/include/cuda/experimental/__sharded/shard.cuhcudax/include/cuda/experimental/__sharded/sharded_array.cuhcudax/include/cuda/experimental/__sharded/sort.cuhcudax/include/cuda/experimental/__sharded/sort_shared_va.cuhcudax/include/cuda/experimental/__sharded/stream_scope.cuhcudax/include/cuda/experimental/__sharded/transform.cuhcudax/include/cuda/experimental/__sharded/unique.cuhcudax/include/cuda/experimental/places.cuhcudax/include/cuda/experimental/sharded.cuhcudax/test/CMakeLists.txtcudax/test/places/CMakeLists.txtcudax/test/sharded/CMakeLists.txtcudax/test/sharded/algorithms/compaction.cucudax/test/sharded/algorithms/composition_verbs.cucudax/test/sharded/algorithms/count_histogram.cucudax/test/sharded/algorithms/elementwise.cucudax/test/sharded/algorithms/generic_tier.cucudax/test/sharded/algorithms/reduce_into.cucudax/test/sharded/algorithms/reduce_scan.cucudax/test/sharded/algorithms/segmented_reduce.cucudax/test/sharded/algorithms/sort.cucudax/test/sharded/algorithms/zip_transform.cucudax/test/sharded/concepts/foreign_models.cucudax/test/sharded/concepts/models.cucudax/test/sharded/containers/contiguous.cucudax/test/sharded/containers/fork_join.cucudax/test/sharded/containers/sharded_array.cucudax/test/sharded/graph_capture/contiguous_capture.cucudax/test/sharded/graph_capture/elementwise_pipeline.cucudax/test/sharded/graph_capture/graph_owned_memory.cucudax/test/sharded/graph_capture/must_not_capture.cucudax/test/sharded/graph_capture/reduce_scan_capture.cucudax/test/sharded/header_standalone.cucudax/test/sharded/include_only.cucudax/test/sharded/stream_scope.cudocs/cudax/index.rstdocs/cudax/places.rstdocs/cudax/sharded.rst
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.
|
|
||
| std::printf("%zu values -> %zu distinct -> %zu odd, sum = %lld (expected %lld)\n", n, distinct, odds, total, ref_sum); | ||
|
|
||
| if (distinct == 0 || odds != host.size() || total != ref_sum) |
Contributor
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
important: Compare distinct with the host count immediately after std::unique. The current condition accepts every nonzero value. If unique(data) returns an incorrect nonzero count, this example prints the wrong count and still reports PASSED.
Comment on lines
+472
to
+495
| const ::std::vector<cudaStream_t>& get_or_create_streams(const exec_place& place) | ||
| { | ||
| // Locate the cache slot for this place. | ||
| size_t idx = 0; | ||
| for (; idx < places_.size(); idx++) | ||
| { | ||
| if (places_[idx] == place) | ||
| { | ||
| break; | ||
| } | ||
| } | ||
| if (idx >= places_.size()) | ||
| { | ||
| _CCCL_THROW(::std::invalid_argument, "place_group: place does not belong to this group"); | ||
| } | ||
|
|
||
| ::std::lock_guard<::std::mutex> lock(mutex_); | ||
| auto& cache = stream_cache_[idx]; | ||
| if (cache.empty()) | ||
| { | ||
| cache = place.pick_all_streams(*resources_); | ||
| } | ||
| return cache; | ||
| } |
Contributor
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | 🏗️ Heavy lift
important: the slot lookup matches places by VALUE, so duplicate places in a group alias one stream set.
exec_place::device(0) returns a process-global singleton, and exec_place::operator== compares pimpl identity first. Two entries for the same device therefore compare equal. The loop returns the first matching index, so every duplicate place resolves to stream_cache_[0].
Consequences for a group built over repeated places (place_group{exec_place::repeat(exec_place::device(0), N)}, and the two-entry grid in the UNITTEST at Line 517):
get_stream(i, lane)returns the same stream for everyi.envs()hands one identical stream to every shard, so the generic sharded algorithms serialize instead of overlapping.stream_cache_[1..N-1]stays empty, sosync()skips those slots.
Key the cache by place index, and let the index-based get_stream at Line 297 use the index directly instead of searching by value.
Proposed fix
cudaStream_t get_stream(const exec_place& place, size_t lane_id = 0)
{
- const auto& streams = get_or_create_streams(place);
+ const auto& streams = get_or_create_streams(index_of(place));
_CCCL_ASSERT(!streams.empty(), "place has an empty stream pool");
return streams[lane_id % streams.size()];
}
/// `@brief` Get the stream of the idx-th place for a given lane_id.
cudaStream_t get_stream(size_t place_idx, size_t lane_id = 0)
{
- return get_stream(place(place_idx), lane_id);
+ const auto& streams = get_or_create_streams(place_idx);
+ _CCCL_ASSERT(!streams.empty(), "place has an empty stream pool");
+ return streams[lane_id % streams.size()];
}- const ::std::vector<cudaStream_t>& get_or_create_streams(const exec_place& place)
- {
- // Locate the cache slot for this place.
- size_t idx = 0;
- for (; idx < places_.size(); idx++)
- {
- if (places_[idx] == place)
- {
- break;
- }
- }
- if (idx >= places_.size())
- {
- _CCCL_THROW(::std::invalid_argument, "place_group: place does not belong to this group");
- }
-
+ // First index whose place equals `place`; throws when the place is foreign.
+ size_t index_of(const exec_place& place) const
+ {
+ for (size_t idx = 0; idx < places_.size(); idx++)
+ {
+ if (places_[idx] == place)
+ {
+ return idx;
+ }
+ }
+ _CCCL_THROW(::std::invalid_argument, "place_group: place does not belong to this group");
+ }
+
+ // One stream set per cache SLOT, so duplicate places keep distinct streams.
+ const ::std::vector<cudaStream_t>& get_or_create_streams(size_t idx)
+ {
+ _CCCL_ASSERT(idx < places_.size(), "place_group: place index out of range");
::std::lock_guard<::std::mutex> lock(mutex_);
auto& cache = stream_cache_[idx];
if (cache.empty())
{
- cache = place.pick_all_streams(*resources_);
+ cache = places_[idx].pick_all_streams(*resources_);
}
return cache;
}Note that pick_all_streams draws from the place's registry pool, so duplicate places still share the underlying pool streams round-robin. Distinct cache slots at least stop all shards from collapsing onto lane 0 of one slot.
Comment on lines
+131
to
+141
| if constexpr (__env_has_mr) | ||
| { | ||
| auto staging_mr = ::cuda::mr::get_memory_resource(call_env); | ||
| h_new_sizes = | ||
| static_cast<count_type*>(staging_mr.allocate_sync(num_shards * sizeof(count_type), alignof(count_type))); | ||
| } | ||
| else | ||
| { | ||
| h_new_sizes = static_cast<count_type*>(reserved::__pinned_staging(num_shards * sizeof(count_type))); | ||
| } | ||
| ::std::fill(h_new_sizes, h_new_sizes + num_shards, count_type{0}); |
Contributor
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
important: in all four places the host staging buffer taken from the call environment's memory resource is released only on the success path. The intervening cuda_safe_call, barrier, commit_sizes, and kernel-launch calls can throw, so the buffer leaks. Apply the same SCOPE(exit) discipline already used for the device counters.
cudax/include/cuda/experimental/__sharded/copy_if.cuh#L131-L141: releaseh_new_sizesin aSCOPE(exit)armed right after the allocation, and drop the tail release at Line 208.cudax/include/cuda/experimental/__sharded/copy_if.cuh#L275-L285: same change for the out-of-place staging; drop the tail release at Line 349.cudax/include/cuda/experimental/__sharded/adjacent_difference.cuh#L144-L152: guardh_lastand drop the tail release at Line 197.cudax/include/cuda/experimental/__sharded/reduce.cuh#L112-L120: guardh_partials, and also guard the per-shard device scratch allocated at Line 143 so it is freed when a later shard's enqueue throws.
📍 Affects 3 files
cudax/include/cuda/experimental/__sharded/copy_if.cuh#L131-L141(this comment)cudax/include/cuda/experimental/__sharded/copy_if.cuh#L275-L285cudax/include/cuda/experimental/__sharded/adjacent_difference.cuh#L144-L152cudax/include/cuda/experimental/__sharded/reduce.cuh#L112-L120
Comment on lines
+121
to
+129
| if constexpr (__env_has_mr) | ||
| { | ||
| auto staging_mr = ::cuda::mr::get_memory_resource(call_env); | ||
| h_counts = static_cast<size_t*>(staging_mr.allocate_sync(num_shards * sizeof(size_t), alignof(size_t))); | ||
| } | ||
| else | ||
| { | ||
| h_counts = static_cast<size_t*>(reserved::__pinned_staging(num_shards * sizeof(size_t))); | ||
| } |
Contributor
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
important: The synchronous algorithms release host staging and device temporaries only on the success path. Every cuda_safe_call, barrier, and commit_sizes between the allocation and the release throws, so the release is skipped and the memory leaks. unique.cuh already applies the right pattern for its device counters with SCOPE(exit); apply it to every temporary.
cudax/include/cuda/experimental/__sharded/count.cuh#L121-L129: guardh_countswithSCOPE(exit)and release the per-shardd_outfrom Line 143 on the throwing path.cudax/include/cuda/experimental/__sharded/scan.cuh#L120-L128: guardh_totalswithSCOPE(exit)so the Phase 1,barrier, and__generic_mapthrow paths still release it.cudax/include/cuda/experimental/__sharded/histogram.cuh#L111-L119: guardh_histswithSCOPE(exit)and release the per-shardd_histfrom Line 133 on the throwing path.cudax/include/cuda/experimental/__sharded/unique.cuh#L126-L134: add aSCOPE(exit)forh_basenext to the existingd_countsguard.cudax/include/cuda/experimental/__sharded/sort_shared_va.cuh#L422-L423: move thedeferredrelease loop into aSCOPE(exit)and guardd_splitsfrom Line 485 the same way.
📍 Affects 5 files
cudax/include/cuda/experimental/__sharded/count.cuh#L121-L129(this comment)cudax/include/cuda/experimental/__sharded/scan.cuh#L120-L128cudax/include/cuda/experimental/__sharded/histogram.cuh#L111-L119cudax/include/cuda/experimental/__sharded/unique.cuh#L126-L134cudax/include/cuda/experimental/__sharded/sort_shared_va.cuh#L422-L423
Comment on lines
+62
to
+72
| if (__a.__size < __bytes) | ||
| { | ||
| if (__a.__ptr != nullptr) | ||
| { | ||
| places::cuda_safe_call(cudaFreeHost(__a.__ptr)); | ||
| } | ||
| void* __p = nullptr; | ||
| places::cuda_safe_call(cudaMallocHost(&__p, __bytes)); | ||
| __a.__ptr = __p; | ||
| __a.__size = __bytes; | ||
| } |
Contributor
There was a problem hiding this comment.
🩺 Stability & Availability | 🔴 Critical | ⚡ Quick win
critical: the grow path corrupts the cached arena when cudaMallocHost fails. Line 66 frees the old block, then Line 69 throws through cuda_safe_call. __a.__ptr still holds the freed pointer and __a.__size still holds the old size. The next call with __bytes <= __a.__size returns that freed pointer to a combine, and the next grow call, plus the destructor, free it again. A single pinned-allocation failure — the realistic outcome for a grow-on-demand pinned arena — turns into a use-after-free and a double free.
Clear the arena state before allocating.
if (__a.__size < __bytes)
{
if (__a.__ptr != nullptr)
{
places::cuda_safe_call(cudaFreeHost(__a.__ptr));
+ // Drop the stale state before the next allocation: a failing
+ // cudaMallocHost must not leave a freed pointer in the arena.
+ __a.__ptr = nullptr;
+ __a.__size = 0;
}
void* __p = nullptr;
places::cuda_safe_call(cudaMallocHost(&__p, __bytes));
__a.__ptr = __p;
__a.__size = __bytes;
}📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
Suggested change
| if (__a.__size < __bytes) | |
| { | |
| if (__a.__ptr != nullptr) | |
| { | |
| places::cuda_safe_call(cudaFreeHost(__a.__ptr)); | |
| } | |
| void* __p = nullptr; | |
| places::cuda_safe_call(cudaMallocHost(&__p, __bytes)); | |
| __a.__ptr = __p; | |
| __a.__size = __bytes; | |
| } | |
| if (__a.__size < __bytes) | |
| { | |
| if (__a.__ptr != nullptr) | |
| { | |
| places::cuda_safe_call(cudaFreeHost(__a.__ptr)); | |
| // Drop the stale state before the next allocation: a failing | |
| // cudaMallocHost must not leave a freed pointer in the arena. | |
| __a.__ptr = nullptr; | |
| __a.__size = 0; | |
| } | |
| void* __p = nullptr; | |
| places::cuda_safe_call(cudaMallocHost(&__p, __bytes)); | |
| __a.__ptr = __p; | |
| __a.__size = __bytes; | |
| } |
Comment on lines
+94
to
+106
| auto group = place_group{make_locality_domain_grid()}; | ||
| const size_t P = group.size(); | ||
|
|
||
| std::vector<unsigned*> d_smids(P); | ||
| std::vector<float*> d_data(P); | ||
| std::vector<cudaStream_t> streams(P); | ||
| for (size_t i = 0; i < P; i++) | ||
| { | ||
| cuda_safe_call(cudaMalloc(&d_smids[i], NB * sizeof(unsigned))); | ||
| cuda_safe_call(cudaMalloc(&d_data[i], N * sizeof(float))); | ||
| cuda_safe_call(cudaMemset(d_data[i], 0, N * sizeof(float))); | ||
| streams[i] = group.get_stream(i); | ||
| } |
Contributor
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | 🏗️ Heavy lift
important: both tests allocate device memory with device 0 current, then hand the shards to environments that may live on other devices. On a multi-GPU runner the per-shard work executes on device i over device-0 memory, which needs peer mapping, so these tests only hold on single-GPU machines.
cudax/test/sharded/stream_scope.cu#L94-L106: pass the device explicitly (make_locality_domain_grid(0)), or allocated_smids[i]/d_data[i]inside anexec_place_scope(group.place(i)).cudax/test/sharded/concepts/foreign_models.cu#L290-L312: allocate one buffer per place inside anexec_place_scope, or build the spans from a single-device group instead ofexec_place::all_devices().
📍 Affects 2 files
cudax/test/sharded/stream_scope.cu#L94-L106(this comment)cudax/test/sharded/concepts/foreign_models.cu#L290-L312
Comment on lines
+184
to
+194
| sharded::transform(arr, op); // envs derived | ||
| double r = sharded::reduce(arr, cuda::std::plus<>{}, 0.0); | ||
|
|
||
| // explicit environments (any sharded_view, foreign structures included) | ||
| auto envs = sharded::default_envs(arr); | ||
| sharded::transform(arr, envs, op); | ||
|
|
||
| // asynchronous: ordered against a caller stream, no host synchronization | ||
| const auto sp = cuda::std::execution::prop{cuda::get_stream, stream_ref{s}}; | ||
| const auto env = cuda::std::execution::env{sp}; | ||
| sharded::transform(arr, op, env); |
Contributor
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
🔎 Supported by static analysis
🤖 get_repo_knowledge executed:
get_repo_knowledge NVIDIA/cccl /tmp/coderabbit-repo-knowledge/nvidia-cccl-2f773195/learnings /tmp/coderabbit-repo-knowledge/nvidia-cccl-2f773195/conventions
Length of output: 47692
🏁 Script executed:
#!/bin/bash
set -e
printf '%s\n' '--- target file excerpts ---'
sed -n '1,220p' docs/cudax/sharded.rst
printf '%s\n' '--- later referenced lines ---'
sed -n '300,345p' docs/cudax/sharded.rst
printf '%s\n' '--- namespace declarations and sharded references ---'
rg -n -C 2 'using namespace|namespace sharded|sharded::(transform|reduce)|cuda::experimental::sharded' docs/cudax/sharded.rstRepository: NVIDIA/cccl
Length of output: 15502
🏁 Script executed:
#!/bin/bash
set -e
cat -n docs/cudax/sharded.rst | sed -n '320,334p'Repository: NVIDIA/cccl
Length of output: 937
🏁 Script executed:
#!/bin/bash
set -e
printf '%s\n' '--- stream_ref declarations ---'
rg -n -m 20 'class stream_ref|struct stream_ref|using stream_ref|stream_ref[[:space:]]*[({]' cudax/include docs test 2>/dev/null || true
printf '%s\n' '--- nearby namespace usage in documentation ---'
rg -n -C 3 'stream_ref|default_envs|zip_transform' docs/cudax/sharded.rst | head -120Repository: NVIDIA/cccl
Length of output: 9270
🏁 Script executed:
#!/bin/bash
set -e
printf '%s\n' '--- stream_ref declaration context ---'
cat -n cudax/include/cuda/experimental/__stream/stream_ref.cuh | sed -n '30,60p'
printf '%s\n' '--- sharded namespace exports ---'
rg -n -C 3 'stream_ref|using namespace|using ::cuda|_CCCL_BEGIN_NAMESPACE_CUDA|namespace sharded' cudax/include/cuda/experimental/sharded.cuh cudax/include/cuda/experimental/__sharded 2>/dev/null | head -180Repository: NVIDIA/cccl
Length of output: 17271
🏁 Script executed:
#!/bin/bash
set -e
cat -n cudax/include/cuda/experimental/sharded.cuh | sed -n '1,120p'
printf '%s\n' '--- exact namespace declarations in the public sharded header ---'
rg -n 'namespace cuda::experimental::sharded|using .*stream_ref|stream_ref' cudax/include/cuda/experimental/sharded.cuh cudax/include/cuda/experimental/__sharded/*.cuhRepository: NVIDIA/cccl
Length of output: 11994
important: Make each code block self-contained.
The earlier using namespace cuda::experimental::sharded; directive does not define sharded, so the sharded::... calls are unresolved. stream_ref is also declared in cuda::experimental, not re-exported by sharded; qualify it in the examples, including lines 327–328.
Source: Path instructions
This comment has been minimized.
This comment has been minimized.
…ssing Multi-GPU capture crashed with cudaErrorInvalidResourceHandle at event-record time in __wait_stream_on (first seen: 2-GPU run of reduce_into and elementwise_pipeline). Root cause: get_device_from_stream refused to query a capturing stream — believing the query illegal — and guessed the current device instead, so transient fork/join events were created in the capture origin's context while recording on another device's lane stream (cuEventRecord requires event and stream from the same context). Single-device machines made the guess always right, which is why every prior run was green. The premise is false: probed on CTK 13.4, cudaStreamGetDevice succeeds during thread-local, relaxed and global capture, on device and green-context streams alike, without invalidating the capture (so do cuStreamGetCtx and cuStreamGetId). Drop the capture fallback on the CTK >= 12.8 path and just ask; keep it, documented as single-device-only, on the unverified pre-12.8 driver path. get_stream_id keeps reporting unknown under capture, now for the real reason: sync-skip caching must never elide a cudaStreamWaitEvent under capture (the wait IS the edge). stream_scope's capture note rewritten to the corrected premise. Verified single-GPU (GB300): reduce_into, elementwise_pipeline, must_not_capture all green; standalone gates pass. Multi-GPU validation on a 2-GPU node to follow. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Contributor
There was a problem hiding this comment.
Actionable comments posted: 2
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 7cc63628-2f01-448e-b9ab-9a1dd9ec8936
📒 Files selected for processing (2)
cudax/include/cuda/experimental/__places/stream_pool.cuhcudax/include/cuda/experimental/__sharded/stream_scope.cuh
🚧 Files skipped from review as they are similar to previous changes (1)
- cudax/include/cuda/experimental/__sharded/stream_scope.cuh
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.
| * `cudaStreamGetDevice`, ...) would be rejected with | ||
| * `cudaErrorStreamCaptureUnsupported` and would *invalidate* the in-flight | ||
| * capture. | ||
| * capturing). |
Contributor
There was a problem hiding this comment.
🩺 Stability & Availability | 🟡 Minor | ⚡ Quick win
🔎 Supported by static analysis
🏁 Script executed:
#!/bin/bash
set -eu
file='cudax/include/cuda/experimental/__places/stream_pool.cuh'
printf '%s\n' '--- target excerpt ---'
sed -n '1,140p' "$file"
printf '%s\n' '--- related bindings/usages ---'
rg -n -C 3 'is_stream_capturing|cudaStreamIsCapturing' cudax/include/cuda/experimentalRepository: NVIDIA/cccl
Length of output: 17348
🤖 get_repo_knowledge executed:
get_repo_knowledge NVIDIA/cccl /tmp/coderabbit-repo-knowledge/nvidia-cccl-2f773195/learnings /tmp/coderabbit-repo-knowledge/nvidia-cccl-2f773195/conventions
Length of output: 47840
🌐 Web query:
site:docs.nvidia.com/cuda cudaStreamIsCapturing cudaErrorStreamCaptureImplicit legacy stream blocking stream capture
💡 Result:
The cudaErrorStreamCaptureImplicit error occurs when calling cudaStreamIsCapturing on the legacy default stream (cudaStreamLegacy) while a blocking stream on the same device is actively capturing a CUDA graph [1][2][3]. In this state: - The legacy stream enters an unusable state until the blocking stream's capture sequence is terminated [1][4]. - The legacy stream itself does not support stream capture [1][5]. - Any attempted use of the legacy stream while a blocking stream is capturing would inherently create an implicit dependency on the capturing stream(s), which is not permitted [1][4]. - If cudaStreamIsCapturing is called on the legacy stream under these conditions, the function returns cudaErrorStreamCaptureImplicit, and the output pointer for the capture status is left in an unspecified state [1][2]. The capture process of the blocking stream is not invalidated by this error [1][4]. To avoid this, ensure that CUDA graph capture is not performed on streams that synchronize with the legacy stream unless they are explicitly created as non-blocking streams (using the cudaStreamNonBlocking flag) [1][5]. The legacy stream is an implicit stream that synchronizes with all other streams in the same context, except for those specifically marked as non-blocking [5].
Citations:
- 1: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__STREAM.html
- 2: https://docs.nvidia.com/cuda/archive/13.2.1/cuda-runtime-api/group__CUDART__STREAM.html
- 3: https://docs.nvidia.com/cuda/archive/13.1.0/cuda-runtime-api/group__CUDART__STREAM.html
- 4: https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__STREAM.html
- 5: https://docs.nvidia.com/cuda/cuda-runtime-api/stream-sync-behavior.html
🏁 Script executed:
#!/bin/bash
set -eu
file='cudax/include/cuda/experimental/__places/stream_pool.cuh'
sed -n '1,140p' "$file"Repository: NVIDIA/cccl
Length of output: 4846
🏁 Script executed:
#!/bin/bash
set -eu
sed -n '1,140p' cudax/include/cuda/experimental/__places/stream_pool.cuhRepository: NVIDIA/cccl
Length of output: 4846
🌐 Web query:
NVIDIA CCCL cudax/include/cuda/experimental/__places/stream_pool.cuh is_stream_capturing
💡 Result:
In the NVIDIA CCCL library, specifically within cudax/include/cuda/experimental/places/stream_pool.cuh, is_stream_capturing (or its functional equivalent) is utilized to manage capture-safety, particularly when interacting with CUDA streams during graph capture [1]. Recent technical changes in the CCCL repository have clarified that certain operations, such as calling cuStreamGetId, are not capture-safe because they can invalidate an in-flight CUDA graph capture [1]. To address this, the implementation in stream_pool.cuh uses cudaStreamIsCapturing as a safeguard [1]. When active capture is detected, the implementation conservatively returns a default value (such as k_no_stream_id) rather than performing operations that would disrupt the capture sequence [1]. This mechanism ensures that stream management within the stream_pool remains robust during graph capture by preventing unsafe API calls when a stream is actively capturing work [1]. The cudax API also exposes is_stream_capturing as part of its public interface to assist users in identifying the capture status of streams [2].
Citations:
- 1: GitHub pull request 9395 in NVIDIA/cccl (link omitted to avoid creating a cross-reference)
- 2: https://nvidia.github.io/cccl/unstable/cudax/api/index.html
🏁 Script executed:
#!/bin/bash
set -eu
python3 - <<'PY'
from pathlib import Path
p = Path("cudax/include/cuda/experimental/__places/stream_pool.cuh")
lines = p.read_text().splitlines()
for i, line in enumerate(lines[:130], 1):
print(f"{i:4}: {line}")
PYRepository: NVIDIA/cccl
Length of output: 5184
🏁 Script executed:
#!/bin/bash
set -eu
file='cudax/include/cuda/experimental/__stf/utility/cuda_safe_call.cuh'
rg -n -C 8 'cuda_try|template.*cudaStreamIsCapturing|cuda_safe_call' "$file"Repository: NVIDIA/cccl
Length of output: 9691
important: Handle the legacy stream before calling cudaStreamIsCapturing.
is_stream_capturing passes nullptr to cuda_try<cudaStreamIsCapturing>, although its documentation promises false. CUDA returns cudaErrorStreamCaptureImplicit for this query while a blocking stream is capturing, and cuda_try propagates the error. Return false before the query, or remove the unconditional false guarantee.
Source: MCP tools
…GPU aware composition_verbs.cu, segmented_reduce.cu and sort.cu existed but were not in the test source list, so they never ran as ctest targets. Registering them surfaced two single-device assumptions on a 4x GB200 node: - composition_verbs: test_lane_wait sized its arrays with a hard-coded pair of shards; size one shard per place of the group instead. - sort: the shared-address-space engine refuses shards on different devices by contract. On a multi-device grid the test now asserts that refusal (array left untouched), then runs the engine tests on make_locality_domain_grid(0). Single-device behaviour is unchanged. All three pass on 4 GPUs; the full sharded|places suite is 91/91 there. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
… fabric-friendly
The multi-sequence selection that computes the exact splitters had a
merge-path fast path for two runs and, for P > 2, a nested binary search:
one thread per boundary, O(P^2 log^2 n) dependent loads. Measured on a
4x GB200 grid (runs on different devices, peer-mapped pools) it cost
2.8-4.8 ms per sort regardless of size — 0.85-1.5 ms even with every run
local — and made P = 4 slower than P = 2 below ~2^29 keys.
Replace it with two exact selections on local copies:
1. gather every run's every-stride-th key to the selecting place (scattered
but fully parallel loads), and select the sample of rank floor(R/stride).
That sample is inside the rank-R prefix and within (P+1)*stride ranks of
its end, and the samples bracket each run's count-before within stride,
so every run's split lies in a window of at most (P+2)*stride elements;
2. bulk-copy those windows local (contiguous, coalesced) and select again
with the rank offset by the window starts.
Both use rank_select_kernel: one thread per element, each computing its own
rank with P independent binary searches and answering any target of that
rank. No serial chain of dependent loads remains, and none crosses the
fabric. stride is the power of two nearest sqrt(N / (T*P*(P+2))), balancing
sample count against window volume. Trivial boundaries from empty shards are
answered on the host; duplicate sample ranks (boundaries collapsing onto one
sample) are handled.
P = 4, 4x GB200, uniform uint32 keys (GKeys/s): 2^24 5.1 -> 29.4,
2^27 25.4 -> 84.9, 2^28 36.4 -> 103.7, 2^30 97.7 -> 118.2; the selection
phase itself 4.83 -> 0.13 ms at 2^28. P = 1 and P = 2 unchanged.
sort.cu gains test_many_places: P = 3, 4, 5, 7 copies of one place through
the public API — random, heavy ties with a descending custom comparator,
all-equal, uneven with empty shards, one key per shard, fewer keys than
shards, single key, presorted and reverse-sorted. Full sharded suite 43/43.
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
caugonnet
commented
Sep 5, 2026
| test_repeated_runs(group); | ||
| test_many_places(group.place(0)); | ||
|
|
||
| printf("sort: all tests passed\n"); |
Contributor
Author
There was a problem hiding this comment.
should remove this
…sharded views
generate_uniform / generate_normal(data, envs, seed, ...): the fill is
bitwise identical to a whole-array generation for ANY sharding of the
index space. Counter-based Philox with non-DYNAMIC ordering; per shard,
generation starts at the shard's global element index. Poisson excluded
by contract (data-dependent draw count breaks the position mapping).
Path selection per element type picks the mechanisms measured exact:
float uses the host API with curandSetGeneratorOffset (bitwise drop-in
for a stock whole-array run, verified at odd offsets and lengths);
double uses a per-element device-API kernel (curand_init, subsequence =
global index) as a WORKAROUND — the host API's FP64 offset positioning
measured inconsistent at large offsets on current toolkits, so the FP64
sequence of this interface is defined by the position-pure device
mapping instead.
Opt-in vendor tier: random.cuh is not part of the sharded.cuh umbrella;
the test links CUDA::curand. Driven through __generic_map (synchronous
convenience contract: capture-refusing, joins all lanes); per-call
generators live in an explicit RAII holder destroyed after the join.
Certification: test/sharded/algorithms/random_fill.cu — 16/16 bitwise
(uniform/normal x f32/f64 x {1, even, prime-odd, per-domain} shardings)
on GB300/CTK 13.4; the A/B is the per-toolkit gate for the contract.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…r, alignment, kernel signature wrapping) Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Contributor
Author
|
/ok to test 3625825 |
Contributor
There was a problem hiding this comment.
Actionable comments posted: 1
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 66125dee-0726-4cf8-a0c0-3b40443fa43e
📒 Files selected for processing (6)
cudax/include/cuda/experimental/__sharded/default_envs.cuhcudax/include/cuda/experimental/__sharded/random.cuhcudax/include/cuda/experimental/__sharded/sort_shared_va.cuhcudax/test/sharded/CMakeLists.txtcudax/test/sharded/algorithms/random_fill.cucudax/test/sharded/algorithms/sort.cu
🚧 Files skipped from review as they are similar to previous changes (3)
- cudax/test/sharded/algorithms/sort.cu
- cudax/include/cuda/experimental/__sharded/default_envs.cuh
- cudax/include/cuda/experimental/__sharded/sort_shared_va.cuh
Included review availability: Your plan provides up to 12 included reviews per hour; 10 remain after this review.
…normal fill (review r3940236061) Pseudo-random normal generation documents an even-count requirement (CURAND_STATUS_LENGTH_NOT_MULTIPLE); CTK 13.4 happens to accept odd lengths, which is why the invariance A/B was green, but the header promises portability, not toolkit luck. Decompose each shard into an even-offset/even-count main run plus up to two boundary elements taken from 2-element runs at even offsets (via a small stream-ordered scratch, so the neighbor shard's element is never written) — every host-API call is contract-conforming and the global sequence is bitwise unchanged. Test grows a 'tiny + odd' sharding (size-1 shards at odd offsets) that pins the head-only and tail-only decompositions: 20/20 bitwise on GB300. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…n, no trailing copy The pairwise merge tree alternated between the destination and one scratch buffer starting from the destination, so for 3-4 (and 7-8, ...) runs the result ended in scratch and was copied back: one extra whole-shard pass per sort. Pick the first level's buffer from the level count so the last level writes the destination; assert the parity instead of copying. Measured on 4x GB200, P=4, uint32: 2^28 2.59 -> 2.46 ms (109 GKeys/s), 2^30 8.95 -> 8.59 ms (125 GKeys/s); P=3 at 2^28 3.24 -> 3.03 ms. Also measured and rejected here: a one-pass k-way merge (splitter samples -> consistent chunks -> one block per chunk, cub::BlockMergeSort tile) was 3x slower (1.85 vs 0.65 ms per destination at 2^26 keys/GPU) — the tile sort runs at radix-sort speed and eats the pass it saves; and bulk-staging the remote runs with cudaMemcpyAsync (700 GB/s ingest) before a local tree ties the direct tree (0.27 + 0.31 vs 0.42 + 0.14 ms). The remote-read level is already near the fabric's bandwidth; the local levels are cheap. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…in the fill tier The float path of generate_uniform/generate_normal created one host-API Philox generator per shard per call and destroyed it after the join. Two of those steps are host-blocking and serialize across shards: the first generate on a fresh generator pays a lazy state setup (0.2-4 ms measured) and curandDestroyGenerator is an implicit device sync (0.2-1 ms). The fill therefore got SLOWER with every place added — on 4x GB200, 2^28 floats: 1.04 ms at one place, 2.64 ms at four — while the double path (device-API kernel, no generator object) scaled normally. Seed and offset are host-side state applied at the next generate, so the sharding-invariance contract depends only on them, never on a generator's history. Keep generators in a process-wide cache keyed by (device, stream): created once, reseeded and repositioned per call. Keyed by stream because the offset is per object and shards on distinct lanes generate concurrently; sequential reuse on one stream is safe. The cache is deliberately never destroyed (static teardown races the runtime's). After, uniform float, 4x GB200 (P=1 / 2 / 3 / 4, Gelem/s): 2^26: 644 / 802 / 759 / 709 (launch-bound at this size) 2^28: 820 / 1329 / 1560 / 1728 2^30: 880 / 1665 / 2208 / 2767 (3.14x at P=4; P=1 equals stock cuRAND) Normal float 2^28: 468 / 842 / - / 1232. Bitwise identity with stock whole-array cuRAND holds at every P; random_fill invariance test passes. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
generate_uniform (cuRAND tier) -> zip_transform (in-circle test) -> reduce (hit count): the stages compose through the sharded structures alone. Because generation is sharding-invariant and the 0/1-flag sum is integer-exact, the estimate is bitwise reproducible across shardings — the example computes pi twice with different shard boundaries (one odd, prime cut) and checks the hit counts match exactly. The caveat is stated in the header on purpose: production Monte Carlo fuses generation into the consumer and never materializes samples; this example materializes deliberately to show the API seams, and points at the lazy-input iterator relaxation (concepts.cuh v1 notes) as the fused spelling. Registered in the places examples with a CUDA::curand link. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ialized one pi is exactly the workload where a reviewer should ask 'why materialize at all?' — so the example now answers instead of inviting the question: a per-shard fused kernel (positional device-API generation + in-circle test + block reduction in one pass over the index space, zero bytes of samples materialized) runs next to the three-stage materialized pipeline. Each spelling is bitwise reproducible across shardings (both checked); the header states when each is the right choice (fuse for single consumption, materialize when the data must persist and be re-read) and points at the lazy-input iterator relaxation as what makes the fused form generic over foreign structures. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
This comment has been minimized.
This comment has been minimized.
…ragged compaction The typical-cuRAND-use example beyond fills: draw (x, u) candidate pairs as one invariant stream (a float2 view over an even-boundary sharded float array — pair k is always draws 2k/2k+1, element grouping stated as a sharding contract), copy_if the pairs accepted under f(x) = 4x(1-x) into a co-partitioned destination with DATA-DEPENDENT per-shard sizes (the ragged commit), then zip_transform + reduce the moments over the ragged result. Target p(x) = 6x(1-x) gives three exact checks (acceptance 2/3, E[x] = 1/2, E[x^2] = 3/10; measured 0.66680 / 0.49992 / 0.29990 at 8M pairs). The property compaction inherits from generation: a stable filter of a sharding-invariant stream is itself sharding-invariant — the example runs the pipeline under two different shard boundaries and verifies the concatenated accepted sequence is bitwise identical (5593482 pairs). Registered next to sharded_pi in the cuRAND-linked examples. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…icitly The product is the sample set itself — ragged, already distributed and placed for downstream sharded consumers — not a reduced scalar; the moments are verification. Names it as the complementary pole to the pi example (fuse when samples are intermediate; ragged materialization is intrinsic when the dataset IS the output). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…density) Answering the natural question 'and then one uses this x population?': the example now consumes it the way it normally would be consumed — histogram_even over the ragged accepted set builds the empirical density, checked per bin against the analytic mass of p(x) = 6x(1-x) (max bin error 3.1e-4 at 5.6M samples). Downstream algorithms take the data-dependent result unchanged; functional estimates and the density are both shown as consumers. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…tioned scratch The example hand-harvested per-shard sizes into a vector to allocate arrays co-partitioned with the ragged compaction result; allocate_like has done exactly that since the container landed (same sizes, placements AND reference streams — so the consumers stay on the producer's lanes). Same output, less API to misdemonstrate. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…0 of the bench-revisit plan) sharded_csr (vendor-free, in the umbrella): the validated torq-lineage container — rebased self-contained per-shard sub-CSRs, nnz-balanced default split, time_balanced_boundaries rebalance model, row-partitioned output factories, fork/join ordering — re-housed under this branch's discipline: - THREE construction grades, adopt-first: NEW zero-copy adopt() aliases the caller's device colinds/values at row-aligned nnz slices (only the rebased offsets are owned; the caller's arrays stay the storage and keep their placement — closes the recorded adopt-from-views gap and is the cuOpt-faithful entry); from_device/host ctor remain the placed- copy grades that buy the measured locality win. - NO library state on the container: the type-erased lib_state slots are gone (both matrix- and group-side). - fork_from uses get_device_from_stream (capture-safe post-1eec09c27e) instead of hand-rolled capture guessing. sparse.cuh (opt-in vendor tier, random.cuh's model): spmv/spmm and the shard-time rebalance utilities over EXPLICIT caller-held state — cusparse_handles (place-bound, lazy per-place creation under the exec scope, shared across matrices) and spmv_plan/spmm_plan (matrix-bound descriptors/workspace/preprocessed plans, lazy build, pointer-rebind reuse). The per-shard engine pipelines (CSR_ALG2/ALG3, warm-up launch, stream rebinds) are the validated originals, verbatim. This resolves the paused lib_state-vs-explicit-helpers decision in the direction the audit recorded, with the cuRAND tier's lifecycle lesson designed in. Tests: the four suites ported to the new API (handle_lifecycle rewritten for explicit-handles semantics; container suite gains test_adopt pinning alias identity, mutation-through-caller-arrays and container-death safety, replacing the retired lib_state test). All four green on GB300; standalone gates pass for both headers. CMake: container test in the main list, product tests linked against CUDA::cusparse. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…vailability The header-hygiene glob was picking up random.cuh and sparse.cuh unconditionally: on an install without the cuRAND/cuSPARSE dev headers that breaks the BUILD, not just the tests. Vendor headers now leave the glob and get their own hygiene targets inside per-library blocks gated on the toolkit components (TARGET CUDA::curand / CUDA::cusparse), along with their test registrations — the recorded per-library gating pattern (the PR-7 precedent), so library-less installs skip cleanly. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Contributor
Author
|
/ok to test a5e5d4f |
Contributor
😬 CI Workflow Results🟥 Finished in 1h 09m: Pass: 34%/63 | Total: 8h 48m | Max: 51m 21s | Hits: 100%/2652See results here. AI failure analysis1. cudax sharded header compilation cannot find cusparse.h · 37 jobsExplanation: The PR adds every `cuda/experimental/__sharded/*.cuh` file to the unconditional sharded header-test target, causing the opt-in `sparse.cuh` header to compile in build images without cuSPARSE development headers. The dedicated sparse tests already handle this dependency correctly by checking `TARGET CUDA::cusparse` and linking `CUDA::cusparse`. Evidence: Copy this prompt into a coding agentJobs:
|
The integration-facing shape: a consumer library's outputs are its own whole-device buffers, so each shard writes its disjoint row block at y + row_begin (C + row_begin * n_cols) into one plain pointer — no sharded_array on the output side. Bitwise-verified against the sharded_array path in the product tests. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Consumer libraries (raft-based solvers like cuOpt) run their cuSPARSE handles in device pointer mode with alpha/beta living on the device; the host-value API would force an 8-byte D2H sync per call at the seam. Add run_device_scalars to the shard plans (binds POINTER_MODE_DEVICE for the launch, restores host mode; build/warm-up stay host-mode) and the matching contiguous-output spmv/spmm overloads taking const _Tp* d_alpha/d_beta. Bitwise-verified against the host-scalar path. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
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
2 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
closes
Checklist