CUCo: An Agentic Framework for Compute and Communication Co-design
Organizations: UT Austin
Abstract
Computation and communication in distributed LLM training and inference are traditionally optimized in isolation; expert-crafted systems such as DeepEP, FLUX, and TokenWeave show the potential of co-design but require deep systems expertise and hardware-specific tuning; CUCo is an agentic framework that automates compute-communication co-design of CUDA kernels by combining a structured design-space formalization with a correctness-first fast-path agent for reliable baselines and an evolution-driven slow-path agent for high-performance strategies, achieving up to 1.57x speedup across four multi-GPU workloads and discovering a two-stream overlap strategy on a DeepSeek-V3 MoE layer that hides dispatch behind local compute at an LLM inference cost under $10 per workload.
Figures & tables
| System | Multi-GPU | Comm Support | No Training | Device-Init |
|---|---|---|---|---|
| CUDAForge Zhang et al. (2025) | ✗ | ✗ | ✓ | ✗ |
| STARK Dong et al. (2025) | ✗ | ✗ | ✓ | ✗ |
| Kevin-32B Baronio et al. (2025) | ✗ | ✗ | ✗ | ✗ |
| KernelFalcon Wang et al. (2025) | ✗ | ✗ | ✓ | ✗ |
| CUCo | ✓ | ✓ | ✓ | ✓ |
| Concrete API Dimensions | Allowed Values |
|---|---|
| : Backend | GIN | LSA | Hybrid |
| : Sync Mechanism | Barrier | Signal | SignalShadow | Counter |
| : Issuer | Thread | Warp | WarpSpan | Tile⟨N⟩ | CTA |
| : Ordering | Relaxed | Acquire | Release | AcqRel |
| : Contexts | | | |
| Intent-Based Dimensions | Allowed Values |
| System | ||||||||
|---|---|---|---|---|---|---|---|---|
| DeepEP (NVL) | LSA | Barrier | Deferred | Local | CTA | Per-peer | Release | – |
| DeepEP (IB) | GIN | Signal | Deferred | World | CTA | Per-peer | Acquire | 1 |
| FLUX | LSA | Barrier | Tile-fused | Local | Warp | Per-tile | AcqRel | – |
| Workload | Dominant bottleneck | Discovered strategy | Key design choices |
|---|---|---|---|
| Flash Attention | Pipeline bubbles between ring exchange rounds | Persistent cooperative kernel with per-tile pipelining; comm block polls tile readiness via atomics | : tile-fused; : counter; : CTA (1 comm + 80 compute) |
| MoE Dispatch/ Combine | Sequential comm–compute phases; padding overhead on imbalanced ranks | Self/remote compute split with two-stream dispatch overlap; variable-size per-peer transfers | : GIN (inter-node); : stream-split; : per-peer |
| KV-Cache Transfer | Compute-to-send idle gap from host-launched sends | GPU-triggered chained send: K GEMM send K V GEMM send V with signal | : device; : release; : signal |
| GEMM + AllGather | Post-compute idle before collective | Immediate device-initiated broadcast fused after GEMM completion | : deferred; : LSA (intra) / GIN (inter) |
Appendix figures & tables3 assets
Supplementary material from the paper’s appendix.
Appendix
| Field | Content |
|---|---|
| Purpose | Device-side peer memory access and barrier synchronization within local GPU teams via direct load/store. |
| Core API | ncclLsaBarrierSession (provides synchronization), ncclGetLsaPointer (team-indexed access), ncclGetLocalPointer (local access), … |
| Sync Scope | ncclTeamLsa (intra-node local team), ncclTeamWorld (all ranks); tighter scope reduces overhead but limits overlap potential. |
| Invariants | Barrier index must be unique per CTA ( blockIdx.x ), … |
| Canonical Pattern | Instantiate ncclLsaBarrierSession with blockIdx.x ; perform memory operations; call sync(ncclCoopCta(), memory_order_release) ; access peers via ncclGetLsaPointer . |
| Workload | Fast-Path Iter. | Slow-Path Gen. | Wall Time | LLM Cost |
|---|---|---|---|---|
| Flash Attention | 3 | 20 | 1.8 hr | $10.96 |
| MoE (GIN) | 2 | 18 | 1.5 hr | $10.20 |
| KV-Cache | 2 | 18 | 1.5 hr | $9.74 |
| GEMM + AllGather | 2 | 10 | 51 min | $4.98 |
GPU Kernel Optimization
GPU kernel optimization improves the speed and memory efficiency of computations by tailoring code to GPU hardware and workloads.
See all 209 papers