tilelang2ascend-tilelang-designer

Generates TileLang kernel designs and implementations for complex Ascend NPU operators.

Updated Sep 15, 2026
One-click install
npx skills add https://github.com/WangWindow/CANN-BatchMatMulMaxsum --skill tilelang2ascend-tilelang-designer-wangwindow
Or copy as Structured Prompt for Agent
Please help me install this Agent Skill.
Skill: tilelang2ascend-tilelang-designer
Source: https://github.com/WangWindow/CANN-BatchMatMulMaxsum/tree/main/.agents/skills/tilelang2ascend-tilelang-designer
Command: npx skills add https://github.com/WangWindow/CANN-BatchMatMulMaxsum --skill tilelang2ascend-tilelang-designer-wangwindow

SYSTEM DOCUMENTATION & REQUIREMENTS

💡 This Skill requires torch, torch_npu, tilelang, and includes scripts (resource) and references (resource) components.

What problem does it solve? Developing high-performance fused operators (Attention, MatMul variants, Norm, Sort/TopK) for Ascend NPUs requires coordinating block-level task partitioning, tile-level implementation, and measurement-driven performance tuning, which is error-prone when done ad hoc. ## Core Features & Use Cases - Orchestrated Design Workflow: Coordinates design methodology and coding-convention references to produce block_level/ designs, tile_level/ TileLang kernels, and a model_new_tilelang.py implementation from a PyTorch reference model. - Mandatory Pattern Routing: Enforces gated checklists for Attention, reduction, shuffle, Sort/TopK, and Norm-fusion operator families before any design code is written. - Verification and Performance Iteration: Ships scripts for TileLang accuracy verification, PyTorch-regression static detection, and msprof-based performance tuning with a geomean >= 0.6x target versus PyTorch reference kernels. - Use Case: Given a BatchMatmulMaxSum PyTorch model, produce a fused single-kernel TileLang design with CV-fusion pipelining, pass accuracy validation, and iterate on performance before handing off to the AscendC translator. ## Quick Start Run this skill on my output_dir containing model.py to generate the block-level and tile-level TileLang design plus model_new_tilelang.py for my fused operator.

Frequently Asked Questions about tilelang2ascend-tilelang-designer

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

FAQPage Schema
How do I generate a TileLang kernel design from a PyTorch model?

Provide an output directory containing model.py with the reference PyTorch model. The skill produces design/block_level/ and design/tile_level/ TileLang kernels plus model_new_tilelang.py, then runs accuracy verification and performance iteration.

What operator types does TileLang kernel design support on Ascend NPU?

It targets complex operators: Attention variants (FlashAttention, GQA, sparse), fused MatMul, Norm variants (RMSNorm, LayerNorm), Sort/TopK, and multi-input fused ops. Simple operators like Gather or Scatter use a separate direct-invoke workflow instead.

How is TileLang kernel accuracy verified against PyTorch?

The evaluate_tilelang.sh script runs verification_tilelang.py, which loads model.py and model_new_tilelang.py, executes both on NPU with identical inputs, and compares outputs with atol/rtol tolerances. Accuracy must pass before any performance optimization is allowed.

Can model_new_tilelang.py use torch operators for part of the computation?

No. All core computation must be fused into a single TileLang kernel; torch is only allowed for tensor creation and layout transforms. The validate_tilelang_impl.py script statically detects PyTorch fallback, uncalled kernels, and scalar for-loop regressions.

What is the performance target for TileLang kernels versus PyTorch?

The target is geomean speedup >= 0.6x of TileLang kernel time versus PyTorch reference kernel time, measured with msprof over at least 20 runs. If the baseline misses the target, the skill iterates on optimization points with an upper-bound analysis before delivery.