Borrowing it
Nothing to install: this file belongs to tile-ai/tilelang-ascend. Take a copy, put it at the same path in your own repository, and replace the rules that are about this project with yours.
curl -O https://raw.githubusercontent.com/tile-ai/tilelang-ascend/ascendc_pto/.agents/skills/npu-arch/SKILL.mdgit clone --depth 1 https://github.com/tile-ai/tilelang-ascendWrote this? Show the measurements
A badge with what this costs and how it scanned, read live from this page, so it follows the numbers instead of freezing them. Markdown for a README, HTML for a documentation site or a project page.
[](https://agentmods.dev/skills/tile-ai/tilelang-ascend/npu-arch)<a href="https://agentmods.dev/skills/tile-ai/tilelang-ascend/npu-arch"><img src="https://agentmods.dev/badge/skills/tile-ai/tilelang-ascend/npu-arch/github.svg" alt="Measured on agentmods" height="20"></a>Or the 80×15 button, for a site that already has a row of RSS and ATOM ones. Only the verdict fits; the numbers stay here.
<a href="https://agentmods.dev/skills/tile-ai/tilelang-ascend/npu-arch"><img src="https://agentmods.dev/badge/skills/tile-ai/tilelang-ascend/npu-arch.svg" alt="Reviewed on agentmods" width="80" height="20"></a>What it costs to keep this loaded
Counted locally with the o200k_base tokenizer, which is exact for GPT models; Claude uses its own tokenizer and its counts differ. Treat this as one consistent yardstick across the catalogue rather than a bill. Prices are per million input tokens.
| Model | Per session | Once invoked |
|---|---|---|
| Fable 5.1 | $0.00081 | $0.01459 |
| Opus 5.5 | $0.00032 | $0.00584 |
| Sonnet 5.5 | $0.00016 | $0.00292 |
| Haiku 4.5 | $0.00008 | $0.00146 |
Grade A, and why
npu-arch scanned grade A with 0 findings against 26 rules in 11 categories — prompt injection, anti-refusal, data exfiltration, privilege escalation, supply chain, agent snooping, system-prompt leakage, SSRF and excessive agency — measured today.
A static scan of the body, not an audit. Every finding is printed with the line that produced it so you can judge whether it matters here. A mod is markdown that instructs an agent; that is exactly why what it instructs is worth reading.
Nothing flagged
None of the 26 patterns this scan looks for appear in this file: no shell pipes, no recursive deletes, no credential paths, no hidden text, no instruction-override or anti-refusal phrasing, no agent-config snooping. That is not a guarantee, it is the absence of the things that are checkable.
How it starts
The opening of the file, as written. The whole thing — 88 lines — stays where its author put it; the contents beside it link to each section on GitHub.
Ascend NPU 架构知识
架构代际概述
| 概念 | 说明 |
|---|---|
| NpuArch | 芯片架构号,定义指令集和微架构,运行时通过 GetCurNpuArch() 获取 |
| SocVersion | 片上系统版本,软件命名标识,运行时通过 GetSocVersion() 获取 |
| NPU_ARCH | Device 侧编译宏,四位数值,用于条件编译 |
| archXX | 算子仓目录简写,取 DAV 编号前两位(如 DAV_3510 → arch35) |
| DAV_C310 | 构建系统内部宏,等价于 NpuArch::DAV_3510 / arch35 / __NPU_ARCH__=3510,不可按数值推断 |
完整映射表
完整产品系列 / SocVersion / NpuArch / 芯片型号映射见 npu-hardware-params.md §0 产品映射表。
获取当前架构
#include "utils/tiling/platform/platform_ascendc.h"
auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
NpuArch npuArch = ascendcPlatform.GetCurNpuArch(); // DAV_2201 / DAV_3510 / ...
platform_ascendc::SocVersion socVer = ascendcPlatform.GetSocVersion();
GetCurNpuArch() 失败返回 NpuArch::DAV_RESV,GetSocVersion() 失败返回 SocVersion::RESERVED_VERSION。
DAV_3510 相对 DAV_2201 的关键变化
详细硬件参数真值见
references/npu-hardware-params.md。运行时必须以PlatformAscendC接口获取实际值,禁止硬编码。
Buffer(同架构内通常一致)
| Buffer | DAV_2201 | DAV_3510 |
|---|---|---|
| L0C | 128 KB | 256 KB |
| UB | 192 KB | 248 KB |
| BT | 1 KB | 4 KB |
频率/核数/L2/Memory(因子型号/形态而异)
| 参数 | Ascend910B2 (DAV_2201) | Ascend950PR PCIE | Ascend950PR Server |
|---|---|---|---|
| Cube 核数 | 24 | 28 | 32 |
| 频率 | 1.8 GHz | 1.65 GHz | 1.65 GHz |
| L2 | 192 MB | 112 MB | 128 MB |
| Memory | 64 GB | 112 GB | 128 GB |
详见 典型SKU示例。
Memory 带宽:950PR 1.6/1.4 TB/s 已由 950 白皮书表3-1 证实,910B2 仍为公开资料与经验值,详见
npu-hardware-params.md§4。
指令集与微架构
| 类别 | 变化 |
|---|---|
| 数据格式 | 新增 FP8 / MXFP8 / MXFP4 / HiF8 Cube MMAD |
| CV 直通通路 | 新增 L0C→UB、UB→L1、SSBuffer 消息 |
| 同步机制 | BufferID 替代 set/wait 强配对 |
| 编程模型 | 新增 SIMT、SIMD-Regbase、NDDMA、CCU 通算融合 |
详细文档索引
references/ 按需加载:
npu-hardware-params.md— 硬件参数参考:各架构子型号一致参数、典型SKU示例、基于公开资料与经验值的规格simt-arch-guide.md— SIMT 架构概念参考:线程层次结构、Warp调度机制、内存空间、UB分区、数据类型、函数调用层级、LAUNCH_BOUND与寄存器数量、SIMT与SIMD核心差异npu-arch-guide.md:- 典型硬件参数与获取方式:核数 / Buffer 容量典型值,
GetCoreNumA*/GetCoreMemSize/aclrtGetDeviceInfo接口用法(含硬编码反例) - DAV_3510 微架构:AI Core Buffer 层级、MTE 引擎、CV 数据通路改动、BufferID 同步
- SIMT vs SIMD:硬件能力差异与适用场景
- SIMD-Regbase:寄存器组、核心优势、代码形态变化
- 数据格式扩展:完整数据类型表、C++ 类型名映射、MXFP Tiling 约束
- NDDMA / CCU:高维 DMA 用法、CCU 三种通信范式、KFC 调度模型
- 架构兼容性检查清单 / 参考信息来源
- 典型硬件参数与获取方式:核数 / Buffer 容量典型值,
What ships with it
5 files beside SKILL.md in the same directory: the scripts, references and assets a skill reads on demand. Not counted in the per-session cost; read them before you install if any of them is executable.
What this file has done since we first saw it
Hashed on every crawl. A supply-chain change to an agent config is a question of when, not whether, so the history is kept rather than the latest state alone.
- today First seen · 88 lines · 81 tokens per session scan A 112c588afddb
npu-arch is a skill published in the GitHub repository tile-ai/tilelang-ascend (427 stars, last pushed yesterday), licensed MIT. It adds 81 tokens to every session and 1,459 once invoked, about $0.0003 per session on Opus 5.5. A static security scan graded it A with 0 findings. No closer match exists in the catalogue, so it is treated as the original; first seen 2026-11-06.
npu-arch compared with skills from other repositories
| Mod | Type | Stars | Last push | Licence | Scan | Tokens | Works with |
|---|---|---|---|---|---|---|---|
| npu-archtile-ai/tilelang-ascend | Skill | 427 | yesterday | MIT | A | 81 | Claude Code, Codex |
| gke-compute-classesgoogle/skills | Skill | 22k | yesterday | Apache-2.0 | A | 83 | Claude Code, Codex |
| jetson-diagnosticNVIDIA/skills | Skill | 3.9k | yesterday | Apache-2.0 | A | 30 | Claude Code, Codex |
| holoscan-install-wheelNVIDIA/skills | Skill | 3.9k | yesterday | Apache-2.0 | A | 34 | Claude Code, Codex |
| offensive-z-waveSnailSploit/Claude-Red | Skill | 10k | 25d ago | MIT | A | 113 | Claude Code, Codex |
| doca-socket-relayNVIDIA/skills | Skill | 3.9k | yesterday | Apache-2.0 | A | 236 | Claude Code, Codex |
| hsb-flashNVIDIA/skills | Skill | 3.9k | yesterday | Apache-2.0 | B | 94 | Claude Code |