triton_ascend入门学习三性能调优

以_layer_norm_fwd_fused算子为例

msprof op --kernel-name=_layer_norm_fwd_fused python3 ./triton-ascend/third_party/ascend/tutorials/03-layer-norm.py

export LD_LIBRARY_PATH=$ASCEND_TOOLKIT_HOME/tools/simulator/Ascend910B4/lib:$LD_LIBRARY_PATH

msprof op simulator --kernel-name=_layer_norm_fwd_fused --soc-version=Ascend910B4 --core-id=5 python3 ./triton-ascend/third_party/ascend/tutorials/03-layer-norm.py  
#4c8g会跑的相当慢 换16c第1个数据类型跑完后也很慢 而且第1个数据类型打印的tensor都是0,感觉不太对 指定core-id有bug? 但不指定也是0。 可能是不应该用cann8.5自带的msprof! 不对，msopgen才是有开源仓的，msprof还没有。未解。论坛达人解释仿真不必看数据准确性。

#解决simulator后到mindstudio insight里面无source的问题：
export LD_LIBRARY_PATH=$ASCEND_TOOLKIT_HOME/tools/simulator/Ascend910B4/lib:$LD_LIBRARY_PATH

TRITON_DISABLE_LINE_INFO=0 msprof op simulator --kernel-name=_layer_norm_fwd_fused --soc-version=Ascend910B4 --core-id=5 python3 /home/atomgit/triton-ascend/third_party/ascend/tutorials/03-layer-norm.py 

方法一：通过上板Profiling分析流水情况（某csv文件）

每条流水线的利用率理想情况下为100%，不到就可能有提升空间（理论不要太当真~）。

的Vector流水的利用率aiv_vec_ratio小于10%，判断未充分发挥算力；

Scalar流水的利用率aiv_scalar_ratio大于60%，判断Scalar是最长的流水：算子源码中是否对标量值进行复杂的运算，SIMD微架构更适合多数据并行计算；另一种可能是，一部分指令在硬件上不支持特定的数据类型，Triton将向量计算退化为标量计算。

MTE2搬运场景：三个输入的shape分别为(128,128)、(128,1)、(128,1)，数据类型为float16。当前算法为Two-pass方法，因此有三次X的搬入，以及W、B的各一次搬入，由此可以计算出总共需要搬运的数据量，计算出理论值为 sizeof(float16) * (128 * 128 * 3 + 128 + 128) / 1.8 TB/s ≈ 0.1991 us （按照1 TB = $10^{12}$ Byte来计算），与实际性能数据aiv_mte2_time存在比较大的差距。经分析，输入数据的总大小小于UB的空间（A2型号为192KB）。因此MTE2时间过长可能是Tiling计算得到的基本块太小，导致发射冗余的搬运指令。

方法二：通过上板Profiling分析Tiling情况

硬件有48个Vector核，发了过多的Block（Block Dim > 48），造成Host调度开销过大。

方法三：通过仿真流水图分析流水情况

Vector核的SCALAR以及FLOWCTRL的指令饱和，是否存在过多标量计算以及不支持的向量化操作。另一方面，Vector核的相关流水（veccore0的MTE2、VECTOR等）有规律性的断流现象，即大量无操作的空白段，是否存在基本块切分过小等因素导致断流。优化方向为流水优化，其次结合Tiling优化和内存优化等手段进一步提升Vector流水利用率。

方法四：通过代码热点分析情况

左侧源码 load接口语句 x = tl.load(X + cols, mask=cols < N, other=0.).to(tl.float32) 对应到右侧的一组汇编指令中（已筛选仅显示代码行相关指令，并按运行拍数Cycles降序排序），标量指令占比高 (ST_XD_XN_IMM dtype:B8,XD:X13,XN:X31,IMM)，不符合load作为访存接口应该MTE占比高的情况。

示例：i64/i32 的compare在npu上无法启用vector导致向量计算转为标量计算

【描述】i64/i32 的cmp在npu上无法启用vector，退化为scalar计算效率降低；通过转化为fp32来利用vec_cast和vec_cmp实现vector操作加速。

【注意】在tl.load和tl.save中的mask使用cmp功能，大部分情况下编译器可以自动优化为vec操作，本例中tl.where则需要手动转换。

核心原因：tl.where 中的整数比较退化为标量计算

第55行正常（计算 mean 循环）：

a = tl.load(X + cols, mask=cols < N, other=0.).to(tl.float32)
_mean += a  # 直接累加

第62-63行（计算 variance 循环）：

x = tl.load(X + cols, mask=cols < N, other=0.).to(tl.float32)
x = tl.where(cols < N, x - mean, 0.)  # ← 问题在这里！

• 第63行的 tl.where(cols < N, ...) 中的比较操作退化为 scalar 计算
• cols 是 tl.arange(0, BLOCK_SIZE) 返回的 int64 类型
• N 是标量参数
• 这个 int64 的比较在 NPU 上不支持向量化，导致大量 scalar 指令

解决方案：将 cols 转换为 float32 类型，使比较操作能够向量化：

# 计算方差
_var = tl.zeros([BLOCK_SIZE], dtype=tl.float32)
for off in range(0, N, BLOCK_SIZE):
    cols = off + tl.arange(0, BLOCK_SIZE)
    x = tl.load(X + cols, mask=cols < N, other=0.).to(tl.float32)
    # 优化：将 cols 转换为 float32 以启用向量化比较
    cols_f32 = cols.to(tl.float32)
    x = tl.where(cols_f32 < N, x - mean, 0.)
    _var += x * x

总结