[STF] Add bundles: grouped logical data as a single task dependency - #10837
[STF] Add bundles: grouped logical data as a single task dependency#10837caugonnet wants to merge 15 commits into
Conversation
A bundle ties several logical data together behind one object (a CSR matrix's three arrays, a graph's topology) submitted as ONE dependency: the construct expands it into ordinary per-field dependencies, and the user lambda receives one tuple of views per bundle. Every field remains a first-class logical data (a bundle owns no data and no tracking state), so bundle-level and bare-handle dependencies interoperate by meeting at the same logical data. Fields declared 'constant' carry a read-only ceiling: whole-bundle modes distribute per field as the strongest admitted mode (rw() on a bundle clamps constant fields to read), and constant fields are const-qualified in every view via the existing readonly_type_of mapping. write() also clamps constant fields to read (a writer needs the structure fetched to interpret what it writes). Mechanism: header-only adapter layer. bundle_dep expands at the context entry points (context + backend_ctx overloads, SFINAE-gated on any_bundle_dep_v so existing overloads are untouched); a named host-device functor regroups the flat views into per-bundle tuples in front of the user function, with its call operator constrained to the wrapped function's applicability so convention probing keeps working; the extended-lambda classification in parallel_for/launch looks through the adapter at the wrapped lambda (identity for everything else). Examples: cg_csr.cu's csr_matrix becomes a bundle (SpMV: 5 deps -> 3, structure constant); pagerank.cu groups the CSR topology as a constant bundle next to its reduce dependency. Test: interface/bundle.cu covers adopt/create constructors, mode distribution and const views (compile-time), mixed bundle+bare-leaf use, task/parallel_for/ host_launch, stream and graph backends. Verified on GB300 sm_103 / CUDA 13.4: new test + both examples build and run under c++17 and c++20; launch/reduce/token regression files recompile clean. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…for overloads Tokens (void_interface deps) are filtered out of the lambda argument list by the constructs, so the regrouping adapter now computes USER-VISIBLE arities per slot (0 for token deps, per-non-void-field counts for bundles) and assembles the call with per-slot tuple pieces flattened by tuple_cat — zero-arity slots vanish, arity-1 slots stay forwarded references (reduction accumulators), bundles stay one tuple. Mixing tokens and bundles in one construct now works and is tested. Also adds the bundle-aware overloads for launch (spec+place, place, bare) and partitioner parallel_for on both context and backend_ctx. The ths-without-place launch form is omitted (ambiguous against the spec+place form; spell the place explicitly). Test additions: token+bundle parallel_for, launch with a bundle. Verified GB300 sm_103 / CUDA 13.4, c++17+c++20, stream+graph. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
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:
📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesCUDASTF now supports non-owning bundles of logical data. Bundle dependencies expand into individual fields and reconstruct grouped arguments for task callables. Execution APIs, lambda classification, C++ examples, Python bindings, and integration tests support this model. CUDASTF bundle dependencies
Suggested reviewers: Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 2
🧹 Nitpick comments (5)
cudax/examples/stf/linear_algebra/cg_csr.cu (1)
31-33: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Use uniform initialization for bundle construction.
cudax/examples/stf/linear_algebra/cg_csr.cu#L31-L33: initialize thebundlebase with braces.cudax/examples/stf/graph_algorithms/pagerank.cu#L75: constructgraphwith braces.cudax/test/stf/interface/bundle.cu#L55: constructBwith braces.cudax/test/stf/interface/bundle.cu#L114: constructCwith braces.As per coding guidelines, “Use uniform initialization for class constructors and compile-time conversions.”
Source: Coding guidelines
cudax/test/stf/interface/bundle.cu (2)
31-37: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Use the CCCL host-device API macro.
Replace
__host__ __device__with_CCCL_HOST_DEVICE_APIoncheck_rw_view_types.Based on learnings, CUDAX test CUDA sources use
_CCCL_HOST_DEVICE_APIfor host-and-device helpers.Sources: Coding guidelines, Learnings
41-41: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Declare
Nas constexpr.
Nis a compile-time value. Useconstexpr size_t N{64};.As per coding guidelines, “All variables that can be evaluated at compile time must be declared
constexpr.”Source: Coding guidelines
cudax/include/cuda/experimental/__stf/internal/bundle.cuh (2)
79-90: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: fix the trait name in the documentation and reject unknown traits.
Line 83 documents the trait as
constant_t, but the tag isconstant. Unknown trait types are also folded away silently, so a misspelled trait produces a mutable field with no diagnostic.proposed change
- * `@tparam` Traits optional traits (`constant_t`) + * `@tparam` Traits optional traits (`constant`) */ template <typename T, typename... Traits> struct field { + static_assert((::cuda::std::is_same_v<Traits, constant> && ...), "unknown field trait"); using type = T; static constexpr bool is_constant = (::cuda::std::is_same_v<Traits, constant> || ... || false); };
324-339: 🩺 Stability & Availability | 🔵 Trivial | ⚡ Quick winsuggestion: Add host-device annotations to
arity_ofandoffset_of. When clang-cuda compiles this code,slot_piececannot call these unannotated functions. Add_CCCL_HOST_DEVICEto both functions.Source: Path instructions
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 2a7f9cb5-3cf1-4489-8b2a-c29cffc6b31d
📒 Files selected for processing (9)
cudax/examples/stf/graph_algorithms/pagerank.cucudax/examples/stf/linear_algebra/cg_csr.cucudax/include/cuda/experimental/__stf/internal/backend_ctx.cuhcudax/include/cuda/experimental/__stf/internal/bundle.cuhcudax/include/cuda/experimental/__stf/internal/context.cuhcudax/include/cuda/experimental/__stf/internal/launch.cuhcudax/include/cuda/experimental/__stf/internal/parallel_for_scope.cuhcudax/test/stf/CMakeLists.txtcudax/test/stf/interface/bundle.cu
…kernel(_chain) CodeRabbit round: (1) the bundle partitioner parallel_for and launch overloads on context now live inside the same CUDASTF_DISABLE_CODE_GENERATION / CUDA-compilation guard as their non-bundle counterparts (verified by compiling with the macro defined); backend_ctx has no such guard convention, unchanged. (2) interface/bundle.cu now exercises cuda_kernel and cuda_kernel_chain with bundle dependencies on both backends. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
There was a problem hiding this comment.
🧹 Nitpick comments (1)
cudax/test/stf/interface/bundle.cu (1)
29-35: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winsuggestion: Make the CUDA-kernel tests observable.
bundle_kerneladds0.0 * (...)toout. The assertion at Line 120 cannot distinguish successful execution from skipped execution, incorrect bundle-field wiring, or an omitted descriptor incuda_kernel_chain. Use separate output logical data for these calls, write a deterministic non-zero result, and assert one contribution fromcuda_kerneland two contributions fromcuda_kernel_chain.Verify that the targeted stream and graph tests fail when either descriptor is removed.
As per path instructions, “Focus on correctness, lifetime/resource ownership, stream ordering, host/device annotations, experimental API clarity, tests, and compatibility with the supported CUDA toolchains.” The PR objective requires coverage for both APIs on both backends.
Also applies to: 104-113
Source: Path instructions
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 9300a8f6-310b-4dfc-95de-01f7dc6d6e86
📒 Files selected for processing (2)
cudax/include/cuda/experimental/__stf/internal/context.cuhcudax/test/stf/interface/bundle.cu
🚧 Files skipped from review as they are similar to previous changes (1)
- cudax/include/cuda/experimental/__stf/internal/context.cuh
@caugonnet, there was an error processing your request: See the following link for more information: https://docs.gha-runners.nvidia.com/cpr/e/2/ |
|
/ok to test 06012e2 |
N is constexpr and read without odr-use; Clang14's -Wunused-lambda-capture rejects the explicit capture under -Werror (CTK12.0 Clang14 CI lane). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
/ok to test f534b2f |
This comment has been minimized.
This comment has been minimized.
The header-hygiene lane forbids bare 'I' in CCCL headers; bundle.cuh's field-index template parameters are now 'Idx'. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
/ok to test 8d4e1fe |
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
…d code cuda.stf._experimental.bundles adds bundle / constant / bundle_dep / bundle_task over the flat bindings: bundle deps flatten into ordinary deps before the Cython task constructor, and bundle_task.get(i) regroups per-field views with each bundle counting as one slot (mirroring the numba_task interop wrapper idiom). The conformance semantics match the C++ side case for case: whole-bundle modes distribute per field (rw()/write() clamp constant fields to read), explicit requests above a field's ceiling raise, unspecified fields in per-field spellings default to read. Registering a CUDA-Array-Interface object now infers the device data place from cudaPointerGetAttributes (previously an opaque host-pinning assertion). tests/stf/test_bundles.py implements the shared conformance checklist: mode distribution, ceiling errors, slot counting, adoption vs registration, mixed bundle + bare-field use, device-array inference. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
/ok to test ce5408e |
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: 51147ad5-d5f9-4a79-8fbc-51243f6fb1f4
📒 Files selected for processing (3)
python/cuda_stf/cuda/stf/_experimental/__init__.pypython/cuda_stf/cuda/stf/_experimental/bundles.pypython/cuda_stf/tests/stf/test_bundles.py
| "bundle": ".bundles", | ||
| "bundle_dep": ".bundles", | ||
| "bundle_task": ".bundles", | ||
| "constant": ".bundles", |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
important: Add bundle, bundle_dep, bundle_task, and constant to __all__. The lazy loader supports direct imports, but from cuda.stf._experimental import * omits all four new public APIs.
__all__ = [
"AccessMode",
"CudaStream",
"DeviceArray",
"TaskGraph",
+ "bundle",
+ "bundle_dep",
+ "bundle_task",
+ "constant",| err, attr = _rt.cudaPointerGetAttributes(cai["data"][0]) | ||
| if int(err) == 0 and attr.type == _rt.cudaMemoryType.cudaMemoryTypeDevice: | ||
| return ctx.logical_data(value, _b.data_place.device(attr.device)) |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/usr/bin/env bash
set -euo pipefail
python - <<'PY'
from cuda.bindings import runtime as rt
assert hasattr(rt.cudaMemoryType, "cudaMemoryTypeDevice")
assert hasattr(rt.cudaMemoryType, "cudaMemoryTypeManaged")
PY
rg -n -C 5 'cudaPointerGetAttributes|cudaMemoryTypeManaged|logical_data\(value\)' \
python/cuda_stf/cuda/stf/_experimental/bundles.pyRepository: NVIDIA/cccl
Length of output: 266
🏁 Script executed:
#!/usr/bin/env bash
set -euo pipefail
printf '%s\n' '--- target file ---'
sed -n '1,130p' python/cuda_stf/cuda/stf/_experimental/bundles.py
printf '%s\n' '--- related symbols and tests ---'
rg -n -C 4 'cudaPointerGetAttributes|cudaMemoryType(Device|Managed)|_register|bundle\(' \
python/cuda_stf tests python 2>/dev/null | head -n 300
printf '%s\n' '--- repository declarations for CUDA Python bindings ---'
rg -n -C 3 'cuda-python|cuda\.bindings|cudaMemoryTypeManaged|cudaMemoryTypeDevice' \
pyproject.toml setup.cfg setup.py requirements* python CMakeLists.txt 2>/dev/null | head -n 250Repository: NVIDIA/cccl
Length of output: 28202
🌐 Web query:
CUDA Python bindings cudaPointerGetAttributes cudaMemoryTypeManaged cudaPointerAttributes device managed memory documentation
💡 Result:
In the NVIDIA CUDA Python bindings (cuda-python), you can determine the type of memory associated with a pointer using the cudaPointerGetAttributes function, which maps directly to the underlying CUDA runtime API [1][2][3]. When you query a pointer, the cudaPointerAttributes structure returned contains a type field, which indicates the memory type using the cudaMemoryType enum [4][2][3]. Managed memory is identified by the value cudaMemoryTypeManaged [2][3]. The cudaPointerAttributes structure includes the following relevant fields: - type: An enum (cudaMemoryType) identifying the memory type (e.g., cudaMemoryTypeManaged, cudaMemoryTypeDevice, cudaMemoryTypeHost, or cudaMemoryTypeUnregistered) [4][3]. - device: The device against which the memory was allocated [4][3]. - devicePointer: A device-side pointer alias (if accessible) [3][5]. - hostPointer: A host-side pointer alias (if accessible) [3][5]. For a higher-level object-oriented interface, the cuda-python library also provides cuda.core.ManagedBuffer, which features an is_managed property to directly check if a buffer is managed [6]. Note that when working with low-level memory resources or specific driver/hardware configurations, relying solely on driver pointer attributes can sometimes be ambiguous; for instance, the library's ManagedMemoryResource uses a combination of driver signals and resource-level information to accurately report managed status [7][8]. To access these features in your code: 1. Use cuda.bindings.runtime.cudaPointerGetAttributes(ptr) to retrieve the cudaPointerAttributes object [2][9]. 2. Compare the resulting type attribute against cuda.bindings.runtime.cudaMemoryType.cudaMemoryTypeManaged [2]. 3. If using the high-level cuda.core API, simply check the .is_managed property on your buffer object [6].
Citations:
- 1: https://nvidia.github.io/cuda-python/cuda-bindings/latest/api.html
- 2: https://github.com/NVIDIA/cuda-python/blob/2a2186c8/cuda_bindings/docs/source/module/runtime.rst
- 3: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__UNIFIED.html
- 4: https://docs.nvidia.com/cuda/cuda-runtime-api/structcudaPointerAttributes.html
- 5: https://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/html/group__CUDART__UNIFIED_gccb4831aa37562c0af3e6b6712e0f12c.html
- 6: https://nvidia.github.io/cuda-python/cuda-core/latest/generated/cuda.core.ManagedBuffer.html
- 7: Fix is_managed reporting for pool-allocated managed memory cuda-python#1924
- 8: https://github.com/NVIDIA/cuda-python/blob/2a2186c8/cuda_core/cuda/core/_memory/_managed_memory_resource.pyx
- 9: https://github.com/NVIDIA/cuda-python/blob/2a2186c8/cuda_bindings/cuda/bindings/cyruntime.pyx.in
🌐 Web query:
site:docs.nvidia.com/cuda/cuda-runtime-api cudaPointerAttributes device cudaMemoryTypeManaged allocated
💡 Result:
When querying cudaPointerAttributes for a pointer that has been allocated with cudaMemoryTypeManaged (such as via cudaMallocManaged), the device field indicates the device that was current at the time the memory was allocated or registered [1][2]. This differs from memory categorized as cudaMemoryTypeDevice, where the device field identifies the specific physical device where the memory resides [1][2]. For managed memory, because the Unified Memory system automatically migrates data between host and device as needed, the device attribute in cudaPointerAttributes serves as a reference to the initial context or "home" device state rather than a static physical location [1][2]. If the device associated with this initial allocation is deinitialized, the managed allocation is also impacted, as it depends on that device's state [1][2].
Citations:
- 1: https://docs.nvidia.com/cuda/cuda-runtime-api/structcudaPointerAttributes.html
- 2: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__UNIFIED.html
important: Handle cudaMemoryTypeManaged like cudaMemoryTypeDevice. Otherwise, managed CUDA Array Interface values fall through to ctx.logical_data(value) without a device data place and can trigger the host-pinning assertion. Add a regression test for a managed-memory CAI value passed to bundle.
cg.py ported cg_csr_stackable.cu but 'simplified to use a dense matrix' — spelling three dependencies per task for a CSR matrix was the friction. cg_csr_bundle.py is the honest sparse port: the CSR matrix is one bundle (values tracked, structure constant), every task takes it as a single argument, numba kernels launch on the task stream, and the solve verifies A @ x == b. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
/ok to test 825d681 |
This comment has been minimized.
This comment has been minimized.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
/ok to test 27e9971 |
Review: the separate bundle_task wrapper forked the API for no reason. context.task now flattens bundle dependencies itself (several flat deps, one slot) and the task object gains get(slot) — uniform for plain and bundle dependencies alike: a plain dep's get(i) returns its CUDA Array Interface view, a bundle's returns a namespace of per-field views. bundle_task is gone; the example reads 'with ctx.task(A.read(), lx.read(), ly.rw()) as t'. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
/ok to test 99f0222 |
This comment has been minimized.
This comment has been minimized.
Review: passing three separate arrays to the SpMV kernel lost the
bundle abstraction at the device boundary. task.get(slot) now returns
a cached namedtuple instead of a namespace — numba types namedtuples
of arrays, so the whole bundle view crosses the kernel boundary as ONE
argument (the device-side analog of the C++ tuple view):
@cuda.jit
def _spmv_kernel(a, x, y):
... a.vals[k] * x[a.colind[k]] ...
Also: tests and the example now use the public package surface
(stf.AccessMode / stf.context / stf.data_place) instead of reaching
into _stf_bindings, and compare IntFlag members directly instead of
unwrapping .value.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
/ok to test 6e52392 |
This comment has been minimized.
This comment has been minimized.
Review: bundle(ctx, **fields) was a method spelled as a free
constructor. context.bundle(**fields) now mirrors ctx.logical_data,
so the whole surface anchors on the context:
A = ctx.bundle(vals=vals, colind=constant(colind))
with ctx.task(A.rw(), ly.rw()) as t:
a = t.get(0)
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
/ok to test fa12691 |
This comment has been minimized.
This comment has been minimized.
Review: passing raw arrays (vals=np.zeros(...)) conflated registration with grouping and hid the data-place policy inside the bundle (which is what forced the pointer-attribute inference). Bundles now group handles only — register arrays first with ctx.logical_data, which owns the placement policy explicitly. Raw values raise a TypeError pointing there. Duck-typed on read()/rw() so stackable logical data qualify. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
/ok to test 9df192e |
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
Bundles model objects (invariant-bound fields, library descriptors), not loose collections: a solver workspace whose tasks touch varying subsets with varying modes (Krylov vectors) should keep bare deps — grouping would over-declare and serialize tasks that share no data. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
/ok to test d30994e |
Review: bundling adds no synchronization of its own — a bundle dep flattens to exactly the leaf deps you would write by hand. The docs now say so precisely: the cost of whole-object spellings and the read-default is acquiring (ordering + transferring) fields a task never touches — right for objects, meaningless for workspaces. AccessMode.NONE completes the mode set: an excluded field contributes no dependency and no transfer, and its view is None (the namedtuple keeps the field for shape stability). Also legitimate for object use: structure-only access to a matrix without acquiring its values. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
/ok to test a7c3c4e |
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
stackable_context.bundle() / stackable_context.task() accepting bundle
deps / stackable_task.get(slot) — the symmetric counterparts of the
plain-context integration, so bundles compose with graph scopes and
launchable_graph_scope record/replay. Exercised by a Newton physics
phase-x-group prototype where per-State token bundles ({rigid, soft})
replace hand-maintained token dicts: whole-bundle modes distribute,
sensors use dep(rigid=READ, soft=NONE), and the recorded task graph
replays at parity with native capture while overlapping sensor work
the coarse-token layout serializes.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
/ok to test d04ac17ec4696802ed922186973c3719a20933c9 |
@caugonnet, there was an error processing your request: See the following link for more information: https://docs.gha-runners.nvidia.com/cpr/e/2/ |
|
/ok to test d04ac17 |
⏱️ CCCL compile-time benchmark comparison: Public headers compile-time benchResult: 0 regression row(s), 2 improvement row(s) above threshold.
Artifacts: reports and traces Direct file processing
🟢 Direct file processing — Improvements
|
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
🥳 CI Workflow Results🟩 Finished in 5h 38m: Pass: 100%/148 | Total: 1d 15h | Max: 57m 16s | Hits: 100%/37016See results here. |
|
|
||
| // The CSR graph topology is one object: a bundle of the two constant arrays. | ||
| // Tasks depend on `graph` as a whole and receive one tuple of (const) views. | ||
| bundle<field<slice<int>, constant>, field<slice<int>, constant>> graph(loffsets, lnonzeros); |
There was a problem hiding this comment.
Not sure about the capabilities and constraints, but as a putative user I'd expect:
bundle<const slice<int>, slice<int>> graph(loffsets, lnonzeros);Also, having template deduction take care of things would be massively nicer:
auto graph = bundle(as_const(loffsets), as_const(lnonzeros));2 participants
Description
Adds bundles: non-owning groups of logical data that tasks can depend
on with a single argument, while every constituent field remains an
ordinary logical data usable on its own.
Many pieces of data are one object at the level users reason about, but
several logical data at the level STF tracks: a CSR/BSR matrix
(values + column indices + row offsets), a graph (offsets + nonzeros),
paired operands that always travel together. Today each task spells every
array as a separate dependency and receives one lambda argument per array
(
cg_csr.cudeclared 5 deps and 5 lambda args for SpMV; everygraph_algorithms/example re-declares the same two topology deps).Key properties:
copies — no data ownership, no tracking state. Bundle-level and
bare-handle dependencies interoperate by meeting at the same logical
data (tested: bundle dep + bare dep on its own field in one task).
constantis a per-bundle read ceiling. Whole-bundle modesdistribute per field as the strongest admitted mode:
rw()clampsconstant fields to
read, and their views are const-qualified in everyspelling (via the existing
readonly_type_of), checked at compiletime.
write()also clamps constant fields toread(a writer needsthe structure fetched to interpret what it writes). This is a promise
about bundle users, not global immutability: constant fields generate
ordinary read deps, which is what serializes correctly against
legitimate writers through the bare handle.
deps at the context entry points (new SFINAE-gated overloads on
contextandbackend_ctx; existing overloads untouched); a namedhost-device functor regroups the flat views into one tuple per bundle
in front of the user function, with its call operator constrained to
the wrapped function's own applicability so convention probing keeps
working. The extended-lambda classification in
parallel_for/launchlooks through the adapter at the wrapped lambda (identity otherwise).
Token (
void_interface) deps produce no lambda argument, and theadapter's grouping is computed in user-visible arities so tokens and
bundles mix freely.
Covered constructs:
task,parallel_for(incl. partitioner form),launch,host_launch,cuda_kernel,cuda_kernel_chain.Examples updated
linear_algebra/cg_csr.cu:csr_matrixbecomes a bundle (SpMV5 deps → 3; structure declared constant).
graph_algorithms/pagerank.cu: CSR topology grouped as a constantbundle next to its reduce dependency.
Python front end
cuda.stf._experimental.bundlesimplements the same feature over thePython bindings with zero shared code:
ctx.bundle(...)(mirroringctx.logical_data) andconstantproducebundle dependencies that
ctx.task(...)accepts directly (several flatdeps, one slot), and
task.get(slot)returns per-slot views — a plaindependency's view, or a namedtuple of per-field views for a bundle,
which numba types as a single kernel argument: the bundle stays one object
all the way into device code (
a.vals[k] * x[a.colind[k]]). The two front ends implement one conformance checklist casefor case: mode distribution against ceilings, loud errors on explicit excess,
unspecified-fields-default-read, one submitted dependency = one
getslot.Registering CUDA-Array-Interface objects now infers the device data place
from the pointer attributes. Tested in
tests/stf/test_bundles.py, plus a worked example:tests/stf/examples/cg_csr_bundle.py— the sparse CG that the earliercg.pyport avoided ("simplified to use a dense matrix"): the CSR matrix isone bundle, every task takes it as a single argument, numba kernels launch on
the task stream, and the solve verifies
A @ x == b.Testing
New
test/stf/interface/bundle.cu: adopt/create constructors, modedistribution and const views (compile-time asserts), mixed bundle +
bare-leaf use, token + bundle in one construct,
task/parallel_for/launch/host_launch, stream and graph backends. Verified on GB300(sm_103, CUDA 13.4) under c++17 and c++20; launch/reduce/token
regression files recompile clean.
Not in this PR (planned follow-ups)
t.get<view>(i)for bundles (slot-table generalization).wp.sparse.BsrMatrixadapter consuming the Python layer (planned as a Warp-side PR stacked on Warp's STF integration).stackable_ctxintegration: bundle deps currently fail loudly there (the adopting constructor takeslogical_data, and stackable task paths assert onstackable_task_dep). The planned integration distributesresolve_dep/validation over bundle members and adds a stackable bundle variant.🤖 Generated with Claude Code