已合并
feat: 迁入 evolution 策略库 161 策略卡 + 142 Playbook 至 contrib/ascendc #103
sijiaali创建于 26 天前
feat: 迁入 evolution 策略库 161 策略卡 + 142 Playbook 至 contrib/ascendc #103
已合并
共 401 个文件变更+41362-0
| @@ -1,3 +1,4 @@ | |||
| 1 | # 社区贡献的 Ascend C 知识 | 1 | # 社区贡献的 Ascend C 知识 |
| 2 | 2 | ||
| 3 | - [优化经验](optimizations/index.md)：从受控实验沉淀的 Ascend C 算子优化经验。 | 3 | - [优化经验](optimizations/index.md)：从受控实验沉淀的 Ascend C 算子优化经验。 |
| 4 | +- [Runbook](runbooks/index.md)：Ascend C 算子优化与落地的分步流程。 | ||
| @@ -0,0 +1,63 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'A1: FP32 Intermediate Computation (FP32中间计算)' | ||
| 4 | +description: 参考实现在DAG中显式地进行了类型转换：输入数据类型U (BF16/FP16) 被转换为T (FP32) 后进行所有计算，最后转换回U输出。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- paradigm.elementwise | ||
| 8 | +- topic.mixture_of_experts | ||
| 9 | +- topic.compute.normalization | ||
| 10 | +- topic.optimizer | ||
| 11 | +status: stable | ||
| 12 | +sources: | ||
| 13 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/A1_fp32_intermediate.md | ||
| 14 | + role: primary | ||
| 15 | +created_at: '2026-10-09T00:00:00Z' | ||
| 16 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 17 | +platforms: | ||
| 18 | +- a2 | ||
汪 | |||
| 19 | +--- | ||
| 20 | + | ||
| 21 | +# A1: FP32 Intermediate Computation (FP32中间计算) | ||
| 22 | + | ||
| 23 | +## 核心思想 | ||
| 24 | +参考实现在DAG中显式地进行了类型转换：输入数据类型U (BF16/FP16) 被转换为T (FP32) 后进行所有计算，最后转换回U输出。这种设计确保了累加精度（AdaGrad中accum的累加操作在FP32上进行，避免BF16/FP16的精度损失）、除法精度（sqrt(accum)和除法操作在FP32上进行，减少数值误差）、梯度精度（梯度平方和乘法在更高精度下进行）。参考实现全程使用float32，虽然精度足够但失去了BF16/FP16的内存带宽优势。 | ||
| 25 | + | ||
| 26 | +## 代码骨架 | ||
| 27 | + | ||
| 28 | +// === 改造前（基线）=== | ||
| 29 | +```cpp | ||
| 30 | +// 参考实现仅FP32，无需转换 | ||
| 31 | +AscendC::Add(accumLocal, accumLocal, inputLocal, C); | ||
| 32 | +``` | ||
| 33 | + | ||
| 34 | +// === 改造后（优化实现）=== | ||
| 35 | +```cpp | ||
| 36 | +// 升精度累加 | ||
| 37 | +if constexpr (std::is_same_v<T, float>) { | ||
| 38 | + Add(sumBufLocal, sumBufLocal, inputLocal, len); | ||
| 39 | +} else { | ||
| 40 | + LocalTensor<float> castBufLocal = castBuf.Get<float>(); | ||
| 41 | + Cast(castBufLocal, inputLocal, RoundMode::CAST_NONE, len); | ||
| 42 | + Add(sumBufLocal, sumBufLocal, castBufLocal, len); | ||
| 43 | +} | ||
| 44 | + | ||
| 45 | +// 降精度输出 | ||
| 46 | +Cast(outputLocal, sumBufLocal, RoundMode::CAST_RINT, len); | ||
| 47 | +``` | ||
| 48 | + | ||
| 49 | +## 关键修改点 | ||
| 50 | + | ||
| 51 | +1. 预期收益: 有效避免FP16/BF16多次累加产生的精度损失，数值稳定性提升 | ||
| 52 | + | ||
| 53 | +## 常见陷阱 | ||
| 54 | + | ||
| 55 | +⚠️ 需要额外的castBuf和sumBuf(FP32)存储空间 | ||
| 56 | +⚠️ 需要在Tiling阶段计算每个输出点的kernel大小 | ||
| 57 | +⚠️ 增加了类型转换操作的开销，但通常远小于内存带宽节省的收益 | ||
| 58 | + | ||
| 59 | +## 代码搜索关键词 | ||
| 60 | + | ||
| 61 | +```bash | ||
| 62 | +grep -n "BUFFER_NUM\|InitBuffer\|TQue\|tileSize\|ubFactor\|Tiling\|BLOCK_DIM\|GetBlockNum" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 63 | +``` | ||
| @@ -0,0 +1,56 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'A2: Welford Numerically Stable Mean/Var (Welford数值稳定算法)' | ||
| 4 | +description: 在训练模式下，BatchNorm需要计算当前batch的均值和方差用于前向传播，同时更新running_mean和running_var。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.compute.normalization | ||
| 8 | +status: stable | ||
| 9 | +sources: | ||
| 10 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/A2_welford_algorithm.md | ||
| 11 | + role: primary | ||
| 12 | +created_at: '2026-10-09T00:00:00Z' | ||
| 13 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 14 | +platforms: | ||
| 15 | +- a2 | ||
| 16 | +--- | ||
| 17 | + | ||
| 18 | +# A2: Welford Numerically Stable Mean/Var (Welford数值稳定算法) | ||
| 19 | + | ||
| 20 | +## 核心思想 | ||
| 21 | +在训练模式下，BatchNorm需要计算当前batch的均值和方差用于前向传播，同时更新running_mean和running_var。参考实现正确处理了样本方差（除以N-1）与总体方差（除以N）的区别：1）Batch统计量：使用batchVarScale系数将总体方差转换为样本方差（无偏估计），当patternR0 * patternR1 > 1时，batchVarScale = N / (N-1)；2）Running统计量：running_mean和running_var使用指数移动平均（EMA）更新，不涉及N-1校正；3）特殊值处理：当patternR0 * patternR1 == 1时（单样本场景），避免除以0，直接设置batchVarScale = 1.0。 | ||
| 22 | + | ||
| 23 | +## 代码骨架 | ||
| 24 | + | ||
| 25 | +// === 改造后（优化实现）=== | ||
| 26 | +```cpp | ||
| 27 | +if (amsgrad_) { | ||
| 28 | + PipeBarrier<PIPE_V>(); | ||
| 29 | + Max(dataOutLocal[maxGradOutOffset_], dataLocal[maxGradNormOffset_], | ||
| 30 | + dataOutLocal[expAvgSqOffset_], dataCount); | ||
| 31 | + PipeBarrier<PIPE_V>(); | ||
| 32 | + Sqrt(dataLocal[varOffset_], dataOutLocal[maxGradOutOffset_], dataCount); | ||
| 33 | + PipeBarrier<PIPE_V>(); | ||
| 34 | + Muls(dataLocal[expAvgOffset_], dataLocal[varOffset_], biasCorrection2Sqrt_, dataCount); | ||
| 35 | +} else { | ||
| 36 | + Sqrt(dataLocal[varOffset_], dataOutLocal[expAvgSqOffset_], dataCount); | ||
| 37 | + // ... | ||
| 38 | +} | ||
| 39 | +``` | ||
| 40 | + | ||
| 41 | +## 关键修改点 | ||
| 42 | + | ||
| 43 | +1. 用细粒度同步替代粗粒度 SyncAll | ||
| 44 | +2. 预期收益: 解决Adam发散问题；提升训练稳定性；支持Adam/AMSGrad切换 | ||
| 45 | + | ||
| 46 | +## 常见陷阱 | ||
| 47 | + | ||
| 48 | +⚠️ 额外的Max操作；需要维护max_grad_norm状态 | ||
| 49 | +⚠️ 需要额外的乘法操作；需要特殊处理单样本场景 | ||
| 50 | +⚠️ 算法复杂度增加，需要维护额外的累加器状态 | ||
| 51 | + | ||
| 52 | +## 代码搜索关键词 | ||
| 53 | + | ||
| 54 | +```bash | ||
| 55 | +grep -n "BUFFER_NUM\|InitBuffer\|TQue\|tileSize\|ubFactor\|Tiling\|BLOCK_DIM" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 56 | +``` | ||
| @@ -0,0 +1,64 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'A3: Rounding Mode Control (舍入模式控制)' | ||
| 4 | +description: 参考实现在FP16版本中采用了多层次类型转换策略来平衡精度与性能。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.attention | ||
| 8 | +- topic.programming_model.execution_unit.cube | ||
| 9 | +- topic.compute.quantization | ||
| 10 | +status: stable | ||
| 11 | +sources: | ||
| 12 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/A3_rounding_mode.md | ||
| 13 | + role: primary | ||
| 14 | +created_at: '2026-10-09T00:00:00Z' | ||
| 15 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 16 | +platforms: | ||
| 17 | +- a2 | ||
| 18 | +--- | ||
| 19 | + | ||
| 20 | +# A3: Rounding Mode Control (舍入模式控制) | ||
| 21 | + | ||
| 22 | +## 核心思想 | ||
| 23 | +参考实现在FP16版本中采用了多层次类型转换策略来平衡精度与性能。计算流程为：1) 将FP16输入Cast到FP32进行除法运算（避免FP16除法精度损失）；2) 在FP32域进行Muls和Adds操作；3) 通过Cast到int32并指定RoundMode::CAST_RINT实现四舍五入；4) 将结果Cast回FP32进行后续计算；5) 最终输出时Cast到FP16。这种精细的类型转换链确保了关键计算步骤的精度，同时减少了不必要的精度转换开销。相比之下，参考实现使用了Round指令（AscendC::Round(xqRoundedLocal, xScaledLocal, this->tileSize)），虽然简单但可能无法提供与Cast+RoundMode::CAST_RINT相同的精度和灵活性。 | ||
| 24 | + | ||
| 25 | +## 代码骨架 | ||
| 26 | + | ||
| 27 | +// === 改造前（基线）=== | ||
| 28 | +```cpp | ||
| 29 | +// 统一使用CAST_NONE | ||
| 30 | +AscendC::Cast(addedLocal, xLocal, AscendC::RoundMode::CAST_NONE, this->tileLength); | ||
| 31 | +AscendC::Cast(outputLocal, addedLocal, AscendC::RoundMode::CAST_NONE, this->tileLength); | ||
| 32 | +``` | ||
| 33 | + | ||
| 34 | +// === 改造后（优化实现）=== | ||
| 35 | +```cpp | ||
| 36 | +// 精细化舍入控制 | ||
| 37 | +// BF16: 使用RINT保证精度 | ||
| 38 | +Cast(x1_fp32, x1Local, RoundMode::CAST_NONE, numCol); | ||
| 39 | +Cast(x2_fp32, x2Local, RoundMode::CAST_NONE, numCol); | ||
| 40 | +PipeBarrier<PIPE_V>(); | ||
| 41 | +Add(x1_fp32, x1_fp32, x2_fp32, numCol); | ||
| 42 | +PipeBarrier<PIPE_V>(); | ||
| 43 | +Cast(xLocal, x1_fp32, RoundMode::CAST_RINT, numCol); // 输出用RINT | ||
| 44 | + | ||
| 45 | +// FP16: 输入输出都用NONE | ||
| 46 | +Cast(yLocal, xFp32, RoundMode::CAST_NONE, numCol); | ||
| 47 | +``` | ||
| 48 | + | ||
| 49 | +## 关键修改点 | ||
| 50 | + | ||
| 51 | +1. 用细粒度同步替代粗粒度 SyncAll | ||
| 52 | +2. 预期收益: 中间计算使用截断减少指令开销，最终输出使用四舍五入提高精度 | ||
| 53 | + | ||
| 54 | +## 常见陷阱 | ||
| 55 | + | ||
| 56 | +⚠️ 需要针对不同数据类型选择不同策略，增加代码复杂度 | ||
| 57 | +⚠️ 增加4条指令和同步开销 | ||
| 58 | +⚠️ 需要额外的条件分支，可能略微影响性能 | ||
| 59 | + | ||
| 60 | +## 代码搜索关键词 | ||
| 61 | + | ||
| 62 | +```bash | ||
| 63 | +grep -n "tileSize\|ubFactor\|Tiling\|BLOCK_DIM\|GetBlockNum\|coreNum\|blockIdx\|SplitCore" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 64 | +``` | ||
| @@ -0,0 +1,51 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'A4: SetFlag/WaitFlag Event Sync (事件同步保证精度)' | ||
| 4 | +description: 参考实现在关键计算节点插入了PipeBarrier<PIPE_V>()进行向量管道同步，确保数据转换和计算操作的正确顺序。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.attention | ||
| 8 | +- paradigm.elementwise | ||
| 9 | +status: stable | ||
| 10 | +sources: | ||
| 11 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/A4_pipeline_barrier.md | ||
| 12 | + role: primary | ||
| 13 | +created_at: '2026-10-09T00:00:00Z' | ||
| 14 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 15 | +platforms: | ||
| 16 | +- a2 | ||
| 17 | +--- | ||
| 18 | + | ||
| 19 | +# A4: SetFlag/WaitFlag Event Sync (事件同步保证精度) | ||
| 20 | + | ||
| 21 | +## 核心思想 | ||
| 22 | +参考实现在关键计算节点插入了PipeBarrier<PIPE_V>()进行向量管道同步，确保数据转换和计算操作的正确顺序。这种细粒度的同步控制在BF16的Cast-Compute-Cast流程中尤为重要，防止了数据竞争和乱序执行导致的精度问题。PipeBarrier确保前一步操作完全完成后才执行下一步，虽然会引入一定的性能开销，但对于精度敏感的场景是必要的。参考实现完全没有使用pipeline barrier，在复杂场景下可能出现同步问题。 | ||
| 23 | + | ||
| 24 | +## 代码骨架 | ||
| 25 | + | ||
| 26 | +// === 改造后（优化实现）=== | ||
| 27 | +```cpp | ||
| 28 | +PipeBarrier<PIPE_V>(); | ||
| 29 | +Cast(float32Tensor, dataLocal[index * maxCastDataCount], RoundMode::CAST_NONE, dataCount); | ||
| 30 | +PipeBarrier<PIPE_V>(); | ||
| 31 | +op(float32Tensor[offset], float32Tensor, scalarVal, dataCount); | ||
| 32 | +PipeBarrier<PIPE_V>(); | ||
| 33 | +Cast(outLocal[index * maxCastDataCount], float32Tensor[offset], RoundMode::CAST_RINT, dataCount); | ||
| 34 | +``` | ||
| 35 | + | ||
| 36 | +## 关键修改点 | ||
| 37 | + | ||
| 38 | +1. 用细粒度同步替代粗粒度 SyncAll | ||
| 39 | +2. 预期收益: 确保计算顺序正确性，防止数据竞争，保证数值精度 | ||
| 40 | + | ||
| 41 | +## 常见陷阱 | ||
| 42 | + | ||
| 43 | +⚠️ 引入同步开销，可能降低流水线效率 | ||
| 44 | +⚠️ 增加了同步开销，但精度收益显著 | ||
| 45 | +⚠️ 同步会引入轻微性能损失，但对于正确性是必需的 | ||
| 46 | + | ||
| 47 | +## 代码搜索关键词 | ||
| 48 | + | ||
| 49 | +```bash | ||
| 50 | +grep -n "SyncAll\|PipeBarrier\|ExecuteTask\|PRELOAD\|SyncAll\|SetFlag\|WaitFlag\|PipeBarrier" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 51 | +``` | ||
| @@ -0,0 +1,58 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'A5: Numerical Safety & Special Value Handling (数值安全与特殊值处理)' | ||
| 4 | +description: 参考实现在计算输出值时，通过Compare+Select组合实现了对NaN（Not a Number）的特殊处理。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.attention | ||
| 8 | +- topic.data_movement.memory_access | ||
| 9 | +status: stable | ||
| 10 | +sources: | ||
| 11 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/A5_softmax_stability.md | ||
| 12 | + role: primary | ||
| 13 | +created_at: '2026-10-09T00:00:00Z' | ||
| 14 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 15 | +platforms: | ||
| 16 | +- a2 | ||
| 17 | +--- | ||
| 18 | + | ||
| 19 | +# A5: Numerical Safety & Special Value Handling (数值安全与特殊值处理) | ||
| 20 | + | ||
| 21 | +## 核心思想 | ||
| 22 | +参考实现在计算输出值时，通过Compare+Select组合实现了对NaN（Not a Number）的特殊处理。具体逻辑是：Compare(maskTemp, xLocal, xLocal, CMPMODE::EQ, calCount) - 如果x是NaN，则x != x，比较结果为false；然后通过Select(curTemp, maskTemp, curTemp, 0.0f, SELMODE::VSEL_TENSOR_SCALAR_MODE, ...)，当mask为false（即x为NaN）时，输出0.0f。这种处理确保了量化结果在遇到NaN输入时不会传播NaN，而是输出一个确定值（0.0f）。这对于神经网络的量化训练尤为重要，因为NaN的传播可能导致训练不稳定。 | ||
| 23 | + | ||
| 24 | +## 代码骨架 | ||
| 25 | + | ||
| 26 | +// === 改造前（基线）=== | ||
| 27 | +```cpp | ||
| 28 | +// 基础形状检查，无索引范围检查 | ||
| 29 | +``` | ||
| 30 | + | ||
| 31 | +// === 改造后（优化实现）=== | ||
| 32 | +```cpp | ||
| 33 | +OP_CHECK_IF((xDimNum != NCDHW_DIM_NUM) || (gradDimNum != NCDHW_DIM_NUM), | ||
| 34 | + OP_LOGE(context_->GetNodeName(), "Input dim num should equal = %lu", NCDHW_DIM_NUM), | ||
| 35 | + return false); | ||
| 36 | + | ||
| 37 | +OP_CHECK_IF(maxPoolGradParams.diDim * maxPoolGradParams.hiDim * maxPoolGradParams.wiDim > MAX_INT32, | ||
| 38 | + OP_LOGE(context_->GetNodeName(), "Shape too big"), | ||
| 39 | + return ge::GRAPH_FAILED); | ||
| 40 | + | ||
| 41 | +maxPoolGradParams.dGcd = Gcd(maxPoolGradParams.doDim, maxPoolGradParams.diDim); | ||
| 42 | +``` | ||
| 43 | + | ||
| 44 | +## 关键修改点 | ||
| 45 | + | ||
| 46 | +1. 预期收益: 保证极端输入下的数值稳定性和正确性 | ||
| 47 | + | ||
| 48 | +## 常见陷阱 | ||
| 49 | + | ||
| 50 | +⚠️ 增加tiling阶段的开销 | ||
| 51 | +⚠️ 需要更多的UB内存用于FP32缓冲区 | ||
| 52 | +⚠️ 增加UB内存占用(4字节 vs 2字节) | ||
| 53 | + | ||
| 54 | +## 代码搜索关键词 | ||
| 55 | + | ||
| 56 | +```bash | ||
| 57 | +grep -n "BUFFER_NUM\|InitBuffer\|TQue\|tileSize\|ubFactor\|Tiling\|BLOCK_DIM\|GetBlockNum" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 58 | +``` | ||
| @@ -0,0 +1,69 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'A6: High-Precision rsqrt via Newton-Raphson (高精度rsqrt)' | ||
| 4 | +description: Layer Norm的归一化需要计算1/sqrt(var + eps)。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.compute.normalization | ||
| 8 | +status: stable | ||
| 9 | +sources: | ||
| 10 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/A6_high_precision_rsqrt.md | ||
| 11 | + role: primary | ||
| 12 | +created_at: '2026-10-09T00:00:00Z' | ||
| 13 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 14 | +platforms: | ||
| 15 | +- a2 | ||
| 16 | +--- | ||
| 17 | + | ||
| 18 | +# A6: High-Precision rsqrt via Newton-Raphson (高精度rsqrt) | ||
| 19 | + | ||
| 20 | +## 核心思想 | ||
| 21 | +Layer Norm的归一化需要计算1/sqrt(var + eps)。参考实现使用Newton-Raphson迭代提高rsqrt精度：初始近似y0 = rsqrt(x)（硬件指令，精度约8bit），第一次迭代y1 = y0 * (1.5 - 0.5 * x * y0^2)，第二次迭代y2 = y1 + 0.5 * y1 * (1 - x * y1^2)。相比单次sqrt+div，迭代方法避免了除法指令的高延迟，同时获得更高的精度（接近FP32全精度）。 | ||
| 22 | + | ||
| 23 | +## 代码骨架 | ||
| 24 | + | ||
| 25 | +// === 改造前（基线）=== | ||
| 26 | +```cpp | ||
| 27 | +// 参考实现 - 传统方法 | ||
| 28 | +AscendC::Adds(rmsLocal, meanSqLocal, this->eps, 1); | ||
| 29 | +AscendC::Sqrt(rmsLocal, rmsLocal, 1); | ||
| 30 | +float rmsVal = rmsLocal.GetValue(0); | ||
| 31 | +AscendC::Reciprocal(rmsLocal, rmsLocal, 1); | ||
| 32 | +float invRms = rmsLocal.GetValue(0); | ||
| 33 | +``` | ||
| 34 | + | ||
| 35 | +// === 改造后（优化实现）=== | ||
| 36 | +```cpp | ||
| 37 | +// 参考实现 - 牛顿迭代 | ||
| 38 | +static constexpr float SCALAR1 = -0.5; | ||
| 39 | +static constexpr float SCALAR2 = 1.5; | ||
| 40 | + | ||
| 41 | +Div(r, one, var, pregLoop); // 初始近似 | ||
| 42 | +Sqrt(y, r, pregLoop); | ||
| 43 | +Muls(t, var, SCALAR1, pregLoop); | ||
| 44 | +Mul(t, t, y, pregLoop); | ||
| 45 | +Mula(t1, t, y, pregLoop); | ||
| 46 | +Mul(rstd, y, t1, pregLoop); | ||
| 47 | + | ||
| 48 | +// 特殊值处理 | ||
| 49 | +CompareScalar(cmpRegZero, var, POS_INF, pregLoop); | ||
| 50 | +Select(rstd, scalarZero, rstd, cmpRegZero); | ||
| 51 | +CompareScalar(cmpRegInf, var, float(0.0), pregLoop); | ||
| 52 | +Select(rstd, scalarInf, rstd, cmpRegInf); | ||
| 53 | +``` | ||
| 54 | + | ||
| 55 | +## 关键修改点 | ||
| 56 | + | ||
| 57 | +1. 预期收益: 更高的精度和更好的性能，特殊值处理确保数值稳定性; 更高的数值精度和指令吞吐，正确处理特殊边界值 | ||
| 58 | + | ||
| 59 | +## 常见陷阱 | ||
| 60 | + | ||
| 61 | +⚠️ 代码更复杂，需要理解牛顿迭代原理 | ||
| 62 | +⚠️ 增加约6条Vector指令，适合训练场景 | ||
| 63 | +⚠️ 增加约8-10条Vector指令 | ||
| 64 | + | ||
| 65 | +## 代码搜索关键词 | ||
| 66 | + | ||
| 67 | +```bash | ||
| 68 | +grep -n "SCALAR1\|SCALAR2" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 69 | +``` | ||
| @@ -0,0 +1,52 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'A7: Index & Boundary Safety (索引与边界安全处理)' | ||
| 4 | +description: 参考实现提供了完整的InferShapeForMultiScaleDeformableAttentionGrad函数，精确计算三个输出(grad_value, | ||
| 5 | + grad_sampling_loc, grad_attn_weight)的shape。 | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- topic.compute.indexing | ||
| 9 | +status: stable | ||
| 10 | +sources: | ||
| 11 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/A7_index_boundary.md | ||
| 12 | + role: primary | ||
| 13 | +created_at: '2026-10-09T00:00:00Z' | ||
| 14 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 15 | +platforms: | ||
| 16 | +- a2 | ||
| 17 | +--- | ||
| 18 | + | ||
| 19 | +# A7: Index & Boundary Safety (索引与边界安全处理) | ||
| 20 | + | ||
| 21 | +## 核心思想 | ||
| 22 | +参考实现提供了完整的InferShapeForMultiScaleDeformableAttentionGrad函数，精确计算三个输出(grad_value, grad_sampling_loc, grad_attn_weight)的shape。这确保了在动态shape场景下，输出tensor的shape能够正确推导，避免了运行时错误。相比之下，参考实现的shape推导只是简单地复制输入shape，可能导致shape不匹配的错误。 | ||
| 23 | + | ||
| 24 | +## 代码骨架 | ||
| 25 | + | ||
| 26 | +// === 改造前（基线）=== | ||
| 27 | +```cpp | ||
| 28 | +this->Input("indices") | ||
| 29 | + .DataType({ge::DT_INT64}) | ||
| 30 | +``` | ||
| 31 | + | ||
| 32 | +// === 改造后（优化实现）=== | ||
| 33 | +```cpp | ||
| 34 | +this->Input("argmax") | ||
| 35 | + .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32}) | ||
| 36 | +``` | ||
| 37 | + | ||
| 38 | +## 关键修改点 | ||
| 39 | + | ||
| 40 | +1. 预期收益: 减少50%索引数据传输带宽，提升内存受限场景性能 | ||
| 41 | + | ||
| 42 | +## 常见陷阱 | ||
| 43 | + | ||
| 44 | +⚠️ 索引范围限制在2^31-1以内（足够覆盖实际场景） | ||
| 45 | +⚠️ 每次迭代增加一次比较操作 | ||
| 46 | +⚠️ 需要上层框架保证索引合法性 | ||
| 47 | + | ||
| 48 | +## 代码搜索关键词 | ||
| 49 | + | ||
| 50 | +```bash | ||
| 51 | +grep -n "tileSize\|ubFactor\|Tiling\|BLOCK_DIM\|GetBlockNum\|coreNum\|blockIdx\|SplitCore" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 52 | +``` | ||
| @@ -0,0 +1,72 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'A8: Quantization-Specific Precision (量化专用精度处理)' | ||
| 4 | +description: 参考实现通过Maxs和Mins指令精确控制量化边界，确保量化值严格限制在[quant_min, quant_max]范围内。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.compute.quantization | ||
| 8 | +status: stable | ||
| 9 | +sources: | ||
| 10 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/A8_quant_precision.md | ||
| 11 | + role: primary | ||
| 12 | +created_at: '2026-10-09T00:00:00Z' | ||
| 13 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 14 | +platforms: | ||
| 15 | +- a2 | ||
| 16 | +--- | ||
| 17 | + | ||
| 18 | +# A8: Quantization-Specific Precision (量化专用精度处理) | ||
| 19 | + | ||
| 20 | +## 核心思想 | ||
| 21 | +参考实现通过Maxs和Mins指令精确控制量化边界，确保量化值严格限制在[quant_min, quant_max]范围内。具体实现为：Maxs(curInt32Temp, curInt32Temp, static_cast<int32_t>(this->quantMin), calCount)后接Mins(curInt32Temp, curInt32Temp, static_cast<int32_t>(this->quantMax), calCount)。这种实现确保了即使在浮点计算误差的情况下，量化结果也不会超出指定范围。此外，参考实现还在Cast操作中指定了不同的RoundMode（CAST_RINT用于四舍五入，CAST_ROUND用于向最近偶数舍入），根据不同的计算阶段选择最合适的舍入模式，进一步优化数值精度。 | ||
| 22 | + | ||
| 23 | +## 代码骨架 | ||
| 24 | + | ||
| 25 | +// === 改造前（基线）=== | ||
| 26 | +```cpp | ||
| 27 | +// 参考实现：单一计算路径（未区分bias类型） | ||
| 28 | +AscendC::Cast(inputHalfLocal, inputLocal, AscendC::RoundMode::CAST_NONE, tileSize); | ||
| 29 | +AscendC::Cast(inputF32Local, inputHalfLocal, AscendC::RoundMode::CAST_NONE, tileSize); | ||
| 30 | +AscendC::Adds(dequantLocal, inputF32Local, -zeroPointVal, tileSize); | ||
| 31 | +AscendC::Muls(dequantLocal, dequantLocal, scaleVal, tileSize); | ||
| 32 | +AscendC::Add(outputLocal, dequantLocal, biasLocal, tileSize); | ||
| 33 | +``` | ||
| 34 | + | ||
| 35 | +// === 改造后（优化实现）=== | ||
| 36 | +```cpp | ||
| 37 | +// 参考实现：INT32 bias的计算路径（先加后转） | ||
| 38 | +__aicore__ inline void ComputeDequantWithBiasInt32(...) { | ||
| 39 | + for (int64_t i = 0; i < inRows; i++) { | ||
| 40 | + Add(xLocal[i * nAlign_], xLocal[i * nAlign_], biasLocal_, nAlign_); // INT32加法 | ||
| 41 | + } | ||
| 42 | + PipeBarrier<PIPE_V>(); | ||
| 43 | + LocalTensor<float> xLocalFp32 = xLocal.template ReinterpretCast<float>(); | ||
| 44 | + Cast(xLocalFp32, xLocal, RoundMode::CAST_NONE, inRows * nAlign_); // 转为FP32 | ||
| 45 | + // ... 后续乘法 | ||
| 46 | +} | ||
| 47 | + | ||
| 48 | +// 其他类型bias的计算路径（先转后加） | ||
| 49 | +__aicore__ inline void ComputeDequantWithBiasFloat(...) { | ||
| 50 | + LocalTensor<float> xLocalFp32 = xLocal.template ReinterpretCast<float>(); | ||
| 51 | + Cast(xLocalFp32, xLocal, RoundMode::CAST_NONE, inRows * nAlign_); // 先转为FP32 | ||
| 52 | + // ... 乘法后再加bias | ||
| 53 | + Add(xLocalFp32[i * nAlign_], xLocalFp32[i * nAlign_], biasLocal_, nAlign_); | ||
| 54 | +} | ||
| 55 | +``` | ||
| 56 | + | ||
| 57 | +## 关键修改点 | ||
| 58 | + | ||
| 59 | +1. 用细粒度同步替代粗粒度 SyncAll | ||
| 60 | +2. 预期收益: 保证INT32 bias场景的数值精度，避免舍入误差累积 | ||
| 61 | + | ||
| 62 | +## 常见陷阱 | ||
| 63 | + | ||
| 64 | +⚠️ 需要维护两套计算路径 | ||
| 65 | +⚠️ 每行需要独立计算统计量，增加计算开销 | ||
| 66 | +⚠️ 增加类型转换指令；需要理解不同RoundMode的语义 | ||
| 67 | + | ||
| 68 | +## 代码搜索关键词 | ||
| 69 | + | ||
| 70 | +```bash | ||
| 71 | +grep -n "GetBlockNum\|coreNum\|blockIdx\|SplitCore\|ComputeDequantWithBiasInt32\|ComputeDequantWithBiasFloat" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 72 | +``` | ||
| @@ -0,0 +1,69 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'D1: Mixed Precision Architecture (混合精度架构)' | ||
| 4 | +description: 参考实现仅支持 float32 输入类型，这在实际 AI 推理场景中是不足的。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.compute.composite | ||
| 8 | +status: stable | ||
| 9 | +sources: | ||
| 10 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/D1_mixed_precision.md | ||
| 11 | + role: primary | ||
| 12 | +created_at: '2026-10-09T00:00:00Z' | ||
| 13 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 14 | +platforms: | ||
| 15 | +- a2 | ||
| 16 | +--- | ||
| 17 | + | ||
| 18 | +# D1: Mixed Precision Architecture (混合精度架构) | ||
| 19 | + | ||
| 20 | +## 核心思想 | ||
| 21 | +参考实现仅支持 float32 输入类型，这在实际 AI 推理场景中是不足的。现代大语言模型和推理框架通常使用 FP16 或 BF16 进行计算以减少内存带宽和存储需求。参考实现通过模板参数 typename T 支持 half (FP16) 和 bfloat16_t (BF16) 两种输入类型。关键技术细节包括：类型特化的计算精度（当输入为 half 时，计算类型 calcType 提升为 float，避免 FP16 精度损失；当输入为 bfloat16_t 时，直接以 BF16 计算）、条件类型推导使用 std::conditional 实现编译期类型选择、MicroAPI 自动处理中 FP16 输入会自动通过 Cast 转换为 FP32 进行计算而 BF16 保持原格式。 | ||
| 22 | + | ||
| 23 | +## 代码骨架 | ||
| 24 | + | ||
| 25 | +// === 改造前（基线）=== | ||
| 26 | +```cpp | ||
| 27 | +this->Input("x") | ||
| 28 | + .ParamType(REQUIRED) | ||
| 29 | + .DataType({ge::DT_FLOAT}) | ||
| 30 | + .Format({ge::FORMAT_ND}) | ||
| 31 | + .UnknownShapeFormat({ge::FORMAT_ND}); | ||
| 32 | +this->Output("y") | ||
| 33 | + .ParamType(REQUIRED) | ||
| 34 | + .DataType({ge::DT_FLOAT}) | ||
| 35 | + .Format({ge::FORMAT_ND}) | ||
| 36 | + .UnknownShapeFormat({ge::FORMAT_ND}); | ||
| 37 | +``` | ||
| 38 | + | ||
| 39 | +// === 改造后（优化实现）=== | ||
| 40 | +```cpp | ||
| 41 | +this->Input("x") | ||
| 42 | + .ParamType(DYNAMIC) | ||
| 43 | + .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_BF16}) | ||
| 44 | + .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 45 | + .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 46 | + .AutoContiguous(); | ||
| 47 | +this->Input("scalar") | ||
| 48 | + .ParamType(REQUIRED) | ||
| 49 | + .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32, ge::DT_FLOAT}) | ||
| 50 | + .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 51 | + .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 52 | + .AutoContiguous(); | ||
| 53 | +``` | ||
| 54 | + | ||
| 55 | +## 关键修改点 | ||
| 56 | + | ||
| 57 | +1. 预期收益: 支持更广泛的业务场景，减少数据类型转换开销，提高端到端性能; 覆盖更广的精度需求，支持2-4倍内存带宽优化，适应不同硬件平台和精度要求; 一个算子支持多种数据类... | ||
| 58 | + | ||
| 59 | +## 常见陷阱 | ||
| 60 | + | ||
| 61 | +⚠️ 代码复杂度增加，需要维护多类型模板实例 | ||
| 62 | +⚠️ 模板代码复杂度增加，编译时间可能增加 | ||
| 63 | +⚠️ 需要在运行时检查属性，略微增加开销 | ||
| 64 | + | ||
| 65 | +## 代码搜索关键词 | ||
| 66 | + | ||
| 67 | +```bash | ||
| 68 | +grep -n "BUFFER_NUM\|InitBuffer\|TQue\|tileSize\|ubFactor\|Tiling\|BLOCK_DIM\|GetBlockNum" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 69 | +``` | ||
| @@ -0,0 +1,68 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'D2: Template Kernel Type Dispatch (模板化内核类型分发)' | ||
| 4 | +description: 参考实现通过C++模板机制实现了对多种数据类型的统一支持。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- paradigm.elementwise | ||
| 8 | +- topic.compute.normalization | ||
| 9 | +- topic.compute.composite | ||
| 10 | +- topic.programming_model.execution_unit.scalar | ||
| 11 | +status: stable | ||
| 12 | +sources: | ||
| 13 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/D2_template_kernel.md | ||
| 14 | + role: primary | ||
| 15 | +created_at: '2026-10-09T00:00:00Z' | ||
| 16 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 17 | +platforms: | ||
| 18 | +- a2 | ||
| 19 | +--- | ||
| 20 | + | ||
| 21 | +# D2: Template Kernel Type Dispatch (模板化内核类型分发) | ||
| 22 | + | ||
| 23 | +## 核心思想 | ||
| 24 | +参考实现通过C++模板机制实现了对多种数据类型的统一支持。在Kernel端，SparseToDenseSimt类使用三个模板参数：IDX_T（indices数据类型）、Y_T（values/output数据类型）、COMP_T（内部计算使用的数据类型）。这种设计允许灵活组合不同的数据类型，如indices使用int32或int64，而values可以是float16、float32、bfloat16、int8~int64、bool等多种类型。Host端通过VALUE_DTYPE和INDICES_DTYPE集合进行严格的类型校验，确保只有支持的类型组合才能通过编译。这种设计不仅提高了算子的通用性，还为不同精度需求的场景提供了优化空间——例如在推理场景下可以使用float16减少内存带宽，在训练场景下使用float32保证精度。 | ||
| 25 | + | ||
| 26 | +## 代码骨架 | ||
| 27 | + | ||
| 28 | +// === 改造前（基线）=== | ||
| 29 | +```cpp | ||
| 30 | +// 参考实现仅支持float32 | ||
| 31 | +AscendC::GlobalTensor<float> varGm; | ||
| 32 | +AscendC::GlobalTensor<float> accumGm; | ||
| 33 | +// ... 所有tensor都是float类型 | ||
| 34 | +``` | ||
| 35 | + | ||
| 36 | +// === 改造后（优化实现）=== | ||
| 37 | +```cpp | ||
| 38 | +// 参考实现 - 模板化类型处理 | ||
| 39 | +template <typename U, typename T = float> | ||
| 40 | +struct ApplyAdagradDUpdateSlots { | ||
| 41 | + using OpCopyInVar = Bind<Vec::CopyIn<U>, Placeholder::In0<U>>; | ||
| 42 | + using OpVarCast = Bind<Vec::Cast<T, U, 0>, OpCopyInVar>; | ||
| 43 | + // ... 在T类型上进行计算 | ||
| 44 | +}; | ||
| 45 | + | ||
| 46 | +// Tiling中针对不同类型的实例化 | ||
| 47 | +if (this->varDtype_ == ge::DT_FLOAT16) { | ||
| 48 | + eleBaseTiling.DoTiling<ApplyAdagradDOp::ApplyAdagradDUpdateSlots<half>::OpDag>(...); | ||
| 49 | +} else if (this->varDtype_ == ge::DT_BF16) { | ||
| 50 | + eleBaseTiling.DoTiling<ApplyAdagradDOp::ApplyAdagradDUpdateSlots<bfloat16_t>::OpDag>(...); | ||
| 51 | +} | ||
| 52 | +``` | ||
| 53 | + | ||
| 54 | +## 关键修改点 | ||
| 55 | + | ||
| 56 | +1. 预期收益: 支持BF16/FP16/FP32三种数据类型，内存带宽节省50%(BF16)或25%(FP16)，同时通过FP32中间计算保持精度 | ||
| 57 | + | ||
| 58 | +## 常见陷阱 | ||
| 59 | + | ||
| 60 | +⚠️ 增加了模板代码复杂度，编译时间可能增加，需要为每种类型生成独立的二进制代码 | ||
| 61 | +⚠️ 增加了代码重复（两个DAG定义），需要更多编译时间生成多个模板实例 | ||
| 62 | +⚠️ 模板元编程增加编译时间；代码可读性降低 | ||
| 63 | + | ||
| 64 | +## 代码搜索关键词 | ||
| 65 | + | ||
| 66 | +```bash | ||
| 67 | +grep -n "BUFFER_NUM\|InitBuffer\|TQue\|tileSize\|ubFactor\|Tiling\|BLOCK_DIM\|GetBlockNum" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 68 | +``` | ||
| @@ -0,0 +1,65 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'D3: TilingKey-Driven Type Dispatch (TilingKey驱动类型分发)' | ||
| 4 | +description: 参考实现通过模板类FakeQuantAffineCachemaskFp32和FakeQuantAffineCachemaskFp16实现了对FP16和FP32两种数据类型的完整支持。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.programming_model.execution_unit.cube | ||
| 8 | +- topic.mixture_of_experts | ||
| 9 | +- topic.compute.quantization | ||
| 10 | +- topic.compute.composite | ||
| 11 | +status: stable | ||
| 12 | +sources: | ||
| 13 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/D3_tilingkey_dispatch.md | ||
| 14 | + role: primary | ||
| 15 | +created_at: '2026-10-09T00:00:00Z' | ||
| 16 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 17 | +platforms: | ||
| 18 | +- a2 | ||
| 19 | +--- | ||
| 20 | + | ||
| 21 | +# D3: TilingKey-Driven Type Dispatch (TilingKey驱动类型分发) | ||
| 22 | + | ||
| 23 | +## 核心思想 | ||
| 24 | +参考实现通过模板类FakeQuantAffineCachemaskFp32和FakeQuantAffineCachemaskFp16实现了对FP16和FP32两种数据类型的完整支持。这种设计允许根据输入数据类型自动选择最优的计算路径，充分利用昇腾芯片对不同精度计算的硬件支持。在算子定义文件中，通过配置多组DataType实现了运行时数据类型的灵活选择。Tiling阶段通过SetTilingKeyMode函数根据数据类型设置不同的Tiling Key（FP16_MODE=2, FP32_MODE=1），从而在kernel入口处分发到对应的模板实例。这种策略的优势在于：1) 避免了运行时的类型判断开销；2) 允许针对每种数据类型进行专门的优化；3) 通过编译期多态实现零开销抽象。值得注意的是，FP16和FP32的实现细节有所不同，例如FP16实现中使用了更多的Cast操作进行精度转换，而FP32实现则尽量保持原生精度计算。 | ||
| 25 | + | ||
| 26 | +## 代码骨架 | ||
| 27 | + | ||
| 28 | +// === 改造前（基线）=== | ||
| 29 | +```cpp | ||
| 30 | +// 参考实现仅支持FP32 | ||
| 31 | +this->Input("x") | ||
| 32 | + .DataType({ge::DT_FLOAT}); | ||
| 33 | +``` | ||
| 34 | + | ||
| 35 | +// === 改造后（优化实现）=== | ||
| 36 | +```cpp | ||
| 37 | +// 三种数据类型支持 | ||
| 38 | +this->Input("x") | ||
| 39 | + .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16}); | ||
| 40 | + | ||
| 41 | +// 差异化RoundMode处理 | ||
| 42 | +if constexpr (std::is_same_v<T, float>) { | ||
| 43 | + DataCopy(outputLocal, sumBufLocal, len); | ||
| 44 | +} else if constexpr (std::is_same_v<T, half>) { | ||
| 45 | + Cast(outputLocal, sumBufLocal, RoundMode::CAST_NONE, len); | ||
| 46 | +} else { | ||
| 47 | + Cast(outputLocal, sumBufLocal, RoundMode::CAST_RINT, len); | ||
| 48 | +} | ||
| 49 | +``` | ||
| 50 | + | ||
| 51 | +## 关键修改点 | ||
| 52 | + | ||
| 53 | +1. 预期收益: 支持FP16/BF16/FP32三种精度，通过低精度计算获得2x性能提升，同时保持数值稳定性 | ||
| 54 | + | ||
| 55 | +## 常见陷阱 | ||
| 56 | + | ||
| 57 | +⚠️ 代码复杂度增加，需要模板编程和编译期分支处理 | ||
| 58 | +⚠️ Tiling逻辑复杂度增加 | ||
| 59 | +⚠️ 增加代码量和维护成本 | ||
| 60 | + | ||
| 61 | +## 代码搜索关键词 | ||
| 62 | + | ||
| 63 | +```bash | ||
| 64 | +grep -n "BUFFER_NUM\|InitBuffer\|TQue\|tileSize\|ubFactor\|Tiling\|BLOCK_DIM\|GetBlockNum" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 65 | +``` | ||
| @@ -0,0 +1,61 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'D4: FP8/INT4 Quantization Conversion (FP8/INT4量化输出类型)' | ||
| 4 | +description: 参考实现 使用固定的 CAST_ROUND 模式进行类型转换，这可能在某些场景下引入系统性偏差。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.compute.quantization | ||
| 8 | +- topic.memory | ||
| 9 | +status: stable | ||
| 10 | +sources: | ||
| 11 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/D4_fp8_int4_conversion.md | ||
| 12 | + role: primary | ||
| 13 | +created_at: '2026-10-09T00:00:00Z' | ||
| 14 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 15 | +platforms: | ||
| 16 | +- a2 | ||
| 17 | +--- | ||
| 18 | + | ||
| 19 | +# D4: FP8/INT4 Quantization Conversion (FP8/INT4量化输出类型) | ||
| 20 | + | ||
| 21 | +## 核心思想 | ||
| 22 | +参考实现 使用固定的 CAST_ROUND 模式进行类型转换，这可能在某些场景下引入系统性偏差。参考实现支持多种 rounding mode：MODE_RINT (默认，四舍五入到最近偶数/Banker's rounding)、MODE_ROUND (标准四舍五入)、MODE_FLOOR (向下取整)、MODE_TRUNC (向零取整)、MODE_CEIL (向上取整)。不同场景需要不同的 rounding mode：训练量化通常使用 rint 以减少累积误差、推理量化可能需要 floor 或 trunc 以满足特定硬件要求、FP8 输出强制使用 MODE_RINT 以符合标准。 | ||
| 23 | + | ||
| 24 | +## 代码骨架 | ||
| 25 | + | ||
| 26 | +// === 改造前（基线）=== | ||
| 27 | +```cpp | ||
| 28 | +this->Input("x1").DataType({ge::DT_FLOAT16}); | ||
| 29 | +this->Output("output").DataType({ge::DT_INT8}); | ||
| 30 | +``` | ||
| 31 | + | ||
| 32 | +// === 改造后（优化实现）=== | ||
| 33 | +```cpp | ||
| 34 | +static const std::vector<ge::DataType> xDataType91095 = { | ||
| 35 | + ge::DT_FLOAT16, ge::DT_BF16, ... | ||
| 36 | +}; | ||
| 37 | +static const std::vector<ge::DataType> yDataType91095 = { | ||
| 38 | + ge::DT_INT8, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E5M2, ge::DT_HIFLOAT8 | ||
| 39 | +}; | ||
| 40 | +if constexpr (is_same<T, half>::value) { | ||
| 41 | + Cast(xOut, xLocalFp32, RoundMode::CAST_NONE, elementCount); | ||
| 42 | +} else { // BF16 | ||
| 43 | + Cast(xOut, xLocalFp32, RoundMode::CAST_RINT, elementCount); | ||
| 44 | +} | ||
| 45 | +``` | ||
| 46 | + | ||
| 47 | +## 关键修改点 | ||
| 48 | + | ||
| 49 | +1. 预期收益: 支持多种数据类型组合，适配不同硬件和场景需求，避免算子碎片化 | ||
| 50 | + | ||
| 51 | +## 常见陷阱 | ||
| 52 | + | ||
| 53 | +⚠️ 增加编译时间和代码复杂度，需要更多测试覆盖 | ||
| 54 | +⚠️ 代码复杂度增加，需要维护多个模板类实现 | ||
| 55 | +⚠️ 需要额外的条件判断和地址计算 | ||
| 56 | + | ||
| 57 | +## 代码搜索关键词 | ||
| 58 | + | ||
| 59 | +```bash | ||
| 60 | +grep -n "tileSize\|ubFactor\|Tiling\|BLOCK_DIM\|GetBlockNum\|coreNum\|blockIdx\|SplitCore" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 61 | +``` | ||
Aknowledge/contrib/ascendc/optimizations/evolution_strategies/cards/d5_bf16_special_handling.md+56-0
| @@ -0,0 +1,56 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'D5: BF16 & Platform-Specific Handling (BF16/多平台处理)' | ||
| 4 | +description: 参考实现针对不同的Ascend芯片（910B, 910_93, 910_95, 310P, kirinx90等）提供了特定的配置支持。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.compute.composite | ||
| 8 | +- topic.implementation.pipeline | ||
| 9 | +status: stable | ||
| 10 | +sources: | ||
| 11 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/D5_bf16_special_handling.md | ||
| 12 | + role: primary | ||
| 13 | +created_at: '2026-10-09T00:00:00Z' | ||
| 14 | +updated_at: '2026-10-13T00:00:00Z' | ||
| 15 | +platforms: | ||
| 16 | +- a2 | ||
| 17 | +--- | ||
| 18 | + | ||
| 19 | +# D5: BF16 & Platform-Specific Handling (BF16/多平台处理) | ||
| 20 | + | ||
| 21 | +## 核心思想 | ||
| 22 | +参考实现针对不同的Ascend芯片（910B, 910_93, 910_95, 310P, kirinx90等）提供了特定的配置支持。这种多平台适配策略不仅体现在编译配置上，还深入到了Shape校验和平台特性检测层面。例如，在Ascend910_95上，算子支持2D shape (g, n)用于分组矩阵乘法场景，而在其他平台上只支持(1, n)或(n,)。这种差异化支持允许算子在不同的硬件能力和使用场景下都能获得最优性能。平台特性检测通过GetChipFeature函数实现，动态检测芯片是否支持特定功能（如Intrinsic_data_move_l12bt的bf16支持），从而实现运行时适配。 | ||
| 23 | + | ||
| 24 | +> 注：本卡实测平台为 a2（910B 系）；910_95/310P/kirinx90 等其他平台支持描述源自参考实现，为可迁移推断，未实测。 | ||
| 25 | + | ||
| 26 | +## 代码骨架 | ||
| 27 | + | ||
| 28 | +// === 改造前（基线）=== | ||
| 29 | +```cpp | ||
| 30 | +this->AICore().AddConfig("ascend910b"); | ||
| 31 | +``` | ||
| 32 | + | ||
| 33 | +// === 改造后（优化实现）=== | ||
| 34 | +```cpp | ||
| 35 | +this->AICore().AddConfig("ascend910_95"); | ||
| 36 | +this->AICore().AddConfig("ascend910_93"); | ||
| 37 | +this->AICore().AddConfig("ascend910b"); | ||
| 38 | +OpAICoreConfig config_kirin = GetKirinCoreConfig(); | ||
| 39 | +this->AICore().AddConfig("kirinx90", config_kirin); | ||
| 40 | +``` | ||
| 41 | + | ||
| 42 | +## 关键修改点 | ||
| 43 | + | ||
| 44 | +1. 预期收益: 算子可移植性强，适配多种硬件; 算子可移植性强，支持多种华为Ascend芯片; 提升算子的可移植性和适用范围 | ||
| 45 | + | ||
| 46 | +## 常见陷阱 | ||
| 47 | + | ||
| 48 | +⚠️ 需要维护多个平台配置 | ||
| 49 | +⚠️ 需要额外的FP32 workspace内存和cast操作的开销 | ||
| 50 | +⚠️ 需要维护多套配置，增加测试负担 | ||
| 51 | + | ||
| 52 | +## 代码搜索关键词 | ||
| 53 | + | ||
| 54 | +```bash | ||
| 55 | +grep -n "BUFFER_NUM\|InitBuffer\|TQue\|tileSize\|ubFactor\|Tiling\|BLOCK_DIM\|GetBlockNum" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 56 | +``` | ||
| @@ -0,0 +1,194 @@ | |||
| 1 | +# 进化策略卡（P/D/A/R 系列） | ||
| 2 | + | ||
| 3 | +AscendC 算子性能进化优化的精选策略库，161 张卡按原文件名平铺（原 evolution-strategies `references/cards/`，文件名小写化），索引按算子族分组。 | ||
| 4 | + | ||
| 5 | +## Matmul 系（42 张） | ||
| 6 | + | ||
| 7 | +- [P102: 标量/向量模拟矩阵运算 → MIX CV 融合重写](p102_scalar_matmul_to_cv_fusion_rewrite.md)：当 kernel 用**标量循环或纯 Vector 指令模拟矩阵运算**（dot product、matmul、三角求… | ||
| 8 | +- [P103: Cube Neumann 乘积式三角求逆（替代标量 forward substitution）](p103_cube_neumann_triangular_solve.md)：下三角线性系统求解（forward substitution / back substitution）是 O(C²)… | ||
| 9 | +- [P104: Matmul 多核切分扩展（Multi-Core Tiling Expansion）](p104_matmul_multi_core_tiling_expansion.md)：`MatmulApiTiling::GetTiling(TCubeTiling&)` 自动给出的 `singleCor… | ||
| 10 | +- [P105: Matmul L2 Cache Tile 尺寸规划（L2 Cache Tile Sizing）](p105_matmul_l2_cache_tile_sizing.md)：大矩阵乘中，若一次性把所有 base 块顺序下发，A/B 子块在 L2 中会被后续 tile 反复挤出，造成 MTE2… | ||
| 11 | +- [P106: Matmul Batch 维度合并到 M（Batch Merge to M）](p106_matmul_batch_merge_to_m.md)：在 BatchMatMul 中，当 B 矩阵无 batch（`batchB == 1`）且 A 未转置（`!trans… | ||
| 12 | +- [P107: Matmul Serpentine 反向遍历（Serpentine Traversal）](p107_matmul_serpentine_traversal.md)：在 L2 级 Tiling 中，若按行优先正向遍历 N 方向 tile，每切换一行 M 就要重新从 L2 加载远端 B… | ||
| 13 | +- [P108: Matmul 平台路由（Platform Routing）](p108_matmul_platform_routing.md)：Ascend C Matmul Tiling 设计强烈依赖硬件特征（`NpuArch`、`GetCoreNumAic`… | ||
| 14 | +- [P128: FP32-via-BF16 Hi/Lo 拆分 3-GEMM（Cube 不吃 fp32 时的保精度路径）](p128_fp32_via_bf16_hilo_3gemm.md)：910B Cube 的 Mmad 不直接消费 fp32 输入。 | ||
| 15 | +- [P129: Fused img2col LoadData3D 随路展开（含 band 窗口化突破 L1 容量）](p129_fused_img2col_loaddata3d.md)：卷积的 im2col GM 物化路径流量为 `特征图→UB→GM→L1→L0`，GM 往返可占整算子 70%。 | ||
| 16 | +- [P130: 1×1 Conv 翻转 GEMM（原始布局即操作数，免 wrepack/im2col）](p130_1x1conv_flipped_gemm.md)：1×1 卷积的 im2col 是恒等变换：NCHW 原始布局的 x 本身就是 `[K=Ci, M=H*W]` 的 GE… | ||
| 17 | +- [P131: Host 侧权重预打包结果 LRU 缓存（Weight Repack Cache）](p131_weight_repack_lru_cache.md)：权重布局变换（K-major repack、fractal 重排、量化打包等）的输出是（weight 内容, shap… | ||
| 18 | +- [P16: Cost-Driven Core Partition (基于代价模型的多粒度分核优化)](p16_cost_driven_core_partition.md)：与简单的按元素/行数均分不同，本策略通过构建计算代价模型（考虑对齐开销、轴权重等硬件特性），对每个计算块评估实际代价… | ||
| 19 | +- [P17: CV Parallel Block Sizing (CV并行基本块大小设计)](p17_cv_parallel_block_sizing.md)：在 Flash Attention 等需要 Cube 和 Vector 并行执行的融合算子中，合理设计 Cube 基本… | ||
| 20 | +- [P21: 二维矩阵缓冲策略（Matrix2x2BufferPolicy）](p21_matrix2x2_buffer_policy.md)：针对 Matmul 切 M 和切 K 两个维度同时需要流水的场景，用 2x2=4 个 buffer 组成二维矩阵，按列… | ||
| 21 | +- [P24: ND↔NZ 格式转换搬运](p24_nd2nz_format_copy.md)：在 GM→L1 搬运阶段同时完成 ND→NZ 格式转换，避免额外 TransData 指令。 | ||
| 22 | +- [P26: Stride 搬运模式（转置与列提取）](p26_stride_fused_transpose.md)：通过 srcStride=0（连读）+ dstStride=跳步（跳写）实现零拷贝转置，将 UB 上连续数据写回 GM… | ||
| 23 | +- [P28: 自定义 HardEvent 生产者-消费者同步](p28_hardevent_producer_consumer.md)：用编译期确定的 HardEvent 类型（MTE2_MTE1, MTE1_M, M_FIX 等）建立精确的生产者-消费… | ||
| 24 | +- [P29: 跨核同步（AIC-AIV CrossCore Sync）](p29_crosscore_aic_aiv_sync.md)：在 AIC（Cube 核）和 AIV（Vector 核）之间传递数据时，通过 CrossCoreSetFlag/Cro… | ||
| 25 | +- [P44: Transpose Buffer 常驻复用](p44_transpose_buffer_resident.md)：gradTransposeBuf_ 作为转置后的梯度数据常驻 UB，在整个 numIters 迭代循环中被 colNo… | ||
| 26 | +- [P46: MatmulImpl 高阶矩阵乘 API (MatmulImpl High-Level API)](p46_matmulimpl_highlevel_api.md)：用 `matmul::MatmulImpl` 高阶 API + MDL 配置替代手写 L0/L1 缓冲管理，将 ~16… | ||
| 27 | +- [P47: 对角线块调度 (Diagonal Block Scheduling)](p47_diagonal_block_scheduling.md)：用 MNBlockIdxCompute 对角线分配替代线性行优先分块，优化多核场景下的 L2 cache 复用和负载均… | ||
| 28 | +- [P48: 多量化模式编译期分发 (Multi-Quantization Mode Dispatch)](p48_multi_quant_dispatch.md)：通过 TilingKey + 编译期宏分发支持 A8W8/A4W4/A16W8/A8W4 MSD 等多种量化路径。 | ||
| 29 | +- [P49: 硬件加速反量化 (Hardware-Accelerated Dequantization)](p49_hw_dequantization.md)：用 AscendDequant 硬件融合指令（int32→float32 * scale 一步完成）+ Brcb 广播… | ||
| 30 | +- [P50: SwiGLU 融合流水线 (SwiGLU Fusion Pipeline)](p50_swiglu_fusion.md)：用高阶 SwiGLU API + Ping-Pong 双缓冲实现 AIC-AIV 流水线，将 5 步手动 sigmoi… | ||
| 31 | +- [P51: 动态 AIC/AIV 核配比 (Dynamic AIC/AIV Core Ratio)](p51_dynamic_aic_aiv_ratio.md)：根据 K 维大小动态选择 AIC:AIV 比例（1:1 vs 1:2 vs CUBE_ONLY），在 Cube 密集和… | ||
| 32 | +- [P52: L2 Cache Hint 优化 (L2 Cache Hint Optimization)](p52_l2_cache_hint.md)：通过 SetL2CacheHint 在多核竞争场景下主动禁用 L2 缓存，避免缓存抖动（cache thrashing… | ||
| 33 | +- [P53: L1 常驻复用策略 (L1 Resident Reuse)](p53_l1_resident_reuse.md)：在 FlashAttention/IFA 等 Cube 密集型算子中，将 Q 矩阵（左矩阵）在首次 N 方向迭代时加载… | ||
| 34 | +- [P63: Iterate 异步接口避免 AIC/AIV 同步依赖 (Iterate Async Interface)](p63_iterate_async_interface.md)：在 MIX（AIC+AIV 混合编程）场景中，使用 Matmul Iterate<false> 或 IterateAl… | ||
| 35 | +- [P70: FP Buffer 随路量化 (FixPipe Buffer Inline Quantization)](p70_fp_buffer_inline_quant.md)：对矩阵乘结果进行量化时，反例做法是将 CO1 结果搬到 workspace 再搬到 UB 进行 Vector 量化计算。 | ||
| 36 | +- [P71: Matmul IBShare L1 共享矩阵 (Matmul IBShare L1 Sharing)](p71_matmul_ibshare_l1.md)：MIX 场景下多个 AIV 核使用相同的 A 或 B 矩阵数据时，默认每个 AIV 都独立从 GM 搬运到 L1。 | ||
| 37 | +- [P72: Matmul 多核切 K (Matmul Multi-Core Split-K)](p72_matmul_split_k.md)：当矩阵 M 和 N 较小无法有效切分多核时，可沿 K 轴切分实现多核并行。 | ||
| 38 | +- [P73: CV 并行 AIC:AIV 比例与多 Workspace 流水 (CV Parallel AIC:AIV Ratio & Multi-Workspace](p73_cv_parallel_aic_aiv_ratio.md)：GroupedMatmul 等 Cube+Vector 融合算子中，当 Vector 计算为主要瓶颈时，可将 AIC:… | ||
| 39 | +- [P75: 双 AIV M/S1 轴工作量分裂 (Dual AIV M/S1-Axis Workload Split)](p75_dual_aiv_workload_split.md)：利用 Ascend 910B 的 1:2 AIC:AIV 核比例，将 Vector 阶段的 M 轴或 S1 轴工作量在… | ||
| 40 | +- [P78: 片上缓存加速 Matmul (On-Chip Cache Matmul Acceleration)](p78_onchip_cache_matmul_acceleration.md)：利用 TSCM（Tensor Shared Cache Memory）或 L1 Carry 路径将 Matmul 的输… | ||
| 41 | +- [P79: Grad Cube2+Cube3 融合计算 (Grad Cube2+Cube3 Fusion for d=64)](p79_grad_cube23_fusion.md)：在 Flash Attention 反向传播中，Cube2（dS*K→dQ）和 Cube3（dS^T*Q→dK, S^… | ||
| 42 | +- [P83: L1 Resident Reuse with Multi-Buffer Partitioning (L1 常驻复用与多 Buffer 分区)](p83_l1_resident_reuse_partition.md)：在 FlashAttention、IFA、MLA 等 Cube 密集型算子中，可将 Q 或左矩阵在首次迭代时搬入 L1… | ||
| 43 | +- [P86: Matmul Internal Data-Path Efficiency Tuning (Matmul 高阶 API 内部数据通路效率调优)](p86_matmul_api_transport_efficiency_tun.md)：该策略统一覆盖四类同源优化：选择更合适的调度模板（MDL/NBuffer33/Preload）、放大 base blo… | ||
| 44 | +- [P87: Manual Mmad Pipeline Timing Control (手动 Mmad 流水线时序控制)](p87_manual_mmad_pipeline_timing_control.md)：当 AscendC 的 Matmul 高阶 API 无法满足特定时序需求时，可以完全绕过 Matmul 库，手动控制… | ||
| 45 | +- [P88: Multi-Phase Compute Phase Decomposition (多阶段计算拆分与资源重分配)](p88_multiphase_compute_phase_decomposit.md)：当算子内部存在前置预处理、主计算、后处理或某个会被主循环重复引用的子计算时，可将其拆成多个彼此串行的独立阶段，并在阶段… | ||
| 46 | +- [P90: 条件性 PipeBarrier 基于矩阵规模 (Conditional PipeBarrier by Matrix Size)](p90_conditional_pipebarrier.md)：在手动 Matmul 流水线中，当 L0C 处于累加模式（`isL0CAccum=true`）时，根据矩阵乘的 M/N… | ||
| 47 | +- [P92: Matmul 基本块计算访存比优化 (Matmul Base Block Compute-to-Memory Ratio)](p92_matmul_block_compute_ratio.md)：Matmul 基本块参数 `[baseM, baseN, baseK]` 直接影响 Cube 计算访存比和 MTE2… | ||
| 48 | +- [P93: Matmul K 轴错峰访问内存 (Matmul K-Axis Staggered Memory Access)](p93_matmul_k_stagger_access.md)：多核执行 Matmul 时，若输入矩阵 A 或 B 的 GM 地址相同，多核同时访问相同 GM 地址会导致地址冲突，M… | ||
| 49 | + | ||
| 50 | +## Attention 系（38 张） | ||
| 51 | + | ||
| 52 | +- [P115: FlashAttention Tiling 路由](p115_flashattention_tiling_routing.md)：FlashAttention 类算子采用四层文档架构（L1 路由 → L2 地基 → L3 设计流 → L4 变体）。 | ||
| 53 | +- [P116: GQA Grouped-M 头合并（G 个 Query Head 合并进 Matmul M 维）](p116_gqa_grouped_m_head_merge.md)：GQA 中 `N_q = N_kv * G`，同一 KV head 被连续 G 个 query head 共享。 | ||
| 54 | +- [P117: 双 BMM 链 MatmulApiTiling Host 配方（SetBufferSpace + SetFixSplit + depthA1](p117_matmul_api_tiling_host_recipe.md)：Attention 类算子的 QK^T / PV 是两个 chained matmul。 | ||
| 55 | +- [P118: Causal Mask 独立 AIV Kernel 并行预生成（按需 Dense 替代通用压缩表）](p118_mask_precompute_parallel_aiv.md)：Causal mask 不该占用 attention 主路径的计算/搬运时间。 | ||
| 56 | +- [P119: Two-Pass Softmax → Online Softmax 单遍 KV 流式改写](p119_two_pass_to_online_softmax.md)：两遍 softmax——pass1 扫全部 KV 求行 max，pass2 再扫一遍算 `exp(s - max)`… | ||
| 57 | +- [P120: FA 向量点积模拟 Matmul → Cube mmad 重写（QK^T / PV 上 Cube）](p120_fa_vec_dot_to_cube_mmad.md)：kernel 内层循环用 `Mul + ReduceSum` 逐 `(qi, kj)` 对算 QK^T 点积、用标量… | ||
| 58 | +- [P121: FA SplitFuse CV 流水拓扑（AIC QK+PV / AIV softmax+rescaleO，GM workspace 握手槽](p121_fa_splitfuse_cv_pipeline.md)：FlashAttention 单核内的四个阶段天然分属两种核：`QK^T`(cube)、`softmax update… | ||
| 59 | +- [P122: KV 堆叠大块化 + L1 容量自适应 Q/KV 划分](p122_kv_stack_large_tile_l1_adaptive.md)：KV tile 是 32/128 小块且尺寸硬编码时，每次 QK^T 的 N 维、PV 的 K 维喂不饱 mmad，且… | ||
| 60 | +- [P123: Causal/Sparse KV 范围跳过 + 边界 Tile Mask 特化 + dtype×mask 编译期分发](p123_causal_kv_skip_boundary_mask.md)：causal 的朴素实现是"全量 KV 循环 + 逐元素兜底"：KV 循环永远跑满 `kvLen/BN` 轮，每轮对… | ||
| 61 | +- [P124: MLA nope+rope 拼接进单 GEMM K 维（embed-split K 维流水，一次 L0C 累加）](p124_mla_nope_rope_concat_gemm_k.md)：MLA 的 QK^T 中 Q = [Q_nope | Q_rope]、K = [K_nope | K_rope] 是两… | ||
| 62 | +- [P125: 跨任务 Cube 预取（任务边界提前发射下一任务首个 QK^T，消除切换空泡）](p125_cross_task_cube_prefetch.md)：CV 流水（P121）的 PRE_LAUNCH 预发射只覆盖**任务内** KV 维；任务边界处流水照样断——当前任务… | ||
| 63 | +- [P126: GM Tiling 表驱动变长任务调度（任务列表拍平进 GM，一次 launch 跑全 batch decode）](p126_gm_tiling_table_task_schedule.md)：decode 场景每个 batch 的 kvSeqlen 都不同，老派做法是 kernel 内按 `totalQ /… | ||
| 64 | +- [P137: Token 级稀疏 Attention 的 workspace 紧凑化走 Cube（top-k 不再回退 Vector）](p137_sparse_topk_compact_gather_to_cube.md)：Token 级稀疏（top-k）FA 被误认为"不规则索引无法用 Cube"。 | ||
| 65 | +- [P138: host tiling 查询 L0A/L0C 容量并按需自适应降维（跨硬件稳定性）](p138_host_tiling_l0_capacity_guard.md)：Cube 侧 host tiling 常用**当前 SoC 实测值写死**（如 `groupTile` / `base… | ||
| 66 | +- [P139: 多槽流水 workspace 全槽零初始化（防 padded 行 NaN → 偶发崩溃）](p139_multi_slot_workspace_zero_init.md)：P14 式 CV 预发射用 `kPreloadNum=2` 双槽流水时，一次性零初始化（`ZeroWorkspaceR… | ||
| 67 | +- [P140: MLA（kvHeads=1）的 M-tile 放大到 qHeads（UB/L0 自适应）](p140_mla_large_m_tile.md)：FA 里单个 Cube GEMM 的 M 维上限 = **共享同一份 K/V 的查询行数**。 | ||
| 68 | +- [P142: Attention UB 常驻 O 累加器](p142_attention_ub_o_accumulator.md)：Flash Attention 每处理一个 KV tile 都要更新输出 O：… | ||
| 69 | +- [P143: Attention host 端 launch 合并 + per-shape tiling 缓存](p143_attention_host_launch_merge.md)：910B2C 上小 shape / fp32 路径的 attention 容易被 host 侧固定开销主导：每次调用重… | ||
| 70 | +- [P144: host 端计算全局池化与 BN 训练统计，标量传参替代设备归约/往返同步](p144_host_pool_bnstats_scalar_pass.md)：CV-Attn（CP8 门控广播）与含 `BN(training)` 的算子，其数学链常是：… | ||
| 71 | +- [P145: 多 launch 纯 AIV 流水架构（避免同 launch GM 写后读）](p145_multilaunch_no_same_launch_gm_rw.md)：**910B2C 同一次 kernel launch 内，MTE3 写 GM → MTE2 读 GM 的可见性不可靠*… | ||
| 72 | +- [P146: 910B2C 非对齐 DMA 与标量通道加固套件](p146_910b2c_alignment_scalar_hardening.md)：910B2C 上小空间维（H×W ≤ 23×23）算子的搬运/标量行为与常规假设差异很大，违反即… | ||
| 73 | +- [P147: NPU 卷积输入 12-sig-bit 舍入模拟（host 侧复现 F.conv2d 数值）](p147_npu_conv_12sig_rounding_emulation.md)：910B2C 的 `F.conv2d` 参考实现会先把**激活与卷积核**都转成 **12 有效位（12-sig-bi… | ||
| 74 | +- [P14: CV 流水预发射 (CV Pipeline Preload)](p14_cv_pipeline_preload.md)：Cube 核和 Vector 核交替执行不同迭代阶段的任务，用环形任务缓存交错编排，消除跨核空等。 | ||
| 75 | +- [P18: L1 7-Buffer Resident Partitioning (L1 七缓冲常驻分区)](p18_l1_7buf_resident_partition.md)：针对 Cube 核上 GM→L1→L0→Cube 的多级数据搬运流水线，将 L1 总容量（512KB）静态划分为 7… | ||
| 76 | +- [P38: Softmax 状态 Buffer 跨 S2 循环常驻](p38_softmax_state_resident.md)：Flash Attention 的 online softmax 需要在 S2 方向多次循环中累积 softmaxMa… | ||
| 77 | +- [P45: 稀疏 Attention 离散块处理](p45_topk_index_cache_iterator.md)：将 topK 选出的离散 block indices 预取到片上 cache 数组中，封装为 L1ChunkItera… | ||
| 78 | +- [P55: BlockReduceMax 替代 DataCopy slice (Block-level Reduction Optimization)](p55_block_reduce_max_optimization.md)：当需要从 UB 中按固定间隔提取元素（如每 8 个元素取 1 个）并写回 GM 时，使用 BlockReduceMax… | ||
| 79 | +- [P57: FlashDecode G 分核归约 (FlashDecode G-Partition Reduction)](p57_flashdecode_g_partition.md)：在 FlashDecode 场景下，当 B 和 N2 都很小而 G 很大时，传统归约方案只使用很少的核数。 | ||
| 80 | +- [P58: TND 负载均衡分核 (TND Load-Balanced Partition)](p58_tnd_load_balance_partition.md)：在 TND（T=B*S 合轴）场景下，通过线段抽象和贪心算法实现跨核负载均衡分配。 | ||
| 81 | +- [P60: KV NZ 格式优化 (KV NZ Format Optimization)](p60_kv_nz_format_optimization.md)：将 KV 矩阵从 ND 格式改为 NZ（Channel-first）格式，使数据布局与 Cube 计算的 L0 输入格… | ||
| 82 | +- [P61: L2 Cache 优化 (L2 Cache Optimization)](p61_l2_cache_optimization.md)：通过关闭双页表功能（开启 KV 的 L2 Cache），减少 GM 访问延迟。 | ||
| 83 | +- [P62: Draft: Sparse/Dense Dual-Path Dispatch by Reachable Coverage（按可达覆盖度切换](p62_sparse_dense_dual_path_dispatch_dra.md)：部分 sparse attention 训练算子并不会无条件执行 gather-merge sparse path。 | ||
| 84 | +- [P68: 低延迟归约指令组合 (Low-Latency Reduction Instruction Combination)](p68_low_latency_reduction.md)：归约操作中，WholeReduceSum 等归约指令延迟约为 Add 指令的 2-5 倍。 | ||
| 85 | +- [P76: V_TEMPLATE KV 预合并 Workspace (V_TEMPLATE KV Pre-Merge Workspace)](p76_v_template_kv_premerge.md)：在 Sparse Attention 的 V_TEMPLATE 模式下，由 Vector 核在 Cube 计算之前将分… | ||
| 86 | +- [P77: L1 Chunk Iterator 稀疏块自适应聚合 (L1 Chunk Iterator Sparse Block Adaptive Aggregation)](p77_l1_chunk_iterator_sparse_aggregatio.md)：在 Sparse Attention 推理场景中，多个不连续的 sparse block 可能无法单独填满一个 L1… | ||
| 87 | +- [P80: Sink Token 流水线注入与 S2 块跳过 (Sink Task Pipeline Injection and S2 Block Skip)](p80_sink_task_pipeline_injection.md)：在 Streaming LLM 的 Attention Sink 场景中，将 sink token（前 N 个 tok… | ||
| 88 | +- [P81: Resident UB/TBuf Buffer Reuse (UB/TBuf 常驻 Buffer 复用)](p81_resident_ub_tbuf_reuse.md)：将小尺寸参数、跨迭代状态、梯度累加器或迭代式中间结果一次性搬入或初始化到 UB/TBuf 常驻 buffer，并在整个… | ||
| 89 | +- [P91: Softmax 梯度预计算独立阶段 (Softmax Gradient Pre-Computation Phase)](p91_softmax_grad_precompute.md)：在 Flash Attention 反向传播中，将 softmax gradient 的计算（`sum_i = sum… | ||
| 90 | + | ||
| 91 | +## 通用策略（26 张） | ||
| 92 | + | ||
| 93 | +- [A1: FP32 Intermediate Computation (FP32中间计算)](a1_fp32_intermediate.md)：参考实现在DAG中显式地进行了类型转换：输入数据类型U (BF16/FP16) 被转换为T (FP32) 后进行所有计… | ||
| 94 | +- [A2: Welford Numerically Stable Mean/Var (Welford数值稳定算法)](a2_welford_algorithm.md)：在训练模式下，BatchNorm需要计算当前batch的均值和方差用于前向传播，同时更新running_mean和ru… | ||
| 95 | +- [A3: Rounding Mode Control (舍入模式控制)](a3_rounding_mode.md)：参考实现在FP16版本中采用了多层次类型转换策略来平衡精度与性能。 | ||
| 96 | +- [A4: SetFlag/WaitFlag Event Sync (事件同步保证精度)](a4_pipeline_barrier.md)：参考实现在关键计算节点插入了PipeBarrier<PIPE_V>()进行向量管道同步，确保数据转换和计算操作的正确顺… | ||
| 97 | +- [A5: Numerical Safety & Special Value Handling (数值安全与特殊值处理)](a5_softmax_stability.md)：参考实现在计算输出值时，通过Compare+Select组合实现了对NaN（Not a Number）的特殊处理。 | ||
| 98 | +- [A6: High-Precision rsqrt via Newton-Raphson (高精度rsqrt)](a6_high_precision_rsqrt.md)：Layer Norm的归一化需要计算1/sqrt(var + eps)。 | ||
| 99 | +- [A7: Index & Boundary Safety (索引与边界安全处理)](a7_index_boundary.md)：参考实现提供了完整的InferShapeForMultiScaleDeformableAttentionGrad函数… | ||
| 100 | +- [A8: Quantization-Specific Precision (量化专用精度处理)](a8_quant_precision.md)：参考实现通过Maxs和Mins指令精确控制量化边界，确保量化值严格限制在[quant_min, quant_max]范… | ||
| 101 | +- [D1: Mixed Precision Architecture (混合精度架构)](d1_mixed_precision.md)：参考实现仅支持 float32 输入类型，这在实际 AI 推理场景中是不足的。 | ||
| 102 | +- [D2: Template Kernel Type Dispatch (模板化内核类型分发)](d2_template_kernel.md)：参考实现通过C++模板机制实现了对多种数据类型的统一支持。 | ||
| 103 | +- [D3: TilingKey-Driven Type Dispatch (TilingKey驱动类型分发)](d3_tilingkey_dispatch.md)：参考实现通过模板类FakeQuantAffineCachemaskFp32和FakeQuantAffineCachem… | ||
| 104 | +- [D4: FP8/INT4 Quantization Conversion (FP8/INT4量化输出类型)](d4_fp8_int4_conversion.md)：参考实现 使用固定的 CAST_ROUND 模式进行类型转换，这可能在某些场景下引入系统性偏差。 | ||
| 105 | +- [D5: BF16 & Platform-Specific Handling (BF16/多平台处理)](d5_bf16_special_handling.md)：参考实现针对不同的Ascend芯片（910B, 910_93, 910_95, 310P, kirinx90等）提供了… | ||
| 106 | +- [P10: Vectorized Data Copy & Instructions (向量化数据与指令)](p10_vectorized_copy.md)：参考实现在mask生成逻辑中采用了Compare+Select的组合指令，而非简单的CompareScalar。 | ||
| 107 | +- [P11: Tail Block Handling (尾块与边界处理)](p11_tail_block_handling.md)：参考实现精心设计了尾块处理机制以解决数据量不能被核心数整除的问题：1) 双轨制计算：同时计算正常核（normCoreH… | ||
| 108 | +- [P12: Broadcast & Mask Operations (广播与掩码操作)](p12_broadcast_mask.md)：参考实现使用SelectWithBytesMask高阶API实现mask应用，替代手动if-else或Add操作。 | ||
| 109 | +- [P13: Special Algorithms & High-Level AscendC APIs (特殊算法与高阶API)](p13_special_algorithms.md)：参考实现采用了面向对象的设计模式，定义了FakeQuantAffineCachemaskBase基类和FakeQuan… | ||
| 110 | +- [P1: 双缓冲机制 (Double Buffering)](p1_double_buffer.md)：用两个 buffer 交替工作：一个用于计算时，另一个用于 DMA 搬运。 | ||
| 111 | +- [P2: Adaptive Tiling (自适应分块策略)](p2_adaptive_tiling.md)：参考实现的Tiling不是固定值，而是根据平台UB大小和数据类型动态计算。 | ||
| 112 | +- [P3: Small-D Multi-Row Merging (小D多行合并优化)](p3_small_d_optimization.md)：当embedding_dim ≤ 512时，参考实现采用完全不同的算法：索引排序+连续处理。 | ||
| 113 | +- [P4: 多核负载均衡 (Multi-core Load Balancing)](p4_load_balance.md)：将数据均匀分配到所有 NPU 核，避免某些核空闲、某些核过载。 | ||
| 114 | +- [P5: 流水线同步优化 (Pipeline Sync)](p5_pipeline_sync.md)：用细粒度的 PipeBarrier/SetFlag/WaitFlag 替代粗粒度的 SyncAll，只在真正有数据依赖… | ||
| 115 | +- [P6: Multi-Algorithm Adaptive Selection (多算法自适应选择)](p6_multi_algo_selection.md)：Welford算法是BatchNorm训练场景的核心算法，相比传统的两遍计算（先算均值再算方差），Welford算法单… | ||
| 116 | +- [P7: 32B Alignment + DataCopyPad (数据对齐与填充)](p7_data_alignment.md)：参考实现充分利用了Ascend C的向量化指令（如Adds、Cast、Duplicate等）进行高效计算。 | ||
| 117 | +- [P8: UB Memory Partitioning (UB内存分区管理)](p8_ub_memory_mgmt.md)：针对BF16（BFloat16）数据类型，参考实现采用了高精度中间计算策略。 | ||
| 118 | +- [P9: Deterministic Output via Workspace (确定性输出)](p9_deterministic_output.md)：在多核并行场景下，如果多个核同时更新同一权重行，由于执行顺序不确定，最终结果虽然是正确的（原子加），但浮点累加顺序不同… | ||
| 119 | + | ||
| 120 | +## 小族与数据通路（26 张） | ||
| 121 | + | ||
| 122 | +- [P127: Conversion 模式特化路由（Transpose/Permute 禁止通用 gather 主路径）](p127_conversion_pattern_specialization_routing.md)：Layout-transform 类算子（permute / transpose / reshape 触发拷贝 / s… | ||
| 123 | +- [P20: 三缓冲轮转（BuffersPolicy3buff）](p20_triple_buffer_rotation.md)：在 AIC/AIV 混合核架构下，三缓冲允许搬运、Cube 计算、Vector 计算三个阶段同时进行，形成三级流水。 | ||
| 124 | +- [P25: DataCopyPad 带 Padding 的搬运](p25_datacopy_pad_alignment.md)：当源数据列数不对齐 32B 时，使用 DataCopyPad + DataCopyPadExtParams 在搬运过程… | ||
| 125 | +- [P32: Scatter Block Update 的 stride 寻址](p32_scatter_block_stride.md)：支持输入张量 dim-0 非连续（通过 inputStride0_ 和 inputStride1_），实现 input… | ||
| 126 | +- [P33: Gather 偏移表驱动的结构化数据提取](p33_gather_offset_table.md)：通过预计算字节偏移表，使用 Gather 指令一次性从混合矩阵中提取多个子张量，偏移表只计算一次常驻 UB。 | ||
| 127 | +- [P37: 梯度累加器跨 batch 常驻](p37_grad_accumulator_cross_batch.md)：gradWeight 需要在所有 batch 和 seq 上累加。 | ||
| 128 | +- [P40: Workspace 双缓冲中间结果常驻](p40_workspace_double_buffer.md)：Cube 和 Vector 之间通过 workspace GM 传递中间结果，每个 core 分配 2 份 works… | ||
| 129 | +- [P42: workspace buffer 多区域分时复用](p42_workspace_zone_reuse.md)：一块大的 TBuf<VECCALC> 被划分为多个逻辑区域，供 RmsNorm 和 RoPE 交替使用。 | ||
| 130 | +- [P54: 头开销优化策略 (Launch Overhead Optimization)](p54_launch_overhead_optimization.md)：头开销定义为算子启动到开始搬运 Tensor 数据（MTE2）的耗时，这段代码主要由 Scalar 流水执行，完成计算… | ||
| 131 | +- [P56: DataCopy UB2GM 合并输出 (Batch DataCopy Output)](p56_datacopy_batch_output.md)：将多次小批量 DataCopy UB2GM 操作合并为一次大批量写回，减少 MTE3 指令发射次数，提升写回效率。 | ||
| 132 | +- [P59: 输出 Transpose 融合 (Output Transpose Fusion)](p59_output_transpose_fusion.md)：在 IFA 等算子输出阶段，通过 DataCopy stride 参数实现输出格式转换（Transpose），避免额外… | ||
| 133 | +- [P64: 避免 GM 同地址访问冲突 (GM Address Conflict Avoidance)](p64_gm_address_conflict_avoidance.md)：MTE2/MTE3/Scalar 访问 GM 时按 512 字节粒度对齐处理，多核同时访问连续 512 字节范围内的地… | ||
| 134 | +- [P65: UB Bank 冲突规避 (UB Bank Conflict Avoidance)](p65_ub_bank_conflict_avoidance.md)：UB（192KB）划分为 48 个 bank（16 个 bank group，每组 3 个 bank），每 bank… | ||
| 135 | +- [P66: GM 地址 512B 对齐优化带宽 (GM 512B Alignment Bandwidth)](p66_gm_512b_alignment.md)：AI 处理器内部设计约束下，GM 地址 512B 对齐可最有效发挥带宽效率。 | ||
| 136 | +- [P74: L2 Cache 选择性禁用 (Selective L2 Cache Disable)](p74_selective_l2_cache_disable.md)：在 KV cache 数据量远大于 L2 容量的推理场景中，通过 `SetL2CacheHint(CacheMode:… | ||
| 137 | +- [P82: BufferPolicy Class for Multi-Stage Pipeline (手写 BufferPolicy 类替代 TQue](p82_manual_multi_buffer_pipeline_policy.md)：当标准 `TQue` 或单一双缓冲无法表达复杂的 buffer 访问路径时，可以手写 `BufferPolicy` 类… | ||
| 138 | +- [P85: On-Chip Buffer Zone Reuse (片上 Buffer 分区与分时复用)](p85_on_chip_buffer_zone_reuse.md)：当算子包含多个串行计算阶段，且各阶段的临时 buffer 生命周期不重叠时，可将它们映射到同一块 UB TBuf 或… | ||
| 139 | +- [P96: 硬件排序引擎替代标量排序/选择 (Hardware Sort Engine Replacement)](p96_hardware_sort_replacement.md)：检测 kernel 中的标量排序（bubble sort、selection sort）或 top-k 选择算法，用… | ||
| 140 | +- [R1: VF 函数融合](r1_vf_fusion.md)：将多个独立的 `__simd_vf__` 函数合并为一个，减少中间结果的 Load/Store 操作，最大化寄存器复用。 | ||
| 141 | +- [R2: 寄存器复用](r2_register_reuse.md)：通过布尔代数简化、指令重排和原地复用减少 RegTensor 使用量，避免超过 32 个限制导致的 spill/fil… | ||
| 142 | +- [R3: 指令双发射](r3_dual_issue.md)：通过循环展开和指令重排，暴露无数据依赖的指令对，最大化 A5 Vector Core 的双发射能力。 | ||
| 143 | +- [R4: Hardware Loop 规范化](r4_hw_loop.md)：确保 VF 内的循环满足 Hardware Loop 编码要求，避免退化为 Software Loop。 | ||
| 144 | +- [R5: 非对齐访问优化](r5_unalign_opt.md)：使用 UnalignRegForLoad/Store 优化非 32B 对齐地址的数据访问。 | ||
| 145 | +- [R6: 低延迟归约](r6_low_latency_reduce.md)：选择最优的归约指令组合，减少 Reduce 类操作的延迟。 | ||
| 146 | +- [R7: SIMD/SIMT 混合编程](r7_simd_simt_hybrid.md)：对同一算子中的不同操作，根据访问模式选择 SIMD 或 SIMT 模型。 | ||
| 147 | +- [R8: Mutex 细粒度同步](r8_mutex_sync.md)：使用 Mutex 替代粗粒度 SetFlag/WaitFlag，减少流水线停顿。 | ||
| 148 | + | ||
| 149 | +## Elementwise（10 张） | ||
| 150 | + | ||
| 151 | +- [P100: GatherMask 替代逐元素 SetValue/GetValue (GatherMask Replacement)](p100_gathermask_replacement.md)：当需要从连续 tensor 中按固定间隔取元素重组数据时（如 RoPE 奇偶分拆、解交织、矩阵转置等），优化前用 `f… | ||
| 152 | +- [P101: GatherMask 替代 Host 侧列拆分 (GatherMask Replaces Host-Side Column Split)](p101_gathermask_host_column_split.md)：当输入为 interleaved 多分量格式（如 `(N,4)` 的 `[x1,y1,x2,y2, x1,y1,x2,… | ||
| 153 | +- [P112: Elementwise 向量化 Tile 尺寸设计](p112_elementwise_vectorized_tile_sizing.md)：Elementwise 算子无跨元素依赖，应优先展平为 1D，并按 Vector 单元宽度与 UB 容量切分 tile。 | ||
| 154 | +- [P22: TQueBind 双向复用队列](p22_tquebind_bidirectional_reuse.md)：对于既需要从 GM 搬入又需要搬出到 GM 的数据（如 conv_states），使用 TQueBind 将 VECI… | ||
| 155 | +- [P30: PipeBarrier 密集插入的纯 Vector 同步](p30_dense_pipebarrier_vector.md)：在递推类算子（y0→y1→y2 依赖链）中，每条 Vector 指令后插入 PipeBarrier<PIPE_V>，确… | ||
| 156 | +- [P43: 反向梯度 4路PingPong 常驻与 Scatter-Add 延迟写回](p43_grad_pingpong_deferred_scatter.md)：反向传播中 dK/dV 梯度累积在 workspace 的 4 路 PingPong buffer 中（selectd… | ||
| 157 | +- [P69: UB 融合连续 Vector 计算 (UB Fused Vector Chain)](p69_ub_fused_vector_chain.md)：多次 Vector 计算且前一次输出是后一次输入时，反例做法是每次计算后将结果从 UB 搬到 GM 再搬回 UB。 | ||
| 158 | +- [P84: Vector Compute Efficiency Patterns (UB 内 Vector 计算效率优化模式)](p84_vector_compute_efficiency_patterns.md)：在纯 Vector 或 Vector 主导的算子中，性能瓶颈常来自三类 UB 内部效率损失：Scalar 侧为每个 V… | ||
| 159 | +- [P95: 标量逐元素运算向量化 (Scalar Element-wise → Adds/Muls Vector Chain)](p95_scalar_elementwise_to_vector.md)：检测 kernel 中逐元素标量算术循环（`for(i) y[i] = f(x[i])`），用 AscendC 标量-… | ||
| 160 | +- [P97: 零拷贝子张量切片 + 批量向量运算 (Zero-Copy Sub-Tensor Slicing)](p97_zero_copy_subtensor.md)：当需要从 UB 大 tensor 中按偏移量取子集做批量运算时（如 NMS 的逐候选框 IoU、TopK 的逐行合并）… | ||
| 161 | + | ||
| 162 | +## Reduction 系（10 张） | ||
| 163 | + | ||
| 164 | +- [P109: Reduction 多轴 Tiling 系统化设计](p109_reduction_multi_axis_tiling.md)：将任意 N 维 shape + axes 先合轴为 A/R 交替序列，再按单轴归约（AR/ARA）逐维处理。 | ||
| 165 | +- [P110: Reduction UB 累加器复用](p110_reduction_ub_accumulator_reuse.md)：在 AR/ARA FullLoad 模式下，整行/整块数据一次驻留 UB，归约后的中间结果直接复用 UB buffer… | ||
| 166 | +- [P19: 自定义多级 PingPong 双缓冲](p19_custom_pingpong_double_buffer.md)：绕过 TQue 标准队列机制，用手写 ping/pong 轮转实现 L0A/L0B/L1 等多级存储的双缓冲。 | ||
| 167 | +- [P34: Weight 常驻与预计算复用](p34_weight_load_once_reuse.md)：将 weight/scale 等小尺寸、跨 loop 不变的数据在 Process 开头搬入 UB 并 cast 到… | ||
| 168 | +- [P35: TBuf 常驻中间累加器](p35_tbuf_resident_accumulator.md)：使用 TBuf<VECCALC> 分配的 buffer 不参与队列管理，在 Init 中分配后在整个 Process… | ||
| 169 | +- [P41: gamma/scale/offset 常驻 UB](p41_gamma_scale_resident.md)：gamma（RmsNorm 权重）、quantization scale/offset 等小张量在整个 kernel… | ||
| 170 | +- [P89: UB 高阶 API 临时 Buffer 共享复用 (High-Level API Temp Buffer Sharing)](p89_ub_api_temp_buffer_sharing.md)：当算子使用需要临时 Buffer 的高阶 API（如 SoftMax）时，该临时空间会挤占其他计算空间。 | ||
| 171 | +- [P94: 标量归约向量化 (Scalar Reduction → Vector Reduce API)](p94_scalar_reduction_to_vector.md)：检测 kernel 中逐元素的标量归约循环（`for(i) sum += x[i]`、`for(i) if (x[i]… | ||
| 172 | +- [P98: 混合标量/向量路径自适应调度 (Hybrid Scalar/Vector Dispatch)](p98_hybrid_scalar_vector_dispatch.md)：向量化并非在所有 workload 大小上都优于标量。 | ||
| 173 | +- [P99: BroadCast + 批量向量运算替代逐行标量循环 (BroadCast + Batch Vector Ops)](p99_broadcast_batch_vector_ops.md)：当小向量（如 `rstd[gRows]`）需要广播到大矩阵（`[gRows, D]`）做逐元素运算时，优化前用 `fo… | ||
| 174 | + | ||
| 175 | +## SSM/Scan 系（5 张） | ||
| 176 | + | ||
| 177 | +- [P132: 线性递归 chunk-scan 三阶段分解（块内预计算 / 块间状态扫描 / 输出投影）](p132_chunk_scan_three_stage_decomposition.md)：线性递归 / 状态空间类算子（Gated Delta Rule、Gated Linear Attention、Mamb… | ||
| 178 | +- [P133: 块内预计算 AIC/AIV CV 流水（cube 算 KKT/求逆、vector 算 cumsum/exp/前代，6-flag 定向握手）](p133_ssm_stage1_cv_pipeline.md)：P132 三阶段分解后，Stage1 块内预计算是**多级 matmul + 向量预处理交错**的序列：prep →… | ||
| 179 | +- [P134: chunk-group 任务划分（group 间串行 + group 内 (nv, chunk) 多核并行）](p134_ssm_chunk_group_partitioning.md)：P132 三阶段后，块间状态递推（Stage2）强制 **chunk 之间串行**——第 k+1 个 chunk 的状… | ||
| 180 | +- [P135: delta-rule 矩阵求逆分治（对角向量前代 + 非对角 cube MM）](p135_delta_rule_inverse_divide_conquer.md)：Delta Rule（GDN/KDA）类算子块内预计算要解一个 `C×C` 下三角线性系统（GDR：`attn_1 =… | ||
| 181 | +- [P136: SSM 中间量 bf16 量化 + 状态 UB 常驻（head grouping）](p136_ssm_intermediate_bf16_quant.md)：三阶段 chunk-scan（P132）会物化大量中间量：`v_beta / q_prime / attn_1 / k… | ||
| 182 | + | ||
| 183 | +## Broadcast（2 张） | ||
| 184 | + | ||
| 185 | +- [P113: Broadcast 对齐 Tiling 设计](p113_broadcast_aligned_tiling.md)：Broadcast 算子先通过 DimensionCollapse 把多轴问题降维，再判断是 OneDim（合轴后单维… | ||
| 186 | +- [P67: Vector Counter 模式 (Vector Counter Mode)](p67_vector_counter_mode.md)：Normal 模式下 Vector API 需手动计算 repeatTimes、区分主块/尾块、多次设置 mask，涉… | ||
| 187 | + | ||
| 188 | +## Sort（1 张） | ||
| 189 | + | ||
| 190 | +- [P111: Sort 两级归并 Tiling 系统化设计](p111_sort_two_level_merge_tiling.md)：当 `N > tileSize × coreNum` 时，Sort / TopK 必须走 Pattern C：先在各核… | ||
| 191 | + | ||
| 192 | +## Conversion（1 张） | ||
| 193 | + | ||
| 194 | +- [P114: Transpose 多维对齐 Tiling 设计](p114_transpose_nd_tiling_with_alignment.md)：small-channel transpose 类算子先把问题统一建模为 `[C, N] → [N, C]`，其中 `… | ||
Aknowledge/contrib/ascendc/optimizations/evolution_strategies/cards/p100_gathermask_replacement.md+95-0
| @@ -0,0 +1,95 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P100: GatherMask 替代逐元素 SetValue/GetValue (GatherMask Replacement)' | ||
| 4 | +description: 当需要从连续 tensor 中按固定间隔取元素重组数据时（如 RoPE 奇偶分拆、解交织、矩阵转置等），优化前用 `for` 循环逐元素 | ||
| 5 | + `src.GetValue(i)` → `dst.SetValue(j, val)` 操作，每条调用仅处理 1 个元素（O(n) 次 Scalar ↔ Vecto… | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- paradigm.elementwise | ||
| 9 | +- topic.compute.composite | ||
| 10 | +- topic.compute.indexing | ||
| 11 | +- topic.attention | ||
| 12 | +- topic.programming_model.execution_unit.scalar | ||
| 13 | +status: stable | ||
| 14 | +sources: | ||
| 15 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P100_gathermask_replacement.md | ||
| 16 | + role: primary | ||
| 17 | +created_at: '2026-10-09T00:00:00Z' | ||
| 18 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 19 | +platforms: | ||
| 20 | +- a2 | ||
| 21 | +--- | ||
| 22 | + | ||
| 23 | +# P100: GatherMask 替代逐元素 SetValue/GetValue (GatherMask Replacement) | ||
| 24 | + | ||
| 25 | +## 核心思想 | ||
| 26 | +当需要从连续 tensor 中按固定间隔取元素重组数据时（如 RoPE 奇偶分拆、解交织、矩阵转置等），优化前用 `for` 循环逐元素 `src.GetValue(i)` → `dst.SetValue(j, val)` 操作，每条调用仅处理 1 个元素（O(n) 次 Scalar ↔ Vector 数据通路切换）。优化后用 AscendC `GatherMask` API，在 Vector 单元上一次完成间隔采集，O(1) 条向量指令替代 O(n) 次标量操作，且零拷贝（直接从源 tensor 的偏移视图上操作）。 | ||
| 27 | + | ||
| 28 | +## 代码骨架 | ||
| 29 | + | ||
| 30 | +```cpp | ||
| 31 | +// === 改造前（基线）：逐元素 SetValue/GetValue 解交织 === | ||
| 32 | +// 从 kvF32[rms_size_:] 中分离偶数位和奇数位 | ||
| 33 | +int half = rope_size_ / 2; | ||
| 34 | +LocalTensor<float> kF32 = ropeBuf_.Get<float>(); | ||
| 35 | +for (int i = 0; i < half; ++i) { | ||
| 36 | + kF32.SetValue(i, kvF32.GetValue(rms_size_ + 2 * i)); // 偶数位 | ||
| 37 | +} | ||
| 38 | +for (int i = 0; i < half; ++i) { | ||
| 39 | + kF32.SetValue(half + i, kvF32.GetValue(rms_size_ + 2 * i + 1)); // 奇数位 | ||
| 40 | +} | ||
| 41 | +// 2 × half 次 SetValue + GetValue Scalar 调用，每次仅处理 1 个元素 | ||
| 42 | + | ||
| 43 | +// === 改造后：GatherMask 解交织（2 条指令完成） === | ||
| 44 | +uint64_t rsvdCnt = 0; | ||
| 45 | +LocalTensor<float> ropeSrc = kvF32[rms_size_]; // 零拷贝偏移视图 | ||
| 46 | + | ||
| 47 | +// deinterleave=1: 取偶数索引位 | ||
| 48 | +GatherMask(kF32, ropeSrc, | ||
| 49 | + static_cast<uint8_t>(1), true, | ||
| 50 | + static_cast<uint32_t>(rope_size_), {1, 1, 8, 0}, rsvdCnt); | ||
| 51 | + | ||
| 52 | +// deinterleave=2: 隔 1 取 1，取奇数索引位 → 写入 kF32[half] | ||
| 53 | +GatherMask(kF32[half], ropeSrc, | ||
| 54 | + static_cast<uint8_t>(2), true, | ||
| 55 | + static_cast<uint32_t>(rope_size_), {1, 1, 8, 0}, rsvdCnt); | ||
| 56 | +PipeBarrier<PIPE_V>(); | ||
| 57 | +``` | ||
| 58 | + | ||
| 59 | +## GatherMask 参数说明 | ||
| 60 | + | ||
| 61 | +```cpp | ||
| 62 | +GatherMask(dst, src, deinterleave, isReverse, count, repeatMaskParams, rsvdCnt); | ||
| 63 | +// dst: 输出 tensor（可用 dst[half] 切片拆分两段） | ||
| 64 | +// src: 源 tensor（可用零拷贝偏移视图 src[offset]） | ||
| 65 | +// deinterleave: uint8_t 间隔步长。1=连续取（偶数位）, 2=隔1取1（奇数位）, N=隔N-1取1 | ||
| 66 | +// isReverse: bool，默认传 true | ||
| 67 | +// count: uint32_t 要收集的元素总数（非 half） | ||
| 68 | +// repeatMaskParams: 常用默认值 {1, 1, 8, 0}（单 repeat，dst/src block stride=1/8） | ||
| 69 | +// rsvdCnt: uint64_t& 保留参数，必须声明并初始化为 0 | ||
| 70 | +``` | ||
| 71 | + | ||
| 72 | +## 关键修改点 | ||
| 73 | + | ||
| 74 | +1. **零拷贝偏移**：`LocalTensor<T> ropeSrc = src[offset]` 创建视图，GatherMask 直接在原 tensor 上操作（配合 P97） | ||
| 75 | +2. **PipeBarrier**：GatherMask 后加 `PipeBarrier<PIPE_V>()`，确保 Vector 写入完成后再被后续指令读取 | ||
| 76 | +3. **dst 拆分**：利用 `dst[half]` 切片将两段 GatherMask 结果写入同一 tensor 的不同区域 | ||
| 77 | +4. **预期收益**：元素数 ≥ 128 时 5-20x 加速；< 64 时收益递减但通常仍优于标量 | ||
| 78 | + | ||
| 79 | +## 常见陷阱 | ||
| 80 | + | ||
| 81 | +⚠️ GatherMask 后必须加 `PipeBarrier<PIPE_V>()`，否则 dst 数据可能被后续指令读为旧值 | ||
| 82 | +⚠️ `deinterleave` 是 uint8_t 类型，传入 int 需显式 cast：`static_cast<uint8_t>(1)` | ||
| 83 | +⚠️ `count` 传元素总数而非 half；两段 GatherMask 的 count 通常相同（各取一半元素） | ||
| 84 | +⚠️ `rsvdCnt` 必须声明为 `uint64_t` 并初始化为 0 | ||
| 85 | +⚠️ GatherMask 从 src 当前位置开始取，需确保 `ropeSrc = src[offset]` 偏移正确 | ||
| 86 | + | ||
| 87 | +## 代码搜索关键词 | ||
| 88 | + | ||
| 89 | +```bash | ||
| 90 | +grep -nE "GatherMask|deinterleave|rsvdCnt|GatherMaskParams" op_kernel/*.cpp | ||
| 91 | +``` | ||
| 92 | + | ||
| 93 | +## 来源 | ||
| 94 | + | ||
| 95 | +- tmp.md 优化点2：GatherMask 替代 SetValue/GetValue 逐元素解交织（RoPE 场景），2 条向量指令替代 2×half 次标量操作 | ||
| @@ -0,0 +1,156 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P101: GatherMask 替代 Host 侧列拆分 (GatherMask Replaces Host-Side Column Split)' | ||
| 4 | +description: 当输入为 interleaved 多分量格式（如 `(N,4)` 的 `[x1,y1,x2,y2, x1,y1,x2,y2, ...]`）时，传统做法是 | ||
| 5 | + Host 侧 `.select(1,i).contiguous()` 拆成独立 1D 张量传入 kernel。 | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- paradigm.elementwise | ||
| 9 | +- topic.compute.composite | ||
| 10 | +- topic.compute.indexing | ||
| 11 | +- topic.attention | ||
| 12 | +- topic.data_movement.memory_access | ||
| 13 | +status: stable | ||
| 14 | +sources: | ||
| 15 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P101_gathermask_host_column_split.md | ||
| 16 | + role: primary | ||
| 17 | +created_at: '2026-10-09T00:00:00Z' | ||
| 18 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 19 | +platforms: | ||
| 20 | +- a2 | ||
| 21 | +--- | ||
| 22 | + | ||
| 23 | +# P101: GatherMask 替代 Host 侧列拆分 (GatherMask Replaces Host-Side Column Split) | ||
| 24 | + | ||
| 25 | +## 核心思想 | ||
| 26 | + | ||
| 27 | +当输入为 interleaved 多分量格式（如 `(N,4)` 的 `[x1,y1,x2,y2, x1,y1,x2,y2, ...]`）时，传统做法是 Host 侧 `.select(1,i).contiguous()` 拆成独立 1D 张量传入 kernel。GatherMask 内置固定模式 (built-in patterns 3/4/5/6) 可以在 kernel 内一条 VEC 指令拆分一个分量，消除 Host 侧拷贝和多余 DMA： | ||
| 28 | + | ||
| 29 | +``` | ||
| 30 | +传统: Host 8×.select().contiguous() → Kernel 8+ 路 DMA → Compute | ||
| 31 | +P101: Host 传原始 interleaved 张量 → Kernel 1 路 DMA → 4×GatherMask → Compute | ||
| 32 | +``` | ||
| 33 | + | ||
| 34 | +GatherMask 的 4 条 VEC 指令开销远小于 Host 侧拷贝 + 多路 DMA 的累计 overhead。 | ||
| 35 | + | ||
| 36 | +## 代码骨架 | ||
| 37 | + | ||
| 38 | +```cpp | ||
| 39 | +// === 改造前：Host 侧列拆分 + 多路 DMA === | ||
| 40 | +// Host: | ||
| 41 | +at::Tensor x1 = bboxes.select(1, 0).contiguous(); // 4 路 .select().contiguous() | ||
| 42 | +at::Tensor y1 = bboxes.select(1, 1).contiguous(); | ||
| 43 | +at::Tensor x2 = bboxes.select(1, 2).contiguous(); | ||
| 44 | +at::Tensor y2 = bboxes.select(1, 3).contiguous(); | ||
| 45 | +// Kernel: 4 路 DMA | ||
| 46 | +CopyIn(x1Tile, x1Gm[tileStart], tileLen); // 4 次 DataCopyPad | ||
| 47 | +CopyIn(y1Tile, y1Gm[tileStart], tileLen); | ||
| 48 | +CopyIn(x2Tile, x2Gm[tileStart], tileLen); | ||
| 49 | +CopyIn(y2Tile, y2Gm[tileStart], tileLen); | ||
| 50 | + | ||
| 51 | +// === 改造后：Kernel 内 GatherMask 解交织 === | ||
| 52 | +// Host: 直接传原始 (N,4) 张量，不做列拆分 | ||
| 53 | +// Kernel: 1 路 DMA + 4×GatherMask | ||
| 54 | +CopyIn(interleaved, bboxesGm[tileStart * 4], tileLen * 4); // 1 次 DataCopyPad | ||
| 55 | +PipeBarrier<PIPE_ALL>(); | ||
| 56 | + | ||
| 57 | +// Normal 模式: 每次 repeat 处理 256B = 64 float32 → 16 dst per repeat | ||
| 58 | +// pattern 3 (00010001): 每 4 个取第 1 个 → x1 | ||
| 59 | +// pattern 4 (00100010): 每 4 个取第 2 个 → y1 | ||
| 60 | +// pattern 5 (01000100): 每 4 个取第 3 个 → x2 | ||
| 61 | +// pattern 6 (10001000): 每 4 个取第 4 个 → y2 | ||
| 62 | +int32_t alignedLen = (tileLen / 16) * 16; | ||
| 63 | +uint16_t repeatTimes = alignedLen / 16; | ||
| 64 | +GatherMaskParams gmParams(1, repeatTimes, 8, 0); // stride=8 (256B) 不可写成 1 | ||
| 65 | + | ||
| 66 | +uint64_t rsvdCnt = 0; | ||
| 67 | +GatherMask(x1, src, static_cast<uint8_t>(3), false, 0, gmParams, rsvdCnt); | ||
| 68 | +GatherMask(y1, src, static_cast<uint8_t>(4), false, 0, gmParams, rsvdCnt); | ||
| 69 | +GatherMask(x2, src, static_cast<uint8_t>(5), false, 0, gmParams, rsvdCnt); | ||
| 70 | +GatherMask(y2, src, static_cast<uint8_t>(6), false, 0, gmParams, rsvdCnt); | ||
| 71 | +PipeBarrier<PIPE_ALL>(); | ||
| 72 | + | ||
| 73 | +// 尾块: tileLen 非 16 倍数时标量补齐 alignedLen..tileLen-1 | ||
| 74 | +for (int32_t j = alignedLen; j < tileLen; ++j) { | ||
| 75 | + int32_t off = j * 4; | ||
| 76 | + x1.SetValue(j, src.GetValue(off)); | ||
| 77 | + y1.SetValue(j, src.GetValue(off + 1)); | ||
| 78 | + x2.SetValue(j, src.GetValue(off + 2)); | ||
| 79 | + y2.SetValue(j, src.GetValue(off + 3)); | ||
| 80 | +} | ||
| 81 | +``` | ||
| 82 | + | ||
| 83 | +## 关键修改点 | ||
| 84 | + | ||
| 85 | +1. **解交织维度放外层循环**：解交织后 UB 常驻，内层维度全量复用；放内层则每 tile 都要重 DMA + 重 GatherMask，DMA 复杂度 O(外层×内层/tile²) | ||
| 86 | +2. **Tile 填到 UB 上限**：tile 过小会导致 DMA 次数随 N 线性增长，压过 GatherMask 收益。tile 填到 UB 上限（910B2C 约 2048）。 | ||
| 87 | + | ||
| 88 | +3. **Normal 模式 + `src0RepeatStride=8`**：Normal 模式每次 repeat 处理 256B（64 float32），stride=8 DataBlocks 保证 repeat 间不重叠。stride=1 会导致连续 repeat 重叠 224B → 全 case 精度 fail | ||
| 89 | +4. **16 对齐 + 标量尾块**：Normal 模式每次 repeat 输出 16 个 dst 元素，tileLen 向下对齐到 16 的倍数交 GatherMask，剩余 0-15 个元素走标量补齐 | ||
| 90 | +5. **只解交织需要复用的大维度**：小维度用标量 GM 访问即可，两侧都做 GatherMask 增加 UB 压力和指令数但无额外复用收益 | ||
| 91 | + | ||
| 92 | +## 适用性检测 (grep) | ||
| 93 | + | ||
| 94 | +```bash | ||
| 95 | +# 检测 Host 侧列拆分（本策略的替代目标） | ||
| 96 | +grep -nE "\.select\(.*\)\.contiguous\(\)" op_host/*.cpp | ||
| 97 | + | ||
| 98 | +# 检测 interleaved 多分量输入格式 (N,4) (N,3) (N,2) | ||
| 99 | +grep -nE "size\(1\)\s*==\s*[234]" op_host/*.cpp | ||
| 100 | + | ||
| 101 | +# 检测 GatherMask 是否已使用 built-in patterns 3/4/5/6 | ||
| 102 | +grep -nE "GatherMask.*static_cast<uint8_t>\([3-6]\)" op_kernel/*.cpp | ||
| 103 | +``` | ||
| 104 | + | ||
| 105 | +## 必然推论: 被消除的 GM Tensor 所承载的计算必须迁移到 Kernel 侧 | ||
| 106 | + | ||
| 107 | +当 Host 侧列拆分被消除后，GM 输入数量从 N 路减少到 1 路（interleaved）。**任何原先对这 N 路 GM tensor 做的 Host 侧预处理或预计算，其结果不再有独立的 GM 通道传入 kernel**，必须满足以下二选一： | ||
| 108 | + | ||
| 109 | +| 情形 | 判定条件 | 处理方式 | | ||
| 110 | +|------|---------|---------| | ||
| 111 | +| 可以直接丢弃 | 预计算结果仅用于减少 kernel 计算量，kernel 有原始分量后可在 UB 上重算 | **在 kernel 内用向量指令重算**（通常 2-3 条 VEC 指令，成本远低于单独 DMA） | | ||
| 112 | +| 不能丢弃 | 预计算结果依赖运行时才能确定的数据（如索引表、动态阈值） | 预计算结果保留为独立 GM 输入，但需评估：剩余 DMA 开销是否超过了 P101 节省的 DMA 开销？若超过，P101 不可行 | | ||
| 113 | + | ||
| 114 | +**实施检查清单**（采用 P101 时必须逐项通过，否则为无效实现）： | ||
| 115 | + | ||
| 116 | +```bash | ||
| 117 | +# 1. 列出基线 op_host 中所有 GM 输入 tensor（包括预计算产物） | ||
| 118 | +grep -nE "EXEC_KERNEL_CMD|SetGlobalBuffer" op_host/*.cpp op_kernel/*.cpp | ||
| 119 | + | ||
| 120 | +# 2. 对每个被消除的 GM 输入，回答： | ||
| 121 | +# - 这个 tensor 是原始数据还是预计算结果？ | ||
| 122 | +# - 如果是预计算结果：它的输入数据量是什么？能迁移到 kernel UB 上算吗？ | ||
| 123 | +# - 如果不能迁移：它单独占一条 DMA——P101 省下的 DMA 够抵消吗？ | ||
| 124 | + | ||
| 125 | +# 3. 改造后验证：kernel 的 SetGlobalBuffer 数量应严格 ≤ 改造前 - 被消除的原始分量数 | ||
| 126 | +``` | ||
| 127 | + | ||
| 128 | +**典型反模式**： | ||
| 129 | +- 消除了输入张量的 C 路列拆分 DMA | ||
| 130 | +- **但没有消除**对拆分后各列做的预计算结果 DMA（如各列差值乘积，完全可以在 UB 上用 2-3 条向量指令重算） | ||
| 131 | +- 结果：节省的 DMA 收益被残留的预计算结果 DMA 抵消，P101 退化为无效 | ||
| 132 | + | ||
| 133 | +**正确做法**： | ||
| 134 | +- 列拆分 → GatherMask：消除 C 路 DMA | ||
| 135 | +- 与拆分列相关的预计算 → kernel UB 内若干条向量指令重算：消除对应 DMA | ||
| 136 | +- 与拆分列无关的预计算（原本走标量路径）→ 保持不变 | ||
| 137 | +- 最终 GM 输入数量应严格等于：未拆分的原始 tensor 数 + 与列拆分无关的辅助 tensor 数 + output | ||
| 138 | + | ||
| 139 | +## 常见陷阱 | ||
| 140 | + | ||
| 141 | +⚠️ **解交织维度必须外层循环**：内层循环会导致每 tile 重 DMA + 重 GatherMask，DMA 复杂度从 O(N) 变为 O(M×N) | ||
| 142 | +⚠️ **tile 过小导致收益被 DMA 次数抵消**：tile 越小，DMA 次数随 N 线性增长，大 N 时退化严重 | ||
| 143 | +⚠️ **残留的 GM 预计算结果会抵消 P101 收益**：见上方"必然推论"——每一条未迁移到 kernel 侧计算的预计算结果都有独立 DMA 开销 | ||
| 144 | + | ||
| 145 | +> GatherMask API 参数约束（stride=8、16 对齐、PipeBarrier 等）见 `api-gathermask.md`。 | ||
| 146 | + | ||
| 147 | +## 代码搜索关键词 | ||
| 148 | + | ||
| 149 | +```bash | ||
| 150 | +grep -nE "GatherMask|GatherMaskParams|src0RepeatStride|interleaved|deinterleave" op_kernel/*.cpp | ||
| 151 | +``` | ||
| 152 | + | ||
| 153 | +## 来源 | ||
| 154 | + | ||
| 155 | +- 31_IOU GatherMask vs Host 预拆分对照实验 | ||
| 156 | +- P100 (GatherMask 替代标量解交织) 覆盖 RoPE 奇偶拆分等步长取元素场景，本卡片覆盖列分量分离场景，两者互补 | ||
| @@ -0,0 +1,85 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P102: 标量/向量模拟矩阵运算 → MIX CV 融合重写' | ||
| 4 | +description: 当 kernel 用**标量循环或纯 Vector 指令模拟矩阵运算**（dot product、matmul、三角求解、链式矩阵乘加）时，P94/P95 | ||
| 5 | + 的向量化收益会封顶——Vector 单元本就不是为矩阵运算设计的。 | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- topic.programming_model.execution_unit.cube | ||
| 9 | +- topic.attention | ||
| 10 | +- topic.compute.composite | ||
| 11 | +- topic.programming_model.execution_unit.scalar | ||
| 12 | +status: stable | ||
| 13 | +sources: | ||
| 14 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P102_scalar_matmul_to_cv_fusion_rewrite.md | ||
| 15 | + role: primary | ||
| 16 | +created_at: '2026-10-09T00:00:00Z' | ||
| 17 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 18 | +platforms: | ||
| 19 | +- a2 | ||
| 20 | +--- | ||
| 21 | + | ||
| 22 | +# P102: 标量/向量模拟矩阵运算 → MIX CV 融合重写 | ||
| 23 | + | ||
| 24 | +## 核心思想 | ||
| 25 | + | ||
| 26 | +当 kernel 用**标量循环或纯 Vector 指令模拟矩阵运算**（dot product、matmul、三角求解、链式矩阵乘加）时，P94/P95 的向量化收益会封顶——Vector 单元本就不是为矩阵运算设计的。此时唯一数量级收益的路径是**架构级重写为 MIX CV 融合**：Cube（AIC）承担全部 matmul（Mmad），Vector（AIV）承担 elementwise/decay/mask/状态更新，中间结果经 GM workspace 交接，`SyncAll<false>()` 做全局 barrier。 | ||
| 27 | + | ||
| 28 | +实证（gated_delta_rule，910B2C）：scalar/VEC 实现优化到 3.3x 后仍比参考慢 17x（scalar 做矩阵运算是原罪）→ MIX_AIC_1_1 重写后 **geomean 8.85x vs PyTorch 参考**，单 stage 对比：AIV 标量三角求解 98µs/chunk → Cube Neumann 链 11 个 Mmad（见 P103）。 | ||
| 29 | + | ||
| 30 | +## 适用判定（何时选这张卡而不是 P94/P95） | ||
| 31 | + | ||
| 32 | +``` | ||
| 33 | +kernel 内层循环是以下形态之一？ | ||
| 34 | + ├─ for(i) for(j) acc += A[i][k] * B[k][j] (标量 matmul/dot) | ||
| 35 | + ├─ for(i) for(j) for(k) ... (嵌套标量乘加，维度 ≥ 64) | ||
| 36 | + ├─ for(i) for(j<i) solve/linsolve 递推 (三角求解、前向替换) | ||
| 37 | + └─ Vector 指令拼出的"伪 matmul"（Mul+ReduceSum 逐行） | ||
| 38 | + │ | ||
| 39 | + ├─ 是 → 矩阵维度 ≥ 64×64？ | ||
| 40 | + │ ├─ 是 → **选 P102（CV 融合重写）**，P94/P95 只做收尾的 elementwise 部分 | ||
| 41 | + │ └─ 否（< 64）→ P94/P95 向量化即可，Cube 启动开销不划算 | ||
| 42 | + └─ 否（纯 elementwise/归约）→ P94/P95 | ||
| 43 | +``` | ||
| 44 | + | ||
| 45 | +## 执行入口（SOP 在 Playbook，调试手册在 references） | ||
| 46 | + | ||
| 47 | +- **完整实操 SOP（Step 0-6，强制执行）**：`playbooks/P102_scalar_matmul_to_cv_fusion_rewrite.md` | ||
| 48 | + - Step 0：必读 mix-cv-fusion-guide.md | ||
| 49 | + - Step 1-2：定位 + 改造计划表 | ||
| 50 | + - Step 3：**Phase 0 纯 AIC 单路径 derisk（禁止跳过直接写 MIX）** + canonical template + 自检 | ||
| 51 | + - Step 4：MIX 骨架 + 配对探针（五条铁律） | ||
| 52 | + - Step 5：逐 stage 增量 + 端到端验证 | ||
| 53 | + - Step 6：失败归因纪律（R1-R6 分类 + "平台不支持"结论的证据标准） | ||
| 54 | +- **MIX 机制与 hang 排查手册**：`references/mix-cv-fusion-guide.md`（R1-R6 根因分类、四阶段流程、双缓冲决策矩阵、调试工具箱） | ||
| 55 | + | ||
| 56 | +## 关键约束（节点描述必带） | ||
| 57 | + | ||
| 58 | +- **四阶段不可跳步**：AIC_ONLY 单 matmul 验证 → MIX 骨架 → 配对探针 → 逐 stage 增量。直接写全双工 MIX 是最高频失败模式。 | ||
| 59 | +- **hang/NaN ≠ 平台不支持**：MIX 在 910B2C 已实证可行（8.85x）。下"环境不支持"结论前必须满足 playbook Step 6 的四条证据标准（最小探针失败 + 清 build 缓存 + 换干净 NPU + 具体报错原文）。 | ||
| 60 | +- **build 缓存陷阱**：改代码必须 `rm -rf build`，否则"改了没生效"造成假诊断。 | ||
| 61 | +- **单 chunk 不算验证**：状态依赖路径必须 T≥128 多 chunk 用例。 | ||
| 62 | + | ||
| 63 | +## 常见陷阱 | ||
| 64 | + | ||
| 65 | +⚠️ **改了两样东西（launch type + 参数）时，hang 不一定归因于最后改的那个**——AIC_ONLY 编译运行正常 ⇒ kernel 逻辑无问题，launch 模式是病根，不要回退参数空转 | ||
| 66 | +⚠️ **build 缓存**：setup.py/构建脚本检测到 .so 存在会跳过 cmake——**改代码必须 `rm -rf build`**，否则"改了没生效"造成假诊断 | ||
| 67 | +⚠️ **带病 NPU**：被前一次 hang kernel 污染的设备会让 cross-core barrier 失效，后续输出非确定性 NaN（位置逐 run 变化）——同 binary 换一块干净 NPU PASS ⇒ 卡问题，与代码无关 | ||
| 68 | +⚠️ **单 chunk 不算验证**：状态依赖路径（S 更新、跨 chunk scale）在单 chunk 下恰好正确，必须 T≥128 多 chunk 用例 | ||
| 69 | +⚠️ **诊断前先 torch 算法仿真**：纯 torch 复刻 kernel 的分解公式再对参考，不花编译时间就能抓公式错误 | ||
| 70 | + | ||
| 71 | +## 代码搜索关键词 | ||
| 72 | + | ||
| 73 | +```bash | ||
| 74 | +# 检测标量/向量模拟矩阵运算（P102 适用信号） | ||
| 75 | +grep -nE "for\s*\(.*\+\+.*\)\s*\{?\s*for\s*\(" kernel/*.cpp # 嵌套标量循环 | ||
| 76 | +grep -nE "reinterpret_cast.*GetPhyAddr|\.GetPhyAddr\(\)" kernel/*.cpp # raw pointer 标量计算 | ||
| 77 | +grep -nE "Mul\(.*\).*ReduceSum|ReduceSum" kernel/*.cpp # Vector 拼伪 matmul | ||
| 78 | + | ||
| 79 | +# 检测是否已是 MIX 架构（已重写则不适用） | ||
| 80 | +grep -nE "KERNEL_TYPE_MIX|ASCEND_IS_AIC|Mmad" kernel/*.cpp | ||
| 81 | +``` | ||
| 82 | + | ||
| 83 | +## 来源 | ||
| 84 | + | ||
| 85 | +- gated_delta_rule CV 融合重写（a5_ops_slim/workspace/gated_delta_rule_cv）：scalar/VEC 3.3x 且比参考慢 17x → MIX_AIC_1_1 重写 geomean **8.85x**（T=64 时 14.0x），精度全 shape PASS（MERE < 9.766e-04）。方法论见该目录 docs/MIX_DEBUG_WORKFLOW.md | ||
| @@ -0,0 +1,88 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P103: Cube Neumann 乘积式三角求逆（替代标量 forward substitution）' | ||
| 4 | +description: 下三角线性系统求解（forward substitution / back substitution）是 O(C²) 的**标量串行链**：每一行依赖前面所有行的结果，无法向量化，是 | ||
| 5 | + MIX 算子中最长的串行尾。 | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- topic.programming_model.execution_unit.cube | ||
| 9 | +- topic.attention | ||
| 10 | +- topic.compute.composite | ||
| 11 | +- topic.programming_model.execution_unit.scalar | ||
| 12 | +status: stable | ||
| 13 | +sources: | ||
| 14 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P103_cube_neumann_triangular_solve.md | ||
| 15 | + role: primary | ||
| 16 | +created_at: '2026-10-09T00:00:00Z' | ||
| 17 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 18 | +platforms: | ||
| 19 | +- a2 | ||
| 20 | +--- | ||
| 21 | + | ||
| 22 | +# P103: Cube Neumann 乘积式三角求逆（替代标量 forward substitution） | ||
| 23 | + | ||
| 24 | +## 核心思想 | ||
| 25 | + | ||
| 26 | +下三角线性系统求解（forward substitution / back substitution）是 O(C²) 的**标量串行链**：每一行依赖前面所有行的结果，无法向量化，是 MIX 算子中最长的串行尾。数学上等价的替代： | ||
| 27 | + | ||
| 28 | +``` | ||
| 29 | +求解 (I + A) x = b，其中 A 严格下三角（A^C = 0，C 为块维） | ||
| 30 | +⇔ 求 (I - X)⁻¹，X = -A（X 严格下三角，X^C = 0，Neumann 级数有限截断） | ||
| 31 | +⇔ (I - X)⁻¹ = (I + X)(I + X²)(I + X⁴)(I + X⁸)(I + X¹⁶)(I + X³²) (C = 64) | ||
| 32 | +``` | ||
| 33 | + | ||
| 34 | +整条链 = **5 次平方 Mmad + 6 次乘积 Mmad = 11 个 Mmad**，全部在 Cube 上完成，替代 C(C-1)/2 ≈ 2016 次标量迭代。fp16 中间精度误差 ~1.2e-4（在 GDR 精度预算内）。 | ||
| 35 | + | ||
| 36 | +## 代码骨架（C=64 实例，来自 gated_delta_rule_cv） | ||
| 37 | + | ||
| 38 | +```cpp | ||
| 39 | +// 输入: X16 (x = -A2) 与 F0 (I + x)，均为 fp16 GM workspace slot | ||
| 40 | +// 输出: A16 (solved, 即 (I+X)^{-1} 的乘积结果) fp16 slot | ||
| 41 | + | ||
| 42 | +// ---- 平方链: x2 = x@x, x4 = x2@x2, ..., x32（各自 fp32→fp16 全往返）---- | ||
| 43 | +MMf16(X16, X16, X2, 64, 64, 64); PipeBarrier<PIPE_ALL>(); | ||
| 44 | +MMf16(X2, X2, X4, 64, 64, 64); PipeBarrier<PIPE_ALL>(); | ||
| 45 | +MMf16(X4, X4, X8, 64, 64, 64); PipeBarrier<PIPE_ALL>(); | ||
| 46 | +MMf16(X8, X8, X16P,64, 64, 64); PipeBarrier<PIPE_ALL>(); | ||
| 47 | +MMf16(X16P,X16P,X32, 64, 64, 64); PipeBarrier<PIPE_ALL>(); | ||
| 48 | + | ||
| 49 | +// ---- 乘积链: R 从 F0 开始，R ← R @ (I + x^{2^i}) = R + R @ x^{2^i} ---- | ||
| 50 | +// c1L (L0C) 整条链持有：首个 Mmad cmatrixInitVal=1（覆盖）， | ||
| 51 | +// 后续 cmatrixInitVal=0（L0C 内累加 R@P 到 R 上） | ||
| 52 | +auto c1L = c1_.AllocTensor<float>(); | ||
| 53 | +MmadParams mp{}; mp.m = 64; mp.n = 64; mp.k = 64; | ||
| 54 | +// i=0: R = F0（Mmad(F0, I16, init=1)），随后同一步折叠 F0@x2（init=0） | ||
| 55 | +// i=1..4: R += R @ x^{2^{i+1}} | ||
| 56 | +// 每步结束 Fixpipe 回 GM slot（除最后一步写 A16 外写 R16），PipeBarrier<PIPE_ALL>() | ||
| 57 | +c1_.FreeTensor(c1L); | ||
| 58 | +``` | ||
| 59 | + | ||
| 60 | +## 关键修改点 | ||
| 61 | + | ||
| 62 | +1. **L0C 驻留累加**：乘积链的 `R + R@P` 利用 `cmatrixInitVal=0` 在 L0C 内直接累加，避免 fp32 中间结果往返 GM——这是把 6 次乘降压到单次 L0C 驻留的关键 | ||
| 63 | +2. **F0@x2 折叠**：`i==0` 步在同一个 L0C 驻留窗口内连做两个 Mmad（先 `F0@I` 初始化，再 `F0@x2` 累加），省一次 Fixpipe 往返 | ||
| 64 | +3. **平方链必须 fp32→fp16 全往返**：每级平方 Fixpipe 回 fp16 GM slot 再加载，防止 fp16 下三角矩阵在 L0 上精度雪崩（实测 ~1.2e-4 总误差） | ||
| 65 | +4. **identity 由 host 预写**：`I16` 放 GM workspace 固定 slot，host 侧 `at::eye(64)` 一次性写入，kernel 内不构造 | ||
| 66 | +5. **同步纪律**：Mmad→Fixpipe、Fixpipe→下次 GM 读之间必须 `PipeBarrier<PIPE_ALL>`（L0 队列事件同步缺失会触发 aicore exception） | ||
| 67 | + | ||
| 68 | +## 适用性检测 (grep) | ||
| 69 | + | ||
| 70 | +```bash | ||
| 71 | +# 检测标量三角求解 / 递推循环（P103 适用信号） | ||
| 72 | +grep -nE "for\s*\(.*i.*\+\+.*\)\s*\{?\s*for\s*\(.*j\s*=\s*0.*j\s*<\s*i" op_kernel/*.cpp | ||
| 73 | +grep -nE "solve|substitution|forward_sub|linsolve" op_kernel/*.cpp | ||
| 74 | + | ||
| 75 | +# 前提：已是 MIX/Cube 架构（否则先走 P102） | ||
| 76 | +grep -nE "KERNEL_TYPE_MIX|ASCEND_IS_AIC" op_kernel/*.cpp | ||
| 77 | +``` | ||
| 78 | + | ||
| 79 | +## 常见陷阱 | ||
| 80 | + | ||
| 81 | +⚠️ **fp16 精度预算**：Neumann 链是 fp16 中间表示，误差随平方层级积累（~1.2e-4 @ C=64）；精度预算紧的算子需先用 torch 仿真验证该误差可接受，再进 kernel | ||
| 82 | +⚠️ **严格下三角前提**：X 必须严格下三角（对角为 0）才有 X^C=0 的有限截断；含对角的 mask 用错（inclusive vs strict）会导致结果系统性偏差且**单 chunk 测试无法暴露** | ||
| 83 | +⚠️ **C 维上限**：示例为 C=64（6 级乘积）；C=128 需 7 级（x64），workspace 与 L0C 占用重估 | ||
| 84 | +⚠️ **队列完整性**：每个 Mmad 操作数必须 Alloc→Load→EnQue→DeQue→Mmad→Free 完整走 a2_/b2_ 队列，缺事件同步 = aicore exception（非 hang，有报错） | ||
| 85 | + | ||
| 86 | +## 来源 | ||
| 87 | + | ||
| 88 | +- gated_delta_rule CV 融合重写（a5_ops_slim/workspace/gated_delta_rule_cv）：AIV 标量 forward substitution 98µs/chunk → Cube Neumann 链 11 Mmad，是 geomean 3.79x → 8.85x 的关键一跃 | ||
| @@ -0,0 +1,90 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P104: Matmul 多核切分扩展（Multi-Core Tiling Expansion）' | ||
| 4 | +description: '`MatmulApiTiling::GetTiling(TCubeTiling&)` 自动给出的 `singleCoreM / singleCoreN` | ||
| 5 | + 可能直接将整行/整列矩阵分配给一个核，导致 `totalBlock = 1`、`blockDim = 1`，多核完全未生效。' | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- topic.programming_model.execution_unit.cube | ||
| 9 | +- topic.parallelism | ||
| 10 | +- topic.implementation.tiling | ||
| 11 | +status: stable | ||
| 12 | +sources: | ||
| 13 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P104_matmul_multi_core_tiling_expansion.md | ||
| 14 | + role: primary | ||
| 15 | +created_at: '2026-10-09T00:00:00Z' | ||
| 16 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 17 | +platforms: | ||
| 18 | +- a2 | ||
| 19 | +--- | ||
| 20 | + | ||
| 21 | +# P104: Matmul 多核切分扩展（Multi-Core Tiling Expansion） | ||
| 22 | + | ||
| 23 | +## 核心思想 | ||
| 24 | + | ||
| 25 | +`MatmulApiTiling::GetTiling(TCubeTiling&)` 自动给出的 `singleCoreM / singleCoreN` 可能直接将整行/整列矩阵分配给一个核，导致 `totalBlock = 1`、`blockDim = 1`，多核完全未生效。P104 在 Host 侧**手动缩小 `singleCoreM / singleCoreN`**，使 `totalBlock ≥ aicCoreNum`，同时保持 16 字节对齐且不低于 `baseM / baseN`，从而把任务扩散到所有 AIC 核。 | ||
| 26 | + | ||
| 27 | +## 代码骨架 | ||
| 28 | + | ||
| 29 | +```cpp | ||
| 30 | +// === 改造前（基线）：GetTiling 后直接序列化，可能只用到 1 个核 === | ||
| 31 | +MatmulApiTiling::GetTiling(cubeTiling); | ||
| 32 | +// cubeTiling.singleCoreM 可能 == M，totalBlock == 1 | ||
| 33 | + | ||
| 34 | +// === 改造后：多核切分扩展 === | ||
| 35 | +constexpr uint32_t ALIGNED_H = 16; | ||
| 36 | +auto singleCoreM = cubeTiling.singleCoreM; | ||
| 37 | +auto singleCoreN = cubeTiling.singleCoreN; | ||
| 38 | +auto baseM = cubeTiling.baseM; | ||
| 39 | +auto baseN = cubeTiling.baseN; | ||
| 40 | +auto aicCoreNum = platform.GetCoreNumAic(); | ||
| 41 | + | ||
| 42 | +int64_t mTotalCnt = CeilDiv(M, static_cast<int64_t>(singleCoreM)); | ||
| 43 | +int64_t nTotalCnt = CeilDiv(N, static_cast<int64_t>(singleCoreN)); | ||
| 44 | +int64_t totalBlock = mTotalCnt * nTotalCnt; | ||
| 45 | + | ||
| 46 | +if (totalBlock < aicCoreNum) { | ||
| 47 | + while (totalBlock < aicCoreNum && singleCoreM > baseM) { | ||
| 48 | + singleCoreM = std::max(baseM, (singleCoreM / 2) / ALIGNED_H * ALIGNED_H); | ||
| 49 | + mTotalCnt = CeilDiv(M, static_cast<int64_t>(singleCoreM)); | ||
| 50 | + totalBlock = mTotalCnt * nTotalCnt; | ||
| 51 | + } | ||
| 52 | + while (totalBlock < aicCoreNum && singleCoreN > baseN) { | ||
| 53 | + singleCoreN = std::max(baseN, (singleCoreN / 2) / ALIGNED_H * ALIGNED_H); | ||
| 54 | + nTotalCnt = CeilDiv(N, static_cast<int64_t>(singleCoreN)); | ||
| 55 | + totalBlock = mTotalCnt * nTotalCnt; | ||
| 56 | + } | ||
| 57 | +} | ||
| 58 | + | ||
| 59 | +cubeTiling.singleCoreM = singleCoreM; | ||
| 60 | +cubeTiling.singleCoreN = singleCoreN; | ||
| 61 | +cubeTiling.usedCoreNum = std::min(totalBlock, aicCoreNum); | ||
| 62 | + | ||
| 63 | +int32_t mBaseTail = M - (mTotalCnt - 1) * singleCoreM; | ||
| 64 | +int32_t nBaseTail = N - (nTotalCnt - 1) * singleCoreN; | ||
| 65 | +``` | ||
| 66 | + | ||
| 67 | +## 关键修改点 | ||
| 68 | + | ||
| 69 | +1. **强制对齐**：缩小后的 `singleCoreM / singleCoreN` 必须是 `ALIGNED_H = 16` 的整数倍。 | ||
| 70 | +2. **下限保护**：不得小于 `baseM / baseN`，否则 cube 一次发射的尺寸无法保证。 | ||
| 71 | +3. **同步更新派生字段**：`mTotalCnt / nTotalCnt / totalBlock / usedCoreNum / mBaseTail / nBaseTail` 全部重新计算。 | ||
| 72 | +4. **Kernel 尾块守卫**：Kernel 端按 `(mIdx == mTotalCnt - 1) ? mBaseTail : singleCoreM` 调用 `SetSingleShape`。 | ||
| 73 | + | ||
| 74 | +## 常见陷阱 | ||
| 75 | + | ||
| 76 | +⚠️ 缩小后未重新计算尾块，导致最后一块越界写 GM。 | ||
| 77 | +⚠️ `usedCoreNum` 未取 `min(totalBlock, aicCoreNum)`，富余核无任务空转。 | ||
| 78 | +⚠️ Kernel 入口缺少 `if (GetBlockIdx() >= totalBlock) return;`，导致越界。 | ||
| 79 | +⚠️ 将 `singleCoreM / N` 缩小到低于 `baseM / N`，Cube 发射异常或性能崩溃。 | ||
| 80 | + | ||
| 81 | +## 代码搜索关键词 | ||
| 82 | + | ||
| 83 | +```bash | ||
| 84 | +grep -nE "MatmulApiTiling|GetTiling|singleCoreM|singleCoreN|baseM|baseN|totalBlock|usedCoreNum|GetCoreNumAic" \ | ||
| 85 | + op_host/*_tiling.cpp op_kernel/*.cpp | ||
| 86 | +``` | ||
| 87 | + | ||
| 88 | +## 来源 | ||
| 89 | + | ||
| 90 | +- `ops/ascendc-tiling-design/references/matmul/ascendc-api-matmul-tiling.md` §多核切分扩展 | ||
| @@ -0,0 +1,73 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P105: Matmul L2 Cache Tile 尺寸规划（L2 Cache Tile Sizing）' | ||
| 4 | +description: 大矩阵乘中，若一次性把所有 base 块顺序下发，A/B 子块在 L2 中会被后续 tile 反复挤出，造成 MTE2 重读。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.programming_model.execution_unit.cube | ||
| 8 | +- topic.data_movement.memory_access | ||
| 9 | +status: stable | ||
| 10 | +sources: | ||
| 11 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P105_matmul_l2_cache_tile_sizing.md | ||
| 12 | + role: primary | ||
| 13 | +created_at: '2026-10-09T00:00:00Z' | ||
| 14 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 15 | +platforms: | ||
| 16 | +- a2 | ||
| 17 | +--- | ||
| 18 | + | ||
| 19 | +# P105: Matmul L2 Cache Tile 尺寸规划（L2 Cache Tile Sizing） | ||
| 20 | + | ||
| 21 | +## 核心思想 | ||
| 22 | + | ||
| 23 | +大矩阵乘中，若一次性把所有 base 块顺序下发，A/B 子块在 L2 中会被后续 tile 反复挤出，造成 MTE2 重读。P105 在 Host 侧增加 **L2 级 Tiling**：把 M/N 方向切分为若干 L2 tile，每个 L2 tile 内的 base 块共享 L2 中的 A 行 / B 列，并保证 `mTileBlock * nTileBlock * dtypeSize` 不超过 L2 容量的一半，从而提升数据局部性、降低 MTE2 压力。 | ||
| 24 | + | ||
| 25 | +## 代码骨架 | ||
| 26 | + | ||
| 27 | +```cpp | ||
| 28 | +// === 改造前：只有 Block/BaseBlock 两级切分 === | ||
| 29 | +L2cacheTilePara l2Para{}; // 未使用或 mTileCntL2 = nTileCntL2 = 1 | ||
| 30 | + | ||
| 31 | +// === 改造后：引入 L2 级切分 === | ||
| 32 | +constexpr uint32_t L2_SIZE = 2 * 1024 * 1024; // 以实际 platform 为准 | ||
| 33 | +L2cacheTilePara l2Para{}; | ||
| 34 | + | ||
| 35 | +// 小算子退化 | ||
| 36 | +if (M * N * dtypeSize > L2_SIZE / 2) { | ||
| 37 | + l2Para.mTileCntL2 = ceil(mTotalCnt, l2Para.mTileBlock); | ||
| 38 | + l2Para.nTileCntL2 = ceil(nTotalCnt, l2Para.nTileBlock); | ||
| 39 | + // 让单个 L2 tile 的 base 块 footprint ≈ L2/2 | ||
| 40 | + // 例如：l2Para.mTileBlock = ceil(mTotalCnt, l2Para.mTileCntL2); | ||
| 41 | + // l2Para.nTileBlock = ceil(nTotalCnt, l2Para.nTileCntL2); | ||
| 42 | +} else { | ||
| 43 | + l2Para.mTileCntL2 = 1; | ||
| 44 | + l2Para.nTileCntL2 = 1; | ||
| 45 | +} | ||
| 46 | + | ||
| 47 | +l2Para.calOrder = (M >= N) ? ROW_FIRST : COL_FIRST; | ||
| 48 | +``` | ||
| 49 | + | ||
| 50 | +## 关键修改点 | ||
| 51 | + | ||
| 52 | +1. **容量 guard**：单个 L2 tile 内 A/B 的工作集应 `≤ L2 / 2`，否则退化为 1。 | ||
| 53 | +2. **calOrder 选择**：`M ≫ N` 选 `ROW_FIRST` 复用 A 行；`N ≫ M` 选 `COL_FIRST` 复用 B 列。 | ||
| 54 | +3. **派生计数**：`mCnt = ceil(mTotalCnt, mTileCntL2)`，`nCnt = ceil(nTotalCnt, nTileCntL2)`，`totalTileCnt = mCnt * nCnt`。 | ||
| 55 | +4. **与 Block 级配合**：Block 级仍按 `totalTileCnt` 分配核，L2 级只改变遍历顺序。 | ||
| 56 | + | ||
| 57 | +## 常见陷阱 | ||
| 58 | + | ||
| 59 | +⚠️ 小算子仍启用 L2 切分，引入不必要的调度开销。 | ||
| 60 | +⚠️ `mTileBlock / nTileBlock` 未对齐到 base 块整数倍，导致 L2 tile 边界尾块处理错误。 | ||
| 61 | +⚠️ `calOrder` 与 serpentine 遍历方向冲突：ROW_FIRST 时仍可用 serpentine，只是优先维度不同。 | ||
| 62 | +⚠️ L2 tile footprint 计算未包含双缓冲 / L1 占用，实际可用容量比 L2_SIZE / 2 更小。 | ||
| 63 | + | ||
| 64 | +## 代码搜索关键词 | ||
| 65 | + | ||
| 66 | +```bash | ||
| 67 | +grep -nE "L2cacheTilePara|mTileCntL2|nTileCntL2|mTileBlock|nTileBlock|calOrder|ROW_FIRST|COL_FIRST" \ | ||
| 68 | + op_host/*_tiling.cpp op_kernel/*.cpp | ||
| 69 | +``` | ||
| 70 | + | ||
| 71 | +## 来源 | ||
| 72 | + | ||
| 73 | +- `ops/ascendc-tiling-design/references/matmul/ascendc-api-matmul-tiling.md` §L2 级 Tiling | ||
Aknowledge/contrib/ascendc/optimizations/evolution_strategies/cards/p106_matmul_batch_merge_to_m.md+66-0
| @@ -0,0 +1,66 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P106: Matmul Batch 维度合并到 M（Batch Merge to M）' | ||
| 4 | +description: 在 BatchMatMul 中，当 B 矩阵无 batch（`batchB == 1`）且 A 未转置（`!transA`）时，可以把 A | ||
| 5 | + 的 batch 维度折叠进 M 维度：`M' = batchA * M`、`batchA' = 1`。 | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- topic.programming_model.execution_unit.cube | ||
| 9 | +- topic.parallelism | ||
| 10 | +- topic.implementation.tiling | ||
| 11 | +status: stable | ||
| 12 | +sources: | ||
| 13 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P106_matmul_batch_merge_to_m.md | ||
| 14 | + role: primary | ||
| 15 | +created_at: '2026-10-09T00:00:00Z' | ||
| 16 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 17 | +platforms: | ||
| 18 | +- a2 | ||
| 19 | +--- | ||
| 20 | + | ||
| 21 | +# P106: Matmul Batch 维度合并到 M（Batch Merge to M） | ||
| 22 | + | ||
| 23 | +## 核心思想 | ||
| 24 | + | ||
| 25 | +在 BatchMatMul 中，当 B 矩阵无 batch（`batchB == 1`）且 A 未转置（`!transA`）时，可以把 A 的 batch 维度折叠进 M 维度：`M' = batchA * M`、`batchA' = 1`。这样原先把每个 batch 当作独立调度单元的问题，转化为更大的 M 维度问题，更容易触发多核切分扩展（P104），减少 `IterateBatch` 的跨 batch 调度开销。 | ||
| 26 | + | ||
| 27 | +## 代码骨架 | ||
| 28 | + | ||
| 29 | +```cpp | ||
| 30 | +// === 改造前：batchA 与 M 分开处理 === | ||
| 31 | +uint32_t batchA = aShape.size() > 2 ? aShape[aShape.size() - 3] : 1; | ||
| 32 | +uint32_t batchB = bShape.size() > 2 ? bShape[bShape.size() - 3] : 1; | ||
| 33 | +uint32_t M = aShape[aShape.size() - 2]; | ||
| 34 | + | ||
| 35 | +// === 改造后：在合法条件下把 batchA 合并进 M === | ||
| 36 | +if (batchB == 1 && !transA) { | ||
| 37 | + M = batchA * M; | ||
| 38 | + batchA = 1; | ||
| 39 | + // 后续 tiling 以新的 M 为基准 | ||
| 40 | +} | ||
| 41 | +``` | ||
| 42 | + | ||
| 43 | +## 关键修改点 | ||
| 44 | + | ||
| 45 | +1. **前提严格判断**：仅当 `batchB == 1 && !transA` 时合并；`transA` 为 true 时 A 的 batch 与 M 在物理上不连续，不能折叠。 | ||
| 46 | +2. **Broadcast 表更新**：合并后 C 的 batch 由 `batchC` 重新按广播规则计算，避免维度错位。 | ||
| 47 | +3. **Kernel offset 不变**：因为 batch 已合并进 M，Kernel 端仍按合并后的 `(M', N, K)` 计算 `blockIdx → (mIdx, nIdx)`，无需额外 batch 索引。 | ||
| 48 | +4. **Bias 处理**：若带 bias 且 bias 原本按 batch 排布，需确认合并后 bias 广播语义仍成立。 | ||
| 49 | + | ||
| 50 | +## 常见陷阱 | ||
| 51 | + | ||
| 52 | +⚠️ `transA == true` 时仍合并，导致 A 的 stride 错误，结果错位。 | ||
| 53 | +⚠️ 合并后未更新 `batchC`，Host/KERNEL 参数不一致。 | ||
| 54 | +⚠️ 合并前 `batchA * M` 溢出 `uint32_t`，应使用 `int64_t` 中间计算。 | ||
| 55 | +⚠️ Bias 或 Scale 带有 batch 维度时，合并后广播失败。 | ||
| 56 | + | ||
| 57 | +## 代码搜索关键词 | ||
| 58 | + | ||
| 59 | +```bash | ||
| 60 | +grep -nE "batchB|batchA|transA|MultiBatchInfo|IterateBatch|batchC" \ | ||
| 61 | + op_host/*_tiling.cpp op_kernel/*.cpp | ||
| 62 | +``` | ||
| 63 | + | ||
| 64 | +## 来源 | ||
| 65 | + | ||
| 66 | +- `ops/ascendc-tiling-design/references/matmul/ascendc-api-matmul-tiling.md` §Batch 合并到 M | ||
| @@ -0,0 +1,69 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P107: Matmul Serpentine 反向遍历（Serpentine Traversal）' | ||
| 4 | +description: 在 L2 级 Tiling 中，若按行优先正向遍历 N 方向 tile，每切换一行 M 就要重新从 L2 加载远端 B 行。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.programming_model.execution_unit.cube | ||
| 8 | +- topic.data_movement.memory_access | ||
| 9 | +status: stable | ||
| 10 | +sources: | ||
| 11 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P107_matmul_serpentine_traversal.md | ||
| 12 | + role: primary | ||
| 13 | +created_at: '2026-10-09T00:00:00Z' | ||
| 14 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 15 | +platforms: | ||
| 16 | +- a2 | ||
| 17 | +--- | ||
| 18 | + | ||
| 19 | +# P107: Matmul Serpentine 反向遍历（Serpentine Traversal） | ||
| 20 | + | ||
| 21 | +## 核心思想 | ||
| 22 | + | ||
| 23 | +在 L2 级 Tiling 中，若按行优先正向遍历 N 方向 tile，每切换一行 M 就要重新从 L2 加载远端 B 行。P107 采用 **serpentine 遍历**：偶数行 M 正向走 N，奇数行 M 反向走 N，使相邻两行在 L2 tile 边界处共享同一组 B 行，提升 L2 命中率并减少 MTE2 读 B。 | ||
| 24 | + | ||
| 25 | +## 代码骨架 | ||
| 26 | + | ||
| 27 | +```cpp | ||
| 28 | +// === 改造前：正向双层循环 === | ||
| 29 | +for (uint64_t mTile = 0; mTile < mTileCntL2; ++mTile) { | ||
| 30 | + for (uint64_t nTile = 0; nTile < nTileCntL2; ++nTile) { | ||
| 31 | + ProcessL2Tile(mTile, nTile); | ||
| 32 | + } | ||
| 33 | +} | ||
| 34 | + | ||
| 35 | +// === 改造后：Serpentine 反向遍历 === | ||
| 36 | +bool reverse = true; | ||
| 37 | +for (uint64_t mTile = 0; mTile < mTileCntL2; ++mTile) { | ||
| 38 | + reverse = !reverse; | ||
| 39 | + for (uint64_t t = 0; t < nTileCntL2; ++t) { | ||
| 40 | + uint64_t nTile = reverse ? (nTileCntL2 - t - 1) : t; | ||
| 41 | + ProcessL2Tile(mTile, nTile); | ||
| 42 | + } | ||
| 43 | +} | ||
| 44 | +``` | ||
| 45 | + | ||
| 46 | +## 关键修改点 | ||
| 47 | + | ||
| 48 | +1. **作用域**：仅当 `nTileCntL2 > 1` 时生效；只有 1 列 tile 时无需反向。 | ||
| 49 | +2. **与 calOrder 配合**：`calOrder` 决定 L2 tile 内 base 块遍历方向；serpentine 决定 L2 tile 之间的遍历方向，二者正交。 | ||
| 50 | +3. **同步语义不变**：每个 L2 tile 内部仍按原有 Iterate/SetSingleShape 流程执行，只是下发顺序改变。 | ||
| 51 | +4. **负载均衡**：serpentine 不改变 `totalTileCnt`，也不改变 Block 级核间分配。 | ||
| 52 | + | ||
| 53 | +## 常见陷阱 | ||
| 54 | + | ||
| 55 | +⚠️ 反向遍历导致 kernel 端依赖 tile 下发顺序（如 Atomic Add 顺序敏感）时结果非确定性。 | ||
| 56 | +⚠️ `nTileCntL2 == 1` 时仍加入 serpentine，引入无意义分支。 | ||
| 57 | +⚠️ 与 `COL_FIRST` 组合时，优先复用 B 列而非 B 行，需重新评估收益。 | ||
| 58 | +⚠️ 若 `ProcessL2Tile` 内部有 workspace 状态跨 tile 依赖，serpentine 会打破顺序。 | ||
| 59 | + | ||
| 60 | +## 代码搜索关键词 | ||
| 61 | + | ||
| 62 | +```bash | ||
| 63 | +grep -nE "serpentine|reverse\s*=|mTileCntL2|nTileCntL2|ProcessL2Tile|calOrder" \ | ||
| 64 | + op_host/*_tiling.cpp op_kernel/*.cpp | ||
| 65 | +``` | ||
| 66 | + | ||
| 67 | +## 来源 | ||
| 68 | + | ||
| 69 | +- `ops/ascendc-tiling-design/references/matmul/ascendc-api-matmul-tiling.md` §Serpentine 反向遍历 | ||
Aknowledge/contrib/ascendc/optimizations/evolution_strategies/cards/p108_matmul_platform_routing.md+69-0
| @@ -0,0 +1,69 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P108: Matmul 平台路由（Platform Routing）' | ||
| 4 | +description: Ascend C Matmul Tiling 设计强烈依赖硬件特征（`NpuArch`、`GetCoreNumAic`、L1/L0 容量、Cube | ||
| 5 | + 数据通路）。 | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- topic.programming_model.execution_unit.cube | ||
| 9 | +status: stable | ||
| 10 | +sources: | ||
| 11 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P108_matmul_platform_routing.md | ||
| 12 | + role: primary | ||
| 13 | +created_at: '2026-10-09T00:00:00Z' | ||
| 14 | +updated_at: '2026-10-13T00:00:00Z' | ||
| 15 | +platforms: | ||
| 16 | +- a2 | ||
| 17 | +--- | ||
| 18 | + | ||
| 19 | +# P108: Matmul 平台路由（Platform Routing） | ||
| 20 | + | ||
| 21 | +## 核心思想 | ||
| 22 | + | ||
| 23 | +Ascend C Matmul Tiling 设计强烈依赖硬件特征（`NpuArch`、`GetCoreNumAic`、L1/L0 容量、Cube 数据通路）。`DAV_2201`（Atlas A2 / A3）适用 `MatmulImpl` + `MatmulApiTiling` 路径；而 `DAV_3510`（Ascend 950）应走 Blaze / tensor_api 路径。P108 要求 Host 在编译期或运行期按 `NpuArch` 路由到对应的 Tiling / Kernel 实现，避免把 A2/A3 的 `TCubeTiling` 假设套用到不兼容平台。 | ||
| 24 | + | ||
| 25 | +> 注：本卡实测平台为 a2；`DAV_3510`（Ascend 950）的 Blaze / tensor_api 路径为设计约定与可迁移推断，未在 950 上实测。 | ||
| 26 | + | ||
| 27 | +## 代码骨架 | ||
| 28 | + | ||
| 29 | +```cpp | ||
| 30 | +// === Host 端平台路由示例 === | ||
| 31 | +auto platform = PlatformAscendC(context); | ||
| 32 | +auto socVersion = platform.GetSocVersion(); // 或从 attr 读取 | ||
| 33 | + | ||
| 34 | +if (socVersion == "Ascend910B" || socVersion == "Ascend910_93") { | ||
| 35 | + // DAV_2201 路径 | ||
| 36 | + MatmulApiTiling::GetTiling(cubeTiling); | ||
| 37 | + // 应用 P104/P105/P106/P107 | ||
| 38 | +} else if (socVersion == "Ascend950") { | ||
| 39 | + // DAV_3510 路径：走 Blaze / tensor_api | ||
| 40 | + // 不应用本 guide 中的 TCubeTiling 假设 | ||
| 41 | +} else { | ||
| 42 | + // 未知平台：兜底报错或走通用路径 | ||
| 43 | +} | ||
| 44 | +``` | ||
| 45 | + | ||
| 46 | +## 关键修改点 | ||
| 47 | + | ||
| 48 | +1. **版本来源**：优先从 `PlatformAscendC` 获取；若 CANN 版本差异大，可用 `ascendc-docs-search` 确认字段名。 | ||
| 49 | +2. **条件编译 vs 运行时分支**：对性能敏感路径可用 TilingKey / 模板特化；调试期可用运行时 `if`。 | ||
| 50 | +3. **隔离平台差异**：不同平台使用独立的 Tiling 结构体，避免共用 `TCubeTiling` 导致字段错位。 | ||
| 51 | +4. **扩展约定**：新增架构出现时，按 `NpuArch` 新增 reference 文件或更新平台适配说明，禁止在原有 A2/A3 路径中硬编码新平台逻辑。 | ||
| 52 | + | ||
| 53 | +## 常见陷阱 | ||
| 54 | + | ||
| 55 | +⚠️ 在 `DAV_3510` 上仍使用 `MatmulApiTiling`，编译或运行时报 unsupported。 | ||
| 56 | +⚠️ 用字符串比较 `socVersion` 时未覆盖别名（如 `Ascend910B2C`、`Ascend910B3` 等）。 | ||
| 57 | +⚠️ 平台分支后未重新计算 `blockDim`、`usedCoreNum` 等派生参数。 | ||
| 58 | +⚠️ 把 `DAV_2201` 的 `baseM=256` 直接套用到不支持该 Cube 粒度的架构。 | ||
| 59 | + | ||
| 60 | +## 代码搜索关键词 | ||
| 61 | + | ||
| 62 | +```bash | ||
| 63 | +grep -nE "PlatformAscendC|GetSocVersion|NpuArch|DAV_2201|DAV_3510|Ascend950|Ascend910B" \ | ||
| 64 | + op_host/*_tiling.cpp op_kernel/*.cpp | ||
| 65 | +``` | ||
| 66 | + | ||
| 67 | +## 来源 | ||
| 68 | + | ||
| 69 | +- `ops/ascendc-tiling-design/references/matmul/ascendc-api-matmul-tiling.md` §适用平台 / §平台路由 | ||
| @@ -0,0 +1,61 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P109: Reduction 多轴 Tiling 系统化设计' | ||
| 4 | +description: 将任意 N 维 shape + axes 先合轴为 A/R 交替序列，再按单轴归约（AR/ARA）逐维处理。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.compute.reduction | ||
| 8 | +- topic.compute.normalization | ||
| 9 | +- topic.implementation.tiling | ||
| 10 | +- topic.programming_model.execution_unit.scalar | ||
| 11 | +status: stable | ||
| 12 | +sources: | ||
| 13 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P109_reduction_multi_axis_tiling.md | ||
| 14 | + role: primary | ||
| 15 | +created_at: '2026-10-09T00:00:00Z' | ||
| 16 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 17 | +platforms: | ||
| 18 | +- a2 | ||
| 19 | +--- | ||
| 20 | + | ||
| 21 | +# P109: Reduction 多轴 Tiling 系统化设计 | ||
| 22 | + | ||
| 23 | +## 核心思想 | ||
| 24 | + | ||
| 25 | +将任意 N 维 shape + axes 先合轴为 A/R 交替序列，再按单轴归约（AR/ARA）逐维处理。避免在多维场景下直接手写嵌套归约循环，而是把问题分解为： | ||
| 26 | + | ||
| 27 | +1. **合轴**：消除 size=1 维度，合并相邻同类型轴。 | ||
| 28 | +2. **模式选择**：A0=1 走 AR（尾轴连续），A0>1 走 ARA（非尾轴块）。 | ||
| 29 | +3. **多轴展开**：ARARA 等场景使用 Shape 三步变换，把每个 R 维单独归约。 | ||
| 30 | + | ||
| 31 | +## 代码骨架 | ||
| 32 | + | ||
| 33 | +```cpp | ||
| 34 | +// === ARA 模式：非尾轴归约 === | ||
| 35 | +uint32_t alignedCols = ((tileA0Len * sizeof(float) + 31) / 32) * 32 / sizeof(float); | ||
| 36 | +uint32_t srcShape[] = {R, alignedCols}; | ||
| 37 | +ReduceMax<float, Pattern::Reduce::RA>(resultLocal, xLocal, tmpLocal, srcShape, true); | ||
| 38 | + | ||
| 39 | +// === AR 模式：尾轴归约 === | ||
| 40 | +ReduceSum<float, true>(scalarLocal, xLocal, reduceTmp, rLength); | ||
| 41 | +``` | ||
| 42 | + | ||
| 43 | +## 关键修改点 | ||
| 44 | + | ||
| 45 | +1. **合轴函数**：host/tiling 侧先把 `shape` + `axes` 转成 A/R 序列，再决定 kernel 循环结构。 | ||
| 46 | +2. **A0 判断**：若合轴后首维为 1，说明 R 是尾轴，用 Level 2 Reduce API；否则用 `Pattern::Reduce::RA`。 | ||
| 47 | +3. **多轴归约**：每处理完一个 R 维，输出 shape 降维，继续处理下一个 R。 | ||
| 48 | +4. **`rLengthAlign` 用于 buffer 分配和 offset，`rLength` 用于 Reduce API count 和 DataCopyPad blockLen**。 | ||
| 49 | + | ||
| 50 | +## 常见陷阱 | ||
| 51 | + | ||
| 52 | +- ❌ 不归约的维度被误合并到 R 中 → 导致结果 shape 错误。 | ||
| 53 | +- ❌ AR/ARA 选择只看 `axis` 是否尾轴，未考虑合轴后的 A0。 | ||
| 54 | +- ❌ `alignedCols` 未按 32B 对齐 → `Pattern::Reduce::RA` 行为异常。 | ||
| 55 | +- ❌ 多轴场景直接一次性归约所有轴 → UB 放不下，应逐级处理。 | ||
| 56 | + | ||
| 57 | +## 代码搜索关键词 | ||
| 58 | + | ||
| 59 | +```bash | ||
| 60 | +grep -nE "ReduceMax|ReduceSum|Pattern::Reduce::RA|alignedCols|rLength|rLengthAlign" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 61 | +``` | ||
| @@ -0,0 +1,66 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P10: Vectorized Data Copy & Instructions (向量化数据与指令)' | ||
| 4 | +description: 参考实现在mask生成逻辑中采用了Compare+Select的组合指令，而非简单的CompareScalar。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.compute.composite | ||
| 8 | +- topic.data_movement.memory_access | ||
| 9 | +- topic.implementation.pipeline | ||
| 10 | +status: stable | ||
| 11 | +sources: | ||
| 12 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P10_vectorized_copy.md | ||
| 13 | + role: primary | ||
| 14 | +created_at: '2026-10-09T00:00:00Z' | ||
| 15 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 16 | +platforms: | ||
| 17 | +- a2 | ||
| 18 | +--- | ||
| 19 | + | ||
| 20 | +# P10: Vectorized Data Copy & Instructions (向量化数据与指令) | ||
| 21 | + | ||
| 22 | +## 核心思想 | ||
| 23 | +参考实现在mask生成逻辑中采用了Compare+Select的组合指令，而非简单的CompareScalar。这种策略的优势在于：1) Compare指令产生bitmask结果，可以进行向量化操作；2) Select指令根据bitmask进行条件选择，实现了向量化条件判断；3) 通过两次Compare+Select组合，可以高效地实现区间判断（quant_min <= x <= quant_max）。相比之下，参考实现使用了CompareScalar+手动循环AND操作的方式，虽然逻辑清晰，但无法充分利用向量化引擎。参考实现还通过Mul指令（Mul(curHf16Temp, selectTemp, curHf16Temp, calCount)）实现了逻辑AND操作，利用FP16的数值计算代替位运算，这也是一种向量化友好的优化策略。 | ||
| 24 | + | ||
| 25 | +## 代码骨架 | ||
| 26 | + | ||
| 27 | +// === 改造后（优化实现）=== | ||
| 28 | +```cpp | ||
| 29 | +template <typename T> | ||
| 30 | +__aicore__ inline void TransposeBase16M8(LocalTensor<T>& dstUb, LocalTensor<T>& srcUb, | ||
| 31 | + uint64_t rowNum, uint64_t colNum) { | ||
| 32 | + uint64_t srcAddrList[TRANS_ADDR_LEN]; | ||
| 33 | + uint64_t dstAddrList[TRANS_ADDR_LEN]; | ||
| 34 | + for (uint64_t r = 0; r < rowNum / TRANS_ADDR_LEN; r++) { | ||
| 35 | + for (uint64_t i = 0; i < TRANS_ADDR_LEN; i++) { | ||
| 36 | + srcAddrList[i] = (uint64_t)(srcUb[r * TRANS_ADDR_LEN * colNum + i * colNum].GetPhyAddr()); | ||
| 37 | + dstAddrList[i] = (uint64_t)(dstUb[r * TRANS_ADDR_LEN + i / 2 * rowNum + i % 2 * BLOCK_NUM_32].GetPhyAddr()); | ||
| 38 | + } | ||
| 39 | + struct TransDataTo5HDParams transDataParams; | ||
| 40 | + transDataParams.repeatTimes = colNum / BLOCK_NUM_32; | ||
| 41 | + TransDataTo5HD<float>(dstAddrList, srcAddrList, transDataParams); | ||
| 42 | + } | ||
| 43 | +} | ||
| 44 | + | ||
| 45 | +if constexpr (is_same<TGrad, float>::value) { | ||
| 46 | + TransposeBase16M8(gradTranUb, gradUb, params_.singleCoreNc, block_.dohowoAlign8); | ||
| 47 | +} else { | ||
| 48 | + TransposeBase16M16(gradTranUb, gradUb, params_.singleCoreNc, block_.dohowoAlign16); | ||
| 49 | +} | ||
| 50 | +``` | ||
| 51 | + | ||
| 52 | +## 关键修改点 | ||
| 53 | + | ||
| 54 | +1. 预期收益: 向量化转置可将内存访问模式优化为连续访问，减少bank conflict，提升Vector指令效率4-8倍 | ||
| 55 | + | ||
| 56 | +## 常见陷阱 | ||
| 57 | + | ||
| 58 | +⚠️ 需要额外的UB buffer存储转置数据 | ||
| 59 | +⚠️ 增加一层函数调用开销(通常可内联) | ||
| 60 | +⚠️ 对于单次运算略显复杂 | ||
| 61 | + | ||
| 62 | +## 代码搜索关键词 | ||
| 63 | + | ||
| 64 | +```bash | ||
| 65 | +grep -n "BUFFER_NUM\|InitBuffer\|TQue\|SyncAll\|PipeBarrier\|ExecuteTask\|PRELOAD\|SyncAll" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 66 | +``` | ||
| @@ -0,0 +1,64 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P110: Reduction UB 累加器复用' | ||
| 4 | +description: 在 AR/ARA FullLoad 模式下，整行/整块数据一次驻留 UB，归约后的中间结果直接复用 UB buffer 作为累加器，避免重复搬入；在 | ||
| 5 | + Split 模式下，分 chunk 归约后通过 Max/Add 跨 chunk 合并，同样需要合理分配 `tmpBuf` 与 accumulator。 | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- topic.compute.reduction | ||
| 9 | +- topic.compute.normalization | ||
| 10 | +- topic.memory | ||
| 11 | +- topic.programming_model.execution_unit.scalar | ||
| 12 | +status: stable | ||
| 13 | +sources: | ||
| 14 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P110_reduction_ub_accumulator_reuse.md | ||
| 15 | + role: primary | ||
| 16 | +created_at: '2026-10-09T00:00:00Z' | ||
| 17 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 18 | +platforms: | ||
| 19 | +- a2 | ||
| 20 | +--- | ||
| 21 | + | ||
| 22 | +# P110: Reduction UB 累加器复用 | ||
| 23 | + | ||
| 24 | +## 核心思想 | ||
| 25 | + | ||
| 26 | +在 AR/ARA FullLoad 模式下，整行/整块数据一次驻留 UB，归约后的中间结果直接复用 UB buffer 作为累加器，避免重复搬入；在 Split 模式下，分 chunk 归约后通过 Max/Add 跨 chunk 合并，同样需要合理分配 `tmpBuf` 与 accumulator。 | ||
| 27 | + | ||
| 28 | +## 代码骨架 | ||
| 29 | + | ||
| 30 | +```cpp | ||
| 31 | +// === FullLoad：中间结果直接复用 UB === | ||
| 32 | +DataCopyPad(xLocal, gmInput, blockLen); // 一次搬入整行 | ||
| 33 | +ReduceMax<float, true>(outLocal, xLocal, tmpLocal, rLength); | ||
| 34 | +DataCopyPad(gmOutput, outLocal, ...); | ||
| 35 | + | ||
| 36 | +// === Split：分 chunk 合并 === | ||
| 37 | +for (uint32_t chunk = 0; chunk < numChunks; ++chunk) { | ||
| 38 | + DataCopyPad(partLocal, gmInput + chunk * Ar, Ar); | ||
| 39 | + ReduceMax<float, true>(partLocal, partLocal, tmpLocal, Ar); | ||
| 40 | + // 跨 chunk 合并 | ||
| 41 | + if (chunk == 0) { Copy(accLocal, partLocal); } | ||
| 42 | + else { Max(accLocal, accLocal, partLocal, ...); } | ||
| 43 | +} | ||
| 44 | +``` | ||
| 45 | + | ||
| 46 | +## 关键修改点 | ||
| 47 | + | ||
| 48 | +1. **Buffer 预算**：`tmpBufSize` 必须按 `rLengthAlign` 分配，Reduce API workspace 单独计算。 | ||
| 49 | +2. **FullLoad 条件**：UB 至少能放 1 整行 R（AR）或 `R × tileA0Len`（ARA）。 | ||
| 50 | +3. **Split 合并**：跨 chunk 使用与算子语义对应的二元 API（Sum 用 Add，Max 用 Max）。 | ||
| 51 | +4. **精度敏感**：大 R sum 可叠加 Welford Online 或二分累加算法。 | ||
| 52 | + | ||
| 53 | +## 常见陷阱 | ||
| 54 | + | ||
| 55 | +- ❌ FullLoad 误判为 Split → 多轮 DataCopy 浪费带宽。 | ||
| 56 | +- ❌ `tmpLocal` 与 `outLocal` 复用同一块 UB → Reduce API 覆盖输出。 | ||
| 57 | +- ❌ Split 场景 chunk 边界未按 32B 对齐 → DataCopyPad 参数错误。 | ||
| 58 | +- ❌ 未区分 `rLength` 与 `rLengthAlign` → buffer 越界或结果不对。 | ||
| 59 | + | ||
| 60 | +## 代码搜索关键词 | ||
| 61 | + | ||
| 62 | +```bash | ||
| 63 | +grep -nE "tmpBuf|tmpLocal|accumulator|accLocal|rLengthAlign|FullLoad|ColSplit|RowSplit" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 64 | +``` | ||
| @@ -0,0 +1,69 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P111: Sort 两级归并 Tiling 系统化设计' | ||
| 4 | +description: 当 `N > tileSize × coreNum` 时，Sort / TopK 必须走 Pattern C：先在各核内部把多个有序 tile | ||
| 5 | + 归并为 1 路，再跨核归并。 | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- topic.compute.ordering | ||
| 9 | +- topic.memory | ||
| 10 | +- topic.implementation.pipeline | ||
| 11 | +- topic.programming_model.execution_unit.scalar | ||
| 12 | +status: stable | ||
| 13 | +sources: | ||
| 14 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P111_sort_two_level_merge_tiling.md | ||
| 15 | + role: primary | ||
| 16 | +created_at: '2026-10-09T00:00:00Z' | ||
| 17 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 18 | +platforms: | ||
| 19 | +- a2 | ||
| 20 | +--- | ||
| 21 | + | ||
| 22 | +# P111: Sort 两级归并 Tiling 系统化设计 | ||
| 23 | + | ||
| 24 | +## 核心思想 | ||
| 25 | + | ||
| 26 | +当 `N > tileSize × coreNum` 时，Sort / TopK 必须走 Pattern C：先在各核内部把多个有序 tile 归并为 1 路，再跨核归并。通过 workspace 双缓存、每轮归并后的 `SyncAll`、以及 Top-K 截断，将跨核同步次数降到 `ceil(log_M(coreNum)) + 1`。 | ||
| 27 | + | ||
| 28 | +## 代码骨架 | ||
| 29 | + | ||
| 30 | +```cpp | ||
| 31 | +__aicore__ inline void Process() | ||
| 32 | +{ | ||
| 33 | + // Phase 1: 各核并行排序 tile | ||
| 34 | + for (uint32_t t = 0; t < myTiles; ++t) { | ||
| 35 | + SortInSingleCore(curTileSize, globalOffset); | ||
| 36 | + } | ||
| 37 | + // Phase 2/3 | ||
| 38 | + pipe_.Reset(); | ||
| 39 | + InitMergeBuffers(); | ||
| 40 | + Phase2IntraCoreMerge(); | ||
| 41 | + SyncAll(); | ||
| 42 | + Phase3InterCoreMerge(); | ||
| 43 | + // Phase 4 | ||
| 44 | + pipe_.Reset(); | ||
| 45 | + InitOutputBuffers(); | ||
| 46 | + Phase4FinalMergeAndOutput(); | ||
| 47 | +} | ||
| 48 | +``` | ||
| 49 | + | ||
| 50 | +## 关键修改点 | ||
| 51 | + | ||
| 52 | +1. **tileSize 取 4096**：对齐到 32，且通常实测最优。 | ||
| 53 | +2. **四阶段划分**：Phase 1 排序；Phase 2 核内归并；Phase 3 跨核归并；Phase 4 Core 0 输出。 | ||
| 54 | +3. **workspace 双缓存**：`workSpaceFlag` 每轮翻转，输入输出分离。 | ||
| 55 | +4. **截断逻辑**：当 `currentElements × M ≥ K` 时，后续只读取前 K 个元素。 | ||
| 56 | +5. **SyncAll 时机**：Phase 2 结束一次；Phase 3 每轮一次；Phase 4 结束一次。 | ||
| 57 | + | ||
| 58 | +## 常见陷阱 | ||
| 59 | + | ||
| 60 | +- ❌ `truncationFlag` 在归并前设置 → 首次触发时数据不完整。 | ||
| 61 | +- ❌ Phase 3 每 group 后 SyncAll → 过度同步；应在每轮结束后同步。 | ||
| 62 | +- ❌ `UpdateSortInfo()` 未调用 → 死循环。 | ||
| 63 | +- ❌ Workspace 按总元素数而非核数分配 → 尾核越界。 | ||
| 64 | + | ||
| 65 | +## 代码搜索关键词 | ||
| 66 | + | ||
| 67 | +```bash | ||
| 68 | +grep -nE "Sort<|MrgSort|SyncAll|workSpaceFlag|truncationFlag|onceMaxElements|tileSize" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 69 | +``` | ||
| @@ -0,0 +1,62 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P112: Elementwise 向量化 Tile 尺寸设计' | ||
| 4 | +description: Elementwise 算子无跨元素依赖，应优先展平为 1D，并按 Vector 单元宽度与 UB 容量切分 tile。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- paradigm.elementwise | ||
| 8 | +- topic.data_movement.memory_access | ||
| 9 | +- topic.implementation.tiling | ||
| 10 | +status: stable | ||
| 11 | +sources: | ||
| 12 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P112_elementwise_vectorized_tile_sizing.md | ||
| 13 | + role: primary | ||
| 14 | +created_at: '2026-10-09T00:00:00Z' | ||
| 15 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 16 | +platforms: | ||
| 17 | +- a2 | ||
| 18 | +--- | ||
| 19 | + | ||
| 20 | +# P112: Elementwise 向量化 Tile 尺寸设计 | ||
| 21 | + | ||
| 22 | +## 核心思想 | ||
| 23 | + | ||
| 24 | +Elementwise 算子无跨元素依赖，应优先展平为 1D，并按 Vector 单元宽度与 UB 容量切分 tile。通过 `maxElemNum` 公式把 `bufferNum`、`maxDtypeBits`、`minDtypeBits` 纳入预算，同时保证多核对齐（512）与 UB 对齐（256B），实现 Vector 单元满载。 | ||
| 25 | + | ||
| 26 | +## 代码骨架 | ||
| 27 | + | ||
| 28 | +```cpp | ||
| 29 | +// 1D 处理骨架 | ||
| 30 | +for (uint32_t blockIdx = 0; blockIdx < blockNum; ++blockIdx) { | ||
| 31 | + uint32_t offset = blockIdx * blockFormer * sizeof(T); | ||
| 32 | + uint32_t curLen = (blockIdx == blockNum - 1) ? blockTail : blockFormer; | ||
| 33 | + CopyIn(inputUb, gmInput + offset, curLen); | ||
| 34 | + Compute(outputUb, inputUb, curLen); | ||
| 35 | + CopyOut(gmOutput + offset, outputUb, curLen); | ||
| 36 | +} | ||
| 37 | +``` | ||
| 38 | + | ||
| 39 | +## 关键修改点 | ||
| 40 | + | ||
| 41 | +1. **Shape 判定**：输入输出 shape 必须完全相同；否则走 Broadcast 路径。 | ||
| 42 | +2. **dtype 分支**：FP16/BF16 加减链路在 spec 未声明同量级时，应 Cast→FP32→Cast。 | ||
| 43 | +3. **maxElemNum 计算**： | ||
| 44 | + ``` | ||
| 45 | + maxElemNum = (ubSize - extraSize) * 8 / (bufferNum * maxDtypeBits) | ||
| 46 | + maxElemNum = floor_align(maxElemNum, 256 * 8 / minDtypeBits) | ||
| 47 | + ``` | ||
| 48 | +4. **多核切分**：把 `ubSplitAxis` 及外层轴展平为 `fusedProduct`，均分给 `coreNum`。 | ||
| 49 | +5. **核利用率不足**：循环缩小 `maxElemNum`（每次减 `CACHE_LINE`）重算，直到喂满更多核。 | ||
| 50 | + | ||
| 51 | +## 常见陷阱 | ||
| 52 | + | ||
| 53 | +- ❌ 把 broadcast 场景误当 elementwise 处理 → 结果 shape 错误。 | ||
| 54 | +- ❌ FP16 加减未升精度 → 大数吃小数导致精度不达标。 | ||
| 55 | +- ❌ `maxElemNum` 未按 dtype 位宽折算 → UB 溢出。 | ||
| 56 | +- ❌ 尾 block 未单独处理 → 越界或向量 mask 错误。 | ||
| 57 | + | ||
| 58 | +## 代码搜索关键词 | ||
| 59 | + | ||
| 60 | +```bash | ||
| 61 | +grep -nE "maxElemNum|ubFormer|blockFormer|blockTail|ELEM_ALIGN_FACTOR|CACHE_LINE" op_host/*_tiling.cpp op_kernel/*.cpp | ||
| 62 | +``` | ||
Aknowledge/contrib/ascendc/optimizations/evolution_strategies/cards/p113_broadcast_aligned_tiling.md+70-0
| @@ -0,0 +1,70 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P113: Broadcast 对齐 Tiling 设计' | ||
| 4 | +description: Broadcast 算子先通过 DimensionCollapse 把多轴问题降维，再判断是 OneDim（合轴后单维）还是多维广播。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.data_movement.memory_access | ||
| 8 | +- paradigm.elementwise | ||
| 9 | +- topic.memory | ||
| 10 | +- topic.implementation.tiling | ||
| 11 | +status: stable | ||
| 12 | +sources: | ||
| 13 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P113_broadcast_aligned_tiling.md | ||
| 14 | + role: primary | ||
| 15 | +created_at: '2026-10-09T00:00:00Z' | ||
| 16 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 17 | +platforms: | ||
| 18 | +- a2 | ||
| 19 | +--- | ||
| 20 | + | ||
| 21 | +# P113: Broadcast 对齐 Tiling 设计 | ||
| 22 | + | ||
| 23 | +## 核心思想 | ||
| 24 | + | ||
| 25 | +Broadcast 算子先通过 DimensionCollapse 把多轴问题降维，再判断是 OneDim（合轴后单维）还是多维广播。多维场景根据硬件版本与数据特征选择 UB Broadcast（BRC）或 NDDMA 硬件广播，并按 256B 对齐切分 UB，确保多核负载均衡。 | ||
| 26 | + | ||
| 27 | +## 代码骨架 | ||
| 28 | + | ||
| 29 | +```cpp | ||
| 30 | +// DimensionCollapse + stride 计算 | ||
| 31 | +for (int i = shapeLen - 1; i >= 0; --i) { | ||
| 32 | + flag[i] = computeBroadcastFlag(dims, i); | ||
| 33 | + if (i < shapeLen - 1 && flag[i] == flag[i+1]) { | ||
| 34 | + collapsedDims[i] *= collapsedDims[i+1]; | ||
| 35 | + // 合并轴 | ||
| 36 | + } | ||
| 37 | + strides[i] = (dims[j][i] == 1 && outDim > 1) ? 0 : runningStride; | ||
| 38 | + runningStride *= collapsedDims[i]; | ||
| 39 | +} | ||
| 40 | + | ||
| 41 | +// UB 切分：找到 ubSplitAxis | ||
| 42 | +for (int i = shapeLen - 1; i >= 0; --i) { | ||
| 43 | + curProduct *= dims[i]; | ||
| 44 | + if (curProduct > maxElemNum) { ubSplitAxis = i; break; } | ||
| 45 | +} | ||
| 46 | +``` | ||
| 47 | + | ||
| 48 | +## 关键修改点 | ||
| 49 | + | ||
| 50 | +1. **补维 + 合轴**：输入维度不足时左侧补 1；相邻同 flag 轴合并。 | ||
| 51 | +2. **广播轴 stride 置 0**，非广播轴按连续 stride 计算。 | ||
| 52 | +3. **广播方式选择**： | ||
| 53 | + - NLast + 尾轴大 → UB BRC | ||
| 54 | + - INT8/FP16/BF16 + 尾轴 32B 对齐 → UB BRC | ||
| 55 | + - 其他 → NDDMA | ||
| 56 | +4. **NDDMA >5 轴**：外层用 Kernel for-loop，最内 5 轴交给 `DataCopy<T,5,config>`。 | ||
| 57 | +5. **对齐**：OneDim 128B，多维 256B。 | ||
| 58 | + | ||
| 59 | +## 常见陷阱 | ||
| 60 | + | ||
| 61 | +- ❌ 未补维导致 flag 位图错位。 | ||
| 62 | +- ❌ 广播轴 stride 未置 0 → DMA 读到错误数据。 | ||
| 63 | +- ❌ NLast 场景误选 NDDMA → dcache 被反复冲刷。 | ||
| 64 | +- ❌ NDDMA 维度超过 5 未拆分 → 编译失败。 | ||
| 65 | + | ||
| 66 | +## 代码搜索关键词 | ||
| 67 | + | ||
| 68 | +```bash | ||
| 69 | +grep -nE "DimensionCollapse|broadcast|flag|stride|ubSplitAxis|NDDMA|BRC|DataCopy<.*5" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 70 | +``` | ||
| @@ -0,0 +1,62 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P114: Transpose 多维对齐 Tiling 设计' | ||
| 4 | +description: small-channel transpose 类算子先把问题统一建模为 `[C, N] → [N, C]`，其中 `C ≤ 16` 为小通道维，`N` | ||
| 5 | + 为展平长轴。 | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- topic.data_movement.layout_transform | ||
| 9 | +- topic.data_movement.memory_access | ||
| 10 | +- topic.implementation.tiling | ||
| 11 | +status: stable | ||
| 12 | +sources: | ||
| 13 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P114_transpose_nd_tiling_with_alignment.md | ||
| 14 | + role: primary | ||
| 15 | +created_at: '2026-10-09T00:00:00Z' | ||
| 16 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 17 | +platforms: | ||
| 18 | +- a2 | ||
| 19 | +--- | ||
| 20 | + | ||
| 21 | +# P114: Transpose 多维对齐 Tiling 设计 | ||
| 22 | + | ||
| 23 | +## 核心思想 | ||
| 24 | + | ||
| 25 | +small-channel transpose 类算子先把问题统一建模为 `[C, N] → [N, C]`，其中 `C ≤ 16` 为小通道维，`N` 为展平长轴。通过 `TransDataTo5HD` + `vnchwconv` + `Gather` 完成重排，并按 32 元素对齐 `tileNA`，同时满足 `repeats ≤ 255` 与 UB 预算约束。 | ||
| 26 | + | ||
| 27 | +## 代码骨架 | ||
| 28 | + | ||
| 29 | +```cpp | ||
| 30 | +// UB 预算与 tile 尺寸 | ||
| 31 | +uint32_t ubBudget = ubSize - reservedBytes; | ||
| 32 | +uint32_t perElemBytes = 16 * C + 32; | ||
| 33 | +uint32_t tileNMax = AlignDown(ubBudget / perElemBytes, 32); | ||
| 34 | +tileNMax = Min(tileNMax, 255 * 16); | ||
| 35 | + | ||
| 36 | +uint32_t tileN = AlignUp(CeilDiv(N, blockDim), 32); | ||
| 37 | +if (tileN > tileNMax) { /* 重算 tileN 使总 tile 数接近 blockDim 整数倍 */ } | ||
| 38 | + | ||
| 39 | +uint32_t tileNA = AlignUp(tileN, 32); | ||
| 40 | +uint32_t repeats = tileNA / 16; | ||
| 41 | +``` | ||
| 42 | + | ||
| 43 | +## 关键修改点 | ||
| 44 | + | ||
| 45 | +1. **统一建模**：把任意 small-channel transpose 整理为 `[C, N] → [N, C]`。 | ||
| 46 | +2. **对齐规则**：`tileNA = AlignUp(tileN, 32)`；`repeats = tileNA / 16 ≤ 255`。 | ||
| 47 | +3. **UB 预算**：`ubBytes = tileNA * (16 * C + 32)`，含双缓冲、中间 buffer、vnchwconv 输出、offset table。 | ||
| 48 | +4. **offset table**：按 `tileNA` 构建，每个 16-half block 前 `C` 个位置有效。 | ||
| 49 | +5. **多核切分**：`blockDim = min(coreNum, totalTiles)`，`tilesPerCore = ceil(totalTiles / blockDim)`。 | ||
| 50 | + | ||
| 51 | +## 常见陷阱 | ||
| 52 | + | ||
| 53 | +- ❌ 把小 C 轴与长轴展平顺序搞反 → layout 错误。 | ||
| 54 | +- ❌ `repeats` 超过 255 → 编译失败。 | ||
| 55 | +- ❌ offset table 按 `tileN` 而非 `tileNA` 构建 → 越界或 Gather 错位。 | ||
| 56 | +- ❌ 尾块未用 `curN` 区分有效长度 → 写入越界。 | ||
| 57 | + | ||
| 58 | +## 代码搜索关键词 | ||
| 59 | + | ||
| 60 | +```bash | ||
| 61 | +grep -nE "TransDataTo5HD|vnchwconv|Gather|tileNA|repeats|offsetBuff|SetValue" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 62 | +``` | ||
| @@ -0,0 +1,60 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P115: FlashAttention Tiling 路由' | ||
| 4 | +description: FlashAttention 类算子采用四层文档架构（L1 路由 → L2 地基 → L3 设计流 → L4 变体）。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.attention | ||
| 8 | +- topic.implementation.tiling | ||
| 9 | +- topic.programming_model.execution_unit.cube | ||
| 10 | +- topic.implementation.pipeline | ||
| 11 | +status: stable | ||
| 12 | +sources: | ||
| 13 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P115_flashattention_tiling_routing.md | ||
| 14 | + role: primary | ||
| 15 | +created_at: '2026-10-09T00:00:00Z' | ||
| 16 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 17 | +platforms: | ||
| 18 | +- a2 | ||
| 19 | +--- | ||
| 20 | + | ||
| 21 | +# P115: FlashAttention Tiling 路由 | ||
| 22 | + | ||
| 23 | +## 核心思想 | ||
| 24 | + | ||
| 25 | +FlashAttention 类算子采用四层文档架构（L1 路由 → L2 地基 → L3 设计流 → L4 变体）。设计时先选基础形态（GQA / MHA / MLA），再叠加正交 trait（量化 / 稀疏 / feature flags），然后按节点 G→N12 推进，确保 Roofline、多核切分、内存划分、流水编排一致。 | ||
| 26 | + | ||
| 27 | +## 代码骨架 | ||
| 28 | + | ||
| 29 | +```cpp | ||
| 30 | +// 示意：设计节点对应的关键决策映射 | ||
| 31 | +N1 宏观形态: kernel type / UB mode / Cube block / pipe depth | ||
| 32 | +N2 并行策略: task-level vs split-KV reduce | ||
| 33 | +N3 Roofline: mEff / AI 可行域 | ||
| 34 | +N4 Softmax: V1(max-trick) vs V2(online) | ||
| 35 | +N6 多核切分: task 维度 + 负载均衡 | ||
| 36 | +N7 内存划分: UB / L1 / L0 / GM | ||
| 37 | +N8 编译特化: constexpr / macro / runtime-if 三分 | ||
| 38 | +N10 流水编排: slot / SyncAll / HardEvent | ||
| 39 | +``` | ||
| 40 | + | ||
| 41 | +## 关键修改点 | ||
| 42 | + | ||
| 43 | +1. **基础形态三选一**：默认 GQA；`G=1` 退化到 MHA；含 latent absorption 走 MLA。 | ||
| 44 | +2. **正交 trait 叠加**：量化（MXFP8/INT8/AntiQuant）、稀疏、PSE/RoPE/Sink 等 feature flag 按需加载。 | ||
| 45 | +3. **按节点推进**：必须 G→N12 顺序执行，不可跳过。 | ||
| 46 | +4. **组合可行性校核**：未验证组合（如量化 × 稀疏）需先在 `ops/ascendc-tiling-design/references/flashattention/foundation/composition.md` 确认。 | ||
| 47 | +5. **常量语义分离**：跨核握手槽、L1 rotation 槽、V2 自读自写槽用不同常量。 | ||
| 48 | + | ||
| 49 | +## 常见陷阱 | ||
| 50 | + | ||
| 51 | +- ❌ 把 GQA / MHA / MLA 当作互斥子族 → 它们只是基础形态变化。 | ||
| 52 | +- ❌ 跳过 N3 Roofline 直接定 tile → 性能达不到理论极限。 | ||
| 53 | +- ❌ 未校核 trait 组合 → 量化 × 稀疏等组合可能未验证。 | ||
| 54 | +- ❌ 同一算子混用多种基础形态 → 违反架构约束。 | ||
| 55 | + | ||
| 56 | +## 代码搜索关键词 | ||
| 57 | + | ||
| 58 | +```bash | ||
| 59 | +grep -nE "GQA|MHA|MLA|kvHeadNum|gMaxPerTask|splitKV|FlashAttention|online_softmax|SetFlag|WaitFlag" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 60 | +``` | ||
Aknowledge/contrib/ascendc/optimizations/evolution_strategies/cards/p116_gqa_grouped_m_head_merge.md+78-0
| @@ -0,0 +1,78 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P116: GQA Grouped-M 头合并（G 个 Query Head 合并进 Matmul M 维）' | ||
| 4 | +description: GQA 中 `N_q = N_kv * G`，同一 KV head 被连续 G 个 query head 共享。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.attention | ||
| 8 | +- topic.programming_model.execution_unit.cube | ||
| 9 | +- topic.implementation.tiling | ||
| 10 | +status: stable | ||
| 11 | +sources: | ||
| 12 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P116_gqa_grouped_m_head_merge.md | ||
| 13 | + role: primary | ||
| 14 | +created_at: '2026-10-09T00:00:00Z' | ||
| 15 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 16 | +platforms: | ||
| 17 | +- a2 | ||
| 18 | +--- | ||
| 19 | + | ||
| 20 | +# P116: GQA Grouped-M 头合并（G 个 Query Head 合并进 Matmul M 维） | ||
| 21 | + | ||
| 22 | +## 核心思想 | ||
| 23 | + | ||
| 24 | +GQA 中 `N_q = N_kv * G`，同一 KV head 被连续 G 个 query head 共享。decode / 短 Sq（`Sq <= 4`）场景下，未合并实现按 `(b, h_q)` 逐头发射 `M = Sq`（甚至 M=1）的小 matmul——Cube 的 M 维利用率极低，且同一 KV head 被重复加载 G 次。 | ||
| 25 | + | ||
| 26 | +由于 BSH/BNSD 布局下同一 `(b, s)` 的 G 个 query head 在内存中相邻（行 stride = D），可以把 `[B, Sq, N_kv, G, D]` **零拷贝重解释**为 target-private BNSD view `[B*Sq, N_kv, 1, G, D]`：一个 cube task 的 M 维 = G，一次 matmul 算完同一 KV head 下的全部 G 个 query head，K/V 只加载一次。matmul 次数从 `B*Sq*N_q` 降到 `B*Sq*N_kv`，M 从 `Sq<=4` 提升到 `G`（常见 4~16）。 | ||
| 27 | + | ||
| 28 | +合并判定条件：`groupedQueryM = (sq <= 4 && G > 1)`。 | ||
| 29 | + | ||
| 30 | +## 代码骨架 | ||
| 31 | + | ||
| 32 | +```cpp | ||
| 33 | +// === Host tiling：判定 + 维度重映射 === | ||
| 34 | +const int64_t originalG = nq / nkv; | ||
| 35 | +// Sq<=4 时把 [B,Sq,Nkv,G,D] 重解释为 [B*Sq,Nkv,1,G,D]： | ||
| 36 | +// 一个 cube task 处理共享同一 K/V head 的全部 G 个 query head， | ||
| 37 | +// 而不是发射 G 个 M=Sq 的小 matmul。 | ||
| 38 | +const bool groupedQueryM = sq <= 4 && originalG > 1; | ||
| 39 | +const int64_t logicalB = groupedQueryM ? b * sq : b; | ||
| 40 | +const int64_t logicalSq = groupedQueryM ? originalG : sq; // M 维 = G | ||
| 41 | +const int64_t logicalG = groupedQueryM ? 1 : originalG; | ||
| 42 | +// view 内 query 行 stride = d（G 个头相邻）；非合并时保持原 nq*d | ||
| 43 | +const int64_t queryRowStride = groupedQueryM ? d : nq * d; | ||
| 44 | +input.set_layoutType(groupedQueryM ? LAYOUT_BNSD : LAYOUT_BSH); | ||
| 45 | +// mask 行数需随 M 维扩展：同一 (b,s) 的 G 个头 causal 语义相同，行内容重复 G 份 | ||
| 46 | +const int64_t maskRows = isCausal ? sq * (groupedQueryM ? originalG : 1) : 0; | ||
| 47 | + | ||
| 48 | +// === MatmulApiTiling：OrgShape 用重映射后的 stride === | ||
| 49 | +bmm1.SetShape(MinHost(s1Base, logicalSq), bmm1N, d); | ||
| 50 | +bmm1.SetOrgShape(logicalSq, s2Base * nRatio, queryRowStride, kvRowStride); | ||
| 51 | +``` | ||
| 52 | + | ||
| 53 | +## 关键修改点 | ||
| 54 | + | ||
| 55 | +1. **合并判定**：`sq <= 4 && G > 1` 才启用；Sq 较大时 M 本来就够大，合并反而限制并行度。 | ||
| 56 | +2. **维度重映射四元组**：`logicalB / logicalSq / logicalG / queryRowStride` 必须同时切换，漏一个就会读错地址。 | ||
| 57 | +3. **layout 切换**：合并后按 BNSD 语义取数（view 内 head 维即 M 维），K/V 偏移不变（G 维上 K/V 恒定）。 | ||
| 58 | +4. **mask 行扩展**：合并后 M 维每行对应一个 query head，causal mask 需从 `sq` 行扩到 `sq * G` 行（同一 s 下 G 行内容相同），与 P118 的 mask 预生成 kernel 配合。 | ||
| 59 | +5. **输出写回**：view 是纯内存重解释，输出 `[B*Sq, N_kv, 1, G, D]` 与原 `[B, Sq, N_q, D]` 逐字节同构，无需 transpose。 | ||
| 60 | +6. **与 P57 的分工**：P57（G 分核归约）解决"核不够分"的并行度问题，P116 解决"单 task M 太小"的 cube 效率问题；decode 场景先 P116 合并 M，仍吃不饱核时再叠加 P57 拆 G。 | ||
| 61 | + | ||
| 62 | +## 常见陷阱 | ||
| 63 | + | ||
| 64 | +- ❌ 非相邻布局强合并：TND / head-major 布局下 G 个头不连续（stride ≠ D），重解释会读到错误数据——先确认 `query[b, s, h_kv*G + g, :]` 对 g 连续。 | ||
| 65 | +- ❌ `Sq > 4` 仍合并：M 维已够大，合并减少 task 数反而造成核闲置。 | ||
| 66 | +- ❌ mask 行数忘记 ×G：合并后每行 M 都需要独立 mask 行，行数不足会读到越界/脏 mask。 | ||
| 67 | +- ❌ 只改 layout 不改 `SetOrgShape` 的行 stride：matmul 按旧 stride `nq*d` 取数，结果整体错位。 | ||
| 68 | + | ||
| 69 | +## 代码搜索关键词 | ||
| 70 | + | ||
| 71 | +```bash | ||
| 72 | +# 检测 GQA 结构（P116 适用信号） | ||
| 73 | +grep -nE "nq\s*/\s*nkv|headNum\s*/\s*kvHeadNum|gSize|group.*size" op_host/*.cpp op_kernel/*.h | ||
| 74 | +# 检测是否逐头发射小 matmul（未合并证据） | ||
| 75 | +grep -nE "for.*hidx|headsPerBlock|h_q|M\s*=\s*1\b" op_kernel/*.cpp | ||
| 76 | +# 检测是否已做 Grouped-M（已合并则不适用） | ||
| 77 | +grep -nE "groupedQueryM|logicalG|grouped.*[Mm]|G.*merge.*M" op_host/*.cpp op_kernel/*.cpp | ||
| 78 | +``` | ||
| @@ -0,0 +1,87 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P117: 双 BMM 链 MatmulApiTiling Host 配方（SetBufferSpace + SetFixSplit + depthA1 | ||
| 4 | + 流水微调）' | ||
| 5 | +description: Attention 类算子的 QK^T / PV 是两个 chained matmul。 | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- topic.attention | ||
| 9 | +- topic.programming_model.execution_unit.cube | ||
| 10 | +status: stable | ||
| 11 | +sources: | ||
| 12 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P117_matmul_api_tiling_host_recipe.md | ||
| 13 | + role: primary | ||
| 14 | +created_at: '2026-10-09T00:00:00Z' | ||
| 15 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 16 | +platforms: | ||
| 17 | +- a2 | ||
| 18 | +--- | ||
| 19 | + | ||
| 20 | +# P117: 双 BMM 链 MatmulApiTiling Host 配方（SetBufferSpace + SetFixSplit + depthA1 流水微调） | ||
| 21 | + | ||
| 22 | +## 核心思想 | ||
| 23 | + | ||
| 24 | +Attention 类算子的 QK^T / PV 是两个 chained matmul。手写 L1/L0 buffer 管理与 K 维切分（Nd2Nz/LoadData/Mmad/Fixpipe 全套）既易错又拿不到框架的自动流水优化；应在 **host 端用 `matmul_tiling::MatmulApiTiling` 分别为 bmm1/bmm2 生成 TCubeTiling**，kernel 侧交给 Matmul 高阶 API 消费（P46）。 | ||
| 25 | + | ||
| 26 | +三个容易漏掉的配方细节： | ||
| 27 | + | ||
| 28 | +1. **`SetBufferSpace(L1, L0C)`**：显式告知真实片上容量，tiling 才能打满 L1 多缓冲； | ||
| 29 | +2. **`d ∈ (64, 128]` 时 `SetFixSplit(baseM=128, baseN=128, baseK=alignUp(d,16))`**：固定 base block 避免默认切分在典型 head_dim 上产生碎块； | ||
| 30 | +3. **depthA1 2→3 微调**：当 `baseK >= singleCoreK`（一个 A tile 已覆盖 K）且 `depthA1 == 2` 时改为 3，让第三级 A1 流水与当前 cube 计算重叠。默认 `GetTiling` 不做此调整，必须手动改。 | ||
| 31 | + | ||
| 32 | +## 代码骨架 | ||
| 33 | + | ||
| 34 | +```cpp | ||
| 35 | +// === Host：为 bmm1 (QK^T) 生成 tiling === | ||
| 36 | +matmul_tiling::MatmulApiTiling bmm1; | ||
| 37 | +bmm1.SetAType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, inputDtype, false); | ||
| 38 | +bmm1.SetBType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, inputDtype, true); // K 转置 | ||
| 39 | +bmm1.SetCType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, | ||
| 40 | + matmul_tiling::DataType::DT_FLOAT); // fp32 累加输出 | ||
| 41 | +bmm1.SetShape(MinHost(s1Base, sq), bmm1N, d); // single: (M, N, K) | ||
| 42 | +bmm1.SetOrgShape(sq, s2Base * nRatio, queryRowStride, kvRowStride); // org + 真实行 stride | ||
| 43 | +bmm1.SetBias(false); | ||
| 44 | +if (bmm1.SetBufferSpace(l1Size, l0cSize) != 0) return -1; // 显式片上容量 | ||
| 45 | +if (d > 64 && d <= 128) { // 典型 head_dim 固定切分 | ||
| 46 | + bmm1.SetFixSplit(MinHost(128, AlignUp(sq, 16)), | ||
| 47 | + MinHost(128, AlignUp(skv, 16)), AlignUp(d, 16)); | ||
| 48 | +} | ||
| 49 | +bmm1.GetTiling(tiling.bmm1TilingData); | ||
| 50 | +// depthA1 2→3：一个 A tile 覆盖 K 时，第三级 A1 流水与 cube 计算重叠 | ||
| 51 | +if (tiling.bmm1TilingData.baseK >= tiling.bmm1TilingData.singleCoreK && | ||
| 52 | + tiling.bmm1TilingData.depthA1 == 2) { | ||
| 53 | + tiling.bmm1TilingData.depthA1 = 3; | ||
| 54 | +} | ||
| 55 | +tiling.bmm1TilingData.shareMode = 0; | ||
| 56 | +tiling.bmm1TilingData.shareL1Size = l1Size; | ||
| 57 | +tiling.bmm1TilingData.shareL0CSize = l0cSize; | ||
| 58 | + | ||
| 59 | +// bmm2 (PV)：A=P(fp32→cast 回输入 dtype)、B=V 不转置，SetShape(min(s1Base,sq), d, bmm1N) | ||
| 60 | +``` | ||
| 61 | + | ||
| 62 | +## 关键修改点 | ||
| 63 | + | ||
| 64 | +1. **bmm1/bmm2 分别建 tiling**：形状、转置标记、OrgShape stride 都不同，不可复用同一个 MatmulApiTiling 对象的结果。 | ||
| 65 | +2. **SetOrgShape 的行 stride 必须按真实 layout**（BSH 是 `nq*d`，grouped-M view 是 `d`，见 P116），stride 错 = 整体错位。 | ||
| 66 | +3. **CType 用 `DT_FLOAT`**：matmul 输出 fp32 累加，softmax/归一化后再 cast 回输入 dtype（配合 A1）。 | ||
| 67 | +4. **SetFixSplit 只在 `d ∈ (64, 128]`** 收益明确；d≤64 默认切分已足够，d>128 走 dOuter 循环。 | ||
| 68 | +5. **depthA1 微调条件必须同时满足** `baseK >= singleCoreK && depthA1 == 2`，乱改会破坏流水假设。 | ||
| 69 | +6. **TilingData GM→local 拷贝用 8 字节对齐字传输**（`uint64_t` 数组 + 尾部字节补齐），不要逐字节发 GM 事务——direct-launch 场景每次 kernel 启动都要读 tiling，逐字节拷贝是启动开销大头。 | ||
| 70 | + | ||
| 71 | +## 常见陷阱 | ||
| 72 | + | ||
| 73 | +- ❌ bmm1 的 BType 忘记转置标记：QK^T 是 `Q @ K^T`，K 必须 `transpose=true`，否则 shape 对不上或结果错误。 | ||
| 74 | +- ❌ `SetBufferSpace` 失败不检查返回值：L1/L0C 超配时 GetTiling 产出不可用 tiling，必须显式 fallback（缩小 s1Base/s2Base 重试）。 | ||
| 75 | +- ❌ depthA1 无条件改 3：`baseK < singleCoreK`（K 需多轮循环）时改 3 会读越界 A1 buffer。 | ||
| 76 | +- ❌ 两个 bmm 共享一份 TCubeTiling：singleCoreN/baseK 不同，kernel 侧必挂。 | ||
| 77 | + | ||
| 78 | +## 代码搜索关键词 | ||
| 79 | + | ||
| 80 | +```bash | ||
| 81 | +# 检测双 BMM 链（P117 适用信号） | ||
| 82 | +grep -nE "bmm1|bmm2|qk|pv|matmul" op_host/*_tiling.cpp op_kernel/*.cpp | ||
| 83 | +# 检测是否已用 MatmulApiTiling（已用则转向 P86 调优） | ||
| 84 | +grep -nE "MatmulApiTiling|SetBufferSpace|SetFixSplit|GetTiling" op_host/*.cpp | ||
| 85 | +# 检测手写 cube 搬运（未用高阶 API 的证据） | ||
| 86 | +grep -nE "Nd2Nz|LoadData|Mmad|Fixpipe" op_kernel/*.cpp | ||
| 87 | +``` | ||
| @@ -0,0 +1,85 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P118: Causal Mask 独立 AIV Kernel 并行预生成（按需 Dense 替代通用压缩表）' | ||
| 4 | +description: Causal mask 不该占用 attention 主路径的计算/搬运时间。 | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.attention | ||
| 8 | +- topic.data_movement.memory_access | ||
| 9 | +- topic.implementation.pipeline | ||
| 10 | +status: stable | ||
| 11 | +sources: | ||
| 12 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P118_mask_precompute_parallel_aiv.md | ||
| 13 | + role: primary | ||
| 14 | +created_at: '2026-10-09T00:00:00Z' | ||
| 15 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 16 | +platforms: | ||
| 17 | +- a2 | ||
| 18 | +--- | ||
| 19 | + | ||
| 20 | +# P118: Causal Mask 独立 AIV Kernel 并行预生成（按需 Dense 替代通用压缩表） | ||
| 21 | + | ||
| 22 | +## 核心思想 | ||
| 23 | + | ||
| 24 | +Causal mask 不该占用 attention 主路径的计算/搬运时间。两个正交要点： | ||
| 25 | + | ||
| 26 | +1. **按需 dense 替代通用压缩表**：只生成当前 shape 需要的 `s1 × s2` dense right-down mask（byte 级 0/1），而不是预生成一张通用 2048×2048 压缩表再解压——短序列场景下压缩表的生成 + 读取开销占主导地位，dense 直达更省。 | ||
| 27 | +2. **mask 生成下沉独立 vector kernel**：单独的轻量 kernel 用**全部 AIV 核**并行生成 mask 到 GM workspace，attention cube kernel（AIC）直接 DataCopy 消费。mask kernel 计算量极小（每行一次 Duplicate + 边界 word 修正），与主 kernel 串行 launch 也只占零头；异构分工（AIV 生成 / AIC 消费）避免 mask 逻辑挤占 cube kernel 的 UB 与流水。 | ||
| 28 | + | ||
| 29 | +行生成用 word 级技巧：每行先 `Duplicate(0x01010101)` 填满（全屏蔽），再对允许区间的前 `allowed/4` 个 word 写 0（全保留），边界 word 按余数写 `0x01010100 / 0x01010000 / 0x01000000`——纯 vector 两次 Duplicate + 一次 SetValue，无逐元素循环。 | ||
| 30 | + | ||
| 31 | +## 代码骨架 | ||
| 32 | + | ||
| 33 | +```cpp | ||
| 34 | +// === 独立 mask kernel（全 AIV 并行，行 stride 分配）=== | ||
| 35 | +__global__ __aicore__ void causal_mask_kernel( | ||
| 36 | + GM_ADDR mask, int64_t rows, int64_t columns, int64_t groupSize, int64_t numBlocks) | ||
| 37 | +{ | ||
| 38 | + // 每行 wordsPerRow 个 int32，byte 语义：0=保留，1=屏蔽 | ||
| 39 | + const int64_t wordsPerRow = (columns + 3) / 4; | ||
| 40 | + for (int64_t row = GetBlockIdx(); row < rows; row += numBlocks) { | ||
| 41 | + auto local = outQueue.AllocTensor<int32_t>(); | ||
| 42 | + Duplicate(local, 0x01010101, wordsPerRow); // 默认全屏蔽 | ||
| 43 | + const int64_t queryRow = row / groupSize; // GQA 合并 M 时 G 行共享 | ||
| 44 | + int64_t allowed = queryRow + (columns - queryRows) + 1; // right-down 对齐 | ||
| 45 | + allowed = min(max(allowed, 0), columns); | ||
| 46 | + if (allowed / 4 > 0) Duplicate(local, 0, allowed / 4); // 全保留段 | ||
| 47 | + PipeBarrier<PIPE_V>(); | ||
| 48 | + const int32_t boundary[3] = {0x01010100, 0x01010000, 0x01000000}; | ||
| 49 | + if ((allowed & 3) != 0) local.SetValue(allowed / 4, boundary[(allowed & 3) - 1]); | ||
| 50 | + // EnQue → DeQue → DataCopyPad 写出 | ||
| 51 | + } | ||
| 52 | +} | ||
| 53 | + | ||
| 54 | +// === Host tiling：NO_COMPRESS + 行字节数对齐 4 === | ||
| 55 | +input.set_attenMaskCompressMode(NO_COMPRESS_MODE); // 禁用通用 2048x2048 压缩表 | ||
| 56 | +input.set_attenMaskS2Size(AlignUp(skv, 4)); // 行字节数 4B 对齐（word 语义） | ||
| 57 | +maskRows = isCausal ? sq * maskGroupSize : 0; // 配合 P116：合并 M 后 ×G | ||
| 58 | +``` | ||
| 59 | + | ||
| 60 | +## 关键修改点 | ||
| 61 | + | ||
| 62 | +1. **mask 语义先确认**：本配方 `0=保留 / 1=屏蔽`；若 kernel 消费端语义相反，Duplicate 初值与边界字全部取反。 | ||
| 63 | +2. **right-down 对齐**：`allowed = queryRow + (skv - sq) + 1`，Sq≠Skv 时因果边界整体右移，别写成左上对齐。 | ||
| 64 | +3. **groupSize 行共享**：GQA grouped-M（P116）后 mask 行数 = `sq * G`，同一 query row 的 G 个 mask 行内容相同（`queryRow = row / groupSize`）。 | ||
| 65 | +4. **压缩模式关闭**：`attenMaskCompressMode = NO_COMPRESS_MODE`，并保证消费端按 dense 行读取（`attenMaskS2Size = AlignUp(skv, 4)`）。 | ||
| 66 | +5. **launch 时序**：mask kernel 与 attention kernel 同 stream 串行即安全；跨 stream 并行需显式 event 同步。 | ||
| 67 | +6. **超长 skv 分行**：单行 word 数受 UB buffer 限制（参考常量 `kMaxMaskColumns = 8192`），超出需按行分段生成。 | ||
| 68 | + | ||
| 69 | +## 常见陷阱 | ||
| 70 | + | ||
| 71 | +- ❌ 继续生成 2048×2048 通用压缩表：短序列（skv≪2048）下表生成 + 解压读取 dominated，性能反噬。 | ||
| 72 | +- ❌ mask kernel 用单核生成：行数 `sq*G` 可达数千，必须 `row += numBlocks` 全核分摊。 | ||
| 73 | +- ❌ 边界 word 余数表写错：`allowed & 3` ∈ {1,2,3} 分别对应 `0x01010100 / 0x01010000 / 0x01000000`，顺序颠倒会使因果边界偏移 1~2 列，精度挂得很难查。 | ||
| 74 | +- ❌ mask kernel 与 attention 跨 stream 无同步：attention 读到未写完的 mask。 | ||
| 75 | + | ||
| 76 | +## 代码搜索关键词 | ||
| 77 | + | ||
| 78 | +```bash | ||
| 79 | +# 检测 causal mask 现状（P118 适用信号） | ||
| 80 | +grep -nE "causal|attenMask|maskCompress|2048" op_host/*.cpp op_kernel/*.cpp | ||
| 81 | +# 检测 mask 是否在主 kernel 内逐元素生成（下沉信号） | ||
| 82 | +grep -nE "if.*causal.*continue|mask.*SetValue|ComputeMask|CreateMask" op_kernel/*.cpp | ||
| 83 | +# 检测是否已独立 mask kernel（已下沉则不适用） | ||
| 84 | +grep -nE "mask_kernel|MaskKernel|maskOffset|launch.*mask" op_kernel/*.cpp op_host/*.cpp | ||
| 85 | +``` | ||
| @@ -0,0 +1,91 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P119: Two-Pass Softmax → Online Softmax 单遍 KV 流式改写' | ||
| 4 | +description: 两遍 softmax——pass1 扫全部 KV 求行 max，pass2 再扫一遍算 `exp(s - max)` 并累加分子分母——把 | ||
| 5 | + **K 从 GM 读两遍、V 读一遍**，MTE2 流量接近 2×，且两遍之间无任何流水重叠。 | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- topic.attention | ||
| 9 | +- topic.compute.normalization | ||
| 10 | +- topic.data_movement.memory_access | ||
| 11 | +- topic.implementation.pipeline | ||
| 12 | +status: stable | ||
| 13 | +sources: | ||
| 14 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P119_two_pass_to_online_softmax.md | ||
| 15 | + role: primary | ||
| 16 | +created_at: '2026-10-09T00:00:00Z' | ||
| 17 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 18 | +platforms: | ||
| 19 | +- a2 | ||
| 20 | +--- | ||
| 21 | + | ||
| 22 | +# P119: Two-Pass Softmax → Online Softmax 单遍 KV 流式改写 | ||
| 23 | + | ||
| 24 | +## 核心思想 | ||
| 25 | + | ||
| 26 | +两遍 softmax——pass1 扫全部 KV 求行 max，pass2 再扫一遍算 `exp(s - max)` 并累加分子分母——把 **K 从 GM 读两遍、V 读一遍**，MTE2 流量接近 2×，且两遍之间无任何流水重叠。online softmax（flash 范式）把 max/sum/acc 三个状态改为 running 更新，**K/V 各只读一遍**： | ||
| 27 | + | ||
| 28 | +``` | ||
| 29 | +每块 S2 tile： | ||
| 30 | + s = QK^T * scale (cube) | ||
| 31 | + m_new = max(m_old, rowmax(s)) | ||
| 32 | + p = exp(s - m_new) | ||
| 33 | + alpha = exp(m_old - m_new) // rescale 因子 | ||
| 34 | + l = l * alpha + rowsum(p) // 分母 running sum | ||
| 35 | + acc = acc * alpha + p @ V // 分子 running 累加（cube） | ||
| 36 | + m_old = m_new | ||
| 37 | +收尾：out = acc / l | ||
| 38 | +``` | ||
| 39 | + | ||
| 40 | +GM 流量减半只是账面收益；真正的质变是 **QK^T(cube) → softmax update(vector) → PV(cube) 在同一 S2 循环内首尾相接**，配合状态 buffer 常驻（P38）后 cube/vector 流水可以持续重叠。 | ||
| 41 | + | ||
| 42 | +精度说明：fp32 累加 + max-trick 下 online 更新本身数值稳定（A5）。改写后**不要**保留逐 token Kahan 补偿——它每 token 引入 5+ 次额外 vector 往返，会打碎流水。 | ||
| 43 | + | ||
| 44 | +## 代码骨架 | ||
| 45 | + | ||
| 46 | +```cpp | ||
| 47 | +// === 改造前：两遍，K 读两次 === | ||
| 48 | +// pass1: for j in S_kv { load K_j; dot; rowmax } // K 第 1 遍 | ||
| 49 | +// pass2: for j in S_kv { load K_j; dot; exp; load V_j; acc += p*V_j } // K 第 2 遍 + V | ||
| 50 | + | ||
| 51 | +// === 改造后：单遍 online 更新 === | ||
| 52 | +Duplicate(rowMax, -INF, s1); Duplicate(rowSum, 0, s1); Duplicate(acc, 0, s1 * d); | ||
| 53 | +for (int64_t s2 = 0; s2 < skv; s2 += s2Base) { // K/V 各读一次 | ||
| 54 | + // 1) cube: sTile = Q @ K[s2:s2+s2Base]^T * scale | ||
| 55 | + // 2) vector: tileMax = RowMax(sTile); newMax = Max(rowMax, tileMax) | ||
| 56 | + // 3) vector: p = Exp(sTile - newMax); alpha = Exp(rowMax - newMax) | ||
| 57 | + // 4) vector: rowSum = rowSum * alpha + RowSum(p) | ||
| 58 | + // 5) vector: acc *= alpha （broadcast 到 d 维） | ||
| 59 | + // 6) cube: acc += p_cast @ V[s2:s2+s2Base] | ||
| 60 | + // 7) rowMax = newMax | ||
| 61 | +} | ||
| 62 | +// 收尾：out = acc / rowSum（rowSum==0 的全 mask 行输出 0） | ||
| 63 | +``` | ||
| 64 | + | ||
| 65 | +## 关键修改点 | ||
| 66 | + | ||
| 67 | +1. **删掉 pass1**：行 max 不再预扫描，改为 tile 级 running max；K 的 GM 读取次数 2→1。 | ||
| 68 | +2. **三状态常驻 UB**：rowMax/rowSum/acc 跨 S2 循环常驻（P38 的 buffer 组织方式），acc 用 fp32。 | ||
| 69 | +3. **rescale 顺序**：先算 `alpha = exp(m_old - m_new)`，**先 rescale rowSum/acc 再累加当前 tile**，顺序反了结果错。 | ||
| 70 | +4. **p 的 dtype**：PV matmul 前把 p cast 回输入 dtype（fp16/bf16），分子累加仍在 fp32 acc 中。 | ||
| 71 | +5. **全 mask 行兜底**：causal 下存在全屏蔽行（m 保持 -INF、l=0），归一化时 `l == 0 → out = 0`，不要除零。 | ||
| 72 | +6. **causal 逐 tile 裁剪**：visible 范围按 tile 边界判断，配合 P118 的预生成 dense mask 直接乘加，不要在 token 级 `continue`。 | ||
| 73 | +7. **与 cube 链接轨**：单遍结构是 P116/P117 的前置——逐 token 两遍结构无法有效利用 matmul tile。 | ||
| 74 | + | ||
| 75 | +## 常见陷阱 | ||
| 76 | + | ||
| 77 | +- ❌ rescale 与累加顺序颠倒：必须 `acc = acc * alpha + p @ V`，写成 `acc = (acc + p @ V) * alpha` 全错。 | ||
| 78 | +- ❌ `exp(m_old - m_new)` 中 m_old 初值：必须 -INF（`exp(-INF - m_new) = 0`），用 0 初始化第一块会把 acc 缩放错。 | ||
| 79 | +- ❌ 保留逐 token Kahan：online + fp32 已够，Kahan 的每 token 5 次 vector 往返会把流水打碎。 | ||
| 80 | +- ❌ 归一化不防零：全 mask 行 `l == 0` 产生 NaN。 | ||
| 81 | +- ❌ tile 内先减 tileMax 再合并：p 必须减全局 running max（m_new），只减 tile max 会溢出。 | ||
| 82 | + | ||
| 83 | +## 代码搜索关键词 | ||
| 84 | + | ||
| 85 | +```bash | ||
| 86 | +# 检测两遍结构（P119 适用信号） | ||
| 87 | +grep -nE "pass|PASS|for.*S_kv|for.*skv" op_kernel/*.cpp | ||
| 88 | +grep -nE "Exp\\s*\\(|ReduceSum|WholeReduceSum|BlockReduceMax" op_kernel/*.cpp | ||
| 89 | +# 检测是否已是 online（已改写则不适用） | ||
| 90 | +grep -nE "SoftmaxFlash|online|rescale|rowMax|running" op_kernel/*.cpp | ||
| 91 | +``` | ||
| @@ -0,0 +1,66 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P11: Tail Block Handling (尾块与边界处理)' | ||
| 4 | +description: 参考实现精心设计了尾块处理机制以解决数据量不能被核心数整除的问题：1) 双轨制计算：同时计算正常核（normCoreHandleDefaultValues/normCoreHandleSparses）和尾核（tailCoreHandleDefaultValues/tailCoreHandleSpar… | ||
| 5 | +tags: | ||
| 6 | +- task.optimization | ||
| 7 | +- topic.compute.indexing | ||
| 8 | +- topic.implementation.tiling | ||
| 9 | +status: stable | ||
| 10 | +sources: | ||
| 11 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P11_tail_block_handling.md | ||
| 12 | + role: primary | ||
| 13 | +created_at: '2026-10-09T00:00:00Z' | ||
| 14 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 15 | +platforms: | ||
| 16 | +- a2 | ||
| 17 | +--- | ||
| 18 | + | ||
| 19 | +# P11: Tail Block Handling (尾块与边界处理) | ||
| 20 | + | ||
| 21 | +## 核心思想 | ||
| 22 | +参考实现精心设计了尾块处理机制以解决数据量不能被核心数整除的问题：1) 双轨制计算：同时计算正常核（normCoreHandleDefaultValues/normCoreHandleSparses）和尾核（tailCoreHandleDefaultValues/tailCoreHandleSparses）的处理数据量；2) 循环展开：正常循环使用defaultUbFactor批量处理，尾循环单独处理剩余数据；3) Loop/Tail分离：Host端计算normBlockLoop/tailBlockLoop和normBlockTailLoopSize/tailBlockTailLoopSize，Kernel端根据blockIdx_判断自己是正常核还是尾核。这种设计确保了在所有场景下都能实现近完美的负载均衡，避免了因数据量不是核心数倍而导致的性能损失。 | ||
| 23 | + | ||
| 24 | +## 代码骨架 | ||
| 25 | + | ||
| 26 | +// === 改造前（基线）=== | ||
| 27 | +```cpp | ||
| 28 | +// 参考实现: 固定tile处理 | ||
| 29 | +for (uint32_t i = 0; i < this->innerLoops; i++) { | ||
| 30 | + CopyIn(i); | ||
| 31 | + Compute(i); | ||
| 32 | + CopyOut(i); | ||
| 33 | +} | ||
| 34 | +``` | ||
| 35 | + | ||
| 36 | +// === 改造后（优化实现）=== | ||
| 37 | +```cpp | ||
| 38 | +// 参考实现: 余数处理 | ||
| 39 | +for (uint32_t i = 0; i < copyTimes; i++) { | ||
| 40 | + bool isRemainder = (i == copyTimes - 1 && copyTimesRemainder > 0); | ||
| 41 | + uint32_t tempDataCount = isRemainder ? copyTimesRemainder : Base::maxDataCount; | ||
| 42 | + if (isRemainder) { | ||
| 43 | + DataCopyExtParams copyParams{1, static_cast<uint32_t>(dataCount * sizeof(T)), 0, 0, 0}; | ||
| 44 | + DataCopyPadExtParams<T> padParams{false, 0, 0, 0}; | ||
| 45 | + DataCopyPad(dataLocal, inTensorsGM[...], copyParams, padParams); | ||
| 46 | + } else { | ||
| 47 | + DataCopy(dataLocal, inTensorsGM[...], dataCount); | ||
| 48 | + } | ||
| 49 | +} | ||
| 50 | +``` | ||
| 51 | + | ||
| 52 | +## 关键修改点 | ||
| 53 | + | ||
| 54 | +1. 预期收益: 保证数据安全性，避免越界访问，同时主循环保持高效; 确保所有边界情况下的计算正确性 | ||
| 55 | + | ||
| 56 | +## 常见陷阱 | ||
| 57 | + | ||
| 58 | +⚠️ AtomicAdd操作有一定性能开销 | ||
| 59 | +⚠️ 需要使用额外的tmpPattern buffer | ||
| 60 | +⚠️ 需要预计算和存储kernel索引，增加UB使用 | ||
| 61 | + | ||
| 62 | +## 代码搜索关键词 | ||
| 63 | + | ||
| 64 | +```bash | ||
| 65 | +grep -n "BUFFER_NUM\|InitBuffer\|TQue\|tileSize\|ubFactor\|Tiling\|BLOCK_DIM\|GetBlockNum" op_kernel/*.cpp op_host/*_tiling.cpp | ||
| 66 | +``` | ||
Aknowledge/contrib/ascendc/optimizations/evolution_strategies/cards/p120_fa_vec_dot_to_cube_mmad.md+82-0
| @@ -0,0 +1,82 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P120: FA 向量点积模拟 Matmul → Cube mmad 重写（QK^T / PV 上 Cube）' | ||
| 4 | +description: kernel 内层循环用 `Mul + ReduceSum` 逐 `(qi, kj)` 对算 QK^T 点积、用标量 `Muls + Add` | ||
| 5 | + 逐元素累加 PV、逐元素 `GetValue` 读 score——这是用 Vector 单元模拟矩阵运算，性能在架构层面被封顶（Cube 的 fp16 算力… | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- topic.attention | ||
| 9 | +- topic.programming_model.execution_unit.cube | ||
| 10 | +- topic.programming_model.execution_unit.scalar | ||
| 11 | +status: stable | ||
| 12 | +sources: | ||
| 13 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P120_fa_vec_dot_to_cube_mmad.md | ||
| 14 | + role: primary | ||
| 15 | +created_at: '2026-10-09T00:00:00Z' | ||
| 16 | +updated_at: '2026-10-13T00:00:00Z' | ||
| 17 | +platforms: | ||
| 18 | +- a2 | ||
| 19 | +--- | ||
| 20 | + | ||
| 21 | +# P120: FA 向量点积模拟 Matmul → Cube mmad 重写（QK^T / PV 上 Cube） | ||
| 22 | + | ||
| 23 | +## 核心思想 | ||
| 24 | + | ||
| 25 | +kernel 内层循环用 `Mul + ReduceSum` 逐 `(qi, kj)` 对算 QK^T 点积、用标量 `Muls + Add` 逐元素累加 PV、逐元素 `GetValue` 读 score——这是用 Vector 单元模拟矩阵运算，性能在架构层面被封顶（Cube 的 fp16 算力比 Vector 高一个数量级以上）。若同时把 Q/K/V 先 `Cast` 成 fp32 再算，UB 流量还要翻倍。 | ||
| 26 | + | ||
| 27 | +重写为 Cube mmad 的四个关键决策： | ||
| 28 | + | ||
| 29 | +1. **输入 dtype 直接进 Cube**：Q/K/V 保持 fp16/bf16 进 mmad，**不要先 Cast fp32**；fp32 只出现在 L0C 输出 S、online softmax 状态和 O 累加链。 | ||
| 30 | +2. **K 按 ColumnMajor 读取 = 免转置 K^T**：QK^T 的 B 矩阵就是 K 的转置。GM 上 K 以行主序存储，按列主序视图读取，经 `Nd2NzParams` 一次 DMA 转为 nZ 分形进 L1，L1→L0B `ifTranspose=false` 直送——零显式 transpose。 | ||
| 31 | +3. **Q preload 常驻 L1**：每个 task 开头一次性把 Q tile（128 行）搬入 L1 并常驻，KV 循环只搬 K/V；L0A/L0B/L0C 全部双缓冲。 | ||
| 32 | +4. **P 回写输入 dtype 喂第二级 mmad**：online softmax 输出的 P 从 fp32 `DownCast` 回 fp16/bf16（bf16 用 `CAST_RINT`、fp16 用 `CAST_NONE`）写 GM workspace，PV mmad 直接消费，O 以 fp32 在 L0C/GM 累加。 | ||
| 33 | + | ||
| 34 | +与 P102 的关系：P102 是"标量/向量模拟矩阵运算 → MIX CV 融合重写"的**通用判定卡**；本卡是判定通过后在 **attention 场景的具体落地配方**（layout 选型、分形搬运、精度链路、tile 基准值）。 | ||
| 35 | + | ||
| 36 | +## 代码骨架 | ||
| 37 | + | ||
| 38 | +```cpp | ||
| 39 | +// === 改造前（向量点积 + 全量 fp32 Cast 形态）=== | ||
| 40 | +for (int32_t qi = 0; qi < qr; qi++) | ||
| 41 | + for (int32_t kj = 0; kj < kr; kj++) { | ||
| 42 | + Mul(rw, qf[qi*Df], kf[kj*Df], Df); // 逐对向量乘 | ||
| 43 | + ReduceSum(sp[so+kj*P8], rw, rt, Df); // 逐对归约 = 一个 score | ||
| 44 | + } | ||
| 45 | +// PV：标量读 p，逐元素 Muls+Add 累加（同样 Vector 模拟 matmul） | ||
| 46 | + | ||
| 47 | +// === 改造后（Cube mmad + 分形 layout）=== | ||
| 48 | +using LayoutQ = layout::RowMajor; // A=Q: GM→L1 nd2nz→zN | ||
| 49 | +using LayoutK = layout::ColumnMajor; // B=K: 列主序读取→nZ，天然 K^T | ||
| 50 | +using L1TileShapeQK = GemmShape<128, 128, 128>; // L1 tile: M/N/K | ||
| 51 | +using L0TileShapeQK = GemmShape<128, 128, 128>; // L0 tile | ||
| 52 | +using BlockMmadQK = Gemm::Block::BlockMmad<MmadAtlasA2FAIQK<...>, | ||
| 53 | + L1TileShapeQK, L0TileShapeQK, QType, KType, SType>; // S=fp32 | ||
| 54 | +using L1TileShapePV = GemmShape<128, 128, 256>; // PV 的 K 维 = KV stack 256 | ||
| 55 | +using BlockMmadPV = Gemm::Block::BlockMmad<MmadAtlasA2FAIPV<...>, | ||
| 56 | + L1TileShapePV, L0TileShapePV, PType, VType, OTmpType>; // OTmp=fp32 | ||
| 57 | + | ||
| 58 | +blockMmadQK.loadQGM(gQ[gmOffsetQ], layoutQ, rowNum, qNBlockSize, qHeads); // Q 常驻 L1 | ||
| 59 | +// KV 循环内：blockMmadQK(...) 只搬 K；softmax(AIV)；blockMmadPV(...) 消费 P(fp16) | ||
| 60 | +``` | ||
| 61 | + | ||
| 62 | +## 关键修改点 | ||
| 63 | + | ||
| 64 | +1. **删除全部逐对 `Mul+ReduceSum` / 标量 `Muls+Add` 矩阵模拟代码**，替换为 mmad 调用；score 矩阵 S 由 L0C 以 fp32 fixpipe 到 GM workspace。 | ||
| 65 | +2. **精度链路重排**：Q/K/V 进 cube 前不做 fp32 Cast；softmax 在 fp32 的 S 上做；P downcast 回输入 dtype；O 累加 fp32，收尾一次 Cast 输出。 | ||
| 66 | +3. **L1/L0 双缓冲与事件流水**：GM→L1（MTE2）、L1→L0（MTE1）、mmad（M）、L0C→GM（FIX）四级用 `SetFlag/WaitFlag<HardEvent::MTE2_MTE1 / M_MTE1 / MTE1_M / M_FIX / FIX_M>` 衔接。 | ||
| 67 | +4. **小 tile 加 PIPE_M barrier**：`(m/16)*(n/16) < 10` 时 Mmad 后插 `PipeBarrier<PIPE_M>` 防流水冲突。 | ||
| 68 | +5. **nd2nz stride 回退**：GM 行 stride 超过 65535 时按行/列降级为多次单条 DataCopy，避免 DMA 参数非法。 | ||
| 69 | + | ||
| 70 | +## 常见陷阱 | ||
| 71 | + | ||
| 72 | +- ❌ 保留"先 Cast fp32 再进 cube"的老习惯 → UB 流量翻倍且 cube 只能吃 fp32（算力减半或不支持）；fp16/bf16 必须直达 L1。 | ||
| 73 | +- ❌ 对 K 做显式 transpose 再进 L0B → 多一次 UB/L1 往返；用 ColumnMajor 视图 + nd2nz 一次到位。 | ||
| 74 | +- ❌ P 以 fp32 喂 PV mmad → PV 算力减半；必须 downcast 回输入 dtype（注意 bf16 用 CAST_RINT）。 | ||
| 75 | +- ❌ L0C 不双缓冲 → FIX 与 M 互相停等，cube 利用率腰斩。 | ||
| 76 | +- ❌ 矩阵维 < 64 仍强行上 cube → 启动开销不划算，退回 P94/P95 向量化（判定见 P102）。 | ||
| 77 | + | ||
| 78 | +## 代码搜索关键词 | ||
| 79 | + | ||
| 80 | +```bash | ||
| 81 | +grep -nE "ReduceSum|Mmad|BlockMmad|Nd2Nz|ColumnMajor|loadQGM|L0C|Cast\(.*RoundMode" op_kernel/*.cpp op_kernel/*.h | ||
| 82 | +``` | ||
Aknowledge/contrib/ascendc/optimizations/evolution_strategies/cards/p121_fa_splitfuse_cv_pipeline.md+96-0
| @@ -0,0 +1,96 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P121: FA SplitFuse CV 流水拓扑（AIC QK+PV / AIV softmax+rescaleO，GM workspace 握手槽 | ||
| 4 | + + PRE_LAUNCH 预发射）' | ||
| 5 | +description: FlashAttention 单核内的四个阶段天然分属两种核：`QK^T`(cube)、`softmax update`(vector)、`PV`(cube)、`O | ||
| 6 | + rescale`(vector)。 | ||
| 7 | +tags: | ||
| 8 | +- task.optimization | ||
| 9 | +- topic.attention | ||
| 10 | +- topic.programming_model.execution_unit.cube | ||
| 11 | +- topic.implementation.pipeline | ||
| 12 | +status: stable | ||
| 13 | +sources: | ||
| 14 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P121_fa_splitfuse_cv_pipeline.md | ||
| 15 | + role: primary | ||
| 16 | +created_at: '2026-10-09T00:00:00Z' | ||
| 17 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 18 | +platforms: | ||
| 19 | +- a2 | ||
| 20 | +--- | ||
| 21 | + | ||
| 22 | +# P121: FA SplitFuse CV 流水拓扑（AIC QK+PV / AIV softmax+rescaleO，GM workspace 握手槽 + PRE_LAUNCH 预发射） | ||
| 23 | + | ||
| 24 | +## 核心思想 | ||
| 25 | + | ||
| 26 | +FlashAttention 单核内的四个阶段天然分属两种核：`QK^T`(cube)、`softmax update`(vector)、`PV`(cube)、`O rescale`(vector)。串行实现（或 `SyncAll` 全局屏障实现）里 Cube 等 Vector、Vector 等 Cube，两类核各有一半时间在空转。 | ||
| 27 | + | ||
| 28 | +FIA SplitFuse 拓扑把四个阶段铺成 **AIC/AIV 双核生产者-消费者流水**： | ||
| 29 | + | ||
| 30 | +``` | ||
| 31 | +AIC: QK[i] ──qkReady──► ... ──softmaxReady──► PV[i] ──pvReady──► ... | ||
| 32 | +AIV: ... ──wait qkReady──► Softmax[i] ──wait...──► RescaleO[i] ──... | ||
| 33 | +``` | ||
| 34 | + | ||
| 35 | +三个关键机制： | ||
| 36 | + | ||
| 37 | +1. **GM workspace 多槽环 = CV 间数据通路**：S/P/OTmp/Update 全部经 GM workspace 交接，按 `coreIdx * SLOTS + slot % SLOTS` 寻址（`SLOTS = PRE_LAUNCH + 1`，FIA 取 PRE_LAUNCH=2 即 3 槽）。槽数 ≥2 才能让"写 i+1"与"读 i"并行；workspace 预算 = `cores × Q_TILE(128) × KV_STACK(512) × (4B S + 2B P + 4B OTmp + 4B Update) × SLOTS`。 | ||
| 38 | +2. **CrossCoreFlag 三握手**：`qkReady(1) / softmaxReady(2) / pvReady(3)` 三个 flag，`CrossCoreSetFlag<0x2, PIPE_FIX>` / `CrossCoreWaitFlag` 做 AIC↔AIV 定向通知（0x2 = 跨核生产者-消费者模式），**不用 SyncAll**。 | ||
| 39 | +3. **PRE_LAUNCH 预发射错位**：KV 循环上界写成 `kvEnd + preKVNum`，循环体内"当前 i 做 QK+softmax，i-PRE_LAUNCH 做 PV+rescaleO"——Cube 永远领先 Vector 两个阶段，尾部靠多跑的 preKVNum 次迭代排空。配合**初始化时把所有 HardEvent 预置 SetFlag**（M_MTE1/MTE1_MTE2/FIX_M/MTE3_V/...），第一次 Wait 直接通过，流水无冷启动。 | ||
| 40 | + | ||
| 41 | +与 P14 的关系：P14 是通用"环形任务缓存 CV 预发射"模式；本卡是 FA 场景的具体拓扑——双 cube 阶段夹双 epilogue、GM workspace 槽位寻址公式、三 flag 协议、预置 SetFlag 暖机清单。与 P73 的关系：P73 调 AIC:AIV 核数比例；本卡决定流水结构本身。 | ||
| 42 | + | ||
| 43 | +## 代码骨架 | ||
| 44 | + | ||
| 45 | +```cpp | ||
| 46 | +// 初始化暖机：预置全部事件，首次 WaitFlag 直通（AIC/AIV 各自一套） | ||
| 47 | +SetFlag<HardEvent::M_MTE1>(EVENT_ID0..7); SetFlag<HardEvent::FIX_M>(EVENT_ID0/1); ... | ||
| 48 | + | ||
| 49 | +for (uint32_t kvSIdx = kvStart; kvSIdx < kvEnd + preKVNum; kvSIdx++) { | ||
| 50 | + uint32_t slot = stackSeqCount % (PRE_LAUNCH + 1); | ||
| 51 | + uint64_t wsOff = coreIdx * WS_BLOCK * (PRE_LAUNCH+1) + slot * WS_BLOCK; | ||
| 52 | + if (kvSIdx < kvEnd) { // 阶段 1+2：当前 tile | ||
| 53 | +#ifdef CUBE | ||
| 54 | + blockMmadQK(gQ, gK, gS[wsOff], ...); // AIC: S=QK^T → workspace 槽 | ||
| 55 | + CrossCoreSetFlag<0x2, PIPE_FIX>(qkReady); | ||
| 56 | +#else | ||
| 57 | + CrossCoreWaitFlag(qkReady); // AIV: softmax(S)→P 回写同槽 | ||
| 58 | + epilogueOnlineSoftmax(gP[wsOff], gS[wsOff], ...); | ||
| 59 | + CrossCoreSetFlag<0x2, PIPE_MTE3>(softmaxReady); | ||
| 60 | +#endif | ||
| 61 | + } | ||
| 62 | + if (kvSIdx >= kvStart + preKVNum) { // 阶段 3+4：滞后 tile | ||
| 63 | +#ifdef CUBE | ||
| 64 | + blockMmadPV(gP[wsOffL], gV, gOTmp[wsOffL], ..., softmaxReady); | ||
| 65 | + CrossCoreSetFlag<0x2, PIPE_FIX>(pvReady); | ||
| 66 | +#else | ||
| 67 | + CrossCoreWaitFlag(pvReady); // AIV: O rescale/累加 | ||
| 68 | + epilogueRescaleO(gO, gOTmp[wsOffL], gOUpdate[...], ...); | ||
| 69 | +#endif | ||
| 70 | + } | ||
| 71 | + stackSeqCount++; | ||
| 72 | +} | ||
| 73 | +// 任务级负载均衡：taskIdx = coreIdx; taskIdx += coreNum 跨步领取 | ||
| 74 | +``` | ||
| 75 | + | ||
| 76 | +## 关键修改点 | ||
| 77 | + | ||
| 78 | +1. **删除主循环内全部 `SyncAll`**，替换为三 flag 定向握手；flag id 用常量（1/2/3）防冲突。 | ||
| 79 | +2. **workspace 按核分槽**：每核独占 `SLOTS` 个槽，槽内 S/P/OTmp/Update 分区；寻址公式必须核间不重叠。 | ||
| 80 | +3. **循环上界 `+ preKVNum`** 并拆分"当前阶段/滞后阶段"两个 if，保证尾部 PV/rescale 被排空。 | ||
| 81 | +4. **初始化预置 SetFlag 清单**必须与主循环 Wait 的事件一一对应，漏一个就死锁或首拍停等。 | ||
| 82 | +5. **softmax 输出 P 与 QK 输出 S 复用同一槽位偏移**（读完 S 写 P），省一半 workspace。 | ||
| 83 | + | ||
| 84 | +## 常见陷阱 | ||
| 85 | + | ||
| 86 | +- ❌ 槽位数 < PRE_LAUNCH+1 → 写 i+1 覆盖未读完的 i，数据竞争且结果偶发错误（难复现）。 | ||
| 87 | +- ❌ 用 `SyncAll` 替代 flag 握手 → 退化为串行，本策略收益清零。 | ||
| 88 | +- ❌ 忘记初始化预置 SetFlag → 第一次 WaitFlag 永久阻塞（kernel hang）。 | ||
| 89 | +- ❌ AIV 侧 `coreIdx` 忘记除以 `GetSubBlockNum()` → 多 AIV 子核写同一 workspace 槽。 | ||
| 90 | +- ❌ 循环上界不加 preKVNum → 尾部 PRE_LAUNCH 个 tile 的 PV/rescale 丢失，输出缺块。 | ||
| 91 | + | ||
| 92 | +## 代码搜索关键词 | ||
| 93 | + | ||
| 94 | +```bash | ||
| 95 | +grep -nE "CrossCoreSetFlag|CrossCoreWaitFlag|SyncAll|PRE_LAUNCH|workspace|WORKSPACE|GetSubBlockNum|qkReady|softmaxReady|pvReady" op_kernel/*.cpp op_kernel/*.h | ||
| 96 | +``` | ||
| @@ -0,0 +1,89 @@ | |||
| 1 | +--- | ||
| 2 | +type: Optimization | ||
| 3 | +title: 'P122: KV 堆叠大块化 + L1 容量自适应 Q/KV 划分' | ||
| 4 | +description: KV tile 是 32/128 小块且尺寸硬编码时，每次 QK^T 的 N 维、PV 的 K 维喂不饱 mmad，且每 32~128 个 | ||
| 5 | + KV 就要完成一轮"搬运→计算→握手"，Cube 启动开销和同步开销占比随 tile 变小线性放大。 | ||
| 6 | +tags: | ||
| 7 | +- task.optimization | ||
| 8 | +- topic.attention | ||
| 9 | +- topic.programming_model.execution_unit.cube | ||
| 10 | +- topic.data_movement.memory_access | ||
| 11 | +- topic.implementation.pipeline | ||
| 12 | +status: stable | ||
| 13 | +sources: | ||
| 14 | +- resource: https://gitcode.com/cann/cannbot-skills/blob/120ae95ac453606fe8eb69a5d1662382d30fc315/community/ops/evolution-strategies/references/cards/P122_kv_stack_large_tile_l1_adaptive.md | ||
| 15 | + role: primary | ||
| 16 | +created_at: '2026-10-09T00:00:00Z' | ||
| 17 | +updated_at: '2026-10-09T00:00:00Z' | ||
| 18 | +platforms: | ||
| 19 | +- a2 | ||
| 20 | +--- | ||
| 21 | + | ||
| 22 | +# P122: KV 堆叠大块化 + L1 容量自适应 Q/KV 划分 | ||
| 23 | + | ||
| 24 | +## 核心思想 | ||
| 25 | + | ||
| 26 | +KV tile 是 32/128 小块且尺寸硬编码时，每次 QK^T 的 N 维、PV 的 K 维喂不饱 mmad，且每 32~128 个 KV 就要完成一轮"搬运→计算→握手"，Cube 启动开销和同步开销占比随 tile 变小线性放大。 | ||
| 27 | + | ||
| 28 | +把 KV 维**堆叠成大块**（基准 `KV_STACK = 512`，paged 场景把多个 128 物理块 stack 成一个逻辑块）： | ||
| 29 | + | ||
| 30 | +1. **每次 CV 握手覆盖 512 个 KV**（而非 32/128）：跨核 flag 同步频率降到 1/4~1/16，P121 流水的握手开销被摊薄。 | ||
| 31 | +2. **mmad 的 N/K 维拉满**：QK^T 一次算 `[128q × 512kv]`，PV 一次算 `K=512`（L1 tile 256×2 拍），L0 装载次数和启动开销按比例下降。 | ||
| 32 | +3. **L1 容量自适应划分**（不硬编码 tile）：K/V 的 L1 块尺寸由 L1 剩余预算动态推出—— | ||
| 33 | + | ||
| 34 | +```cpp | ||
| 35 | +kDyn = max(RoundUp(embed, 128), 256); // K 维：对齐 128，下限 256 | ||
| 36 | +budget = L1_MAX(512K) - embedV * KV_STACK(512) * dtype; // 先扣掉 V 常驻区 | ||
| 37 | +qSize = Q_TILE(128) * kDyn * dtype; // 再扣 Q 常驻区 | ||
| 38 | +nDyn = (budget - qSize) / kDyn / dtype / DOUBLE_BUFFER; // KV N 维，含双缓冲 | ||
| 39 | +nDyn = min(nDyn, 128); nDyn = RoundDown(nDyn, 32); // 上限 128，对齐 32 | ||
| 40 | +``` | ||
| 41 | + | ||
| 42 | +Q 常驻 + V 常驻 + K 双缓冲三者瓜分 L1，剩余容量全部折成 KV 的 N 维块——embed 越大块越小，全自动适配。 | ||
| 43 | + | ||
| 44 | +与 P18/P83 的关系：P18 是 L1 7-buffer 常驻划分的通用模式，P83 是 L1 常驻复用；本卡给出 FA 场景的具体预算公式与堆叠基准值（512/128/32 对齐）。 | ||
| 45 | + | ||
| 46 | +## 代码骨架 | ||
| 47 | + | ||
| 48 | +```cpp | ||
| 49 | +// === 改造前：小 KV tile，硬编码 === | ||
| 50 | +constexpr int BN = 32; // 每轮 32 个 KV | ||
| 51 | +for (int t = 0; t < kvLen; t += BN) { load K/V tile; compute; sync; } | ||
| 52 | + | ||
| 53 | +// === 改造后：KV 堆叠 512 + L1 预算自适应 === | ||
| 54 | +constexpr uint32_t MAX_KV_STACK_LEN = 512; // 逻辑 KV 块 | ||
| 55 | +constexpr uint32_t Q_TILE_CEIL = 128; | ||
| 56 | +uint32_t kDynNum = RoundUp(embed, 128); kDynNum = max(kDynNum, 256); | ||
| 57 | +uint32_t maxQKPL1 = L1_MAX_SIZE - embedV * MAX_KV_STACK_LEN * sizeof(T); | ||
| 58 | +uint32_t maxQL1 = Q_TILE_CEIL * kDynNum * sizeof(T); | ||
| 59 | +uint32_t nDynNum = (maxQKPL1 - maxQL1) / kDynNum / sizeof(T) / DOUBLE_BUFFER; | ||
| 60 | +nDynNum = min(RoundDown(nDynNum, 32), 128); | ||
| 61 | +blockMmadQK.init(resource, nDynNum, kDynNum, MAX_KV_STACK_LEN); | ||
| 62 | +// paged cache：blockStackNum = ceil(512 / blockSize)，跨物理块拼逻辑块 | ||
| 63 | +for (kvSIdx ...; kvSIdx < kvEnd; kvSIdx++) { // 每轮 512 个 KV | ||
| 64 | + stackSeqTile = min(512, noSkipKvS - kvSIdx * 512); // 尾块取余 | ||
| 65 | + blockMmadQK(...{rowNum, stackSeqTile, embed}...); | ||
| 66 | +} | ||
| 67 | +``` | ||
| 68 | + | ||
| 69 | +## 关键修改点 | ||
| 70 | + | ||
| 71 | +1. **KV 逻辑块从 32/128 提到 512**（或按 UB/L1 预算重算），paged 场景按物理块堆叠 `blockStackNum = ceil(STACK/blockSize)`。 | ||
| 72 | +2. **尾块取余**：最后一个 stack tile 用 `noSkipKvS - idx*512`，配合 P123 的 causal 跳过边界。 | ||
| 73 | +3. **nDyn/kDyn 全部公式化**，禁硬编码；kDyn 下限 256 保证 K 维效率，nDyn 上限 128 防 L0/L1 溢出。 | ||
| 74 | +4. **对齐铁律**：kDyn 对齐 128，nDyn 对齐 32，DOUBLE_BUFFER=2 计入预算。 | ||
| 75 | +5. **PV 的 K 维联动**：`kPV = nDyn * kDyn / PV_L1_M`，QK 与 PV 的 L1 划分共享同一份 K/V 常驻区（PV init 传 `L1_QK_SIZE` 偏移复用）。 | ||
| 76 | + | ||
| 77 | +## 常见陷阱 | ||
| 78 | + | ||
| 79 | +- ❌ 堆叠后忘记改 workspace 槽尺寸（`Q_TILE × STACK`）→ workspace 溢出踩相邻核数据。 | ||
| 80 | +- ❌ nDyn 未对齐 32 / 未除双缓冲 → L1 分配越界，运行时随机踩内存。 | ||
| 81 | +- ❌ embed > 256 时仍按 256 切 kDyn → L1 预算公式溢出为负，nDyn 算成巨大值（无符号回绕）。 | ||
| 82 | +- ❌ causal 场景尾块仍按 512 满载 → 多读多算且 mask 越界；必须配合 P123 的 `noSkipKvS`。 | ||
| 83 | +- ❌ 堆叠块大于实际 kvSeqlen 时不取 min → GM 越界读。 | ||
| 84 | + | ||
| 85 | +## 代码搜索关键词 | ||
| 86 | + | ||
| 87 | +```bash | ||
| 88 | +grep -nE "KV_STACK|STACK_LEN|kDyn|nDyn|L1_MAX|DOUBLE_BUFFER|blockStackNum|RoundUp|RoundDown" op_kernel/*.cpp op_kernel/*.h | ||
| 89 | +``` | ||
这个平台需要明确一下，是否是仅在a2平台上适用