内存访问优化
块大小选择原理
块大小需要根据算子类型和硬件存储层级来平衡:
- VEC 类算子(element-wise、reduce、softmax 等):数据需放入 UB(192KB/VEC),
BLOCK_SIZE * sizeof(dtype)需小于 UB 可用容量,同时兼顾计算并行度。过小并行度不足,过大溢出 UB - CUBE 类算子(matmul、attention 等):左矩阵放 L0A(* KB),右矩阵放 L0B(* KB),结果放 L0C(* KB),具体参考硬件信息文档:
m0 * k0 * sizeof(A.dtype) ≤ * KB(L0A)k0 * n0 * sizeof(B.dtype) ≤ * KB(L0B)m0 * n0 * sizeof(C.dtype) ≤ * KB(L0C)
- 所有数据传输按 256 Bytes 对齐,BLOCK_SIZE 为 32 的倍数最优
2D 数据:优先 tl.make_block_ptr
A_block_ptr = tl.make_block_ptr(
base=A_ptr, shape=(M, K), strides=(stride_am, stride_ak),
offsets=(pid_m * BLOCK_M, 0), block_shape=(BLOCK_M, BLOCK_K), order=(1, 0),
)
a = tl.load(A_block_ptr, boundary_check=(0, 1))
# 移动指针
A_block_ptr = tl.advance(A_block_ptr, (0, BLOCK_K))
连续内存:一维访问
非连续张量先 .contiguous() 转换,再用一维 ptr + offsets 访问:
class ModelNew(torch.nn.Module):
def forward(self, x):
if not x.is_contiguous():
x = x.contiguous()
out = torch.empty_like(x)
n = x.numel()
grid = (triton.cdiv(n, BLOCK_SIZE),)
kernel[grid](x, out, n, BLOCK_SIZE=1024)
return out
一维访问比 stride 计算效率更高,推荐优先使用。
对齐要求
- Ascend 256B 对齐: element-wise / reduce 算子
- Ascend 512B 对齐: MatMul 切分
- 数据搬运带宽上限约 256*256B,据此设计搬运策略
要点
- 优先
.contiguous()+ 一维访问 - 连续内存访问效率远高于 stride 计算开销