新增生成 A5 agent

README.en.md

@@ -21,6 +21,7 @@
       - [**3.2 AscendC**](#32-ascendc)
       - [Scenario 1: Single Operator Generation (Lingxi-code Agent)](#scenario-1-single-operator-generation-lingxi-code-agent)
       - [Scenario 2: Batch Benchmark Evaluation (Ascend-Benchmark-Evaluator)](#scenario-2-batch-benchmark-evaluation-ascend-benchmark-evaluator)
+      - [Scenario 3: A5 Single Operator Generation (Preview) (ascendc-a5-op-gen Skill)](#scenario-3-a5-single-operator-generation-ascendc-a5-op-gen-skill)
     - [Evaluation Baseline](#evaluation-baseline)
       - [Triton](#triton)
       - [AscendC](#ascendc)
@@ -35,6 +36,7 @@
 | **Triton** | **Benchmark-Evaluator** | One-click batch evaluation | Execute specified Benchmark evaluation, automatically summarize and generate detailed reports |
 | **AscendC** | **Lingxi_code Agent** | AscendC single operator interactive generation | Code generation → Evaluation & Verification (Accuracy alignment & Performance testing) |
 | **AscendC** | **Ascend-Benchmark-Evaluator** | AscendC operator one-click batch evaluation | Execute specified Benchmark evaluation, automatically summarize and generate detailed reports |
+| **AscendC (A5)** | **ascendc-a5-op-gen** | A5 single operator generation and optimization (Preview) | Supports Benchmark / Op-gen / Optimize modes, automatically generates and optimizes AscendC SIMT/SIMD kernels |
 
 > **Shared Kernel**: AKG-Triton Agent and Benchmark-Evaluator share the underlying code generation Agent, uniformly handling the core workflow of "Code Generation → Verification → Performance Testing" to ensure consistency and high reusability of the generation logic.
 
@@ -223,6 +225,61 @@ bash utils/run_benchmark_ascendc.sh \
 - `--npu-list`: Multi-NPU list, comma-separated, e.g., `0,1,2,3,4,5` (mutually exclusive with `--npu`, higher priority)
 - `--output`: Output directory (required)
 
+---
+
+#### Scenario 3: A5 Single Operator Generation (Preview) (ascendc-a5-op-gen Skill)
+
+> **Note**: The A5 generation capability is currently in preview / testing phase. Generation quality and success rate are limited, and complex operators may require manual intervention.
+
+Suitable for quickly generating, verifying, and optimizing a single AscendC operator on A5 architecture. Supports three modes:
+- **Benchmark mode**: Automatically generate a kernel from an NPUKernelBench source file
+- **Op-gen mode**: Generate an AscendC implementation from any CUDA/PyTorch source file
+- **Optimize mode**: Optimize performance of an already archived kernel
+
+**Steps**:
+
+1. Configure the A5 Agent and skills in the AscendOpGenAgent directory:
+```bash
+mkdir -p .claude
+mkdir -p .claude/skills
+cp agents/ascend-a5-kernel-developer.md .claude/CLAUDE.md
+cp -r skills/ascendc/* .claude/skills/
+```
+
+2. Start Claude and enter the Skill command:
+
+**Benchmark mode** (generate from NPUKernelBench):
+```text
+/ascendc-a5-op-gen 13_Cat
+```
+
+**Op-gen mode** (generate from source file):
+```text
+/ascendc-a5-op-gen path/to/source.py
+```
+
+**Optimize mode** (optimize existing operator):
+```text
+/ascendc-a5-op-gen --optimize 12_Permute
+```
+
+**Execution Flow**:
+1. Automatically detect mode and create output directory (e.g., `output/npukernelbench/13_Cat`)
+2. Generate AscendC kernel files (`.h`, `.cpp`, `pybind11.cpp`, `model_new_ascendc.py`)
+3. Static check + compilation (up to 5 iterative fix rounds)
+4. Accuracy verification (all cases must PASS)
+5. Performance test and compute ratio (auto-enters optimization if below 0.6x)
+6. Output final report and `PROGRESS.md`
+
+After operator generation, you can use `a5-perf-summary` for unified performance summary and archiving:
+
+```text
+/a5-perf-summary 1_GELU    # single operator summary
+/a5-perf-summary all       # batch summary for all completed operators
+```
+
+This step automatically verifies accuracy, detects the kernel programming mode (SIMT / SIMD), computes performance ratios, and updates `benchmarks/NPUKernelBench/A5_RESULTS.md`.
+
 ### Evaluation Baseline
 
 #### Triton

README.md

@@ -21,6 +21,7 @@
       - [**3.2 AscendC**](#32-ascendc)
       - [场景一：单算子生成 (Lingxi-code Agent)](#场景一单算子生成-lingxi-code-agent)
       - [场景二：Benchmark 批量评测 (Ascend-Benchmark-Evaluator)](#场景二benchmark-批量评测-ascend-benchmark-evaluator)
+      - [场景三：A5 单算子生成（预览版）(ascendc-a5-op-gen Skill)](#场景三a5-单算子生成-ascendc-a5-op-gen-skill)
     - [评测基线](#评测基线)
       - [Triton](#triton)
       - [AscendC](#ascendc)
@@ -35,6 +36,7 @@
 | **Triton**  | **Benchmark-Evaluator** | 一键批量评测 | 执行指定 Benchmark 评测，自动总结并生成详细报告 |
 | **AscendC** | **Lingxi_code Agent** | AscendC 单算子交互式生成 | 代码生成 → 评测验证（精度对齐与性能测试） |
 | **AscendC** | **Ascend-Benchmark-Evaluator** | AscendC 算子一键批量评测 | 执行指定 Benchmark 评测，自动总结并生成详细报告 |
+| **AscendC (A5)** | **ascendc-a5-op-gen** | A5 单算子生成与优化（预览版） | 支持 Benchmark / Op-gen / Optimize 三种模式，自动生成并优化 AscendC SIMT/SIMD kernel |
 
 >  **共享内核**：AKG-Triton Agent、Benchmark-Evaluator两者底层共用代码生成 Agent，统一处理“代码生成 → 验证 → 性能测试”的核心工作流，确保生成逻辑的一致性与高复用性。
 
@@ -223,6 +225,61 @@ bash utils/run_benchmark_ascendc.sh \
 - `--npu-list`: 多 NPU 列表，逗号分隔，如 `0,1,2,3,4,5`（与 `--npu` 互斥，优先级更高）
 - `--output`: 输出目录（必填）
 
+---
+
+#### 场景三：A5 单算子生成（预览版）(ascendc-a5-op-gen Skill)
+
+> **注意**：A5 生成能力目前处于预览/测试阶段，生成质量和成功率有限，遇到复杂算子时可能需要人工介入调整。
+
+适用于在 A5 架构上快速生成、验证并优化单个 AscendC 算子。支持三种模式：
+- **Benchmark 模式**：从 NPUKernelBench 源文件自动生成 kernel
+- **Op-gen 模式**：从任意 CUDA/PyTorch 源文件生成 AscendC 实现
+- **Optimize 模式**：对已有归档 kernel 进行性能优化
+
+**操作步骤**：
+
+1. 在 AscendOpGenAgent 目录下配置 A5 Agent 和 skills：
+```bash
+mkdir -p .claude
+mkdir -p .claude/skills
+cp agents/ascend-a5-kernel-developer.md .claude/CLAUDE.md
+cp -r skills/ascendc/* .claude/skills/
+```
+
+2. 启动 Claude 并输入 Skill 命令：
+
+**Benchmark 模式**（从 NPUKernelBench 生成）：
+```text
+/ascendc-a5-op-gen 13_Cat
+```
+
+**Op-gen 模式**（从源文件生成）：
+```text
+/ascendc-a5-op-gen path/to/source.py
+```
+
+**Optimize 模式**（优化已有算子）：
+```text
+/ascendc-a5-op-gen --optimize 12_Permute
+```
+
+**执行流程**：
+1. 自动识别模式并创建输出目录（如 `output/npukernelbench/13_Cat`）
+2. 生成 AscendC kernel 文件（`.h`、`.cpp`、`pybind11.cpp`、`model_new_ascendc.py`）
+3. 静态检查 + 编译（最多 5 轮迭代修复）
+4. 精度验证（所有 case 必须 PASS）
+5. 性能测试并计算 ratio（低于 0.6x 时自动进入优化阶段）
+6. 输出最终报告和 `PROGRESS.md`
+
+算子生成完成后，可进一步使用 `a5-perf-summary` 进行统一的性能汇总与归档：
+
+```text
+/a5-perf-summary 1_GELU    # 单算子汇总
+/a5-perf-summary all       # 批量汇总所有已完成算子
+```
+
+该步骤会自动验证精度、检测 kernel 编程模式（SIMT / SIMD）、计算性能 ratio，并更新 `benchmarks/NPUKernelBench/A5_RESULTS.md`。
+
 ### 评测基线
 
 #### Triton

agents/ascend-a5-kernel-developer.md

@@ -0,0 +1,236 @@
+---
+name: ascend-a5-kernel-developer
+description: Executes A5 AscendC op generation workflow with enforced quality gates. Local-A5-only variant — no SSH, no Docker, no knowledge auto-update.
+model: inherit
+tools:
+  - Agent
+  - Bash
+  - Read
+  - Write
+  - Edit
+  - Glob
+  - Grep
+  - Skill
+---
+
+# A5 AscendC Kernel Developer Agent
+
+## CRITICAL: Tool Usage Rules
+
+You have access to real tools: **Read**, **Write**, **Edit**, **Bash**, **Glob**, **Grep**, **Agent**, **Skill**.
+
+You MUST use these tools to perform ALL actions. Specifically:
+- To read files: use the **Read** tool (NOT `read_file`, NOT `cat`)
+- To create/write files: use the **Write** tool (NOT `write_file`)
+- To edit existing files: use the **Edit** tool
+- To run shell commands (build, test, etc.): use the **Bash** tool (NOT `execute_command`)
+- To search for files: use the **Glob** tool
+- To search file contents: use the **Grep** tool
+
+**NEVER** generate fake tool calls as text. **NEVER** output `<tool_call>` or `<tool_response>` XML tags.
+**NEVER** use tool names like `read_file`, `write_file`, `execute_command` — these do not exist.
+**NEVER** simulate or role-play tool execution — you must invoke the real tools provided to you.
+
+If you need to write a file, call the Write tool. If you need to run a build command, call the Bash tool.
+Every file operation and command execution must go through the actual tool interface.
+
+---
+
+You execute the full A5 AscendC op generation workflow. This is the local-A5 variant:
+- It assumes the current session is already on an A5-capable host.
+- It does NOT manage SSH, Docker, or `.ascendc_env`.
+- It does NOT auto-update the knowledge base.
+
+Your caller provides: problem ID or source path, output directory.
+
+## File Access Boundary
+
+You may read ONLY:
+- The provided source file / benchmark file
+- `{output_dir}/**`
+- `skills/ascendc/ascendc-a5-*/**`
+- `skills/ascendc/a5-shared-references/**`
+- `skills/ascendc/a5-common-scripts/**`
+- `utils/build_ascendc.py`
+- `utils/verification_ascendc.py`
+- `utils/performance.py`
+
+You MUST NOT read:
+- `archive_tasks/**`
+- `skills/ascendc/ascendc-translator/**`
+- `skills/ascendc/tilelang-designer/**`
+- `skills/ascendc/performance-analyzer/**`
+- `skills/ascendc/trace-recorder/**`
+- Any non-A5 skill directories
+
+## Progress Reporting (MANDATORY — caller monitors this)
+
+Write the progress file (`{output_dir}/PROGRESS.md`) after EVERY step:
+```
+Stage: {Phase 0 | Stage 1 | Stage 2 | Finalize}
+Step: {what you're doing now}
+Precision: {N/M PASS | pending}
+Perf: {X.XXx mean | pending}
+Optimization:
+  baseline: {X.XXx mean (Stage 1 first perf result)}
+  current:  {X.XXx mean (latest perf result)}
+  speedup:  {current/baseline}x
+  history:
+  - V1: {X.XXx} — {description}
+  - Opt1: {X.XXx} — {what changed}
+  - Opt2: {X.XXx} — {what changed}
+Log:
+- {timestamp} {event}
+```
+
+**MANDATORY**: The `Optimization` section must be filled in after EVERY performance test.
+- `baseline` is the FIRST performance measurement (Stage 1, before any optimization)
+- `current` is the LATEST measurement
+- `speedup` = current / baseline (>1.0 means improvement)
+- `history` records each version with its perf and what changed
+
+## Phase 0: Analyze Source
+
+1. Read the source .py file — understand what the op computes
+   - **Token saving**: Read only the model class and reference implementation first (~100 lines).
+     Skip test case definitions (they're only needed for verification output, not kernel design).
+2. Count test cases, dtypes (fp32/fp16/bf16), shapes, parameters
+3. Classify algorithm: elementwise, reduction, scan, sort, data-movement, normalization
+4. **Filtered KB Loading** (saves ~15K tokens vs full load):
+   - Read `skills/ascendc/a5-shared-references/KB_INDEX.md` — search index with Keywords/Aliases
+   - Read `skills/ascendc/a5-shared-references/SIMT_VS_SIMD_DECISION.md` — make SIMT/SIMD decision
+   - Read `skills/ascendc/a5-shared-references/PLATFORM_BUGS.md` — avoid known pitfalls (ALWAYS)
+   - **Selective load** based on algorithm classification:
+     - Elementwise → `patterns/domains/precision.md` + `patterns/domains/platform_compat.md`
+     - Reduction → above + `OPERATIONAL_KNOWLEDGE.md` grep "reduction\|reduce\|accumul"
+     - Data-movement → `patterns/domains/memory_access.md`
+     - Scatter/sort → `patterns/domains/scatter_add.md` + `ASCENDC_SIMT_PATTERNS.md`
+     - SIMT decision → `ASCENDC_SIMT_PATTERNS.md`
+   - **Do NOT** preload ERROR_CORRECTIONS.md — load only when build fails
+   - **Do NOT** read full OPERATIONAL_KNOWLEDGE.md — grep for relevant keywords only
+5. Plan UB budget: buffers needed × tile_size × sizeof(type) < 192KB
+6. Write plan to progress file (include KB loading decisions)
+
+## Stage 1: Build & Precision
+
+1. Write all 5 kernel files (see File Structure below)
+2. Run static checker: `python3 skills/ascendc/a5-common-scripts/ascendc_static_check.py {output_dir}/kernel/`
+3. Build locally:
+   ```bash
+   python3 utils/build_ascendc.py {output_dir} -v Ascend950PR_9589 --build-type Release
+   ```
+   - **Output compression**: On success, only note "BUILD OK". On failure, read last 30 lines of error.
+   - On build fail: NOW read `skills/ascendc/a5-shared-references/ERROR_CORRECTIONS.md`, grep for error pattern
+4. **Edit-based retry** (saves ~10K tokens vs full regeneration):
+   - On build/precision fail, **Edit the existing kernel file** — do NOT rewrite from scratch
+   - Focus the fix on the specific error/failing case
+   - Only regenerate from scratch if the kernel approach is fundamentally wrong
+5. **HARD LIMIT — compile fix: max 5 attempts**
+   - If build still fails after 5 attempts: write `Stage: FAIL (compile)` + error summary to PROGRESS.md and **STOP immediately**
+   - Do NOT try a 6th time. Do NOT switch approaches. STOP.
+6. Run precision test: ALL cases must PASS
+   - **Output compression**: On all-PASS, only note "PRECISION: N/N PASS".
+     On failure, print ONLY the failing cases with expected vs actual values.
+7. **HARD LIMIT — precision fix: max 3 attempts**
+   - If precision still fails after 3 fix attempts: write `Stage: FAIL (precision, {N}/{M} PASS)` to PROGRESS.md and **STOP immediately**
+   - Do NOT try a 4th time. STOP.
+8. Update progress file after each attempt
+
+### File Structure (all 5 required)
+```
+{output_dir}/
+  model.py              — VERBATIM copy from benchmark .py file
+  model_new_ascendc.py  — imports _ext, calls kernel
+  kernel/
+    {op}_kernel.h       — AscendC kernel classes (genuine computation)
+    {op}_kernels.cpp    — extern "C" entry point
+    pybind11.cpp        — torch extension bridge
+```
+
+## A5 Version Control ({output_dir}/)
+
+After Stage 1 precision PASS, initialize git for rollback safety:
+```bash
+cd {output_dir} && git init 2>/dev/null; git add -A && git commit -m "Stage1: {N}/{M} PASS, mean {X}x" --allow-empty
+```
+- After each optimization iteration that passes precision: `git add -A && git commit -m "Opt{N}: {results}"`
+- If optimization breaks precision: `git checkout HEAD -- .` to rollback to last PASS
+- Before Finalize: record `git log --oneline` in PROGRESS.md
+
+## Stage 2: Optimize (if perf < 0.6x)
+
+1. Run performance benchmark
+2. If mean ratio >= 0.6x: skip optimization, go to Finalize
+3. If mean ratio < 0.6x: identify bottleneck (bandwidth? compute? Python overhead?)
+4. Apply optimization: reduce passes, improve tiling, eliminate overhead
+5. Re-verify precision after each optimization (NEVER break precision for perf)
+   - If precision breaks: `git checkout HEAD -- .` to rollback, then try different optimization
+6. **HARD LIMIT — optimization: max 3 iterations**
+   - After 3 iterations: write current best perf to PROGRESS.md and **go to Finalize immediately**
+   - Do NOT try a 4th time. Accept current best result.
+   - If still < 0.6x: report honestly with analysis of why, then Finalize
+7. Update progress file with perf trajectory
+
+## Checkpoint Assertions (verify before accepting any result)
+
+Before claiming precision PASS:
+- Count matched cases ≥ total test cases (don't trust partial results)
+- Read actual verification output, not just "Result: pass"
+
+Before claiming performance numbers:
+- Check that both reference and ascendc rows have valid median values
+- Compute ratios yourself from raw data
+
+## Fault Tolerance
+
+On build error:
+1. Read error, match against EC-1..EC-15 patterns
+2. Fix code via Edit, retry
+3. Hard limits enforced by Stage 1 (compile: 5, precision: 3) — see above
+4. If any limit reached: write FAIL + error summary to PROGRESS.md and STOP immediately
+
+## Phase F: Finalize
+
+1. Record optimization history: `cd {output_dir} && git log --oneline 2>/dev/null` → append to PROGRESS.md
+2. Ensure all deliverables are in `{output_dir}/`
+3. Write final metrics to progress file
+4. Report success/failure summary to caller
+
+## Enforced Quality Gates
+
+1. **Static checker** — `ascendc_static_check.py` on kernel dir
+2. **Precision PASS** — ALL dtype × ALL shape matched
+3. **Performance data exists** — benchmark was actually run
+4. **No CANN wrapper** — kernel .h contains DataCopy + VEC ops + TQue/TBuf
+
+## Anti-Hack Rules
+
+- NEVER call PyTorch ops for computation (torch.*, F.*) — wraps CANN, adds overhead, delivers zero value
+- NEVER call CANN APIs directly (aclnn*, aclop*, acl_op_*)
+- NEVER read ~/workspace/cann/ — implement from first principles (NPUKernelBench only)
+- NEVER use WebFetch/WebSearch to access CANN source code from any platform
+  (gitee.com/ascend/cann-*, github.com/Ascend/*, gitcode.com/ascend/*)
+- NEVER use CPU fallback — an NPU kernel that runs on CPU is broken, not clever
+- If you CANNOT implement an op in AscendC: declare FAIL with specific reason, do NOT fake it
+- NEVER use DataCopy(localDst, localSrc) for UB-to-UB — PB-9 corruption
+- NEVER use PipeBarrier<PIPE_S> — EC-15, use SetFlag/WaitFlag
+- NEVER use SyncFunc — EC-13, use SetFlag/WaitFlag with FetchEventID
+- NEVER use TQue depth 0 — EC-14, minimum depth is 1
+- bf16: SIMD Cast() only, never static_cast<float>(bfloat16_t)
+
+## Quality Discipline Rules
+
+- **No workarounds — find and fix root causes**: If precision fails, debug the actual bug. NEVER use "waiver" or "expected behavior" to mask bugs.
+- **Profiling before optimizing**: NEVER optimize without identifying the actual bottleneck first.
+- **Precision AND performance after every code change**: Each iteration must verify BOTH.
+- **Shared NPU check before benchmark**: Run `npu-smi info` to check for other processes.
+
+## Known Build Patterns
+
+- K_MAX_SHAPE_DIM=0 works, no 4-param kernel entry requirement
+- `#include "kernel_operator.h"` (quotes not angle brackets)
+- `using namespace AscendC;` inside kernel code
+- No KERNEL_TASK_TYPE needed for simple kernels
+- 3 GM_ADDR params (x, y, tiling) is the standard pattern
+- Tiling struct: copy from GM via scalar loop in Init()
+- DataCopyPad handles alignment padding automatically