(7) 张量、NHWC布局与256字节对齐：框架里的”砖块”

——“一个人用AI如何写出比PyTorch更快的自研深度学习框架”系列文章之七

上一篇我们把 Tech-Renaissance 的七大模块串成了一张全局地图，从 DTS 文件读取、CPU 预处理、H2D 异步传输，到前向、反向、梯度通信、优化器更新，再到指标输出，整条流水线已经在你脑子里有了一根主线。

但地图再漂亮，也得一块砖一块砖地砌。从今天开始，我们把这些”砖块”逐一拿起来细看。而所有砖块里最基础、最不起眼、却又无处不在的，就是张量（Tensor）。

张量这个东西，初学者第一次见往往觉得它不过是个”多维数组”：标量是零维，向量是一维，矩阵是二维，batch 图像是四维。但如果你真去做一个高性能深度学习框架，你会发现事情远没有这么简单。张量怎么在内存里排布、通道放在哪个维度、首地址要不要对齐、stride 怎么计算——这些看起来细枝末节的问题，最终会直接决定你卷积能跑多快、显存能省多少、以及 CUDA Graph 能不能稳定捕获。

这篇就来讲讲 Tech-Renaissance 里张量的设计哲学：为什么统一 NHWC、为什么首地址 256 字节对齐、为什么 DTensor 是一个不持有内存的纯描述符，以及 CPU 和 CUDA 路径为什么会有两套 stride。 这些都是后面 MemoryPlan、CUDA Graph、多流并发等篇章的公共底座，必须先把它夯实。

一、张量不只是一块内存

1.1 从标量到张量

在深度学习的语境里，张量（Tensor）是数据的通用容器。一个标量可以看作零阶张量，一个向量是一阶张量，一个矩阵是二阶张量。而深度学习里最常见的批量图像数据——比如一个 batch 的 224×224 RGB 图片——就是一个四阶张量，形状通常写作 [N, H, W, C]：N 是 batch 大小，H 和 W 是高和宽，C 是通道数。

但”形状”只是张量的一个属性。一个训练框架里的张量，至少要携带以下信息：

• 形状（Shape）：每个维度的大小；
• 数据类型（DType）：FP32、FP16、INT8、INT32 等；
• 内存布局（Layout）：NCHW 还是 NHWC，通道在最内层还是外层；
• Stride：沿每个维度走一格需要跳过多少元素；
• 所在设备：CPU 内存还是 GPU 显存；
• 是否需要梯度：是否参与反向传播；
• 指向的计算图节点：动态图里还要记录它是哪个算子产生的。

这些信息里，Shape、DType 和 Layout 是最基础的，它们共同决定了”这块内存里到底装着什么”。而 Stride 则告诉你”怎么按逻辑顺序访问它”。

1.2 NCHW 与 NHWC：两种世界观

同样一张 [N, H, W, C] 的图像 batch，在内存里可以有两种完全不同的排法。

NCHW 把通道放在高宽之前。逻辑上你先遍历 batch N，再遍历通道 C，然后才是高 H、宽 W。内存里的顺序大致是：先把第一张图的所有 R 通道像素（整张 H×W 平面）放完，再放所有 G 通道像素，然后 B 通道，接着第二张图…… 因此，同一通道的数据在内存里是连续存放的，但同一个像素位置的不同通道值之间隔了一整张 H×W 平面。

NHWC 则把通道放在最内层。内存顺序是：先放第一个空间位置（H=0, W=0）的所有通道值（R、G、B），再放下一个空间位置的所有通道值…… 因此，同一个像素位置的不同通道值在内存里是挨着的，但相邻空间位置的同一通道数据在内存里并不连续。

这两种布局没有绝对的数学优劣，但访存特性差异巨大。

NCHW 适合做”通道级”操作：比如你要对整张特征图做全局池化，或者按通道做归一化，NCHW 下同一通道的数据是连续的，访问局部性好。NHWC 则更适合”空间级”操作：比如卷积核在 H-W 平面上滑动时，每滑到一个位置都要同时读多个通道，NHWC 下这些通道值在内存里紧挨着，一次读取就能拿到一个像素位置的所有通道。而在 NCHW 下，同一个像素的 C 个通道跨越了整个 H×W 平面，需要多次分散访问。

主流框架的选择也很有趣：

• PyTorch 默认使用 NCHW。这主要与 cuDNN 早期对 NCHW 路径支持最成熟的生态有关。cuDNN 最初的许多算子对 NCHW 支持最成熟，PyTorch 沿用了这一传统。
• TensorFlow 早期默认使用 NHWC，原因之一是 TensorFlow 最初在 CPU 上开发，NHWC 对 CPU cache 更友好。但 TensorFlow 也支持 data_format='NCHW'，在 GPU 上常常更快。
• ONNX Runtime、TensorRT 等推理框架则会根据后端和硬件自动做布局转换。

所以长期以来，社区里有一种说法：”GPU 用 NCHW，CPU 用 NHWC。”这句话在几年前大体成立，但近些年情况已经发生变化——cuDNN 对 NHWC 的支持越来越好，Tensor Core 在很多场景下反而更偏好 NHWC。NVIDIA 官方文档和不少评测都表明，对于 ResNet、VGG 这类典型 CNN，NHWC 配合 FP16 Tensor Core 常常能获得比 NCHW 更高的吞吐。

1.3 内存对齐：为什么不是”随便放”

除了布局，另一个容易被忽视的问题是对齐（Alignment）。

GPU 访问全局内存时，是以 warp（32 个线程）为单位发起请求的。如果 warp 里的 32 个线程访问的地址是连续且对齐的，GPU 硬件可以把这些请求合并成一次或几次大的内存事务（Coalesced Access），最大限度地利用显存带宽。反之，如果地址分散或不对齐，同一个 warp 的访问可能被拆成几十次小事务，带宽利用率大幅下降。

现代 NVIDIA GPU 的全局内存访问以 32 字节 sector 为基础，若干 sector 再组合成更大的内存事务。如果 warp 访问的地址连续且对齐，硬件就能把请求合并成少数几次大事务；反之则会被拆成多次小事务，带宽利用率大幅下降。CUDA 的 cudaMalloc 返回的地址默认至少 256 字节对齐，很多高性能计算库（cuDNN、cuBLAS）内部也会要求输入指针满足一定的对齐条件，否则可能直接拒绝或者回退到慢路径。

所以”首地址对齐”不是强迫症，而是硬件层面的硬性要求。256 字节对齐比 128 字节更严格，可以确保：

• 对齐边界足够宽，不容易出现跨边界的分裂访问；
• cuDNN 的某些 kernel 在拿到 256 对齐地址时可以直接走最优实现；
• 多流、多卡场景下，地址规律一致，便于做批量传输和通信优化。

二、Tech-Renaissance 的张量分层

在 Tech-Renaissance 里，张量不是一个类打天下，而是被明确地分成了三个层次：

• Shape / DType：最基础的类型系统，定义在 include/renaissance/core/types.h；
• Tensor：CPU 端数据容器，只用于主机-设备数据搬运；
• DTensor（DistributedTensor）：设备端计算的一等公民，是一个不持有内存的纯描述符。

这种分层看起来多此一举，实则是经过深思熟虑的。Tensor 负责”数据在 CPU 上长什么样”，DTensor 负责”数据在 GPU 上怎么被引用”，MemoryPlan 负责”数据在显存里放在哪里”，DeviceContext 负责”把描述符翻译成真实指针”。 四个角色各司其职，才能把静态图、CUDA Graph、多卡并行串起来。

2.1 Shape 与 DType：被固定为 4D NHWC

先看最基础的 Shape：

// include/renaissance/core/types.h
struct Shape {
    int dims[4] = {1, 1, 1, 1};  // NHWC格式：dims[0]=N, dims[1]=H, dims[2]=W, dims[3]=C
    ...
    int n() const { return dims[0]; }
    int h() const { return dims[1]; }
    int w() const { return dims[2]; }
    int c() const { return dims[3]; }
    int64_t numel() const noexcept { ... }
};

注意这个设计：Shape 永远是 4 维，语义固定为 [N, H, W, C]。全连接层的权重在这个框架里也会被表达成 [O, 1, 1, I] 的 4D NHWC 形状，偏置是 [O, 1, 1, 1]。这样做的好处是所有算子都按同一套维度约定处理，不需要为全连接、卷积、池化分别写形状推导。

Shape 的构造函数里还有一个值得一提的细节：由于 MSVC 在 /O2 优化级别下曾对简单 POD 结构体的初始化列表产生错误优化，框架里使用 TR_NOINLINE 配合 volatile 写入来强制生成实际内存写入指令，规避这个编译器层面的坑。这个细节说明，做底层基础设施软件，你面对的不只是算法和架构，还有编译器的”脾气”。

DType 则很简单：

// include/renaissance/core/types.h
enum class DType : uint8_t { FP32, FP16, INT8, INT32 };

2.2 Tensor：CPU 端的”二等公民”

Tensor 类的基本情况是这样：

// include/renaissance/tensor/tensor.h
/**
 * @class Tensor
 * @brief CPU 端数据容器（移动语义，禁用拷贝）
 *
 * 用于从主机向设备（H2D）或设备向主机（D2H）搬运数据。
 * ...
 * V4.20重要变更：
 * - Tensor类强制紧凑布局（无padding）
 * - 所有数据按NHWC顺序连续存储
 * - 移除所有行步幅对齐要求
 * - 对齐要求仅保留：首地址256字节对齐
 * - 与DTensor的non-compact布局转换在transfer时自动处理
 */

这里的关键词是：强制紧凑、无 padding、NHWC、首地址 256 字节对齐。

为什么 Tensor 要强制紧凑？因为它只负责 CPU 端的数据搬运。无论是从磁盘读到内存、从内存拷到显存，还是从显存拷回内存，CPU 侧的数据都希望是”所见即所得”：N×H×W×C 个元素，每个元素占对应 dtype 的字节数，总大小就是 numel() * sizeof(dtype)，不多不少。如果 CPU 侧也做 padding，那文件读取、CRC 校验、预处理管线都会变复杂。

但 Tensor 的首地址仍然要求 256 字节对齐。实现上，GPU 模式用 cudaMallocHost 分配页锁定内存（CUDA 保证返回地址至少 256 对齐），CPU 模式用 mi_malloc_aligned(..., 256)，并做防御性校验：

// src/tensor/tensor.cpp
Tensor::Tensor(const Shape& shape, DType dtype) : shape_(shape), dtype_(dtype) {
    ...
    nbytes_ = static_cast<size_t>(num) * elem_size;
    elem_size_ = elem_size;

    constexpr size_t base_alignment = 256;

#ifdef TR_USE_CUDA
    cudaError_t err = cudaMallocHost(&ptr_, nbytes_);
#else
    ptr_ = mi_malloc_aligned(nbytes_, base_alignment);
#endif

    if (reinterpret_cast<uintptr_t>(ptr_) % base_alignment != 0) { ... }
}

Tensor 还禁用拷贝、只支持移动语义。原因很直接：在训练热路径上随意复制一块 CPU 内存是不可接受的，框架鼓励你用 std::move 转移所有权，实现零开销传递。

2.3 DTensor：设备端的纯描述符

如果说 Tensor 是”拿着一块内存的容器”，那 DTensor 就是”描述一块内存的图纸”。

// include/renaissance/tensor/distributed_tensor.h
/**
 * @brief 分布式张量 — "一张图纸，八卡共享"的内存抽象
 *
 * DTensor 是一个纯虚拟概念：只存形状/偏移量/stride，不持有内存、不存指针。
 * 同一个DTensor指代所有卡上相同的物理内存区域，布局完全相同，但数据可以有别。
 *
 * 职责分离：
 *   DTensor       — 描述单张量在多卡上的统一内存视图（本文件）
 *   MemoryPlan    — 分配 DTensor、计算偏移、确保无间隙和 256B 对齐
 *   DeviceContext — 将 DTensor.offset 解析为真实 GPU/CPU 指针
 *
 * 关键约定：
 *   - 全部 4D NHWC 物理布局，Shape 存逻辑维度（padding前）
 *   - 首地址必定 256 字节对齐，slot_bytes() 必定为 256 的整数倍
 *   - 创建后 id/shape/dtype/region 不可变
 *   - offset_ 由 MemoryPlan::finalize 赋值，此前恒为 -1（哨兵）
 *   - stride 由 shape + dtype 计算，创建后不可变
 */

DTensor 里不存指针，只存：

• id：全局唯一编号；
• shape / dtype / region：形状、类型、所属显存语义区；
• offset_：在 MemoryPlan 里的字节偏移；
• slot_bytes_：在显存里占用的槽位大小；
• 两套 stride：CUDA 对齐 stride 和 CPU 紧凑 stride。

运行时需要真实指针怎么办？通过 DeviceContext::ptr_at(dtensor_id) 动态解析：

// src/backend/device_context.cpp
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()));
}

这个设计的威力在于：同一张 MemoryPlan 图纸可以在多张 GPU 上共享。八张卡上同一个 DTensor id 拥有相同的 offset 和 slot_bytes，只是分别绑定到各自 Arena 的基地址。多卡并行时不需要为每张卡维护一份独立的张量描述，MemoryPlan 直接复用。

三、为什么 Tech-Renaissance 统一采用 NHWC

3.1 GPU 训练，布局要为 GPU 服务

Tech-Renaissance 的定位是高吞吐训练框架，主战场在 NVIDIA GPU。在这个前提下，统一 NHWC 是自然而然的决定。

原因并不玄乎：

第一，cuDNN 对 NHWC 的优化已经非常成熟。cuDNN Frontend 的 Graph API 允许调用方通过 stride 显式表达 NHWC，NVIDIA 在 Ampere、Hopper 及后续架构上为 NHWC 做了大量针对性的融合 kernel。对于 Conv + BN + ReLU 这种经典组合，NHWC 配合 FP16 Tensor Core 往往能达到最优吞吐。

第二，NHWC 与数据管线天然匹配。图像预处理本质上是在 H-W 平面上对 RGB 像素做操作，预处理完后直接产出 NHWC 格式最自然，不需要额外做一次 NHWC → NCHW 的转置。

第三，避免布局转换的开销。PyTorch 里如果你用 channels_last，框架会在 NCHW 和 NHWC 之间做转换；转换本身是一次显存读写，对于高分辨率特征图并不便宜。Tech-Renaissance 从第一层输入开始就统一 NHWC，所有卷积、池化、BN、激活都在同一种布局下完成，消除了这种来回倒腾。

当然，统一 NHWC 不是没有代价。CPU 端某些通道级操作的局部性会变差，所以我们在 CPU 算子里用 SIMD 和紧凑 stride 做补偿。但这是一个清醒的 trade-off：GPU 训练吞吐才是核心指标，CPU 只是辅助搬运。

3.2 cuDNN 桥接：用 NCHW 顺序传 dim，用 stride 表达 NHWC

一个有趣的细节是，Tech-Renaissance 调用 cuDNN 时，维度数组仍然按 NCHW 顺序传入，但 stride 给出 c_stride = 1，这样 cuDNN 就知道物理布局其实是 NHWC。这种”表里不一”是 cuDNN API 历史包袱造成的。

// include/renaissance/tensor/distributed_tensor.h
/**
 * @brief 以 cuDNN 的 NCHW API 顺序返回逻辑维度
 *
 * dim_index: 0→N, 1→C, 2→H, 3→W
 *
 * 关键：cuDNN API 期望 NCHW 顺序，但物理布局是 NHWC。
 * 此方法填正确值即可——cuDNN 通过 cudnn_stride(1)==1 自动识别 NHWC。
 */
int64_t cudnn_dim(int dim_index) const noexcept { ... }

int64_t cudnn_stride(int dim_index) const noexcept {
    switch (dim_index) {
        case 0: return n_stride_cuda_;
        case 1: return c_stride_cuda_;  // = 1
        case 2: return h_stride_cuda_;
        case 3: return w_stride_cuda_;
        default: return 0;
    }
}

cudnn_stride(1) == 1 这个信号一给，cuDNN 就明白”C 维度是最内层”，于是按 NHWC 物理布局调度对应的 kernel。这是一种优雅的”隐式协商”——框架不需要显式地告诉 cuDNN”我用的是 NHWC”，cuDNN 自己从 stride 信息中推断出来，并选择最优的执行计划。

四、256 字节对齐：从工具函数到全局契约

4.1 一个不起眼的工具函数

整个框架的 256 字节对齐都基于这个函数：

// include/renaissance/core/types.h
constexpr inline size_t align_up_256(size_t size) noexcept {
    return (size + 255) & ~static_cast<size_t>(255);
}

它本身只是一行位运算，但它背后是一套贯穿全局的契约：Tensor 分配要 256 对齐、DTensor 的 slot 要 256 对齐、MemoryPlan 的 offset 要 256 对齐、workspace 要 256 对齐、TSR 文件 RAW 模式的数据块也要 256 对齐。

4.2 DTensor 的 slot_bytes：对齐 + 预留

每个 DTensor 在显存里占用的物理槽位不是简单的 numel * sizeof(dtype)，而是：

// include/renaissance/tensor/distributed_tensor.h
static uint64_t compute_slot_bytes(const Shape& shape, DType dtype, Region region) noexcept {
    ...
    int64_t pc = align_up(static_cast<int64_t>(shape.c()),
                          static_cast<int64_t>(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);  // FP16: 2 bytes per element
    } else if (dtype == DType::INT8) {
        return utils::align_up_256(elems * 1 + 16);  // INT8: 1 byte per element
    } else {
        return 2 * utils::align_up_256(elems * 2 + 16);  // FP32 基准 = 2×FP16
    }
}

注意几个细节：

• 先按 region 决定的 alignment 对 C 通道做 padding；
• 计算 padded 元素数后，再按 dtype 换算字节；
• 最后 +16 预留，再 align_up_256。

那个 +16 不是拍脑袋。它为某些算子在边界处理、向量化读写或 CPU 侧 XNNPACK 等路径上提供了安全边际。而 256 对齐则保证了 MemoryPlan 线性累加时，下一个 DTensor 的 offset 仍然落在 256 边界上。

另一个值得注意的地方是 FP32 槽位 = 2 × FP16 等效槽位。这个倍率关系不是巧合：它让同一张量在不同精度下的显存位置保持字节级对应，后续就可以用单个 range op 完成整区 FP16↔FP32 的转换，而不必逐层、逐张量地写循环。

4.3 MemoryPlan 如何保证每个 offset 都 256 对齐

MemoryPlan::finalize() 做一遍线性累加：

// src/graph/memory_plan.cpp
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;

    validate_region_order();
    validate_contiguity();
    validate_layer_correspondence();
    validate_alignment();
    finalized_ = true;
}

因为每个 slot_bytes() 都是 256 的倍数，所以 cursor 永远停留在 256 边界上，每个 DTensor 的 offset_ 自然满足 256 对齐。最后 validate_alignment() 再做一次防御性校验：

// src/graph/memory_plan.cpp
void MemoryPlan::validate_alignment() const {
    for (const auto& entry : entries_) {
        TR_CHECK(entry.dt.offset() % 256 == 0, ValueError,
             "Alignment violated: DTensor " << entry.dt.id << " offset=" << entry.dt.offset());
    }
}

这种”构造时保证 + 结束时校验”的双重机制，让 CUDA Graph 全捕获所需的地址稳定性有了一个坚实基础。后面讲 CUDA Graph 时会看到，地址稳定是多么重要。

五、双轨 stride：CUDA 对齐与 CPU 紧凑

这是 Tech-Renaissance 张量系统里最有个性的设计之一。

前面说过，CPU 端的 Tensor 强制紧凑无 padding；但 GPU 端为了适应 cuDNN 和 Tensor Core 的向量化要求，某些 Region 的 FP16 张量需要在 C 通道做 padding。比如输入缓冲区按 4 对齐，特征图区按 8 对齐。这就导致同一个逻辑形状在 CPU 和 CUDA 两条路径上的 stride 不一样。

于是 DTensor 直接提供了两套 stride：

// include/renaissance/tensor/distributed_tensor.h
/**
 * @brief CUDA 对齐 stride — 由 padded_c()（cuda_alignment）推导
 */
int64_t n_stride_cuda() const noexcept { return n_stride_cuda_; }
int64_t h_stride_cuda() const noexcept { return h_stride_cuda_; }
int64_t w_stride_cuda() const noexcept { return w_stride_cuda_; }
int64_t c_stride_cuda() const noexcept { return c_stride_cuda_; }

/**
 * @brief CPU 紧凑 stride — 框架保证 CPU 上所有 DTensor 必定紧凑
 */
int64_t n_stride_cpu()  const noexcept { return n_stride_cpu_; }
int64_t h_stride_cpu()  const noexcept { return h_stride_cpu_; }
int64_t w_stride_cpu()  const noexcept { return w_stride_cpu_; }
int64_t c_stride_cpu()  const noexcept { return c_stride_cpu_; }

CUDA 路径上，C 通道可能 padding：

// include/renaissance/tensor/distributed_tensor.h
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:
#ifdef TR_USE_CUDA
                return 8;
#else
                return 1;
#endif
            default: return 1;
        }
    }
    ...
    return 1;
}

int64_t padded_c() const noexcept {
    return align_up(static_cast<int64_t>(shape.c()),
                    static_cast<int64_t>(cuda_alignment()));
}

构造时两套 stride 分别计算：

// include/renaissance/tensor/distributed_tensor.h
DistributedTensor(int32_t i, Shape s, DType d, Region r)
    : id(i), shape(s), dtype(d), region(r), ... {
    slot_bytes_ = compute_slot_bytes(s, d, r);
    int64_t ac = padded_c();
    c_stride_cuda_ = 1;
    w_stride_cuda_ = ac;
    h_stride_cuda_ = ac * s.w();
    n_stride_cuda_ = ac * s.w() * s.h();

    c_stride_cpu_ = 1;
    w_stride_cpu_ = s.c();
    h_stride_cpu_ = s.w() * s.c();
    n_stride_cpu_ = s.h() * s.w() * s.c();
}

CPU 路径始终紧凑，w_stride_cpu = C；CUDA 路径在需要时 w_stride_cuda = padded_c。两者谁该用哪套，由 capture 和算子实现自行选择。例如 CPU capture 用紧凑 stride：

// src/graph/capture_cpu.cpp
op_ctx->n_stride = dt_in.n_stride_cpu();
op_ctx->h_stride = dt_in.h_stride_cpu();
op_ctx->w_stride = dt_in.w_stride_cpu();
op_ctx->c_stride = dt_in.c_stride_cpu();

CUDA 卷积算子用对齐 stride：

// src/backend/ops/dtensor/conv_op_impl.cpp
inline std::vector<int64_t> make_nhwc_stride(const DTensor& dt) {
    return {dt.n_stride_cuda(), dt.c_stride_cuda(),
            dt.h_stride_cuda(), dt.w_stride_cuda()};
}

如果 DTensor 是非紧凑的，CPU 与 GPU 之间搬运时就会触发一次 layout conversion，把 padding 的通道补上或去掉。这个转换由 transfer_to_rank / fetch_from_rank 自动处理，用户无感知。

六、与 PyTorch 的对比

说到这里，可能有人会问：PyTorch 不是也有 channels_last 吗？它不也支持 NHWC 吗？

是的，但路径完全不同。

PyTorch 默认是 NCHW。torch.channels_last 是在 NCHW 主路径之上提供的一种可选内存格式，用户需要显式调用 .to(memory_format=torch.channels_last)。它带来了几个问题：

• 不是所有算子都原生支持 channels_last。某些算子在 NHWC 输入下会内部转回 NCHW，产生额外的 transpose；
• 与 eager 的交互复杂。动态图里算子一个接一个地跑，布局转换点不固定，性能波动大；
• 静态优化受限。torch.compile 可以做布局传播，但图断裂（graph break）和动态形状会让它难以保证端到端 NHWC。

Tech-Renaissance 的选择是：从第一层开始就只认 NHWC。没有可选格式，没有运行时转换，所有算子、所有预处理、所有显存布局都围绕 NHWC 设计。这不是说 NHWC 普遍优于 NCHW，而是在我们的目标场景——GPU 上追求极致训练吞吐——下，统一 NHWC 是最简洁、最可控的方案。

七、小结

张量看起来是深度学习框架里最”无聊”的部件，但正是这些”无聊”的细节，决定了框架能不能跑得快、能不能稳、能不能扩展到多卡。

Tech-Renaissance 的张量设计可以概括为四句话：

1. Shape 固定 4D NHWC，统一所有算子的维度语义；
2. Tensor 是 CPU 端紧凑容器，首地址 256 对齐，只负责搬运；
3. DTensor 是设备端纯描述符，不持内存，通过 offset 在多卡间共享同一份布局图纸；
4. 双轨 stride 分离 CUDA 对齐与 CPU 紧凑，让 GPU 算子走最优路径，CPU 搬运保持简单。

而这些设计的共同底座，就是 MemoryPlan——那个把显存按语义分区、一遍线性布局、保证 256 字节对齐的”显存图纸系统”。DTensor 的 offset 从哪来、slot_bytes 怎么定、多卡布局如何一致，最终都要落到 MemoryPlan 上。

下一篇，我们把视角转向数据侧，看看 Tech-Renaissance 如何用同一套抽象兼容 MNIST、CIFAR 与 ImageNet。