skills.rest
WangWindow
WangWindowCOMMUNITY
Explore all skills by this author →
WebsiteGitHub

pypto-op-develop

Generates PyPTO operator kernel implementation files following Layer A-L design conventions.

Updated Sep 15, 2026
Search Tags:#code generation#tiling#jit#pypto#cann#ascend npu#operator kernel
One-click install
npx skills add https://github.com/WangWindow/CANN-BatchMatMulMaxsum --skill pypto-op-develop-wangwindow
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: pypto-op-develop
Source: https://github.com/WangWindow/CANN-BatchMatMulMaxsum/tree/main/.agents/skills/pypto-op-develop
Command: npx skills add https://github.com/WangWindow/CANN-BatchMatMulMaxsum --skill pypto-op-develop-wangwindow

SYSTEM DOCUMENTATION & REQUIREMENTS

💡 This Skill includes scripts (resource) and references (resource) components.

What problem does it solve? Writing PyPTO operator kernels for Ascend NPU hardware requires strict adherence to layered architecture rules, tile shape constraints, dynamic axis handling, and JIT entry conventions. This Skill guides a coder agent through per-Phase incremental construction of operator implementation files, preventing common compile-time and runtime failures. ## Core Features & Use Cases - Per-Phase Module Generation: Produces one <op>_module<k>_impl.py per orchestrator dispatch based on DESIGN.md and module_interfaces.yaml contracts, then a cleanup pass that consolidates into <op>_impl.py plus a Chinese README.md. - Layer G-K Structure Enforcement: Builds implementations on the impl_template.py skeleton with strict separation between cache bridge, PyPTO sub-kernels, loop-owning kernel impl, JIT entry, and host wrapper. - Constraint & Error Guidance: Ships execution-constraints references, error-code troubleshooting flows, and shell scripts for CANN environment setup and idle NPU chip detection. - Use Case: An orchestrator assigns active_module M2 for a fused attention operator; the coder agent reads the design docs, emits the applicable constraint checklist, and generates exactly one compliant module impl file before stopping. ## Quick Start Ask the agent to implement the current active module of your operator using the pypto-op-develop skill with the existing SPEC.md, DESIGN.md, and module_interfaces.yaml.

Frequently Asked Questions about pypto-op-develop

High-intent search queries and answers about installing and using this skill.

FAQPage Schema
How do I implement a PyPTO operator kernel for Ascend NPU?▼

Read the design documents and module_interfaces.yaml, then generate one module impl file per Phase using the impl_template.py skeleton with Layers G through K. Set tile shapes per stage, use pypto.loop for dynamic axes, and write outputs back via assemble or slice assignment.

What inputs are required before coding a PyPTO operator?▼

You need the operator spec (name, math formula, input/output shapes and dtypes), the design plan (API mapping, tiling strategy, loop structure), and the golden reference implementation. If any are missing, report back to the orchestrator instead of guessing.

Why does PyPTO matmul fail with dynamic shape dimensions?▼

PyPTO matmul requires concrete compile-time shapes, so DYNAMIC axes appear as -1 and are rejected. Use the 2D reshape plus nested pypto.loop plus concrete tile pattern, where view shapes are Python ints and only offsets carry symbolic values.

Can the host wrapper use a Python for loop to call the kernel per chunk?▼

No. Layer K must call the JIT entry exactly once; chunk iteration belongs inside Layer I as pypto.loop with pypto.view offsets. Python loops in the wrapper collapse the kernel into single-chunk launches and fail verification gates.

How do I fix pypto.zeros device keyword errors in the host wrapper?▼

Layer K runs in host context, so allocate outputs with torch.empty or torch.zeros with explicit dtype and device, then pass them to the JIT entry. pypto.zeros and similar creation APIs only work inside the JIT graph.

When is pypto.loop(1) allowed in a kernel?▼

Only when the kernel has no other pypto.loop call and a vector-pipe op needs to satisfy the layout check. If a real pypto.loop(N) already exists, wrapping it in pypto.loop(1) is forbidden.

Related Skills

Explore More Software Engineering Skills→
AscendAscend
31

ascendc-operator-design

Generate AscendC operator design documents with tiling and performance plans.

Official
Advanced
guqiong96guqiong96
105

add-jit-kernel

Add custom CUDA kernels to the SGLang JIT kernel module.

Community
Advanced
hw-native-syshw-native-sys
101

feature-dev

Develops PyPTO IR features including operators, passes, and transformations.

Official
Intermediate
xchang1121xchang1121
6

tilelang-ascend-attention

Guides implementation of Attention operators using TileLang Cube and Vector fusion on Ascend hardware.

Community
Advanced