# T-CCL: Resource Efficient and Performant Collective Communication using Tensor Memory Accelerator

arXiv: https://arxiv.org/abs/2610.07098 · HTML: https://mlsanity.com/paper/2610.07098/

## Key facts

- No citing papers in the map yet (data through 14 Oct 2026).
- Builds on 12 of its 25 references that are in the map.
- First posted on arXiv 5 Oct 2026.

## Builds on (most cited first, 12 of 12)

- [Efficient Memory Management for Large Language Model Serving with PagedAttention](https://mlsanity.com/paper/2309.06180/) (2023)
- [Qwen2.5 Technical Report](https://mlsanity.com/paper/2412.15115/) (2024)
- [Splitwise: Efficient generative LLM inference using phase splitting](https://mlsanity.com/paper/2311.18677/) (2023)
- [NanoFlow: Towards Optimal Large Language Model Serving Throughput](https://mlsanity.com/paper/2408.12757/) (2024)
- [ThunderKittens: Simple, Fast, and Adorable AI Kernels](https://mlsanity.com/paper/2410.20399/) (2024)
- [Demystifying NCCL: An In-depth Analysis of GPU Communication Protocols and Algorithms](https://mlsanity.com/paper/2507.04786/) (2025)
- [T3: Transparent Tracking & Triggering for Fine-grained Overlap of Compute & Collectives](https://mlsanity.com/paper/2401.16677/) (2024)
- [ParallelKittens: Systematic and Practical Simplification of Multi-GPU AI Kernels](https://mlsanity.com/paper/2511.13940/) (2025)
- [MSCCL++: Rethinking GPU Communication Abstractions for AI Inference](https://mlsanity.com/paper/2504.09014/) (2025)
- [Efficient and Adaptable Overlapping for Computation and Communication via Signaling and Reordering](https://mlsanity.com/paper/2504.19519/) (2025)
- [GC3: An Optimizing Compiler for GPU Collective Communication](https://mlsanity.com/paper/2201.11840/) (2022)
- [An Efficient, Reliable and Observable Collective Communication Library in Large-scale GPU Training Clusters](https://mlsanity.com/paper/2510.00991/) (2025)

## Abstract (arXiv)

Large transformer-based models increasingly depend on multi-GPU execution, which requires frequent collective communication among GPUs. Existing communication libraries often rely on many GPU threads to achieve high bandwidth or low latency, resulting in a large streaming multiprocessor (SM)-side resource footprint. This footprint can limit the resources available to other GPU work, particularly when communication and computation execute concurrently. Thus, efficient collective communication should not only achieve high collective performance but also reduce its SM-side resource usage. This paper presents T-CCL, a resource-efficient collective communication library based on the Tensor Memory Accelerator (TMA) for intra-node communication. T-CCL offloads both data movement and reduction operations to TMA and executes each collective as a pipelined series of asynchronous TMA operations, reducing the SM resources required for collective communication while maintaining high bandwidth. Evaluated across AllReduce, AllGather, and ReduceScatter collectives, T-CCL outperforms NCCL by up to 2.4x with unrestricted communication resources and up to 3.42x under restricted resource budgets, remains competitive with NCCL's recent symmetric-memory kernels, and occupies the same or fewer SMs in profiled cases. In a GEMM-collective overlap case study, switching the communication backend from NCCL to T-CCL raises the average operator-level speedup over a sequential baseline from 1.12x to 1.25x on two GPUs and from 1.04x to 1.14x on four GPUs, as T-CCL uses fewer SMs for communication, leaving more SMs available to the overlapped GEMM. Integrated into vLLM as a communication backend, T-CCL improves end-to-end inference throughput over vLLM's automatic backend dispatch by up to 1.31x, outperforming it at every evaluated batch size on both the conversation and decode-heavy workloads.