基于Atlas 800I A2平台的MSELoss 性能优化实践与经验总结
·
作者:昇腾实战派
知识地图:https://blog.csdn.net/Lumos_Lovegood/article/details/161601003
背景概述
在深度学习模型训练中,MSELoss 作为常用的损失函数,其计算性能直接影响整体训练效率。本文针对大规模向量(如 33.5M 元素)下的 MSELoss 计算,从基线实现中存在的性能瓶颈出发,逐步优化,最终在昇腾硬件上实现接近甚至超越 PyTorch 原生实现的吞吐性能。文章记录了优化思路、关键改进及性能数据,旨在为类似场景下的算子优化提供参考。
基线问题
原始实现(mse_loss-base.py)采用变长 Grid 策略:
grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']), )
每个 program 仅处理一个 block,并直接通过 tl.atomic_add 将结果写入输出。当 N=33.5M 时,共启动 32768 个 program,每个 program 均竞争同一个 atomic_add 地址,导致严重的写冲突和性能下降。
优化历程
1. Grid-Strided Loop(核心优化)
| 维度 | 基线 | 优化后 |
|---|---|---|
| Grid 策略 | 变长 ceil(N/BLOCK_SIZE) |
固定 num_vector_cores |
| Atomic 竞争 | 32768 次竞争 | 56 次(每个 program 循环累加后一次写入) |
| 耗时占比 | 向量等待 81.5% | 向量等待 38~66% |
本质优化:将数千次 kernel 启动和数万次 atomic_add 合并为 1 次 kernel 启动 + 56 次写入,大幅降低竞争开销。
2. BLOCK_SIZE 自适应
| N | 基线 | 优化后 |
|---|---|---|
| ≤16K | 256 | 1024 |
| ≤128K | 512 | 2048 |
| ≤4M | 512 | 4096 |
| >4M | 1024 | 8192 |
依据 UB 192KB 容量上限,最大化每次 DMA 搬运粒度,摊薄固定开销,提升整体吞吐。
3. tl.store 替代 atomic_add
将 tl.atomic_add 改为 tl.store(output + pid, ...) + host 侧 sum()。L2 write hit 从 0% 提升到 100%,但由于输出数据仅占总数据量的 0.002%,对总吞吐影响可忽略。
Profiler 指标解读
| 指标 | 含义 | 理想值 |
|---|---|---|
mte2_ratio |
L1 ↔ UB 搬运占时间比 | > 80% |
mte3_ratio |
GM ↔ UB 搬运占时间比 | < 1%(数据在 L1 中复用) |
vec_ratio |
向量计算占时间比 | < 30%(MSELoss 是 memory-bound) |
vec_wait |
向量单元等待比例 | < 40%(越高越空闲) |
L2 write hit |
L2 写命中率 | 越高越好 |
优化瓶颈
当前 Atlas 800I A2 上最大瓶颈是 MTE2(DMA 管道)占空比:
可尝试方向:
- 增大 BLOCK_SIZE 至 16384/32768(受 UB 限制,需确认)
- 双缓冲/软件流水线(Triton DSL 中需编译器支持)
- 使用向量 load 指令替代 DMA(需 Triton Ascend 后端支持)
更多推荐


所有评论(0)