——“一个人用AI如何写出比PyTorch更快的自研深度学习框架”系列文章之十六
上一篇我们讲了 MemoryPlan 如何把显存按语义切成一个个 Region,并在编译期把每张张量该放在哪里、占多少字节、生命周期多长,全部安排得明明白白。但 MemoryPlan 解决的是“单卡内部”的布局问题;一旦进入分布式训练,八张卡甚至更多卡同时工作,问题就上升了一个维度:如何让所有 GPU 对同一张“内存图纸”达成一致?
在回答这个问题之前,我们不妨先退一步,看看一个深度学习框架里的“张量”到底在扮演什么角色。
一、张量的两种身份
在深度学习框架中,一个张量其实承载着两种截然不同的身份。
第一种身份是数据容器。 它持有实际的内存缓冲区,存储着浮点数或整数,你可以往里面写数据,也可以从中读数据。当我们说“把一个 batch 的图像加载到 GPU 显存”,我们谈论的是数据容器的身份。
第二种身份是内存描述符。 它描述了一块内存的逻辑属性:形状是什么、数据类型是什么、在显存的哪个位置、stride 是多少。当我们说“卷积层的输入张量形状是 [128, 224, 224, 3]”,我们谈论的是描述符的身份。
主流框架通常把这两种身份合二为一。PyTorch 的 torch.Tensor 同时持有元数据(shape、stride、dtype、device)和实际数据缓冲区(storage)。这种设计的好处是直观——你创建一个 Tensor,它自动分配内存,然后你就能直接操作它。这种“所见即所得”的体验非常适合研究和原型开发,也是 PyTorch 能在学术界迅速流行的关键原因。
但劣势同样明显:元数据和数据绑定在一起,意味着你无法在不触及具体数据的前提下去规划整个显存布局。你必须先创建 Tensor、分配内存,才能确定它的 shape 和 stride,这就把“先规划后执行”的可能性堵死了。在分布式训练中,8 张 GPU 上各有各的 Tensor 实例,它们的元数据虽然相同,但都是独立维护的副本;每次进程启动时 cudaMalloc 返回的地址也可能不同,给 CUDA Graph 的地址稳定性带来挑战。
Tech-Renaissance 的选择是把这两种身份拆开。
Tensor 是数据容器,它只在 CPU 端存在,专门负责主机与设备之间的数据搬运(H2D/D2H)。它持有页锁定内存,支持各种初始化方式,但它不参与任何 GPU 端的运算。
DTensor(DistributedTensor)是纯内存描述符。它只存储形状、偏移量、stride 和 Region 信息,不持有任何内存指针,不分配任何缓冲区。它是一个“虚拟张量”——它告诉你“有这样一个张量,它的形状是 [128, 224, 224, 3],数据类型是 FP16,位于显存偏移量 X 处,stride 是 [HWC, WC, C, 1]”,但它自己并不拥有那块内存。
这种分离带来了一个深远的结果:同一份 DTensor 描述,在所有 GPU 上对应相同的内存偏移位置,布局完全相同,但各卡显存中存储的实际数据可以有别。 这就是“一张图纸,八卡共享”的真正含义。
有了这个概念,我们就可以进入 DTensor 的设计细节。
二、分布式训练为什么需要统一的张量抽象
现代深度学习模型很少只跑在一张 GPU 上。无论是为了缩短训练时间,还是为了装下更大的模型,多卡协同都是常态。多卡训练最主流的模式是数据并行:每张卡都持有完整模型的副本,各自处理不同的数据子集,算出梯度后再通过 AllReduce 把所有卡上的梯度加和平均,保证各卡参数始终一致。
在这个模式下,每张卡执行的计算图本质上是同一张图,只是输入数据不同。既然如此,理想情况就是:同一段代码、同一张计算图、同一种显存布局,在所有卡上共享执行,而不是为每张卡单独构图、单独分配、单独调优。
但这需要解决几个关键问题。
第一,地址必须稳定。CUDA Graph 捕获要求重放时的张量地址与捕获时完全一致;如果每张卡各自动态分配,地址就难以保证一致,图也无法复用。
第二,布局必须一致。所有卡上同一个逻辑张量——比如第 3 层的卷积权重——必须位于各自显存池的相同偏移处,否则同一段 kernel 代码无法在所有卡上同样解析。
第三,变体必须兼容。训练过程中可能遇到最后一个不完整的 batch、渐进式分辨率切换、验证与训练不同尺寸等情况。如果每种情况都导致偏移漂移,那图捕获和跨卡通信都会乱套。
第四,通信必须可表达。数据并行最终要落实在 NCCL AllReduce 这类集合通信上;框架需要知道“哪些字节属于哪个张量、哪些张量需要被一起归约”。
在主流框架里,这些问题的解法各不相同。
PyTorch DTensor 是一个基于 Placement 的张量抽象。它通过 DeviceMesh、Shard(dim)、Replicate()、Partial() 等概念描述张量在多个设备上的分布方式,然后由框架自动进行分片传播和通信生成。这是一个非常通用的设计,可以支持数据并行、张量并行、流水线并行等多种策略,但相应地,它的实现复杂度高,运行时也需要做大量的分片计算和调度决策。
JAX 通过 pjit 和 NamedSharding 把同样的思想做进了函数式编译管线:用户用 jax.device_put 或 with_sharding_constraint 标注张量如何放置,XLA 在编译时决定通信并生成 SPMD 程序。
TensorFlow DTensor 也类似,围绕 Layout 和 Mesh 构建,强调跨框架的一致性。
这些方案的共同点是:把“张量如何分布到设备”作为一等公民。它们强大、通用,但通用性本身也带来了开销。
Tech-Renaissance 的选择则更为直接:我们当前聚焦的是同质 GPU 集群上的高速数据并行训练。在这个场景下,模型可以放进单张 GPU,每张卡持有完整副本,通信主体只有梯度 AllReduce 和 BN 统计量同步。于是我们不需要一个完整的分片/布局系统;我们需要的是一个跨卡一致的内存视图描述符——这就是 DTensor。
三、DTensor:纯描述符,不持有内存
在 include/renaissance/tensor/distributed_tensor.h 中,DTensor 被定义为一个非常轻量的结构体:
struct DistributedTensor {
int32_t id = -1;
Shape shape; // [N, H, W, C] 逻辑维度
DType dtype = DType::FP32;
Region region = Region::DEFAULT;
int32_t n_ = 1, h_ = 1, w_ = 1, c_ = 1;
InitConfig init_config;
int64_t numel() const noexcept;
uint64_t nbytes() const noexcept;
uint64_t padded_bytes() const noexcept;
uint64_t slot_bytes() const noexcept;
bool valid() const noexcept;
bool is_compact() const noexcept;
uint8_t cuda_alignment() const noexcept;
int64_t padded_c() const noexcept;
static uint64_t compute_slot_bytes(const Shape& shape, DType dtype, Region region) noexcept;
uint64_t offset() const;
DistributedTensor(int32_t i, Shape s, DType d, Region r);
DistributedTensor(int32_t i, Shape s, DType d, Region r, uint64_t sb);
};
using DTensor = DistributedTensor;
注意它的核心特征:DTensor 是一个纯虚拟概念,只存形状、偏移量、stride 和内存区域信息,不持有任何实际内存,也不保存设备指针。 同一个 DTensor 描述符在所有卡上对应相同的偏移位置,布局完全相同,但各卡显存中存储的实际数据可以不同。
这个设计把“张量是什么”和“张量存在哪里”解耦了。DTensor 描述的是“图纸上的位置”——它有 id、shape、dtype、region、offset、slot_bytes;而真正的物理内存由后端的 ArenaKeeper 管理,每张卡有一块独立的显存池。运行时,DeviceContext::ptr_at(dtensor_id) 才把 DTensor 解析为当前 rank 上的实际指针。
Tech-Renaissance 对 DTensor 有一个核心约定:所有 rank 的张量行为必须同步——只允许数据不同,执行的所有操作必须完全相同。 这不是一个软性的“约定”,而是由 DTensor 的架构从根本上保证的硬性约束。因为所有 rank 共享同一份 MemoryPlan 和同一份 ComputationGraph,所以“对 DTensor #42 做卷积,结果写入 DTensor #43”这条指令在所有卡上完全一致;唯一的区别只是 ptr_at() 返回的物理地址指向不同 GPU 的显存。
这与 PyTorch 的 torch.Tensor 或 JAX 的 DeviceArray 很不一样。后两者通常直接持有数据或指向数据的指针,并附带设备信息;而 DTensor 更像是蓝图里的一个编号和坐标。你可以把它理解为建筑图纸上的“第 15 号柱子在 3 轴交 B 轴、截面 400×400”,至于工地现场第 15 号柱子具体由哪根钢筋水泥浇筑,那是施工队按图纸去找的事。
框架里还有一个 Tensor 类(include/renaissance/tensor/tensor.h),它专门负责 CPU 端的数据搬运:加载数据集、做预处理、H2D 传输前的暂存。Tensor 是紧凑的、 owning 的内存容器;DTensor 则从不 owning 内存。两者职责清晰:Tensor 管“数据从哪里来”,DTensor 管“数据在 GPU 上按什么规则存放”。
四、MemoryPlan + ArenaKeeper:图纸如何落地到每张卡
DTensor 能跨卡共享的关键,是 MemoryPlan 在所有 rank 上完全一致。
在编译阶段,Compiler 会为每个训练变体生成一个 MemoryPlan。MemoryPlan 通过语义化接口分配 DTensor,例如:
DTensor MemoryPlan::alloc_fc_weight(const Shape& shape); DTensor MemoryPlan::alloc_deep_conv_weight(const Shape& shape); DTensor MemoryPlan::alloc_feature(const Shape& shape, DType dtype);
每个接口内部都调用 alloc_impl,给张量分配一个全局唯一的 id,并计算它的 slot_bytes。当所有 DTensor 都分配完毕后,调用 MemoryPlan::finalize():
void MemoryPlan::finalize() {
uint64_t cursor = 0;
for (size_t ri = 0; ri < static_cast<size_t>(Region::NUM_REGIONS); ++ri) {
auto& info = region_infos_[ri];
info.base_offset = cursor;
for (int32_t dt_id : region_dt_ids_[ri]) {
auto& entry = entries_[id_to_idx_.at(dt_id)];
entry.dt.offset_ = cursor;
cursor += entry.dt.slot_bytes();
}
info.total_bytes = cursor - info.base_offset;
}
total_bytes_ = cursor;
}
这段代码做的事情很简单:按 Region 顺序线性遍历,把每个 DTensor 的 offset_ 设为当前 cursor,然后 cursor 前进 slot_bytes()。最终所有 DTensor 的偏移被一次性确定,没有空隙、没有碎片、没有对运行时分配器的依赖。
这里的 Region 是显存的语义分区。当前版本的 Region 枚举定义了 69 个分区(NUM_REGIONS = 69),按 B(BN 统计量)、W(主权重)、E(EMA 权重)、A(AMP FP16 权重)、G(梯度)、R(结果区)、M(一阶动量)、V(二阶动量)、N(LARS 范数)、I(输入缓冲)、F(特征图与梯度槽)、S(标量与掩码)、T(临时张量)等字母系列组织。每个系列内部再按张量类型细分,例如 W 系列包含 W_BN_BIAS、W_BN_WEIGHT、W_FC_BIAS、W_FC_WEIGHT、W_FIRST_CONV、W_DEEP_CONV 等。这种命名不是装饰,而是直接服务于 RANGE OP:当框架需要“把 FP32 主权重全部转成 FP16”或“把深层卷积梯度做一次 AllReduce”时,编译器只需给出起始 Region 和结束 Region,就能覆盖整个连续区间。
重要的是:每个 rank 拿到的是同一份 MemoryPlan 的副本——同一张图纸。于是第 15 号 DTensor 在 rank 0 上位于 offset X,在 rank 7 上也位于 offset X,只是它们对应的物理显存池不同。
物理显存池由 ArenaKeeper 管理。初始化时,框架调用:
ArenaKeeper::initialize(true, gpu_ids, total_bytes_per_device);
为每张 GPU 独立 cudaMalloc 一块大小完全相同的显存。运行时,DeviceContext::ptr_at() 的代码只有两行:
void* DeviceContext::ptr_at(int dtensor_id) const noexcept {
const DTensor& dt = current_mp_->get_dtensor(dtensor_id);
return ArenaKeeper::instance().ptr_at(rank_for_context_, static_cast<size_t>(dt.offset()));
}
ArenaKeeper::ptr_at 更纯粹:
void* ArenaKeeper::ptr_at(int rank, size_t offset) const noexcept {
return static_cast<char*>(arenas_[rank]->base_ptr()) + offset;
}
热路径上就是一次基址加偏移。没有哈希查找,没有运行时分配,没有跨 rank 的分片计算。这正是数据并行场景下最干净的抽象。
五、跨变体一致性:slot_bytes 与 stride 的分离
训练中常常需要切换形状变体。例如最后一个 batch 的大小可能不是 128 而是 37;渐进式分辨率训练可能先用 224×224 再用 288×288;验证时通常不用数据增强、batch 尺寸也可能不同。如果每种变体都导致 DTensor 偏移变化,那 CUDA Graph 捕获就得每种变体重新做一遍。
Tech-Renaissance 的解法是:跨变体取最大 slot_bytes,但 stride 仍按各自 shape 计算。
DTensor 有两个构造函数:
// 标准构造:slot_bytes 从 shape 推导 DistributedTensor(int32_t i, Shape s, DType d, Region r); // 变体构造:slot_bytes 显式传入 DistributedTensor(int32_t i, Shape s, DType d, Region r, uint64_t sb);
在 Compiler 的第二阶段,框架会对每个 (layer, tensor) 位置跨所有变体计算最大 slot 字节数:
uint64_t max_slot = DTensor::compute_slot_bytes(shape_max, dtype, region);
然后在第三阶段为每个变体创建 MemoryPlan 时,用同一个 max_slot 去构造 DTensor:
DTensor dt = alloc_impl(shape_variant, dtype, region, max_slot);
这样做的结果是:不同变体下同一个 DTensor id 的 offset 完全相同,因为 MemoryPlan::finalize() 只认 slot_bytes()。与此同时,DTensor 的 stride 是在构造时根据自己的 shape 算出来的,所以每个变体运行时使用的 stride 仍然是正确的。
这是一种非常精巧的解耦:slot_bytes 决定“图纸上的地盘大小”,保证偏移一致;stride 决定“实际怎么走路径”,保证数据访问正确。两者互不干扰。
六、NHWC、对齐与 padding 的工程细节
DTensor 的 stride 计算不是简单地把 shape.c() 当作最内层维度,而是要根据 cuda_alignment() 做 C 通道 padding。
uint8_t cuda_alignment() const noexcept {
if (dtype == DType::FP16) {
switch (region) {
case Region::I_A_DATA:
case Region::I_B_DATA: return 4;
case Region::F_FEATURE_FP16:
case Region::F_GRAD_SLOT_FP16:
return 8;
default: return 1;
}
}
if (dtype == DType::INT8 && region == Region::S_MASK) {
return 8;
}
return 1;
}
padded_c() 把原始 C 通道向上对齐到 cuda_alignment 的倍数:
// 变量说明: // C = shape.c() // 原始通道数 // A = cuda_alignment() // 对齐因子(1/4/8) int64_t padded_c = align_up(C, A);
对于 FP16 输入缓冲区,C 通道按 4 对齐;对于 FP16 特征图和梯度槽,按 8 对齐。这是为了充分利用 Tensor Core 的向量化访问模式:当通道数是 8 的倍数时,半精度卷积和矩阵乘可以更高效地读取和写入。
slot_bytes() 的计算公式也很有讲究:
static uint64_t compute_slot_bytes(const Shape& shape, DType dtype, Region region) noexcept {
// ... alignment 推断 ...
int64_t pc = align_up(static_cast<int64_t>(shape.c()), alignment);
uint64_t elems = static_cast<uint64_t>(shape.n()) * shape.h() * shape.w() * pc;
if (dtype == DType::FP16) {
return utils::align_up_256(elems * 2 + 16);
} else if (dtype == DType::INT8) {
return utils::align_up_256(elems * 1 + 16);
} else {
return 2 * utils::align_up_256(elems * 2 + 16); // FP32 = 2 × FP16
}
}
这里有几个值得注意的设计点。
第一,末尾预留 16 字节。 这不是浪费,而是为了兼容 XNNPACK 等底层库的边界要求。某些 kernel 在处理最后一行时可能会做越界读取或向量化存取,预留 16 字节可以避免非法访问风险,同时不影响逻辑数据。
第二,slot 大小向上取整到 256 字节。 这与框架“首地址 256 字节对齐”的原则一致,确保所有张量的起始地址满足 GPU 的对齐要求,也有利于内存合并访问。
第三,FP32 的 slot 是 FP16 等效 slot 的两倍。 这不是随意定的。框架在 AMP 训练时需要在 FP32 主权重和 FP16 计算权重之间做整区转换;如果 FP32 区域的总字节数与对应 FP16 区域总字节数始终保持严格的 2:1 关系,那么从“FP32 权重区”到“FP16 权重区”的批量类型转换就可以用一个简单的 RANGE OP 一次性完成——一个 kernel 遍历所有层的所有参数,按固定比例缩放即可。这比自己写逐层转换快得多,也是框架速度快的根基之一。
DTensor 还显式区分了 CUDA 对齐 stride 和 CPU 紧凑 stride:
int64_t n_stride_cuda() const noexcept; // 基于 padded_c int64_t n_stride_cpu() const noexcept; // 基于 shape.c(),恒紧凑
CPU 上的 DTensor 永远紧凑,与 Tensor 的内存排布一致,方便 H2D/D2H 传输;GPU 上则允许 padding 以换取卷积效率。当 is_compact() 返回 false 时,传输路径会做必要的 layout conversion。
DTensor 构造时同时算出两套 stride。以 CUDA 对齐 stride 为例:
// 变量说明: // ac = padded_c() // 对齐后的 C 通道数 // W = shape.w() // 宽度 // H = shape.h() // 高度 int64_t c_stride_cuda_ = 1; int64_t w_stride_cuda_ = ac; int64_t h_stride_cuda_ = ac * W; int64_t n_stride_cuda_ = ac * W * H;
CPU 紧凑 stride 则用原始 shape.c() 计算,不引入 padding:
// 变量说明: // C = shape.c() // 原始通道数 // W = shape.w() // H = shape.h() int64_t c_stride_cpu_ = 1; int64_t w_stride_cpu_ = C; int64_t h_stride_cpu_ = C * W; int64_t n_stride_cpu_ = C * W * H;
这些对齐规则看似琐碎,实则共同服务于一个目标:让批量操作(RANGE OP)成为可能。DTensor 的 Region 分类、slot 对齐、FP32 与 FP16 之间的 2:1 字节关系,都不是孤立的工程细节,而是为了让编译器可以用一次 kernel launch 覆盖整个语义分区。
以 AMP 训练的权重同步为例。模型同时维护 FP32 主权重(W 系列 Region)和 FP16 计算权重(A 系列 Region)。由于每个 FP32 张量的 slot_bytes 恒等于其对应 FP16 张量 slot 的两倍,所有层的主权重在 W 区连续排列,所有层的 AMP 权重在 A 区连续排列,W 区总字节数与 A 区总字节数也保持严格的 2:1 关系。于是“FP32→FP16”的整区转换只需要一个 RANGE CAST 算子:一个 kernel 从 W 区起点遍历到 W 区终点,按 2:1 比例把数据写入 A 区。同理,优化器更新所有 BN bias、所有 Conv weight、所有 FC weight,也可以按 Region 分组,用 1~2 个 kernel 完成,而不是逐层遍历。
换句话说,DTensor 的对齐与分区规则,是把“模型有多少层”这个复杂度从运行时抹掉的关键。无论模型是 19 层的 VGG16 还是 152 层的 ResNet,只要同一类张量落在同一个 Region 里,批量操作的开销就不会随层数线性增长。
七、DTensor 如何支撑分布式图执行
DTensor 不只是内存描述符,它也是计算图里的“地址符号”。
在 ComputationGraph 中,每个算子节点只保存 DTensor 的 id,不保存指针,也不保存形状。例如一个卷积节点会记录 input_ids = {3, 5}、output_ids = {7};运行时,DeviceContext::ptr_at(3)、ptr_at(5)、ptr_at(7) 把这些 id 解析为当前 rank 上的实际指针。
由于所有 rank 的 MemoryPlan 相同,同一个 DTensor id 在所有 rank 上解析出的偏移相同;由于 ArenaKeeper 为每张卡分配了独立的池,解析出的指针指向各自卡上的数据。于是同一张计算图可以在所有 rank 上无修改地执行,这正是数据并行训练的基础。
对于集合通信,DTensor 的 Region 语义进一步发挥了作用。梯度 AllReduce 不是逐张量做的,而是按 Region 批量做的。在 src/backend/ops/range/allreduce_op.cpp 中:
static void launch_allreduce_cuda_impl(
const GraphNode& node,
const MemoryPlan& mp,
const DeviceContext& ctx,
MultiStreamCaptureState& state)
{
cudaStream_t s = static_cast<cudaStream_t>(ctx.stream(StreamKind::UPDATE));
// ...
for (size_t i = 0; i < node.input_ranges.size(); ++i) {
auto [src_off, src_sz] = mp.resolve_region_bounds(
static_cast<Region>(node.input_ranges[i].start_region_id),
static_cast<Region>(node.input_ranges[i].end_region_id));
auto [dst_off, dst_sz] = mp.resolve_region_bounds(
static_cast<Region>(node.output_ranges[i].start_region_id),
static_cast<Region>(node.output_ranges[i].end_region_id));
void* src = ArenaKeeper::instance().ptr_at(ctx.rank_for_context(), src_off);
void* dst = ArenaKeeper::instance().ptr_at(ctx.rank_for_context(), dst_off);
size_t count = std::min(src_sz, dst_sz) / sizeof(float);
if (src != dst) cudaMemcpyAsync(dst, src, count * sizeof(float),
cudaMemcpyDeviceToDevice, s);
ncclAllReduce(dst, dst, count, ncclFloat32, ncclSum,
static_cast<ncclComm_t>(ctx.nccl_comm()), s);
if (do_mean && world_size > 1) {
float inv = 1.0f / static_cast<float>(world_size);
launch_tr_scale_fp32_kernel(static_cast<float*>(dst), inv, count, s);
}
}
}
这里 GraphNode 的 input_ranges 和 output_ranges 是 Region 范围。MemoryPlan::resolve_region_bounds 把 Region 区间解析成 (offset, size),然后 ArenaKeeper::ptr_at 把 offset 变成当前 rank 的指针。一次 ncclAllReduce 可以覆盖整个梯度桶——例如 G_DEEP_CONV 区所有深层卷积梯度、或 G_FC_WEIGHT 区所有全连接权重梯度。
因为 DTensor 的 Region 是编译期确定的,所以这些范围在图捕获阶段就解析好了;CUDA Graph 捕获时,NCCL 调用的指针地址也是固定的。运行时只需要 cudaGraphLaunch,不需要再做任何 Region 到指针的转换。
八、与主流张量抽象的对比
有人可能会问:Tech-Renaissance 的 DTensor 为什么不支持 PyTorch 那种 Shard(dim) / Replicate() / Partial() 的通用分片语义?
答案与定位有关。PyTorch DTensor、JAX 的 NamedSharding、TensorFlow DTensor 都把“张量如何分布到设备”作为一等公民,目标是支持数据并行、张量并行、流水线并行等多种策略。这种通用性很强大,但运行时也需要做分片传播、布局计算和通信生成。
Tech-Renaissance 当前聚焦的是同质 GPU 集群上的高速数据并行训练:模型能放进单张 GPU,每张卡持有完整副本,通信模式固定为“本地算梯度 → 全局 AllReduce → 本地更新”。在这个场景下,我们不需要一个完整的 Placement 系统;我们需要的是一个跨卡一致的内存视图描述符。
所以 Tech-Renaissance 的 DTensor 是一个“面向数据并行的内存布局描述符”,而不是“面向通用并行的分片张量”。它的优势在于:
- 指针解析极简:基址加偏移,一次加法。
- 图共享无歧义:同一张 ComputationGraph 可以在所有 rank 上复用。
- CUDA Graph 友好:地址稳定、Region 范围静态确定、通信可捕获。
- 变体切换零偏移漂移:跨变体最大 slot 保证同一 id 在不同 shape 下 offset 一致。
它的代价是不支持模型并行、张量并行、流水线并行。如果未来要把单个大模型拆到多张卡上,这个抽象就需要扩展。但就当前任务而言——让能放进单卡的模型在八卡上跑得更快——它已经足够好。
从另一个角度看,Tech-Renaissance 的 DTensor 与 PyTorch 的 torch.Tensor 代表了两种设计哲学。PyTorch 的 Tensor 是状态型对象:元数据和数据绑定,创建即分配,销毁即释放,直观但把“布局规划”留在了运行时。DTensor 是纯描述符:不持有内存,编译期即可确定 offset、stride、slot_bytes,牺牲了一部分直接可操作性,但换来了布局确定性、跨卡一致性和批量操作效率。两者没有绝对的优劣,只是为不同的执行模型做了不同的取舍。
九、小结
DTensor 是 Tech-Renaissance 里连接“单卡内存规划”与“多卡分布式执行”的桥梁。
它不持有内存,只描述内存:id、shape、dtype、region、offset、slot_bytes、stride。MemoryPlan 负责在所有 rank 上生成同一份图纸;ArenaKeeper 为每张卡分配独立但大小相同的显存池;DeviceContext 在运行时把 DTensor id 解析为当前 rank 的实际指针。因为图纸一致,所以同一张计算图可以八卡共享;因为 slot_bytes 跨变体取最大,所以不同 batch 大小、不同分辨率切换时偏移不会漂移。
这套设计的核心思想可以概括为一句话:一张图纸,多卡共享;数据不同,执行相同。
下一篇,我们将进入运行时阶段,先讲 Tech-Renaissance 的多流并发架构——三条计算流、一条传输流、一条更新流如何协同工作,把数据搬运、梯度通信和计算最大限度地重叠起来。
