القائمة
Triton Ascend 性能优化实战经验。从批量自适应搜索中提炼的通用优化模式，覆盖 tile 调优方法论、内存加载优化、reduction 优化、隐式广播、多 Pass 合并、数据访问重构等。
يُحدَّد Prompt الذي يراجع المصدر أولًا بشكل افتراضي. يمكنك التبديل إلى أمر مباشر أو تنزيل نسخة محلية.
اقرأ SKILL.md وأي ملفات مرافقة يعرضها SkillsMP قبل أن تقرر التثبيت.
التثبيت باستخدام Codex أو Claude انسخ هذا Prompt والصقه في Codex أو Claude أو مساعد آخر ليراجع صفحة Skill ويثبّتها لك.
يتجاوز الأمر المباشر Prompt المخصّص للمراجعة. افحص المصدر قبل تشغيله.
npx skills add https://github.com/mindspore-ai/akg --skill triton-ascend-performance-improvementيبقى الأمر في سطر واحد. مرّر أفقيًا لمراجعته كاملًا قبل النسخ.
تفضّل نسخة محلية؟ نزّل الملفات المتاحة حاليًا لدى SkillsMP.
矩阵乘法矩阵乘法 A[M, K] @ B[K, N] = C[M, N]中，大K维度矩阵乘法(K>>M,N)优化：针对M/N较小但K极大(如M=N=256,K=131072)的场景，Split-K切分K维度并行化、Workspace+Reduce替代全局同步，实现显著性能提升
Triton Ascend hard API restrictions and forbidden syntax. MUST-follow rules that apply to every kernel: forbidden control flow (return/break/continue/lambda/while), tensor slice/index restrictions, scalar conversion rules, BLOCK_SIZE upper bound. Violating any of these produces a compile or runtime error on Ascend.
Triton Ascend 性能优化通用策略: BLOCK_SIZE 选择 (1024-2048 for elementwise, must be <65536), grid configuration (use VEC_CORE_NUM / CUBE_CORE_NUM, 2D/3D grid for matmul / conv / reduce, 1D grid + inner loop for elementwise / pointwise), 256B alignment for memory transfers, autotune block-size patterns, fp16 / fp32 precision conversion. Bind via keywords like matmul, elementwise, reduce, block_size, grid, autotune, alignment, fp16, fp32, tile, interleaved-loop, cube-core, vec-core.
استنادا إلى تصنيف SOC المهني
عرض SKILL.md
| name | triton-ascend-performance-improvement |
| description | Triton Ascend 性能优化实战经验。从批量自适应搜索中提炼的通用优化模式，覆盖 tile 调优方法论、内存加载优化、reduction 优化、隐式广播、多 Pass 合并、数据访问重构等。 |
| category | improvement |
| version | 1.0.0 |
| metadata | {"case_type":"improvement","backend":"ascend","dsl":"triton_ascend"} |
tile 尺寸需满足硬件存储约束（具体容量参考传入的硬件信息文档）：
CUBE 路径 (matmul / tl.dot)：tile 必须能放入 L0A/L0B/L0C
BLOCK_M × BLOCK_K × sizeof(dtype) ≤ L0A 容量VEC 路径 (elementwise / reduce)：所有活跃 tensor 需放入 UB
BLOCK_SIZE × sizeof(dtype) × 活跃tensor数 × multi_buffer系数 ≤ UB 容量tl.where 产生的临时缓冲）也占用 UB调优策略：从较大 tile 开始尝试，遇到 ub overflow / cbuf overflow 编译错误时逐级缩小。配合 @triton.autotune 自动选优。
在 tl.load 时直接应用 mask 和填充值，而非先无条件加载再用 tl.where 筛选：
# 次优：先加载再 where（多一次中间操作，可能触发 vsel 编译错误）
tile = tl.zeros([BLOCK_SIZE], dtype=tl.float32)
tile = tl.where(mask, tl.load(ptr + offsets), 0.0)
# 推荐：直接在 load 时应用 mask
tile = tl.load(ptr + offsets, mask=mask, other=0.0)
需要将低维 tensor 广播到高维时，用维度扩展（[:, None]）利用 Triton 的隐式广播，避免 tl.broadcast_to 创建临时矩阵：
# 次优：显式展开为完整矩阵
a_broadcast = tl.broadcast_to(a_tile[:, None], (BLOCK_M, BLOCK_N))
c_tile = a_broadcast * b_tile
# 推荐：隐式广播
c_tile = a_tile[:, None] * b_tile
core_sum = 0.0 局部累加，避免 tensor 索引问题tl.atomic_add(out_ptr, core_sum) 一次写入在 host 端预先计算 tensor 的 stride 并作为参数传入 kernel，而非在 kernel 内部计算，可以减少传参错误并辅助编译器优化：
# host 端
stride_am, stride_ak = A.stride(0), A.stride(1)
kernel[grid](A_ptr, ..., stride_am, stride_ak, ...)
当算子对同一数据进行多次独立遍历时（如 softmax 的 max→exp_sum→normalize，或 topk 的多次扫描），应合并为单次遍历。每减少一次 HBM 读取，理论上可获得等比例加速。多 pass 合并对归约维度较小（能放入单个 BLOCK）的场景效果最佳。
当算子涉及非连续的 strided 访问（如需要访问相邻元素的配对计算），应重构数据加载策略：
flat_idx % D + 跨 stride 的随机访问Ascend 硬件对非连续访问有显著性能惩罚，重构后性能差距可达数倍至数十倍。
对于 LayerNorm/RMSNorm 类算子，两阶段方案（第一遍统计 mean/var，第二遍归一化）通常优于单 pass（在线统计）。 原因：单 pass 循环体内活跃 tensor 增多，UB 压力增大，编译器流水线优化效率下降。实测中 2-pass 比单 pass 快约 20%。