版本v0.9草案状态开放设计草案 / FRAND 授权意向定位封装内双计算 Die 一致性缝合互连对标 NVIDIA NV-HBI发布日期2026-09-30免责与设计声明NV-HBI 为 NVIDIA 私有协议其内部实现细节未公开。本规范以公开数据为锚点NV-HBI 10 TB/s 双向带宽、UCIe 3.0 PHY 参数、AMBA CHI C2C / CXL 一致性工程经验、MCM-GPU 学术研究是一套工程上自洽的开放规范设计不代表任何已流片产品。目录设计目标与适用范围术语与缩写总体架构与分层CoreX-P物理层CoreX-L链路层CoreX-C一致性协议层CoreX-M内存模型映射性能信封与 NV-HBI 对标RAS 与可管理性功耗管理验证、合规与 KGD 流程开放性与授权与 NV-HBI 的诚实差距声明路线图参考文献与依据1. 设计目标与适用范围1.1 目标项目指标备注定位同一封装内两颗对等计算 Die的缓存一致性缝合互连点对点、双节点总带宽≥ 10 TB/s 双向对标 Blackwell NV-HBI 公开值单次跨 Die 事务附加延迟≤ 10 ns含 PHY 链路 协议处理PHY TxRx 延迟≤ 2 ns对齐 UCIe 3.0 指标能效0.5–0.75 pJ/bit对齐 UCIe 3.0 64 GT/s 档语义目标两 Die 呈现为单一设备统一地址空间、统一逻辑 L2、硬件一致性软件零感知空载功耗较满载降 ≥ 85%动态 lane 关断1.2 非目标明确的边界声明不设计多 Die2路由与目录广播——留给 v1.x双节点约束是本规范省电、省面积、低延迟的前提不定义封装外一致性传播——sys 作用域操作越过本封装后的传播PCIe/C2C/NVLink 域由对应接口规范负责不规定对端 Die 的内部微架构——本规范只管接口上的事务语义不管 SM/L1/L2 内部怎么实现不追求 CPU 式全层级一致性——GPU 的一致性点在 L2L1 不参与一致性状态机见 6.2。1.3 设计哲学三条贯穿全文的原则单点串行化优于目录广播每个地址有且仅有一个 home slice天然全序免广播、免监听均匀化即最优数据局部性不可调度时用地址哈希把跨 Die 流量统计均匀化调度器零复杂度为 GPU 语义雕刻事务集不为通用性付税——事务集直接对应 CUDA/PTX 内存操作不做 CPU 式 MESI 的全集。2. 术语与缩写术语含义Home Slice某地址的唯一管辖 L2 分片一致性串行化点Peer Die对端计算 DieFar Atomic在 home slice 的 ALU 执行的远端原子读-改-写Flit链路层定长传输单元64 BSectorL2 缓存行子粒度32 BVC虚拟通道Virtual ChannelDSM分布式共享内存cluster 内 block 间互访TMATensor Memory Accelerator异步批量拷贝引擎KGDKnown Good Die已知良好裸片ShorelineDie 边缘可用于 PHY 布线的海岸线长度FRAND公平、合理、无歧视授权3. 总体架构与分层┌─────────────────────────────────────────────────────┐ │ CoreX-M 内存模型映射CUDA/PTX → 事务 强制映射表 │ ├─────────────────────────────────────────────────────┤ │ CoreX-C 一致性协议层 │ │ 地址治理 / Home Slice / 三态一致性 / │ │ Far Atomic / Fence / DSM / TMA / Barrier │ ├─────────────────────────────────────────────────────┤ │ CoreX-L 链路层 │ │ 64B Flit / 4×VC / CRC-32重传 / 边带 │ ├─────────────────────────────────────────────────────┤ │ CoreX-P 物理层 │ │ 64 GT/s NRZ 单端 / 45µm 凸点 / 硅桥走线 │ └─────────────────────────────────────────────────────┘ ↕ 封装内局部硅桥LSI走线 ≤ 2 mm ┌─────────────────────────────────────────────────────┐ │ Peer Die对称镜像栈 │ └─────────────────────────────────────────────────────┘分层间的可替换性声明CoreX-P 可直接采用商用 UCIe 3.0 PHY IP先进封装档CoreX-L 与 UCIe 链路层在 flit 概念上兼容但不逐比特兼容CoreX-C/M 为本规范独有无任何现成开放协议可替代——这是实现方的主要工作量所在。4. CoreX-P物理层4.1 电气与封装参数参数取值依据信令制式NRZ 单端64 GT/sUCIe 3.0 最高档调制/编码无前向纠错必选可选轻量 FEC 留给 v1.x延迟优先凸点间距45 µm先进封装档允许 25–55 µmUCIe-A 范围通道走线距离≤ 2 mm位于局部硅桥LSI之下与 CoWoS-L 同构避免整板硅中介层成本单 lane 单向带宽8 GB/s64 Gbps目标 lane 数≥ 625 lane/方向典型配置 700含冗余5 TB/s 单向 ÷ 8 GB/s海岸线预算双向合计 ≥ 5 mm64G 下带宽密度约 2600 GB/s/mm5 TB/s 需约 2 mm留 ≥25% 冗余BER≤ 10⁻¹²UCIe 3.0 64G 目标PHY 延迟TxRx ≤ 2 ns目标值能效0.5–0.75 pJ/bitUCIe 3.0 64G 档4.2 Lane 组织与备援Lane 按簇Cluster组织每簇 64 lane 2 条备援 lane支持降级运行单簇内坏 lane ≤ 2 时由备援顶替超过时整簇降速运行48 GT/s 模式训练序列上电时逐簇训练deskew、眼图居中训练结果存入本地寄存器供运行时可重校准复用。4.3 时钟与初始化每方向独立 forwarded clock随路时钟源同步初始化顺序边带握手 → 时钟锁定 → 逐簇训练 → 链路层 credit 交换 → 协议层上线冷启动到协议可用目标 ≤ 10 ms。4.4 边带通道独立低速边带open-drain双 Lane 冗余承载训练/重校准控制遥测误码计数、眼图裕量、温度联动紧急关断Emergency Shutdown与快速节流Fast Throttle——热失控场景下确定性低延迟信令不占用主数据通道。5. CoreX-L链路层5.1 Flit 格式定长64 B512 bitFlit63 0 ┌──────────┬────────────┬──────────────┬──────────────┐ │ Header │ Payload │ Payload │ CRC-32 / │ │ (4 B) │ A (28 B) │ B (28 B) │ Ctrl (4 B) │ └──────────┴────────────┴──────────────┴──────────────┘Header4 B位级定义位段字段说明[31:29]VC虚拟通道号0–3[28:24]OpCode事务操作码见 6.4[23:22]Scope内存作用域cta/cluster/gpu/sys[21:18]FlitSeq多 flit 事务的序号128 B 行 2 flit[17:12]SrcID源 slice/SM 标识[11:0]Credit/Tag事务标签乱序返回匹配5.2 虚拟通道与流控VC名称承载设计理由VC0Request读/写/原子请求与响应分离防请求-响应循环依赖VC1Response数据返回、写确认数据体量最大独立 credit 池VC2Coherence无效化、确认、FenceToken一致性维护不被数据流阻塞VC3Latency-CriticalDSM 消息、barrier、TMA DoneToken32 B 小消息最高优先级仲裁每条 VC 独立 credit 流控接收方按 flit 粒度发放 credit死锁避免VC 间无依赖环路请求不等待请求协议层强制「响应可在任意请求未完成时发出」仲裁VC3 VC2 VC0 VC1严格优先级 防饥饿老化计数器。5.3 可靠性每 flit CRC-32接收方检错 → NACK → 发送方从 replay buffer 重传Replay buffer 深度 64 flit覆盖 ≤ 2 个往返延迟的带宽积可选Raw 模式旁路 CRC/重传仅用于调试与一致性协议形式化验证。5.4 链路功耗状态状态行为进入/退出延迟L0全速—L0p按簇关断 lane带宽减半/再减半 20 ns无训练L1全链路休眠边带保活退出 500 ns重校准复用训练结果6. CoreX-C一致性协议层本层是 CoreX-HBI 的「灵魂」也是对标 NV-HBI 的真正难点。PHY 可以买 IP本层必须自己实现。6.1 地址治理缓存行128 B扇区32 Bsector 填充粒度对齐 GPU L2 行为Home 分配默认按 128 B 行粒度做 XOR 哈希分治到两颗 Die 的 L2 home slicehome_die XOR(addr[47:7] 的若干折叠位) 0x1可配置4 KB 页粒度模式模式 P供局部性敏感负载对比实验运行时切换需 flush 全链路分配器约束显存分配器强制页交错分布到全部 HBM 堆栈保证跨 Die 流量统计均匀均匀化即最优原则。6.2 一致性模型L2 唯一一致点 L1 作用域失效GPU 不需要 CPU 式全层级 MESI本规范明确L2 home slice 是唯一一致点所有写与原子操作在此串行化L1 不写分配、全局写穿透不持有一致性状态——不存在 L1↔L1 监听L1 的失效由作用域化失效事务触发InvAll(scope)而非逐行监听行状态精简为三态状态含义EDExclusive-Dirtyhome 持有最新值HBM 为旧值对端无副本SCShared-Clean与 HBM 一致对端 L2 可能持有只读副本模式 BIInvalid本 slice 不持有有效副本6.3 Home Slice 目录模式 B 用每 home slice 配粗粒度目录按行记录「对端 Die 是否持有只读副本」1 bit/行目录查找与 L2 tag 查找并行不增加关键路径目录溢出策略溢出即失效对端副本保守正确性能损失有界。6.4 事务集OpCode事务语义对应 GPU 场景0x00RdS读共享返回数据状态→SCSM 全局加载0x01RdO读独占返回数据状态→ED即将写的加载0x02WrFull整行写128 B全局存储整行0x03WrPtl字节掩码部分写掩码精度到字节分散全局存储0x04AtomRMWFar Atomichome slice ALU 执行读-改-写返回旧值atomicAdd/atomicCAS等0x05FenceToken屏障令牌排空指定作用域在途事务并回执__threadfence()系0x06FlushWB写回所有 ED 行至 HBMkernel 边界/DMA 前0x07InvAll按作用域失效本端/对端 L1 及模式 B 副本上下文切换、显存重映射0x08BulkCopy异步批量拷贝描述符源/目的/长度/strideTMA 跨 Die 拷贝0x09DoneToken异步操作完成通知TMA/barrier 完成0x0ABarrierArr/ 0x0BBarrierWait硬件屏障网络卸载cooperative launch 全局 barrier0x0CDSMMsg32 B 低延迟簇内消息VC3cluster 分布式共享内存0x0D–0x1F保留—扩展空间6.5 Far Atomic 子规范执行位置数据 home slice 的原子 ALU单点串行化的物理体现操作集add / min / max / and / or / xor / exch / CAS位宽与类型4/8/16 B 整数f32/f64/bf16 浮点浮点原子按 IEEE 754 语义非结合性由单点串行保证全局一致顺序吞吐目标≥ 1 op/clk/slice 2 GHz同地址原子严格全序作用域gpu 作用域原子仅在双 Die 域内定序sys 作用域原子在 home slice 额外触发外域回执外域传播协议不在本规范内。6.6 定序规则形式化同地址规则同一 128 B 行的一切操作在 home slice 严格全序宽松序缺省异地址操作无定序承诺除非被 fence 约束Fence 排水规则FenceToken(scope)从发起方出发沿途排空该线程按 SrcID 追踪在 scope 内的全部在途事务到达 home slice 完成确认后原路回执回执到达发起方 fence 语义成立作用域层级cta ⊂ cluster ⊂ gpu ⊂ sys低作用域 fence 不得排空高作用域事务性能保护高作用域 fence 必须包含低作用域Acquire/Release 映射release 写 WrPtl FenceToken(scope)前置排水acquire 读 RdS 本地依赖栅栏禁止后续访存越过该读的返回点。6.7 NUMA 缓解双模式模式 A默认模式 BPeer Caching远端读每次穿 CoreX-HBI 取数只读行可缓存于本端 L2目录不需要home 目录追踪对端副本写路径直写 home写时单点失效对端副本适用流式/一次性数据重复读的共享只读数据权重、常量软件可见性无无运行时按页统计自动切换6.8 典型事务时序跨 Die 读模式 AL2 missDie0 SM → Die0 L2(miss, 非本端 home) → [CoreX-C RdS] → CoreX-L(VC0) → CoreX-P → Die1 L2 home slice(命中 SC/ED 或读 HBM) → [数据] → CoreX-L(VC1) → Die0 L2 → Die0 SM 附加延迟预算PHY 2ns 链路 2ns 协议处理 ≤ 6ns ≈ 总计 ≤ 10ns跨 Die 原子Far AtomicDie0 SM atomicAdd(addr home 在 Die1) → CoreX-C AtomRMW(VC0) → Die1 home slice ALU 执行 → 返回旧值(VC1) 定序确认(VC2) → Die0 SMgpu 作用域 fence线程执行 fence.acq_rel.gpu → FenceToken(gpu) 广播至两端所有在途 home slice → 各 slice 排空该 SrcID 在途事务 → 回执汇聚 → 全部回执到达 → fence 完成后续访存放行7. CoreX-M内存模型映射规范附带强制映射表PTX 8.x 参考实现实现方按自家 ISA 调整但必须通过等价性验证PTX/CUDA 操作CoreX-C 事务序列ld.relaxed.gpuRdSld.acquire.gpuRdS 本地依赖栅栏st.relaxed.gpuWrPtlst.release.gpuFenceToken(gpu)→WrPtlred.relaxed.gpu.addAtomRMW(add, gpu)无需返回值atom.acq_rel.sys.casFenceToken(sys)→AtomRMW(cas, sys) 外域回执fence.sc.gpuFenceToken(gpu)全量排水__threadfence_block()无 CoreX 事务CTA 内本地处理cluster.barrierBarrierArr/Waitcluster 作用域cp.async.bulkTMABulkCopyDoneTokenkernel 边界隐式同步FlushWBInvAll(gpu)等价性验证要求实现方须对映射表跑 PTX 内存模型 litmus 测试集含 message-passing、store-buffering、read-to-write causality 等经典模式全部通过方可宣称 CoreX-M 合规。8. 性能信封与 NV-HBI 对标指标NV-HBI公开值/反推CoreX-HBI v0.9 目标达成手段双向带宽10 TB/s10 TB/s约 8 TB/s 有效载荷 协议开销冗余64 GT/s × ~700 lane/向PHY 延迟未公开≤ 2 ns TxRxUCIe 3.0 档指标跨 Die 读附加延迟未公开NUMA 惩罚 1–3% 反推≤ 10 ns精简 flit、单点串行免广播能效未公开0.5–0.75 pJ/bitUCIe 3.0 64G 档满载 PHY 功耗未公开~40–60 W80 Tbps × 0.6 pJ/bit 估算均值远低于此L0p/L1一致性语义CUDA 内存模型CoreX-M 映射表 litmus 验证第 7 节封装CoWoS-L LSI同构硅桥 ≤ 2 mm 走线第 4 节9. RAS 与可管理性检错纠错flit 级 CRC-32 滑动窗口重传PHY 可选 FEC 留给 v1.xLane 备援每簇 2 条 spare超备援能力时整簇降速64G→48G可靠性目标FIT 1遥测边带承载误码计数、眼图裕量、每 VC credit 水位、每 slice 跨 Die 流量统计供 NUMA 模式切换决策调试flit 级追踪口Raw 模式旁路一致性协议状态机快照导出紧急关断边带 Emergency Shutdown确定性延迟 100 ns。10. 功耗管理L0p 动态宽度缩放按 VC credit 水位预测流量逐簇关断/唤醒 20 ns 无训练切换L1 全链路休眠空闲超时进入边带保活退出 500 ns运行时可重校准复用冷启动训练结果空载功耗目标较满载降 ≥ 85%热联动Die 温度传感器 → 边带 Fast Throttle → 对端降速防止封装热点。11. 验证、合规与 KGD 流程11.1 分层验证层验证手段CoreX-PPHY BIST、眼图/BER 扫描、环回测试CoreX-L链路环回、CRC 注入测试、credit 压力测试CoreX-C形式化验证定序规则模型检验 协议状态机 fuzzCoreX-MPTX litmus 测试集强制见 7系统级双 Die FPGA/仿真原型跑真实 kernel 负载11.2 KGD 流程封装前晶圆测试 → PHY BIST → 链路环回 → 一致性协议自检内置 litmus 微引擎 → 分级full/lane-degraded→ KGD 出厂11.3 合规性测试套件金色参考 RTL事务级模型作为仲裁基准互操作 plugfest不同实现方的 PHY/链路层对插合规标识CoreX-HBI-ReadyPHY链路/ CoreX-HBI-Coherent含 CoreX-C/M。12. 开放性与授权规范本体CoreX-P/L/C/M开放授权FRAND金色参考 RTL 与 litmus 套件开源Apache 2.0 意向PHY 实现可自由选用商用 UCIe 3.0 IP 或自研明确声明本规范只解决双 Die 缝合多 Die2需引入路由与目录广播属 v1.x 范围。13. 与 NV-HBI 的诚实差距声明本设计能在带宽、延迟、一致性功能三个维度对标 NV-HBI但有三处无法真正 1:1协同设计深度NV-HBI 与 NVIDIA 的 L2 仲裁粒度、TMA、DSM 是同一支团队联合雕刻的开放规范只能规定接口语义规定不了对端 Die 内部的微架构配合度能效尾部标准 PHY 的合规电路训练、CRC、边带、备援带来约 10–20% 的固定功耗与面积税验证成熟度NV-HBI 已经历两代量产Blackwell、RubinCoreX-HBI 从 v0.9 到流片验证最乐观估计 24–30 个月。现实定位本规范的用途不是「替代 NVIDIA」而是给任何想做双 Die GPU 的厂商国产 GPU、AI ASIC 创业公司一张可直接开工的图纸——其中 CoreX-P/L 可直接购买 UCIe 3.0 IP 实现真正需要自己写的是 CoreX-C 的事务引擎与 CoreX-M 的映射表。14. 路线图版本内容目标v0.9本版双 Die 缝合全集设计冻结评审v1.0补形式化证明附录、参考 RTL、litmus 套件首个 FPGA 原型v1.5可选 PHY FEC、模式 B 目录细化、功耗状态细化流片验证v2.0多 Die≤4扩展路由 分级目录对标下一代多芯粒 GPU15. 参考文献与依据NVIDIA Blackwell 架构公开资料NV-HBI 10 TB/s 双向带宽NVIDIA Rubin 架构公开资料双计算 Die NV-HBI HBM4UCIe 3.0 规范公开摘要48/64 GT/s、45µm 凸点、0.5–0.75 pJ/bit、运行时可重校准、边带紧急关断Arm AMBA CHI C2C 规范IHI0098与 Neoverse CSS 芯粒实践MCM-GPU 一致性研究HMG、CPElide 等学术协议PTX ISA 内存一致性模型文档作用域与 fence 语义OCP BoW / OpenHBI 公开规范PHY 层对比基线CoreX-HBI v0.9 — 为双计算 Die 缝合而生的开放一致性互连规范。