What is tilelang-op-profiling?

A performance profiling workflow for TileLang kernels running on real NPU hardware. An NPU is a processor designed for neural-network workloads.

Use it to collect and compare kernel timing data, check performance against a target, locate bottlenecks, and verify changes after optimization.

It replaces guesswork with measurements from the hardware, helping identify why an operator is slow and whether an optimization worked.

How tile-ai/tilelang-ascend configures Claude Code and Codex

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/tilelang-op-profiling/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.

README.md

[![agentmods](https://agentmods.dev/badge/skills/tile-ai/tilelang-ascend/tilelang-op-profiling/github.svg)](https://agentmods.dev/skills/tile-ai/tilelang-ascend/tilelang-op-profiling)

Your own site

<a href="https://agentmods.dev/skills/tile-ai/tilelang-ascend/tilelang-op-profiling"><img src="https://agentmods.dev/badge/skills/tile-ai/tilelang-ascend/tilelang-op-profiling/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.

Your own site · 80×15

<a href="https://agentmods.dev/skills/tile-ai/tilelang-ascend/tilelang-op-profiling"><img src="https://agentmods.dev/badge/skills/tile-ai/tilelang-ascend/tilelang-op-profiling.svg" alt="Reviewed on agentmods" width="80" height="20"></a>

Per session 93 Skills are progressive disclosure: only the name and description are preloaded; the body loads when the skill is used.

When invoked 7,931 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.

Model	Per session	Once invoked
Fable 5.1	$0.00093	$0.07931
Opus 5.5	$0.00037	$0.03172
Sonnet 5.5	$0.00019	$0.01586
Haiku 4.5	$0.00009	$0.00793

Measured 7d ago against content hash 21070ef80f0d, method: parsed. Prices are Anthropic first-party input rates as of 2026-10-07, from the pricing page.

Security

Grade A, and why

tilelang-op-profiling 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 7d ago.

The scan reads SKILL.md. This mod also ships 1 executable file (scripts/perf_summary.py), listed below but not scanned — reading those needs a real analyzer, not pattern matching.

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/tilelang-op-profiling/SKILL.md · 453 lines

How it starts

The opening of the file, as written. The whole thing — 453 lines — stays where its author put it; the contents beside it link to each section on GitHub.

TileLang 算子上板性能采集与调优

在真实 NPU 上采集 TileLang kernel 性能数据，系统化解读性能指标文件，判定性能是否达标，定位瓶颈类型，并给出可操作的优化建议。

后端

上游传入 target 时原样使用；否则优先复用现有入口已经显式指定的 target。两者都未指定时默认使用 AscendC：现有 JIT 入口不传 target 时会使用 target="auto"，当前仓库将 auto 解析为 AscendC；新建临时 profiling 入口时，在 @tilelang.jit 或 tilelang.compile / tilelang.engine.lower 中显式传入 target="ascendc"。AOT 入口的 lower 与 LibraryGenerator 使用同一 target，预运行和 msprof 采集也保持该 target。

职责边界

本 Skill 负责：运行入口识别、正确性门禁、msprof op 采集、kernel/case/launch 映射校验、原始数据归档、性能判定、瓶颈诊断、修改代码、正确性复验和优化前后对比。
优先复用算子已有 target、运行入口和参数。
do_bench 可用于建立稳定延迟基线，但最终硬件流水、缓存、冲突等结论以 msprof op 数据为准。报告必须标明数据来源和单位。

适用场景

场景	说明
算子开发完成后的性能验收	确认算子达到预期性能水平
性能问题定位	通过 CSV 指标精确定位瓶颈
优化效果验证	对比优化前后的 CSV 数据
Agent team 测试阶段	tester/developer 调用，自动化性能分析

工作流

多 launch 算子、混合 AIC/AIV 计时、AOT/AclGraph 入口或用户要求 pytest 最终验收时，先读整链计时与采集。下面一个 case 一次 kernel 的命令用于单 kernel 入口；多 kernel 链路按实际名称集合与 launch 次数调整。

Step 1: 准备并校验 TileLang profiling 入口
    ↓
Step 2: msprof op 采集指标
    ├── 成功 → 继续
    └── 仅 DBI/插桩/导出故障且直调正常 → 普通 msprof fallback
    ↓
Step 3: 归档数据 + 生成统计摘要
    ↓
Step 4: 读摘要 → 对照性能标准判定（需要时读原始 CSV）
    ↓
Step 5: 如不达标，查阅瓶颈优化速查表
    ↓
Step 6: 修改代码 → 回到 Step 2 重新采集（数据自动归档为新一轮）

Step 1: 准备并校验 TileLang profiling 入口

先检查目标算子目录，按以下优先级选择入口：

已有 perf_*.py、--perf、--profiler msprof 或 bench.sh 入口；
已有单文件 example 的 __main__ 入口；
AOT 入口或 AclGraph 的公开 replay 入口；
只有上述入口无法隔离目标 kernel 时，才在临时目录生成最小 profiling wrapper。

按以下顺序确定算子输出根目录 <variant_output_dir>：

上游已传入 variant_output_dir 时原样使用；
未传入时，从目标 kernel 的生产源码路径确定 <operator_name>：
- examples/operators/<name>/... 使用 <name>；
- 其他 examples/.../<name>/<source_file> 使用生产源码所在的实际算子目录 <name>；
- 源码直接位于 examples/ 时，从文件名去掉 example_ 前缀和扩展名；
默认令 <variant_output_dir> 为仓库根目录下的 examples/operators/<operator_name>，所有算子的 profiling 输出统一放在 examples/operators/；
临时 profiling wrapper 只作为运行入口，不能用临时目录名推断算子名；
无法唯一确定生产源码或 <operator_name> 时停止并询问，不能猜测另一个算子目录。

Read the full file on GitHub · 453 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.

7d ago First seen · 453 lines · 93 tokens per session scan A 21070ef80f0d

Subscribe to this mod's changes

tilelang-op-profiling is a skill published in the GitHub repository tile-ai/tilelang-ascend (427 stars, last pushed yesterday), licensed MIT. It adds 93 tokens to every session and 7,931 once invoked, about $0.0004 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-10-30.

Analytics and metrics

tilelang-op-profiling next to 6 skills from other repositories
Mod	Type	Stars	Last push	Licence	Scan	Tokens	Works with
tilelang-op-profilingtile-ai/tilelang-ascend	Skill	427	yesterday	MIT	A	93	Claude Code, Codex
imagegenopenai/codex	Skill	133k	yesterday	Apache-2.0	A	113	Codex
playwright-climicrosoft/playwright	Skill	99k	yesterday	Apache-2.0	A	19	Claude Code
local-ai-agentsmicrosoft/ai-agents-for-beginners	Skill	79k	4d ago	MIT	A	200	Claude Code, Codex
systematic-debuggingobra/superpowers	Skill	313k	3d ago	MIT	A	21	Claude Code
next-cache-components-adoptionvercel/next.js	Skill	143k	today	MIT	A	95	Claude Code
next-partial-prefetching-adoptionvercel/next.js	Skill	143k	today	MIT	A	103	Claude Code

tilelang-ascend: Skill for Claude Code