mlsanityData to 14 Oct 2026

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

arXiv 2610.07098 ↗First posted 5 Oct 2026cs.DC cs.AIcited by 0

Not on the map. Why?

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.

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

12 of 25 references.

References · 12

  1. 2023Efficient Memory Management for Large Language Model Serving with PagedAttentioncited by 5,944
  2. 2024Qwen2.5 Technical Reportcited by 4,942
  3. 2023Splitwise: Efficient generative LLM inference using phase splittingcited by 558
  4. 2024NanoFlow: Towards Optimal Large Language Model Serving Throughputcited by 90
  5. 2024ThunderKittens: Simple, Fast, and Adorable AI Kernelscited by 36
  6. 2025Demystifying NCCL: An In-depth Analysis of GPU Communication Protocols and Algorithmscited by 20
  7. 2024T3: Transparent Tracking & Triggering for Fine-grained Overlap of Compute & Collectivescited by 16
  8. 2025ParallelKittens: Systematic and Practical Simplification of Multi-GPU AI Kernelscited by 12
  9. 2025MSCCL++: Rethinking GPU Communication Abstractions for AI Inferencecited by 10
  10. 2025Efficient and Adaptable Overlapping for Computation and Communication via Signaling and Reorderingcited by 10
  11. 2022GC3: An Optimizing Compiler for GPU Collective Communicationcited by 7
  12. 2025An Efficient, Reliable and Observable Collective Communication Library in Large-scale GPU Training Clusterscited by 4

Built on by

Not cited yet.