已合并
feat: 迁移 ascendc-port-orchestrator OKF 卡并按审查处置收敛 #36
ariel_zjl创建于 9月16日
feat: 迁移 ascendc-port-orchestrator OKF 卡并按审查处置收敛 #36
已合并
共 620 个文件变更+44337-3
| @@ -322,7 +322,7 @@ cases: | |||
| 322 | expected: | 322 | expected: |
| 323 | - common/platforms/concepts/ascend_910b | 323 | - common/platforms/concepts/ascend_910b |
| 324 | - id: common_platform_unmatched_terms_search | 324 | - id: common_platform_unmatched_terms_search |
| 325 | - query: platform mapping | 325 | + query: thermal throttling policy |
| 326 | filters: | 326 | filters: |
| 327 | domain: common | 327 | domain: common |
| 328 | technology: platforms | 328 | technology: platforms |
| @@ -1,9 +1,12 @@ | |||
| 1 | # 硬件参考 | 1 | # 硬件参考 |
| 2 | 2 | ||
| 3 | -Ascend 910B/910_95 硬件规格与强制约束速查。生成任何 Triton Ascend kernel 之前必须过一遍约束检查清单。 | 3 | +Ascend 910B/910_95/950PR/910C 硬件规格与强制约束速查。生成任何 Triton Ascend kernel 之前必须过一遍约束检查清单。 |
| 4 | 4 | ||
| 5 | - [ascend_910b](ascend_910b.md)：910B2/910B3 核心架构、Grid 语义、内存层次、编译选项 | 5 | - [ascend_910b](ascend_910b.md)：910B2/910B3 核心架构、Grid 语义、内存层次、编译选项 |
| 6 | - [ascend_910_95](ascend_910_95.md)：910_95/950PR/950DT（dav-c310，Reg-based A5 架构）核数、UB/L1/L0C 容量、专有数据通路与 910B 差异 | 6 | - [ascend_910_95](ascend_910_95.md)：910_95/950PR/950DT（dav-c310，Reg-based A5 架构）核数、UB/L1/L0C 容量、专有数据通路与 910B 差异 |
| 7 | - [NPU 内存层次与数据通路模型](npu_memory_model.md)：Ascend NPU 的 GM/L1/UB/L0A/L0B/L0C 存储层次、容量与对齐约束、MTE/FIX 数据通路及 tl.load/tl.store/tl.dot 的硬件映射。 | 7 | - [NPU 内存层次与数据通路模型](npu_memory_model.md)：Ascend NPU 的 GM/L1/UB/L0A/L0B/L0C 存储层次、容量与对齐约束、MTE/FIX 数据通路及 tl.load/tl.store/tl.dot 的硬件映射。 |
| 8 | - [constraints_cheatsheet](constraints_cheatsheet.md)：强制约束速查（UB/Gird/对齐） + 性能约束 + autotune 搜索空间 | 8 | - [constraints_cheatsheet](constraints_cheatsheet.md)：强制约束速查（UB/Gird/对齐） + 性能约束 + autotune 搜索空间 |
| 9 | - [ascend_patterns](ascend_patterns.md)：Ascend 特有的编程模式与 GPU→Ascend 迁移差异 | 9 | - [ascend_patterns](ascend_patterns.md)：Ascend 特有的编程模式与 GPU→Ascend 迁移差异 |
| 10 | +- [Ascend950PR (David V100) -- Target Platform](target_ascend950pr.md)：950PR（DAV_3510/arch35）完整平台规格：核数、UB/L1/L0 容量、6 管线与 Scalar 单元、SIMT+SIMD 模型、原子操作、跨核同步、Cube 数值行为与已知缺陷（手册+白皮书+实测）。 | ||
| 11 | +- [Ascend910C (DaVinci V220, dual-die) — Target Platform](target_ascend910c.md)：910C（A3）平台规格：SIMD-only、80~96 AIV、UB 192KB/AIV、L0C 128KB、SOC 字符串与构建参数、与 A5 移植差异及待实测项。 | ||
| 12 | +- [Ascend910B (DaVinci V220) — Target Platform](target_ascend910b.md)：910B（A2）平台规格：SIMD-only、B1–B4 SOC 型号、UB 192KB/AIV、L0C 128KB、构建参数与待上机验证项。 | ||
| @@ -0,0 +1,124 @@ | |||
| 1 | +--- | ||
| 2 | +type: Concept | ||
| 3 | +title: "Ascend910B (DaVinci V220) — Target Platform" | ||
| 4 | +description: "Ascend910B（DaVinci V220 单 die，A2）目标平台规格：SIMD-only 编程模型、B1–B4 SOC 型号、24/20 AICore、UB 192KB/AIV、L0C 128KB、CMake 构建参数与待硬件实测确认项（前瞻性规格，尚未上机验证）。" | ||
| 5 | +tags: [task.lookup, topic.architecture, topic.memory] | ||
| 6 | +status: draft | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/runbooks/hardware/target-ascend910b.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-07-12T16:00:00Z" | ||
| 11 | +updated_at: "2026-09-16T07:31:58Z" | ||
| 12 | +platforms: [a2] | ||
| 13 | +aliases: [hw/ascend910b] | ||
| 14 | +--- | ||
| 15 | + | ||
| 16 | +# Ascend910B (DaVinci V220) — Target Platform | ||
| 17 | + | ||
| 18 | +> 原 OKF v1 卡号：**hw/ascend910b**（原 OKF v1 kind: guide；confidence: single_run）。 | ||
| 19 | + | ||
| 20 | +> 本卡规格未经真机核验（forward-looking spec），保持 draft 待硬件验证。 | ||
| 21 | + | ||
| 22 | +> **NOT YET VERIFIED ON HARDWARE.** This file is a forward-looking spec | ||
| 23 | +> derived from CANN source + public references. When an A2 server becomes | ||
| 24 | +> available, run smoke tests (build + `acl.get_soc_name()` + a SIMD `Add` | ||
| 25 | +> kernel) and update `verified: true` + `date_verified`. | ||
| 26 | +> | ||
| 27 | +> A2 = single die DaVinci V220. CANN treats it as `ascend910b_list` family, | ||
| 28 | +> compiles with `BUILD_MODE=c220`, arch dir `arch22`. Same arch as A3 but | ||
| 29 | +> single-die (so half the AIV count of A3). | ||
| 30 | +> | ||
| 31 | +> **SIMD-only.** Same SIMT-restriction list as A3 — see [target_ascend910c](target_ascend910c.md). | ||
| 32 | + | ||
| 33 | +## SOC version reference | ||
| 34 | + | ||
| 35 | +| Variant | SOC string | Marketed as | | ||
| 36 | +|---------|-----------|-------------| | ||
| 37 | +| B1 | `Ascend910B1` | 910B1 | | ||
| 38 | +| B2 | `Ascend910B2` / `Ascend910B2C` | 910B2 | | ||
| 39 | +| B3 | `Ascend910B3` | 910B3 | | ||
| 40 | +| B4 | `Ascend910B4` / `Ascend910B4-1` | 910B4 | | ||
| 41 | + | ||
| 42 | +`acl.get_soc_name()` returns the active SOC string at runtime — query before | ||
| 43 | +choosing build flags. Most public A2 deployments are B3. | ||
| 44 | + | ||
| 45 | +## Compute (per chip) | ||
| 46 | + | ||
| 47 | +| Parameter | B2 | B3/B4 | Source | | ||
| 48 | +|-----------|----|----|--------| | ||
| 49 | +| AICore total | 24 | 20 | HAMi guide, public refs | | ||
| 50 | +| AIV per AICore | 2 | 2 | CANN `GetAicAivTaskRation()` | | ||
| 51 | +| **AIV total** | **48** | **40** | derived | | ||
| 52 | +| AIC total (Cube) | 24 | 20 | derived | | ||
| 53 | +| **Programming model** | SIMD only | SIMD only | DavidV100 manual §25.1.1 | | ||
| 54 | +| `threadIdx` / `blockDim` SIMT | NOT available | NOT available | arch22 | | ||
| 55 | +| `Simt::` namespace | NOT available | NOT available | arch22 | | ||
| 56 | + | ||
| 57 | +What works / fails on A2 = same as A3 (see [target_ascend910c](target_ascend910c.md)). | ||
| 58 | + | ||
| 59 | +## Memory (per chip) | ||
| 60 | + | ||
| 61 | +| Parameter | B2 | B3/B4 | Source | | ||
| 62 | +|-----------|----|----|--------| | ||
| 63 | +| HBM type | HBM2e | HBM3e (B3/B4) | public refs | | ||
| 64 | +| **HBM capacity** | 64 GB | 32 GB (B4) / 64 GB | HAMi guide | | ||
| 65 | +| **HBM bandwidth** | ~400 GB/s | ~800 GB/s | public refs | | ||
| 66 | +| **UB / AIV** | **192 KB** | **192 KB** | CANN `hardware.h` | | ||
| 67 | +| L1 / AIC | 512 KB | 512 KB | CANN `arch.h` | | ||
| 68 | +| L0A / AIC | 64 KB | 64 KB | CANN `arch.h` | | ||
| 69 | +| L0B / AIC | 64 KB | 64 KB | CANN `arch.h` | | ||
| 70 | +| **L0C / AIC** | **128 KB** | **128 KB** | CANN `arch.h` (`__NPU_ARCH__==3510`→256, else→128 — **was "3101", corrected 2026-06-18 vs arch.h:29**) | | ||
| 71 | +| **BT (bias table) / AIC** | **1 KB (1024)** | **1 KB (1024)** | cannbot npu-arch §2.3 DAV_2201 `bt_size` (A5/950PR=4KB — differs) | | ||
| 72 | +| **Structured sparsity 4:2** | **supported** (`sparsity=1`) | **supported** | cannbot npu-arch §2.3 DAV_2201 (A5/950PR `sparsity=0` — NOT supported; don't assume 4:2 on A5) | | ||
| 73 | +| L2 cache | TBD | TBD | needs probe | | ||
| 74 | + | ||
| 75 | +## Atomics | ||
| 76 | + | ||
| 77 | +Same TBD list as A3 — see [target_ascend910c](target_ascend910c.md) § Atomics. A2 and A3 share the | ||
| 78 | +V220 architecture, so atomic capability should be identical, but verify | ||
| 79 | +on first A2 port. | ||
| 80 | + | ||
| 81 | +## Cross-platform porting summary | ||
| 82 | + | ||
| 83 | +| What | A5 (950PR) | A3 (910C) | A2 (910B) | | ||
| 84 | +|------|-----------|-----------|-----------| | ||
| 85 | +| Programming model | SIMT + SIMD | SIMD only | SIMD only | | ||
| 86 | +| `__NPU_ARCH__` | 3510 | 2201 | 2201 | | ||
| 87 | +| arch dir | arch35 | arch22 | arch22 | | ||
| 88 | +| AIV count | 56 | 80~96 | 40~48 | | ||
| 89 | +| UB per AIV | 256 KB | 192 KB | 192 KB | | ||
| 90 | +| L0C per AIC | 256 KB | 128 KB | 128 KB | | ||
| 91 | +| HBM bandwidth | 1.6 TB/s | ~1.6 TB/s (dual-die) | 0.4–0.8 TB/s | | ||
| 92 | + | ||
| 93 | +## Build flags (CMake) | ||
| 94 | + | ||
| 95 | +```cmake | ||
| 96 | +set(SOC_VERSION "Ascend910B3" CACHE STRING "...") # or B1/B2/B4 to match your chip | ||
| 97 | +set(ASCEND_CANN_PACKAGE_PATH "/usr/local/Ascend/cann" CACHE STRING "...") | ||
| 98 | +include(${ASCEND_CANN_PACKAGE_PATH}/aarch64-linux/tikcpp/ascendc_kernel_cmake/ascendc.cmake) | ||
| 99 | +``` | ||
| 100 | + | ||
| 101 | +## Open verification tasks (when A2 server becomes available) | ||
| 102 | + | ||
| 103 | +1. Confirm SOC string via `acl.get_soc_name()` and update this doc's | ||
| 104 | + `soc_version` frontmatter | ||
| 105 | +2. Confirm AIV count via `GetCoreNumAiv()` (40, 48, or other) | ||
| 106 | +3. Smoke test: build the SIMD `Add` test kernel with the right SOC and run it | ||
| 107 | +4. Probe: FP16/BF16 atomicAdd support (same probe as A3) | ||
| 108 | +5. Probe: HBM bandwidth measurement (compare against the public 400 GB/s figure) | ||
| 109 | +6. Update `verified: true` and `date_verified` in frontmatter | ||
| 110 | + | ||
| 111 | +## References | ||
| 112 | + | ||
| 113 | +- 源仓 `docs/design/PLUGIN_PARADIGM_NOTES.md#ascend-chip-comparison`（插件仓内部文档，未迁移） — A2/A3/A5 comparison | ||
| 114 | +- CANN source (`code_channel_infer.py`, `arch.h`, `hardware.h`) | ||
| 115 | +- HAMi Ascend 910B virtualization guide (public) | ||
| 116 | +- HiAscend AscendC manual 9.0.0 (Chinese version) | ||
| 117 | +- Hi1980B&C V100 Davinci Cloud 芯片 用户指南（本地 docx 文档，未迁移） — | ||
| 118 | + primary hardware doc covering both Hi1980B (A2) and Hi1980C (A3) | ||
| 119 | + | ||
| 120 | +## 相关卡片 | ||
| 121 | + | ||
| 122 | +- [Ascend910C (DaVinci V220, dual-die) — Target Platform](target_ascend910c.md)：A3 双 die 同族平台（SIMT 限制清单与 Atomics 待测项见该卡）。 | ||
| 123 | +- [Ascend950PR (David V100) -- Target Platform](target_ascend950pr.md)：A5 目标平台规格。 | ||
| 124 | +- [Ascend 910B2/910B3 硬件规格](ascend_910b.md)：同芯片的 Triton 视角硬件速查（Grid 语义、编译选项），与本卡来源不同、互为补充。 | ||
| @@ -0,0 +1,180 @@ | |||
| 1 | +--- | ||
| 2 | +type: Concept | ||
| 3 | +title: "Ascend910C (DaVinci V220, dual-die) — Target Platform" | ||
| 4 | +description: "Ascend910C（DaVinci V220 双 die 封装，A3）目标平台规格：SIMD-only 编程模型、40~48 AICore、UB 192KB/AIV、L0C 128KB、SOC 字符串与 CMake 构建参数、与 A5 的跨平台移植差异及待实测确认项。" | ||
| 5 | +tags: [task.lookup, topic.architecture, topic.memory] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/runbooks/hardware/target-ascend910c.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-07-12T16:00:00Z" | ||
| 11 | +updated_at: "2026-07-12T16:00:00Z" | ||
| 12 | +platforms: [a3] | ||
| 13 | +aliases: [hw/ascend910c] | ||
| 14 | +--- | ||
| 15 | + | ||
| 16 | +# Ascend910C (DaVinci V220, dual-die) — Target Platform | ||
| 17 | + | ||
| 18 | +> 原 OKF v1 卡号：**hw/ascend910c**（原 OKF v1 kind: guide；confidence: single_run）。 | ||
| 19 | + | ||
| 20 | +> A3 platform = two `Ascend910B` dies in one package. CANN treats it as a member | ||
| 21 | +> of the `ascend910b_list` family (NOT a separate `ascend910c_list`); compiler | ||
| 22 | +> selects `BUILD_MODE=c220` and arch dir `arch22`, same as A2/910B. | ||
| 23 | +> | ||
| 24 | +> **SIMD-only.** No SIMT (no `threadIdx`, no `Simt::` namespace, no | ||
| 25 | +> `WarpShflSync`/`ThreadBarrier`/`LAUNCH_BOUND`). Kernels written for A5 (950PR) | ||
| 26 | +> in SIMT mode will NOT compile for A3 — they must be rewritten using the | ||
| 27 | +> `TPipe / TQue / DataCopy` SIMD pattern. | ||
| 28 | +> | ||
| 29 | +> Verified data on 192.0.2.210 container `npu-a3` (CANN 9.0.0_beta.2). | ||
| 30 | +> Many memory/atomic numbers below are inherited from CANN source + | ||
| 31 | +> public references and **need empirical confirmation on actual A3 hardware** | ||
| 32 | +> (probe queue tracked out-of-band; 源仓流程为运行 `aog-hardware-probe` 把发现沉淀到 | ||
| 33 | +> `kb/okf/runbooks/hardware/`)。 | ||
| 34 | + | ||
| 35 | +## SOC version reference | ||
| 36 | + | ||
| 37 | +| Variant | SOC string | Notes | | ||
| 38 | +|---------|-----------|-------| | ||
| 39 | +| 9392 | `Ascend910_9392` | **The chip on our verified A3 server** (npu-smi `Ascend910`) | | ||
| 40 | +| Other 93xx | `Ascend910_9391`, `Ascend910_9381`, `Ascend910_9382`, `Ascend910_9372`, `Ascend910_9362` | Other A3 sub-models — same `c220` build path | | ||
| 41 | + | ||
| 42 | +`acl.get_soc_name()` returns the active SOC string at runtime — query before | ||
| 43 | +choosing build flags. | ||
| 44 | + | ||
| 45 | +## Compute (per package = 2 dies) | ||
| 46 | + | ||
| 47 | +| Parameter | Value | Source | | ||
| 48 | +|-----------|-------|--------| | ||
| 49 | +| AICore total | **40 ~ 48** (= 2× 20-24 per die) | comparison doc; verify with npu-smi | | ||
| 50 | +| AIV per AICore | 2 | CANN `GetAicAivTaskRation()` | | ||
| 51 | +| **AIV total** | **80 ~ 96** | AICore × 2 | | ||
| 52 | +| AIC total (Cube) | 40 ~ 48 | AICore × 1 | | ||
| 53 | +| **Programming model** | **SIMD only** | DavidV100 manual §25.1.1; arch22 has no SIMT path | | ||
| 54 | +| `threadIdx` / `blockDim` (SIMT) | **NOT available** | arch22 lacks SIMT path | | ||
| 55 | +| `Simt::` namespace | **NOT available** | (ditto) | | ||
| 56 | +| Warp / LAUNCH_BOUND | **N/A** | (ditto) | | ||
| 57 | +| `GetBlockNum()` / `GetBlockIdx()` | available (block-level only) | universal AscendC | | ||
| 58 | +| Block scheduling | Software time-slicing | shared with A2/A5 family | | ||
| 59 | +| Clock frequency | TBD | needs verification | | ||
| 60 | + | ||
| 61 | +### What WORKS on A3 (SIMD path) | ||
| 62 | + | ||
| 63 | +- `TPipe`, `TQue<QuePosition::VECIN/VECOUT>`, `TBuf` | ||
| 64 | +- `DataCopy(LocalTensor, GlobalTensor, len)` | ||
| 65 | +- Vector ops: `Add`, `Mul`, `Muls`, `Adds`, `Cast`, `Exp`, `Reduce*` (SIMD form) | ||
| 66 | +- Cube ops: `Mmad` family | ||
| 67 | +- Block-level parallelism via `GetBlockNum() / GetBlockIdx()` | ||
| 68 | +- Atomic ops on int32/uint32 (CAS, EXCH) | ||
| 69 | + | ||
| 70 | +### What FAILS on A3 (SIMT-only patterns) | ||
| 71 | + | ||
| 72 | +- Any `Simt::WarpShflSync`, `Simt::WarpReduceAddSync`, etc. | ||
| 73 | +- `ThreadBarrier()`, `ThreadFence()` | ||
| 74 | +- `__syncthreads()` equivalents | ||
| 75 | +- Per-thread allocation patterns relying on `threadIdx.x` | ||
| 76 | +- `LAUNCH_BOUND(...)` annotations | ||
| 77 | + | ||
| 78 | +If an arch35 implementation uses warp shuffle, scatter-add via per-thread atomics, | ||
| 79 | +or block-wide barriers, the A3 port MUST replace those patterns with SIMD | ||
| 80 | +equivalents (e.g. block-wide reduction via UB scratchpad + DataCopy, or | ||
| 81 | +sort-then-segment-reduce instead of unsorted scatter-add). | ||
| 82 | + | ||
| 83 | +## Memory (per package) | ||
| 84 | + | ||
| 85 | +| Parameter | Value | Source | | ||
| 86 | +|-----------|-------|--------| | ||
| 87 | +| HBM type | HBM2e / HBM3 (mixed by SKU) | public; comparison doc | | ||
| 88 | +| **HBM capacity** | **128 GB** (dual-die total; varies by SKU) | comparison doc | | ||
| 89 | +| **HBM bandwidth (peak)** | **~1.6 TB/s** (sum of two dies) | comparison doc | | ||
| 90 | +| **L2 cache** | ~64 MB per die (TBD if shared cross-die) | comparison doc; **needs probe** | | ||
| 91 | +| **UB / AIV** | **192 KB** | CANN `hardware.h` (V220 default) | | ||
| 92 | +| UB / AICore | 384 KB (= 2 × 192 KB) | derived | | ||
| 93 | +| **L1 Buffer / AIC** | **512 KB** | CANN `arch.h` L1_SIZE (universal) | | ||
| 94 | +| L0A Buffer / AIC | 64 KB | CANN `arch.h` | | ||
| 95 | +| L0B Buffer / AIC | 64 KB | CANN `arch.h` | | ||
| 96 | +| **L0C Buffer / AIC** | **128 KB** (vs A5's 256 KB — half size!) | CANN `arch.h`: `__NPU_ARCH__==3510`→256KB, else→128KB (**was "3101", corrected 2026-06-18 vs arch.h:29**) | | ||
| 97 | +| **BT (bias table) / AIC** | **1 KB (1024)** | cannbot npu-arch §2.3 DAV_2201 `bt_size` (A5/950PR=4KB — differs) | | ||
| 98 | +| **Structured sparsity 4:2** | **supported** (`sparsity=1`) | cannbot npu-arch §2.3 DAV_2201 (A5/950PR `sparsity=0` — NOT supported; don't assume 4:2 on A5) | | ||
| 99 | +| Per-AICore AXI | TBD (probe required) | comparison doc | | ||
| 100 | +| Cross-die interconnect | TBD | comparison doc | | ||
| 101 | + | ||
| 102 | +> **Cross-platform delta vs A5:** UB is 192 KB (not 256 KB) and L0C is 128 KB | ||
| 103 | +> (not 256 KB). Tiling code that assumes 256 KB UB **will OOM at compile time** | ||
| 104 | +> on A3. Always derive tile sizes from `GetUBSizeInBytes()` (runtime) or | ||
| 105 | +> a `target` parameter (compile-time). | ||
| 106 | + | ||
| 107 | +## Atomics | ||
| 108 | + | ||
| 109 | +| Operation | Status | Source | | ||
| 110 | +|-----------|--------|--------| | ||
| 111 | +| `AtomicCas` U32/S32 | Supported | universal V220 | | ||
| 112 | +| atomicAdd S32/U32 | Supported | universal V220 | | ||
| 113 | +| **atomicAdd FP32** | **TBD — needs probe** | comparison doc; not yet verified | | ||
| 114 | +| **atomicAdd FP16** | **TBD — needs probe** | comparison doc; high-risk for FP16 scatter-add ports | | ||
| 115 | +| **atomicAdd BF16** | **TBD — needs probe** | comparison doc | | ||
| 116 | +| atomicMax / atomicMin FP32 | TBD — needs probe | comparison doc | | ||
| 117 | + | ||
| 118 | +> **Scatter-add risk:** A5 has FP16/BF16 atomicAdd. If A3 lacks them, ports of | ||
| 119 | +> source scatter-add backwards (e.g. SparseGather) need a cast-to-FP32 wrapper or | ||
| 120 | +> a sort-then-reduce restructure. **First A3 port that touches scatter-add | ||
| 121 | +> MUST run a probe to confirm.** | ||
| 122 | + | ||
| 123 | +## Pipelines | ||
| 124 | + | ||
| 125 | +A3 inherits the V220 pipeline structure (CUBE, VEC, MTE1, MTE2, MTE3, FIXP + | ||
| 126 | +independent Scalar Unit). The [A5 (950PR) doc](target_ascend950pr.md)'s pipeline guidance applies *except*: | ||
| 127 | + | ||
| 128 | +- No SIMT pipeline / SIMT DCache / SIMT shared memory (those are A5-only) | ||
| 129 | +- Cube path (MMAD) and Vector path (SIMD) are present and behave like A5 | ||
| 130 | + | ||
| 131 | +## Cross-platform porting summary | ||
| 132 | + | ||
| 133 | +| What | A5 (950PR) | A3 (910C) | | ||
| 134 | +|------|-----------|-----------| | ||
| 135 | +| Programming model | SIMT + SIMD | **SIMD only** | | ||
| 136 | +| `__NPU_ARCH__` | 3510 | 2201 | | ||
| 137 | +| arch dir | arch35 | arch22 | | ||
| 138 | +| AIV count | 56 | 80 ~ 96 | | ||
| 139 | +| UB per AIV | 256 KB | **192 KB** | | ||
| 140 | +| L0C per AIC | 256 KB | **128 KB** | | ||
| 141 | +| FP16/BF16 atomicAdd | supported | **TBD** | | ||
| 142 | +| HBM bandwidth | 1.6 TB/s | ~1.6 TB/s (dual-die sum) | | ||
| 143 | +| Block scheduling | software time-slicing | software time-slicing | | ||
| 144 | + | ||
| 145 | +## Build flags (CMake) | ||
| 146 | + | ||
| 147 | +```cmake | ||
| 148 | +# In your project's CMakeLists.txt: | ||
| 149 | +set(SOC_VERSION "Ascend910_9392" CACHE STRING "...") # or other 93xx variant | ||
| 150 | +set(ASCEND_CANN_PACKAGE_PATH "/usr/local/Ascend/cann" CACHE STRING "...") | ||
| 151 | +# CANN 9.0 layout — note the new aarch64-linux/ prefix: | ||
| 152 | +include(${ASCEND_CANN_PACKAGE_PATH}/aarch64-linux/tikcpp/ascendc_kernel_cmake/ascendc.cmake) | ||
| 153 | +``` | ||
| 154 | + | ||
| 155 | +The CANN 8.x path (`compiler/tikcpp/...`) does NOT exist in CANN 9.0 — the | ||
| 156 | +multi-target CMakeLists.txt searches both layouts and picks the one that exists. | ||
| 157 | + | ||
| 158 | +## Open verification questions (PROBE QUEUE) | ||
| 159 | + | ||
| 160 | +1. FP16/BF16 atomicAdd support — write a kernel that calls atomicAdd on a FP16 | ||
| 161 | + GM tensor and verify result vs reference (源仓 probe `Q_atomic_fp16_a3.md`，未迁移) | ||
| 162 | +2. L2 cache total size and cross-die sharing semantics | ||
| 163 | +3. Cross-die HBM access latency / bandwidth penalty | ||
| 164 | +4. Actual frequency, register file size per AIV | ||
| 165 | +5. SIMD vector unit width (likely 256B × 2, same as A5/A2; needs confirmation) | ||
| 166 | + | ||
| 167 | +## References | ||
| 168 | + | ||
| 169 | +- 源仓 `docs/design/PLUGIN_PARADIGM_NOTES.md#ascend-chip-comparison`（插件仓内部文档，未迁移） — authoritative A2/A3/A5 comparison | ||
| 170 | +- CANN source (`code_channel_infer.py`, `arch.h`, `hardware.h`) | ||
| 171 | +- HiAscend AscendC manual 9.0.0 (Chinese version) — see TODO probe queue for | ||
| 172 | + extracted A3-specific items | ||
| 173 | +- Hi1980B&C V100 Davinci Cloud 芯片 用户指南 / 硬件指南（本地 docx 文档，未迁移） — | ||
| 174 | + primary hardware doc (Chinese) | ||
| 175 | + | ||
| 176 | +## 相关卡片 | ||
| 177 | + | ||
| 178 | +- [Ascend950PR (David V100) -- Target Platform](target_ascend950pr.md)：A5 目标平台规格（pipeline 指导、SIMT 能力来源卡）。 | ||
| 179 | +- [Ascend910B (DaVinci V220) — Target Platform](target_ascend910b.md)：A2 单 die 同族平台（SIMT 限制、Atomics 待测项与本卡相同）。 | ||
| 180 | +- [Ascend 910B2/910B3 硬件规格](ascend_910b.md)：同族芯片的 Triton 视角硬件速查，与本卡来源不同、互为补充。 | ||
| @@ -0,0 +1,91 @@ | |||
| 1 | +--- | ||
| 2 | +type: Concept | ||
| 3 | +title: "Q1 Probe Report: TBuf<TPosition::A1> in pure AIV kernel" | ||
| 4 | +description: "Ascend950PR 实测（CANN 9.0.0 / bisheng 2026-03-21）：纯 AIV kernel 中 `TBuf<TPosition::A1>` + UB↔L1 `DataCopy` 编译零告警但运行时 aivec error 259 非法指令；351x UB↔L1 硬通道未暴露给纯 AIV，L1-spill 改写在该工具链上不可行。" | ||
| 5 | +tags: [task.debugging, topic.memory, topic.data_movement] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/runbooks/hardware/probe-2026-04-21-q-l1-scratch-op11-kind2.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-07-12T16:00:00Z" | ||
| 11 | +updated_at: "2026-07-12T16:00:00Z" | ||
| 12 | +platforms: ["950"] | ||
| 13 | +aliases: [hw/2026-04-21_Q_l1_scratch_op11_kind2] | ||
| 14 | +--- | ||
| 15 | + | ||
| 16 | +# Q1 Probe Report: TBuf<TPosition::A1> in pure AIV kernel | ||
| 17 | + | ||
| 18 | +> 原 OKF v1 卡号：**hw/2026-04-21_Q_l1_scratch_op11_kind2**（原 OKF v1 kind: guide；confidence: single_run）。 | ||
| 19 | + | ||
| 20 | +## Verdict | ||
| 21 | + | ||
| 22 | +**ACCEPT_MISCOMPILE** — bisheng compiles the kernel with no errors or warnings, but the emitted code triggers an `aivec` illegal-instruction fault at runtime (error 259, runtime 507035) the first time either `DataCopy(UB, L1)` or `DataCopy(L1, UB)` executes on the AIV core. | ||
| 23 | + | ||
| 24 | +## Environment | ||
| 25 | + | ||
| 26 | +Captured 2026-04-22 from A5 container `npu_dev3`: | ||
| 27 | + | ||
| 28 | +- **OS**: openEuler 22.03 (LTS-SP4) | ||
| 29 | +- **Bisheng/ccec compiler**: clang 15.0.5 (clang-5c68a1cb1231 flang-5c68a1cb1231), build stamp `2026-03-21T17:07:34+08:00` | ||
| 30 | +- **CANN**: 9.0.0, innerversion V100R001C10SPC001B218, path `/usr/local/Ascend/cann-9.0.0` | ||
| 31 | +- **NPU-SMI**: 25.7.rc1.b087, driver Version 25.7.rc1.b087 | ||
| 32 | +- **SOC**: Ascend950PR (per npu-smi), built with `-DSOC_VERSION=Ascend950PR_9589` | ||
| 33 | +- **NPU used**: device 1 (AICore 0%, HBM 5235 / 131072 MB — least-loaded of 3 visible NPUs, 0 running processes) | ||
| 34 | +- **torch / torch_npu**: not cleanly importable standalone in the container shell (libhccl.so missing on PATH during bare-shell probe), but import works from inside the `run_probe.py` script via the launcher env. | ||
| 35 | + | ||
| 36 | +## Probe design | ||
| 37 | + | ||
| 38 | +4 KB (1024 fp32) round-trip: GM → UB → L1 (`TBuf<TPosition::A1>`) → UB → GM. One block, pure AIV (`KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY)`), no Cube / Mmad / Load2D. | ||
| 39 | + | ||
| 40 | +Source: `workspace/probe_l1_scratch/kernel/probe_kernel.h` — see `l1Buf_` declaration `TBuf<TPosition::A1> l1Buf_;` and the UB↔L1 `DataCopy` calls. | ||
| 41 | + | ||
| 42 | +## Iterations run (2 / 2 allowed) | ||
| 43 | + | ||
| 44 | +### Iter 1 — PipeBarrier<PIPE_ALL> between UB→L1 and L1→UB | ||
| 45 | + | ||
| 46 | +- **Build**: OK. No warnings mentioning L1/A1/TPosition. `libkernels.a` and `_probe_ext.cpython-311-x86_64-linux-gnu.so` both linked successfully. | ||
| 47 | +- **Runtime**: Vector core exception, error 259. | ||
| 48 | + | ||
| 49 | +Excerpt: | ||
| 50 | +``` | ||
| 51 | +EZ9999: ... there is an aivec error exception, core id is 0, error code = 259, | ||
| 52 | +dump info: pc start: 0x120041000000, current: 0x12004100030c, ... | ||
| 53 | +extend info: errcode:(259) errorStr: Illegal instruction, which is usually | ||
| 54 | +caused by unaligned UUB addresses. subErrType: 0x4. | ||
| 55 | +Kernel task happen error, retCode=0x31, [vector core exception]. | ||
| 56 | +``` | ||
| 57 | + | ||
| 58 | +### Iter 2 — explicit MTE3→MTE1 SetFlag / WaitFlag sync (replacing the barrier) | ||
| 59 | + | ||
| 60 | +Hypothesis: the fault was a missing cross-pipe ordering (UB→L1 goes on MTE3, L1→UB goes on MTE1), not the L1 access itself. | ||
| 61 | + | ||
| 62 | +- **Build**: OK. | ||
| 63 | +- **Runtime**: identical failure mode — error code 259, subErrType 0x4, same "Illegal instruction on aivec core". Current-PC shifted to `0x12004100031c` (the added SetFlag/WaitFlag extended the prologue by ~16 bytes); the trap still fires at the first `DataCopy` touching L1. | ||
| 64 | + | ||
| 65 | +Interpretation: the issue is not a pipe-sync bug. It is an instruction-level incompatibility — bisheng emits some opcode / operand that the AIV scalar/MTE unit rejects when the source or destination is an L1 (`A1`) LocalTensor rather than a VECIN/VECOUT UB tensor. Sync does not change which opcodes are emitted. | ||
| 66 | + | ||
| 67 | +## What this tells us | ||
| 68 | + | ||
| 69 | +1. **Syntactic acceptance ≠ runtime support.** Both `TBuf<TPosition::A1>` / `TPipe::InitBuffer(l1Buf_, bytes)` and the `DataCopy(L1_tensor, UB_tensor)` / `DataCopy(UB_tensor, L1_tensor)` overloads compile cleanly (no warnings), but the instruction stream they produce is illegal on an AIV-only kernel on this CANN 9.0.0 / bisheng 2026-03-21 build. | ||
| 70 | +2. **The generic `DataCopy` overload is not the UB↔L1 bridge for AIV on 351x.** The 351x public doc describes a hardware UB↔L1 path (MTE3 UB→L1, MTE1 L1→UB), and the AscendC API ref's `TPosition` enum does list A1 as a valid position, but the user-facing intrinsic to drive that hardware path from pure AIV code is either (a) not the generic `DataCopy(LocalTensor, LocalTensor)` template, (b) gated on a Cube/AIC context being present in the same kernel, or (c) not yet exposed in this CANN version. Probe cannot distinguish (a) / (b) / (c), but all three have the same consequence for Op#11. | ||
| 71 | +3. **No miscompile of pointer math / alignment that we caused.** The probe buffer is 4 KB, the `LocalTensor<float>` is 32 B-aligned by construction (`InitBuffer` with 4096 B), and the `DataCopy` count is 1024 elements which is a multiple of the fp32 block size (8). The "unaligned UUB" string in the error is misleading — it is the generic `errcode:259` message, not proof that our pointer was unaligned. | ||
| 72 | + | ||
| 73 | +## Recommendation to orchestrator (3 lines) | ||
| 74 | + | ||
| 75 | +1. **Do NOT pursue the Kind-2 rewrite of Op#11 DequantSwigluQuant that assumed `TBuf<TPosition::A1>` as UB spill.** On this CANN 9.0.0 / bisheng toolchain, the compiler silently accepts the construct but runtime fails — any kernel built on this assumption will crash at deploy, after wasting a worker's 5-iter budget. | ||
| 76 | +2. **Unblock Op#11's 0.54x via a different axis**: reduce per-tile UB footprint (smaller tile, fused buffer reuse, or fp16 intermediate storage with fp32 compute), or split the kernel into two launches so each pass has smaller UB. 源仓债项 DEBT-030（未迁移）should be closed as "A1 scratch not viable on this CANN" with this probe report as evidence. | ||
| 77 | +3. **Escalate upstream only if the 2x gain is critical**: file a ticket/question with CANN asking whether a dedicated UB↔L1 intrinsic (e.g. `Copy`, `CopyUbufToL1`, or an AIV-scope `LoadData` variant) exists on 351x and is not yet in the public API ref. Do not block Op#11 on that answer. | ||
| 78 | + | ||
| 79 | +## Evidence files | ||
| 80 | + | ||
| 81 | +- `workspace/probe_l1_scratch/kernel/probe_kernel.h` (with explicit MTE3→MTE1 sync, iter 2) | ||
| 82 | +- `workspace/probe_l1_scratch/kernel/probe_kernels.cpp` | ||
| 83 | +- `workspace/probe_l1_scratch/kernel/pybind11.cpp` | ||
| 84 | +- `workspace/probe_l1_scratch/kernel/probe_tiling.h` | ||
| 85 | +- `workspace/probe_l1_scratch/model_new_ascendc.py` | ||
| 86 | +- `/home/npu_user/workspace/AscendOpGenAgent/current_task/run_probe.py` (runner, copies also visible on A5 at `/root/AscendOpGenAgent/current_task/run_probe.py`) | ||
| 87 | +- Full stderr from both iters captured above. | ||
| 88 | + | ||
| 89 | +## 相关卡片 | ||
| 90 | + | ||
| 91 | +- [PB-16 `DataCopy(UB, TBuf<TPosition::A1>)` in pure-AIV kernel](../runbooks/compilation/pb_16_datacopy_ub_tbuf_tposition_a1_in_pure_aiv_kernel_s.md)：同一发现的 Runbook 处置卡（症状/规避/升级状态），其 Evidence 即本卡对应的 probe 报告；两卡互补，未合并。 | ||
| @@ -0,0 +1,219 @@ | |||
| 1 | +--- | ||
| 2 | +type: Concept | ||
| 3 | +title: "PROBE REPORT — V351 DataCopyPad UB→GM non-aligned tail" | ||
| 4 | +description: "Ascend950PR（V351，`__CCE_AICORE__ == 300`）实测（CANN 9.0.0 / bisheng 15.0.5）：DataCopyPad UB→GM 非 32B 对齐 blockLen（31/33/47/63 字节）精确写入请求字节数、不越界、不报错；V220 的 EC-23 规避不适用于 V351，OL-167 / P-P98 在 V351 成立。" | ||
| 5 | +tags: [task.implementation, topic.data_movement.data_copy, topic.alignment] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/runbooks/hardware/probe-2026-05-18-adhoc-datacopypad-v300-tail.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-07-12T16:00:00Z" | ||
| 11 | +updated_at: "2026-07-12T16:00:00Z" | ||
| 12 | +platforms: ["950"] | ||
| 13 | +aliases: [hw/2026-05-18_adhoc_datacopypad_v300_tail] | ||
| 14 | +--- | ||
| 15 | + | ||
| 16 | +# PROBE REPORT — V351 DataCopyPad UB→GM non-aligned tail | ||
| 17 | + | ||
| 18 | +> 原 OKF v1 卡号：**hw/2026-05-18_adhoc_datacopypad_v300_tail**（原 OKF v1 kind: guide；confidence: single_run）。 | ||
| 19 | + | ||
| 20 | +## Verdict | ||
| 21 | +ACCEPT_CORRECT | ||
| 22 | + | ||
| 23 | +**One-line summary**: On Ascend950PR (V351, `__CCE_AICORE__ == 300`) with | ||
| 24 | +CANN 9.0.0 + bisheng 15.0.5, `DataCopyPad` UB→GM with a non-32B-aligned | ||
| 25 | +byte-level `blockLen` (tested 31, 33, 47, 63) writes **exactly the | ||
| 26 | +requested number of bytes**, does **not overflow** into adjacent GM, and | ||
| 27 | +does **not raise** any runtime error. [EC-23](../runbooks/compilation/ec_23_datacopypad_ub_gm_crashes_on_ascend950pr_507035.md) (V220 crash) does NOT apply | ||
| 28 | +to V351. [OL-167](../runbooks/compilation/ol_167_datacopy_count_form_silently_truncates_non_32b_ali.md) / [P-P98](../optimizations/p_p98_non_aligned_tail_write_to_gm_via_datacopypad_dat.md) are confirmed valid on V351. | ||
| 29 | + | ||
| 30 | +## Environment | ||
| 31 | + | ||
| 32 | +``` | ||
| 33 | +- date: 2026-05-19T04:48Z | ||
| 34 | +- host: 198.51.100.35 (npu_dev3 container) | ||
| 35 | +- kernel: 6.6.110-42.13.tl4.x86_64 | ||
| 36 | +- OS: openEuler 22.03 (LTS-SP4) | ||
| 37 | +- CANN: /data/cann_b103/cann-9.0.0 (sourced via bin/setenv.bash) | ||
| 38 | +- bisheng (ccec): clang version 15.0.5 (clang-5c68a1cb1231 flang-5c68a1cb1231) | ||
| 39 | + build 2026-04-03T18:03:06+08:00 | ||
| 40 | + Target: x86_64-unknown-linux-gnu | ||
| 41 | +- npu-smi: 25.7.rc1 | ||
| 42 | +- torch: 2.7.1.post4.dev20260312 | ||
| 43 | +- torch_npu: shipped with /data/cann_b103/cann-9.0.0 | ||
| 44 | +- SOC_VERSION: Ascend950PR_9579 (V351, __CCE_AICORE__ == 300) | ||
| 45 | +- NPU 4 utilization before run: 0%; HBM 4568/114688 MB; no other processes | ||
| 46 | +- NPU 4 utilization after run: 0%; HBM 4780 MB; no orphaned processes | ||
| 47 | +``` | ||
| 48 | + | ||
| 49 | +## Question | ||
| 50 | + | ||
| 51 | +On V351/Ascend950PR with bisheng A5, does `DataCopyPad` UB→GM for a | ||
| 52 | +non-32B-aligned tail (blockLen=47 bytes from a fp16 UB) write cleanly, | ||
| 53 | +or does it crash like V220 (per [EC-23](../runbooks/compilation/ec_23_datacopypad_ub_gm_crashes_on_ascend950pr_507035.md) "DataCopyPad UB→GM crashes on | ||
| 54 | +V220 — use plain DataCopy + manual tail pad")? | ||
| 55 | + | ||
| 56 | +## Hypothesis vs Observation | ||
| 57 | + | ||
| 58 | +- **Hypothesis** (per 源仓任务简报 BRIEF.md，未迁移): Verdict (a) `ACCEPT_CORRECT` — V351 is | ||
| 59 | + newer than V220 and `DataCopyPad` is the documented primitive for | ||
| 60 | + non-aligned writes. | ||
| 61 | +- **Observation**: Confirmed. All four non-aligned byte counts (31, 33, | ||
| 62 | + 47, 63) wrote exactly `blockLen` bytes; the immediately adjacent guard | ||
| 63 | + bytes preserved the 0xCAFE sentinel intact; no `507035 / EZ9999 / | ||
| 64 | + EE9999` runtime errors; no kernel launch failure; NPU 4 remained | ||
| 65 | + idle and healthy after each invocation. | ||
| 66 | + | ||
| 67 | +## Evidence | ||
| 68 | + | ||
| 69 | +### Probe design | ||
| 70 | + | ||
| 71 | +- **GM buffer**: 256 bytes (128 × fp16), host-allocated then pre-filled | ||
| 72 | + with raw bit pattern `0xCAFE` per fp16 slot (little-endian bytes | ||
| 73 | + `0xFE 0xCA 0xFE 0xCA ...`) as overflow sentinel. | ||
| 74 | +- **UB scratch**: 64 × fp16 = 128 bytes; first 24 slots set to bit | ||
| 75 | + pattern `0xBEEF` (little-endian bytes `0xEF 0xBE 0xEF 0xBE ...`); the | ||
| 76 | + remaining 40 slots set to `0xDEAD` (bytes `0xAD 0xDE`) — these are | ||
| 77 | + the "would-overflow-if-bug" bytes we'd see leak into GM if the API | ||
| 78 | + rounded `blockLen` up to a 32B-aligned write. | ||
| 79 | +- **Kernel**: single AIV, `KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY)`, | ||
| 80 | + one call to `DataCopyPad(gmDst, ubLocal, {1, blockLen, 0, 0, 0})`. | ||
| 81 | +- **Host check** (after `torch.npu.synchronize()`): | ||
| 82 | + - bytes `[0, blockLen)` must equal the expected payload sliced from | ||
| 83 | + `(0xBEEF×24 || 0xDEAD×40)` little-endian byte stream | ||
| 84 | + - bytes `[blockLen, GM_BYTES)` must each remain at their pre-init | ||
| 85 | + sentinel (`0xFE` even index, `0xCA` odd index) | ||
| 86 | + - any Python-level exception or non-zero `aclrtlaunch_*` return code | ||
| 87 | + fails the run | ||
| 88 | + | ||
| 89 | +### Build | ||
| 90 | + | ||
| 91 | +`build_ascendc.py` configured CMake under | ||
| 92 | +`/root/probe_work/probe_datacopypad_v300_tail/kernel/build/` with | ||
| 93 | +`-DSOC_VERSION=Ascend950PR_9579 -DASCEND_CANN_PACKAGE_PATH=/data/cann_b103/cann-9.0.0`. | ||
| 94 | + | ||
| 95 | +- First attempt: compile error — `KERNEL_TASK_TYPE_DEFAULT` was | ||
| 96 | + mistakenly prefixed with `AscendC::` namespace; macro lives in | ||
| 97 | + global scope. One-line fix → unprefixed. | ||
| 98 | +- Second attempt: **clean build**. No warnings of substance. Library | ||
| 99 | + produced at | ||
| 100 | + `/root/probe_work/probe_datacopypad_v300_tail/kernel/build/_probe_datacopypad_ext.cpython-311-x86_64-linux-gnu.so`. | ||
| 101 | + | ||
| 102 | +### Runtime | ||
| 103 | + | ||
| 104 | +`run_results.json` (full, captured from container): | ||
| 105 | + | ||
| 106 | +```json | ||
| 107 | +[ | ||
| 108 | + { | ||
| 109 | + "blockLen": 47, | ||
| 110 | + "exception": null, | ||
| 111 | + "valid_match": true, | ||
| 112 | + "overflow_detected": false, | ||
| 113 | + "first_64_bytes_hex": "efbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefcafecafecafecafecafecafecafecafeca" | ||
| 114 | + }, | ||
| 115 | + { | ||
| 116 | + "blockLen": 33, | ||
| 117 | + "exception": null, | ||
| 118 | + "valid_match": true, | ||
| 119 | + "overflow_detected": false, | ||
| 120 | + "first_64_bytes_hex": "efbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefcafecafecafecafecafecafecafecafecafecafecafecafecafecafecafeca" | ||
| 121 | + }, | ||
| 122 | + { | ||
| 123 | + "blockLen": 63, | ||
| 124 | + "exception": null, | ||
| 125 | + "valid_match": true, | ||
| 126 | + "overflow_detected": false, | ||
| 127 | + "first_64_bytes_hex": "efbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeaddeaddeaddeaddeaddeaddeaddeadca" | ||
| 128 | + }, | ||
| 129 | + { | ||
| 130 | + "blockLen": 31, | ||
| 131 | + "exception": null, | ||
| 132 | + "valid_match": true, | ||
| 133 | + "overflow_detected": false, | ||
| 134 | + "first_64_bytes_hex": "efbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefbeefcafecafecafecafecafecafecafecafecafecafecafecafecafecafecafecafeca" | ||
| 135 | + } | ||
| 136 | +] | ||
| 137 | +``` | ||
| 138 | + | ||
| 139 | +### Byte-by-byte trace for blockLen=47 (the headline case) | ||
| 140 | + | ||
| 141 | +``` | ||
| 142 | +byte index: 0 1 2 3 ... 44 45 46 | 47 48 49 50 51 ... | ||
| 143 | +expected : EF BE EF BE ... EF BE EF | CA FE CA FE CA ... <- valid 47B then sentinel | ||
| 144 | +observed : ef be ef be ... ef be ef | ca fe ca fe ca ... | ||
| 145 | + ^^ | ||
| 146 | + first byte past the 47-byte tail is the SENTINEL | ||
| 147 | + (0xCA = high byte of 0xCAFE, odd index 47). | ||
| 148 | + If DataCopyPad had over-written to the next 32B | ||
| 149 | + boundary (byte 63 inclusive), we would instead | ||
| 150 | + see 'be' (low byte of 0xBEEF[23]) or 'ad/de' | ||
| 151 | + (0xDEAD for UB slots >=24) here. | ||
| 152 | +``` | ||
| 153 | + | ||
| 154 | +For blockLen=63 the pattern transitions from 0xBEEF region into the | ||
| 155 | +0xDEAD region within the valid window (consistent with UB layout — | ||
| 156 | +slot 24 onwards is 0xDEAD), then resumes 0xCAFE sentinel at byte 63: | ||
| 157 | +`...efbe[23rd 0xBEEF] addead addead ad de ad ca` — the final `de ad` | ||
| 158 | +pair is the UB[31] low/high bytes (bytes 62/63 of payload — wait, byte | ||
| 159 | +63 is past the 63-byte tail: index 63 is the 64th byte). Re-decoding | ||
| 160 | +the 63 case: | ||
| 161 | + | ||
| 162 | +``` | ||
| 163 | +bytes 0..45 : 0xBEEF × 23 elements = 'efbe' × 23 (46 bytes) | ||
| 164 | +byte 46 : low byte of UB[23] (0xBEEF) = 0xEF | ||
| 165 | +byte 47 : high byte of UB[23] (0xBEEF) = 0xBE | ||
| 166 | +bytes 48..61: UB[24..30] (0xDEAD × 7) = 'addead'×... actually | ||
| 167 | + fp16 size 2 → 14 bytes 'ad de ad de ad de ad de ad de ad de ad de' | ||
| 168 | +byte 62 : low byte of UB[31] (0xDEAD) = 0xAD | ||
| 169 | + ^ | ||
| 170 | + 63rd byte of write | ||
| 171 | +bytes 63.. : 0xCAFE sentinel, first byte at odd index 63 = 0xCA ✓ | ||
| 172 | +``` | ||
| 173 | + | ||
| 174 | +That precisely matches `...adde adde adde ad ca fe ca fe...` in the | ||
| 175 | +hex dump (low-nibble first per byte): tail of `...addeaddeaddeaddeadca` | ||
| 176 | +where the final `ad` is byte 62 (last byte written) and `ca` is byte 63 | ||
| 177 | +(first untouched sentinel byte). | ||
| 178 | + | ||
| 179 | +### Raw files | ||
| 180 | + | ||
| 181 | +- Kernel: `workspace/probe_datacopypad_v300_tail/kernel/probe_kernel.cpp` | ||
| 182 | +- Pybind: `workspace/probe_datacopypad_v300_tail/kernel/pybind11.cpp` | ||
| 183 | +- Host driver: `workspace/probe_datacopypad_v300_tail/run_probe.py` | ||
| 184 | +- Run log (full stdout/stderr): `workspace/probe_datacopypad_v300_tail/run.log` | ||
| 185 | +- Machine-readable: `workspace/probe_datacopypad_v300_tail/run_results.json` | ||
| 186 | + | ||
| 187 | +## Recommendation for orchestrator | ||
| 188 | + | ||
| 189 | +- **Amend [OL-167](../runbooks/compilation/ol_167_datacopy_count_form_silently_truncates_non_32b_ali.md) + [P-P98](../optimizations/p_p98_non_aligned_tail_write_to_gm_via_datacopypad_dat.md)** with `verified_on: Ascend950PR (V351, CANN | ||
| 190 | + 9.0.0, bisheng 15.0.5)` plus a pointer to this probe report. The | ||
| 191 | + byte-exact `blockLen` semantic is confirmed empirically — `DataCopyPad` | ||
| 192 | + IS the correct primitive for non-aligned UB→GM writes on V351. | ||
| 193 | +- **Confirm [EC-23](../runbooks/compilation/ec_23_datacopypad_ub_gm_crashes_on_ascend950pr_507035.md)'s scope as V220-only** (no widening needed). The | ||
| 194 | + EC-23 `applies_to: __CCE_AICORE__ == 220` tag is correct; do NOT | ||
| 195 | + propagate the V220 crash carve-out into V351 kernels. | ||
| 196 | +- **Unblock foreach_neg cold-start (源仓 task #22，未迁移)**: the `aog-kernel-worker` | ||
| 197 | + that cited EC-23 as a conflict for V351 was over-applying a | ||
| 198 | + V220-specific defect rule. Direct instruction: "use `DataCopyPad` | ||
| 199 | + with byte-level `blockLen` per P-P98; no pybind narrow+contiguous | ||
| 200 | + workaround required on V351". | ||
| 201 | +- **No new PB entry needed.** No new finalize-pipeline gate needed. | ||
| 202 | + | ||
| 203 | +## Caveats | ||
| 204 | + | ||
| 205 | +- Verified at `__CCE_AICORE__ == 300` only — not at 310/320/220 etc. | ||
| 206 | +- Verified for single-AIV launch, single-row `blockCount=1`, | ||
| 207 | + `srcStride=dstStride=0`. Pathological cases (`blockCount > 1` with | ||
| 208 | + small `blockLen`; sub-32B GM offsets via `gm_dst + odd_offset`) were | ||
| 209 | + NOT exercised by this probe — they remain unverified. | ||
| 210 | +- bisheng build date 2026-04-03; if bisheng <2026-04 is in use the | ||
| 211 | + result may not transfer (DataCopyPad codegen has changed across | ||
| 212 | + versions historically — e.g. EC-23 itself is a regression note). | ||
| 213 | +- The GM buffer was a fresh `torch::empty + view-as-int16-fill` | ||
| 214 | + allocation; pre-existing dirty GM scenarios were not tested. | ||
| 215 | +- No multi-stream / multi-AIV contention scenarios were tested. | ||
| 216 | + | ||
| 217 | +## 相关卡片 | ||
| 218 | + | ||
| 219 | +- [V351 DataCopyPad 尾 tile 按 32B 上取整多读](../runbooks/precision/v351_datacopypad_32b_roundup_oob_001.md)：相反方向（GM→UB 载入、行跨步寻址源、CANN 9.2.0）观测到的上取整多读污染；本卡结论仅覆盖 UB→GM 写出、连续 GM、CANN 9.0.0，两者方向与版本不同、互不矛盾，勿用本卡否定该载入侧缺陷。 | ||
| @@ -0,0 +1,12 @@ | |||
| 1 | +# 概念 | ||
| 2 | + | ||
| 3 | +Ascend C 算子开发的稳定概念、编程模型与决策框架。 | ||
| 4 | + | ||
| 5 | +- [Roofline Model for Ascend950PR (A5)](roofline_model.md)：Ascend950PR roofline 模型：VEC 峰值 fp16 约 56 TFLOPS、实测 CUBE 峰值 fp16 约 373 TFLOPS、HBM 1.5 TB/s；matmul 须按 CUBE 峰值定天花板。 | ||
| 6 | +- [SIMT vs SIMD Decision Framework](simt_vs_simd_decision.md)：SIMT 与 SIMD 选型决策框架：原子加/散写/间接寻址/group-local 计算选 SIMT，连续读加向量计算、批量搬运选 SIMD；含精度约束与四个实测案例。 | ||
| 7 | +- [PROBE REPORT — Q5 Sort/Reduce instruction cycles](sort_reduce_instruction_cycles.md)：950PR 实测：WholeReduceMax/BlockReduceMax 64×fp32 ≈100 cyc/call，MrgSort 成本随输出对数线性（≈0.55 cyc/pair + ~42 cyc 启动），宜单次长队列调用。 | ||
| 8 | +- [Q1 Probe Report: TBuf<TPosition::A1> in pure AIV kernel](aiv_ub_l1_datacopy_unsupported.md)：950PR 实测：纯 AIV kernel 的 UB↔L1 DataCopy 编译通过但运行时 aivec error 259；351x UB↔L1 硬通道未暴露给纯 AIV，L1-spill 改写不可行。 | ||
| 9 | +- [PROBE REPORT — Q_mte2_parallel (Q4): MTE2 parallel channels](mte2_parallel_channels.md)：950PR 实测：单 AIV 并发双 DataCopy 仅 ~10% 重叠收益（K2/K3=1.106×），双通道假设证伪；保持 1 DataCopy + 双缓冲计算重叠。 | ||
| 10 | +- [PROBE REPORT — Q3 Scalar broadcast sync-bypass (K_base vs K_muls_flex vs K_brcb)](scalar_broadcast_sync_bypass.md)：950PR 实测：逐行标量缩放用一次性 Brcb+宽 Mul 比 GetValue+Muls 快 25.3×；灵活标量 Muls 无同步旁路优势（假设证伪）。 | ||
| 11 | +- [PROBE REPORT — Q2 UB bank count / width / conflict behavior](ub_bank_conflict_behavior.md)：950PR UB stride 扫描：dense 比任意 strided 快 ~22%，mod-4 plateau ~6%，排除单端口 K=8/16 假设；bank 数/宽度未能反推（INCONCLUSIVE_PARTIAL）。 | ||
| 12 | +- [PROBE REPORT — V351 DataCopyPad UB→GM non-aligned tail](datacopypad_ub_to_gm_unaligned_tail.md)：950PR（V351）实测：DataCopyPad UB→GM 非 32B 对齐 blockLen 精确写入、不越界不报错；V220 的 EC-23 规避不适用于 V351。 | ||
| @@ -0,0 +1,116 @@ | |||
| 1 | +--- | ||
| 2 | +type: Concept | ||
| 3 | +title: "PROBE REPORT — Q_mte2_parallel (Q4): MTE2 parallel channels" | ||
| 4 | +description: "Ascend950PR 实测（CANN 9.0.0 / bisheng 2026-03-21）：单 AIV 并发两个 DataCopy 仅 ~10% 重叠收益（K2/K3=1.106×），“双 MTE2 通道打满”模式证伪（SHARED_CHANNEL）；应保持 1 DataCopy + 双缓冲 TQue 计算重叠。" | ||
| 5 | +tags: [task.design, topic.data_movement, topic.performance.analysis] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/runbooks/hardware/probe-2026-04-21-q-mte2-parallel.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-07-12T16:00:00Z" | ||
| 11 | +updated_at: "2026-07-12T16:00:00Z" | ||
| 12 | +platforms: ["950"] | ||
| 13 | +aliases: [hw/2026-04-21_Q_mte2_parallel] | ||
| 14 | +--- | ||
| 15 | + | ||
| 16 | +# PROBE REPORT — Q_mte2_parallel (Q4): MTE2 parallel channels | ||
| 17 | + | ||
| 18 | +> 原 OKF v1 卡号：**hw/2026-04-21_Q_mte2_parallel**（原 OKF v1 kind: guide；confidence: single_run）。 | ||
| 19 | + | ||
| 20 | +## Verdict | ||
| 21 | + | ||
| 22 | +**ACCEPT_CORRECT** — all 3 kernels compile and execute bit-exact. | ||
| 23 | + | ||
| 24 | +**MTE2 interpretation: SHARED_CHANNEL with marginal overlap** (ratio K2/K3 = 1.106×; K3/K1 = 1.761×). | ||
| 25 | + | ||
| 26 | +**One-line summary**: Two concurrent `DataCopyPad`/`DataCopy` issues from a single AIV do NOT achieve the ≥1.5× dual-channel threshold the template hypothesised; there is only ~10% arbitration/prefetch benefit from issuing back-to-back vs serialized. The pattern "schedule 2 concurrent DataCopy on critical path for 2× MTE2 throughput" is NOT valid on this CANN/bisheng build — schedule `1 DataCopy + compute overlap` instead. | ||
| 27 | + | ||
| 28 | +## Environment | ||
| 29 | + | ||
| 30 | +- **date**: 2026-04-22T07:50Z | ||
| 31 | +- **host**: npu_dev3 container on A5 (198.51.100.35) | ||
| 32 | +- **NPU used**: NPU 1 (AICore 0%, idle before+after; HBM 5235/131072 MB; only NPU 0 had pytest traffic) | ||
| 33 | +- **OS (container)**: openEuler 22.03 (LTS-SP4) | ||
| 34 | +- **SOC**: Ascend950PR (per npu-smi), built with `-DSOC_VERSION=Ascend950PR_9589` | ||
| 35 | +- **npu-smi**: 25.7.rc1.b087 | ||
| 36 | +- **CANN**: 9.0.0, innerversion `V100R001C10SPC001B218`, path `/usr/local/Ascend/cann-9.0.0` | ||
| 37 | +- **bisheng/ccec**: `2026-03-21T17:07:34+08:00 clang version 15.0.5 (clang-5c68a1cb1231 flang-5c68a1cb1231)` | ||
| 38 | +- **python / torch**: Python 3.11.13 at `/root/python3.11.13`, torch import works under CANN `set_env.sh` + `LD_LIBRARY_PATH` prepending `cann-9.0.0/x86_64-linux/lib64` (libhccl.so path) | ||
| 39 | + | ||
| 40 | +## Question | ||
| 41 | + | ||
| 42 | +Do two concurrent `DataCopy` (GM→UB) operations issued from a single AIV saturate 2× read bandwidth (documented as "per-AICore AXI 2×128B read + 1×128B write"), or do they share one channel? | ||
| 43 | + | ||
| 44 | +## Hypothesis vs Observation | ||
| 45 | + | ||
| 46 | +- **Hypothesis** (template): partial saturation, ~1.6×, not full 2×. | ||
| 47 | +- **Observation**: K2/K3 = **1.106×**. K3 (two concurrent) is only ~10% faster than K2 (two sequential) for the same total bytes moved. K3/K1 = **1.761×** (K3 moves 2× the bytes of K1 in 1.76× the time, i.e., ~12% savings versus hypothetical full serialization). | ||
| 48 | +- The observed parallel-overlap benefit (~10%) is well below the 1.5× "dual-channel" and 1.2× "partial-parallel" thresholds; it falls in the **SHARED_CHANNEL** bin per the template's verdict rules. | ||
| 49 | + | ||
| 50 | +Interpretation: either (a) bisheng serialises the two `DataCopy` issues despite back-to-back placement, or (b) the hardware has one effective MTE2 issue slot per AIV so the documented "2 read channels per AICore" are split across the AIC + paired AIV rather than available within one AIV. The 10% overlap is consistent with instruction-pipeline / outstanding-request overlap (one copy's tail overlapping the other's head), not with true dual-channel saturation. | ||
| 51 | + | ||
| 52 | +## Evidence | ||
| 53 | + | ||
| 54 | +### Kernels (3 variants, all AIV-only, single block) | ||
| 55 | + | ||
| 56 | +- **K1**: `REPEAT_N=64` iterations of: 1× `DataCopy(GM→UB, 64KB)` + MTE2→MTE3 flag + 1× `DataCopy(UB→GM, 64KB)` + MTE3→MTE2 flag. | ||
| 57 | +- **K2**: same REPEAT_N, but each iteration does **two sequential halves** (first half read+write, then second half read+write) — enforces serialization via explicit MTE2↔MTE3 sync. | ||
| 58 | +- **K3**: same REPEAT_N, but each iteration **issues two reads back-to-back** into two separate `TBuf<VECCALC>` regions (64KB each) before one joint MTE2→MTE3 sync, then **two writes back-to-back** + joint sync. | ||
| 59 | + | ||
| 60 | +Per-buffer size 64 KB (2 × 64 KB = 128 KB fits within AIV UB). Loop factor 64× chosen so total DMA work (4 MB for K1, 8 MB for K2/K3) dominates kernel-launch overhead. | ||
| 61 | + | ||
| 62 | +### Build | ||
| 63 | + | ||
| 64 | +`/root/AscendOpGenAgent/utils/build_ascendc.py current_task -v Ascend950PR_9589 --clean` — build succeeded with no kernel-level warnings. (Clock-skew warnings unrelated; they come from tar-preserved future mtimes.) | ||
| 65 | + | ||
| 66 | +``` | ||
| 67 | +[100%] Built target pybind11_lib | ||
| 68 | +[build_ascendc] Build completed: /root/AscendOpGenAgent/current_task/kernel/build | ||
| 69 | +``` | ||
| 70 | + | ||
| 71 | +Full log in `workspace/probe_mte2_parallel/build.log`. | ||
| 72 | + | ||
| 73 | +### Runtime (on NPU 1, warmup=10, iters=30, aclrtEvent timing) | ||
| 74 | + | ||
| 75 | +| Kernel | Bytes moved / launch (total read+write) | Median (us) | Min (us) | Stdev (us) | | ||
| 76 | +|---|---|---|---|---| | ||
| 77 | +| K1 (1× copy ×64) | 8 MB (4 MB rd + 4 MB wr) | 58.91 | 58.54 | 0.39 | | ||
| 78 | +| K2 (2× seq ×64) | 16 MB (8 MB rd + 8 MB wr) | 114.74 | 114.28 | 0.59 | | ||
| 79 | +| K3 (2× concurrent ×64) | 16 MB | 103.76 | 103.31 | 0.42 | | ||
| 80 | + | ||
| 81 | +**Stdev under 1%** on all three — measurement noise is negligible; NPU 1 was idle throughout (re-checked post-run). | ||
| 82 | + | ||
| 83 | +Ratios: | ||
| 84 | +- **K2 / K3 = 1.106×** → only ~10% gain from issuing two copies concurrently vs serializing them. **Falls in SHARED_CHANNEL bin per template** (<1.2×). | ||
| 85 | +- **K3 / K1 = 1.761×** → K3 moves 2× the bytes of K1; ratio of 1.76× (not 2.0×) shows ~12% overlap between the two concurrent transfers, not full parallelism. | ||
| 86 | +- **K2 / K1 = 1.948×** → K2's two sequential copies scale almost exactly 2× (as expected for serialized DMAs). | ||
| 87 | + | ||
| 88 | +### Raw files | ||
| 89 | + | ||
| 90 | +- Kernel: `workspace/probe_mte2_parallel/kernel/probe_kernel.h`, `probe_kernels.cpp`, `probe_tiling.h`, `pybind11.cpp` | ||
| 91 | +- Runner: `workspace/probe_mte2_parallel/run_probe.py`, `model_new_ascendc.py` | ||
| 92 | +- Build log: `workspace/probe_mte2_parallel/build.log` | ||
| 93 | +- Run log: `workspace/probe_mte2_parallel/run.log` | ||
| 94 | +- Per-iter timings: `workspace/probe_mte2_parallel/timings.csv` (30 rows, all 3 kernels) | ||
| 95 | +- On A5: `/root/AscendOpGenAgent/current_task/` (kernel build artifacts) | ||
| 96 | + | ||
| 97 | +## Recommendation for orchestrator | ||
| 98 | + | ||
| 99 | +1. **Do NOT promote "schedule two concurrent DataCopy to saturate dual MTE2 channels" as a KB performance pattern for Ascend950PR single-AIV kernels.** The 10% benefit is too small to justify the added register pressure / UB fragmentation / sync complexity. Flag to [target_ascend950pr](../../../common/platforms/concepts/target_ascend950pr.md) §Memory as "MTE2 parallel channels: observed only ~10% arbitration overlap within a single AIV (probe 2026-04-22, K2/K3=1.106×). Pattern 'concurrent DataCopy for 2× MTE2' is FALSIFIED on CANN 9.0.0 / bisheng 2026-03-21." | ||
| 100 | +2. **Keep "1 DataCopy + compute overlap via double-buffered TQue"** as the canonical UB-bound optimization. The K2 vs K1 ratio (1.948×) confirms MTE2 does not magically scale — halving load is the actual win. | ||
| 101 | +3. **One open question** worth a follow-up probe: is the dual-channel possibly available **across the AIC + AIV pair** (not within a single AIV)? Would need a mixed Cube+Vector kernel to test. Low priority — most AIV-only kernels can't exploit that anyway. | ||
| 102 | + | ||
| 103 | +## Caveats | ||
| 104 | + | ||
| 105 | +- **Buffer size**: 64 KB per buffer was chosen to fit 2 in UB. Larger concurrent transfers were not testable (2× 256 KB would exceed AIV UB). The ~10% overlap measured at 64 KB may be different at smaller sizes (where launch overhead dominates) or if UB were larger. The probe does NOT distinguish "shared channel" from "dual channel gated on larger transfer granularity". | ||
| 106 | +- **REPEAT_N=64** reuses the same GM region each iteration, so HBM row buffer / L2-like caches may be warm. True cold-HBM behavior per iteration would need non-overlapping GM strides; not tested here. This does NOT invalidate the ratio (K2 and K3 have identical access pattern, both benefit equally from any caching). | ||
| 107 | +- **CANN version sensitivity**: result may change with a future CANN / bisheng release — record re-probe on any bisheng upgrade. | ||
| 108 | +- **Cross-AIV**: within a single AICore there are 2 AIVs sharing the AXI; this probe uses only 1 block (1 AIV). Saturation across 2 AIVs (2-block launch) was not tested and may behave differently. | ||
| 109 | +- **Methodology lesson applied**: the [Q1 L1-scratch finding](aiv_ub_l1_datacopy_unsupported.md) warned "build success + no warnings ≠ runtime correctness". We hit exactly that trap at iter 1 — two earlier attempts had a live kernel crash (K3 write OOB) and silent wrong output (buffer-sizing mismatch). Only the repeat-loop + matched Python/kernel size convention gave bit-exact output on all 3 kernels. | ||
| 110 | + | ||
| 111 | +## Iterations run (4 iters; 1 compile, 3 runtime-semantic) | ||
| 112 | + | ||
| 113 | +1. **Iter 1** — initial kernels with 256 KB per buffer + 2 TQues for K3. Build OK; K3 crashed with `aivec error 263 "MTE write address out of range"`; root cause: 2× 256 KB exceeded AIV UB (192 KB). | ||
| 114 | +2. **Iter 2** — refactored K3 to use TBufs (raw VECCALC) instead of TQues. Build OK; K3 crashed with identical error. Confirmed the issue was UB size, not TQue semantics. | ||
| 115 | +3. **Iter 3** — shrunk buffers to 64 KB each (2× = 128 KB, fits in UB); fixed Python-side `HALF` constant to match kernel's `PROBE_ELEMS_HALF`. All kernels ran, K1 and K3 correct, K2 had one half zero (likely off-by-one in double-blocked TQue, not the question's focus). Timings 4–5 us each → launch overhead dominated → no conclusive signal. | ||
| 116 | +4. **Iter 4** — loop-amplified each kernel 64× inside a single launch; replaced TQues with TBuf+explicit MTE2↔MTE3 flags (fixes K2's correctness issue from iter 3). All 3 bit-exact; per-launch time 58–115 us; stdev <1%. **Final data reported above.** | ||
| @@ -0,0 +1,118 @@ | |||
| 1 | +--- | ||
| 2 | +type: Concept | ||
| 3 | +title: "Roofline Model for Ascend950PR (A5)" | ||
| 4 | +description: "Ascend950PR (A5) roofline 模型：56 个 AI Vector Core、VEC 峰值 fp16 约 56 TFLOPS、实测 CUBE 峰值 fp16 约 373 TFLOPS、HBM 1.5 TB/s、UB 每核 248KB；matmul 须按 CUBE 峰值定天花板。" | ||
| 5 | +tags: [task.optimization, topic.performance.analysis, topic.architecture] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/reference/porter/handbook/roofline_model.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-08-05T14:23:52Z" | ||
| 11 | +updated_at: "2026-09-04T23:59:25Z" | ||
| 12 | +platforms: ["950"] | ||
| 13 | +--- | ||
| 14 | + | ||
| 15 | +# Roofline Model for Ascend950PR (A5) | ||
| 16 | + | ||
| 17 | +Use this model to determine if a kernel is compute-bound or memory-bound, and to set realistic performance targets. | ||
| 18 | + | ||
| 19 | +## Hardware Parameters | ||
| 20 | + | ||
| 21 | +| Parameter | Value | Source | | ||
| 22 | +|-----------|-------|--------| | ||
| 23 | +| AI Vector Cores | 56 | [target_ascend950pr](../../../common/platforms/concepts/target_ascend950pr.md) | | ||
| 24 | +| VEC Peak (FP32) | ~28 TFLOPS | 56 cores × 512 FLOPS/cycle × 1GHz | | ||
| 25 | +| VEC Peak (FP16) | ~56 TFLOPS | 2x FP32 (half-precision packing) | | ||
| 26 | +| **CUBE Peak (FP16)** | **~373 TFLOPS** | **Measured 2026-06-12, .171 NPU1 957b, torch.matmul 8192³** | | ||
| 27 | +| **CUBE Peak (BF16)** | **~368 TFLOPS** | **Measured (same run)** | | ||
| 28 | +| **CUBE Peak (FP32)** | **~24 TFLOPS** | **Measured — cube is fp16/bf16-optimized, fp32 not favored** | | ||
| 29 | +| HBM Bandwidth | 1.5 TB/s | Measured (theoretical ~2 TB/s) | | ||
| 30 | +| UB per Core | **248 KB usable** (256 KB phys) | `GetCoreMemSize(UB)`=253952 (262144 − 8KB framework reserve); cannbot npu-arch §2.5 canonical. SIMT-effective ≈208KB (−40KB DCache). Never hardcode — 见 [target_ascend950pr](../../../common/platforms/concepts/target_ascend950pr.md) | | ||
| 31 | +| L2 Cache | 48 MB shared | Read-only cache | | ||
| 32 | + | ||
| 33 | +> ⚠ **CUBE vs VEC peak — do not conflate.** matmul / FlashAttention / any | ||
| 34 | +> op dominated by `Matmul`/`Mmad` runs on the **CUBE unit** (peak ~373 TFLOPS | ||
| 35 | +> fp16), NOT the vector unit (~56 TFLOPS fp16). Using the VEC peak as the | ||
| 36 | +> ceiling for a cube-bound op understates it ~6.7× → an op at 24% of its real | ||
| 37 | +> cube ceiling looks like ~80%+ "efficient" and the perf gate wrongly stops | ||
| 38 | +> optimizing. Elementwise / reduction / softmax are vector-unit ops → use the | ||
| 39 | +> VEC peak. 源仓的 `roofline_eval.py` 按 `peak_cube_*` 字段 + `_peak_tflops(op_type)` | ||
| 40 | +> 编码了这个拆分，ridge 按单元分别计算（见下）。 | ||
| 41 | + | ||
| 42 | +## Operational Intensity Calculation | ||
| 43 | + | ||
| 44 | +``` | ||
| 45 | +OI = FLOPs / Bytes_transferred (FLOP/byte) | ||
| 46 | + | ||
| 47 | +For SG Forward (per token, top_k experts): | ||
| 48 | + FLOPs = 2 * top_k * hdim (multiply + add per element per expert) | ||
| 49 | + Bytes = top_k * hdim * sizeof(T) (load expert) + hdim * sizeof(T) (store output) | ||
| 50 | + OI = 2 * top_k * hdim / ((top_k + 1) * hdim * sizeof(T)) | ||
| 51 | + = 2 * top_k / ((top_k + 1) * sizeof(T)) | ||
| 52 | + | ||
| 53 | + fp32, K=4: OI = 8 / (5*4) = 0.4 FLOP/byte → memory-bound | ||
| 54 | + fp16, K=4: OI = 8 / (5*2) = 0.8 FLOP/byte → memory-bound | ||
| 55 | +``` | ||
| 56 | + | ||
| 57 | +## Ridge Point | ||
| 58 | + | ||
| 59 | +``` | ||
| 60 | +Ridge OI = Peak_FLOPS / Peak_BW | ||
| 61 | + | ||
| 62 | +VEC-bound ops (elementwise / reduction / softmax): | ||
| 63 | + fp32: 28 TFLOPS / 1.5 TB/s ≈ 18.7 FLOP/byte | ||
| 64 | + fp16: 56 TFLOPS / 1.5 TB/s ≈ 37.3 FLOP/byte | ||
| 65 | + | ||
| 66 | +CUBE-bound ops (matmul / attention): | ||
| 67 | + fp16: 373 TFLOPS / 1.5 TB/s ≈ 248.7 FLOP/byte | ||
| 68 | + fp32: 24 TFLOPS / 1.5 TB/s ≈ 16.0 FLOP/byte | ||
| 69 | +``` | ||
| 70 | + | ||
| 71 | +The ridge for a cube-bound op is ~6.7× higher than the vec ridge — many | ||
| 72 | +matmul/FA shapes that look "compute-bound" against the vec ridge are still | ||
| 73 | +well below the cube ridge (i.e. memory-bound on the cube unit). | ||
| 74 | + | ||
| 75 | +SG kernels (OI < 1) are **deeply memory-bound**. Optimizations should focus on: | ||
| 76 | +1. Reducing GM reads (data reuse, loop reorder) | ||
| 77 | +2. Maximizing MTE2/VEC overlap (TQue prefetch) | ||
| 78 | +3. NOT on compute optimizations (VEC is not the bottleneck) | ||
| 79 | + | ||
| 80 | +## Theoretical Performance Bounds | ||
| 81 | + | ||
| 82 | +### SG Forward (Memory-Bound) | ||
| 83 | + | ||
| 84 | +``` | ||
| 85 | +min_time = total_bytes / bandwidth | ||
| 86 | +total_bytes = T * K * hdim * sizeof(T) # expert reads (dominant) | ||
| 87 | + + T * hdim * sizeof(T) # output writes | ||
| 88 | + | ||
| 89 | +Example: prod_a (T=8192, K=4, H=256, fp32): | ||
| 90 | + read_bytes = 8192 * 4 * 256 * 4 = 32 MB | ||
| 91 | + write_bytes = 8192 * 256 * 4 = 8 MB | ||
| 92 | + total = 40 MB | ||
| 93 | + min_time = 40 MB / 1.5 TB/s = 0.027 ms | ||
| 94 | + | ||
| 95 | + With expert reuse (expert-major, 64 experts): | ||
| 96 | + read_bytes = 64 * 256 * 4 = 64 KB (experts) + 8192 * 4 * 4 = 128 KB (indices) | ||
| 97 | + total ≈ 8.2 MB (output + indices + expert data) | ||
| 98 | + min_time = 8.2 MB / 1.5 TB/s = 0.005 ms | ||
| 99 | +``` | ||
| 100 | + | ||
| 101 | +## Using the Model | ||
| 102 | + | ||
| 103 | +### Before Exploration | ||
| 104 | +1. Calculate OI for the kernel | ||
| 105 | +2. Determine if compute-bound or memory-bound | ||
| 106 | +3. Calculate theoretical min_time | ||
| 107 | +4. Compare with actual time → efficiency = theoretical / actual | ||
| 108 | + | ||
| 109 | +### Setting Targets | ||
| 110 | +- If efficiency < 30%: significant optimization opportunity exists | ||
| 111 | +- If efficiency 30-60%: moderate improvements possible | ||
| 112 | +- If efficiency > 60%: close to hardware limits, diminishing returns | ||
| 113 | +- If efficiency > 80%: near optimal, explore algorithmic changes only | ||
| 114 | + | ||
| 115 | +### Early Stopping | ||
| 116 | +If after an optimization: | ||
| 117 | +- The kernel is >60% of theoretical peak → consider stopping | ||
| 118 | +- Further improvement requires >2x reduction in data movement → may need algorithm redesign (escalate to human) | ||
| @@ -0,0 +1,129 @@ | |||
| 1 | +--- | ||
| 2 | +type: Concept | ||
| 3 | +title: "PROBE REPORT — Q3 Scalar broadcast sync-bypass (K_base vs K_muls_flex vs K_brcb)" | ||
| 4 | +description: "Ascend950PR 实测（CANN 9.0.0 / bisheng 15.0.5）：多行逐行标量缩放用一次性 Brcb + 单次宽 Mul 比 GetValue+Muls 快 25.3×；灵活标量位置 Muls 无同步旁路优势（0.974×，假设证伪）。" | ||
| 5 | +tags: [task.implementation, topic.programming_model.execution_unit.scalar, topic.performance.analysis] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/runbooks/hardware/probe-2026-04-21-q-scalar-broadcast.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-07-12T16:00:00Z" | ||
| 11 | +updated_at: "2026-07-12T16:00:00Z" | ||
| 12 | +platforms: ["950"] | ||
| 13 | +aliases: [hw/2026-04-21_Q_scalar_broadcast] | ||
| 14 | +--- | ||
| 15 | + | ||
| 16 | +# PROBE REPORT — Q3 Scalar broadcast sync-bypass (K_base vs K_muls_flex vs K_brcb) | ||
| 17 | + | ||
| 18 | +> 原 OKF v1 卡号：**hw/2026-04-21_Q_scalar_broadcast**（原 OKF v1 kind: guide；confidence: single_run）。 | ||
| 19 | + | ||
| 20 | +## Verdict | ||
| 21 | + | ||
| 22 | +**ACCEPT_CORRECT** — all 3 kernels compile and execute bit-exactly (within fp32 rounding tolerance over 512 cumulative multiplies; rel error 1.21e-5). | ||
| 23 | + | ||
| 24 | +**One-line summary**: `K_muls_flex/K_base = 0.974x` (flexible scalar form is **NOT** faster than baseline scalar Muls — hypothesis A **FALSIFIED**); `K_brcb/K_base = 25.3x` (one-time Brcb + single vector Mul **dominates** both scalar variants by a wide margin). Best variant for multi-row scalar-per-row scaling = **K_brcb**. | ||
| 25 | + | ||
| 26 | +## Environment | ||
| 27 | + | ||
| 28 | +- **date**: 2026-04-22T00:25Z | ||
| 29 | +- **host**: A5 (198.51.100.35), container `npu_dev3` | ||
| 30 | +- **NPU used**: NPU 1 (AICore 0% before and after run; HBM 5236/131072 MB; NPU 0 had ongoing pytest traffic — avoided) | ||
| 31 | +- **OS (container)**: openEuler 22.03 (LTS-SP4) | ||
| 32 | +- **SOC**: Ascend950PR (per npu-smi), built with `-DSOC_VERSION=Ascend950PR_9589` (→ `__NPU_ARCH__ == 3510`, which is the arch gate for the `Muls (灵活标量位置)` / flexible-scalar overload) | ||
| 33 | +- **npu-smi**: 25.7.rc1.b087 | ||
| 34 | +- **CANN**: 9.0.0, path `/usr/local/Ascend/cann-9.0.0` | ||
| 35 | +- **bisheng**: `clang version 15.0.5 (clang-5c68a1cb1231 flang-5c68a1cb1231)` (installed 2025-09-23, used inside container — matches Q4's toolchain) | ||
| 36 | +- **python / torch**: Python 3.11.13 at `/root/python3.11.13`; torch import succeeds after sourcing `set_env.sh` + prepending `cann-9.0.0/x86_64-linux/lib64` to `LD_LIBRARY_PATH` | ||
| 37 | + | ||
| 38 | +## Question | ||
| 39 | + | ||
| 40 | +Given the `GetValue(k_ub) → Muls(work, k_scalar, N)` pattern in op#11 serializes on the Scalar pipe (MTE2_S → S_V → V_S → S_V sync chain), which of these bypasses that serialization? | ||
| 41 | + | ||
| 42 | +- **K_base**: `GetValue(k_ub[r]) → Muls(work[r], work[r], k_scalar, H)` per row (current op#11 pattern) | ||
| 43 | +- **K_muls_flex**: `Muls(work[r], work[r], k_ub[r], H)` per row — flexible-scalar overload where `src1` is a LocalTensor slice in UB instead of a scalar variable (no explicit `GetValue` in source) | ||
| 44 | +- **K_brcb**: once-off `Brcb(k_full, k_ub, N_ROWS/8, {1, 8})` expands `k_ub[0..63]` into a 512-element tensor (each k replicated to 8 fp32 = one 32 B block), then a single `Mul(work, work, k_full, 512)` | ||
| 45 | + | ||
| 46 | +## Hypothesis vs Observation | ||
| 47 | + | ||
| 48 | +- **Hypothesis A**: `K_muls_flex ≥ 1.3× K_base` (flexible scalar bypasses S_V sync because src comes from LocalTensor rather than scalar register). | ||
| 49 | + - **Observation**: `K_muls_flex / K_base = 0.974×` → K_muls_flex is **~2.6% slower** than K_base. **Hypothesis FALSIFIED** on CANN 9.0.0 + bisheng 15.0.5. The compiler/runtime does not materialise any win from the flexible-scalar overload in this shape; the per-row 64× sequential Muls calls pay essentially the same sync cost either way (likely because the flex overload's single-point LocalTensor read still flows through the Scalar pipe via a different opcode that preserves the chain). | ||
| 50 | +- **Hypothesis B**: `K_brcb` similar or better for multi-row scenarios. | ||
| 51 | + - **Observation**: `K_brcb / K_base = 25.3×`. Not merely "similar" — dominant. The Brcb+single-Mul pattern eliminates 64 per-row issue-and-sync pairs entirely and replaces them with one tensor broadcast (handled by the VEC engine in cycles measured in the low hundreds) + one wide Mul over 512 elements (one VEC instruction of 8 repeats). **Hypothesis CONFIRMED and then some.** | ||
| 52 | + | ||
| 53 | +## Evidence | ||
| 54 | + | ||
| 55 | +### Kernels (3 variants, all AIV-only, single block; see `kernel/probe_kernel.h`) | ||
| 56 | + | ||
| 57 | +Scenario: `N_ROWS=64, H=8 fp32 per row, TOTAL=512 fp32 (2 KB)`. H=8 chosen to match Brcb's native 32B-block output width (each k replicated to exactly one block). Inner work amplified by outer `REPEAT_N=512` so each launch does 64×512=32,768 scalar Muls (K_base/K_muls_flex) or 512 Brcb + 512 wide Mul (K_brcb). | ||
| 58 | + | ||
| 59 | +- **K_base**: outer REPEAT_N × inner {for r in 0..63: `k = kT.GetValue(r); Muls(work[r*8], work[r*8], k, 8);`} | ||
| 60 | +- **K_muls_flex**: outer REPEAT_N × inner {for r in 0..63: `Muls(work[r*8], work[r*8], kT[r], 8);`} — flexible scalar via single-point LocalTensor slice (exercises the `Muls<U,S,V>(dst, src0, src1, count)` overload gated on `__NPU_ARCH__ == 3510 || 5102` in `kernel_operator_vec_binary_scalar_intf.h`) | ||
| 61 | +- **K_brcb**: outer REPEAT_N × inner {`Brcb(kFull, kT, /*repeatTime=*/8, BrcbRepeatParams(1, 8));` + `Mul(work, work, kFull, 512);`} | ||
| 62 | + | ||
| 63 | +All three take the same GM inputs `work_gm[512]` and `k_gm[64]`, and all three write the same output `out_gm[512] = work[r*8:(r+1)*8] * k[r]^REPEAT_N`. Host sets `work=1` and `k[r] = 1 + (r-32)*5e-6` so `k^512 ∈ [0.919, 1.082]` — stays well inside fp32 normal range. | ||
| 64 | + | ||
| 65 | +### Build | ||
| 66 | + | ||
| 67 | +Command: `python3 utils/build_ascendc.py current_task -v Ascend950PR_9589 --clean` (with `source /usr/local/Ascend/cann-9.0.0/set_env.sh` + `LD_LIBRARY_PATH` prepending `cann-9.0.0/x86_64-linux/lib64`). | ||
| 68 | + | ||
| 69 | +All three kernels compile clean. Key lines from `build.log`: | ||
| 70 | + | ||
| 71 | +``` | ||
| 72 | +[ 95%] Linking CXX static library lib/libkernels.a | ||
| 73 | +[100%] Built target kernels | ||
| 74 | +[ 97%] Building CXX object CMakeFiles/pybind11_lib.dir/.../pybind11.cpp.o | ||
| 75 | +[ 22%] Linking CXX shared library _probe_ext.cpython-311-x86_64-linux-gnu.so | ||
| 76 | +[ 25%] Built target pybind11_lib | ||
| 77 | +[build_ascendc] Build completed | ||
| 78 | +``` | ||
| 79 | + | ||
| 80 | +**No kernel-level warnings.** The flexible-scalar Muls overload compiled without any arch-gate error, confirming `Ascend950PR_9589` does activate `__NPU_ARCH__ == 3510` and unlocks that overload. Clock-skew warnings (tar-preserved future mtimes) are unrelated; worked around by `find … -exec touch +` before cmake. | ||
| 81 | + | ||
| 82 | +### Runtime (on NPU 1, WARMUP=10, ITERS=25, aclrtEvent timing) | ||
| 83 | + | ||
| 84 | +| Kernel | Work per launch | Median (us) | Min (us) | Stdev (us) | | ||
| 85 | +|---|---|---|---|---| | ||
| 86 | +| **K_base** (per-row `GetValue+Muls`, 64×512=32768 scalar Muls) | 32,768 scalar Muls | **758.53** | 757.79 | 0.66 | | ||
| 87 | +| **K_muls_flex** (per-row flex-scalar Muls, same 32768 calls) | 32,768 flex-scalar Muls | **778.52** | 777.92 | 0.51 | | ||
| 88 | +| **K_brcb** (512× Brcb + 512× single wide Mul) | 512 Brcb + 512 wide Mul | **29.95** | 29.60 | 0.31 | | ||
| 89 | + | ||
| 90 | +Stdev < 1 us across all three — measurement noise negligible; NPU 1 idle pre- and post-run (verified). | ||
| 91 | + | ||
| 92 | +Precision (after 512 cumulative multiplies): all three output bit-exactly match the CPU reference `work * k^512` within `rel_err = 1.21e-5` (cumulative rounding bound is ≈ `512 · 1.19e-7 = 6.1e-5` so all three are at 1/5 of the rigorously allowed fp32 error, i.e. identical rounding behavior). | ||
| 93 | + | ||
| 94 | +Ratios: | ||
| 95 | +- **K_base / K_muls_flex = 0.974×** (K_base is 2.6% FASTER than the flex-scalar variant → hypothesis A falsified) | ||
| 96 | +- **K_base / K_brcb = 25.32×** (K_brcb is 25× faster than the baseline) | ||
| 97 | +- **K_muls_flex / K_brcb = 26.0×** | ||
| 98 | + | ||
| 99 | +### Raw files | ||
| 100 | + | ||
| 101 | +- Kernels: `workspace/probe_scalar_broadcast/kernel/probe_kernel.h`, `probe_kernels.cpp`, `probe_tiling.h`, `pybind11.cpp` | ||
| 102 | +- Runner: `workspace/probe_scalar_broadcast/run_probe.py`, `model_new_ascendc.py` | ||
| 103 | +- Build log: `workspace/probe_scalar_broadcast/build.log` | ||
| 104 | +- Run log: `workspace/probe_scalar_broadcast/run.log` | ||
| 105 | +- Per-iter timings: `workspace/probe_scalar_broadcast/timings.csv` (25 rows × 3 columns) | ||
| 106 | +- On A5: `/root/AscendOpGenAgent/current_task/` (kernel build artifacts) | ||
| 107 | + | ||
| 108 | +## Recommendation for orchestrator | ||
| 109 | + | ||
| 110 | +1. **For op#11 and any similar multi-row × per-row-scalar scaling pattern, adopt the K_brcb idiom**: once per outer step, `Brcb(k_full, k_ub, N_ROWS/8, {1,8})` to expand the scalar-per-row tensor into a full-width per-element tensor, then do a single `Mul(work, work, k_full, N_ROWS*H_block)`. This eliminates the entire MTE2_S → S_V → V_S scalar-pipe serialization chain and collapses 64 sequential Muls into 1 wide Mul. Observed **25.3× speedup** in this shape. If H per row > 8, replicate `k_full` blocks across H/8 sub-blocks (or use `Mul` with repeat-strides src1BlkStride=0). The "k fits in 8-elem blocks via Brcb" assumption is cheap: Brcb repeatTime=ceil(N_ROWS/8). | ||
| 111 | +2. **Do NOT promote `Muls (灵活标量位置)` (flexible single-point-LocalTensor scalar src) as a sync-bypass optimisation.** This probe shows it performs **marginally worse** than the straightforward `GetValue + Muls` pattern. The documentation may suggest cleaner source code, but it does **not** reduce runtime in this scalar-per-row scenario on CANN 9.0.0 + bisheng 15.0.5. Flag to KB [target_ascend950pr](../../../common/platforms/concepts/target_ascend950pr.md) § Scalar pipe: "Muls flex-scalar overload has no sync-bypass advantage over GetValue+Muls scalar form (probe 2026-04-22, ratio 0.974x)." | ||
| 112 | +3. **Anchor a new template `Q_scalar_broadcast.md` in KB**（源仓 probe 模板，未迁移）with the K_base vs K_brcb contrast as the canonical example, and a note that flex-scalar Muls is **not** a sync-bypass solution. | ||
| 113 | +4. **Apply immediately to op#11's inner loop** if op#11's per-row scalar Muls is on the critical path. Projected per-iter saving depends on op#11's H (for small H per row, the 25× ratio is the ceiling; for H=2048 per row the ratio will compress, but a scalar-only GetValue cost scales with N_ROWS and is always ≥ the Brcb+Mul cost for N_ROWS≥8). | ||
| 114 | + | ||
| 115 | +## Caveats | ||
| 116 | + | ||
| 117 | +- **Shape sensitivity**: this probe uses `H=8` per row to exactly match Brcb's 32B block output. For op#11's real H=2048 per row, K_brcb needs a secondary stage (either replicate k_full via VEC Copy with stride-0, or use `Mul` with the `UnaryRepeatParams` src1BlkStride=0 pattern). The 25× speedup is an **upper bound** — a realistic op#11-shaped version may see 3–8× once H is large (because the wide `Mul` cost grows with H while scalar-Muls cost grows with N_ROWS × (H/VEC_WIDTH)). The **direction** (K_brcb always wins for N_ROWS ≥ 8) is robust; the **magnitude** needs a shape-matched follow-up probe before landing in op#11 KB guidance. | ||
| 118 | +- **Cumulative multiplies**: the probe intentionally runs k^512 to amplify per-launch time past launch overhead; this means fp32 rounding accumulates. 5e-4 relative tolerance was adopted; observed 1.2e-5 is well below. If a future variant diverges (e.g. truncation vs round-to-even), this probe will catch it as a miscompile. | ||
| 119 | +- **CANN/bisheng version sensitivity**: bisheng may refactor scalar-pipe scheduling in future releases. Re-run this probe on any bisheng upgrade; the flex-scalar overload could start bypassing sync in a later compiler. | ||
| 120 | +- **Single-AIV single-block**: dual-AIV / multi-block parallelism not tested. K_brcb's 25× win is within one AIV; across multiple blocks each block independently gets the same win, so the ratio should hold in an N-block launch of the same shape. | ||
| 121 | +- **Not a dedicated template prior to this probe**: this anchors the future `Q_scalar_broadcast.md` template — the sync-bypass mechanism identified (Brcb+Mul over per-row scalar Muls) is the definition of that template's question. | ||
| 122 | + | ||
| 123 | +## Iterations run (2 iters total; 1 build-fix, 1 run) | ||
| 124 | + | ||
| 125 | +1. **Build iter 1**: script auto-selected `/usr/local/Ascend/cann-9.0.0.T501` as CANN path, which lacks `tikcpp/ascendc_kernel_cmake` subtree → cmake configure FAIL. Fixed by exporting `ASCEND_HOME_PATH=/usr/local/Ascend/cann-9.0.0` + `source /usr/local/Ascend/cann-9.0.0/set_env.sh` + `LD_LIBRARY_PATH` prepend for `libhccl.so` (needed by `import torch_npu` at cmake-configure time for TORCH_PATH detection). | ||
| 126 | +2. **Build iter 2**: clean build; encountered "clock skew" style failure (`cmake -E touch: failed to update`) caused by WSL-side files having future mtime after tar extract; `find -exec touch +` + full rm -rf build, then rebuild. Succeeded with no kernel warnings. All 3 kernels + pybind produced. | ||
| 127 | +3. **Run iter 1**: residual Q4 `run_probe.py` (for `run_wrm` wholeReduceMax) was still present; initial run crashed with `AttributeError: module '_probe_ext' has no attribute 'run_wrm'`. Fixed by `docker cp` of the new `run_probe.py` + `model_new_ascendc.py` (the tar path mapping didn't overwrite these at the container root; probably extracted into current_task/ subdir but not replaced top-level). Subsequent run produced the final numbers above in a single pass — no timing iterations needed, stdev < 1% on first attempt. | ||
| 128 | + | ||
| 129 | +(Well inside the 2-compile-fix + 30-min budget: total wall ≈ 20 min, including env capture and report.) | ||
| @@ -0,0 +1,174 @@ | |||
| 1 | +--- | ||
| 2 | +type: Concept | ||
| 3 | +title: "SIMT vs SIMD Decision Framework" | ||
| 4 | +description: "SIMT 与 SIMD 选型决策框架（P-P9 权威参考）：给出按算子特征的速查表与完整决策树——原子加/散写/间接寻址/group-local 计算选 SIMT，连续读加向量计算、批量搬运选 SIMD；含精度约束 OL-30 与 Sparse-Gather、MXFP4 四个实测案例。" | ||
| 5 | +tags: [task.design, topic.programming_model.execution_model.simt, topic.programming_model.execution_model.simd] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/reference/porter/handbook/simt_vs_simd_decision.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-08-05T14:23:52Z" | ||
| 11 | +updated_at: "2026-09-16T07:31:58Z" | ||
| 12 | +platforms: ["950"] | ||
| 13 | +--- | ||
| 14 | + | ||
| 15 | +# SIMT vs SIMD Decision Framework | ||
| 16 | + | ||
| 17 | +> When to use SIMT, when to use SIMD, and why. Based on verified production experience | ||
| 18 | +> with Sparse-Gather (E8-E14) and MXFP4 quantization (2026-04-07)，实测在 Ascend950PR（A5）上完成。 | ||
| 19 | +> | ||
| 20 | +> **This is the authoritative reference for P-P9 (SIMT vs SIMD choice)**（原 OKF v1 卡号 P-P9 的权威参考）。 | ||
| 21 | +> | ||
| 22 | +> **2026-09-16 修订**：旧规则「间接寻址 → SIMT」已废弃——间接寻址须区分 addressing layer | ||
| 23 | +> 与 data layer，地址间接但数据连续时（如 SG Forward 的 expert 行）SIMD 反而快 2–7×。 | ||
| 24 | +> 见下方 Quick Decision Table 与 Decision Tree Step 2 的更新，修订依据见文末 | ||
| 25 | +> [SIMD vs SIMT selection（P-P9 修订卡）](../optimizations/p_p9_simd_vs_simt_selection.md)。 | ||
| 26 | + | ||
| 27 | +## Quick Decision Table | ||
| 28 | + | ||
| 29 | +| 算子特征 | SIMT | SIMD | 原因 | | ||
| 30 | +|----------|:----:|:----:|------| | ||
| 31 | +| atomicAdd / scatter-write | ✅ | ❌ | SIMD 无 atomicAdd（SetAtomicAdd 要求对齐且功能受限） | | ||
| 32 | +| ~~间接寻址 (arr[index[i]])~~ **旧规则已废弃** | ⚠️ | ⚠️ | 须区分 addressing layer / data layer，见 Decision Tree Step 2 | | ||
| 33 | +| 间接读 + 加权求和（地址间接、数据连续，如 SG Forward） | ❌ | ✅ | DataCopy 块搬运 + 4 流水并行，比 SIMT 散读快 2–7× | | ||
| 34 | +| 散读 + 散写混合（两端都不规则） | ✅ | ❌ | SIMD 无法编排两端不规则访问 | | ||
| 35 | +| 连续读 + 向量计算 (Add/Muls/Cast) | ❌ | ✅ | MTE2/VEC pipeline 重叠 1.6-2.3x | | ||
| 36 | +| **Group-local 计算 (group_size < tile_size)** | ✅ | ❌ | **SIMD per-group 循环慢于 SIMT 并行** | | ||
| 37 | +| Per-element 异构操作 (不同 shift/branch) | ✅ | ⚠️ | SIMD 可用 Compare+Select，但开销大 | | ||
| 38 | +| 纯位操作 (reinterpret float↔int) | ⚠️ | ⚠️ | 两者都可，看并行度 | | ||
| 39 | +| 批量数据搬运 + 简单计算 | ❌ | ✅ | MTE2 DMA 带宽 >> dcache | | ||
| 40 | + | ||
| 41 | +## Decision Tree (完整版) | ||
| 42 | + | ||
| 43 | +``` | ||
| 44 | +Step 1: 是否有 atomicAdd / scatter-write? | ||
| 45 | + YES → SIMT (SIMD 无法做原子操作) | ||
| 46 | + NO → Step 2 | ||
| 47 | + | ||
| 48 | +Step 2: 是否有间接寻址 (arr[index[i]])? | ||
| 49 | + 【2026-09-16 修订】旧规则 "间接寻址 → SIMT" 已废弃。 | ||
| 50 | + 新规则: 区分 addressing layer 与 data layer — | ||
| 51 | + a) addressing layer 间接 (选哪一行/哪个 expert)， | ||
| 52 | + 但 data layer 连续 (目标行本身是连续内存，如 SG Forward 的 expert row) | ||
| 53 | + → SIMD: scalar 侧算地址，DataCopy/MTE2 批量搬运数据层， | ||
| 54 | + 比 SIMT 线程标量散读快 2-7x (msprof 实测) | ||
| 55 | + b) 散读 + 散写混合 (两端都不规则) | ||
| 56 | + → SIMT (SIMD 无法编排) | ||
| 57 | + 例外: 如果 index 有序且可 batch，考虑 sort-to-reuse (P-P24，链接见文末) | ||
| 58 | + 原因 (msprof 证据): SIMT 模式下 GM 访问只走 VEC pipe 的 load/store 单元， | ||
| 59 | + MTE2/MTE3 DMA 引擎不参与 (硬件限制，非代码问题)，4 流水并行不可达； | ||
| 60 | + SIMT 的 VEC pipe 被 GM 读延迟打满 (vec=0.95+, mte2=0.000)， | ||
| 61 | + 而 SIMD 可同时跑 4 条流水 (vec+scl+mte2+mte3 各 30-90%)。 | ||
| 62 | + NO → Step 3 | ||
| 63 | + | ||
| 64 | +Step 3: 计算是否有 group-local 依赖? | ||
| 65 | + (即: 每 N 个元素共享一个参数，N < tile_size) | ||
| 66 | + YES → Step 3a | ||
| 67 | + NO → Step 4 (所有元素独立，纯 element-wise) | ||
| 68 | + | ||
| 69 | +Step 3a: group_size >= 256? | ||
| 70 | + YES → SIMD (group 足够大，per-group 循环开销可接受) | ||
| 71 | + 使用 TQue<VECIN,4> pipeline + per-group VEC 操作 | ||
| 72 | + NO → Step 3b | ||
| 73 | + | ||
| 74 | +Step 3b: per-group 计算是否全部可向量化? | ||
| 75 | + (即: group 内每个元素执行完全相同的指令序列) | ||
| 76 | + YES → SIMD 可能有优势，但需验证: | ||
| 77 | + - 如果 group=32, tile=1024: 需要 32 次 per-group 循环 | ||
| 78 | + - SIMT 128 threads 并行 128 groups = 4096 elements/dispatch | ||
| 79 | + - 经验: MXFP4(group=32) SIMD V3 比 SIMT 慢 6x | ||
| 80 | + → **默认 SIMT，除非 profiling 证明 SIMD 更快** | ||
| 81 | + NO → **SIMT** (per-element 异构计算，SIMT 线程并行最优) | ||
| 82 | + 例: MXFP4 per-element x_exp + variable shift → SIMT 快 6-20x | ||
| 83 | + | ||
| 84 | +Step 4: 计算是否可用 SIMD 向量指令表达? | ||
| 85 | + (Abs, Add, Muls, Cast, Compare+Select — 同一操作应用到所有元素) | ||
| 86 | + YES → **SIMD** (MTE2/VEC pipeline 重叠) | ||
| 87 | + 例: SG forward (DataCopy + Muls + Add) → SIMD TQue 1.6-2.3x | ||
| 88 | + NO → **SIMT** (需要 per-element 分支/位操作) | ||
| 89 | +``` | ||
| 90 | + | ||
| 91 | +## ⚠️ 精度约束 (OL-30) | ||
| 92 | + | ||
| 93 | +**SIMD 优化不能以精度降级为代价。** | ||
| 94 | + | ||
| 95 | +常见陷阱: 为了 SIMD 性能，将 per-group 操作改为 tile-wide 操作。 | ||
| 96 | +- ❌ MXFP4: 1024 元素共享一个 exponent（spec 要求每 32 个）→ **精度 bug** | ||
| 97 | +- ❌ A3 手写 SIMD: BATCH=512 共享 exponent → **精度 bug（已确认）** | ||
| 98 | +- ✅ SG forward: per-token 计算，每个 token 独立 → 无精度问题 | ||
| 99 | + | ||
| 100 | +**规则**: 如果优化改变了算法的 group/block precision 语义， | ||
| 101 | +则不能作为 production kernel。必须标注 "approximate" 并提供精确版本。 | ||
| 102 | + | ||
| 103 | +## Verified Case Studies | ||
| 104 | + | ||
| 105 | +### Case 1: Sparse-Gather Forward (SIMD wins) | ||
| 106 | + | ||
| 107 | +| 特征 | 值 | | ||
| 108 | +|------|---| | ||
| 109 | +| 计算 | DataCopy + Muls + Add (全向量化) | | ||
| 110 | +| Group 依赖 | 无 (per-token 独立) | | ||
| 111 | +| 数据访问 | 连续 (expert embedding) | | ||
| 112 | +| **结果** | SIMD TQue<4> **1.6-2.3x** faster than PipeBarrier | | ||
| 113 | + | ||
| 114 | +为什么 SIMD 赢: 每个元素执行完全相同的 Muls+Add，无 group-local 依赖，MTE2 prefetch 有效。 | ||
| 115 | + | ||
| 116 | +### Case 2: Sparse-Gather Backward Sorted (SIMD wins) | ||
| 117 | + | ||
| 118 | +| 特征 | 值 | | ||
| 119 | +|------|---| | ||
| 120 | +| 计算 | DataCopy + Muls + Add (累加到 TQue<VECOUT>) | | ||
| 121 | +| Group 依赖 | per-expert 累加，但 expert run 够长 | | ||
| 122 | +| 数据访问 | 连续 | | ||
| 123 | +| **结果** | SIMD TQue **1-10%** faster, 7 PipeBarrier→0 | | ||
| 124 | + | ||
| 125 | +为什么 SIMD 赢（但幅度小）: scalar pipe 是真瓶颈 (0.648)，MTE2/VEC 重叠收益有限。 | ||
| 126 | + | ||
| 127 | +### Case 3: MXFP4 Quantization (SIMT wins) | ||
| 128 | + | ||
| 129 | +| 特征 | 值 | | ||
| 130 | +|------|---| | ||
| 131 | +| 计算 | per-element log2 + floor + pow2 + round (异构) | | ||
| 132 | +| Group 依赖 | **per-32-element shared exponent** | | ||
| 133 | +| 数据访问 | 连续但 group-local | | ||
| 134 | +| **结果** | SIMT **6x** faster than SIMD V3, **20x** faster than SIMD V1 | | ||
| 135 | + | ||
| 136 | +为什么 SIMT 赢: | ||
| 137 | +1. group_size=32 太小 → SIMD 需要 per-group 循环 → 串行 | ||
| 138 | +2. SIMT 128 threads 并行处理 128 groups (4096 elements) → 天然并行 | ||
| 139 | +3. 消除 per-group 循环 (SIMD V4 fast) 性能提升但 **精度不符 spec** | ||
| 140 | +4. Per-element x_exp 导致 Compare+Select 开销（SIMD V3），SIMT 线程各自独立计算无开销 | ||
| 141 | + | ||
| 142 | +### Case 4: MXFP4 Tile-Wide Approximate (SIMD wins, but wrong precision) | ||
| 143 | + | ||
| 144 | +| 特征 | 值 | | ||
| 145 | +|------|---| | ||
| 146 | +| 计算 | tile-wide Abs + Muls + Cast(FLOOR) + Select | | ||
| 147 | +| Group 依赖 | **无** (放弃了 per-group，改为 tile-wide) | | ||
| 148 | +| 精度 | ⚠️ **不符合 MXFP4 spec** | | ||
| 149 | +| **结果** | SIMD **1.08x** faster than SIMT on 4K | | ||
| 150 | + | ||
| 151 | +为什么 SIMD 赢: 消除 per-group 循环后，全部是 tile-wide SIMD 向量操作，MTE2 pipeline 生效。 | ||
| 152 | +**但不可用于 production** — 精度降级不可接受。 | ||
| 153 | + | ||
| 154 | +## Pattern P-P9 Summary | ||
| 155 | + | ||
| 156 | +**默认选择**: | ||
| 157 | +- 有 atomicAdd / scatter-write → **SIMT** | ||
| 158 | +- 间接寻址先分 addressing/data 两层：地址间接但数据连续 → **SIMD**；两端都不规则 → **SIMT** | ||
| 159 | +- 有 group-local 依赖 (group < 256) → **SIMT** | ||
| 160 | +- 连续读 + 纯向量计算 + 无 group 依赖 → **SIMD** | ||
| 161 | + | ||
| 162 | +**验证必须**: | ||
| 163 | +- 选择后必须 A/B benchmark 确认（OL-27: 同 NPU 同 session） | ||
| 164 | +- SIMD 优化不得降低精度（OL-30） | ||
| 165 | +- 如果 SIMD 比 SIMT 慢 → 直接用 SIMT，不需要继续优化 SIMD | ||
| 166 | + | ||
| 167 | +（OL-27 / OL-30 / P-P24 原源仓 runbook/patterns 卡号，均已迁入本仓： | ||
| 168 | +[OL-27](../runbooks/performance/ol_27_performance_claims_must_be_based_on_same_condition.md)、 | ||
| 169 | +[OL-30](../runbooks/precision/ol_30_simd_performance_optimization_must_not_come_at_the.md)、 | ||
| 170 | +[P-P24](../optimizations/p_p24_sort_to_reuse_eliminate_gm_read_amplification_fr.md)。） | ||
| 171 | + | ||
| 172 | +## 相关卡片 | ||
| 173 | + | ||
| 174 | +- [SIMD vs SIMT selection（P-P9 修订卡）](../optimizations/p_p9_simd_vs_simt_selection.md)：废弃「间接寻址 → SIMT」旧规则的 msprof 证据与 addressing layer / data layer 区分准则的原始出处。 | ||
| @@ -0,0 +1,140 @@ | |||
| 1 | +--- | ||
| 2 | +type: Concept | ||
| 3 | +title: "PROBE REPORT — Q5 Sort/Reduce instruction cycles" | ||
| 4 | +description: "Ascend950PR 实测（CANN 9.0.0 / bisheng 2026-03-21）：WholeReduceMax/BlockReduceMax 64×fp32 ≈100 cyc/call（≤64 fp32 不必再切分），MrgSort 成本随输出对数线性 ≈0.55 cyc/pair + ~42 cyc 启动开销，宜单次长队列调用。" | ||
| 5 | +tags: [task.design, topic.performance.analysis, topic.compute.reduction, topic.compute.ordering] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/runbooks/hardware/probe-2026-04-21-q-instruction-cycles.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-07-12T16:00:00Z" | ||
| 11 | +updated_at: "2026-07-12T16:00:00Z" | ||
| 12 | +platforms: ["950"] | ||
| 13 | +aliases: [hw/2026-04-21_Q_instruction_cycles] | ||
| 14 | +--- | ||
| 15 | + | ||
| 16 | +# PROBE REPORT — Q5 Sort/Reduce instruction cycles | ||
| 17 | + | ||
| 18 | +> 原 OKF v1 卡号：**hw/2026-04-21_Q_instruction_cycles**（原 OKF v1 kind: guide；confidence: single_run）。 | ||
| 19 | + | ||
| 20 | +## Verdict | ||
| 21 | +ACCEPT_CORRECT | ||
| 22 | + | ||
| 23 | +**One-line summary**: On Ascend950PR_9589 / CANN 9.0.0 / bisheng 15.0.5 (2026-03-21), | ||
| 24 | +`WholeReduceMax<float>(64)` ≈ 102 cyc/call, `BlockReduceMax<float>(64)` ≈ 99 cyc/call, | ||
| 25 | +`MrgSort<float>` scales near-linearly with output length (≈ 0.55 cyc per output pair | ||
| 26 | +at Q≥256, with ~100 cyc/call fixed overhead). | ||
| 27 | + | ||
| 28 | +## Environment | ||
| 29 | +- date: 2026-04-22 (UTC, host clock ~18 min ahead of WSL — cosmetic) | ||
| 30 | +- host: 198.51.100.35 → container `npu_dev3` | ||
| 31 | +- NPU: `/dev/davinci1` (NPU 1, Ascend950PR, idle pre + post, AICore 0%) | ||
| 32 | +- CANN: `Ascend-cann-toolkit 9.0.0 V100R001C10SPC001B218` | ||
| 33 | +- bisheng / ccec: `clang 15.0.5 (clang-5c68a1cb1231 flang-5c68a1cb1231)` build stamp `2026-03-21T17:07:34+08:00` | ||
| 34 | +- SOC target at build: `-DSOC_VERSION=Ascend950PR_9589` | ||
| 35 | +- frequency assumption for cycle derivation: **1.8 GHz** (AIV clock on 950PR; re-scale | ||
| 36 | + if vendor spec differs — time columns in timings.csv are authoritative). | ||
| 37 | + | ||
| 38 | +## Question | ||
| 39 | +Per-repeat cycle latency and steady-state throughput of: | ||
| 40 | +- `WholeReduceMax<float>` (max 64 fp32 per call in fp32 mask-count mode) | ||
| 41 | +- `BlockReduceMax<float>` (32-byte block = 8 fp32 per block, 8 blocks = 64 fp32) | ||
| 42 | +- `MrgSort<float>` (4-way merge of pre-sorted (score,index) queues, varied queue length) | ||
| 43 | + | ||
| 44 | +## Hypothesis vs Observation | ||
| 45 | +- Hypothesis: no strong prior; record numbers. | ||
| 46 | +- Observation: all three instructions produced stable, reproducible timings | ||
| 47 | + (stdev < 1 µs on ≈ 230–5200 µs means); MrgSort is linear in queue length | ||
| 48 | + (tree-log model is **not** the dominant cost — output size dominates). | ||
| 49 | + | ||
| 50 | +## Evidence | ||
| 51 | + | ||
| 52 | +### Results (REPEAT_N = 4096 calls, 1.8 GHz assumed) | ||
| 53 | + | ||
| 54 | +| primitive | elements_per_call | median_us | cycles_per_repeat | throughput (elts/cycle) | | ||
| 55 | +|---|---|---|---|---| | ||
| 56 | +| WholeReduceMax | 64 (fp32) | 232.30 | 102.1 | 0.63 | | ||
| 57 | +| BlockReduceMax | 64 (fp32, 8×8-blk) | 225.18 | 99.0 | 0.65 | | ||
| 58 | +| MrgSort_q64 | 4×64 = 256 in / out | 413.94 | 181.9 | 1.41 | | ||
| 59 | +| MrgSort_q256 | 4×256 = 1024 in / out | 1366.99| 600.7 | 1.70 | | ||
| 60 | +| MrgSort_q1024 | 4×1024 = 4096 in / out | 5179.60| 2276.2 | 1.80 | | ||
| 61 | + | ||
| 62 | +### Derived MrgSort model | ||
| 63 | + | ||
| 64 | +Regression on 3 MrgSort points (output pairs N = 256, 1024, 4096): | ||
| 65 | +- Slope: ≈ (2276−182) cyc / (4096−256) pairs = **0.545 cyc/output-pair** | ||
| 66 | +- Intercept: ≈ 42 cyc (per-call setup) | ||
| 67 | +- Throughput asymptotically **~1.83 output-pairs/cycle** as queues grow | ||
| 68 | +- Confirms **linear (not super-linear)** scaling → hardware is not internally | ||
| 69 | + serializing on queue length; the "tree-log" ideal does not apply because | ||
| 70 | + MrgSort must emit every output element regardless of comparator tree depth. | ||
| 71 | + | ||
| 72 | +### Reduce instructions | ||
| 73 | + | ||
| 74 | +WholeReduceMax and BlockReduceMax over 64 fp32 are nearly identical in cost | ||
| 75 | +(~100 cyc); the BlockReduce variant is marginally cheaper (~3%). The ≈ 100-cyc | ||
| 76 | +floor likely includes pipeline setup and the result-writeback; the per-element | ||
| 77 | +"useful work" is sub-cycle. | ||
| 78 | + | ||
| 79 | +**Practical guidance**: for scan-then-reduce patterns over ≤ 64 fp32, the | ||
| 80 | +single-instruction cost is the floor; further chunking will not help. | ||
| 81 | + | ||
| 82 | +### Build | ||
| 83 | + | ||
| 84 | +Clean build succeeded with one non-fatal warning: | ||
| 85 | + | ||
| 86 | +``` | ||
| 87 | +[WARNING]: Multiple kernel functions are detected. It is recommended to define | ||
| 88 | +only one kernel function per file. | ||
| 89 | +``` | ||
| 90 | +(cosmetic — 5 `__global__` kernels share `probe_kernels.cpp`; does not affect correctness.) | ||
| 91 | + | ||
| 92 | +### Runtime | ||
| 93 | + | ||
| 94 | +Two independent runs produced medians within 0.3% of each other (232.3 vs 232.8 | ||
| 95 | +µs for WRM; 225.18 vs 225.14 µs for BRM; MrgSort rows within 0.05 µs), | ||
| 96 | +confirming the loop-amplification + discard-first-5 methodology removes I-cache | ||
| 97 | +cold-start noise. NPU 1 verified idle (AICore 0%, no foreign process) before | ||
| 98 | +AND after the measurement set. | ||
| 99 | + | ||
| 100 | +### Raw files | ||
| 101 | +- Kernel: `workspace/probe_instruction_cycles/kernel/probe_kernel.h` (ProbeWholeMax, ProbeBlockMax, ProbeMrgSort<QLEN>) | ||
| 102 | +- Kernel exports: `workspace/probe_instruction_cycles/kernel/probe_kernels.cpp` | ||
| 103 | +- Timing harness: `workspace/probe_instruction_cycles/kernel/pybind11.cpp` | ||
| 104 | +- Runner: `workspace/probe_instruction_cycles/run_probe.py` | ||
| 105 | +- Build log: `workspace/probe_instruction_cycles/build.log` (196 lines, 1 cosmetic warning) | ||
| 106 | +- Run log: `workspace/probe_instruction_cycles/run.log` | ||
| 107 | +- Timings CSV: `workspace/probe_instruction_cycles/timings.csv` | ||
| 108 | + | ||
| 109 | +## Recommendation for orchestrator | ||
| 110 | + | ||
| 111 | +1. **KB entry**（[target_ascend950pr](../../../common/platforms/concepts/target_ascend950pr.md) §"Sort/Reduce cycle data, bisheng 2026-03-21 / CANN 9.0.0"）: | ||
| 112 | + record the table above. Tag version-stamp; **not** cross-version portable. | ||
| 113 | +2. **Op-gen heuristic for reductions over ≤ 64 fp32**: a single WholeReduceMax / | ||
| 114 | + BlockReduceMax is ≈ 100 cyc. Do not chunk below this size; cost is amortized | ||
| 115 | + by the pipeline setup, not by per-element work. | ||
| 116 | +3. **Op-gen heuristic for MrgSort**: expected cost ≈ `0.55 × 4 × QLEN + 100 cyc` | ||
| 117 | + per invocation. For a fixed input size, shorter queues × more MrgSort calls | ||
| 118 | + is *slower* than one long-queue call (per-call overhead is constant). Confirms | ||
| 119 | + the existing pattern in 源仓算子目录 `9_TopKTopP`（未迁移）of preferring Sort over iterative | ||
| 120 | + pair-merges when possible. | ||
| 121 | +4. **Follow-up (not urgent)**: re-probe after any bisheng upgrade; validate | ||
| 122 | + 1.8 GHz assumption (or switch to msprof cycle counters) if cycle counts | ||
| 123 | + diverge from the ×1.8e9 scaling by >5%. | ||
| 124 | + | ||
| 125 | +## Caveats | ||
| 126 | + | ||
| 127 | +- **Frequency**: cycle numbers depend on assumed 1.8 GHz AIV clock. The raw | ||
| 128 | + `median_us` column in `timings.csv` is authoritative and re-scalable. | ||
| 129 | +- **Loop body DCE risk**: the compiler might have hoisted / CSE'd some of the | ||
| 130 | + repeated WholeReduceMax calls; however, the observed cycle count (~100/call) | ||
| 131 | + is far above a `nop` loop cost, and measurements scale with REPEAT_N (halved | ||
| 132 | + REPEAT_N → halved time; sanity-confirmed during iteration on build cache). | ||
| 133 | +- **bisheng-specific**: results will not transfer to a different bisheng | ||
| 134 | + codegen version. | ||
| 135 | +- **Single-block**: all timings are with `blockDim=1` (single AIV). Multi-block | ||
| 136 | + throughput should scale linearly until memory/issue-slot saturation; not probed here. | ||
| 137 | +- **MrgSort input format**: used fp32 score + fp32 "index slot" (8 B / pair). The | ||
| 138 | + index field is reinterpret-cast at the hardware level; correctness of the sorted | ||
| 139 | + output was not bit-checked here (probe is timing-only; sort engine functional | ||
| 140 | + correctness is covered by CANN itself). | ||
| @@ -0,0 +1,190 @@ | |||
| 1 | +--- | ||
| 2 | +type: Concept | ||
| 3 | +title: "PROBE REPORT — Q2 UB bank count / width / conflict behavior" | ||
| 4 | +description: "Ascend950PR UB stride 扫描实测（CANN 9.0.0 / bisheng 2026-03-21）：dense-sequential 比任意 strided 访问快 ~22%，mod-4 plateau 仅 ~6%，排除单端口 K=8/16×32B 交织假设；bank 数量与每 bank 宽度无法从 timing 反推（Verdict INCONCLUSIVE_PARTIAL）。" | ||
| 5 | +tags: [task.optimization, topic.memory, topic.performance.analysis] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/runbooks/hardware/probe-2026-04-21-q-ub-bank-count.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-07-12T16:00:00Z" | ||
| 11 | +updated_at: "2026-07-12T16:00:00Z" | ||
| 12 | +platforms: ["950"] | ||
| 13 | +aliases: [hw/2026-04-21_Q_ub_bank_count] | ||
| 14 | +--- | ||
| 15 | + | ||
| 16 | +# PROBE REPORT — Q2 UB bank count / width / conflict behavior | ||
| 17 | + | ||
| 18 | +> 原 OKF v1 卡号：**hw/2026-04-21_Q_ub_bank_count**（原 OKF v1 kind: guide；confidence: single_run）。 | ||
| 19 | + | ||
| 20 | +## Verdict | ||
| 21 | +INCONCLUSIVE_PARTIAL | ||
| 22 | + | ||
| 23 | +**One-line summary**: The stride-sweep methodology successfully revealed a | ||
| 24 | +**dense-sequential fast path** (stride=1 at 12.26 cyc/repeat) vs a **strided | ||
| 25 | +slow zone** (14.7–15.5 cyc/repeat, +20–27%). But the observed spread is too | ||
| 26 | +compressed to nail down bank count K or per-bank width W directly. No | ||
| 27 | +catastrophic K× cliff was observed at any stride up to 32, which **rules out | ||
| 28 | +a naïve K=8 banks × 32B interleave** (that would predict ~8× penalty at S=8) | ||
| 29 | +but does not rule out a more forgiving bank network with ≥ 2 R-ports per bank | ||
| 30 | +group (consistent with the 351x public doc claim of "2R+2W per bank group"). | ||
| 31 | + | ||
| 32 | +## Environment | ||
| 33 | +- date: 2026-04-21 (UTC) | ||
| 34 | +- host: 198.51.100.35 → container `npu_dev3` | ||
| 35 | +- NPU: `/dev/davinci1` (NPU 1, Ascend950PR, idle pre+post, AICore 0%, no | ||
| 36 | + foreign process; NPU 0 busy with op#11 pytest, avoided) | ||
| 37 | +- CANN: `Ascend-cann-toolkit 9.0.0 V100R001C10SPC001B218` | ||
| 38 | +- bisheng: `clang 15.0.5 (clang-5c68a1cb1231 ...)` build stamp `2026-03-21T17:07:34+08:00` | ||
| 39 | +- SOC target at build: `-DSOC_VERSION=Ascend950PR_9589` | ||
| 40 | +- AIV frequency assumed for cycle derivation: **1.8 GHz** (same assumption as [Q5](sort_reduce_instruction_cycles.md)) | ||
| 41 | +- Python env: torch + torch_npu (requires `setenv.bash` to load libhccl.so) | ||
| 42 | + | ||
| 43 | +## Question | ||
| 44 | +How many banks does the 256KB UB have, what's the per-bank width, and what | ||
| 45 | +stride/alignment avoids bank conflicts? | ||
| 46 | + | ||
| 47 | +## Hypothesis vs Observation | ||
| 48 | +- **Hypothesis**: common options are 8 or 16 banks × 32B or 64B, interleaved | ||
| 49 | + at native-load granularity (likely 32B × 8 banks on a 256-byte VEC register | ||
| 50 | + chip). | ||
| 51 | +- **Observation**: Stride-sweep of `Mul<float>(mask=64, repeat=8, src0BlkStride=S)`: | ||
| 52 | + - S=1 (dense): **12.26 cyc/repeat** — baseline / fast path | ||
| 53 | + - S=2,3,5,6,7: **14.7–15.0 cyc/repeat** (+20–22%) | ||
| 54 | + - S=4,8,16: **15.5 cyc/repeat** (+27%) — mild upper plateau | ||
| 55 | + - S=32: 15.3 cyc/repeat — no further degradation | ||
| 56 | + - **No catastrophic cliff** anywhere. The dynamic range between best and | ||
| 57 | + worst non-unit stride is only ~6%, which is too compressed to cleanly | ||
| 58 | + back-solve bank count. | ||
| 59 | + | ||
| 60 | +## Evidence | ||
| 61 | + | ||
| 62 | +### Results (REPEAT_N = 4096 Mul calls, INNER_REPEATS = 8 per call, 1.8 GHz) | ||
| 63 | + | ||
| 64 | +| src0BlkStride | median µs | min µs | stdev | cyc/call | cyc/repeat | | ||
| 65 | +|---:|---:|---:|---:|---:|---:| | ||
| 66 | +| 1 | 223.25 | 222.85 | 0.49 | 98.1 | **12.26** | | ||
| 67 | +| 2 | 272.90 | 272.24 | 0.41 | 119.9 | 14.99 | | ||
| 68 | +| 3 | 268.26 | 267.49 | 0.91 | 117.9 | 14.74 | | ||
| 69 | +| 4 | 282.89 | 282.46 | 0.38 | 124.3 | **15.54** | | ||
| 70 | +| 5 | 268.03 | 267.44 | 0.37 | 117.8 | 14.72 | | ||
| 71 | +| 6 | 272.81 | 272.06 | 0.40 | 119.9 | 14.99 | | ||
| 72 | +| 7 | 268.15 | 267.48 | 0.33 | 117.8 | 14.73 | | ||
| 73 | +| 8 | 282.84 | 282.36 | 0.41 | 124.3 | **15.54** | | ||
| 74 | +| 16 | 282.75 | 282.41 | 0.25 | 124.3 | 15.53 | | ||
| 75 | +| 32 | 278.08 | 277.37 | 0.38 | 122.2 | 15.28 | | ||
| 76 | + | ||
| 77 | +### Interpretation / Inference attempts | ||
| 78 | + | ||
| 79 | +**Clear signal #1 — dense is 22% faster than any non-unit stride.** | ||
| 80 | +S=1 is the only stride where the 8 × 32B blocks of a mask-64 repeat land in | ||
| 81 | +strictly consecutive UB addresses. The ~22% penalty for any S≥2 says the vector | ||
| 82 | +unit *can* burst-read 256B contiguous in a single cycle but must serialize / | ||
| 83 | +arbitrate on any scatter. Practical guidance: **contiguous UB layout is a | ||
| 84 | +first-order performance invariant**, bigger than any secondary bank effect. | ||
| 85 | + | ||
| 86 | +**Clear signal #2 — mod-4 parity in the even strides.** | ||
| 87 | +S ∈ {4, 8, 16} all sit at 15.53 cyc. S ∈ {2, 6} at 14.99 cyc. S=32 slightly | ||
| 88 | +lower (15.28). Odd strides all at ~14.73. This is consistent with some | ||
| 89 | +bank-group alignment where access patterns that are multiples of 4 blocks | ||
| 90 | +(128B) align all 8 accesses to 2 bank groups, while non-4-multiples spread | ||
| 91 | +across 4–8 bank groups, slightly reducing the per-group contention. The | ||
| 92 | +**"conflict" is only ~6% between the aligned-bad (S=4,8,16) and misaligned- | ||
| 93 | +better (S=3,5,7)** cases, far below a naïve "K-way conflict = K× slowdown". | ||
| 94 | + | ||
| 95 | +**What we can rule out**: | ||
| 96 | +- **K=8 banks × 32B interleave, single-port** — would predict ~8× slowdown | ||
| 97 | + at S=8. Observed: 1.27×. Rejected. | ||
| 98 | +- **K=16 banks × 32B interleave, single-port** — same style reasoning. Rejected. | ||
| 99 | +- **Any single-port bank layout** where stride-S collision triggers K-way | ||
| 100 | + serialization is inconsistent with the 6% spread. | ||
| 101 | + | ||
| 102 | +**What is consistent with observation**: | ||
| 103 | +- Multi-ported bank groups (the public doc's "2R+2W per group" statement). | ||
| 104 | + With 2 read ports per group, a 2-way conflict is absorbed for free; only | ||
| 105 | + a 3-way or worse conflict costs cycles, and even then the port structure | ||
| 106 | + throttles rather than stalls. That matches a small ~20% plateau rather | ||
| 107 | + than a cliff. | ||
| 108 | +- The fast path (S=1) exploits a separate "full-width burst" mode of the | ||
| 109 | + vector register-file that bypasses bank arbitration entirely. | ||
| 110 | + | ||
| 111 | +**What we CANNOT determine from this probe**: | ||
| 112 | +- Exact bank count K. Our stride grid (1,2,3,4,5,6,7,8,16,32) does not | ||
| 113 | + produce a distinguishable fingerprint for K ∈ {8, 16, 32}. | ||
| 114 | +- Per-bank width W (32B vs 64B). Our smallest access granularity was 32B | ||
| 115 | + (one block); distinguishing 32B vs 64B banks would require sub-block | ||
| 116 | + strided tests that AscendC BinaryRepeatParams doesn't expose. | ||
| 117 | +- Read-vs-write conflict behavior. We only probed RR (two src reads into | ||
| 118 | + one dst write); didn't isolate WW or RW conflicts. | ||
| 119 | + | ||
| 120 | +### Build | ||
| 121 | +Clean build succeeded (after an initial env setup hurdle: `libhccl.so` | ||
| 122 | +requires `source /usr/local/Ascend/cann-9.0.0/x86_64-linux/bin/setenv.bash` | ||
| 123 | +before invoking the cmake-based build, otherwise TORCH_PATH resolves to | ||
| 124 | +empty string and `<torch/extension.h>` cannot be found). One cosmetic warning: | ||
| 125 | + | ||
| 126 | +``` | ||
| 127 | +[WARNING]: Multiple kernel functions are detected. It is recommended to define | ||
| 128 | +only one kernel function per file. | ||
| 129 | +``` | ||
| 130 | +(10 stride-variant `__global__` kernels share `probe_kernels.cpp` — does not | ||
| 131 | +affect correctness.) | ||
| 132 | + | ||
| 133 | +### Runtime | ||
| 134 | +Single run produced stable timings across all 10 stride variants (stdev < | ||
| 135 | +1 µs on medians of 223–283 µs, i.e. < 0.4% relative). Discard-first-5 | ||
| 136 | +methodology from [Q5](sort_reduce_instruction_cycles.md) carried over cleanly. | ||
| 137 | + | ||
| 138 | +### Raw files | ||
| 139 | +- Kernel: `workspace/probe_ub_bank_count/kernel/{probe_kernels.cpp,probe_kernel.h,probe_tiling.h,pybind11.cpp}` | ||
| 140 | +- Build log: `workspace/probe_ub_bank_count/build.log` | ||
| 141 | +- Run log: `workspace/probe_ub_bank_count/run.log` | ||
| 142 | +- Timings: `workspace/probe_ub_bank_count/timings.csv` | ||
| 143 | + | ||
| 144 | +## Practical recommendations for orchestrator / future ops | ||
| 145 | + | ||
| 146 | +- **Rule of thumb for UB access**: keep the primary access dense-sequential | ||
| 147 | + (src0BlkStride=1 / src1BlkStride=1). Any strided access will cost ~22% | ||
| 148 | + regardless of stride value — there is **no magic "safe" stride** that | ||
| 149 | + matches the dense-sequential speed. | ||
| 150 | +- **Don't worry about specific stride-K choices for bank avoidance.** The | ||
| 151 | + worst observed penalty (S=4,8,16) is +27% over baseline, only +6% over | ||
| 152 | + the best non-unit stride. Fine-tuning stride for bank structure yields | ||
| 153 | + diminishing returns. | ||
| 154 | +- **If a kernel already uses stride≥2**, pick odd strides (3, 5, 7) if free | ||
| 155 | + — they measure ~5% faster than even-power strides (4, 8, 16). Not worth | ||
| 156 | + reshaping the kernel over, but worth knowing. | ||
| 157 | +- **For aggressive optimization**, prefer restructuring to make UB access | ||
| 158 | + fully contiguous (e.g. transpose-via-DataCopy before VEC, rather than | ||
| 159 | + VEC-with-stride). The 22% dense→strided cliff is the real prize. | ||
| 160 | + | ||
| 161 | +## Caveats | ||
| 162 | +- The stride grid is sparse — {9,10,11,12,14,24} untested; a full power-of-2 | ||
| 163 | + + nearby-offset sweep (e.g. 9,17,33) might reveal finer structure. | ||
| 164 | +- Methodology measures one Mul with src0 strided and src1+dst dense. A | ||
| 165 | + symmetric two-sided stride probe (both src0 and src1 strided) might produce | ||
| 166 | + cleaner conflict signatures. | ||
| 167 | +- Single NPU, single bisheng version; result may shift with compiler | ||
| 168 | + updates that change VEC scheduling. | ||
| 169 | +- Frequency assumption (1.8 GHz) affects absolute cycle numbers but not | ||
| 170 | + relative cycle ratios (which are what matters for conflict analysis). | ||
| 171 | +- **This verdict is indirect inference only.** Without access to an msprof | ||
| 172 | + `aiv_ub_bank_conflict` counter (checked via `msprof --help` — not available | ||
| 173 | + in CANN 9.0.0 user-facing interface), the true bank count cannot be | ||
| 174 | + definitively stated from timing data alone. | ||
| 175 | + | ||
| 176 | +## Suggested next-step probes (if budget permits) | ||
| 177 | + | ||
| 178 | +1. **Write-conflict probe**: two concurrent VEC stores with stride, to | ||
| 179 | + isolate WW conflicts from RR conflicts. | ||
| 180 | +2. **Sub-block probe**: use `Mul` with `dstBlkStride` and `src0BlkStride` | ||
| 181 | + differing at sub-block level via element-offset tricks — may expose | ||
| 182 | + 32B vs 64B per-bank width. | ||
| 183 | +3. **msprof counter enumeration**: run `msprof --info counters` on 950PR to | ||
| 184 | + see if `bank_conflict` or similar hardware counters are exposed; if so, | ||
| 185 | + re-run the existing stride sweep while counting — that would produce a | ||
| 186 | + definitive K,W answer. | ||
| 187 | + | ||
| 188 | +## 相关卡片 | ||
| 189 | + | ||
| 190 | +- [OL-50 Avoid UB bank conflicts](../runbooks/compilation/ol_50_avoid_ub_bank_conflicts_consecutive_vec_instructio.md)：通用 bank-conflict 规避建议（draft，平台未定）；本卡为 950PR 上的实测 stride 扫描——实测显示 stride 微调收益 ≤6%，与"按 bank 精细调 offset"的预期收益量级不同，消费时以本卡的实测边界为准。 | ||
| @@ -0,0 +1,183 @@ | |||
| 1 | +--- | ||
| 2 | +type: Example | ||
| 3 | +title: "a3 MIX many-small-matmul cube primitives — P-P119 (hand-rolled cube vs IterateAll) + P-P120 (Neumann (I+P) L0C-fold)" | ||
| 4 | +description: "a3/arch22 上 MIX 多个小矩阵乘的两条 cube 决策：≤128 的小 matmul 用手写单 block cube 原语而非 MatmulImpl::IterateAll（实测受 aic_scalar 限制）；用 Mmad 的 cmatrixInitVal 把 (I+P) 折进 L0C。" | ||
| 5 | +tags: [task.implementation, topic.programming_model.execution_unit.cube, topic.attention] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/reference/porter/patterns/a3_mix_small_matmul_cube.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-08-05T14:23:52Z" | ||
| 11 | +updated_at: "2026-09-04T23:59:25Z" | ||
| 12 | +platforms: [a3] | ||
| 13 | +--- | ||
| 14 | + | ||
| 15 | +# a3 MIX many-small-matmul cube primitives — P-P119 (hand-rolled cube vs IterateAll) + P-P120 (Neumann (I+P) L0C-fold) | ||
| 16 | + | ||
| 17 | +> 原 OKF v1 卡号：**P-P119**、**P-P120**。 | ||
| 18 | +> 适用范围：soc=Ascend910_9382（a3 / arch22，Atlas A3），CANN 9.1.0；两条决策在 Ascend910_9382 | ||
| 19 | +> 上设备实证（gated_delta_rule fwd cube-solve perf-opt，2026-07-20），**在 Ascend950PR 上未验证**。 | ||
| 20 | +> | ||
| 21 | +> These are the two cube-side decisions that separated our first pipeline-generated | ||
| 22 | +> `gated_delta_rule` kernel from the leaner customer cv-reference (both correct; cv-ref | ||
| 23 | +> leaner because the pipeline lacked exactly these two facts). Found by white-box | ||
| 24 | +> comparison, then device-measured during the fwd cube-solve perf-opt (2026-07-20). | ||
| 25 | +> These sit ON TOP of the a3 MIX chunked-recurrence algorithm pattern **P-P117** | ||
| 26 | +> ([gated_delta_rule_a3_recurrence](gated_delta_rule_a3_recurrence.md)) and the a3 MIX | ||
| 27 | +> cube+vector SYNC template **P-P116** ([fa_class_a3_mix_template](fa_class_a3_mix_template.md)). | ||
| 28 | +> The cube-primitive param values they emit are grounded in the source repo's | ||
| 29 | +> `fa_class/cv_reference_concrete_params.md` (`matmul_primitive` decision, not migrated). | ||
| 30 | + | ||
| 31 | +--- | ||
| 32 | + | ||
| 33 | +## P-P119 — many-small-matmul a3 MIX: emit a hand-rolled single-block cube primitive, NOT `MatmulImpl::IterateAll` | ||
| 34 | + | ||
| 35 | +### Trigger / applies_to | ||
| 36 | +a3 (Ascend910_9382, arch22) MIX op in the **chunked-recurrence / gated-linear-attention** | ||
| 37 | +family (P-P117) where EACH chunk performs **MANY small matmuls** — dimensions ≤128, so | ||
| 38 | +each matmul is a single L0C tile (`M,N,K ≤ 128`, one Mmad pass, no inner K-loop). A | ||
| 39 | +gated_delta_rule chunk runs ~8+ such matmuls (`Acc=kb@knT`, the 6-step Neumann product, | ||
| 40 | +`Am`, `U`, `W`, `WS`, `o_inter`, `S`-update) per (head, chunk), times `Nc` chunks times | ||
| 41 | +heads → hundreds-to-thousands of tiny matmuls per op. | ||
| 42 | + | ||
| 43 | +### The finding (why the library primitive is the wrong primitive here) | ||
| 44 | +`MatmulImpl<...>::IterateAll<sync=true>` (the P-P68 NON-KFC cube) is the correct, safe | ||
| 45 | +choice for a **few large** GEMMs. But for MANY SMALL matmuls its **per-call tiling scalar | ||
| 46 | +setup DOMINATES runtime** — every `IterateAll` re-runs the library's tiling/address | ||
| 47 | +bookkeeping (and on a reused object you additionally pay `SetOrgShape` per call, CAND-GDR-3). | ||
| 48 | +That scalar work is fixed-cost per call and does not shrink with the tiny `≤128³` MAC work, | ||
| 49 | +so at high call-count the kernel becomes **aic_scalar-bound, not MAC-bound**. | ||
| 50 | + | ||
| 51 | +**msprof evidence (a3/Ascend910_9382, gated_delta_rule fwd, IterateAll version):** | ||
| 52 | +`aic_scalar` **≈32%** / `aic_mac` **≈3.7%** / `cube_util` **≈17%**. The cube is ~9× more | ||
| 53 | +time in scalar tiling setup than in actual MACs — a textbook per-call-overhead-bound | ||
| 54 | +signature. | ||
| 55 | + | ||
| 56 | +### The fix — emit a hand-rolled single-block cube primitive per matmul | ||
| 57 | +Replace each `IterateAll` call with a direct, minimal cube pipeline for one L0C tile: | ||
| 58 | + | ||
| 59 | +``` | ||
| 60 | +Nd2Nz (GM/UB → L1, ND → NZ fractal layout) | ||
| 61 | + → LoadData / LoadData3D (L1 → L0A / L0B, no ifTranspose games; see below) | ||
| 62 | + → Mmad (single pass, M,N,K ≤ 128, one L0C tile) | ||
| 63 | + → Fixpipe[F322F16] (L0C fp32 accumulator → fp16 GM/UB, with the fp32→fp16 cast) | ||
| 64 | +``` | ||
| 65 | + | ||
| 66 | +There is NO tiling scalar setup per call — the block dims are compile-time-known | ||
| 67 | +(`≤128³`), the load/Mmad params are constants, so the aic_scalar cost per matmul collapses | ||
| 68 | +to just the fixed load+Mmad+Fixpipe issue. This is the SAME primitive family already | ||
| 69 | +grounded for FA-class cube QK^T in the source repo's `fa_class/cv_reference_concrete_params.md` | ||
| 70 | +(`matmul_primitive`: `Mmad` 4-arg accumulate form + `Fixpipe` per-context `srcStride` | ||
| 71 | +units); reuse those exact param rules — DO NOT re-derive `srcStride` (per-context: `/C0` | ||
| 72 | +for intermediate L0C→workspace copies, ELEMENTS for final L0C→GM writes; a wrong unit → | ||
| 73 | +507015 ECC read). | ||
| 74 | + | ||
| 75 | +### Measured result (device, a3/Ascend910_9382, 2026-07-20) | ||
| 76 | +This change ALONE closed a **~1.3× compute-bound deficit** vs the cv-reference to **parity**: | ||
| 77 | +**T4096 −35%**, **T1024 −30%** device-time, **precision unchanged** (still 16/16 @ the | ||
| 78 | +customer fp64 gate). No algorithm change — same matmul math, leaner cube emit. | ||
| 79 | + | ||
| 80 | +### Decision rule (put in the worker's hands) | ||
| 81 | +- **Few large GEMMs (one/few per op, M·N·K ≫ 128³)** → `MatmulImpl::IterateAll` (P-P68). The | ||
| 82 | + library tiling amortizes; hand-rolling buys nothing. | ||
| 83 | +- **Many small matmuls (≤128 each, hundreds+ per op — chunked-recurrence)** → hand-rolled | ||
| 84 | + single-block cube primitive. Confirm with msprof: `aic_scalar ≫ aic_mac` (e.g. 32% vs | ||
| 85 | + 3.7%) IS the "per-call tiling setup dominates" signature that says switch. | ||
| 86 | +- **⚠ CORRECTNESS gate on the standalone op-gen build — hand-rolled is REQUIRED from the | ||
| 87 | + FIRST pass, not a perf swap-in.** On the standalone `build_ascendc.py` path (no CANN | ||
| 88 | + op-framework), `MatmulImpl<> + IterateAll<sync=true>` under `MIX_AIC_1_1` needs the | ||
| 89 | + framework KFC bootstrap (`REGIST_MATMUL_OBJ` + msg-ring) the standalone build does NOT | ||
| 90 | + provide → the kernel LAUNCHES but throws a **507057 MIX-kernel task exception at | ||
| 91 | + execution (no output tensor)**. The hand-rolled single-block primitive has NO KFC | ||
| 92 | + dependency (raw `Nd2Nz→LoadData3D→Mmad→Fixpipe`), so it runs standalone. **Therefore do | ||
| 93 | + NOT use the "start with `MatmulImpl` for a correctness pass, swap to hand-rolled for perf" | ||
| 94 | + strategy on the standalone build — that pass 507057s and never reaches correctness. Emit | ||
| 95 | + the hand-rolled primitive from the start.** Device-proven: gated_delta_rule fwd hand-rolled | ||
| 96 | + = **16/16 @ fp64** (standalone build) vs `MatmulImpl` standalone = **0/16 / 507057** | ||
| 97 | + (cold-start re-gen 2026-07-21). Cross-ref OL-235 (KFC standalone wall) / PB-35 / | ||
| 98 | + CAND-KFC-standalone-bootstrap-teardown (source-repo runbook IDs, not migrated). | ||
| 99 | +- Precision is invariant across the two primitives (same Mmad math). The primitive choice is | ||
| 100 | + a PERF lever ONLY where BOTH are runnable (few-large, or a framework-integrated build); on | ||
| 101 | + the standalone many-small a3-MIX path it is a **CORRECTNESS gate** (MatmulImpl unrunnable → | ||
| 102 | + hand-rolled mandatory). | ||
| 103 | + | ||
| 104 | +--- | ||
| 105 | + | ||
| 106 | +## P-P120 — Neumann / (I+P) triangular-solve: fold identity+product on the cube via accumulate-mode Mmad (`cmatrixInitVal`), no separate +I vector-add | ||
| 107 | + | ||
| 108 | +### Trigger / applies_to | ||
| 109 | +a3 (Ascend910_9382, arch22) cube computing an **iterative (I+P)-style matmul chain** — | ||
| 110 | +canonically the **Neumann power-product triangular solve** used to invert | ||
| 111 | +`(I + strict-lower)` in chunked gated-delta-rule / gated-linear-attention: | ||
| 112 | +`A = (I + x)^-1 = Π_i (I + x^{2^i})`, i.e. a chain of steps each of the form | ||
| 113 | +`R_next = R @ (I + P)` where `P = x^{2^i}`. Any `R@(I+P)` / `R + R@P` fold qualifies. | ||
| 114 | + | ||
| 115 | +### The finding (the naive emit inserts a redundant vector +I) | ||
| 116 | +The naive lowering of one step is: (1) cube `R@P`, (2) vector `Add` the identity/additive | ||
| 117 | +term (`R + R@P`, or `I + P` built with a vector `Duplicate`-diagonal). Each step then costs | ||
| 118 | +a cube Mmad PLUS an AIV add PLUS the barriers separating them. Over a 5–6 step Neumann chain | ||
| 119 | +that is ~9 AIV ops + their MIX cross-core fences — pure overhead, since the cube can carry | ||
| 120 | +the additive term itself. | ||
| 121 | + | ||
| 122 | +### The fix — inject the additive term through the L0C accumulator | ||
| 123 | +`Mmad`'s `MmadParams.cmatrixInitVal` selects whether the C-matrix is initialized fresh | ||
| 124 | +(`true` / `=1` on the first tile) or ACCUMULATED onto the existing L0C content | ||
| 125 | +(`false` / `=0`). Exploit it to compute `R@(I+P) = R + R@P` in ONE L0C pass: | ||
| 126 | + | ||
| 127 | +1. **First Mmad, `cmatrixInitVal = 1`** — seed the L0C accumulator with the identity/additive | ||
| 128 | + term (the `R` of `R + R@P`; equivalently `R@I`). This lands the additive part IN L0C with | ||
| 129 | + NO vector op. | ||
| 130 | +2. **Second Mmad, `cmatrixInitVal = 0`** — accumulate `R@P` directly onto that same L0C | ||
| 131 | + accumulator. Result in L0C is `R + R@P = R@(I+P)`. | ||
| 132 | +3. Hold ONE L0C accumulator across the product chain; only the final step Fixpipes out. | ||
| 133 | + | ||
| 134 | +No separate `+I` vector-add, no diagonal `Duplicate`, no extra AIV↔AIC barrier for the add. | ||
| 135 | +This is the SAME `cmatrixInitVal=(ki==0)` accumulate idiom the K-tile loop uses in | ||
| 136 | +`fa_class/cv_reference_concrete_params.md` (`matmul_primitive`) — here repurposed to inject | ||
| 137 | +an ALGEBRAIC additive term (the I / R) rather than to accumulate K-tiles. | ||
| 138 | + | ||
| 139 | +### Measured result (structure + op-count) | ||
| 140 | +Cuts AIV ops + barriers to **~6 vs ~9** for a 5-step Neumann chain (the redundant per-step | ||
| 141 | +vector add + its MIX fence removed). Precision unchanged (the additive term is exact in the | ||
| 142 | +fp32 L0C accumulator — arguably MORE faithful than an fp16 round-tripped vector add). | ||
| 143 | +Reference STRUCTURE: the cv-reference `gated_delta_rule` `NeumannSolve`. | ||
| 144 | + | ||
| 145 | +**Cross-backend witness (independent, PR #200 / customer liuyu15819):** the customer's | ||
| 146 | +Kimi-K3 + a5_ops gated_delta_rule CV-fusion run reached the SAME `cmatrixInitVal` L0C-fold | ||
| 147 | +from scratch on **910B2C / arch 220x / CANN 8.5.1** — replacing a 2016-iteration AIV | ||
| 148 | +scalar-AXPY solve (~98µs/chunk) with an 11-Mmad cube chain (~28µs/chunk): T=4096 | ||
| 149 | +8422→2137µs, geomean **3.79×→8.85×** vs PyTorch, MERE ≤9e-4, 16/16. Same mechanism reached | ||
| 150 | +independently on a different backend + SOC-variant (same 220x arch as our Ascend910_9382) = | ||
| 151 | +strong cross-confirmation that this fold is backend-spontaneous, not KB-dependent. See the | ||
| 152 | +reconciled `OL-278` (guarded-MIX base) + `OL-279` (the concrete Fixpipe/Nd2Nz stride units | ||
| 153 | +that make the hand-cube emit runnable) in the source repo (not migrated). | ||
| 154 | + | ||
| 155 | +### Decision rule | ||
| 156 | +- Any `R@(I+P)` / `R + R@P` / `(I+x)^-1` product-chain step on the cube → seed the additive | ||
| 157 | + term with `cmatrixInitVal=1`, accumulate the product with `cmatrixInitVal=0`, on ONE L0C | ||
| 158 | + accumulator held across the chain. Do NOT emit a separate vector `+I` add. | ||
| 159 | +- Watch the Neumann solve SIGN (CAND-GDR-1): the accumulator must realize the reference's | ||
| 160 | + `(I + strict)^-1`, not `(I - strict)^-1` — the fold mechanism is sign-agnostic, so the | ||
| 161 | + sign bug (if present) is orthogonal and must still be handled at the operand level. | ||
| 162 | + | ||
| 163 | +--- | ||
| 164 | + | ||
| 165 | +## Cross-references | ||
| 166 | + | ||
| 167 | +- **P-P117** ([gated_delta_rule_a3_recurrence](gated_delta_rule_a3_recurrence.md)) — the a3 MIX | ||
| 168 | + chunked-recurrence ALGORITHM these two cube decisions optimize; the many-small-matmul | ||
| 169 | + call-site (§2 decomposition) and the Neumann product solve (`T=(I+strict)^-1`) are exactly | ||
| 170 | + where P-P119 and P-P120 apply. | ||
| 171 | +- **P-P116** ([fa_class_a3_mix_template](fa_class_a3_mix_template.md)) — the a3 MIX cube+vector | ||
| 172 | + SYNC template both sit on. | ||
| 173 | +- **P-P68** — single-AIC GEMM via `MatmulImpl` + `IterateAll<sync=true>` (the library primitive | ||
| 174 | + P-P119 says to REPLACE for the many-small case, but KEEP for few-large; source-repo card, not migrated). | ||
| 175 | +- **`fa_class/cv_reference_concrete_params.md`** (`matmul_primitive` decision) — the concrete | ||
| 176 | + `Mmad` 4-arg accumulate form + `cmatrixInitVal=(ki==0)` + per-context `Fixpipe srcStride` | ||
| 177 | + units the hand-rolled primitive (P-P119) and the L0C-fold (P-P120) both emit. Grounding | ||
| 178 | + for the exact param values; do not re-derive (source-repo file, not migrated). | ||
| 179 | +- **CAND-GDR-1** — Neumann solve SIGN gotcha (orthogonal to the P-P120 fold; source-repo ID). | ||
| 180 | +- **CAND-GDR-3** — reused `MatmulImpl` needs `SetOrgShape` per call: another facet of the | ||
| 181 | + per-call library overhead P-P119 sheds by hand-rolling (source-repo ID). | ||
| 182 | +- [cube_vector_fusion](cube_vector_fusion.md) — the a5/arch35 sibling manual-`Mmad` cube-MIX | ||
| 183 | + note (different chip; P-P119/P-P120 are the a3/arch22 chunked-recurrence-specific decisions). | ||
| @@ -0,0 +1,108 @@ | |||
| 1 | +--- | ||
| 2 | +type: Example | ||
| 3 | +title: "Cube-Vector Fusion (MIX) — port_a3 cube-class canonical pattern" | ||
| 4 | +description: "P-P102：A5/Ascend950PR cube 类算子移植的原生 MIX 范式——Mmad cube + vector 收尾 + WorkspaceQueue 跨核同步；纯 VEC kernel 顶替属架构性作弊，finalize 门会判 ARCHITECTURAL_HACK。含 worked example 与可迁移的 WHY。" | ||
| 5 | +tags: [task.implementation, topic.programming_model.execution_unit.cube, topic.implementation.synchronization, topic.attention] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/reference/porter/patterns/cube_vector_fusion.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-08-05T14:23:52Z" | ||
| 11 | +updated_at: "2026-09-16T07:31:58Z" | ||
| 12 | +platforms: ["950"] | ||
| 13 | +--- | ||
| 14 | + | ||
| 15 | +# Cube-Vector Fusion (MIX) — port_a3 cube-class canonical pattern | ||
| 16 | + | ||
| 17 | +> 原 OKF v1 卡号：**P-P102**。 | ||
| 18 | +> 适用范围：soc=Ascend950PR（A5 / V351 / arch35）。本模板教的是构建 A5/V351 MIX kernel—— | ||
| 19 | +> 正文自述 "Cube-class A5 ports use native MIX"；220x/a3 一侧只是被移植的**来源**算法， | ||
| 20 | +> 本模板缝合的设备代码与跨核握手都是 A5 侧。`verified_on: flash_attention_score A5` | ||
| 21 | +> （decision_manifest 2026-05-29）。 | ||
| 22 | + | ||
| 23 | +## Principle (the WHAT) | ||
| 24 | + | ||
| 25 | +**Cube-class A5 ports use native MIX (Mmad cube + vector epilogue + WorkspaceQueue | ||
| 26 | +cross-core sync), never pure-vec.** A pure-VEC kernel generated for a cube-required | ||
| 27 | +op is a HACK (OL-188, same anti-cheat tier as CPU fallback) — the finalize gate | ||
| 28 | +`_check_architecture_class` rejects it as `ARCHITECTURAL_HACK` (PR #316). | ||
| 29 | + | ||
| 30 | +- **applies_to**: `port_a3_to_a5; op_class CUBE_MIX (matmul / attention / conv / rnn / | ||
| 31 | + gmm / ffn families); A5/V351`. | ||
| 32 | +- **verified_on**: `flash_attention_score A5` (`output/a3_to_a5_port/src/kernels/ | ||
| 33 | + flash_attention_score/op_kernel/`, decision_manifest.jsonl 2026-05-29). The concrete | ||
| 34 | + anchors below are distilled from that verified archive so a worker reads them HERE, | ||
| 35 | + not from the archive. | ||
| 36 | + | ||
| 37 | +## Concrete anchor (the worked example, distilled) | ||
| 38 | + | ||
| 39 | +- **File split**: `<op>_cube.h` (cube class) + `<op>_vec.h` (vec class) + `<op>_kernel.h`/ | ||
| 40 | + `.cpp` orchestrator. Class names contain the literal `Cube`/`Vec` (the finalize | ||
| 41 | + pre-build regex + arch-class classifier look for cube markers in these files). | ||
| 42 | +- **Cube primitive**: manual `AscendC::Mmad`. FA used 9 `Mmad` calls; | ||
| 43 | + `MmadParams.cmatrixInitVal=true` for a single K-tile, `cmatrixSource=false`. | ||
| 44 | + **EXECUTABLE params (this anchor is contract-level; the code-level detail lives in | ||
| 45 | + the [`fa-concrete-lowering-params-507015`](../optimizations/fa_concrete_lowering_params_507015.md) | ||
| 46 | + runbook，已迁入本仓 optimizations/):** | ||
| 47 | + the `Mmad` signature, full `MmadParams` with `cmatrixInitVal=(ki==0)` for the K-tile accumulate | ||
| 48 | + loop (not always-true — that's the single-tile case only), the L0A/L0B Nz-fractal layout (16×C0, | ||
| 49 | + C0=16 fp16), the QK-load `ifTranspose` form, the Fixpipe `srcStride` PER-CONTEXT units (`/C0` for | ||
| 50 | + L0C→workspace vs ELEMENTS for L0C→GM), and the 16/C0 block-tile-alignment rule. An agent writing | ||
| 51 | + the cube needs that file — this anchor stops short of the code. | ||
| 52 | + `matmul::Matmul<>` / `MatmulImpl<>` is the library alternative, but FA chose `Mmad` | ||
| 53 | + after V220 evidence that `matmul::Matmul<>` was numerically wrong ~500× — the A5 | ||
| 54 | + path picks the safer verified `Mmad`. (Cross-ref: the standalone `matmul::Matmul<>`/ | ||
| 55 | + KFC path is separately blocked on V220 — CAND-KFC-standalone-bootstrap-teardown, | ||
| 56 | + rendezvous deadlock — so `Mmad` is also the unblocked choice.) | ||
| 57 | +- **Vec epilogue**: the op's vector reduction/activation. FA: `SoftmaxFlashV2` (online | ||
| 58 | + softmax). General cube-class epilogue = the op's reduction/elementwise tail. | ||
| 59 | +- **Cross-core sync**: `WorkspaceQueue<T, RING_SLOTS=3>` ring buffer, **ONE per | ||
| 60 | + producer↔consumer direction** (FA: 3 — S cube→vec, P vec→cube, O cube→vec), with | ||
| 61 | + paired flag IDs and **raw `PIPE_FIX`/`PIPE_MTE3`/`PIPE_MTE2` literals** (CANN 9.0.0 | ||
| 62 | + forbids templated `pipe_t`). **NOT** file-scope inline `CrossCoreSetFlag` in the loop | ||
| 63 | + (Antipattern A → 507015). (Cross-ref OL-190 WorkspaceQueue.) | ||
| 64 | +- **Task type**: `KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2)` — V220-native | ||
| 65 | + (PB-28 FALSIFIED 2026-05-25). Do NOT arch-guard it out. | ||
| 66 | +- **Tiling**: regbase per-AIV sub-block (`vec_num=2` for MIX_AIC_1_2 attention), under | ||
| 67 | + the V351 248 KiB UB budget. | ||
| 68 | + | ||
| 69 | +## The WHY (so a worker can ADAPT, not copy) | ||
| 70 | + | ||
| 71 | +Per the study-CANN-understand-WHY direction: the MIX structure exists because cube | ||
| 72 | +(AIC) and vector (AIV) are **separate cores** that must be kept both busy. The cross- | ||
| 73 | +core sync is a **counting-semaphore** handshake (`WorkspaceQueue` flag IDs) pairing one | ||
| 74 | +AIC producer with the AIV consumer per data direction; getting the AIC/AIV pairing and | ||
| 75 | +the producer-PIPE↔consumer-PIPE matching right is what avoids both deadlock (507014- | ||
| 76 | +class rendezvous) and the lockstep pipeline-collapse (OL-200: cube/vec each ~half idle | ||
| 77 | +when scheduled in lockstep instead of `i%K`-staggered). A worker adapting this to a new | ||
| 78 | +cube op must reproduce the *contract* (one queue per direction, correct PIPE pairing, | ||
| 79 | +staggered schedule), not the FA-specific tile sizes. | ||
| 80 | + | ||
| 81 | +## Generation methodology — understand→KB→research→regenerate | ||
| 82 | + | ||
| 83 | +(NOT line-port V220, NOT pure-vec fallback. Codified in the Layer-2 brief + here.) | ||
| 84 | + | ||
| 85 | +1. **Understand**: read the A3/V220 algorithm source (`op_kernel/*.h` top-level + | ||
| 86 | + `op_host/*.h`) for the algorithm — what cube op, what epilogue, what dataflow. | ||
| 87 | +2. **KB**: consult this pattern + L-tier RegBase MicroAPI KB (OL-143/144) for the A5 | ||
| 88 | + mapping. | ||
| 89 | +3. **Research (only if you cannot generate)**: read the A5 reference for UNDERSTANDING | ||
| 90 | + — allowed per the study-CANN direction; the `ARCH35_WRAP_CHEAT` gate + copy-shape | ||
| 91 | + scanners remain the anti-copy net (understand, never copy verbatim into the kernel). | ||
| 92 | +4. **Generate from understanding** using the scaffold above. | ||
| 93 | +5. **Regenerate** on gate rejection / precision-perf gap by looping back to | ||
| 94 | + understand/KB/research — **NEVER** to a pure-VEC fallback (the gate rejects it). | ||
| 95 | + | ||
| 96 | +## Anti-pattern cross-refs | ||
| 97 | + | ||
| 98 | +- **OL-188** — pure-VEC for a cube-required op = HACK (the finalize gate this pattern | ||
| 99 | + satisfies). | ||
| 100 | +- **OL-190** — WorkspaceQueue ring-buffer cross-core sync. | ||
| 101 | +- **OL-200** — MIX pipeline collapse: lockstep cube/vec leaves each ~half idle; stagger | ||
| 102 | + by `i%K` (precision-correct ≠ perf-correct). | ||
| 103 | +- **EC-57 / EC-58** — per-source `SPLIT_CORE_CUBE`/`SPLIT_CORE_VEC` defines + | ||
| 104 | + `matmul_intf.h` include discipline for the MIX compile. | ||
| 105 | +- **PB-35** — `event_t(0)` cross-Path deadlock. | ||
| 106 | +- **PB-28 FALSIFIED** — `KERNEL_TYPE_MIX_AIC_1_2` is V220-native; do not arch-guard out. | ||
| 107 | + | ||
| 108 | +（OL-* / EC-* / PB-* 原源仓 runbook 卡号，已迁入 ops/ascendc/runbooks/ 与 ops/ascendc/optimizations/。） | ||
| @@ -0,0 +1,113 @@ | |||
| 1 | +--- | ||
| 2 | +type: Example | ||
| 3 | +title: "a3 c220 GM-comm-slot LOOPED bidirectional MIX probe — RUN_LOG" | ||
| 4 | +description: "a3/220x 上 GM 通信槽式双向 MIX 跨核握手的实测日志：共享广播 flag 12 + 按 block_idx 分槽 + dcci/dsb 能否让 8 轮循环的 cube↔vec 握手免死锁（对照每对不同 flag id 的循环已死锁 507014），含 CANN 9.0.0 a3 同步头文件实测细节。" | ||
| 5 | +tags: [task.debugging, topic.implementation.synchronization, topic.attention] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/reference/porter/patterns/fa_a3_multicore_gmslot_run_log.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-08-05T14:23:52Z" | ||
| 11 | +updated_at: "2026-09-04T23:59:25Z" | ||
| 12 | +platforms: [a3] | ||
| 13 | +--- | ||
| 14 | + | ||
| 15 | +# a3 c220 GM-comm-slot LOOPED bidirectional MIX probe — RUN_LOG | ||
| 16 | + | ||
| 17 | +> 适用范围：soc=Ascend910_9382（a3 / 220x / arch22），CANN 9.0.0 a3 容器。一次负面结果实测： | ||
| 18 | +> GM-slot 配方**没有**解开 LOOPED 双向 MIX 死锁。 | ||
| 19 | + | ||
| 20 | +DECISIVE QUESTION: does the shipped c220 GM-comm-slot mechanism (single shared broadcast | ||
| 21 | +flag + per-block GM slot + dcci/dsb) make a LOOPED (>=8 rounds) bidirectional multi-core | ||
| 22 | +MIX cube<->vec handshake run DEADLOCK-FREE on a3/220x, where the naive per-sub DISTINCT | ||
| 23 | +flag-id loop DEADLOCKED (mode3, 507014)? | ||
| 24 | + | ||
| 25 | +Env: host a3-host, a3 CANN 9.0.0 container, NPU **the a3 NPU ONLY** (ASCEND_RT_VISIBLE_DEVICES=<n>). | ||
| 26 | +the a3 NPU (the a3 NPU, ) verified CLEAN at start: AICore 0%, HBM 3098MB, no proc. | ||
| 27 | + | ||
| 28 | +## Prior (mode4_probe/RUN_LOG.md, device-measured) | ||
| 29 | +- mode0 1-shot distinct-id: PASS deadlock-free | ||
| 30 | +- mode1 1-shot shared-id: PASS deadlock-free | ||
| 31 | +- mode3 LOOPED distinct-id (8 rounds): DEADLOCK 507014 (device latched AICore=100%) | ||
| 32 | +- mode2 LOOPED shared-id: NOT RUN (device latched first) | ||
| 33 | + | ||
| 34 | +## KB findings this session (cann-9.0.0 a3 headers) — measured | ||
| 35 | +### pto/npu/a2a3/custom/TSyncCVID.hpp — TSYNC_CVID (c2v ONLY) | ||
| 36 | + CV_COMM_CTRL=12 (13,14,15 reserved). kCvCommSlotBytes=512. kCvMaxCores=25. | ||
| 37 | + CVSyncMode: C_ALL_CORE_SYNC=0, V_SUBCORES_SYNC=1, CV_CORE_SYNC=2. | ||
| 38 | + _getFFTSMsg(mode,flag,base=1) = (base&0xf) + ((mode&3)<<4) + ((flag&0xf)<<8). | ||
| 39 | + cube: write comm_slot -> GM[block_idx*512]; dcci(slot,SINGLE_CACHE_LINE); dsb(DSB_DDR); | ||
| 40 | + ffts_cross_core_sync(PIPE_MTE2, msg(CV_CORE_SYNC=2, CV_COMM_CTRL=12)). | ||
| 41 | + vec: dcci(slot); wait_flag_dev(12); read GM slot. | ||
| 42 | + => single SHARED broadcast flag id 12, disambiguation by GM SLOT (block_idx), NOT distinct ids. | ||
| 43 | + | ||
| 44 | +### pto/npu/a2a3/custom/TSync_Custom.hpp — bidirectional record/wait/allocate/free | ||
| 45 | + c2v record: cube ffts(PIPE_FIX, msg(CV_CORE_SYNC, flag)); vec wait_flag_dev(flag). | ||
| 46 | + v2c record: vec ffts(PIPE_MTE3, msg(CV_CORE_SYNC, flag)); cube wait_flag_dev(flag). | ||
| 47 | + backward allocate/free use flag+1. All SHARED broadcast (CV_CORE_SYNC=2). | ||
| 48 | + | ||
| 49 | +### pto/npu/a2a3/TPush.hpp — full shipped FIFO (canonical FA bidirectional) | ||
| 50 | + "both" dir uses 4 flags: FlagID (c2v fwd), FlagID+1 (c2v free), FlagID+2 (v2c fwd), FlagID+3 (v2c free). | ||
| 51 | + Each a SINGLE SHARED CV_CORES_SYNC=2 broadcast. SyncPeriod credit-batching (shouldWaitFree/shouldNotifyFree) | ||
| 52 | + => syncs only every SyncPeriod tiles, NOT every round -> avoids per-round orphan-credit accumulation. | ||
| 53 | + KEY INSIGHT: shipped recipe = shared broadcast per direction + GM-slot data + PERIODIC (batched) credit sync. | ||
| 54 | + | ||
| 55 | +## Design of GM-slot looped mode (this probe) | ||
| 56 | +Reuse mode4_probe/kernel base. Add mode 5 = LOOPED bidirectional, GM-slot recipe: | ||
| 57 | + per round r in 0..T-1 (T>=8): | ||
| 58 | + c2v: cube does trivial matmul -> GM data slot; dcci+dsb; ffts(PIPE_MTE2, msg(CV_CORE_SYNC, 12)); vec wait_flag_dev(12). | ||
| 59 | + vec: trivial vec op; reverse: intra-vec sub barrier, subblock0 ffts(PIPE_MTE3, msg(CV_CORE_SYNC, 13)); cube wait_flag_dev(13). | ||
| 60 | + Single shared broadcast flags 12 (c2v) + 13 (v2c). Both AIVs consume the c2v broadcast (balanced fanout-2); | ||
| 61 | + reverse aggregated to 1 via subblock0-only issue (balanced 1:1). GM-slot per block_idx*512. | ||
| 62 | +Keep mode3 (distinct-id looped) as known-DEADLOCK CONTROL. | ||
| 63 | + | ||
| 64 | +## Timeline | ||
| 65 | +23:47:58Z RUN_LOG created, headers grepped | ||
| 66 | + | ||
| 67 | +$(date -u +%H:%M:%SZ) Kernel design finalized. Building mode 5. | ||
| 68 | +KEY INSIGHT (grounded in headers): the DECISIVE anti-deadlock variable is single SHARED | ||
| 69 | +broadcast flag id per direction with BALANCED consumption: | ||
| 70 | + - c2v: cube ONE set (CV_CORE_SYNC=2 broadcast) -> BOTH subs wait same id (broadcast delivers | ||
| 71 | + a credit to each recipient; mode1 1-shot proved 1 set satisfies 2 waits). Zero orphan. | ||
| 72 | + - v2c: each sub sets same id -> cube waits it N(=2) times. Balanced. | ||
| 73 | + vs mode3 DISTINCT ids: cube's broadcast of id X physically reaches BOTH subs but only 1 waits X | ||
| 74 | + -> sibling orphan accumulates each round -> FFTS stall. That is the 507014 root. | ||
| 75 | +GM-slot's role = lets you use a SHARED flag (data disambiguated by block_idx*512 slot, not by | ||
| 76 | +distinct flag id). Faithful mode5 = shared flags CV_COMM_CTRL=12 (c2v) + 13 (v2c) + per-block | ||
| 77 | +GM slot write + DataCacheCleanAndInvalid (AscendC API; raw dcci/dsb are bisheng c220 builtins, | ||
| 78 | +using the portable AscendC cache API instead). mode2 (shared, no gm-slot) = flag-equivalent control. | ||
| 79 | +RUN ORDER (device budget: hang LATCHES): candidate mode5 FIRST on clean device. If PASS -> | ||
| 80 | +decisive positive. Then mode3 control LAST (expected to latch = accepted end state). | ||
| 81 | + | ||
| 82 | +23:56:17Z BUILD SUCCESS: mode5 GM-slot kernel compiles clean on bisheng arch22 (Ascend910_9382). | ||
| 83 | + DataCacheCleanAndInvalid<int32_t,SINGLE_CACHE_LINE,CACHELINE_OUT>, CV_COMM_CTRL=12/13, GM slot all OK. | ||
| 84 | + .so = _mix2_ext.cpython-311-aarch64-linux-gnu.so (9.2MB). Running mode5 (CANDIDATE) on clean the a3 NPU. | ||
| 85 | + | ||
| 86 | +## RE-RUN verdict (gmslot-rerun-kw-2) | ||
| 87 | +Host=a3-host , a3 CANN 9.0.0 container. | ||
| 88 | +Probe files synced at <scratch>/mode4_probe (build_gmslot.sh, test_mix2.py, kernel/ present, Jul 20 07:54). | ||
| 89 | +STEP1 the a3 NPU clean gate: npu-smi the a3 NPU () AICore=0%, HBM=3097/65536MB (baseline), | ||
| 90 | + "No process in device", no python/aclrt/test_mix procs. => the a3 NPU CLEAN/IDLE, NOT latched. Proceeding. | ||
| 91 | +STEP2 build: BUILD SUCCESS. _mix2_ext.cpython-311-aarch64-linux-gnu.so (9276648 bytes) at Jul 20 08:19. | ||
| 92 | +STEP3 running mode5 CANDIDATE (hard 90s timeout), ASCEND_RT_VISIBLE_DEVICES=<n>... | ||
| 93 | + | ||
| 94 | +### FINAL RE-RUN VERDICT | ||
| 95 | +STEP3 mode5 CANDIDATE: **HANG** (EXIT_CODE=124, timeout 90s). stdout captured: | ||
| 96 | + "torch_npu DEV=npu:0 visible=2 mode=5 (LOOPED GM-comm-slot ...) CANDIDATE" | ||
| 97 | + "[mode=5] launching blockDim=2 (2 MIX groups)..." | ||
| 98 | + -> then NO further output; no "RESULT mode=5" line. Host sync blocked >90s = deadlock signature. | ||
| 99 | + max_rel=n/a, latency=n/a (never returned). | ||
| 100 | +the a3 NPU end-state AFTER python kill: AICore=0%, HBM=3098MB (baseline), "No process in device" | ||
| 101 | + => device RECOVERED cleanly on process teardown this run (no persistent 507014 latch observed; | ||
| 102 | + the timeout-kill of python tore down the aclrt context and freed the device). No >120s recovery needed. | ||
| 103 | +STEP4 mode3 control: **SKIPPED** (brief: only run if mode5 PASSED; it hung). | ||
| 104 | + | ||
| 105 | +INTERPRETATION: mode5 GM-comm-slot recipe (single SHARED broadcast flag 12/13 per direction + | ||
| 106 | +per-block GM slot + DataCacheCleanAndInvalid, LOOPED >=8 rounds bidirectional, blockDim=2) ALSO | ||
| 107 | +DEADLOCKS on a3/220x (Ascend910_9382). The shared-flag + GM-slot disambiguation does NOT escape the | ||
| 108 | +LOOPED-bidirectional MIX deadlock limit. Consistent with prior device witness: single-shot concurrent | ||
| 109 | +MIX works, LOOPED bidirectional (full-FA class) hangs. => GM-slot recipe is NOT the missing key; | ||
| 110 | +the deadlock is deeper than distinct-vs-shared flag-id orphan-credit. Full-FA multicore on a3 via this | ||
| 111 | +path remains UNACHIEVED. Device did recover cleanly (good for budget), but verdict is negative. | ||
| 112 | + | ||
| 113 | +STRUCTURED: MODE5=HANG; max_rel=n/a; latency_ms=n/a; MODE3_CONTROL=SKIPPED; device2_end_state=clean(recovered-post-kill) | ||
| @@ -0,0 +1,96 @@ | |||
| 1 | +--- | ||
| 2 | +type: Example | ||
| 3 | +title: "RESULT — a3 multi-core FA via library/KFC path" | ||
| 4 | +description: "a3/arch22 多核 FlashAttention 走 library/KFC 路线的结果：MatmulImpl IterateAll<sync=true> 加每对一次 MODE2 握手扩到多核，精度 20/20 PASS 且与单核逐位一致；性能为 vendor 的 0.186 倍（20 核）。" | ||
| 5 | +tags: [task.implementation, topic.attention, topic.parallelism] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/reference/porter/patterns/fa_a3_multicore_result.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-08-05T14:23:52Z" | ||
| 11 | +updated_at: "2026-09-04T23:59:25Z" | ||
| 12 | +platforms: [a3] | ||
| 13 | +--- | ||
| 14 | + | ||
| 15 | +# RESULT — a3 multi-core FA via library/KFC path | ||
| 16 | + | ||
| 17 | +> 适用范围：arch22 / Ascend910_9382 / CANN 9.0.0 / the a3 NPU。Started 2026-07-19。 | ||
| 18 | +> 源仓中本卡的精度与性能原始数据分别记录在同目录 `fa_a3_multicore_mc_verify.json`（20 行 | ||
| 19 | +> 精度明细）与 `fa_a3_multicore_mc_perf.json`（性能计时）两个证据文件里；目标仓不提交该类 | ||
| 20 | +> 数据资产，关键数值已全部并入下文正文。 | ||
| 21 | + | ||
| 22 | +## Structured verdict (fill as stages complete) | ||
| 23 | +``` | ||
| 24 | +STUDY=done | ||
| 25 | +BUILD=PASS | ||
| 26 | +npu_fusion_attention_on_a3=YES | ||
| 27 | +precision=PASS 20/20 (multi-core, deterministic, bit-identical to 1-core) | ||
| 28 | +perf_vs_vendor=0.186x multi-core (20 cores); 0.0146x single-core | ||
| 29 | +verdict=multi-core library-cube FA runs deadlock-free on a3; perf gap structural (non-flash) | ||
| 30 | +``` | ||
| 31 | + | ||
| 32 | +## Stage log | ||
| 33 | + | ||
| 34 | +### STUDY = done (2026-07-19) | ||
| 35 | +See STUDY.md. Key findings: | ||
| 36 | +- The brief's "vendor FA = 0-CrossCore async-KFC library" path is documented-BLOCKED on standalone a3 | ||
| 37 | + (507014; CAND-FA-A3-PERF-STRUCTURAL-1 / CAND-KFC-standalone-bootstrap / PB-53 / PB-54). Not attempted. | ||
| 38 | +- The VERIFIED-working library route = `MatmulImpl IterateAll<sync=true>` + ONE per-pair PB-55 handshake | ||
| 39 | + (single-core base, 17/17 PASS). This IS a genuine library cube (not hand-Mmad, not a hand-rolled | ||
| 40 | + cross-core ring). Phase-1 plan = extend it to multi-core (blockDim>1, one head-slice/core), keeping | ||
| 41 | + only the minimal sanctioned per-pair MODE2 (CV_CORE_SYNC) handshake — NO added cross-core flags. | ||
| 42 | +- earlier_fa_prototype "vendor" files are actually a hand-rolled CrossCore ring (the anti-pattern) measured at | ||
| 43 | + 0.25x vendor — used only as a negative reference. | ||
| 44 | +- npu_fusion_attention DID run on an a3 device (58-61us, S=1024, 2026-05-30 per CAND); re-verify on the a3 NPU. | ||
| 45 | + | ||
| 46 | +### BUILD = PASS (2026-07-19) | ||
| 47 | +- Host a3-host, a3 CANN 9.0.0 container, CANN 9.0.0, SOC=Ascend910_9382, bisheng. | ||
| 48 | +- `cmake -B build -DSOC_VERSION=Ascend910_9382 -DASCEND_CANN_PACKAGE_PATH=$CANN_PATH | ||
| 49 | + then `cmake --build build -j8`. BUILD_OK **first try**, no compile-fix iters. | ||
| 50 | +- Produced `kernel/build/_fusion_attention_ext.cpython-311-aarch64-linux-gnu.so`. | ||
| 51 | +- Multi-core kernel = single-core base extended to blockDim>1 (one head-slice/core), | ||
| 52 | + keeping ONLY the per-pair MODE2=CV_CORE_SYNC PB-55 handshake; no cross-core flags. | ||
| 53 | + | ||
| 54 | +### DEVICE-CLEAN (the a3 NPU) = PASS | ||
| 55 | +- npu-smi the a3 NPU (the a3 NPU, both chips): AICore 0%, "No process in device". Clean. | ||
| 56 | + | ||
| 57 | +### npu_fusion_attention on a3 = YES (2026-07-19) | ||
| 58 | +- the a3 NPU, BNSD B4N8S512D64 fp16. `torch_npu.npu_fusion_attention(q,k,v,N,"BNSD",scale,keep_prob=1.0)` | ||
| 59 | + RAN and produced output. 配套 perf 证据数据记录 median 102.3us / min 95.4us，单核探针中为 115us。 | ||
| 60 | +- Settles the contested "a3 has no vendor FA baseline" question: vendor CANN FA DOES run on a3 today. | ||
| 61 | + | ||
| 62 | +### PRECISION = PASS 20/20 (multi-core, deadlock-free) | ||
| 63 | +- 精度验证脚本跑 5 shapes × 4 core-counts {1,2,4,auto}，20 行明细全部记录在配套 verify 证据数据中： | ||
| 64 | + 全部 20 行 prec=PASS, det=True, eq1core=True。 | ||
| 65 | + - cos 0.999998..1.000015, max_abs 3.3e-4..5.7e-4 vs fp32 dense-SDPA golden。 | ||
| 66 | + - determinism: every call bit-exact on re-run。 | ||
| 67 | + - core-count equivalence: multi-core output BIT-IDENTICAL to nc=1 base at every shape。 | ||
| 68 | +- **The §5 looped-deadlock did NOT materialize.** blockDim>1 (nc=2,4,20) ran clean. Confirms the | ||
| 69 | + STUDY hypothesis: MODE2=CV_CORE_SYNC scopes FLAG_S/FLAG_P to each core's OWN AIC<->AIV pair → | ||
| 70 | + zero cross-core flag interaction → each core is an independent, already-verified single-core FA. | ||
| 71 | + (The §5 deadlocks were HAND-ROLLED cross-core looped rings; this design has none.) | ||
| 72 | + | ||
| 73 | +### PERF vs vendor = measured (mine slower; honest) | ||
| 74 | +Shape B4N8S512D64 (BN=32), the a3 NPU, median device-time (sync-bracketed): | ||
| 75 | +- vendor npu_fusion_attention : ~102us | ||
| 76 | +- mine multi-core auto (20 AIC cores) : 550.8us → ratio vendor/mine = **0.186x** (mine 5.4x slower) | ||
| 77 | +- mine single-core (nc=1) : 7866us → ratio vendor/mine = **0.0146x** (mine 68x slower) | ||
| 78 | +- Multi-core parallelization gave **~14.3x** over single-core (7866→550us; 32 heads over 20 cores). | ||
| 79 | +- Gap to vendor is STRUCTURAL, not a handshake bug: this kernel is NON-flash — it materializes the | ||
| 80 | + full [S,S] score + [S,S] prob matrices to GM per head and does a row-at-a-time scalar softmax. | ||
| 81 | + The genuine library cube (IterateAll) beat the earlier hand-Mmad port (0.25x) on the cube side, | ||
| 82 | + but the full-matrix-to-GM traffic + row-serial softmax dominate. Vendor flash-attention keeps | ||
| 83 | + scores in L1/UB and fuses the online softmax — a different algorithm class. | ||
| 84 | + | ||
| 85 | +## STRUCTURED VERDICT | ||
| 86 | +``` | ||
| 87 | +BUILD=PASS (first try, bisheng SOC=Ascend910_9382) | ||
| 88 | +npu_fusion_attention_on_a3=YES (~102us, B4N8S512D64) | ||
| 89 | +single_core_vs_vendor=0.0146x (7866us vs 115us) | ||
| 90 | +multicore=PASS 20/20 cos>=0.999998 ratio 0.186x (550us vs 102us; deadlock-free at blockDim=2/4/20; bit-identical to 1-core) | ||
| 91 | +verdict=Library-cube (sync IterateAll) multi-core FA BUILDS + RUNS deadlock-free on a3 standalone | ||
| 92 | + pybind, precision 20/20, deterministic. The per-pair MODE2 handshake replicated across | ||
| 93 | + cores does NOT hit the §5 looped-deadlock (that was a hand-rolled cross-core ring, avoided | ||
| 94 | + here). Perf 0.186x vendor — gap is structural (non-flash, full-matrix-to-GM + row softmax), | ||
| 95 | + not a sync defect; multi-core still recovered 14x over single-core. | ||
| 96 | +``` | ||
| @@ -0,0 +1,207 @@ | |||
| 1 | +--- | ||
| 2 | +type: Example | ||
| 3 | +title: "FlashAttention-class MIX Template Domain (a3 / arch22) — GENERATE a cube+vector attention op (kernel + host)" | ||
| 4 | +description: "P-P116：a3/arch22 专用的 FlashAttention 类 MIX_AIC_1_2 注意力算子生成模板（kernel + host），教结构与关键行而非整体拷贝；MIX 派发不加 KERNEL_TASK_TYPE_DEFAULT 宏，动笔前先跑通最小 AIC↔AIV SYNC witness。禁止用于 a5。" | ||
| 5 | +tags: [task.implementation, topic.attention, topic.implementation.synchronization, topic.programming_model.execution_unit.cube] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/reference/porter/patterns/fa_class_a3_mix_template.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-08-05T14:23:52Z" | ||
| 11 | +updated_at: "2026-09-04T23:59:25Z" | ||
| 12 | +platforms: [a3] | ||
| 13 | +--- | ||
| 14 | + | ||
| 15 | +# FlashAttention-class MIX Template Domain (a3 / arch22) — GENERATE a cube+vector attention op (kernel + host) | ||
| 16 | + | ||
| 17 | +> 原 OKF v1 卡号：**P-P116**。`op_class=attention-fwd / CUBE_MIX (flash-attention family, a3 hand-authored)`。 | ||
| 18 | + | ||
| 19 | +> ## ⚠ ARCH SCOPE — this template is `arch22` / `220x` / Atlas A3 **ONLY** | ||
| 20 | +> `arch_scope: arch22 · 220x · soc=Ascend910_9382 · product=Atlas A3 · dav-2201 · cann=9.0.0`。 | ||
| 21 | +> MIX 派发（不加 KERNEL_TASK_TYPE_DEFAULT 宏）与 AIC↔AIV 跨核握手在 Ascend910_9382（dav-2201, | ||
| 22 | +> arch22）上设备实证，**从未在 a5 上验证**。 | ||
| 23 | +> **NOT for an `arch35` / Ascend950PR / Atlas A5 FA target.** The a5 side has its OWN templates — | ||
| 24 | +> [cube_vector_fusion](cube_vector_fusion.md) (**P-P102**) and [fa_class_template](fa_class_template.md) | ||
| 25 | +> (**P-P103**). Their recipes (arch35 §4 sync mode-4, the `KERNEL_TASK_TYPE_DEFAULT` macro, C310 | ||
| 26 | +> Fixpipe, MicroAPI regbase softmax) are verified on Ascend950PR and are WRONG for a3. | ||
| 27 | + | ||
| 28 | +## What this template is (and how to use it) | ||
| 29 | + | ||
| 30 | +This is a **REUSABLE, GENERATION-TEACHING** template, not a copy source. It teaches the *pattern* of an a3 cube+vector MIX attention op — the structure, the load-bearing exact lines, the axes you generalize along, and the host side — so a worker GENERATES its own kernel + host for its own op. It is deliberately the OPPOSITE of a whole-body copy: there is no liftable op body here. | ||
| 31 | + | ||
| 32 | +**How a worker uses it:** | ||
| 33 | + | ||
| 34 | +1. **READ** this template to learn the MIX pattern + the host-generation guidance. | ||
| 35 | +2. **BUILD + RUN** the minimal a3 MIX SYNC-witness (see §6) once, to *witness the AIC↔AIV handshake close deadlock-free* — proof beats prose before you author. | ||
| 36 | +3. **GENERATE** your op's kernel + host from the knowledge here, **using shipped-library primitives** (see the rule below). Do NOT copy the witness body. | ||
| 37 | + | ||
| 38 | +### ★ Reference SHIPPED libraries; generate only the op-glue | ||
| 39 | + | ||
| 40 | +The customer's build environment has the CANN library and catlass. **USE them — do not reinvent the wheel:** | ||
| 41 | + | ||
| 42 | +| need | USE this SHIPPED primitive (encouraged) | do NOT hand-roll | | ||
| 43 | +|---|---|---| | ||
| 44 | +| cube matmul (`S=Q@Kᵀ`, `O=P@V`) | CANN `MatmulImpl<>` + `IterateAll<sync=true>` (`#include "lib/matmul_intf.h"`) | a hand-written GEMM, or a pure-vector emulation (OL-188, forbidden) | | ||
| 45 | +| row-wise softmax | CANN library `AscendC::SoftMax<T>(...)` + `SoftMaxTiling` / `GetSoftMaxMinTmpSize` (memory-base-vector lib) | our `MixSoftmax` + `RowReduceMax`/`RowReduceSum` wrappers | | ||
| 46 | +| low-level row reduce (if you hand-fold) | library intrinsics `WholeReduceMax` / `WholeReduceSum` (fp32) | a bespoke reduce tree you present as reusable | | ||
| 47 | +| AIC↔AIV handshake | CANN `CrossCoreSetFlag<MODE,PIPE>()` / `CrossCoreWaitFlag()` (catlass `catlass/arch/cross_core_sync.hpp`) | a hand-rolled scratch-flag polling loop | | ||
| 48 | + | ||
| 49 | +**BAN (the customer will NOT have these):** CANN *internal* `.cpp`, headers NOT in the shipped CANN library, and **OUR example's hand-written helpers** — `RowReduceMax` / `RowReduceSum` / `Align8` / `FloorPow2` / `MakeMixCfg` / the `MixCubeQK` / `MixCubePV` / `MixSoftmax` wrappers. These are **op-glue you GENERATE yourself** (thin wrappers over the library APIs above), NOT liftable library code. | ||
| 50 | + | ||
| 51 | +> **The agent host needs NO CANN.** Compilation happens on the NPU server (which has CANN + catlass). Your generated code freely `#include`s CANN / catlass headers (`lib/matmul_intf.h`, the softmax lib, `catlass/arch/cross_core_sync.hpp`); they resolve at BUILD time on the server. Reference library headers without hesitation. | ||
| 52 | + | ||
| 53 | +## verified_on | ||
| 54 | + | ||
| 55 | +`verified_on: soc=Ascend910_9382; cann=9.0.0; 2026-07-18; DS famix (single-head FA core) + famix_mh (multi-head)`. Device re-confirmed on the CANN **9.0.0** container: MIX builds + runs deadlock-free + precision-passes (cosine 0.99999905, `allclose` True, deterministic). | ||
| 56 | + | ||
| 57 | +--- | ||
| 58 | + | ||
| 59 | +## 1. Arch scope (restate in your generated op) | ||
| 60 | + | ||
| 61 | +- Target `arch22` / `dav-2201`, `SOC_VERSION=Ascend910_9382`, Atlas A3. | ||
| 62 | +- Reusable TEMPLATE, NOT a copy — you generate the op-specific kernel + host. | ||
| 63 | +- This closes the a3 half of **DEBT-222** (the a3 FA-class starting knowledge the KB previously lacked). | ||
| 64 | + | ||
| 65 | +## 2. The MIX pattern — PSEUDOCODE (the structure you generate) | ||
| 66 | + | ||
| 67 | +One `__global__ __aicore__` entry. The default MIX launch reaches BOTH the AIC and the AIV cores; `if ASCEND_IS_AIC` / `if ASCEND_IS_AIV` partition the work inside it. The two cores rendezvous through two cross-core flags: | ||
| 68 | + | ||
| 69 | +``` | ||
| 70 | +__global__ mix_attn(q, k, v, o, s_scratch, p_scratch, seq, d, scale): | ||
| 71 | + | ||
| 72 | + if AIC: # cube core | ||
| 73 | + cube1: S = Q @ Kᵀ via library MatmulImpl, IterateAll<sync=true> | ||
| 74 | + CrossCoreSetFlag(FLAG_S) # publish S — BROADCAST, releases both AIVs | ||
| 75 | + CrossCoreWaitFlag(FLAG_P) # block until AIV has written P | ||
| 76 | + cube2: O = P @ V via library MatmulImpl, IterateAll<sync=true> | ||
| 77 | + | ||
| 78 | + if AIV: # vector core (1:2 paired to the AIC) | ||
| 79 | + CrossCoreWaitFlag(FLAG_S) # block until S ready | ||
| 80 | + P = softmax(scale * S) row-wise, fp32 accumulate via library SoftMax | ||
| 81 | + CrossCoreSetFlag(FLAG_P) # publish P — see §3(a): BOTH subblocks must set | ||
| 82 | +``` | ||
| 83 | + | ||
| 84 | +- `S`,`P` are `[seq,seq]` scratch tensors in **GM** (allocated host-side, passed in). Softmax works one row at a time in UB. | ||
| 85 | +- The forward flag (`FLAG_S`, AIC→AIV) is **broadcast**; the reverse flag (`FLAG_P`, AIV→AIC) is **per-subblock-counted** — this asymmetry is the whole trap (§3a). | ||
| 86 | + | ||
| 87 | +## 3. Minimal NON-NEGOTIABLE snippets ONLY | ||
| 88 | + | ||
| 89 | +These exact lines are **load-bearing** — use the CANN/catlass library APIs shown; everything around them you generate. Grounded verbatim in the device-proven `famix` probe / the `examples/a3_mix_fa_min` witness. | ||
| 90 | + | ||
| 91 | +**(a) PB-55 — BOTH AIV subblocks set the reverse flag (uniquely correct; single-setter DEADLOCKS).** The reverse AIV→AIC handshake in `MIX_AIC_1_2` is per-subblock-COUNTED: the single AIC `CrossCoreWaitFlag(FLAG_P)` requires a set from EVERY AIV subblock of the 1:2 pair. Both subblocks compute the (identical → benign) softmax and BOTH set the flag: | ||
| 92 | + | ||
| 93 | +```cpp | ||
| 94 | +if ASCEND_IS_AIV { | ||
| 95 | + CrossCoreWaitFlag(FLAG_S); // forward flag is broadcast: one AIC set releases both AIVs | ||
| 96 | + /* ...softmax P = softmax(scale*S), via library SoftMax... */ | ||
| 97 | + CrossCoreSetFlag<MIX_SYNC_MODE2, PIPE_MTE3>(FLAG_P); // BOTH subblocks reach here — NO `if (GetSubBlockIdx()==0)` guard | ||
| 98 | +} | ||
| 99 | +``` | ||
| 100 | + | ||
| 101 | +Raising `FLAG_P` from only subblock 0 → AIC waits forever (measured DEADLOCK, no fault code). `PIPE_MTE3` because the producer of P is the `DataCopy` that drains it to GM. | ||
| 102 | + | ||
| 103 | +**(b) `MatmulImpl<>` + `IterateAll<sync=true>` — the genuine (non-async) library cube invocation shape.** Both matmuls use this; NEVER the async KfcServer `Iterate()`/`GetTensor()` path (PB-34 deadlock): | ||
| 104 | + | ||
| 105 | +```cpp | ||
| 106 | +MatmulImpl<AT, BT, CT, BiasT, MM_CFG> mm; // AT/BT/CT = MatmulType<GM, ND, half, /*ISTRANS=*/false> | ||
| 107 | +mm.Init(&tiling, &pipe); | ||
| 108 | +mm.SetTensorA(aG, /*isTransposeA=*/false); | ||
| 109 | +mm.SetTensorB(bG, /*isTransposeB=*/true); // S=Q@Kᵀ: transpose is a RUNTIME bool on K; the static ISTRANS flag stays false on all three operands | ||
| 110 | +mm.SetSingleShape(M, N, K); | ||
| 111 | +mm.template IterateAll<true>(outG, 0, false, false, false); // sync=true, no atomic | ||
| 112 | +mm.End(); | ||
| 113 | +``` | ||
| 114 | + | ||
| 115 | +(`MM_CFG` / `tiling` are op-glue you fill from your shapes — generate them, don't lift our `MakeMixCfg`.) | ||
| 116 | + | ||
| 117 | +**(c) arch22 dispatch — NO task-type macro.** Do NOT emit `KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2)`. That macro is arch35-only and is REJECTED on Ascend910_9382 → `ACL_ERROR_RT_PARAM_INVALID = 107000` at `RegisterAscendBinary` (PB-28). The default MIX dispatch already reaches both cores: | ||
| 118 | + | ||
| 119 | +```cpp | ||
| 120 | +extern "C" __global__ __aicore__ void mix_attn(/* ... */) { | ||
| 121 | + if ASCEND_IS_AIC { /* cubes + FLAG_S set / FLAG_P wait */ } | ||
| 122 | + if ASCEND_IS_AIV { /* softmax + FLAG_S wait / FLAG_P set */ } | ||
| 123 | +} | ||
| 124 | +``` | ||
| 125 | + | ||
| 126 | +> **Macro coupling (load-bearing — the GDR 507014 trap, PB-57).** The "NO task-type macro" | ||
| 127 | +> rule above is correct **for THIS template's path** — a *scoped* `CrossCoreSetFlag`/`WaitFlag` | ||
| 128 | +> MODE2 handshake (§(d) below), which is built around the 1:2 default asymmetry. It does NOT | ||
| 129 | +> generalize to a *whole-device* `SyncAll<false>()` barrier: that path REQUIRES the 1:1 macro | ||
| 130 | +> `KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_1)`, because with the no-macro 1:2 default the | ||
| 131 | +> idle AIV cores skip the per-work-item loop and never issue their `SyncAll` → 507014 deadlock | ||
| 132 | +> (PB-57). So `KERNEL_TYPE_MIX_AIC_1_1` **IS a valid arch22 macro** (PB-28 declares it | ||
| 133 | +> out-of-scope of the 107000 ban); only `MIX_AIC_1_2` / `*_ONLY` are arch35-only. Decision rule: | ||
| 134 | +> **scoped CrossCore MODE2 → no macro (this template); whole-device `SyncAll` → 1:1 macro | ||
| 135 | +> (PB-57 / P-P117). NEVER mix "no macro (1:2)" with whole-device `SyncAll`.** | ||
| 136 | +> | ||
| 137 | +> **⚠️ Generated scoped-1:2 determinism caveat (PB-57 (b')).** The "scoped MODE2 → | ||
| 138 | +> 1:2, no macro" rule above is DEADLOCK-correct, but a generated kernel using | ||
| 139 | +> `MIX_AIC_1_2` plus many per-id `CrossCoreSetFlag<0x2>` handshakes became run-to-run | ||
| 140 | +> nondeterministic above roughly 12 concurrent MIX blocks. For generated FA-class kernels that | ||
| 141 | +> need determinism, select the balanced `KERNEL_TYPE_MIX_AIC_1_1` template. No whole-device | ||
| 142 | +> `SyncAll` is needed: device evidence (N=30) showed all 16 cases deterministic after changing | ||
| 143 | +> only the AIC:AIV pairing balance. Keep the per-id flags and WAR barriers unchanged; see PB-57 | ||
| 144 | +> (b') for the full evidence. | ||
| 145 | + | ||
| 146 | +**(d) Distinct flag ids, NEVER 0 (PB-35).** `event_t(0)` collides with the cube-internal pipe-sync chain → silent hang. Use the user range 1..7, one id per direction: | ||
| 147 | + | ||
| 148 | +```cpp | ||
| 149 | +constexpr uint8_t MIX_SYNC_MODE2 = 2; // mode 2 suffices on arch22 (NOT the arch35 §4 mode-4 disjoint-id recipe) | ||
| 150 | +constexpr uint16_t FLAG_S = 4; // AIC→AIV | ||
| 151 | +constexpr uint16_t FLAG_P = 5; // AIV→AIC | ||
| 152 | +``` | ||
| 153 | + | ||
| 154 | +The AIC raises `FLAG_S` with `PIPE_FIX` (its producer is the cube fixpipe): `CrossCoreSetFlag<MIX_SYNC_MODE2, PIPE_FIX>(FLAG_S);`. | ||
| 155 | + | ||
| 156 | +## 4. Generalization — the axes you extend along | ||
| 157 | + | ||
| 158 | +The proven skeleton is single-head, `seq≤384`, `d=64`, fp16, non-flash, no mask. Your op generalizes along these axes (prose + tiny illustrative pseudocode — you write the full impl): | ||
| 159 | + | ||
| 160 | +- **Multi-head / batch** — a device-side loop over `nheads = B*H`, per-head Q/K/V/O offset `bh*seq*d`, reusing ONE `[seq,seq]` S/P scratch pair across all heads. The `FLAG_S`/`FLAG_P` chain serializes access to that shared scratch, so it is safe for any `nheads` (iter *i+1*'s cube#1 overwrites S only after AIV read it; AIV overwrites P only after cube#2(*i*) consumed it): | ||
| 161 | + ``` | ||
| 162 | + for bh in 0..nheads: qh=q+bh*seq*d; ...; # same AIC/AIV chain per head, one reused scratch pair | ||
| 163 | + ``` | ||
| 164 | +- **Causal / `sparse_mode`** — add the mask INSIDE the softmax (before rowmax): set masked positions to `-inf` (or a large negative), then the library `SoftMax` / your reduce handles the rest. Not wired in the proven skeleton. | ||
| 165 | +- **Layouts (BSH / SBH / BSND / BNSD)** — differ only in the index math that computes each head's base offset and the row/col strides you pass to `SetGlobalBuffer` / the matmul shapes. The compute pattern is identical; generate the offset arithmetic per layout. | ||
| 166 | +- **Dtypes (fp16 / bf16)** — I/O dtype is the `MatmulType` `T`; softmax ALWAYS accumulates in **fp32** (Cast in → reduce/exp in fp32 → Cast out with `CAST_RINT`), independent of I/O dtype. PV accumulate stays fp32 in the cube. | ||
| 167 | +- **Shapes** — single-pass **non-flash** softmax is correct while a full `[seq]` fp32 row fits UB (fine at moderate seq, since S lives in GM and only one row is UB-resident). For very long seq, switch to **flash online-softmax + KV-block tiling** (running-max / running-sum / expMax rescale) — codify that delta from **P-P101** (source-repo card, not migrated), not here. | ||
| 168 | + | ||
| 169 | +## 5. HOST-side generation guidance (generate this too, with standard CANN APIs) | ||
| 170 | + | ||
| 171 | +Author the host with **standard CANN / torch_npu APIs** — no exotic wiring. Pattern, not a copyable file: | ||
| 172 | + | ||
| 173 | +- **Tiling struct** — the op-specific scalar args the kernel needs (`seq`, `d`, `scale`, plus `nheads` / layout strides / mask flags for the general op). For the `IterateAll<sync=true>` cube path, the `TCubeTiling` is filled from device-side scalar args (`M/N/Ka/Kb`, `singleCoreM/N/K`) — so the host does NOT need the matmul host-tiling lib on its link line (OL-235). If you use the library `SoftMax`, generate its `SoftMaxTiling` (via `GetSoftMaxMinTmpSize`/`GetSoftMaxMaxTmpSize` for the tmp-buffer size). | ||
| 174 | +- **Scratch allocation** — allocate the `[seq,seq]` (or `[nheads,seq,seq]`, or one reused `[seq,seq]`) S / P GM scratch tensors host-side and pass their pointers in (they are NOT op outputs). | ||
| 175 | +- **Launch** — the standard `aclrtlaunch_<kernel>(blockDim, stream, ...ptrs..., scalars)`. `blockDim = 1` (one AIC block; the runtime pairs the AIV(s) for the 1:2 MIX). **The standard `aclrtlaunch` runtime already supplies the FFTS descriptor** (via `rtGetC2cCtrlAddr`) — do **NOT** hand-wire FFTS; that is what makes the user `CrossCore*` handshake work under the default launch. | ||
| 176 | +- **pybind marshaling** — keep the pybind shim PURE: device/dtype/dim `TORCH_CHECK`s, `.contiguous()`, `torch::empty` for output + scratch, compute `scale = 1/√d`, call the `extern "C" aclrtlaunch_<kernel>` shim, return the output. NO torch math in the shim (all compute is in the kernel). | ||
| 177 | +- **Build** — `SOC_VERSION=Ascend910_9382`, the standard AscendC `aclrtlaunch` flow (separate AIC+AIV device objects; pybind links the auto-generated `extern "C" aclrtlaunch_<kernel>`). Compiles on the NPU server; the generated code `#include`s CANN/catlass headers resolved there. | ||
| 178 | + | ||
| 179 | +## 6. Witness step — BUILD + RUN the SYNC-witness BEFORE you author | ||
| 180 | + | ||
| 181 | +Before authoring your kernel, build and run the minimal a3 MIX SYNC-witness at: | ||
| 182 | + | ||
| 183 | +`src/skills/references/target/ascendc/examples/a3_mix_fa_min/`（源仓代码资产路径，未迁移） | ||
| 184 | + | ||
| 185 | +(main is slimming this into a handshake-only witness — **main owns its final form**; reference it by that path.) Build it on an a3 CANN 9.0.0 container and run it once: the point is to watch `torch.npu.synchronize()` **return** (not hang) — that is the device proof the AIC↔AIV handshake closes deadlock-free, which beats any prose assurance. Do NOT author or fork the witness (main's lane) and do NOT construct your own; READ + RUN it, then GENERATE your op. | ||
| 186 | + | ||
| 187 | +## 7. Cross-references | ||
| 188 | + | ||
| 189 | +- **PB-55** — the `MIX_AIC_1_2` AIC↔AIV handshake is direction-ASYMMETRIC (reverse per-subblock-counted; single-setter deadlocks); §3(a) implements it. | ||
| 190 | +- **PB-34** — `MatmulImpl<>` + async KfcServer path deadlocks in this MIX mode → §3(b) mandates `IterateAll<sync=true>`. | ||
| 191 | +- **PB-35** — `event_t(0)` collides with the cube-internal pipe-sync chain → §3(d) forbids flag id 0. | ||
| 192 | +- **DEBT-222** — the a3 FA-class domain-template gap; this file closes its a3 half. | ||
| 193 | +- **PB-28** — `KERNEL_TASK_TYPE_DEFAULT` is arch35-only (`107000` on arch22) → §3(c) omits it. | ||
| 194 | +- **P-P101** — de-scalarized online-softmax; the delta to add for flash KV-block tiling on long seq (§4; source-repo card, not migrated). | ||
| 195 | +- **P-P102 / P-P103** — the a5/arch35 native-MIX counterparts ([cube_vector_fusion](cube_vector_fusion.md) / [fa_class_template](fa_class_template.md)); disjoint arch scope (this a3 template is their sibling, not a substitute). | ||
| 196 | +- **OL-188** — pure-VEC for a cube-required op = forbidden architectural hack (§ library-rule). | ||
| 197 | +- **OL-235** — the harness pybind link line has no host matmul-tiling lib → `IterateAll<sync=true>` is the buildable cube path (§5). | ||
| 198 | + | ||
| 199 | +（PB-* / OL-* / DEBT-* 为源仓 runbook 或跟踪卡号，未迁移。） | ||
| 200 | + | ||
| 201 | +## 8. HONEST SCOPE | ||
| 202 | + | ||
| 203 | +**Device-proven** (`Ascend910_9382`, CANN 9.0.0, 2026-07-18): the MIX **pattern** — dispatch, the 2-flag AIC↔AIV handshake, genuine library-cube matmuls, fp32 softmax — at seq ≤ 384, d = 64, fp16 I/O, single-head + multi-head (device loop, one reused scratch pair). cosine 0.99999905, deterministic. | ||
| 204 | + | ||
| 205 | +**NOT proven (do not claim as covered):** flash online-softmax + KV-block tiling (long seq — use P-P101); causal / attention mask; **perf** (single-pass non-flash, not tuned — but a3 **DOES** have a vendor baseline: `npu_fusion_attention` runs ~102µs, and the source repo's `cross_core_sync.md` §5 / PB-56 measure this template's single-core at 0.0146× and a per-head-independent multi-core extension at 0.186× vendor — see [fa_a3_multicore_result](fa_a3_multicore_result.md). **Corrects the earlier "a3 has no vendor FA baseline" claim — it does; perf is measurable, just uncompetitive until a flash rework**); **no claim on a5** (`unverified_on: soc=Ascend950PR` — a5 uses P-P102 / P-P103). | ||
| 206 | + | ||
| 207 | +Each op generated from this template is verified SEPARATELY on device (owner-scheduled); this template's existence enables generation, it does not certify any specific generated op. | ||
| @@ -0,0 +1,195 @@ | |||
| 1 | +--- | ||
| 2 | +type: Example | ||
| 3 | +title: "Chunked Gated-Delta-Rule / Gated-Linear-Attention Forward — a3 MIX multi-chunk recurrence (P-P117 full body)" | ||
| 4 | +description: "P-P117：a3/arch22 上分块 gated-delta-rule / 门控线性注意力前向的生成范式——每头顺序分块递推、kernel AIV 内做逐块 decay-fold（禁 host 侧代算）、matmul 走 AIC 的 MIX 执行。非 softmax 注意力。" | ||
| 5 | +tags: [task.implementation, topic.attention, topic.programming_model.execution_unit.cube] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/reference/porter/patterns/gated_delta_rule_a3_recurrence.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-08-05T14:23:52Z" | ||
| 11 | +updated_at: "2026-09-04T23:59:25Z" | ||
| 12 | +platforms: [a3] | ||
| 13 | +--- | ||
| 14 | + | ||
| 15 | +# Chunked Gated-Delta-Rule / Gated-Linear-Attention Forward — a3 MIX multi-chunk recurrence (P-P117 full body) | ||
| 16 | + | ||
| 17 | +> 原 OKF v1 卡号：**P-P117**。 | ||
| 18 | +> 适用范围：soc=Ascend910_9382（a3 / arch22，Atlas A3），CANN 9.1.0。每头顺序分块递推、 | ||
| 19 | +> 逐块 decay-fold 在 kernel AIV 内完成（禁 host 侧 torch 代算）、matmul 走 AIC 的 MIX 执行、 | ||
| 20 | +> 整设备 SyncAll（MIX_AIC_1_1，PB-57）均在 Ascend910_9382（dav-2201, arch22）设备实证， | ||
| 21 | +> **在 Ascend950PR 上未验证**。 | ||
| 22 | +> | ||
| 23 | +> Chip: **a3 / arch22 / Ascend910_9382**, CANN 9.1.0. Device-verified 16/16 @ the customer | ||
| 24 | +> fp64 gate (rtol=0.02, ≥17× headroom, deterministic, NaN-free), 9–28× device-time vs | ||
| 25 | +> torch_npu (gen3, 2026-07-20). | ||
| 26 | +> This is the ALGORITHM pattern; it sits on top of the a3 MIX cube+vector SYNC template | ||
| 27 | +> **P-P116** ([fa_class_a3_mix_template](fa_class_a3_mix_template.md)). Read both. | ||
| 28 | + | ||
| 29 | +## 1. Scope | ||
| 30 | + | ||
| 31 | +A chunked **gated-delta-rule** (a.k.a. gated-linear-attention / GDN) forward pass on a3: | ||
| 32 | +arbitrary sequence length `T`, grouped-query attention (GQA), multi-batch. Reference is the | ||
| 33 | +Qwen-team fp64 chunked GDR (`ref_gdr.py::chunk_gated_delta_rule_fwd`). The kernel is genuine | ||
| 34 | +MIX cube+vector — every matmul runs on AIC via the library, elementwise on AIV. This is NOT a | ||
| 35 | +flash/online-softmax attention; there is no softmax. It is a linear-attention recurrence over | ||
| 36 | +chunks of the sequence. | ||
| 37 | + | ||
| 38 | +## 2. The decomposition — every kernel matmul is a plain `A@B` | ||
| 39 | + | ||
| 40 | +Per head, loop chunks `c = 0 .. ceil(T/64)-1` **SEQUENTIALLY** with a persistent state | ||
| 41 | +`S[128,128]` in GM (init 0). The decomposition keeps **every kernel matmul a plain `A@B`** by | ||
| 42 | +pre-folding the per-chunk decay bookkeeping into the matmul operands. **That fold is computed | ||
| 43 | +ON-DEVICE, in the AscendC kernel** (an AIV pre-stage, "Stage 0") — NOT on the host. | ||
| 44 | + | ||
| 45 | +> ⚠️ **Compute the decay ON-DEVICE. Do NOT fold on the host with torch.** The no-delegation rule | ||
| 46 | +> (*ALL computation must use AscendC primitives*) forbids computing the gate cumsum / | ||
| 47 | +> exp / operand fold with host `torch` ops (`gc.cumsum()`, `eg.exp()`, `k*eg`, …). Those are real | ||
| 48 | +> transcendental + elementwise compute and belong in the kernel. The sibling | ||
| 49 | +> `recurrent_gated_delta_rule` op proves feasibility (it does its whole gate decay on-device with | ||
| 50 | +> 0 delegation via the `Exp()` vector primitive). An earlier version of THIS archive folded on the | ||
| 51 | +> host and was caught by `scan_delegation_cheating.py` with 4 violations (torch `cumsum`+`exp`). | ||
| 52 | +> The host may do ONLY layout marshaling: GQA head-expand, zero-pad, chunk-reshape, dtype cast, | ||
| 53 | +> and permute to head-major (incl. providing a raw transposed `kT` — a permute, not compute). | ||
| 54 | + | ||
| 55 | +Per-chunk operand fold (`gc = cumsum(g)`, `eg = exp(gc)`); the kernel produces these in scratch: | ||
| 56 | + | ||
| 57 | +``` | ||
| 58 | +kb = k * beta * eg | ||
| 59 | +knT = (k / eg)^T = (k * exp(-gc))^T | ||
| 60 | +qs = scale * q * eg | ||
| 61 | +vb = v * beta | ||
| 62 | +kdT = (k * exp(gc_last - gc))^T | ||
| 63 | +sc = exp(gc_last) = eg[C-1] # per-head scalar decay applied to state (no extra exp) | ||
| 64 | +``` | ||
| 65 | + | ||
| 66 | +**On-device Stage 0 (AIV), per (head, chunk):** | ||
| 67 | +- `gc = cumsum(g)` over the chunk — an in-chunk **scalar prefix-sum** (`C = 64`; `gc[i]+=gc[i-1]`), | ||
| 68 | + matching the reference `x.cumsum(dim=2)`. (No AscendC `Cumsum` API on this CANN; the 64-wide | ||
| 69 | + scalar scan is trivial vs the matmuls.) | ||
| 70 | +- `eg = Exp(gc)`, `egInv = Exp(-gc)`, `kd = Exp(gc_last - gc)` via the `Exp` vector primitive; | ||
| 71 | + `sc = eg[C-1]` (reuse — `exp(gc_last)` is already the last `eg`). | ||
| 72 | +- **Row-broadcast fold** for `kb`/`qs`/`vb` (`[C,D]`, scale per chunk-position row): `Broadcast<T,2,1>` | ||
| 73 | + the `[C]` row-scale to `[C,D]` then `Mul`. Compute in fp32 UB (`Cast` half→fp32 in, fp32→half | ||
| 74 | + `CAST_RINT` out) so precision matches the host fold. | ||
| 75 | +- **Col-broadcast fold** for the transposed `knT`/`kdT` (`[D,C]`, scale per chunk-position column): | ||
| 76 | + the repeat-`Mul` broadcast idiom — `Mul(dst,src,vec, mask=C, repeat=D, {1,1,1,C/8,C/8,0})` with | ||
| 77 | + `src1RepStride=0` so the `[C]` vec is reused for every row. (This is the sibling | ||
| 78 | + `recurrent_gated_delta_rule::MatVecMul` idiom — studied for structure, not lifted.) | ||
| 79 | +- Fold results land in dedicated scratch slots (`S_KB/S_QS/S_VB/S_KNT/S_KDT`); stages 1–8 read them | ||
| 80 | + exactly as if host-provided. A `SyncAll<false>()` separates Stage 0 (AIV write) from Stage 1 | ||
| 81 | + (AIC read). | ||
| 82 | + | ||
| 83 | +**Zero-pad `T → Nc*64`** and take the per-chunk cumsum: the padded tail then carries `gc_last` | ||
| 84 | +automatically (cumsum of the zero-padded gate), and the padded rows/cols vanish under the causal | ||
| 85 | +mask — so there is NO explicit last-chunk fill to special-case. | ||
| 86 | + | ||
| 87 | +Within-chunk (all matmuls plain `A@B`, all on AIC): | ||
| 88 | + | ||
| 89 | +``` | ||
| 90 | +Acc = kb @ knT | ||
| 91 | +L = -strict(Acc) # strict = strictly-lower triangle; SIGN matters, see CAND-GDR-1 | ||
| 92 | +T = (I + strict(Acc))^-1 # Neumann power product prod(I + L^{2^k}), k=0..5 | ||
| 93 | +Am = incl(qs @ knT) # incl = inclusive-lower (with diagonal) | ||
| 94 | +U = T @ vb | ||
| 95 | +``` | ||
| 96 | + | ||
| 97 | +Cross-chunk (persistent state `S`): | ||
| 98 | + | ||
| 99 | +``` | ||
| 100 | +W = T @ kb | ||
| 101 | +WS = W @ S | ||
| 102 | +o_inter = qs @ S | ||
| 103 | +vn = U - WS | ||
| 104 | +o = Am @ vn + o_inter # chunk output | ||
| 105 | +S = S * sc + kdT @ vn # state update for next chunk | ||
| 106 | +``` | ||
| 107 | + | ||
| 108 | +`Nc == 1` (single chunk, `S = 0`) collapses cleanly to the single-chunk case. GQA is handled by a | ||
| 109 | +host `repeat_interleave` head-expand before the per-head loop. | ||
| 110 | + | ||
| 111 | +## 3. MIX execution + sync (a3-specific — the load-bearing part) | ||
| 112 | + | ||
| 113 | +- **Matmuls on AIC** via `MatmulImpl<...>` + `IterateAll<sync=true>` — the arch22-safe NON-KFC | ||
| 114 | + cube (P-P68). NEVER the async `matmul::Matmul` + KfcServer path (PB-34 / a3-standalone-blocked). | ||
| 115 | +- **One reused `MatmulImpl` object across DIFFERENT shapes** (K=128 KKT / K=64 Neumann chain / | ||
| 116 | + N=128 for U,W,O): call `mm.SetOrgShape(M,N,K)` **per matmul** and `Init` the object at the MAX | ||
| 117 | + shape (M=N=K=128). `SetSingleShape` alone leaves stale org dims → K-mismatch garbage | ||
| 118 | + (**CAND-GDR-3**). | ||
| 119 | +- **Elementwise on AIV.** Reusing UB `LocalTensor`s across consecutive AIV ops needs a | ||
| 120 | + `PipeBarrier<PIPE_ALL>()` at the helper entry — the prior op's V-read races the next op's MTE2 | ||
| 121 | + overwrite otherwise (**CAND-GDR-2**, an INTRA-core WAR; cross-core sync does NOT cover it). | ||
| 122 | +- **Cross-core barrier = whole-device `SyncAll<false>()`.** On a3 this REQUIRES pinning | ||
| 123 | + `KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_1)` (1:1). With the no-macro 1:2 default, idle | ||
| 124 | + AIV cores skip the per-work-item loop and never issue their `SyncAll` → 507014 deadlock | ||
| 125 | + (**PB-57**). This is the OPPOSITE macro choice from the P-P116 scoped-CrossCore MODE2 path | ||
| 126 | + (which keeps the 1:2 default and emits NO macro). Never mix "no macro (1:2)" with whole-device | ||
| 127 | + `SyncAll`. | ||
| 128 | +- **⚠️ DETERMINISM (generated scoped-1:2 MIX, PB-57(b')).** The scoped per-id | ||
| 129 | + `CrossCoreSetFlag`/`WaitFlag` MODE2 path can be deadlock-free yet run-to-run nondeterministic at | ||
| 130 | + high parallelism because the FFTS flag pool aliases above roughly 12 concurrent MIX blocks. | ||
| 131 | + Select the balanced `MIX_AIC_1_1` template for deterministic generated GDR kernels. Keep the | ||
| 132 | + per-id flags and WAR barriers unchanged; pairing balance alone removed the nondeterminism in | ||
| 133 | + 12/12 canonical runs, with all 16 cases deterministic and no whole-device `SyncAll` required. | ||
| 134 | + | ||
| 135 | +## 4. GM slot layout | ||
| 136 | + | ||
| 137 | +`S_NUM = 32` GM scratch slots, including a 2-slot persistent `S` (double-buffered across the | ||
| 138 | +chunk loop) and a 2-slot `Sdelta`. The remaining slots hold the **on-device Stage-0 folded operands** | ||
| 139 | +(`S_KB`/`S_QS`/`S_VB`/`S_KNT`/`S_KDT` — 5 slots; written by Stage 0, read by stages 1–8), the | ||
| 140 | +within-chunk intermediates (`Acc`/`L`/`T`/`Am`/`U`), and the cross-chunk temporaries | ||
| 141 | +(`W`/`WS`/`o_inter`/`vn`). Persistent `S` survives across chunk iterations; everything else is | ||
| 142 | +recycled per chunk. (The earlier host-fold variant had `S_NUM = 27`; moving the fold on-device | ||
| 143 | +added the 5 fold slots.) | ||
| 144 | + | ||
| 145 | +## 5. The 4 codegen bugs found + fixed (cross-refs) | ||
| 146 | + | ||
| 147 | +All four were device-pinned during gen3 (each by a device run or a CPU-vs-device intermediate diff): | ||
| 148 | + | ||
| 149 | +1. **CAND-GDR-1** — Neumann solve SIGN: the kernel's `(I-L)^-1` must equal the reference's | ||
| 150 | + `(I+L)^-1`; negate the host strict mask. (Corrects an earlier WRONG "fp16-overflow → fp32 | ||
| 151 | + intermediates" thesis — CPU fp16-sim passes; it was the sign.) | ||
| 152 | +2. **CAND-GDR-2** — intra-AIV UB-buffer WAR fence (`PipeBarrier<PIPE_ALL>`); cross-core sync does | ||
| 153 | + NOT fix it. (Corrects an earlier WRONG "scoped CrossCore is the race fix" thesis.) | ||
| 154 | +3. **CAND-GDR-3** — reused `MatmulImpl` needs `SetOrgShape(M,N,K)` per call + Init at max shape. | ||
| 155 | +4. **CAND-GDR-4** — precision-bar threshold conflation: score against the customer fp64 gate | ||
| 156 | + (16/16, ≥17× headroom), not `model.py`'s strict-fp32 9.766e-4 bar (14/16, under-rates via a | ||
| 157 | + forced `.to(float32)`). | ||
| 158 | + | ||
| 159 | +**Review-stage methodology fix (delegation)** — the first archive computed the gate cumsum + exp + | ||
| 160 | +operand fold on the HOST with torch (`gc.cumsum()`, `eg.exp()`), caught by | ||
| 161 | +`scan_delegation_cheating.py` (4 violations). Fixed by moving the whole fold ON-DEVICE into the | ||
| 162 | +Stage-0 AIV pre-stage (§2). Re-verified 16/16 @ fp64 (max_err byte-identical to the host-fold | ||
| 163 | +version → numerically faithful) + scanner 0. **Lesson: fold the decay in the kernel from the | ||
| 164 | +start; do not reach for host torch cumsum/exp — that is delegation, not a shortcut.** | ||
| 165 | + | ||
| 166 | +## 6. Capability note (why this is worth codifying) | ||
| 167 | + | ||
| 168 | +The a5ops backend emitted these 4 codegen-correctness bugs on this op that a different backend did | ||
| 169 | +not — same pipeline, different backend ⇒ a codegen-DISCIPLINE short-board, NOT missing knowledge. | ||
| 170 | +Git history shows the relevant sync KB (P-P116 + PB-55, `cross_core_sync.md`) was in place BEFORE | ||
| 171 | +generation and never deleted — NOT a KB regression. The load-bearing clarity gaps that DID | ||
| 172 | +contribute (fa_class §(c) coupling unspelt; `MIX_AIC_1_1` buried) are addressed by the PB-57 entry | ||
| 173 | +and the [fa_class_a3_mix_template](fa_class_a3_mix_template.md) §(c) / `cross_core_sync.md` §6 | ||
| 174 | +amendments shipped with this pattern. | ||
| 175 | + | ||
| 176 | +## 7. Cross-references | ||
| 177 | + | ||
| 178 | +- **P-P116** ([fa_class_a3_mix_template](fa_class_a3_mix_template.md)) — the a3 MIX cube+vector | ||
| 179 | + SYNC template this sits on. | ||
| 180 | +- **P-P68** — single-AIC GEMM with static tiling (the NON-KFC cube route; source-repo card, not migrated). | ||
| 181 | +- **PB-57** — whole-device `SyncAll` MIX needs `KERNEL_TYPE_MIX_AIC_1_1` (1:1; source-repo runbook ID). | ||
| 182 | +- **CAND-GDR-1/2/3/4** — the 4 device-pinned bugs above (source-repo IDs). | ||
| 183 | +- **`fa_class/cross_core_sync.md` §6** — cross-core sync does NOT cover an intra-core UB WAR (source-repo file, not migrated). | ||
| 184 | +- **P-P119** ([a3_mix_small_matmul_cube](a3_mix_small_matmul_cube.md)) — the CUBE-EMIT decision for §2's many-small | ||
| 185 | + matmuls: emit a hand-rolled single-block cube primitive (`Nd2Nz→LoadData3D→Mmad→Fixpipe`), | ||
| 186 | + NOT `MatmulImpl::IterateAll` (per-call tiling scalar setup dominates → aic_scalar-bound; | ||
| 187 | + closed a ~1.3× deficit to cv-ref parity, precision-invariant). | ||
| 188 | +- **P-P120** ([a3_mix_small_matmul_cube](a3_mix_small_matmul_cube.md)) — the Neumann `T=(I+strict)^-1` solve: fold | ||
| 189 | + `R@(I+P)=R+R@P` into ONE L0C pass via accumulate-mode Mmad (`cmatrixInitVal=1` seeds the | ||
| 190 | + additive term, `cmatrixInitVal=0` accumulates the product) — no separate +I vector-add. | ||
| 191 | +- **`recurrent_gated_delta_rule`** (sibling op) — on-device gate decay via the `Exp()` vector | ||
| 192 | + primitive + `MatVecMul` repeat-broadcast `Mul` idiom; the structural reference for the on-device | ||
| 193 | + Stage-0 fold (§2). 0 delegation. | ||
| 194 | +- [gated_delta_rule_bwd_a3_recurrence](gated_delta_rule_bwd_a3_recurrence.md) — the BACKWARD companion (P-P118). | ||
| 195 | +- [ChunkGatedDeltaRule A2/A3 ACLNN 算子知识](../operators/aclnn/transformer/attention/chunk_gated_delta_rule.md) — 官方 ACLNN 算子链路卡；本卡为手写 MIX 生成范式。 | ||
| @@ -0,0 +1,135 @@ | |||
| 1 | +--- | ||
| 2 | +type: Example | ||
| 3 | +title: "Chunked Gated-Delta-Rule / Gated-Linear-Attention BACKWARD — a3 MIX 3-pass reverse-recurrence (P-P118 full body)" | ||
| 4 | +description: "P-P118：a3/arch22 上分块 gated-delta-rule 反向的三趟结构——并行 PASS A、求 dstate 的逆向递推 PASS B、逐块 PASS C，产出 dq/dk/dv/dg/db/dh0 六个梯度；含 T 非 64 倍数的尾块处理。" | ||
| 5 | +tags: [task.implementation, topic.attention, topic.precision] | ||
| 6 | +status: stable | ||
| 7 | +sources: | ||
| 8 | + - resource: https://gitcode.com/cann/cannbot-skills/blob/bb96c9c595c7a8e96745d365291769dc48319c22/plugins-community/ascendc-port-orchestrator/kb/okf/reference/porter/patterns/gated_delta_rule_bwd_a3_recurrence.md | ||
| 9 | + role: primary | ||
| 10 | +created_at: "2026-08-05T14:23:52Z" | ||
| 11 | +updated_at: "2026-09-04T23:59:25Z" | ||
| 12 | +platforms: [a3] | ||
| 13 | +--- | ||
| 14 | + | ||
| 15 | +# Chunked Gated-Delta-Rule / Gated-Linear-Attention BACKWARD — a3 MIX 3-pass reverse-recurrence (P-P118 full body) | ||
| 16 | + | ||
| 17 | +> 原 OKF v1 卡号：**P-P118**。 | ||
| 18 | +> 适用范围：soc=Ascend910_9382（a3 / arch22，Atlas A3），CANN 9.1.0。三趟结构（并行 PASS A、 | ||
| 19 | +> 逆向递推 PASS B 求 dstate、逐块 PASS C）、ON-DEVICE decay-fold + reverse-cumsum(dg) + GQA | ||
| 20 | +> head-sum（kernel AIV，禁 host torch 代算）、matmul-with-ones 归约、整设备 SyncAll | ||
| 21 | +> （MIX_AIC_1_1，PB-57）均在 Ascend910_9382（dav-2201, arch22）设备实证，**在 Ascend950PR 上未验证**。 | ||
| 22 | +> | ||
| 23 | +> Chip: **a3 / arch22 / Ascend910_9382**, CANN 9.1.0. Device-verified **11/11** @ the customer | ||
| 24 | +> fp64 gate (rtol=0.02 per-gradient, ~14-18× headroom, deterministic, NaN-free, input-mutation-safe), | ||
| 25 | +> **2.22–11.25× vs the PyTorch-NPU reference** (2026-07-20). | ||
| 26 | +> This is the BACKWARD companion to **P-P117** ([gated_delta_rule_a3_recurrence](gated_delta_rule_a3_recurrence.md)). | ||
| 27 | +> It sits on top of the a3 MIX cube+vector SYNC template **P-P116** | ||
| 28 | +> ([fa_class_a3_mix_template](fa_class_a3_mix_template.md)). Read all three. | ||
| 29 | + | ||
| 30 | +## 1. Scope | ||
| 31 | + | ||
| 32 | +The BACKWARD of a chunked **gated-delta-rule** (a.k.a. gated-linear-attention / GDN) pass on a3: | ||
| 33 | +arbitrary `T`, GQA, multi-batch. It produces **six gradients** `dq, dk, dv, dg, db, dh0` from the | ||
| 34 | +forward intermediates `(q, k, v, g, beta, A, h, do, dht, initial_state)`. Reference is the Qwen-team | ||
| 35 | +fp64 chunked-GDR backward (`ref_gdr.py::chunk_gated_delta_rule_bwd`). Genuine MIX cube+vector — every | ||
| 36 | +matmul on AIC, elementwise on AIV. Customer gate = FQA `_assert_relative(rtol=0.02)` applied | ||
| 37 | +**per-gradient** (all 6 must pass). | ||
| 38 | + | ||
| 39 | +## 2. Two math simplifications that make it a clean 3-pass matmul assembly | ||
| 40 | + | ||
| 41 | +- **Use the passed-in recurrent state `h` directly.** The reference recomputes `h` identically in | ||
| 42 | + fp32, so there is **NO in-kernel forward recurrence sweep**. Per chunk `vn = u - w@h[c]`. This makes | ||
| 43 | + the whole forward-direction sweep (PASS A) **parallel** per-chunk. | ||
| 44 | +- **Only step 4 (`dstate`) is a cross-chunk recurrence — and it runs in REVERSE chunk order.** Steps | ||
| 45 | + 1, 3, 5, 6 are per-chunk (parallel). | ||
| 46 | + | ||
| 47 | +## 3. The 3 passes (per head `bh = b*Hv + hv`, blockdim = nHead, MIX_AIC_1_1) | ||
| 48 | + | ||
| 49 | +``` | ||
| 50 | +PASS A (forward c): w = A @ kbg | ||
| 51 | + u = A @ vb | ||
| 52 | + vn = u - w @ h[c] | ||
| 53 | + AmT = mUpInc ⊙ (kn @ qsT) # transposed attn (avoids in-kernel Am transpose) | ||
| 54 | + dv = AmT @ do | ||
| 55 | + dsi = qs^T @ do | ||
| 56 | + store w, vn, dv, dsi | ||
| 57 | + | ||
| 58 | +PASS B (REVERSE c): dh[c] = dstate | ||
| 59 | + dv[c] += kd @ dstate | ||
| 60 | + dstate = dstate*sc + dsi - w^T @ dv # fp32 accumulation in GM | ||
| 61 | + dh0 = dstate # after last (c=0) reverse step | ||
| 62 | + | ||
| 63 | +PASS C (forward c): step5 (dq5, dk5, dw5, dg5) + step6 (dk, dv, db, dg) per-chunk matmul assembly | ||
| 64 | + verified line-by-line vs model.py step5/6 — signs, masks, transposes, and the | ||
| 65 | + ORDER dependence: dg5 is taken from dq5 BEFORE the ds@k term is added. | ||
| 66 | +``` | ||
| 67 | + | ||
| 68 | +## 4. ON-DEVICE fold / reductions / cumsum (all compute in AscendC — do NOT host-fold) | ||
| 69 | + | ||
| 70 | +> ⚠️ **Compute the decay, the reductions, and the reverse-cumsum ON-DEVICE. Do NOT fold on the host | ||
| 71 | +> with torch.** The no-delegation rule (*ALL computation must use AscendC primitives*) | ||
| 72 | +> forbids host `torch` `exp`/`cumsum`/`sum` for the operand fold, the `dg/db` reductions, or the | ||
| 73 | +> `reverse-cumsum(dg)`. An earlier host-fold bwd variant scored delegation violations pinpointed in | ||
| 74 | +> `op_host` (the `#204` scanner reads `op_host`); it was corrected to the on-device fold. Same lesson | ||
| 75 | +> as the forward (P-P117) — this is the backward restatement. | ||
| 76 | + | ||
| 77 | +- **Per-chunk operand fold** (`kb/kn/qs/qsT/kd/doeg/kbgT/begK/eglgK/egK` + decay masks `dsMask/dsMask2` | ||
| 78 | + + `sc`): folded in kernel AIV stages via the real `Exp()` vector primitive. `g` arrives **already | ||
| 79 | + per-chunk cumsum'd** (from the `verify_cases.py` padding contract), so the kernel does `exp` only — | ||
| 80 | + no in-kernel forward cumsum for the operands. | ||
| 81 | +- **Reductions `dg`, `db`**: done as **matmul-with-ones** producing `[C,16]` fp32 accumulators (host | ||
| 82 | + takes col0). A uniform matmul+elementwise path — no vector-reduce API risk, and R16=16 fan-out | ||
| 83 | + avoids a degenerate `N=1` matmul. | ||
| 84 | +- **`reverse-cumsum(dg)`**: an in-chunk scalar prefix-sum on UB (`dg[c] = sum_{j>=c} dg[j]`), | ||
| 85 | + `SetValue/GetValue`. | ||
| 86 | +- **GQA head-sum(dq, dk)**: a device compaction stage (reduces Hv-expanded grads back to Hk). | ||
| 87 | +- **fp32 accumulation**: `dstate/dg/db/dh0` accumulate in fp32 GM (5th-lesson discipline). | ||
| 88 | +- Transposes via runtime `SetTensorA/B(bool)` (P-P69); the host pre-transposes host-input operands | ||
| 89 | + (`AT/hT/knT/qsT/...`) — a layout permute, not compute. | ||
| 90 | +- **Host does layout marshaling only**: GQA head-expand (`repeat_interleave`), zero-pad `T→Nc*64`, | ||
| 91 | + chunk-reshape, dtype cast, pre-transpose permutes, then `dg/db` col0 extract + reshape. | ||
| 92 | + | ||
| 93 | +## 5. Partial chunk (T not a multiple of 64) | ||
| 94 | + | ||
| 95 | +No `fill_last_chunk` and no knowledge of the original `T` is needed: the benchmark's only partial | ||
| 96 | +case (T=200) has `dht=0`, which makes the last-chunk `sc`-scaling multiply 0 (irrelevant), and causal | ||
| 97 | +masking + zero-padding vanish the padded contributions. CPU-prototype confirmed fill on/off is | ||
| 98 | +identical. (Caveat: a case combining **nonzero dht AND a partial chunk** would need an explicit | ||
| 99 | +last-chunk fill — none exists in the 11-case set.) | ||
| 100 | + | ||
| 101 | +## 6. MIX execution + sync | ||
| 102 | + | ||
| 103 | +Same substrate as the forward (P-P116/P-P117): matmuls on AIC (`MatmulImpl` + `IterateAll<sync=true>`, | ||
| 104 | +arch22-safe NON-KFC cube per P-P68 — one reused `mm` object across shapes needs `SetOrgShape` per call, | ||
| 105 | +CAND-GDR-3), elementwise + the on-device folds/reductions on AIV (intra-AIV UB reuse needs a | ||
| 106 | +`PipeBarrier<PIPE_ALL>` fence, CAND-GDR-2), whole-device `SyncAll<false>()` (PB-57: needs | ||
| 107 | +`KERNEL_TYPE_MIX_AIC_1_1`). | ||
| 108 | + | ||
| 109 | +## 7. Backward-specific gotchas (device-proven) | ||
| 110 | + | ||
| 111 | +- **CAND-GDR-BWD-1** — the A-matrix backward decay mask `dsMask2`: the reference `decay_mask.swapaxes(-2,-1)` | ||
| 112 | + on a `[B,N,C,C,Hv]` tensor swaps the last `C` with the **HEAD** axis (head relocation), **NOT** the | ||
| 113 | + two `C` axes. The correct per-head 2D matrix is `exp(g_c - g_d)` masked **strict-lower** (`c>d`). A | ||
| 114 | + literal `[C,C]` transpose (`exp(g_d - g_c)`, `d>c`) fails dk/dg/db (~0.12-0.27) while dq/dv/dh0 stay | ||
| 115 | + clean — the failing grads share the dsMask2 chain, which localizes it. | ||
| 116 | +- **CAND-GDR-BWD-2** — `dstate` init must be an independent copy (`.clone()`) of the input `dht`. A | ||
| 117 | + `dht_.to(kFloat).reshape().contiguous()` chain is a no-op **VIEW** on fp32-contiguous `dht`; PASS-B | ||
| 118 | + updates `dstate` **in-place** → corrupts the caller's `dht_` (a real training-loop input-mutation | ||
| 119 | + bug) → dk/dg/db/dv drift on a 2nd call with the same `dht` (dq/dh0 stay clean — they don't depend on | ||
| 120 | + `dstate`). Negative-control proven load-bearing. | ||
| 121 | +- **CAND-GDR-BWD-3** — building the passed-in state `h` in the TEST harness must thread `initial_state` | ||
| 122 | + (the kernel trusts the passed-in `h`; the reference recomputes it WITH init internally). Omitting it | ||
| 123 | + fails only the nonzero-`h0` case — a test-harness bug that masquerades as a kernel precision failure. | ||
| 124 | + | ||
| 125 | +## 8. `applies_to` / provenance | ||
| 126 | + | ||
| 127 | +`applies_to: soc=Ascend910_9382 (a3/arch22); cann=9.1.0; op_class=gated-linear-attention-backward/CUBE_MIX; | ||
| 128 | +verified_on=gated_delta_rule bwd a3 (2026-07-20); unverified_on: Ascend950PR`. | ||
| 129 | +Generate-not-copy: authored from the model.py 8-step backward math (CPU-prototype-first: `proto.py` | ||
| 130 | +fp32+fp16 11/11 before the device port), NOT lifted from any shipped bwd kernel. Cross-ref P-P117 | ||
| 131 | +(forward, [gated_delta_rule_a3_recurrence](gated_delta_rule_a3_recurrence.md)), P-P116 (a3 MIX sync | ||
| 132 | +template, [fa_class_a3_mix_template](fa_class_a3_mix_template.md)), P-P68 (NON-KFC cube), P-P69 | ||
| 133 | +(runtime transpose), PB-57 (1:1 macro), CAND-GDR-2/3, CAND-GDR-BWD-1/2/3（后四者为源仓 runbook 卡号，未迁移）。 | ||
| 134 | + | ||
| 135 | +相关：[ChunkGatedDeltaRule A2/A3 ACLNN 算子知识](../operators/aclnn/transformer/attention/chunk_gated_delta_rule.md)（官方 ACLNN 前向算子链路卡；本卡为手写 MIX 反向生成范式）。 | ||