tilelang-ascend: Skill for Claude Code

.agents/skills/npu-arch/SKILL.md

npu-arch is a skill for Claude Code, Codex from tile-ai/tilelang-ascend. It costs 81 tokens per session (1,459 once invoked), scanned A, original, MIT.

A reference skill for Ascend NPU hardware architecture and its software identifiers. An NPU is a processor designed for neural-network workloads; the skill maps chip models, architecture generations, and compile-time feature names.

In plain words
What is it for?
It helps identify an NPU architecture, compare supported features, map chip and software versions, and choose conditional compilation paths.
Why use it?
It helps avoid choosing the wrong hardware branch or assuming that similarly named architecture values mean the same thing.

Skill for Claude CodeCodex

Written for no agent in particular: nothing here depends on one. Also seen: installed under .agents/ (shared by several agents).

This is tile-ai/tilelang-ascend's own configuration. It tells Claude Code and Codex how to work on tilelang-ascend itself, so it is not a mod to install elsewhere. Copy it as a starting point and replace the rules that are about this project. Everything tilelang-ascend configures →

Reuse

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.

Copy the file
curl -O https://raw.githubusercontent.com/tile-ai/tilelang-ascend/ascendc_pto/.agents/skills/npu-arch/SKILL.md
Clone the repo
git clone --depth 1 https://github.com/tile-ai/tilelang-ascend

Made for: Claude Code, Codex.

Wrote 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.

agentmods badge for npu-arch

README.md
[![agentmods](https://agentmods.dev/badge/skills/tile-ai/tilelang-ascend/npu-arch/github.svg)](https://agentmods.dev/skills/tile-ai/tilelang-ascend/npu-arch)
Your own site
<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.

agentmods 80×15 button for npu-arch

Your own site · 80×15
<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>
Per session 81 Skills are progressive disclosure: only the name and description are preloaded; the body loads when the skill is used.
When invoked 1,459 The whole file, excluding the scripts and references it only reads on demand.
Security scan A 0 findings. A grade says what 26 rules found in the file — not that it is safe.
Origin original No closer match found in the catalogue.
Token cost

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.

ModelPer sessionOnce 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

Measured today against content hash 112c588afddb, method: parsed. Prices are Anthropic first-party input rates as of 2026-10-07, from the pricing page.

Security

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.

.agents/skills/npu-arch/SKILL.md · 88 lines

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 调度模型
    • 架构兼容性检查清单 / 参考信息来源

Read the full file on GitHub · 88 lines

Files

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.

Changes

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.

  1. today First seen · 88 lines · 81 tokens per session scan A 112c588afddb

Subscribe to this mod's changes

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.

Related

npu-arch compared with skills from other repositories

npu-arch next to 6 skills from other repositories
ModTypeStarsLast pushLicenceScanTokensWorks with
npu-archtile-ai/tilelang-ascendSkill427yesterdayMITA81Claude Code, Codex
gke-compute-classesgoogle/skillsSkill22kyesterdayApache-2.0A83Claude Code, Codex
jetson-diagnosticNVIDIA/skillsSkill3.9kyesterdayApache-2.0A30Claude Code, Codex
holoscan-install-wheelNVIDIA/skillsSkill3.9kyesterdayApache-2.0A34Claude Code, Codex
offensive-z-waveSnailSploit/Claude-RedSkill10k25d agoMITA113Claude Code, Codex
doca-socket-relayNVIDIA/skillsSkill3.9kyesterdayApache-2.0A236Claude Code, Codex
hsb-flashNVIDIA/skillsSkill3.9kyesterdayApache-2.0B94Claude Code