xllm: Skill for Claude Code

TileLang Ascend Kernel

When to use

Use this skill when the task involves any of the following in the xLLM repo:

• xllm/python/kernels_npu/tilelang/*.py
• xllm/compiler/tilelang/targets/ascend/aot/*.py
• xllm/core/kernels/npu/tilelang/*_wrapper.cpp
• xllm/core/kernels/npu/tilelang/CMakeLists.txt
• generated TileLang artifacts such as manifest.json, registry.inc, or specialization .cpp

Run build and test commands inside the NPU container.

Run TileLang commands from the xLLM repo root, not from an installed-package environment.

Entry points and TL_ROOT

• Prefer python xllm/compiler/tilelang_launcher.py ... for end-to-end TileLang compile flows.
• From the xLLM repo root, use export TL_ROOT=$PWD/third_party/tilelang-ascend for xLLM TileLang tooling and verify test -f "$TL_ROOT/tilelang/__init__.py".
• Before any raw script does import tilelang, run export TL_ROOT=$PWD/third_party/tilelang-ascend && source third_party/tilelang-ascend/set_env.sh, then execute the script.
• Do not run kernel files directly with python rope.py; use module execution because these files rely on relative imports.
• For direct kernel-script debugging, run them as modules and pass required CLI args:

cd xllm
export PYTHONPATH="$PWD:$PWD/xllm${PYTHONPATH:+:$PYTHONPATH}"
python -m compiler.tilelang.targets.ascend.aot.rope \
  --output .tmp/rope.cpp
# Expected: [INFO] RoPE output matches torch reference

• The same module-style rule applies to other AOT descriptors under xllm/compiler/tilelang/targets/ascend/aot/.

Primary Reference And Mode Preference

Primary reference:

• third_party/tilelang-ascend/docs/TileLang-Ascend Programming Guide.md

Use third_party/tilelang-ascend/.agents/skills/tilelang-custom-skill/tilelang-api-best-practices/references/api-tile-ops.md when the task depends on T.tile.xxx semantics such as compare, select, cast, or other vector intrinsics.

Default to Expert mode for xLLM Ascend kernels:

• prefer T.tile.xxx, explicit UB/shared allocation, and explicit T.copy
• prefer explicit T.serial control for row/block traversal
• do not introduce Developer mode T.Parallel unless the kernel is a clearly tile-local element-wise expression and the change does not reduce control over UB usage, temporary buffers, or exact runtime semantics
• when translating Triton kernels, preserve the Triton runtime semantics first, then choose the smallest Expert-mode lowering that matches them