akg

矩阵乘法矩阵乘法 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.

通过 adaptive_search 或 evolve 搜索式 workflow 生成优化算子。 后台 silent mode 执行，轮询监控进度。

适用于归约(reduce)类算子和含归约子步骤的复合算子（如归一化）的优化指南。典型算子包括：sum, mean, max, min, prod, argmax, argmin, cumsum, cumprod, softmax, logsoftmax, layernorm, rmsnorm, groupnorm, instancenorm, batchnorm, l1norm, l2norm, frobeniusnorm, var, std, average_pooling, sum_pooling 等。特别重要：当归约维度不是最后一维（如 dim=1 归约 shape=[B,F,D1,D2]），需要正确处理多维索引和两阶段归约。包含 PyTorch normalized_shape 多轴归一化语义说明。不适用于纯逐元素运算或矩阵乘法。如果算子是损失函数（先逐元素计算再全局归约），应选择 elementwise-reduce-fused 指南。

CPU C++ 算子核心概念、标准结构模式、KernelBench 代码规范和内嵌扩展方法

ARM CPU 架构性能优化技巧、NEON SIMD 向量化、数值稳定性和调试策略

x64 CPU 架构性能优化技巧、SIMD/AVX 向量化、数值稳定性和调试策略

akg_agents 环境准备。缓存优先，检查 akg_cli + LLM 连通性， 未安装则引导安装，采集硬件/Framework/DSL。 FULL_SETUP 模式额外包含参数确认和运行时依赖安装。

算子算法草图设计 Skill — 负责根据任务需求设计高质量的算法草图（sketch），提供伪代码形式的算法方案、优化建议和实现策略。 支持多种 DSL：triton_cuda、triton_ascend、cpp、cuda_c、tilelang_cuda、pypto。 支持 Hint 模式（参数空间配置）。

算子内核代码生成 Skill — 负责算子实现的全部智力工作：方案讨论、代码生成、基于反馈修改。 支持多种 DSL：triton_cuda、triton_ascend、cpp、cuda_c、tilelang_cuda、pypto。

算子代码验证 Skill — 静态代码检查 + 精度对比验证。 包含两阶段验证：先做零成本静态检查（语法、编译、import、DSL 合规性）， 通过后再对比框架实现与生成实现的输出一致性。 支持多框架（torch / mindspore）、多后端（cuda / ascend / cpu）。

从用户提供的代码或自然语言描述，构建标准化的单文件自包含任务代码。

非reduce轴很小、reduce轴很大的归约优化：将reduce轴映射到多核（而非常规的非reduce轴），使用原子操作跨线程块归约，通过二次切分避免超UB，适用于极端shape比例（M<<N如16×262144）的归约场景

原子操作归约（amin）优化：非reduce轴很小时将reduce轴映射多核，提供循环内/外两种原子操作方案（减少存储vs减少竞争），通过二次切分+计算重组提升性能，适用于M<<N（如16×262144）的极端shape比例场景

Triton Ascend 编程基础，包括核心概念（program_id、block、grid）、内核函数结构、装饰器用法和标准代码模式。适用使用 Triton Ascend、需要了解基本语法结构的任意内核代码生成场景

Triton CUDA 编程基础，包括核心概念（program_id、block、grid）、内核函数结构、装饰器用法和标准代码模式。适用于使用 Triton CUDA、需要了解基本语法结构的任意 CUDA 内核代码生成场景

逐元素算子(element-wise)优化策略，包括 add/mul/relu/sigmoid/tanh/gelu/exp/log 等操作的向量化实现和融合技巧。适用于实现激活函数、逐元素运算、广播操作等向量模式算子的 CUDA 内核代码生成场景

矩阵乘法算子(matmul/bmm/linear)优化策略，包括分块 Tiling、共享内存缓存、Tensor Core 利用和大矩阵处理技巧。适用于实现 GEMM、批量矩阵乘、全连接层等矩阵运算的 CUDA 内核代码生成场景

CUDA GPU 内存访问优化策略，包括共享内存利用、合并访存、Bank Conflict 避免和数据布局优化技巧。适用于内存带宽受限、需要优化全局内存访问效率、或处理大规模数据的 CUDA 内核性能优化场景

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

提交前代码自审工具。检查 rebase 冲突、代码规范（ruff/bandit）、危险函数、SPEC.md 合规性。 生成审查报告写入 .tmp/review/。当用户输入 /akg-review 或要求代码审查、提交前检查时使用。

【仅适用于 vLLM-MindSpore，不适用于 PTA/vLLM-Ascend 等其他框架】稀疏量化模型加载与适配 - 在 vLLM-MindSpore 中加载稀疏量化模型、适配代码实现或修复常见 Bug。Triggers: '稀疏量化', '加载稀疏量化模型', 'W8A8SC', 'sparse quantization', '稀疏量化适配', '稀疏量化加载', 'sparse-quant'。

通用模型算子融合分析 - 分析任意vllm-ascend支持模型的算子路径及融合优化策略，支持生成新的融合算子

后处理优化 - 对vllm-ascend计算结果进行优化处理，包括Triton Kernel优化、NPU算子优化等。

后处理优化 - 对计算结果进行优化处理，包括缓存、并行化、向量化等。

triton-ascend 常见错误修复：UB/CBUF溢出、BiShengIR编译失败、语法限制违反、数值正确性、多维索引分解错误、张量连续性

Ascend NPU 上融合算子的深度优化方法论。覆盖性能天花板分析框架、多 Pass 合并策略、数据访问模式重构、Normalization 两阶段决策、NPU 原生算子评估方法论。适用于 elementwise 融合、归一化融合、softmax+topk 融合、matmul+activation 融合等场景。

Triton Ascend 性能优化实战经验。从批量自适应搜索中提炼的通用优化模式，覆盖 tile 调优方法论、内存加载优化、reduction 优化、隐式广播、多 Pass 合并、数据访问重构等。

triton-cuda常见错误及修复方法，用于代码生成时避免同类问题

Ascend 硬件约束与编译器限制速查。涵盖 CUBE/VEC 存储层级预算计算方法、bishengIR 编译器已知限制、strided access 性能特征。适用于所有 Triton Ascend 算子生成和调试场景。

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

Ascend NPU 内存访问优化策略，包括 UB（统一缓冲区）利用、数据布局优化、合并访存和预取技巧。适用于内存带宽受限、需要优化数据搬运效率、或处理大规模数据的内核代码性能优化场景

适用于同时包含逐元素计算和全局归约两个阶段的复合算子。典型算子包括：损失函数（MSELoss, HuberLoss, HingeLoss, SmoothL1Loss, CrossEntropyLoss, KLDivLoss, CosineSimilarityLoss, TripletMarginLoss 等）、以及自定义的先逐元素变换再全局聚合的算子。这类算子的计算模式是：第一步对张量每个元素独立执行变换（差值、平方、clamp 等），第二步对变换结果做全局或按维度归约（sum/mean）得到标量或低维结果。与纯 elementwise 或纯 reduce 不同，这类算子需要在同一个 kernel 中融合两个阶段以避免中间结果的额外 GM 读写。

适用于纯逐元素(element-wise)类算子的优化指南。当算子的核心计算是对张量每个元素独立执行相同操作、无跨元素依赖时应选择此指南，典型算子包括：relu, sigmoid, tanh, gelu, selu, leaky_relu, elu, swish, softplus, hardsigmoid, hardtanh, softsign, exp, log, sqrt, pow, add, mul, sub, div, abs, neg, clamp, cast(类型转换), where, fill, copy 等。也适用于涉及标量广播(broadcast)的运算。不适用于需要跨元素归约(如 sum/mean/max)或矩阵乘法的算子。如果算子同时包含逐元素计算和全局归约（如损失函数 MSELoss、HuberLoss、HingeLoss），应选择 elementwise-reduce-fused 指南。

适用于矩阵乘法(matmul)类算子的优化指南。当算子的核心计算涉及二维或更高维的矩阵乘法时应选择此指南，典型算子包括：matmul, mm, bmm, linear, gemm, outer_product, einsum(含矩阵乘), conv(转为矩阵乘实现)等。涵盖 Cube Core 使用、分块(tiling)策略、Swizzle 优化、大 K 维处理等关键技巧。不适用于纯逐元素运算或纯归约运算。对于 attention 机制中的 QK^T 和 score*V 矩阵乘，若算子整体是注意力计算，应优先选择 attention 指南。

生成符合 AKG 项目规范的 Issue 描述文件（Bug Report / RFC / Task）。 支持两种模式：用户描述生成、基于分支 diff 自动生成。 生成的 .md 和 .json 文件写入 .tmp/issue/，经规范校验后供用户确认和提交。

基于当前分支与目标分支的 diff，自动生成符合 AKG 项目规范的 PR 描述文件。 生成的 .md 和 .json 文件写入 .tmp/pr/，经规范校验后供用户确认和提交。

双内核调用模式的 Triton Ascend 实现示例。展示在 forward 中先后调用两个 kernel 的标准写法：中间结果缓冲区分配、两次 kernel 启动。适用于需要分阶段计算的融合算子。

LayerNorm 归约算子的完整 Triton Ascend 实现示例。展示两阶段归约模式（统计量计算 → 归一化输出）、标量累加器、分块遍历等技巧。当生成 reduce/normalize 类算子时可参考此示例的代码结构。