Partitioning Work Across Devices and the Cost of Every Dependency

Adding a second device to a computation adds a second processor and a new class of cost. The processor is easy: a device that can execute the same kernel on half the data roughly halves the compute time, provided the kernel is the kind that scales. The cost is coordination, and it appears wherever the partition boundary requires the two halves to agree on something: a reduction, a halo exchange, a dependency between stages, a shared data structure.

The design question is not how to use all the devices. It is how to choose a partition whose boundaries require the fewest agreements, and then how to make each agreement cost as little as possible. Getting that wrong produces a program that uses more devices and runs slower, which is a common enough outcome that it deserves a name: the parallel slowdown.

Three partition shapes and what each one costs

The simplest partition is disjoint: each device owns a slice of the output and reads whatever input it needs, with no cross-device communication at all. A per-pixel image filter, an elementwise transform, an independent simulation of many small systems: all of these divide cleanly. The cost is the input replication: if every device needs the whole input, the input is copied N times, and the transfer volume scales with the device count rather than the problem size.

The second shape is a halo partition: each device owns a contiguous region plus a border of neighbouring data, and after each step the borders are exchanged. A stencil computation is the canonical case. The transfer volume per step is proportional to the perimeter of each region, while the compute is proportional to its volume. For a three-dimensional region split across N devices, the perimeter-to-volume ratio grows as the cube root of N in the best case, so the communication fraction grows slowly; for a two-dimensional region it grows as the square root, and for one dimension linearly. The shape of the partition therefore matters as much as its count.

The third shape is a pipeline: each device runs one stage of a sequence, and data flows from stage to stage. The cost is that the slowest stage sets the rate for the whole pipeline, and the pipeline must be filled and drained, which adds latency proportional to the number of stages. A pipeline of N stages with unequal work has a steady-state throughput set by the maximum stage time, not the average, so balancing is not optional.

/* One five-point stencil step over the interior rows of a device's
   block, reading `src` and writing `dst`. The two must be distinct: the
   update reads rows r-1 and r+1, so an in-place pass would let some
   threads read values their neighbours had already overwritten. Rows 0
   and rows+1 are halos; `cols` is also the block size. */
__global__ void stencil(const float *src, float *dst, int rows, int cols)
{
    int r = blockIdx.x + 1;              /* interior row, 1..rows */
    int c = threadIdx.x;
    if (c == 0 || c == cols - 1) return; /* column halos are fixed */
    dst[r * cols + c] = 0.25f * (src[(r - 1) * cols + c] + src[(r + 1) * cols + c] +
                                 src[r * cols + c - 1] + src[r * cols + c + 1]);
}

/* One exchange step across two devices: compute into `nxt`, hand the
   new boundary rows across, and wait before the next step reads them. */
void step(float *cur[2], float *nxt[2], int rows, int cols, cudaStream_t s[2])
{
    cudaEvent_t ready[2];
    for (int k = 0; k < 2; k++) {
        cudaSetDevice(k);
        stencil<<<rows, cols, 0, s[k]>>>(cur[k], nxt[k], rows, cols);
        cudaEventCreateWithFlags(&ready[k], cudaEventDisableTiming);
        cudaEventRecord(ready[k], s[k]);   /* source row is now final */
    }

    /* Row `rows` of device 0 becomes device 1's new top halo; row 1 of
       device 1 becomes device 0's new bottom halo. Each peer copy is
       queued on the destination's stream, so an event has to order it
       behind the kernel on the source device. */
    cudaSetDevice(1);
    cudaStreamWaitEvent(s[1], ready[0], 0);
    cudaMemcpyPeerAsync(nxt[1], 1, nxt[0] + rows * cols, 0,
                        cols * sizeof(float), s[1]);
    cudaSetDevice(0);
    cudaStreamWaitEvent(s[0], ready[1], 0);
    cudaMemcpyPeerAsync(nxt[0] + (rows + 1) * cols, 0, nxt[1] + cols, 1,
                        cols * sizeof(float), s[0]);

    /* The next step must not start until both halos have landed. */
    for (int k = 0; k < 2; k++) {
        cudaSetDevice(k);
        cudaStreamSynchronize(s[k]);
    }
}

Why the slowest participant sets the cost

Every cross-device dependency is a join, and a join completes when its last input arrives. If device 0 finishes its share in 8 ms and device 1 in 12 ms, the join costs 12 ms, and device 0 has been idle for 4 ms. The idle time is not recovered anywhere; it is simply lost. Over many steps the loss accumulates in proportion to the imbalance, not to the total work.

This makes load balancing a correctness-of-performance issue rather than a nicety. A static partition that splits the data evenly is optimal only if the work per element is uniform. When it is not, as with a sparse structure of uneven density, a simulation where some regions become active and others go quiet, or a sort whose partitions diverge in cost, the even split produces a load imbalance that grows over time.

The usual responses are overdecomposition and work stealing. Overdecomposition splits the data into many more chunks than there are devices and assigns chunks dynamically as devices become free, so imbalance is bounded by the size of one chunk rather than by the size of one device’s share. Work stealing goes further, letting an idle device take queued work from a busy one, at the cost of the bookkeeping and the atomics needed to coordinate the queue. Both trade scheduling overhead for balance, and both are worth it when the imbalance is large relative to the chunk cost.

Synchronization frequency is the multiplier

The cost of a dependency is the latency of the mechanism used to enforce it, and that latency is on the order of microseconds for an event or a host-side wait. A computation that synchronizes once has paid a few microseconds. A computation that synchronizes once per iteration of a loop with a thousand iterations has paid a few milliseconds, which may be comparable to the compute.

The multiplier explains why strong scaling, meaning the same problem on more devices, saturates long before the devices are all busy. As the problem is split more finely, the work per device shrinks but the synchronization latency does not, so the synchronization fraction rises. The point at which it exceeds the compute fraction is the practical limit on the device count for that decomposition, and it moves only if the synchronization gets cheaper or the work per synchronization gets larger.

Batching is the standard lever. Instead of synchronizing after each element, accumulate a batch and synchronize after the batch. Instead of exchanging halos every step, exchange every k steps if the stencil’s dependency distance allows it, which it does when the stencil reaches only one cell, so k steps can proceed with a halo of width k. The trade is memory for latency: a wider halo costs more storage and more transfer per exchange, but the exchange happens less often.

Where the devices are, and why it matters twice

The physical topology affects both the transfer cost and the synchronization cost. A peer copy between devices under the same PCIe switch is cheaper than a copy that traverses the root complex, as described in peer-to-peer transfers and link topology. But topology also determines whether a host-side synchronization has to cross sockets, and on a two-socket machine a device on socket 0 waiting on a device on socket 1 pays the inter-socket latency on every join.

The mapping from logical partitions to physical devices should therefore be topology-aware. A pipeline whose stages have tight dependencies should place adjacent stages on devices with a short path between them, while independent stages can be spread for thermal or power reasons. The runtime’s topology query provides the information; using it requires deciding the partition before assigning devices, which is the opposite of the usual order of operations.

Memory capacity enters as a hard constraint rather than a soft one. A partition that does not fit in a device’s memory does not run, and the workarounds, managed memory with migration or peer access to another device’s memory, reintroduce the very costs the partition was meant to avoid. The capacity arithmetic that governs managed memory, covered in residency hints and thrashing under oversubscription, applies unchanged to a partition.

The accounting that decides whether to add a device

A partition is worth making when the time saved on compute exceeds the time added on coordination. Writing that out: the compute time for the whole problem is T, the fraction that parallelizes is p, and the coordination cost per step is c. Splitting across N devices gives roughly T(1−p) + Tp/N + cN in the optimistic case where coordination scales with the device count. The derivative with respect to N turns positive when c exceeds Tp/N², which is to say the optimum device count is finite and depends on the coordination cost per device.

The formula is a sketch rather than a prediction, because real coordination costs are not linear in N, since a reduction across N devices is a tree of depth log N and a broadcast is a fan-out, but the shape is right. Coordination costs that grow with the device count eventually dominate, and the crossover arrives sooner for fine-grained decompositions.

The practical version of the accounting is a measurement rather than a formula: run the same problem on one device and on two, and compare the wall-clock time. If two devices are not meaningfully faster, the partition’s boundaries are too expensive, and the next step is to widen them with fewer, larger exchanges rather than to add a third device.

Where the model breaks

The accounting above assumes devices are interchangeable. They are not. Devices of different generations or from different vendors have different memory capacities, different instruction sets, and different peak throughputs. A partition across heterogeneous devices must either split work in proportion to capability or accept that the weakest device sets the rate, and the programming interfaces for the two are frequently different enough that the kernel itself must be duplicated.

The model also assumes the host is idle during device work. In a pipeline with host-side stages such as I/O, decompression, and serialization, the host is a participant with its own latency, and a join that includes a host stage pays the host’s scheduling latency, which is larger and less predictable than the device’s. Systems that put the host in the loop need the host’s work to be asynchronous with respect to the devices, or the join costs hundreds of microseconds rather than tens.

Finally, the accounting ignores the cost of getting the data to the devices in the first place. For a computation that runs once on a large input, the transfer is a fixed cost that the partition does not change, and it may dominate. For a computation that runs many times on resident data, the transfer amortizes to nothing and the partition is judged on compute and coordination alone. Which regime a program is in determines whether the whole exercise is worth starting.

The ordering guarantees that make a join correct, as opposed to merely fast, are those of stream ordering and submission semantics, and the flag-based coordination that a peer topology makes possible without host involvement is described in system-scope atomics and memory ordering.