Places#

Places are abstractions that represent where code executes and where data resides across the non-uniform memory of a CUDA system. They provide a unified interface for managing execution affinity, stream pools, memory allocation, and device context switching – independently of any task-based programming model.

Places come in two flavors:

  • Execution places (exec_place) determine where code is executed.

  • Data places (data_place) specify where data is located in memory.

The places API is part of the cuda::experimental::places C++ namespace and can be used standalone via the cuda/experimental/places.cuh header, without pulling in the full CUDASTF task-graph framework. For backward compatibility, all places types are also available in the cuda::experimental::stf namespace.

Execution places#

An execution place describes a location where computation can occur. The following factory methods create the most common execution places:

  • exec_place::device(id) – a specific CUDA device

  • exec_place::host() – the host CPU

  • exec_place::current_device() – the CUDA device that is currently active

  • exec_place::cuda_context(ctx, devid) – an externally-owned CUDA driver context; the device ordinal is derived from the context when devid is omitted

  • exec_place::locality_domain(devid, domain) – one locality domain of a device (see Locality domain places (CUDA 13.4+))

When an execution place is activated, it sets the appropriate CUDA context (e.g. calls cudaSetDevice). Each execution place also has an affine data place: the memory location naturally associated with it. For a device execution place the affine data place is the device’s global memory; for the host it is pinned host memory (RAM).

A CUDA-context execution place is non-owning. The caller must keep the CUcontext alive while the place and any streams obtained from it are in use.

Data places#

A data place describes a memory location where data can reside. The following factory methods are available:

  • data_place::device(id) – global memory of a specific CUDA device

  • data_place::host() – pinned host memory

  • data_place::managed() – CUDA managed (unified) memory

  • data_place::affine() – the data place naturally associated with the current execution place

  • data_place::locality_domain(devid, domain) – memory localized to one locality domain of a device (see Locality domain places (CUDA 13.4+))

The affine data place is the default: when no data place is specified, data is placed in the memory that is local to the execution place. For example, a task running on device 0 will access data in device 0’s global memory by default.

Non-affine placement is also supported: data can be placed on a different device or in host memory regardless of where the computation runs. This is useful for sparse accesses (leveraging CUDA Unified Memory page faulting) or for addressing memory capacity constraints. Non-affine placement assumes the hardware and OS support such accesses (NVLINK, UVM, etc.).

Places as container keys#

Both exec_place and data_place can be used as keys in standard associative containers. The library provides the required comparison and hash support:

  • ``std::map`` and ``std::set`` use operator< (strict weak ordering) for keys. Both place types implement operator<, so they can be used as ordered map or set keys.

  • ``std::unordered_map`` and ``std::unordered_set`` require a hash function and equality. The library specializes cuda::experimental::stf::hash for both place types (accessible from both the stf and places namespaces), and both implement operator==.

This allows, for example, maintaining per-place handles (e.g. CUBLAS or CUSOLVER handles keyed by exec_place) or per-place caches keyed by data_place, using either ordered or hash-based containers as needed. The following snippet shows lazy creation of a CUBLAS handle per execution place using an std::unordered_map keyed by exec_place:

#include <cuda/experimental/places.cuh>
#include <cublas_v2.h>

using namespace cuda::experimental::places;

cublasHandle_t& get_cublas_handle(const exec_place& ep = exec_place::current_device())
{
  static std::unordered_map<exec_place, cublasHandle_t, hash<exec_place>> handles;
  auto& h = handles[ep];
  if (h == cublasHandle_t{})
  {
    exec_place_scope scope(ep);
    cuda_safe_call(cublasCreate(&h));
  }
  return h;
}

Locality domain places (CUDA 13.4+)#

Some devices partition their multiprocessors and memory into locality domains (see the CU_DEVICE_ATTRIBUTE_LOCALITY_DOMAIN_COUNT device attribute). Compute throughput and memory bandwidth are higher within a domain than across domains, so pinning both execution and data to the same domain can improve locality. Locality domain places expose this capability:

  • locality_domain_count(devid) – number of domains the device exposes (never 0: a device without locality-domain support reports a single whole-device domain; invalid device ordinals are rejected with an exception)

  • exec_place::locality_domain(devid, domain[, split]) – an execution place backed by an SM partition for the requested domain (a green context built from an SM split by locality domain); the optional split selects the SM split method (see below)

  • data_place::locality_domain(devid, domain) – a data place whose allocations are localized to the requested domain (stream-ordered memory pools and VMM physical handles)

  • make_locality_domain_grid(devid[, split]) – a grid with one execution place per domain of the device, all built with the same SM split method

  • locality_domain_helper – enumerates the domains of a device, mirroring green_context_helper; hands out locality_domain_view identity tokens accepted by both factories

exec_place::locality_domain(d, i) and data_place::locality_domain(d, i) share the same domain ordinal and are therefore co-located: the execution place’s affine data place is the matching locality-domain data place.

SM split methods#

Not every SM of a device necessarily participates in a strict per-domain SM split: SMs may not be assigned to any domain, and the default co-scheduling alignment can leave a domain’s incomplete SM groups unassigned. The locality_domain_sm_split enumeration therefore lets execution-place construction choose how the per-domain SM partitions are carved out of the device (with cuDevSmResourceSplit):

  • locality_domain_sm_split::backfill (the default): every domain place is sized to an even share of the device total and backfilled by the driver – with the target domain’s SMs first, then SMs not assigned to any domain, then SMs from other domains – so the domain places together cover the whole device. The backfilled SMs may sit outside the place’s domain (they have no memory affinity with it), and the partition uses the finest co-scheduling granularity, which does not support launching thread-block clusters.

  • locality_domain_sm_split::aligned – only SMs of the domain that form complete co-scheduled groups at the device’s default alignment. Every SM of the place is affine to the place’s domain and thread-block cluster launches remain available, but incomplete groups and SMs outside any domain are left out of the partition.

  • locality_domain_sm_split::fine – all of the domain’s SMs, grouped at the finest co-scheduling granularity (groups of 2). Every SM of the place is affine to the place’s domain, at the cost of thread-block cluster launches.

backfill is the least surprising default: work spread over the domain places uses the whole device. The strictly per-domain methods trade that coverage for affinity – when work is partitioned by data affinity, SMs backfilled from outside a domain execute against remote memory, which can offset the locality benefit the partitioning was meant to capture. Prefer aligned or fine when per-place SM/memory affinity matters more than whole-device coverage.

The split method only affects the execution side: data places take no method, and places built with different methods for the same (device, domain) are distinct places sharing the same (equal) affine data place. Backends without native locality-domain support (pre-13.4 toolkits, the whole-device degrade, the fake-topology override) accept and ignore the method.

All three methods rest on the CUDA 13.4 driver surface that locality-domain places already require (splitting by locality domain with cuDevSmResourceSplit); the extra pieces backfill and fine use predate it, so no method needs a toolkit newer than 13.4.

// Strictly per-domain partitions for affinity-partitioned work
auto grid = make_locality_domain_grid(dev, locality_domain_sm_split::fine);

The following schematic example assumes the usual CUDASTF setup (a context ctx, a logical data lX and a kernel, as in the STF introduction):

const int dev = 0;

// Never 0: a device without locality-domain support reports a single
// whole-device domain, so the same code runs everywhere.
const unsigned int n = locality_domain_count(dev);

// One task per domain (with CUDASTF)
for (unsigned int i = 0; i < n; i++) {
    ctx.task(exec_place::locality_domain(dev, i), lX.rw())
        ->*[](cudaStream_t s, auto x) { kernel<<<16, 128, 0, s>>>(x); };
}

// Or distribute a parallel_for over all domains at once
auto grid = make_locality_domain_grid(dev);
ctx.parallel_for(blocked_partition(), grid, lX.shape(), lX.rw())
    ->*[] __device__(size_t i, auto x) { x(i) *= 2.0; };

Like data_place::device(id), a locality-domain place is identified by a (device, domain) pair that acts as an identity token: construction of a data place performs no existence check, and the ordinals are validated lazily when the place is actually used. Places for distinct domains hash and compare as distinct values, so they can be used as container keys.

Behavior on older toolkits (before CUDA 13.4):

The same API compiles and runs, with whole-device semantics: every valid device reports a single locality domain (locality_domain_count() == 1), data places allocate plain device memory, and execution places activate the whole device. The domain ordinal is still carried through hashing, comparison and to_string(), so distinct ordinals remain distinguishable as labels. This lets code name a locality domain precisely even when the underlying implementation is the whole device, and keeps a single code path across toolkits. The same degrade applies at runtime on a CUDA 13.4+ toolkit whose driver cannot answer the locality-domain query: the count is never 0, and such a device reports exactly one whole-device domain.

Environment variables:

  • CUDASTF_DISABLE_LOCALIZED_MEMORY – locality-domain data places hand out plain device memory instead of domain-localized memory, while execution still runs on per-domain SM partitions. This is an A/B knob to measure the effect of memory localization independently of execution confinement.

  • CUDASTF_FAKE_LOCALITY_DOMAINS=N – forces N domains per device, backed by an even green-context SM split (granularity-aware) with plain device memory, on any green-context-capable toolkit (CUDA 12.4+). This is the converse ablation control – SM confinement without memory localization – and also lets locality-domain code paths be exercised on devices with a single domain. The override is strict: when the device cannot provide N domains (SM budget, group granularity, or no green-context support at runtime), locality-domain queries and factories throw instead of silently reporting a smaller topology.

Setting the current device or context#

The exec_place::activate() method provides a generic alternative to cudaSetDevice() that works uniformly across different execution place types. This is useful when you want to set the current CUDA device or context without using tasks.

The method returns an exec_place representing the previous state, which can be used to restore the original device or context.

Behavior by execution place type:

  • Device places (exec_place::device(id)): Calls cudaSetDevice(id)

  • Green context places: Sets the current CUDA driver context via cuCtxSetCurrent()

  • Host places: No-op

Basic usage with devices:

exec_place place = exec_place::device(1);
exec_place prev = place.activate();  // Switch to device 1

// ... perform operations on device 1 ...

place.deactivate(prev);  // Restore previous device

Alternative restoration pattern:

You can also restore by calling activate() on the returned place:

exec_place place = exec_place::device(1);
exec_place prev = place.activate();

// ... work on device 1 ...

prev.activate();  // Equivalent to place.deactivate(prev)

Usage with green contexts (CUDA 12.4+):

Green contexts provide SM-level partitioning of GPU resources. The activate()/deactivate() methods handle the underlying driver context management:

// Create green contexts with 8 SMs each
green_context_helper gc(8, device_id);
auto view = gc.get_view(0);

exec_place gc_place = exec_place::green_ctx(view);
exec_place prev = gc_place.activate();  // Sets green context as current

// ... GPU work runs with SM affinity ...

gc_place.deactivate(prev);  // Restore original context

RAII scope for scoped activation:

For exception-safe code or when you want automatic restoration, use the exec_place_scope RAII helper:

{
    exec_place_scope scope(exec_place::device(1));
    // Device 1 is now active
    // ... perform operations on device 1 ...
}
// Previous device is automatically restored when scope goes out of scope

The guard automatically restores the previous execution place when it goes out of scope, making it useful for exception-safe code.

Stream management with execution places#

Execution places can be used independently of any task system to manage CUDA streams in a structured way. This is useful when you want to use place abstractions (devices, green contexts) for stream management without the full task-based programming model.

Stream pools for pooled places (device(N), host()) live in an exec_place_resources registry that the caller owns. Pass the registry to exec_place::pick_stream to get a CUDA stream; the per-place pool inside the registry is created lazily on first request and is destroyed when the registry is destroyed.

The method accepts an optional for_computation hint (defaults to true) that may select between computation and data transfer stream pools to improve overlapping. This is purely a performance hint, and it does not affect correctness. Not all execution places enforce it.

#include <cuda/experimental/places.cuh>
using namespace cuda::experimental::places;

// Standalone use: own the registry yourself.
exec_place_resources resources;

// Get a stream from the current device
exec_place place = exec_place::current_device();
cudaStream_t stream = place.pick_stream(resources);

// Use the stream for CUDA operations
myKernel<<<grid, block, 0, stream>>>(d_data);

// Get streams from specific devices (sharing the same registry)
cudaStream_t stream_dev0 = exec_place::device(0).pick_stream(resources);
cudaStream_t stream_dev1 = exec_place::device(1).pick_stream(resources);

Inside a CUDASTF context, the context’s async_resources_handle already holds an exec_place_resources registry. Convenience overloads accept the handle directly so call sites do not have to dereference it:

cudaStream_t stream = place.pick_stream(ctx.async_resources());

Stream pools are populated lazily – CUDA streams are only created when first requested via pick_stream(resources) (or pick_stream(ctx.async_resources()) inside CUDASTF). Self-contained places (exec_place::cuda_stream(s), green-context places) ignore the registry and return their own embedded pool instead, so the user-provided cudaStream_t / CUgreenCtx must outlive any place that wraps it.

Memory allocation with data places#

Data places provide a unified interface for memory allocation that works across different memory types (host, device, managed) and place extensions (green contexts, user-defined places). This allows you to allocate memory while benefiting from the place abstraction.

The data_place::allocate() and data_place::deallocate() methods provide raw memory allocation. The stream parameter defaults to nullptr, which is convenient for non-stream-ordered allocations (host, managed) where the stream is ignored:

#include <cuda/experimental/places.cuh>
using namespace cuda::experimental::places;

// Allocate on host (pinned memory) - stream defaults to nullptr
void* host_ptr = data_place::host().allocate(1024);
// ... use host_ptr ...
data_place::host().deallocate(host_ptr, 1024);

// Allocate on a specific device (stream-ordered)
cudaStream_t stream;
cudaStreamCreate(&stream);
void* dev_ptr = data_place::device(0).allocate(1024, stream);
// ... use dev_ptr with stream ...
data_place::device(0).deallocate(dev_ptr, 1024, stream);
cudaStreamDestroy(stream);

// Allocate managed memory - stream defaults to nullptr
void* managed_ptr = data_place::managed().allocate(1024);
// ... use managed_ptr from host or device ...
data_place::managed().deallocate(managed_ptr, 1024);

Stream-ordered vs immediate allocations:

Different data places have different allocation behaviors:

  • Host (data_place::host()): Uses cudaMallocHost() / cudaFreeHost() - immediate, stream parameter is ignored

  • Managed (data_place::managed()): Uses cudaMallocManaged() / cudaFree() - immediate, stream parameter is ignored (note: cudaFree may introduce implicit synchronization)

  • Device (data_place::device(id)): Uses cudaMallocAsync() / cudaFreeAsync() - stream-ordered

  • Extensions (green contexts, etc.): Behavior depends on the extension implementation

You can query whether a place uses stream-ordered allocation with allocation_is_stream_ordered():

data_place place = data_place::device(0);
if (place.allocation_is_stream_ordered()) {
    // Allocation is stream-ordered - synchronize via the stream
    void* ptr = place.allocate(size, stream);
    myKernel<<<grid, block, 0, stream>>>(ptr);
    place.deallocate(ptr, size, stream);
    cudaStreamSynchronize(stream);
} else {
    // Allocation is immediate - stream is ignored, safe to use right away
    void* ptr = place.allocate(size);
    // ... use ptr ...
    place.deallocate(ptr, size);
}

This abstraction is particularly useful when writing generic code that needs to work with different types of places, including custom place extensions.

Some places need to know the shape of the tensor being allocated, not just its size: a composite data place distributes the allocation according to a partitioner that maps element coordinates to places. allocate_nd() takes the tensor extents (dimension 0 varying fastest) and the element size:

// 2-D tensor of nx x ny doubles, distributed by the place's partitioner
void* ptr = place.allocate_nd(dim4(nx, ny), sizeof(double));
// ...
place.deallocate(ptr, nx * ny * sizeof(double));

For most places this is equivalent to allocate(prod(dims) * elemsize). For composite places it is required: the byte-count allocate() throws there, since a byte count alone cannot carry the geometry the partitioner needs. A caller that genuinely has untyped bytes states that explicitly with allocate_nd(dim4(nbytes), 1), which distributes the buffer with byte granularity. This raw-byte form applies to composite places built from scale-free partitioners only; a composite place backed by a structured partition (see Structured partitions) accepts exactly the extents of the tensor the partition was built for and rejects anything else, including a flat byte count.

For advanced use cases involving CUDA’s Virtual Memory Management (VMM) API, data_place also provides the mem_create() method. This is a lower-level interface used internally by localized arrays (composite_slice) to create physical memory segments that are then mapped into a contiguous virtual address space.

Unlike allocate(), which returns a usable pointer directly, mem_create() returns a CUmemGenericAllocationHandle that must be subsequently mapped with cuMemMap() before use:

#include <cuda/experimental/places.cuh>
using namespace cuda::experimental::places;

// Create a physical memory handle for device 0
CUmemGenericAllocationHandle handle;
data_place::device(0).mem_create(&handle, size);

// The handle must be mapped to a virtual address before use
// (see CUDA VMM documentation for cuMemMap, cuMemSetAccess, etc.)

When to use each method:

  • Use allocate() for most cases - it provides ready-to-use memory with stream-ordered semantics where applicable.

  • Use mem_create() only when you need explicit control over virtual memory mapping, such as creating localized arrays that span multiple devices with a unified virtual address space.

Limitations of mem_create:

  • Only supports device memory and host memory (pinned)

  • Managed memory is not supported by the VMM API

  • The returned handle requires additional VMM API calls to be usable

Custom place extensions can override mem_create() to provide specialized VMM allocation behavior (e.g., memory localization for hardware partitions).

Grid of places#

It is possible to manipulate places which are a collection of multiple places. In particular, it is possible to define an execution place which corresponds to multiple device execution places.

A grid of execution places is an exec_place that contains multiple underlying places. Grids are created with the make_grid free function:

// Create a 1D grid from a vector of places
exec_place grid = make_grid(std::vector<exec_place>{
    exec_place::device(0), exec_place::device(1)
});

The exec_place::all_devices() helper creates a grid of all available CUDA devices:

exec_place all = exec_place::all_devices();

Similarly, exec_place::n_devices(n) creates a grid from the first n devices:

exec_place first_four = exec_place::n_devices(4);

It is possible to retrieve the total number of elements in a grid using the size() method, and individual places with get_place(i):

exec_place grid = exec_place::all_devices();
for (size_t i = 0; i < grid.size(); i++) {
    exec_place dev = grid.get_place(i);
    // ...
}

Grids of places need not be 1D arrays. They can be structured as a multi-dimensional grid described with a dim4 class by passing it to make_grid or n_devices:

// Create a shaped grid: 8 devices arranged as a 2x2x2 cube
exec_place cube = exec_place::n_devices(8, dim4(2, 2, 2));

// Or from an explicit vector
exec_place shaped = make_grid(my_places, dim4(4, 2));

Note that the total size of the dim4 must match the number of places.

It is possible to query the shape of the grid using get_dims(), which returns a dim4 object. Individual places can be accessed by multi-dimensional position using get_place(pos4).

An existing grid can be viewed with different dimensions using reshape(). The new dimensions must contain exactly the same number of places:

exec_place cube = make_grid(my_places, dim4(2, 3, 4));
exec_place flat = cube.reshape(dim4(24));

Reshaping changes only the grid coordinate system. It preserves dimension-0- fastest linear order, so flat.get_place(i) == cube.get_place(i) for every linear index i. It does not reorder, replicate, or remove places.

collapse_axes(first, last) is a convenience operation that combines a contiguous inclusive range of axes. The collapsed extent is the product of the selected extents; later axes shift left and trailing extents become one:

exec_place grid = make_grid(my_places, dim4(2, 3, 4));

exec_place grid_6x4 = grid.collapse_axes(0, 1); // dim4(6, 4)
exec_place grid_2x12 = grid.collapse_axes(1, 2); // dim4(2, 12)
exec_place grid_24 = grid.collapse_axes(0, 3); // dim4(24)

These operations are useful when a partition should consume several axes of a processor grid as one logical axis. They are coordinate transformations, not partitioning: the latter decomposes a place into constituent resources.

The place_partition class partitions an execution place at a given granularity. This is useful for splitting a multi-device grid into its constituent devices, or for partitioning a device into locality domains, green contexts or CUDA streams.

The partitioning granularity is specified by place_partition_scope:

  • place_partition_scope::cuda_device – partition into individual devices

  • place_partition_scope::locality_domain – partition into locality domains (devices without locality-domain support contribute a single whole-device domain; an optional locality_domain_sm_split argument selects the SM split method)

  • place_partition_scope::green_context – partition into green contexts (CUDA 12.4+)

  • place_partition_scope::cuda_stream – partition into CUDA streams

Partitioning exec_place::all_devices() at locality_domain scope is the machine-wide form: it yields every locality domain of every device. The single-device helper make_locality_domain_grid(dev_id) is convenience sugar over this mechanism.

exec_place grid = exec_place::all_devices();

// Partition into individual devices
place_partition devices(grid, place_partition_scope::cuda_device);
for (auto& dev : devices) {
    // dev is an exec_place for a single device
}

// Convert back to an exec_place grid
exec_place new_grid = devices.to_exec_place();

The exec_place::partition_by_scope() method provides a shorthand that returns a new exec_place grid directly:

exec_place grid = exec_place::all_devices();
exec_place by_device = grid.partition_by_scope(place_partition_scope::cuda_device);

When using a grid of places with CUDASTF constructs such as parallel_for, data partitioning policies express how data and index spaces are dispatched over the different places of a grid.

class MyPartition : public partitioner_base {
public:
    template <typename S_out, typename S_in>
    static const S_out apply(const S_in& in, pos4 position, dim4 grid_dims);

    void get_executor(pos4* result, pos4 data_coords, dim4 data_dims, dim4 grid_dims);
};

A partitioning class must implement an apply method which takes:

  • a reference to a shape of type S_in

  • a position within a grid of execution places, described using an object of type pos4

  • the dimension of this grid expressed as a dim4 object

apply returns a shape which corresponds to the subset of the in shape associated to this entry of the grid. Note that the output shape type S_out may be different from the S_in type of the input shape.

To support different types of shapes, appropriate overloads of the apply method should be implemented.

This apply method is typically used by the parallel_for construct in order to dispatch indices over the different places.

A partitioning class must also implement the get_executor virtual method which allows localized data allocators. This method indicates, for each entry of a shape, on which place this entry should preferably be allocated.

get_executor writes a pos4 coordinate in the execution place grid into *result, and its input arguments are:

  • a coordinate within the shape described as a pos4 object

  • the dimension of the shape expressed as a dim4 object

  • the dimension of the execution place grid expressed as a dim4 object

Defining the get_executor makes it possible to map a piece of data over an execution place grid. The get_executor method of a partitioning policy in an execution place grid therefore defines the affine data place of a logical data accessed on that grid.

Predefined partitioning policies#

There are currently two policies readily available:

  • tiled_partition<TILE_SIZE> dispatches entries of a shape using a tiled layout. For multi-dimensional shapes, the outermost dimension is dispatched into contiguous tiles of size TILE_SIZE.

  • blocked_partition dispatches entries of the shape using a blocked layout, where each entry of the grid of places receives approximately the same contiguous portion of the shape, dispatched along the outermost dimension.

This illustrates how a 2D shape is dispatched over 3 places using the blocked layout:

 __________________________________
|           |           |         |
|           |           |         |
|           |           |         |
|    P 0    |    P 1    |   P 2   |
|           |           |         |
|           |           |         |
|___________|___________|_________|

This illustrates how a 2D shape is dispatched over 3 places using a tiled layout, where the dimension of the tiles is indicated by the TILE_SIZE parameter:

 ________________________________
|     |     |     |     |     |  |
|     |     |     |     |     |  |
|     |     |     |     |     |  |
| P 0 | P 1 | P 2 | P 0 | P 1 |P2|
|     |     |     |     |     |  |
|     |     |     |     |     |  |
|_____|_____|_____|_____|_____|__|

Structured partitions#

The classic partitioning policies above are scale-free: blocked_partition splits whatever shape it is handed, knows nothing about the tensor it will be applied to, and always dispatches along the outermost dimension. A structured partition (cute_partition) is the complementary tool: it describes, dimension by dimension, how one specific tensor maps onto a grid of places.

using namespace cuda::experimental::places;

// A 3-D tensor: dimension 1 blocked over the places of the grid,
// dimensions 0 and 2 not distributed
auto part = make_partition(
    dim4(nx, ny, nz),
    partition_spec{whole, blocked<0>, whole},
    grid.get_dims());

Each entry in partition_spec selects a policy for the corresponding tensor dimension: whole (not distributed), blocked<axis>, cyclic<axis>, or block_cyclic<axis>(block_size). Rank, policy, mesh-axis, and leaf counts are preserved in the C++ type; tensor extents, strides, and block sizes remain runtime values. This is strictly more expressive than the classic policies – splitting dimension 1 of a 3-D tensor, or mixing policies across dimensions, cannot be stated with blocked_partition.

The first argument of make_partition is the tensor’s extents: unlike a classic policy, a structured partition is bound to one reference shape, and remains the authority on it. This is a deliberate trade, and the source of most of the type’s properties:

  • Split dimensions are padded up to divisibility (a 10-element dimension blocked over 3 places is treated as 12, in chunks of 4). Padding makes the underlying layout exact and bijective, which is what keeps every query closed-form: validation is a linear pass over the layout, and the owner of a coordinate is a chain of divisions and modulos.

  • Coordinates beyond the true extents (the padding phantoms) own no bytes and do no work: consumers discard them by comparing coordinates against the true extents. This is the predication idiom of CUTLASS/CuTe (“partition the rounded-up shape, predicate the boundary”) rather than per-place clamping, which would break the layout’s uniformity.

Ownership can be queried directly, and – more importantly – a candidate mapping can be scored before any memory is committed:

pos4 owner = part.owner(pos4(x, y, z));   // grid position owning (x,y,z)

// Dry run: same block-majority decision procedure as a real allocation
localized_stats stats = evaluate_localized_placement(grid, part, sizeof(double));
// stats.bytes_per_place, stats.accuracy() (estimated fraction of local bytes),
// stats.nallocs, ... -- tune the spec, then allocate

A structured partition can back a composite data place. Because the partition is bound to one tensor, such a place is per-tensor – allocate with the partition’s exact extents (compare with the classic composite place, which is a reusable shape-free policy):

data_place dp = make_composite_data_place(grid, part);
void* ptr     = dp.allocate_nd(dim4(nx, ny, nz), sizeof(double));
// physical pages land on the place owning them, per the partition
dp.deallocate(ptr, nx * ny * nz * sizeof(double));

Two structured composite places built from equal partitions compare equal, so they denote the same data placement wherever data places are compared.

  • Extents follow the dimension-0-fastest linearization of dim4::get_index() (the convention of STF slices). A row-major front-end must present its whole description in this order – the extents, the per-dimension partition_spec, and any coordinates passed to owner() reverse together, since reversing only the extents would silently re-target each policy at the wrong axis.

  • At most 4 tensor dimensions (the pos4/dim4 domain).

  • Typed partitions and their kernel-facing sub-shapes store exactly their layout leaves. Runtime interfaces (including C/Python opaque handles) erase them to a canonical descriptor only at the data-place boundary.

  • The partition object is trivially copyable and its queries are host/device callable.

The partitioned_axpy example shows the intended workflow end to end: express the partition once, evaluate it, run tasks over data placed by it, and perform a raw geometry-aware allocation.

The same parallel_for entry point that accepts the classic policies accepts a structured partition instance, which then decides both the per-place kernel decomposition and (through the task’s affine data place) the placement of the data those kernels touch – one object, both sides:

// Every place computes exactly the coordinates it owns
ctx.parallel_for(part, grid, lX.shape(), lX.write())
    ->*[] __device__(size_t x, size_t y, size_t z, auto X) { ... };

The shape argument may also be a box describing a region within the tensor the partition was built for (validated by containment) – e.g. the interior of a stencil domain. Each place still enumerates its own coordinates; those outside the region (like the padding phantoms of uneven extents) are skipped by a per-coordinate predicate, so iteration stays aligned with data ownership rather than re-splitting the region:

box interior({1ul, nx - 1}, {1ul, ny - 1}, {1ul, nz - 1});
ctx.parallel_for(part, grid, interior, lX.rw())->*...;

Predication has a cost proportional to the rejected fraction of the enumerated coordinates, which makes it the right tool for regions that are dense in their bounds (interiors: the rejected boundary shell is a surface-to-volume fraction) and the wrong tool for thin regions. For boundary-style updates – a face of the domain, say – prefer one of:

  • fuse the boundary handling into the volumetric kernel’s body when the condition is cheap (application-dependent);

  • iterate the face with a classic scale-free policy (tight, no rejected coordinates) while an explicit dependency keeps placement on the partition’s composite place:

    auto dist = make_composite_data_place(grid, part);
box face({0ul, nx}, {0ul, ny}, {0ul, 1ul});
ctx.parallel_for(blocked_partition(), grid, face, lX.rw(dist))->*...;

    The face’s few remote writes (places computing parts of a face another place owns) are typically negligible against the volumetric traffic.

The fdtd_mgpu example demonstrates the full pattern: a single make_partition call decides which dimension splits for every task – initialization over the full shape, updates over interior boxes, a point source – and places the fields’ data, so changing the distribution of the whole simulation is editing one partition_spec entry.