摘要
本文系统拆解张量布局(Tensor Layout)的本质、分类与工程实现。我们从“多维索引 → 一维内存地址”的映射函数出发,对比 NumPy、PyTorch ATen 以及 CUTLASS/CUTE 等开源实现,解释不同布局的共性、设计原理、数学表示与适用场景。你将看到:为什么 Strided Layout 最通用、分块布局为何更贴近缓存、Swizzle 如何缓解 GPU 共享内存 bank 冲突,以及矩阵乘法专用布局为什么“强但不通用”。
目录
- 引言
- 张量布局的本质
- 张量布局的分类
- Strided Layout:最常用的布局
- 分块布局(Blocked Layout)
- 稀疏布局
- Packed Layout(打包布局)
- Swizzled Layout
- 矩阵乘法专用布局
- 性能分析与比较
- 实现考虑
- 总结
1. 引言
在深度学习与高性能计算中,张量(Tensor)是多维数据的基本载体。张量布局定义了这些多维元素在内存中的排列方式,它会直接影响:
- 计算性能(带宽、吞吐、并行访问效率)
- 内存利用率(连续/非连续、对齐、填充)
- 硬件效率(缓存命中、预取、GPU 合并访问、shared memory bank 冲突)
想优化计算密集型应用,先把“数据到底怎么躺在内存里”搞清楚,往往比盲目改算子更有效。
2. 张量布局的本质
2.1 数学定义
张量布局的核心是一个 映射函数:把 N 维索引空间映射到一维内存偏移空间。
给定形状为 [S=(d_0,d1,\dots,d{N-1})] 的 N 维张量,对任意索引 [i=(i_0,i1,\dots,i{N-1})](其中 [0\le i_k<d_k]),布局函数定义为:
[
\mathrm{offset}=f(i_0,i1,\dots,i{N-1})
]
真实内存地址为:
[
\mathrm{addr}(i)=\mathrm{base}+f(i)\times \mathrm{elem_bytes}
]
其中 [\mathrm{elem_bytes}] 是元素字节数(例如 float32 为 4 字节)。
公式逐项解读:
- [f(i)]:从张量起始位置到目标元素需要跨过的“元素个数”(偏移量),是相对位置而非绝对地址
- [\mathrm{base}]:数据起始地址(绝对地址)
- [\mathrm{elem_bytes}]:单元素字节数
- [\mathrm{addr}(i)]:物理地址 = 基地址 + 偏移 × 元素大小
所有布局(strided、blocked、swizzle、packed…)本质上都是在实现不同的 [f(\cdot)]。
2.2 物理本质
从体系结构视角看,张量布局可以理解为三件事:
-
内存访问模式的抽象
布局规定了数据在内存中的分布规律。规律越强,硬件越容易“帮你优化”(缓存、预取、合并访问等)。
-
连接算法与硬件的桥梁
CPU/GPU/TPU 的内存层次与并行访问特性差异很大。布局要让“算法想怎么读写”和“硬件擅长怎么读写”尽量一致。
-
性能优化的杠杆
很多优化并不是减少计算,而是通过重排数据改变缓存行为、带宽利用率与并行冲突情况,从而把瓶颈从“访存”拉回“计算”。
2.3 布局函数的可逆性
可逆性指:能否从偏移量唯一恢复多维索引。形式化表达为:
对任意偏移量 [\mathrm{offset}],若存在唯一索引 [i] 使得 [f(i)=\mathrm{offset}],则布局在其索引域上可逆。
数学基础:单射(injective)
若 [f(i)=f(j)\Rightarrow i=j],则 [f] 为单射。单射保证“不同索引不会撞到同一偏移”。
但工程上“好用的逆映射”往往需要更强条件:
- 显式公式可计算:能直接算出 [f^{-1}],而不是依赖复杂搜索/查表
- 效率高:逆映射复杂度应接近正向映射
- 边界稳定:对 0、最大偏移等边界有清晰定义
可逆性的实际价值:
- 调试与验证:越界/错读时,从地址反推张量元素更容易定位问题
- 序列化/反序列化:保存数据 + 布局信息即可重建结构
- 反向传播:自动微分需要把梯度内存位置映射回原张量位置
- 内存检查:用于越界检测与一致性验证
现实限制:
- 并非所有布局都需要可逆
- 稀疏布局通常不可逆(“不存储的零”不对应唯一偏移)
- 大多数稠密布局(如 strided)可逆,且对开发很有帮助
2.4 通用性质
无论布局怎么变化,通常都需要这些基本要素:
| 要素 |
描述 |
示例 |
| 索引空间 |
各维度取值范围 |
[0\le i_k<d_k] |
| 映射函数 |
索引 → 偏移的规则 |
[\mathrm{offset}=f(i)] |
| 内存覆盖 |
布局访问的内存区域 |
[\text{覆盖字节数}] |
| 遍历顺序 |
以布局定义的顺序访问元素 |
行优先、列优先、块内优先 |
这些共同构件决定了:只要你能定义 [f(\cdot)],你就定义了一种布局。
3. 张量布局的分类
3.1 按内存连续性分类
| 类型 |
描述 |
适用场景 |
| 连续布局 |
元素在内存中连续排列 |
顺序访问、向量化 |
| 非连续布局 |
元素在内存中有间隔 |
视图操作、转置 |
3.2 按访问模式分类
| 类型 |
描述 |
示例 |
| 正则布局 |
偏移是索引的线性函数 |
Strided Layout |
| 不规则布局 |
需要查表或更复杂计算 |
稀疏布局、Swizzled |
| 分块布局 |
数据按块组织 |
分块矩阵乘法 |
3.3 按硬件优化分类
| 类型 |
描述 |
目标硬件 |
| CPU 优化布局 |
关注缓存行、预取 |
CPU |
| GPU 优化布局 |
关注合并访问、bank 冲突 |
GPU |
| 专用硬件布局 |
匹配特定计算单元 |
TPU、NPU |
分类的意义在于:你选布局不是“哪个好”,而是“哪个更贴合你的访问模式与硬件”。
4. Strided Layout:最常用的布局
4.1 基本定义
Strided Layout 最常见、最通用。它的偏移量计算为:
[
\mathrm{offset}=\sum_{k=0}^{N-1} i_k \cdot s_k
]
其中 [s_k] 是第 [k] 维的步长(stride),通常以“元素个数”计(而非字节)。
解读:
- [i_k]:第 [k] 维索引(从 0 开始)
- [s_k]:沿该维移动 1 个位置,在内存中要跳过的元素数
- 总偏移是各维索引的“加权和”,权重就是 stride
4.2 内存连续性条件
对于形状 [(d0,\dots,d{N-1})] 的张量:
- 行优先(C-order)连续:
[
sk=\prod{t=k+1}^{N-1} d_t
]
- 列优先(F-order)连续:
[
sk=\prod{t=0}^{k-1} d_t
]
两者都保证相邻元素在内存中按某个维度连续排列。
4.3 可逆性及其应用
Strided Layout 通常是单射,并且可用显式公式从偏移反推索引。一个常见的逆映射算法如下:
def offset_to_indices(offset, strides, shape):
"""
将偏移量转换回多维索引
参数:
offset: 内存偏移量(以元素个数为单位)
strides: 各维度步长列表
shape: 张量形状
返回:
多维索引元组
"""
indices = []
remaining = offset
# 从最高维开始处理(步长最大的维度)
for k in range(len(strides)):
# 计算当前维度的索引
index_k = remaining // strides[k]
indices.append(index_k)
# 更新剩余偏移量
remaining = remaining % strides[k]
return tuple(indices)
算法要点:
- remaining 从 offset 开始
- 用整除拿到当前维索引
- 用取余更新 remaining
- 逐维拆解直到结束
应用场景(可逆性带来的工程收益):
| 应用场景 |
具体用途 |
可逆性收益 |
| 内存调试 |
定位越界/错读位置 |
快速确定问题元素 |
| 视图操作 |
reshape、unflatten |
无需复制数据 |
| 序列化 |
保存/加载张量 |
存储结构更清晰 |
4.4 示例:二维矩阵
考虑一个 2×3 的矩阵。
行优先(C-order)布局
- 形状:[(2,3)]
- 步长(按元素):[(3,1)]
- 内存排列:[[a00,a01,a02,a10,a11,a12]]
- 偏移量:[\mathrm{offset}(i,j)=i\cdot 3 + j]
列优先(F-order)布局
- 形状:[(2,3)]
- 步长(按元素):[(1,2)]
- 内存排列:[[a00,a10,a01,a11,a02,a12]]
- 偏移量:[\mathrm{offset}(i,j)=i + j\cdot 2]
两者对比:
| 特性 |
行优先 |
列优先 |
| 内存排列 |
逐行存储 |
逐列存储 |
| 行遍历 |
缓存友好 |
缓存不友好 |
| 列遍历 |
缓存不友好 |
缓存友好 |
| 典型生态 |
C/C++/Python |
Fortran/Matlab |
4.5 PyTorch ATen 实现
PyTorch ATen 中,strided layout 的核心信息就是 sizes/strides:
// 简化版的TensorImpl结构
struct TensorImpl {
// 数据指针
void* data_ptr_;
// 布局信息
std::vector<int64_t> sizes_; // 各维度大小
std::vector<int64_t> strides_; // 各维度步长
// 计算元素偏移
int64_t offset(const std::vector<int64_t>& indices) const {
int64_t offset = 0;
for (int64_t i = 0; i < indices.size(); ++i) {
offset += indices[i] * strides_[i];
}
return offset;
}
// 从偏移量计算索引(需要形状信息)
std::vector<int64_t> indices_from_offset(int64_t offset) const {
std::vector<int64_t> indices(sizes_.size());
int64_t remaining = offset;
for (int64_t i = 0; i < sizes_.size(); ++i) {
indices[i] = remaining / strides_[i];
remaining = remaining % strides_[i];
}
return indices;
}
};
阅读提示:
data_ptr_:数据起始地址
sizes_:形状
strides_:步长
offset():实现 [\sum i_k s_k]
indices_from_offset():体现“可逆性”接口
4.6 转置操作的布局变化
转置通常是 零拷贝:只改变 strides/shape,不复制数据。
import numpy as np
# 创建原始矩阵
A = np.array([[1, 2, 3], [4, 5, 6]], dtype=np.float32)
# 转置操作
A_t = A.T
要点:
- 原始形状:
(2, 3);原始步长:(12, 4) 字节(float32=4 字节)
- 转置后形状:
(3, 2);转置后步长:(4, 12) 字节
- 不复制数据,只改变访问方式
这种“视图”在工程里极其常见:快,但可能导致后续算子因非连续访问而变慢,因此常配合 ascontiguousarray 或框架内部的 contiguous 转换。
5. 分块布局(Blocked Layout)
5.1 设计动机
当矩阵很大时,简单的连续行优先/列优先会让缓存命中率下降:你正在处理的那一小块数据可能很快被冲掉。分块布局把矩阵切成小块(tile),提高空间局部性。
5.2 二维分块布局
对矩阵 [A\in \mathbb{R}^{M\times N}],按块大小 [B_m\times B_n] 划分为多个块。
5.3 内存排列
分块布局常见两种组织方式(块内/块间顺序都可选):
- 块内行优先,块间行优先
- 块内行优先,块间列优先
核心思想是:先把“同一块内的元素”尽量放得更近。
5.4 偏移量计算(二维示例)
设块大小为 [B_m\times B_n],对全局索引 [(i,j)]:
- 块坐标:
[
b_i=\left\lfloor \frac{i}{B_m}\right\rfloor,\quad b_j=\left\lfloor \frac{j}{B_n}\right\rfloor
]
- 块内坐标:
[
r=i\bmod B_m,\quad c=j\bmod B_n
]
- 每行块数:[
n_b=\left\lceil \frac{N}{B_n}\right\rceil
]
若“块间行优先、块内行优先”,则可写成:
[
\mathrm{offset}=(b_i\cdot n_b + b_j)\cdot (B_m\cdot B_n) + r\cdot B_n + c
]
解读:
- 先把 [(i,j)] 拆成“在哪个块 + 块内位置”
- 块的线性编号决定“块起始偏移”
- 再加上块内行优先偏移
5.5 CUTLASS 示例
CUTLASS 中常用分块布局优化 GEMM:
// 定义分块布局
using BlockLayout = cutlass::layout::RowMajorBlocked<32, 4>;
// 使用示例
float data[M * N];
BlockLayout layout(cutlass::MatrixCoord(M, N));
// 计算元素(i, j)的偏移
int offset = layout({i, j}); // 返回元素索引
分块布局的特殊性:
- 主要用于二维矩阵(可推广到高维,但优化收益最稳定的是二维)
- 块大小必须谨慎选择:匹配 CPU 缓存行(常见 64B)或 GPU shared memory 组织
- 小矩阵可能不划算:分块的地址计算与边界处理会带来开销
| 分块布局优势 |
具体表现 |
| 缓存友好 |
局部性更强、命中率更高 |
| 并行友好 |
更适合多线程/多核或 GPU tile 计算 |
| 硬件适配 |
便于对齐缓存行与共享内存 |
6. 稀疏布局
COO 用坐标 + 值存储非零元素:
# COO格式表示稀疏矩阵
indices = torch.tensor([[0, 1, 2], # 行索引
[1, 2, 0]]) # 列索引
values = torch.tensor([3.0, 4.0, 5.0])
shape = (3, 3)
# 对应的稠密矩阵
# [[0, 3, 0],
# [0, 0, 4],
# [5, 0, 0]]
COO 的数学表示可以看作存储三元组 [(r_t, c_t, v_t)],共有 [\mathrm{nnz}] 个非零元素。
6.2 CSR(Compressed Sparse Row)
CSR 用 row_ptr 压缩行边界,用 col_ind 标识列索引:
# CSR格式表示
row_ptr = [0, 1, 2, 3] # 每行的起始位置
col_ind = [1, 2, 0] # 列索引
values = [3.0, 4.0, 5.0]
shape = (3, 3)
对行 [r],其非零元素范围是 [\mathrm{row_ptr}[r]] 到 [\mathrm{row_ptr}[r+1])。
6.3 稀疏布局的内存效率
| 格式 |
存储开销 |
随机访问 |
行操作 |
列操作 |
| COO |
较高(存坐标) |
O(nnz) |
O(nnz) |
O(nnz) |
| CSR |
中等 |
O(log(nnz_row)) |
O(1)(遍历行) |
O(nnz) |
| CSC |
中等 |
O(log(nnz_col)) |
O(nnz) |
O(1)(遍历列) |
其中 [\mathrm{nnz}] 为非零数量。
稀疏布局的特殊性:
- 非单射性:零元素不存储,多个零索引对应“无存储状态”
- 格式选择依赖算子:CSR 擅长按行,CSC 擅长按列
- 压缩收益依赖稀疏度:不够稀疏反而更浪费
- 构建成本高:收集非零并排序/压缩需要额外开销
可逆性分析:
稀疏布局通常 不是单射,因此不可逆:仅凭 offset 很难唯一恢复 [(i,j)],必须依赖额外索引数组。
7. Packed Layout(打包布局)
7.1 量化数据布局
量化网络中常见 8-bit、4-bit 甚至更低精度。Packed Layout 把多个低精度值打包进一个字节/字中,以节省内存与带宽。
7.2 4 位量化示例
两个 4-bit 值打包成 1 字节:
struct Packed4BitLayout {
// 假设矩阵大小为 M×N
int M, N;
int stride; // 以字节为单位的行步长
uint8_t* data; // 数据指针
// 获取元素(i, j)
uint8_t get(int i, int j) {
// 计算字节偏移
int byte_offset = i * stride + (j / 2);
uint8_t packed = data[byte_offset];
// 提取4位值
if (j % 2 == 0) {
return packed & 0x0F; // 低4位
} else {
return (packed >> 4) & 0x0F; // 高4位
}
}
// 设置元素(i, j)
void set(int i, int j, uint8_t value) {
// 计算字节偏移
int byte_offset = i * stride + (j / 2);
uint8_t& packed = data[byte_offset];
// 设置4位值
if (j % 2 == 0) {
packed = (packed & 0xF0) | (value & 0x0F);
} else {
packed = (packed & 0x0F) | ((value & 0x0F) << 4);
}
}
};
要点:
get()/set() 都依赖位运算
- 奇偶列分别落在字节高/低 4 位
- 内存用量显著下降,但访问更复杂
7.3 偏移量计算(示意)
若每个存储单元打包 [p] 个元素,二维情况下可理解为:
- 存储单元偏移:[
\mathrm{byte_offset}=i\cdot \mathrm{stride} + \left\lfloor\frac{j}{p}\right\rfloor
]
- 单元内位置:[
\mathrm{lane}=j\bmod p
]
打包布局的特殊性:
- 位操作开销不可忽视
- 对齐与边界处理更敏感
- 不同硬件对低比特操作支持差异大
- 低精度会引入量化误差(属于算法层面的代价)
| 打包布局优势 |
具体表现 |
| 内存节省 |
4-bit 可显著降低存储需求 |
| 带宽优化 |
相同数据量占用更少传输带宽 |
| 能效提升 |
访存减少通常意味着能耗降低 |
8. Swizzled Layout
8.1 基本概念
Swizzled Layout 通过重映射地址的某些低位,让访问分布更均匀。它主要用于解决 GPU 共享内存(shared memory)bank 冲突。
8.2 Bank 冲突问题
GPU shared memory 通常划分为多个 bank(常见 32)。当同一 warp 的多个线程访问落在同一 bank 的不同地址时,会发生 bank 冲突,导致访问串行化,从而拉低吞吐。
8.3 简单 Swizzle:XOR 重写 bank 位
// 基本的Swizzle函数
inline int swizzle_address(int addr) {
// 假设32个bank,bank索引由地址的5-9位决定
int bank_index = (addr >> 5) & 0x1F;
int swizzled_bank = bank_index ^ ((addr >> 2) & 0x1F);
return (addr & ~(0x1F << 5)) | (swizzled_bank << 5);
}
解读:
- 提取原始 bank 索引(某段地址位)
- 用 XOR 混入其他低位,打散 bank 分布
- 替换回地址,得到 swizzled 地址
- 目标是减少同 warp 内 bank 撞车
8.4 CUTLASS 中的 Swizzle
// 使用Swizzle布局
using SwizzleLayout = cutlass::layout::TensorSwizzle<
cutlass::layout::Swizzle<256, 8>>;
// Swizzle函数通常形式
int swizzle_func(int addr) {
// 多种实现方式,如XOR、置换等
return addr ^ ((addr >> 5) & 0x1F);
}
8.5 数学表示
可把 Swizzle 看成对原始偏移做一个双射变换:
[
\mathrm{offset}' = g(\mathrm{offset})
]
其中 [g] 设计为双射(bijection),确保地址一一对应。
8.6 可逆性分析
Swizzled Layout 往往可逆。对于 XOR,逆就是它本身:
// XOR Swizzle的逆函数(XOR的逆就是自身)
inline int unswizzle_address(int addr) {
// 由于 XOR 的逆就是自身
return swizzle_address(addr);
}
但要注意:可逆不等于“有用”。Swizzle 的主要价值在于提升 shared memory 并行访问效率,运行时很少需要频繁反解地址。
| Swizzle 优势 |
具体表现 |
| 减少 bank 冲突 |
提升共享内存并行度 |
| 提高带宽利用率 |
更有效的访存分布 |
| 硬件适配 |
针对 GPU 内存结构优化 |
9. 矩阵乘法专用布局
9.1 背景
GEMM 优化同时受多因素影响:
- 数据重用(Arithmetic Intensity)
- 全局内存访问合并(Coalesced Access)
- 缓存/共享内存层次利用
布局往往与分块策略、warp 级加载方式、Tensor Core 形状要求一起被“整体设计”。
9.2 CUTLASS 中的 GEMM 布局
// 定义矩阵乘法的布局
using LayoutA = cutlass::layout::ColumnMajorInterleaved<32>;
using LayoutB = cutlass::layout::RowMajorInterleaved<32>;
using LayoutC = cutlass::layout::RowMajor;
// 这些布局优化了:
// 1. 全局内存到共享内存的加载
// 2. 共享内存到寄存器的加载
// 3. Tensor Core的数据布局要求
直观理解:
- A 用列主序交错:让加载更贴合某些 tile/warp 的读取方式
- B 用行主序交错:配合另一个操作数的访问模式
- C 通常用行主序:写回更简单或更匹配后续算子
9.3 Volta Tensor Core 布局(概念说明)
Tensor Core 往往要求操作数以特定片段(fragment)形式组织。布局会强制:
- 每个线程处理多个元素
- warp 级协同加载
- shared memory 排布与 fragment 映射一致
这种布局的特殊性:
- 高度专用:只为特定硬件 + 特定算子服务
- 参数复杂:tile、交错因子、对齐、swizzle 往往耦合
- 移植性差:换架构可能要重调甚至重写
- 实现门槛高:必须理解硬件访存与执行模型
10. 性能分析与比较
10.1 内存访问模式对比
| 布局类型 |
顺序访问 |
随机访问 |
缓存友好 |
向量化友好 |
| 行连续 |
优秀 |
O(1) |
优秀 |
优秀 |
| 列连续 |
优秀(按列) |
O(1) |
一般 |
一般 |
| 分块 |
优秀(块内) |
O(1) |
优秀 |
优秀 |
| 稀疏 COO |
差 |
O(nnz) |
差 |
差 |
| 稀疏 CSR |
一般(按行) |
O(log nnz_row) |
一般 |
一般 |
10.2 计算复杂度比较
| 操作 |
稠密行优先 |
稠密列优先 |
CSR 稀疏 |
COO 稀疏 |
| 获取元素 (i,j) |
O(1) |
O(1) |
O(log nnz_row) |
O(nnz) |
| 矩阵向量乘 |
O(MN) |
O(MN) |
O(nnz) |
O(nnz) |
| 转置 |
O(1)(布局变化) |
O(1)(布局变化) |
O(nnz) |
O(nnz) |
复杂度提示:
O(1):恒定时间,最理想
O(MN):稠密矩阵很难绕开
O(nnz):稀疏典型复杂度
O(log nnz_row):通常靠二分查找实现
11. 实现考虑
11.1 布局泛化
现代张量库需要把布局作为可替换组件(模板/策略类),以统一接口支持多种 layout:
template <typename Layout>
class Tensor {
private:
void* data_;
std::vector<int64_t> shape_;
Layout layout_;
public:
// 泛化的偏移计算
int64_t offset(const std::vector<int64_t>& indices) {
return layout_.offset(indices, shape_);
}
// 泛化的迭代器
Iterator begin() {
return layout_.begin(data_, shape_);
}
// 可逆布局需要实现的接口
std::vector<int64_t> indices_from_offset(int64_t offset) {
return layout_.indices_from_offset(offset, shape_);
}
};
设计原则:
- 模板化/策略化:布局可插拔
- 接口统一:调用方不关心内部细节
- 可逆性可选:不是所有布局都必须实现
indices_from_offset
11.2 布局转换
布局转换常见两种策略:
- 惰性转换:只记录布局变化,不复制数据(典型:转置视图)
- 急切转换:复制数据到新布局(典型:contiguous 化)
import numpy as np
# 惰性转换示例(视图)
x = np.array([[1, 2, 3], [4, 5, 6]])
x_t = x.T # 仅改变strides,不复制数据
# 急切转换示例(复制)
x_contiguous = np.ascontiguousarray(x_t) # 实际复制数据
选择哪种策略,通常取决于:你是更在意“当前操作的开销”,还是更在意“后续大量算子的连续访存收益”。
11.3 布局兼容性检查
逐元素操作不仅要看广播规则,还要看布局是否能高效协同(尤其是不同类型布局混算时):
bool are_layouts_compatible(
const Layout& lhs,
const Layout& rhs,
const std::vector<int64_t>& shape) {
// 检查是否都是strided layout
if (is_strided(lhs) && is_strided(rhs)) {
// 检查广播兼容性
return can_broadcast(lhs.shape(), rhs.shape(), shape);
}
// 其他布局类型需要特殊处理
return false;
}
12. 总结
张量布局是“多维索引在内存中的落点规则”,其本质是索引空间到内存地址空间的映射函数。尽管布局类型多样,它们通常都具备这些共性:
- 定义索引到偏移的映射(核心能力)
- 持有形状信息(决定索引域)
- 支持遍历(迭代访问所有元素)
- 决定性能特性(缓存、预取、并行冲突、带宽利用率)
各布局特殊性回顾:
- Strided Layout:最通用,任意维度,常有显式逆公式
- 分块布局:二维矩阵缓存优化最常见,块大小选择关键
- 稀疏布局:为节省存储而生,通常不可逆,格式强依赖算子
- 打包布局:为量化/带宽优化服务,位操作与对齐更复杂
- Swizzled 布局:GPU 专用,重点解决 shared memory bank 冲突
- 矩阵乘法专用布局:为特定硬件/算子量身定做,强但不通用
“实用逆映射”的标准:
- 逆映射有显式公式(或近似同成本计算)
- 计算复杂度接近正向映射
- 在调试、序列化、视图等场景确实用得上
- 易实现、易验证、边界清晰
如果你正在做张量算子、推理引擎或 GPU kernel 调优,建议把布局当作一等公民来设计:访问模式明确了吗?目标硬件的瓶颈是什么?要不要为了后续算子把视图变成连续?更多相关讨论也可以在 云栈社区 继续延伸。