【学习】 NVIDIA CFT（Compute Fabric Transport）特性分析

NVIDIA CFTCompute Fabric Transport特性分析版本2026-10 素材来源NVIDIA 官方文档NCCL 2.32.3 Device API / Usage Guide、CUDA Toolkit 13.4 官方博客、NCCL Release Notes、PyTorch SymmetricMemory 文档定位面向通信库/高级应用开发者的传输为中心transport-centric设备端数据移动机制0. 一句话概述CFTCompute Fabric Transport计算织物传输是 NVIDIA 从 NCCL 2.31 / CUDA 13.3 开始提供的设备端通信能力它把 CUDA fabric 的逻辑端点logical endpoint暴露给 GPU 内核让内核直接用 fabric 指令发起异步 put / get / 归约 / barrier操作跨 NVLink 织物fabric移动数据而无需把远端 GPU 内存映射进本地进程的虚拟地址空间。硬件要求BlackwellSM_100及更新架构软件要求CUDA Toolkit 13.3、驱动 ≥ 610.43.02、NCCL 2.31定位仅通过 CUDA Driver API 使用面向通信库开发者NCCL / NVSHMEM / PyTorch 等普通应用层开发者应使用 NCCL / NVSHMEM 等高阶库1. 背景与动机1.1 规模扩展带来的地址空间压力AI 工厂AI factory时代超节点规模持续扩大以 72-GPU Vera Rubin NVL72 为例单域 NVLink 提供每 GPU 3.6 TB/s 双向带宽、机架级 260 TB/s。GPU 数量越多每个进程都把参与通信的远端 GPU 分配映射进自己的虚拟地址空间的传统做法就越难以为继虚拟地址VA压力多 GPU 多进程 大显存HBM3e 上百 GB/卡peer 映射 / IPC 映射数量随 rank 数平方增长64 位 VA 空间也会被吃紧映射管理开销每次分配/映射、IPC 句柄交换都需要 host 参与且映射粒度粗、无法精细控制通信与计算分离传统 NCCL 集合通信是独立的 kernel/库调用难以在自定义内核内部直接算完就发计算-通信重叠依赖库层调度。1.2 硬件前提NVLink NVSwitch NVLink SHARPNVLink 是 NVIDIA 专为 AI 工厂打造的 scale-up 网络第 4 代 NVSwitchHopper开始内置SHARPScalable Hierarchical Aggregation and Reduction Protocol引擎即NVLink SHARPNVLSGPU 可发出 multimem PTX load/store 指令访问多播地址交换机在转发的同时完成数据复制与归约Blackwell 延续并强化了该能力multimem 指令、多播归约、fabric 指令集CFT 正是建立在fabric 指令 logical endpoint这套硬件抽象之上。1.3 软件演进路线传统 IPC / peer mapping映射每个远端分配 ↓ NVLS / SymmetricMemory window多播 交换机内归约仍需 window 地址空间 ↓ NCCL Device APIncclDevComm内核内集合通信原语 ↓ CFTtransport-centriclogical endpoint fabric 指令内核直接 put/get/red/barrier2. 核心原理2.1 transport-centric 模型CFT 的核心转变是从内存映射为中心到传输为中心维度传统方式映射为中心CFT传输为中心寻址远端内存映射成本地 VA 指针命名逻辑端点(le_id, le_offset)数据移动load/store / 拷贝fabric 指令异步 put/get/red通信方host 编排 kernel 参与内核直接发起无需 host 逐次介入地址空间每个远端分配占一份 VA不占本地 VA压力大幅下降模式单播为主单播 多播multicast2.2 三个关键抽象对称内存窗口Symmetric Memory Window参与通信的 rank 各自分配、大小一致、对齐一致的缓冲区注册为 window 后所有 rank 看到的是同一份对称地址布局——这是 CFT 寻址的基础逻辑端点Logical EndpointCFT 数据移动操作都作用在端点 ID 端点内偏移上。(le_id, le_offset)是不透明对由 NCCL 提供 host/device 两侧的查询 API把(window, offset, peer)翻译成(le_id, le_offset)CFT Team用 team单播组 / 多播组而不是 world rank 来寻址对端支持构建层级化通信模式flat / hierarchical。2.3 操作原语操作方向说明put/putMultimem本地 smem → 远端异步写multimem 为多播变体get远端 → 本地 smem异步读red/redMultimem本地 smem → 远端归约归约在 fabric 路径上完成复用 SHARP 能力pullRed多播端点 → 本地 smem拉取 归约需 warp 级 coopbarrier全 team跨 CFT team 的设备线程同步putCounted/redCounted/waitCounted带计数器去掉数据 fence flag模式用用户管理的计数器做完成通知降低同步延迟归约操作类型Sum/And/Xor/Or/Min/Max具体数据类型的支持以 CFT PTX 文档为准支持unicast 与 multicast两种端点multicast 需在创建 device communicator 时请求NCCL_CFT_MULTIMEM完成与错误上报flush可返回错误报告cudaFabricOpErrorStatusGet/Count解码应用可据此检测、重试或改道失败的 fabric 传输。2.4 与既有方案的关系vs NVLSNVLink SHARPNVLS 是 NCCL 内基于 symmetric window multimem PTX 的集合通信实现CFT 更底层直接暴露 logical endpoint 与设备端操作 API可被 NVLS 以外的自定义内核使用vs IPC / peer mappingCFT 不需要远端内存映射进本地 VA解决大规模多 GPU 的 VA 压力vs NVSHMEMNVSHMEM 是更上层的 PGAS 编程模型CFT 是其可依赖的更底层传输原语NCCL 2.31.2 已提供 nccl4rust / NIIN 等 NVSHMEM 兼容接口构建在 Device API 原语之上。3. 如何实现API 流程3.1 前置条件硬件Blackwell 或更新SM_100所有参与 rank 都能创建兼容的 CUDA fabric logical endpoints软件CUDA Toolkit 13.3编译 fabric PTX 指令、驱动 ≥ 610.43.02、NCCL 2.31内核必须为支持 fabric 指令的架构编译。关键约束如果某个 communicator 上并非所有 rank 都支持 CFTncclDevCommCreate()直接失败全有或全无。3.2 Host 端设置流程// 1) 分配并注册对称内存窗口ncclMemAlloc(buffer,memSize);ncclCommWindowRegister(comm,buffer,memSize,win,NCCL_WIN_COLL_SYMMETRIC);// 可加 NCCL_WIN_CFT_COUNTED// 2) 创建带 CFT 能力的 device communicatorncclDevComm devComm;ncclDevCommRequirements reqsNCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;reqs.cftCapsNCCL_CFT|NCCL_CFT_MULTIMEM;// 单播 多播reqs.cftBarrierCountnCTAs;// barrier slot 数每 CTA 一个NCCLCHECK(ncclDevCommCreate(comm,reqs,devComm));3.3 Team 与 Endpoint 查询// device 侧取单播 CFT team并按模式过滤FLAT / HIER_MULTIMEM / HIER_LSAncclTeam cftTeamncclTeamCft(devComm,NCCL_CFT_TEAM_FLAT);// 把 (window, offset, peer) 翻译成逻辑端点ncclCftLeId leId;size_tleOffset;ncclGetCftLeInfo(win,byteOffset,peerCft,cftTeam,devComm,leId,leOffset);host 侧对应ncclGetCftDeviceLeInfo()/ncclGetPeerDeviceLeInfo()/ncclGetMultimemDeviceLeInfo()若需在创建 CFT device communicator之前查询端点可在初始化 communicator 时配置hostCftMode。3.4 内核内操作典型 put 流程stage 数据到共享内存 (smem, 16B 对齐) → 构造 ncclCftncclCoopCta 对象共享 ncclCftSmem 状态 → cft.put(coop, leId, leOffset, smem, bytes) → cft.submit(coop) // 提交使操作开始推进 → cft.flushSmem(coop) // 等待 smem 被消费后可复用更早的回收点 → cft.flush(coop) // 等待所有操作完成含错误报告__global__ void cftPutKernel(ncclDevComm devComm, ncclWindow_t win) { __shared__ ncclCftSmem cftSmem; __shared__ alignas(16) char smem[128]; ncclCoopCta coop; ncclCftncclCoopCta cft{coop, cftSmem}; ncclTeam cftTeam ncclTeamCft(devComm); int peer (cftTeam.rank 1) % cftTeam.nRanks; ncclCftLeId leId; size_t leOffset; ncclGetCftLeInfo(win, 0, peer, cftTeam, devComm, leId, leOffset); cft.put(coop, leId, leOffset, smem, sizeof(smem)); cft.submit(coop); cft.flush(coop); }3.5 Barrier__global__ void cftBarrierKernel(ncclDevComm devComm) { ncclCoopCta coop; ncclCftBarrierSessionncclCoopCta bar{coop, devComm, blockIdx.x}; bar.sync(coop, cuda::memory_order_acq_rel, ncclMemProxyType::Generic, ncclMemProxyType::Fabric); }多播 barrier 需在创建 communicator 时带NCCL_CFT_MULTIMEM并传multimemtrue。3.6 Counted 操作延迟优化Counted 操作把传统的数据 内存 fence flag 更新同步模式替换为用户管理的计数器窗口用NCCL_WIN_CFT_COUNTED注册计数器的更新被硬件保证排在数据落盘之后发送方cft.putCounted(coop, leId, dataOffset, counterOffset, smem, bytes)接收方cft.waitCounted(coop, acquire, Generic, counterPtr, expectedBytes, nullptr)poll 计数器达到期望值即返回waitCounted封装了协作线程同步、内存序、abort/timeout目前 abort 传nullptr。3.7 内存序Cross-proxy Fence数据与 flag 之间需要建立 happens-before 关系时使用跨 proxy 的 fence// 发送数据后用 Fabric→Generic 的 release fence 保证 flag 在数据之后可见 ncclMemFence(coop, cuda::memory_order_release, ncclMemProxyType::Fabric, ncclMemProxyType::Generic, ncclMemFenceScope::Sys); cft.put(coop, leId, leOffset flagOffset, flag, 16); cft.submit(coop); cft.flush(coop);另一个典型场景Generic proxy 写入 smem 后用 fence 让 Fabric proxy 可见再发起 put。4. 已知问题与限制类别限制硬件/软件门槛仅 BlackwellSM_100、CUDA 13.3、驱动 ≥ 610.43.02、NCCL 2.31内核需为支持 fabric 指令的架构编译全有或全无任一 rank 不支持 CFTncclDevCommCreate()即失败无法降级对齐/尺寸smem 源/目标必须 16B 对齐传输字节数必须是 16 的倍数中转要求数据必须经共享内存中转smemSource/smemDestination不适合直接 global→global 的原生路径in-flight 上限单个ncclCftSmem最多跟踪 16MB 非 fetching 操作put/red 及多播变体、1MB fetching 操作get/pullRedcounter 约束每 GPU 并发活跃 counter ≤ 32 个为最佳性能counter 建议 256B 对齐同步易错submit / flushSmem / flush、cross-proxy fence、counted 的 memory order 全部需要手工正确编排错误使用会导致数据竞争或悬挂barrier 配额cftBarrierCount必须覆盖内核用到的全部 barrier index通常每 CTA 一个生态成熟度新特性2026 年随 CUDA 13.4 正式宣传官方示例目前仅 barrier 示例文档/工具链仍在完善使用面仅 Driver API、面向通信库作者普通应用无直接收益应走 NCCL/NVSHMEMabort 支持waitCounted的 abortFlag 目前需传nullptr无运行时中止路径归约类型受 CFT PTX 指令集约束Sum/And/Xor/Or/Min/Max类型覆盖以 PTX 文档为准5. 生态与展望NCCL 2.31.2新增 CFT host/device API支持 window 内存注册到 CUDA logical endpoints以及 device 侧 Put / Get / Red / NVLS 操作PyTorch SymmetricMemory已把 CFT 作为底层传输之一——CFT 把 peer 的 window 注册内存暴露为(le_id, le_offset)自定义内核可直接用它发起 put/get/reduce无需构造nccl内核CUDA Toolkit 13.4正式把 CFT 列为新特性与 locality domain、CDMM、Rubin 预览等并列并明确reports completion and error status so applications can detect, retry, or reroute failed fabric transfers方向CFT 是 NVIDIAGPU 即网络端点路线的底层基石之一未来 Rubincompute capability 107与 NVLink 6 上logical endpoint fabric 指令这套抽象预计会成为 in-kernel 通信、计算-通信重叠、集合通信卸载的主要通道。6. 参考链接NCCL Device API - CFThttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/api/device_cft.htmlNCCL Usage - Compute Fabric Transporthttps://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/cft.htmlNCCL Release 2.31.2 Noteshttps://docs.nvidia.com/deeplearning/nccl/release-notes/rel_2-31-2.htmlCUDA Toolkit 13.4 官方博客CFT 章节https://developer.nvidia.com/blog/?p121255NVIDIA NVLink: The Scale-Up Network for AI Factorieshttps://developer.nvidia.com/blog/nvidia-nvlink-the-scale-up-network-for-ai-factories/PyTorch Symmetric MemoryCFT 说明https://docs.pytorch.org/docs/main/symmetric_memory.htmlNVLS / NVLink SHARP 背景TokenWeavearXivhttps://arxiv.org/html/2505.11329v1/本文为技术分析文档所有功能描述与版本信息以 NVIDIA 官方文档为准未声明处不构成对 CFT 可用性、性能或兼容性的承诺。