skills.rest
fla-org
fla-orgOFFICIAL
Explore all skills by this author →
GitHub

fla-ascend-performance

Profile, diagnose, and optimize Triton-Ascend NPU kernels in the flash-linear-attention repository.

5.7k|685|Updated Dec 20, 2023
Search Tags:#performance tuning#torch_npu#kernel optimization#ascend npu#Triton#Profiling#flash-linear-attention
One-click install
npx skills add https://github.com/fla-org/flash-linear-attention --skill fla-ascend-performance
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: fla-ascend-performance
Source: https://github.com/fla-org/flash-linear-attention/tree/main/.agents/skills/fla-ascend-performance
Command: npx skills add https://github.com/fla-org/flash-linear-attention --skill fla-ascend-performance

SYSTEM DOCUMENTATION & REQUIREMENTS

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

What problem does it solve?

Optimizing Triton kernels on Ascend NPUs requires navigating profiler CSVs, pipe utilization metrics, UB memory budgets, and Ascend-specific compiler traps that differ sharply from CUDA Triton. This Skill provides a repeatable profile-diagnose-optimize-verify workflow so kernel performance work on the flash-linear-attention repo is systematic instead of ad hoc.

Core Features & Use Cases

  • Generic NPU profiling scripts: profile_npu.py wraps any workload() callable with torch_npu profiler collection (PipeUtilization, MemoryUB, L2Cache), and analyze_profile.py parses op_statistic and kernel_details CSVs to surface hotspots.
  • Bottleneck diagnosis tables: Maps profiler signals (Cube/MAC ratio, Vector ratio, MTE2/MTE3, scalar, UB bandwidth) to concrete bottleneck classes and matching optimization levers such as UB tiling, grid splits, fusion, and G_T_CONTIG gate loading.
  • Ascend-specific trap catalog: Documents int32 address overflow, tl.dot left-operand clobber, constexpr DMA-path splitting, make_block_ptr int32 offsets, and the prohibition on num_warps/num_stages, with per-kernel case notes.
  • Use Case: A developer notices a Gated DeltaNet backward kernel is slow on Ascend 910. They freeze the baseline, run the generic profiler, identify a stride-HV gather on gate g via high MTE2, apply the G_T_CONTIG host transpose pattern, and verify with the frozen pytest gate plus a synchronized benchmark.

Quick Start

Ask the assistant to profile and optimize a specific Triton-Ascend kernel in the fla repo, for example: profile the causal_conv1d backward kernel on NPU, diagnose its bottleneck, and optimize it following the fla-ascend-performance workflow.

Frequently Asked Questions about fla-ascend-performance

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

FAQPage Schema
How do I profile a Triton kernel on Ascend NPU with torch_npu?▼

Run the generic profile_npu.py script with an exec file that defines a workload() function, selecting one aic_metrics such as PipeUtilization per run. Then parse the output with analyze_profile.py to inspect op_statistic and kernel_details CSVs for the target kernel.

How do I diagnose Ascend NPU kernel bottlenecks from profiler CSVs?▼

Read kernel_details sorted by Duration, then map pipe columns to bottleneck classes: high aic_mac_ratio means Cube-bound, high aiv_vec_ratio means Vector-bound, high mte2/mte3 with low compute means memory-move-bound, and high scalar_ratio means scalar-bound. Each class has matching optimization levers like tiling, fusion, or contiguous loads.

Can I use num_warps or num_stages in Ascend Triton kernels?▼

No. Ascend Triton does not support num_warps or num_stages, so they must never appear in kernel launches, autotune configs, or wrappers. Tune performance through tile sizes, grid configuration, layout, fusion or splitting, and UB budget instead.

Why does my Ascend kernel give wrong results with no compile error?▼

A common cause is the Ascend tl.dot left-operand clobber, where tl.dot(lhs, rhs) overwrites lhs in UB, unlike CUDA Triton. Fix it by reloading the tile from global memory or copying with tile + 0.0 before the first dot that uses it as the left operand.

Why do strided gate loads slow down Ascend kernels?▼

Loading gate g laid out as [B, T, HV] along the time axis creates stride-HV gathers that Ascend MTE handles poorly, often 10x to 35x slower. The fix is a host-side transpose to [B, HV, T] with a G_T_CONTIG constexpr flag so kernel loads are stride-1.

How do I avoid int32 overflow in Triton-Ascend address math?▼

Cast runtime indices to int64 before multiplying by strides using tl.cast, for example tl.cast(i_t, tl.int64) * BT, since .to(tl.int64) fails on specialized constexpr arguments. Keep make_block_ptr offsets in int32 and reserve int64 for flattened pointer arithmetic.

Related Skills

Explore More Software Engineering Skills→
AscendAscend
31

simple-vector-triton-gpu-to-npu

Migrate Triton GPU kernels to Ascend NPU with 1D grid mapping.

Official
Advanced
hatch3rhatch3r
26

hatch3r-pr-creation

Create platform-native pull requests or merge requests following repository conventions.

Official
Advanced
0xmariowu0xmariowu
40

autosearch:reflective-search-loop

Coordinate multi-round research tasks with explicit loop state and budget checks.

Community
Advanced
xchang1121xchang1121
6

triton-cuda-optimization

Optimize Triton CUDA kernels with memory access and block configuration guidance.

Community
Advanced