Skip to content

SelectNDStrategy — ND 集体算法选择器

此页面上的所有地址适用于 libtpu-0.0.40-cp314 轮子中的 libtpu.so(构建 ID 89edbbe81c5b328a958fe628a9f2207d)。其他版本会有所不同。 .text VMA等于文件偏移量;所有地址均为 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 对象大小、它们的构造函数以及 StrategyND ctor 的区别参数([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>=3LogicalDevicesPerChip==1)
终端类别StrategySubgroupND (0x638)、StrategyND (0x5f0)、TwistedTorusND (0x610)
StrategyND ctor0x137c2f40ComputeColorDimensions 0x137c3ba0BuildStrategy 0x137c4660
降级轴减速机xla::jellyfish::GetDegradedAxis @ 0x1c894c20 (0x16b = 363 B)
颜色暗淡重新映射BaseStrategyND::InitColorDimensionsDegraded @ 0x137c6580
弹性门UseResilientAlgorithmBase @ 0x1c894da0(env[0x1116]==1NumNetworkDimensions==3GetDegradedAxis≠-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)和损坏的符号修复了完整的签名:

c
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)。 在任何分支之前,跨模块标志与操作码测试一起折叠:

c
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):

text
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"),不是 kAllReducekAllReduce 是操作码 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:

text
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(StrategyND ND 平面臂)发射 "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 == 00x137c6bff)且设备分配组件计数时为 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 >= 3LogicalDevicesPerChip == 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 / 姓名级建造(尺寸、结构)
Aenable_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-icross_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-iiiUseStridedStrategyND(单片、NumNetDims>=3LDPC==1)"Enable StridedStrategyND" (0x960d0b9)StrategyND 跨步 (0x5f0)
D其他(无VLOG)StrategyND 默认ND环(0x5f0)

env = ObjectView<TpuCompilationEnvironment>;括号内的值是标志结构中的字节偏移量。每个 StrategyND 进一步解析为 BuildStrategy 内的一维与 N 维环(下一节)。


策略枚举 — 五个终端类

operator newCtor角色
StrategySubgroupND0x6380x137d4c00子平面(2-D-of-3-D-torus)全归约;每个子组环然后在环上
StrategyND(ND环)0x5f00x137c2f40伞 1-D / N-D 环; BuildStrategy 通过 [obj+0xa8] 解决
StrategyND(N路)0x5f00x137c2f402-/4-路模型平行环(环类代码7)
StrategyND(跨步)0x5f00x137c2f40步进ND环(stride0/stride1) 非矩形 ICI 链路图上的
TwistedTorusND0x6100x137d0040扭曲圆环

五个终端“类”中的三个是 相同 C++ 类型 StrategyND,仅通过 ctor 参数向量进行区分(ND 环、N 路、跨步)。对象大小(0x5f0)相同;只是区别参数有所不同。

StrategyND 构造函数

StrategyND::StrategyND(0x137c2f40,12个参数)记录区分参数,然后导出颜色维度表并构建环:

c
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-环选择器(==1UniDirectionNDRingStrategy,否则 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) 上门控 UniDirection1DRingStrategyUniDirectionNDRingStrategyBuildStrategy 内部的每颜色 RingLocation 结构未完全解码(低);它发出的合作伙伴计划记录在 二项式/递归加倍 上。


降级轴容错重映射

环面轴在切片启动时可能有部分失效的 ICI 链接。 “弹性”集体路径通过将该轴降级到主环之外来围绕它布线。三个函数实现了这一点:每轴减速器、重映射和弹性门。

GetResourceFromIciResource — ICI 插槽映射 (0x1c894c00)

7 条指令体将 IciResource 枚举值映射到打包的 {slot, valid} 对上。逐字反编译:

c
__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} 是成本模型 ResourceVector3 个圆环尺寸 × 2 个环方向(请参见 SPMD 链路计数成本)。高位双字(valid = 1,来自常量 0x10000000D)是每个资源的“此插槽是真正的 ICI 链接”标记,在累积链接计数/存款周期时使用。

Target::Is{X,Y,Z}Degraded — 每轴失败链接标志

Target 结构中的三个相邻字节,每个字节都是布尔值“此维度有部分/失败的 ICI 链接,围绕它进行路由”:

配件地址源字节
Target::IsXDegraded()0x1d615940Target[+0x3f8]
Target::IsYDegraded()0x1d615960Target[+0x3f9]
Target::IsZDegraded()0x1d615980Target[+0x3fa]

这些是从切片拓扑描述符 / TpuDegradedAxesProto 填充的 - 请参阅 降级轴摄取 了解原型→目标摄取。

GetDegradedAxis — 将每轴标志减少到一个轴或 -1 (0x1c894c20)

输入:Target* targetbitset<3> usable_mask(bit0=X、bit1=Y、bit2=Z 可用)。它从 [Target+0x3b8](反编译中的 *(this+119);Xext = +0x58,Yext = +0x5c,Zext = +0x60)读取拓扑范围,并针对每个轴应用谓词:

text
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 反编译:

text
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) != -1

env[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)。反编译步骤:

  1. degraded_axis = GetDegradedAxis(target, mask) (0x137c659a)。

  2. 推导两个健康轴(反编译v5/v6):

    c
    healthy_a (v5) = (degraded_axis == 0) ? 1 : 0;     // Y if X bad, else X
    healthy_b (v6) = (healthy_a ^ 3) - degraded_axis;

给出,每个退化轴:

degraded_axishealthy_ahealthy_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)
  1. color_dims[c][0..2] 填充为 c = 0 .. min(num_colors, 6) - 1(反编译使用 a2 != 1..6 防护展开所有六行;最终的 ud1>6 无法到达的陷阱):

    text
    even 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 解码。

偏移标志(推断)读取者效果
0xe1fxla_pf_enable_nd_allreduceA路门启用 ND 子平面全归约(文档 “在 3-D 环面的子平面上使用 2-D 算法”)
0x10152D算法/ND平面启用IsGroupNDPlane,路径 B启用二维平面ND环
0x1116xla_tpu_use_resilient_collective_emitterUseResilientAlgorithmBase启用降级轴弹性路径
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::SelectNDStrategy0x137c78e0拾取器 (1999 B)
IsCrossReplicaAllReduce0x1e5a0020入门折叠
ReplicaGroupsOnNDPlane0x1c890960路径 A 单ND平面测试(平面=2)
BaseStrategyND::IsGroupNDPlane0x137c6700路径 B 二维平面门
BaseStrategyND::UseSpecialStrategyNDNWay0x137c6be0C-i N 向闸门
BaseStrategyND::UseStridedStrategyND0x137c72e0C-iii 跨步闸门
StrategyND::StrategyND0x137c2f40ND环元件(12个参数)
StrategyND::BuildStrategy0x137c4660[obj+0xa8] 1-D-vs-ND 环
BaseStrategyND::ComputeColorDimensions0x137c3ba0[6][3] 颜色暗淡生产者
BaseStrategyND::GetColorCount0x137c6260env[0xf45] 颜色计数
StrategySubgroupND::StrategySubgroupND0x137d4c00子平面向量
TwistedTorusND::TwistedTorusND0x137d0040扭曲环面
GetResourceFromIciResource0x1c894c00ICI插槽图
GetDegradedAxis0x1c894c20每轴减速机(363B)
Target::IsXDegraded / IsYDegraded / IsZDegraded0x1d615940 / …960 / …980每轴标志字节 Target+0x3f8..3fa
UseResilientAlgorithmBase0x1c894da0弹性门
UseResilientAlgorithmTwistedTorus0x1c894fc0扭转弹性门
BaseStrategyND::InitColorDimensionsDegraded0x137c6580[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 字段布局和平面计数语义是开放的。低的。
  • GetDegradedAxis X 臂中的精确位折叠。 每轴降级语义已得到证实;无分支 X 轴索引写入的文字 SSA 被注释为 impl 详细信息。低的。

交叉引用