Tilelang Mixcv Skill

TileLang MixCV Skill

Mandatory routing rule

Before answering, follow AGENTS.md section "Docs Auto Routing Rules (Mandatory)".

Operator baseline rule (Mandatory)

• Before writing a new MixCV operator, first check examples/ and testing/npuir/.
• Prefer adapting an existing operator case rather than writing from scratch.

Focus

• mixed kernels with both Cube and Vector stages
• staged producer-consumer synchronization
• flash-attention-like patterns

Developer mode identification rule (Mandatory)

• In Developer mode, classify as MixCV when both conditions are true in the same kernel:
  1. Cube-side compute contains T.gemm.
  2. Vector-side compute contains at least one v-prefix op (for example T.vmul, T.vadd, T.vexp, T.vcast, T.vbrc).
• If both conditions hold, route to this skill even if the user does not explicitly say "mixcv".

Key primitives

• T.Scope("Cube") and T.Scope("Vector")
• T.rs("PIPE_FIX") and other pipe regions
• T.sync_block_set and T.sync_block_wait
• Pipelined loops where suitable

References

• references/pipeline.md
• references/flash-attn-pattern.md
• references/flash-attn-dev.md

Official docs to consult

• docs/Tilelang.language/同步管道操作/T.sync_block_set.md
• docs/Tilelang.language/同步管道操作/T.sync_block_wait.md
• docs/Tilelang.language/同步管道操作/T.pipe_barrier.md
• docs/Tilelang.language/线性代数操作/T.gemm.md
• docs/Tilelang.language/数学操作/T.vexp.md

Related skills

• tilelang-cube-skill
• tilelang-vector-skill
• tilelang-debug-helper