cudax places: place_group — a grid of places plus its execution resources - #3
Open
caugonnet wants to merge 14 commits into
Open
cudax places: place_group — a grid of places plus its execution resources#3caugonnet wants to merge 14 commits into
caugonnet wants to merge 14 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>
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
1 participant
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.
What this adds
cuda::experimental::places::place_group: a set of execution places together with the resources it takes to execute on them — lazily created per-place stream pools and per-place memory resources.The one-liner: a grid of places names where things can run; a
place_groupowns what it takes to run there. A grid (orstd::vector<exec_place>) is a pure value — copyable, transient, derivable fromplace_partition, no lifetime. Aplace_groupis a resource scope: pools, memory resources, teardown order. Two groups over the same grid are two deliberately distinct isolation scopes. This mirrorsMPI_Group(membership) vsMPI_Comm(membership + attached state), with explicit construction of the second from the first. Construction stays a one-liner:place_group{grid}.Contents
place_group: constructors from a vector of places / anexec_placegrid (flattened), factoriesby_devices(ids = {})andby_locality_domains(ids = {})(whole-device fallback where domains are unsupported)exec_place_resourcesregistry; a group can instead borrow the registry of an STFasync_resources_handle(place_group{places, ctx.async_resources()}), so exactly one pool owner exists when both layers coexist. The borrowed handle is kept alive by the group.place_memory_resource: acuda::mrresource overdata_place(device / host / managed / green-context / locality-domain places), plusplace_group::env(...)building the stream + memory-resource environments accepted by CUB single-call algorithms — temporaries land on the place that runs the workall_device_ids,places_from_devices,places_from_grid,places_from_locality_domains,make_stream_wait_for(place, color)streams from the per-place pools, with round-robin assignment for independent workTests
cudax/test/places/place_group/place_group.cu: construction and factories; lazy per-place stream pools (stability per (place, color), distinct colors = distinct streams, kernels run on every place); isolation (two groups over the same places own distinct pools); memory resources (stream-ordered + sync paths, equality by place); borrowing (the borrowed group's pool is the handle's pool, an owning group's is not; the raw-registry seam); move semantics. Also registered as a unittested header.Verified on a 2-locality-domain GB300 node (CUDA 13.4): all places tests pass (17/17), including the new ones.
Sequenced next
The sharded container/algorithm tier that consumes
place_groupis stacked on this PR (next in the series).Opened on the fork for design review with Andrei — API shape, naming, and the ownership/borrowing seam are the review points.
🤖 Generated with Claude Code