diff --git a/.gitmodules b/.gitmodules new file mode 100644 index 00000000..1646f273 --- /dev/null +++ b/.gitmodules @@ -0,0 +1,9 @@ +[submodule "skills/triton/latency-optimizer/references/docs_triton_IR/third_party"] + path = skills/triton/latency-optimizer/references/docs_triton_IR/third_party + url = https://github.com/Ascend/AscendNPU-IR.git +[submodule "skills/triton/latency-optimizer/references/docs_triton_IR/third_party_triton-ascend"] + path = skills/triton/latency-optimizer/references/docs_triton_IR/third_party_triton-ascend + url = https://github.com/Ascend/triton-ascend.git +[submodule "skills/triton/latency-optimizer/references/docs_triton_IR/thrid_party_AscendNPU-IR"] + path = skills/triton/latency-optimizer/references/docs_triton_IR/thrid_party_AscendNPU-IR + url = https://github.com/Ascend/AscendNPU-IR.git diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/AGENTS.md b/skills/triton/latency-optimizer/references/docs_triton_IR/AGENTS.md new file mode 100644 index 00000000..ac16b267 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/AGENTS.md @@ -0,0 +1,200 @@ +# AGENTS.md + +## Project Purpose + +Migrate GPU Triton kernels to Ascend 910_95 NPU and iteratively optimize performance. The agent follows a minimal-change, verify-after-each-step workflow. + +## Target Hardware + +| Parameter | Value | +|-----------|-------| +| Architecture | `dav-c310` (Reg-based) | +| AI Core | 1 Cube + 2 Vector | +| UB Capacity | 248 KB (256KB - 8KB reserved) | +| L0C Capacity | 256 KB | +| L1 Capacity | 512 KB | +| UB Alignment | 32B | +| L0C Alignment | 512B | + +## First Action: Read Architecture Docs + +Before any migration task, read these files to understand the 910_95 hardware: + +1. `docs_ascendnpu_ir/00-Architecture/01-npu-hardware-overview.md` +2. `docs_ascendnpu_ir/00-Architecture/02-memory-hierarchy.md` +3. `docs_ascendnpu_ir/00-Architecture/03-pipeline-execution-model.md` +4. `docs_ascendnpu_ir/00-Architecture/04-data-layout.md` + +## Migration Workflow + +### Phase 1: Make It Run (minimal changes only) + +Apply these 5 changes in order. Do NOT add any optimization at this stage. Even if `docs_for_triton_agent/` contains a complete migration example for the target kernel (e.g., Flash Attention, Fused MatMul), do NOT apply all the changes from the example at once — only apply the minimal changes needed to make the kernel compile and produce correct results. + +1. **Device replacement**: `cuda` → `npu`, add `import torch_npu` before any triton import +2. **Remove GPU-only parameters**: delete `num_warps`, `num_stages`, `cache_modifier`, `eviction_policy` +3. **Data type replacement**: fp64 → fp32, uint series → int series +4. **Grid adjustment**: Grid size aligned to physical core count, prefer 1D +5. **Add accuracy verification**: add `if __name__ == "__main__"` block comparing Triton output against PyTorch reference + +After Phase 1, invoke **ascend-triton-data-collector** subagent with the script path and target kernel names to verify correctness, extract IR, and collect baseline performance. + +If accuracy check fails at this stage, the migration itself has a bug — fix it before proceeding. + +### Phase 2: Iterative Optimization (one change at a time) + +After baseline is established, iteratively identify and apply optimizations. Consult `docs_for_triton_agent/` to find optimization strategies matching the current bottleneck. + +**Critical rule: Apply ONLY ONE optimization per iteration.** Even if `docs_for_triton_agent/` contains a complete migration example for the target kernel (e.g., Flash Attention, Fused MatMul), do NOT apply all changes at once. First apply only the minimal changes needed to pass accuracy verification, then add optimizations one at a time. If accuracy fails or performance regresses, it must be clear which change caused it. + +For each iteration: + +``` +1. Analyze msprof data and IR from ascend-triton-data-collector to identify the biggest bottleneck +2. Consult docs_for_triton_agent/ to find ONE optimization that addresses the bottleneck +3. Apply the optimization to the kernel code +4. Invoke ascend-triton-data-collector subagent → verify accuracy + collect IR + profile performance +5. If accuracy fails OR performance regresses: + a. Revert the change + b. Analyze the failure — can the optimization be corrected or adjusted? + c. If yes → apply the corrected version, re-verify + d. If no → move on to the next optimization candidate +6. If improved → proceed to next bottleneck +``` + +Stop iterating when: all key metrics are in healthy range, or no remaining optimization improves performance, or accuracy cannot be maintained after correction attempts. + +### Phase 3: 910_95 Specific (after Phase 2 converges) + +After Phase 2 converges, consult `docs_for_triton_agent/` for 910_95-exclusive optimizations (L0C→UB direct path, FP8, SIMT mode, etc.). Apply the same one-at-a-time rule with analysis on failure. + +## Accuracy Verification Template + +Every migrated script must include an `if __name__ == "__main__"` block: + +```python +if __name__ == "__main__": + dtype = torch.float16 + x = torch.randn((1024, 1024), dtype=dtype, device="npu").requires_grad_() + + # --- Triton kernel --- + y = YourTritonFunction.apply(x, ...) + dy = torch.randn_like(y) + y.backward(dy) + triton_dx, x.grad = x.grad.clone(), None + + # --- PyTorch reference --- + ref = your_pytorch_reference(x, ...) + ref.backward(dy) + torch_dx, x.grad = x.grad.clone(), None + + # --- Compare --- + atol, rtol = 1e-3, 1e-3 + if torch.allclose(y, ref, atol=atol, rtol=rtol): + print("✅ [Fwd]Triton and Torch match") + else: + print("❌ [Fwd]Triton and Torch differ") + if torch.allclose(triton_dx, torch_dx, atol=atol, rtol=rtol): + print("✅ [Bwd]Triton and Torch match") + else: + print("❌ [Bwd]Triton and Torch differ") +``` + +Key rules: +- Reference must be pure PyTorch (no Triton calls) +- Clone gradients immediately after `.backward()`, reset `x.grad` to `None` +- Use `atol=1e-3, rtol=1e-3` for fp16/bf16; `atol=1e-4, rtol=1e-4` for fp32 +- The `✅`/`❌` format is important — data collector reports these + +## Performance Analysis Guide + +After each ascend-triton-data-collector run, **prioritize IR analysis** over msprof metrics. IR reveals the actual computation structure, tiling strategy, and pipeline overlap; msprof ratios alone can be misleading (e.g., high scalar_ratio may be hidden by pipeline overlap and not a real bottleneck). + +### Step 1: IR Analysis (primary) + +Examine the compiler IR output to identify: +- **Scalar-heavy patterns**: excessive `arith` ops (int64 comparisons, type conversions) that are NOT absorbed by pipeline overlap +- **Memory access patterns**: non-continuous loads/stores, redundant copies, misaligned tiling +- **Pipeline structure**: whether Vector and Cube ops are properly overlapped; whether scalar ops sit on the critical path or are hidden behind data transfers +- **Sync/barrier density**: unnecessary barriers between independent operations + +### Step 2: msprof Metrics (supplementary, validate IR findings) + +| Metric | Ideal | Caveat | +|--------|-------|--------| +| `aiv_vec_ratio` | > 80% | Low ratio may indicate Vector pipeline stalled OR that Cube/transfer dominates — cross-check with IR | +| `aiv_mte2_ratio` | < 50% | High ratio confirms memory bottleneck seen in IR | +| `aiv_scalar_ratio` | < 20% | **High ratio does NOT always mean scalar-bound** — if scalar ops are hidden by pipeline overlap (MTE2 wait time covers scalar execution), they are not the bottleneck. Only treat as scalar degradation when IR confirms scalar ops are on the critical path | +| `aic_cube_ratio` | > 80% | Low ratio may indicate Cube pipeline stalled — cross-check with IR tiling | +| `aic_mte1_ratio` | moderate | High ratio confirms L1→L0A/L0B transfer bottleneck seen in IR | + +### Bottleneck → Action + +| Bottleneck | Action | +|------------|--------| +| Compute-bound (high vec/cube ratio) | Increase tiling, CV fusion | +| Memory-bound (high mte2 ratio) | MultiBuffer, care_padding=False, continuous access | +| Scalar degradation (IR-confirmed critical-path scalar ops) | Type conversion (int64→int32, int cmp→fp32 cmp), where→get_element+insert_slice | +| Pipeline-hidden scalar (high scalar_ratio but IR shows overlap) | No action needed — scalar ops are masked by data transfer latency | +| Sync overhead | sync_solver, reduce barrier count | + +## Knowledge Base Hierarchy + +Consult documentation in this priority order: + +### Priority 1: Agent Migration Guide (problem-oriented, code-driven) + +Path: `docs_for_triton_agent/` + +| Scenario | Document | +|----------|----------| +| Hardware specs / memory / alignment | `00-hardware-quick-ref.md` | +| First migration / architecture differences | `01-migration-overview.md` | +| API not working / need alternatives | `02-api-differences.md` | +| BLOCK_SIZE / Grid contraction | `03-tiling-and-grid.md` | +| Memory access optimization | `04-memory-access-patterns.md` | +| tl.dot + vector ops co-optimization | `05-cv-pipeline-optimization.md` | +| Performance far below expected | `06-scalar-degradation-avoidance.md` | +| multibuffer / enable_mixed_cv / sync_solver | `07-compile-params.md` | +| Type conversion / precision | `08-data-type-precision.md` | +| tl.make_block_ptr | `09-block-pointer-migration.md` | +| Autotune on NPU | `10-autotune-on-npu.md` | +| tl.dot + bias fusion | `11-fixpipe-and-bias-fusion.md` | +| Multiple tl.store merge | `12-store-merge.md` | +| Single-position tl.where | `13-where-optimization.md` | +| compile_hint / extension APIs | `14-compile-hint-and-extension.md` | +| Cube-Vector synchronization | `15-sync-and-barrier.md` | +| Compilation / runtime errors | `16-debugging-common-errors.md` | +| Flash Attention migration | `17-flash-attention-migration.md` | +| Fused MatMul migration | `18-fused-matmul-migration.md` | +| Fused SwiGLU migration | `19-fused-swiglu-migration.md` | +| RoPE migration | `20-rope-migration.md` | +| Softcap migration | `21-softcap-migration.md` | +| Advanced optimization | `22-advanced-optimization.md` | + +### Priority 2: Triton-Ascend Developer Docs (API reference, compilation flow) + +Path: `docs_triton_ascend/` + +Consult when Priority 1 docs don't cover the detail needed: API parameters, compilation pipeline internals, extension usage examples. + +### Priority 3: AscendNPU-IR Compiler Docs (IR semantics, pass definitions) + +Path: `docs_ascendnpu_ir/` + +Consult when analyzing extracted IR or understanding compiler behavior: HIVM dialect operations, pass semantics, memory management, data layout. + +### Priority 4: Source Code (last resort) + +- `triton-ascend/` — Triton-Ascend implementation +- `AscendNPU-IR/` — bishengir compiler source + +Only consult when documentation cannot resolve the question. + +## Boundaries + +- Never apply more than one optimization per iteration — if accuracy fails, you must know exactly which change caused it +- Never skip accuracy verification after any change +- Never modify files under `triton-ascend/third_party/ascend/` or `AscendNPU-IR/bishengir/` +- Do not use parameters or constants from 910B/A2 architectures — all values must be for 910_95 (dav-c310) +- If accuracy fails or performance regresses after an optimization, revert first, then analyze the failure — try to correct if possible, otherwise move to the next optimization diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/00-Architecture/01-npu-hardware-overview.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/00-Architecture/01-npu-hardware-overview.md new file mode 100644 index 00000000..927b8cee --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/00-Architecture/01-npu-hardware-overview.md @@ -0,0 +1,391 @@ +# NPU 硬件架构总览 + +> 关键词：AI Core, Cube, Vector, NPUTargetSpec, AddressSpace, TCoreType, VFMode, dav-c220, dav-c310, Reg-based, Mem-based + +## 概述 + +华为昇腾 NPU 是面向 AI 计算的专用处理器，其核心设计理念是通过多级存储层次和专用计算单元实现高吞吐的矩阵与向量计算。AscendNPU-IR 项目通过 HACC 和 HIVM 两个 Dialect 将 NPU 硬件特征映射到 MLIR IR 中，使得编译器能够精确描述和优化针对特定硬件的计算逻辑。 + +每个 AI Core 是 NPU 的基本计算单元，内部包含 Cube（矩阵乘法）和 Vector（向量计算）两类核心。不同 NPU 型号在 AI Core 数量、存储容量、数据通路等方面存在差异，这些差异通过 [NPUTargetSpec.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td) 中的 TableGen 定义精确描述，并在编译时通过 `hacc.DeviceSpec` 属性传递给 IR。 + +本文档从源码中精确提取所有 NPU 型号规格、核心类型枚举和硬件概念在 IR 中的映射关系，为理解 AscendNPU-IR 的后续文档奠定基础。 + +## AI Core 架构 + +每个 AI Core 包含以下组件： + +| 组件 | 功能 | 说明 | +|------|------|------| +| **AIC (AI Cube)** | 包含 Cube 单元，用于矩阵乘法计算 | 每个 AI Core 有 1 个 AIC | +| **AIV (AI Vector)** | 包含 Vector 单元，用于向量计算 | 每个 AI Core 有 2 个 AIV | +| **Scalar** | 标量计算单元，拥有 DCache 和 ICache | 每个 AIC/AIV 都有自己的 Scalar 单元 | + +**核心数量关系**： +- 每个 AI Core = 1 个 Cube Core + 2 个 Vector Core +- 即 `VectorCoreCount = 2 * CubeCoreCount = 2 * AiCoreCount` + +## NPU 型号规格 + +### Ascend910B 系列 + +源文件：[NPUTargetSpec.td:64-98](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td#L64-L98) + +基础规格（`Ascend910B_BaseSpec`）：UB=192KB, L1=512KB, L0A=64KB, L0B=64KB, L0C=128KB, 架构代号=`dav-c220` + +| 型号 | AI Core | Cube Core | Vector Core | UB | L1 | L0A | L0B | L0C | Arch | +|------|---------|-----------|-------------|-----|-----|-----|-----|-----|------| +| Ascend910B1 | 24 | 24 | 48 | 192KB | 512KB | 64KB | 64KB | 128KB | dav-c220 | +| Ascend910B2 | 24 | 24 | 48 | 192KB | 512KB | 64KB | 64KB | 128KB | dav-c220 | +| Ascend910B3 | 20 | 20 | 40 | 192KB | 512KB | 64KB | 64KB | 128KB | dav-c220 | +| Ascend910B4 | 20 | 20 | 40 | 192KB | 512KB | 64KB | 64KB | 128KB | dav-c220 | + +### Ascend910_93 系列 + +源文件：[NPUTargetSpec.td:104-150](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td#L104-L150) + +基础规格（`Ascend910_93_BaseSpec`）：与 910B 系列相同，UB=192KB, L1=512KB, L0A=64KB, L0B=64KB, L0C=128KB, 架构代号=`dav-c220` + +| 型号 | AI Core | Cube Core | Vector Core | UB | L1 | L0A | L0B | L0C | Arch | +|------|---------|-----------|-------------|-----|-----|-----|-----|-----|------| +| Ascend910_9362 | 20 | 20 | 40 | 192KB | 512KB | 64KB | 64KB | 128KB | dav-c220 | +| Ascend910_9372 | 20 | 20 | 40 | 192KB | 512KB | 64KB | 64KB | 128KB | dav-c220 | +| Ascend910_9381 | 24 | 24 | 48 | 192KB | 512KB | 64KB | 64KB | 128KB | dav-c220 | +| Ascend910_9382 | 24 | 24 | 48 | 192KB | 512KB | 64KB | 64KB | 128KB | dav-c220 | +| Ascend910_9391 | 24 | 24 | 48 | 192KB | 512KB | 64KB | 64KB | 128KB | dav-c220 | +| Ascend910_9392 | 24 | 24 | 48 | 192KB | 512KB | 64KB | 64KB | 128KB | dav-c220 | + +### Ascend310B 系列 + +源文件：[NPUTargetSpec.td:156-190](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td#L156-L190) + +基础规格（`Ascend310B_BaseSpec`）：UB=256KB, L1=1024KB, L0A=64KB, L0B=64KB, L0C=128KB, 架构代号=`dav-m300` + +| 型号 | AI Core | Cube Core | Vector Core | UB | L1 | L0A | L0B | L0C | Arch | +|------|---------|-----------|-------------|-----|------|-----|-----|-----|------| +| Ascend310B1 | 1 | 1 | 1 | 256KB | 1024KB | 64KB | 64KB | 128KB | dav-m300 | +| Ascend310B2 | 1 | 1 | 1 | 256KB | 1024KB | 64KB | 64KB | 128KB | dav-m300 | +| Ascend310B3 | 1 | 1 | 1 | 256KB | 1024KB | 64KB | 64KB | 128KB | dav-m300 | +| Ascend310B4 | 1 | 1 | 1 | 256KB | 1024KB | 64KB | 64KB | 128KB | dav-m300 | + +### Ascend910_95 系列 + +源文件：[NPUTargetSpec.td:196-262](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td#L196-L262) + +基础规格（`Ascend950_BaseSpec`）：UB=248KB（预留 8KB 给编译器）, DCache=32KB~120KB, L1=512KB, L0A=64KB, L0B=64KB, L0C=256KB, 架构代号=`dav-c310` + +| 型号 | AI Core | Cube Core | Vector Core | UB | DCache | L1 | L0A | L0B | L0C | Arch | +|------|---------|-----------|-------------|------|--------|-----|-----|-----|-----|------| +| Ascend910_950z | 4 | 4 | 8 | 248KB | 32~120KB | 512KB | 64KB | 64KB | 256KB | dav-c310 | +| Ascend910_9579 | 28 | 28 | 56 | 248KB | 32~120KB | 512KB | 64KB | 64KB | 256KB | dav-c310 | +| Ascend910_957b | 28 | 28 | 56 | 248KB | 32~120KB | 512KB | 64KB | 64KB | 256KB | dav-c310 | +| Ascend910_957d | 28 | 28 | 56 | 248KB | 32~120KB | 512KB | 64KB | 64KB | 256KB | dav-c310 | +| Ascend910_9581 | 32 | 32 | 64 | 248KB | 32~120KB | 512KB | 64KB | 64KB | 256KB | dav-c310 | +| Ascend910_9589 | 32 | 32 | 64 | 248KB | 32~120KB | 512KB | 64KB | 64KB | 256KB | dav-c310 | +| Ascend910_958a | 32 | 32 | 64 | 248KB | 32~120KB | 512KB | 64KB | 64KB | 256KB | dav-c310 | +| Ascend910_958b | 32 | 32 | 64 | 248KB | 32~120KB | 512KB | 64KB | 64KB | 256KB | dav-c310 | +| Ascend910_9599 | 36 | 36 | 72 | 248KB | 32~120KB | 512KB | 64KB | 64KB | 256KB | dav-c310 | + +### Ascend950PR 系列 + +源文件：[NPUTargetSpec.td:264-346](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td#L264-L346) + +基础规格与 Ascend910_95 系列相同（`Ascend950_BaseSpec`）。 + +| 型号 | AI Core | Cube Core | Vector Core | +|------|---------|-----------|-------------| +| Ascend950PR_950z | 4 | 4 | 8 | +| Ascend950PR_9579 | 28 | 28 | 56 | +| Ascend950PR_957a | 28 | 28 | 56 | +| Ascend950PR_957b | 28 | 28 | 56 | +| Ascend950PR_957c | 28 | 28 | 56 | +| Ascend950PR_957d | 28 | 28 | 56 | +| Ascend950PR_9589 | 32 | 32 | 64 | +| Ascend950PR_958a | 32 | 32 | 64 | +| Ascend950PR_958b | 32 | 32 | 64 | +| Ascend950PR_958c | 32 | 32 | 64 | +| Ascend950PR_958d | 32 | 32 | 64 | +| Ascend950PR_9599 | 36 | 36 | 72 | +| Ascend950PR_959a | 36 | 36 | 72 | +| Ascend950PR_959b | 36 | 36 | 72 | + +### Ascend950DT 系列 + +源文件：[NPUTargetSpec.td:348-490](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td#L348-L490) + +基础规格与 Ascend910_95 系列相同（`Ascend950_BaseSpec`）。 + +| 型号 | AI Core | Cube Core | Vector Core | +|------|---------|-----------|-------------| +| Ascend950DT_950x | 8 | 8 | 16 | +| Ascend950DT_950y | 8 | 8 | 16 | +| Ascend950DT_9571 | 28 | 28 | 56 | +| Ascend950DT_9572 | 28 | 28 | 56 | +| Ascend950DT_9573 | 28 | 28 | 56 | +| Ascend950DT_9574 | 28 | 28 | 56 | +| Ascend950DT_9575 | 28 | 28 | 56 | +| Ascend950DT_9576 | 28 | 28 | 56 | +| Ascend950DT_9577 | 28 | 28 | 56 | +| Ascend950DT_9578 | 28 | 28 | 56 | +| Ascend950DT_9581 | 32 | 32 | 64 | +| Ascend950DT_9582 | 32 | 32 | 64 | +| Ascend950DT_9583 | 32 | 32 | 64 | +| Ascend950DT_9584 | 32 | 32 | 64 | +| Ascend950DT_9585 | 32 | 32 | 64 | +| Ascend950DT_9586 | 32 | 32 | 64 | +| Ascend950DT_9587 | 32 | 32 | 64 | +| Ascend950DT_9588 | 32 | 32 | 64 | +| Ascend950DT_9591 | 36 | 36 | 72 | +| Ascend950DT_9592 | 36 | 36 | 72 | +| Ascend950DT_9595 | 36 | 36 | 72 | +| Ascend950DT_9596 | 36 | 36 | 72 | +| Ascend950DT_95A1 | 36 | 36 | 72 | +| Ascend950DT_95A2 | 36 | 36 | 72 | + +## 存储层次概览 + +Ascend NPU 采用多级存储架构，包含通用可寻址存储空间和专用硬件缓冲区两类。 + +### 通用可寻址存储空间 + +| 存储空间 | IR 标识符 | 说明 | 910B/910_93 | 910_95/950PR/950DT | 310B | +|----------|-----------|------|-------------|---------------------|------| +| GM | `gm` | 全局内存 (HBM/L2)，设备外部存储 | - | - | - | +| L1 | `cbuf` | Cube 单元的一级缓存 | 512KB | 512KB | 1024KB | +| L0A | `ca` | 矩阵 A 输入缓存 | 64KB | 64KB | 64KB | +| L0B | `cb` | 矩阵 B 输入缓存 | 64KB | 64KB | 64KB | +| L0C | `cc` | 矩阵乘法结果缓存 | 128KB | 256KB | 128KB | +| UB | `ub` | 统一缓冲区，Vector 单元使用 | 192KB | 248KB | 256KB | + +### 对齐要求 + +源文件：[NPUTargetSpec.td:64-74](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td#L64-L74) 和 [NPUTargetSpec.td:196-208](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td#L196-L208) + +| 存储空间 | 对齐要求（所有架构一致） | 源码值（bits） | +|----------|------------------------|---------------| +| UB | 32B | 256 bits | +| L1 | 32B | 256 bits | +| L0C | 512B | 4096 bits | + +### 专用硬件缓冲区 + +| 缓冲区 | 大小 | 对齐 | 说明 | 访问方式 | +|--------|------|------|------|----------| +| BT Buffer (BiasTable) | 1KB | 64B | 存放矩阵乘法的 Bias 数据 | 通过 `copy_cbuf_to_bt` 从 L1 拷贝 | +| FP Buffer (FixPipe) | 7KB | 128B | FixPipe 流水线的中间缓冲区 | 通过 `hivm.fixpipe` 隐式使用 | + +### Ascend950 架构特有存储 + +| 存储空间 | 大小范围 | 说明 | +|----------|---------|------| +| DCache (SIMT) | 32KB ~ 120KB | SIMT Vector 的数据缓存，大小可配置 | + +## 核心类型枚举 + +源文件：[HIVMAttrs.td:298-317](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L298-L317) + +### TCoreType — 操作级核心类型 + +用于标注 HIVM 操作在哪种核心上执行。 + +| 枚举值 | C++ 符号 | 数值 | 说明 | +|--------|---------|------|------| +| CUBE | `TCoreType::CUBE` | 1 | 操作在 Cube 核心上执行 | +| VECTOR | `TCoreType::VECTOR` | 2 | 操作在 Vector 核心上执行 | +| CUBE_OR_VECTOR | `TCoreType::CUBE_OR_VECTOR` | 3 | 操作可在 Cube 或 Vector 核心上执行 | +| CUBE_AND_VECTOR | `TCoreType::CUBE_AND_VECTOR` | 4 | 操作需要在 Cube 和 Vector 核心上同时执行 | + +IR 语法示例： + +```mlir +#hivm.tcore_type +#hivm.tcore_type +``` + +### TFuncCoreType — 函数级核心类型 + +源文件：[HIVMAttrs.td:250-269](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L250-L269) + +| 枚举值 | C++ 符号 | 数值 | 说明 | +|--------|---------|------|------| +| AIC | `TFuncCoreType::AIC` | 1 | 函数运行在 AI Cube 核心上 | +| AIV | `TFuncCoreType::AIV` | 2 | 函数运行在 AI Vector 核心上 | +| MIX | `TFuncCoreType::MIX` | 3 | 函数混合使用 Cube 和 Vector 核心 | +| AIC_OR_AIV | `TFuncCoreType::AIC_OR_AIV` | 4 | 函数可在 Cube 或 Vector 核心上运行 | + +### TModuleCoreType — 模块级核心类型 + +源文件：[HIVMAttrs.td:271-292](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L271-L292) + +| 枚举值 | C++ 符号 | 数值 | 说明 | +|--------|---------|------|------| +| AIC | `TModuleCoreType::AIC` | 1 | 模块内所有函数均为 AIC 类型 | +| AIV | `TModuleCoreType::AIV` | 2 | 模块内所有函数均为 AIV 类型 | +| MIX | `TModuleCoreType::MIX` | 3 | 模块内函数混合使用 AIC 和 AIV | + +推断规则（源自源码注释）： +- 若模块内所有函数的 `func_core_type` 均为 `AIV`，则模块核心类型为 `AIV` +- 若模块内所有函数的 `func_core_type` 均为 `AIC`，则模块核心类型为 `AIC` +- 否则，模块核心类型为 `MIX` + +## VFMode 枚举 + +源文件：[HIVMAttrs.td:948-960](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L948-L960) + +VFMode 用于描述 Vector Function 的执行模式。 + +| 枚举值 | C++ 符号 | 数值 | 说明 | +|--------|---------|------|------| +| SIMD | `VFMode::SIMD` | 0 | 单指令多数据模式，传统 Vector 执行方式 | +| SIMT | `VFMode::SIMT` | 1 | 单指令多线程模式，类似 GPU 的线程级并行 | +| MIX | `VFMode::MIX` | 2 | 混合模式，同时使用 SIMD 和 SIMT | + +IR 语法示例： + +```mlir +#hivm.vf_mode +#hivm.vf_mode +#hivm.vf_mode +``` + +## 硬件概念在 IR 中的映射表 + +| 硬件概念 | IR 表示 | 源文件位置 | +|----------|---------|-----------| +| NPU 型号规格 | `hacc.DeviceSpec` 属性 | [NPUTargetSpec.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td) | +| 内存空间 | `#hivm.address_space` | [HIVMAttrs.td:171-197](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L171-L197) | +| 操作核心类型 | `#hivm.tcore_type` | [HIVMAttrs.td:298-317](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L298-L317) | +| 函数核心类型 | `#hivm.func_core_type` | [HIVMAttrs.td:250-269](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L250-L269) | +| 模块核心类型 | `#hivm.module_core_type` | [HIVMAttrs.td:271-292](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L271-L292) | +| Vector 执行模式 | `#hivm.vf_mode` | [HIVMAttrs.td:948-960](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L948-L960) | +| 执行流水线 | `#hivm.pipe` | [HIVMAttrs.td:203-244](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L203-L244) | +| 紧耦合缓冲区 | `#hivm.tightly_coupled_buffer` | [HIVMAttrs.td:1010-1017](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L1010-L1017) | +| 数据布局 | `#hivm.data_layout` | [HIVMAttrs.td:84-165](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L84-L165) | + +## 架构规格对比汇总 + +| 特性 | Ascend910B / 910_93 | Ascend310B | Ascend910_95 / 950PR / 950DT | +|------|---------------------|------------|-------------------------------| +| 架构代号 | dav-c220 | dav-m300 | dav-c310 | +| UB 大小 | 192KB | 256KB | 248KB（预留 8KB） | +| L1 大小 | 512KB | 1024KB | 512KB | +| L0A 大小 | 64KB | 64KB | 64KB | +| L0B 大小 | 64KB | 64KB | 64KB | +| L0C 大小 | 128KB | 128KB | 256KB | +| UB 对齐 | 32B | 32B | 32B | +| L1 对齐 | 32B | 32B | 32B | +| L0C 对齐 | 512B | 512B | 512B | +| DCache | 无 | 无 | 32KB ~ 120KB（SIMT 可配置） | +| 紧耦合缓冲区 | 不支持 | 不支持 | 支持（MoveToUb / MoveToL1） | +| L0C -> UB 通路 | 不支持 | 不支持 | 支持 | +| Fixpipe Dual Dst | 不支持 | 不支持 | 支持（ROW_SPLIT / COLUMN_SPLIT） | + +## 架构分类：Reg-based 与 Mem-based + +Ascend NPU 的编译器将硬件架构分为 **Reg-based（寄存器基）** 和 **Mem-based（内存基）** 两大类。这一分类的根本依据是**是否支持 SIMT（Single Instruction Multiple Threads）VF 模式**：Reg-based 架构支持 SIMT，向量操作可基于寄存器进行；Mem-based 架构不支持 SIMT，向量操作全部基于内存（UB 缓冲区）。两种架构在同步机制、算子降级、内存规划等方面也存在系统性差异。 + +源码参考：[Utils.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HACC/Utils/Utils.cpp)、[Utility.h](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/Transforms/GraphSyncSolver/Utility.h) + +### 分类定义 + +```cpp +// Architecture is memory based (A2/A3). +const bool isMemBasedArch; + +// Architecture is register based (A5). +const bool isRegBasedArch; +``` + +| 架构类型 | 代号 | 对应芯片系列 | 架构代号 | +|---------|------|------------|---------| +| **Mem-based**（内存基） | A2/A3 | Ascend910B, Ascend910_93 | dav-c220 | +| **Reg-based**（寄存器基） | A5 | Ascend310B, Ascend950 | dav-m300, dav-c310 | + +### 判断函数 + +编译器通过 `isRegBasedArch` 和 `isMemBasedArch` 两个函数判断目标架构类型： + +```cpp +bool isRegBasedArch(TargetDevice targetDevice) { + return isAscend310B(targetDevice) || isAscend950(targetDevice); +} + +bool isMemBasedArch(TargetDevice targetDevice) { + return isAscend910B(targetDevice) || isAscend910_93(targetDevice); +} +``` + +当 IR 中未指定目标设备时，默认按 910B（Mem-based）处理。 + +### 核心区别：SIMT 支持 + +两种架构最根本的区别在于**是否支持 SIMT VF 模式**： + +| 特性 | Mem-based (A2/A3) | Reg-based (A5) | +|------|-------------------|-----------------| +| SIMT VF 模式 | **不支持** | **支持**（Ascend310B 和 Ascend950 均支持） | +| VFMode 推断 | 不运行 InferVFMode | 运行 InferVFMode，推断 SIMD/SIMT/MIX | +| 数据访问模型 | 全部基于 UB 缓冲区（SIMD 模式） | SIMT 基于寄存器，SIMD 基于 UB | +| SIMT 编译路径 | 无 | SIMT VF 拆分后走 Triton GPU 编译路径 | +| SIMD/SIMT 混合 | 不适用 | 通过 `--enable-simd-simt-mix-compile` 启用 | +| DCache | 无 | 有（950 系列 32-120KB） | + +SIMT VF 模式下，向量操作基于寄存器进行，每个线程独立执行标量操作，需要 `LocalLoadOp`/`LocalStoreOp` 在 UB 和寄存器之间转移数据。编译器仅在 Reg-based 架构上运行 `InferVFModePass` 和 `InsertInferVFModeFuncPass`，Mem-based 架构直接跳过这些 Pass。 + +### 附加区别：同步机制 + +两种架构的核间同步实现方式也不同： + +| 特性 | Mem-based (A2/A3) | Reg-based (A5) | +|------|-------------------|-----------------| +| 同步方式 | 基于内存的 FFTS（Fast Flag Transmit Storage） | 基于寄存器的 SetFlag/WaitFlag 指令 | +| 跨核同步 | 需要设置 FFTS base addr，使用 `SetCrossCoreInstrOp` | 使用 `SetFlagOp`/`WaitFlagOp` 寄存器级指令 | +| 块同步降低 | 使用 `SetCrossCoreInstrOp` | 使用 `IntraBlockSet`/`IntraBlockRegInstrOp` | +| Pipe Barrier | 对所有 Pipe 生成 barrier | 跳过 PIPE_V 的 barrier | + +### 编译器行为差异 + +架构类型在以下编译 Pass 中产生不同的行为： + +| 编译 Pass | Mem-based 行为 | Reg-based 行为 | +|-----------|---------------|----------------| +| **InjectSync** | 对所有 Pipe 设置 barrier | 跳过 PIPE_V 的 barrier | +| **CrossCoreGSS** | 需要获取并设置 FFTS base addr | 不需要 FFTS | +| **MmadL1 同步** | 标准 SetFlag/WaitFlag | 额外注入 PIPE_M → PIPE_MTE1 的 SetFlag/WaitFlag | +| **VReduceOp 标量降级** | 基本归约仅 i64 降级；argmax/argmin 条件更多 | 除 argmax/argmin 内存对齐问题外，基本归约不降级 | +| **VReduceOp Extra Buffer** | 需要额外临时缓冲区 | 不需要额外缓冲区 | +| **Normalize** | 应用 CmpVne 规范化（vcmp NE → vnot(vcmp EQ)） | 不应用 CmpVne 规范化 | +| **Stride 对齐** | 标准 stride 对齐规则 | 对非单位最后维度 stride 有更严格的对齐要求 | +| **内存规划** | SIMT/MIX 模式下不需要动态调整 UB | SIMT/MIX 模式下需要动态调整 UB 空间（考虑 DCache） | +| **入口内核配置** | `configureEntryForMembaseArch` | `configureEntryForRegbaseArch` | + +### 对 Triton 算子优化的影响 + +- **归约操作**：Reg-based 架构（A5）对基本归约（sum/prod/max/min）有更好的向量硬件支持，除 argmax/argmin 的内存对齐问题外不会标量降级。Mem-based 架构（A2/A3）下 i64 归约和整数 argmax/argmin 会标量降级 +- **比较操作规范化**：Mem-based 架构会将 `vcmp(NE)` 规范化为 `vnot(vcmp(EQ))` 以正确处理 NaN；Reg-based 架构不做此规范化 +- **SIMT 模式**：仅 Reg-based 架构（Ascend310B 和 Ascend950）支持 SIMT VF 模式，Mem-based 架构（910B/910_93）不支持 + +> 详细的标量降级差异见 [11-scalar-lowering.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/11-scalar-lowering.md) + +## 常见问题 + +**Q: Ascend910B1 和 Ascend910B2 的硬件规格有什么区别？** +A: 从 [NPUTargetSpec.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td) 的定义来看，两者的 AI Core 数量（24/24/48）和存储规格完全相同。区别可能在于芯片频率、HBM 容量等 TableGen 中未描述的特性。 + +**Q: 为什么 950 系列的 UB 是 248KB 而不是 256KB？** +A: 源码注释明确标注 `UbSize = 2031616; // bits = 248KB, reserve 8KB for compiler`，即预留了 8KB 给编译器内部使用。 + +**Q: CUBE_OR_VECTOR 和 CUBE_AND_VECTOR 有什么区别？** +A: `CUBE_OR_VECTOR` 表示操作可以在 Cube 或 Vector 任一核心上执行（选择其一），`CUBE_AND_VECTOR` 表示操作需要 Cube 和 Vector 核心同时参与执行。 + +**Q: VFMode 的 SIMT 模式在哪些设备上可用？** +A: SIMT VF 模式是 Reg-based 架构（A5 代）的核心特性，Ascend310B 和 Ascend950 均支持。Mem-based 架构（A2/A3 代：Ascend910B/910_93）不支持 SIMT，编译器在这些设备上会跳过 `InferVFModePass` 等相关 Pass。 + +**Q: Reg-based 和 Mem-based 架构对 Triton 算子编写有什么影响？** +A: 最显著的影响在归约操作上：Reg-based 架构（A5 代：Ascend310B/950）对基本归约有更好的向量硬件支持，而 Mem-based 架构（A2/A3 代：Ascend910B/910_93）下 i64 归约和整数 argmax/argmin 会标量降级，性能损失较大。此外，Mem-based 架构会将 `vcmp(NE)` 规范化为 `vnot(vcmp(EQ))` 以正确处理 NaN。详见 [11-scalar-lowering.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/11-scalar-lowering.md)。 + +## 相关文档 + +- 源码参考：[NPUTargetSpec.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td) +- 源码参考：[HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) +- 下一节：[02-memory-hierarchy.md](./02-memory-hierarchy.md) — 内存层次详解 +- 下一节：[03-pipeline-execution-model.md](./03-pipeline-execution-model.md) — Pipeline 执行模型 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/00-Architecture/02-memory-hierarchy.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/00-Architecture/02-memory-hierarchy.md new file mode 100644 index 00000000..22f97708 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/00-Architecture/02-memory-hierarchy.md @@ -0,0 +1,386 @@ +# 内存层次详解 + +> 关键词：AddressSpace, GM, L1, L0A, L0B, L0C, UB, MTE2, MTE1, MTE3, FIX, 数据通路, 随路操作, 紧耦合缓冲区 + +## 概述 + +Ascend NPU 的内存层次是其计算性能的关键基础。与通用处理器不同，NPU 的各级存储空间与特定的计算单元和执行流水线紧密绑定，数据在不同存储层次之间的搬运由专用的 DMA 引擎（MTE 系列 Pipeline）完成，而非通过统一的缓存层次自动管理。 + +理解内存层次对于编写高效的 HIVM IR 至关重要：每个 `memref` 必须通过 `#hivm.address_space<...>` 属性标注其所在的存储空间，每个数据搬运操作必须使用正确的 Pipeline，且不同存储空间之间的数据通路存在严格的硬件约束。 + +本文档从 [NPUTargetSpec.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td) 和 [HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) 中精确提取各级存储的容量、对齐要求和数据通路信息。 + +## 完整内存层次 + +### 通用可寻址存储空间 + +| 层次 | IR 标识符 | 枚举值 | 说明 | 所属计算单元 | +|------|-----------|--------|------|-------------| +| GM | `gm` | 1 | 全局内存 (HBM/L2)，设备外部存储 | 所有单元共享 | +| L1 | `cbuf` | 2 | 一级缓存 | Cube 单元 | +| L0A | `ca` | 3 | 矩阵 A 输入缓存 | Cube A 端 | +| L0B | `cb` | 4 | 矩阵 B 输入缓存 | Cube B 端 | +| L0C | `cc` | 5 | 矩阵乘法结果缓存 | Cube C 端 | +| UB | `ub` | 6 | 统一缓冲区 | Vector 单元 | + +### 专用硬件缓冲区 + +| 缓冲区 | 大小 | 对齐 | 说明 | 访问方式 | +|--------|------|------|------|----------| +| BT Buffer (BiasTable) | 1KB | 64B | 存放矩阵乘法的 Bias 数据 | 通过 `copy_cbuf_to_bt` 从 L1 拷贝 | +| FP Buffer (FixPipe) | 7KB | 128B | FixPipe 流水线的中间缓冲区 | 通过 `hivm.fixpipe` 隐式使用 | + +## 各级存储的容量和对齐要求 + +### Ascend910B / 910_93 系列 + +源文件：[NPUTargetSpec.td:64-74](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td#L64-L74) + +| 存储空间 | 大小 | 对齐要求 | 源码值（bits） | +|----------|------|----------|---------------| +| UB | 192KB | 32B | UbSize=1572864, UbAlignSize=256 | +| L1 | 512KB | 32B | L1Size=4194304, L1AlignSize=256 | +| L0A | 64KB | - | L0aSize=524288 | +| L0B | 64KB | - | L0bSize=524288 | +| L0C | 128KB | 512B | L0cSize=1048576, L0cAlignSize=4096 | + +### Ascend310B 系列 + +源文件：[NPUTargetSpec.td:156-166](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td#L156-L166) + +| 存储空间 | 大小 | 对齐要求 | 源码值（bits） | +|----------|------|----------|---------------| +| UB | 256KB | 32B | UbSize=2097152, UbAlignSize=256 | +| L1 | 1024KB | 32B | L1Size=8388608, L1AlignSize=256 | +| L0A | 64KB | - | L0aSize=524288 | +| L0B | 64KB | - | L0bSize=524288 | +| L0C | 128KB | 512B | L0cSize=1048576, L0cAlignSize=4096 | + +### Ascend910_95 / 950PR / 950DT 系列 + +源文件：[NPUTargetSpec.td:196-208](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td#L196-L208) + +| 存储空间 | 大小 | 对齐要求 | 源码值（bits） | +|----------|------|----------|---------------| +| UB | 248KB（预留 8KB） | 32B | UbSize=2031616, UbAlignSize=256 | +| DCache | 32KB ~ 120KB | - | MinimalDCacheSize=262144, MaximumDCacheSize=983040 | +| L1 | 512KB | 32B | L1Size=4194304, L1AlignSize=256 | +| L0A | 64KB | - | L0aSize=524288 | +| L0B | 64KB | - | L0bSize=524288 | +| L0C | 256KB | 512B | L0cSize=2097152, L0cAlignSize=4096 | + +## 数据通路详解 + +### AddressSpace 枚举与硬件存储的完整映射表 + +源文件：[HIVMAttrs.td:171-197](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L171-L197) + +| 枚举值 | C++ 符号 | 数值 | IR 标识符 | 硬件存储 | 说明 | +|--------|---------|------|-----------|---------|------| +| Zero | `AddressSpace::Zero` | 0 | `zero` | - | 默认/零地址空间 | +| GM | `AddressSpace::GM` | 1 | `gm` | HBM/L2 | 全局内存 | +| L1 | `AddressSpace::L1` | 2 | `cbuf` | L1 Cache | Cube 一级缓存 | +| L0A | `AddressSpace::L0A` | 3 | `ca` | L0A Buffer | 矩阵 A 输入缓存 | +| L0B | `AddressSpace::L0B` | 4 | `cb` | L0B Buffer | 矩阵 B 输入缓存 | +| L0C | `AddressSpace::L0C` | 5 | `cc` | L0C Buffer | 矩阵乘法结果缓存 | +| UB | `AddressSpace::UB` | 6 | `ub` | UB | 统一缓冲区 | + +IR 使用示例： + +```mlir +memref> +memref> +memref<256x256xf16, #hivm.address_space> +``` + +### 源-目标地址空间到 Pipeline 的映射 + +源文件：[HIVMDMAOps.cpp:616-622](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L616-L622) + +| 源地址空间 | 目标地址空间 | Pipeline | IR 操作 | 说明 | +|-----------|-------------|----------|---------|------| +| GM | L1 | PIPE_MTE2 | `hivm.nd2nz` | GM 到 L1，支持 ND->NZ 转换 | +| GM | UB | PIPE_MTE2 | `hivm.load` | GM 到 UB，支持 Padding | +| L1 | GM | PIPE_MTE2 | `copy_cbuf_to_gm` | L1 到 GM，支持 NZ->ND 转换 | +| L1 | L0A | PIPE_MTE1 | 内部指令 | L1 到矩阵 A 缓存 | +| L1 | L0B | PIPE_MTE1 | 内部指令 | L1 到矩阵 B 缓存 | +| L1 | BT Buffer | PIPE_MTE1 | `copy_cbuf_to_bt` | L1 到 Bias Table 缓存 | +| L1 | UB | PIPE_MTE1 | `hivm.l12ub` | L1 到 UB | +| L0A/L0B | L0C | PIPE_M | Cube 计算 | 矩阵乘法 | +| L0C | GM | PIPE_FIX | `hivm.fixpipe` | L0C 到全局内存 | +| L0C | L1 | PIPE_FIX | `hivm.fixpipe` | L0C 到 L1 缓存 | +| L0C | UB | PIPE_FIX | `hivm.fixpipe` | L0C 到 UB（仅 950 系列） | +| UB | UB | PIPE_V | `hivm.copy` | UB 内复制 | +| UB | GM | PIPE_MTE3 | `hivm.store` | UB 到全局内存 | +| UB | L1 | PIPE_MTE3 | `hivm.copy` | UB 到 L1（仅 950 系列） | + +### Cube 数据通路 + +Cube 计算路径涉及从 GM 加载数据到 L1，再从 L1 加载到 L0A/L0B，经 Cube 计算后结果写入 L0C，最后通过 FixPipe 输出。 + +``` +GM ──[MTE2]──▶ L1 ──[MTE1]──▶ L0A/L0B/BT Buffer ──[M]──▶ L0C ──[FIX]──▶ GM/L1/UB +``` + +详细步骤： + +1. **MTE2**: 从 GM 加载矩阵 A、矩阵 B 数据到 L1 +2. **MTE2**: 从 GM 加载 Bias 数据到 L1 +3. **MTE1**: 从 L1 加载矩阵 A 数据到 L0A +4. **MTE1**: 从 L1 加载矩阵 B 数据到 L0B +5. **MTE1**: 从 L1 加载 Bias 数据到 BT Buffer +6. **M**: Cube 执行矩阵乘法，结果写入 L0C +7. **FIX**: L0C 数据通过 FixPipe 输出到 GM/L1/UB + +### Vector 数据通路 + +Vector 计算路径从 GM 加载数据到 UB，在 UB 中完成向量计算后写回 GM。 + +``` +GM ──[MTE2]──▶ UB ──[V]──▶ UB ──[MTE3]──▶ GM +``` + +### 910_95 特殊通路 + +Ascend950 架构引入了两条特殊数据通路： + +**L0C -> UB 直通通路**：Cube 计算结果可直接通过 FixPipe 输出到 UB，无需经过 GM 中转。这使得 Cube-Vector 混合计算路径更高效： + +``` +910_95: GM ──[MTE2]──▶ L1 ──[MTE1]──▶ L0A/L0B ──[M]──▶ L0C ──[FIX]──▶ UB ──[V]──▶ UB ──[MTE3]──▶ GM +非910_95: GM ──[MTE2]──▶ L1 ──[MTE1]──▶ L0A/L0B ──[M]──▶ L0C ──[FIX]──▶ GM ──[MTE2]──▶ UB ──[V]──▶ UB ──[MTE3]──▶ GM +``` + +**UB -> L1 通路**：Vector 处理后的数据可从 UB 搬运到 L1，供后续 Cube 操作使用。 + +## 随路操作汇总表 + +随路操作是指在数据搬运过程中，由硬件 DMA 引擎自动完成的附加操作，无需额外的计算指令。 + +| Pipeline | 数据流向 | 支持的随路操作 | IR 属性/操作 | 备注 | +|----------|---------|---------------|-------------|------| +| **MTE1** | L1 -> L0A/L0B | 矩阵转置 | `a_transpose`/`b_transpose` | 在 `mmadL1` 等操作中设置 | +| **MTE1** | L1 -> L0A/L0B | 布局转换 | zN <-> nZ | 支持格式互转 | +| **MTE2** | GM -> L1 | ND -> NZ 布局转换 | `hivm.nd2nz` | 将 ND 格式转为 NZ 格式 | +| **MTE2** | L1 -> GM | NZ -> ND 布局转换 | `hivm.nz2nd` | 将 NZ 格式转为 ND 格式 | +| **MTE2** | GM -> UB | Padding | `pad_mode`, `pad_value`, `left_padding_num`, `right_padding_num` | 在 `hivm.load` 中设置 | +| **MTE2** | GM -> UB | 隐式转置 | `may_implicit_transpose_with_last_axis` | 在 `hivm.load` 中设置 | +| **MTE3** | UB -> GM | 原子操作 | `atomic_kind` | 支持 add, max, min, and, or, xor, CAS, XCHG | +| **MTE3** | UB -> GM | 隐式转置 | `may_implicit_transpose_with_last_axis` | 在 `hivm.store` 中设置 | +| **FIX** | L0C -> GM/L1/UB | 预量化 | `pre_quant` | FP32->FP16, FP32->BF16, INT32->INT8 | +| **FIX** | L0C -> GM/L1/UB | 预激活 | `pre_relu` | ReLU, Leaky ReLU, P-ReLU | +| **FIX** | L0C -> GM/L1/UB | 布局转换 | `dma_mode` | NZ2ND, NZ2DN, NZ2NZ | +| **FIX** | L0C -> UB | 双目标模式 | `dual_dst_mode` | ROW_SPLIT, COLUMN_SPLIT（仅 950 系列） | +| **FIX** | L0C -> GM/L1/UB | Channel Split | `channel_split` | 仅 950 系列 | + +## 紧耦合缓冲区（TightlyCoupledBuffer） + +源文件：[HIVMAttrs.td:1010-1017](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L1010-L1017) + +紧耦合缓冲区是 **Ascend950 架构特有** 的 CV（Cube-Vector）通信机制，用于在 Cube 操作和 Vector 操作之间高效传递数据，无需经过全局内存中转。 + +**IR 表示**： + +```mlir +#hivm.tightly_coupled_buffer> +``` + +**工作原理**： + +紧耦合缓冲区通过 `InsertCVTightCoupledBuffer` Pass 在 Fixpipe 和 Vector 操作之间插入，支持两种数据搬运模式： + +| 模式 | 数据流向 | 说明 | +|------|---------|------| +| MoveToUb | L0C -> UB | 将 Cube 计算结果从 L0C 直接搬运到 UB | +| MoveToL1 | UB -> L1 | 将 Vector 处理后的数据从 UB 搬运到 L1 | + +**Pipeline 选择逻辑**： + +``` +if (isAscend950(target)) { + if (enableLayoutOptimization) { + InsertCVDataMovement // A5 新布局优化路径 + } else { + InsertCVTightCoupledBuffer // 传统紧耦合缓冲区路径 + } +} else { + InsertLoadStoreForMixCV // 非 950 设备的混合 CV 路径 +} +``` + +## 数据流 ASCII 图 + +### Ascend910B / 910_93 架构 + +``` ++-----------------------------------------------------------------------------+ +| Global Memory (GM / HBM) | ++---------------------------------------+-------------------------------------+ + | + +-------------------+-------------------+ + | | + +-----v-----+ +-----v-----+ + | MTE2 | | MTE2 | + | GM -> L1 | | GM -> UB | + | (双向) | | (单向) | + +-----+-----+ +-----+-----+ + | | + v v + +-----------------------+ +-----------------+ + | L1 | | UB | + | (cbuf, 512KB) | | (ub, 192KB) | + | Cube输入缓存 | | Vector工作区 | + +-----------+-----------+ +--------+--------+ + | | + +-----------+-----------+ | + | | | | + +-----v-----+ +---v---+ +-----v-----+ | + | MTE1 | | MTE1 | | MTE1 | | + | L1 -> L0A | |L1->L0B| |L1 -> BT Buf| | + +-----+-----+ +---+---+ +-----+-----+ | + | | | | + v v v | + +-----------+ +-----------+ +-----------+ | + | L0A | | L0B | | BT Buffer | | + | (ca,64KB) | | (cb,64KB) | | (1KB) | | + | 矩阵A输入 | | 矩阵B输入 | | Bias数据 | | + +-----+-----+ +-----+-----+ +-----+-----+ | + | | | | + +-------------+-------------+ | + | | + v | + +------------------+ | + | Cube | | + | (MatMul) | | + +--------+---------+ | + | | + v | + +------------------+ | + | L0C | | + | (cc, 128KB) | | + | 矩阵乘法结果 | | + +--------+---------+ | + | | + +-----------+-----------+ | + | | | | + +-----v-----+ +---v---+ | + | FIX | | FIX | | + | L0C -> GM | |L0C->L1| | + +-----+-----+ +---+---+ | + | | | + v v | + +-----------+ +-----------+ | + | GM | | L1 | | + +-----------+ +-----------+ | + | + +-------------------------------+ + | + +-----v-----+ + | MTE3 | + | UB -> GM | + | (单向) | + +-----+-----+ + | + v + +-----------+ + | GM | + +-----------+ +``` + +### Ascend910_95 / 950PR / 950DT 架构 + +``` ++-----------------------------------------------------------------------------+ +| Global Memory (GM / HBM) | ++---------------------------------------+-------------------------------------+ + | + +-------------------+-------------------+ + | | + +-----v-----+ +-----v-----+ + | MTE2 | | MTE2 | + | GM -> L1 | | GM -> UB | + | (双向) | | (单向) | + +-----+-----+ +-----+-----+ + | | + v v + +-----------------------+ +----------------------+ + | L1 | | UB | + | (cbuf, 512KB) | | (ub, 248KB,预留8KB) | + | Cube输入缓存 | | Vector工作区 | + +-----------+-----------+ +--------+-------------+ + | | + +-----------+-----------+ | + | | | | + +-----v-----+ +---v---+ +-----v-----+ | + | MTE1 | | MTE1 | | MTE1 | | + | L1 -> L0A | |L1->L0B| |L1 -> BT Buf| | + +-----+-----+ +---+---+ +-----+-----+ | + | | | | + v v v | + +-----------+ +-----------+ +-----------+ | + | L0A | | L0B | | BT Buffer | | + | (ca,64KB) | | (cb,64KB) | | (1KB) | | + | 矩阵A输入 | | 矩阵B输入 | | Bias数据 | | + +-----+-----+ +-----+-----+ +-----+-----+ | + | | | | + +-------------+-------------+ | + | | + v | + +------------------+ | + | Cube | | + | (MatMul) | | + +--------+---------+ | + | | + v | + +------------------+ | + | L0C | | + | (cc, 256KB) | | + | 矩阵乘法结果 | | + +--------+---------+ | + | | + +-----------+-----------+-----------+ | + | | | | | + +-----v-----+ +---v---+ +-----v-----+ | | + | FIX | | FIX | | FIX | | | + | L0C -> GM | |L0C->L1| | L0C -> UB |<----+ | + | | | | | (950特有) | | + +-----+-----+ +---+---+ +-----+-----+ | + | | | | + v v v | + +-----------+ +-----------+ +----------------------+ | + | GM | | L1 | | UB (紧耦合缓冲区) |<---------+ + +-----------+ +-----------+ +----------+-----------+ | + | | + +-----v-----+ | + | MTE3 | | + | UB -> GM | | + +-----+-----+ | + | | + v | + +-----------+ | + | GM | | + +-----------+ | +``` + +## 常见问题 + +**Q: 为什么 L0A/L0B 没有对齐要求字段？** +A: 在 [NPUTargetSpec.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td) 的 `TargetSpec` 基类中，只定义了 `UbAlignSize`、`L1AlignSize` 和 `L0cAlignSize` 三个对齐字段，L0A/L0B 的对齐要求未在 TableGen 中显式描述。 + +**Q: 950 系列的 UB -> L1 通路使用哪个 Pipeline？** +A: 根据源码 [HIVMDMAOps.cpp:616-622](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L616-L622)，`{AddressSpace::UB, AddressSpace::L1}` 映射到 `PIPE::PIPE_MTE3`。 + +**Q: `hivm.copy` 操作支持哪些地址空间组合？** +A: 根据源码 [HIVMDMAOps.cpp:440-448](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L440-L448)，`copy` 支持的组合为：`UB -> UB` 和 `GM -> L1`。对于 950 系列还额外支持 `UB -> L1`。 + +**Q: 紧耦合缓冲区和普通数据通路有什么区别？** +A: 紧耦合缓冲区是 950 架构特有的 CV 通信机制，它允许 Cube 和 Vector 之间直接传递数据而无需经过 GM 中转。普通数据通路中，非 950 设备的 Cube 结果必须先写回 GM，再由 Vector 从 GM 加载到 UB。 + +## 相关文档 + +- 源码参考：[NPUTargetSpec.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Targets/NPUTargetSpec.td) +- 源码参考：[HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) +- 源码参考：[HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) +- 源码参考：[HIVMDMAOps.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp) +- 上一节：[01-npu-hardware-overview.md](./01-npu-hardware-overview.md) — NPU 硬件架构总览 +- 下一节：[03-pipeline-execution-model.md](./03-pipeline-execution-model.md) — Pipeline 执行模型 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/00-Architecture/03-pipeline-execution-model.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/00-Architecture/03-pipeline-execution-model.md new file mode 100644 index 00000000..30e45dcf --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/00-Architecture/03-pipeline-execution-model.md @@ -0,0 +1,347 @@ +# Pipeline 执行模型 + +> 关键词：PIPE, OpPipeTrait, MacroOpPipeTrait, SinglePipeOpTrait, set_flag, wait_flag, pipe_barrier, 同步, 流水线 + +## 概述 + +Ascend NPU 的 AI Core 内部采用多 Pipeline 并行执行架构。每个 Pipeline 对应一个硬件执行单元（如向量计算单元、矩阵计算单元、DMA 引擎等），不同 Pipeline 之间可以并行工作，但同一 Pipeline 内的操作是顺序执行的。 + +HIVM IR 通过 `PIPE` 枚举和一系列 TableGen Trait（`OpPipeTrait`、`MacroOpPipeTrait`、`SinglePipeOpTrait`）将每个操作绑定到特定的 Pipeline。编译器利用这些信息进行同步分析（`InjectSync` Pass），在需要时插入 `set_flag`/`wait_flag`/`pipe_barrier` 同步操作，确保数据依赖关系得到满足。 + +理解 Pipeline 执行模型对于编写正确的 HIVM IR 和调试同步问题至关重要。本文档从 [HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) 和 [HIVMTraits.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td) 中精确提取所有 Pipeline 定义和 Trait 机制。 + +## Pipe 枚举完整列表 + +源文件：[HIVMAttrs.td:203-236](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L203-L236) + +| 枚举值 | C++ 符号 | 数值 | IR 标识符 | 硬件执行单元 | 说明 | +|--------|---------|------|-----------|-------------|------| +| PIPE_S | `PIPE::PIPE_S` | 0 | `PIPE_S` | Scalar 单元 | 标量流水线，执行标量计算和控制流 | +| PIPE_V | `PIPE::PIPE_V` | 1 | `PIPE_V` | Vector 单元 | 向量计算流水线，执行向量运算和 UB 内数据搬运 | +| PIPE_M | `PIPE::PIPE_M` | 2 | `PIPE_M` | Cube 单元 | 矩阵计算流水线，执行矩阵乘法 | +| PIPE_MTE1 | `PIPE::PIPE_MTE1` | 3 | `PIPE_MTE1` | MTE1 DMA 引擎 | L1 到 L0A/L0B/BT Buffer 的单向数据通路 | +| PIPE_MTE2 | `PIPE::PIPE_MTE2` | 4 | `PIPE_MTE2` | MTE2 DMA 引擎 | GM 与 L1/UB 之间的双向数据通路 | +| PIPE_MTE3 | `PIPE::PIPE_MTE3` | 5 | `PIPE_MTE3` | MTE3 DMA 引擎 | UB 到 GM 的单向数据通路 | +| PIPE_ALL | `PIPE::PIPE_ALL` | 6 | `PIPE_ALL` | 所有单元 | 所有流水线的统称，用于同步操作 | +| PIPE_MTE4 | `PIPE::PIPE_MTE4` | 7 | `PIPE_MTE4` | MTE4 DMA 引擎 | 额外的数据传输通路 | +| PIPE_MTE5 | `PIPE::PIPE_MTE5` | 8 | `PIPE_MTE5` | MTE5 DMA 引擎 | 额外的数据传输通路 | +| PIPE_V2 | `PIPE::PIPE_V2` | 9 | `PIPE_V2` | 第二 Vector 单元 | 第二向量流水线 | +| PIPE_FIX | `PIPE::PIPE_FIX` | 10 | `PIPE_FIX` | FixPipe 单元 | FixPipe 数据通路，L0C 到 GM/L1/UB | +| VIRTUAL_PIPE_MTE2_L1A | `PIPE::VIRTUAL_PIPE_MTE2_L1A` | 11 | `VIRTUAL_PIPE_MTE2_L1A` | 虚拟 Pipeline | 虚拟 MTE2 L1A 通路，用于编译器内部区分 | +| VIRTUAL_PIPE_MTE2_L1B | `PIPE::VIRTUAL_PIPE_MTE2_L1B` | 12 | `VIRTUAL_PIPE_MTE2_L1B` | 虚拟 Pipeline | 虚拟 MTE2 L1B 通路，用于编译器内部区分 | +| PIPE_NUM | `PIPE::PIPE_NUM` | 13 | `PIPE_NUM` | - | Pipeline 总数计数，非实际 Pipeline | +| PIPE_UNASSIGNED | `PIPE::PIPE_UNASSIGNED` | 99 | `PIPE_UNASSIGNED` | - | 未分配 Pipeline，用于无 Pipe 属性的操作 | + +### Pipe 分类 + +**物理 Pipeline**（对应实际硬件执行单元）：PIPE_S, PIPE_V, PIPE_M, PIPE_MTE1, PIPE_MTE2, PIPE_MTE3, PIPE_MTE4, PIPE_MTE5, PIPE_V2, PIPE_FIX + +**虚拟 Pipeline**（编译器内部使用，用于更细粒度的同步控制）：VIRTUAL_PIPE_MTE2_L1A, VIRTUAL_PIPE_MTE2_L1B + +**特殊值**：PIPE_ALL（同步用）, PIPE_NUM（计数用）, PIPE_UNASSIGNED（未分配） + +## 每个 Pipe 值与硬件执行单元的对应关系 + +| Pipe | 硬件执行单元 | 数据流 | 典型操作 | +|------|-------------|--------|---------| +| PIPE_S | Scalar | 标量计算 | 循环控制、条件判断 | +| PIPE_V | Vector (AIV) | UB -> UB | `hivm.vadd`, `hivm.vmul`, `hivm.vcast` 等 | +| PIPE_M | Cube (AIC) | L0A/L0B -> L0C | `hivm.mmadL1` 中的矩阵乘法部分 | +| PIPE_MTE1 | MTE1 DMA | L1 -> L0A/L0B/BT | `hivm.mmadL1` 中的数据加载部分 | +| PIPE_MTE2 | MTE2 DMA | GM <-> L1/UB | `hivm.load`, `hivm.nd2nz` | +| PIPE_MTE3 | MTE3 DMA | UB -> GM/L1 | `hivm.store`, `hivm.nz2nd` | +| PIPE_FIX | FixPipe | L0C -> GM/L1/UB | `hivm.fixpipe` | +| PIPE_V2 | 第二 Vector | UB -> UB | 第二向量单元操作 | +| PIPE_MTE4 | MTE4 DMA | 额外传输通路 | 预留 | +| PIPE_MTE5 | MTE5 DMA | 额外传输通路 | 预留 | + +## IR 操作的 Pipe 属性映射表 + +### DMA 操作 + +源文件：[HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) + +| 操作 | Pipe | Trait 声明 | 核心类型 | +|------|------|-----------|---------| +| `hivm.load` | PIPE_MTE2 | `OpPipeTrait<"PIPE::PIPE_MTE2">` | 可推断 | +| `hivm.store` | PIPE_MTE3 | `OpPipeTrait<"PIPE::PIPE_MTE3">` | 可推断 | +| `hivm.copy` | 动态推断 | `getPipe()` 方法 | 可推断 | +| `hivm.fixpipe` | PIPE_FIX | `OpPipeTrait<"PIPE::PIPE_FIX">` | CUBE | +| `hivm.nd2nz` | PIPE_MTE2 | `OpPipeTrait<"PIPE::PIPE_MTE2">` | CUBE | +| `hivm.nz2nd` | PIPE_MTE3 | `OpPipeTrait<"PIPE::PIPE_MTE3">` | CUBE | +| `hivm.l12ub` | PIPE_MTE1 | `OpPipeTrait<"PIPE::PIPE_MTE1">` | CUBE | + +**`hivm.copy` 的 Pipe 推断逻辑**： + +源文件：[HIVMDMAOps.cpp:616-622](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L616-L622) + +| 源地址空间 | 目标地址空间 | 推断的 Pipe | +|-----------|-------------|------------| +| UB | UB | PIPE_V | +| L0C | GM | PIPE_FIX | +| GM | L1 | PIPE_MTE2 | +| UB | L1 | PIPE_MTE3 | + +### 向量操作 + +源文件：[HIVMVectorOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td) + +所有向量操作默认继承自 `HIVM_VectorOp`，其 Pipe 和核心类型通过基类 Trait 定义： + +```tablegen +class HIVM_VectorOp traits = [], + list vecSpecialTraits=[OpPipeTrait<"PIPE::PIPE_V">, VectorCoreTypeTrait]> : + HIVM_StructuredOp +``` + +即所有向量操作默认 Pipe 为 `PIPE_V`，核心类型为 `VECTOR`。 + +| 操作类别 | 代表操作 | Pipe | 核心类型 | +|---------|---------|------|---------| +| 一元运算 | `vexp`, `vabs`, `vln`, `vrelu`, `vrsqrt`, `vsqrt`, `vtanh`, `vsin`, `vcos`, `verf`, `vrec`, `vnot`, `vcast` | PIPE_V | VECTOR | +| 二元运算 | `vadd`, `vsub`, `vmul`, `vdiv`, `vmax`, `vmin`, `vor`, `vand`, `vxor`, `vshl`, `vshr`, `vcmp`, `vpow`, `vmod`, `vmodui` | PIPE_V | VECTOR | +| 三元运算 | `vsel` | PIPE_V | VECTOR | +| 广播 | `vbrc` | PIPE_V | VECTOR | +| 规约 | `vreduce` | PIPE_V | VECTOR | +| 转置 | `vtranspose` | PIPE_V | VECTOR | +| 其他 | `varange`, `vinterleave`, `vdeinterleave`, `vflip`, `vmulextended`, `vpad`, `vconcat`, `vgather`, `vcumprod`, `vcumsum`, `vsort` | PIPE_V | VECTOR | + +### 宏操作 + +源文件：[HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td) + +宏操作使用 `MacroOpPipeTrait` 标注其输入和输出 Pipeline。 + +| 操作 | InPipe | OutPipe | Trait 声明 | 核心类型 | +|------|--------|---------|-----------|---------| +| `hivm.mmadL1` | PIPE_MTE1 | PIPE_M | `MacroOpPipeTrait<"PIPE::PIPE_MTE1, PIPE::PIPE_M">` | CUBE | +| `hivm.batchMmadL1` | PIPE_MTE1 | PIPE_M | `MacroOpPipeTrait<"PIPE::PIPE_MTE1, PIPE::PIPE_M">` | CUBE | +| `hivm.matmul` | PIPE_MTE2 | PIPE_MTE3 | `MacroOpPipeTrait<"PIPE::PIPE_MTE2, PIPE::PIPE_MTE3">` | 可推断 | +| `hivm.mix_matmul` | PIPE_MTE2 | PIPE_MTE3 | `MacroOpPipeTrait<"PIPE::PIPE_MTE2, PIPE::PIPE_MTE3">` | 可推断 | +| `hivm.mix_group_matmul` | PIPE_MTE2 | PIPE_MTE3 | `MacroOpPipeTrait<"PIPE::PIPE_MTE2, PIPE::PIPE_MTE3">` | 可推断 | + +## OpPipeTrait / MacroOpPipeTrait / SinglePipeOpTrait 的 TableGen 定义 + +源文件：[HIVMTraits.td:157-173](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td#L157-L173) + +### SinglePipeOpTrait + +标记操作拥有单一 Pipeline 属性。 + +```tablegen +def SinglePipeOpTrait : NativeOpTrait<"SinglePipeOpTrait">; +``` + +### OpPipeTrait + +参数化 Trait，将操作绑定到特定的单一 Pipeline。继承自 `SinglePipeOpTrait`。 + +```tablegen +class OpPipeTrait + : ParamNativeOpTrait<"OpPipeTrait", Pipe, [SinglePipeOpTrait]>; +``` + +使用示例： + +```tablegen +OpPipeTrait<"PIPE::PIPE_MTE2"> // 操作属于 MTE2 Pipeline +OpPipeTrait<"PIPE::PIPE_FIX"> // 操作属于 FIX Pipeline +``` + +### MacroOpTrait + +标记操作为宏操作（包含多个子操作，涉及多个 Pipeline）。 + +```tablegen +def MacroOpTrait : NativeOpTrait<"MacroOpTrait">; +``` + +### MacroOpPipeTrait + +参数化 Trait，标注宏操作的输入 Pipeline 和输出 Pipeline。继承自 `MacroOpTrait`。 + +```tablegen +class MacroOpPipeTrait + : ParamNativeOpTrait<"MacroOpPipeTrait", InOutPipes, + [MacroOpTrait]>; +``` + +使用示例： + +```tablegen +MacroOpPipeTrait<"PIPE::PIPE_MTE1, PIPE::PIPE_M"> // mmadL1: 输入=MTE1, 输出=M +MacroOpPipeTrait<"PIPE::PIPE_MTE2, PIPE::PIPE_MTE3"> // matmul: 输入=MTE2, 输出=MTE3 +``` + +### OpPipeInterface + +源文件：[OpPipeInterface.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/Interfaces/OpPipeInterface.td) + +`OpPipeInterface` 是统一的操作接口，提供以下方法： + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `isSinglePipeOp()` | `bool` | 判断是否为单 Pipe 操作 | +| `isMacroOp()` | `bool` | 判断是否为宏操作 | +| `getPipe()` | `mlir::hivm::PIPE` | 获取单 Pipe 操作的 Pipeline（非单 Pipe 返回 `PIPE_UNASSIGNED`） | +| `getInPipe()` | `mlir::hivm::PIPE` | 获取宏操作的输入 Pipeline（非宏操作返回 `PIPE_UNASSIGNED`） | +| `getOutPipe()` | `mlir::hivm::PIPE` | 获取宏操作的输出 Pipeline（非宏操作返回 `PIPE_UNASSIGNED`） | + +## Pipeline 同步在 IR 中的表示 + +源文件：[HIVMSynchronizationOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMSynchronizationOps.td) + +### set_flag + +通知硬件某个 Pipeline 的操作已完成，设置事件标志。 + +```mlir +hivm.hir.set_flag [#hivm.pipe, #hivm.pipe, EVENT_ID0] +``` + +参数说明： +- `set_pipe`：发出通知的 Pipeline（数据生产方） +- `wait_pipe`：等待通知的 Pipeline（数据消费方） +- `static_event_id` / `dynamic_event_id`：事件 ID（0-7） + +### wait_flag + +等待某个 Pipeline 的事件标志，确保数据依赖满足。 + +```mlir +hivm.hir.wait_flag [#hivm.pipe, #hivm.pipe, EVENT_ID0] +``` + +参数与 `set_flag` 一致。 + +### pipe_barrier + +同一 Pipeline 内的屏障操作，确保该 Pipeline 之前的所有操作完成。 + +```mlir +hivm.hir.pipe_barrier [#hivm.pipe] +``` + +参数： +- `pipe`：需要执行屏障的 Pipeline + +### 同步操作参数的 Pipe 含义 + +| 同步操作 | `set_pipe` 含义 | `wait_pipe` 含义 | +|---------|----------------|-----------------| +| `set_flag` | 完成数据生产的 Pipeline | 需要被通知的 Pipeline | +| `wait_flag` | 产生数据的 Pipeline | 等待数据的 Pipeline | + +典型同步模式： + +``` +MTE2 加载数据 -> set_flag[MTE2, V, event0] -> V 使用数据前 -> wait_flag[MTE2, V, event0] +``` + +## 紧耦合缓冲区的 Pipeline 选择逻辑 + +紧耦合缓冲区是 950 架构特有的 CV 通信机制，其 Pipeline 选择取决于编译选项： + +``` +if (isAscend950(target)) { + if (enableLayoutOptimization) { + InsertCVDataMovement // A5 新布局优化路径 + } else { + InsertCVTightCoupledBuffer // 传统紧耦合缓冲区路径 + } +} else { + InsertLoadStoreForMixCV // 非 950 设备的混合 CV 路径 +} +``` + +在不同路径下，Cube 和 Vector 之间的数据传递使用不同的 Pipeline 组合： + +| 路径 | 数据流 | 使用的 Pipeline | +|------|--------|----------------| +| InsertCVTightCoupledBuffer | L0C -> UB | PIPE_FIX (fixpipe) + PIPE_V (vector) | +| InsertCVDataMovement | L0C -> UB | PIPE_FIX (fixpipe) + PIPE_V (vector) | +| InsertLoadStoreForMixCV | L0C -> GM -> UB | PIPE_FIX (fixpipe) + PIPE_MTE2 (load) + PIPE_V (vector) | + +## 执行顺序约束与同步要求 + +### Pipeline 间并行 + +不同 Pipeline 的操作可以并行执行。例如，MTE2 加载数据到 L1 的同时，V 可以在 UB 上执行向量计算。 + +### Pipeline 内顺序 + +同一 Pipeline 内的操作严格按程序顺序执行，无需额外同步。 + +### 跨 Pipeline 数据依赖 + +当数据生产者和消费者位于不同 Pipeline 时，必须插入同步操作： + +| 场景 | 生产 Pipeline | 消费 Pipeline | 同步方式 | +|------|-------------|-------------|---------| +| GM 加载到 UB 后 V 计算 | PIPE_MTE2 | PIPE_V | set_flag/wait_flag | +| V 计算后写回 GM | PIPE_V | PIPE_MTE3 | set_flag/wait_flag | +| L1 加载到 L0A 后 M 计算 | PIPE_MTE1 | PIPE_M | set_flag/wait_flag | +| M 计算后 FIX 输出 | PIPE_M | PIPE_FIX | set_flag/wait_flag | +| FIX 输出到 UB 后 V 计算 | PIPE_FIX | PIPE_V | set_flag/wait_flag | + +### Event ID 资源 + +源文件：[HIVMAttrs.td:479-498](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L479-L498) + +硬件提供 8 个 Event ID（EVENT_ID0 ~ EVENT_ID7），用于区分不同的同步事件。编译器的 `SyncEventIdAllocation` Pass 负责分配和复用 Event ID。 + +### Unit Flag 模式 + +源文件：[HIVMAttrs.td:512-534](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L512-L534) + +Unit Flag 是一种条件同步机制，用于循环中依赖首次迭代才能确定同步的场景。 + +| 模式 | C++ 符号 | 数值 | 说明 | +|------|---------|------|------| +| DISABLED | `UNIT_FLAG_DISABLED` | 0 | 禁用 Unit Flag | +| RESERVED | `UNIT_FLAG_RESERVED` | 1 | 保留 | +| ENABLED_WITHOUT_UPDATE | `ENABLED_WITHOUT_UPDATE` | 2 | 启用但不更新 | +| ENABLED_WITH_UPDATE | `ENABLED_WITH_UPDATE` | 3 | 启用并更新 | + +### SyncBlock 模式 + +源文件：[HIVMAttrs.td:540-563](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L540-L563) + +SyncBlock 用于不同 AI Core 之间的同步。 + +| 模式 | 说明 | +|------|------| +| ALL_CUBE | 所有 Cube 核心同步到同一点 | +| ALL_VECTOR | 所有 Vector 核心同步到同一点 | +| ALL_SUB_VECTOR | 所有子 Vector 核心同步到同一点 | +| BARRIER_CUBE | Cube-Cube 同步，降低为 barrier.pipe_all | +| BARRIER_VECTOR | Vector-Vector 同步，降低为 barrier.pipe_all | +| ALL | 所有 AIC/AIV 同步到同一点 | + +## 常见问题 + +**Q: 为什么 `hivm.copy` 没有静态的 OpPipeTrait？** +A: `hivm.copy` 支持多种地址空间组合（UB->UB, UB->L1 等），其 Pipe 需要根据源和目标地址空间动态推断，因此通过 `getPipe()` 方法实现而非静态 Trait。 + +**Q: VIRTUAL_PIPE_MTE2_L1A 和 VIRTUAL_PIPE_MTE2_L1B 的用途是什么？** +A: 这两个虚拟 Pipeline 用于编译器内部更细粒度的同步控制。当 MTE2 同时向 L1 的不同区域（A 矩阵区域和 B 矩阵区域）搬运数据时，虚拟 Pipeline 允许编译器区分这两条数据流，实现更精确的同步。 + +**Q: 宏操作（如 mmadL1）的 InPipe 和 OutPipe 如何影响同步？** +A: 宏操作内部包含多个子操作（如 mmadL1 包含 MTE1 数据加载和 M 矩阵计算）。InPipe 表示宏操作的第一个子操作所属的 Pipeline，OutPipe 表示最后一个子操作所属的 Pipeline。同步分析使用这些信息确定宏操作与其他操作之间的依赖关系。 + +**Q: PIPE_MTE4 和 PIPE_MTE5 目前有操作使用吗？** +A: 从当前源码来看，PIPE_MTE4 和 PIPE_MTE5 已在枚举中定义，但尚未有 IR 操作显式使用。它们为未来硬件扩展预留。 + +## 相关文档 + +- 源码参考：[HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) +- 源码参考：[HIVMTraits.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td) +- 源码参考：[HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) +- 源码参考：[HIVMVectorOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td) +- 源码参考：[HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td) +- 源码参考：[HIVMSynchronizationOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMSynchronizationOps.td) +- 源码参考：[OpPipeInterface.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/Interfaces/OpPipeInterface.td) +- 上一节：[02-memory-hierarchy.md](./02-memory-hierarchy.md) — 内存层次详解 +- 下一节：[04-data-layout.md](./04-data-layout.md) — 数据布局详解 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/00-Architecture/04-data-layout.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/00-Architecture/04-data-layout.md new file mode 100644 index 00000000..5070290f --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/00-Architecture/04-data-layout.md @@ -0,0 +1,323 @@ +# 数据布局详解 + +> 关键词：DataLayout, ND, NZ, zN, nZ, Fractal, DOTA_ND, DOTB_ND, DOTC_ND, nd2nz, nz2nd, DataLayoutAttr, fractalSizes + +## 概述 + +Ascend NPU 的 Cube 单元（矩阵乘法引擎）对输入数据的布局有严格要求：矩阵 A 和矩阵 B 必须以特定的 Fractal（分形）格式存储在 L0A/L0B 中，计算结果 L0C 也以 Fractal 格式存储。而用户数据通常以 ND（N-Dimensional，行优先）格式存储在全局内存中。因此，数据在 GM 和 L1 之间搬运时需要进行布局转换（ND <-> NZ），这是 NPU 编程中一个核心且独特的概念。 + +HIVM IR 通过 `DataLayout` 枚举和 `DataLayoutAttr` 参数化属性精确描述数据布局信息，通过 `nd2nz`/`nz2nd` 操作显式表示布局转换。理解数据布局对于正确编写和优化 HIVM IR 至关重要，因为错误的布局会导致计算结果错误或硬件异常。 + +本文档从 [HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) 和 [HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) 中精确提取所有布局定义和转换操作。 + +## DataLayout 枚举完整列表 + +源文件：[HIVMAttrs.td:84-101](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L84-L101) + +| 枚举值 | C++ 符号 | 数值 | IR 标识符 | 物理含义 | +|--------|---------|------|-----------|---------| +| DOTA_ND | `DataLayout::DOTA_ND` | 1 | `dotA_ND` | 矩阵 A 的 ND 布局，用于 Cube 乘法的左操作数 | +| DOTB_ND | `DataLayout::DOTB_ND` | 2 | `dotB_ND` | 矩阵 B 的 ND 布局，用于 Cube 乘法的右操作数 | +| DOTC_ND | `DataLayout::DOTC_ND` | 3 | `dotC_ND` | 矩阵 C 的 ND 布局，用于 Cube 乘法的结果 | +| nZ | `DataLayout::nZ` | 4 | `nZ` | nZ 布局（列优先 Fractal），N 维在低地址连续 | +| zN | `DataLayout::zN` | 5 | `zN` | zN 布局（行优先 Fractal），Z 维在低地址连续 | +| ND | `DataLayout::ND` | 6 | `ND` | ND 布局（标准 N 维行优先），用户数据的标准格式 | +| Fractal | `DataLayout::Fractal` | 7 | `Fractal` | 通用 Fractal 布局，通过 fractalSizes 参数描述具体分形 | + +### 布局物理含义详解 + +#### ND 布局 + +ND（N-Dimensional）是标准的行优先多维数组布局，也是用户数据在全局内存中的自然存储格式。在 ND 布局中，最后一个维度（最右维度）在内存中连续存储。 + +``` +ND 布局示例（2x4 矩阵）: +[a00, a01, a02, a03, a10, a11, a12, a13] +``` + +#### zN 布局 + +zN（Z-major N-minor）是 Cube 单元使用的 Fractal 布局之一。在 zN 布局中，数据按 Fractal 块组织，每个块内 Z 维度（行方向）在低地址连续。zN 布局是矩阵 A 输入到 L0A 时所需的格式。 + +``` +zN 布局示意（Fractal 块大小 16x16）: +按 16x16 分块，块内行优先 +[block(0,0), block(1,0), block(0,1), block(1,1), ...] +每个 block 内: [row0, row1, ..., row15] +``` + +#### nZ 布局 + +nZ（N-major Z-minor）是另一种 Fractal 布局。在 nZ 布局中，数据按 Fractal 块组织，每个块内 N 维度（列方向）在低地址连续。nZ 布局是矩阵 B 输入到 L0B 时所需的格式。 + +``` +nZ 布局示意（Fractal 块大小 16x16）: +按 16x16 分块，块内列优先 +[block(0,0), block(0,1), block(1,0), block(1,1), ...] +每个 block 内: [col0, col1, ..., col15] +``` + +#### DOTA_ND / DOTB_ND / DOTC_ND + +这三种布局是 Cube 矩阵乘法中各操作数的 ND 格式标记，它们在语义上区分了矩阵乘法中 A、B、C 三个操作数的角色。`transpose` 属性仅对 DOTA_ND 和 DOTB_ND 有效且必须提供。 + +| 布局 | 角色 | transpose 属性 | 说明 | +|------|------|---------------|------| +| DOTA_ND | 矩阵 A | 必须提供 | 标记矩阵 A 是否需要转置 | +| DOTB_ND | 矩阵 B | 必须提供 | 标记矩阵 B 是否需要转置 | +| DOTC_ND | 矩阵 C | 不适用 | 矩阵乘法结果的 ND 格式 | + +#### Fractal 布局 + +通用 Fractal 布局，通过 `fractalSizes` 参数描述具体的分形块大小。`fractalSizes` 是一个包含两个 int64 值的数组，分别表示 Fractal 块的行方向和列方向的大小。 + +## DataLayoutAttr 参数化属性定义 + +源文件：[HIVMAttrs.td:103-165](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L103-L165) + +### TableGen 定义 + +```tablegen +def HIVM_DataLayoutAttr : HIVM_Attr<"DataLayout", "data_layout"> { + let parameters = (ins + EnumParameter:$data_layout, + OptionalParameter<"BoolAttr">:$transpose, + OptionalParameter<"DenseI64ArrayAttr">:$fractalSizes + ); + let description = [{ + HIVM data layout mapping attribute. Maps to DOTA_ND, DOTB_ND, DOTC_ND, zN, nZ and ND. + - `transpose`: Indicates that the layout is transposed. + Only valid and must be present for DOTA_ND and DOTB_ND layout. + }]; + let assemblyFormat = "`<` $data_layout (`,` struct($transpose, $fractalSizes)^)? `>`"; +} +``` + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `data_layout` | `DataLayout` 枚举 | 是 | 布局类型 | +| `transpose` | `BoolAttr` | 否（DOTA_ND/DOTB_ND 时必须） | 是否转置 | +| `fractalSizes` | `DenseI64ArrayAttr` | 否 | Fractal 块大小 [row_size, col_size] | + +### 额外方法 + +源码中定义了以下辅助方法： + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getFractalSizesArray()` | `std::optional>` | 获取 fractalSizes 为 SmallVector | +| `getTransposeValue()` | `std::optional` | 获取 transpose 值 | +| `isNDLayout()` | `bool` | 判断是否为 ND 类布局（DOTA_ND/DOTB_ND/DOTC_ND/ND） | +| `getFractalBlockSizes()` | `FailureOr` | 提取恰好 2 个 fractal 块大小 | + +### IR 语法示例 + +```mlir +#hivm.data_layout +#hivm.data_layout +#hivm.data_layout +#hivm.data_layout +#hivm.data_layout +#hivm.data_layout +#hivm.data_layout +#hivm.data_layout +``` + +## 各布局的使用场景 + +| 布局 | 使用场景 | 存储位置 | 说明 | +|------|---------|---------|------| +| ND | 用户数据在 GM 中的标准格式 | GM | 所有用户数据默认格式 | +| zN | Cube 矩阵 A 输入 | L1, L0A | 矩阵 A 在 L1 中以 zN 格式存储 | +| nZ | Cube 矩阵 B 输入 | L1, L0B | 矩阵 B 在 L1 中以 nZ 格式存储 | +| DOTA_ND | 标记矩阵 A 的 ND 格式及转置信息 | - | 用于 `mmadL1` 的布局接口 | +| DOTB_ND | 标记矩阵 B 的 ND 格式及转置信息 | - | 用于 `mmadL1` 的布局接口 | +| DOTC_ND | 标记矩阵 C 的 ND 格式 | - | 用于 `mmadL1` 的布局接口 | +| Fractal | 通用分形布局描述 | - | 通过 fractalSizes 参数化描述 | + +### 操作与布局的对应关系 + +| 操作 | 输入布局 | 输出布局 | 说明 | +|------|---------|---------|------| +| `hivm.nd2nz` | ND | zN/nZ | GM -> L1 时将 ND 转为 Fractal 格式 | +| `hivm.nz2nd` | zN/nZ | ND | L1 -> GM 时将 Fractal 转为 ND 格式 | +| `hivm.mmadL1` | zN (A), nZ (B) | Fractal (C) | Cube 矩阵乘法要求特定输入布局 | +| `hivm.fixpipe` | Fractal (L0C) | ND/NZ/DN | 输出时可进行布局转换 | + +## ND <-> NZ 转换在 IR 中的表示 + +### nd2nz 操作 + +源文件：[HIVMDMAOps.td:328-368](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td#L328-L368) + +将数据从 ND 格式转换为 NZ 格式，同时从 GM 搬运到 L1。属于 PIPE_MTE2 Pipeline。 + +```mlir +hivm.nd2nz ins(%src : memref<256x256xf16, #hivm.address_space>) + outs(%dst : memref>) +``` + +属性： +- `dst_continuous`：可选 UnitAttr，标记目标数据是否连续存储 +- `init_out_buffer`：是否初始化输出缓冲区 +- `pad_value`：填充值 + +### nz2nd 操作 + +源文件：[HIVMDMAOps.td:370-386](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td#L370-L386) + +将数据从 NZ 格式转换为 ND 格式，同时从 L1 搬运到 GM。属于 PIPE_MTE3 Pipeline。 + +```mlir +hivm.nz2nd ins(%src : memref>) + outs(%dst : memref<256x256xf16, #hivm.address_space>) +``` + +### l12ub 操作 + +源文件：[HIVMDMAOps.td:388-404](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td#L388-L404) + +将数据从 L1 搬运到 UB，同时进行 NZ 到 ND 的布局转换。属于 PIPE_MTE1 Pipeline。 + +```mlir +hivm.l12ub ins(%src : memref>) + outs(%dst : memref>) +``` + +## Fractal 布局详解 + +### Cube 矩阵乘法的输入输出布局要求 + +Cube 单元执行矩阵乘法 `C = A * B` 时，对数据的布局有严格要求： + +| 操作数 | 存储位置 | 要求布局 | 说明 | +|--------|---------|---------|------| +| 矩阵 A | L0A | zN | 行优先 Fractal 格式 | +| 矩阵 B | L0B | nZ | 列优先 Fractal 格式 | +| 矩阵 C | L0C | Fractal | 结果以 Fractal 格式存储 | +| Bias | BT Buffer | 特定格式 | 通过 `copy_cbuf_to_bt` 加载 | + +### fractalSizes 参数含义 + +`fractalSizes` 参数是一个包含两个 int64 值的数组 `[blockM, blockN]`，描述 Fractal 块的大小： + +- `blockM`：Fractal 块在 M 维度（行方向）的大小 +- `blockN`：Fractal 块在 N 维度（列方向）的大小 + +典型的 Fractal 块大小为 16x16（对于 FP16/FP32 数据类型）。 + +### mmadL1 的布局接口 + +源文件：[HIVMMacroOps.td:152-172](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td#L152-L172) + +`MmadL1Op` 实现了 `OpLayoutInterface`，提供以下布局查询方法： + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getOperandALayout()` | `FailureOr` | 获取矩阵 A 的目标布局 | +| `getOperandBLayout()` | `FailureOr` | 获取矩阵 B 的目标布局 | +| `getOperandCLayout()` | `FailureOr` | 获取矩阵 C 的目标布局 | +| `getOperandBiasLayout()` | `FailureOr` | 获取 Bias 的目标布局 | +| `getOperandsTargetFractalLayout()` | 接口方法 | 获取所有操作数的目标 Fractal 布局 | + +### Fixpipe 的布局转换 + +源文件：[HIVMAttrs.td:847-859](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L847-L859) + +Fixpipe 支持三种布局转换模式（`dma_mode` 属性）： + +| 模式 | IR 标识符 | 数值 | 说明 | +|------|-----------|------|------| +| NZ2ND | `nz2nd` | 0 | 将 NZ 格式转为 ND 格式 | +| NZ2DN | `nz2dn` | 1 | 将 NZ 格式转为 DN 格式（仅 950 系列） | +| NZ2NZ | `normal` | 2 | 保持 NZ 格式不变（默认值） | + +## 910B vs 950 架构的数据流 + +### Ascend910B 架构数据流 + +``` +GM (ND格式) + | + | [MTE2 + nd2nz] + v +L1 (NZ格式) + | + | [MTE1] + v +L0A (zN格式) L0B (nZ格式) + | | + +------+-------+ + | + | [M: Cube矩阵乘法] + v + L0C (Fractal格式) + | + | [FIX + dma_mode=NZ2ND] + v + GM (ND格式) +``` + +### Ascend950 架构数据流 + +``` +GM (ND格式) + | + | [MTE2 + nd2nz] + v +L1 (NZ格式) + | + | [MTE1] + v +L0A (zN格式) L0B (nZ格式) + | | + +------+-------+ + | + | [M: Cube矩阵乘法] + v + L0C (Fractal格式, 256KB) + | + +----+----+----+ + | | | + | [FIX] | [FIX] | [FIX] + | NZ2ND | NZ2ND | NZ2ND/NZ2DN + v v v + GM (ND) L1 (NZ) UB (ND) ----[V]---- UB ----[MTE3]---- GM + ^ + | + 紧耦合缓冲区/InsertCVDataMovement + (950特有: L0C直通UB) +``` + +950 架构的关键差异： +1. L0C 容量增大到 256KB（910B 为 128KB） +2. 支持 L0C -> UB 直通通路，Fixpipe 可直接将结果输出到 UB +3. 支持 NZ2DN 布局转换模式 +4. 支持紧耦合缓冲区，Cube 和 Vector 之间可高效传递数据 +5. 支持 Dual Dst 模式（ROW_SPLIT / COLUMN_SPLIT），将 Cube 结果切分给两个 Vector 单元 + +## 常见问题 + +**Q: 为什么 Cube 单元需要 Fractal 布局而不是 ND 布局？** +A: Fractal 布局将矩阵按固定大小的块（如 16x16）重新排列，使得 Cube 单元可以高效地按块加载和计算。这种布局优化了 L0A/L0B 缓存的数据局部性，减少了缓存未命中。 + +**Q: DOTA_ND 和 ND 有什么区别？** +A: ND 是通用的行优先布局标记，DOTA_ND 是专门用于矩阵乘法中矩阵 A 的布局标记，它额外携带 `transpose` 属性来标记是否需要转置。`isNDLayout()` 方法对两者都返回 true。 + +**Q: nd2nz 和 fixpipe 的 NZ2ND 有什么区别？** +A: `nd2nz` 是 GM -> L1 方向的布局转换（ND 转 NZ），属于 PIPE_MTE2。`fixpipe` 的 NZ2ND 是 L0C -> GM/L1/UB 方向的布局转换（NZ 转 ND），属于 PIPE_FIX。两者方向相反，且在不同的 Pipeline 上执行。 + +**Q: fractalSizes 参数在什么情况下需要设置？** +A: 当使用 `Fractal` 布局类型时，需要通过 `fractalSizes` 指定具体的分形块大小。对于 `zN`、`nZ`、`ND` 等标准布局，fractalSizes 通常不需要设置。 + +**Q: 950 系列的 NZ2DN 模式是什么？** +A: NZ2DN 是 950 架构新增的布局转换模式，将 NZ 格式转为 DN（Diagonal-N）格式。源码中标注此模式仅在 Ascend950 上支持。 + +## 相关文档 + +- 源码参考：[HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) +- 源码参考：[HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) +- 源码参考：[HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td) +- 上一节：[03-pipeline-execution-model.md](./03-pipeline-execution-model.md) — Pipeline 执行模型 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/00-overview.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/00-overview.md new file mode 100644 index 00000000..5ebb9905 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/00-overview.md @@ -0,0 +1,318 @@ +# HIVM 方言总览 + +> 关键词：HIVM、Hybrid Intelligence Virtual Machine、方言定义、操作分类 + +## 概述 + +HIVM（Hybrid Intelligence Virtual Machine，混合智能虚拟机）是 AscendNPU-IR 项目中的核心 IR 方言，用于描述华为 Ascend NPU 上的计算操作。HIVM 方言定义在 BiShengIR 编译栈中，作为硬件无关与硬件相关 IR 之间的桥梁，承载了从高级语义到低级硬件指令的映射。 + +HIVM 方言的设计目标是： +- 提供统一的 IR 抽象，覆盖 Ascend NPU 的 Cube（矩阵计算）、Vector（向量计算）和 Mix（混合计算）三种核心类型 +- 支持 DMA 数据搬运、向量计算、矩阵乘法宏操作、同步控制、自定义操作和底层内建指令 +- 通过 Pipeline 和 Address Space 属性精确描述硬件资源约束 + +## 方言定义 + +### 基本属性 + +| 属性 | 值 | +|------|------| +| 方言名称 | `hivm` | +| C++ 命名空间 | `::mlir::hivm` | +| 操作前缀 | `hir.` | +| TableGen 定义文件 | [HIVMBase.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMBase.td) | + +### 依赖方言 + +HIVM 方言声明了以下依赖方言（定义于 [HIVMBase.td:L33-L37](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMBase.td#L33-L37)）： + +| 依赖方言 | 用途 | +|----------|------| +| `arith::ArithDialect` | 算术运算基础 | +| `bishengir::memref_ext::MemRefExtDialect` | MemRef 扩展操作 | +| `math::MathDialect` | 数学运算 | +| `memref::MemRefDialect` | 内存引用操作 | +| `hacc::HACCDialect` | 硬件加速器配置 | +| `tensor::TensorDialect` | 张量操作 | + +### 操作基类 + +HIVM 定义了两个核心操作基类（[HIVMBase.td:L52-L68](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMBase.td#L52-L68)）： + +- **`HIVM_Op`**：基础操作基类，所有 HIVM 操作的前缀为 `hir.` +- **`HIVM_StructuredOp`**：结构化操作基类，继承 `HIVM_Op` 并实现 `HIVMStructuredOpInterface`、`MemoryEffectsOpInterface`、`FlattenInterface`、`LibraryFunctionOpInterface` 等接口 + +## 操作分类总表 + +### DMA 操作（数据搬运） + +定义文件：[HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) + +| 操作名 | IR 语法 | 一句话描述 | +|--------|---------|-----------| +| `LoadOp` | `hir.load` | 从 GM 加载数据到本地缓冲区（UB） | +| `StoreOp` | `hir.store` | 从本地缓冲区（UB）存储数据到 GM | +| `CopyOp` | `hir.copy` | 本地内存层级间的数据拷贝 | +| `FixpipeOp` | `hir.fixpipe` | L0C 到其他内存层级的搬运，支持随路量化/激活 | +| `ND2NZOp` | `hir.nd2nz` | ND 到 NZ 布局转换的数据搬运 | +| `NZ2NDOp` | `hir.nz2nd` | L1 到 GM 的 NZ 到 ND 布局转换搬运 | +| `L12UBOp` | `hir.l12ub` | L1 到 UB 的数据搬运 | +| `AtomicCasOp` | `hir.atomic_cas` | 原子比较并交换操作 | +| `AtomicXchgOp` | `hir.atomic_xchg` | 原子交换操作 | + +### 向量操作 + +定义文件：[HIVMVectorOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td) + +| 操作名 | IR 语法 | 一句话描述 | +|--------|---------|-----------| +| `VExpOp` | `hir.vexp` | 逐元素指数运算 | +| `VAbsOp` | `hir.vabs` | 逐元素绝对值运算 | +| `VLnOp` | `hir.vln` | 逐元素自然对数运算 | +| `VReluOp` | `hir.vrelu` | 逐元素 ReLU 激活 | +| `VRsqrtOp` | `hir.vrsqrt` | 逐元素平方根倒数 | +| `VSqrtOp` | `hir.vsqrt` | 逐元素平方根 | +| `VTanhOp` | `hir.vtanh` | 逐元素双曲正切 | +| `VSinOp` | `hir.vsin` | 逐元素正弦 | +| `VCosOp` | `hir.vcos` | 逐元素余弦 | +| `VErfOp` | `hir.verf` | 逐元素误差函数 | +| `VRecOp` | `hir.vrec` | 逐元素倒数 | +| `VNotOp` | `hir.vnot` | 逐元素取反 | +| `VCastOp` | `hir.vcast` | 逐元素类型转换 | +| `VAddOp` | `hir.vadd` | 逐元素加法 | +| `VMulOp` | `hir.vmul` | 逐元素乘法 | +| `VMulExtOp` | `hir.vmulext` | 逐元素乘法取高 32 位 | +| `VSubOp` | `hir.vsub` | 逐元素减法 | +| `VDivOp` | `hir.vdiv` | 逐元素除法 | +| `VMaxOp` | `hir.vmax` | 逐元素最大值 | +| `VMinOp` | `hir.vmin` | 逐元素最小值 | +| `VOrOp` | `hir.vor` | 逐元素或运算 | +| `VAndOp` | `hir.vand` | 逐元素与运算 | +| `VXorOp` | `hir.vxor` | 逐元素异或运算 | +| `VModOp` | `hir.vmod` | 逐元素取模 | +| `VModUIOp` | `hir.vmodui` | 逐元素无符号取模 | +| `VShLOp` | `hir.vshl` | 逐元素左移 | +| `VShROp` | `hir.vshr` | 逐元素右移 | +| `VCmpOp` | `hir.vcmp` | 逐元素比较 | +| `VPowOp` | `hir.vpow` | 逐元素幂运算 | +| `VSelOp` | `hir.vsel` | 逐元素选择 | +| `VBrcOp` | `hir.vbrc` | 向量广播 | +| `VReduceOp` | `hir.vreduce` | 向量规约 | +| `VTransposeOp` | `hir.vtranspose` | 向量转置 | +| `VArangeOp` | `hir.varange` | 向量范围生成 | +| `VInterleaveOp` | `hir.vinterleave` | 向量交织 | +| `VDeinterleaveOp` | `hir.vdeinterleave` | 向量解交织 | +| `VFlipOp` | `hir.vflip` | 向量翻转 | +| `VMulextendedOp` | `hir.vmulextended` | 扩展乘法（高低 16 位） | +| `VPadOp` | `hir.vpad` | 向量填充 | +| `VConcatOp` | `hir.vconcat` | 向量拼接 | +| `VGatherOp` | `hir.vgather` | 向量收集 | +| `VCumprodOp` | `hir.vcumprod` | 累积乘积 | +| `VCumsumOp` | `hir.vcumsum` | 累积求和 | +| `VSortOp` | `hir.vsort` | 向量排序 | + +### 宏操作（矩阵乘法） + +定义文件：[HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td) + +| 操作名 | IR 语法 | 一句话描述 | +|--------|---------|-----------| +| `MmadL1Op` | `hir.mmadL1` | L1 输入的矩阵乘加操作 | +| `BatchMmadL1Op` | `hir.batchMmadL1` | L1 输入的批量矩阵乘加操作 | +| `MatmulOp` | `hir.matmul` | GM 输入的全局矩阵乘法 | +| `MixMatmulOp` | `hir.mix_matmul` | 支持后向量融合的混合矩阵乘法 | +| `MixGroupMatmulOp` | `hir.mix_group_matmul` | 支持分组专家的混合矩阵乘法 | + +### 同步操作 + +定义文件：[HIVMSynchronizationOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMSynchronizationOps.td) + +| 操作名 | IR 语法 | 一句话描述 | +|--------|---------|-----------| +| `SetFlagOp` | `hir.set_flag` | 设置同步标志 | +| `WaitFlagOp` | `hir.wait_flag` | 等待同步标志 | +| `PipeBarrierOp` | `hir.pipe_barrier` | Pipeline 屏障 | +| `SyncBlockOp` | `hir.sync_block` | 块间同步 | +| `SyncBlockSetOp` | `hir.sync_block_set` | 设置块同步标志 | +| `SyncBlockWaitOp` | `hir.sync_block_wait` | 等待块同步标志 | +| `CreateSyncBlockLockOp` | `hir.create_sync_block_lock` | 创建块同步锁 | +| `SyncBlockLockOp` | `hir.sync_block_lock` | 获取块同步锁 | +| `SyncBlockUnlockOp` | `hir.sync_block_unlock` | 释放块同步锁 | + +### Custom / 基础操作 + +定义文件：[HIVMOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMOps.td) + +| 操作名 | IR 语法 | 一句话描述 | +|--------|---------|-----------| +| `GetBlockIdxOp` | `hir.get_block_idx` | 获取当前线程块索引 | +| `GetBlockNumOp` | `hir.get_block_num` | 获取线程块数量 | +| `GetSubBlockIdxOp` | `hir.get_sub_block_idx` | 获取子块索引 | +| `GetSubBlockNumOp` | `hir.get_sub_block_num` | 获取子块数量 | +| `SetAtomicOp` | `hir.set_atomic` | 设置原子操作模式 | +| `SetMaskNormOp` | `hir.set_mask_norm` | 设置掩码规范化 | +| `SetCtrlOp` | `hir.set_ctrl` | 设置控制位 | +| `LoadScalarOp` | `hir.load_scalar` | 加载标量 | +| `DCCIOp` | `hir.dcci` | 数据缓存清理/无效化 | +| `ConvertLayoutOp` | `hir.convert_layout` | 布局转换 | +| `PointerCastOp` | `hir.pointer_cast` | 指针类型转换 | +| `BitcastOp` | `hir.bitcast` | 位重解释 | +| `SetFFTSBaseAddrOp` | `hir.set_ffts_base_addr` | 设置 FFTS 基地址 | +| `CustomOp` | `hir.custom` | 自定义操作 | +| `CustomMacroOp` | `hir.custom_macro` | 自定义宏操作 | +| `GatherLoadOp` | `hir.gather_load` | 稀疏内存加载 | +| `ScatterStoreOp` | `hir.scatter_store` | 稀疏内存存储 | +| `LocalLoadOp` | `hir.local_load` | 从 UB 加载张量（SIMD 到 SIMT） | +| `LocalStoreOp` | `hir.local_store` | 存储张量到 UB（SIMT 到 SIMD） | +| `IndirectLoadOp` | `hir.indirect_load` | 间接内存加载 | +| `IndirectStoreOp` | `hir.indirect_store` | 间接内存存储 | +| `GatherTOp` | `hir.gatherT` | 按轴收集操作 | +| `IndexPutOp` | `hir.index_put` | 按轴散列写操作 | +| `ScatterTOp` | `hir.scatterT` | 按轴散列存储操作 | +| `EmbeddingGatherOp` | `hir.embedding_gather` | 嵌入查找操作 | +| `DebugOp` | `hir.debug` | 设备端调试 | +| `InitDebugOp` | `hir.init_debug` | 初始化调试 | +| `FinishDebugOp` | `hir.finish_debug` | 结束调试 | + +### 内建指令 + +定义文件：[HIVMIntrinOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMIntrinOps.td) + +| 操作名 | IR 语法 | 一句话描述 | +|--------|---------|-----------| +| `GetBlockIdxInstrOp` | `hir.intr.hivm.GET.BLOCK.IDX` | 获取块索引指令 | +| `GetBlockNumInstrOp` | `hir.intr.hivm.GET.BLOCK.NUM` | 获取块数量指令 | +| `GetSubBlockIdxInstrOp` | `hir.intr.hivm.GET.SUBBLOCKID` | 获取子块索引指令 | +| `GetSubBlockNumInstrOp` | `hir.intr.hivm.GET.SUBBLOCKDIM` | 获取子块维度指令 | +| `SetFlagImmInstrOp` | `hir.intr.hivm.SET.FLAG.IMM` | 立即数设置标志指令 | +| `WaitFlagImmInstrOp` | `hir.intr.hivm.WAIT.FLAG.IMM` | 立即数等待标志指令 | +| `SetFlagRegInstrOp` | `hir.intr.hivm.SET.FLAG.REG` | 寄存器设置标志指令 | +| `WaitFlagRegInstrOp` | `hir.intr.hivm.WAIT.FLAG.REG` | 寄存器等待标志指令 | +| `PipeBarrierInstrOp` | `hir.intr.hivm.BARRIER` | Pipeline 屏障指令 | +| `SetFftsBaseAddrInstrOp` | `hir.intr.hivm.SET.FFTS.BASE.ADDR` | 设置 FFTS 基地址指令 | +| `SetCrossCoreInstrOp` | `hir.intr.hivm.SET.CROSS.CORE` | FFTS 跨核同步指令 | +| `WaitFlagDevInstrOp` | `hir.intr.hivm.WAIT.FLAG.DEV.REG` | FFTS 块/子块同步指令 | +| `SetMaskNormInstrOp` | `hir.intr.hivm.SET.MASK.NORM` | 设置掩码规范化指令 | +| `GetCtrlInstrOp` | `hir.intr.hivm.GET.CTRL` | 获取控制位指令 | +| `SetCtrlInstrOp` | `hir.intr.hivm.SET.CTRL` | 设置控制位指令 | +| `SBitSet0InstrOp` | `hir.intr.hivm.SBITSET0` | 设置状态位为 0 | +| `SBitSet1InstrOp` | `hir.intr.hivm.SBITSET1` | 设置状态位为 1 | +| `DCCIDstInstrOp` | `hir.intr.hivm.DCCI.DST` | 数据缓存清理 GM 指令 | +| `DCCIDstUBInstrOp` | `hir.intr.hivm.DCCI.DST.UB` | 数据缓存清理 UB 指令 | + +## 核心属性与枚举 + +### Address Space（地址空间） + +定义于 [HIVMAttrs.td:L170-L197](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L170-L197) + +| 枚举值 | IR 字面量 | 含义 | +|--------|----------|------| +| `Zero` | `zero` | 默认地址空间 | +| `GM` | `gm` | 全局内存 | +| `L1` | `cbuf` | L1 缓存（Cube Buffer） | +| `L0A` | `ca` | L0A 缓存（Cube A 矩阵） | +| `L0B` | `cb` | L0B 缓存（Cube B 矩阵） | +| `L0C` | `cc` | L0C 缓存（Cube C 矩阵/累加器） | +| `UB` | `ub` | 统一缓冲区（Vector Buffer） | + +### PIPE（Pipeline） + +定义于 [HIVMAttrs.td:L203-L244](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L203-L244) + +| 枚举值 | 含义 | +|--------|------| +| `PIPE_S` | Scalar Pipeline | +| `PIPE_V` | Vector Pipeline | +| `PIPE_M` | Cube Pipeline | +| `PIPE_MTE1` | MTE1 Pipeline（L1 搬运） | +| `PIPE_MTE2` | MTE2 Pipeline（GM 到本地搬运） | +| `PIPE_MTE3` | MTE3 Pipeline（本地到 GM 搬运） | +| `PIPE_FIX` | Fix Pipeline（L0C 搬出） | +| `PIPE_MTE4` | MTE4 Pipeline | +| `PIPE_MTE5` | MTE5 Pipeline | + +### Core Type（核心类型） + +定义于 [HIVMAttrs.td:L298-L317](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L298-L317) + +| 枚举值 | 含义 | +|--------|------| +| `CUBE` | Cube 核心（矩阵计算） | +| `VECTOR` | Vector 核心（向量计算） | +| `CUBE_OR_VECTOR` | Cube 或 Vector 核心 | +| `CUBE_AND_VECTOR` | Cube 和 Vector 核心（混合） | + +## 与 HFusion / Triton 方言的关系 + +``` +Triton Python API + | + v + Triton Dialect (tt) -- 前端 IR，描述 Triton 语义 + | + v + HFusion Dialect -- 融合 IR，描述算子融合与 Tiling + | + v + HIVM Dialect (hir) -- 硬件映射 IR，描述 NPU 操作 + | + v + LLVM / CCE 指令 -- 最终代码生成 +``` + +- **Triton 方言**：用户编写的 Triton Python 代码首先被编译为 Triton 方言 IR，描述高级语义（如 `tt.load`、`tt.store`、`tt.dot`） +- **HFusion 方言**：Triton 方言经过 Tiling 和融合优化后转换为 HFusion 方言，描述算子融合策略和 Tile 级别的计算 +- **HIVM 方言**：HFusion 方言进一步 lowering 为 HIVM 方言，精确映射到 Ascend NPU 的硬件操作，包括 DMA 搬运、Pipeline 分配、同步控制等 + +## 方言在编译栈中的位置 + +``` ++--------------------------------------------------+ +| Triton Python Source | ++--------------------------------------------------+ + | + v ++--------------------------------------------------+ +| Triton Dialect (tt) | +| tt.load, tt.store, tt.dot, tt.make_range ... | ++--------------------------------------------------+ + | + v ++--------------------------------------------------+ +| HFusion Dialect | +| 融合策略、Tiling、Bufferization | ++--------------------------------------------------+ + | + v ++--------------------------------------------------+ +| HIVM Dialect (hir) | +| hir.load, hir.store, hir.mmadL1, hir.vadd ... | +| Pipeline 分配、同步插入、内存规划 | ++--------------------------------------------------+ + | + v ++--------------------------------------------------+ +| CCE / LLVM IR | +| 库函数调用、内建指令 | ++--------------------------------------------------+ + | + v ++--------------------------------------------------+ +| Ascend NPU 二进制 | ++--------------------------------------------------+ +``` + +HIVM 方言处于编译栈的中间层，是连接高级语义与底层硬件的关键桥梁。在此阶段： +- 数据搬运被映射为具体的 DMA 操作（Load/Store/Copy/Fixpipe 等） +- 计算操作被分配到 Cube 或 Vector 核心 +- Pipeline 归属被确定，同步操作被插入 +- 内存层级（GM/L1/UB/L0A/L0B/L0C）被显式标注 + +## 相关文档 + +- DMA 操作详解：[01-DMA-Operations/00-overview.md](01-DMA-Operations/00-overview.md) +- 源码参考： + - [HIVMBase.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMBase.td) - 方言基础定义 + - [HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) - 属性与枚举定义 + - [HIVMInterfaces.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMInterfaces.td) - 接口定义 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/00-overview.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/00-overview.md new file mode 100644 index 00000000..9855b4e0 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/00-overview.md @@ -0,0 +1,181 @@ +# DMA 操作总览 + +> 关键词：DMA、数据搬运、Pipeline、Address Space、MTE2、MTE3、FIX + +## 概述 + +DMA（Direct Memory Access）操作是 HIVM 方言中最核心的操作类别之一，负责在 Ascend NPU 的不同内存层级之间搬运数据。DMA 操作直接映射到硬件的 DMA 引擎，通过不同的 Pipeline 执行，是实现高效数据流的关键。 + +Ascend NPU 的内存层级和 DMA 操作构成了一个层次化的数据通路，理解这些通路对于编写高性能内核至关重要。 + +## 内存层级与数据通路 + +### 内存层级 + +``` ++----------------------------------------------------------+ +| GM (Global Memory) | +| 片外全局内存，大容量高延迟 | ++----------------------------------------------------------+ + | ^ + | MTE2 (Load/ND2NZ) | MTE3 (Store/NZ2ND) + v | ++----------------------------------------------------------+ +| L1 (Cube Buffer) | +| 片上缓存，Cube 专用 | ++----------------------------------------------------------+ + | ^ + | MTE1 (L12UB) | + v | ++----------------------------------------------------------+ +| UB (Unified Buffer) | +| 统一缓冲区，Vector 专用 | ++----------------------------------------------------------+ + | ^ + | | V (Copy UB->UB) + v v ++----------------------------------------------------------+ +| UB (Unified Buffer) | ++----------------------------------------------------------+ + ++----------------------------------------------------------+ +| L0A / L0B | +| Cube 矩阵输入缓存（由 MMAD 隐式使用） | ++----------------------------------------------------------+ + ^ + | M (MMAD 写入 L0C) + | ++----------------------------------------------------------+ +| L0C | +| Cube 矩阵累加器缓存 | ++----------------------------------------------------------+ + | + | FIX (Fixpipe) + v ++----------------------------------------------------------+ +| GM / UB / L1 | +| Fixpipe 输出目标 | ++----------------------------------------------------------+ +``` + +### 数据通路图 + +``` + GM + / | \ + / | \ + MTE2 / | \ MTE2 + Load / | \ ND2NZ + v | v + UB L1 L1 + | ^ ^ + | Copy | | + v | MTE3 | + UB ------+ | + | / + MTE3 | / + Store | / Copy (Ascend950) + v v + GM UB + + +L0C ---- FIX (Fixpipe) ----> GM / UB / L1 + 支持随路量化/激活 +``` + +## Pipeline 归属表 + +每个 DMA 操作都归属于特定的硬件 Pipeline，这决定了操作的执行单元和同步方式。 + +| 操作 | Pipeline | 枚举值 | 说明 | +|------|----------|--------|------| +| `hir.load` | MTE2 | `PIPE_MTE2` | GM 到本地缓冲区的数据加载 | +| `hir.nd2nz` | MTE2 | `PIPE_MTE2` | GM 到 L1 的 ND 到 NZ 布局转换加载 | +| `hir.store` | MTE3 | `PIPE_MTE3` | 本地缓冲区到 GM 的数据存储 | +| `hir.nz2nd` | MTE3 | `PIPE_MTE3` | L1 到 GM 的 NZ 到 ND 布局转换存储 | +| `hir.l12ub` | MTE1 | `PIPE_MTE1` | L1 到 UB 的数据搬运 | +| `hir.fixpipe` | FIX | `PIPE_FIX` | L0C 到其他层级的搬运 | +| `hir.copy` | 动态 | 取决于 src/dst | 根据源和目标地址空间动态确定 | +| `hir.atomic_cas` | 无固定 | - | 原子比较并交换 | +| `hir.atomic_xchg` | 无固定 | - | 原子交换 | + +### Copy 操作的动态 Pipeline + +`hir.copy` 操作的 Pipeline 根据源和目标地址空间动态确定（定义于 [HIVMDMAOps.cpp:L607-L632](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L607-L632)）： + +| 源地址空间 | 目标地址空间 | Pipeline | +|-----------|-------------|----------| +| UB | UB | `PIPE_V` | +| L0C | GM | `PIPE_FIX` | +| GM | L1 | `PIPE_MTE2` | +| UB | L1 | `PIPE_MTE3` | + +## 同步要求概述 + +DMA 操作涉及多个 Pipeline 之间的数据依赖，需要通过同步机制保证数据一致性。HIVM 提供了以下同步原语： + +### Pipeline 内同步 + +- **`hir.set_flag` / `hir.wait_flag`**：Pipeline 间的标志同步，用于生产者-消费者模式 +- **`hir.pipe_barrier`**：Pipeline 内屏障，确保同一 Pipeline 内操作的顺序性 + +### 块间同步 + +- **`hir.sync_block`**：不同 Kernel 间的同步 +- **`hir.sync_block_set` / `hir.sync_block_wait`**：细粒度块间标志同步 + +### 典型同步模式 + +``` +hir.load (MTE2) ---- set_flag[MTE2, V] --> +hir.vadd (V) ---- wait_flag[MTE2, V] --> + set_flag[V, MTE3] --> +hir.store (MTE3) ---- wait_flag[V, MTE3] --> +``` + +## DMA 操作一览 + +| 操作 | IR 语法 | 数据通路 | Pipeline | 布局转换 | 随路功能 | +|------|---------|---------|----------|---------|---------| +| `hir.load` | `hir.load ins(...) outs(...)` | GM -> UB | MTE2 | 无 | Padding、Eviction | +| `hir.store` | `hir.store ins(...) outs(...)` | UB -> GM | MTE3 | 无 | 原子操作 | +| `hir.copy` | `hir.copy ins(...) outs(...)` | UB->UB, GM->L1, UB->L1 | 动态 | 无 | Padding | +| `hir.fixpipe` | `hir.fixpipe ins(...) outs(...)` | L0C->GM/UB/L1 | FIX | NZ2ND/NZ2DN/NZ2NZ | 量化、ReLU、双目标 | +| `hir.nd2nz` | `hir.nd2nz ins(...) outs(...)` | GM -> L1 | MTE2 | ND -> NZ | 初始化缓冲区 | +| `hir.nz2nd` | `hir.nz2nd ins(...) outs(...)` | L1 -> GM | MTE3 | NZ -> ND | 无 | +| `hir.l12ub` | `hir.l12ub ins(...) outs(...)` | L1 -> UB | MTE1 | NZ -> ND | 无 | +| `hir.atomic_cas` | `hir.atomic_cas ins(...) outs(...)` | GM <-> UB | 无 | 无 | 原子 CAS | +| `hir.atomic_xchg` | `hir.atomic_xchg ins(...) outs(...)` | GM <-> UB | 无 | 无 | 原子交换 | + +## 支持的数据通路汇总 + +| 通路 | 操作 | 硬件约束 | +|------|------|---------| +| GM -> UB | `hir.load` | 所有 Ascend 型号 | +| GM -> L1 | `hir.nd2nz` | 所有 Ascend 型号 | +| UB -> GM | `hir.store` | 所有 Ascend 型号 | +| L1 -> GM | `hir.nz2nd` | 所有 Ascend 型号 | +| L1 -> UB | `hir.l12ub` | 所有 Ascend 型号 | +| UB -> UB | `hir.copy` | 所有 Ascend 型号 | +| GM -> L1 | `hir.copy` | 所有 Ascend 型号 | +| UB -> L1 | `hir.copy` | 仅 Ascend950 系列 | +| L0C -> GM | `hir.fixpipe` | 所有 Ascend 型号 | +| L0C -> L1 | `hir.fixpipe` | 所有 Ascend 型号 | +| L0C -> UB | `hir.fixpipe` | 仅 Ascend950 系列 | + +## 相关文档 + +- [hir.load 详解](01-load.md) +- [hir.store 详解](02-store.md) +- [hir.nd2nz 详解](03-nd2nz.md) +- [hir.nz2nd 详解](04-nz2nd.md) +- [hir.copy 详解](05-copy.md) +- [hir.fixpipe 详解](06-fixpipe.md) +- [原子操作详解](07-atomic.md) +- [Gather/Scatter 详解](08-gather-scatter.md) +- [间接访问详解](09-indirect-access.md) +- [随路功能详解](10-padding-quantization.md) +- 源码参考： + - [HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) - DMA 操作 TableGen 定义 + - [HIVMDMAOps.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp) - DMA 操作实现 + - [dma-ops.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/dma-ops.mlir) - DMA 操作测试用例 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/01-load.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/01-load.md new file mode 100644 index 00000000..d0b1d8f4 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/01-load.md @@ -0,0 +1,319 @@ +# hir.load + +> 关键词：Load、DMA、MTE2、GM 到 UB、Padding、Eviction Policy + +## 概述 + +`hir.load` 是 HIVM 方言中的核心 DMA 操作，用于将数据从全局内存（GM）加载到本地缓冲区（目前仅支持加载到统一缓冲区 UB）。该操作映射到硬件的 MTE2 Pipeline，是大多数内核中最频繁执行的 DMA 操作。 + +`hir.load` 支持丰富的随路功能，包括 Padding（填充）、Eviction Policy（驱逐策略）和缓冲区初始化，使其能够处理边界不齐、数据对齐等常见场景。 + +> Python API 对应：tl.load -- 详见 [docs_triton_ascend/02-Core-API/01-memory-ops.md](../../../docs_triton_ascend/02-Core-API/01-memory-ops.md) + +## IR 操作定义 + +### TableGen 定义 + +来源：[HIVMDMAOps.td:L53-L136](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td#L53-L136) + +```tablegen +def LoadOp : HIVM_DmaOp<"load", [ + AttrSizedOperandSegments, StaticMaxRankTrait<3>, + SinglePipeOpTrait, OpPipeTrait<"PIPE::PIPE_MTE2">, + DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + DeclareOpInterfaceMethods, + OperElemTypeConstraints<[0], [I8, UI8, I16, UI16, F16, BF16, + I32, UI32, F32, UI64, I64, F8E4M3FN, F8E5M2]>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, +]> { + let summary = "HIVM data load operation"; + let description = [{ + Loads the data from the global memory to the local buffer. + Currently only support loading to the unified buffer. + + Examples: + ```mlir + hivm.load ins(%src : memref<16x16xf16, #hivm.address_space>) outs(%dst : memref<16x16xf16, #hivm.address_space>) + ``` + + Constraints: + - `src` and `dst` are expected to have the same element type. + - If `pad_mode` is not set, `src` and `dst` shape should be the same. + - Supports both left and right padding. + - `pad_value` should have the same element type as `src` and `dst`. + }]; + let arguments = (ins TensorOrMemref:$src, + TensorOrMemref:$dst, + OptionalAttr:$pad_mode, + Optional:$pad_value, + Optional:$left_padding_num, + Optional:$right_padding_num, + DefaultValuedOptionalAttr:$init_out_buffer, + Optional:$init_condition, + OptionalAttr:$eviction_policy + ); + let results = (outs Optional:$result_tensor); + let assemblyFormat = [{ + `ins` `(` $src `:` type($src) `)` + `outs` `(` $dst `:` type($dst) `)` + attr-dict + (`pad_mode` `=` $pad_mode^)? + (`pad_value` `=` $pad_value^ `:` type($pad_value))? + (`left_padding_num` `=` $left_padding_num^ `:` type($left_padding_num))? + (`init_out_buffer` `=` $init_out_buffer^ )? + (`right_padding_num` `=` $right_padding_num^ `:` type($right_padding_num))? + (`init_condition` `=` $init_condition^ `:` type($init_condition))? + (`eviction_policy` `=` $eviction_policy^)? + (`->` type($result_tensor)^)? + }]; + let hasFolder = 1; + let hasVerifier = 1; + let extraClassDeclaration = DmaOpBaseDecl; +} +``` + +### MLIR 语法 + +```mlir +hivm.hir.load ins(%src : memref>) + outs(%dst : memref>) + +hivm.hir.load ins(%src : tensor) + outs(%dst : tensor) + pad_mode = #hivm.padmode + pad_value = %val : f16 + -> tensor +``` + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| `src` | TensorOrMemref | 是 | 源数据缓冲区 | 必须为 GM 地址空间（memref 语义下） | +| `dst` | TensorOrMemref | 是 | 目标数据缓冲区 | 必须为 UB 地址空间（memref 语义下） | +| `pad_mode` | HIVM_PadModeAttr | 否 | 填充模式 | PadNull / PadFirstElem / PadValue | +| `pad_value` | AnyType | 否 | 填充值 | 类型须与 src/dst 元素类型相同 | +| `left_padding_num` | Index | 否 | 左侧填充元素数 | 仅在 pad_mode 设置时有意义 | +| `right_padding_num` | AnyType | 否 | 右侧填充元素数 | 仅在 pad_mode 设置时有意义 | +| `init_out_buffer` | BoolAttr | 否 | 是否初始化输出缓冲区 | 默认 false | +| `init_condition` | AnyType | 否 | 初始化条件 | 用于条件化初始化 | +| `eviction_policy` | HIVM_EvictionPolicyAttr | 否 | 缓存驱逐策略 | EvictFirst / EvictLast | + +### 结果说明 + +| 结果 | 类型 | 说明 | +|------|------|------| +| `result_tensor` | Optional | Tensor 语义下的结果张量 | + +### 属性说明 + +| 属性 | 类型 | 默认值 | 说明 | 可选值 | +|------|------|--------|------|--------| +| `pad_mode` | HIVM_PadModeAttr | 无 | 填充模式 | `PadNull`(0)、`PadFirstElem`(1)、`PadValue`(2) | +| `init_out_buffer` | BoolAttr | `false` | 是否在加载前初始化输出缓冲区 | `true` / `false` | +| `eviction_policy` | HIVM_EvictionPolicyAttr | 无 | 数据缓存驱逐策略 | `EvictFirst`(0)、`EvictLast`(1) | + +### 数据类型约束 + +来源：`OperElemTypeConstraints<[0], [...]>`，约束操作数索引 0（src）的元素类型。 + +| 支持的元素类型 | 说明 | +|---------------|------| +| `i8`, `ui8` | 8 位整数 | +| `i16`, `ui16` | 16 位整数 | +| `f16` | 半精度浮点 | +| `bf16` | BFloat16 | +| `i32`, `ui32` | 32 位整数 | +| `f32` | 单精度浮点 | +| `ui64`, `i64` | 64 位整数 | +| `f8E4M3FN` | 8 位浮点 E4M3 格式 | +| `f8E5M2` | 8 位浮点 E5M2 格式 | + +### PadMode 枚举 + +定义于 [HIVMAttrs.td:L330-L349](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L330-L349) + +| 枚举值 | 数值 | IR 字面量 | 说明 | +|--------|------|----------|------| +| `PadNull` | 0 | `PadNull` | 不填充 | +| `PadFirstElem` | 1 | `PadFirstElem` | 使用第一个元素填充 | +| `PadValue` | 2 | `PadValue` | 使用指定值填充 | + +### EvictionPolicy 枚举 + +定义于 [HIVMAttrs.td:L356-L372](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L356-L372) + +| 枚举值 | 数值 | IR 字面量 | 说明 | +|--------|------|----------|------| +| `EvictFirst` | 0 | `EvictFirst` | 优先驱逐 | +| `EvictLast` | 1 | `EvictLast` | 最后驱逐 | + +## IR 示例 + +### 基础加载 + +最简单的 GM 到 UB 数据加载，源和目标形状相同： + +```mlir +func.func @hivm_memref_load_gm_to_ub() { + %src = memref.alloc() : memref<16x16xf16, #hivm.address_space> + %dst = memref.alloc() : memref<16x16xf16, #hivm.address_space> + hivm.hir.load ins(%src : memref<16x16xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + return +} +``` + +### 带 Padding 的加载 + +当源数据形状小于目标缓冲区时，使用 Padding 填充多余位置： + +```mlir +func.func @hivm_memref_copy_gm_to_ub_pad_value() { + %val = arith.constant 10.0 : f16 + %src = memref.alloc() : memref<16x16xf16, #hivm.address_space> + %dst = memref.alloc() : memref<16x16xf16, #hivm.address_space> + hivm.hir.load ins(%src : memref<16x16xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + pad_mode = #hivm.padmode + pad_value = %val : f16 + return +} +``` + +使用 PadFirstElem 模式（用第一个元素填充）： + +```mlir +func.func @hivm_memref_copy_gm_to_ub_pad_first() { + %src = memref.alloc() : memref<16x15xf16, #hivm.address_space> + %dst = memref.alloc() : memref<16x16xf16, #hivm.address_space> + hivm.hir.load ins(%src : memref<16x15xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + pad_mode = #hivm.padmode + return +} +``` + +带左侧填充数和仅指定 pad_value（自动推断 PadValue 模式）： + +```mlir +func.func @hivm_memref_copy_gm_to_ub_pad_value_only() { + %val = arith.constant 10.0 : f16 + %src = memref.alloc() : memref<16x16xf16, #hivm.address_space> + %dst = memref.alloc() : memref<16x16xf16, #hivm.address_space> + hivm.hir.load ins(%src : memref<16x16xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + pad_value = %val : f16 + return +} +``` + +### 带 Eviction Policy 的加载 + +```mlir +func.func @hivm_memref_load_with_eviction() { + %src = memref.alloc() : memref<16x16xf16, #hivm.address_space> + %dst = memref.alloc() : memref<16x16xf16, #hivm.address_space> + hivm.hir.load ins(%src : memref<16x16xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + eviction_policy = #hivm.evictionpolicy + return +} +``` + +### Tensor 语义 + +```mlir +func.func @hivm_tensor_load() -> tensor<16x16xf32> { + %src = tensor.empty() : tensor<16x16xf32> + %dst = tensor.empty() : tensor<16x16xf32> + %res = hivm.hir.load ins(%src : tensor<16x16xf32>) + outs(%dst : tensor<16x16xf32>) + -> tensor<16x16xf32> + return %res : tensor<16x16xf32> +} +``` + +## IR 层约束与验证 + +来源：[HIVMDMAOps.cpp:L219-L272](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L219-L272) + +`hir.load` 的验证器执行以下检查： + +1. **元素类型一致性**：`src` 和 `dst` 的元素类型必须相同 +2. **Rank 一致性**：`src` 和 `dst` 必须具有相同的 rank +3. **形状兼容性**：如果未设置 `pad_mode`，`src` 和 `dst` 的形状必须兼容（相同） +4. **PadValue 必要性**：如果 `pad_mode` 为 `PadValue`，则 `pad_value` 必须提供 +5. **PadValue 类型一致性**：`pad_value` 的类型必须与 `dst` 的元素类型相同 +6. **地址空间约束**（memref 语义下）： + - `src` 必须为 GM 地址空间 + - `dst` 不能为 GM 地址空间 +7. **Tensor 语义约束**： + - `result_tensor` 的元素类型必须与 `dst` 相同 + - `result_tensor` 的 rank 必须与 `dst` 相同 + - 如果未设置 `pad_mode`，`result_tensor` 的形状必须与 `dst` 兼容 + +### 验证错误示例 + +``` +error: element types of dst and src should be the same! +error: src and dst should have the same dimensions! +error: if pad_mode is not set, src and dst shape should be the same! +error: if padmode is PadValue, pad_value is required! +error: dtype of pad_value and element type of dst/src should be the same! +error: only support src == gm and dst != gm currently! +``` + +## 与其他 IR 操作的关系 + +### 从 Triton 到 HIVM + +``` +tt.load --> hivm.hir.load + (带 padding 参数映射) +``` + +Triton 的 `tt.load` 中的 `padding_option` 和 `other` 参数映射到 HIVM 的 `pad_mode` 和 `pad_value`。 + +### 从 HFusion 到 HIVM + +HFusion 方言中的加载操作在 lowering 到 HIVM 时，会根据目标地址空间和布局需求生成 `hir.load` 或 `hir.nd2nz`。 + +### 后续 Lowering + +`hir.load` 最终被 lowering 为库函数调用，函数名格式为： + +``` +load_gm_to_ubuf_{rank}d_{datatype} +``` + +## 常见问题 + +### Q: 为什么 load 只支持 GM -> UB？ + +A: 这是当前硬件的限制。如果需要将数据加载到 L1，应使用 `hir.nd2nz`；如果需要在本地层级间搬运，应使用 `hir.copy`。 + +### Q: pad_mode 和 pad_value 的关系是什么？ + +A: `pad_mode` 决定填充策略。如果仅指定 `pad_value` 而不指定 `pad_mode`，编译器会自动推断为 `PadValue` 模式。如果 `pad_mode` 为 `PadFirstElem`，则不需要 `pad_value`。 + +### Q: init_out_buffer 的作用是什么？ + +A: 当设置为 `true` 时，会在加载数据之前先将目标缓冲区初始化为零。这在某些需要确保缓冲区干净的场景下有用，例如当加载的数据不能完全覆盖目标缓冲区时。 + +### Q: eviction_policy 如何选择？ + +A: `EvictFirst` 表示该数据在缓存中优先被驱逐，适用于只使用一次的数据；`EvictLast` 表示该数据尽量保留在缓存中，适用于会被反复访问的数据。 + +## 相关文档 + +- Python API：[docs_triton_ascend/02-Core-API/01-memory-ops.md](../../../docs_triton_ascend/02-Core-API/01-memory-ops.md) +- DMA 操作总览：[00-overview.md](00-overview.md) +- 随路功能详解：[10-padding-quantization.md](10-padding-quantization.md) +- 源码参考： + - [HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) - TableGen 定义 + - [HIVMDMAOps.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp) - 验证逻辑实现 + - [dma-ops.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/dma-ops.mlir) - IR 测试用例 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/02-store.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/02-store.md new file mode 100644 index 00000000..125f53c8 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/02-store.md @@ -0,0 +1,237 @@ +# hir.store + +> 关键词：Store、DMA、MTE3、UB 到 GM、原子操作 + +## 概述 + +`hir.store` 是 HIVM 方言中的核心 DMA 操作，用于将数据从本地缓冲区（UB）存储到全局内存（GM）。该操作映射到硬件的 MTE3 Pipeline，是大多数内核中数据写回的主要方式。 + +`hir.store` 支持原子操作模式，可以在存储时执行原子加、原子最大值、原子最小值等操作，用于实现 Triton 中的原子存储语义。 + +> Python API 对应：tl.store -- 详见 [docs_triton_ascend/02-Core-API/01-memory-ops.md](../../../docs_triton_ascend/02-Core-API/01-memory-ops.md) + +## IR 操作定义 + +### TableGen 定义 + +来源：[HIVMDMAOps.td:L138-L190](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td#L138-L190) + +```tablegen +def StoreOp : HIVM_DmaOp<"store", [ + StaticMaxRankTrait<3>, + SinglePipeOpTrait, OpPipeTrait<"PIPE::PIPE_MTE3">, + DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + DeclareOpInterfaceMethods, + OperElemTypeConstraints<[0], [I8, UI8, I16, UI16, F16, BF16, + I32, UI32, F32, UI64, I64, F8E4M3FN, F8E5M2]>, + DeclareOpInterfaceMethods, +]> { + let summary = "HIVM data store operation"; + let description = [{ + Stores the data on local buffer to global memory. + Currently only support storing data on the unified buffer. + + Examples: + ```mlir + hivm.store ins(%src : memref<16x16xf16, #hivm.address_space>) outs(%dst : memref<16x16xf16, #hivm.address_space>) + ``` + + Constraints: + - `src` and `dst` are expected to have the same element type. + - If `atomic_kind` is set, the kind is one of `add`, `max`, `min`. + }]; + let arguments = (ins TensorOrMemref:$src, + TensorOrMemref:$dst, + OptionalAttr:$atomic_kind + ); + let results = (outs Optional:$result_tensor); + let builders = [ + OpBuilder<(ins "TypeRange":$res, "Value":$src, "Value":$dst)> + ]; + let assemblyFormat = [{ + `ins` `(` $src `:` type($src) `)` + `outs` `(` $dst `:` type($dst) `)` + attr-dict + (`atomic` `=` $atomic_kind^)? + (`->` type($result_tensor)^)? + }]; + let hasFolder = 1; + let hasVerifier = 1; + let extraClassDeclaration = DmaOpBaseDecl # [{ + // Return whether atomic store is enabled. + bool isAtomic(); + + // Return whether block-sync atomic store enabled is implemented by hardware. + bool isHWAtomic(); + + // Return whether block-sync atomic store enabled is implemented by software. + bool isSWAtomic(); + }]; +} +``` + +### MLIR 语法 + +```mlir +hivm.hir.store ins(%src : memref>) + outs(%dst : memref>) + +hivm.hir.store ins(%src : memref>) + outs(%dst : memref>) + atomic = #hivm.atomic_kind +``` + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| `src` | TensorOrMemref | 是 | 源数据缓冲区 | 必须为 UB 地址空间（memref 语义下） | +| `dst` | TensorOrMemref | 是 | 目标数据缓冲区 | 必须为 GM 地址空间（memref 语义下） | +| `atomic_kind` | HIVM_AtomicKindAttr | 否 | 原子操作类型 | NONE / ADD / MAX / MIN / AND / OR / XOR / CAS / XCHG | + +### 结果说明 + +| 结果 | 类型 | 说明 | +|------|------|------| +| `result_tensor` | Optional | Tensor 语义下的结果张量 | + +### 属性说明 + +| 属性 | 类型 | 默认值 | 说明 | 可选值 | +|------|------|--------|------|--------| +| `atomic_kind` | HIVM_AtomicKindAttr | 无 | 原子操作类型 | 见下表 | + +#### AtomicKind 枚举 + +定义于 [HIVMAttrs.td:L637-L664](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L637-L664) + +| 枚举值 | 数值 | IR 字面量 | 说明 | 实现方式 | +|--------|------|----------|------|---------| +| `NONE` | 0 | `none` | 非原子操作 | - | +| `ADD` | 1 | `add` | 原子加 | 硬件实现 | +| `MAX` | 2 | `max` | 原子最大值 | 硬件实现 | +| `MIN` | 3 | `min` | 原子最小值 | 硬件实现 | +| `AND` | 4 | `and` | 原子与 | 软件实现 | +| `OR` | 5 | `or` | 原子或 | 软件实现 | +| `XOR` | 6 | `xor` | 原子异或 | 软件实现 | +| `CAS` | 7 | `cas` | 原子比较并交换 | 软件实现 | +| `XCHG` | 8 | `xchg` | 原子交换 | 软件实现 | + +### 数据类型约束 + +来源：`OperElemTypeConstraints<[0], [...]>` + +| 支持的元素类型 | 说明 | +|---------------|------| +| `i8`, `ui8` | 8 位整数 | +| `i16`, `ui16` | 16 位整数 | +| `f16` | 半精度浮点 | +| `bf16` | BFloat16 | +| `i32`, `ui32` | 32 位整数 | +| `f32` | 单精度浮点 | +| `ui64`, `i64` | 64 位整数 | +| `f8E4M3FN` | 8 位浮点 E4M3 格式 | +| `f8E5M2` | 8 位浮点 E5M2 格式 | + +## IR 示例 + +### 基础存储 + +最简单的 UB 到 GM 数据存储： + +```mlir +func.func @hivm_memref_store_ub_to_gm() { + %src = memref.alloc() : memref<16x16xf16, #hivm.address_space> + %dst = memref.alloc() : memref<16x16xf16, #hivm.address_space> + hivm.hir.store ins(%src : memref<16x16xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + return +} +``` + +### 原子加存储 + +```mlir +func.func @hivm_memref_atomic_add_store() { + %src = memref.alloc() : memref<16x16xf32, #hivm.address_space> + %dst = memref.alloc() : memref<16x16xf32, #hivm.address_space> + hivm.hir.store ins(%src : memref<16x16xf32, #hivm.address_space>) + outs(%dst : memref<16x16xf32, #hivm.address_space>) + atomic = #hivm.atomic_kind + return +} +``` + +### 原子最大值存储 + +```mlir +func.func @hivm_memref_atomic_max_store() { + %src = memref.alloc() : memref<16x16xf16, #hivm.address_space> + %dst = memref.alloc() : memref<16x16xf16, #hivm.address_space> + hivm.hir.store ins(%src : memref<16x16xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + atomic = #hivm.atomic_kind + return +} +``` + +## IR 层约束与验证 + +来源：[HIVMDMAOps.cpp:L348-L376](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L348-L376) + +1. **元素类型一致性**：`src` 和 `dst` 的元素类型必须相同 +2. **Rank 一致性**：`src` 和 `dst` 必须具有相同的 rank +3. **地址空间约束**（memref 语义下）： + - `src` 必须为 UB 地址空间 + - `dst` 必须为 GM 地址空间 +4. **Tensor 语义约束**： + - `result_tensor` 的元素类型必须与 `dst` 相同 + - `result_tensor` 的 rank 必须与 `dst` 相同 + +### 原子操作分类 + +来源：[HIVMDMAOps.cpp:L378-L394](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L378-L394) + +- **硬件原子操作**（`isHWAtomic()`）：ADD、MAX、MIN -- 由硬件直接支持 +- **软件原子操作**（`isSWAtomic()`）：AND、OR、XOR、CAS、XCHG -- 通过软件锁机制实现 + +## 与其他 IR 操作的关系 + +### 从 Triton 到 HIVM + +``` +tt.store --> hivm.hir.store +tt.store (atomic) --> hivm.hir.store atomic = #hivm.atomic_kind +``` + +### 后续 Lowering + +`hir.store` 最终被 lowering 为库函数调用，函数名格式为： + +``` +store_ubuf_to_gm_{rank}d_{datatype} +``` + +## 常见问题 + +### Q: 为什么 store 只支持 UB -> GM？ + +A: 这是当前硬件的限制。如果需要将 L1 数据写回 GM，应使用 `hir.nz2nd`；如果需要将 L0C 数据写出，应使用 `hir.fixpipe`。 + +### Q: 硬件原子和软件原子有什么区别？ + +A: 硬件原子（ADD/MAX/MIN）由 Ascend NPU 的 DMA 引擎直接支持，性能更高。软件原子（AND/OR/XOR/CAS/XCHG）需要通过软件锁机制实现，开销较大。在性能敏感场景中，应优先使用硬件原子操作。 + +### Q: 如何在 store 之前设置全局原子模式？ + +A: 可以使用 `hir.set_atomic` 操作设置全局原子模式，后续的 store 操作将自动启用原子语义。使用 `hir.set_atomic kind= #hivm.atomic_kind` 重置。 + +## 相关文档 + +- Python API：[docs_triton_ascend/02-Core-API/01-memory-ops.md](../../../docs_triton_ascend/02-Core-API/01-memory-ops.md) +- DMA 操作总览：[00-overview.md](00-overview.md) +- 原子操作详解：[07-atomic.md](07-atomic.md) +- 源码参考： + - [HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) - TableGen 定义 + - [HIVMDMAOps.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp) - 验证逻辑实现 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/03-nd2nz.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/03-nd2nz.md new file mode 100644 index 00000000..5e246241 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/03-nd2nz.md @@ -0,0 +1,219 @@ +# hir.nd2nz + +> 关键词：ND2NZ、布局转换、MTE2、GM 到 L1、Fractal 布局 + +## 概述 + +`hir.nd2nz` 是 HIVM 方言中的 DMA 操作，用于将数据从全局内存（GM）搬运到 L1 缓存，同时执行 ND（Normal Data）到 NZ（NzFormat/Z-Fractals）的布局转换。该操作映射到硬件的 MTE2 Pipeline。 + +ND 到 NZ 的布局转换是 Ascend NPU 矩阵计算的关键前置步骤。Cube 核心的矩阵乘法操作（MMAD）要求输入数据以 NZ（Fractal）格式存储在 L1 缓存中，因此 `hir.nd2nz` 是从 GM 加载矩阵数据到 Cube 计算流水线的必经之路。 + +> Python API 对应：无直接对应，由编译器自动插入 + +## IR 操作定义 + +### TableGen 定义 + +来源：[HIVMDMAOps.td:L328-L368](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td#L328-L368) + +```tablegen +def ND2NZOp : HIVM_DmaOp<"nd2nz", [ + AttrSizedOperandSegments, SinglePipeOpTrait, OpPipeTrait<"PIPE::PIPE_MTE2">, + HIVMCoreTypeInterface, CubeCoreTypeTrait, NoMaxRankTrait, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods +]> { + let summary = "HIVM data copy operation with on-the-fly ND to NZ layout transformation"; + let description = [{ + - `dst_continuous`: if present, signify that the source data is stored continuously + in the destination buffer. This must be set in order for this op to be converted to + library function call. + Constraints: + - if `init_out_buffer` is true, `pad_value` should have value. + }]; + let arguments = (ins AnyShaped:$src, + AnyShaped:$dst, + OptionalAttr:$dst_continuous, + DefaultValuedOptionalAttr:$init_out_buffer, + Optional:$pad_value, + Optional:$init_condition + ); + let results = (outs Variadic:$result_tensor); + let assemblyFormat = [{ + attr-dict + `ins` `(` $src `:` type($src) `)` + `outs` `(` $dst `:` type($dst) `)` + (`init_out_buffer` `=` $init_out_buffer^ )? + (`pad_value` `=` $pad_value^ `:` type($pad_value))? + (`init_condition` `=` $init_condition^ `:` type($init_condition))? + (`->` type($result_tensor)^)? + }]; + let builders = [ + OpBuilder<(ins "TypeRange" : $res, "Value" : $src, "Value" : $dst, + "UnitAttr" : $dst_continuous)>, + OpBuilder<(ins "TypeRange" : $res, "Value" : $src, "Value" : $dst, + "UnitAttr" : $dst_continuous, "bool" : $init_out_buffer, + "Value" : $pad_value)>, + ]; + let extraClassDeclaration = DmaOpBaseDecl; +} +``` + +### MLIR 语法 + +```mlir +hivm.hir.nd2nz ins(%src : memref) outs(%dst : memref) + +hivm.hir.nd2nz {dst_continuous} + ins(%src : memref) outs(%dst : memref) + +hivm.hir.nd2nz {dst_continuous} + ins(%src : memref>) + outs(%dst : memref<8x16x16x8xf32, #hivm.address_space>) + init_out_buffer = true pad_value = %cst : f32 +``` + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| `src` | AnyShaped | 是 | 源数据缓冲区 | 通常为 GM 地址空间 | +| `dst` | AnyShaped | 是 | 目标数据缓冲区 | 通常为 L1 地址空间 | +| `dst_continuous` | UnitAttr | 否 | 目标数据是否连续存储 | 设置后才能转换为库函数调用 | +| `init_out_buffer` | BoolAttr | 否 | 是否初始化输出缓冲区 | 默认 false；若为 true 则 pad_value 必须有值 | +| `pad_value` | AnyType | 否 | 填充值 | 类型须与 src/dst 元素类型相同 | +| `init_condition` | AnyType | 否 | 初始化条件 | 用于条件化初始化 | + +### 结果说明 + +| 结果 | 类型 | 说明 | +|------|------|------| +| `result_tensor` | Variadic | Tensor 语义下的结果张量 | + +### 属性说明 + +| 属性 | 类型 | 默认值 | 说明 | 可选值 | +|------|------|--------|------|--------| +| `dst_continuous` | UnitAttr | 无 | 目标数据连续存储标志 | 存在即表示连续 | +| `init_out_buffer` | BoolAttr | `false` | 是否在搬运前初始化输出缓冲区 | `true` / `false` | + +## IR 示例 + +### 基础 ND 到 NZ 转换 + +```mlir +func.func @test_nd2nz() { + %gmA = memref.alloc() : memref<1024x2048xf16> + %gmASubview = memref.subview %gmA[0, 0][256, 128][1, 1] + : memref<1024x2048xf16> to + memref<256x128xf16, strided<[2048, 1], offset: 0>> + %l1A = memref.alloc() : memref<256x128xf16> + hivm.hir.nd2nz ins(%gmASubview : memref<256x128xf16, strided<[2048, 1], offset: 0>>) + outs(%l1A: memref<256x128xf16>) + return +} +``` + +### 带连续存储标志 + +```mlir +hivm.hir.nd2nz {dst_continuous} + ins(%gmASubview : memref<256x128xf16, strided<[2048, 1], offset: 0>>) + outs(%l1A: memref<256x128xf16>) +``` + +### 带缓冲区初始化 + +```mlir +func.func @test_nd2nz_tensor_init_out_buffer(%arg0: memref>) { + %cst = arith.constant 1.000000e+00 : f32 + %c0 = arith.constant 0 : index + %alloc = memref.alloc() : memref<8x16x16x8xf32, #hivm.address_space> + hivm.hir.nd2nz {dst_continuous} + ins(%arg0 : memref>) + outs(%alloc : memref<8x16x16x8xf32, #hivm.address_space>) + init_out_buffer = true pad_value = %cst : f32 + return +} +``` + +### Tensor 语义 + +```mlir +func.func @test_nd2nz_tensor() { + %gmA = tensor.empty() : tensor<1024x2048xf16> + %gmASubview = tensor.extract_slice %gmA[0, 0][256, 128][1, 1] + : tensor<1024x2048xf16> to tensor<256x128xf16> + %l1A = tensor.empty() : tensor<256x128xf16> + %ret = hivm.hir.nd2nz ins(%gmASubview : tensor<256x128xf16>) + outs(%l1A: tensor<256x128xf16>) -> tensor<256x128xf16> + return +} +``` + +## IR 层约束与验证 + +1. **元素类型一致性**：`src` 和 `dst` 的元素类型必须相同 +2. **init_out_buffer 与 pad_value**：如果 `init_out_buffer` 为 true，则 `pad_value` 必须提供 +3. **Core Type**：该操作属于 Cube 核心类型（`CubeCoreTypeTrait`） +4. **Pipeline**：固定为 MTE2（`OpPipeTrait<"PIPE::PIPE_MTE2">`） +5. **无最大 Rank 限制**：`NoMaxRankTrait` 表示不限制操作数的最大 rank + +### 库函数命名 + +`hir.nd2nz` 的库函数名会根据下游操作自动调整。如果下游是 `MmadL1Op` 且当前操作提供 per-channel bias，则函数名会添加 `_forbias` 后缀（[HIVMDMAOps.cpp:L660-L671](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L660-L671)）。 + +## 与其他 IR 操作的关系 + +### ND 与 NZ 布局 + +``` +ND 布局 (Normal Data): NZ 布局 (Fractal): ++---+---+---+---+ +===+===+ +| 0 | 1 | 2 | 3 | |0,1|2,3| ++---+---+---+---+ +===+===+ +| 4 | 5 | 6 | 7 | --> |4,5|6,7| ++---+---+---+---+ +===+===+ +| 8 | 9 |10 |11 | |8,9|10,11| ++---+---+---+---+ +===+===+ + +行优先连续存储 按 Fractal Block 重新排列 + 适配 Cube 矩阵计算引擎 +``` + +### 典型使用模式 + +``` +hir.nd2nz (GM -> L1, ND -> NZ) + | + v +hir.mmadL1 (L1 -> L0A/L0B -> L0C, 矩阵乘加) + | + v +hir.fixpipe (L0C -> GM/UB, 可选 NZ -> ND) +``` + +## 常见问题 + +### Q: 什么时候需要使用 nd2nz 而不是 load？ + +A: 当数据需要被 Cube 核心用于矩阵乘法时，必须使用 `nd2nz` 将数据从 GM 搬运到 L1 并转换为 NZ 格式。如果数据只是被 Vector 核心使用，则应使用 `hir.load` 将数据搬运到 UB。 + +### Q: dst_continuous 标志的作用是什么？ + +A: `dst_continuous` 表示目标缓冲区中的数据是连续存储的。设置此标志后，操作才能被转换为优化的库函数调用。如果未设置，编译器可能需要生成更通用的搬运代码。 + +### Q: 为什么 nd2nz 属于 Cube 核心类型？ + +A: 因为 ND 到 NZ 的布局转换是 Cube 矩阵计算流水线的专用操作，转换后的 NZ 格式数据仅供 Cube 核心的 MMAD 操作使用。 + +## 相关文档 + +- DMA 操作总览：[00-overview.md](00-overview.md) +- hir.nz2nd：[04-nz2nd.md](04-nz2nd.md) +- hir.fixpipe：[06-fixpipe.md](06-fixpipe.md) +- 源码参考： + - [HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) - TableGen 定义 + - [HIVMDMAOps.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp) - 实现代码 + - [dma-ops.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/dma-ops.mlir) - 测试用例 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/04-nz2nd.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/04-nz2nd.md new file mode 100644 index 00000000..856f1d58 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/04-nz2nd.md @@ -0,0 +1,160 @@ +# hir.nz2nd + +> 关键词：NZ2ND、布局转换、MTE3、L1 到 GM、Fractal 到 ND + +## 概述 + +`hir.nz2nd` 是 HIVM 方言中的 DMA 操作，用于将数据从 L1 缓存搬运到全局内存（GM），同时执行 NZ（NzFormat/Z-Fractals）到 ND（Normal Data）的布局转换。该操作映射到硬件的 MTE3 Pipeline。 + +`hir.nz2nd` 是 `hir.nd2nz` 的逆操作，通常用于将 Cube 计算产生的 NZ 格式中间结果从 L1 写回 GM，同时恢复为 ND 格式以便后续处理。 + +> Python API 对应：无直接对应，由编译器自动插入 + +## IR 操作定义 + +### TableGen 定义 + +来源：[HIVMDMAOps.td:L370-L386](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td#L370-L386) + +```tablegen +def NZ2NDOp : HIVM_DmaOp<"nz2nd", + [StaticMaxRankTrait<2>, + SinglePipeOpTrait, OpPipeTrait<"PIPE::PIPE_MTE3">, + HIVMCoreTypeInterface, CubeCoreTypeTrait + ]> { + let summary = "HIVM data copy operation from L1 to Global Memory with NZ2ND conversion"; + let description = [{ NZ2ND does data movement from L1 to OUT with NZ2ND conversion. }]; + let arguments = (ins TensorOrMemref:$src, TensorOrMemref:$dst); + let results = (outs Optional:$result_tensor); + let assemblyFormat = [{ + attr-dict + `ins` `(` $src `:` type($src) `)` + `outs` `(` $dst `:` type($dst) `)` + (`->` type($result_tensor)^)? + }]; + let extraClassDeclaration = DmaOpBaseDecl; +} +``` + +### MLIR 语法 + +```mlir +hivm.hir.nz2nd ins(%src : memref>) + outs(%dst : memref>) +``` + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| `src` | TensorOrMemref | 是 | 源数据缓冲区 | 必须为 L1 地址空间 | +| `dst` | TensorOrMemref | 是 | 目标数据缓冲区 | 必须为 GM 地址空间 | + +### 结果说明 + +| 结果 | 类型 | 说明 | +|------|------|------| +| `result_tensor` | Optional | Tensor 语义下的结果张量 | + +### 属性说明 + +`hir.nz2nd` 没有额外属性，是最简洁的 DMA 操作之一。 + +### 数据类型约束 + +`hir.nz2nd` 继承自 `HIVM_DmaOp`，受 `HIVM_StructuredOp` 的通用类型约束。源和目标的元素类型必须相同。 + +### Rank 约束 + +`StaticMaxRankTrait<2>` 限制操作数的最大 rank 为 2。 + +## IR 示例 + +### 基础 NZ 到 ND 转换 + +```mlir +func.func @hivm_nz2nd_l1_to_gm() { + %src = memref.alloc() : memref<256x128xf16, #hivm.address_space> + %dst = memref.alloc() : memref<256x128xf16, #hivm.address_space> + hivm.hir.nz2nd ins(%src : memref<256x128xf16, #hivm.address_space>) + outs(%dst : memref<256x128xf16, #hivm.address_space>) + return +} +``` + +### Tensor 语义 + +```mlir +func.func @hivm_nz2nd_tensor() -> tensor<256x128xf16> { + %src = tensor.empty() : tensor<256x128xf16> + %dst = tensor.empty() : tensor<256x128xf16> + %res = hivm.hir.nz2nd ins(%src : tensor<256x128xf16>) + outs(%dst : tensor<256x128xf16>) + -> tensor<256x128xf16> + return %res : tensor<256x128xf16> +} +``` + +## IR 层约束与验证 + +1. **元素类型一致性**：`src` 和 `dst` 的元素类型必须相同 +2. **地址空间约束**（memref 语义下）： + - `src` 必须为 L1 地址空间 + - `dst` 必须为 GM 地址空间 +3. **Rank 约束**：最大 rank 为 2 +4. **Core Type**：属于 Cube 核心类型（`CubeCoreTypeTrait`） +5. **Pipeline**：固定为 MTE3（`OpPipeTrait<"PIPE::PIPE_MTE3">`） + +### 库函数命名 + +`hir.nz2nd` 的库函数名格式为： + +``` +nz2nd_{src_rank}d_to_{dst_rank}d_{datatype} +``` + +来源：[HIVMDMAOps.cpp:L693-L714](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L693-L714) + +## 与其他 IR 操作的关系 + +### 与 nd2nz 的关系 + +``` +hir.nd2nz (GM -> L1, ND -> NZ) hir.nz2nd (L1 -> GM, NZ -> ND) + | ^ + v | + [Cube 计算] [Cube 结果写出] + hir.mmadL1 | + | | + v | + hir.fixpipe (L0C -> L1, NZ) -----------+ +``` + +### 与 fixpipe 的关系 + +`hir.nz2nd` 仅执行 L1 到 GM 的布局转换搬运。如果需要从 L0C 搬运数据并同时执行 NZ 到 ND 的布局转换，应使用 `hir.fixpipe` 并设置 `dma_mode = NZ2ND`。 + +### 典型使用场景 + +`hir.nz2nd` 通常用于以下场景： +1. 将 Cube 计算的中间结果从 L1 写回 GM +2. 在多轮矩阵乘法之间，将 NZ 格式的部分和写回 GM 暂存 + +## 常见问题 + +### Q: nz2nd 和 fixpipe 的 NZ2ND 模式有什么区别？ + +A: `hir.nz2nd` 是独立的 L1 -> GM 搬运操作，仅执行布局转换。`hir.fixpipe` 的 NZ2ND 模式从 L0C 搬运数据到 GM，同时支持随路量化和激活。如果数据源是 L0C，应优先使用 `hir.fixpipe`；如果数据源已经是 L1 中的 NZ 格式数据，则使用 `hir.nz2nd`。 + +### Q: 为什么 nz2nd 的最大 rank 限制为 2？ + +A: 这是由硬件约束决定的。NZ 到 ND 的布局转换在 2D 矩阵上最为常见，更高维度的数据通常需要先 reshape 为 2D 再进行转换。 + +## 相关文档 + +- DMA 操作总览：[00-overview.md](00-overview.md) +- hir.nd2nz：[03-nd2nz.md](03-nd2nz.md) +- hir.fixpipe：[06-fixpipe.md](06-fixpipe.md) +- 源码参考： + - [HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) - TableGen 定义 + - [HIVMDMAOps.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp) - 实现代码 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/05-copy.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/05-copy.md new file mode 100644 index 00000000..273ee6e7 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/05-copy.md @@ -0,0 +1,243 @@ +# hir.copy + +> 关键词：Copy、DMA、本地搬运、UB 到 UB、Padding + +## 概述 + +`hir.copy` 是 HIVM 方言中的 DMA 操作，用于在本地内存层级之间拷贝数据。与 `hir.load` 和 `hir.store` 不同，`hir.copy` 的 Pipeline 归属是动态的，取决于源和目标地址空间的组合。 + +当前支持的拷贝通路包括： +- UB 到 UB（Vector Pipeline） +- GM 到 L1（MTE2 Pipeline） +- UB 到 L1（MTE3 Pipeline，仅 Ascend950 系列） + +`hir.copy` 支持随路 Padding 功能，可以在拷贝过程中填充目标缓冲区的多余位置。 + +> Python API 对应：无直接对应，由编译器在 bufferization 阶段自动生成 + +## IR 操作定义 + +### TableGen 定义 + +来源：[HIVMDMAOps.td:L192-L244](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td#L192-L244) + +```tablegen +def CopyOp : HIVM_DmaOp<"copy", [ + SinglePipeOpTrait, StaticMaxRankTrait<3>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + DeclareOpInterfaceMethods, + OperElemTypeConstraints<[0], [I1, I8, UI8, I16, UI16, F16, BF16, + I32, UI32, F32, UI64, I64, F8E4M3FN, F8E5M2]>, + DeclareOpInterfaceMethods, +]> { + let summary = "HIVM data copy operation"; + let description = [{ + Copy the data between local memory hierarchies. + Currently support: + - UB to UB + - UB to L1 (for Ascend950 series) + + Examples: + ```mlir + hivm.hir.copy ins(%src : memref<16x16xf16, #hivm.address_space>) outs(%dst : memref<16x16xf16, #hivm.address_space>) + ``` + + Constraints: + - `src` and `dst` are expected to have the same element type. + - If `pad_mode` is not set, `src` and `dst` shape should be the same. + - Only support left padding. + - `pad_value` should have the same element type as `src` and `dst`. + }]; + let arguments = (ins TensorOrMemref:$src, + TensorOrMemref:$dst, + OptionalAttr:$pad_mode, + Optional:$pad_value + ); + let results = (outs Optional:$result_tensor); + let builders = [ + OpBuilder<(ins "TypeRange":$res, "Value":$src, "Value":$dst)> + ]; + let assemblyFormat = [{ + `ins` `(` $src `:` type($src) `)` + `outs` `(` $dst `:` type($dst) `)` + attr-dict + (`pad_mode` `=` $pad_mode^)? + (`pad_value` `=` $pad_value^ `:` type($pad_value))? + (`->` type($result_tensor)^)? + }]; + let hasFolder = 1; + let hasVerifier = 1; + let extraClassDeclaration = DmaOpBaseDecl # [{ + // Declare functions necessary for SinglePipeOpTrait. + PIPE getPipe(); + }]; +} +``` + +### MLIR 语法 + +```mlir +hivm.hir.copy ins(%src : memref<16x16xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + +hivm.hir.copy ins(%src : tensor<16x16xf32>) + outs(%dst : tensor<16x16xf32>) + -> tensor<16x16xf32> +``` + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| `src` | TensorOrMemref | 是 | 源数据缓冲区 | 取决于通路 | +| `dst` | TensorOrMemref | 是 | 目标数据缓冲区 | 取决于通路 | +| `pad_mode` | HIVM_PadModeAttr | 否 | 填充模式 | PadNull / PadFirstElem / PadValue | +| `pad_value` | AnyType | 否 | 填充值 | 类型须与 src/dst 元素类型相同 | + +### 结果说明 + +| 结果 | 类型 | 说明 | +|------|------|------| +| `result_tensor` | Optional | Tensor 语义下的结果张量 | + +### 属性说明 + +| 属性 | 类型 | 默认值 | 说明 | 可选值 | +|------|------|--------|------|--------| +| `pad_mode` | HIVM_PadModeAttr | 无 | 填充模式 | `PadNull`(0)、`PadFirstElem`(1)、`PadValue`(2) | + +### 数据类型约束 + +来源：`OperElemTypeConstraints<[0], [...]>` + +| 支持的元素类型 | 说明 | +|---------------|------| +| `i1` | 布尔类型 | +| `i8`, `ui8` | 8 位整数 | +| `i16`, `ui16` | 16 位整数 | +| `f16` | 半精度浮点 | +| `bf16` | BFloat16 | +| `i32`, `ui32` | 32 位整数 | +| `f32` | 单精度浮点 | +| `ui64`, `i64` | 64 位整数 | +| `f8E4M3FN` | 8 位浮点 E4M3 格式 | +| `f8E5M2` | 8 位浮点 E5M2 格式 | + +### 动态 Pipeline 归属 + +来源：[HIVMDMAOps.cpp:L607-L632](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L607-L632) + +| 源地址空间 | 目标地址空间 | Pipeline | +|-----------|-------------|----------| +| UB | UB | `PIPE_V` | +| L0C | GM | `PIPE_FIX` | +| GM | L1 | `PIPE_MTE2` | +| UB | L1 | `PIPE_MTE3` | + +## IR 示例 + +### UB 到 UB 拷贝 + +```mlir +func.func @hivm_memref_copy_ub_to_ub() { + %src = memref.alloc() : memref<16x16xf16, #hivm.address_space> + %dst = memref.alloc() : memref<16x16xf16, #hivm.address_space> + hivm.hir.copy ins(%src : memref<16x16xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + return +} +``` + +### UB 到 L1 拷贝（Ascend950） + +```mlir +module attributes {hacc.target = #hacc.target<"Ascend950PR_9589">} { + func.func @hivm_memref_copy_ub_to_l1() { + %src = memref.alloc() : memref<16x16xf16, #hivm.address_space> + %dst = memref.alloc() : memref<16x16xf16, #hivm.address_space> + hivm.hir.copy ins(%src : memref<16x16xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + return + } +} +``` + +### Tensor 语义拷贝 + +```mlir +func.func @hivm_tensor_copy() -> tensor<16x16xf32> { + %src = tensor.empty() : tensor<16x16xf32> + %dst = tensor.empty() : tensor<16x16xf32> + %res = hivm.hir.copy ins(%src : tensor<16x16xf32>) + outs(%dst : tensor<16x16xf32>) + -> tensor<16x16xf32> + return %res : tensor<16x16xf32> +} +``` + +## IR 层约束与验证 + +来源：[HIVMDMAOps.cpp:L552-L605](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L552-L605) + +1. **元素类型一致性**：`src` 和 `dst` 的元素类型必须相同 +2. **Rank 一致性**：`src` 和 `dst` 必须具有相同的 rank +3. **形状兼容性**：如果未设置 `pad_mode`，`src` 和 `dst` 的形状必须兼容 +4. **PadValue 必要性**：如果 `pad_mode` 为 `PadValue`，则 `pad_value` 必须提供 +5. **PadValue 类型一致性**：`pad_value` 的类型必须与 `dst` 的元素类型相同 +6. **地址空间约束**（memref 语义下）： + - 仅支持以下通路：UB->UB, GM->L1, UB->L1（Ascend950） +7. **Padding 限制**：仅支持左侧 Padding + +### 支持的地址空间组合 + +来源：[HIVMDMAOps.cpp:L440-L448](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L440-L448) + +| 组合 | 支持情况 | +|------|---------| +| UB -> UB | 所有型号 | +| GM -> L1 | 所有型号 | +| UB -> L1 | 仅 Ascend950 | + +## 与其他 IR 操作的关系 + +### 与 load/store 的区别 + +| 操作 | 数据通路 | Pipeline | Padding | 布局转换 | +|------|---------|----------|---------|---------| +| `hir.load` | GM -> UB | MTE2 | 支持（左+右） | 无 | +| `hir.store` | UB -> GM | MTE3 | 不支持 | 无 | +| `hir.copy` | 本地 <-> 本地 | 动态 | 仅左侧 | 无 | + +### 后续 Lowering + +`hir.copy` 最终被 lowering 为库函数调用，函数名格式为： + +``` +copy_{src_space}_to_{dst_space}_{rank}d_{datatype} +``` + +## 常见问题 + +### Q: copy 和 load 的 Padding 有什么区别？ + +A: `hir.copy` 仅支持左侧 Padding，而 `hir.load` 同时支持左侧和右侧 Padding。这是因为 copy 操作的硬件实现在 Padding 方面的限制。 + +### Q: 什么时候使用 copy 而不是 load/store？ + +A: 当数据搬运不涉及 GM 时（如 UB 到 UB），应使用 `hir.copy`。当涉及 GM 时，应使用 `hir.load`（GM 到本地）或 `hir.store`（本地到 GM）。 + +### Q: 为什么 UB 到 L1 的拷贝仅支持 Ascend950？ + +A: 这是硬件约束。Ascend950 系列的 Vector 核心可以直接写入 L1 缓存，而其他型号不支持此功能。 + +## 相关文档 + +- DMA 操作总览：[00-overview.md](00-overview.md) +- hir.load：[01-load.md](01-load.md) +- hir.store：[02-store.md](02-store.md) +- 随路功能详解：[10-padding-quantization.md](10-padding-quantization.md) +- 源码参考： + - [HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) - TableGen 定义 + - [HIVMDMAOps.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp) - 实现代码 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/06-fixpipe.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/06-fixpipe.md new file mode 100644 index 00000000..746e8f5f --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/06-fixpipe.md @@ -0,0 +1,341 @@ +# hir.fixpipe + +> 关键词：Fixpipe、DMA、FIX Pipeline、L0C、量化、ReLU、双目标、NZ2ND + +## 概述 + +`hir.fixpipe` 是 HIVM 方言中最复杂的 DMA 操作，负责将数据从 L0C（Cube 累加器缓存）搬运到其他内存层级（GM/UB/L1），同时支持丰富的随路功能： + +- **布局转换**：NZ 到 ND、NZ 到 DN、NZ 到 NZ（无转换） +- **随路量化**：F32 到 F16、F32 到 I8、F32 到 BF16 等类型转换 +- **随路激活**：ReLU、Leaky ReLU、P-ReLU +- **双目标输出**：将 L0C 数据拆分到两个 UB 缓冲区 + +`hir.fixpipe` 映射到硬件的 FIX Pipeline，属于 Cube 核心类型，是矩阵乘法计算结果写出的核心操作。 + +> Python API 对应：al.fixpipe -- 详见 [docs_triton_ascend/03-Ascend-Extensions/03-fixpipe.md](../../../docs_triton_ascend/03-Ascend-Extensions/03-fixpipe.md) + +## IR 操作定义 + +### TableGen 定义 + +来源：[HIVMDMAOps.td:L246-L326](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td#L246-L326) + +```tablegen +def FixpipeOp : HIVM_DmaOp<"fixpipe", [ + AttrSizedOperandSegments, + SinglePipeOpTrait, OpPipeTrait<"PIPE::PIPE_FIX">, + HIVMCoreTypeInterface, CubeCoreTypeTrait, NoMaxRankTrait, + HIVMUnitFlagEnabledInterface, DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods +]> { + let summary = + "HIVM data copy operation from L0C to other memory hierarchies."; + let description = [{ + Fixpipe is pipeline that performing data movement from L0C to other memory hierarchies, + with on-the-fly fixed function of pre-stage quantization, + pre-stage ReLU, element-wise add, post-stage ReLU, post-stage quantization. + Currently support: + - L0C to OUT + - L0C to L1 + - L0C to UB (for Ascend950 series) + + Additionally, Fixpipe is also capable of layout transform. + + ### Attributes + + #### 'dma_mode' + HIVM data movement model from L0C to destination, There are three values: NZ2DN, NZ2ND, and NZ2NZ(normal). + + #### `dual_dst_mode` + HIVM dual destination mode control. dual destination mode can be enabled only when nz2nd or normal data + movement mode is enabled, and data movement is being performed from L0C to UB. Only supported on Ascend950 series. + }]; + let arguments = (ins AnyShaped:$src, + AnyShaped:$dst, + Arg, [{ + An optional condition to enable unit-flag mode, + useful if there is a dependency on a for loop to run at least once. + }]>:$unit_flag_cond, + DefaultValuedAttr:$dma_mode, + DefaultValuedOptionalAttr:$dual_dst_mode, + DefaultValuedOptionalAttr:$pre_quant, + DefaultValuedOptionalAttr:$pre_relu, + DefaultValuedOptionalAttr:$channel_split, + OptionalAttr:$unit_flag_mode, + Optional:$quant_scale + ); + let results = (outs Optional:$result_tensor); + let assemblyFormat = [{ + attr-dict + `ins` `(` $src `:` type($src) `)` + `outs` `(` $dst `:` type($dst) `)` + (`quant_scale` `=` $quant_scale^ `:` type($quant_scale) )? + (`dual_dst_mode` `=` $dual_dst_mode^)? + (`unit_flag_mode` `(` $unit_flag_mode^ `)` )? + (`unit_flag_cond` `(` $unit_flag_cond^ `)` )? + (`->` type($result_tensor)^)? + }]; + let hasVerifier = 1; + let extraClassDeclaration = DmaOpBaseDecl # [{ + int getFixpipeState(); + int needFixpipePreFuse(); + bool hasStore(); + }]; +} +``` + +### MLIR 语法 + +```mlir +hivm.hir.fixpipe ins(%src : memref<256x128xf16>) + outs(%dst : memref<256x128xf16>) + +hivm.hir.fixpipe {dma_mode = #hivm.dma_mode} + ins(%src : memref<256x128xf16>) + outs(%dst : memref<256x128xf16>) + +hivm.hir.fixpipe {pre_quant = #hivm.fixpipe_pre_quant_mode} + ins(%src : tensor<256x128xf32>) + outs(%dst : tensor<256x128xf16>) + -> tensor<256x128xf16> +``` + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| `src` | AnyShaped | 是 | 源数据缓冲区（L0C） | 通常为 L0C 地址空间 | +| `dst` | AnyShaped | 是 | 目标数据缓冲区 | GM/UB/L1 | +| `unit_flag_cond` | Variadic | 否 | Unit-flag 模式条件 | 用于循环依赖 | +| `quant_scale` | AnyFloat | 否 | 量化缩放因子 | 仅在量化模式下使用 | + +### 结果说明 + +| 结果 | 类型 | 说明 | +|------|------|------| +| `result_tensor` | Optional | Tensor 语义下的结果张量 | + +### 属性说明 + +| 属性 | 类型 | 默认值 | 说明 | +|------|------|--------|------| +| `dma_mode` | HIVM_FixpipeDMAModeAttr | `NZ2NZ` | DMA 搬运模式 | +| `dual_dst_mode` | HIVM_FixpipeDualDstModeAttr | `NO_DUAL` | 双目标模式 | +| `pre_quant` | HIVM_FixpipePreQuantModeAttr | `NO_QUANT` | 随路量化模式 | +| `pre_relu` | HIVM_FixpipePreReluModeAttr | `NO_RELU` | 随路激活模式 | +| `channel_split` | BoolAttr | `false` | 通道拆分标志 | +| `unit_flag_mode` | UnitFlagArrayAttr | 无 | Unit-flag 模式数组 | + +## 枚举详解 + +### FixpipeDMAMode（DMA 搬运模式） + +定义于 [HIVMAttrs.td:L847-L859](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L847-L859) + +| 枚举值 | 数值 | IR 字面量 | 说明 | +|--------|------|----------|------| +| `NZ2ND` | 0 | `nz2nd` | NZ 格式转换为 ND 格式后搬运 | +| `NZ2DN` | 1 | `nz2dn` | NZ 格式转换为 DN 格式后搬运（仅 Ascend950） | +| `NZ2NZ` | 2 | `normal` | 保持 NZ 格式直接搬运（默认） | + +### FixpipeDualDstMode（双目标模式） + +定义于 [HIVMAttrs.td:L821-L845](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L821-L845) + +| 枚举值 | 数值 | IR 字面量 | 说明 | +|--------|------|----------|------| +| `NO_DUAL` | 0 | `NO_DUAL` | 单目标模式（默认） | +| `ROW_SPLIT` | 1 | `ROW_SPLIT` | 按 M 维度拆分，M/2 x N 写入每个 UB，M 须为 2 的倍数 | +| `COLUMN_SPLIT` | 2 | `COLUMN_SPLIT` | 按 N 维度拆分，M x N/2 写入每个 UB，N 须为 32 的倍数 | + +双目标模式的约束： +- 仅在 `dma_mode` 为 NZ2ND 或 NZ2NZ 时可用 +- 仅在数据从 L0C 搬运到 UB 时可用 +- 仅在 Ascend950 系列上支持 + +### FixpipePreQuantMode（随路量化模式） + +定义于 [HIVMAttrs.td:L783-L801](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L783-L801) + +| 枚举值 | 数值 | IR 字面量 | 说明 | +|--------|------|----------|------| +| `NO_QUANT` | 0 | `NO_QUANT` | 不执行量化（默认） | +| `F322F16` | 1 | `F322F16` | F32 转 F16 量化 | +| `S322I8` | 9 | `S322I8` | F32 转 I8 量化 | +| `QF322F32_PRE` | 15 | `QF322F32_PRE` | 带 scale 的 F32 到 F32 预量化 | +| `F322BF16` | 16 | `F322BF16` | F32 转 BF16 量化 | + +### FixpipePreReluMode（随路激活模式） + +定义于 [HIVMAttrs.td:L803-L819](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L803-L819) + +| 枚举值 | 数值 | IR 字面量 | 说明 | +|--------|------|----------|------| +| `NO_RELU` | 0 | `NO_RELU` | 不执行激活（默认） | +| `NORMAL_RELU` | 1 | `NORMAL_RELU` | 标准 ReLU：max(0, x) | +| `LEAKY_RELU` | 2 | `LEAKY_RELU` | Leaky ReLU | +| `P_RELU` | 3 | `P_RELU` | Parametric ReLU | + +## IR 示例 + +### 基础搬运（NZ2NZ 模式） + +```mlir +func.func @test_fixpipe() { + %gmC = memref.alloc() : memref<1024x2048xf16> + %gmCSubview = memref.subview %gmC[0, 0][256, 128][1, 1] + : memref<1024x2048xf16> to + memref<256x128xf16, strided<[2048, 1], offset: 0>> + %l0c = memref.alloc() : memref<256x128xf16> + hivm.hir.fixpipe ins(%l0c : memref<256x128xf16>) + outs(%gmCSubview : memref<256x128xf16, strided<[2048, 1], offset: 0>>) + return +} +``` + +### NZ2ND 布局转换 + +```mlir +hivm.hir.fixpipe {dma_mode = #hivm.dma_mode} + ins(%l0c : memref<256x128xf16>) + outs(%gmCSubview : memref<256x128xf16, strided<[2048, 1], offset: 0>>) +``` + +### NZ2DN 布局转换（仅 Ascend950） + +```mlir +hivm.hir.fixpipe {dma_mode = #hivm.dma_mode} + ins(%l0c : memref<256x128xf16>) + outs(%gmCSubview : memref<256x128xf16, strided<[2048, 1], offset: 0>>) +``` + +### L0C 到 UB 搬运（Ascend950） + +```mlir +module attributes {hacc.target = #hacc.target<"Ascend950PR_9579">} { + func.func @test_fixpipe_l0c_to_ub() { + %alloc = memref.alloc() : memref<1024x2048xf16, #hivm.address_space> + %subview = memref.subview %alloc[0, 0] [256, 128] [1, 1] + : memref<1024x2048xf16, #hivm.address_space> + to memref<256x128xf16, strided<[2048, 1]>, #hivm.address_space> + %alloc_0 = memref.alloc() : memref<256x128xf16, #hivm.address_space> + hivm.hir.fixpipe ins(%alloc_0 : memref<256x128xf16, #hivm.address_space>) + outs(%subview : memref<256x128xf16, strided<[2048, 1]>, #hivm.address_space>) + return + } +} +``` + +### 随路量化（F322F16） + +```mlir +%l0c1 = tensor.empty() : tensor<256x128xf32> +%ret2 = hivm.hir.fixpipe {pre_quant = #hivm.fixpipe_pre_quant_mode} + ins(%l0c1 : tensor<256x128xf32>) + outs(%gmCSubview : tensor<256x128xf16>) + -> tensor<256x128xf16> +``` + +### 随路激活（Leaky ReLU） + +```mlir +%ret3 = hivm.hir.fixpipe {pre_relu = #hivm.fixpipe_pre_relu_mode} + ins(%l0c : tensor<256x128xf16>) + outs(%gmCSubview : tensor<256x128xf16>) + -> tensor<256x128xf16> +``` + +### 双目标模式 + +```mlir +%l0c1 = memref.alloc() : memref<16x16xf16, #hivm.address_space> +%ub = memref.alloc() : memref<16x16xf16, #hivm.address_space> +hivm.hir.fixpipe ins(%l0c1 : memref<16x16xf16, #hivm.address_space>) + outs(%ub : memref<16x16xf16, #hivm.address_space>) + dual_dst_mode = #hivm.fixpipe_dual_dst_mode +``` + +## IR 层约束与验证 + +来源：[HIVMDMAOps.cpp:L848-L891](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L848-L891) + +1. **NZ2DN 模式限制**：NZ2DN 仅在 Ascend950 上支持 +2. **dst=UB 限制**：目标为 UB 仅在 Ascend950 上支持 +3. **双目标模式约束**： + - `dma_mode` 不能为 NZ2DN + - 仅在 Ascend950 上支持 + - 数据搬运必须从 L0C 到 UB +4. **Core Type**：属于 Cube 核心类型（`CubeCoreTypeTrait`） +5. **Pipeline**：固定为 FIX（`OpPipeTrait<"PIPE::PIPE_FIX">`） + +### FixpipeState + +来源：[HIVMDMAOps.cpp:L808-L842](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp#L808-L842) + +Fixpipe 的状态由 `getFixpipeState()` 返回： + +| 状态 | 值 | 含义 | +|------|---|------| +| `Init` | -1 | 初始状态 | +| `QuantOrActivation` | 0 | 需要随路量化或激活 | +| `End` | 1 | 包含存储操作，已完成 | + +### 库函数命名 + +``` +fixpipe_{dma_mode}_{dual?}{src_type}_to_{dst_type}_{src_rank}d_to_{dst_rank}d_{dst_space} +``` + +## 与其他 IR 操作的关系 + +### 在矩阵乘法流水线中的位置 + +``` +hir.nd2nz (GM -> L1) hir.nd2nz (GM -> L1) + | | + v v +hir.mmadL1 (L1 -> L0A/L0B -> L0C) + | + v +hir.fixpipe (L0C -> GM/UB/L1) +``` + +### 与 nz2nd 的关系 + +`hir.fixpipe` 的 NZ2ND 模式与 `hir.nz2nd` 的区别： +- `hir.fixpipe` 从 L0C 搬运，支持随路量化/激活 +- `hir.nz2nd` 从 L1 搬运，仅执行布局转换 + +## 常见问题 + +### Q: 什么时候使用 fixpipe 而不是 nz2nd？ + +A: 当数据源是 L0C（矩阵乘法结果）时，应使用 `hir.fixpipe`，因为它可以利用随路量化和激活功能，减少额外的计算操作。当数据源是 L1 中的 NZ 格式数据时，使用 `hir.nz2nd`。 + +### Q: pre_quant 和 pre_relu 可以同时使用吗？ + +A: 可以。Fixpipe 支持同时执行随路量化和随路激活，这是其核心优势之一。 + +### Q: dual_dst_mode 的实际用途是什么？ + +A: 双目标模式用于将 L0C 中的矩阵拆分到两个 UB 缓冲区，常用于混合核心（Mix）模式下的 Cube-Vector 协同计算。ROW_SPLIT 按行拆分，COLUMN_SPLIT 按列拆分。 + +### Q: quant_scale 的作用是什么？ + +A: `quant_scale` 用于 `QF322F32_PRE` 量化模式，提供量化缩放因子。在其他量化模式下不需要此参数。 + +## 相关文档 + +- Python API：[docs_triton_ascend/03-Ascend-Extensions/03-fixpipe.md](../../../docs_triton_ascend/03-Ascend-Extensions/03-fixpipe.md) +- DMA 操作总览：[00-overview.md](00-overview.md) +- 随路功能详解：[10-padding-quantization.md](10-padding-quantization.md) +- 源码参考： + - [HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) - TableGen 定义 + - [HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) - 枚举定义 + - [HIVMDMAOps.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/HIVMDMAOps.cpp) - 验证逻辑实现 + - [dma-ops.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/dma-ops.mlir) - IR 测试用例 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/07-atomic.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/07-atomic.md new file mode 100644 index 00000000..b05c8cbf --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/07-atomic.md @@ -0,0 +1,277 @@ +# hir.atomic_cas / hir.atomic_xchg + +> 关键词：Atomic、CAS、Compare-And-Swap、XCHG、Exchange、原子操作 + +## 概述 + +HIVM 方言提供了两个独立的原子操作：`hir.atomic_cas`（原子比较并交换）和 `hir.atomic_xchg`（原子交换）。这两个操作用于在全局内存（GM）上执行不可分割的读-修改-写操作，是实现并发同步原语的基础。 + +与 `hir.store` 的原子模式不同，`atomic_cas` 和 `atomic_xchg` 是独立的操作，不依赖于 `hir.set_atomic` 的全局原子模式设置。 + +> Python API 对应：tl.atomic_cas / tl.atomic_xchg -- 详见 [docs_triton_ascend/02-Core-API/01-memory-ops.md](../../../docs_triton_ascend/02-Core-API/01-memory-ops.md) + +## hir.atomic_cas + +### 概述 + +原子比较并交换（Compare-And-Swap, CAS）操作。该操作读取内存位置 V 的值，如果等于期望值 A，则将其更新为新值 B，并返回 V 的原始值。整个过程是原子的，不会被其他线程中断。 + +### TableGen 定义 + +来源：[HIVMDMAOps.td:L406-L443](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td#L406-L443) + +```tablegen +def AtomicCasOp : HIVM_Op<"atomic_cas", + [NoLibraryFunctionTrait, + OperElemTypeConstraints<[0], [I8, I16, I32, F16, F32, I64]> + ]> { + let summary = "Atomic Compare-And-Swap (CAS) Op"; + let description = [{ + Compare-And-Swap (CAS) is an atomic operation that consists of three operands: + Memory location (V), Expected old value (A), New value (B). + The semantics of the operation are: the value of V is updated to B, + only if the value of memory location V is equal to the expected old value A. + The operation returns the original value of V regardless of whether it is updated or not. + + Constraints: + 1. The input memref and output memref must have the same rank + and the same element type. + + Arguments: + * `src0`: expected old value + * `src1`: new value + * `dst`: memory location in GM + + Examples: + ```mlir + hivm.hir.atomic_cas ins(%src0, %src1 : memref, memref) outs(%dst : memref) + %result = hivm.hir.atomic_cas ins(%src0, %src1 : tensor, tensor) outs(%dst : tensor) -> tensor + ``` + }]; + let arguments = (ins Variadic:$src, + TensorOrMemref:$dst + ); + let results = (outs Optional:$result_tensor); + let assemblyFormat = [{ + attr-dict + `ins` `(` $src `:` type($src) `)` + `outs` `(` $dst `:` type($dst) `)` + (`->` type($result_tensor)^)? + }]; +} +``` + +### MLIR 语法 + +```mlir +hivm.hir.atomic_cas ins(%src0, %src1 : memref, memref) + outs(%dst : memref) + +%result = hivm.hir.atomic_cas ins(%src0, %src1 : tensor, tensor) + outs(%dst : tensor) -> tensor +``` + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| `src` | Variadic | 是 | 源操作数（2 个） | src0=期望旧值, src1=新值 | +| `dst` | TensorOrMemref | 是 | 目标内存位置 | GM 地址空间 | + +### 结果说明 + +| 结果 | 类型 | 说明 | +|------|------|------| +| `result_tensor` | Optional | 返回 V 的原始值（Tensor 语义） | + +### 数据类型约束 + +来源：`OperElemTypeConstraints<[0], [...]>` + +| 支持的元素类型 | 说明 | +|---------------|------| +| `i8` | 8 位有符号整数 | +| `i16` | 16 位有符号整数 | +| `i32` | 32 位有符号整数 | +| `i64` | 64 位有符号整数 | +| `f16` | 半精度浮点 | +| `f32` | 单精度浮点 | + +### IR 示例 + +#### Memref 语义 + +```mlir +func.func @test_atomic_cas_memref() { + %src0 = memref.alloc() : memref<16xf32, #hivm.address_space> + %src1 = memref.alloc() : memref<16xf32, #hivm.address_space> + %dst = memref.alloc() : memref<16xf32, #hivm.address_space> + hivm.hir.atomic_cas ins(%src0, %src1 : memref<16xf32, #hivm.address_space>, memref<16xf32, #hivm.address_space>) + outs(%dst : memref<16xf32, #hivm.address_space>) + return +} +``` + +#### Tensor 语义 + +```mlir +func.func @test_atomic_cas_tensor() -> tensor<16xi32> { + %src0 = tensor.empty() : tensor<16xi32> + %src1 = tensor.empty() : tensor<16xi32> + %dst = tensor.empty() : tensor<16xi32> + %result = hivm.hir.atomic_cas ins(%src0, %src1 : tensor<16xi32>, tensor<16xi32>) + outs(%dst : tensor<16xi32>) -> tensor<16xi32> + return %result : tensor<16xi32> +} +``` + +## hir.atomic_xchg + +### 概述 + +原子交换（Exchange）操作。该操作读取内存位置的当前值，写入新值，并返回旧值。整个过程是原子的。 + +### TableGen 定义 + +来源：[HIVMDMAOps.td:L445-L481](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td#L445-L481) + +```tablegen +def AtomicXchgOp : HIVM_Op<"atomic_xchg", + [NoLibraryFunctionTrait, + OperElemTypeConstraints<[0], [I8, I16, I32, F16, F32, I64]> + ]> { + let summary = "Atomic Exchange Op"; + let description = [{ + Atomic exchange is an atomic operation that consists of three steps: + 1. Read the current value of the specified memory address + 2. Write the new value to the memory address + 3. Return the old value read previously + The whole process is atomic, that is, it will not be interrupted by other threads during the operation. + + Constraints: + 1. The input memref and output memref must have the same rank + and the same element type. + + Arguments: + * `src`: new value + * `dst`: memory location in GM + + Examples: + ```mlir + hivm.hir.atomic_xchg ins(%src : memref) outs(%dst : memref) + %result = hivm.hir.atomic_xchg ins(%src : tensor) outs(%dst : tensor) -> tensor + ``` + }]; + let arguments = (ins Variadic:$src, + TensorOrMemref:$dst + ); + let results = (outs Optional:$result_tensor); + let assemblyFormat = [{ + attr-dict + `ins` `(` $src `:` type($src) `)` + `outs` `(` $dst `:` type($dst) `)` + (`->` type($result_tensor)^)? + }]; +} +``` + +### MLIR 语法 + +```mlir +hivm.hir.atomic_xchg ins(%src : memref) + outs(%dst : memref) + +%result = hivm.hir.atomic_xchg ins(%src : tensor) + outs(%dst : tensor) -> tensor +``` + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| `src` | Variadic | 是 | 源操作数（1 个） | 新值 | +| `dst` | TensorOrMemref | 是 | 目标内存位置 | GM 地址空间 | + +### 结果说明 + +| 结果 | 类型 | 说明 | +|------|------|------| +| `result_tensor` | Optional | 返回 V 的原始值（Tensor 语义） | + +### 数据类型约束 + +与 `atomic_cas` 相同：`i8`, `i16`, `i32`, `i64`, `f16`, `f32` + +### IR 示例 + +```mlir +func.func @test_atomic_xchg_memref() { + %src = memref.alloc() : memref<16xf32, #hivm.address_space> + %dst = memref.alloc() : memref<16xf32, #hivm.address_space> + hivm.hir.atomic_xchg ins(%src : memref<16xf32, #hivm.address_space>) + outs(%dst : memref<16xf32, #hivm.address_space>) + return +} +``` + +## IR 层约束与验证 + +### atomic_cas 约束 + +1. **操作数数量**：`src` 必须包含恰好 2 个操作数（期望旧值和新值） +2. **元素类型一致性**：输入和输出的元素类型必须相同 +3. **Rank 一致性**：输入和输出的 rank 必须相同 +4. **NoLibraryFunctionTrait**：不生成库函数调用，直接映射到硬件指令 + +### atomic_xchg 约束 + +1. **操作数数量**：`src` 必须包含恰好 1 个操作数（新值） +2. **元素类型一致性**：输入和输出的元素类型必须相同 +3. **Rank 一致性**：输入和输出的 rank 必须相同 +4. **NoLibraryFunctionTrait**：不生成库函数调用，直接映射到硬件指令 + +## 与 store 原子操作的区别 + +| 特性 | `hir.store` + atomic | `hir.atomic_cas` | `hir.atomic_xchg` | +|------|---------------------|-------------------|-------------------| +| 操作类型 | ADD/MAX/MIN/AND/OR/XOR | CAS | XCHG | +| 依赖 set_atomic | 是 | 否 | 否 | +| 返回旧值 | 否 | 是 | 是 | +| 实现方式 | 硬件/软件 | 硬件指令 | 硬件指令 | +| 数据类型 | I8~F8E5M2 | I8/I16/I32/I64/F16/F32 | I8/I16/I32/I64/F16/F32 | + +## 与其他 IR 操作的关系 + +### 从 Triton 到 HIVM + +``` +tt.atomic_cas --> hivm.hir.atomic_cas +tt.atomic_xchg --> hivm.hir.atomic_xchg +``` + +### 与 set_atomic 的关系 + +`hir.atomic_cas` 和 `hir.atomic_xchg` 是独立的原子操作，不需要通过 `hir.set_atomic` 设置全局原子模式。它们直接映射到硬件的原子指令。 + +## 常见问题 + +### Q: atomic_cas 和 store 的 CAS 原子模式有什么区别？ + +A: `hir.store` 的 CAS 模式需要先通过 `hir.set_atomic` 设置全局原子模式，且 CAS 属于软件原子实现。`hir.atomic_cas` 是独立的操作，直接映射到硬件指令，性能更高。 + +### Q: atomic_cas 的返回值有什么用？ + +A: `atomic_cas` 返回内存位置的原始值，可以用于判断 CAS 是否成功（如果返回值等于期望旧值，说明更新成功；否则说明有并发修改）。 + +### Q: 为什么 atomic_cas/xchg 不支持 BF16 和 F8 类型？ + +A: 这是硬件约束。原子操作的数据类型支持由硬件决定，当前 Ascend NPU 的原子指令不支持 BF16 和 F8 类型。 + +## 相关文档 + +- Python API：[docs_triton_ascend/02-Core-API/01-memory-ops.md](../../../docs_triton_ascend/02-Core-API/01-memory-ops.md) +- DMA 操作总览：[00-overview.md](00-overview.md) +- hir.store：[02-store.md](02-store.md) +- 源码参考： + - [HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) - TableGen 定义 + - [HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) - AtomicKind 枚举 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/08-gather-scatter.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/08-gather-scatter.md new file mode 100644 index 00000000..483610e2 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/08-gather-scatter.md @@ -0,0 +1,296 @@ +# hir.gather_load / hir.scatter_store + +> 关键词：Gather、Scatter、稀疏内存访问、间接索引、Mask + +## 概述 + +`hir.gather_load` 和 `hir.scatter_store` 是 HIVM 方言中的稀疏内存访问操作，用于按索引张量从全局内存中收集数据或向全局内存散列写数据。这两个操作支持掩码（Mask）机制，可以条件化地控制哪些元素被加载或存储。 + +与 `hir.load`/`hir.store` 的连续内存访问不同，Gather/Scatter 操作支持非连续的、按索引的内存访问模式，适用于稀疏数据、嵌入查找等场景。 + +> Python API 对应：tl.load (带 index) / tl.store (带 index) -- 详见 [docs_triton_ascend/02-Core-API/01-memory-ops.md](../../../docs_triton_ascend/02-Core-API/01-memory-ops.md) + +## hir.gather_load + +### 概述 + +稀疏内存加载操作。从源内存缓冲区中按索引张量指定的偏移位置收集元素，生成输出张量。支持掩码和回退值。 + +### TableGen 定义 + +来源：[HIVMOps.td:L389-L426](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMOps.td#L389-L426) + +```tablegen +def GatherLoadOp : HIVM_Op<"gather_load", [ + DeclareOpInterfaceMethods, + SameVariadicOperandSize, +]> { + let summary = [{ + Perform sparse memory loading with optional mask and other + }]; + + let description = [{ + This operation takes a source memory buffer and a tensor of offsets, + and produces an output tensor by gathering elements from the source + at the specified offset positions. The operation supports masking + and provides fallback values for masked-out positions. + }]; + let arguments = (ins AnyMemRef:$base, + RankedTensorOf<[I32, I64]>:$indices, + AnyTypeOf<[I32, I64]>:$burst_len, + Optional>:$mask, + Optional>:$other, + + OptionalAttr:$boundaryCheck, + OptionalAttr:$padding, + OptionalAttr:$cache, + OptionalAttr:$evict, + OptionalAttr:$isVolatile + ); + let results = (outs AnyRankedTensor:$result); + + let assemblyFormat = [{ + `ins` `(` $base `:` type($base) `,` $indices `:` type($indices) `,` + $burst_len `:` type($burst_len) + (`,` $mask `:` type($mask)^)? (`,` $other `:` type($other)^)? `)` + attr-dict + `->` type($result) + }]; + + let hasVerifier = 1; +} +``` + +### MLIR 语法 + +```mlir +%result = hivm.hir.gather_load ins(%base : memref, %indices : tensor<16xi32>, %burst_len : i32) + -> tensor<16xf32> + +%result = hivm.hir.gather_load ins(%base : memref, %indices : tensor<16xi32>, %burst_len : i32, %mask : tensor<16xi1>, %other : f32) + -> tensor<16xf32> +``` + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| `base` | AnyMemRef | 是 | 源内存缓冲区基地址 | GM 地址空间 | +| `indices` | RankedTensorOf<[I32, I64]> | 是 | 偏移索引张量 | I32 或 I64 | +| `burst_len` | AnyTypeOf<[I32, I64]> | 是 | 突发传输长度 | I32 或 I64 | +| `mask` | RankedTensorOf<[I1]> | 否 | 掩码张量 | 布尔类型 | +| `other` | AnyTypeOf<[AnyInteger, AnyFloat]> | 否 | 掩码位置的回退值 | 与结果元素类型相同 | + +### 结果说明 + +| 结果 | 类型 | 说明 | +|------|------|------| +| `result` | AnyRankedTensor | 收集到的输出张量 | + +### 属性说明 + +| 属性 | 类型 | 默认值 | 说明 | 可选值 | +|------|------|--------|------|--------| +| `boundaryCheck` | DenseI32ArrayAttr | 无 | 边界检查维度 | 维度索引数组 | +| `padding` | HIVM_PaddingOptionAttr | 无 | 填充选项 | `PAD_ZERO`(1) / `PAD_NAN`(2) | +| `cache` | HIVM_CacheModifierAttr | 无 | 缓存修改器 | NONE/CA/CG/WB/CS/WT/CV | +| `evict` | HIVM_EvictionPolicyAttr | 无 | 驱逐策略 | EvictFirst / EvictLast | +| `isVolatile` | BoolAttr | 无 | 是否为 volatile 访问 | true / false | + +### PaddingOption 枚举 + +定义于 [HIVMAttrs.td:L920-L928](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L920-L928) + +| 枚举值 | 数值 | IR 字面量 | 说明 | +|--------|------|----------|------| +| `PAD_ZERO` | 1 | `zero` | 用零填充 | +| `PAD_NAN` | 2 | `nan` | 用 NaN 填充 | + +### CacheModifier 枚举 + +定义于 [HIVMAttrs.td:L930-L942](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L930-L942) + +| 枚举值 | 数值 | IR 字面量 | 说明 | +|--------|------|----------|------| +| `NONE` | 1 | `none` | 无缓存修改 | +| `CA` | 2 | `ca` | Cache All | +| `CG` | 3 | `cg` | Cache Global | +| `WB` | 4 | `wb` | Write Back | +| `CS` | 5 | `cs` | Cache Streaming | +| `WT` | 6 | `wt` | Write Through | +| `CV` | 7 | `cv` | Cache Volatile | + +### IR 示例 + +#### 基础 Gather Load + +```mlir +func.func @test_gather_load(%base: memref, %indices: tensor<16xi32>, %burst_len: i32) -> tensor<16xf32> { + %result = hivm.hir.gather_load ins(%base : memref, %indices : tensor<16xi32>, %burst_len : i32) + -> tensor<16xf32> + return %result : tensor<16xf32> +} +``` + +#### 带 Mask 和 Other 的 Gather Load + +```mlir +func.func @test_gather_load_masked(%base: memref, %indices: tensor<16xi32>, %burst_len: i32, %mask: tensor<16xi1>, %other: f32) -> tensor<16xf32> { + %result = hivm.hir.gather_load ins(%base : memref, %indices : tensor<16xi32>, %burst_len : i32, %mask : tensor<16xi1>, %other : f32) + -> tensor<16xf32> + return %result : tensor<16xf32> +} +``` + +## hir.scatter_store + +### 概述 + +稀疏内存存储操作。将源张量中的元素按索引张量指定的偏移位置写入目标 GM 缓冲区。支持掩码机制。 + +### TableGen 定义 + +来源：[HIVMOps.td:L459-L492](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMOps.td#L459-L492) + +```tablegen +def ScatterStoreOp : HIVM_Op<"scatter_store", [ + DeclareOpInterfaceMethods, + SameVariadicOperandSize, +]> { + let summary = [{ + Performs sparse memory storing with optional mask + }]; + + let description = [{ + This operation takes a source tensor and a tensor of offsets from UB, + and stores elements from the source into the destination GM buffer + at the specified offset positions. The operation + supports masking to conditionally control which elements are stored. + }]; + let arguments = (ins AnyMemRef:$base, + RankedTensorOf<[I32, I64]>:$indices, + AnyRankedTensor:$data, + AnyTypeOf<[I32, I64]>:$burst_len, + Optional>:$mask, + + OptionalAttr:$boundaryCheck, + OptionalAttr:$cache, + OptionalAttr:$evict + ); + + let assemblyFormat = [{ + `ins` `(` $base `:` type($base) `,` $indices `:` type($indices) `,` + $data `:` type($data) `,` $burst_len `:` type($burst_len) + (`,` $mask `:` type($mask)^)? `)` + attr-dict + }]; + + let hasVerifier = 1; +} +``` + +### MLIR 语法 + +```mlir +hivm.hir.scatter_store ins(%base : memref, %indices : tensor<16xi32>, %data : tensor<16xf32>, %burst_len : i32) + +hivm.hir.scatter_store ins(%base : memref, %indices : tensor<16xi32>, %data : tensor<16xf32>, %burst_len : i32, %mask : tensor<16xi1>) +``` + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| `base` | AnyMemRef | 是 | 目标内存缓冲区基地址 | GM 地址空间 | +| `indices` | RankedTensorOf<[I32, I64]> | 是 | 偏移索引张量 | I32 或 I64 | +| `data` | AnyRankedTensor | 是 | 源数据张量 | UB 中的数据 | +| `burst_len` | AnyTypeOf<[I32, I64]> | 是 | 突发传输长度 | I32 或 I64 | +| `mask` | RankedTensorOf<[I1]> | 否 | 掩码张量 | 布尔类型 | + +### 属性说明 + +| 属性 | 类型 | 默认值 | 说明 | 可选值 | +|------|------|--------|------|--------| +| `boundaryCheck` | DenseI32ArrayAttr | 无 | 边界检查维度 | 维度索引数组 | +| `cache` | HIVM_CacheModifierAttr | 无 | 缓存修改器 | 同 gather_load | +| `evict` | HIVM_EvictionPolicyAttr | 无 | 驱逐策略 | EvictFirst / EvictLast | + +### IR 示例 + +#### 基础 Scatter Store + +```mlir +func.func @test_scatter_store(%base: memref, %indices: tensor<16xi32>, %data: tensor<16xf32>, %burst_len: i32) { + hivm.hir.scatter_store ins(%base : memref, %indices : tensor<16xi32>, %data : tensor<16xf32>, %burst_len : i32) + return +} +``` + +#### 带 Mask 的 Scatter Store + +```mlir +func.func @test_scatter_store_masked(%base: memref, %indices: tensor<16xi32>, %data: tensor<16xf32>, %burst_len: i32, %mask: tensor<16xi1>) { + hivm.hir.scatter_store ins(%base : memref, %indices : tensor<16xi32>, %data : tensor<16xf32>, %burst_len : i32, %mask : tensor<16xi1>) + return +} +``` + +## IR 层约束与验证 + +### gather_load 约束 + +1. **索引类型**：`indices` 必须为 I32 或 I64 类型的张量 +2. **掩码类型**：`mask` 必须为 I1 类型的张量 +3. **回退值类型**：`other` 的类型必须与输出张量的元素类型相同 +4. **SameVariadicOperandSize**：可变操作数的大小必须一致 + +### scatter_store 约束 + +1. **索引类型**：`indices` 必须为 I32 或 I64 类型的张量 +2. **掩码类型**：`mask` 必须为 I1 类型的张量 +3. **SameVariadicOperandSize**：可变操作数的大小必须一致 + +## 与其他 IR 操作的关系 + +### Gather/Scatter vs Load/Store + +| 特性 | hir.load / hir.store | hir.gather_load / hir.scatter_store | +|------|---------------------|-------------------------------------| +| 访问模式 | 连续 | 按索引非连续 | +| 索引方式 | 隐式（连续偏移） | 显式（索引张量） | +| 掩码支持 | 通过 Padding | 通过 Mask 张量 | +| Pipeline | MTE2/MTE3 | 无固定 Pipeline | +| 适用场景 | 密集张量操作 | 稀疏数据、嵌入查找 | + +### 与 indirect_load/indirect_store 的区别 + +| 特性 | gather_load / scatter_store | indirect_load / indirect_store | +|------|---------------------------|-------------------------------| +| 数据通路 | GM <-> UB | GM <-> UB | +| 索引方式 | 偏移量 | 偏移量 | +| Pipeline | 无固定 | PIPE_V | +| 接口风格 | ins/outs 分离 | ins/outs 统一 | +| 掩码 | 可选 mask 张量 | 可选 mask 张量 | + +## 常见问题 + +### Q: gather_load 和 indirect_load 有什么区别？ + +A: `gather_load` 是更通用的稀疏加载操作，支持缓存修改器和边界检查等属性。`indirect_load` 是更结构化的间接加载操作，使用 Destination Style 接口，属于 Vector Pipeline。在编译流程中，`gather_load` 可能被分解为 `indirect_load`。 + +### Q: burst_len 参数的作用是什么？ + +A: `burst_len` 指定每次内存访问的突发传输长度，影响 DMA 引擎的传输效率。较大的突发长度可以提高带宽利用率，但需要更多的缓冲区空间。 + +### Q: 什么时候应该使用 gather/scatter 而不是 load/store？ + +A: 当数据访问模式是非连续的（如按索引访问数组元素、嵌入表查找）时，应使用 gather/scatter。当数据访问是连续的时，使用 load/store 更高效。 + +## 相关文档 + +- Python API：[docs_triton_ascend/02-Core-API/01-memory-ops.md](../../../docs_triton_ascend/02-Core-API/01-memory-ops.md) +- DMA 操作总览：[00-overview.md](00-overview.md) +- 间接访问：[09-indirect-access.md](09-indirect-access.md) +- 源码参考： + - [HIVMOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMOps.td) - TableGen 定义 + - [HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) - 属性枚举 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/09-indirect-access.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/09-indirect-access.md new file mode 100644 index 00000000..bd01981f --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/09-indirect-access.md @@ -0,0 +1,332 @@ +# hir.indirect_load / hir.indirect_store + +> 关键词：Indirect、间接访问、SIMT、Vector Pipeline、Mask + +## 概述 + +`hir.indirect_load` 和 `hir.indirect_store` 是 HIVM 方言中的间接内存访问操作，用于按偏移张量从全局内存（GM）中加载数据或向 GM 存储数据。这两个操作使用 SIMT（Single Instruction Multiple Thread）模板执行，归属于 Vector Pipeline（PIPE_V），支持 1D 到 5D 的数据访问。 + +与 `hir.gather_load`/`hir.scatter_store` 类似，`indirect_load`/`indirect_store` 也支持非连续的按索引内存访问，但它们使用 Destination Style 接口，具有更结构化的操作语义。 + +> Python API 对应：无直接对应，由编译器在 lowering 阶段自动生成 + +## hir.indirect_load + +### 概述 + +间接内存加载操作。从源内存缓冲区中按偏移张量指定的位置收集元素到目标张量，支持掩码和回退值。 + +### TableGen 定义 + +来源：[HIVMOps.td:L524-L592](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMOps.td#L524-L592) + +```tablegen +def IndirectLoadOp : HIVM_Op<"indirect_load",[ + StaticMaxRankTrait<5>, + DestinationStyleOpInterface, + OpPipeInterface, + SinglePipeOpTrait, OpPipeTrait<"PIPE::PIPE_V">, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + AttrSizedOperandSegments + ]> { + let summary = [{ + Performs indirect memory loading with masking and fallback values. + }]; + + let description = [{ + This operation takes a source memory buffer and a tensor of offsets, + and produces an output tensor by gathering elements from the source + at the specified offset positions. The operation supports masking + and provides fallback values for masked-out positions. + This operation supports 1D-5D. + + For each position in the output tensor: + 1D: + dst[i] = mask[i] ? src[offsets[i]] : other[i] + 2D: + dst[i][j] = mask[i][j] ? src[offsets[i][j]] : other[i][j] + 3D: + dst[i][j][k] = mask[i][j][k] ? src[offsets[i][j][k]] : other[i][j][k] + 4D: + dst[i][j][k][l] = mask[i][j][k][l] ? src[offsets[i][j][k][l]] : other[i][j][k][l] + 5D: + dst[i][j][k][l][m] = mask[i][j][k][l][m] ? src[offsets[i][j][k][l][m]] : other[i][j][k][l][m] + + Where: + - src: source memory buffer to load from + - offsets: indices specifying positions in the source buffer + - mask: boolean mask controlling which elements to load + - other: fallback values used when mask is false + - dst: destination tensor specifying output shape and type + + This operation is useful for sparse data access patterns and + gather operations with conditional loading semantics. + }]; + + let arguments = (ins AnyMemRef:$src, + TensorOrMemref:$offsets, + TensorOrMemref:$dst, + Optional:$mask, + Optional:$other + ); + let results = (outs Optional:$result); + + let assemblyFormat = [{ + `ins` `(` $src `:` type($src) `,` $offsets `:` type($offsets) + (`,` $mask^ `:` type($mask))? + (`,` $other^ `:` type($other))? `)` + `outs` `(` $dst `:` type($dst) `)` + attr-dict + (`->` type($result)^)? + }]; + + let extraClassDeclaration = [{ + static StringRef getOpName() { return "indirect_load"; } + ::mlir::MutableOperandRange getDpsInitsMutable() { + return getDstMutable(); + } + }]; + + let hasVerifier = 1; +} +``` + +### MLIR 语法 + +```mlir +hivm.hir.indirect_load ins(%src : memref, %offsets : tensor<16xi32>) + outs(%dst : tensor<16xf32>) + +hivm.hir.indirect_load ins(%src : memref, %offsets : tensor<16xi32>, %mask : tensor<16xi1>, %other : tensor<16xf32>) + outs(%dst : tensor<16xf32>) -> tensor<16xf32> +``` + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| `src` | AnyMemRef | 是 | 源内存缓冲区 | GM 地址空间 | +| `offsets` | TensorOrMemref | 是 | 偏移索引张量 | 指定加载位置 | +| `dst` | TensorOrMemref | 是 | 目标张量 | 指定输出形状和类型 | +| `mask` | TensorOrMemref | 否 | 布尔掩码张量 | 控制哪些位置加载 | +| `other` | TensorOrMemref | 否 | 掩码位置的回退值 | 与 dst 元素类型相同 | + +### 结果说明 + +| 结果 | 类型 | 说明 | +|------|------|------| +| `result` | Optional | Tensor 语义下的结果张量 | + +### 语义说明 + +``` +dst[i] = mask[i] ? src[offsets[i]] : other[i] +``` + +当 mask 为 true 时，从 src 的 offsets[i] 位置加载数据；当 mask 为 false 时，使用 other[i] 作为回退值。 + +## hir.indirect_store + +### 概述 + +间接内存存储操作。将源张量中的元素按偏移张量指定的位置写入目标 GM 缓冲区，支持掩码。 + +### TableGen 定义 + +来源：[HIVMOps.td:L598-L660](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMOps.td#L598-L660) + +```tablegen +def IndirectStoreOp : HIVM_Op<"indirect_store",[ + StaticMaxRankTrait<5>, + DestinationStyleOpInterface, + OpPipeInterface, + SinglePipeOpTrait, OpPipeTrait<"PIPE::PIPE_V">, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + ]> { + let summary = [{ + Performs indirect memory storing with SIMT template. + }]; + + let description = [{ + This operation takes a source tensor and a tensor of offsets from UB, + and stores elements from the source into the destination GM buffer + at the specified offset positions with SIMT template. The operation + supports masking to conditionally control which elements are stored. + This operation supports 1D-5D. + + For each position in the source tensor: + 1D: + if (mask[i]) dst[offsets[i]] = src[i] + 2D: + if (mask[i][j]) dst[offsets[i][j]] = src[i][j] + 3D: + if (mask[i][j][k]) dst[offsets[i][j][k]] = src[i][j][k] + 4D: + if (mask[i][j][k][l]) dst[offsets[i][j][k][l]] = src[i][j][k][l] + 5D: + if (mask[i][j][k][l][m]) dst[offsets[i][j][k][l][m]] = src[i][j][k][l][m] + + Where: + - src: source tensor containing values to store from UB + - offsets: indices specifying positions in the destination buffer + - dst: destination memory buffer to store into GM + - mask: optional boolean mask controlling which elements to store + + When no mask is provided, all elements from the source tensor are stored + to the corresponding offset positions in the destination buffer. + }]; + + let arguments = (ins AnyMemRef:$dst, + TensorOrMemref:$offsets, + TensorOrMemref:$src, + Optional:$mask + ); + + let assemblyFormat = [{ + `ins` `(` $src `:` type($src) `,` $offsets `:` type($offsets) + (`,` $mask^ `:` type($mask))? `)` + `outs` `(` $dst `:` type($dst) `)` + attr-dict + }]; + + let extraClassDeclaration = [{ + static StringRef getOpName() { return "indirect_store"; } + ::mlir::MutableOperandRange getDpsInitsMutable() { + return getDstMutable(); + } + }]; + + let hasVerifier = 1; +} +``` + +### MLIR 语法 + +```mlir +hivm.hir.indirect_store ins(%src : tensor<16xf32>, %offsets : tensor<16xi32>) + outs(%dst : memref) + +hivm.hir.indirect_store ins(%src : tensor<16xf32>, %offsets : tensor<16xi32>, %mask : tensor<16xi1>) + outs(%dst : memref) +``` + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| `dst` | AnyMemRef | 是 | 目标内存缓冲区 | GM 地址空间 | +| `offsets` | TensorOrMemref | 是 | 偏移索引张量 | 指定存储位置 | +| `src` | TensorOrMemref | 是 | 源数据张量 | UB 中的数据 | +| `mask` | TensorOrMemref | 否 | 布尔掩码张量 | 控制哪些位置存储 | + +### 语义说明 + +``` +if (mask[i]) dst[offsets[i]] = src[i] +``` + +当 mask 为 true 时，将 src[i] 存储到 dst 的 offsets[i] 位置。当无 mask 时，所有元素都被存储。 + +## IR 示例 + +### indirect_load 基础用法 + +```mlir +func.func @test_indirect_load(%src: memref, %offsets: tensor<16xi32>) -> tensor<16xf32> { + %dst = tensor.empty() : tensor<16xf32> + %result = hivm.hir.indirect_load ins(%src : memref, %offsets : tensor<16xi32>) + outs(%dst : tensor<16xf32>) -> tensor<16xf32> + return %result : tensor<16xf32> +} +``` + +### indirect_load 带 Mask 和 Other + +```mlir +func.func @test_indirect_load_masked(%src: memref, %offsets: tensor<16xi32>, %mask: tensor<16xi1>, %other: tensor<16xf32>) -> tensor<16xf32> { + %dst = tensor.empty() : tensor<16xf32> + %result = hivm.hir.indirect_load ins(%src : memref, %offsets : tensor<16xi32>, %mask : tensor<16xi1>, %other : tensor<16xf32>) + outs(%dst : tensor<16xf32>) -> tensor<16xf32> + return %result : tensor<16xf32> +} +``` + +### indirect_store 基础用法 + +```mlir +func.func @test_indirect_store(%dst: memref, %offsets: tensor<16xi32>, %src: tensor<16xf32>) { + hivm.hir.indirect_store ins(%src : tensor<16xf32>, %offsets : tensor<16xi32>) + outs(%dst : memref) + return +} +``` + +### indirect_store 带 Mask + +```mlir +func.func @test_indirect_store_masked(%dst: memref, %offsets: tensor<16xi32>, %src: tensor<16xf32>, %mask: tensor<16xi1>) { + hivm.hir.indirect_store ins(%src : tensor<16xf32>, %offsets : tensor<16xi32>, %mask : tensor<16xi1>) + outs(%dst : memref) + return +} +``` + +## IR 层约束与验证 + +### indirect_load 约束 + +1. **Rank 限制**：最大 rank 为 5（`StaticMaxRankTrait<5>`） +2. **Pipeline**：固定为 Vector Pipeline（`OpPipeTrait<"PIPE::PIPE_V">`） +3. **Destination Style**：实现 `DestinationStyleOpInterface` +4. **AttrSizedOperandSegments**：可变长度操作数段的大小由属性控制 +5. **元素类型一致性**：`src`、`dst`、`other` 的元素类型必须一致 + +### indirect_store 约束 + +1. **Rank 限制**：最大 rank 为 5（`StaticMaxRankTrait<5>`） +2. **Pipeline**：固定为 Vector Pipeline（`OpPipeTrait<"PIPE::PIPE_V">`） +3. **Destination Style**：实现 `DestinationStyleOpInterface` +4. **元素类型一致性**：`src` 和 `dst` 的元素类型必须一致 + +## 与其他 IR 操作的关系 + +### indirect_load/store vs gather_load/scatter_store + +| 特性 | indirect_load/store | gather_load/scatter_store | +|------|-------------------|--------------------------| +| Pipeline | PIPE_V | 无固定 | +| 接口风格 | Destination Style | 非 Destination Style | +| Rank 支持 | 1D-5D | 无限制 | +| 库函数 | 有（LibraryFunctionOpInterface） | 无 | +| 缓存属性 | 不支持 | 支持 cache/evict | +| 边界检查 | 不支持 | 支持 boundaryCheck | + +### 典型转换关系 + +``` +hir.gather_load --> hir.indirect_load (分解后) +hir.gather_load --> hir.custom name="__builtin_gather_load" (内置 Custom Op) +``` + +## 常见问题 + +### Q: indirect_load 和 gather_load 应该使用哪个？ + +A: `indirect_load` 是更结构化的操作，属于 Vector Pipeline，有对应的库函数实现。`gather_load` 是更通用的操作，支持缓存属性和边界检查。在编译流程中，`gather_load` 可能被分解为 `indirect_load` 或 Custom Op。 + +### Q: indirect_load/store 为什么属于 Vector Pipeline？ + +A: 因为间接内存访问使用 SIMT（Single Instruction Multiple Thread）模板执行，每个线程独立处理一个索引位置，这是 Vector 核心的执行模式。 + +### Q: 5D 限制是否足够？ + +A: 对于大多数深度学习场景，5D 已经足够覆盖常见的张量维度（batch, seq_len, head, height, width）。如果需要更高维度的间接访问，可能需要先 reshape 数据。 + +## 相关文档 + +- DMA 操作总览：[00-overview.md](00-overview.md) +- Gather/Scatter：[08-gather-scatter.md](08-gather-scatter.md) +- 源码参考： + - [HIVMOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMOps.td) - TableGen 定义 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/10-padding-quantization.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/10-padding-quantization.md new file mode 100644 index 00000000..625451c3 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/01-DMA-Operations/10-padding-quantization.md @@ -0,0 +1,358 @@ +# 随路 Padding / 量化 / 激活 + +> 关键词：PadMode、FixpipePreQuantMode、FixpipePreReluMode、随路功能、量化、ReLU + +## 概述 + +HIVM 方言的 DMA 操作支持多种随路（on-the-fly）功能，这些功能在数据搬运过程中同时执行，无需额外的计算操作。随路功能是 Ascend NPU 硬件的重要特性，可以显著减少计算开销和内存访问次数。 + +本文档汇总所有 DMA 操作的随路功能，包括： +- **Padding（填充）**：在数据搬运时填充目标缓冲区 +- **Quantization（量化）**：在 Fixpipe 搬运时执行类型转换 +- **Activation（激活）**：在 Fixpipe 搬运时执行激活函数 + +## Padding（填充） + +### 概述 + +Padding 功能允许在 DMA 搬运过程中，将目标缓冲区中未被源数据覆盖的位置填充为指定值。这对于处理边界不齐的数据非常有用，例如在 Tiling 后的边界 Tile 中，源数据可能小于目标缓冲区。 + +### PadMode 枚举 + +定义于 [HIVMAttrs.td:L329-L349](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L329-L349) + +| 枚举值 | 数值 | IR 字面量 | 说明 | +|--------|------|----------|------| +| `PadNull` | 0 | `PadNull` | 不填充，源和目标形状必须相同 | +| `PadFirstElem` | 1 | `PadFirstElem` | 使用源数据的第一个元素填充 | +| `PadValue` | 2 | `PadValue` | 使用指定的 pad_value 填充 | + +### IR 表示 + +```mlir +#hivm.padmode +#hivm.padmode +#hivm.padmode +``` + +### 支持 Padding 的操作 + +| 操作 | 左侧 Padding | 右侧 Padding | PadFirstElem | PadValue | +|------|-------------|-------------|-------------|---------| +| `hir.load` | 支持 | 支持 | 支持 | 支持 | +| `hir.copy` | 支持 | 不支持 | 支持 | 支持 | + +### 使用模式 + +#### PadValue 模式 + +```mlir +%val = arith.constant 0.0 : f16 +hivm.hir.load ins(%src : memref<16x15xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + pad_mode = #hivm.padmode + pad_value = %val : f16 +``` + +当仅指定 `pad_value` 而不指定 `pad_mode` 时，编译器自动推断为 `PadValue` 模式： + +```mlir +hivm.hir.load ins(%src : memref<16x15xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + pad_value = %val : f16 +``` + +#### PadFirstElem 模式 + +```mlir +hivm.hir.load ins(%src : memref<16x15xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + pad_mode = #hivm.padmode +``` + +#### 左侧 Padding + +```mlir +%c0 = arith.constant 0 : index +hivm.hir.load ins(%src : memref<16x16xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + pad_mode = #hivm.padmode + pad_value = %val : f16 + left_padding_num = %c0 : index +``` + +### Padding 与 init_out_buffer 的关系 + +`init_out_buffer` 属性控制是否在搬运前初始化整个目标缓冲区。当 Padding 无法覆盖所有位置时（例如右侧 Padding 之外的区域），`init_out_buffer` 可以确保缓冲区的干净状态。 + +```mlir +hivm.hir.load ins(%src : memref<16x15xf16, #hivm.address_space>) + outs(%dst : memref<16x16xf16, #hivm.address_space>) + init_out_buffer = true + pad_value = %val : f16 +``` + +### PaddingOption 枚举（Gather/Scatter 专用） + +定义于 [HIVMAttrs.td:L920-L928](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L920-L928) + +| 枚举值 | 数值 | IR 字面量 | 说明 | +|--------|------|----------|------| +| `PAD_ZERO` | 1 | `zero` | 用零填充越界位置 | +| `PAD_NAN` | 2 | `nan` | 用 NaN 填充越界位置 | + +此枚举仅用于 `hir.gather_load` 操作的 `padding` 属性。 + +## Quantization（量化） + +### 概述 + +量化功能是 `hir.fixpipe` 操作的独有随路功能，在数据从 L0C 搬运到其他内存层级时，同时执行类型转换。这避免了额外的向量计算操作，可以显著提升性能。 + +### FixpipePreQuantMode 枚举 + +定义于 [HIVMAttrs.td:L783-L801](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L783-L801) + +| 枚举值 | 数值 | IR 字面量 | 源类型 | 目标类型 | 说明 | +|--------|------|----------|--------|---------|------| +| `NO_QUANT` | 0 | `NO_QUANT` | - | - | 不执行量化（默认） | +| `F322F16` | 1 | `F322F16` | F32 | F16 | F32 到 F16 转换 | +| `S322I8` | 9 | `S322I8` | F32 | I8 | F32 到 I8 量化 | +| `QF322F32_PRE` | 15 | `QF322F32_PRE` | F32 | F32 | 带 scale 的 F32 预量化 | +| `F322BF16` | 16 | `F322BF16` | F32 | BF16 | F32 到 BF16 转换 | + +### IR 表示 + +```mlir +#hivm.fixpipe_pre_quant_mode +#hivm.fixpipe_pre_quant_mode +#hivm.fixpipe_pre_quant_mode +#hivm.fixpipe_pre_quant_mode +#hivm.fixpipe_pre_quant_mode +``` + +### 量化模式详解 + +#### F322F16 + +将 F32 累加结果转换为 F16，适用于混合精度训练和推理： + +```mlir +%l0c = tensor.empty() : tensor<256x128xf32> +%dst = tensor.empty() : tensor<256x128xf16> +%result = hivm.hir.fixpipe {pre_quant = #hivm.fixpipe_pre_quant_mode} + ins(%l0c : tensor<256x128xf32>) + outs(%dst : tensor<256x128xf16>) + -> tensor<256x128xf16> +``` + +#### S322I8 + +将 F32 累加结果量化为 I8，适用于量化推理： + +```mlir +%l0c = tensor.empty() : tensor<256x128xf32> +%dst = tensor.empty() : tensor<256x128xi8> +%result = hivm.hir.fixpipe {pre_quant = #hivm.fixpipe_pre_quant_mode} + ins(%l0c : tensor<256x128xf32>) + outs(%dst : tensor<256x128xi8>) + -> tensor<256x128xi8> +``` + +#### QF322F32_PRE + +带缩放因子的 F32 到 F32 预量化，需要提供 `quant_scale` 参数： + +```mlir +%scale = arith.constant 0.125 : f32 +%l0c = tensor.empty() : tensor<256x128xf32> +%dst = tensor.empty() : tensor<256x128xf32> +%result = hivm.hir.fixpipe {pre_quant = #hivm.fixpipe_pre_quant_mode} + ins(%l0c : tensor<256x128xf32>) + outs(%dst : tensor<256x128xf32>) + quant_scale = %scale : f32 + -> tensor<256x128xf32> +``` + +#### F322BF16 + +将 F32 累加结果转换为 BF16： + +```mlir +%l0c = tensor.empty() : tensor<256x128xf32> +%dst = tensor.empty() : tensor<256x128xbf16> +%result = hivm.hir.fixpipe {pre_quant = #hivm.fixpipe_pre_quant_mode} + ins(%l0c : tensor<256x128xf32>) + outs(%dst : tensor<256x128xbf16>) + -> tensor<256x128xbf16> +``` + +### 量化与激活的组合 + +量化与激活可以同时使用，执行顺序为：先量化，后激活： + +```mlir +%l0c = tensor.empty() : tensor<256x128xf32> +%dst = tensor.empty() : tensor<256x128xf16> +%result = hivm.hir.fixpipe {pre_quant = #hivm.fixpipe_pre_quant_mode, + pre_relu = #hivm.fixpipe_pre_relu_mode} + ins(%l0c : tensor<256x128xf32>) + outs(%dst : tensor<256x128xf16>) + -> tensor<256x128xf16> +``` + +## Activation（激活） + +### 概述 + +激活功能是 `hir.fixpipe` 操作的独有随路功能，在数据从 L0C 搬运时同时执行激活函数。与量化类似，这避免了额外的向量计算操作。 + +### FixpipePreReluMode 枚举 + +定义于 [HIVMAttrs.td:L803-L819](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L803-L819) + +| 枚举值 | 数值 | IR 字面量 | 公式 | 说明 | +|--------|------|----------|------|------| +| `NO_RELU` | 0 | `NO_RELU` | - | 不执行激活（默认） | +| `NORMAL_RELU` | 1 | `NORMAL_RELU` | max(0, x) | 标准 ReLU | +| `LEAKY_RELU` | 2 | `LEAKY_RELU` | x > 0 ? x : alpha * x | Leaky ReLU | +| `P_RELU` | 3 | `P_RELU` | x > 0 ? x : p * x | Parametric ReLU | + +### IR 表示 + +```mlir +#hivm.fixpipe_pre_relu_mode +#hivm.fixpipe_pre_relu_mode +#hivm.fixpipe_pre_relu_mode +#hivm.fixpipe_pre_relu_mode +``` + +### 激活模式详解 + +#### NORMAL_RELU + +标准 ReLU 激活，将负值截断为 0： + +```mlir +%l0c = tensor.empty() : tensor<256x128xf16> +%dst = tensor.empty() : tensor<256x128xf16> +%result = hivm.hir.fixpipe {pre_relu = #hivm.fixpipe_pre_relu_mode} + ins(%l0c : tensor<256x128xf16>) + outs(%dst : tensor<256x128xf16>) + -> tensor<256x128xf16> +``` + +#### LEAKY_RELU + +Leaky ReLU 激活，负值乘以一个小的斜率因子： + +```mlir +%l0c = tensor.empty() : tensor<256x128xf16> +%dst = tensor.empty() : tensor<256x128xf16> +%result = hivm.hir.fixpipe {pre_relu = #hivm.fixpipe_pre_relu_mode} + ins(%l0c : tensor<256x128xf16>) + outs(%dst : tensor<256x128xf16>) + -> tensor<256x128xf16> +``` + +#### P_RELU + +Parametric ReLU，负值乘以可学习的参数： + +```mlir +%l0c = tensor.empty() : tensor<256x128xf16> +%dst = tensor.empty() : tensor<256x128xf16> +%result = hivm.hir.fixpipe {pre_relu = #hivm.fixpipe_pre_relu_mode} + ins(%l0c : tensor<256x128xf16>) + outs(%dst : tensor<256x128xf16>) + -> tensor<256x128xf16> +``` + +## 随路功能组合 + +### Fixpipe 随路功能执行顺序 + +``` +L0C 数据 + | + v +[Pre-Quant] (可选) -- 类型转换：F32->F16/F32->I8/F32->BF16/F32->F32(scale) + | + v +[Pre-ReLU] (可选) -- 激活函数：ReLU/LeakyReLU/P-ReLU + | + v +[DMA Mode] -- 布局转换：NZ2ND/NZ2DN/NZ2NZ + | + v +[Dual Dst] (可选) -- 双目标拆分：ROW_SPLIT/COLUMN_SPLIT + | + v +目标缓冲区 (GM/UB/L1) +``` + +### 组合约束 + +| 组合 | 是否支持 | 说明 | +|------|---------|------| +| 量化 + 激活 | 支持 | 先量化后激活 | +| 量化 + NZ2ND | 支持 | 先量化/激活后布局转换 | +| 量化 + NZ2DN | 支持 | 仅 Ascend950 | +| 量化 + 双目标 | 支持 | 仅 Ascend950，dst=UB | +| 激活 + NZ2ND | 支持 | - | +| 激活 + 双目标 | 支持 | 仅 Ascend950，dst=UB | +| NZ2DN + 双目标 | 不支持 | 互斥 | + +### Eviction Policy（缓存驱逐策略） + +定义于 [HIVMAttrs.td:L356-L372](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L356-L372) + +| 枚举值 | 数值 | IR 字面量 | 说明 | +|--------|------|----------|------| +| `EvictFirst` | 0 | `EvictFirst` | 优先驱逐，适用于一次性数据 | +| `EvictLast` | 1 | `EvictLast` | 最后驱逐，适用于反复访问数据 | + +Eviction Policy 仅用于 `hir.load` 操作，控制加载的数据在缓存中的保留策略。 + +## 随路功能在各操作中的支持矩阵 + +| 功能 | hir.load | hir.store | hir.copy | hir.fixpipe | hir.nd2nz | +|------|---------|----------|---------|------------|----------| +| PadValue | 支持 | - | 支持 | - | - | +| PadFirstElem | 支持 | - | 支持 | - | - | +| 左侧 Padding | 支持 | - | 支持 | - | - | +| 右侧 Padding | 支持 | - | - | - | - | +| init_out_buffer | 支持 | - | - | - | 支持 | +| Eviction Policy | 支持 | - | - | - | - | +| 量化 (PreQuant) | - | - | - | 支持 | - | +| 激活 (PreReLU) | - | - | - | 支持 | - | +| 布局转换 | - | - | - | 支持 | - | +| 双目标 | - | - | - | 支持 | - | +| 原子操作 | - | 支持 | - | - | - | + +## 常见问题 + +### Q: 为什么 copy 只支持左侧 Padding 而 load 支持双侧？ + +A: 这是硬件 DMA 引擎的限制。Copy 操作的硬件实现在 Padding 方面功能较弱，仅支持左侧填充。 + +### Q: 量化和激活可以同时使用吗？ + +A: 可以。Fixpipe 支持同时执行随路量化和随路激活，执行顺序为先量化后激活。这是 Fixpipe 的核心优势之一。 + +### Q: QF322F32_PRE 量化模式中 quant_scale 的作用是什么？ + +A: `quant_scale` 是一个缩放因子，在 F32 到 F32 的预量化过程中，将结果乘以该缩放因子。这常用于量化感知训练中，需要在保持 F32 精度的同时模拟量化效果。 + +### Q: PadFirstElem 的典型使用场景是什么？ + +A: `PadFirstElem` 常用于需要用边界值填充的场景，例如在卷积的 Padding 中，用边缘像素值填充边界，而不是用零值。 + +## 相关文档 + +- hir.load：[01-load.md](01-load.md) +- hir.copy：[05-copy.md](05-copy.md) +- hir.fixpipe：[06-fixpipe.md](06-fixpipe.md) +- 源码参考： + - [HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) - 枚举定义 + - [HIVMDMAOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMDMAOps.td) - DMA 操作定义 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/00-overview.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/00-overview.md new file mode 100644 index 00000000..1ca6c0be --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/00-overview.md @@ -0,0 +1,245 @@ +# HIVM 向量操作总览 + +> 关键词：HIVM, Vector Operations, PIPE_V, Elementwise, TableGen, AscendNPU + +## 概述 + +HIVM（Hybrid Intelligence Virtual Machine）方言的向量操作是 AscendNPU 向量计算单元（AIV）的核心 IR 抽象。所有向量操作在硬件层面映射到 Vector Pipe（PIPE_V），由向量计算核心执行。这些操作遵循 MLIR 的 DestinationStyleOpInterface，同时支持 tensor 和 memref 两种语义。 + +> Python API 对应：Triton Ascend 的 `tl` 原子操作最终通过编译流水线降级为 HIVM 向量操作。详见 docs_triton_ascend 相关文档。 + +## 操作继承层次 + +HIVM 向量操作采用 TableGen 多层继承定义，层次结构如下： + +``` +HIVM_Op -- 所有 HIVM 操作的基类（前缀 "hir."） + └── HIVM_StructuredOp -- 结构化操作基类（实现 HIVMStructuredOpInterface 等） + └── HIVM_VectorOp -- 向量操作基类（PIPE_V, VectorCoreTypeTrait） + ├── HIVM_ElementwiseNaryOp -- 逐元 N-ary 操作模板 + │ ├── HIVM_ElementwiseUnaryOp -- 一元操作（N=1） + │ ├── HIVM_ElementwiseBinaryOp -- 二元操作（N=2） + │ └── HIVM_ElementwiseTernaryOp -- 三元操作（N=3） + └── [独立向量操作] -- vbrc, vreduce, vtranspose, varange 等 +``` + +### 基类定义 + +| 基类 | 关键特性 | 源码位置 | +|------|---------|---------| +| [HIVM_Op](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMBase.td#L52-L58) | 操作前缀 `hir.`，命名空间 `::mlir::hivm` | HIVMBase.td | +| [HIVM_StructuredOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMBase.td#L60-L68) | 实现 HIVMStructuredOpInterface, MemoryEffectsOpInterface, FlattenInterface, LibraryFunctionOpInterface | HIVMBase.td | +| [HIVM_VectorOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L24-L40) | PIPE_V, VectorCoreTypeTrait, AlwaysSpeculatable, SinglePipeOpTrait, DestinationStyleOpInterface | HIVMVectorOps.td | +| [HIVM_ElementwiseNaryOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L46-L95) | AttrSizedOperandSegments, HIVMOpSameOperandsAndResultRank, UniformReassociationFlattenTrait, CollapsibleConsecutiveTargetDimsTrait, TransposableOTF | HIVMVectorOps.td | + +## Pipeline 归属 + +所有 HIVM 向量操作归属于 **PIPE_V**（Vector Pipe），对应硬件的向量计算单元。这在 `HIVM_VectorOp` 基类中通过 `OpPipeTrait<"PIPE::PIPE_V">` 静态指定。 + +``` +PIPE 枚举值 含义 适用操作类别 +────────────────────────────────────────────── +PIPE_S Scalar Pipe 标量操作 +PIPE_V Vector Pipe 向量操作（本文档范围） +PIPE_M Matrix Pipe 矩阵乘操作 +PIPE_MTE1 数据搬入 L1 Load 类操作 +PIPE_MTE2 数据搬入 L0A/L0B Load 类操作 +PIPE_MTE3 数据搬出 Store 类操作 +``` + +源码参考：[HIVMAttrs.td - Pipe 枚举定义](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L203-L236) + +## Traits 分类 + +HIVM 向量操作使用以下关键 Traits 进行约束： + +### 核心 Traits + +| Trait | 说明 | 依赖 | 源码位置 | +|-------|------|------|---------| +| ElementwiseNaryOpTrait\ | 逐元 N-ary 操作约束：N 个输入，1 个输出，相同 rank | HIVMStructuredOpInterface, HIVMOpSameOperandsAndResultRank | [HIVMTraits.td#L53-L56](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td#L53-L56) | +| CommutativeOpTrait | 交换律：输入操作数可交换 | 无 | [HIVMTraits.td#L76](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td#L76) | +| OperElemTypeConstraints\ | 操作数元素类型约束 | 无 | [HIVMTraits.td#L90-L106](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td#L90-L106) | +| VectorOnlyTrait\ | 指定操作数只能为向量类型 | 无 | [HIVMTraits.td#L80-L81](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td#L80-L81) | +| ScalarOnlyTrait\ | 指定操作数只能为标量类型 | 无 | [HIVMTraits.td#L87](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td#L87) | +| StaticMaxRankTrait\ | 静态已知最大 rank 限制 | 无 | [HIVMTraits.td#L64](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td#L64) | +| InferMaxRankTrait | 运行时推断最大 rank | 无 | [HIVMTraits.td#L67](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td#L67) | + +### OTF（On-The-Fly）Traits + +| Trait | 说明 | 源码位置 | +|-------|------|---------| +| BroadcastableOTF | 支持 OTF 广播：在计算时自动扩展指定维度 | [HIVMTraits.td#L130-L132](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td#L130-L132) | +| TransposableOTF | 支持 OTF 转置：在计算时自动重排维度 | [HIVMTraits.td#L153-L155](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td#L153-L155) | + +### Flatten 相关 Traits + +| Trait | 说明 | 源码位置 | +|-------|------|---------| +| UniformReassociationFlattenTrait | 所有操作数和结果可以使用相同的维度重关联进行展平 | [HIVMTraits.td#L183](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td#L183) | +| CollapsibleConsecutiveTargetDimsTrait | 标记操作在展平时必须保持目标维度的独立性和秩 | [HIVMTraits.td#L191-L192](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td#L191-L192) | + +## 与 HFusion 逐元操作的对应关系 + +HIVM 向量操作在编译流水线中会降级到 HFusion 方言的逐元操作。主要对应关系： + +| HIVM 操作 | HFusion 操作 | 说明 | +|-----------|-------------|------| +| hir.vrelu | hfusion.elemwise_unary {fun = relu} | ReLU 激活 | +| hir.vnot | hfusion.elemwise_unary {fun = vnot} | 按位取反 | +| hir.vcast | hfusion.cast {round_mode = ...} | 类型转换 | +| hir.vcmp | hfusion.compare {compare_fn = veq/vne/vlt/vle/vgt/vge} | 比较操作 | +| hir.vreduce | hfusion.reduce / hfusion.reduce_with_index | 归约操作 | +| hir.varange | hfusion.arange | 范围生成 | +| hir.vinterleave | hfusion.interleave | 交错合并 | +| hir.vdeinterleave | hfusion.deinterleave | 交错分离 | + +对于标准的算术操作（vadd, vsub, vmul, vdiv, vabs, vexp, vln 等），HIVM 操作降级到 `linalg` 方言的对应操作（linalg.add, linalg.sub, linalg.abs, linalg.exp 等）。 + +## 所有向量操作总表 + +### 一元运算（Unary Ops） + +| 操作 | 助记符 | 说明 | 支持元素类型 | 最大 Rank | +|------|--------|------|-------------|-----------| +| VExpOp | hir.vexp | 逐元指数运算 | F16, F32 | 3 | +| VAbsOp | hir.vabs | 逐元绝对值 | F16, F32, I8, I16, I32, I64 | 3 | +| VLnOp | hir.vln | 逐元自然对数 | F16, F32 | 3 | +| VReluOp | hir.vrelu | 逐元 ReLU | F16, F32, I32 | 3 | +| VRsqrtOp | hir.vrsqrt | 逐元倒数平方根 | F16, F32 | 3 | +| VSqrtOp | hir.vsqrt | 逐元平方根 | F16, F32 | 3 | +| VTanhOp | hir.vtanh | 逐元双曲正切 | 浮点类型 | - | +| VSinOp | hir.vsin | 逐元正弦 | 浮点类型 | - | +| VCosOp | hir.vcos | 逐元余弦 | 浮点类型 | - | +| VErfOp | hir.verf | 逐元误差函数 | 浮点类型 | - | +| VRecOp | hir.vrec | 逐元倒数 | F16, F32 | 3 | +| VNotOp | hir.vnot | 逐元按位取反 | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32 | 3 | +| VCastOp | hir.vcast | 逐元类型转换 | 多种（见 04-cast-ops.md） | 2 | + +### 二元运算（Binary Ops） + +| 操作 | 助记符 | 说明 | 支持元素类型 | 最大 Rank | 交换律 | +|------|--------|------|-------------|-----------|--------| +| VAddOp | hir.vadd | 逐元加法 | I8, I16, I32, F16, F32, I64 | 3 | 是 | +| VSubOp | hir.vsub | 逐元减法 | I8, I16, I32, F16, F32, I64 | 3 | 否 | +| VMulOp | hir.vmul | 逐元乘法 | I16, I32, F16, F32, I64 | 3 | 是 | +| VDivOp | hir.vdiv | 逐元除法 | F16, F32, I16, I32, I64 | 3 | 否 | +| VMaxOp | hir.vmax | 逐元最大值 | I16, I32, F16, F32, I64 | 3 | 是 | +| VMinOp | hir.vmin | 逐元最小值 | I16, I32, F16, F32, I64 | 3 | 是 | +| VOrOp | hir.vor | 逐元按位或 | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32 | 3 | 是 | +| VAndOp | hir.vand | 逐元按位与 | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32 | 3 | 是 | +| VXorOp | hir.vxor | 逐元按位异或 | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64 | 2 | 否 | +| VModOp | hir.vmod | 逐元取模（有符号） | I16, I32, I64 | 1 | 否 | +| VModUIOp | hir.vmodui | 逐元取模（无符号） | I16, I32, I64 | 1 | 否 | +| VPowOp | hir.vpow | 逐元幂运算 | I32 | 1 | 否 | +| VShLOp | hir.vshl | 逐元左移 | I16, I32, I64 | 3 | 否 | +| VShROp | hir.vshr | 逐元右移 | I16, I32, I64 | 3 | 否 | +| VCmpOp | hir.vcmp | 逐元比较 | F16, F32, I8, I16, I32, I64（输入）; I1, I8（输出） | 1 | 否 | +| VMulExtOp | hir.vmulext | 逐元乘法高32位 | I32 | 3 | 否 | + +### 三元运算（Ternary Ops） + +| 操作 | 助记符 | 说明 | 支持元素类型 | 最大 Rank | +|------|--------|------|-------------|-----------| +| VSelOp | hir.vsel | 逐元条件选择 | 条件: I1/I8; 数据: I1, AnyI8, AnyI16, F16, BF16, AnyI32, F32, I64 | 1 | + +### 归约运算（Reduction Ops） + +| 操作 | 助记符 | 说明 | 支持元素类型 | +|------|--------|------|-------------| +| VReduceOp | hir.vreduce | 向量归约 | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, F32 | + +### 数据搬移（Data Movement Ops） + +| 操作 | 助记符 | 说明 | 支持元素类型 | +|------|--------|------|-------------| +| VBrcOp | hir.vbrc | 向量广播 | I8, UI8, I16, F16, UI16, I32, F8E4M3FN, F8E5M2, F32, UI32, BF16, I64, UI64, I1 | +| VTransposeOp | hir.vtranspose | 维度转置 | AnyI8, AnyI16, AnyI32, F16, BF16, F32, I64, UI64, F8E4M3FN, F8E5M2 | +| VInterleaveOp | hir.vinterleave | 交错合并 | I16, F16, UI16, I32, F32, UI32, BF16, I64, UI64, F8E4M3FN, F8E5M2 | +| VDeinterleaveOp | hir.vdeinterleave | 交错分离 | I8, I16, F16, UI16, I32, F32, UI32, BF16, I64, UI64 | +| VFlipOp | hir.vflip | 维度翻转 | I8, UI8, I16, I32, UI16, UI32, I64, UI64, F16, F32, BF16 | +| VPadOp | hir.vpad | 填充 | 无类型约束（由 pad_value 决定） | +| VConcatOp | hir.vconcat | 拼接 | SameOperandsElementType | +| VGatherOp | hir.vgather | 按索引收集 | 数据: I1, I8, I16, UI16, I32, UI32, F16, BF16, F32, F8E4M3FN, F8E5M2; 索引: I32 | + +### 累积与排序（Cumulative & Sort Ops） + +| 操作 | 助记符 | 说明 | 支持元素类型 | 最大 Rank | +|------|--------|------|-------------|-----------| +| VCumsumOp | hir.vcumsum | 累积求和 | I1, I8, I16, I32, I64, F16, F32, BF16 | 2 | +| VCumprodOp | hir.vcumprod | 累积求积 | I1, I8, I16, I32, I64, F16, F32, BF16 | 1 | +| VSortOp | hir.vsort | 排序 | F16, F32, I32, I64 | 1 | + +### 特殊操作（Special Ops） + +| 操作 | 助记符 | 说明 | 支持元素类型 | 最大 Rank | +|------|--------|------|-------------|-----------| +| VArangeOp | hir.varange | 范围序列生成 | I16, I32, F16, F32, I64 | 3 | +| VMulextendedOp | hir.vmulextended | 扩展乘法（高低位） | I16 | 1 | +| VMulExtOp | hir.vmulext | 乘法高32位 | I32 | 3 | + +## 通用操作数结构 + +### Elementwise N-ary 操作通用参数 + +所有继承自 `HIVM_ElementwiseNaryOp` 的操作共享以下参数结构： + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| $src | Variadic\ 或 Variadic\ | 是 | 输入操作数（支持向量-向量或向量-标量） | +| $dst | Variadic\ | 是 | 输出操作数（DestinationStyle） | +| $temp_buffer | Optional\ | 否 | 临时缓冲区（部分操作需要） | +| $transpose | DenseI64ArrayAttr (默认 {}) | 否 | OTF 转置维度排列 | +| $broadcast | DenseI64ArrayAttr (默认 {}) | 否 | OTF 广播维度索引 | +| $result | Variadic\ | 是（tensor 语义） | 结果类型 | + +### 通用 Assembly Format + +``` +hir. attr-dict + ins( : ) + outs( : ) + [temp_buffer( : )] + [broadcast = ] + [transpose = ] + [-> ] +``` + +## 标量降级 + +部分 HIVM 向量操作在特定条件下会被降级为标量循环执行，而非使用硬件向量指令。这通常发生在硬件向量计算单元不支持某些数据类型和操作组合时（如 i64 类型的算术运算、整数大小比较等）。标量降级会显著影响性能。 + +实现标量降级的操作均声明了 `ImplByScalarOpInterface`，通过 `shouldLowerToScalarLoops()` 方法判断是否需要降级。 + +**详细文档**：[11-scalar-lowering.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/11-scalar-lowering.md) + +### 支持标量降级的操作一览 + +| 操作 | 助记符 | 标量降级主要触发条件 | +|------|--------|-------------------| +| VAddOp | hir.vadd | i64 或 SIMT VF | +| VSubOp | hir.vsub | i64 或 SIMT VF | +| VMulOp | hir.vmul | i64 或 SIMT VF | +| VMinOp | hir.vmin | i64 或 SIMT VF | +| VMaxOp | hir.vmax | i64 或 SIMT VF | +| VAbsOp | hir.vabs | i64 或 SIMT VF | +| VShLOp | hir.vshl | i64 或 SIMT VF | +| VShROp | hir.vshr | i64 或 SIMT VF | +| VInterleaveOp | hir.vinterleave | i64 或 SIMT VF | +| VDeinterleaveOp | hir.vdeinterleave | i64 或 SIMT VF | +| VCmpOp | hir.vcmp | 整数类型 且 (非 i32 或 非 EQ/NE) | +| VMulExtOp | hir.vmulext | i32 或 i64（即始终降级） | +| VCumsumOp | hir.vcumsum | i64 或 累积维度为最后维度 | +| VCumprodOp | hir.vcumprod | i64 或 累积维度为最后维度 | +| VReduceOp | hir.vreduce | 依架构和 reduceOp 类型 | + +## 相关文档 + +- 标量降级详解：[11-scalar-lowering.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/11-scalar-lowering.md) +- Python API：docs_triton_ascend 中的原子操作文档 +- 源码参考： + - [HIVMVectorOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td) — 向量操作 TableGen 定义 + - [HIVMTraits.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td) — Traits 定义 + - [HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) — 枚举属性定义 + - [HIVMBase.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMBase.td) — 基类定义 + - [convert-hivm-to-upstream.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/ExecutionEngine/convert-hivm-to-upstream.mlir) — IR 测试示例 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/01-unary-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/01-unary-ops.md new file mode 100644 index 00000000..f1d11006 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/01-unary-ops.md @@ -0,0 +1,514 @@ +# HIVM 一元向量运算 + +> 关键词：HIVM, Unary, vexp, vabs, vln, vrelu, vrsqrt, vsqrt, vtanh, vsin, vcos, verf, vrec, vnot + +## 概述 + +HIVM 一元向量运算继承自 `HIVM_ElementwiseUnaryOp`，对输入向量的每个元素独立执行运算，产生一个相同形状的结果向量。所有一元运算满足 `ElementwiseNaryOpTrait<1>`，即 1 个输入操作数、1 个输出结果。 + +一元运算可分为以下类别： +- **数学函数**：vexp, vln, vsqrt, vrsqrt, vrec, vsin, vcos, vtanh, verf +- **数值操作**：vabs, vrelu +- **位操作**：vnot + +> Python API 对应：`tl.math.exp()`, `tl.math.log()`, `tl.abs()`, `tl.math.rsqrt()`, `tl.math.sqrt()`, `tl.math.tanh()`, `tl.math.sin()`, `tl.math.cos()`, `tl.math.erf()`, `tl.math.reciprocal()` 等。 + +## IR 操作定义 + +### 基类：HIVM_ElementwiseUnaryOp + +```tablegen +class HIVM_ElementwiseUnaryOp traits = []> : + HIVM_ElementwiseNaryOp], traits)>; +``` + +源码参考：[HIVMVectorOps.td#L101-L103](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L101-L103) + +--- + +### hir.vexp — 逐元指数运算 + +#### TableGen 定义 + +```tablegen +def VExpOp : HIVM_ElementwiseUnaryOp<"vexp", + [SameOperandsElementType, StaticMaxRankTrait<3>, + VectorOnlyTrait<0>, + OperElemTypeConstraints<[0, 1], [F16, F32]>, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> { + let summary = "Elementwise Vector Exponential Op"; + let description = baseClassDescription # [{ + Additional constraints: + 1. The input/init operands and result have the same element type. + }]; + let arguments = (ins Variadic:$src, + Variadic:$dst, + Optional:$temp_buffer, + DefaultValuedAttr:$transpose, + DefaultValuedAttr:$broadcast + ); +} +``` + +源码参考：[HIVMVectorOps.td#L105-L136](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L105-L136) + +#### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src | Variadic\ | 是 | 输入向量 | VectorOnly, 元素类型 F16/F32 | +| $dst | Variadic\ | 是 | 输出向量（DestinationStyle） | 与 src 相同元素类型 | +| $temp_buffer | Optional\ | 否 | 临时缓冲区 | ExtraBufferOpInterface | +| $transpose | DenseI64ArrayAttr (默认 {}) | 否 | OTF 转置维度 | TransposableOTF 约束 | +| $broadcast | DenseI64ArrayAttr (默认 {}) | 否 | OTF 广播维度 | BroadcastableOTF 约束 | +| $result | Variadic\ | 是（tensor 语义） | 结果 | 与 src 相同元素类型 | + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vexp | F16, F32 | OperElemTypeConstraints<[0, 1], [F16, F32]> | + +#### IR 示例 + +```mlir +%result = hivm.hir.vexp ins(%src : tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) -> tensor<5x?x10xf32> + +hivm.hir.vexp ins(%src : memref<5x?x10xf32>) outs(%dst : memref<5x?x10xf32>) + +%result = hivm.hir.vexp ins(%src : tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) broadcast = [0] -> tensor<5x?x10xf32> +``` + +--- + +### hir.vabs — 逐元绝对值 + +#### TableGen 定义 + +```tablegen +def VAbsOp : HIVM_ElementwiseUnaryOp<"vabs", + [SameOperandsElementType, StaticMaxRankTrait<3>, + VectorOnlyTrait<0>, + OperElemTypeConstraints<[0], [F16, F32, I8, I16, I32, I64]>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> +``` + +源码参考：[HIVMVectorOps.td#L138-L171](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L138-L171) + +#### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src | Variadic\ | 是 | 输入向量 | VectorOnly, 元素类型 F16/F32/I8/I16/I32/I64 | +| $dst | Variadic\ | 是 | 输出向量 | 与 src 相同元素类型 | +| $temp_buffer | Optional\ | 否 | 临时缓冲区 | - | +| $transpose | DenseI64ArrayAttr (默认 {}) | 否 | OTF 转置维度 | - | +| $broadcast | DenseI64ArrayAttr (默认 {}) | 否 | OTF 广播维度 | - | + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vabs | F16, F32, I8, I16, I32, I64 | OperElemTypeConstraints<[0], [F16, F32, I8, I16, I32, I64]> | + +#### 特殊接口 + +- **ImplByScalarOpInterface**：支持标量实现路径 +- **VectorizableOpInterface**：支持向量化 + +#### IR 示例 + +```mlir +%result = hivm.hir.vabs ins(%src : tensor<1x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) broadcast = [0] -> tensor<5x?x10xf32> + +hivm.hir.vabs ins(%src : memref<1x?x10xf32>) outs(%dst : memref<5x?x10xf32>) broadcast = [0] +``` + +--- + +### hir.vln — 逐元自然对数 + +#### TableGen 定义 + +```tablegen +def VLnOp : HIVM_ElementwiseUnaryOp<"vln", + [SameOperandsElementType, StaticMaxRankTrait<3>, + VectorOnlyTrait<0>, + OperElemTypeConstraints<[0], [F16, F32]>, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> +``` + +源码参考：[HIVMVectorOps.td#L173-L204](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L173-L204) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vln | F16, F32 | OperElemTypeConstraints<[0], [F16, F32]> | + +#### IR 示例 + +```mlir +%result = hivm.hir.vln ins(%src : tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) -> tensor<5x?x10xf32> + +hivm.hir.vln ins(%src : memref<5x?x10xf32>) outs(%dst : memref<5x?x10xf32>) +``` + +--- + +### hir.vrelu — 逐元 ReLU + +#### TableGen 定义 + +```tablegen +def VReluOp : HIVM_ElementwiseUnaryOp<"vrelu", + [SameOperandsElementType, StaticMaxRankTrait<3>, + VectorOnlyTrait<0>, + OperElemTypeConstraints<[0], [F16, F32, I32]>, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> +``` + +源码参考：[HIVMVectorOps.td#L206-L237](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L206-L237) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vrelu | F16, F32, I32 | OperElemTypeConstraints<[0], [F16, F32, I32]> | + +#### IR 示例 + +```mlir +%result = hivm.hir.vrelu ins(%src : tensor) outs(%dst : tensor<5x?x10xf32>) transpose = [1, 0, 2] -> tensor<5x?x10xf32> + +hivm.hir.vrelu ins(%src : memref<5x1x10xi32>) outs(%dst : memref<5x?x10xi32>) broadcast = [1] +``` + +降级到 HFusion：`hfusion.elemwise_unary {fun = #hfusion.unary_fn}` + +--- + +### hir.vrsqrt — 逐元倒数平方根 + +#### TableGen 定义 + +```tablegen +def VRsqrtOp : HIVM_ElementwiseUnaryOp<"vrsqrt", + [SameOperandsElementType, StaticMaxRankTrait<3>, + VectorOnlyTrait<0>, + OperElemTypeConstraints<[0], [F16, F32]>, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> +``` + +源码参考：[HIVMVectorOps.td#L239-L270](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L239-L270) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vrsqrt | F16, F32 | OperElemTypeConstraints<[0], [F16, F32]> | + +#### IR 示例 + +```mlir +%result = hivm.hir.vrsqrt ins(%src : tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) -> tensor<5x?x10xf32> + +hivm.hir.vrsqrt ins(%src : memref<5x?x10xf32>) outs(%dst : memref<5x?x10xf32>) +``` + +--- + +### hir.vsqrt — 逐元平方根 + +#### TableGen 定义 + +```tablegen +def VSqrtOp : HIVM_ElementwiseUnaryOp<"vsqrt", + [SameOperandsElementType, StaticMaxRankTrait<3>, + VectorOnlyTrait<0>, + OperElemTypeConstraints<[0], [F16, F32]>, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> +``` + +源码参考：[HIVMVectorOps.td#L272-L303](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L272-L303) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vsqrt | F16, F32 | OperElemTypeConstraints<[0], [F16, F32]> | + +#### IR 示例 + +```mlir +%result = hivm.hir.vsqrt ins(%src : tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) -> tensor<5x?x10xf32> +``` + +--- + +### hir.vtanh — 逐元双曲正切 + +#### TableGen 定义 + +```tablegen +def VTanhOp : HIVM_ElementwiseUnaryOp<"vtanh", + [SameOperandsElementType, + VectorOnlyTrait<0>, NoLibraryFunctionTrait + ]> +``` + +源码参考：[HIVMVectorOps.td#L305-L315](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L305-L315) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vtanh | 浮点类型 | SameOperandsElementType（无显式 OperElemTypeConstraints） | + +注意：vtanh 使用 `NoLibraryFunctionTrait`，表示不通过库函数实现，而是使用内联实现路径。 + +#### IR 示例 + +```mlir +%result = hivm.hir.vtanh ins(%src : tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) -> tensor<5x?x10xf32> + +hivm.hir.vtanh ins(%src : memref<5x?x10xf32>) outs(%dst : memref<5x?x10xf32>) +``` + +--- + +### hir.vsin — 逐元正弦 + +#### TableGen 定义 + +```tablegen +def VSinOp : HIVM_ElementwiseUnaryOp<"vsin", + [SameOperandsElementType, + VectorOnlyTrait<0>, NoLibraryFunctionTrait + ]> +``` + +源码参考：[HIVMVectorOps.td#L317-L327](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L317-L327) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vsin | 浮点类型 | SameOperandsElementType（无显式 OperElemTypeConstraints） | + +#### IR 示例 + +```mlir +%result = hivm.hir.vsin ins(%src : tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) -> tensor<5x?x10xf32> +``` + +--- + +### hir.vcos — 逐元余弦 + +#### TableGen 定义 + +```tablegen +def VCosOp : HIVM_ElementwiseUnaryOp<"vcos", + [SameOperandsElementType, + VectorOnlyTrait<0>, NoLibraryFunctionTrait + ]> +``` + +源码参考：[HIVMVectorOps.td#L329-L339](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L329-L339) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vcos | 浮点类型 | SameOperandsElementType（无显式 OperElemTypeConstraints） | + +#### IR 示例 + +```mlir +%result = hivm.hir.vcos ins(%src : tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) -> tensor<5x?x10xf32> +``` + +--- + +### hir.verf — 逐元误差函数 + +#### TableGen 定义 + +```tablegen +def VErfOp : HIVM_ElementwiseUnaryOp<"verf", + [SameOperandsElementType, + VectorOnlyTrait<0>, NoLibraryFunctionTrait + ]> +``` + +源码参考：[HIVMVectorOps.td#L341-L351](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L341-L351) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| verf | 浮点类型 | SameOperandsElementType（无显式 OperElemTypeConstraints） | + +#### IR 示例 + +```mlir +%result = hivm.hir.verf ins(%src : tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) -> tensor<5x?x10xf32> + +hivm.hir.verf ins(%src : memref<5x?x10xf32>) outs(%dst : memref<5x?x10xf32>) +``` + +--- + +### hir.vrec — 逐元倒数 + +#### TableGen 定义 + +```tablegen +def VRecOp : HIVM_ElementwiseUnaryOp<"vrec", + [SameOperandsElementType, + VectorOnlyTrait<0>, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0], [F16, F32]>, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> +``` + +源码参考：[HIVMVectorOps.td#L353-L384](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L353-L384) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vrec | F16, F32 | OperElemTypeConstraints<[0], [F16, F32]> | + +#### IR 示例 + +```mlir +%result = hivm.hir.vrec ins(%src : tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) -> tensor<5x?x10xf32> + +hivm.hir.vrec ins(%src : memref<5x?x10xf32>) outs(%dst : memref<5x?x10xf32>) +``` + +--- + +### hir.vnot — 逐元按位取反 + +#### TableGen 定义 + +```tablegen +def VNotOp : HIVM_ElementwiseUnaryOp<"vnot", + [SameOperandsElementType, + VectorOnlyTrait<0>, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0], + [I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32]>, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> +``` + +源码参考：[HIVMVectorOps.td#L386-L418](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L386-L418) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vnot | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32 | OperElemTypeConstraints<[0], [I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32]> | + +注意：vnot 支持浮点类型的按位取反（即对浮点数的位模式取反），这在某些数值算法中有特殊用途。 + +#### IR 示例 + +```mlir +hivm.hir.vnot ins(%src : memref<5x1x10xi32>) outs(%dst : memref<5x?x10xi32>) broadcast = [1] +``` + +降级到 HFusion：`hfusion.elemwise_unary {fun = #hfusion.unary_fn}` + +## 数据类型约束汇总 + +| 操作 | 支持的元素类型 | 最大 Rank | OTF 广播 | OTF 转置 | Extra Buffer | 标量实现 | 向量化 | +|------|--------------|-----------|---------|---------|-------------|---------|--------| +| vexp | F16, F32 | 3 | 是 | 是 | 是 | - | - | +| vabs | F16, F32, I8, I16, I32, I64 | 3 | 是 | 是 | 是 | 是 | 是 | +| vln | F16, F32 | 3 | 是 | 是 | 是 | - | - | +| vrelu | F16, F32, I32 | 3 | 是 | 是 | 是 | - | - | +| vrsqrt | F16, F32 | 3 | 是 | 是 | 是 | - | - | +| vsqrt | F16, F32 | 3 | 是 | 是 | 是 | - | - | +| vtanh | 浮点 | - | - | - | - | - | - | +| vsin | 浮点 | - | - | - | - | - | - | +| vcos | 浮点 | - | - | - | - | - | - | +| verf | 浮点 | - | - | - | - | - | - | +| vrec | F16, F32 | 3 | 是 | 是 | 是 | - | - | +| vnot | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32 | 3 | 是 | 是 | 是 | - | - | + +## IR 层约束与验证 + +1. **元素类型一致性**：所有一元操作要求输入和输出具有相同的元素类型（`SameOperandsElementType`） +2. **Rank 一致性**：输入和输出必须具有相同的 rank（`HIVMOpSameOperandsAndResultRank`） +3. **VectorOnly 约束**：第一个输入操作数（$src[0]）必须为向量类型 +4. **OTF 广播约束**： + - broadcast 数组中的维度必须唯一 + - 所有广播维度 d 满足 `0 <= d < rank(dst)` + - 广播维度上 src 的 size 为 1 或等于 dst 的 size +5. **OTF 转置约束**： + - transpose 必须是 `range(rank(dst))` 的排列 + - `transpose[rank(dst) - 1] = rank(dst) - 1`（最后一维不可转置） +6. **temp_buffer**：需要 ExtraBufferOpInterface 的操作在硬件执行时需要额外的临时存储空间 + +## 与其他 IR 操作的关系 + +| HIVM 操作 | 上游 linalg 降级 | HFusion 降级 | +|-----------|-----------------|-------------| +| vexp | linalg.exp | - | +| vabs | linalg.abs | - | +| vln | linalg.log | - | +| vrelu | - | hfusion.elemwise_unary {fun = relu} | +| vrsqrt | linalg.rsqrt | - | +| vsqrt | linalg.sqrt | - | +| vtanh | linalg.tanh | - | +| vrec | linalg.reciprocal | - | +| verf | linalg.erf | - | +| vnot | - | hfusion.elemwise_unary {fun = vnot} | + +## 常见问题 + +**Q: vtanh, vsin, vcos, verf 为什么没有显式的 OperElemTypeConstraints？** +A: 这些操作使用 `NoLibraryFunctionTrait`，表示不通过预编译库函数实现。它们的类型约束由 `SameOperandsElementType` 隐式保证（输入输出同类型），具体支持的类型由硬件实现决定。 + +**Q: vnot 为什么支持浮点类型？** +A: vnot 执行的是按位取反操作，对浮点数的位模式取反在某些数值算法中有特殊用途（如快速生成 NaN 掩码等）。 + +**Q: 什么是一元操作的 temp_buffer？** +A: 部分一元操作（如 vexp, vln, vrsqrt 等）在硬件执行时需要额外的临时存储空间来保存中间计算结果。temp_buffer 参数用于提供这个空间。 + +## 相关文档 + +- Python API：docs_triton_ascend 中的 `tl.math` 模块文档 +- 源码参考： + - [HIVMVectorOps.td - Unary Ops](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L105-L418) + - [HIVMTraits.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td) + - [convert-hivm-to-upstream.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/ExecutionEngine/convert-hivm-to-upstream.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/02-binary-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/02-binary-ops.md new file mode 100644 index 00000000..29683d48 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/02-binary-ops.md @@ -0,0 +1,547 @@ +# HIVM 二元向量运算 + +> 关键词：HIVM, Binary, vadd, vsub, vmul, vdiv, vmax, vmin, vor, vand, vxor, vmod, vmodui, vpow + +## 概述 + +HIVM 二元向量运算继承自 `HIVM_ElementwiseBinaryOp`，对两个输入向量的对应元素执行运算，产生一个相同形状的结果向量。所有二元运算满足 `ElementwiseNaryOpTrait<2>`，即 2 个输入操作数、1 个输出结果。 + +二元运算可分为以下类别： +- **算术运算**：vadd, vsub, vmul, vdiv, vmod, vmodui, vpow +- **极值运算**：vmax, vmin +- **位运算**：vor, vand, vxor + +> Python API 对应：`tl.abs()`, Triton 的 `+`, `-`, `*`, `/`, `%`, `//` 等运算符，以及 `tl.minimum()`, `tl.maximum()` 等。 + +## IR 操作定义 + +### 基类：HIVM_ElementwiseBinaryOp + +```tablegen +class HIVM_ElementwiseBinaryOp traits = []> : + HIVM_ElementwiseNaryOp], traits)>; +``` + +源码参考：[HIVMVectorOps.td#L506-L508](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L506-L508) + +--- + +### hir.vadd — 逐元加法 + +#### TableGen 定义 + +```tablegen +def VAddOp : HIVM_ElementwiseBinaryOp<"vadd", + [SameOperandsElementType, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0, 1], [I8, I16, I32, F16, F32, I64]>, + CommutativeOpTrait, VectorOnlyTrait<0>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> { + let summary = "Elementwise Binary Vector Addition Op"; + let description = baseClassDescription # [{ + Additional constraints: + 1. The input/init operands and result have the same element type. + 2. Support both Vector-Vector and Vector-Scalar operation. + }]; + let arguments = (ins Variadic:$src, Variadic:$dst, + Optional:$temp_buffer, + DefaultValuedAttr:$transpose, + DefaultValuedAttr:$broadcast); +} +``` + +源码参考：[HIVMVectorOps.td#L510-L542](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L510-L542) + +#### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src | Variadic\ | 是 | 两个输入操作数（支持向量+向量或向量+标量） | 元素类型 I8/I16/I32/F16/F32/I64 | +| $dst | Variadic\ | 是 | 输出向量 | 与 src 相同元素类型 | +| $temp_buffer | Optional\ | 否 | 临时缓冲区 | - | +| $transpose | DenseI64ArrayAttr (默认 {}) | 否 | OTF 转置维度 | - | +| $broadcast | DenseI64ArrayAttr (默认 {}) | 否 | OTF 广播维度 | - | + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vadd | I8, I16, I32, F16, F32, I64 | OperElemTypeConstraints<[0, 1], [I8, I16, I32, F16, F32, I64]> | + +#### 特殊特性 + +- **CommutativeOpTrait**：满足交换律，`a + b = b + a` +- **ImplByScalarOpInterface**：支持向量-标量操作 +- **VectorizableOpInterface**：支持向量化 +- **$src 类型为 AnyType**：允许标量输入（非 AnyShaped），支持 Vector-Scalar 模式 + +#### IR 示例 + +```mlir +%result = hivm.hir.vadd ins(%a, %b : tensor, tensor) outs(%dst : tensor<5x?x10xf32>) transpose = [1, 0, 2] -> tensor<5x?x10xf32> + +hivm.hir.vadd ins(%vec, %scalar : tensor<23x77xi32>, i32) outs(%dst : tensor<23x77xi32>) -> tensor<23x77xi32> + +hivm.hir.vadd ins(%a, %b : memref<16xf16, #hivm.address_space>, memref<16xf16, #hivm.address_space>) outs(%c : memref<16xf16, #hivm.address_space>) +``` + +--- + +### hir.vsub — 逐元减法 + +#### TableGen 定义 + +```tablegen +def VSubOp + : HIVM_ElementwiseBinaryOp< + "vsub", [SameOperandsElementType, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0, 1], [I8, I16, I32, F16, F32, I64]>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + BroadcastableOTF]> +``` + +源码参考：[HIVMVectorOps.td#L596-L633](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L596-L633) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vsub | I8, I16, I32, F16, F32, I64 | OperElemTypeConstraints<[0, 1], [I8, I16, I32, F16, F32, I64]> | + +注意：vsub 不具有 CommutativeOpTrait（减法不满足交换律）。 + +#### IR 示例 + +```mlir +%result = hivm.hir.vsub ins(%a, %b : tensor<5x?x10xf32>, tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) -> tensor<5x?x10xf32> + +%result = hivm.hir.vsub ins(%a, %b : tensor<64x1xf32>, tensor<1x64xf32>) outs(%dst : tensor<64x64xf32>) broadcast = [0, 1] -> tensor<64x64xf32> +``` + +--- + +### hir.vmul — 逐元乘法 + +#### TableGen 定义 + +```tablegen +def VMulOp : HIVM_ElementwiseBinaryOp<"vmul", + [SameOperandsElementType, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0, 1], [I16, I32, F16, F32, I64]>, + CommutativeOpTrait, VectorOnlyTrait<0>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> +``` + +源码参考：[HIVMVectorOps.td#L544-L576](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L544-L576) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vmul | I16, I32, F16, F32, I64 | OperElemTypeConstraints<[0, 1], [I16, I32, F16, F32, I64]> | + +注意：vmul 不支持 I8 类型（与 vadd/vsub 不同）。 + +#### IR 示例 + +```mlir +%result = hivm.hir.vmul ins(%a, %b : tensor<5x?x10xf32>, tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) -> tensor<5x?x10xf32> +``` + +--- + +### hir.vdiv — 逐元除法 + +#### TableGen 定义 + +```tablegen +def VDivOp : HIVM_ElementwiseBinaryOp< + "vdiv", [SameOperandsElementType, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0, 1], [F16, F32, I16, I32, I64]>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + BroadcastableOTF]> { + let arguments = (ins Variadic:$src, + Variadic:$dst, + Optional:$temp_buffer, + DefaultValuedAttr:$isSigned, + DefaultValuedAttr:$transpose, + DefaultValuedAttr:$broadcast + ); +} +``` + +源码参考：[HIVMVectorOps.td#L635-L674](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L635-L674) + +#### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src | Variadic\ | 是 | 两个输入操作数 | 元素类型 F16/F32/I16/I32/I64 | +| $dst | Variadic\ | 是 | 输出向量 | 与 src 相同元素类型 | +| $temp_buffer | Optional\ | 否 | 临时缓冲区 | - | +| $isSigned | BoolAttr (默认 true) | 否 | 是否为有符号除法 | 仅对整数类型有效 | +| $transpose | DenseI64ArrayAttr (默认 {}) | 否 | OTF 转置维度 | - | +| $broadcast | DenseI64ArrayAttr (默认 {}) | 否 | OTF 广播维度 | - | + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vdiv | F16, F32, I16, I32, I64 | OperElemTypeConstraints<[0, 1], [F16, F32, I16, I32, I64]> | + +注意：vdiv 不支持 I8 类型，且仅支持 Vector-Vector 操作（无 ImplByScalarOpInterface）。 + +#### IR 示例 + +```mlir +%result = hivm.hir.vdiv ins(%a, %b : tensor<5x?x10xf32>, tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) -> tensor<5x?x10xf32> + +hivm.hir.vdiv ins(%a, %b : memref<5x?x10xf32>, memref<5x?x10xf32>) outs(%dst : memref<5x?x10xf32>) +``` + +--- + +### hir.vmax — 逐元最大值 + +#### TableGen 定义 + +```tablegen +def VMaxOp : HIVM_ElementwiseBinaryOp<"vmax", + [SameOperandsElementType, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0, 1], [I16, I32, F16, F32, I64]>, + CommutativeOpTrait, VectorOnlyTrait<0>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> +``` + +源码参考：[HIVMVectorOps.td#L676-L708](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L676-L708) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vmax | I16, I32, F16, F32, I64 | OperElemTypeConstraints<[0, 1], [I16, I32, F16, F32, I64]> | + +#### IR 示例 + +```mlir +%result = hivm.hir.vmax ins(%vec, %scalar : tensor<23x77xi32>, i32) outs(%dst : tensor<23x77xi32>) -> tensor<23x77xi32> + +hivm.hir.vmax ins(%a, %b : memref<5x?x10xf32>, memref<5x?x10xf32>) outs(%dst : memref<5x?x10xf32>) +``` + +--- + +### hir.vmin — 逐元最小值 + +#### TableGen 定义 + +```tablegen +def VMinOp : HIVM_ElementwiseBinaryOp<"vmin", + [SameOperandsElementType, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0, 1], [I16, I32, F16, F32, I64]>, + CommutativeOpTrait, VectorOnlyTrait<0>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> +``` + +源码参考：[HIVMVectorOps.td#L710-L742](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L710-L742) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vmin | I16, I32, F16, F32, I64 | OperElemTypeConstraints<[0, 1], [I16, I32, F16, F32, I64]> | + +#### IR 示例 + +```mlir +%result = hivm.hir.vmin ins(%vec, %scalar : tensor<23x77xi32>, i32) outs(%dst : tensor<23x77xi32>) -> tensor<23x77xi32> +``` + +--- + +### hir.vor — 逐元按位或 + +#### TableGen 定义 + +```tablegen +def VOrOp : HIVM_ElementwiseBinaryOp<"vor", + [SameOperandsElementType, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0, 1], + [I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32]>, + VectorOnlyTrait<0>, VectorOnlyTrait<1>, CommutativeOpTrait, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> +``` + +源码参考：[HIVMVectorOps.td#L744-L775](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L744-L775) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vor | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32 | OperElemTypeConstraints<[0, 1], [I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32]> | + +注意：vor 两个输入均为 VectorOnly，仅支持 Vector-Vector 操作。 + +#### IR 示例 + +```mlir +%result = hivm.hir.vor ins(%a, %a : tensor, tensor) outs(%dst : tensor<5x?x10xf32>) transpose = [1, 0, 2] -> tensor<5x?x10xf32> + +hivm.hir.vor ins(%a, %b : memref<5x1x10xi32>, memref<5x1x10xi32>) outs(%dst : memref<5x?x10xi32>) broadcast = [1] +``` + +--- + +### hir.vand — 逐元按位与 + +#### TableGen 定义 + +```tablegen +def VAndOp : HIVM_ElementwiseBinaryOp<"vand", + [SameOperandsElementType, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0, 1], + [I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32]>, + VectorOnlyTrait<0>, VectorOnlyTrait<1>, CommutativeOpTrait, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> +``` + +源码参考：[HIVMVectorOps.td#L777-L808](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L777-L808) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vand | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32 | OperElemTypeConstraints<[0, 1], [I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32]> | + +#### IR 示例 + +```mlir +%result = hivm.hir.vand ins(%a, %b : tensor<23x77xi32>, tensor<23x77xi32>) outs(%dst : tensor<23x77xi32>) -> tensor<23x77xi32> + +hivm.hir.vand ins(%a, %b : memref<5x1x10xi32>, memref<5x1x10xi32>) outs(%dst : memref<5x?x10xi32>) broadcast = [1] +``` + +--- + +### hir.vxor — 逐元按位异或 + +#### TableGen 定义 + +```tablegen +def VXorOp : HIVM_ElementwiseBinaryOp<"vxor", + [SameOperandsElementType, StaticMaxRankTrait<2>, + OperElemTypeConstraints<[0, 1], + [I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64]>, + VectorOnlyTrait<0>, VectorOnlyTrait<1>, + DeclareOpInterfaceMethods + ]> +``` + +源码参考：[HIVMVectorOps.td#L810-L843](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L810-L843) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vxor | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64 | OperElemTypeConstraints<[0, 1], [I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64]> | + +注意： +- vxor 不支持浮点类型（与 vor/vand 不同） +- vxor 最大 Rank 为 2（与 vor/vand 的 3 不同） +- vxor 不具有 CommutativeOpTrait +- vxor 不支持 BroadcastableOTF + +#### IR 示例 + +```mlir +hivm.hir.vxor ins(%a, %b : memref<5x?x10xi32>, memref<5x?x10xi32>) outs(%dst : memref<5x?x10xi32>) +``` + +--- + +### hir.vmod — 逐元取模（有符号） + +#### TableGen 定义 + +```tablegen +def VModOp : HIVM_ElementwiseBinaryOp<"vmod", + [SameOperandsElementType, StaticMaxRankTrait<1>, + VectorOnlyTrait<0>, + OperElemTypeConstraints<[0, 1], + [I16, I32, I64]> + ]> +``` + +源码参考：[HIVMVectorOps.td#L845-L857](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L845-L857) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vmod | I16, I32, I64 | OperElemTypeConstraints<[0, 1], [I16, I32, I64]> | + +注意：vmod 最大 Rank 仅为 1，且使用默认的 ElementwiseNaryOp 参数（无 temp_buffer, 无 OTF 广播/转置）。 + +#### IR 示例 + +```mlir +%result = hivm.hir.vmod ins(%a, %b : tensor<32xi64>, i64) outs(%dst : tensor<32xi64>) -> tensor<32xi64> +``` + +降级到上游：`arith.remsi` + +--- + +### hir.vmodui — 逐元取模（无符号） + +#### TableGen 定义 + +```tablegen +def VModUIOp : HIVM_ElementwiseBinaryOp<"vmodui", + [SameOperandsElementType, StaticMaxRankTrait<1>, + VectorOnlyTrait<0>, + OperElemTypeConstraints<[0, 1], + [I16, I32, I64]> + ]> +``` + +源码参考：[HIVMVectorOps.td#L859-L871](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L859-L871) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vmodui | I16, I32, I64 | OperElemTypeConstraints<[0, 1], [I16, I32, I64]> | + +#### IR 示例 + +```mlir +%result = hivm.hir.vmodui ins(%a, %b : tensor<32xi64>, i64) outs(%dst : tensor<32xi64>) -> tensor<32xi64> +``` + +降级到上游：`arith.remui` + +--- + +### hir.vpow — 逐元幂运算 + +#### TableGen 定义 + +```tablegen +def VPowOp : HIVM_ElementwiseBinaryOp<"vpow", + [SameOperandsElementType, StaticMaxRankTrait<1>, + OperElemTypeConstraints<[0, 1], [I32]>, + VectorOnlyTrait<0>, VectorOnlyTrait<1>, + DeclareOpInterfaceMethods + ]> +``` + +源码参考：[HIVMVectorOps.td#L981-L1010](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L981-L1010) + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vpow | I32 | OperElemTypeConstraints<[0, 1], [I32]> | + +注意：vpow 仅支持 I32 类型，两个输入均为 VectorOnly。 + +#### IR 示例 + +```mlir +%result = hivm.hir.vpow ins(%base, %exp : tensor<32xi32>, tensor<32xi32>) outs(%dst : tensor<32xi32>) -> tensor<32xi32> +``` + +## 数据类型约束汇总 + +| 操作 | 支持的元素类型 | 最大 Rank | 交换律 | Vec-Scalar | OTF 广播 | Extra Buffer | +|------|--------------|-----------|--------|-----------|---------|-------------| +| vadd | I8, I16, I32, F16, F32, I64 | 3 | 是 | 是 | 是 | 是 | +| vsub | I8, I16, I32, F16, F32, I64 | 3 | 否 | 是 | 是 | 是 | +| vmul | I16, I32, F16, F32, I64 | 3 | 是 | 是 | 是 | 是 | +| vdiv | F16, F32, I16, I32, I64 | 3 | 否 | 否 | 是 | 是 | +| vmax | I16, I32, F16, F32, I64 | 3 | 是 | 是 | 是 | 是 | +| vmin | I16, I32, F16, F32, I64 | 3 | 是 | 是 | 是 | 是 | +| vor | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32 | 3 | 是 | 否 | 是 | 是 | +| vand | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, BF16, F32 | 3 | 是 | 否 | 是 | 是 | +| vxor | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64 | 2 | 否 | 否 | 否 | 是 | +| vmod | I16, I32, I64 | 1 | 否 | 否 | 否 | 否 | +| vmodui | I16, I32, I64 | 1 | 否 | 否 | 否 | 否 | +| vpow | I32 | 1 | 否 | 否 | 否 | 是 | + +## IR 层约束与验证 + +1. **元素类型一致性**：所有二元操作要求输入和输出具有相同的元素类型（`SameOperandsElementType`），vcmp 除外 +2. **Rank 一致性**：输入和输出必须具有相同的 rank +3. **VectorOnly 约束**：标记为 VectorOnly 的操作数必须为向量类型 +4. **ScalarOnly 约束**：标记为 ScalarOnly 的操作数必须为标量类型（如 vshl 的第二个操作数） +5. **vdiv 的 isSigned 属性**：对整数除法，isSigned 决定使用有符号还是无符号除法 +6. **CommutativeOpTrait**：具有交换律的操作允许编译器交换输入操作数以优化性能 + +## 与其他 IR 操作的关系 + +| HIVM 操作 | 上游 linalg 降级 | 说明 | +|-----------|-----------------|------| +| vadd | linalg.add | - | +| vsub | linalg.sub | - | +| vmul | linalg.mul | - | +| vdiv | linalg.div | - | +| vmax | linalg.max | - | +| vmin | linalg.min | - | +| vor | linalg.map { arith.ori } | 浮点类型需要 bitcast | +| vand | linalg.map { arith.andi } | 浮点类型需要 bitcast | +| vxor | linalg.map { arith.xori } | - | +| vmod | arith.remsi | 有符号取模 | +| vmodui | arith.remui | 无符号取模 | + +## 常见问题 + +**Q: vdiv 的 isSigned 默认值是什么？** +A: 默认为 `true`，即有符号除法。对于无符号整数除法，需要显式设置 `isSigned = false`。 + +**Q: 为什么 vor/vand 支持浮点类型而 vxor 不支持？** +A: vor/vand 对浮点类型执行的是位级别的或/与操作（先 bitcast 到整数类型，执行位运算，再 bitcast 回浮点），这在某些掩码操作中有用。vxor 的硬件实现不支持浮点位操作。 + +**Q: vmod 和 vmodui 的区别是什么？** +A: vmod 执行有符号取模（对应 C 语言的 `%` 运算符，降级为 `arith.remsi`），vmodui 执行无符号取模（对应 `arith.remui`）。 + +**Q: vpow 为什么只支持 I32？** +A: 这是硬件约束。vpow 的硬件实现仅针对 I32 类型设计，浮点幂运算需要通过其他方式实现。 + +## 相关文档 + +- Python API：docs_triton_ascend 中的算术运算符文档 +- 源码参考： + - [HIVMVectorOps.td - Binary Ops](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L506-L1010) + - [convert-hivm-to-upstream.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/ExecutionEngine/convert-hivm-to-upstream.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/03-ternary-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/03-ternary-ops.md new file mode 100644 index 00000000..9ff3e90e --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/03-ternary-ops.md @@ -0,0 +1,152 @@ +# HIVM 三元向量运算 + +> 关键词：HIVM, Ternary, vsel, select, conditional + +## 概述 + +HIVM 三元向量运算继承自 `HIVM_ElementwiseTernaryOp`，对三个输入操作数执行逐元条件选择运算。目前仅有一个操作 `hir.vsel`，它根据条件向量的值从两个数据源向量中选择元素。 + +> Python API 对应：`tl.where(condition, x, y)` — 条件选择操作。 + +## IR 操作定义 + +### 基类：HIVM_ElementwiseTernaryOp + +```tablegen +class HIVM_ElementwiseTernaryOp traits = []> : + HIVM_ElementwiseNaryOp], traits)>; +``` + +源码参考：[HIVMVectorOps.td#L1016-L1018](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1016-L1018) + +--- + +### hir.vsel — 逐元条件选择 + +#### TableGen 定义 + +```tablegen +def VSelOp : HIVM_ElementwiseTernaryOp<"vsel", + [StaticMaxRankTrait<1>, + OperElemTypeConstraints<[/*condition=*/0], [I1,I8]>, + OperElemTypeConstraints<[/*src0=*/1, /*src0=*/2], [I1, AnyI8, AnyI16, F16, BF16, AnyI32, F32, I64]>, + DeclareOpInterfaceMethods, + BroadcastableOTF]> { + let summary = "Elementwise Vector Selection Op"; + let description = baseClassDescription # [{ + Select elements from two source vector according to the binary `condition` vector. + If the corresponding bit of the indicator is 1, select `src0`. Otherwise, + select `src1`. + + Additional constraints: + 1. The input vectors and output vector must have the same ranks. + 2. The element type of indicator vector must be bool. + }]; + let arguments = (ins Variadic:$src, + Variadic:$dst, + Optional:$temp_buffer, + DefaultValuedAttr:$transpose, + DefaultValuedAttr:$broadcast + ); + let assemblyFormat = [{ + attr-dict `ins` `(` $src `:` type($src) `)` + `outs` `(` $dst `:` type($dst) `)` + (`temp_buffer` `(` $temp_buffer^ `:` type($temp_buffer) `)`)? + (`->` type($result)^)? + }]; +} +``` + +源码参考：[HIVMVectorOps.td#L1020-L1050](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1020-L1050) + +#### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src[0] | AnyType | 是 | 条件向量（condition/indicator） | 元素类型 I1 或 I8 | +| $src[1] | AnyType | 是 | 数据源 src0（条件为真时选择） | 元素类型 I1/AnyI8/AnyI16/F16/BF16/AnyI32/F32/I64 | +| $src[2] | AnyType | 是 | 数据源 src1（条件为假时选择） | 与 src0 相同元素类型 | +| $dst | Variadic\ | 是 | 输出向量 | 与 src0/src1 相同元素类型 | +| $temp_buffer | Optional\ | 否 | 临时缓冲区 | ExtraBufferOpInterface | +| $transpose | DenseI64ArrayAttr (默认 {}) | 否 | OTF 转置维度 | - | +| $broadcast | DenseI64ArrayAttr (默认 {}) | 否 | OTF 广播维度 | - | + +#### 数据类型约束 + +| 操作数位置 | 语义 | 支持的元素类型 | 约束来源 | +|-----------|------|--------------|----------| +| $src[0] | 条件 | I1, I8 | OperElemTypeConstraints<[0], [I1, I8]> | +| $src[1] | 数据源0 | I1, AnyI8, AnyI16, F16, BF16, AnyI32, F32, I64 | OperElemTypeConstraints<[1, 2], [I1, AnyI8, AnyI16, F16, BF16, AnyI32, F32, I64]> | +| $src[2] | 数据源1 | I1, AnyI8, AnyI16, F16, BF16, AnyI32, F32, I64 | OperElemTypeConstraints<[1, 2], ...> | + +#### 语义说明 + +``` +dst[i] = condition[i] ? src0[i] : src1[i] +``` + +当条件向量的对应位为 1 时，选择 src0 的元素；否则选择 src1 的元素。 + +#### IR 示例 + +```mlir +%result = hivm.hir.vsel ins(%cond, %src0, %src1 : tensor<23x77xi1>, f32, tensor<23x77xf32>) outs(%dst : tensor<23x77xf32>) -> tensor<23x77xf32> + +%result = hivm.hir.vsel ins(%cond, %a, %b : i1, tensor<5x?x10xf32>, tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) -> tensor<5x?x10xf32> +``` + +完整使用示例（来自测试文件）： + +```mlir +%cond = hivm.hir.vcmp ins(%a, %b : tensor<23x77xf32>, f32) + outs(%init : tensor<23x77xi1>) compare_mode = -> tensor<23x77xi1> +%result = hivm.hir.vsel ins(%cond, %val_true, %val_false : tensor<23x77xi1>, f32, tensor<23x77xf32>) + outs(%dst : tensor<23x77xf32>) -> tensor<23x77xf32> +``` + +## IR 层约束与验证 + +1. **条件类型约束**：条件向量（$src[0]）的元素类型必须为 I1 或 I8 +2. **数据源类型一致性**：src0 和 src1 必须具有相同的元素类型 +3. **Rank 一致性**：所有输入向量和输出向量必须具有相同的 rank +4. **最大 Rank 限制**：StaticMaxRankTrait<1>，仅支持 1 维 +5. **BroadcastableOTF**：支持 OTF 广播，允许条件或数据源在指定维度上进行广播 + +## 与其他 IR 操作的关系 + +| HIVM 操作 | 上游降级 | HFusion 降级 | 说明 | +|-----------|---------|-------------|------| +| vsel | linalg.select | - | 语义等价于 linalg.select | + +vsel 通常与 vcmp 配合使用，形成条件选择模式： + +``` +vcmp → vsel (比较后选择) +``` + +降级示例： +```mlir +%cond = hivm.hir.vcmp ... compare_mode = -> tensor +%result = hivm.hir.vsel ins(%cond, %a, %b : ...) -> tensor + +%result = linalg.select ins(%cond, %a, %b : tensor, tensor, tensor) -> tensor +``` + +## 常见问题 + +**Q: vsel 的条件向量为什么支持 I8 而不仅仅是 I1？** +A: 硬件实现中，条件判断可以基于 I8 类型的非零值，不仅仅是布尔值。这提供了更大的灵活性，允许直接使用比较结果或掩码向量。 + +**Q: vsel 支持标量条件吗？** +A: 是的。从 IR 示例可以看到，$src 支持 AnyType（而非仅 AnyShaped），因此条件可以是标量 `i1` 值。当条件为标量时，所有元素使用相同的条件值。 + +**Q: vsel 的最大 Rank 为什么只有 1？** +A: 这是当前硬件实现的限制。对于多维条件选择，需要先展平为 1 维操作，或在编译阶段通过循环分解处理。 + +## 相关文档 + +- Python API：docs_triton_ascend 中的 `tl.where()` 文档 +- 源码参考： + - [HIVMVectorOps.td - VSelOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1020-L1050) + - [convert-hivm-to-upstream.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/ExecutionEngine/convert-hivm-to-upstream.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/04-cast-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/04-cast-ops.md new file mode 100644 index 00000000..04553128 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/04-cast-ops.md @@ -0,0 +1,272 @@ +# HIVM 类型转换操作 + +> 关键词：HIVM, vcast, round_mode, cast, bitcast, type conversion + +## 概述 + +`hir.vcast` 是 HIVM 方言的逐元类型转换操作，支持在不同数据类型之间进行转换。它是 HIVM 中最复杂的向量操作之一，提供了丰富的舍入模式（round_mode）和转换方式（cast）属性来控制转换行为。 + +> Python API 对应：`tl.cast()`, `tl.to()`, 以及隐式类型转换。 + +## IR 操作定义 + +### hir.vcast — 逐元类型转换 + +#### TableGen 定义 + +```tablegen +def VCastOp : HIVM_ElementwiseUnaryOp<"vcast", + [StaticMaxRankTrait<2>, + VectorOnlyTrait<0>, HIVMOpSameOperandsAndResultRank, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> { + let summary = "Elementwise Vector Type Conversion Op"; + let description = baseClassDescription # [{ + Additional constraints: + 1. Supports the following conversions: + + | src | dst | roundingmode | + |------|------|---------------------------------------------------| + | f32 | f32 | round, rint, floor, ceil, trunc | + | f32 | f16 | round, rint, floor, ceil, trunc, odd | + | f32 | i64 | round, rint, floor, ceil, trunc | + | f32 | i32 | round, rint, floor, ceil, trunc | + | f32 | i16 | round, rint, floor, ceil, trunc | + | f32 | s64 | round, rint, floor, ceil, trunc | + | f32 | bf16 | round, rint, floor, ceil, trunc | + | f16 | f32 | rint | + | f16 | i32 | round, rint, floor, ceil, trunc | + | f16 | i16 | round, rint, floor, ceil, trunc | + | f16 | i8 | round, rint, floor, ceil, trunc | + | f16 | ui8 | round, rint, floor, ceil, trunc | + | f16 | i4 | round, rint, floor, ceil, trunc | + | bf16 | f32 | rint | + | bf16 | i32 | round, rint, floor, ceil, trunc | + | ui8 | f16 | rint | + | i8 | f16 | rint | + | i8 | i1 | rint | + | i16 | f16 | round, rint, floor, ceil, trunc | + | i16 | f32 | rint | + | i32 | f32 | round, rint, floor, ceil, trunc | + | i32 | i64 | rint | + | i32 | i16 | rint | + | i64 | i32 | rint | + | i64 | f32 | round, rint, floor, ceil, trunc | + | i4 | f16 | rint | + | i1 | f16 | rint | + | i1 | f32 | rint | + }]; + let arguments = (ins Variadic:$src, + Variadic:$dst, + Optional:$temp_buffer, + DefaultValuedAttr:$round_mode, + DefaultValuedAttr:$cast, + DefaultValuedAttr:$transpose, + DefaultValuedAttr:$broadcast + ); + let hasVerifier = 1; +} +``` + +源码参考：[HIVMVectorOps.td#L420-L500](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L420-L500) + +#### 参数说明 + +| 参数 | 类型 | 必选 | 默认值 | 说明 | 约束 | +|------|------|------|--------|------|------| +| $src | Variadic\ | 是 | - | 输入向量 | VectorOnly | +| $dst | Variadic\ | 是 | - | 输出向量 | 与 src 相同 rank | +| $temp_buffer | Optional\ | 否 | - | 临时缓冲区 | - | +| $round_mode | HIVM_RoundModeAttr | 否 | RINT | 舍入模式 | 见下方枚举表 | +| $cast | HIVM_TypeFnAttr | 否 | cast_signed | 转换方式 | 见下方枚举表 | +| $transpose | DenseI64ArrayAttr | 否 | {} | OTF 转置维度 | - | +| $broadcast | DenseI64ArrayAttr | 否 | {} | OTF 广播维度 | - | + +## round_mode 属性 + +`round_mode` 控制 vcast 在精度降低时的舍入行为。定义在 [HIVMAttrs.td#L378-L406](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L378-L406)。 + +| 枚举值 | 字符串表示 | 说明 | C 语言等价 | +|--------|-----------|------|-----------| +| RINT (0) | rint | 向最近整数舍入，偶数优先 | `rint()` | +| ROUND (1) | round | 向最近整数舍入，远离零优先 | `round()` | +| FLOOR (2) | floor | 向负无穷方向舍入 | `floor()` | +| CEIL (3) | ceil | 向正无穷方向舍入 | `ceil()` | +| TRUNC (4) | trunc | 向零方向舍入 | `trunc()` | +| ODD (5) | odd | 向奇数舍入（Von Neumann 舍入） | - | +| TRUNCWITHOVERFLOW (6) | truncwithoverflow | 截断并允许溢出 | - | + +### 舍入模式选择指南 + +| 场景 | 推荐模式 | 说明 | +|------|---------|------| +| 默认转换 | rint | 最常用，IEEE 754 标准舍入 | +| 精确向下取整 | floor | 用于计算索引 | +| 精确向上取整 | ceil | 用于分配大小 | +| F32 → F16 降精度 | odd | 保持最大精度，避免舍入偏差 | +| 截断小数部分 | trunc | 直接丢弃小数部分 | + +## cast 属性 + +`cast` 属性控制整数类型转换时的符号处理方式。定义在 [HIVMAttrs.td#L431-L446](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L431-L446)。 + +| 枚举值 | 值 | 说明 | +|--------|---|------| +| cast_signed | 0 | 有符号转换（默认） | +| cast_unsigned | 1 | 无符号转换 | +| bitcast | 2 | 位转换（不改变位模式） | + +### cast 模式说明 + +- **cast_signed**：将源整数视为有符号数进行转换。例如 `i16 → f32` 时，负数会被正确解释。 +- **cast_unsigned**：将源整数视为无符号数进行转换。例如 `ui8 → f16` 时，所有值被视为正数。 +- **bitcast**：不改变底层位模式，仅重新解释类型。例如 `f32 → i32` 的 bitcast 保持 32 位不变。 + +## unsigned_mode 属性（隐含） + +虽然 vcast 本身使用 `cast` 属性，但 HIVM 还定义了 `UnsignedMode` 枚举，用于更细粒度的有符号/无符号转换控制： + +| 枚举值 | 字符串表示 | 说明 | +|--------|-----------|------| +| SI2SI (0) | si2si | 有符号 → 有符号 | +| SI2UI (1) | si2ui | 有符号 → 无符号 | +| UI2SI (2) | ui2si | 无符号 → 有符号 | +| UI2UI (3) | ui2ui | 无符号 → 无符号 | + +源码参考：[HIVMAttrs.td#L408-L425](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L408-L425) + +## 支持的转换路径 + +### 浮点 → 浮点 + +| 源类型 | 目标类型 | 支持的 round_mode | +|--------|---------|------------------| +| f32 | f32 | round, rint, floor, ceil, trunc | +| f32 | f16 | round, rint, floor, ceil, trunc, odd | +| f32 | bf16 | round, rint, floor, ceil, trunc | +| f16 | f32 | rint | +| bf16 | f32 | rint | + +### 浮点 → 整数 + +| 源类型 | 目标类型 | 支持的 round_mode | +|--------|---------|------------------| +| f32 | i64 | round, rint, floor, ceil, trunc | +| f32 | i32 | round, rint, floor, ceil, trunc | +| f32 | i16 | round, rint, floor, ceil, trunc | +| f16 | i32 | round, rint, floor, ceil, trunc | +| f16 | i16 | round, rint, floor, ceil, trunc | +| f16 | i8 | round, rint, floor, ceil, trunc | +| f16 | ui8 | round, rint, floor, ceil, trunc | +| f16 | i4 | round, rint, floor, ceil, trunc | +| bf16 | i32 | round, rint, floor, ceil, trunc | + +### 整数 → 浮点 + +| 源类型 | 目标类型 | 支持的 round_mode | +|--------|---------|------------------| +| i8 | f16 | rint | +| ui8 | f16 | rint | +| i16 | f16 | round, rint, floor, ceil, trunc | +| i16 | f32 | rint | +| i32 | f32 | round, rint, floor, ceil, trunc | +| i64 | f32 | round, rint, floor, ceil, trunc | +| i4 | f16 | rint | +| i1 | f16 | rint | +| i1 | f32 | rint | + +### 整数 → 整数 + +| 源类型 | 目标类型 | 支持的 round_mode | +|--------|---------|------------------| +| i8 | i1 | rint | +| i32 | i64 | rint | +| i32 | i16 | rint | +| i64 | i32 | rint | + +## IR 示例 + +### 基本类型转换 + +```mlir +hivm.hir.vcast ins(%src : memref<2x16xbf16>) outs(%dst : memref<2x16xf32>) +``` + +### 指定舍入模式 + +```mlir +hivm.hir.vcast ins(%src : memref<2x16xbf16>) outs(%dst : memref<2x16xf32>) + round_mode = #hivm.round_mode + +hivm.hir.vcast ins(%src : memref<2x16xbf16>) outs(%dst : memref<2x16xi32>) + round_mode = #hivm.round_mode + +hivm.hir.vcast ins(%src : memref<2x16xbf16>) outs(%dst : memref<2x16xi32>) + round_mode = #hivm.round_mode + +hivm.hir.vcast ins(%src : memref<2x16xbf16>) outs(%dst : memref<2x16xi32>) + round_mode = #hivm.round_mode +``` + +### Tensor 语义 + +```mlir +%result = hivm.hir.vcast ins(%src : tensor<23x77xi32>) outs(%dst : tensor<23x77xf32>) -> tensor<23x77xf32> +``` + +### 指定转换方式 + +```mlir +hivm.hir.vcast ins(%src : memref<2x16xi32>) outs(%dst : memref<2x16xf32>) + cast = #hivm.cast + +hivm.hir.vcast ins(%src : memref<2x16xi32>) outs(%dst : memref<2x16xf32>) + cast = #hivm.cast +``` + +## IR 层约束与验证 + +1. **Rank 一致性**：输入和输出必须具有相同的 rank（`HIVMOpSameOperandsAndResultRank`） +2. **最大 Rank 限制**：StaticMaxRankTrait<2>，最多支持 2 维 +3. **转换路径验证**：`hasVerifier = 1`，操作包含自定义验证器，根据硬件版本验证参数合法性 +4. **round_mode 与转换路径兼容性**：不是所有 round_mode 都适用于所有转换路径，详见上方"支持的转换路径"表 +5. **元素类型可以不同**：vcast 是唯一一个不要求 SameOperandsElementType 的 ElementwiseUnaryOp + +## 与其他 IR 操作的关系 + +| HIVM 操作 | 上游降级 | HFusion 降级 | 说明 | +|-----------|---------|-------------|------| +| vcast | - | hfusion.cast {round_mode = ...} | round_mode 属性直接映射 | + +降级示例： +```mlir +hivm.hir.vcast ins(%src : memref<2x16xbf16>) outs(%dst : memref<2x16xf32>) + round_mode = #hivm.round_mode + +%result = hfusion.cast {round_mode = #hfusion.round_mode} %src : memref<2x16xbf16> -> memref<2x16xf32> +``` + +## 常见问题 + +**Q: vcast 和 hivm.hir.bitcast 有什么区别？** +A: `hivm.hir.bitcast` 是独立的位转换操作，不改变位模式仅重新解释类型，不需要 round_mode。`hir.vcast` 使用 `cast = bitcast` 时语义类似，但 vcast 还支持需要舍入的类型转换。 + +**Q: F32 → F16 为什么支持 odd 模式而其他转换不支持？** +A: F32 → F16 是最常见的降精度场景，odd 舍入（Von Neumann 舍入）可以避免统计偏差，在量化场景中特别重要。硬件仅为此路径实现了 odd 模式。 + +**Q: TRUNCWITHOVERFLOW 和 TRUNC 有什么区别？** +A: TRUNC 在溢出时行为未定义或饱和，而 TRUNCWITHOVERFLOW 允许溢出发生而不做饱和处理。后者在需要检测溢出的场景中使用。 + +**Q: 为什么有些转换路径只支持 rint？** +A: 从低精度到高精度的转换（如 f16 → f32, i16 → f32）是精确的，不需要舍入，因此只使用默认的 rint 模式。从高精度到低精度的转换才需要指定舍入策略。 + +## 相关文档 + +- Python API：docs_triton_ascend 中的 `tl.cast()` 文档 +- 源码参考： + - [HIVMVectorOps.td - VCastOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L420-L500) + - [HIVMAttrs.td - RoundMode 枚举](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L378-L406) + - [HIVMAttrs.td - TypeFn 枚举](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L431-L446) + - [HIVMAttrs.td - UnsignedMode 枚举](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L408-L425) + - [convert-hivm-to-upstream.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/ExecutionEngine/convert-hivm-to-upstream.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/05-compare-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/05-compare-ops.md new file mode 100644 index 00000000..a217e555 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/05-compare-ops.md @@ -0,0 +1,247 @@ +# HIVM 比较运算 + +> 关键词：HIVM, vcmp, compare, EQ, NE, LT, GT, GE, LE + +## 概述 + +`hir.vcmp` 是 HIVM 方言的逐元比较操作，对两个输入向量的对应元素执行比较运算，产生一个布尔类型的结果向量。比较模式通过 `compare_mode` 属性指定，支持六种比较关系。 + +> Python API 对应：Triton 的比较运算符 `==`, `!=`, `<`, `>`, `<=`, `>=`，以及 `tl.math` 中的相关函数。 + +## IR 操作定义 + +### hir.vcmp — 逐元比较 + +#### TableGen 定义 + +```tablegen +def VCmpOp : HIVM_ElementwiseBinaryOp<"vcmp", + [StaticMaxRankTrait<1>, + OperElemTypeConstraints<[0, 1], [F16, F32, I8, I16, I32, I64]>, + OperElemTypeConstraints<[/*dstIdx=*/2], [I1, I8]>, + VectorOnlyTrait<0>, + DeclareOpInterfaceMethods + ]> { + let summary = "Elementwise Binary Vector Comparison Op"; + let description = baseClassDescription # [{ + Compare elements from two source vector. If the comparison result is true, + the corresponding bit of `dst` is 1 or 8. + + Additional constraints: + 1. The input vectors and output vector must have the same ranks + 2. The element type of `dst` must be bool + 3. The input is vector-only. + 4. Supports the following data type: + + | compare mode | element type | + |-------------------|-------------------------| + | GE/GT/LE/LT/NE/EQ | f16, f32, i16, i32, i64 | + }]; + let arguments = (ins Variadic:$src, + Variadic:$dst, + DefaultValuedAttr:$compare_mode, + DefaultValuedAttr:$transpose, + DefaultValuedAttr:$broadcast + ); + let assemblyFormat = [{ + attr-dict `ins` `(` $src `:` type($src) `)` + `outs` `(` $dst `:` type($dst) `)` + (`compare_mode` `=` $compare_mode^)? + (`->` type($result)^)? + }]; +} +``` + +源码参考：[HIVMVectorOps.td#L944-L979](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L944-L979) + +#### 参数说明 + +| 参数 | 类型 | 必选 | 默认值 | 说明 | 约束 | +|------|------|------|--------|------|------| +| $src[0] | AnyType | 是 | - | 第一个输入向量 | VectorOnly, 元素类型 F16/F32/I8/I16/I32/I64 | +| $src[1] | AnyType | 是 | - | 第二个输入向量/标量 | 元素类型 F16/F32/I8/I16/I32/I64 | +| $dst | Variadic\ | 是 | - | 输出向量 | 元素类型 I1 或 I8 | +| $compare_mode | HIVM_CmpModeAttr | 否 | EQ | 比较模式 | 见下方枚举表 | +| $transpose | DenseI64ArrayAttr | 否 | {} | OTF 转置维度 | - | +| $broadcast | DenseI64ArrayAttr | 否 | {} | OTF 广播维度 | - | + +注意：vcmp 没有 `temp_buffer` 参数，也没有 `SameOperandsElementType` 约束（因为输入和输出的元素类型不同）。 + +## compare_mode 属性 + +`compare_mode` 控制 vcmp 的比较语义。定义在 [HIVMAttrs.td#L452-L473](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L452-L473)。 + +| 枚举值 | 值 | 字符串表示 | 语义 | Python 等价 | +|--------|---|-----------|------|------------| +| EQ | 0 | eq | 等于 | `a == b` | +| NE | 1 | ne | 不等于 | `a != b` | +| LT | 2 | lt | 小于 | `a < b` | +| GT | 3 | gt | 大于 | `a > b` | +| GE | 4 | ge | 大于等于 | `a >= b` | +| LE | 5 | le | 小于等于 | `a <= b` | + +## 数据类型约束 + +vcmp 使用两组独立的 OperElemTypeConstraints： + +| 操作数位置 | 语义 | 支持的元素类型 | 约束来源 | +|-----------|------|--------------|----------| +| $src[0], $src[1] | 输入 | F16, F32, I8, I16, I32, I64 | OperElemTypeConstraints<[0, 1], [F16, F32, I8, I16, I32, I64]> | +| $dst (idx=2) | 输出 | I1, I8 | OperElemTypeConstraints<[2], [I1, I8]> | + +### 比较模式与数据类型对应 + +| 比较模式 | 支持的输入元素类型 | +|---------|------------------| +| EQ, NE, LT, GT, GE, LE | F16, F32, I16, I32, I64 | + +注意：虽然 OperElemTypeConstraints 允许 I8 输入，但 TableGen 描述中的约束表仅列出 f16, f32, i16, i32, i64。I8 的支持可能取决于具体硬件版本。 + +## IR 示例 + +### 等于比较 + +```mlir +%eq = hivm.hir.vcmp + ins(%a, %b : tensor<4xf32>, tensor<4xf32>) + outs(%init : tensor<4xi1>) + compare_mode = #hivm.compare_mode -> tensor<4xi1> +``` + +### 不等于比较 + +```mlir +%ne = hivm.hir.vcmp + ins(%a, %b : tensor<4xf32>, tensor<4xf32>) + outs(%init : tensor<4xi1>) + compare_mode = #hivm.compare_mode -> tensor<4xi1> +``` + +### 小于比较 + +```mlir +%lt = hivm.hir.vcmp + ins(%a, %b : tensor<4xf32>, tensor<4xf32>) + outs(%init : tensor<4xi1>) + compare_mode = #hivm.compare_mode -> tensor<4xi1> +``` + +### 向量-标量比较 + +```mlir +%result = hivm.hir.vcmp ins(%vec, %scalar : tensor<23x77xf32>, f32) + outs(%dst : tensor<23x77xi1>) compare_mode = -> tensor<23x77xi1> +``` + +### Memref 语义 + +```mlir +hivm.hir.vcmp ins(%a, %b : memref<4xf32>, memref<4xf32>) + outs(%dst : memref<4xi1>) compare_mode = #hivm.compare_mode +``` + +### 与 vsel 配合使用的典型模式 + +```mlir +%cond = hivm.hir.vcmp ins(%a, %zero : tensor, f32) + outs(%init : tensor) compare_mode = -> tensor +%result = hivm.hir.vsel ins(%cond, %val_true, %val_false : tensor, f32, tensor) + outs(%dst : tensor) -> tensor +``` + +## IR 层约束与验证 + +1. **输出类型约束**：输出向量的元素类型必须为 I1 或 I8 +2. **最大 Rank 限制**：StaticMaxRankTrait<1>，仅支持 1 维 +3. **VectorOnly 约束**：第一个输入操作数必须为向量类型 +4. **ImplByScalarOpInterface**：第二个输入可以是标量，支持向量-标量比较。此外，整数类型的 vcmp 在特定条件下会被降级为标量循环（详见下方「标量降级」章节） +5. **Rank 一致性**：输入向量和输出向量必须具有相同的 rank +6. **无 SameOperandsElementType**：输入和输出的元素类型不同（输入为数值类型，输出为布尔类型） + +## 与其他 IR 操作的关系 + +| HIVM 操作 | compare_mode | HFusion 降级 | 说明 | +|-----------|-------------|-------------|------| +| vcmp | eq | hfusion.compare {compare_fn = veq} | 等于 | +| vcmp | ne | hfusion.compare {compare_fn = vne} | 不等于 | +| vcmp | lt | hfusion.compare {compare_fn = vlt} | 小于 | +| vcmp | le | hfusion.compare {compare_fn = vle} | 小于等于 | +| vcmp | gt | hfusion.compare {compare_fn = vgt} | 大于 | +| vcmp | ge | hfusion.compare {compare_fn = vge} | 大于等于 | + +vcmp 的典型使用模式是与 vsel 配合，形成条件选择： + +``` +vcmp (compare_mode) → vsel (条件选择) +``` + +## 标量降级 + +vcmp 实现了 `ImplByScalarOpInterface`，在特定条件下会被降级为标量循环（`scf.for` + `arith.CmpIOp`），而非使用硬件向量比较指令。 + +### 降级条件 + +vcmp 在以下条件**全部满足**时降级为标量循环： + +1. 操作具有纯 buffer 语义（`hasPureBufferSemantics()` 为 true） +2. 第一个操作数为 MemRefType 或 TensorType +3. 元素类型为整数类型 +4. **且**满足以下任一条件： + - 元素类型不是 i32（如 i8, i16, i64） + - 元素类型是 i32，但比较模式不是 EQ 且不是 NE + +### 降级判断矩阵 + +| 元素类型 | EQ | NE | LT | GT | LE | GE | +|---------|----|----|----|----|----|-----| +| f16 | 向量 ✅ | 向量 ✅ | 向量 ✅ | 向量 ✅ | 向量 ✅ | 向量 ✅ | +| f32 | 向量 ✅ | 向量 ✅ | 向量 ✅ | 向量 ✅ | 向量 ✅ | 向量 ✅ | +| i8 | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | +| i16 | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | +| i32 | 向量 ✅ | 向量 ✅ | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | +| i64 | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | **标量** ⚠️ | + +### 根本原因 + +硬件向量比较指令对整数类型的支持有限： +- 浮点类型（f16/f32）：全部 6 种比较模式均有向量指令支持 +- i32 的 EQ/NE：硬件向量指令支持相等/不等比较 +- i32 的 LT/GT/LE/GE 及其他整数宽度：无对应向量指令，只能退化为标量循环 + +### 降级实现细节 + +标量降级时，vcmp 被分解为嵌套 `scf.for` 循环，循环体内逐元素执行 `arith.CmpIOp`（有符号比较谓词：slt/sgt/sle/sge/eq/ne）。由于 store 操作不支持 i1 类型，比较结果（i1）会先通过 `arith.ExtUIOp` 零扩展为 i8 再存储。 + +此外，在 `HIVMDecomposeOp` Pass 中有预处理步骤：将 vcmp 的输出从 i1 转为 i8 临时 buffer，之后用 `VCastOp` 转回 i1。 + +### 优化建议 + +- **优先使用浮点比较**：f16/f32 的所有比较模式都走向量路径，性能最优 +- **整数相等/不等比较使用 i32**：i32 的 EQ/NE 是唯一能走向量路径的整数比较 +- **避免整数大小比较**：i32 的 LT/GT/LE/GE 以及 i8/i16/i64 的所有比较都会标量降级 +- **替代策略**：如果业务逻辑允许，将整数比较转为浮点比较（先 cast 再 compare），可避免标量降级 + +> 完整的标量降级文档：[11-scalar-lowering.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/11-scalar-lowering.md) + +## 常见问题 + +**Q: vcmp 的输出为什么支持 I8 而不仅仅是 I1？** +A: 硬件实现中，比较结果可以存储为 I8 类型（非零表示真，零表示假），这在某些后续操作中更高效。I1 是最紧凑的表示，但 I8 在内存对齐方面更有优势。 + +**Q: vcmp 支持浮点比较的 NaN 处理吗？** +A: HIVM 遵循 IEEE 754 标准的浮点比较语义。NaN 与任何值（包括自身）的比较结果均为 false（EQ 和 NE 除外：NaN != NaN 为 true）。 + +**Q: 为什么 vcmp 没有 temp_buffer？** +A: 比较操作是简单的逐元操作，不需要额外的临时存储空间。 + +**Q: vcmp 的默认 compare_mode 为什么是 EQ？** +A: 等于比较是最常用的比较模式，作为默认值可以简化最常见的使用场景。 + +## 相关文档 + +- 标量降级详解：[11-scalar-lowering.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/11-scalar-lowering.md) +- Python API：docs_triton_ascend 中的比较运算符文档 +- 源码参考： + - [HIVMVectorOps.td - VCmpOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L944-L979) + - [HIVMAttrs.td - CmpMode 枚举](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L452-L473) + - [convert-hivm-to-upstream.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/ExecutionEngine/convert-hivm-to-upstream.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/06-shift-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/06-shift-ops.md new file mode 100644 index 00000000..952b214b --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/06-shift-ops.md @@ -0,0 +1,201 @@ +# HIVM 移位运算 + +> 关键词：HIVM, vshl, vshr, shift, round + +## 概述 + +HIVM 移位运算包含左移（`hir.vshl`）和右移（`hir.vshr`）两个操作。两者均属于 `HIVM_ElementwiseBinaryOp`，但具有特殊的操作数约束：第一个操作数为向量，第二个操作数为标量，即仅支持 Vector-Scalar 模式。 + +> Python API 对应：Triton 的 `<<` 和 `>>` 运算符。 + +## IR 操作定义 + +### hir.vshl — 逐元左移 + +#### TableGen 定义 + +```tablegen +def VShLOp : HIVM_ElementwiseBinaryOp<"vshl", + [SameOperandsElementType, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0, 1], [I16, I32, I64]>, + VectorOnlyTrait<0>, ScalarOnlyTrait<1>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> { + let summary = "Elementwise Binary Vector Shift Left Op"; + let description = baseClassDescription # [{ + Additional constraints: + 1. The input vector and result have the same element type. + 2. Support only Vector - Scalar operation. + }]; + let arguments = (ins Variadic:$src, Variadic:$dst, + Optional:$temp_buffer, + DefaultValuedAttr:$transpose, + DefaultValuedAttr:$broadcast); +} +``` + +源码参考：[HIVMVectorOps.td#L873-L904](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L873-L904) + +#### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src[0] | AnyType | 是 | 输入向量 | VectorOnly, 元素类型 I16/I32/I64 | +| $src[1] | AnyType | 是 | 移位量（标量） | ScalarOnly, 元素类型 I16/I32/I64 | +| $dst | Variadic\ | 是 | 输出向量 | 与 src 相同元素类型 | +| $temp_buffer | Optional\ | 否 | 临时缓冲区 | - | +| $transpose | DenseI64ArrayAttr (默认 {}) | 否 | OTF 转置维度 | - | +| $broadcast | DenseI64ArrayAttr (默认 {}) | 否 | OTF 广播维度 | - | + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vshl | I16, I32, I64 | OperElemTypeConstraints<[0, 1], [I16, I32, I64]> | + +#### 特殊约束 + +- **VectorOnlyTrait<0>**：第一个操作数（被移位向量）必须为向量 +- **ScalarOnlyTrait<1>**：第二个操作数（移位量）必须为标量 +- 仅支持 Vector-Scalar 操作模式 + +#### IR 示例 + +```mlir +%result = hivm.hir.vshl ins(%vec, %shift : tensor<32xi32>, i32) outs(%dst : tensor<32xi32>) -> tensor<32xi32> + +hivm.hir.vshl ins(%vec, %shift : memref<32xi32>, i32) outs(%dst : memref<32xi32>) +``` + +--- + +### hir.vshr — 逐元右移 + +#### TableGen 定义 + +```tablegen +def VShROp : HIVM_ElementwiseBinaryOp<"vshr", + [SameOperandsElementType, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0, 1], [I16, I32, I64]>, + VectorOnlyTrait<0>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + BroadcastableOTF + ]> { + let summary = "Elementwise Binary Vector Shift Right Op"; + let description = baseClassDescription # [{ + Additional constraints: + 1. The input vector and result have the same element type. + 2. Support only Vector - Scalar operation. + 3. If `round` is set to true, rounding is applied during arithmetic + shift right. + }]; + let arguments = (ins Variadic:$src, + Variadic:$dst, + Optional:$temp_buffer, + DefaultValuedOptionalAttr:$round, + DefaultValuedAttr:$transpose, + DefaultValuedAttr:$broadcast + ); + let assemblyFormat = [{ + attr-dict `ins` `(` $src `:` type($src) `)` + `outs` `(` $dst `:` type($dst) `)` + (`temp_buffer` `(` $temp_buffer^ `:` type($temp_buffer) `)`)? + (`round` `:` $round^ )? + (`->` type($result)^)? + }]; +} +``` + +源码参考：[HIVMVectorOps.td#L906-L942](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L906-L942) + +#### 参数说明 + +| 参数 | 类型 | 必选 | 默认值 | 说明 | 约束 | +|------|------|------|--------|------|------| +| $src[0] | AnyType | 是 | - | 输入向量 | VectorOnly, 元素类型 I16/I32/I64 | +| $src[1] | AnyType | 是 | - | 移位量（标量） | 元素类型 I16/I32/I64 | +| $dst | Variadic\ | 是 | - | 输出向量 | 与 src 相同元素类型 | +| $temp_buffer | Optional\ | 否 | - | 临时缓冲区 | - | +| $round | BoolAttr | 否 | true | 是否在算术右移时进行舍入 | - | +| $transpose | DenseI64ArrayAttr | 否 | {} | OTF 转置维度 | - | +| $broadcast | DenseI64ArrayAttr | 否 | {} | OTF 广播维度 | - | + +#### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vshr | I16, I32, I64 | OperElemTypeConstraints<[0, 1], [I16, I32, I64]> | + +#### round 属性说明 + +`round` 属性是 vshr 独有的，控制算术右移时的舍入行为： + +- **round = true（默认）**：在算术右移时进行舍入。如果移出的最高位为 1，则结果加 1。这相当于 `(src + (1 << (shift - 1))) >> shift`，实现更精确的除以 2^n 运算。 +- **round = false**：不进行舍入，直接截断。相当于标准的算术右移。 + +#### 特殊约束 + +- **VectorOnlyTrait<0>**：第一个操作数必须为向量 +- 注意：vshr 没有 `ScalarOnlyTrait<1>`，但描述中说明仅支持 Vector-Scalar 操作 +- **ImplByScalarOpInterface**：支持标量实现路径 + +#### IR 示例 + +```mlir +%result = hivm.hir.vshr ins(%vec, %shift : tensor<32xi32>, i32) outs(%dst : tensor<32xi32>) -> tensor<32xi32> + +hivm.hir.vshr ins(%vec, %shift : memref<32xi32>, i32) outs(%dst : memref<32xi32>) + +%result = hivm.hir.vshr ins(%vec, %shift : tensor<32xi32>, i32) outs(%dst : tensor<32xi32>) round : true -> tensor<32xi32> + +%result = hivm.hir.vshr ins(%vec, %shift : tensor<32xi32>, i32) outs(%dst : tensor<32xi32>) round : false -> tensor<32xi32> +``` + +## 数据类型约束汇总 + +| 操作 | 支持的元素类型 | 最大 Rank | 操作模式 | OTF 广播 | Extra Buffer | round 属性 | +|------|--------------|-----------|---------|---------|-------------|-----------| +| vshl | I16, I32, I64 | 3 | Vector-Scalar | 是 | 是 | 无 | +| vshr | I16, I32, I64 | 3 | Vector-Scalar | 是 | 是 | 是（默认 true） | + +## IR 层约束与验证 + +1. **元素类型一致性**：输入向量和输出向量必须具有相同的元素类型 +2. **VectorOnly 约束**：第一个输入操作数必须为向量类型 +3. **ScalarOnly 约束（vshl）**：第二个输入操作数必须为标量类型 +4. **仅支持整数类型**：I16, I32, I64，不支持浮点类型 +5. **移位量语义**：移位量为标量，所有元素使用相同的移位量 + +## 与其他 IR 操作的关系 + +| HIVM 操作 | 上游降级 | 说明 | +|-----------|---------|------| +| vshl | arith.shli | 逻辑左移 | +| vshr (round=false) | arith.shrsi | 算术右移（有符号） | +| vshr (round=true) | 需要额外的加法和移位组合 | 舍入右移 | + +## 常见问题 + +**Q: vshl 和 vshr 为什么只支持 Vector-Scalar 模式？** +A: 硬件向量移位指令的设计是所有元素使用相同的移位量。如果需要逐元素不同移位量，需要通过循环或其他方式实现。 + +**Q: vshr 的 round 属性有什么实际用途？** +A: round 属性在将右移用作除以 2^n 的近似时特别有用。例如，`x >> 1` 等价于 `x / 2`（截断），而 `round: true` 的 `(x + 1) >> 1` 更接近数学上的四舍五入除法。 + +**Q: vshr 是算术右移还是逻辑右移？** +A: vshr 执行算术右移（保留符号位）。逻辑右移（无符号右移）需要先将数据视为无符号类型。 + +**Q: vshl 为什么有 ScalarOnlyTrait 而 vshr 没有？** +A: 这是一个实现细节差异。vshr 虽然没有显式的 ScalarOnlyTrait<1>，但其描述中明确说明仅支持 Vector-Scalar 操作。ImplByScalarOpInterface 也暗示了标量操作数的支持。 + +## 相关文档 + +- Python API：docs_triton_ascend 中的位运算文档 +- 源码参考： + - [HIVMVectorOps.td - VShLOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L873-L904) + - [HIVMVectorOps.td - VShROp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L906-L942) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/07-reduction-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/07-reduction-ops.md new file mode 100644 index 00000000..7e960888 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/07-reduction-ops.md @@ -0,0 +1,236 @@ +# HIVM 归约运算 + +> 关键词：HIVM, vreduce, reduction, sum, prod, max, min, reduce_dims + +## 概述 + +`hir.vreduce` 是 HIVM 方言的向量归约操作，沿指定维度对输入向量执行归约运算。它支持 12 种归约操作，从基本的求和/求积到带索引的最大/最小值归约。vreduce 是 HIVM 中功能最丰富的向量操作之一，直接映射到硬件的向量归约指令。 + +> Python API 对应：`tl.reduce()`, `tl.sum()`, `tl.max()`, `tl.min()`, `tl.argmax()`, `tl.argmin()` 等。 + +## IR 操作定义 + +### hir.vreduce — 向量归约 + +#### TableGen 定义 + +```tablegen +def VReduceOp : HIVM_VectorOp<"vreduce", + [AttrSizedOperandSegments, + OperElemTypeConstraints<[0, 1], [I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, F32]>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + InferMaxRankTrait, DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + CollapsibleConsecutiveTargetDimsTrait, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods + ]> { + let summary = "Vector Reduction Op"; + let description = [{ + Recuce one or more axes of the source vector according to + the reduction axes array, starting from an init value. + + Constraints: + 1. The input vector and output vector must have the same rank + and the same element type. + 2. For the output operand, the size of the reduced axis must be 1. + 3. The reduction indices array can not be empty, + nor can be larger than the ranks of the input vector. + 4. The reduced indices must be in `[0, RankOfDstVec)`. + + Examples: + ```mlir + hivm.hir.vreduce ins(%src : memref) outs(%dst : memref<1xf32>) reduce_dims : [1] + %result = hivm.hir.vreduce ins(%src : tensor) outs(%dst : tensor<1xf32>) reduce_dims : [0] -> tensor<1xf32> + ``` + }]; + let arguments = (ins TensorOrMemref:$src, + Variadic:$dst, + Optional:$temp_buffer, + HIVM_ReduceOpAttr:$arith, + BoolAttr:$unsigned_src, + OptionalAttr:$tie_break_left, + DenseI64ArrayAttr:$reduce_dims + ); + let results = (outs Variadic:$result); + let assemblyFormat = [{ + attr-dict $arith `ins` `(` $src `:` type($src) `)` + `outs` `(` $dst `:` type($dst) `)` + (`temp_buffer` `(` $temp_buffer^ `:` type($temp_buffer) `)`)? + `unsigned_src` `=` $unsigned_src + (`tie_break_left` `=` $tie_break_left^)? + `reduce_dims` `=` $reduce_dims + (`->` type($result)^)? + }]; +} +``` + +源码参考：[HIVMVectorOps.td#L1129-L1216](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1129-L1216) + +#### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src | TensorOrMemref | 是 | 输入向量 | 元素类型 I1/I8/UI8/I16/UI16/I32/UI32/I64/UI64/F16/F32 | +| $dst | Variadic\ | 是 | 输出向量（归约值） | 与 src 相同 rank 和元素类型 | +| $temp_buffer | Optional\ | 否 | 临时缓冲区 | ExtraBufferOpInterface | +| $arith | HIVM_ReduceOpAttr | 是 | 归约操作类型 | 见下方枚举表 | +| $unsigned_src | BoolAttr | 是 | 源操作数是否为无符号 | - | +| $tie_break_left | OptionalAttr\ | 否 | 归约冲突时偏向左侧 | 仅 max_with_index/min_with_index | +| $reduce_dims | DenseI64ArrayAttr | 是 | 归约维度数组 | 不可为空，索引在 [0, rank) 内 | + +## reduce_operation 属性（$arith） + +`$arith` 属性指定归约操作的类型。定义在 [HIVMAttrs.td#L596-L631](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L596-L631)。 + +| 枚举值 | 值 | 字符串表示 | 语义 | 输出数量 | +|--------|---|-----------|------|---------| +| sum | 1 | \ | 求和 | 1 | +| prod | 2 | \ | 求积 | 1 | +| max | 3 | \ | 最大值 | 1 | +| min | 4 | \ | 最小值 | 1 | +| max_with_index | 5 | \ | 最大值及其索引 | 2 | +| min_with_index | 6 | \ | 最小值及其索引 | 2 | +| any | 7 | \ | 任意为真 | 1 | +| all | 8 | \ | 全部为真 | 1 | +| xori | 9 | \ | 异或归约 | 1 | +| ori | 10 | \ | 或归约 | 1 | +| andi | 11 | \ | 与归约 | 1 | +| none | 0 | \ | 无归约（占位符） | - | + +### 归约操作分类 + +| 类别 | 操作 | 说明 | +|------|------|------| +| 算术归约 | sum, prod | 数值累加/累乘 | +| 极值归约 | max, min | 最大/最小值 | +| 带索引极值 | max_with_index, min_with_index | 最大/最小值及其位置索引 | +| 逻辑归约 | any, all | 逻辑或/与归约 | +| 位归约 | xori, ori, andi | 位运算归约 | + +## unsigned_src 属性 + +`unsigned_src` 指示源操作数是否应被视为无符号数： + +| 值 | 说明 | +|---|------| +| true | 源操作数视为无符号数 | +| false | 源操作数视为有符号数 | + +这对整数类型的归约操作（特别是 max/min）有影响，因为无符号和有符号整数的比较语义不同。 + +## tie_break_left 属性 + +`tie_break_left` 仅对 `max_with_index` 和 `min_with_index` 归约有效。当多个元素具有相同的最大/最小值时： + +| 值 | 说明 | +|---|------| +| true | 返回最左侧（最小索引）的位置 | +| false | 返回最右侧（最大索引）的位置 | +| 未设置 | 使用默认行为 | + +## reduce_dims 参数 + +`reduce_dims` 指定沿哪些维度进行归约： + +- 必须为非空数组 +- 索引必须在 `[0, rank(dst))` 范围内 +- 被归约维度在输出中的大小必须为 1 +- 未被归约维度在输入和输出中的大小必须相同 + +## 数据类型约束 + +| 操作数位置 | 语义 | 支持的元素类型 | 约束来源 | +|-----------|------|--------------|----------| +| $src (idx=0) | 输入 | I1, I8, UI8, I16, UI16, I32, UI32, I64, UI64, F16, F32 | OperElemTypeConstraints<[0, 1], [...]> | +| $dst (idx=1) | 输出 | 与 src 相同 | OperElemTypeConstraints<[0, 1], [...]> | + +## IR 示例 + +### 基本求和归约 + +```mlir +hivm.hir.vreduce ins(%src : memref) outs(%dst : memref<1xf32>) unsigned_src = false reduce_dims = [1] +``` + +### 最大值归约（tensor 语义） + +```mlir +%result = hivm.hir.vreduce ins(%src : tensor) outs(%dst : tensor<1xf32>) unsigned_src = false reduce_dims = [0] -> tensor<1xf32> +``` + +### 带索引的最大值归约 + +```mlir +%val, %idx = hivm.hir.vreduce ins(%src : tensor<23x77xf32>) outs(%dst_val, %dst_idx : tensor<23x1xf32>, tensor<23x1xi32>) unsigned_src = false tie_break_left = true reduce_dims = [1] -> tensor<23x1xf32>, tensor<23x1xi32> +``` + +### 逻辑归约 + +```mlir +%result = hivm.hir.vreduce ins(%src : tensor) outs(%dst : tensor<1xi1>) unsigned_src = false reduce_dims = [0] -> tensor<1xi1> +``` + +## IR 层约束与验证 + +1. **Rank 一致性**：输入和输出向量必须具有相同的 rank +2. **元素类型一致性**：输入和输出必须具有相同的元素类型 +3. **归约维度约束**：输出在归约维度上的大小必须为 1 +4. **reduce_dims 非空**：归约维度数组不能为空 +5. **reduce_dims 范围**：索引必须在 `[0, rank(dst))` 范围内 +6. **max_with_index/min_with_index 输出**：需要两个输出操作数（值和索引） +7. **hasVerifier = 1**：包含自定义验证器 +8. **hasCanonicalizer = 1**：包含自定义规范化器 + +## 与其他 IR 操作的关系 + +| HIVM 操作 | $arith | HFusion 降级 | 说明 | +|-----------|--------|-------------|------| +| vreduce | sum | hfusion.reduce \ | 求和 | +| vreduce | max | hfusion.reduce \ | 最大值 | +| vreduce | min | hfusion.reduce \ | 最小值 | +| vreduce | max_with_index | hfusion.reduce_with_index \ | 带索引最大值 | +| vreduce | min_with_index | hfusion.reduce_with_index \ | 带索引最小值 | + +降级示例： +```mlir +%val, %idx = hivm.hir.vreduce ins(%src : tensor<23x77xf32>) + outs(%v, %i : tensor<23x1xf32>, tensor<23x1xi32>) + unsigned_src = false tie_break_left = true reduce_dims = [1] + -> tensor<23x1xf32>, tensor<23x1xi32> + +%val, %idx = hfusion.reduce_with_index {tie_break_left = true, unsigned_src = false} + %src : tensor<23x77xf32> -> tensor<23x1xf32>, tensor<23x1xi32> +``` + +## 常见问题 + +**Q: vreduce 的输出为什么与输入具有相同的 rank？** +A: 这是 DestinationStyleOpInterface 的要求。归约维度在输出中的大小为 1，但 rank 保持不变。这与 NumPy 的 `keepdims=True` 行为类似。 + +**Q: max_with_index 的索引输出类型是什么？** +A: 索引输出的元素类型为 I32，与输入的数据类型无关。 + +**Q: unsigned_src 对浮点类型有影响吗？** +A: 对浮点类型没有影响，unsigned_src 仅对整数类型的归约操作有意义。 + +**Q: 可以同时对多个维度进行归约吗？** +A: 可以。reduce_dims 是一个数组，可以指定多个归约维度。例如 `reduce_dims = [0, 2]` 同时归约第 0 维和第 2 维。 + +**Q: vreduce 和 vcumsum/vcumprod 有什么区别？** +A: vreduce 将归约维度压缩为大小 1，输出 rank 不变但维度减小。vcumsum/vcumprod 保持所有维度不变，输出与输入形状相同，只是每个元素变为从起始到当前位置的累积值。 + +## 相关文档 + +- Python API：docs_triton_ascend 中的 `tl.reduce()`, `tl.sum()`, `tl.max()` 等文档 +- 源码参考： + - [HIVMVectorOps.td - VReduceOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1129-L1216) + - [HIVMAttrs.td - ReduceOp 枚举](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L596-L631) + - [convert-hivm-to-upstream.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/ExecutionEngine/convert-hivm-to-upstream.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/08-data-movement.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/08-data-movement.md new file mode 100644 index 00000000..4da7bad3 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/08-data-movement.md @@ -0,0 +1,466 @@ +# HIVM 数据搬移操作 + +> 关键词：HIVM, vbrc, vtranspose, vinterleave, vdeinterleave, vflip, vpad, vconcat, vgather + +## 概述 + +HIVM 数据搬移操作负责在向量级别进行数据的重排、广播、拼接、收集等操作。这些操作不执行计算，而是改变数据的布局和形状，是向量计算的重要辅助操作。 + +> Python API 对应：`tl.broadcast_to()`, `tl.trans()`, `tl.flip()`, `tl.pad()`, `tl.concat()`, `tl.gather()` 等。 + +## hir.vbrc — 向量广播 + +### TableGen 定义 + +```tablegen +def VBrcOp : HIVM_VectorOp<"vbrc", + [SameOperandsElementType, + OperElemTypeConstraints< + [0], [I8, UI8, I16, F16, UI16, I32, F8E4M3FN, F8E5M2, F32, UI32, BF16, I64, UI64, I1]>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + CollapsibleConsecutiveTargetDimsTrait, + DeclareOpInterfaceMethods, + InferMaxRankTrait, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + ], []> { + let summary = "Vector Broadcast Op"; + let arguments = (ins AnyType:$src, + TensorOrMemref:$dst, + Optional:$temp_buffer, + DefaultValuedAttr:$broadcast_dims + ); +} +``` + +源码参考：[HIVMVectorOps.td#L1056-L1123](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1056-L1123) + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src | AnyType | 是 | 输入（标量或向量） | 元素类型见下方约束表 | +| $dst | TensorOrMemref | 是 | 输出向量 | 与 src 相同元素类型 | +| $temp_buffer | Optional\ | 否 | 临时缓冲区 | - | +| $broadcast_dims | DenseI64ArrayAttr (默认 {}) | 否 | 广播维度数组 | 标量输入时必须为空 | + +### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vbrc | I8, UI8, I16, F16, UI16, I32, F8E4M3FN, F8E5M2, F32, UI32, BF16, I64, UI64, I1 | OperElemTypeConstraints<[0], [...]> | + +### 约束规则 + +1. 输入和输出必须具有相同的 rank 和元素类型 +2. 对于输入向量，被广播维度的 size 必须为 1 +3. 向量输入时 broadcast_dims 不能为空 +4. 标量输入时 broadcast_dims 必须为空 +5. broadcast_dims 中的索引必须在 `[0, rank(src))` 范围内 +6. I1 类型的输出尾轴需要对齐到 16 + +### IR 示例 + +```mlir +%result = hivm.hir.vbrc ins(%scalar : i32) outs(%dst : tensor<23x77xi32>) -> tensor<23x77xi32> + +%result = hivm.hir.vbrc ins(%src : tensor<1xi32>) outs(%dst : tensor) broadcast_dims = [0] -> tensor + +hivm.hir.vbrc ins(%scalar : f32) outs(%dst : memref) +``` + +--- + +## hir.vtranspose — 维度转置 + +### TableGen 定义 + +```tablegen +def VTransposeOp : HIVM_VectorOp<"vtranspose", + [OperElemTypeConstraints<[0], [AnyI8, AnyI16, AnyI32, F16, BF16, F32, I64, UI64, F8E4M3FN, F8E5M2]>, + DeclareOpInterfaceMethods, + InferMaxRankTrait, + UniformReassociationFlattenTrait, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + ]> { + let summary = "Vector Transpose Op"; + let arguments = (ins TensorOrMemref:$src, + TensorOrMemref:$dst, + Optional:$temp_buffer, + DefaultValuedAttr:$permutation + ); +} +``` + +源码参考：[HIVMVectorOps.td#L1222-L1272](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1222-L1272) + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src | TensorOrMemref | 是 | 输入向量 | 元素类型见下方约束表 | +| $dst | TensorOrMemref | 是 | 输出向量 | 与 src 相同 rank 和元素类型 | +| $temp_buffer | Optional\ | 否 | 临时缓冲区 | - | +| $permutation | DenseI64ArrayAttr (默认 {}) | 否 | 维度排列 | 必须是 range(rank) 的排列 | + +### 语义 + +``` +dim(dst, i) = dim(src, permutation[i]) +``` + +### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vtranspose | AnyI8, AnyI16, AnyI32, F16, BF16, F32, I64, UI64, F8E4M3FN, F8E5M2 | OperElemTypeConstraints<[0], [...]> | + +### IR 示例 + +```mlir +%result = hivm.hir.vtranspose ins(%src : tensor<32x8xf32>) outs(%dst : tensor<8x32xf32>) permutation = [1, 0] -> tensor<8x32xf32> + +hivm.hir.vtranspose ins(%src : memref<32x8xf32>) outs(%dst : memref<8x32xf32>) permutation = [1, 0] + +%result = hivm.hir.vtranspose ins(%src : tensor) outs(%dst : tensor<5x?x10xf32>) permutation = [1, 0, 2] -> tensor<5x?x10xf32> +``` + +--- + +## hir.vinterleave — 交错合并 + +### TableGen 定义 + +```tablegen +def VInterleaveOp : HIVM_VectorOp<"vinterleave", + [SameOperandsElementType, AttrSizedOperandSegments, + HIVMOpSameOperandsAndResultRank, StaticMaxRankTrait<1>, + OperElemTypeConstraints<[0], + [I16, F16, UI16, I32, F32, UI32, BF16, I64, UI64, F8E4M3FN, F8E5M2]>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + DeclareOpInterfaceMethods, + ]> { + let summary = "Vetor Interleave Op"; + let arguments = (ins Variadic:$src, + TensorOrMemref:$dst, + Optional:$temp_buffer, + DefaultValuedAttr:$interleave_channel_nums + ); +} +``` + +源码参考：[HIVMVectorOps.td#L1336-L1377](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1336-L1377) + +### 参数说明 + +| 参数 | 类型 | 必选 | 默认值 | 说明 | 约束 | +|------|------|------|--------|------|------| +| $src | Variadic\ | 是 | - | 输入向量列表 | 所有向量形状相同 | +| $dst | TensorOrMemref | 是 | - | 输出向量 | 最后一维 = src 最后一维 * channel_nums | +| $temp_buffer | Optional\ | 否 | - | 临时缓冲区 | - | +| $interleave_channel_nums | I64Attr | 否 | 2 | 交错通道数 | 必须等于 $src 的数量 | + +### 语义 + +将 N 个张量沿最后一维交错合并。例如，两个张量 `[a0, a1, a2]` 和 `[b0, b1, b2]` 交错后为 `[a0, b0, a1, b1, a2, b2]`。 + +### IR 示例 + +```mlir +%result = hivm.hir.vinterleave ins(%a, %b : tensor<2x16xf32>, tensor<2x16xf32>) outs(%c : tensor<2x32xf32>) interleave_channel_nums = 2 -> tensor<2x32xf32> +``` + +降级到 HFusion：`hfusion.interleave` + +--- + +## hir.vdeinterleave — 交错分离 + +### TableGen 定义 + +```tablegen +def VDeinterleaveOp : HIVM_VectorOp<"vdeinterleave", + [SameOperandsElementType, HIVMOpSameOperandsAndResultRank, + OperElemTypeConstraints<[0], + [I8, I16, F16, UI16, I32, F32, UI32, BF16, I64, UI64]>, + DeclareOpInterfaceMethods, + InferMaxRankTrait, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + DeclareOpInterfaceMethods, + ]> { + let arguments = (ins TensorOrMemref:$src, + Variadic:$dst, + DefaultValuedOptionalAttr:$channel_num, + DefaultValuedOptionalAttr:$index_mode + ); +} +``` + +源码参考：[HIVMVectorOps.td#L1383-L1423](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1383-L1423) + +### 参数说明 + +| 参数 | 类型 | 必选 | 默认值 | 说明 | +|------|------|------|--------|------| +| $src | TensorOrMemref | 是 | - | 输入向量 | +| $dst | Variadic\ | 是 | - | 输出向量列表 | +| $channel_num | I64Attr | 否 | 2 | 通道数 | +| $index_mode | HIVM_DeinterleaveModeAttr | 否 | ALL_CHANNELS | 分离模式 | + +### index_mode 属性 + +| 枚举值 | 值 | 说明 | +|--------|---|------| +| CHANNEL_0 | 0 | 仅提取通道 0 | +| CHANNEL_1 | 1 | 仅提取通道 1 | +| ALL_CHANNELS | 999 | 提取所有通道 | + +### IR 示例 + +```mlir +%result = hivm.hir.vdeinterleave ins(%src : tensor<32xf32>) outs(%dst : tensor<16xf32>) index_mode = -> tensor<16xf32> +``` + +降级到 HFusion：`hfusion.deinterleave %src channel<0>` + +--- + +## hir.vflip — 维度翻转 + +### TableGen 定义 + +```tablegen +def VFlipOp : HIVM_VectorOp<"vflip", + [SameOperandsElementType, StaticMaxRankTrait<1>, + OperElemTypeConstraints<[0, 1], [I8, UI8, I16, I32, UI16, UI32, I64, UI64, F16, F32, BF16]>, + DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + DeclareOpInterfaceMethods + ]> { + let arguments = (ins TensorOrMemref:$src, + TensorOrMemref:$dst, + I64Attr:$flip_axis + ); +} +``` + +源码参考：[HIVMVectorOps.td#L1429-L1457](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1429-L1457) + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| $src | TensorOrMemref | 是 | 输入向量 | +| $dst | TensorOrMemref | 是 | 输出向量 | +| $flip_axis | I64Attr | 是 | 翻转的轴 | + +### IR 示例 + +```mlir +%result = hivm.hir.vflip ins(%src : tensor<10xf32>) outs(%dst : tensor<10xf32>) flip_axis = 0 -> tensor<10xf32> +``` + +--- + +## hir.vpad — 填充 + +### TableGen 定义 + +```tablegen +def VPadOp : HIVM_VectorOp<"vpad", + [HIVMOpSameOperandsAndResultRank, + AttrSizedOperandSegments, NoLibraryFunctionTrait, + DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods]> { + let arguments = (ins TensorOrMemref:$src, + TensorOrMemref:$dst, + AnyType:$pad_value, + Variadic:$low, + Variadic:$high, + DenseI64ArrayAttr:$static_low, + DenseI64ArrayAttr:$static_high + ); +} +``` + +源码参考：[HIVMVectorOps.td#L1502-L1559](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1502-L1559) + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| $src | TensorOrMemref | 是 | 输入张量 | +| $dst | TensorOrMemref | 是 | 输出张量（bufferization） | +| $pad_value | AnyType | 是 | 填充值 | +| $low | Variadic\ | 是 | 每维起始方向的填充长度 | +| $high | Variadic\ | 是 | 每维末尾方向的填充长度 | +| $static_low | DenseI64ArrayAttr | 是 | 静态起始填充长度 | +| $static_high | DenseI64ArrayAttr | 是 | 静态末尾填充长度 | + +### IR 示例 + +```mlir +hivm.hir.vpad ins(%src : tensor<2x16xf32>) outs(%dst : tensor) + low[%first_dim_low, 0] high[%first_dim_high, 0] + pad_value %pad_value : f32 + -> tensor +``` + +--- + +## hir.vconcat — 拼接 + +### TableGen 定义 + +```tablegen +def VConcatOp : HIVM_VectorOp<"vconcat", + [SameOperandsElementType, NoLibraryFunctionTrait, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + DeclareOpInterfaceMethods + ]> { + let arguments = (ins I64Attr:$dim, + Variadic:$src, + TensorOrMemref:$dst + ); +} +``` + +源码参考：[HIVMVectorOps.td#L1565-L1604](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1565-L1604) + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| $dim | I64Attr | 是 | 拼接维度 | +| $src | Variadic\ | 是 | 输入张量列表 | +| $dst | TensorOrMemref | 是 | 输出张量 | + +### 语义 + +沿指定维度拼接多个张量。拼接维度的大小等于所有输入在该维度大小之和，其他维度大小必须相同。 + +### IR 示例 + +```mlir +%result = hivm.hir.vconcat dim(0) ins(%a, %b : tensor<5x?x10xf32>, tensor) outs(%c : tensor) -> tensor + +hivm.hir.vconcat dim(1) ins(%0, %1 : tensor<136x2048xf32>, tensor<136x2048xf32>) outs(%2 : tensor<136x4096xf32>) -> tensor<136x4096xf32> +``` + +降级到上游：`tensor.concat` + +--- + +## hir.vgather — 按索引收集 + +### TableGen 定义 + +```tablegen +def VGatherOp : HIVM_VectorOp<"vgather", + [HIVMOpSameOperandsAndResultRank, + StaticMaxRankTrait<1>, + OperElemTypeConstraints<[0], [I1, I8, I16, UI16, I32, UI32, F16, BF16, F32, F8E4M3FN, F8E5M2]>, + OperElemTypeConstraints<[1], [I32]>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + DeclareOpInterfaceMethods + ]> { + let arguments = (ins TensorOrMemref:$src, + TensorOrMemref:$indices, + TensorOrMemref:$dst, + Optional:$temp_buffer + ); +} +``` + +源码参考：[HIVMVectorOps.td#L1610-L1654](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1610-L1654) + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src | TensorOrMemref | 是 | 源数据 | 元素类型见上方约束 | +| $indices | TensorOrMemref | 是 | 索引向量 | 元素类型 I32 | +| $dst | TensorOrMemref | 是 | 输出向量 | 与 src 相同元素类型 | +| $temp_buffer | Optional\ | 否 | 临时缓冲区 | - | + +### 语义 + +根据索引向量从源数据中收集元素，存储到输出向量中。收集轴为最后一维。 + +``` +dst[i] = src[indices[i]] +``` + +### IR 示例 + +```mlir +%result = hivm.hir.vgather ins(%src : tensor<100xf32>) indices(%idx : tensor<10xi32>) outs(%dst : tensor<10xf32>) -> tensor<10xf32> +``` + +## 数据类型约束汇总 + +| 操作 | 支持的元素类型 | 最大 Rank | +|------|--------------|-----------| +| vbrc | I8, UI8, I16, F16, UI16, I32, F8E4M3FN, F8E5M2, F32, UI32, BF16, I64, UI64, I1 | 推断 | +| vtranspose | AnyI8, AnyI16, AnyI32, F16, BF16, F32, I64, UI64, F8E4M3FN, F8E5M2 | 推断 | +| vinterleave | I16, F16, UI16, I32, F32, UI32, BF16, I64, UI64, F8E4M3FN, F8E5M2 | 1 | +| vdeinterleave | I8, I16, F16, UI16, I32, F32, UI32, BF16, I64, UI64 | 推断 | +| vflip | I8, UI8, I16, I32, UI16, UI32, I64, UI64, F16, F32, BF16 | 1 | +| vpad | 无显式约束（由 pad_value 决定） | - | +| vconcat | SameOperandsElementType | - | +| vgather | 数据: I1, I8, I16, UI16, I32, UI32, F16, BF16, F32, F8E4M3FN, F8E5M2; 索引: I32 | 1 | + +## 常见问题 + +**Q: vbrc 和 Elementwise 操作的 broadcast 属性有什么区别？** +A: vbrc 是独立的广播操作，生成新的张量。Elementwise 操作的 broadcast 属性是 OTF（On-The-Fly）广播，在计算的同时进行广播，不需要额外的广播步骤。 + +**Q: vinterleave 和 vconcat 有什么区别？** +A: vinterleave 是沿最后一维交错合并多个张量（如 `[a0,b0,a1,b1,...]`），而 vconcat 是沿指定维度简单拼接（如 `[a0,a1,...,b0,b1,...]`）。 + +**Q: vgather 的索引范围有约束吗？** +A: 索引值必须在源数据的最后一维大小范围内，即 `0 <= indices[i] < dim(src, last_dim)`。越界访问的行为未定义。 + +## 相关文档 + +- Python API：docs_triton_ascend 中的数据操作文档 +- 源码参考： + - [HIVMVectorOps.td - Data Movement Ops](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1056-L1654) + - [convert-hivm-to-upstream.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/ExecutionEngine/convert-hivm-to-upstream.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/09-cumulative-sort.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/09-cumulative-sort.md new file mode 100644 index 00000000..01d3b95d --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/09-cumulative-sort.md @@ -0,0 +1,320 @@ +# HIVM 累积与排序操作 + +> 关键词：HIVM, vcumsum, vcumprod, vsort, cumulative, sort + +## 概述 + +HIVM 累积与排序操作包括累积求和（`hir.vcumsum`）、累积求积（`hir.vcumprod`）和排序（`hir.vsort`）。累积操作沿指定维度计算从起始到当前位置的累积值，排序操作沿指定维度对元素进行排序。 + +> Python API 对应：`tl.cumsum()`, `tl.cumprod()`, `tl.sort()` 等。 + +## hir.vcumsum — 累积求和 + +### TableGen 定义 + +```tablegen +def VCumsumOp : HIVM_VectorOp<"vcumsum", + [SameOperandsElementType, + StaticMaxRankTrait<2>, + OperElemTypeConstraints<[0], [I1, I8, I16, I32, I64, F16, F32, BF16]>, + DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods + ]> { + let summary = "Vector Cumsum Op"; + let description = [{ + Calculate the cumulative sum of each element along the specified axis of + `src`. Each element along the specified axis in the output of cumsum + contains the sum of all elements from the first element to the current + position in the original `src`. + + Constraints: + 1. The input vector and output vector must have the same rank + and the same element type. + + Arguments: + * `src`: the tensor/memref from which to calculate the cumulative sum + * `dst`: the tensor/memref to store elements + * `cum_dims`: specifies the dimension along which to calculate the + cumulative sum. + + Examples: + ```mlir + hivm.hir.vcumsum ins(%src : memref) outs(%dst : memref) cum_dims : [0] + %result = hivm.hir.vcumsum ins(%src : tensor) outs(%dst : tensor) cum_dims : [0] -> tensor + ``` + }]; + let arguments = (ins TensorOrMemref:$src, + TensorOrMemref:$dst, + ConfinedAttr]>:$cum_dims, + BoolAttr:$reverse + ); + let results = (outs Variadic:$result); + let assemblyFormat = [{ + attr-dict `ins` `(` $src `:` type($src) `)` + `outs` `(` $dst `:` type($dst) `)` + `cum_dims` `=` $cum_dims + `reverse` `=` $reverse + (`->` type($result)^)? + }]; + let hasCanonicalizer = 1; + let hasVerifier = 1; +} +``` + +源码参考：[HIVMVectorOps.td#L1719-L1772](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1719-L1772) + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src | TensorOrMemref | 是 | 输入向量 | 元素类型 I1/I8/I16/I32/I64/F16/F32/BF16 | +| $dst | TensorOrMemref | 是 | 输出向量 | 与 src 相同 rank 和元素类型 | +| $cum_dims | DenseI64ArrayAttr | 是 | 累积维度 | 严格排序，Confined 约束 | +| $reverse | BoolAttr | 是 | 是否反向累积 | - | + +### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vcumsum | I1, I8, I16, I32, I64, F16, F32, BF16 | OperElemTypeConstraints<[0], [I1, I8, I16, I32, I64, F16, F32, BF16]> | + +### 语义 + +正向累积（reverse = false）： +``` +dst[i] = sum(src[0:i+1]) // 沿 cum_dims 维度 +``` + +反向累积（reverse = true）： +``` +dst[i] = sum(src[i:end]) // 沿 cum_dims 维度 +``` + +### IR 示例 + +```mlir +hivm.hir.vcumsum ins(%src : memref<5x?x10xi32>) outs(%dst : memref<5x?x10xi32>) cum_dims = [1] reverse = false + +%result = hivm.hir.vcumsum ins(%src : tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) cum_dims = [0] reverse = false -> tensor<5x?x10xf32> +``` + +降级到上游：`linalg.generic` 包含 `arith.addi`（整数）或 `arith.addf`（浮点） + +--- + +## hir.vcumprod — 累积求积 + +### TableGen 定义 + +```tablegen +def VCumprodOp : HIVM_VectorOp<"vcumprod", + [SameOperandsElementType, + StaticMaxRankTrait<1>, + OperElemTypeConstraints<[0], [I1, I8, I16, I32, I64, F16, F32, BF16]>, + DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods + ]> { + let summary = "Vector Cumprod Op"; + let description = [{ + Calculate the cumulative product of each element along the specified axis + of `src`. Each element along the specified axis in the output of cumprod + contains the product of all elements from the first element to the current + position in the original `src`. + }]; + let arguments = (ins TensorOrMemref:$src, + TensorOrMemref:$dst, + ConfinedAttr]>:$cum_dims, + BoolAttr:$reverse + ); +} +``` + +源码参考：[HIVMVectorOps.td#L1660-L1713](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1660-L1713) + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src | TensorOrMemref | 是 | 输入向量 | 元素类型 I1/I8/I16/I32/I64/F16/F32/BF16 | +| $dst | TensorOrMemref | 是 | 输出向量 | 与 src 相同 rank 和元素类型 | +| $cum_dims | DenseI64ArrayAttr | 是 | 累积维度 | 严格排序，Confined 约束 | +| $reverse | BoolAttr | 是 | 是否反向累积 | - | + +### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vcumprod | I1, I8, I16, I32, I64, F16, F32, BF16 | OperElemTypeConstraints<[0], [I1, I8, I16, I32, I64, F16, F32, BF16]> | + +### 语义 + +正向累积（reverse = false）： +``` +dst[i] = prod(src[0:i+1]) // 沿 cum_dims 维度 +``` + +反向累积（reverse = true）： +``` +dst[i] = prod(src[i:end]) // 沿 cum_dims 维度 +``` + +### IR 示例 + +```mlir +%result = hivm.hir.vcumprod ins(%src : tensor<5x?x10xf32>) outs(%dst : tensor<5x?x10xf32>) cum_dims = [0] reverse = false -> tensor<5x?x10xf32> + +hivm.hir.vcumprod ins(%src : memref<5x?x10xi32>) outs(%dst : memref<5x?x10xi32>) cum_dims = [1] reverse = false +``` + +降级到上游：`linalg.generic` 包含 `arith.muli`（整数）或 `arith.mulf`（浮点） + +--- + +## hir.vsort — 排序 + +### TableGen 定义 + +```tablegen +def VSortOp : HIVM_VectorOp<"vsort", + [StaticMaxRankTrait<1>, + AttrSizedOperandSegments, + OperElemTypeConstraints<[0, 1], [F16, F32, I32, I64]>, + DeclareOpInterfaceMethods + ]> { + let summary = "Vector Sort Op"; + let description = [{ + Sort the sorting axis of `src` in ascending or descending order, and output + the sorted value and the index corresponding to the value. + + Constraints: + 1. The input vector and output vector must have the same rank. + 2. Currently only tail axis sorting is supported. + + Arguments: + * `src`: the tensor/memref from which to be sorted + * `dst_value`: the tensor/memref to store the sorted value + * `dst_index`: the tensor/memref to store the index corresponding to dst_value + * `descending`: determines whether to sort in ascending or descending + order. The default is false, which means ascending order + * `sort_axis`: Axis to be sorted + }]; + let arguments = (ins TensorOrMemref:$src, + Variadic:$dst, + Optional:$temp_buffer, + DefaultValuedOptionalAttr:$descending, + DefaultValuedAttr:$sort_axis + ); +} +``` + +源码参考：[HIVMVectorOps.td#L1778-L1839](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1778-L1839) + +### 参数说明 + +| 参数 | 类型 | 必选 | 默认值 | 说明 | 约束 | +|------|------|------|--------|------|------| +| $src | TensorOrMemref | 是 | - | 输入向量 | 元素类型 F16/F32/I32/I64 | +| $dst | Variadic\ | 是 | - | 输出（排序值 + 索引） | 2 个输出 | +| $temp_buffer | Optional\ | 否 | - | 临时缓冲区 | - | +| $descending | BoolAttr | 否 | false | 是否降序排序 | - | +| $sort_axis | I64Attr | 否 | -1 | 排序轴 | 目前仅支持尾轴 | + +### 数据类型约束 + +| 操作数位置 | 语义 | 支持的元素类型 | 约束来源 | +|-----------|------|--------------|----------| +| $src (idx=0) | 输入 | F16, F32, I32, I64 | OperElemTypeConstraints<[0, 1], [F16, F32, I32, I64]> | +| $dst (idx=1) | 排序值输出 | 与 src 相同 | OperElemTypeConstraints<[0, 1], [...]> | + +注意：索引输出的元素类型为 I32。 + +### 语义 + +``` +dst_value[i] = src[sorted_indices[i]] +dst_index[i] = sorted_indices[i] +``` + +- `descending = false`：升序排序（默认） +- `descending = true`：降序排序 + +### IR 示例 + +```mlir +hivm.hir.vsort ins(%src : memref) outs(%dst : memref) descending = true sort_axis = 0 + +%result = hivm.hir.vsort ins(%src : tensor) outs(%dst : tensor) descending = true sort_axis = 0 -> tensor +``` + +### 特殊接口 + +```cpp +Value getDstValue(); // 获取排序值输出 +Value getDstIndex(); // 获取排序索引输出 +int64_t getSignedSortAxis(); // 获取有符号排序轴 +``` + +## 数据类型约束汇总 + +| 操作 | 支持的元素类型 | 最大 Rank | 输出数量 | reverse 属性 | descending 属性 | +|------|--------------|-----------|---------|-------------|----------------| +| vcumsum | I1, I8, I16, I32, I64, F16, F32, BF16 | 2 | 1 | 是 | - | +| vcumprod | I1, I8, I16, I32, I64, F16, F32, BF16 | 1 | 1 | 是 | - | +| vsort | F16, F32, I32, I64 | 1 | 2 | - | 是 | + +## IR 层约束与验证 + +### 通用约束 + +1. **Rank 一致性**：输入和输出必须具有相同的 rank +2. **元素类型一致性**：输入和输出必须具有相同的元素类型（vsort 索引输出除外） +3. **cum_dims 约束**：必须严格排序（`DenseArrayStrictlySorted`） + +### vsort 特有约束 + +1. **仅支持尾轴排序**：当前硬件实现仅支持沿最后一个轴排序 +2. **两个输出**：排序值和排序索引 +3. **sort_axis 默认值 -1**：表示最后一个轴 +4. **hasVerifier = 1**：包含自定义验证器 + +## 与其他 IR 操作的关系 + +| HIVM 操作 | 上游降级 | 说明 | +|-----------|---------|------| +| vcumsum | linalg.generic { arith.addi/addf } | 累积求和 | +| vcumprod | linalg.generic { arith.muli/mulf } | 累积求积 | +| vsort | - | 无直接上游对应 | + +## 常见问题 + +**Q: vcumsum 和 vreduce\ 有什么区别？** +A: vcumsum 保持输出与输入相同的形状，每个位置存储从起始到当前位置的累积和。vreduce\ 将归约维度压缩为 1，输出形状与输入不同。 + +**Q: vsort 为什么只支持尾轴排序？** +A: 这是当前硬件实现的限制。硬件排序指令仅支持沿最后一维排序。如果需要沿其他维度排序，需要先转置再排序再转置回来。 + +**Q: cum_dims 的 DenseArrayStrictlySorted 约束是什么意思？** +A: cum_dims 数组中的索引必须严格递增排列，不允许重复。例如 [0, 2] 合法，[2, 0] 或 [0, 0] 不合法。 + +**Q: reverse 属性在累积操作中的实际用途是什么？** +A: reverse = true 时，累积从数组末尾向起始方向计算。例如，对于序列 [1, 2, 3]，正向累积和为 [1, 3, 6]，反向累积和为 [6, 5, 3]。 + +## 相关文档 + +- Python API：docs_triton_ascend 中的 `tl.cumsum()`, `tl.sort()` 文档 +- 源码参考： + - [HIVMVectorOps.td - VCumsumOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1719-L1772) + - [HIVMVectorOps.td - VCumprodOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1660-L1713) + - [HIVMVectorOps.td - VSortOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1778-L1839) + - [convert-hivm-to-upstream.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/ExecutionEngine/convert-hivm-to-upstream.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/10-special-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/10-special-ops.md new file mode 100644 index 00000000..70d7948c --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/10-special-ops.md @@ -0,0 +1,288 @@ +# HIVM 特殊操作 + +> 关键词：HIVM, varange, vmulextended, vmulext, arange, mul_extended + +## 概述 + +HIVM 特殊操作包括范围序列生成（`hir.varange`）、扩展乘法（`hir.vmulextended`）和乘法高32位（`hir.vmulext`）。这些操作不属于标准的逐元运算分类，但在特定场景中不可或缺。 + +> Python API 对应：`tl.arange()`, 以及 Triton 中的扩展精度运算。 + +## hir.varange — 范围序列生成 + +### TableGen 定义 + +```tablegen +def VArangeOp + : HIVM_VectorOp< + "varange", [AttrSizedOperandSegments, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0], [I16, I32, F16, F32, I64]>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods]> { + let summary = "Vector Arange Op"; + let description = [{ + Fill a vector with range 0,1,2... based on strides and offset. + e.g. offset = 1, strides = [1, 2], tensor/memref shape = [2x4xi32], + the result is [[1, 3, 5, 7, + 2, 4, 6, 8]]. + + Constraints: + 1. Must have at least one stride. + 2. Default offset is 0. + + Examples: + ```mlir + hivm.hir.varange offset[%o] strides[%s0, %s1] outs(%dst : memref<32xf32>) + %result = hivm.hir.varange offset[%o] strides[%s0, %s1] outs(%dst : tensor<32xf32>) + -> tensor<32xf32> + ``` + }]; + let arguments = (ins TensorOrMemref:$dst, + Optional:$offset, + Variadic:$strides + ); + let results = (outs Optional:$result); + let assemblyFormat = [{ + attr-dict + (`offset` `[` $offset^ `]`)? + `strides` `[` $strides `]` + `outs` `(` $dst `:` type($dst) `)` + (`->` type($result)^)? + }]; +} +``` + +源码参考：[HIVMVectorOps.td#L1278-L1330](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1278-L1330) + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $dst | TensorOrMemref | 是 | 输出向量（同时指定形状） | 元素类型 I16/I32/F16/F32/I64 | +| $offset | Optional\ | 否 | 起始偏移量 | 默认为 0 | +| $strides | Variadic\ | 是 | 每维步长 | 至少一个步长 | + +### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| varange | I16, I32, F16, F32, I64 | OperElemTypeConstraints<[0], [I16, I32, F16, F32, I64]> | + +### 语义 + +varange 根据偏移量和步长生成一个范围序列，填充到输出张量中。 + +对于多维张量，每个维度有一个步长。计算公式： + +``` +value[i0, i1, ..., in] = offset + i0 * strides[0] + i1 * strides[1] + ... + in * strides[n] +``` + +示例（来自 TableGen 描述）： +- offset = 1, strides = [1, 2], shape = [2x4] +- 结果：[[1, 3, 5, 7], [2, 4, 6, 8]] + +### IR 示例 + +```mlir +%result = hivm.hir.varange offset[] strides[%c0, %c3, %c2] outs(%dst : tensor<5x?x10xi64>) -> tensor<5x?x10xi64> + +hivm.hir.varange offset[%c3] strides[%c1, %c1, %c1] outs(%dst : memref<5x?x10xi32>) + +%result = hivm.hir.varange strides[%c1] outs(%dst : tensor<1011xi32>) -> tensor<1011xi32> +``` + +降级到 HFusion：`hfusion.arange` + +--- + +## hir.vmulextended — 扩展乘法 + +### TableGen 定义 + +```tablegen +def VMulextendedOp : HIVM_VectorOp<"vmulextended", + [AttrSizedOperandSegments, HIVMOpSameOperandsAndResultRank, + StaticMaxRankTrait<1>, OperElemTypeConstraints<[0], [I16]>, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods, + UniformReassociationFlattenTrait, + DeclareOpInterfaceMethods + ]> { + let summary = "Vector Mulextended Op"; + let description = [{ + Do vmul on two tensors. Get both high and low 16-bits. + }]; + let arguments = (ins Variadic:$src, + Variadic:$dst, + Optional:$temp_buffer + ); +} +``` + +源码参考：[HIVMVectorOps.td#L1463-L1496](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1463-L1496) + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src | Variadic\ | 是 | 两个输入向量 | 元素类型 I16 | +| $dst | Variadic\ | 是 | 两个输出向量（高位和低位） | - | +| $temp_buffer | Optional\ | 否 | 临时缓冲区 | - | + +### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vmulextended | I16 | OperElemTypeConstraints<[0], [I16]> | + +### 语义 + +vmulextended 对两个 I16 向量执行乘法，产生两个 I16 输出：乘法结果的高 16 位和低 16 位。 + +``` +product = src0[i] * src1[i] // 32-bit result +dst_high[i] = (product >> 16) & 0xFFFF // 高 16 位 +dst_low[i] = product & 0xFFFF // 低 16 位 +``` + +### IR 示例 + +```mlir +%high, %low = hivm.hir.vmulextended ins(%a, %b : tensor<32xi16>, tensor<32xi16>) outs(%dh, %dl : tensor<32xi16>, tensor<32xi16>) -> tensor<32xi16>, tensor<32xi16> +``` + +--- + +## hir.vmulext — 乘法高32位 + +### TableGen 定义 + +```tablegen +def VMulExtOp : HIVM_ElementwiseBinaryOp<"vmulext", + [SameOperandsElementType, StaticMaxRankTrait<3>, + OperElemTypeConstraints<[0, 1], [I32]>, + VectorOnlyTrait<0>, + DeclareOpInterfaceMethods + ]> { + let summary = [{ + Elementwise Binary Vector Multiplication that Calculates + the Most Significant 32-bits. + }]; + let description = baseClassDescription # [{ + Additional constraints: + 1. The input/init operands and result have the same element type. + 2. Support Vector-Vector operation. + }]; +} +``` + +源码参考：[HIVMVectorOps.td#L578-L594](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L578-L594) + +### 参数说明 + +| 参数 | 类型 | 必选 | 说明 | 约束 | +|------|------|------|------|------| +| $src[0] | AnyType | 是 | 第一个输入向量 | VectorOnly, 元素类型 I32 | +| $src[1] | AnyType | 是 | 第二个输入向量/标量 | 元素类型 I32 | +| $dst | Variadic\ | 是 | 输出向量 | 元素类型 I32 | + +注意：vmulext 使用默认的 ElementwiseNaryOp 参数（无 temp_buffer, 无 OTF 广播/转置）。 + +### 数据类型约束 + +| 操作 | 支持的元素类型 | 约束来源 | +|------|--------------|----------| +| vmulext | I32 | OperElemTypeConstraints<[0, 1], [I32]> | + +### 语义 + +vmulext 计算两个 I32 值乘积的最高有效 32 位。这等价于 64 位乘法结果的高 32 位。 + +``` +product_64 = src0[i] * src1[i] // 64-bit result +dst[i] = (product_64 >> 32) & 0xFFFFFFFF // 高 32 位 +``` + +### IR 示例 + +```mlir +%result = hivm.hir.vmulext ins(%a, %b : tensor<32xi32>, tensor<32xi32>) outs(%dst : tensor<32xi32>) -> tensor<32xi32> +``` + +## vmulextended 与 vmulext 的区别 + +| 特性 | vmulextended | vmulext | +|------|-------------|---------| +| 输入类型 | I16 | I32 | +| 输出数量 | 2（高位 + 低位） | 1（仅高位） | +| 操作类别 | HIVM_VectorOp（独立操作） | HIVM_ElementwiseBinaryOp | +| 最大 Rank | 1 | 3 | +| 乘法宽度 | I16 * I16 → I32 | I32 * I32 → I64 | +| 输出内容 | 高 16 位 + 低 16 位 | 高 32 位 | +| 标量支持 | 否 | 是（ImplByScalarOpInterface） | + +## 数据类型约束汇总 + +| 操作 | 支持的元素类型 | 最大 Rank | 输出数量 | +|------|--------------|-----------|---------| +| varange | I16, I32, F16, F32, I64 | 3 | 1 | +| vmulextended | I16 | 1 | 2 | +| vmulext | I32 | 3 | 1 | + +## IR 层约束与验证 + +### varange + +1. **至少一个步长**：$strides 不能为空 +2. **步长数量与 rank**：步长数量应与输出张量的 rank 匹配 +3. **默认偏移**：offset 可选，默认为 0 +4. **hasVerifier = 1**：包含自定义验证器 + +### vmulextended + +1. **仅支持 I16**：输入和输出元素类型均为 I16 +2. **两个输出**：高位结果和低位结果 +3. **最大 Rank 1**：仅支持 1 维 + +### vmulext + +1. **仅支持 I32**：输入和输出元素类型均为 I32 +2. **VectorOnly**：第一个输入必须为向量 +3. **ImplByScalarOpInterface**：第二个输入可以是标量 + +## 与其他 IR 操作的关系 + +| HIVM 操作 | 上游降级 | HFusion 降级 | 说明 | +|-----------|---------|-------------|------| +| varange | - | hfusion.arange | 范围序列 | +| vmulextended | - | - | 扩展乘法（无直接上游对应） | +| vmulext | - | - | 乘法高32位（无直接上游对应） | + +## 常见问题 + +**Q: varange 的步长数量必须等于张量的 rank 吗？** +A: 是的。每个维度对应一个步长值，因此步长数量应该等于输出张量的 rank。 + +**Q: varange 支持浮点步长吗？** +A: varange 的步长参数类型为 Index（整数），但输出可以是浮点类型。步长值在计算时会被转换为输出类型。 + +**Q: vmulextended 和 vmulext 的典型应用场景是什么？** +A: 这些操作用于实现扩展精度运算。当标准乘法的位宽不够时，可以使用 vmulextended 获取完整的乘法结果（高位+低位），或使用 vmulext 获取高位部分。这在定点数运算和量化场景中特别有用。 + +**Q: vmulext 为什么继承自 ElementwiseBinaryOp 而不是 VectorOp？** +A: vmulext 的语义与二元逐元操作一致（两个输入，一个输出），且支持标量操作数。将其归类为 ElementwiseBinaryOp 可以复用逐元操作的通用基础设施。 + +## 相关文档 + +- Python API：docs_triton_ascend 中的 `tl.arange()` 文档 +- 源码参考： + - [HIVMVectorOps.td - VArangeOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1278-L1330) + - [HIVMVectorOps.td - VMulextendedOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L1463-L1496) + - [HIVMVectorOps.td - VMulExtOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMVectorOps.td#L578-L594) + - [convert-hivm-to-upstream.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/ExecutionEngine/convert-hivm-to-upstream.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/11-scalar-lowering.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/11-scalar-lowering.md new file mode 100644 index 00000000..f44876f3 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/11-scalar-lowering.md @@ -0,0 +1,380 @@ +# HIVM 向量操作标量降级 + +> 关键词：HIVM, ImplByScalarOpInterface, shouldLowerToScalarLoops, lowerToLoops, scalar, i64, SIMT VF + +## 概述 + +HIVM 向量操作在编译过程中，部分操作在特定条件下会被**降级为标量循环**（Scalar Lowering）执行，而非使用硬件向量指令。这是因为 AscendNPU 的向量计算单元对某些数据类型和操作组合缺乏硬件级支持，编译器通过 `ImplByScalarOpInterface` 接口自动将这类操作退化为逐元素的标量循环。 + +标量降级会显著影响性能：向量指令可以一次处理多个数据元素，而标量循环逐元素执行，吞吐量大幅降低。因此，在编写 Triton 算子时，理解哪些操作在什么条件下会被标量降级，对于性能优化至关重要。 + +> 本文档面向辅助人类编写和优化 Triton 算子的 AI agent，帮助识别和避免标量降级导致的性能瓶颈。 + +## 标量降级机制 + +### 接口定义 + +`ImplByScalarOpInterface` 定义在 [ImplByScalarOpInterface.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/Interfaces/ImplByScalarOpInterface.td)，声明了两个关键方法： + +| 方法 | 说明 | +|------|------| +| `shouldLowerToScalarLoops()` | 判断当前操作是否应该降级为标量循环 | +| `lowerToLoops(RewriterBase &b)` | 执行实际的标量循环降级 | + +### 降级流程 + +``` +HIVM 向量操作 + │ + ├── shouldLowerToScalarLoops() 返回 true？ + │ │ + │ ├── 是 → HIVMLowerToLoopsPass 调用 lowerToLoops() + │ │ │ + │ │ └── 创建 scf.for 嵌套循环 + │ │ 循环体内逐元素执行 arith 标量操作 + │ │ + │ └── 否 → 保持向量操作，后续由硬件向量指令执行 + │ │ + │ ├── Execution Engine 路径：→ hfusion 操作 + │ └── TritonGPU 路径：→ arith 向量操作 +``` + +### 共同前提条件 + +所有标量降级的共同前提是 **`hasPureBufferSemantics()` 为 true**，即操作必须具有纯 buffer 语义（memref 语义，而非 tensor 语义）。在 bufferization 阶段之后，大部分操作都会满足此条件。 + +### 降级实现 + +当 `shouldLowerToScalarLoops()` 返回 true 时，[LowerToLoops.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/ImplByScalarOpInterface/LowerToLoops.cpp) 会将向量操作分解为嵌套 `scf.for` 循环，循环体内逐元素执行对应的 `arith` 标量操作： + +| HIVM 操作 | 标量降级后的 arith 操作 | +|-----------|----------------------| +| VAddOp | arith.AddIOp / arith.AddFOp | +| VSubOp | arith.SubIOp | +| VMulOp | arith.MulIOp / arith.MulFOp | +| VMinOp | arith.MinSIOp | +| VMaxOp | arith.MaxSIOp | +| VAbsOp | math.AbsIOp | +| VShLOp | arith.ShLIOp | +| VShROp | arith.ShRSIOp | +| VCmpOp | arith.CmpIOp + arith.ExtUIOp (i1→i8) | +| VMulExtOp | arith.MulUIExtendedOp | +| VCumsumOp | arith.AddIOp / arith.AddFOp（含前一轮累积值） | +| VCumprodOp | arith.MulIOp / arith.MulFOp（含前一轮累积值） | +| VReduceOp | 依 reduceOp 不同：MinSI/MaxSI/AddI/MulI/XOrI 及 argmin/argmax 复合逻辑 | + +### 对其他编译 Pass 的影响 + +标量降级判断不仅影响 `HIVMLowerToLoopsPass`，还会影响以下 Pass 的行为： + +| Pass | 影响 | +|------|------| +| [OptMemPlanForPipeline](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/Transforms/OptMemPlanForPipeline.cpp#L18-L23) | 标量降级的操作使用不同的 buffer 规划策略 | +| [AdjustAlignUtil](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/Transforms/AlignBuffer/AdjustAlignUtil.cpp#L414-L419) | 标量降级的操作跳过 stride 对齐调整（标量操作不需要向量对齐） | +| [HIVMDecomposeOp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/Transforms/HIVMDecomposeOp.cpp#L655-L712) | 对 VCmpOp 做特殊的 i1→i8 输出转换预处理 | + +## 第 1 组：通用算术操作 + +**适用操作**：VAddOp, VSubOp, VMulOp, VMinOp, VMaxOp, VAbsOp, VShLOp, VShROp, VInterleaveOp, VDeinterleaveOp + +这 10 个操作共享相同的 `shouldLowerToScalarLoops` 逻辑，通过 `ENABLE_DEFAULT_OP_SHOULD_LOWER_TO_SCALAR_LOOPS_IMPL` 宏生成。 + +源码参考：[ShouldLowerToScalarLoops.cpp:56-64](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/ImplByScalarOpInterface/ShouldLowerToScalarLoops.cpp#L56-L64) + +### 降级条件 + +``` +hasPureBufferSemantics() == true +AND (isSIMTVF() == true OR elemType == i64) +``` + +| 条件 | 说明 | +|------|------| +| hasPureBufferSemantics() | 必须具有纯 buffer 语义 | +| isSIMTVF() | 操作处于 SIMT VF 模式（带有 `VFMode::SIMT` 属性） | +| elemType.isInteger(64) | 第一个操作数的元素类型为 i64 | + +### 降级判断矩阵 + +| 元素类型 | SIMD VF 模式 | SIMT VF 模式 | +|---------|-------------|-------------| +| f16 | 不降级 ✅ | **降级** ⚠️ | +| f32 | 不降级 ✅ | **降级** ⚠️ | +| i8 | 不降级 ✅ | **降级** ⚠️ | +| i16 | 不降级 ✅ | **降级** ⚠️ | +| i32 | 不降级 ✅ | **降级** ⚠️ | +| i64 | **降级** ⚠️ | **降级** ⚠️ | + +### 各操作受影响的数据类型 + +| 操作 | IR 支持的元素类型 | 会被标量降级的类型 | +|------|-----------------|------------------| +| VAddOp | I8, I16, I32, F16, F32, I64 | I64（SIMD）/ 全部（SIMT） | +| VSubOp | I8, I16, I32, F16, F32, I64 | I64（SIMD）/ 全部（SIMT） | +| VMulOp | I16, I32, F16, F32, I64 | I64（SIMD）/ 全部（SIMT） | +| VMinOp | I16, I32, F16, F32, I64 | I64（SIMD）/ 全部（SIMT） | +| VMaxOp | I16, I32, F16, F32, I64 | I64（SIMD）/ 全部（SIMT） | +| VAbsOp | F16, F32, I8, I16, I32, I64 | I64（SIMD）/ 全部（SIMT） | +| VShLOp | I16, I32, I64 | I64（SIMD）/ 全部（SIMT） | +| VShROp | I16, I32, I64 | I64（SIMD）/ 全部（SIMT） | +| VInterleaveOp | I16, F16, I32, F32, BF16, I64 等 | I64（SIMD）/ 全部（SIMT） | +| VDeinterleaveOp | I8, I16, F16, I32, F32, BF16, I64 等 | I64（SIMD）/ 全部（SIMT） | + +### 优化建议 + +- **避免 i64 类型的算术运算**：i64 是最常见的标量降级触发条件。如果精度允许，优先使用 i32 类型 +- **SIMT VF 模式下所有算术操作都会降级**：SIMT 模式下向量操作退化为标量循环是设计如此，因为 SIMT 模式本身就是标量线程模型 +- **浮点运算不受影响**（SIMD 模式下）：f16/f32 的加减乘、最大最小值在 SIMD 模式下均走向量路径 + +## 第 2 组：比较操作（VCmpOp） + +源码参考：[ShouldLowerToScalarLoops.cpp:92-114](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/ImplByScalarOpInterface/ShouldLowerToScalarLoops.cpp#L92-L114) + +### 降级条件 + +``` +hasPureBufferSemantics() == true +AND src[0] 是 MemRefType 或 TensorType +AND elemType 是整数类型 +AND (elemType != i32 OR compare_mode ∉ {EQ, NE}) +``` + +### 降级判断矩阵 + +| 元素类型 | EQ | NE | LT | GT | LE | GE | +|---------|----|----|----|----|----|-----| +| f16 | 不降级 ✅ | 不降级 ✅ | 不降级 ✅ | 不降级 ✅ | 不降级 ✅ | 不降级 ✅ | +| f32 | 不降级 ✅ | 不降级 ✅ | 不降级 ✅ | 不降级 ✅ | 不降级 ✅ | 不降级 ✅ | +| i8 | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | +| i16 | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | +| i32 | 不降级 ✅ | 不降级 ✅ | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | +| i64 | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | **降级** ⚠️ | + +### 特殊处理 + +VCmpOp 标量降级时有额外的 i1→i8 扩展处理： + +1. **HIVMDecomposeOp 预处理**：将 vcmp 的输出从 i1 转为 i8 临时 buffer（因为 store 不支持 i1），之后用 `VCastOp` 转回 i1 +2. **LowerToLoops 后处理**：`arith.CmpIOp` 产生 i1 结果，通过 `arith.ExtUIOp` 零扩展为 i8 后再存储 + +### 优化建议 + +- **浮点比较始终走向量路径**：f16/f32 的所有比较模式都由硬件向量指令执行，性能最优 +- **i32 的相等/不等比较走向量路径**：i32 的 EQ/NE 是唯一能走向量路径的整数比较 +- **避免整数大小比较**：i32 的 LT/GT/LE/GE 以及 i8/i16/i64 的所有比较都会标量降级 +- **整数比较的替代策略**：如果业务逻辑允许，将整数比较转为浮点比较（先 cast 再 compare），可避免标量降级 +- **vcmp + vsel 模式**：vcmp 的典型使用模式是与 vsel 配合，整数比较的标量降级会使整个条件选择链路性能下降 + +## 第 3 组：扩展乘法（VMulExtOp） + +源码参考：[ShouldLowerToScalarLoops.cpp:120-126](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/ImplByScalarOpInterface/ShouldLowerToScalarLoops.cpp#L120-L126) + +### 降级条件 + +``` +hasPureBufferSemantics() == true +AND (elemType == i32 OR elemType == i64) +``` + +### 降级判断矩阵 + +| 元素类型 | 是否降级 | +|---------|---------| +| i32 | **降级** ⚠️ | +| i64 | **降级** ⚠️ | + +注意：VMulExtOp 的 IR 定义仅支持 I32 类型，因此实际上 **vmulext 在所有情况下都会被标量降级**。 + +### 优化建议 + +- **vmulext 始终走标量路径**：该操作没有向量硬件支持，应尽量避免使用 +- **替代方案**：如果需要高 32 位乘法结果，考虑使用 vmulextended（I16 输入）或手动拆分乘法 + +## 第 4 组：累积操作（VCumsumOp, VCumprodOp） + +源码参考：[ShouldLowerToScalarLoops.cpp:22-48](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/ImplByScalarOpInterface/ShouldLowerToScalarLoops.cpp#L22-L48) + +### 降级条件 + +``` +hasPureBufferSemantics() == true +AND cumDims.size() == 1 +AND (elemType == i64 OR flattenedCumDims[0] == flattenedRank - 1) +``` + +### 降级判断详解 + +| 条件 | 说明 | +|------|------| +| cumDims.size() == 1 | 累积维度只能有 1 个，多个累积维度不降级 | +| elemType == i64 | 目标元素类型为 i64 时直接降级 | +| flattenedCumDims[0] == flattenedRank - 1 | 累积维度在 flatten 后是最后一个维度时降级 | + +### 降级判断矩阵 + +| 元素类型 | 累积维度 = 最后维度 | 累积维度 ≠ 最后维度 | 多个累积维度 | +|---------|-------------------|-------------------|------------| +| f16/f32/bf16 | **降级** ⚠️ | 不降级 ✅ | 不降级 ✅ | +| i8/i16/i32 | **降级** ⚠️ | 不降级 ✅ | 不降级 ✅ | +| i64 | **降级** ⚠️ | **降级** ⚠️ | 不降级 ✅ | + +### 优化建议 + +- **避免 i64 累积操作**：i64 类型的累积操作无论维度如何都会标量降级 +- **注意累积维度位置**：非 i64 类型下，累积维度为最后维度时会触发降级。如果可能，调整数据布局使累积维度不在最后 +- **多个累积维度不降级**：当 cumDims > 1 时不会标量降级，但这种情况可能触发其他限制 + +## 第 5 组：归约操作（VReduceOp） + +源码参考：[ShouldLowerToScalarLoops.cpp:132-281](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/ImplByScalarOpInterface/ShouldLowerToScalarLoops.cpp#L132-L281) + +VReduceOp 的标量降级逻辑是最复杂的，取决于硬件架构类型和归约操作类型。 + +### 降级条件 + +``` +hasPureBufferSemantics() == true +AND shouldVReduceOpDecomposeToScalarImpl() == true +``` + +### Reg-based 架构（A5 代：Ascend310B, Ascend950） + +Reg-based（寄存器基）架构的核间同步通过寄存器级指令（SetFlag/WaitFlag）实现，在归约操作上有更好的硬件向量支持。 + +| reduceOp | 降级条件 | +|----------|---------| +| max_with_index | 内存访问对齐不合法时降级 | +| min_with_index | 内存访问对齐不合法时降级 | +| 其他（sum/prod/max/min/xori 等） | **不降级** ✅ | + +内存访问对齐合法性由 [isLegalAccessAlignment](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/ImplByScalarOpInterface/ShouldLowerToScalarLoops.cpp#L139-L219) 判断：检查操作数和结果的步幅布局是否满足对齐要求。如果缺少 `StridedLayoutAttr`，默认判定为不合法，需要降级。 + +### Mem-based 架构（A2/A3 代：Ascend910B, Ascend910_93） + +Mem-based（内存基）架构的核间同步依赖 FFTS（Fast Flag Transmit Storage）内存机制，归约操作的标量降级条件更多。 + +| reduceOp | 降级条件 | +|----------|---------| +| sum / prod / max / min / xori | 元素类型为 i64 | +| max_with_index / min_with_index | i64/i32/i16 → 降级；f16/f32/bf16 且 flatten 后 rank > 2 → 降级；其他 → 不降级 | +| any / all / ori / andi / none | **不降级** ✅ | + +### 降级判断矩阵（Mem-based 架构） + +#### 基本归约（sum/prod/max/min/xori） + +| 元素类型 | 是否降级 | +|---------|---------| +| f16 | 不降级 ✅ | +| f32 | 不降级 ✅ | +| i8/i16/i32 | 不降级 ✅ | +| i64 | **降级** ⚠️ | + +#### 带索引极值归约（max_with_index/min_with_index） + +| 元素类型 | flatten rank ≤ 2 | flatten rank > 2 | +|---------|-----------------|-----------------| +| f16 | 不降级 ✅ | **降级** ⚠️ | +| f32 | 不降级 ✅ | **降级** ⚠️ | +| bf16 | 不降级 ✅ | **降级** ⚠️ | +| i16 | **降级** ⚠️ | **降级** ⚠️ | +| i32 | **降级** ⚠️ | **降级** ⚠️ | +| i64 | **降级** ⚠️ | **降级** ⚠️ | + +### 优化建议 + +- **基本归约优先使用 f32**：sum/prod/max/min/xori 在 f32 下始终走向量路径 +- **避免 i64 归约**：i64 的基本归约和带索引极值归约都会标量降级 +- **argmax/argmin 注意维度**：f16/f32/bf16 的 max_with_index/min_with_index 在高维（flatten rank > 2）时会降级，尽量保持低维 +- **整数 argmax/argmin 在 Mem-based 架构上始终降级**：i16/i32/i64 的带索引极值归约无法走向量路径 +- **Reg-based 架构（A5 代）优势**：Ascend310B/950 等芯片对基本归约有更好的硬件支持，但 argmax/argmin 仍需关注内存对齐 + +## 完整汇总表 + +| 序号 | 操作 | 助记符 | 标量降级触发条件 | 最常见触发因素 | +|------|------|--------|----------------|-------------| +| 1 | VAddOp | hir.vadd | i64 或 SIMT VF | i64 类型 | +| 2 | VSubOp | hir.vsub | i64 或 SIMT VF | i64 类型 | +| 3 | VMulOp | hir.vmul | i64 或 SIMT VF | i64 类型 | +| 4 | VMinOp | hir.vmin | i64 或 SIMT VF | i64 类型 | +| 5 | VMaxOp | hir.vmax | i64 或 SIMT VF | i64 类型 | +| 6 | VAbsOp | hir.vabs | i64 或 SIMT VF | i64 类型 | +| 7 | VShLOp | hir.vshl | i64 或 SIMT VF | i64 类型 | +| 8 | VShROp | hir.vshr | i64 或 SIMT VF | i64 类型 | +| 9 | VInterleaveOp | hir.vinterleave | i64 或 SIMT VF | i64 类型 | +| 10 | VDeinterleaveOp | hir.vdeinterleave | i64 或 SIMT VF | i64 类型 | +| 11 | VCmpOp | hir.vcmp | 整数类型 且 (非 i32 或 非 EQ/NE) | 整数大小比较 | +| 12 | VMulExtOp | hir.vmulext | i32 或 i64（即始终降级） | 无向量硬件支持 | +| 13 | VCumsumOp | hir.vcumsum | i64 或 累积维度为最后维度 | 累积维度位置 | +| 14 | VCumprodOp | hir.vcumprod | i64 或 累积维度为最后维度 | 累积维度位置 | +| 15 | VReduceOp | hir.vreduce | 依架构和 reduceOp 类型（见上方详解） | i64 / argmax/argmin | + +## Triton 算子优化速查 + +以下是从 Triton 算子编写角度的快速优化参考，帮助避免标量降级导致的性能问题。 + +### 数据类型选择 + +| 场景 | 推荐类型 | 避免类型 | 原因 | +|------|---------|---------|------| +| 算术运算（加减乘、最大最小） | f32, i32 | i64 | i64 触发标量降级 | +| 比较运算 | f32 | i8, i16, i64 | 整数比较大部分会标量降级 | +| 整数相等/不等比较 | i32 | i8, i16, i64 | 仅 i32 的 EQ/NE 走向量路径 | +| 归约运算 | f32 | i64 | i64 归约标量降级 | +| argmax/argmin | f32（低维） | i16, i32, i64 | 整数 argmax/argmin 始终降级 | + +### 操作模式选择 + +| 场景 | 推荐做法 | 避免做法 | 原因 | +|------|---------|---------|------| +| 整数大小比较 | 先 cast 为浮点再比较 | 直接整数 LT/GT/LE/GE | 整数大小比较标量降级 | +| 累积操作 | 累积维度不在最后维度 | 累积维度为最后维度 | 最后维度累积会标量降级 | +| 高精度乘法 | vmulextended (I16) | vmulext (I32) | vmulext 始终标量降级 | + +### 常见性能陷阱 + +1. **i64 陷阱**：i64 是最常见的标量降级触发因素。几乎所有向量操作在 i64 类型下都会降级为标量循环。如果业务逻辑允许，应尽量避免使用 i64 +2. **整数比较陷阱**：除 i32 的 EQ/NE 外，所有整数比较都会标量降级。这在 `where` 条件选择模式（vcmp + vsel）中尤其影响性能 +3. **vmulext 陷阱**：vmulext 在 IR 层面仅支持 I32，而 I32 恰好触发标量降级，导致该操作实际上始终走标量路径 +4. **累积维度陷阱**：非 i64 类型的累积操作在累积维度为最后维度时仍会标量降级，需要关注数据布局 + +## 常见问题 + +**Q: 标量降级对性能的影响有多大？** +A: 标量降级将向量操作退化为逐元素的标量循环，性能损失通常在 10x-100x 量级，具体取决于向量宽度和操作复杂度。对于计算密集型算子，避免标量降级是最重要的优化手段之一。 + +**Q: 如何判断我的 Triton 算子是否触发了标量降级？** +A: 可以通过查看编译后的 HIVM IR，检查是否存在 `scf.for` 循环包裹 `arith` 标量操作的模式。也可以在编译时启用调试日志，观察 `shouldLowerToScalarLoops` 的判断结果。 + +**Q: SIMT VF 模式下为什么所有算术操作都降级？** +A: SIMT（Single Instruction Multiple Threads）模式本身就是标量线程模型，每个线程独立执行标量操作。在这种模式下，向量操作退化为标量循环是设计如此，不存在向量指令可用。 + +**Q: 为什么 i64 类型总是触发标量降级？** +A: AscendNPU 的向量计算单元对 i64 类型的硬件支持有限。大部分向量指令不支持 i64 数据类型，编译器只能退化为标量循环逐元素处理。 + +**Q: VCmpOp 的整数比较为什么只有 i32 的 EQ/NE 能走向量路径？** +A: 硬件向量比较指令仅支持 i32 的相等/不等判断和浮点类型的全部比较模式。i32 的大小比较（LT/GT/LE/GE）以及其他整数宽度的比较没有对应的向量指令。 + +**Q: 标量降级后结果是否正确？** +A: 是的。标量降级是功能等价的变换，仅影响性能，不影响计算结果的正确性。标量循环逐元素执行与向量指令批量执行在数学上产生相同的结果。 + +**Q: 可以通过编译选项禁用标量降级吗？** +A: 不建议这样做。标量降级是因为硬件不支持对应的向量操作，如果强制禁用，编译会在后续阶段失败。正确的做法是调整数据类型或操作模式，避免触发标量降级条件。 + +## 相关文档 + +- 各操作的详细文档： + - [01-unary-ops.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/01-unary-ops.md)（VAbsOp） + - [02-binary-ops.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/02-binary-ops.md)（VAddOp, VSubOp, VMulOp, VMinOp, VMaxOp） + - [05-compare-ops.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/05-compare-ops.md)（VCmpOp） + - [06-shift-ops.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/06-shift-ops.md)（VShLOp, VShROp） + - [07-reduction-ops.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/07-reduction-ops.md)（VReduceOp） + - [08-data-movement.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/08-data-movement.md)（VInterleaveOp, VDeinterleaveOp） + - [09-cumulative-sort.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/09-cumulative-sort.md)（VCumsumOp, VCumprodOp） + - [10-special-ops.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/01-HIVM-Dialect/02-Vector-Operations/10-special-ops.md)（VMulExtOp） +- 编译流水线文档：[04-hivm-transforms.md](file:///d:/项目/trae/triton_a5/docs_ascendnpu_ir/06-Compilation-Pipeline/04-hivm-transforms.md)（HIVMLowerToLoopsPass） +- 源码参考： + - [ImplByScalarOpInterface.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/Interfaces/ImplByScalarOpInterface.td) — 接口定义 + - [ShouldLowerToScalarLoops.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/ImplByScalarOpInterface/ShouldLowerToScalarLoops.cpp) — 降级判断逻辑 + - [LowerToLoops.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/IR/ImplByScalarOpInterface/LowerToLoops.cpp) — 降级实现 + - [HIVMLowerToLoops.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/Transforms/HIVMLowerToLoops.cpp) — 降级 Pass 入口 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/00-overview.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/00-overview.md new file mode 100644 index 00000000..3179df1e --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/00-overview.md @@ -0,0 +1,92 @@ +# HIVM 宏操作总览 + +> 关键词：MacroOp, MacroOpTrait, MacroOpPipeTrait, 跨 Pipeline, UnitFlag, 矩阵乘加 + +## 概述 + +HIVM 宏操作（Macro Operations）是 AscendNPU-IR 中表达跨 Pipeline 复合计算的核心抽象。与单 Pipe 操作不同，宏操作涉及多个硬件 Pipeline 之间的数据流动和同步协调，通常对应 NPU 上 Cube Core 与 Vector Core 之间的协作计算模式。 + +宏操作在 IR 层面通过 `MacroOpTrait` 和 `MacroOpPipeTrait` 两个关键 Trait 进行标记和约束： + +- **MacroOpTrait**：标识该操作为跨 Pipeline 的宏操作，编译器在同步分析、Pipe 分配等 Pass 中会特殊处理此类操作。 +- **MacroOpPipeTrait**：参数化 Trait，声明宏操作涉及的具体输入/输出 Pipeline 组合。例如 `MacroOpPipeTrait<"PIPE::PIPE_MTE1, PIPE::PIPE_M">` 表示该操作从 MTE1 Pipe 输入、从 M Pipe 输出。 + +宏操作体系在 HIVM 中分为两大类： + +1. **本地矩阵乘加（Local MMAD）**：数据在片上存储层次（L1 → L0C）间流动，包括 `mmadL1` 和 `batchMmadL1`。 +2. **全局矩阵乘（Global MMAD）**：数据从全局内存（GM）直接参与计算，包括 `matmul`、`mix_matmul`、`mix_group_matmul`。 + +> Python API 对应：Triton 的 `tl.dot` / `torch.matmul` 等操作在编译时可能被映射为 HIVM 宏操作。 + +## 宏操作类层次 + +``` +HIVM_MacroOp (基类) +├── HIVM_LocalMmadOp (本地 MMAD 基类) +│ ├── MmadL1Op -- hir.mmadL1 +│ └── BatchMmadL1Op -- hir.batchMmadL1 +└── HIVM_GlobalMmadOp (全局 MMAD 基类) + ├── MatmulOp -- hir.matmul + ├── MixMatmulOp -- hir.mix_matmul + └── MixGroupMatmulOp -- hir.mix_group_matmul +``` + +## 核心 Trait 说明 + +### MacroOpTrait + +标识操作为宏操作。编译器在以下场景中识别此 Trait： + +- **同步注入（InjectSync）**：宏操作需要跨 Pipe 同步，InjectSync Pass 会自动在宏操作前后插入 `set_flag`/`wait_flag`。 +- **Pipe 分配**：宏操作的 Pipe 信息由 `MacroOpPipeTrait` 提供，不同于单 Pipe 操作的 `OpPipeTrait`。 +- **GraphSyncSolver**：基于图的同步求解器会将宏操作建模为多节点依赖图。 + +### MacroOpPipeTrait + +参数化 Trait，格式为 `MacroOpPipeTrait<"PIPE::PIPE_IN, PIPE::PIPE_OUT">`，声明宏操作涉及的 Pipeline 组合： + +| 操作 | MacroOpPipeTrait 参数 | 含义 | +|------|----------------------|------| +| mmadL1 / batchMmadL1 | `PIPE::PIPE_MTE1, PIPE::PIPE_M` | 数据从 L1 加载（MTE1），在 Cube Core 计算（M） | +| matmul | `PIPE::PIPE_MTE2, PIPE::PIPE_MTE3` | 数据从 GM 加载（MTE2），结果写回 GM（MTE3） | +| mix_matmul | `PIPE::PIPE_MTE2, PIPE::PIPE_MTE3` | 同 matmul，额外支持 Vector 后处理 | +| mix_group_matmul | `PIPE::PIPE_MTE2, PIPE::PIPE_MTE3` | 同 matmul，支持分组和 Vector 后处理 | + +## UnitFlag 同步机制 + +宏操作支持 UnitFlag 同步模式，用于处理循环中"至少执行一次"的依赖场景。UnitFlag 有四种模式： + +| 模式 | 值 | 说明 | +|------|---|------| +| DISABLED | 0 | 禁用 UnitFlag | +| RESERVED | 1 | 保留 | +| ENABLED_WITHOUT_UPDATE | 2 | 启用但不更新标志 | +| ENABLED_WITH_UPDATE | 3 | 启用并更新标志 | + +在本地 MMAD 操作中，`unit_flag_cond` 参数提供可选的 i1 条件值，`unit_flag_mode` 属性指定每个输出 Tensor 的 UnitFlag 模式。 + +## 宏操作与单 Pipe 操作的对比 + +| 特性 | 单 Pipe 操作 | 宏操作 | +|------|------------|--------| +| Pipe 数量 | 单个 | 多个 | +| 标记 Trait | `SinglePipeOpTrait` + `OpPipeTrait` | `MacroOpTrait` + `MacroOpPipeTrait` | +| 同步需求 | Pipe 内同步 | 跨 Pipe 同步 | +| 典型操作 | load, store, vadd | mmadL1, matmul | +| DestinationStyleOpInterface | 支持 | 支持 | + +## 操作列表 + +| 操作 | 助记符 | 说明 | 详细文档 | +|------|--------|------|---------| +| MmadL1Op | `hir.mmadL1` | 本地矩阵乘加（L1→L0C） | [01-mmad-l1.md](01-mmad-l1.md) | +| BatchMmadL1Op | `hir.batchMmadL1` | 批量本地矩阵乘加 | [02-batch-mmad-l1.md](02-batch-mmad-l1.md) | +| MatmulOp | `hir.matmul` | 全局矩阵乘（GM→GM） | [03-matmul.md](03-matmul.md) | +| MixMatmulOp | `hir.mix_matmul` | 混合 Cube+Vector 矩阵乘 | [04-mix-matmul.md](04-mix-matmul.md) | +| MixGroupMatmulOp | `hir.mix_group_matmul` | 分组矩阵乘（MoE） | [05-mix-group-matmul.md](05-mix-group-matmul.md) | + +## 相关文档 + +- 源码参考：[HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td) +- 同步体系：[04-Synchronization/00-overview.md](../04-Synchronization/00-overview.md) +- 属性类型：[06-Attributes-Types/01-enumerations.md](../06-Attributes-Types/01-enumerations.md) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/01-mmad-l1.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/01-mmad-l1.md new file mode 100644 index 00000000..e9bd3971 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/01-mmad-l1.md @@ -0,0 +1,166 @@ +# hir.mmadL1 — 本地矩阵乘加（L1→L0C） + +> 关键词：mmadL1, Matrix Multiply and Add, L1, L0C, Cube Core, UnitFlag, Fractal Layout + +## 概述 + +`hir.mmadL1` 是 HIVM 方言中的本地矩阵乘加操作，在 Cube Core 上执行。该操作从 L1 存储层次读取矩阵 A 和 B，在 L0C 中执行乘加运算，结果写回 L0C。计算语义为 `C = C + A x B + (optional) channel_bias`。 + +该操作是 HIVM 宏操作体系中最基础的矩阵计算单元，对应硬件上的 mma_tile 指令。它涉及 MTE1（L1 数据搬运）和 M（Cube 矩阵计算）两个 Pipeline，需要跨 Pipe 同步。 + +mmadL1 支持转置加载（a_transpose/b_transpose）、HF32 加速模式、per-channel bias、以及 UnitFlag 同步条件等高级特性。 + +> Python API 对应：Triton 的 `tl.dot` 操作在 Split-K 场景下可能被映射为 mmadL1。 + +## IR 操作定义 + +从 [HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td#L152-L172) 提取： + +``` +def MmadL1Op : HIVM_LocalMmadOp<"mmadL1", [ + NoMaxRankTrait, + DeclareOpInterfaceMethods, + DeclareOpInterfaceMethods +]> +``` + +基类 `HIVM_LocalMmadOp` 定义（[HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td#L49-L150)）： + +``` +class HIVM_LocalMmadOp traits = []> : + HIVM_MacroOp, + ], traits)> +``` + +## 参数说明 + +### 输入操作数（ins） + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$a` | TensorOrMemref | 是 | 矩阵 A，形状为 `[M, K]`，rank 必须为 2 | +| `$b` | TensorOrMemref | 是 | 矩阵 B，形状为 `[K, N]`，rank 必须为 2 | +| `$init_condition` | I1 | 是 | L0C 数据清零条件：为 true 时清零 L0C 后再使用 | +| `$real_m` | Index | 是 | M 维度的实际数据大小 | +| `$real_k` | Index | 是 | K 维度的实际数据大小 | +| `$real_n` | Index | 是 | N 维度的实际数据大小 | +| `$c` | TensorOrMemref | 是 | 矩阵 C（输出/累加），形状为 `[M, N]` | +| `$per_channel_bias` | TensorOrMemref | 否 | Per-channel bias，形状为 `[N]` | +| `$sync_related_args` | Variadic\ | 否 | 同步相关参数，由 InjectSync Pass 自动管理 | +| `$unit_flag_cond` | Variadic\ | 否 | UnitFlag 启用条件，用于循环依赖场景 | + +### 输出操作数（outs） + +| 参数 | 类型 | 说明 | +|------|------|------| +| `$result_tensors` | Variadic\ | 结果 Tensor | + +### 属性 + +| 属性 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$a_transpose` | UnitAttr | 否 | 矩阵 A 在加载前转置 | +| `$b_transpose` | UnitAttr | 否 | 矩阵 B 在加载前转置 | +| `$enable_HF32` | UnitAttr | 否 | 启用 HF32 模式：FP32 数据在 CUBE 计算前舍入为 HF32，性能翻倍但精度降低 | +| `$unit_flag_mode` | UnitFlagArrayAttr | 否 | 每个输出 Tensor 的 UnitFlag 模式 | + +### 额外类方法 + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getOpName()` | StringRef | 返回 `"mma_tile"` | +| `isInitConstant(opt)` | bool | 查询/设置 init_condition 是否为常量 | +| `setInitCondition(Value)` | void | 设置 init 条件值 | +| `getMatmulBiasMode()` | MatmulBiasMode | 获取 bias 模式 | +| `shouldDecomposeBiasByElementAdd()` | bool | 判断是否应将 bias 分解为逐元素加法 | +| `getNumSyncRelatedArgs()` | int | 获取同步参数数量 | +| `getInputOperands(bool)` | SmallVector\ | 获取输入操作数 | +| `getOperandALayout()` | FailureOr\ | 获取 A 的 Fractal Layout | +| `getOperandBLayout()` | FailureOr\ | 获取 B 的 Fractal Layout | +| `getOperandCLayout()` | FailureOr\ | 获取 C 的 Fractal Layout | +| `getOperandBiasLayout()` | FailureOr\ | 获取 Bias 的 Fractal Layout | + +## IR 示例 + +### 基本用法 + +```mlir +%ma = memref.alloc() : memref<256x128xf16> +%mb = memref.alloc() : memref<128x256xf16> +%mc = memref.alloc() : memref<256x256xf32> +%c256 = arith.constant 256 : index +%c128 = arith.constant 128 : index +%init = arith.constant 1 : i1 +hivm.hir.mmadL1 ins(%ma, %mb, %init, %c256, %c128, %c256 : + memref<256x128xf16>, memref<128x256xf16>, i1, index, index, index) + outs(%mc : memref<256x256xf32>) +``` + +### 带转置 + +```mlir +%ma_t = memref.alloc() : memref<128x256xf16> +hivm.hir.mmadL1 {a_transpose} + ins(%ma_t, %mb, %init, %c256, %c128, %c256 : + memref<128x256xf16>, memref<128x256xf16>, i1, index, index, index) + outs(%mc : memref<256x256xf32>) +``` + +### Split-K 循环中的条件初始化 + +```mlir +%mc = memref.alloc() : memref<256x256xf32> +%start = arith.constant 0 : index +%end = arith.constant 1024 : index +%step = arith.constant 128 : index +scf.for %arg0 = %start to %end step %step { + %ma = memref.alloc() : memref<256x128xf16> + %mb = memref.alloc() : memref<128x256xf16> + %init_condition = arith.cmpi eq, %arg0, %start : index + hivm.hir.mmadL1 ins(%ma, %mb, %init_condition, %c256, %c128, %c256 : + memref<256x128xf16>, memref<128x256xf16>, i1, index, index, index) + outs(%mc : memref<256x256xf32>) +} +``` + +### Tensor 语义 + +```mlir +%mc = tensor.empty() : tensor<256x256xf32> +%res = hivm.hir.mmadL1 ins(%ma, %mb, %init_condition, %c256, %c128, %c256 : + tensor<256x128xf16>, tensor<128x256xf16>, i1, index, index, index) + outs(%mC_iter : tensor<256x256xf32>) -> tensor<256x256xf32> +``` + +## IR 层约束与验证 + +1. **Rank 约束**：矩阵 A、B、C 的 rank 必须为 2（batchMmadL1 为 3）。 +2. **Core Type**：操作必须在 Cube Core 上执行（`CubeCoreTypeTrait`）。 +3. **Pipeline**：操作涉及 MTE1 和 M 两个 Pipeline（`MacroOpPipeTrait<"PIPE::PIPE_MTE1, PIPE::PIPE_M">`）。 +4. **init_condition**：在 Split-K 循环中，首次迭代应设置 init_condition 为 true 以清零 L0C，后续迭代为 false 以累加。 +5. **a_transpose / b_transpose**：当设置转置时，输入矩阵的逻辑形状不变，但数据在加载时按转置方式读取。 +6. **enable_HF32**：仅对 FP32 数据有效，将 FP32 舍入为 HF32 后计算，精度降低但性能提升。 +7. **per_channel_bias**：如果提供，形状必须为 `[N]`，与矩阵 B 的列维度匹配。 +8. **Fractal Layout**：mmadL1 实现了 `OpLayoutInterface` 的 `getOperandsTargetFractalLayout` 方法，用于确定操作数的 Fractal 布局。 + +## 常见问题 + +**Q: mmadL1 和 matmul 的区别是什么？** +A: mmadL1 是本地操作，数据从 L1 读取到 L0C 计算；matmul 是全局操作，数据直接从 GM 读取。mmadL1 需要用户手动管理 L1 数据搬运和同步，matmul 由编译器自动处理。 + +**Q: init_condition 什么时候设为 true？** +A: 当需要清零 L0C 累加器时设为 true，通常在 Split-K 循环的第一次迭代。后续迭代设为 false 以累加部分和。 + +**Q: HF32 模式适用于什么场景？** +A: HF32 将 FP32 舍入为 19-bit（10-bit 尾数），适用于对精度不敏感但需要 FP32 吞吐量的场景。 + +## 相关文档 + +- 源码参考：[HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td#L49-L172) +- 测试用例：[ops.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/ops.mlir) +- UnitFlag 详解：[04-Synchronization/03-unit-flag.md](../04-Synchronization/03-unit-flag.md) +- Fractal Layout：[06-Attributes-Types/02-parameterized-attrs.md](../06-Attributes-Types/02-parameterized-attrs.md) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/02-batch-mmad-l1.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/02-batch-mmad-l1.md new file mode 100644 index 00000000..b65e3631 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/02-batch-mmad-l1.md @@ -0,0 +1,110 @@ +# hir.batchMmadL1 — 批量本地矩阵乘加 + +> 关键词：batchMmadL1, Batch Matrix Multiply and Add, L1, L0C, Cube Core + +## 概述 + +`hir.batchMmadL1` 是 HIVM 方言中的批量本地矩阵乘加操作，在 Cube Core 上执行。与 `mmadL1` 类似，但支持批量维度，适用于批量矩阵乘法场景。计算语义为 `C[i] = C[i] + A[i] x B[i] + (optional) channel_bias`。 + +该操作从 L1 存储层次读取批量矩阵 A 和 B，在 L0C 中执行乘加运算。矩阵 A、B、C 的 rank 必须为 3，其中第 0 维为批量维度。 + +> Python API 对应：Triton 的 `tl.dot` 在 batch 维度场景下可能被映射为 batchMmadL1。 + +## IR 操作定义 + +从 [HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td#L174-L184) 提取： + +``` +def BatchMmadL1Op : HIVM_LocalMmadOp<"batchMmadL1", + [NoLibraryFunctionTrait]> { + let summary = [{ + Batch Matrix Multiply and Add Op with inputs from L1 memory hierarchy. + }]; + let description = localMmadBaseDes # [{ + Note: the rank of A, B, and C Matrix must be three, where the 0-th dimension + being the batch dimension. + }]; +} +``` + +## 参数说明 + +### 输入操作数（ins） + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$a` | TensorOrMemref | 是 | 矩阵 A，形状为 `[Batch, M, K]`，rank 必须为 3 | +| `$b` | TensorOrMemref | 是 | 矩阵 B，形状为 `[Batch, K, N]`，rank 必须为 3 | +| `$init_condition` | I1 | 是 | L0C 数据清零条件 | +| `$real_m` | Index | 是 | M 维度的实际数据大小 | +| `$real_k` | Index | 是 | K 维度的实际数据大小 | +| `$real_n` | Index | 是 | N 维度的实际数据大小 | +| `$c` | TensorOrMemref | 是 | 矩阵 C（输出/累加），形状为 `[Batch, M, N]` | +| `$per_channel_bias` | TensorOrMemref | 否 | Per-channel bias | +| `$sync_related_args` | Variadic\ | 否 | 同步相关参数，由 InjectSync Pass 管理 | +| `$unit_flag_cond` | Variadic\ | 否 | UnitFlag 启用条件 | + +### 输出操作数（outs） + +| 参数 | 类型 | 说明 | +|------|------|------| +| `$result_tensors` | Variadic\ | 结果 Tensor | + +### 属性 + +| 属性 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$a_transpose` | UnitAttr | 否 | 矩阵 A 在加载前转置 | +| `$b_transpose` | UnitAttr | 否 | 矩阵 B 在加载前转置 | +| `$enable_HF32` | UnitAttr | 否 | 启用 HF32 模式 | +| `$unit_flag_mode` | UnitFlagArrayAttr | 否 | 每个输出 Tensor 的 UnitFlag 模式 | + +### 与 mmadL1 的差异 + +| 特性 | mmadL1 | batchMmadL1 | +|------|--------|-------------| +| 矩阵 rank | 2 | 3（第 0 维为 batch） | +| A 形状 | `[M, K]` | `[Batch, M, K]` | +| B 形状 | `[K, N]` | `[Batch, K, N]` | +| C 形状 | `[M, N]` | `[Batch, M, N]` | +| OpLayoutInterface | 实现 `getOperandsTargetFractalLayout` | 未实现 | +| LibraryFunctionTrait | 默认支持 | `NoLibraryFunctionTrait` | +| OpName | `"mma_tile"` | 无（使用默认） | + +## IR 示例 + +### 基本用法 + +```mlir +%ma = memref.alloc() : memref<2x256x128xf16> +%mb = memref.alloc() : memref<2x128x256xf16> +%mc = memref.alloc() : memref<2x256x256xf32> +%c256 = arith.constant 256 : index +%c128 = arith.constant 128 : index +%init = arith.constant 1 : i1 +hivm.hir.batchMmadL1 ins(%ma, %mb, %init, %c256, %c128, %c256 : + memref<2x256x128xf16>, memref<2x128x256xf16>, i1, index, index, index) + outs(%mc : memref<2x256x256xf32>) +``` + +## IR 层约束与验证 + +1. **Rank 约束**：矩阵 A、B、C 的 rank 必须为 3，第 0 维为批量维度。 +2. **Core Type**：操作必须在 Cube Core 上执行（`CubeCoreTypeTrait`）。 +3. **Pipeline**：涉及 MTE1 和 M 两个 Pipeline。 +4. **NoLibraryFunctionTrait**：该操作没有预定义的库函数实现，不支持直接 lowering 到标准库调用。 +5. **批量维度一致性**：A、B、C 的批量维度大小应一致。 + +## 常见问题 + +**Q: batchMmadL1 和多次 mmadL1 有什么区别？** +A: batchMmadL1 在单个操作中处理整个批量，硬件可以利用批量间的数据局部性。多次 mmadL1 需要循环展开，每次处理一个批量元素，可能产生更多的同步开销。 + +**Q: 为什么 batchMmadL1 没有 OpLayoutInterface？** +A: 当前实现中 batchMmadL1 未实现 `getOperandsTargetFractalLayout`，其 Fractal Layout 推导使用默认逻辑。 + +## 相关文档 + +- 源码参考：[HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td#L174-L184) +- mmadL1 详解：[01-mmad-l1.md](01-mmad-l1.md) +- 测试用例：[ops.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/ops.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/03-matmul.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/03-matmul.md new file mode 100644 index 00000000..fad84068 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/03-matmul.md @@ -0,0 +1,157 @@ +# hir.matmul — 全局矩阵乘（GM→GM） + +> 关键词：matmul, Matrix Multiply, Global Memory, Descale, Swizzle, Tiling + +## 概述 + +`hir.matmul` 是 HIVM 方言中的全局矩阵乘操作，直接从全局内存（GM）读取输入矩阵并执行矩阵乘法，结果写回 GM。计算语义为 `C = A * B`（无 bias/descale 时）或 `C = descale * (A * B + bias)`（有 bias/descale 时）。 + +该操作涉及 MTE2（GM 数据加载）和 MTE3（GM 数据写回）两个 Pipeline，由编译器自动管理 L1/L0 层次的数据搬运和同步。用户只需提供 GM 地址和 Tiling 参数，无需手动管理片上存储。 + +matmul 支持反量化（descale）、bias、转置、Swizzle 优化等特性，适用于大矩阵乘法场景。 + +> Python API 对应：Triton 的 `tl.dot` 在非 Split-K 场景下通常被映射为 matmul。 + +## IR 操作定义 + +从 [HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td#L236-L318) 提取： + +``` +def MatmulOp : HIVM_GlobalMmadOp<"matmul"> { + let summary = "HIVM Matrix Multiply Op with inputs from global memory"; + let arguments = (ins AnyShaped:$a, + AnyShaped:$b, + Optional:$tilingParams, + Optional:$bias, + Optional:$descale, + OptionalAttr:$aTranspose, + OptionalAttr:$bTranspose, + OptionalAttr:$descaleMode, + Variadic:$blockSizes, + Variadic:$processSizes, + Optional:$swizzleOffset, + Optional:$swizzleDirection, + Optional:$epiloguePTiles, + AnyShaped:$c); +} +``` + +## 参数说明 + +### 输入操作数（ins） + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$a` | AnyShaped | 是 | 矩阵 A，`m x k` | +| `$b` | AnyShaped | 是 | 矩阵 B，`k x n` | +| `$tilingParams` | AnyShaped | 否 | Tiling 参数 | +| `$bias` | AnyShaped | 否 | Bias 向量，形状为 `[n]` | +| `$descale` | AnyShaped | 否 | 反量化缩放因子，形状取决于 descaleMode | +| `$c` | AnyShaped | 是 | 矩阵 C（输出），`m x n` | + +### 输出操作数（outs） + +| 参数 | 类型 | 说明 | +|------|------|------| +| `$result` | Variadic\ | 结果 Tensor | + +### 属性 + +| 属性 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$aTranspose` | UnitAttr | 否 | 矩阵 A 转置加载 | +| `$bTranspose` | UnitAttr | 否 | 矩阵 B 转置加载 | +| `$descaleMode` | HIVM_DescaleModeAttr | 否 | 反量化模式 | + +### I64 操作数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$blockSizes` | Variadic\ | 否 | M/N/K 维度在 L1 层次处理的数据块大小 | +| `$processSizes` | Variadic\ | 否 | M/N/K 维度在 L0 层次处理的数据块大小 | +| `$swizzleOffset` | I64 | 否 | Swizzle 调度的连续块编号 | +| `$swizzleDirection` | I64 | 否 | Swizzle 调度的块方向 | +| `$epiloguePTiles` | I64 | 否 | Epilogue 阶段一次处理的 P tile 数量 | + +### DescaleMode 说明 + +| 模式 | 值 | descale 形状 | 说明 | +|------|---|-------------|------| +| DescaleNull | 0 | 无 | 不使用反量化 | +| DescalePerChannel | 1 | `[n]` | 按 Channel 反量化，形状等于 N | +| DescalePerTensor | 2 | `[1]` | 按 Tensor 反量化，形状为 1 | + +### 额外类方法 + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getOpName()` | StringRef | 返回 `"matmul"` | +| `getDpsInitsMutable()` | MutableOperandRange | DestinationStyleOpInterface 所需 | + +## IR 示例 + +### 基本矩阵乘 + +```mlir +func.func @test_matmul_basic(%A_gm : memref<16x16xf16, #hivm.address_space>, + %B_gm : memref<16x16xf16, #hivm.address_space>, + %res_gm : memref<16x16xf16, #hivm.address_space>) { + hivm.hir.matmul + ins(%A_gm, %B_gm: + memref<16x16xf16, #hivm.address_space>, memref<16x16xf16, #hivm.address_space>) + outs(%res_gm : memref<16x16xf16, #hivm.address_space>) + descale_mode = #hivm.descale_mode + return +} +``` + +### 带 Per-Channel Descale 和 Bias + +```mlir +hivm.hir.matmul + ins(%A_gm, %B_gm: + memref<16x16xf16, #hivm.address_space>, memref<16x16xf16, #hivm.address_space>) + outs(%res_gm : memref<16x16xf16, #hivm.address_space>) + bias = %bias_gm : memref<16xf16, #hivm.address_space> + descale = %descale_perchannel_gm : memref<16xf16, #hivm.address_space> + descale_mode = #hivm.descale_mode +``` + +### 带 Per-Tensor Descale + +```mlir +hivm.hir.matmul + ins(%A_gm, %B_gm: + memref<16x16xf16, #hivm.address_space>, memref<16x16xf16, #hivm.address_space>) + outs(%res_gm : memref<16x16xf16, #hivm.address_space>) + bias = %bias_gm : memref<16xf16, #hivm.address_space> + descale = %descale_pertensor_gm : memref<1xf16, #hivm.address_space> + descale_mode = #hivm.descale_mode +``` + +## IR 层约束与验证 + +1. **Core Type**：操作在 Cube Core 上执行，通过 `HIVMInferCoreTypeInterface` 推断。 +2. **Pipeline**：涉及 MTE2 和 MTE3 两个 Pipeline。 +3. **NoMaxRankTrait**：不限制操作数的最大 rank。 +4. **Address Space**：输入/输出操作数通常需要 `#hivm.address_space` 标记。 +5. **Descale 一致性**：当提供 descale 操作数时，descaleMode 必须与 descale 形状一致。 +6. **Bias 形状**：bias 的形状必须为 `[n]`，与矩阵 B 的列维度匹配。 +7. **blockSizes / processSizes**：通常为 3 个 I64 值，分别对应 M、N、K 维度的块大小。 + +## 常见问题 + +**Q: matmul 和 mmadL1 的主要区别？** +A: matmul 直接从 GM 读写数据，编译器自动管理 L1/L0 搬运和同步；mmadL1 需要用户手动管理 L1 数据搬运。matmul 适合端到端矩阵乘法，mmadL1 适合需要精细控制数据流的 Split-K 场景。 + +**Q: Swizzle 参数的作用？** +A: Swizzle 用于优化 GM 访存的 bank conflict，通过改变数据块的读取顺序来避免冲突。`swizzleOffset` 指定起始块编号，`swizzleDirection` 指定遍历方向。 + +**Q: descale 的计算公式？** +A: `C = descale * (A * B + bias)`，其中 descale 是反量化缩放因子，用于量化推理场景。 + +## 相关文档 + +- 源码参考：[HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td#L236-L318) +- 测试用例：[ops.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/ops.mlir) +- DescaleMode 枚举：[06-Attributes-Types/01-enumerations.md](../06-Attributes-Types/01-enumerations.md) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/04-mix-matmul.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/04-mix-matmul.md new file mode 100644 index 00000000..4ebb8975 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/04-mix-matmul.md @@ -0,0 +1,147 @@ +# hir.mix_matmul — 混合 Cube+Vector 矩阵乘 + +> 关键词：mix_matmul, Mix Matrix Multiply, Post-Vector Function, Workspace, Communication + +## 概述 + +`hir.mix_matmul` 是 HIVM 方言中的混合矩阵乘操作，在 Cube Core 执行矩阵乘法后，支持在 Vector Core 上执行后处理函数（post-vector function）。这种 Cube+Vector 的混合模式允许在 tile 级别融合后处理操作（如激活函数、类型转换等），避免额外的 GM 读写开销。 + +计算语义与 matmul 相同：`C = descale * (A * B + bias)`，但额外支持： +- `post_vector_func_ins`：Vector 后处理函数的输入参数 +- `workspace_ins`：工作空间缓冲区 +- `comm_params`：通信参数（用于融合通信操作，如 AllReduce） + +> Python API 对应：Triton 的 `tl.dot` + 后续 elementwise 操作在编译时可能被融合为 mix_matmul。 + +## IR 操作定义 + +从 [HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td#L320-L422) 提取： + +``` +def MixMatmulOp : HIVM_GlobalMmadOp<"mix_matmul"> { + let summary = "HIVM (Mix) Matrix Multiply Op with inputs from global memory"; + let arguments = (ins AnyShaped:$a, + AnyShaped:$b, + Variadic:$postVecFuncIns, + Variadic:$workspaceIns, + Optional:$tilingParams, + Optional:$commParams, + Optional:$bias, + Optional:$descale, + OptionalAttr:$aTranspose, + OptionalAttr:$bTranspose, + OptionalAttr:$descaleMode, + Variadic:$blockSizes, + Variadic:$processSizes, + Optional:$swizzleOffset, + Optional:$swizzleDirection, + Optional:$epiloguePTiles, + AnyShaped:$c); +} +``` + +## 参数说明 + +### 输入操作数（ins） + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$a` | AnyShaped | 是 | 矩阵 A，`m x k` | +| `$b` | AnyShaped | 是 | 矩阵 B，`k x n` | +| `$postVecFuncIns` | Variadic\ | 否 | Vector 后处理函数的输入参数 | +| `$workspaceIns` | Variadic\ | 否 | 工作空间缓冲区输入 | +| `$tilingParams` | AnyShaped | 否 | Tiling 参数 | +| `$commParams` | AnyShaped | 否 | 通信相关参数（拓扑、通信器、group 等） | +| `$bias` | AnyShaped | 否 | Bias 向量 | +| `$descale` | AnyShaped | 否 | 反量化缩放因子 | +| `$c` | AnyShaped | 是 | 矩阵 C（输出） | + +### 输出操作数（outs） + +| 参数 | 类型 | 说明 | +|------|------|------| +| `$result` | Variadic\ | 结果 Tensor | + +### 属性 + +| 属性 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$aTranspose` | UnitAttr | 否 | 矩阵 A 转置加载 | +| `$bTranspose` | UnitAttr | 否 | 矩阵 B 转置加载 | +| `$descaleMode` | HIVM_DescaleModeAttr | 否 | 反量化模式 | +| `post_vector_func` | StrAttr | 否 | Vector 后处理函数名称（通过属性指定） | + +### I64 操作数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$blockSizes` | Variadic\ | 否 | L1 层次 M/N/K 块大小 | +| `$processSizes` | Variadic\ | 否 | L0 层次 M/N/K 块大小 | +| `$swizzleOffset` | I64 | 否 | Swizzle 起始块编号 | +| `$swizzleDirection` | I64 | 否 | Swizzle 方向 | +| `$epiloguePTiles` | I64 | 否 | Epilogue P tile 数量 | + +### 与 matmul 的差异 + +| 特性 | matmul | mix_matmul | +|------|--------|------------| +| 后处理融合 | 不支持 | 支持 `post_vector_func_ins` | +| 工作空间 | 不支持 | 支持 `workspace_ins` | +| 通信参数 | 不支持 | 支持 `comm_params` | +| OpName | `"matmul"` | `"mix_matmul"` | + +## IR 示例 + +### 基本用法 + +```mlir +hivm.hir.mix_matmul + ins(%A_gm, %B_gm: + memref<16x16xf16, #hivm.address_space>, memref<16x16xf16, #hivm.address_space>) + outs(%res_gm : memref<16x16xf16, #hivm.address_space>) + tiling_params = %tiling_params_gm : memref<16xf16, #hivm.address_space> + comm_params = %comm_params_gm : memref<16xi64, #hivm.address_space> +``` + +### 带 Post-Vector Function 和 Workspace + +```mlir +hivm.hir.set_ffts_base_addr %arg0 +hivm.hir.mix_matmul {post_vector_func = "bishengir_gen_vector_epilogue_func"} + ins(%arg1, %arg2 : + memref<1024x1024xf16, #hivm.address_space>, memref<1024x1024xf16, #hivm.address_space>) + post_vector_func_ins(%arg3 : memref<1024x1024xf16, #hivm.address_space>) + workspace_ins(%arg4 : memref<1024x1024xf16, #hivm.address_space>) + outs(%arg5 : memref<1024x1024xf16, #hivm.address_space>) + block_sizes(%c128_i64, %c256_i64, %c256_i64 : i64, i64, i64) + process_sizes(%c128_i64, %c256_i64, %c64_i64 : i64, i64, i64) + swizzle_offset = %c1_i64 : i64 + swizzle_direction = %c0_i64 : i64 + epilogue_p_tiles = %c4_i64 : i64 +``` + +## IR 层约束与验证 + +1. **Core Type**：通过 `HIVMInferCoreTypeInterface` 推断，通常为 Cube Core 执行矩阵乘法部分。 +2. **Pipeline**：涉及 MTE2 和 MTE3 两个 Pipeline。 +3. **Post-Vector Function**：`post_vector_func` 属性指定 Vector 后处理函数名称，`postVecFuncIns` 提供其输入参数。 +4. **Workspace**：`workspaceIns` 提供工作空间缓冲区，用于中间计算结果存储。 +5. **Communication**：`commParams` 用于融合通信操作（如 AllReduce），包含拓扑、通信器等信息。 +6. **MIX Kernel**：使用 mix_matmul 的函数通常标记为 `hivm.func_core_type = #hivm.func_core_type` 或 `MIX`。 + +## 常见问题 + +**Q: mix_matmul 的 post_vector_func 是什么？** +A: 它是 Vector Core 上执行的后处理函数，可以对矩阵乘结果进行激活函数、类型转换等操作。函数名通过 `post_vector_func` 属性指定，输入通过 `postVecFuncIns` 传入。 + +**Q: workspace_ins 的用途？** +A: 工作空间缓冲区用于存储中间计算结果，例如在融合 AllReduce 时需要临时存储部分和。 + +**Q: 什么时候应该用 mix_matmul 而不是 matmul？** +A: 当需要在矩阵乘后立即执行 Vector 后处理（如激活、量化）或需要融合通信操作时，使用 mix_matmul 可以避免额外的 GM 读写。 + +## 相关文档 + +- 源码参考：[HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td#L320-L422) +- 测试用例：[ops.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/ops.mlir) +- matmul 详解：[03-matmul.md](03-matmul.md) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/05-mix-group-matmul.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/05-mix-group-matmul.md new file mode 100644 index 00000000..247d9029 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/03-Macro-Operations/05-mix-group-matmul.md @@ -0,0 +1,150 @@ +# hir.mix_group_matmul — 分组矩阵乘（MoE 场景） + +> 关键词：mix_group_matmul, Group Matmul, MoE, tokens_per_expert, Post-Vector Function + +## 概述 + +`hir.mix_group_matmul` 是 HIVM 方言中的分组矩阵乘操作，专为 Mixture-of-Experts（MoE）场景设计。在 MoE 中，不同的 token 被路由到不同的 expert（权重矩阵），`mix_group_matmul` 通过 `tokens_per_expert` 参数指定每个 expert 处理的 token 数量，实现高效的分组矩阵乘法。 + +计算语义与 matmul 相同：`C = descale * (A * B + bias)`，但额外支持： +- `tokens_per_expert`：1D 向量，指定每个 expert 的 token 数量 +- `post_vector_func_ins` / `post_vector_func_outs`：Vector 后处理函数的输入/输出 +- `workspace_ins`：工作空间缓冲区 +- `comm_params`：通信参数 + +> Python API 对应：MoE 模型中的分组矩阵乘操作。 + +## IR 操作定义 + +从 [HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td#L424-L535) 提取： + +``` +def MixGroupMatmulOp : HIVM_GlobalMmadOp<"mix_group_matmul"> { + let summary = "HIVM (Mix) Matrix Group Multiply Op with inputs from global memory"; + let arguments = (ins AnyShaped:$a, // weight, 3D + AnyShaped:$b, // tokens, 2D + AnyShaped:$tokens_per_expert, // tokens_per_expert, 1D + Variadic:$postVecFuncIns, + Variadic:$postVecFuncOuts, + Variadic:$workspaceIns, + Optional:$tilingParams, + Optional:$commParams, + Optional:$bias, + Optional:$descale, + OptionalAttr:$aTranspose, + OptionalAttr:$bTranspose, + OptionalAttr:$descaleMode, + Variadic:$blockSizes, + Variadic:$processSizes, + Optional:$swizzleOffset, + Optional:$swizzleDirection, + Optional:$epiloguePTiles, + AnyShaped:$c); +} +``` + +## 参数说明 + +### 输入操作数（ins） + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$a` | AnyShaped | 是 | 权重矩阵，3D 形状 `[num_experts, K, N]` | +| `$b` | AnyShaped | 是 | Token 矩阵，2D 形状 `[total_tokens, K]` | +| `$tokens_per_expert` | AnyShaped | 是 | 每个 expert 的 token 数量，1D 形状 `[num_experts]` | +| `$postVecFuncIns` | Variadic\ | 否 | Vector 后处理函数输入 | +| `$postVecFuncOuts` | Variadic\ | 否 | Vector 后处理函数输出 | +| `$workspaceIns` | Variadic\ | 否 | 工作空间缓冲区 | +| `$tilingParams` | AnyShaped | 否 | Tiling 参数 | +| `$commParams` | AnyShaped | 否 | 通信参数 | +| `$bias` | AnyShaped | 否 | Bias 向量 | +| `$descale` | AnyShaped | 否 | 反量化缩放因子 | +| `$c` | AnyShaped | 是 | 输出矩阵 C | + +### 输出操作数（outs） + +| 参数 | 类型 | 说明 | +|------|------|------| +| `$result` | Variadic\ | 结果 Tensor | + +### 属性 + +| 属性 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$aTranspose` | UnitAttr | 否 | 权重矩阵转置加载 | +| `$bTranspose` | UnitAttr | 否 | Token 矩阵转置加载 | +| `$descaleMode` | HIVM_DescaleModeAttr | 否 | 反量化模式 | + +### I64 操作数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$blockSizes` | Variadic\ | 否 | L1 层次 M/N/K 块大小 | +| `$processSizes` | Variadic\ | 否 | L0 层次 M/N/K 块大小 | +| `$swizzleOffset` | I64 | 否 | Swizzle 起始块编号 | +| `$swizzleDirection` | I64 | 否 | Swizzle 方向 | +| `$epiloguePTiles` | I64 | 否 | Epilogue P tile 数量 | + +### 与 mix_matmul 的差异 + +| 特性 | mix_matmul | mix_group_matmul | +|------|-----------|-----------------| +| 权重矩阵 A | 2D `[M, K]` | 3D `[num_experts, K, N]` | +| Token 矩阵 B | 2D `[K, N]` | 2D `[total_tokens, K]` | +| tokens_per_expert | 不支持 | 必需，1D `[num_experts]` | +| post_vector_func_outs | 不支持 | 支持 | +| OpName | `"mix_matmul"` | `"mix_group_matmul"` | + +## IR 示例 + +### 基本用法 + +```mlir +hivm.hir.mix_group_matmul + ins(%A_gm, %B_gm, %tokens_per_expert_gm: + memref<16x16x16xf16, #hivm.address_space>, + memref<16x16xf16, #hivm.address_space>, + memref<16xi64, #hivm.address_space>) + outs(%res_gm : memref<16x16xf16, #hivm.address_space>) +``` + +### 带 Post-Vector Function 和通信参数 + +```mlir +hivm.hir.mix_group_matmul + ins(%A_gm, %B_gm, %tokens_per_expert_gm: + memref<16x16x16xf16, #hivm.address_space>, + memref<16x16xf16, #hivm.address_space>, + memref<16xi64, #hivm.address_space>) + post_vector_func_ins(%post_vector_func_ins : memref<1024x1024xf16, #hivm.address_space>) + post_vector_func_outs(%post_vector_func_outs : memref<1024x1024xf16, #hivm.address_space>) + outs(%res_gm : memref<16x16xf16, #hivm.address_space>) + tiling_params = %tiling_params_gm : memref<16xf16, #hivm.address_space> + comm_params = %comm_params_gm : memref<16xi64, #hivm.address_space> +``` + +## IR 层约束与验证 + +1. **Core Type**：通过 `HIVMInferCoreTypeInterface` 推断。 +2. **Pipeline**：涉及 MTE2 和 MTE3 两个 Pipeline。 +3. **权重矩阵维度**：A 必须为 3D，形状 `[num_experts, K, N]`。 +4. **tokens_per_expert**：必须为 1D，长度等于 num_experts，所有元素之和应等于 total_tokens。 +5. **Post-Vector Function**：同时支持输入和输出参数，允许后处理函数产生额外输出。 +6. **NoMaxRankTrait**：不限制操作数的最大 rank。 + +## 常见问题 + +**Q: tokens_per_expert 的含义？** +A: 它是一个 1D 向量，`tokens_per_expert[i]` 表示第 i 个 expert 需要处理的 token 数量。编译器据此将 token 分配到不同的 expert 进行分组矩阵乘。 + +**Q: 为什么权重矩阵 A 是 3D 的？** +A: 在 MoE 场景中，每个 expert 有独立的权重矩阵。A 的第 0 维是 expert 数量，后两维是每个 expert 的权重矩阵。 + +**Q: mix_group_matmul 和多次 mix_matmul 的区别？** +A: mix_group_matmul 在单个操作中处理所有 expert 的分组计算，编译器可以优化数据搬运和同步。多次 mix_matmul 需要循环展开，可能产生更多开销。 + +## 相关文档 + +- 源码参考：[HIVMMacroOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMMacroOps.td#L424-L535) +- 测试用例：[ops.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/ops.mlir) +- mix_matmul 详解：[04-mix-matmul.md](04-mix-matmul.md) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/00-overview.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/00-overview.md new file mode 100644 index 00000000..5a3853d3 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/00-overview.md @@ -0,0 +1,111 @@ +# HIVM 同步体系总览 + +> 关键词：Synchronization, Pipe Sync, Block Sync, FFTS, UnitFlag, Event ID, InjectSync, GraphSyncSolver + +## 概述 + +HIVM 同步体系是 AscendNPU-IR 中协调多 Pipeline、多 Core、多 Block 之间数据依赖的核心机制。由于 NPU 采用多 Pipeline 并行执行架构（MTE1/MTE2/MTE3/M/V 等），不同 Pipeline 上的操作之间存在数据依赖，需要通过同步原语来保证执行顺序的正确性。 + +HIVM 同步体系分为三个层次： + +1. **管道同步（Pipe Synchronization）**：同一 Block 内不同 Pipeline 之间的同步，通过 Event ID 机制实现。这是最基础、最常用的同步层次。 +2. **跨核同步（Block Synchronization）**：不同 Block 之间的同步，通过 FFTS（Fast Flag Transfer System）机制实现。用于多 Block 协作计算场景。 +3. **UnitFlag 同步**：处理循环中"至少执行一次"依赖的特殊同步模式，通过 UnitFlag 条件控制。 + +## 同步操作层次 + +``` +┌─────────────────────────────────────────────────┐ +│ 跨核同步 (Block Sync) │ +│ sync_block / sync_block_set / sync_block_wait │ +│ create_sync_block_lock / sync_block_lock/unlock │ +├─────────────────────────────────────────────────┤ +│ 管道同步 (Pipe Sync) │ +│ set_flag / wait_flag / pipe_barrier │ +├─────────────────────────────────────────────────┤ +│ UnitFlag 同步 │ +│ unit_flag_cond / unit_flag_mode │ +│ (嵌入在宏操作中) │ +└─────────────────────────────────────────────────┘ +``` + +## 管道同步 + +管道同步是 HIVM 中最细粒度的同步机制，用于同一 Block 内不同 Pipeline 之间的协调。核心操作： + +| 操作 | 助记符 | 说明 | +|------|--------|------| +| SetFlagOp | `hir.set_flag` | 在 set_pipe 上设置 Event Flag | +| WaitFlagOp | `hir.wait_flag` | 在 wait_pipe 上等待 Event Flag | +| PipeBarrierOp | `hir.pipe_barrier` | Pipeline 屏障，等待该 Pipe 所有操作完成 | + +### Event ID 机制 + +每个 Pipe 有 8 个 Event ID（EVENT_ID0 ~ EVENT_ID7），用于区分不同的同步点。`set_flag` 和 `wait_flag` 通过 Event ID 配对使用： + +- `set_flag[set_pipe, wait_pipe, event_id]`：在 set_pipe 上发送信号 +- `wait_flag[set_pipe, wait_pipe, event_id]`：在 wait_pipe 上等待信号 + +Event ID 可以是静态的（编译时确定的 `#hivm.event`）或动态的（运行时计算的 i64 值）。 + +## 跨核同步 + +跨核同步用于不同 Block 之间的协调，基于 FFTS 机制。核心操作： + +| 操作 | 助记符 | 说明 | +|------|--------|------| +| SyncBlockOp | `hir.sync_block` | 高层跨核同步，支持多种模式 | +| SyncBlockSetOp | `hir.sync_block_set` | 发送跨核同步信号 | +| SyncBlockWaitOp | `hir.sync_block_wait` | 等待跨核同步信号 | +| CreateSyncBlockLockOp | `hir.create_sync_block_lock` | 创建跨核锁 | +| SyncBlockLockOp | `hir.sync_block_lock` | 获取跨核锁 | +| SyncBlockUnlockOp | `hir.sync_block_unlock` | 释放跨核锁 | + +### SyncBlockMode + +| 模式 | 值 | 说明 | +|------|---|------| +| ALL_CUBE | 0 | 所有 Cube Core 同步 | +| ALL_VECTOR | 1 | 所有 Vector Core 同步 | +| ALL_SUB_VECTOR | 2 | 所有 Sub-Vector 同步 | +| BARRIER_CUBE | 3 | Cube-Cube 屏障同步 | +| BARRIER_VECTOR | 4 | Vector-Vector 屏障同步 | +| ALL | 5 | 所有 AIC/AIV 同步 | + +### FFTS 机制 + +FFTS（Fast Flag Transfer System）是 AscendNPU 的硬件同步机制。每个 Block 有一个 FFTS 基地址（`ffts_base_addr`），FFTS 收集特定 flag_id 后将其设置回 Block 组中的 Block，实现同步。 + +在 Ascend910B 及以上平台上，`ffts_base_addr` 必须通过 `hir.set_ffts_base_addr` 设置。 + +## UnitFlag 同步 + +UnitFlag 是嵌入在宏操作（如 mmadL1）中的同步机制，用于处理循环中"至少执行一次"的依赖场景。详见 [03-unit-flag.md](03-unit-flag.md)。 + +## 自动同步注入 + +HIVM 编译器提供两个自动同步注入 Pass： + +1. **InjectSync Pass**：基于启发式规则，在操作前后自动插入 `set_flag`/`wait_flag`。 +2. **GraphSyncSolver Pass**：基于图的同步求解器，构建操作依赖图后求解最优同步方案。 + +详见 [04-sync-injection.md](04-sync-injection.md)。 + +## 操作列表 + +| 操作 | 助记符 | 详细文档 | +|------|--------|---------| +| SetFlagOp | `hir.set_flag` | [01-pipe-sync.md](01-pipe-sync.md) | +| WaitFlagOp | `hir.wait_flag` | [01-pipe-sync.md](01-pipe-sync.md) | +| PipeBarrierOp | `hir.pipe_barrier` | [01-pipe-sync.md](01-pipe-sync.md) | +| SyncBlockOp | `hir.sync_block` | [02-block-sync.md](02-block-sync.md) | +| SyncBlockSetOp | `hir.sync_block_set` | [02-block-sync.md](02-block-sync.md) | +| SyncBlockWaitOp | `hir.sync_block_wait` | [02-block-sync.md](02-block-sync.md) | +| CreateSyncBlockLockOp | `hir.create_sync_block_lock` | [02-block-sync.md](02-block-sync.md) | +| SyncBlockLockOp | `hir.sync_block_lock` | [02-block-sync.md](02-block-sync.md) | +| SyncBlockUnlockOp | `hir.sync_block_unlock` | [02-block-sync.md](02-block-sync.md) | + +## 相关文档 + +- 源码参考：[HIVMSynchronizationOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMSynchronizationOps.td) +- 测试用例：[sync-ops.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/sync-ops.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/01-pipe-sync.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/01-pipe-sync.md new file mode 100644 index 00000000..0bf33be7 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/01-pipe-sync.md @@ -0,0 +1,168 @@ +# 管道同步操作 — set_flag / wait_flag / pipe_barrier + +> 关键词：Pipe Sync, Event ID, set_flag, wait_flag, pipe_barrier, Pipeline Synchronization + +## 概述 + +管道同步操作是 HIVM 中最细粒度的同步机制，用于同一 Block 内不同 Pipeline 之间的协调。NPU 的多 Pipeline 架构中，MTE1/MTE2/MTE3/M/V 等 Pipeline 可以并行执行，但存在数据依赖时需要通过 Event Flag 机制保证执行顺序。 + +三个核心操作： +- `hir.set_flag`：在 set_pipe 上设置 Event Flag，通知 wait_pipe 数据已就绪 +- `hir.wait_flag`：在 wait_pipe 上等待 Event Flag，阻塞直到数据就绪 +- `hir.pipe_barrier`：Pipeline 屏障，等待该 Pipe 上所有先前操作完成 + +## IR 操作定义 + +### SetFlagOp + +从 [HIVMSynchronizationOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMSynchronizationOps.td#L34-L49) 提取： + +``` +def SetFlagOp : HIVM_SynchronizationOp<"set_flag"> { + let arguments = (ins HIVM_PipeAttr:$set_pipe, + HIVM_PipeAttr:$wait_pipe, + OptionalAttr:$static_event_id, + Optional:$dynamic_event_id); + let assemblyFormat = [{ + `[` + $set_pipe + `,` $wait_pipe + `,` custom($static_event_id, $dynamic_event_id) + `]` attr-dict + }]; +} +``` + +### WaitFlagOp + +从 [HIVMSynchronizationOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMSynchronizationOps.td#L51-L66) 提取： + +``` +def WaitFlagOp : HIVM_SynchronizationOp<"wait_flag"> { + let arguments = (ins HIVM_PipeAttr:$set_pipe, + HIVM_PipeAttr:$wait_pipe, + OptionalAttr:$static_event_id, + Optional:$dynamic_event_id); + let assemblyFormat = [{ + `[` + $set_pipe + `,` $wait_pipe + `,` custom($static_event_id, $dynamic_event_id) + `]` attr-dict + }]; +} +``` + +### PipeBarrierOp + +从 [HIVMSynchronizationOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMSynchronizationOps.td#L68-L72) 提取： + +``` +def PipeBarrierOp : HIVM_SynchronizationOp<"pipe_barrier"> { + let arguments = (ins HIVM_PipeAttr:$pipe); + let assemblyFormat = "`[` $pipe `]` attr-dict"; +} +``` + +## 参数说明 + +### SetFlagOp / WaitFlagOp 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$set_pipe` | HIVM_PipeAttr | 是 | 发送信号的 Pipeline | +| `$wait_pipe` | HIVM_PipeAttr | 是 | 接收信号的 Pipeline | +| `$static_event_id` | HIVM_EventAttr | 否 | 静态 Event ID（编译时确定） | +| `$dynamic_event_id` | I64 | 否 | 动态 Event ID（运行时计算） | + +### PipeBarrierOp 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$pipe` | HIVM_PipeAttr | 是 | 需要屏障的 Pipeline | + +### Event ID 约束 + +- 每个 Pipe 有 8 个 Event ID（EVENT_ID0 ~ EVENT_ID7） +- `static_event_id` 和 `dynamic_event_id` 二选一，不能同时存在 +- 同一对 set_flag/wait_flag 必须使用相同的 (set_pipe, wait_pipe, event_id) 三元组 +- Event ID 在同一 (set_pipe, wait_pipe) 对中不能重复使用（除非前一对已完成 wait） + +### 典型 Pipe 组合 + +| set_pipe | wait_pipe | 场景 | +|----------|-----------|------| +| PIPE_MTE1 | PIPE_M | L1 数据加载完成后通知 Cube 计算 | +| PIPE_MTE2 | PIPE_M | GM 数据加载完成后通知 Cube 计算 | +| PIPE_M | PIPE_MTE3 | Cube 计算完成后通知 GM 写回 | +| PIPE_M | PIPE_FIX | Cube 计算完成后通知 Fixpipe | +| PIPE_M | PIPE_V | Cube 计算完成后通知 Vector 后处理 | + +## IR 示例 + +### 静态 Event ID + +```mlir +hivm.hir.set_flag [#hivm.pipe, #hivm.pipe, #hivm.event] +hivm.hir.wait_flag [#hivm.pipe, #hivm.pipe, #hivm.event] +``` + +### 动态 Event ID + +```mlir +%eventId = arith.constant 1 : i64 +hivm.hir.set_flag [#hivm.pipe, #hivm.pipe, %eventId] +hivm.hir.wait_flag [#hivm.pipe, #hivm.pipe, %eventId] +``` + +### Pipe Barrier + +```mlir +hivm.hir.pipe_barrier [#hivm.pipe] +``` + +### 完整的 Load → Compute → Store 同步模式 + +```mlir +// 加载数据到 L1 +hivm.hir.load ins(%src_gm : ...) outs(%dst_l1 : ...) +// 设置 flag 通知 M Pipe 数据已就绪 +hivm.hir.set_flag [#hivm.pipe, #hivm.pipe, #hivm.event] + +// 等待数据就绪后执行矩阵乘 +hivm.hir.wait_flag [#hivm.pipe, #hivm.pipe, #hivm.event] +hivm.hir.mmadL1 ins(...) outs(...) + +// 设置 flag 通知 MTE3 Pipe 计算结果已就绪 +hivm.hir.set_flag [#hivm.pipe, #hivm.pipe, #hivm.event] + +// 等待计算完成后写回 GM +hivm.hir.wait_flag [#hivm.pipe, #hivm.pipe, #hivm.event] +hivm.hir.store ins(...) outs(...) +``` + +## IR 层约束与验证 + +1. **Event ID 配对**：set_flag 和 wait_flag 必须使用相同的 (set_pipe, wait_pipe, event_id) 三元组。 +2. **Event ID 唯一性**：同一 (set_pipe, wait_pipe) 对中，正在使用的 Event ID 不能重复。 +3. **静态/动态互斥**：`static_event_id` 和 `dynamic_event_id` 不能同时存在。 +4. **Pipe 合法性**：set_pipe 和 wait_pipe 必须是合法的 PIPE 枚举值。 +5. **顺序约束**：set_flag 必须在对应的 wait_flag 之前执行（否则 wait_flag 会阻塞）。 +6. **PipeBarrier**：pipe_barrier 会等待指定 Pipe 上所有先前操作完成，是比 set_flag/wait_flag 更重的同步操作。 + +## 常见问题 + +**Q: 什么时候用静态 Event ID，什么时候用动态 Event ID？** +A: 静态 Event ID 适用于编译时可以确定同步点的场景（如单次 load→compute→store）。动态 Event ID 适用于循环中需要复用 Event ID 的场景，运行时根据循环变量计算 Event ID。 + +**Q: Event ID 数量不够用怎么办？** +A: 每个 Pipe 只有 8 个 Event ID。在复杂场景中，InjectSync Pass 会自动管理 Event ID 的分配和复用。GraphSyncSolver Pass 可以更优化地使用 Event ID。 + +**Q: pipe_barrier 和 set_flag/wait_flag 的区别？** +A: pipe_barrier 是重量级同步，等待整个 Pipe 完成；set_flag/wait_flag 是轻量级同步，只同步特定的数据依赖。通常优先使用 set_flag/wait_flag。 + +## 相关文档 + +- 源码参考：[HIVMSynchronizationOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMSynchronizationOps.td#L34-L72) +- 测试用例：[sync-ops.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/sync-ops.mlir) +- Event 枚举：[06-Attributes-Types/01-enumerations.md](../06-Attributes-Types/01-enumerations.md) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/02-block-sync.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/02-block-sync.md new file mode 100644 index 00000000..b22c3bff --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/02-block-sync.md @@ -0,0 +1,206 @@ +# 跨核同步操作 — sync_block / sync_block_set / sync_block_wait + +> 关键词：Block Sync, FFTS, SyncBlockMode, sync_block, sync_block_set, sync_block_wait, Lock + +## 概述 + +跨核同步操作用于不同 Block 之间的协调，基于 FFTS（Fast Flag Transfer System）硬件机制实现。在多 Block 协作计算场景中（如大矩阵分块乘法、AllReduce 等），不同 Block 需要同步以避免数据竞争。 + +HIVM 提供两种跨核同步接口： +1. **高层接口**：`sync_block`，封装了 FFTS 的细节，通过 SyncBlockMode 指定同步模式 +2. **低层接口**：`sync_block_set` / `sync_block_wait`，直接控制 FFTS 的 set/wait 操作 + +此外，还提供基于锁的同步原语：`create_sync_block_lock` / `sync_block_lock` / `sync_block_unlock`。 + +## IR 操作定义 + +### SyncBlockOp + +从 [HIVMSynchronizationOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMSynchronizationOps.td#L78-L118) 提取： + +``` +def SyncBlockOp : HIVM_SynchronizationOp<"sync_block"> { + let arguments = (ins HIVM_SyncBlockModeAttr:$sync_block_mode, + OptionalAttr:$flag_id, + Optional:$ffts_base_addr, + OptionalAttr:$tcube_pipe, + OptionalAttr:$tvector_pipe); + let assemblyFormat = [{ + attr-dict `[` $sync_block_mode (`,` $flag_id^)?`]` + (`ffts_base_addr` `=` $ffts_base_addr^)? + (`tcube_pipe` `=` $tcube_pipe^)? + (`tvector_pipe` `=` $tvector_pipe^)? + }]; +} +``` + +### SyncBlockSetOp + +从 [HIVMSynchronizationOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMSynchronizationOps.td#L120-L148) 提取： + +``` +def SyncBlockSetOp : HIVM_SynchronizationOp<"sync_block_set", [AttrSizedOperandSegments]> { + let arguments = (ins HIVM_TCoreTypeAttr:$tcore_type, + HIVM_PipeAttr:$tpipe, + HIVM_PipeAttr:$pipe, + OptionalAttr:$static_flag_id, + Optional:$dynamic_flag_id, + Optional:$ffts_base_addr, + DefaultValuedOptionalAttr:$tsync_instr_mode); +} +``` + +### SyncBlockWaitOp + +从 [HIVMSynchronizationOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMSynchronizationOps.td#L150-L175) 提取： + +``` +def SyncBlockWaitOp : HIVM_SynchronizationOp<"sync_block_wait"> { + let arguments = (ins HIVM_TCoreTypeAttr:$tcore_type, + HIVM_PipeAttr:$tpipe, + HIVM_PipeAttr:$pipe, + OptionalAttr:$static_flag_id, + Optional:$dynamic_flag_id, + DefaultValuedOptionalAttr:$tsync_instr_mode); +} +``` + +## 参数说明 + +### SyncBlockOp 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$sync_block_mode` | HIVM_SyncBlockModeAttr | 是 | 同步模式 | +| `$flag_id` | IntegerAttr | 否 | Flag ID（静态） | +| `$ffts_base_addr` | I64 | 否 | FFTS 基地址 | +| `$tcube_pipe` | HIVM_PipeAttr | 否 | Cube Core 等待的 Pipe | +| `$tvector_pipe` | HIVM_PipeAttr | 否 | Vector Core 等待的 Pipe | + +### SyncBlockSetOp / SyncBlockWaitOp 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$tcore_type` | HIVM_TCoreTypeAttr | 是 | Core 类型（CUBE/VECTOR） | +| `$tpipe` | HIVM_PipeAttr | 是 | 目标 Core 的 Pipe | +| `$pipe` | HIVM_PipeAttr | 是 | 当前操作的 Pipe | +| `$static_flag_id` | IntegerAttr | 否 | 静态 Flag ID | +| `$dynamic_flag_id` | I64 | 否 | 动态 Flag ID | +| `$ffts_base_addr` | I64 | 否 | FFTS 基地址（仅 SetOp） | +| `$tsync_instr_mode` | HIVM_SyncBlockInstrModeAttr | 否 | 同步指令模式，默认 INTRA_BLOCK_SYNCHRONIZATION | + +### SyncBlockMode 枚举 + +| 模式 | 值 | 说明 | 必需参数 | +|------|---|------|---------| +| ALL_CUBE | 0 | 所有 Cube Core 同步到同一点 | `tcube_pipe` | +| ALL_VECTOR | 1 | 所有 Vector Core 同步到同一点 | `tvector_pipe` | +| ALL_SUB_VECTOR | 2 | 所有 Sub-Vector 同步 | - | +| BARRIER_CUBE | 3 | Cube-Cube 屏障，lowering 为 barrier.pipe_all，仅复制到 AIC kernel | - | +| BARRIER_VECTOR | 4 | Vector-Vector 屏障，lowering 为 barrier.pipe_all，仅复制到 AIV kernel | - | +| ALL | 5 | 所有 AIC/AIV 同步到同一点 | `tvector_pipe` | + +### SyncBlockInstrMode 枚举 + +| 模式 | 值 | 说明 | +|------|---|------| +| INTER_BLOCK_SYNCHRONIZATION | 0 | 跨 Block 同步 | +| INTER_SUBBLOCK_SYNCHRONIZATION | 1 | 跨 Sub-Block 同步 | +| INTRA_BLOCK_SYNCHRONIZATION | 2 | Block 内同步（默认） | + +### 锁操作参数 + +| 操作 | 参数 | 说明 | +|------|------|------| +| create_sync_block_lock | `lockArg`: Optional\ | 创建锁内存区域，返回 memref\<1xi64\> | +| sync_block_lock | `lock_var`: MemRefRankOf<[I64], [1]> | 获取锁，阻塞直到 lock_var == block_idx | +| sync_block_unlock | `lock_var`: MemRefRankOf<[I64], [1]> | 释放锁，lock_var 递增并释放 | + +## IR 示例 + +### SyncBlockOp — ALL 模式 + +```mlir +%ffts_base_addr = arith.constant 0 : i64 +hivm.hir.sync_block[#hivm.sync_block_mode, 1 : i16] + ffts_base_addr = %ffts_base_addr + tcube_pipe=#hivm.pipe + tvector_pipe=#hivm.pipe +``` + +### SyncBlockOp — ALL_CUBE 模式 + +```mlir +hivm.hir.sync_block[#hivm.sync_block_mode, 1 : i16] + ffts_base_addr = %ffts_base_addr + tcube_pipe=#hivm.pipe +``` + +### SyncBlockSetOp — 静态 Flag ID + +```mlir +%ffts_base_addr = arith.constant 0 : i64 +hivm.hir.sync_block_set[#hivm.tcore_type, #hivm.pipe, #hivm.pipe] + flag = 1 + ffts_base_addr = %ffts_base_addr + sync_instr_mode = #hivm.sync_block_instr_mode +``` + +### SyncBlockSetOp — 动态 Flag ID + +```mlir +%flag_id = arith.constant 0 : i64 +hivm.hir.sync_block_set[#hivm.tcore_type, #hivm.pipe, #hivm.pipe] + flag = %flag_id + ffts_base_addr = %ffts_base_addr + sync_instr_mode = #hivm.sync_block_instr_mode +``` + +### SyncBlockWaitOp + +```mlir +hivm.hir.sync_block_wait[#hivm.tcore_type, #hivm.pipe, #hivm.pipe] flag = 1 +``` + +### 锁操作 + +```mlir +%lock = hivm.hir.create_sync_block_lock() : memref<1xi64> +hivm.hir.sync_block_lock lock_var(%lock : memref<1xi64>) +// ... 临界区代码 ... +hivm.hir.sync_block_unlock lock_var(%lock : memref<1xi64>) +``` + +### 锁操作 — 从外部内存创建 + +```mlir +%lock = hivm.hir.create_sync_block_lock() from %arg : from memref to memref<1xi64> +``` + +## IR 层约束与验证 + +1. **FFTS 基地址**：在 Ascend910B 及以上平台，`ffts_base_addr` 必须设置。 +2. **sync_block 限制**：只能在数据已搬运到 GM 后使用。 +3. **Flag ID**：FFTS 收集特定 flag_id 后将其设置回 Block 组中的 Block，实现同步。 +4. **tcube_pipe / tvector_pipe**：ALL_CUBE 模式需要 `tcube_pipe`，ALL_VECTOR/ALL 模式需要 `tvector_pipe`。 +5. **锁操作**：`sync_block_lock` 阻塞直到 `lock_var == block_idx`，`sync_block_unlock` 递增并释放 `lock_var`。 +6. **sync_block_set/wait 配对**：必须使用相同的 (tcore_type, tpipe, pipe, flag_id) 组合。 + +## 常见问题 + +**Q: sync_block 和 sync_block_set/wait 的区别？** +A: sync_block 是高层封装，自动处理 set/wait 配对；sync_block_set/wait 是低层接口，需要手动配对。通常推荐使用 sync_block。 + +**Q: 什么时候需要设置 ffts_base_addr？** +A: 在 Ascend910B 及以上平台，跨 Block 同步必须设置 FFTS 基地址。通常通过 `hir.set_ffts_base_addr` 在函数入口设置。 + +**Q: 锁操作的使用场景？** +A: 锁操作用于需要互斥访问共享资源的场景，如多个 Block 写入同一 GM 区域。`create_sync_block_lock` 分配锁内存，`sync_block_lock` 获取锁，`sync_block_unlock` 释放锁。 + +## 相关文档 + +- 源码参考：[HIVMSynchronizationOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMSynchronizationOps.td#L78-L244) +- 测试用例：[sync-ops.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/sync-ops.mlir) +- SyncBlockMode 枚举：[06-Attributes-Types/01-enumerations.md](../06-Attributes-Types/01-enumerations.md) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/03-unit-flag.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/03-unit-flag.md new file mode 100644 index 00000000..8eb2e18d --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/03-unit-flag.md @@ -0,0 +1,126 @@ +# UnitFlag 同步机制 + +> 关键词：UnitFlag, DISABLED, RESERVED, ENABLED_WITHOUT_UPDATE, ENABLED_WITH_UPDATE, UnitFlagEnabledInterface + +## 概述 + +UnitFlag 是 HIVM 中嵌入在宏操作（如 mmadL1）内的同步机制，专门处理循环中"至少执行一次"的依赖场景。在典型的 Split-K 矩阵乘法中，L0C 累加器需要在首次迭代清零、后续迭代累加。但如果循环可能不执行（循环次数为 0），则 set_flag/wait_flag 的配对会被打破，导致同步错误。 + +UnitFlag 通过条件化同步解决了这个问题：当 `unit_flag_cond` 为 true 时，即使循环未执行，同步也能正确完成。 + +## UnitFlag 四种模式 + +从 [HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L512-L534) 提取： + +| 模式 | 枚举值 | 二进制 | 说明 | +|------|--------|--------|------| +| DISABLED | 0 | 0b00 | 禁用 UnitFlag，不使用条件同步 | +| RESERVED | 1 | 0b01 | 保留，当前未使用 | +| ENABLED_WITHOUT_UPDATE | 2 | 0b10 | 启用 UnitFlag 但不更新标志计数器 | +| ENABLED_WITH_UPDATE | 3 | 0b11 | 启用 UnitFlag 并更新标志计数器 | + +### 模式详解 + +#### DISABLED（0b00） + +默认模式，不使用 UnitFlag 同步。操作按常规方式参与同步分析。 + +#### RESERVED（0b01） + +保留模式，当前未使用，预留给未来扩展。 + +#### ENABLED_WITHOUT_UPDATE（0b10） + +启用 UnitFlag 但不更新标志计数器。适用于以下场景： +- 循环可能不执行，但需要保证同步正确性 +- UnitFlag 条件由外部控制，不需要操作自身更新 + +#### ENABLED_WITH_UPDATE（0b11） + +启用 UnitFlag 并更新标志计数器。适用于以下场景： +- 循环至少执行一次时，UnitFlag 条件为 true +- 操作自身需要参与标志计数器的更新 + +## UnitFlagEnabledInterface + +从 [HIVMInterfaces.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMInterfaces.td#L680-L761) 提取： + +实现 `UnitFlagEnabledInterface` 的操作可以访问以下方法： + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getUnitFlagModes()` | `optional>` | 获取 UnitFlag 模式数组 | +| `setUnitFlagModes(SmallVector)` | void | 设置 UnitFlag 模式 | +| `getUnitFlagConditions()` | `optional>` | 获取 UnitFlag 条件值 | +| `setUnitFlagConditions(SmallVector)` | void | 设置 UnitFlag 条件值 | +| `getUnitFlagModeLibValue(PatternRewriter&)` | Value | 获取传递给标准库调用的 UnitFlag 模式值（0b00/0b10/0b11） | + +## 在 mmadL1 中的使用 + +mmadL1 操作通过以下参数支持 UnitFlag： + +### unit_flag_cond + +```mlir +// Variadic 类型，可选 +// 提供 i1 条件值，用于控制 UnitFlag 是否启用 +// 通常为循环是否至少执行一次的条件 +hivm.hir.mmadL1 ins(...) + outs(...) + unit_flag_cond(%cond) +``` + +### unit_flag_mode + +```mlir +// UnitFlagArrayAttr 类型，可选 +// 指定每个输出 Tensor 的 UnitFlag 模式 +hivm.hir.mmadL1 ins(...) + outs(...) + unit_flag_mode([#hivm.unit_flag]) +``` + +## IR 示例 + +### Split-K 循环中的 UnitFlag + +```mlir +%mc = memref.alloc() : memref<256x256xf32> +%start = arith.constant 0 : index +%end = arith.constant 1024 : index +%step = arith.constant 128 : index +scf.for %arg0 = %start to %end step %step { + %ma = memref.alloc() : memref<256x128xf16> + %mb = memref.alloc() : memref<128x256xf16> + %init_condition = arith.cmpi eq, %arg0, %start : index + hivm.hir.mmadL1 ins(%ma, %mb, %init_condition, %c256, %c128, %c256 : + memref<256x128xf16>, memref<128x256xf16>, i1, index, index, index) + outs(%mc : memref<256x256xf32>) + unit_flag_mode([#hivm.unit_flag]) + unit_flag_cond(%loop_alive_cond) +} +``` + +## IR 层约束与验证 + +1. **UnitFlagEnabledInterface 验证**：实现该接口的操作必须正确声明 `unit_flag_mode` 和 `unit_flag_cond` 参数。 +2. **unit_flag_mode 数组长度**：应与输出 Tensor 数量一致。 +3. **unit_flag_cond**：为 i1 类型的 Variadic 操作数，通常提供 0 或 1 个条件值。 +4. **库调用值**：`getUnitFlagModeLibValue()` 返回的值只能是 0b00、0b10 或 0b11 之一，RESERVED 模式不传递给库调用。 + +## 常见问题 + +**Q: 什么时候需要使用 UnitFlag？** +A: 当宏操作（如 mmadL1）在循环中使用，且循环可能不执行时，需要 UnitFlag 来保证同步正确性。InjectSync Pass 在启用 `enable-unit-flag` 选项时会自动插入 UnitFlag。 + +**Q: ENABLED_WITHOUT_UPDATE 和 ENABLED_WITH_UPDATE 的区别？** +A: ENABLED_WITH_UPDATE 会让操作更新标志计数器，适用于操作确实执行的场景；ENABLED_WITHOUT_UPDATE 不更新计数器，适用于操作可能不执行但需要保持同步配对的场景。 + +**Q: UnitFlag 如何影响 InjectSync Pass？** +A: 当 `enable-unit-flag=true` 时，InjectSync Pass 会在宏操作中自动设置 UnitFlag 模式和条件，确保循环不执行时同步仍然正确。 + +## 相关文档 + +- 源码参考：[HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L512-L534) +- 接口定义：[HIVMInterfaces.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMInterfaces.td#L680-L761) +- 宏操作：[03-Macro-Operations/01-mmad-l1.md](../03-Macro-Operations/01-mmad-l1.md) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/04-sync-injection.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/04-sync-injection.md new file mode 100644 index 00000000..569a841a --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/04-Synchronization/04-sync-injection.md @@ -0,0 +1,183 @@ +# 同步注入 Pass — InjectSync / GraphSyncSolver + +> 关键词：InjectSync, GraphSyncSolver, SyncMode, Event ID Allocation, IR 变换 + +## 概述 + +HIVM 编译器提供自动同步注入机制，在用户编写的 IR 中自动插入 `set_flag`/`wait_flag` 等同步操作，保证多 Pipeline 执行的正确性。用户通常不需要手动编写同步操作，编译器会根据操作间的数据依赖自动注入。 + +HIVM 提供两种同步注入 Pass： + +1. **InjectSync Pass**：基于启发式规则的同步注入，按操作顺序分析数据依赖并插入同步操作。 +2. **GraphSyncSolver Pass**：基于图的同步求解器，构建操作依赖图后求解最优同步方案，通常能生成更优的同步代码。 + +## InjectSync Pass + +### Pass 定义 + +从 [Passes.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/Transforms/Passes.td#L163-L183) 提取： + +``` +def InjectSync : Pass<"hivm-inject-sync", "func::FuncOp"> { + let summary = "auto inject sync"; + let options = [ + Option<"syncMode", "sync-mode", "hivm::SyncMode", + "hivm::SyncMode::NORMAL", + "inject sync mode">, + Option<"enableUnitFlag", "enable-unit-flag", "bool", "false", + "Enable unit-flag modes for synchronization">, + Option<"assumeAliveLoops", "assume-alive-loops", "bool", "false", + "Assume that all loops will execute at least once."> + ]; +} +``` + +### Pass 选项 + +| 选项 | 类型 | 默认值 | 说明 | +|------|------|--------|------| +| `sync-mode` | SyncMode | NORMAL | 同步模式：NORMAL（正常注入）或 BARRIERALL（注入 pipe_barrier，仅用于调试） | +| `enable-unit-flag` | bool | false | 启用 UnitFlag 模式，处理循环可能不执行的同步场景 | +| `assume-alive-loops` | bool | false | 假设所有循环至少执行一次，简化同步分析 | + +### SyncMode 枚举 + +| 模式 | 说明 | +|------|------| +| NORMAL | 正常模式，根据数据依赖插入 set_flag/wait_flag | +| BARRIERALL | 调试模式，在每个操作前后插入 pipe_barrier | + +### IR 变换效果 + +#### 变换前 + +```mlir +func.func @example(%src : memref<16x16xf16, #hivm.address_space>, + %dst_l1 : memref<16x16xf16>, + %dst_l0c : memref<16x16xf32>) { + hivm.hir.load ins(%src : ...) outs(%dst_l1 : ...) + hivm.hir.mmadL1 ins(%a, %b, %init, ...) outs(%dst_l0c : ...) + hivm.hir.fixpipe ins(%dst_l0c : ...) outs(%res : ...) + return +} +``` + +#### 变换后（NORMAL 模式） + +```mlir +func.func @example(%src : memref<16x16xf16, #hivm.address_space>, + %dst_l1 : memref<16x16xf16>, + %dst_l0c : memref<16x16xf32>) { + hivm.hir.load ins(%src : ...) outs(%dst_l1 : ...) + hivm.hir.set_flag [#hivm.pipe, #hivm.pipe, #hivm.event] + hivm.hir.wait_flag [#hivm.pipe, #hivm.pipe, #hivm.event] + hivm.hir.mmadL1 ins(%a, %b, %init, ...) outs(%dst_l0c : ...) + hivm.hir.set_flag [#hivm.pipe, #hivm.pipe, #hivm.event] + hivm.hir.wait_flag [#hivm.pipe, #hivm.pipe, #hivm.event] + hivm.hir.fixpipe ins(%dst_l0c : ...) outs(%res : ...) + return +} +``` + +#### 变换后（BARRIERALL 模式，调试用） + +```mlir +func.func @example(...) { + hivm.hir.load ins(%src : ...) outs(%dst_l1 : ...) + hivm.hir.pipe_barrier [#hivm.pipe] + hivm.hir.pipe_barrier [#hivm.pipe] + hivm.hir.mmadL1 ins(%a, %b, %init, ...) outs(%dst_l0c : ...) + hivm.hir.pipe_barrier [#hivm.pipe] + hivm.hir.pipe_barrier [#hivm.pipe] + hivm.hir.fixpipe ins(%dst_l0c : ...) outs(%res : ...) + return +} +``` + +## GraphSyncSolver Pass + +### Pass 定义 + +从 [Passes.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/Transforms/Passes.td#L195-L213) 提取： + +``` +def GraphSyncSolver : Pass<"hivm-graph-sync-solver", "func::FuncOp"> { + let summary = "graph sync solver"; + let options = [ + Option<"enableUnitFlag", "enable-unit-flag", "bool", "false", + "Enable unit-flag modes for synchronization">, + Option<"enableTesterMode", "enable-tester-mode", "bool", "false", + "Enable sync-tester mode">, + ListOption<"syncTesterOptions", "sync-tester-options", "int64_t", + "Sync-tester options"> + ]; +} +``` + +### Pass 选项 + +| 选项 | 类型 | 默认值 | 说明 | +|------|------|--------|------| +| `enable-unit-flag` | bool | false | 启用 UnitFlag 模式 | +| `enable-tester-mode` | bool | false | 启用同步测试模式 | +| `sync-tester-options` | list\ | [] | 同步测试选项（num_runs, init_seed, num_ops, num_ptrs, enable-multibuffer） | + +### 工作原理 + +GraphSyncSolver 的工作流程： + +1. **IR 转换**：将 MLIR IR 转换为内部同步图表示（SyncSolverIR） +2. **图构建**：构建操作依赖图，节点为操作，边为数据依赖 +3. **求解**：在图上求解最优同步方案，包括 Event ID 分配和同步点选择 +4. **代码生成**：将求解结果转换回 MLIR IR，插入 set_flag/wait_flag 操作 + +### 与 InjectSync 的对比 + +| 特性 | InjectSync | GraphSyncSolver | +|------|-----------|----------------| +| 算法 | 启发式规则 | 图求解 | +| Event ID 使用 | 可能较多 | 更优化 | +| 同步粒度 | 操作级 | 可优化到更细粒度 | +| 编译速度 | 较快 | 较慢 | +| 同步质量 | 基本正确 | 更优，减少不必要的同步 | +| 调试支持 | BARRIERALL 模式 | Tester 模式 | + +## Pipeline 中的选择 + +从 [HIVMPipelines.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/Pipelines/HIVMPipelines.cpp#L81-L98) 可以看到，Pipeline 中的选择逻辑： + +- 如果 `enableHIVMGraphSyncSolver=true` 且 `enableHIVMInjectBarrierAllSync=false`：使用 GraphSyncSolver +- 否则：使用 InjectSync + +## InjectBlockSync Pass + +除了 Pipe 内同步，还有跨 Block 同步注入 Pass： + +``` +def InjectBlockSync : Pass<"hivm-inject-block-sync", "func::FuncOp"> { + let options = [ + Option<"blockAllSync", "block-all-sync", "bool", "false", + "Enable inject all block sync">, + Option<"assumeAliveLoops", "assume-alive-loops", "bool", "false", + "Assume that all loops will execute at least once."> + ]; +} +``` + +## 常见问题 + +**Q: 应该使用 InjectSync 还是 GraphSyncSolver？** +A: 默认使用 InjectSync。如果需要更优的同步方案（减少同步开销），可以启用 GraphSyncSolver。GraphSyncSolver 编译时间较长，但生成的同步代码通常更高效。 + +**Q: assume-alive-loops 选项的作用？** +A: 当设置为 true 时，编译器假设所有循环至少执行一次，可以简化同步分析，避免插入不必要的 UnitFlag。但如果不满足假设，可能导致同步错误。 + +**Q: BARRIERALL 模式有什么用？** +A: BARRIERALL 是调试模式，在每个操作前后插入 pipe_barrier。虽然性能很差，但可以排除同步问题导致的错误，帮助定位问题。 + +## 相关文档 + +- 源码参考：[Passes.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/Transforms/Passes.td#L163-L213) +- InjectSync 实现：[InjectSync.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/Transforms/InjectSync/InjectSync.cpp) +- GraphSyncSolver 实现：[GraphSyncSolver.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/Transforms/GraphSyncSolver/GraphSyncSolver.cpp) +- Pipeline 配置：[HIVMPipelines.cpp](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/lib/Dialect/HIVM/Pipelines/HIVMPipelines.cpp) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/05-Custom-Operations/00-overview.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/05-Custom-Operations/00-overview.md new file mode 100644 index 00000000..2b1abff2 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/05-Custom-Operations/00-overview.md @@ -0,0 +1,76 @@ +# Custom 操作机制总览 + +> 关键词：Custom Op, BuiltinInfo, __builtin_gather_load, __builtin_index_select, 自定义操作 + +## 概述 + +HIVM Custom 操作机制提供了一种通用的扩展接口，允许用户在现有操作无法满足需求时编写自定义实现。Custom 操作支持两种粒度： + +1. **hir.custom**：单 Pipe 自定义操作，在单个 Pipeline 上执行 +2. **hir.custom_macro**：跨 Pipe 自定义宏操作，涉及多个 Pipeline + +Custom 操作通过 `name` 属性标识操作类型，编译器根据 name 查找对应的实现。部分 name 以 `__builtin` 开头的是内置操作，编译器自带实现；其他 name 需要用户提供实现。 + +## Builtin 机制 + +内置操作（Builtin）是编译器自带的 Custom 操作实现，name 以 `__builtin` 开头。当前支持的内置操作： + +| Builtin Name | 说明 | Core Type | Pipe | +|-------------|------|-----------|------| +| `__builtin_gather_load` | Gather Load 操作 | VECTOR | PIPE_V | +| `__builtin_index_select` | Index Select 操作 | VECTOR | PIPE_V | + +### BuiltinInfo 结构 + +每个 Builtin 操作都有对应的 `BuiltinInfo` 结构，包含： + +| 字段 | CustomOp | CustomMacroOp | +|------|----------|---------------| +| coreType | TCoreType | TCoreType | +| pipe / inPipe | PIPE | PIPE | +| - | - | outPipe (PIPE) | +| vfMode | VFMode | VFMode | +| getOpLibraryCallName | function\ | function\ | +| gmAddrArgsIndices | SmallVector\ | SmallVector\ | + +## Custom 操作属性 + +Custom 操作通过属性传递必要信息： + +### 必需属性 + +| 属性 | 说明 | +|------|------| +| `hivm.tcore_type` | 执行的 Core 类型，参见 TCoreTypeAttr | +| `hivm.vf_mode` | Vector 单元运行模式，参见 VFModeAttr（Cube Core 时忽略） | + +### CustomOp 特有属性 + +| 属性 | 说明 | +|------|------| +| `hivm.pipe` | 执行的 Pipe，参见 PipeAttr | + +### CustomMacroOp 特有属性 + +| 属性 | 说明 | +|------|------| +| `hivm.pipe_in` | 输入 Pipe | +| `hivm.pipe_out` | 输出 Pipe | + +### 可选属性 + +| 属性 | 说明 | +|------|------| +| `gm_addr_args_indices` | I32 数组，指示哪些参数包含 GM 纯地址，lowering 时保留地址值 | + +## 操作列表 + +| 操作 | 助记符 | 详细文档 | +|------|--------|---------| +| CustomOp | `hir.custom` | [01-custom-op.md](01-custom-op.md) | +| CustomMacroOp | `hir.custom_macro` | [02-custom-macro-op.md](02-custom-macro-op.md) | + +## 相关文档 + +- 源码参考：[HIVMOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMOps.td#L904-L1167) +- 测试用例：[custom-op.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/custom-op.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/05-Custom-Operations/01-custom-op.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/05-Custom-Operations/01-custom-op.md new file mode 100644 index 00000000..7445dd64 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/05-Custom-Operations/01-custom-op.md @@ -0,0 +1,160 @@ +# hir.custom — 自定义单 Pipe 操作 + +> 关键词：custom, CustomOp, Builtin, __builtin_gather_load, __builtin_index_select, SinglePipe + +## 概述 + +`hir.custom` 是 HIVM 方言中的自定义单 Pipe 操作，允许用户在现有操作无法满足需求时编写自定义实现。该操作在单个 Pipeline 上执行，通过 `name` 属性标识操作类型。 + +Custom 操作适用于以下场景： +1. 现有操作无法实现所需功能 +2. 现有操作可以实现但性能不佳 +3. 需要私有操作 + +内置操作（name 以 `__builtin` 开头）由编译器自带实现，用户无需指定属性即可使用。 + +> Python API 对应：Triton 的自定义操作可以通过 `tl.extern` 等机制映射为 hir.custom。 + +## IR 操作定义 + +从 [HIVMOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMOps.td#L1069-L1103) 提取： + +``` +def CustomOp : HIVM_CustomOp<"custom", [SinglePipeOpTrait]> { + let arguments = (ins StrAttr:$name, + Variadic:$inputs, + Variadic:$outputs); + let results = (outs Variadic:$results); +} +``` + +基类 `HIVM_CustomOp` 定义（[HIVMOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMOps.td#L904-L967)）： + +``` +class HIVM_CustomOp traits = []> + : HIVM_StructuredOp, + HIVMInferCoreTypeInterface], + traits)> { + dag args = (ins StrAttr:$name, Variadic:$inputs, + Variadic:$outputs); + dag res = (outs Variadic:$results); +} +``` + +## 参数说明 + +### 输入操作数（ins） + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$name` | StrAttr | 是 | 操作名称，`__builtin_` 前缀为内置操作 | +| `$inputs` | Variadic\ | 是 | 输入参数（可变数量） | +| `$outputs` | Variadic\ | 是 | 输出/初始化参数（DestinationStyle） | + +### 输出结果（outs） + +| 参数 | 类型 | 说明 | +|------|------|------| +| `$results` | Variadic\ | 结果值 | + +### 必需属性 + +| 属性 | 类型 | 说明 | +|------|------|------| +| `hivm.tcore_type` | TCoreTypeAttr | 执行的 Core 类型 | +| `hivm.pipe` | PipeAttr | 执行的 Pipeline | +| `hivm.vf_mode` | VFModeAttr | Vector 运行模式（Cube Core 时忽略） | + +### 可选属性 + +| 属性 | 类型 | 说明 | +|------|------|------| +| `gm_addr_args_indices` | DenseI32ArrayAttr | 指示哪些参数包含 GM 纯地址 | + +### 额外类方法 + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getCoreType()` | `optional` | 获取 Core 类型 | +| `setCoreType(TCoreType)` | void | 设置 Core 类型 | +| `getPipe()` | PIPE | 获取 Pipe | +| `setPipe(PIPE)` | void | 设置 Pipe | +| `getVFMode()` | `optional` | 获取 VF 模式 | +| `setVFMode(VFMode)` | void | 设置 VF 模式 | +| `getGMAddrArgsIndices()` | `optional>` | 获取 GM 地址参数索引 | +| `setGMAddrArgsIndices(SmallVector)` | void | 设置 GM 地址参数索引 | +| `isBuiltin()` | bool | 判断是否为内置操作 | +| `getDpsInitsMutable()` | MutableOperandRange | DestinationStyleOpInterface | + +### 内置操作常量 + +```cpp +static constexpr StringLiteral kBuiltinGatherLoadName = "__builtin_gather_load"; +static constexpr StringLiteral kBuiltinIndexSelectName = "__builtin_index_select"; +``` + +## IR 示例 + +### 自定义操作 + +```mlir +%empty = tensor.empty() : tensor<3x3xf32> +%0 = hivm.hir.custom + { hivm.tcore_type = #hivm.tcore_type, hivm.pipe = #hivm.pipe, hivm.vf_mode = #hivm.vf_mode } + "my_custom_op" + ins(%arg0, %arg1, %c4_i64, %c0_i32, %c2_i64, %c1_i64, %c2_i32, %c2_i32, %c0_i32, %c0_i32 + : memref, tensor<3x3xi64>, i64, i32, i64, i64, i32, i32, i32, i32) + outs(%empty : tensor<3x3xf32>) -> tensor<3x3xf32> +``` + +### 内置操作（无需指定属性） + +```mlir +%empty = tensor.empty() : tensor<3x3xf32> +%0 = hivm.hir.custom + "__builtin_gather_load" + ins(%arg0, %arg1, %c4_i64, %c0_i32, %c2_i64, %c1_i64, %c2_i32, %c2_i32, %c0_i32, %c0_i32 + : memref, tensor<3x3xi64>, i64, i32, i64, i64, i32, i32, i32, i32) + outs(%empty : tensor<3x3xf32>) -> tensor<3x3xf32> +``` + +### __builtin_index_select + +```mlir +%0 = tensor.empty() : tensor<1x4x32xf32> +%1 = hivm.hir.custom + {extra_attr = "srcStrideLength=3", hivm.vf_mode = #hivm.vf_mode} + "__builtin_index_select" + ins(%arg0, %arg1, %c0_i32, %c9000_i64, %c0_i32, %c0_i32, %c0_i32, %c0_i32, %c0_i32, %c0_i32, %c32_i32, %c4000_i32, %c32_i32 + : memref, tensor<16x400xi32>, i32, i64, i32, i32, i32, i32, i32, i32, i32, i32, i32) + outs(%0 : tensor<1x4x32xf32>) -> tensor<1x4x32xf32> +``` + +## IR 层约束与验证 + +1. **SinglePipeOpTrait**：操作在单个 Pipeline 上执行。 +2. **AttrSizedOperandSegments**：inputs 和 outputs 的数量由属性决定。 +3. **DestinationStyleOpInterface**：outputs 作为 DPS init 操作数。 +4. **Core Type**：必须通过 `hivm.tcore_type` 属性指定（内置操作可省略，编译器自动推断）。 +5. **Pipe**：必须通过 `hivm.pipe` 属性指定（内置操作可省略）。 +6. **VFMode**：必须通过 `hivm.vf_mode` 属性指定（内置操作可省略，Cube Core 时忽略）。 +7. **Builtin 验证**：内置操作编译器会检查参数正确性并规范化属性。 + +## 常见问题 + +**Q: 自定义操作如何提供实现？** +A: 当前自定义操作的用户实现机制尚未完全开放（标记为 TODO）。内置操作由编译器自动链接到模板库。 + +**Q: __builtin_gather_load 和 hir.gather_load 的区别？** +A: `__builtin_gather_load` 是通过 custom 操作机制实现的 gather load，而 `hir.gather_load` 是独立的 HIVM 操作。两者功能类似，但 custom 版本更灵活，可以通过属性配置。 + +**Q: gm_addr_args_indices 的用途？** +A: 指示哪些输入参数包含 GM 的纯地址值。Lowering Pass 在处理这些参数时会保留地址值，而不是尝试从 memref 中提取数据。 + +## 相关文档 + +- 源码参考：[HIVMOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMOps.td#L1069-L1103) +- 测试用例：[custom-op.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/custom-op.mlir) +- CustomMacroOp：[02-custom-macro-op.md](02-custom-macro-op.md) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/05-Custom-Operations/02-custom-macro-op.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/05-Custom-Operations/02-custom-macro-op.md new file mode 100644 index 00000000..49ead6d2 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/05-Custom-Operations/02-custom-macro-op.md @@ -0,0 +1,135 @@ +# hir.custom_macro — 自定义跨 Pipe 宏操作 + +> 关键词：custom_macro, CustomMacroOp, MacroOpTrait, InPipe, OutPipe, Builtin + +## 概述 + +`hir.custom_macro` 是 HIVM 方言中的自定义跨 Pipe 宏操作，与 `hir.custom` 类似，但涉及多个 Pipeline。该操作通过 `MacroOpTrait` 标记为宏操作，使用 `hivm.pipe_in` 和 `hivm.pipe_out` 属性分别指定输入和输出 Pipeline。 + +CustomMacroOp 适用于需要跨 Pipeline 协作的场景，如数据加载+计算+后处理的融合操作。 + +## IR 操作定义 + +从 [HIVMOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMOps.td#L1105-L1167) 提取： + +``` +def CustomMacroOp : HIVM_CustomOp<"custom_macro", [MacroOpTrait]> { + let arguments = (ins StrAttr:$name, + Variadic:$inputs, + Variadic:$outputs); + let results = (outs Variadic:$results); +} +``` + +## 参数说明 + +### 输入操作数（ins） + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$name` | StrAttr | 是 | 操作名称 | +| `$inputs` | Variadic\ | 是 | 输入参数 | +| `$outputs` | Variadic\ | 是 | 输出/初始化参数 | + +### 输出结果（outs） + +| 参数 | 类型 | 说明 | +|------|------|------| +| `$results` | Variadic\ | 结果值 | + +### 必需属性 + +| 属性 | 类型 | 说明 | +|------|------|------| +| `hivm.tcore_type` | TCoreTypeAttr | 执行的 Core 类型 | +| `hivm.pipe_in` | PipeAttr | 输入 Pipeline | +| `hivm.pipe_out` | PipeAttr | 输出 Pipeline | +| `hivm.vf_mode` | VFModeAttr | Vector 运行模式 | + +### 可选属性 + +| 属性 | 类型 | 说明 | +|------|------|------| +| `gm_addr_args_indices` | DenseI32ArrayAttr | GM 地址参数索引 | + +### 与 CustomOp 的差异 + +| 特性 | CustomOp | CustomMacroOp | +|------|----------|---------------| +| Trait | SinglePipeOpTrait | MacroOpTrait | +| Pipe 属性 | `hivm.pipe`（单个） | `hivm.pipe_in` + `hivm.pipe_out`（两个） | +| Pipeline 数量 | 单个 | 跨多个 | +| 同步需求 | Pipe 内同步 | 跨 Pipe 同步 | +| BuiltinInfo 字段 | coreType, pipe, vfMode | coreType, inPipe, outPipe, vfMode | + +### 额外类方法 + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getInPipe()` | PIPE | 获取输入 Pipe | +| `setInPipe(PIPE)` | void | 设置输入 Pipe | +| `getOutPipe()` | PIPE | 获取输出 Pipe | +| `setOutPipe(PIPE)` | void | 设置输出 Pipe | +| `getCoreType()` | `optional` | 获取 Core 类型 | +| `setCoreType(TCoreType)` | void | 设置 Core 类型 | +| `getVFMode()` | `optional` | 获取 VF 模式 | +| `setVFMode(VFMode)` | void | 设置 VF 模式 | +| `isBuiltin()` | bool | 判断是否为内置操作 | + +### Pipe 属性名称常量 + +```cpp +static constexpr StringLiteral inPipeName = "hivm.pipe_in"; +static constexpr StringLiteral outPipeName = "hivm.pipe_out"; +``` + +## IR 示例 + +### 自定义宏操作 + +```mlir +%empty = tensor.empty() : tensor<3x3xf32> +%0 = hivm.hir.custom_macro + { hivm.tcore_type = #hivm.tcore_type, hivm.vf_mode = #hivm.vf_mode, + hivm.pipe_in = #hivm.pipe, hivm.pipe_out = #hivm.pipe } + "my_custom_op" + ins(%arg0, %arg1, %c4_i64, %c0_i32, %c2_i64, %c1_i64, %c2_i32, %c2_i32, %c0_i32, %c0_i32 + : memref, tensor<3x3xi64>, i64, i32, i64, i64, i32, i32, i32, i32) + outs(%empty : tensor<3x3xf32>) -> tensor<3x3xf32> +``` + +### 内置宏操作 + +```mlir +%empty = tensor.empty() : tensor<3x3xf32> +%0 = hivm.hir.custom_macro + "__builtin_gather_load" + ins(%arg0, %arg1, %c4_i64, %c0_i32, %c2_i64, %c1_i64, %c2_i32, %c2_i32, %c0_i32, %c0_i32 + : memref, tensor<3x3xi64>, i64, i32, i64, i64, i32, i32, i32, i32) + outs(%empty : tensor<3x3xf32>) -> tensor<3x3xf32> +``` + +## IR 层约束与验证 + +1. **MacroOpTrait**：操作被标记为宏操作，编译器在同步分析中特殊处理。 +2. **跨 Pipe 同步**：由于涉及多个 Pipeline，InjectSync/GraphSyncSolver Pass 会在操作前后插入跨 Pipe 同步操作。 +3. **pipe_in / pipe_out**：必须指定输入和输出 Pipeline，编译器据此生成同步操作。 +4. **DestinationStyleOpInterface**：outputs 作为 DPS init 操作数。 +5. **Builtin 验证**：内置操作编译器会检查参数正确性。 + +## 常见问题 + +**Q: 什么时候用 custom_macro 而不是 custom？** +A: 当操作涉及多个 Pipeline（如数据加载+计算+后处理）时使用 custom_macro。单 Pipeline 操作使用 custom。 + +**Q: pipe_in 和 pipe_out 如何选择？** +A: pipe_in 是数据进入的 Pipeline（如 PIPE_MTE2 表示从 GM 加载），pipe_out 是数据输出的 Pipeline（如 PIPE_V 表示在 Vector Core 处理）。 + +**Q: custom_macro 的同步如何处理？** +A: InjectSync/GraphSyncSolver Pass 会根据 pipe_in 和 pipe_out 自动在操作前后插入 set_flag/wait_flag 同步操作。 + +## 相关文档 + +- 源码参考：[HIVMOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMOps.td#L1105-L1167) +- 测试用例：[custom-op.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/custom-op.mlir) +- CustomOp：[01-custom-op.md](01-custom-op.md) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/00-overview.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/00-overview.md new file mode 100644 index 00000000..0cda58e9 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/00-overview.md @@ -0,0 +1,70 @@ +# HIVM 属性类型系统总览 + +> 关键词：Attributes, Types, Enumerations, DataLayout, AddressSpace, Pipe, CoreType + +## 概述 + +HIVM 属性类型系统是 AscendNPU-IR 中描述硬件特性、存储层次、执行单元等关键信息的核心机制。通过 MLIR 的属性系统，HIVM 将 NPU 的硬件约束编码到 IR 中，使编译器能够在类型检查和 Lowering 阶段正确处理这些约束。 + +HIVM 属性类型系统分为以下几类： + +1. **枚举属性**：描述有限的离散值集合，如 IteratorType、DataLayout、AddressSpace、Pipe 等 +2. **参数化属性**：携带额外参数的复合属性，如 DataLayoutAttr（含 transpose/fractalSizes）、AddressSpaceAttr、BlockMappingAttr 等 +3. **类型约束**：MemRef/Tensor/Vector 在 HIVM 中的特殊约束 +4. **接口**：HIVMStructuredOpInterface、HIVMCoreTypeInterface、HIVMUnitFlagEnabledInterface 等 +5. **Trait**：操作行为约束，如 MacroOpTrait、CoreTypeTrait、ElementwiseNaryOpTrait 等 + +## 属性分类 + +### 枚举属性（30+ 个） + +枚举属性是 HIVM 中最常用的属性类型，每个枚举对应一个硬件特性或操作模式。详见 [01-enumerations.md](01-enumerations.md)。 + +### 参数化属性 + +| 属性 | 助记符 | 参数 | 说明 | +|------|--------|------|------| +| DataLayoutAttr | `data_layout` | data_layout, transpose?, fractalSizes? | 数据布局映射 | +| AddressSpaceAttr | `address_space` | address_space | 地址空间映射 | +| PipeAttr | `pipe` | pipe | Pipeline 标识 | +| TCoreTypeAttr | `tcore_type` | tcoretype | 操作 Core 类型 | +| TFuncCoreTypeAttr | `func_core_type` | funcCoreType | 函数 Core 类型 | +| TModuleCoreTypeAttr | `module_core_type` | moduleCoreType | 模块 Core 类型 | +| EventAttr | `event` | event | Event ID | +| UnitFlagAttr | `unit_flag` | unit_flag | UnitFlag 模式 | +| SyncBlockModeAttr | `sync_block_mode` | sync_mode | 同步 Block 模式 | +| SyncBlockInstrModeAttr | `sync_block_instr_mode` | sync_instr_mode | 同步指令模式 | +| ReduceOpAttr | `reduce_op` | reduce_op | 归约操作类型 | +| BlockMappingAttr | `block` | order? | Block 映射 | +| SubBlockMappingAttr | `sub_block` | sub_block | Sub-Block 映射 | + +详见 [02-parameterized-attrs.md](02-parameterized-attrs.md)。 + +### 标记属性 + +| 属性 | 助记符 | 说明 | +|------|--------|------| +| MultiBufferAttr | `multi_buffer` | 多缓冲标记 | +| TPartOfMixAttr | `part_of_mix` | Mix Kernel 组成部分 | +| VFAttr | `vector_function` | Vector 函数标记 | +| HasAliaScopesAttr | `has_alias_scopes` | 别名作用域标记 | +| TightlyCoupledBufferAttr | `tightly_coupled_buffer` | 紧耦合缓冲区（含 id 参数） | +| MemUniqueAttr | `mem_unique` | 唯一内存规划 | +| ParallelLoopAttr | `parallel_loop` | 可并行化循环标记 | +| UnlikelyConditionAttr | `unlikely_condition` | 不可能条件标记 | +| SharedMemoryAttr | `shared_memory` | SIMT VF 共享内存 | + +## 文档列表 + +| 文档 | 内容 | +|------|------| +| [01-enumerations.md](01-enumerations.md) | 所有枚举属性速查（完整枚举值列表） | +| [02-parameterized-attrs.md](02-parameterized-attrs.md) | 参数化属性详细定义 | +| [03-type-system.md](03-type-system.md) | MemRef/Tensor/Vector 在 HIVM 中的约束 | +| [04-interfaces.md](04-interfaces.md) | HIVM 接口方法列表 | +| [05-traits.md](05-traits.md) | 所有 Trait 定义和适用操作 | + +## 相关文档 + +- 源码参考：[HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) +- 测试用例：[attribute.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/attribute.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/01-enumerations.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/01-enumerations.md new file mode 100644 index 00000000..a6e07824 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/01-enumerations.md @@ -0,0 +1,467 @@ +# HIVM 枚举属性速查 + +> 关键词：Enumeration, IteratorType, DataLayout, AddressSpace, Pipe, CoreType, Event, UnitFlag, SyncBlockMode, ReduceOperation, AtomicKind + +## 概述 + +本文档列出 HIVM 方言中所有枚举属性的完整枚举值，是 Agent 最常查阅的参考。所有枚举值均从 [HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) 精确提取。 + +枚举属性在 MLIR IR 中的格式为 `#hivm.<>`，例如 `#hivm.pipe`。 + +--- + +## IteratorType（12 值） + +源码：[HIVMAttrs.td#L51-L78](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L51-L78) + +| 枚举名 | 值 | 字符串 | 说明 | +|--------|---|--------|------| +| kParallel | 0 | `parallel` | 并行迭代 | +| kBroadcast | 1 | `broadcast` | 广播迭代 | +| kTranspose | 2 | `transpose` | 转置迭代 | +| kReduction | 3 | `reduction` | 归约迭代 | +| kInterleave | 4 | `interleave` | 交织迭代 | +| kDeinterleave | 5 | `deinterleave` | 解交织迭代 | +| kInverse | 6 | `inverse` | 逆序迭代 | +| kPad | 7 | `pad` | 填充迭代 | +| kConcat | 8 | `concat` | 拼接迭代 | +| kGather | 9 | `gather` | 收集迭代 | +| kCumulative | 10 | `cumulative` | 累积迭代 | +| kOpaque | 99 | `opaque` | 不透明迭代 | + +--- + +## DataLayout（7 值） + +源码：[HIVMAttrs.td#L84-L101](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L84-L101) + +| 枚举名 | 值 | 字符串 | 说明 | +|--------|---|--------|------| +| DOTA_ND | 1 | `dotA_ND` | 矩阵 A 的 ND 布局（Cube 矩阵乘） | +| DOTB_ND | 2 | `dotB_ND` | 矩阵 B 的 ND 布局 | +| DOTC_ND | 3 | `dotC_ND` | 矩阵 C 的 ND 布局 | +| nZ | 4 | `nZ` | nZ 格式（列优先分块） | +| zN | 5 | `zN` | zN 格式（行优先分块） | +| ND | 6 | `ND` | ND 格式（标准行优先） | +| Fractal | 7 | `Fractal` | Fractal 格式 | + +--- + +## AddressSpace（7 值） + +源码：[HIVMAttrs.td#L171-L188](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L171-L188) + +| 枚举名 | 值 | 字符串 | 说明 | +|--------|---|--------|------| +| Zero | 0 | `zero` | 默认地址空间 | +| GM | 1 | `gm` | 全局内存 | +| L1 | 2 | `cbuf` | L1 缓存（CBuffer） | +| L0A | 3 | `ca` | L0A 缓存（Cube A 矩阵） | +| L0B | 4 | `cb` | L0B 缓存（Cube B 矩阵） | +| L0C | 5 | `cc` | L0C 缓存（Cube C 矩阵/累加器） | +| UB | 6 | `ub` | 统一缓冲区 | + +--- + +## Pipe（15 值） + +源码：[HIVMAttrs.td#L203-L236](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L203-L236) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| PIPE_S | 0 | Scalar Pipe | +| PIPE_V | 1 | Vector Pipe | +| PIPE_M | 2 | Cube Matrix Pipe | +| PIPE_MTE1 | 3 | Memory Transfer Engine 1（L1 搬运） | +| PIPE_MTE2 | 4 | Memory Transfer Engine 2（GM→片内搬运） | +| PIPE_MTE3 | 5 | Memory Transfer Engine 3（片内→GM 搬运） | +| PIPE_ALL | 6 | 所有 Pipe | +| PIPE_MTE4 | 7 | Memory Transfer Engine 4 | +| PIPE_MTE5 | 8 | Memory Transfer Engine 5 | +| PIPE_V2 | 9 | Vector Pipe 2 | +| PIPE_FIX | 10 | Fixpipe（L0C→UB/GM 数据转换） | +| VIRTUAL_PIPE_MTE2_L1A | 11 | 虚拟 Pipe：MTE2 到 L1A | +| VIRTUAL_PIPE_MTE2_L1B | 12 | 虚拟 Pipe：MTE2 到 L1B | +| PIPE_NUM | 13 | Pipe 数量标记 | +| PIPE_UNASSIGNED | 99 | 未分配 Pipe | + +--- + +## TCoreType（4 值） + +源码：[HIVMAttrs.td#L298-L309](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L298-L309) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| CUBE | 1 | Cube Core（矩阵计算） | +| VECTOR | 2 | Vector Core（向量计算） | +| CUBE_OR_VECTOR | 3 | 可在 Cube 或 Vector 上执行 | +| CUBE_AND_VECTOR | 4 | 需要在 Cube 和 Vector 上同时执行 | + +--- + +## TFuncCoreType（4 值） + +源码：[HIVMAttrs.td#L250-L261](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L250-L261) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| AIC | 1 | AI Cube Kernel | +| AIV | 2 | AI Vector Kernel | +| MIX | 3 | 混合 Cube+Vector Kernel | +| AIC_OR_AIV | 4 | AIC 或 AIV Kernel | + +--- + +## TModuleCoreType（3 值） + +源码：[HIVMAttrs.td#L271-L276](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L271-L276) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| AIC | 1 | 所有函数为 AIC | +| AIV | 2 | 所有函数为 AIV | +| MIX | 3 | 包含 AIC 和 AIV 函数 | + +--- + +## PadMode（3 值） + +源码：[HIVMAttrs.td#L330-L341](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L330-L341) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| PadNull | 0 | 无填充 | +| PadFirstElem | 1 | 使用第一个元素填充 | +| PadValue | 2 | 使用指定值填充 | + +--- + +## EvictionPolicy（2 值） + +源码：[HIVMAttrs.td#L356-L364](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L356-L364) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| EvictFirst | 0 | 先驱逐 | +| EvictLast | 1 | 后驱逐 | + +--- + +## RoundMode（7 值） + +源码：[HIVMAttrs.td#L378-L406](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L378-L406) + +| 枚举名 | 值 | 字符串 | 说明 | +|--------|---|--------|------| +| RINT | 0 | `rint` | 就近舍入， ties to even | +| ROUND | 1 | `round` | 就近舍入， ties away from zero | +| FLOOR | 2 | `floor` | 向负无穷舍入 | +| CEIL | 3 | `ceil` | 向正无穷舍入 | +| TRUNC | 4 | `trunc` | 向零舍入 | +| ODD | 5 | `odd` | 向奇数舍入（Von Neumann） | +| TRUNCWITHOVERFLOW | 6 | `truncwithoverflow` | 截断并溢出 | + +--- + +## UnsignedMode（4 值） + +源码：[HIVMAttrs.td#L408-L425](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L408-L425) + +| 枚举名 | 值 | 字符串 | 说明 | +|--------|---|--------|------| +| SI2SI | 0 | `si2si` | 有符号→有符号 | +| SI2UI | 1 | `si2ui` | 有符号→无符号 | +| UI2SI | 2 | `ui2si` | 无符号→有符号 | +| UI2UI | 3 | `ui2ui` | 无符号→无符号 | + +--- + +## TypeFn / Cast（3 值） + +源码：[HIVMAttrs.td#L431-L446](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L431-L446) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| cast_signed | 0 | 有符号类型转换 | +| cast_unsigned | 1 | 无符号类型转换 | +| bitcast | 2 | 位转换（不改变位模式） | + +--- + +## CompareMode（6 值） + +源码：[HIVMAttrs.td#L452-L473](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L452-L473) + +| 枚举名 | 值 | 字符串 | 说明 | +|--------|---|--------|------| +| EQ | 0 | `eq` | 等于 | +| NE | 1 | `ne` | 不等于 | +| LT | 2 | `lt` | 小于 | +| GT | 3 | `gt` | 大于 | +| GE | 4 | `ge` | 大于等于 | +| LE | 5 | `le` | 小于等于 | + +--- + +## Event（8 值） + +源码：[HIVMAttrs.td#L479-L498](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L479-L498) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| EVENT_ID0 | 0 | Event ID 0 | +| EVENT_ID1 | 1 | Event ID 1 | +| EVENT_ID2 | 2 | Event ID 2 | +| EVENT_ID3 | 3 | Event ID 3 | +| EVENT_ID4 | 4 | Event ID 4 | +| EVENT_ID5 | 5 | Event ID 5 | +| EVENT_ID6 | 6 | Event ID 6 | +| EVENT_ID7 | 7 | Event ID 7 | + +--- + +## UnitFlag（4 值） + +源码：[HIVMAttrs.td#L512-L523](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L512-L523) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| DISABLED | 0 | 禁用 UnitFlag | +| RESERVED | 1 | 保留 | +| ENABLED_WITHOUT_UPDATE | 2 | 启用但不更新标志 | +| ENABLED_WITH_UPDATE | 3 | 启用并更新标志 | + +--- + +## SyncBlockMode（6 值） + +源码：[HIVMAttrs.td#L540-L555](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L540-L555) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| ALL_CUBE | 0 | 所有 Cube Core 同步 | +| ALL_VECTOR | 1 | 所有 Vector Core 同步 | +| ALL_SUB_VECTOR | 2 | 所有 Sub-Vector 同步 | +| BARRIER_CUBE | 3 | Cube-Cube 屏障 | +| BARRIER_VECTOR | 4 | Vector-Vector 屏障 | +| ALL | 5 | 所有 AIC/AIV 同步 | + +--- + +## SyncBlockInstrMode（3 值） + +源码：[HIVMAttrs.td#L569-L578](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L569-L578) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| INTER_BLOCK_SYNCHRONIZATION | 0 | 跨 Block 同步 | +| INTER_SUBBLOCK_SYNCHRONIZATION | 1 | 跨 Sub-Block 同步 | +| INTRA_BLOCK_SYNCHRONIZATION | 2 | Block 内同步（默认） | + +--- + +## ReduceOperation（12 值） + +源码：[HIVMAttrs.td#L596-L623](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L596-L623) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| none | 0 | 无归约 | +| sum | 1 | 求和 | +| prod | 2 | 求积 | +| max | 3 | 最大值 | +| min | 4 | 最小值 | +| max_with_index | 5 | 最大值及索引 | +| min_with_index | 6 | 最小值及索引 | +| any | 7 | 任意为真 | +| all | 8 | 全部为真 | +| xori | 9 | 异或（整数） | +| ori | 10 | 或（整数） | +| andi | 11 | 与（整数） | + +--- + +## AtomicKind（9 值） + +源码：[HIVMAttrs.td#L637-L658](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L637-L658) + +| 枚举名 | 值 | 字符串 | 说明 | +|--------|---|--------|------| +| NONE | 0 | `none` | 无原子操作 | +| ADD | 1 | `add` | 原子加 | +| MAX | 2 | `max` | 原子最大值 | +| MIN | 3 | `min` | 原子最小值 | +| AND | 4 | `and` | 原子与 | +| OR | 5 | `or` | 原子或 | +| XOR | 6 | `xor` | 原子异或 | +| CAS | 7 | `or` | 比较并交换 | +| XCHG | 8 | `xor` | 原子交换 | + +--- + +## AlignKind（3 值） + +源码：[HIVMAttrs.td#L670-L678](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L670-L678) + +| 枚举名 | 值 | 字符串 | 说明 | +|--------|---|--------|------| +| ALIGN | 0 | `align` | 对齐 | +| UNALIGNED | 1 | `unaligned` | 未对齐 | +| UNKNOWN | 2 | `unknown` | 未知 | + +--- + +## AxisKind（3 值） + +源码：[HIVMAttrs.td#L688-L696](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L688-L696) + +| 枚举名 | 值 | 字符串 | 说明 | +|--------|---|--------|------| +| FIRST | 0 | `first` | 第一轴 | +| MIDDLE | 1 | `middle` | 中间轴 | +| LAST | 2 | `last` | 最后轴 | + +--- + +## VFMode（3 值） + +源码：[HIVMAttrs.td#L948-L954](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L948-L954) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| SIMD | 0 | SIMD 模式（向量化） | +| SIMT | 1 | SIMT 模式（多线程） | +| MIX | 2 | 混合 SIMD+SIMT 模式 | + +--- + +## DataCacheKind（4 值） + +源码：[HIVMAttrs.td#L736-L747](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L736-L747) + +| 枚举名 | 值 | 字符串 | 说明 | +|--------|---|--------|------| +| ALL | 0 | `all` | 所有缓存 | +| UB | 1 | `ub` | UB 缓存 | +| OUT | 2 | `out` | 输出缓存 | +| ATOMIC | 3 | `atomic` | 原子缓存 | + +--- + +## DCCIMode（2 值） + +源码：[HIVMAttrs.td#L753-L763](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L753-L763) + +| 枚举名 | 值 | 字符串 | 说明 | +|--------|---|--------|------| +| SINGLE_CACHE_LINE | 0 | `single_cache_line` | 单缓存行 | +| ALL_CACHE_LINES | 1 | `all_cache_lines` | 所有缓存行 | + +--- + +## FixpipePreQuantMode（5 值） + +源码：[HIVMAttrs.td#L783-L796](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L783-L796) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| NO_QUANT | 0 | 不量化 | +| F322F16 | 1 | FP32→FP16 | +| S322I8 | 9 | INT32→INT8 | +| QF322F32_PRE | 15 | 量化 FP32→FP32 预处理 | +| F322BF16 | 16 | FP32→BF16 | + +--- + +## FixpipePreReluMode（4 值） + +源码：[HIVMAttrs.td#L803-L814](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L803-L814) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| NO_RELU | 0 | 不激活 | +| NORMAL_RELU | 1 | 标准 ReLU | +| LEAKY_RELU | 2 | Leaky ReLU | +| P_RELU | 3 | P-ReLU | + +--- + +## FixpipeDualDstMode（3 值） + +源码：[HIVMAttrs.td#L821-L830](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L821-L830) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| NO_DUAL | 0 | 单目的地模式 | +| ROW_SPLIT | 1 | M 维度拆分，M/2 x N 写入每个 UB | +| COLUMN_SPLIT | 2 | N 维度拆分，M x N/2 写入每个 UB | + +--- + +## FixpipeDMAMode（3 值） + +源码：[HIVMAttrs.td#L847-L854](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L847-L854) + +| 枚举名 | 值 | 字符串 | 说明 | +|--------|---|--------|------| +| NZ2ND | 0 | `nz2nd` | nZ 格式→ND 格式 | +| NZ2DN | 1 | `nz2dn` | nZ 格式→DN 格式 | +| NZ2NZ | 2 | `normal` | 正常模式（nZ→nZ） | + +--- + +## DeinterleaveMode（3 值） + +源码：[HIVMAttrs.td#L865-L874](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L865-L874) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| CHANNEL_0 | 0 | 通道 0 | +| CHANNEL_1 | 1 | 通道 1 | +| ALL_CHANNELS | 999 | 所有通道 | + +--- + +## DescaleMode（3 值） + +源码：[HIVMAttrs.td#L886-L895](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L886-L895) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| DescaleNull | 0 | 不使用反量化 | +| DescalePerChannel | 1 | 按 Channel 反量化 | +| DescalePerTensor | 2 | 按 Tensor 反量化 | + +--- + +## MatmulBiasMode（5 值） + +源码：[HIVMAttrs.td#L903-L914](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L903-L914) + +| 枚举名 | 值 | 说明 | +|--------|---|------| +| NoBias | 0 | 无 bias | +| PerChannelAdd | 1 | Per-channel 加法 bias | +| PostPerChannelAddWithSplitK | 2 | Split-K 后 Per-channel 加法 bias | +| ElementwiseAdd | 3 | 逐元素加法 bias | +| MMInitPerChannelAddWithSplitK | 4 | Split-K 初始化 Per-channel 加法 bias | + +--- + +## MemoryEffect（3 值） + +源码：[HIVMAttrs.td#L1046-L1055](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L1046-L1055) + +| 枚举名 | 值 | 字符串 | 说明 | +|--------|---|--------|------| +| READ | 0 | `read` | 读内存 | +| WRITE | 1 | `write` | 写内存 | +| READ_WRITE | 2 | `read_write` | 读写内存 | + +--- + +## 相关文档 + +- 源码参考：[HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) +- 测试用例：[attribute.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/attribute.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/02-parameterized-attrs.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/02-parameterized-attrs.md new file mode 100644 index 00000000..9f8d212a --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/02-parameterized-attrs.md @@ -0,0 +1,270 @@ +# HIVM 参数化属性 + +> 关键词：DataLayoutAttr, AddressSpaceAttr, PipeAttr, BlockMappingAttr, SubBlockMappingAttr, TightlyCoupledBufferAttr + +## 概述 + +参数化属性是携带额外参数的复合属性，用于描述 HIVM 中需要多个参数才能完整表达的信息。与简单枚举属性不同，参数化属性可以包含可选参数、数组参数等。 + +## DataLayoutAttr + +源码：[HIVMAttrs.td#L103-L165](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L103-L165) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$data_layout` | EnumParameter\ | 是 | 数据布局枚举值 | +| `$transpose` | OptionalParameter\ | 否 | 是否转置（DOTA_ND/DOTB_ND 时必须提供） | +| `$fractalSizes` | OptionalParameter\ | 否 | Fractal 块大小 | + +### IR 格式 + +```mlir +#hivm.data_layout +#hivm.data_layout +#hivm.data_layout +#hivm.data_layout +#hivm.data_layout +#hivm.data_layout +#hivm.data_layout +``` + +### 约束 + +- `transpose` 仅对 DOTA_ND 和 DOTB_ND 布局有效且必须提供 +- `fractalSizes` 为可选的 I64 数组，指定 Fractal 块大小 + +### 额外方法 + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getFractalSizesArray()` | `optional>` | 获取 Fractal 大小数组 | +| `getTransposeValue()` | `optional` | 获取 transpose 值 | +| `isNDLayout()` | bool | 判断是否为 ND 类布局（DOTA_ND/DOTB_ND/DOTC_ND/ND） | +| `getFractalBlockSizes()` | `FailureOr` | 获取 2 个 Fractal 块大小 | + +## AddressSpaceAttr + +源码：[HIVMAttrs.td#L190-L197](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L190-L197) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$address_space` | EnumParameter\ | 是 | 地址空间枚举值 | + +### IR 格式 + +```mlir +#hivm.address_space +#hivm.address_space +#hivm.address_space +#hivm.address_space +#hivm.address_space +#hivm.address_space +#hivm.address_space +``` + +### 接口 + +实现了 `DeviceMappingAttrInterface`，用于设备映射。 + +## PipeAttr + +源码：[HIVMAttrs.td#L238-L244](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L238-L244) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$pipe` | EnumParameter\ | 是 | Pipeline 枚举值 | + +### IR 格式 + +```mlir +#hivm.pipe +#hivm.pipe +#hivm.pipe +#hivm.pipe +#hivm.pipe +#hivm.pipe +#hivm.pipe +#hivm.pipe +``` + +## TCoreTypeAttr + +源码：[HIVMAttrs.td#L311-L317](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L311-L317) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$tcoretype` | EnumParameter\ | 是 | Core 类型枚举值 | + +### IR 格式 + +```mlir +#hivm.tcore_type +#hivm.tcore_type +#hivm.tcore_type +#hivm.tcore_type +``` + +## TFuncCoreTypeAttr + +源码：[HIVMAttrs.td#L263-L269](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L263-L269) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$funcCoreType` | EnumParameter\ | 是 | 函数 Core 类型 | + +### IR 格式 + +```mlir +#hivm.func_core_type +#hivm.func_core_type +#hivm.func_core_type +``` + +## TModuleCoreTypeAttr + +源码：[HIVMAttrs.td#L278-L292](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L278-L292) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$moduleCoreType` | EnumParameter\ | 是 | 模块 Core 类型 | + +### 推断规则 + +- 所有函数为 AIV → 模块 Core 类型为 AIV +- 所有函数为 AIC → 模块 Core 类型为 AIC +- 否则 → 模块 Core 类型为 MIX + +## EventAttr + +源码：[HIVMAttrs.td#L500-L506](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L500-L506) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$event` | EnumParameter\ | 是 | Event ID | + +### IR 格式 + +```mlir +#hivm.event +#hivm.event +``` + +## UnitFlagAttr + +源码：[HIVMAttrs.td#L525-L531](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L525-L531) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$unit_flag` | EnumParameter\ | 是 | UnitFlag 模式 | + +### 数组类型 + +`UnitFlagArrayAttr` 是 `UnitFlagAttr` 的数组类型，用于宏操作中多个输出 Tensor 的 UnitFlag 模式。 + +## SyncBlockModeAttr + +源码：[HIVMAttrs.td#L557-L563](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L557-L563) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$sync_mode` | EnumParameter\ | 是 | 同步 Block 模式 | + +## SyncBlockInstrModeAttr + +源码：[HIVMAttrs.td#L580-L590](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L580-L590) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$sync_instr_mode` | EnumParameter\ | 是 | 同步指令模式 | + +### 默认值 + +默认值为 `INTRA_BLOCK_SYNCHRONIZATION`。 + +## BlockMappingAttr + +源码：[HIVMAttrs.td#L712-L720](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L712-L720) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$order` | OptionalParameter\\> | 否 | 线性维度序号 | + +### IR 格式 + +```mlir +#hivm.block +#hivm.block +``` + +### 接口 + +实现了 `DeviceMappingAttrInterface`。 + +## SubBlockMappingAttr + +源码：[HIVMAttrs.td#L722-L730](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L722-L730) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$sub_block` | EnumParameter\ | 是 | Sub-Block 映射 ID | + +### 说明 + +用于 Mix Kernel 中 Cube/Vector Block 维度比例的映射。 + +## TightlyCoupledBufferAttr + +源码：[HIVMAttrs.td#L1010-L1017](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L1010-L1017) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$id` | OptionalParameter\\> | 否 | 紧耦合缓冲区 ID | + +### 说明 + +用于 Cube-Vector 紧耦合缓冲区，允许 Cube 和 Vector Core 共享数据。 + +## MemoryEffectAttr + +源码：[HIVMAttrs.td#L1057-L1063](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td#L1057-L1063) + +### 参数 + +| 参数 | 类型 | 必选 | 说明 | +|------|------|------|------| +| `$effect` | EnumParameter\ | 是 | 内存效果类型 | + +### 说明 + +用于 SIMT VF 模式下的内存效果标记。 + +## 相关文档 + +- 源码参考：[HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) +- 测试用例：[attribute.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/attribute.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/03-type-system.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/03-type-system.md new file mode 100644 index 00000000..c7eabfdb --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/03-type-system.md @@ -0,0 +1,132 @@ +# HIVM 类型系统约束 + +> 关键词：MemRef, Tensor, Vector, Element Type, AddressSpace, DataLayout, ShapedType + +## 概述 + +HIVM 在标准 MLIR 类型系统（MemRef、Tensor、Vector）上增加了特定约束，以表达 NPU 硬件的存储层次、数据布局和元素类型限制。这些约束通过 AddressSpace 属性、DataLayout 属性和元素类型限制来体现。 + +## MemRef 在 HIVM 中的约束 + +### AddressSpace 标注 + +HIVM 中的 MemRef 通常需要通过 `#hivm.address_space<...>` 属性标注其所在的存储层次： + +```mlir +memref<16x16xf16, #hivm.address_space> // 全局内存 +memref<16x16xf16, #hivm.address_space> // L1 缓存 +memref<16x16xf16, #hivm.address_space> // L0A 缓存 +memref<16x16xf16, #hivm.address_space> // L0B 缓存 +memref<16x16xf16, #hivm.address_space> // L0C 缓存 +memref<16x16xf16, #hivm.address_space> // 统一缓冲区 +``` + +### 存储层次与操作约束 + +| 存储层次 | AddressSpace | 可用操作 | 说明 | +|---------|-------------|---------|------| +| GM | `gm` | load, store, matmul, mix_matmul | 全局内存，所有 Block 共享 | +| L1 | `cbuf` | nd2nz, copy, mmadL1 输入 | CBuffer，Block 内共享 | +| L0A | `ca` | mmadL1 内部使用 | Cube A 矩阵缓冲区 | +| L0B | `cb` | mmadL1 内部使用 | Cube B 矩阵缓冲区 | +| L0C | `cc` | fixpipe 输入 | Cube C 矩阵/累加器 | +| UB | `ub` | vector ops, fixpipe 输出 | 统一缓冲区，Vector Core 使用 | + +### Strided MemRef + +HIVM 支持 strided MemRef，用于表达子视图： + +```mlir +memref<256x128xf16, strided<[2048, 1], offset: 0>> +memref<256x128xf16, strided<[2048, 1]>, #hivm.address_space> +``` + +## Tensor 在 HIVM 中的约束 + +### Tensor 语义 + +HIVM 中的 Tensor 遵循 MLIR 标准 Tensor 语义，但增加了以下约束： + +1. **DestinationStyleOpInterface**：HIVM 操作通常实现 DPS 接口，Tensor 通过 `outs` 操作数传入 +2. **RankedTensor**：HIVM 操作要求使用 RankedTensor（有确定形状的 Tensor） +3. **Tensor/MemRef 双语义**：大多数 HIVM 操作同时支持 Tensor 和 MemRef 语义 + +### Tensor 与 MemRef 的选择 + +| 场景 | 推荐类型 | 原因 | +|------|---------|------| +| 函数间传递数据 | MemRef + AddressSpace | 明确存储层次 | +| 操作间传递中间结果 | Tensor | 更安全的类型系统 | +| 需要子视图 | MemRef + subview | Tensor 使用 extract_slice | +| Bufferization 后 | MemRef | Bufferization 将 Tensor 转换为 MemRef | + +## Vector 在 HIVM 中的约束 + +HIVM 中的 Vector 类型主要用于 Vector Core 上的计算操作。约束包括： + +1. **元素类型限制**：Vector 的元素类型必须符合硬件支持的类型 +2. **形状约束**：Vector 的形状必须符合 Vector Core 的计算单元大小 + +## 元素类型约束 + +### 支持的元素类型 + +| 类型 | 说明 | 适用场景 | +|------|------|---------| +| f16 | 半精度浮点 | 矩阵乘法、通用计算 | +| bf16 | BFloat16 | 矩阵乘法、推理 | +| f32 | 单精度浮点 | 累加器、通用计算 | +| f64 | 双精度浮点 | 有限支持 | +| i8 | 8-bit 整数 | 量化计算 | +| i16 | 16-bit 整数 | 通用计算 | +| i32 | 32-bit 整数 | 通用计算、索引 | +| i64 | 64-bit 整数 | 地址、索引 | +| i1 | 布尔 | 条件、mask | + +### 矩阵乘法元素类型组合 + +| A 类型 | B 类型 | C 类型 | 说明 | +|--------|--------|--------|------| +| f16 | f16 | f32 | 标准 FP16 矩阵乘，累加为 FP32 | +| f16 | f16 | f16 | FP16 矩阵乘，累加为 FP16 | +| bf16 | bf16 | f32 | BF16 矩阵乘 | +| i8 | i8 | i32 | INT8 量化矩阵乘 | + +### 类型转换约束 + +HIVM 提供以下类型转换操作： + +- **bitcast**：位模式不变，重新解释类型（要求总位宽相同） +- **cast_signed / cast_unsigned**：有符号/无符号类型转换 +- **FixpipePreQuantMode**：Fixpipe 中的在线类型转换（F322F16、S322I8 等） + +## 类型约束 Trait + +HIVM 定义了以下类型约束相关的 Trait： + +| Trait | 说明 | +|-------|------| +| `HIVMOpSameOperandsAndResultRank` | 所有操作数和结果的 rank 相同（除临时缓冲区） | +| `ElementwiseNaryOpTrait` | N 元素逐元素操作，rank 相同 | +| `StaticMaxRankTrait` | 静态已知最大 rank | +| `NoMaxRankTrait` | 无 rank 限制 | +| `VectorOnlyTrait` | 指定操作数只支持 Vector 类型 | +| `ScalarOnlyTrait` | 指定操作数只支持标量类型 | +| `OperElemTypeConstraints` | 指定操作数的元素类型约束 | + +## 常见问题 + +**Q: 什么时候需要标注 AddressSpace？** +A: 当 MemRef 需要明确存储层次时（如 load/store 操作的 GM 参数），必须标注 AddressSpace。编译器在 Bufferization 后会自动推断和添加 AddressSpace。 + +**Q: Tensor 和 MemRef 可以混用吗？** +A: 大多数 HIVM 操作同时支持 Tensor 和 MemRef 语义。但在同一操作中，输入和输出通常使用相同的语义（全 Tensor 或全 MemRef）。 + +**Q: 为什么 L0C 的 AddressSpace 是 `cc`？** +A: L0C 是 Cube Core 的累加器缓冲区，`cc` 代表 Cube C 矩阵。类似地，`ca` 是 Cube A，`cb` 是 Cube B。 + +## 相关文档 + +- 源码参考：[HIVMAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMAttrs.td) +- 源码参考：[HIVMTraits.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td) +- 测试用例：[ops.mlir](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/test/Dialect/HIVM/IR/ops.mlir) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/04-interfaces.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/04-interfaces.md new file mode 100644 index 00000000..0c2d2780 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/04-interfaces.md @@ -0,0 +1,172 @@ +# HIVM 接口方法列表 + +> 关键词：HIVMStructuredOpInterface, HIVMCoreTypeInterface, HIVMInferCoreTypeInterface, HIVMUnitFlagEnabledInterface, OpLayoutInterface + +## 概述 + +HIVM 定义了多个 OpInterface 和 AttrInterface，用于描述操作的通用行为。这些接口是 HIVM 编译器 Pass（如同步注入、Lowering、Bufferization 等）与具体操作之间的契约。 + +## HIVMStructuredOpInterface + +源码：[HIVMInterfaces.td#L67-L678](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMInterfaces.td#L67-L678) + +HIVM 结构化操作接口，继承自以下接口： +- `DestinationStyleOpInterface`：DPS 目标风格操作 +- `OpPipeInterface`：Pipe 信息 +- `HIVMCoreTypeInterface`：Core 类型查询 +- `FlattenInterface`：维度展平 +- `LibraryFunctionOpInterface`：库函数调用 + +### 方法列表 + +#### 操作属性查询 + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `isElemwiseNaryOp()` | bool | 是否为逐元素 N-ary 操作 | +| `isInlineBroadcastable()` | bool | 是否支持内联广播 | +| `isInlineTransposable()` | bool | 是否支持内联转置 | +| `existInlineBroadcastLoopDims()` | bool | 是否存在内联广播维度 | +| `existInlineTransposeLoopDims()` | bool | 是否存在内联转置维度 | + +#### 循环类型处理 + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getIteratorTypesArray()` | `SmallVector` | 获取迭代器类型数组 | +| `setIteratorTypesArray(IteratorType, DenseI64ArrayAttr&)` | LogicalResult | 设置迭代器类型数组 | +| `getNumLoops()` | unsigned | 获取循环总数 | +| `getNumParallelLoops()` | unsigned | 获取并行循环数 | +| `getParallelLoopDims(SmallVectorImpl&)` | void | 获取并行循环维度 | +| `getReductionLoopDims(SmallVectorImpl&)` | void | 获取归约循环维度 | +| `getBroadcastLoopDims(SmallVectorImpl&)` | void | 获取广播循环维度 | +| `getTransposeLoopDims(SmallVectorImpl&)` | void | 获取转置循环维度 | +| `getPadLoopDims(SmallVectorImpl&)` | void | 获取填充循环维度 | +| `getConcatLoopDims(SmallVectorImpl&)` | void | 获取拼接循环维度 | +| `getGatherLoopDims(SmallVectorImpl&)` | void | 获取收集循环维度 | +| `getPermutationArray()` | `ArrayRef` | 获取转置排列数组 | +| `getBroadcastArray()` | `ArrayRef` | 获取广播数组 | +| `getInlinedBroadcastableAxes(OpOperand*)` | `SmallVector` | 获取内联广播轴 | + +#### Indexing Maps + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getIndexingMaps()` | ArrayAttr | 获取 indexing maps 属性 | +| `getIndexingMapsArray()` | `SmallVector` | 获取 indexing maps 数组 | +| `getLoopsToShapesMap()` | AffineMap | 获取循环到形状的映射 | +| `getShapesToLoopsMap()` | AffineMap | 获取形状到循环的映射 | +| `getMatchingIndexingMap(OpOperand*)` | AffineMap | 获取操作数对应的 indexing map | +| `getIndexingMapMatchingResult(OpResult)` | AffineMap | 获取结果对应的 indexing map | + +#### 形状与 Rank + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getRank(OpOperand*)` | int64_t | 获取操作数 rank | +| `getShape(OpOperand*)` | `ArrayRef` | 获取操作数形状 | +| `getStaticShape()` | `SmallVector` | 获取静态形状 | +| `hasDynamicShape()` | bool | 是否有动态形状 | +| `hasIndexSemantics()` | bool | 是否有索引语义 | + +#### 内存效果与操作数 + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getEffects(SmallVectorImpl<...>&)` | void | 获取内存效果 | +| `getTargetSpaceOperands(AddressSpace, bool)` | `SmallVector` | 获取目标地址空间操作数 | +| `getHIVMOperands(bool)` | `SmallVector` | 获取 HIVM 操作数 | +| `getHIVMOperandTypes(bool)` | `SmallVector` | 获取 HIVM 操作数类型 | +| `getHIVMInputOperands(bool)` | `SmallVector` | 获取 HIVM 输入操作数 | +| `isVectorOnlyOperand(size_t)` | bool | 检查操作数是否仅支持 Vector | +| `getContiguousAxes()` | BitVector | 获取连续轴掩码 | +| `getUnitAxesMask()` | BitVector | 获取单位轴掩码 | +| `getPermutedAxesMask()` | BitVector | 获取置换轴掩码 | + +#### 额外声明 + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `createFlatListOfOperandDims(OpBuilder&, Location)` | `SmallVector` | 创建操作数维度平面列表 | +| `createFlatListOfOperandStaticDims()` | `SmallVector` | 创建静态维度平面列表 | +| `createLoopRanges(OpBuilder&, Location)` | `SmallVector` | 创建循环范围 | +| `computeStaticLoopSizes()` | `SmallVector` | 计算静态循环大小 | +| `reifyResultShapes(OpBuilder&, ReifiedRankedShapedTypeDims&)` | LogicalResult | 具体化结果形状 | +| `getIndexingMapIndex(OpOperand*)` | int64_t | 获取 indexing map 索引 | + +## HIVMCoreTypeInterface + +源码：[HIVMInterfaces.td#L27-L46](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMInterfaces.td#L27-L46) + +Core 类型查询接口，用于确定操作在哪种 Core 上执行。 + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getCoreType()` | `optional` | 获取操作的 Core 类型 | + +Core 类型有两种确定方式： +1. **静态**：通过 `CoreTypeTrait` 附加到操作上 +2. **动态**：通过 `InferCoreTypeInterface` 推断 + +## HIVMInferCoreTypeInterface + +源码：[HIVMInterfaces.td#L48-L65](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMInterfaces.td#L48-L65) + +Core 类型推断接口，用于动态推断操作的 Core 类型。 + +| 方法 | 返回类型 | 默认实现 | 说明 | +|------|---------|---------|------| +| `inferCoreType()` | `optional` | 返回 `std::nullopt` | 推断操作的 Core 类型 | + +## HIVMUnitFlagEnabledInterface + +源码：[HIVMInterfaces.td#L680-L761](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMInterfaces.td#L680-L761) + +UnitFlag 启用接口，用于支持 UnitFlag 同步模式的操作。 + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getUnitFlagModes()` | `optional>` | 获取 UnitFlag 模式数组 | +| `setUnitFlagModes(SmallVector)` | void | 设置 UnitFlag 模式 | +| `getUnitFlagConditions()` | `optional>` | 获取 UnitFlag 条件值 | +| `setUnitFlagConditions(SmallVector)` | void | 设置 UnitFlag 条件值 | +| `getUnitFlagModeLibValue(PatternRewriter&)` | Value | 获取传递给库调用的 UnitFlag 值 | + +### 验证 + +该接口包含验证方法 `verifyUnitFlagEnabledInterface`，确保操作正确声明了 `unit_flag_mode` 和 `unit_flag_cond` 参数。 + +## OpLayoutInterface + +源码：[OpLayoutInterface.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/Interfaces/OpLayoutInterface.td) + +布局接口，用于确定操作数的目标 Fractal Layout。 + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getOperandsTargetFractalLayout()` | 需要操作自行实现 | 获取操作数的目标 Fractal Layout | + +mmadL1 实现了此接口，提供以下方法： + +| 方法 | 返回类型 | 说明 | +|------|---------|------| +| `getOperandALayout()` | `FailureOr` | 获取 A 矩阵布局 | +| `getOperandBLayout()` | `FailureOr` | 获取 B 矩阵布局 | +| `getOperandCLayout()` | `FailureOr` | 获取 C 矩阵布局 | +| `getOperandBiasLayout()` | `FailureOr` | 获取 Bias 布局 | + +## 常见问题 + +**Q: HIVMStructuredOpInterface 和 LinalgOp 的关系？** +A: HIVMStructuredOpInterface 借鉴了 LinalgOp 的设计，但增加了 Pipe、CoreType、Flatten 等 NPU 特有的接口。两者都基于 DestinationStyleOpInterface。 + +**Q: 什么时候需要实现 HIVMInferCoreTypeInterface？** +A: 当操作的 Core 类型无法静态确定时（如同步操作可能在 Cube 或 Vector 上执行），需要实现此接口动态推断。 + +**Q: OpLayoutInterface 的用途？** +A: 用于确定操作数在 NPU 存储层次中的数据布局（如 Fractal 格式），编译器据此生成正确的数据搬运指令。 + +## 相关文档 + +- 源码参考：[HIVMInterfaces.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMInterfaces.td) +- 源码参考：[OpLayoutInterface.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/Interfaces/OpLayoutInterface.td) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/05-traits.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/05-traits.md new file mode 100644 index 00000000..399351d4 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/06-Attributes-Types/05-traits.md @@ -0,0 +1,222 @@ +# HIVM Trait 定义 + +> 关键词：Trait, MacroOpTrait, MacroOpPipeTrait, CoreTypeTrait, ElementwiseNaryOpTrait, BroadcastableOTF, TransposableOTF + +## 概述 + +HIVM Trait 是附加在操作上的行为约束标记，用于编译器 Pass 中的操作分类和验证。Trait 从 [HIVMTraits.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td) 定义。 + +## Rank 相关 Trait + +### HIVMOpSameOperandsAndResultRank + +验证所有操作数和结果类型（除临时缓冲区外）的 rank 相同。 + +- **依赖**：HIVMStructuredOpInterface +- **适用操作**：大多数 Elementwise 操作 + +### OpLibraryMaxRankTrait\ + +指定库函数支持的最大 rank。 + +| 参数 | 说明 | +|------|------| +| MaxRank > 0 | 静态已知最大 rank | +| MaxRank = 0 | 静态未知最大 rank（需推断） | +| MaxRank = -1 | 无 rank 限制 | + +### StaticMaxRankTrait\ + +继承自 `OpLibraryMaxRankTrait`，用于静态已知最大 rank 的操作。 + +| 操作 | MaxRank | 说明 | +|------|---------|------| +| debug | 4 | 调试操作最多 4D | +| embedding_gather | 3 | 嵌入查找最多 3D | +| indirect_load / indirect_store | 5 | 间接加载/存储最多 5D | +| gatherT / scatterT / index_put | 5 | 收集/散射最多 5D | + +### InferMaxRankTrait + +继承自 `OpLibraryMaxRankTrait<0>`，用于静态未知最大 rank 的操作。 + +### NoMaxRankTrait + +继承自 `OpLibraryMaxRankTrait<-1>`，无 rank 限制。 + +| 操作 | 说明 | +|------|------| +| mmadL1 | 矩阵乘加，rank 由操作数决定 | +| matmul / mix_matmul / mix_group_matmul | 全局矩阵乘，无 rank 限制 | + +## Elementwise Trait + +### ElementwiseNaryOpTrait\ + +N 元素逐元素操作 Trait。 + +**语义**：对 N 个输入操作数执行逐元素操作，产生单个结果。 + +**约束**： +1. 实现 DestinationStyleOpInterface +2. 输入操作数数量为 N +3. 所有 shaped 操作数和结果的 rank 相同 +4. Buffer 语义时，最终维度的 stride 相等 + +**依赖**：HIVMStructuredOpInterface, HIVMOpSameOperandsAndResultRank + +## 广播与转置 Trait + +### BroadcastableOTF + +支持内联广播的操作 Trait。 + +**语义**： +``` +for i // <- broadcast dim + for j + dst[i, j] = some_op(src1[0, j], src2[0, j], ..., srcN[0, j]) +``` + +**约束**： +1. 实现 DestinationStyleOpInterface +2. 必须有 `DenseI64ArrayAttr` 名为 `"broadcast"` +3. broadcast 维度唯一 +4. broadcast 维度在 `[0, rank(dst))` 范围内 +5. broadcast 维度上 `dim(src_i, d) = 1 || dim(src_i, d) = dim(dst, d)` +6. 非 broadcast 维度上 `dim(src_i, d) = dim(dst, d)` + +**依赖**：HIVMStructuredOpInterface, HIVMOpSameOperandsAndResultRank + +### TransposableOTF + +支持内联转置的操作 Trait。 + +**语义**： +``` +for i // transpose = [1, 0] + for j + dst[i, j] = some_op(src1[j, i], src2[j, i], ..., srcN[j, i]) +``` + +**约束**： +1. 实现 DestinationStyleOpInterface +2. 必须有 `DenseI64ArrayAttr` 名为 `"transpose"` +3. transpose 是 `range(rank(dst))` 的排列 +4. `transpose[rank(dst) - 1] = rank(dst) - 1` +5. `dim(dst, d) = dim(src_i, transpose[d])` + +**依赖**：HIVMStructuredOpInterface, HIVMOpSameOperandsAndResultRank + +## Pipe 相关 Trait + +### SinglePipeOpTrait + +标识操作为单 Pipe 操作。 + +### OpPipeTrait\ + +参数化 Trait，声明操作在单个 Pipeline 上执行。 + +| 操作 | Pipe | 说明 | +|------|------|------| +| embedding_gather | PIPE_V | Vector Pipe | +| indirect_load | PIPE_V | Vector Pipe | +| indirect_store | PIPE_V | Vector Pipe | +| gatherT | PIPE_V | Vector Pipe | +| scatterT | PIPE_V | Vector Pipe | +| index_put | PIPE_V | Vector Pipe | +| custom | 由属性指定 | 用户指定 | + +**依赖**：SinglePipeOpTrait + +### MacroOpTrait + +标识操作为宏操作（跨 Pipeline）。 + +### MacroOpPipeTrait\ + +参数化 Trait，声明宏操作的输入/输出 Pipeline。 + +| 操作 | InOutPipes | 说明 | +|------|-----------|------| +| mmadL1 / batchMmadL1 | PIPE::PIPE_MTE1, PIPE::PIPE_M | L1 加载 + Cube 计算 | +| matmul / mix_matmul / mix_group_matmul | PIPE::PIPE_MTE2, PIPE::PIPE_MTE3 | GM 加载 + GM 写回 | + +**依赖**：MacroOpTrait + +## Core Type Trait + +### CoreTypeTrait\ + +参数化 Trait，静态声明操作的 Core 类型。 + +| Trait | CoreType | 适用操作 | +|-------|----------|---------| +| VectorCoreTypeTrait | TCoreType::VECTOR | 大多数 Vector 操作 | +| CubeCoreTypeTrait | TCoreType::CUBE | mmadL1, batchMmadL1 | +| CubeVectorCoreTypeTrait | TCoreType::CUBE_OR_VECTOR | get_block_idx, pointer_cast, set_ffts_base_addr | + +## 其他 Trait + +### NoLibraryFunctionTrait + +标识操作没有预定义的库函数实现。 + +| 操作 | 说明 | +|------|------| +| batchMmadL1 | 无库函数，不支持直接 lowering | + +### CommutativeOpTrait + +标识操作的输入操作数可交换。 + +### VectorOnlyTrait\ + +指定操作数只支持 Vector（shaped）类型输入。 + +### ScalarOnlyTrait\ + +指定操作数只支持标量类型输入。 + +### OperElemTypeConstraints\ + +指定操作数的元素类型约束。 + +### UniformReassociationFlattenTrait + +标识操作可以统一展平维度。 + +**依赖**：HIVMStructuredOpInterface + +### CollapsibleConsecutiveTargetDimsTrait + +标识操作在展平时必须保留目标维度的独立性和 rank。 + +**依赖**：HIVMStructuredOpInterface, UniformReassociationFlattenTrait + +## Trait 依赖关系 + +``` +HIVMStructuredOpInterface +├── HIVMOpSameOperandsAndResultRank +│ ├── ElementwiseNaryOpTrait +│ ├── BroadcastableOTF +│ └── TransposableOTF +├── UniformReassociationFlattenTrait +│ └── CollapsibleConsecutiveTargetDimsTrait +└── OpLibraryMaxRankTrait + ├── StaticMaxRankTrait + ├── InferMaxRankTrait + └── NoMaxRankTrait + +SinglePipeOpTrait +└── OpPipeTrait + +MacroOpTrait +└── MacroOpPipeTrait +``` + +## 相关文档 + +- 源码参考：[HIVMTraits.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMTraits.td) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/07-Intrin-Operations/00-overview.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/07-Intrin-Operations/00-overview.md new file mode 100644 index 00000000..4332b5c0 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/01-HIVM-Dialect/07-Intrin-Operations/00-overview.md @@ -0,0 +1,167 @@ +# HIVM 内建指令操作 + +> 关键词：IntrinOp, GET.BLOCK.IDX, SET.FLAG.IMM, WAIT.FLAG.IMM, BARRIER, SET.FFTS.BASE.ADDR, SET.CROSS.CORE, INTRA.BLOCK + +## 概述 + +HIVM 内建指令操作（Intrinsic Operations）是 HIVM 方言中最底层的操作，直接对应 NPU 硬件指令。这些操作从 [HIVMIntrinOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMIntrinOps.td) 定义，继承自 `LLVM_IntrOpBase`，在 Lowering 的最终阶段由高层 HIVM 操作转换而来。 + +内建指令操作通常不由用户直接编写，而是由编译器 Pass 自动生成。了解这些操作有助于理解 HIVM IR 到硬件指令的映射关系。 + +## 指令分类 + +### 1. Block 索引指令 + +| 指令 | 高层操作 | 说明 | +|------|---------|------| +| `hivm.GET.BLOCK.IDX` | `hir.get_block_idx` | 获取当前 Block 索引 | +| `hivm.GET.BLOCK.NUM` | `hir.get_block_num` | 获取 Block 总数 | +| `hivm.GET.SUBBLOCKID` | `hir.get_sub_block_idx` | 获取当前 Sub-Block 索引 | +| `hivm.GET.SUBBLOCKDIM` | `hir.get_sub_block_num` | 获取 Sub-Block 总数 | + +### 2. Pipe 同步指令 + +| 指令 | 参数 | 说明 | +|------|------|------| +| `hivm.SET.FLAG.IMM` | set_pipe(I64Attr), wait_pipe(I64Attr), event_id(I64Attr) | 设置 Event Flag（立即数模式） | +| `hivm.WAIT.FLAG.IMM` | set_pipe(I64Attr), wait_pipe(I64Attr), event_id(I64Attr) | 等待 Event Flag（立即数模式） | +| `hivm.SET.FLAG.REG` | set_pipe(I64Attr), wait_pipe(I64Attr), event_id(I64) | 设置 Event Flag（寄存器模式） | +| `hivm.WAIT.FLAG.REG` | set_pipe(I64Attr), wait_pipe(I64Attr), event_id(I64) | 等待 Event Flag（寄存器模式） | +| `hivm.BARRIER` | pipe(I64Attr) | Pipeline 屏障 | + +### 3. FFTS 同步指令 + +| 指令 | 参数 | 说明 | +|------|------|------| +| `hivm.SET.FFTS.BASE.ADDR` | addr(I<64>) | 设置 FFTS 基地址寄存器 | +| `hivm.SET.CROSS.CORE` | pipe(I64Attr), config(I<64>) | FFTS 跨核同步指令 | +| `hivm.WAIT.FLAG.DEV.REG` | flag_id(I<64>) | FFTS Block/Sub-Block 同步等待（寄存器） | +| `hivm.WAIT.FLAG.DEV.PIPE.IMM` | pipe(I64Attr), flag_id(I64Attr) | FFTS 同步等待（Pipe + Flag 立即数） | +| `hivm.WAIT.FLAG.DEV.PIPE.REG` | pipe(I64Attr), flag_id(I<64>) | FFTS 同步等待（Pipe 立即数 + Flag 寄存器） | + +### 4. Intra-Block 同步指令 + +| 指令 | 参数 | 说明 | +|------|------|------| +| `hivm.SET.INTRA.BLOCK.mode` | pipe(I64Attr), sync_id(I<64>) | Block 内同步设置（寄存器模式） | +| `hivm.WAIT.INTRA.BLOCK.mode` | pipe(I64Attr), sync_id(I<64>) | Block 内同步等待（寄存器模式） | +| `hivm.SET.INTRA.BLOCKI.mode` | pipe(I64Attr), sync_id(I64Attr) | Block 内同步设置（立即数模式） | +| `hivm.WAIT.INTRA.BLOCKI.mode` | pipe(I64Attr), sync_id(I64Attr) | Block 内同步等待（立即数模式） | + +### 5. 控制指令 + +| 指令 | 参数 | 说明 | +|------|------|------| +| `hivm.SET.MASK.NORM` | 无 | 设置 Mask 为正常模式 | +| `hivm.GET.CTRL` | 无 | 获取控制寄存器值 | +| `hivm.SET.CTRL` | config(I<64>) | 设置控制寄存器 | +| `hivm.SBITSET0` | x(I<64>), idx(I<64>) | 设置状态位为 0 | +| `hivm.SBITSET1` | x(I<64>), idx(I<64>) | 设置状态位为 1 | + +### 6. 缓存指令 + +| 指令 | 参数 | 说明 | +|------|------|------| +| `hivm.DCCI.DST` | ptr(LLVMPointer), entire(I<64>), dst(I<64>) | 清除/无效化 GM 数据缓存 | +| `hivm.DCCI.DST.UB` | ptr(LLVMPointer), entire(I<64>), dst(I<64>) | 清除/无效化 UB 数据缓存 | + +## 指令详细说明 + +### SET.FLAG.IMM / WAIT.FLAG.IMM + +立即数模式的 Pipe 同步指令，Event ID 编码在指令中。 + +``` +// 参数 +set_pipe: I64Attr -- 发送信号的 Pipe（立即数） +wait_pipe: I64Attr -- 接收信号的 Pipe（立即数） +event_id: I64Attr -- Event ID（立即数） + +// 无返回值 +``` + +### SET.FLAG.REG / WAIT.FLAG.REG + +寄存器模式的 Pipe 同步指令，Event ID 来自寄存器值。 + +``` +// 参数 +set_pipe: I64Attr -- 发送信号的 Pipe（立即数） +wait_pipe: I64Attr -- 接收信号的 Pipe（立即数） +event_id: I<64> -- Event ID（寄存器值） + +// 无返回值 +``` + +### SET.CROSS.CORE + +FFTS 跨核同步指令，发送数据（包括模式和 Flag ID）到 FFTS 目标地址。 + +``` +// 参数 +pipe: I64Attr -- Pipe 类型（立即数） +config: I<64> -- 配置值（寄存器） + +// config 编码格式： +// config = (0x0001 | ((mode & 0x0f) << 4) | ((flagID & 0x0f) << 8)) +// +// mode: +// 0: inter block synchronization +// 1: inter subblock synchronization +// 2: intra block synchronization +// +// flagID: 8-bit flag ID，每个 flag ID 有一个计数器 +``` + +### INTRA.BLOCK 同步指令 + +Block 内 Cube Core 和 Vector Core 之间的同步指令。 + +**ID 映射规则**（Mix-Mode Block，1 CUBECORE + 2 VECCORE）： + +| SET 来源 | 目标 | ID 映射 | +|----------|------|---------| +| VECCORE0 ID 0-15 | CUBECORE | ID 0-15 | +| VECCORE1 ID 0-15 | CUBECORE | ID 16-31 | +| CUBECORE ID 0-15 | VECCORE0 | ID 0-15 | +| CUBECORE ID 16-31 | VECCORE1 | ID 0-15 | + +**同步机制**： +- 每个 sync ID 有一个 4-bit 计数器 +- SET 操作：递增对方 Core 对应 ID 的计数器，不阻塞当前 Pipeline +- WAIT 操作：如果对应 ID 的计数器为 0 则阻塞，否则递减计数器 + +## 高层操作到内建指令的映射 + +| 高层操作 | 内建指令 | 说明 | +|---------|---------|------| +| `hir.get_block_idx` | `hivm.GET.BLOCK.IDX` | Block 索引查询 | +| `hir.get_block_num` | `hivm.GET.BLOCK.NUM` | Block 数量查询 | +| `hir.set_flag` (静态 Event ID) | `hivm.SET.FLAG.IMM` | 立即数模式同步 | +| `hir.set_flag` (动态 Event ID) | `hivm.SET.FLAG.REG` | 寄存器模式同步 | +| `hir.wait_flag` (静态 Event ID) | `hivm.WAIT.FLAG.IMM` | 立即数模式等待 | +| `hir.wait_flag` (动态 Event ID) | `hivm.WAIT.FLAG.REG` | 寄存器模式等待 | +| `hir.pipe_barrier` | `hivm.BARRIER` | Pipeline 屏障 | +| `hir.set_ffts_base_addr` | `hivm.SET.FFTS.BASE.ADDR` | FFTS 基地址设置 | +| `hir.sync_block_set` | `hivm.SET.CROSS.CORE` / `hivm.SET.INTRA.BLOCKI.mode` | 跨核/核内同步 | +| `hir.sync_block_wait` | `hivm.WAIT.FLAG.DEV.REG` / `hivm.WAIT.INTRA.BLOCKI.mode` | 跨核/核内等待 | +| `hir.set_mask_norm` | `hivm.SET.MASK.NORM` | Mask 模式设置 | +| `hir.set_ctrl` | `hivm.SET.CTRL` | 控制寄存器设置 | +| `hir.dcci` | `hivm.DCCI.DST` / `hivm.DCCI.DST.UB` | 缓存操作 | + +## 常见问题 + +**Q: 为什么有 IMM 和 REG 两种模式？** +A: IMM（Immediate）模式将参数编码在指令中，适用于编译时已知的常量参数，执行效率更高。REG（Register）模式从寄存器读取参数，适用于运行时计算的动态参数。 + +**Q: SET.CROSS.CORE 的 config 如何编码？** +A: `config = 0x0001 | ((mode & 0x0f) << 4) | ((flagID & 0x0f) << 8)`，其中 mode 为同步模式（0=跨Block，1=跨Sub-Block，2=Block内），flagID 为 8-bit Flag ID。 + +**Q: INTRA.BLOCK 同步的 ID 映射规则是什么？** +A: 在 Mix-Mode Block 中，CUBECORE 有 32 个 ID，每个 VECCORE 有 16 个 ID。VECCORE0 的 ID 0-15 映射到 CUBECORE 的 ID 0-15，VECCORE1 的 ID 0-15 映射到 CUBECORE 的 ID 16-31，反之亦然。 + +## 相关文档 + +- 源码参考：[HIVMIntrinOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HIVM/IR/HIVMIntrinOps.td) +- 高层同步操作：[04-Synchronization/01-pipe-sync.md](../04-Synchronization/01-pipe-sync.md) +- 跨核同步：[04-Synchronization/02-block-sync.md](../04-Synchronization/02-block-sync.md) diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/00-overview.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/00-overview.md new file mode 100644 index 00000000..a3b1654b --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/00-overview.md @@ -0,0 +1,81 @@ +# HACC 方言总览 + +## 1. 简介 + +HACC（Heterogeneous Async Computing Call，异构异步计算调用）方言是 AscendNPU-IR 中用于描述 Host-Device 异构计算模型的核心方言。它定义了函数的 Host/Device 归属、NPU 设备规格参数、Kernel 参数类型以及 Host-Device 函数绑定关系，为后续的编译流程提供了异构语义基础。 + +- **方言名称**：`hacc` +- **C++ 命名空间**：`::mlir::hacc` +- **依赖方言**：`mlir::DLTIDialect` + +## 2. 核心概念 + +| 概念 | 说明 | +|------|------| +| HACCFuncType | 函数的 Host/Device 归属分类 | +| DeviceSpecEnum | NPU 硬件规格参数枚举 | +| TargetDeviceSpecAttr | 目标设备规格属性，映射具体 NPU 型号 | +| KernelArgType | Kernel 参数类型分类 | +| HostFuncType | Host 端函数角色分类 | +| HACCFunctionInterface | 异构函数接口，提供查询与设置方法 | + +## 3. 源码位置 + +| 文件 | 路径 | +|------|------| +| 方言基类 | [HACCBase.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/IR/HACCBase.td) | +| 属性与枚举 | [HACCAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/IR/HACCAttrs.td) | +| 接口定义 | [HACCInterfaces.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/IR/HACCInterfaces.td) | +| 变换 Pass | [Passes.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Transforms/Passes.td) | + +## 4. 方言定义 + +```tablegen +def HACC_Dialect : Dialect { + let name = "hacc"; + let description = [{ + Heterogeneous Async Computing Call (HACC) dialect. + }]; + let cppNamespace = "::mlir::hacc"; + let useDefaultAttributePrinterParser = 1; + let dependentDialects = ["mlir::DLTIDialect"]; +} +``` + +## 5. 与其他方言的关系 + +``` +HACC 方言 + ├── 被 HFusion 方言依赖（HFusion 依赖 hacc::HACCDialect） + ├── 被 Scope 方言间接使用（scope.scope 可携带 tcore_type 等属性） + ├── 与 DLTIDialect 协作（设备规格通过 DLTI 机制存储） + └── 为 HIVM 层提供 Host-Device 调用约定基础 +``` + +## 6. 典型 IR 示例 + +```mlir +module { + func.func @host_kernel(%arg0: tensor) + attributes {hacc.function_kind = #hacc.function_kind} { + return + } + + func.func @device_kernel(%arg0: tensor, + %arg1: i64 {hacc.arg_type = #hacc.arg_type}) + attributes {hacc.function_kind = #hacc.function_kind, + hacc.tiling_function = #hacc.tiling_function<@tiling_func>} { + return + } +} +``` + +## 7. 文档索引 + +| 文档 | 内容 | +|------|------| +| [01-function-management.md](01-function-management.md) | HOST/DEVICE 函数类型与 HACCFunctionInterface | +| [02-device-specification.md](02-device-specification.md) | NPU 设备规格参数与型号映射 | +| [03-kernel-args.md](03-kernel-args.md) | Kernel 参数类型完整列表 | +| [04-host-device-binding.md](04-host-device-binding.md) | Host-Device 函数绑定关系 | +| [05-transforms.md](05-transforms.md) | HACC 变换 Pass | diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/01-function-management.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/01-function-management.md new file mode 100644 index 00000000..0fb00250 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/01-function-management.md @@ -0,0 +1,135 @@ +# HACC 函数管理 + +## 1. 概述 + +HACC 方言通过 `HACCFuncType` 枚举和 `HACCFunctionInterface` 接口管理异构函数的归属与属性。每个函数被标记为 HOST 或 DEVICE 类型，并通过接口方法查询和设置函数属性。 + +## 2. HACCFuncType 枚举 + +`HACCFuncType` 枚举定义了函数的异构归属类型。 + +| 枚举值 | 整数值 | 助记符 | 说明 | +|--------|--------|--------|------| +| `HOST` | 1 | `HOST` | Host 端函数，运行在 CPU 上 | +| `DEVICE` | 2 | `DEVICE` | Device 端函数，运行在 NPU 上 | + +### 属性定义 + +```tablegen +def HACC_FuncTypeAttr : HACC_Attr<"HACCFuncType", "function_kind"> { + let parameters = (ins EnumParameter:$function_kind); + let assemblyFormat = "`<` params `>`"; +} +``` + +### MLIR 表示 + +```mlir +hacc.function_kind = #hacc.function_kind +hacc.function_kind = #hacc.function_kind +``` + +## 3. HACCFunctionInterface 接口 + +`HACCFunctionInterface` 继承自 `FunctionOpInterface`，为异构函数提供统一的查询和设置方法。 + +> 源码参考：[HACCInterfaces.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/IR/HACCInterfaces.td) + +### 3.1 查询方法 + +| 方法名 | 返回类型 | 参数 | 说明 | +|--------|----------|------|------| +| `getHACCFuncType()` | `std::optional` | 无 | 返回函数的 HACC 类型，未知时返回 `std::nullopt` | +| `isHost()` | `bool` | 无 | 判断是否为 Host 函数 | +| `isDevice()` | `bool` | 无 | 判断是否为 Device 函数 | +| `isDeviceEntry()` | `bool` | 无 | 判断是否为 Device 入口函数 | +| `getHostFuncType()` | `std::optional` | 无 | 返回 Host 函数类型，非 Host 函数返回 `std::nullopt` | + +### 3.2 设置方法 + +| 方法名 | 返回类型 | 参数 | 说明 | +|--------|----------|------|------| +| `setDevice()` | `void` | 无 | 将函数标记为 Device 函数，自动移除不允许的属性 | +| `setDeviceEntry()` | `void` | 无 | 将函数标记为 Device 入口函数，自动移除不允许的属性 | +| `setHost()` | `void` | 无 | 将函数标记为 Host 函数，自动移除不允许的属性 | +| `setHostFuncType(::mlir::hacc::HostFuncType)` | `void` | `funcType` | 设置 Host 函数的具体角色类型 | + +### 3.3 参数查询方法 + +| 方法名 | 返回类型 | 参数 | 说明 | +|--------|----------|------|------| +| `isKernelArg(int, ::mlir::hacc::KernelArgType)` | `bool` | `argIdx`, `argType` | 判断第 `argIdx` 个参数是否具有指定的 `hacc.arg_type` | + +## 4. HACC To LLVM Translation 属性 + +`HACCToLLVMIRTranslateAttr` 枚举用于标记函数在 LLVM IR 翻译阶段的行为。 + +| 枚举值 | 整数值 | 助记符 | 说明 | +|--------|--------|--------|------| +| `ENTRY` | 0 | `hacc.entry` | Device 入口函数 | +| `MIX_ENTRY` | 1 | `hacc.mix_entry` | 混合 Device 入口函数 | +| `ALWAYS_INLINE` | 2 | `hacc.always_inline` | 始终内联函数 | + +## 5. 辅助属性 + +### 5.1 RenameFuncAttr + +```tablegen +def HACC_RenameFuncAttr : HACC_Attr<"RenameFunc", "rename_func"> { + let parameters = (ins AttrParameter<"::mlir::FlatSymbolRefAttr">:$targetName); +} +``` + +指示当前函数应重命名为目标函数名。`hacc-rename-func` Pass 会据此执行重命名。 + +### 5.2 InputIdxAttr / OutputIdxAttr + +| 属性 | 参数 | 说明 | +|------|------|------| +| `hacc.input_idx` | `unsigned:$argIdx` | 标记函数参数为输入索引 | +| `hacc.output_idx` | `unsigned:$argIdx` | 标记函数参数为输出索引（NPU Kernel 的输出通过输入参数传递） | + +### 5.3 其他辅助属性 + +| 属性名 | 助记符 | 说明 | +|--------|--------|------| +| `ExportAsDAG` | `export_as_dag` | 将函数导出为 DAG | +| `DummyFunc` | `dummy_func` | 虚拟函数标记 | +| `ExternalFunctionPath` | `external_function_path` | 外部函数路径 | +| `CachedIO` | `cached_io` | 标记值已被缓存 IO | +| `NoIOAlias` | `no_io_alias` | 标记函数输入输出严格不别名 | +| `BlockDim` | `block_dim` | 函数的 Block 维度属性 | +| `SIMTModule` | `simt_module` | 标记 SIMT 模块 | + +## 6. 典型使用模式 + +### 6.1 Host 函数 + +```mlir +func.func @tiling_func(%arg0: tensor) -> (i64, i64) + attributes {hacc.function_kind = #hacc.function_kind} { + %c = arith.constant 42 : i64 + return %c, %c : i64, i64 +} +``` + +### 6.2 Device 入口函数 + +```mlir +func.func @kernel_entry(%arg0: tensor, + %arg1: i64 {hacc.arg_type = #hacc.arg_type}) + attributes {hacc.function_kind = #hacc.function_kind, + hacc.tiling_function = #hacc.tiling_function<@tiling_func>} { + return +} +``` + +### 6.3 函数重命名 + +```mlir +func.func @bar() attributes {hacc.rename_func = #hacc.rename_func<@foo>} { + return +} +``` + +经过 `hacc-rename-func` Pass 后，`@bar` 将被重命名为 `@foo`。 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/02-device-specification.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/02-device-specification.md new file mode 100644 index 00000000..aadfc712 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/02-device-specification.md @@ -0,0 +1,181 @@ +# NPU 设备规格参数 + +## 1. 概述 + +HACC 方言通过 `DeviceSpecEnum` 枚举和 `TargetDeviceSpecAttr` 属性描述 NPU 设备的硬件规格。编译器可基于这些规格进行 Tiling、内存分配等优化决策。 + +> 源码参考：[HACCAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/IR/HACCAttrs.td#L241-L289)、[Passes.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Transforms/Passes.td#L51-L110) + +## 2. DeviceSpecEnum 完整列表 + +| 枚举值 | 整数值 | 说明 | +|--------|--------|------| +| `AI_CORE_COUNT` | 0 | AI Core 数量 | +| `CUBE_CORE_COUNT` | 1 | Cube Core 数量（矩阵乘单元） | +| `VECTOR_CORE_COUNT` | 2 | Vector Core 数量（向量计算单元） | +| `UB_SIZE` | 3 | Unified Buffer 大小（字节） | +| `L1_SIZE` | 4 | L1 缓存大小（字节） | +| `L0A_SIZE` | 5 | L0A 缓存大小（字节，矩阵乘 A 矩阵缓冲） | +| `L0B_SIZE` | 6 | L0B 缓存大小（字节，矩阵乘 B 矩阵缓冲） | +| `L0C_SIZE` | 7 | L0C 缓存大小（字节，矩阵乘 C 结果缓冲） | +| `UB_ALIGN_SIZE` | 8 | Unified Buffer 对齐大小（字节） | +| `L1_ALIGN_SIZE` | 9 | L1 缓存对齐大小（字节） | +| `L0C_ALIGN_SIZE` | 10 | L0C 缓存对齐大小（字节） | +| `MINIMAL_D_CACHE_SIZE` | 11 | 最小 D Cache 大小 | +| `MAXIMUM_D_CACHE_SIZE` | 12 | 最大 D Cache 大小 | +| `ARCH` | 13 | 架构版本标识 | + +### TableGen 定义 + +```tablegen +def HACC_DeviceSpecEnum : + HACC_I32Enum<"DeviceSpec", "HACC device spec", [ + I32EnumAttrCase<"AI_CORE_COUNT", 0>, + I32EnumAttrCase<"CUBE_CORE_COUNT", 1>, + I32EnumAttrCase<"VECTOR_CORE_COUNT", 2>, + I32EnumAttrCase<"UB_SIZE", 3>, + I32EnumAttrCase<"L1_SIZE", 4>, + I32EnumAttrCase<"L0A_SIZE", 5>, + I32EnumAttrCase<"L0B_SIZE", 6>, + I32EnumAttrCase<"L0C_SIZE", 7>, + I32EnumAttrCase<"UB_ALIGN_SIZE", 8>, + I32EnumAttrCase<"L1_ALIGN_SIZE", 9>, + I32EnumAttrCase<"L0C_ALIGN_SIZE", 10>, + I32EnumAttrCase<"MINIMAL_D_CACHE_SIZE", 11>, + I32EnumAttrCase<"MAXIMUM_D_CACHE_SIZE", 12>, + I32EnumAttrCase<"ARCH", 13> +]> +``` + +## 3. TargetDeviceSpecAttr + +`TargetDeviceSpecAttr` 用于表示具体 NPU 设备的规格参数集合，基于 DLTI（Data Layout Target Information）机制存储。 + +### 3.1 属性定义 + +```tablegen +def HACC_TargetDeviceSpecAttr : + HACC_Attr<"TargetDeviceSpec", "target_device_spec", + [TargetDeviceSpecTrait, HACCTargetDeviceSpecTrait]> { + let parameters = (ins + ArrayRefParameter<"DataLayoutEntryInterface", "single spec entry">:$entries + ); + let assemblyFormat = "`<` $entries `>`"; +} +``` + +### 3.2 MLIR 表示 + +```mlir +#hacc.target_device_spec< + #dlti.dl_entry<"UB_SIZE", 196608 : i32>> +``` + +### 3.3 HACCTargetDeviceSpecInterface + +该接口继承自 `TargetDeviceSpecInterface`，提供按枚举值查询规格的方法： + +| 方法名 | 返回类型 | 参数 | 说明 | +|--------|----------|------|------| +| `getSpecForIdentifierEnum(DeviceSpec)` | `::mlir::DataLayoutEntryInterface` | `identifier` | 根据枚举值返回对应的规格条目 | + +## 4. TargetAttr + +`TargetAttr` 用于指示目标设备名称。 + +```tablegen +def HACC_TargetAttr : HACC_Attr<"Target", "target"> { + let parameters = (ins AttrParameter<"StringAttr", "target device">:$target); + let assemblyFormat = "`<` $target `>`"; +} +``` + +## 5. NPU 型号映射 + +`hacc-append-device-spec` Pass 通过 `--target` 选项指定 NPU 型号，自动附加设备规格信息。支持的型号列表如下： + +### 5.1 Ascend 910B 系列 + +| 型号标识 | +|----------| +| `Ascend910B1` | +| `Ascend910B2` | +| `Ascend910B3` | +| `Ascend910B4` | + +### 5.2 Ascend 910_93 系列 + +| 型号标识 | +|----------| +| `Ascend910_9362` | +| `Ascend910_9372` | +| `Ascend910_9381` | +| `Ascend910_9382` | +| `Ascend910_9391` | +| `Ascend910_9392` | + +### 5.3 Ascend 310B 系列 + +| 型号标识 | +|----------| +| `Ascend310B1` | +| `Ascend310B2` | +| `Ascend310B3` | +| `Ascend310B4` | + +### 5.4 Ascend 950 系列 + +| 型号标识 | 型号标识 | +|----------|----------| +| `Ascend910_950z` | `Ascend950PR_950z` | +| `Ascend910_9579` | `Ascend950PR_9579` | +| `Ascend910_957b` | `Ascend950PR_957a` | +| `Ascend910_957d` | `Ascend950PR_957b` | +| `Ascend910_9581` | `Ascend950PR_957c` | +| `Ascend910_9589` | `Ascend950PR_957d` | +| `Ascend910_958a` | `Ascend950PR_9589` | +| `Ascend910_958b` | `Ascend950PR_958a` | +| `Ascend910_9599` | `Ascend950PR_958b` | +| | `Ascend950PR_958c` | +| | `Ascend950PR_958d` | +| | `Ascend950PR_9599` | +| | `Ascend950PR_959a` | +| | `Ascend950PR_959b` | + +### 5.5 Ascend 950DT 系列 + +| 型号标识 | +|----------| +| `Ascend950DT_950x` | +| `Ascend950DT_950y` | +| `Ascend950DT_9571` ~ `Ascend950DT_9578` | +| `Ascend950DT_9581` ~ `Ascend950DT_9588` | +| `Ascend950DT_9591`, `Ascend950DT_9592` | +| `Ascend950DT_9595`, `Ascend950DT_9596` | +| `Ascend950DT_95A1`, `Ascend950DT_95A2` | + +## 6. 内存层次示意 + +``` ++-------------------------------------------+ +| GM (Global Memory) | ++-------------------------------------------+ + | | ++-------v--------+ +-------v--------+ +| L1 Cache | | L1 Cache | ++-------+--------+ +-------+--------+ + | | ++-------v--------+ +-------v--------+ +| L0A | L0B | L0C| | L0A | L0B | L0C| ++-----+-----+----+ +-----+-----+----+ + | | ++-------v--------+ +-------v--------+ +| UB (Unified | | UB (Unified | +| Buffer) | | Buffer) | ++----------------+ +----------------+ +``` + +- **GM**：全局内存，Host 与 Device 共享 +- **L1**：AI Core 内部一级缓存 +- **L0A/L0B/L0C**：Cube 单元专用缓冲 +- **UB**：统一缓冲区，Vector 计算单元的工作空间 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/03-kernel-args.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/03-kernel-args.md new file mode 100644 index 00000000..c351aa22 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/03-kernel-args.md @@ -0,0 +1,127 @@ +# Kernel 参数类型 + +## 1. 概述 + +HACC 方言通过 `KernelArgType` 枚举对 Kernel 函数的参数进行语义分类。每个参数可附加 `hacc.arg_type` 属性，标识其在 Kernel 执行中的角色。 + +> 源码参考：[HACCAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/IR/HACCAttrs.td#L77-L104) + +## 2. KernelArgType 完整列表 + +| 枚举值 | 整数值 | 助记符 | 说明 | +|--------|--------|--------|------| +| `kFFTSBaseAddr` | 0 | `ffts_base_address` | FFTS（Fast Function Task Schedule）基地址 | +| `kInput` | 1 | `input` | Kernel 输入数据 | +| `kOutput` | 2 | `output` | Kernel 输出数据 | +| `kInputAndOutput` | 3 | `input_and_output` | 同时作为输入和输出的数据（原地操作） | +| `kWorkspace` | 4 | `workspace` | 工作空间内存 | +| `kSyncBlockLock` | 5 | `sync_block_lock` | 同步 Block 锁 | +| `kTilingKey` | 6 | `tiling_key` | Tiling 键值（用于选择 Tiling 策略） | +| `kTilingData` | 7 | `tiling_data` | Tiling 数据（动态 Tiling 参数） | +| `kTilingStruct` | 8 | `tiling_struct` | Tiling 结构体（打包的 Tiling 参数） | +| `kMeshArg` | 9 | `mesh_arg` | Mesh 参数（多卡通信） | +| `kSanitizerAddr` | 10 | `sanitizer_addr` | Sanitizer 地址（用于调试） | +| `kGMAddr` | 11 | `gm_addr` | 全局内存地址 | + +### TableGen 定义 + +```tablegen +def HACC_KernelArgTypeEnum + : HACC_I32Enum<"KernelArgType", "HACC Kernel Arg Category", + [HACC_kFFTSBaseAddr, HACC_kInput, HACC_kOutput, + HACC_kInputAndOutput, HACC_kWorkspace, HACC_kSyncBlockLock, + HACC_kTilingKey, HACC_kTilingData, HACC_kTilingStruct, + HACC_kMeshArg, HACC_kSanitizerAddr, HACC_kGMAddr]> +``` + +## 3. KernelArgTypeAttr 属性 + +```tablegen +def HACC_KernelArgTypeAttr : HACC_Attr<"KernelArgType", "arg_type"> { + let parameters = (ins EnumParameter:$arg_type); + let assemblyFormat = "`<` params `>`"; +} +``` + +## 4. 各参数类型详解 + +### 4.1 kFFTSBaseAddr（ffts_base_address） + +FFTS 基地址参数，用于 Kernel 启动时的快速函数任务调度。FFTS 机制允许 Host 端预先配置 Kernel 的执行参数，减少启动延迟。 + +### 4.2 kInput / kOutput / kInputAndOutput + +| 类型 | 说明 | +|------|------| +| `input` | 标记参数为只读输入，Kernel 从中读取数据 | +| `output` | 标记参数为只写输出，Kernel 向其写入结果 | +| `input_and_output` | 标记参数为读写，Kernel 既读取又写入（原地操作） | + +NPU Kernel 的调用约定中，输出参数也作为输入参数传入（out-param 模式），`hacc.output_idx` 属性用于标记哪个输入参数对应哪个输出值。 + +### 4.3 kWorkspace + +工作空间内存参数，用于 Kernel 执行过程中需要的临时存储。工作空间大小通过 `hacc.infer_workspace_shape_function` 在 Host 端计算。 + +### 4.4 kSyncBlockLock + +同步 Block 锁参数，用于多 Block 间的原子同步。每个原子操作需要一个 `<1xi64>` 类型的 memref 作为锁，锁的数量通过 `hacc.infer_sync_block_lock_num_function` 推断。 + +### 4.5 kTilingKey / kTilingData / kTilingStruct + +| 类型 | 说明 | +|------|------| +| `tiling_key` | Tiling 策略选择键，Host 端根据输入形状计算 | +| `tiling_data` | 单个 Tiling 参数（如 tile 大小、步长等） | +| `tiling_struct` | 打包的 Tiling 结构体，包含多个 Tiling 参数 | + +`hacc-pack-tiling-data` Pass 可将多个 `tiling_data` 参数打包为单个 `tiling_struct`。 + +### 4.6 kMeshArg + +Mesh 参数，用于多卡通信场景，与 HMAP 方言的集合通信操作配合使用。 + +### 4.7 kSanitizerAddr + +Sanitizer 地址参数，用于 Device 端内存调试和越界检测。 + +### 4.8 kGMAddr + +全局内存地址参数，用于直接传递 GM 地址给 Kernel。 + +## 5. MLIR 示例 + +### 5.1 带 Tiling 数据的 Kernel + +```mlir +func.func @device_kernel( + %input: memref {hacc.arg_type = #hacc.arg_type}, + %output: memref {hacc.arg_type = #hacc.arg_type}, + %tiling_key: i64 {hacc.arg_type = #hacc.arg_type}, + %tiling_data: i64 {hacc.arg_type = #hacc.arg_type} +) attributes {hacc.function_kind = #hacc.function_kind} { + return +} +``` + +### 5.2 带 Workspace 和 SyncBlockLock 的 Kernel + +```mlir +func.func @kernel_with_workspace( + %input: memref {hacc.arg_type = #hacc.arg_type}, + %output: memref {hacc.arg_type = #hacc.arg_type}, + %workspace: memref {hacc.arg_type = #hacc.arg_type}, + %lock: memref<1xi64> {hacc.arg_type = #hacc.arg_type} +) attributes {hacc.function_kind = #hacc.function_kind} { + return +} +``` + +### 5.3 使用 isKernelArg 查询 + +```cpp +// 在 C++ 中查询参数类型 +if (funcOp.isKernelArg(argIdx, hacc::KernelArgType::kInput)) { + // 处理输入参数 +} +``` diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/04-host-device-binding.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/04-host-device-binding.md new file mode 100644 index 00000000..9ea0fc96 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/04-host-device-binding.md @@ -0,0 +1,150 @@ +# Host-Device 函数绑定关系 + +## 1. 概述 + +在 HACC 异构模型中，Device Kernel 的执行需要 Host 端提供多种辅助函数（如 Tiling 计算、形状推断等）。HACC 方言通过 `HostFuncType` 枚举和一系列 `FuncRefAttr` 属性建立 Host-Device 函数间的绑定关系。 + +> 源码参考：[HACCAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/IR/HACCAttrs.td#L110-L225) + +## 2. HostFuncType 完整列表 + +| 枚举值 | 整数值 | 助记符 | 说明 | +|--------|--------|--------|------| +| `kEntry` | 1 | `host_entry` | Host 入口函数 | +| `kTilingFunction` | 2 | `tiling_function` | Tiling 计算函数 | +| `kInferOutputShapeFunction` | 3 | `infer_output_shape_function` | 输出形状推断函数 | +| `kInferWorkspaceShapeFunction` | 4 | `infer_workspace_shape_function` | 工作空间大小推断函数 | +| `kInferSyncBlockLockNumFunction` | 5 | `infer_sync_block_lock_num_function` | 同步锁数量推断函数 | +| `kInferSyncBlockLockInitFunction` | 6 | `infer_sync_block_lock_init_function` | 同步锁初始值推断函数 | +| `kInferVFModeFunction` | 7 | `infer_vf_mode_function` | VF 模式推断函数 | +| `kGetTilingStructSizeFunction` | 8 | `get_tiling_struct_size_function` | Tiling 结构体大小获取函数 | + +### TableGen 定义 + +```tablegen +def HACC_HostFuncTypeEnum + : HACC_I32Enum< + "HostFuncType", "HACC Host function type", + [HACC_kEntry, HACC_kTilingFunction, HACC_kInferOutputShapeFunction, + HACC_kInferWorkspaceShapeFunction, + HACC_kInferSyncBlockLockNumFunction, + HACC_kInferSyncBlockLockInitFunction, HACC_kInferVFModeFunction, + HACC_kGetTilingStructSizeFunction]> +``` + +## 3. HostFuncTypeAttr 属性 + +```tablegen +def HACC_HostFuncTypeAttr : HACC_Attr<"HostFuncType", "host_func_type"> { + let parameters = (ins EnumParameter:$host_func_type); + let assemblyFormat = "`<` params `>`"; +} +``` + +## 4. FuncRefAttr 绑定属性 + +Host 端辅助函数通过 `FuncRefAttr` 系列属性绑定到 Device 函数上。每个属性包含一个 `FlatSymbolRefAttr`，指向对应的 Host 函数符号名。 + +### 4.1 基类定义 + +```tablegen +class HACC_FuncRefAttr + : HACC_Attr { + let parameters = (ins AttrParameter<"::mlir::FlatSymbolRefAttr", + "function symbol name">:$funcName); + let assemblyFormat = "`<` $funcName `>`"; +} +``` + +### 4.2 绑定属性列表 + +| 属性名 | 助记符 | 说明 | +|--------|--------|------| +| `TilingFunctionAttr` | `tiling_function` | 指向 Host 端 Tiling 计算函数 | +| `InferOutputShapeFunctionAttr` | `infer_output_shape_function` | 指向 Host 端输出形状推断函数 | +| `InferWorkspaceShapeFunctionAttr` | `infer_workspace_shape_function` | 指向 Host 端工作空间大小推断函数 | +| `InferSyncBlockLockNumFunctionAttr` | `infer_sync_block_lock_num_function` | 指向 Host 端同步锁数量推断函数 | +| `InferSyncBlockLockInitFunctionAttr` | `infer_sync_block_lock_init_function` | 指向 Host 端同步锁初始值推断函数 | +| `InferVFModeFunctionAttr` | `infer_vf_mode_function` | 指向 Host 端 VF 模式推断函数 | +| `GetTilingStructSizeFunctionAttr` | `get_tiling_struct_size_function` | 指向 Host 端 Tiling 结构体大小获取函数 | + +## 5. 各绑定函数详解 + +### 5.1 TilingFunction + +Tiling 函数在 Host 端执行，根据输入张量的形状和设备规格计算 Tiling 参数（如分块大小、迭代次数等），并将结果传递给 Device Kernel。 + +``` +Host: tiling_func(input_shape, device_spec) -> (tile_m, tile_n, tile_k, ...) +Device: kernel(input, output, tile_m, tile_n, tile_k, ...) +``` + +### 5.2 InferOutputShapeFunction + +推断 Device Kernel 的输出张量形状。对于动态形状的 Kernel，Host 端需要预先知道输出大小以分配内存。 + +### 5.3 InferWorkspaceShapeFunction + +推断 Device Kernel 所需的工作空间大小。工作空间用于存储中间计算结果。 + +### 5.4 InferSyncBlockLockNumFunction + +推断 Kernel 所需的同步锁数量。每个原子操作需要所有 Block 共享一个 `<1xi64>` 类型的 memref 作为锁。 + +### 5.5 InferSyncBlockLockInitFunction + +推断同步锁的初始值。每个锁在 Kernel 执行前需要初始化。 + +### 5.6 InferVFModeFunction + +推断 VF（Vector Function）模式，用于控制向量化执行策略。 + +### 5.7 GetTilingStructSizeFunction + +获取 Tiling 结构体的大小（以 i64 为单位），用于 `hacc-pack-tiling-data` Pass 将多个 Tiling 参数打包为结构体。 + +## 6. MLIR 示例 + +### 6.1 完整的 Host-Device 绑定 + +```mlir +module { + func.func @tiling_func(%arg0: tensor) -> (i64, i64) + attributes {hacc.function_kind = #hacc.function_kind, + hacc.host_func_type = #hacc.host_func_type} { + %c1 = arith.constant 128 : i64 + %c2 = arith.constant 256 : i64 + return %c1, %c2 : i64, i64 + } + + func.func @infer_shape(%arg0: tensor) -> tensor + attributes {hacc.function_kind = #hacc.function_kind, + hacc.host_func_type = #hacc.host_func_type} { + return + } + + func.func @device_kernel( + %input: tensor, + %output: tensor, + %t1: i64 {hacc.arg_type = #hacc.arg_type}, + %t2: i64 {hacc.arg_type = #hacc.arg_type} + ) attributes { + hacc.function_kind = #hacc.function_kind, + hacc.tiling_function = #hacc.tiling_function<@tiling_func>, + hacc.infer_output_shape_function = #hacc.infer_output_shape_function<@infer_shape> + } { + return + } +} +``` + +### 6.2 Host 入口函数 + +```mlir +func.func @host_entry(%arg0: tensor) -> tensor + attributes {hacc.function_kind = #hacc.function_kind, + hacc.host_func_type = #hacc.host_func_type} { + %result = func.call @device_kernel(%arg0, ...) : (tensor, ...) -> tensor + return %result : tensor +} +``` diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/05-transforms.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/05-transforms.md new file mode 100644 index 00000000..546f8402 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/02-HACC-Dialect/05-transforms.md @@ -0,0 +1,121 @@ +# HACC 变换 Pass + +## 1. 概述 + +HACC 方言提供了两个变换 Pass，用于函数重命名和设备规格附加。这些 Pass 在编译流程中为后续的代码生成和运行时调度提供必要的信息。 + +> 源码参考：[Passes.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HACC/Transforms/Passes.td) + +## 2. Pass 列表 + +| Pass 名称 | 作用域 | 构造函数 | +|-----------|--------|----------| +| `hacc-rename-func` | `FuncOp` | `mlir::hacc::createRenameFuncPass()` | +| `hacc-append-device-spec` | `ModuleOp` | `mlir::hacc::createAppendDeviceSpecPass()` | + +## 3. hacc-rename-func + +### 3.1 功能 + +根据 `hacc.rename_func` 属性重命名函数，并更新模块内所有对该函数的引用。 + +### 3.2 变换示例 + +输入： + +```mlir +func.func @bar() attributes {hacc.rename_func = #hacc.rename_func<@foo>} { + return +} + +func.func @caller() { + func.call @bar() : () -> () + return +} +``` + +输出： + +```mlir +func.func @foo() { + return +} + +func.func @caller() { + func.call @foo() : () -> () + return +} +``` + +### 3.3 约束 + +- 目标函数名不能与模块中已有函数名冲突 + +## 4. hacc-append-device-spec + +### 4.1 功能 + +根据指定的 NPU 型号，向模块附加设备规格信息（`#hacc.target_device_spec`）。规格信息基于 DLTI 机制存储，供后续编译 Pass 查询使用。 + +### 4.2 选项 + +| 选项名 | 类型 | 默认值 | 说明 | +|--------|------|--------|------| +| `target` | `::mlir::hacc::TargetDevice` | `Unknown` | 目标设备名称 | + +### 4.3 依赖方言 + +- `hacc::HACCDialect` +- `mlir::DLTIDialect` + +### 4.4 支持的设备型号 + +通过 `--target` 选项指定，完整列表参见 [02-device-specification.md](02-device-specification.md) 第 5 节。主要系列包括： + +- Ascend 910B 系列（910B1/910B2/910B3/910B4） +- Ascend 910_93 系列（910_9362 ~ 910_9392） +- Ascend 310B 系列（310B1/310B2/310B3/310B4） +- Ascend 950 系列（910_950z ~ Ascend950DT_95A2） + +### 4.5 使用示例 + +```bash +bishengir-opt --hacc-append-device-spec="target=Ascend910B4" input.mlir +``` + +### 4.6 变换效果 + +该 Pass 会在模块上附加 `#hacc.target_device_spec` 属性，包含指定设备型号的硬件规格参数： + +```mlir +module attributes { + #dlti.target_device_spec = #hacc.target_device_spec< + #dlti.dl_entry<"AI_CORE_COUNT", 30 : i32>, + #dlti.dl_entry<"UB_SIZE", 196608 : i32>, + #dlti.dl_entry<"L1_SIZE", 1048576 : i32>, + ... + > +} { + ... +} +``` + +## 5. Pass 在编译流程中的位置 + +``` +前端 IR + │ + ├── hacc-append-device-spec ── 附加设备规格 + │ + ├── [HFusion 变换流程] + │ ├── hfusion-fuse-ops + │ ├── hfusion-auto-schedule + │ └── hfusion-auto-vectorize + │ + ├── hacc-rename-func ── 函数重命名 + │ + └── 后端 Lowering +``` + +- `hacc-append-device-spec` 通常在编译流程早期执行，为后续 Pass 提供设备规格信息 +- `hacc-rename-func` 通常在编译流程后期执行，确保函数名称符合运行时要求 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/00-overview.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/00-overview.md new file mode 100644 index 00000000..5a7b3276 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/00-overview.md @@ -0,0 +1,85 @@ +# HFusion 方言总览 + +## 1. 简介 + +HFusion（Hybrid Fusion，混合融合）方言是 AscendNPU-IR 中的张量级融合操作方言，提供丰富的张量操作原语，包括逐元操作、归约、矩阵乘、数据搬移、内存操作和特殊操作。HFusion 操作基于 Linalg 结构化操作范式，支持通过 OpDSL 声明式定义，并可通过变换 Pass 进行自动融合和向量化。 + +- **方言名称**：`hfusion` +- **C++ 命名空间**：`::mlir::hfusion` +- **方言具有 Canonicalizer**：是 + +## 2. 依赖方言 + +| 依赖方言 | 说明 | +|----------|------| +| `hacc::HACCDialect` | 异构计算调用 | +| `hmap::HMAPDialect` | 混合 Mesh 感知并行 | +| `linalg::LinalgDialect` | 结构化操作 | +| `mathExt::MathExtDialect` | 扩展数学操作 | +| `mesh::MeshDialect` | Mesh 并行 | +| `symbol::SymbolDialect` | 符号化形状 | + +## 3. 操作分类 + +| 类别 | 操作 | 数量 | +|------|------|------| +| 逐元操作 | elemwise_unary, elemwise_binary, compare, select, cast, bitcast | 6 | +| 归约操作 | reduce_with_index, cumsum, cumprod | 3 | +| 矩阵乘操作 | matmul_mx, group_matmul | 2 | +| 数据搬移 | load, store, broadcast, transpose, concat, pad, interleave, deinterleave, flip, gather, arange | 11 | +| 内存操作 | gather_load, scatter_store, indirect_load, indirect_store, gatherT, index_put, scatterT, atomic_cas, atomic_xchg | 9 | +| 特殊操作 | print, assert, barrier, sort, histogram, embedding_gather, is_inf, is_nan, is_finite, mulext, symbolic_dim | 11 | + +## 4. 源码位置 + +| 文件 | 路径 | 说明 | +|------|------|------| +| 方言基类 | [HFusionBase.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionBase.td) | 方言定义与属性枚举 | +| 枚举定义 | [HFusionEnums.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionEnums.td) | 函数属性枚举 | +| 属性定义 | [HFusionAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionAttrs.td) | 方言属性 | +| 操作定义 | [HFusionOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionOps.td) | 非结构化操作 | +| 结构化操作 | [HFusionStructuredOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionStructuredOps.td) | 结构化操作 | +| OpDSL 定义 | [HFusionNamedStructuredOps.yaml](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionNamedStructuredOps.yaml) | YAML 声明式操作定义 | +| 变换 Pass | [Passes.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/Transforms/Passes.td) | 变换 Pass 定义 | + +## 5. 类型系统 + +| 类型 | 定义 | 说明 | +|------|------|------| +| `TensorOrMemref` | `AnyTypeOf<[AnyMemRef, AnyRankedTensor]>` | 张量或 MemRef | +| `ShapedTypeOf` | 自定义 | 元素类型受限的 ShapedType | + +## 6. 操作基类 + +| 基类 | 说明 | +|------|------| +| `HFusion_Op` | 非结构化操作基类 | +| `HFusionStructuredBase_Op` | 结构化操作基类，实现 LinalgStructuredInterface 和 DestinationStyleOpInterface | + +## 7. 典型 IR 示例 + +```mlir +%result = hfusion.elemwise_unary ins(%input : tensor<128x256xf32>) + outs(%init : tensor<128x256xf32>) + unary_fn = + +%reduced = hfusion.reduce_with_index ins(%input, %init : tensor<128x256xf32>, tensor<128xf32>) + outs(%out_init, %idx_init : tensor<128xf32>, tensor<128xi64>) + reduce_with_index_kind = + +%matmul_result = hfusion.matmul_mx ins(%a, %b, %sa, %sb : tensor<16x32xf8E4M3FN>, tensor<32x64xf8E4M3FN>, tensor<2x2xui8>, tensor<2x4xui8>) + outs(%acc : tensor<16x64xf32>) -> tensor<16x64xf32> +``` + +## 8. 文档索引 + +| 文档 | 内容 | +|------|------| +| [01-elementwise-ops.md](01-elementwise-ops.md) | 逐元操作 | +| [02-reduction-ops.md](02-reduction-ops.md) | 归约操作 | +| [03-matmul-ops.md](03-matmul-ops.md) | 矩阵乘操作 | +| [04-data-movement-ops.md](04-data-movement-ops.md) | 数据搬移操作 | +| [05-memory-ops.md](05-memory-ops.md) | 内存操作 | +| [06-special-ops.md](06-special-ops.md) | 特殊操作 | +| [07-attributes-enums.md](07-attributes-enums.md) | 属性与枚举速查 | +| [08-transforms.md](08-transforms.md) | HFusion 变换 Pass 总览 | diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/01-elementwise-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/01-elementwise-ops.md new file mode 100644 index 00000000..4687dd31 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/01-elementwise-ops.md @@ -0,0 +1,208 @@ +# 逐元操作 + +## 1. 概述 + +HFusion 方言的逐元操作基于 Linalg 结构化操作范式，通过 OpDSL（YAML）声明式定义。每个操作包含一个 Region，描述逐元计算的标量逻辑。操作支持类型转换（cast）和多种函数选择。 + +> 源码参考：[HFusionNamedStructuredOps.yaml](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionNamedStructuredOps.yaml) + +## 2. 操作列表 + +### 2.1 load + +从输入张量逐元加载数据到输出张量，不执行数值类型转换。 + +| 操作数 | 类型 | 角色 | +|--------|------|------| +| `I` | input_tensor (type_var: T1) | 输入 | +| `O` | output_tensor (type_var: U) | 输出 | + +赋值：`O = I` + +### 2.2 store + +将输入张量逐元存储到输出张量，支持原子操作语义。 + +| 操作数 | 类型 | 角色 | +|--------|------|------| +| `I` | input_tensor (type_var: T1) | 输入 | +| `O` | output_tensor (type_var: U) | 输出 | +| `atomic_kind` | atomic_kind_attr (默认: NONE) | 原子操作类型 | + +赋值：`O = atomic_kind(I)` + +### 2.3 elemwise_unary + +对输入张量逐元应用一元函数。 + +| 操作数 | 类型 | 角色 | +|--------|------|------| +| `I` | input_tensor (type_var: T1) | 输入 | +| `O` | output_tensor (type_var: U) | 输出 | +| `fun` | unary_fn_attr (默认: sqrt) | 一元函数选择 | +| `cast` | type_fn_attr (默认: cast_signed) | 类型转换方式 | + +赋值：`O = fun(cast(I))` + +支持的 `unary_fn` 值： + +| 函数名 | 说明 | +|--------|------| +| `relu` | ReLU 激活 | +| `sqrt` | 平方根 | +| `rsqrt` | 平方根倒数 | +| `rec` | 倒数 | +| `vnot` | 按位取反 | +| `tanh` | 双曲正切 | +| `sin` | 正弦 | +| `cos` | 余弦 | +| `atan` | 反正切 | +| `tan` | 正切 | +| `absi` | 绝对值 | +| `erf` | 误差函数 | +| `log2` | 以 2 为底对数 | +| `log10` | 以 10 为底对数 | +| `log1p` | log(1+x) | +| `exp2` | 2^x | +| `expm1` | e^x - 1 | +| `ilogb` | 指数部分 | + +### 2.4 elemwise_binary + +对两个输入张量逐元应用二元函数。 + +| 操作数 | 类型 | 角色 | +|--------|------|------| +| `lhs` | input_tensor (type_var: T1) | 左操作数 | +| `rhs` | input_tensor (type_var: T2) | 右操作数 | +| `O` | output_tensor (type_var: U) | 输出 | +| `fun` | binary_fn_attr (默认: vand) | 二元函数选择 | +| `cast` | type_fn_attr (默认: cast_signed) | 类型转换方式 | + +赋值：`O = fun(cast(lhs), cast(rhs))` + +支持的 `binary_fn` 值： + +| 函数名 | 说明 | +|--------|------| +| `vor` | 按位或 | +| `vand` | 按位与 | +| `vxor` | 按位异或 | +| `minf` | 浮点最小值 | +| `maxf` | 浮点最大值 | +| `powf` | 浮点幂 | +| `mod` | 取模 | +| `shli` | 左移 | +| `shrsi` | 算术右移 | +| `shrui` | 逻辑右移 | +| `ldexp` | ldexp | +| `ceildivsi` | 有符号向上取整除法 | +| `ceildivui` | 无符号向上取整除法 | +| `floordivsi` | 有符号向下取整除法 | +| `powi` | 整数幂 | +| `minnumf` | IEEE 浮点最小值 | +| `maxnumf` | IEEE 浮点最大值 | +| `modui` | 无符号取模 | +| `divfhp` | 高精度除法 | + +### 2.5 compare + +对两个输入张量逐元执行比较操作。 + +| 操作数 | 类型 | 角色 | +|--------|------|------| +| `lhs` | input_tensor (type_var: T1) | 左操作数 | +| `rhs` | input_tensor (type_var: T1) | 右操作数 | +| `O` | output_tensor (type_var: U) | 输出 | +| `compare_fn` | compare_fn_attr (默认: veq) | 比较函数选择 | + +赋值：`O = compare_fn(lhs, rhs)` + +支持的 `compare_fn` 值： + +| 函数名 | 说明 | +|--------|------| +| `veq` | 等于 | +| `vne` | 不等于 | +| `vle` | 有符号小于等于 | +| `vlt` | 有符号小于 | +| `vge` | 有符号大于等于 | +| `vgt` | 有符号大于 | +| `vule` | 无符号小于等于 | +| `vuge` | 无符号大于等于 | +| `vugt` | 无符号大于 | +| `vult` | 无符号小于 | + +### 2.6 select + +根据条件张量逐元选择值。 + +| 操作数 | 类型 | 角色 | +|--------|------|------| +| `cond` | input_tensor (type_var: U) | 条件 | +| `lhs` | input_tensor (type_var: T1) | 真值 | +| `rhs` | input_tensor (type_var: T1) | 假值 | +| `O` | output_tensor (type_var: T1) | 输出 | + +赋值：`O = select(cond, lhs, rhs)` + +### 2.7 cast + +逐元类型转换，支持舍入模式控制。 + +| 操作数 | 类型 | 角色 | +|--------|------|------| +| `I` | input_tensor (type_var: T1) | 输入 | +| `O` | output_tensor (type_var: U) | 输出 | +| `round_mode` | round_mode_attr (默认: RINT) | 舍入模式 | +| `enable_overflow` | enable_overflow_attr (默认: true) | 溢出检测 | +| `enable_saturate` | enable_saturate_attr (默认: false) | 饱和截断 | +| `cast` | type_fn_attr (默认: cast_signed) | 类型转换方式 | +| `unsigned_mode` | unsigned_mode_attr (默认: SI2SI) | 无符号转换模式 | + +赋值：`O = round(cast(I), round_mode, unsigned_mode)` + +### 2.8 bitcast + +逐元位转换，不改变底层位模式。 + +| 操作数 | 类型 | 角色 | +|--------|------|------| +| `I` | input_tensor (type_var: T1) | 输入 | +| `O` | output_tensor (type_var: U) | 输出 | + +赋值：`O = bitcast(I)` + +## 3. MLIR 示例 + +### 3.1 逐元一元操作 + +```mlir +%result = hfusion.elemwise_unary ins(%input : tensor<128x256xf16>) + outs(%init : tensor<128x256xf32>) + unary_fn = cast = +``` + +### 3.2 逐元二元操作 + +```mlir +%result = hfusion.elemwise_binary ins(%lhs, %rhs : tensor<128x256xf16>, tensor<128x256xf16>) + outs(%init : tensor<128x256xf16>) + binary_fn = cast = +``` + +### 3.3 比较操作 + +```mlir +%result = hfusion.compare ins(%lhs, %rhs : tensor<128x256xf32>, tensor<128x256xf32>) + outs(%init : tensor<128x256xi1>) + compare_fn = +``` + +### 3.4 类型转换 + +```mlir +%result = hfusion.cast ins(%input : tensor<128x256xf32>) + outs(%init : tensor<128x256xi8>) + round_mode = cast = +``` diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/02-reduction-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/02-reduction-ops.md new file mode 100644 index 00000000..449424c3 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/02-reduction-ops.md @@ -0,0 +1,149 @@ +# 归约操作 + +## 1. 概述 + +HFusion 方言提供多种归约操作，包括带索引的归约（reduce_with_index）和累积操作（cumsum/cumprod）。归约操作沿指定维度对张量进行聚合计算。 + +> 源码参考：[HFusionStructuredOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionStructuredOps.td#L60-L142)、[HFusionOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionOps.td#L275-L323) + +## 2. reduce_with_index + +### 2.1 功能 + +使用 max/min 对 AnyShaped 执行归约操作，同时返回归约值和对应的索引。支持两种模式：(1) 接受输入和索引，产生归约值和索引；(2) 仅接受输入，产生归约值和索引。 + +### 2.2 操作签名 + +| 操作数/结果 | 类型 | 说明 | +|-------------|------|------| +| `inputs` | `Variadic` | 输入张量 | +| `inits` | `Variadic` | 初始值 | +| `reduce_kind` | `HFusion_ReduceWithIndexOpAttr` | 归约类型（min/max） | +| `unsigned_src` | `BoolAttr` | 源是否为无符号整数类型 | +| `tie_break_left` | `OptionalAttr` | 平局时取最左索引 | +| `dimensions` | `DenseI64ArrayAttr` | 归约维度 | +| `result` | `Variadic` | 归约结果 | + +### 2.3 归约类型 + +| ReduceWithIndexKind | 助记符 | 说明 | +|---------------------|--------|------| +| `MIN` | `min` | 最小值归约 | +| `MAX` | `max` | 最大值归约 | + +### 2.4 Traits + +- `AttrSizedOperandSegments` +- `ResultOnlyIfTensor` +- `DestinationStyleOpInterface` +- `LinalgStructuredInterface` +- `ReifyRankedShapedTypeOpInterface` + +### 2.5 MLIR 示例 + +```mlir +%val, %idx = hfusion.reduce_with_index + ins(%input, %val_init, %idx_init : tensor<128x256xf32>, tensor<128xf32>, tensor<128xi64>) + outs(%val_out, %idx_out : tensor<128xf32>, tensor<128xi64>) + reduce_with_index_kind = + unsigned_src = false + dimensions = [1] +``` + +## 3. cumsum + +### 3.1 功能 + +沿指定维度计算输入张量的累积和。 + +### 3.2 操作签名 + +| 操作数/结果 | 类型 | 说明 | +|-------------|------|------| +| `input` | `RankedTensorOf<[BF16, F16, F32, I8, I16, I32, I64, F8E4M3FN, F8E5M2]>` | 输入张量 | +| `cum_dims` | `DenseI64ArrayAttr` | 累积维度（严格递增排序） | +| `reverse` | `BoolAttr` | 是否反向累积 | +| `output` | `RankedTensorOf<[BF16, F16, F32, I8, I16, I32, I64, F8E4M3FN, F8E5M2]>` | 输出张量 | + +### 3.3 约束 + +- 当前仅支持单个累积维度 +- `cum_dims` 必须严格递增排序 + +### 3.4 MLIR 示例 + +```mlir +%result = hfusion.cumsum %input cum_dims = [1] reverse = false + : tensor<128x256xf32> -> tensor<128x256xf32> +``` + +## 4. cumprod + +### 4.1 功能 + +沿指定维度计算输入张量的累积积。 + +### 4.2 操作签名 + +| 操作数/结果 | 类型 | 说明 | +|-------------|------|------| +| `input` | `RankedTensorOf<[BF16, F16, F32, I8, I16, I32, I64]>` | 输入张量 | +| `cum_dims` | `DenseI64ArrayAttr` | 累积维度（严格递增排序） | +| `reverse` | `BoolAttr` | 是否反向累积 | +| `output` | `RankedTensorOf<[BF16, F16, F32, I8, I16, I32, I64]>` | 输出张量 | + +### 4.3 约束 + +- 当前仅支持单个累积维度 +- `cum_dims` 必须严格递增排序 + +### 4.4 MLIR 示例 + +```mlir +%result = hfusion.cumprod %input cum_dims = [0] reverse = true + : tensor<128x256xf32> -> tensor<128x256xf32> +``` + +## 5. CumOpType 枚举 + +| 枚举值 | 助记符 | 说明 | +|--------|--------|------| +| `UNDEFINED` | `undefined` | 未定义 | +| `CUMSUM` | `cumsum` | 累积和 | +| `CUMPROD` | `cumprod` | 累积积 | + +## 6. 语义说明 + +### 6.1 reduce_with_index 语义 + +对于 max 归约： +``` +result_value[i] = max(input[i][0], input[i][1], ..., input[i][N-1]) +result_index[i] = argmax(input[i][0], input[i][1], ..., input[i][N-1]) +``` + +当 `tie_break_left = true` 时，如果多个元素具有相同的最大值，返回最左边的索引。 + +### 6.2 cumsum 语义 + +正向累积（reverse = false）： +``` +output[i] = sum(input[0], input[1], ..., input[i]) +``` + +反向累积（reverse = true）： +``` +output[i] = sum(input[i], input[i+1], ..., input[N-1]) +``` + +### 6.3 cumprod 语义 + +正向累积（reverse = false）： +``` +output[i] = product(input[0], input[1], ..., input[i]) +``` + +反向累积（reverse = true）： +``` +output[i] = product(input[i], input[i+1], ..., input[N-1]) +``` diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/03-matmul-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/03-matmul-ops.md new file mode 100644 index 00000000..f162e622 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/03-matmul-ops.md @@ -0,0 +1,109 @@ +# 矩阵乘操作 + +## 1. 概述 + +HFusion 方言提供矩阵乘操作，包括微缩放格式矩阵乘（matmul_mx）和分组矩阵乘（group_matmul）。这些操作是 NPU Cube 计算单元的核心负载。 + +> 源码参考：[HFusionOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionOps.td#L964-L1021)、[HFusionNamedStructuredOps.yaml](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionNamedStructuredOps.yaml#L386-L424) + +## 2. matmul_mx + +### 2.1 功能 + +执行微缩放（Microscaling）格式的矩阵乘法，输入通过缩放因子隐式缩放。常用于 FP8/FP4 等量化数据类型。计算公式：`C = (A * scale_a) dot (B * scale_b)`。 + +微缩放格式遵循 OCP Microscaling Formats (MX) 规范：https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf + +### 2.2 操作签名 + +| 操作数/结果 | 类型 | 说明 | +|-------------|------|------| +| `inputA` | `ShapedTypeOf<[F8E4M3FN, F8E5M2]>` | 左矩阵 | +| `inputB` | `ShapedTypeOf<[F8E4M3FN, F8E5M2]>` | 右矩阵 | +| `scaleA` | `ShapedTypeOf<[UI8, I8]>` | 左矩阵缩放因子 | +| `scaleB` | `ShapedTypeOf<[UI8, I8]>` | 右矩阵缩放因子 | +| `acc` | `ShapedTypeOf<[AnyFloat]>` | 累加器（DPS init） | +| `result` | `ShapedTypeOf<[AnyFloat]>` | 输出结果 | + +### 2.3 Traits + +- `Pure` +- `DestinationStyleOpInterface` +- `BiShengIRAggregatedOpInterface`（支持 decomposeOperation） + +### 2.4 MLIR 示例 + +```mlir +%result = hfusion.matmul_mx + ins(%a, %b, %sa, %sb : + tensor<16x32xf8E4M3FN>, tensor<32x64xf8E4M3FN>, + tensor<2x2xui8>, tensor<2x4xui8>) + outs(%acc : tensor<16x64xf32>) + -> tensor<16x64xf32> +``` + +### 2.5 Builder + +```tablegen +OpBuilder<(ins "Value":$inputA, "Value":$inputB, + "Value":$scaleA, "Value":$scaleB, "Value":$acc)> +``` + +## 3. group_matmul + +### 3.1 功能 + +执行分组矩阵乘法，用于 MoE（Mixture of Experts）等场景。每个 Expert 的权重矩阵与其分配的 Token 进行矩阵乘。 + +### 3.2 操作签名 + +| 操作数/结果 | 类型 | 角色 | +|-------------|------|------| +| `w1` | input_tensor (type_var: T1) | Expert 权重矩阵，shape_map: `(d0, d1, d2) -> (d2, d1, d0)` | +| `tokens` | input_tensor (type_var: T2) | Token 嵌入，shape_map: `(d0, d1, d2) -> (d1, d2)` | +| `tokens_per_expert` | input_tensor (type_var: T3) | 每 Expert 的 Token 数，shape_map: `(d0, d1, d2) -> (1)` | +| `output` | output_tensor (type_var: T4) | 输出，shape_map: `(d0, d1, d2) -> (d0, d0)` | + +### 3.3 Indexing Maps + +```mlir +affine_map<(d0, d1, d2) -> (d2, d1, d0)> // w1 +affine_map<(d0, d1) -> (d0, d1)> // tokens +affine_map<(d0) -> (0)> // tokens_per_expert +affine_map<(d0, d1, d2) -> (d0, d0)> // output +``` + +### 3.4 Iterator Types + +``` +["parallel", "parallel", "reduction"] +``` + +### 3.5 MLIR 示例 + +```mlir +%result = hfusion.group_matmul + ins(%w1, %tokens, %tokens_per_expert : + tensor, tensor, tensor<1xi64>) + outs(%output : tensor) +``` + +## 4. MmMapMode 枚举 + +| 枚举值 | 助记符 | 说明 | +|--------|--------|------| +| `CoreOp` | `core_op` | 核心操作模式 | +| `MacroInstr` | `macro_instr` | 宏指令模式 | + +## 5. 数据格式说明 + +### 5.1 FP8 格式 + +| 类型 | 说明 | +|------|------| +| `F8E4M3FN` | 4 位指数、3 位尾数，无无穷，用于前向传播 | +| `F8E5M2` | 5 位指数、2 位尾数，支持无穷，用于反向传播 | + +### 5.2 微缩放格式 + +微缩放格式将一组元素共享一个缩放因子，典型配置为每 32 个元素共享 1 个 UI8 缩放因子。这种格式在保持精度的同时大幅减少了存储和计算开销。 diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/04-data-movement-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/04-data-movement-ops.md new file mode 100644 index 00000000..b7705053 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/04-data-movement-ops.md @@ -0,0 +1,159 @@ +# 数据搬移操作 + +## 1. 概述 + +HFusion 方言提供丰富的数据搬移操作，包括广播、转置、拼接、填充、交织、翻转和聚集等。这些操作用于在不同形状和布局的张量间进行数据重排。 + +> 源码参考：[HFusionOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionOps.td#L128-L213)、[HFusionStructuredOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionStructuredOps.td#L196-L260) + +## 2. interleave + +### 2.1 功能 + +将 n 个输入张量沿最后一维交织，构造一个输出张量。当前仅支持 n=2。 + +### 2.2 操作签名 + +| 操作数/结果 | 类型 | 说明 | +|-------------|------|------| +| `input` | `Variadic` | 输入张量（2 个，形状相同） | +| `output` | `AnyRankedTensor` | 输出张量 | + +### 2.3 Traits + +- `Pure`, `Commutative` +- `SameOperandsAndResultRank` +- `ReifyRankedShapedTypeOpInterface` + +### 2.4 MLIR 示例 + +```mlir +%output = hfusion.interleave %a, %b : tensor<128x256xf32>, tensor<128x256xf32> -> tensor<128x512xf32> +``` + +## 3. deinterleave + +### 3.1 功能 + +将一个输入张量沿最后一维反交织为两个张量。偶数索引元素进入第一个输出，奇数索引元素进入第二个输出。 + +### 3.2 操作签名 + +| 操作数/结果 | 类型 | 说明 | +|-------------|------|------| +| `input` | `AnyRankedTensor` | 输入张量 | +| `channelIndex` | `I64Attr` | 通道选择：-1=全部, 0=偶数, 1=奇数 | +| `output` | `Variadic` | 输出张量 | + +### 3.3 channelIndex 行为 + +| 值 | 行为 | +|----|------| +| -1 | 输出两个张量（偶数索引 + 奇数索引） | +| 0 | 仅输出偶数索引通道 | +| 1 | 仅输出奇数索引通道 | + +### 3.4 约束 + +- 输入张量最后一维大小必须是 2 的倍数 + +### 3.5 MLIR 示例 + +```mlir +%even, %odd = hfusion.deinterleave %input channel_index = -1 + : tensor<128x512xf32> -> tensor<128x256xf32>, tensor<128x256xf32> + +%even_only = hfusion.deinterleave %input channel_index = 0 + : tensor<128x512xf32> -> tensor<128x256xf32> +``` + +## 4. flip + +### 4.1 功能 + +沿指定维度翻转张量。当前仅支持最后一维。 + +### 4.2 操作签名 + +| 操作数/结果 | 类型 | 说明 | +|-------------|------|------| +| `input` | `AnyRankedTensor` | 输入张量 | +| `flip_axis` | `I64Attr` | 翻转轴 | +| `output` | `AnyRankedTensor` | 输出张量 | + +### 4.3 MLIR 示例 + +```mlir +%output = hfusion.flip %input flip_axis = 1 + : tensor<128x256xf32> -> tensor<128x256xf32> +``` + +## 5. gather + +### 5.1 功能 + +沿指定轴从源张量中聚集元素。对应 `triton.language.gather` 语义。 + +### 5.2 操作签名 + +| 操作数/结果 | 类型 | 说明 | +|-------------|------|------| +| `src` | `AnyShaped` | 源张量 | +| `index` | `AnyShaped` | 索引张量 | +| `init` | `AnyShaped` | 初始值 | +| `axis` | `I64Attr` | 聚集轴 | +| `result` | `Variadic` | 输出 | + +### 5.3 Traits + +- `AllRanksMatch<["src", "index", "init"]>` +- `AllShapesMatch<["index", "init"]>` +- `AllElementTypesMatch<["src", "init"]>` +- `BiShengIRAggregatedOpInterface`（支持 decomposeOperation） + +### 5.4 语义 + +给定 src:tensor<16x16> 和 index:tensor<16x4>，axis=1： +``` +for i in 0 to 16: + for j in 0 to 4: + for k in 0 to 16: + output[i][j] = (index[i][j] == k) ? src[i][k] : output[i][j] +``` + +### 5.5 MLIR 示例 + +```mlir +%result = hfusion.gather ins(%src, %index, %init : tensor<16x16xf32>, tensor<16x4xi32>, tensor<16x4xf32>) + axis = 1 -> tensor<16x4xf32> +``` + +## 6. arange + +### 6.1 功能 + +生成等差数列张量，支持偏移和多维步长。 + +### 6.2 操作签名 + +| 操作数/结果 | 类型 | 说明 | +|-------------|------|------| +| `offset` | `Optional` | 偏移量（默认 0） | +| `strides` | `Variadic` | 各维步长 | +| `init` | `AnyShaped` | 初始张量（决定输出形状） | +| `result_tensor` | `Optional` | 输出张量 | + +### 6.3 语义 + +3D arange：`arange[i, j, k] = offset + stride[0] * i + stride[1] * j + stride[2] * k` + +## 7. 其他数据搬移操作 + +以下操作由 Linalg 命名操作或 HFusion 结构化操作覆盖，通过 OpDSL 或 Linalg 通用操作实现： + +| 操作 | 说明 | 实现方式 | +|------|------|----------| +| `broadcast` | 广播张量到目标形状 | Linalg 广播语义 | +| `transpose` | 转置张量维度 | Linalg 转置语义 | +| `concat` | 沿指定维度拼接张量 | Linalg 拼接语义 | +| `pad` | 填充张量边界 | Linalg 填充语义 | diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/05-memory-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/05-memory-ops.md new file mode 100644 index 00000000..ccd07444 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/05-memory-ops.md @@ -0,0 +1,219 @@ +# 内存操作 + +## 1. 概述 + +HFusion 方言提供丰富的内存操作，包括稀疏加载/存储、间接访存、聚集/散射和原子操作。这些操作支持 GM（Global Memory）与 UB（Unified Buffer）之间的数据搬移，以及原子同步语义。 + +> 源码参考：[HFusionOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionOps.td#L329-L928) + +## 2. gather_load + +### 2.1 功能 + +从源内存缓冲区按偏移量聚集加载元素到输出张量，支持掩码和回退值。 + +### 2.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `base` | `AnyMemRef` | 源内存缓冲区 | +| `indices` | `RankedTensorOf<[I32, I64]>` | 偏移量张量 | +| `burst_len` | `AnyTypeOf<[I32, I64]>` | 突发长度 | +| `mask` | `Optional>` | 掩码 | +| `other` | `Optional>` | 掩码位置的回退值 | +| `boundaryCheck` | `OptionalAttr` | 边界检查 | +| `padding` | `OptionalAttr` | 填充选项 | +| `cache` | `OptionalAttr` | 缓存修饰符 | +| `evict` | `OptionalAttr` | 驱逐策略 | +| `isVolatile` | `OptionalAttr` | 是否 volatile | +| `result` | `AnyRankedTensor` | 输出张量 | + +### 2.3 MLIR 示例 + +```mlir +%result = hfusion.gather_load + ins(%base, %indices, %burst_len, %mask, %other : + memref, tensor<128xi32>, i32, tensor<128xi1>, f32) + -> tensor<128xf32> +``` + +## 3. scatter_store + +### 3.1 功能 + +将源张量中的元素按偏移量散射存储到目标内存缓冲区，支持掩码。 + +### 3.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `base` | `AnyMemRef` | 目标内存缓冲区 | +| `indices` | `RankedTensorOf<[I32, I64]>` | 偏移量张量 | +| `data` | `AnyRankedTensor` | 源数据张量 | +| `burst_len` | `AnyTypeOf<[I32, I64]>` | 突发长度 | +| `mask` | `Optional>` | 掩码 | +| `boundaryCheck` | `OptionalAttr` | 边界检查 | +| `cache` | `OptionalAttr` | 缓存修饰符 | +| `evict` | `OptionalAttr` | 驱逐策略 | + +### 3.3 MLIR 示例 + +```mlir +hfusion.scatter_store + ins(%base, %indices, %data, %burst_len, %mask : + memref, tensor<128xi32>, tensor<128xf32>, i32, tensor<128xi1>) +``` + +## 4. indirect_load + +### 4.1 功能 + +间接内存加载，支持 1D-5D 掩码和回退值。 + +### 4.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `src` | `AnyMemRef` | 源内存缓冲区 | +| `offsets` | `TensorOf<[I32, I64]>` | 偏移量张量 | +| `dst` | `AnyRankedTensor` | 目标张量（决定输出形状和类型） | +| `mask` | `TensorOf<[I1, I8]>` | 掩码 | +| `other` | `TensorOf<[AnyInteger, AnyFloat]>` | 回退值 | +| `result` | `Optional` | 输出 | + +### 4.3 语义 + +``` +dst[i] = mask[i] ? src[offsets[i]] : other[i] +``` + +## 5. indirect_store + +### 5.1 功能 + +间接内存存储，支持 1D-5D 掩码，使用 SIMT 模板。 + +### 5.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `dst` | `AnyMemRef` | 目标内存缓冲区 | +| `offsets` | `TensorOf<[I32, I64]>` | 偏移量张量 | +| `src` | `AnyRankedTensor` | 源数据张量 | +| `mask` | `Optional>` | 掩码 | + +### 5.3 语义 + +``` +if (mask[i]) dst[offsets[i]] = src[i] +``` + +## 6. gatherT + +### 6.1 功能 + +SIMT 模式下的 Gather 操作，沿指定轴从源 GM 缓冲区按索引张量聚集元素。 + +### 6.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `src` | `AnyMemRef` | 源 GM 缓冲区 | +| `index` | `AnyRankedTensor` | 索引张量 | +| `dst` | `AnyRankedTensor` | 目标张量 | +| `bound` | `AnyTypeOf<[I32, I64]>` | gather 维大小 | +| `dim` | `AnyTypeOf<[I32, I64]>` | gather 维度 | +| `src_stride` | `Variadic>` | 源张量步长 | +| `index_shape` | `Variadic>` | 索引张量形状 | +| `offsets` | `Variadic>` | 偏移量 | + +## 7. index_put + +### 7.1 功能 + +SIMT 模式下的 IndexPut 操作，按索引将值写入目标张量的指定位置。 + +### 7.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `dst` | `AnyMemRef` | 目标 GM 缓冲区 | +| `index` | `AnyRankedTensor` | 索引张量 | +| `value` | `AnyRankedTensor` | 值张量 | +| `scatter_dim` | `AnyTypeOf<[I32, I64]>` | 散射维度 | +| `bound` | `AnyTypeOf<[I32, I64]>` | 索引上界 | +| `end_offset` | `Variadic>` | 结束偏移 | +| `start_offset` | `Variadic>` | 起始偏移 | +| `dst_stride` | `Variadic>` | 目标步长 | + +## 8. scatterT + +### 8.1 功能 + +SIMT 模式下的 Scatter 操作，按索引将值写入目标张量的指定位置。 + +### 8.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `dst` | `AnyMemRef` | 目标 GM 缓冲区 | +| `value` | `AnyRankedTensor` | 值张量 | +| `index_tile` | `AnyRankedTensor` | 索引张量 | +| `index_boundary` | `AnyTypeOf<[I32, I64]>` | 索引上界 | +| `dim` | `AnyTypeOf<[I32, I64]>` | 维度 | +| `dst_stride` | `Variadic>` | 目标步长 | +| `index_shape` | `Variadic>` | 索引形状 | +| `offsets` | `Variadic>` | 偏移量 | + +## 9. atomic_cas + +### 9.1 功能 + +原子比较并交换（Compare-And-Swap）操作。 + +### 9.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `input` | `Variadic` | 期望旧值(src0)和新值(src1) | +| `dst` | `TensorOrMemref` | 内存位置 | +| `output` | `Variadic` | 原始值 | + +### 9.3 语义 + +``` +if (V == A) V = B; return old_V; +``` + +### 9.4 MLIR 示例 + +```mlir +hfusion.atomic_cas ins(%src0, %src1 : memref, memref) + outs(%dst : memref) + +%result = hfusion.atomic_cas ins(%src0, %src1 : tensor, tensor) + outs(%dst : tensor) -> tensor +``` + +## 10. atomic_xchg + +### 10.1 功能 + +原子交换操作，将新值写入内存地址并返回旧值。 + +### 10.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `input` | `Variadic` | 新值(src) | +| `dst` | `TensorOrMemref` | 内存位置 | +| `output` | `Variadic` | 旧值 | + +### 10.3 MLIR 示例 + +```mlir +hfusion.atomic_xchg ins(%src : memref) outs(%dst : memref) + +%result = hfusion.atomic_xchg ins(%src : tensor) + outs(%dst : tensor) -> tensor +``` diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/06-special-ops.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/06-special-ops.md new file mode 100644 index 00000000..740e113e --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/06-special-ops.md @@ -0,0 +1,254 @@ +# 特殊操作 + +## 1. 概述 + +HFusion 方言的特殊操作包括调试操作、排序、直方图、嵌入聚集、数值检查和扩展乘法等。 + +> 源码参考：[HFusionOps.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionOps.td#L48-L269) + +## 2. print + +### 2.1 功能 + +Device 端打印操作，用于调试。 + +### 2.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `prefix` | `StrAttr` | 前缀字符串 | +| `hex` | `BoolAttr` | 是否以十六进制打印 | +| `arg` | `AnyTypeOf<[AnyInteger, AnyFloat, AnyRankedTensor]>` | 待打印的值 | + +### 2.3 MLIR 示例 + +```mlir +hfusion.print "value:" hex = true %val : tensor<128xf32> +``` + +## 3. assert + +### 3.1 功能 + +Device 端断言操作，用于调试。 + +### 3.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `msg` | `StrAttr` | 断言消息 | +| `cond` | `AnyTypeOf<[AnyInteger, AnyRankedTensor]>` | 条件值 | + +### 3.3 MLIR 示例 + +```mlir +hfusion.assert "index out of range" %cond : i32 +``` + +## 4. barrier + +### 4.1 功能 + +同步一个 Core 上所有 Pipeline 的执行。 + +### 4.2 操作签名 + +无操作数，无结果。 + +### 4.3 MLIR 示例 + +```mlir +hfusion.barrier +``` + +## 5. sort + +### 5.1 功能 + +沿指定轴对张量排序，输出排序后的值和对应索引。 + +### 5.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `src` | `TensorOrMemref` | 待排序张量 | +| `descending` | `BoolAttr` (默认: false) | 是否降序 | +| `sort_axis` | `I64Attr` | 排序轴 | +| `result` | `Variadic` | 排序结果 | + +### 5.3 约束 + +- 输入和输出必须具有相同的 rank +- 当前仅支持尾轴排序 + +### 5.4 MLIR 示例 + +```mlir +%result = hfusion.sort ins(%src : tensor) + descending = true sort_axis = 0 -> tensor +``` + +## 6. histogram + +### 6.1 功能 + +计算整数张量的直方图，支持可选掩码。 + +### 6.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `input` | `RankedTensorOf<[I8, UI8, I16, UI16, I32, UI32, I64, UI64]>` | 输入张量 | +| `num_bins` | `I64Attr` | 直方图桶数 | +| `mask` | `Optional>` | 掩码 | +| `output` | `RankedTensorOf<[I8, I16, I32, I64]>` | 输出直方图 | + +### 6.3 语义 + +对输入张量的每个元素，递增输出直方图中对应的桶。如果提供掩码，仅统计 mask[i]=true 的元素。输出必须为 1D 张量，长度等于 num_bins。 + +### 6.4 MLIR 示例 + +```mlir +%hist = hfusion.histogram %input, 256, %mask : tensor<1024xi32>, tensor<1024xi1> -> tensor<256xi64> +``` + +## 7. embedding_gather + +### 7.1 功能 + +使用 Gather 语义执行嵌入查找操作。 + +### 7.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `src` | `AnyMemRef` | 2D 嵌入表（GM） | +| `index` | `AnyRankedTensor` | 1D/2D 索引张量 | +| `dst` | `AnyRankedTensor` | 目标张量 | +| `bound` | `AnyTypeOf<[I32, I64]>` | 词汇表大小（边界检查） | +| `offsets` | `Variadic>` | 偏移量 | +| `numels` | `Variadic>` | 元素数量 | +| `result` | `Optional` | 输出 | + +### 7.3 语义 + +``` +result[b][i][d] = src[index[b][i]][d] +``` + +### 7.4 MLIR 示例 + +```mlir +%result = hfusion.embedding_gather + ins(%src, %index, %bound, [], [] : + memref<10000x768xf16>, tensor<128xi32>, i32) + outs(%dst : tensor<128x768xf16>) + -> tensor<128x768xf16> +``` + +## 8. is_inf + +### 8.1 功能 + +判断浮点张量元素是否为正无穷或负无穷。 + +### 8.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `input` | `RankedTensorOf<[BF16, F16, F32]>` | 输入张量 | +| `output` | `RankedTensorOf<[I1]>` | 输出布尔张量 | + +### 8.3 MLIR 示例 + +```mlir +%result = hfusion.isinf %input : tensor<128xf32> -> tensor<128xi1> +``` + +## 9. is_nan + +### 9.1 功能 + +判断浮点张量元素是否为 NaN。 + +### 9.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `input` | `RankedTensorOf<[BF16, F16, F32]>` | 输入张量 | +| `output` | `RankedTensorOf<[I1]>` | 输出布尔张量 | + +### 9.3 MLIR 示例 + +```mlir +%result = hfusion.isnan %input : tensor<128xf32> -> tensor<128xi1> +``` + +## 10. is_finite + +### 10.1 功能 + +判断浮点张量元素是否为有限值（非 NaN 且非无穷）。 + +### 10.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `input` | `RankedTensorOf<[BF16, F16, F32]>` | 输入张量 | +| `output` | `RankedTensorOf<[I1]>` | 输出布尔张量 | + +### 10.3 Traits + +- `BiShengIRAggregatedOpInterface`（支持 decomposeOperation） + +### 10.4 MLIR 示例 + +```mlir +%result = hfusion.isfinite %input : tensor<128xf32> -> tensor<128xi1> +``` + +## 11. mulext + +### 11.1 功能 + +扩展有符号整数乘法，返回 2N 位乘积的低半部分和高半部分。 + +### 11.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `lhs` | `SignlessIntegerLike` | 左操作数 | +| `rhs` | `SignlessIntegerLike` | 右操作数 | +| `low` | `SignlessIntegerLike` | 乘积低半部分 | +| `high` | `SignlessIntegerLike` | 乘积高半部分 | + +### 11.3 Traits + +- `Pure`, `Commutative` +- `AllTypesMatch<["lhs", "rhs", "low", "high"]>` + +### 11.4 MLIR 示例 + +```mlir +%low, %high = hfusion.mulext %a, %b : i32 +``` + +## 12. symbolic_dim + +### 12.1 功能 + +引用符号维度并返回 index 类型值。 + +### 12.2 操作签名 + +| 操作数 | 类型 | 说明 | +|--------|------|------| +| `symbolName` | `SymbolRefAttr` | 符号名称 | +| `result` | `Index` | 结果 | + +### 12.3 MLIR 示例 + +```mlir +%0 = hfusion.symbolic_dim @SymName : index +``` diff --git a/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/07-attributes-enums.md b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/07-attributes-enums.md new file mode 100644 index 00000000..9cef3518 --- /dev/null +++ b/skills/triton/latency-optimizer/references/docs_triton_IR/docs_ascendnpu_ir/03-HFusion-Dialect/07-attributes-enums.md @@ -0,0 +1,246 @@ +# 属性与枚举速查 + +## 1. 概述 + +HFusion 方言定义了丰富的枚举和属性，用于控制逐元操作的函数选择、类型转换、舍入模式、原子操作类型等。本文档提供完整的速查表。 + +> 源码参考：[HFusionEnums.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionEnums.td)、[HFusionBase.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionBase.td)、[HFusionAttrs.td](file:///d:/项目/trae/triton_a5/AscendNPU-IR/bishengir/include/bishengir/Dialect/HFusion/IR/HFusionAttrs.td) + +## 2. UnaryFn（一元函数枚举） + +| 枚举值 | 整数值 | 说明 | +|--------|--------|------| +| `relu` | 0 | ReLU 激活 | +| `sqrt` | 1 | 平方根 | +| `rsqrt` | 2 | 平方根倒数 | +| `rec` | 3 | 倒数 | +| `vnot` | 4 | 按位取反 | +| `tanh` | 5 | 双曲正切 | +| `sin` | 6 | 正弦 | +| `cos` | 7 | 余弦 | +| `atan` | 8 | 反正切 | +| `tan` | 9 | 正切 | +| `absi` | 10 | 绝对值 | +| `erf` | 11 | 误差函数 | +| `log2` | 12 | 以 2 为底对数 | +| `log10` | 13 | 以 10 为底对数 | +| `log1p` | 14 | log(1+x) | +| `exp2` | 15 | 2^x | +| `expm1` | 16 | e^x - 1 | +| `ilogb` | 17 | 指数部分 | + +属性语法：`unary_fn = ` + +## 3. BinaryFn（二元函数枚举） + +| 枚举值 | 整数值 | 说明 | +|--------|--------|------| +| `vor` | 0 | 按位或 | +| `vand` | 1 | 按位与 | +| `vxor` | 2 | 按位异或 | +| `minf` | 3 | 浮点最小值 | +| `maxf` | 4 | 浮点最大值 | +| `powf` | 5 | 浮点幂 | +| `mod` | 6 | 取模 | +| `shli` | 7 | 左移 | +| `shrsi` | 8 | 算术右移 | +| `shrui` | 9 | 逻辑右移 | +| `ldexp` | 10 | ldexp | +| `ceildivsi` | 11 | 有符号向上取整除法 | +| `ceildivui` | 12 | 无符号向上取整除法 | +| `floordivsi` | 13 | 有符号向下取整除法 | +| `powi` | 14 | 整数幂 | +| `minnumf` | 15 | IEEE 浮点最小值 | +| `maxnumf` | 16 | IEEE 浮点最大值 | +| `modui` | 17 | 无符号取模 | +| `divfhp` | 18 | 高精度除法 | + +属性语法：`binary_fn = ` + +## 4. CompareFn（比较函数枚举） + +| 枚举值 | 整数值 | 说明 | +|--------|--------|------| +| `veq` | 0 | 等于 | +| `vne` | 1 | 不等于 | +| `vle` | 2 | 有符号小于等于 | +| `vlt` | 3 | 有符号小于 | +| `vge` | 4 | 有符号大于等于 | +| `vgt` | 5 | 有符号大于 | +| `vule` | 6 | 无符号小于等于 | +| `vuge` | 7 | 无符号大于等于 | +| `vugt` | 8 | 无符号大于 | +| `vult` | 9 | 无符号小于 | + +属性语法：`compare_fn = ` + +## 5. TernaryFn（三元函数枚举） + +| 枚举值 | 整数值 | 说明 | +|--------|--------|------| +| `select` | 0 | 条件选择 | + +属性语法：`ternary_fn =