tilelang-ascend-debugging

Diagnose common coding pattern errors in TileLang-Ascend operator kernels.

6|1|Updated Apr 19, 2026
One-click install
npx skills add https://github.com/xchang1121/op-autoresearch --skill tilelang-ascend-debugging-xchang1121
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: tilelang-ascend-debugging
Source: https://github.com/xchang1121/op-autoresearch/tree/main/skills/tilelang-ascend/fundamentals/tilelang-ascend-debugging
Command: npx skills add https://github.com/xchang1121/op-autoresearch --skill tilelang-ascend-debugging-xchang1121

SYSTEM DOCUMENTATION & REQUIREMENTS

What problem does it solve? Writing TileLang kernels for Ascend NPU hardware involves subtle pitfalls—misused symbolic expressions, incorrect scalar-vector operations, and confused kernel launch semantics—that produce cryptic compile or runtime failures. This Skill provides a structured checklist and targeted fixes for the most frequent TileLang-Ascend coding mistakes. ## Core Features & Use Cases - Pre-flight Checklist: Verify output index alignment, UB buffer capacity, Expert/Developer mode configurations, and fused-operator workspace dataflow before running kernels. - Pattern Fixes: Get concrete correct implementations for dynamic shapes with T.symbolic, scalar-vector arithmetic via T.tile APIs, and proper T.Kernel versus T.serial usage. - Constraint Enforcement: Avoid forbidden patterns such as Python built-in functions (min, max, and, or) on TVM symbolic expressions and low-precision T.tile.reciprocal division. - Use Case: After generating a fused Cube-Vector operator for Ascend 910B, run through the checklist to confirm workspace_idx matches the function signature and AUTO_CV_COMBINE pass configs are enabled before evaluation. ## Quick Start Review my TileLang-Ascend kernel code against the debugging checklist and fix any coding pattern violations.

Frequently Asked Questions about tilelang-ascend-debugging

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

FAQPage Schema
How do I handle dynamic shapes in TileLang kernels?▼

Use T.symbolic to declare dynamic dimensions, for example N = T.symbolic('N', 'int32'). This lets the kernel accept runtime-varying sizes instead of hardcoding static shapes in tensor declarations.

How to perform scalar-vector arithmetic in TileLang-Ascend?▼

Scalar-vector operations must use T.tile APIs, not Python operators like + or /. For example, x + 1.0 becomes T.tile.add(x, x, 1.0), and the scalar must be the second operand since some APIs do not accept scalars in other positions.

Why does my TileLang kernel fail when using Python min or max?▼

TileLang kernel expressions are symbolic TVM Expr objects, so Python built-ins like min(), max(), and, or, not are unsupported. Replace them with T.min and T.max, for example hw_end = T.min(hw_start + block_HW, H * W).

What is the difference between T.Kernel and T.serial in TileLang?▼

T.Kernel(n_num, is_npu=True) launches n_num parallel blocks, each running the kernel body once, while T.serial is a sequential loop inside a single block. Do not nest T.serial(n_num) inside T.Kernel(n_num) with the same count, as that duplicates the partitioning semantics.

Why is T.tile.reciprocal forbidden for division in TileLang-Ascend?▼

T.tile.reciprocal has insufficient precision for division operations. Instead compute 2.0 / x using T.tile.div with a broadcasted numerator, such as T.tile.div(dst, T.broadcast(2.0, shape), x).