mindspore-ai/akg - Page 2

triton-ascend-example-matmul

标准矩阵乘法的完整 Triton Ascend 实现示例。展示 2D 分块(tiling)、K 维循环累加、2D mask 处理、Cube Core 利用等关键模式。当生成 matmul 类算子时可参考此示例的代码结构。

Source text: Chinese

triton-ascend-example-relu

ReLU 逐元素算子的完整 Triton Ascend 实现示例。展示向量化逐元素操作的标准模式：1D 分块遍历、mask 边界处理、交错循环。当生成 elementwise 类算子时可参考此示例的代码结构。

Source text: Chinese

triton-ascend-example-softmax

Softmax 归约算子的完整 Triton Ascend 实现示例。展示三阶段归约模式（求 max → 求 sum(exp) → 归一化）、分块累加、标量累加器精度提升等技巧。当生成 reduce 类算子时可参考此示例的代码结构。

Source text: Chinese

triton-ascend-examples-mindspore

MindSpore 框架下 Triton Ascend 内核的集成示例，展示 MindSpore 自定义算子注册、Primitive 定义、tensor 传入传出等标准写法。当目标框架为 mindspore 时应导入此示例作为代码结构参考。

Source text: Chinese

triton-ascend-grid-config

Grid/Block 配置策略，包括核数选择、并行度调优、二次切分和大 shape 算子处理方案。适用于需要确定 kernel 启动参数、优化多核并行效率、或处理超大规模数据的内核代码生成场景

Source text: Chinese

triton-ascend-attention

适用于注意力(attention)机制类算子的优化指南。当算子的核心计算是 Transformer 风格的注意力运算时应选择此指南，典型算子包括：self_attention, cross_attention, multi_head_attention, flash_attention, scaled_dot_product_attention, causal_attention, masked_attention 等。涵盖 QKV 矩阵乘分块、在线 softmax、因果 mask 处理、Flash…

Source text: Chinese

cuda-c-api

CUDA C 编程接口完整参考手册

Source text: Chinese

cuda-c-basics

CUDA C 核心概念、内核结构和标准编程模式

Source text: Chinese

cuda-c-examples-torch

PyTorch + CUDA C 完整集成示例代码

Source text: Chinese

cuda-c-optimization

CUDA C 性能优化、数值稳定性和调试排查

Source text: Chinese

cuda-c-patterns

CUDA C 三大编程模式：向量操作、归约、矩阵乘法

Source text: Chinese

tilelang-cuda-api

TileLang CUDA API 完整参考手册，适用于需要查阅具体 API 用法、了解函数参数含义的任意 TileLang CUDA 内核代码生成场景

Source text: Chinese

tilelang-cuda-basics

TileLang CUDA 核心概念、内核结构和标准编程模式

Source text: Chinese

tilelang-cuda-examples-torch

PyTorch + TileLang CUDA 完整示例代码

Source text: Chinese

tilelang-cuda-memory

TileLang CUDA 内存访问优化策略，包括 T.alloc_shared/fragment 使用、数据布局优化、合并访存和 Bank Conflict 避免技巧。适用于内存带宽受限、需要优化数据搬运效率的 TileLang 内核性能优化场景

Source text: Chinese

tilelang-cuda-optimization

TileLang CUDA 性能优化通用策略、最佳实践和调试技巧汇总。适用于需要提升 TileLang 内核性能、遇到编译/运行错误需要排查、或需要了解 TileLang 平台限制的内核代码生成和优化场景

Source text: Chinese

tilelang-cuda-patterns

TileLang CUDA 核心编程模式（逐元素、归约、矩阵乘法、GEMV）的标准实现范式和代码模板。适用于需要快速确定算子属于哪种编程模式、或需要了解 TileLang 各模式基本代码结构的内核代码生成场景

Source text: Chinese

tilelang-cuda-synchronization

TileLang CUDA 同步规范，包括 T.sync_threads() 使用规则、线程安全最佳实践和死锁预防策略。适用于编写涉及共享内存访问、多线程协作、或需要避免同步死锁的 TileLang 内核代码生成场景

Source text: Chinese

triton-ascend-case-elemwise-broadcast-2d

2D广播除法优化：小维度完整处理不切分（循环外加载复用），通过固定NUM_BLOCKS实现核间并行（40核），核内SUB_M控制粒度平衡UB利用率，适用于broadcast轴大但非broadcast轴小的2D场景

Source text: Chinese

triton-ascend-case-elemwise-broadcast-3d

跨轴3D广播优化（最后一维很小）：采用两阶段kernel策略（先broadcast展开+reshape为2D，再标准多核处理）提升向量化效率，适用于跨轴broadcast且最后一维特别小（<20）导致向量化效果差的场景

Source text: Chinese

triton-ascend-case-elemwise-cast

大shape类型转换（int8→fp16）优化：通过二次切分（BLOCK_SIZE+TILE_SIZE）提高UB利用率，核数2048时性能最优，适用于shape较大（百万级元素）的elementwise类型转换场景

Source text: Chinese

triton-ascend-case-elemwise-concat

Slice+Concat融合算子优化：通过精确切片加载（只load需要部分）和索引计算拼接（避免cat指令）避免中间结果存储和多次内存访问，适用于多输入需要切片后拼接的融合算子场景

Source text: Chinese

triton-ascend-case-elemwise-zeros

小shape张量创建（zeros/arange/full）优化：通过减少核数避免多核启动和调度开销，单核处理性能优于多核并行，适用于shape较小（千级元素）的elementwise创建张量场景

Source text: Chinese

triton-ascend-case-index-histogram

直方图统计（histogram）优化：预排序+二分查找降低算法复杂度（O(n×m)→O(n log n + m log n)，性能提升19倍），转换为float32调用Vec Core硬件加速排序，适用于大规模统计类操作（50万+元素）

Source text: Chinese

triton-ascend-case-index-put

索引赋值（index_put）优化：批量加载索引数据到UB后循环内通过get_element复用（避免重复访问全局内存），显著降低内存访问延迟，适用于需要在循环中多次访问同一片数据的不规则内存访问场景

Source text: Chinese

triton-ascend-case-matmul-swizzle2d

大矩阵乘法Swizzle2D优化：固定核心数启动（grid=20而非所有块）+Swizzle2D块重排（GROUP_SIZE=4）提升缓存局部性，根据M/N比例自适应选择分组方向，适用于大规模矩阵乘法（千万级元素）的Ascend NPU场景

Source text: Chinese

triton-ascend-case-reduction-amax-medium

中等规模归约（amax）优化：计算重组（循环内累加、循环外归约）减少归约次数，grid=40等于核数时性能最优（25.73us），适用于非reduce轴中等、reduce轴较大（千万级元素）的2D归约场景

Source text: Chinese

triton-ascend-case-reduction-amax-small

极小规模归约（amax）优化：单核处理（grid=1）优于多核并行（2.16us vs 3.51us），避免并行化带来的调度开销，适用于数据规模很小（<1000元素）的归约场景

Source text: Chinese

triton-ascend-case-reduction-amin-large

极大规模1D归约（amin）优化：二次切分避免超UB+计算重组减少归约次数，网格数接近AI Core数量（grid=32）、UB用满、无尾块时性能最优（9.61us），适用于极大规模1D数据（400万级元素）的全量归约场景

Source text: Chinese

triton-ascend-case-reduction-amin-medium

大规模2D归约（amin）reduce轴很大优化：在优先占满UB前提下为reduce轴分配较大切分尺寸（BLOCK_SIZE_N=16384最优），减少循环次数但需权衡单次迭代负载，适用于非reduce轴中等、reduce轴很大（50万级元素）的场景

Source text: Chinese

triton-ascend-case-reduction-amin-small

中等规模1D归约（amin）优化：适中并行度（grid=8时最优2.21us），存在最优平衡点（过小导致单块负载过重、过大引入调度开销），适用于中等规模1D数据（6万级元素）的全量归约场景

Source text: Chinese

triton-ascend-case-reduction-mean-large

大规模reduce最后根轴（mean）行二次切分优化：每个kernel计算多行减少线程块数量、kernel内二次切分避免超UB，grid=40且SUB切分不含尾块时性能最优（16.00us），尾块计算会显著降低性能，适用于非reduce轴中等、reduce轴较大的2D归约场景

Source text: Chinese

triton-ascend-case-reduction-mean-medium

中等规模reduce第一根轴（mean）优化：计算重组减少归约次数，网格规模略小于AI Core数量且避免尾块时性能最佳（grid=32最优9.98us），适用于reduce第一根轴、两轴均中等（百万级元素）的2D归约场景

Source text: Chinese

triton-ascend-case-reduction-prod-small

小规模reduce第一根轴（prod）优化：使用自定义mul函数配合tl.reduce实现连乘（triton无prod接口），最优网格数明显小于AI Core数量（grid=16最优2.15us），过高并行度反而因调度开销降低性能，适用于shape较小（10万级元素）的reduce第一根轴场景

Source text: Chinese

triton-ascend-case-reduction-sum-fused

Reduction+Elementwise融合算子优化：先逐元素操作再归约，行二次切分+计算重组，grid=40且SUB切分不含尾块时性能最优（47.58us），融合优化逻辑以reduce为主，适用于需要先逐元素计算再reduce的融合场景

Source text: Chinese

triton-ascend-case-reduction-sum-large

大规模归约（sum）非reduce轴很大优化：计算重组减少归约次数，在优先占满UB前提下为reduce轴分配较大切分尺寸（BLOCK_SIZE_N=1024最优685.65us），适用于非reduce轴非常大（6万+）、reduce轴中等（千级）的2D归约场景

Source text: Chinese

triton-ascend-case-reduction-weighted-swiglu

3D融合算子（Weighted SwiGLU Backward）优化：Reshape降维将前两维合并简化并行策略，行二次切分避免超UB，在优先占满UB前提下为reduce轴分配较大切分尺寸，grid数较大时可能性能更优，适用于3D张量逐元素+reduce融合的场景

Source text: Chinese

triton-cuda-api

Triton CUDA API 完整参考手册，包括 tl.load/store、tl.reduce、tl.dot、tl.atomic 等核心函数的签名、参数和使用示例。适用于需要查阅具体 API 用法、了解函数参数含义的任意 Triton CUDA 内核代码生成场景

Source text: Chinese

triton-cuda-attention

Attention 算子的 Triton-CUDA 实现指南。包含经过验证的 Flash Attention 完整示例、各变体（Causal/GQA/MQA/RoPE）的差异改法、在线 Softmax 算法和常见错误

Source text: Chinese

triton-cuda-debugging

Triton CUDA 调试排查清单和常见错误速查表，包括编译错误、运行时错误、精度问题和性能问题的诊断方法。适用于 CUDA 内核代码出现错误需要定位原因、或需要验证代码正确性的调试场景

Source text: Chinese