SelectNDStrategy — ND 集体算法选择器
此页面上的所有地址适用于
libtpu-0.0.40-cp314轮子中的libtpu.so(构建 ID89edbbe81c5b328a958fe628a9f2207d)。其他版本会有所不同。.textVMA等于文件偏移量;所有地址均为 VMA。
摘要
xla::jellyfish::BaseStrategyND::SelectNDStrategy(0x137c78e0、0x7cf = 1999 字节)是决定将 ICI all-reduce/reduce-scatter/all-gather 编译成哪种集体环算法的单个函数。给定切片拓扑(1-D / 2-D / 3-D ICI 环面)、HloInstruction 和一堆 TpuCompilationEnvironment 标志字节,它会遍历固定决策树并准确构造五个终端策略对象之一:子平面子组策略、通用 ND 环 StrategyND(其本身解析为 1-D 或构建时的 N-D 单向环)、N 路模型并行环、扭曲圆环环或 跨步 ND 环。这里没有成本比较搜索——选择器是一系列谓词门;所选环的“成本”在链路计数模型中单独计算(参见 SPMD 链路计数成本),并且环遍历的“邻居调度”在 StrategyND::BuildStrategy 中的下游构建。
分层在拾取器顶部的是降级轴容错重新映射:当一个环面轴具有部分故障的 ICI 链路时,弹性路径会将该轴折叠到主环之外。 GetDegradedAxis (0x1c894c20) 将每个轴的三个“此维度已降级”标志减少为单个降级轴索引(或 -1 =“无法隔离,放弃”);然后,InitColorDimensionsDegraded (0x137c6580) 重写 [6][3] 每颜色环尺寸表,以便死轴降级到最里面的环尺寸,并且两个正常轴承载环。 GetResourceFromIciResource (0x1c894c00) 将生成的环尺寸映射到成本模型存入周期的 ICI ResourceVector 插槽上。
了解 MPI 环/递归加倍集体的读者拥有该框架:这是 TPU 在矩形环面上的每轴算法选择器,加上单轴容错逃生舱口。重新实施合同是:
- 决策树。 两个顶级路径在
enable_sub_plane参数上分割,并且在每个有序保护谓词(TpuCompEnv 标志字节、操作码测试、单 ND 平面测试、拓扑形状探测、跨模块折叠)中在五个终端类中进行选择。 - 策略枚举。 五个终端类、它们的
operator new对象大小、它们的构造函数以及StrategyNDctor 的区别参数([obj+0xa8]1-D-vs-ND-ring 选择器由BuildStrategy使用)。 - 降级轴折叠。 每轴降级谓词 (
IsAxisDegraded && extent≥2 && usable)、≥2 ⇒ -1缩减以及将死轴推至内环维度的[6][3]颜色维度重新映射。
| 选择器条目 | xla::jellyfish::BaseStrategyND::SelectNDStrategy @ 0x137c78e0 (0x7cf = 1999 B) |
| 单ND平面测试 | xla::jellyfish::ReplicaGroupsOnNDPlane(…, plane=2, false) @ 0x1c890960 |
| 二维平面门 | BaseStrategyND::IsGroupNDPlane @ 0x137c6700(NumNetworkDimensions>=3,NDtopo 调暗不全为 1,env[0x1015]==1) |
| N 向门 | BaseStrategyND::UseSpecialStrategyNDNWay @ 0x137c6be0(单切片,comp-count ∈ {2,4}) |
| 跨步门 | BaseStrategyND::UseStridedStrategyND @ 0x137c72e0(单片、NumNetworkDimensions>=3、LogicalDevicesPerChip==1) |
| 终端类别 | StrategySubgroupND (0x638)、StrategyND (0x5f0)、TwistedTorusND (0x610) |
StrategyND ctor | 0x137c2f40 → ComputeColorDimensions 0x137c3ba0 → BuildStrategy 0x137c4660 |
| 降级轴减速机 | xla::jellyfish::GetDegradedAxis @ 0x1c894c20 (0x16b = 363 B) |
| 颜色暗淡重新映射 | BaseStrategyND::InitColorDimensionsDegraded @ 0x137c6580 |
| 弹性门 | UseResilientAlgorithmBase @ 0x1c894da0(env[0x1116]==1、NumNetworkDimensions==3、GetDegradedAxis≠-1) |
| ICI 插槽图 | GetResourceFromIciResource @ 0x1c894c00 |
| 信心 | HIGH(所有六个核心功能的反编译验证主体),除非行/标注另有说明 |
这在哪里
SelectNDStrategy 是吊舱集体下降的一级。上游,SPMD 分区器 (集体概述) 已经决定需要一个集合以及哪些副本组; SelectNDStrategy 决定“如何”在物理 ICI 环面上实现该集体。下游:
- 它所选择的环的 成本 由每种循环模型和 ICI 链路计数除数定价 - 请参阅 SPMD 链路计数成本。下面记录的
GetResourceFromIciResource槽位图与该模型共享。 - 递归加倍/二项式发射极循环(每等级合作伙伴安排
StrategyND环运行)位于 二项式/递归加倍 上。 - 扭曲圆环拓扑(
TwistedTorusND分支目标的非矩形 ICI 链路图及其BuildStrategy相序)位于 TwistedTorusND::BuildStrategy 和 扭曲概述 上。 - 降级轴原型摄取 — 运行时切片故障描述符如何写入本页*消耗的
Target降级字节 — 位于 降级轴摄取 上。 - SparseCore 卸载
HierarchicalKind策略(不同的卸载路径)位于 分层种类 上。
此页面拥有 SelectNDStrategy 决策树、五类策略枚举和 IciResource → strategy / 降级轴折叠。
签名和参数布局
反编译的序言(0x137c78f4..0x137c790e)和损坏的符号修复了完整的签名:
xla::jellyfish::StrategyND* BaseStrategyND::SelectNDStrategy( // 0x137c78e0
Target& target, // this/BaseStrategyND
const DeviceAssignment* dev_assign,
bool is_cross_module,
ObjectView<TpuCompilationEnvironment> env, // the flag-byte base
const ProgramSharedRegistry* psr,
bool use_global_ids, // StrategyND ctor param
LloRegionBuilder* lrb,
const HloInstruction* hlo,
bool enable_sub_plane, // GATES the top-level split
optional<vector<net_util::RingLocation>> ring_locs,
bool b11,
bool b12);它返回一个堆分配的StrategyND*(每个终端类都派生自StrategyND,因此即使对于StrategySubgroupND / TwistedTorusND,返回类型也是基指针)。
条目折叠(0x137c7912..0x137c7922)。 在任何分支之前,跨模块标志与操作码测试一起折叠:
is_cross_module &= (hlo != nullptr) ? hlo->IsCrossReplicaAllReduce() : 0;
// decompiled: v14 = IsCrossReplicaAllReduce(hlo) & is_cross_module仅当该指令确实是跨副本全归约(IsCrossReplicaAllReduce @ 0x1e5a0020)时,该指令才被视为“跨模块/特殊”。这个折叠的 is_cross_module (反编译器的 v14)是每个下游分支测试的内容,而不是原始参数。
决策树
两个顶级路径在 enable_sub_plane 参数上分开。每条路径要么构造一个终端策略并返回,要么落到一个共同的尾部。下面是有序分支及其保护谓词、命名 VLOG 字符串、终端类和 operator new 大小。 VLOG 字符串从 .rodata 中按字节精确解码,并在反编译正文中重新确认。
路径 A — enable_sub_plane == true:子平面子组
这是 3D 环面全归约子平面上的 2D 算法。 以下所有必须保持(反编译主体行 67..125):
hlo != nullptr
AND !is_cross_module (folded value v14 != 1)
AND env[0xe1f] != 0 (xla_pf_enable_nd_allreduce; *(a4+3615))
AND hlo->opcode != kAnd (0xd) (decompiled: **((DWORD**)hlo + 11) != 13)
AND cast<HloAllReduceInstructionBase>(hlo)->field[0xe1] == 1 (offset 225; use-global / constrain-layout)
AND ReplicaGroupsOnNDPlane(target, dev_assign, hlo->device_list(), plane=2, false).ok
AND the returned vector<MeshNDInfo> is a SINGLE group (begin == end)成功:VLOG "Enable StrategySubgroupND."(.rodata 0xa0c5fc7,26 B;反编译主体行 89)并构造 StrategySubgroupND(operator new(0x638);ctor StrategySubgroupND::StrategySubgroupND @ 0x137d4c00) (lrb, target, env, psr, dev_assign, hlo, ring_locs)。如果任何一个守卫失败,就会掉到公共尾部。
明白了 — 操作码
0xd解码为kAnd(.rodata 0x86f5a47 "and"),不是kAllReduce。kAllReduce是操作码0x9(.rodata 0x86dfb36 "all-reduce")。反编译后的guard为opcode != 13,即排除kAnd元素操作码;否则,该指令将转换为HloAllReduceInstructionBase。注意 — 这里的
ReplicaGroupsOnNDPlane(plane=2)调用与成本模型使用的平面计数原语相同;SelectNDStrategy仅读取其ok字节和 begin==end 单组测试(在0x137c79bf反编译)。内部MeshNDInfo场布局(“plane=2”意味着超出“ND 子平面”)未解码;参见 SPMD 链路计数成本。xla_pf_enable_nd_allreduce标志携带文档字符串 “在 3-D 环面的子平面上使用 2-D 算法” (.rodata 0x850c910)。
路径 B — enable_sub_plane == false:2-D ND 平面环
致电 IsGroupNDPlane(target, env, dev_assign, hlo, &NDtopo, is_cross_module) (0x137c6700)。当切片是一个 3-D 环面且其集合适合单个 2-D 平面时,它返回 true:
Target::NumNetworkDimensions() >= 3 (decompiled: v9 >= 3)
AND !multi-slice (GetMultiSliceTopology / IsMultiSliceDeviceAssignment)
AND the instruction is an HloAllReduceInstruction (cast guard: (opcode & 0xFD) == 9 ⇒ opcode ∈ {kAllReduce 0x9, kAllReduceStart 0xb}; a kFusion (0x3d) wrapper is unwrapped via ExtractInstruction first)
AND cast.field[0xe1] (offset 225) use-global OR IsCrossModuleReduceInstruction
AND the NDtopo dims are not all 1 (decompiled: v28[23]!=1 && v28[22]!=1 && (v28[24]==1 | …))
AND env[0x1015] == 1 (*(a2 + 4117) == 1; the 2-D-algorithm enable)env[0x1015] 字节读取在反编译中被确认为 *(a2 + 4117) == 1 (4117 = 0x1015)。在IsGroupNDPlane == true上:VLOG "Enabling ND sub-plane allreduce"(.rodata 0x86df4d5,31 B;反编译主体行144)并使用ND环参数集构造**StrategyND**(operator new(0x5f0)) - 反编译的ctor调用传递p7 = 1(ND环)以及 Ring-dim 参数中的 NDtopo 范围。如果是 IsGroupNDPlane == false,则落入公共尾部。
明白了 — 子平面(路径 A)门排除
kAnd (0xd)以及其他目标全归约指令; ND 平面(路径 B /IsGroupNDPlane)门还针对全归约指令(转换为HloAllReduceInstruction、(opcode & 0xFD) == 9)。注意0x9解码为kAllReduce,而不是kAllGather。两个 VLOG 字符串很容易交换:路径 A(StrategySubgroupND臂)发出"Enable StrategySubgroupND.";路径 B(StrategyNDND 平面臂)发射"Enabling ND sub-plane allreduce"。"Skipping because 2D subplane wasn't detected"字符串 (.rodata 0xa212791) 不是在这里发出的——它位于上游SplitAllReducePhases::RunImpl(0x10ff0040) 中。
普通尾部 — N 向/扭曲/跨步/默认
当路径 A 和路径 B 均未触发时到达。按顺序评价:
C-i — N 路模型并行性。 如果设置了折叠 is_cross_module,则调用 UseSpecialStrategyNDNWay(target, dev_assign) (0x137c6be0):当切片为单切片(GetMultiSliceTopology == 0、0x137c6bff)且设备分配组件计数时为 true [grp+0x8] ∈ {2,4} (0x137c6c27)。在 true 上:VLOG "Enable Strategy NDNway" (.rodata 0x84bec52, 22 B) 并使用 N 路参数集构造 StrategyND (0x5f0) - 反编译的 ctor 推送 LogicalDeviceCount (*(int*)(*dev_assign + 8)) 和环类代码7。这是 2-/4 路模型并行交叉副本全归约(字符串 0xa00b8b8)。
C-ii — 扭曲环面。 否则(单模块路径)从 [Target+0x3b8] 读取芯片部件环面变暗(X = +0x58,Y = +0x5c,Z = +0x60)并运行扭曲环面形状探头(0x137c7c39..0x137c7c7d):取三个范围的最小值/中值并测试是否为 2·(smaller dim) == (a dim extent)。在扭曲的形状上:VLOG "AllReduceEmitter: Choosing twisted topology"(.rodata 0x84b9da0,43 B)并用(target, lrb)构建**TwistedTorusND**(operator new(0x610) @ 0x137c7c8d;ctor 0x137d0040)。有关生成的非矩形环,请参阅 TwistedTorusND::BuildStrategy。
C-iii — 跨步 ND。 如果不扭曲,则探测 UseStridedStrategyND(target, &stride0, &stride1, dev_assign, hlo) (0x137c72e0):单切片 (GetMultiSliceTopology == 0)、NumNetworkDimensions >= 3 和 LogicalDevicesPerChip == 1 时为 true。为真时:VLOG "Enable StridedStrategyND" (.rodata 0x960d0b9, 24 B) 并构造 StrategyND (0x5f0),并将计算出的 stride0/stride1 作为环尺寸参数传递。
D — 默认 ND 环。 如果以上均未触发,则构造普通默认 StrategyND (0x5f0,ctor 0x137c2f40)作为 ND 环(反编译主体 LABEL_50,第 282..297 行;环类代码7、ND 环选择器 1)。默认尾部发出不 VLOG 字符串。 ("Enable StrategySubgroupND." 字符串 — .rodata 0xa0c5fc7,26 B,反编译正文行 89 — 属于路径 A,而不是此后备。)
汇总表
| 订单 | 防护装置(所有条件) | VLOG / 姓名 | 级建造(尺寸、结构) |
|---|---|---|---|
| A | enable_sub_plane && hlo && !cross_module && env[0xe1f] && opcode!=kAnd(0xd) && cast.field[0xe1]==1 && ReplicaGroupsOnNDPlane(plane=2).ok && single-group | "Enable StrategySubgroupND." (0xa0c5fc7) | StrategySubgroupND(0x638、0x137d4c00) |
| B | !enable_sub_plane && IsGroupNDPlane && env[0x1015]==1 | "Enabling ND sub-plane allreduce" (0x86df4d5) | StrategyND ND 环(0x5f0、0x137c2f40) |
| C-i | cross_module && UseSpecialStrategyNDNWay(单切片,comp-count ∈ {2,4}) | "Enable Strategy NDNway" (0x84bec52) | StrategyND N 路 (0x5f0) |
| C-ii | 单模块&&扭曲形状(2·a == dim) | "AllReduceEmitter: Choosing twisted topology" (0x84b9da0) | TwistedTorusND(0x610、0x137d0040) |
| C-iii | UseStridedStrategyND(单片、NumNetDims>=3、LDPC==1) | "Enable StridedStrategyND" (0x960d0b9) | StrategyND 跨步 (0x5f0) |
| D | 其他 | (无VLOG) | StrategyND 默认ND环(0x5f0) |
env = ObjectView<TpuCompilationEnvironment>;括号内的值是标志结构中的字节偏移量。每个 StrategyND 进一步解析为 BuildStrategy 内的一维与 N 维环(下一节)。
策略枚举 — 五个终端类
| 级 | operator new | Ctor | 角色 |
|---|---|---|---|
StrategySubgroupND | 0x638 | 0x137d4c00 | 子平面(2-D-of-3-D-torus)全归约;每个子组环然后在环上 |
StrategyND(ND环) | 0x5f0 | 0x137c2f40 | 伞 1-D / N-D 环; BuildStrategy 通过 [obj+0xa8] 解决 |
StrategyND(N路) | 0x5f0 | 0x137c2f40 | 2-/4-路模型平行环(环类代码7) |
StrategyND(跨步) | 0x5f0 | 0x137c2f40 | 步进ND环(stride0/stride1) 非矩形 ICI 链路图上的 |
TwistedTorusND | 0x610 | 0x137d0040 | 扭曲圆环 |
五个终端“类”中的三个是 相同 C++ 类型 StrategyND,仅通过 ctor 参数向量进行区分(ND 环、N 路、跨步)。对象大小(0x5f0)相同;只是区别参数有所不同。
StrategyND 构造函数
StrategyND::StrategyND(0x137c2f40,12个参数)记录区分参数,然后导出颜色维度表并构建环:
StrategyND(Target&, LloRegionBuilder*, ProgramSharedRegistry*, ObjectView<TpuCompEnv>,
bool p6, bool p7, long p8, long p9, ulong p10, long p11, long p12, bool p13)| 参数 | 存储于 | 含义 |
|---|---|---|
p6 | [obj+0xa8] | 由 BuildStrategy 消耗的 1-D-vs-ND-环选择器(==1 ⇒ UniDirectionNDRingStrategy,否则 UniDirection1DRingStrategy) |
p7 | [obj+0x5e1] | “使用 ND 环与 1-D 环”标志; ND 平面分支上的 1,因调用站点而异 |
p8, p9 | — | 两个环变暗/程度,或跨步 stride0/stride1;许多分支在这里推送 7 (ring-kind / all-3-dims-usable 代码) |
p10 | — | 大小/计数 — N 路/ND 平面分支中的 LogicalDeviceCount(来自开发分配组的 movsxd) 来自选取器 LogicalDeviceCount 参数的 |
p13 | [obj+0x5a8] |
该构造函数从常量 xmmword 0xae5cdb0(这是 {1, 1},而不是字符串指针 - "UniDirection1DRingStrategy" 类型名称字符串实际上位于 .rodata 0x84ba569)初始化两个 16 字节槽 [obj+0x5b0] 和 [obj+0x5c8],然后设置[obj+0x5c0]=1、[obj+0x5d8]=1、[obj+0xa8]=p6、[obj+0x5e1]=p7、[obj+0x5a8]=p13,最后调用Target::LogicalDevicesPerChip(0x1d615b00)、ComputeColorDimensions(0x137c3ba0—— [6][3] 调暗生产者)和 StrategyND::BuildStrategy(0x137c4660)。
注意 —
SelectNDStrategy选择类别;BuildStrategy(0x137c4660) 将[6][3]颜色调暗表转换为每颜色RingLocation邻居计划,在[obj+0xa8](0x137c4689) 上门控UniDirection1DRingStrategy与UniDirectionNDRingStrategy。BuildStrategy内部的每颜色RingLocation结构未完全解码(低);它发出的合作伙伴计划记录在 二项式/递归加倍 上。
降级轴容错重映射
环面轴在切片启动时可能有部分失效的 ICI 链接。 “弹性”集体路径通过将该轴降级到主环之外来围绕它布线。三个函数实现了这一点:每轴减速器、重映射和弹性门。
GetResourceFromIciResource — ICI 插槽映射 (0x1c894c00)
7 条指令体将 IciResource 枚举值映射到打包的 {slot, valid} 对上。逐字反编译:
__int64 GetResourceFromIciResource(int ici_resource) { // 0x1c894c00
__int64 e = (unsigned int)(ici_resource - 1);
if (e < 6) return e + 0x10000000DLL; // {slot = 0xd + e, valid = 1}
return 0; // {0, 0}
}对于 ici_resource ∈ [1..6](因此 e = ici_resource - 1 ∈ [0..5])→ 打包对 {slot = 0xd + e = 0xc + ici_resource, valid = 1};否则为 {0, 0}。六个插槽 {0xd, 0xe, 0xf, 0x10, 0x11, 0x12} 是成本模型 ResourceVector 的 3 个圆环尺寸 × 2 个环方向(请参见 SPMD 链路计数成本)。高位双字(valid = 1,来自常量 0x10000000D)是每个资源的“此插槽是真正的 ICI 链接”标记,在累积链接计数/存款周期时使用。
Target::Is{X,Y,Z}Degraded — 每轴失败链接标志
Target 结构中的三个相邻字节,每个字节都是布尔值“此维度有部分/失败的 ICI 链接,围绕它进行路由”:
| 配件 | 地址 | 源字节 |
|---|---|---|
Target::IsXDegraded() | 0x1d615940 | Target[+0x3f8] |
Target::IsYDegraded() | 0x1d615960 | Target[+0x3f9] |
Target::IsZDegraded() | 0x1d615980 | Target[+0x3fa] |
这些是从切片拓扑描述符 / TpuDegradedAxesProto 填充的 - 请参阅 降级轴摄取 了解原型→目标摄取。
GetDegradedAxis — 将每轴标志减少到一个轴或 -1 (0x1c894c20)
输入:Target* target、bitset<3> usable_mask(bit0=X、bit1=Y、bit2=Z 可用)。它从 [Target+0x3b8](反编译中的 *(this+119);Xext = +0x58,Yext = +0x5c,Zext = +0x60)读取拓扑范围,并针对每个轴应用谓词:
dX = IsXDegraded() && (Xext >= 2) && usable_mask.bit0
dY = IsYDegraded() && (Yext >= 2) && usable_mask.bit1
dZ = IsZDegraded() && (Zext >= 2) && usable_mask.bit2
num_degraded_axes = dX + dY + dZ (decompiled accumulator v12)
degraded_axis = index of the single degraded axis (0=X, 1=Y, 2=Z), else 0
return (num_degraded_axes >= 2) ? -1 : degraded_axis;反编译尾部(0x1c894cd2)正是if ((unsigned)v12 >= 2) return -1; return v9;。函数 VLOG "num_degraded_axes = {n}, degraded_axis = {a}"(.rodata 0xa219d1f + 0xa219bb2;来源 group_utils.cc:1728)。
注意 — 弹性算法可以围绕至多一个降级轴进行路由。如果两个或三个轴降级,它将返回
-1并且不采用弹性路径。 X轴臂实现无分支(v6 = a2 & 2; v7 = v6 - 1; v8 = v6 >> 1,则cmov式选择);上面的每轴语义是三个臂计算的。 X 臂索引写入的确切 SSA 是 impl 细节(低),谓词为高。
UseResilientAlgorithmBase — 当采用降级路径时 (0x1c894da0)
门(全部必须保留),在 0x1c894dd2..0x1c894e3b 反编译:
env[0x1116] == 1 (xla_tpu_use_resilient_collective_emitter; *(a2+4374)==1)
AND a usable-axis-mask test (also accepts a bool override arg)
AND Target::NumNetworkDimensions() == 3 (3-D torus only)
AND a symmetric-torus dimension comparison (X==Y, and Y==Z or 2·-related)
AND GetDegradedAxis(target, mask) != -1env[0x1116] 字节被确认为 *(a2 + 4374) == 1 (4374 = 0x1116)。当返回 true 时,ComputeColorDimensions 采用下面的降级颜色暗淡初始化路径。扭曲变体 UseResilientAlgorithmTwistedTorus (0x1c894fc0) 对扭曲环面弹性路径重复使用相同的 env[0x1116] 门 + GetDegradedAxis ≠ -1 测试,将其绑定到 C-ii 选择器分支。
InitColorDimensionsDegraded — [6][3] 环暗重映射 (0x137c6580)
签名 (Target&, long num_colors, long(*color_dims)[6][3], const bitset<3>& mask)。反编译步骤:
degraded_axis = GetDegradedAxis(target, mask)(0x137c659a)。推导两个健康轴(反编译
v5/v6):chealthy_a (v5) = (degraded_axis == 0) ? 1 : 0; // Y if X bad, else X healthy_b (v6) = (healthy_a ^ 3) - degraded_axis;
给出,每个退化轴:
degraded_axis | healthy_a | healthy_b | 内部(第 [2] 列) |
|---|---|---|---|
| 0(X 坏) | 1 (Y) | 2 (Z) | 0 (X) |
| 1(Y 坏) | 0 (X) | 2 (Z) | 1 (Y) |
| 2(Z 坏) | 0 (X) | 1 (Y) | 2 (Z) |
将
color_dims[c][0..2]填充为c = 0 .. min(num_colors, 6) - 1(反编译使用a2 != 1..6防护展开所有六行;最终的ud1是>6无法到达的陷阱):texteven rows c: [ healthy_a, healthy_b, degraded ] = [v5, v6, result] odd rows c: [ healthy_b, healthy_a, degraded ] = [v6, v5, result]
每个颜色行将 降级轴固定到内部(最后一列 [2])环尺寸,并在每个颜色的外部两列上交替两个健康轴,因此连续颜色以相反的顺序遍历健康轴(平衡两个健康的 SerDes 方向)。 VLOG "Using resilient collective, degraded_axis = {a}, color count = {n}"(.rodata 0xa219b98 + 0xa218cee;来源all_reduce_strategies.cc:571)。
重新映射如何改变ICI插槽存款
[6][3] color_dims 工作台的结构与 ComputeColorDimensions 的结构相同,ComputeColorDimensions 为造环机 和 提供动力,驱动每个环面尺寸的 ICI ResourceVector 存款(通过 GetResourceFromIciResource 的插槽 0xd..0x12):
- 正常(无降级轴):每种颜色的外部暗淡是活动环面轴;成本模型将周期存入每个活动调光的 2 个插槽中,并且
num_dims除数 =popcnt(active axes)。 - 降级(一个轴坏):坏轴位于每种颜色的内列中,因此承载减少分散/全聚集流量的外(主)环变暗仅是 2 个 健康 轴。降级轴的两个 ICI 插槽不用作主环变暗 - 每个集合的有效尺寸计数下降
3 → 2,这 (a) 将成本模型num_dims除数减少到 2 轴情况,并且 (b) 将故障轴的链接保留在环之外,以便集合在幸存的圆环链接上完成。
注意 — 最终效果是将 3-D-环面集合转换为 2-健康轴 ND 环,其中死轴降级到最内部维度,通过相同的每轴公式计算成本,但降级轴从主暗淡集中排除。数值成本公式和
num_dims除数在 SPMD 链路计数成本 上有详细说明。
TpuCompilationEnvironment 标志字节
选择器从 ObjectView<TpuCompilationEnvironment> 读取四个标志字节。字节偏移量是二进制文件中的字节精确读取;标志 names 通过相邻的 VLOG/文档字符串和辅助语义固定,而不是通过数字 proto-field-id 解码。
| 偏移 | 标志(推断) | 读取者 | 效果 |
|---|---|---|---|
0xe1f | xla_pf_enable_nd_allreduce | A路门 | 启用 ND 子平面全归约(文档 “在 3-D 环面的子平面上使用 2-D 算法”) |
0x1015 | 2D算法/ND平面启用 | IsGroupNDPlane,路径 B | 启用二维平面ND环 |
0x1116 | xla_tpu_use_resilient_collective_emitter | UseResilientAlgorithmBase | 启用降级轴弹性路径 |
0xf45 | 多色/弹性颜色计数启用 | GetColorCount (0x137c6260) | 门弹性路径前的多色分解 |
GOTCHA —
0x1015的偏移↔名称映射带有一个残余不确定性:它被读取为*(env+4117)==1(偏移量及其门控作用为高),但无论是xla_tpu_use_strided_strategy_nd(.rodata 0x86f5a4b) 还是不同的 2-D 算法标志都不是通过原始字段描述符表(TpuCompEnvReflection::GetFieldValue@0x1d7523a0、_table_@0x21cfa9e0)固定。将 name 视为 MEDIUM;字节读取及其分支效果为高。
功能图
| 功能 | 地址 | 角色 |
|---|---|---|
BaseStrategyND::SelectNDStrategy | 0x137c78e0 | 拾取器 (1999 B) |
IsCrossReplicaAllReduce | 0x1e5a0020 | 入门折叠 |
ReplicaGroupsOnNDPlane | 0x1c890960 | 路径 A 单ND平面测试(平面=2) |
BaseStrategyND::IsGroupNDPlane | 0x137c6700 | 路径 B 二维平面门 |
BaseStrategyND::UseSpecialStrategyNDNWay | 0x137c6be0 | C-i N 向闸门 |
BaseStrategyND::UseStridedStrategyND | 0x137c72e0 | C-iii 跨步闸门 |
StrategyND::StrategyND | 0x137c2f40 | ND环元件(12个参数) |
StrategyND::BuildStrategy | 0x137c4660 | [obj+0xa8] 1-D-vs-ND 环 |
BaseStrategyND::ComputeColorDimensions | 0x137c3ba0 | [6][3] 颜色暗淡生产者 |
BaseStrategyND::GetColorCount | 0x137c6260 | env[0xf45] 颜色计数 |
StrategySubgroupND::StrategySubgroupND | 0x137d4c00 | 子平面向量 |
TwistedTorusND::TwistedTorusND | 0x137d0040 | 扭曲环面 |
GetResourceFromIciResource | 0x1c894c00 | ICI插槽图 |
GetDegradedAxis | 0x1c894c20 | 每轴减速机(363B) |
Target::IsXDegraded / IsYDegraded / IsZDegraded | 0x1d615940 / …960 / …980 | 每轴标志字节 Target+0x3f8..3fa |
UseResilientAlgorithmBase | 0x1c894da0 | 弹性门 |
UseResilientAlgorithmTwistedTorus | 0x1c894fc0 | 扭转弹性门 |
BaseStrategyND::InitColorDimensionsDegraded | 0x137c6580 | [6][3] 降级重新映射 |
未解决的问题
0x1015标志名称。 字节偏移和分支效果为高;规范的 XLA 标志字符串未通过原始字段描述符表固定。中等的。- 扭曲圆环形状谓词。 C-ii 探针 (
0x137c7c39..0x137c7c7d) 简化为“2·(smaller dim) == (a dim extent)”;未导出概括Nx_Ny_Nz_twisted拓扑标签的完整封闭形式扭曲不等式。请参阅 扭曲概述。中等的。 StrategyND::BuildStrategy每颜色RingLocation结构。[obj+0xa8]1-D-vs-ND 门已确认;生成每种颜色的RingLocation集的 per-dim memmove 扇出未转录。请参阅 二项式/递归加倍。低的。ReplicaGroupsOnNDPlane(plane=2)内部结构。SelectNDStrategy只读ok+ begin==end 单组测试;MeshNDInfo字段布局和平面计数语义是开放的。低的。GetDegradedAxisX 臂中的精确位折叠。 每轴降级语义已得到证实;无分支 X 轴索引写入的文字 SSA 被注释为 impl 详细信息。低的。
交叉引用
- 集体概述 — 拾取器位于吊舱集体降低位置的位置
- 二项式/递归加倍 - 每个等级的合作伙伴安排
StrategyND环运行(发射器环路,在那里拥有) - SPMD 链路计数成本 — 成本 /
num_dims除数模型和ResourceVector槽存放IciResource地图源 - 降级轴摄取 —
TpuDegradedAxesProto→Target[+0x3f8..3fa]摄取,生成此页面使用的降级标志 - TwistedTorusND::BuildStrategy 和 扭曲概述 — C-ii 分支目标的扭曲环面拓扑
- 分层种类 — SparseCore 卸载集体策略(不同的路径)
- ICI 概述 — 物理 ICI 环面和链路模型