triton-ascend-debugging

triton-ascend-debugging is a skill for Claude Code, Codex from mindspore-ai/akg. It costs 69 tokens per session (1,583 once invoked), scanned A, original, Apache-2.0.

A troubleshooting checklist for Triton kernels running on Ascend hardware. It covers compilation errors, runtime failures, incorrect results, memory access, launch settings, concurrency, and performance checks.

In plain words
What is it for?
Use it when a Triton Ascend kernel fails to compile, crashes, produces inaccurate output, or runs slowly.
Why use it?
It gives a systematic way to narrow down problems such as out-of-bounds access, invalid control flow, incorrect indexing, grid limits, and unsuitable memory use.

Skill for Claude CodeCodex

Written for no agent in particular: nothing here depends on one.

Good fit Use it when a Triton Ascend kernel fails to compile, crashes, produces inaccurate output, or runs slowly.

Compare 6 skills from other repositories ↓
Install with agentmods
npx agentmods add skills/mindspore-ai/akg/triton-ascend-debugging
Install

Getting it into your agent

One page per mod, every tool's command on it. A separate URL per tool would split the same page into five that compete with each other.

Any agent
npx skills add mindspore-ai/akg --skill triton-ascend-debugging
Clone the repo
git clone --depth 1 https://github.com/mindspore-ai/akg

Made for: Claude Code, Codex.

Wrote this? Show the measurements

A badge with what this costs and how it scanned, read live from this page, so it follows the numbers instead of freezing them. Markdown for a README, HTML for a documentation site or a project page.

agentmods badge for triton-ascend-debugging

README.md
[![agentmods](https://agentmods.dev/badge/skills/mindspore-ai/akg/triton-ascend-debugging/github.svg)](https://agentmods.dev/skills/mindspore-ai/akg/triton-ascend-debugging)
Your own site
<a href="https://agentmods.dev/skills/mindspore-ai/akg/triton-ascend-debugging"><img src="https://agentmods.dev/badge/skills/mindspore-ai/akg/triton-ascend-debugging/github.svg" alt="Measured on agentmods" height="20"></a>

Or the 80×15 button, for a site that already has a row of RSS and ATOM ones. Only the verdict fits; the numbers stay here.

agentmods 80×15 button for triton-ascend-debugging

Your own site · 80×15
<a href="https://agentmods.dev/skills/mindspore-ai/akg/triton-ascend-debugging"><img src="https://agentmods.dev/badge/skills/mindspore-ai/akg/triton-ascend-debugging.svg" alt="Reviewed on agentmods" width="80" height="20"></a>
Per session 69 Skills are progressive disclosure: only the name and description are preloaded; the body loads when the skill is used.
When invoked 1,583 The whole file, excluding the scripts and references it only reads on demand.
Security scan A 0 findings. A grade says what 26 rules found in the file — not that it is safe. Third-party audits
  • NVIDIA SkillSpector pass 7 Sept 2026
How audits are shown
Origin original No closer match found in the catalogue.
Token cost

What it costs to keep this loaded

Counted locally with the o200k_base tokenizer, which is exact for GPT models; Claude uses its own tokenizer and its counts differ. Treat this as one consistent yardstick across the catalogue rather than a bill. Prices are per million input tokens.

ModelPer sessionOnce invoked
Fable 5.1 $0.00069 $0.01583
Opus 5 $0.00034 $0.00792
Sonnet 5 $0.00014 $0.00317
Haiku 4.5 $0.00007 $0.00158

Measured 9d ago against content hash 4b00bc1d6486, method: parsed. Prices are Anthropic first-party input rates as of 2026-09-12, from the pricing page.

Security

Grade A, and why

triton-ascend-debugging scanned grade A with 0 findings against 26 rules in 11 categories — prompt injection, anti-refusal, data exfiltration, privilege escalation, supply chain, agent snooping, system-prompt leakage, SSRF and excessive agency — measured 9d ago.

A static scan of the body, not an audit. Every finding is printed with the line that produced it so you can judge whether it matters here. A mod is markdown that instructs an agent; that is exactly why what it instructs is worth reading.

Nothing flagged

None of the 26 patterns this scan looks for appear in this file: no shell pipes, no recursive deletes, no credential paths, no hidden text, no instruction-override or anti-refusal phrasing, no agent-config snooping. That is not a guarantee, it is the absence of the things that are checkable.

akg_agents/python/akg_agents/op/resources/skills/triton-ascend/fundamentals/triton-ascend-debugging/SKILL.md · 85 lines

What it actually says

调试与排查清单

完整调试清单

内存访问问题

  • 所有 load/store 是否都有 mask 或 boundary_check?
  • stride 参数设置是否正确?
  • 数组索引是否越界?

控制流检查

  • 是否误用了 return/break/continue?
  • 复杂条件是否用 mask 组合实现?
  • tl.constexpr 是否只在内核参数中使用?

Grid 与 Block 配置检查

  • Grid 总大小是否不超过 65535?
  • 对于大 shape 算子,是否采用了交错循环 for i in range(pid, total, core_num)
  • Grid 维度是否为 tuple 类型且不超过 3 维?

并发与原子操作检查

  • 并发写入是否使用了原子操作(tl.atomic_add 等)?
  • 原子操作是否必要(能否避免)?

切片与索引检查

  • 是否避免了Python风格的直接切片(如b[0]b[i:j])?
  • 是否对tl.arange生成的张量误用了tl.get_element
  • 切片操作是否使用了正确的API(tl.get_elementtl.extract_slice等)?

性能优化检查

  • 内存访问是否连续(避免跨步访问)?
  • 是否充分利用了块内并行?
  • 复杂算子是否考虑拆分为多个简单kernel?

禁止使用的语法(Ascend 后端)

禁止写法 替代方案
return / break / continue 使用 mask 控制流程
lambda 表达式 内联函数或 tl.where
链式布尔运算 a and b 分步计算 mask:m1 = ...; m2 = ...; m = m1 & m2
张量直接索引 tensor[i] tl.load(ptr + offset) / tl.store(ptr + offset, val)
Python 切片 b[0] / b[i:j] tl.get_element / tl.extract_slice / tl.insert_slice
tl.arange 结果用 get_element 直接计算索引值
while 循环 for i in range(MAX): if i < n:
range() 混用运行时变量和 constexpr 全 constexpr 的 range(0, N, BLOCK_K) + 循环体内运行时 if
tl.float16(scalar) scalar.to(tl.float16)
tl.constexpr 在 host 侧使用 仅在 kernel 参数中使用
if-else 中负偏移 tl.maximum(offset, 0)
复杂 tl.where 用于内存偏移 拆分为 if-else 静态分支

常见错误速查表

编译错误

错误类型 典型症状 常见原因 解决方案
UB/CBUF 溢出 ub overflow, requires X bits while 1572864 bits available BLOCK 尺寸过大或中间变量过多 缩小 BLOCK 尺寸;减少同时活跃的 tensor 数
HiVM vsel 错误 hivm.hir.vsel: Unsupported op for finding the root alloc 嵌套 mask + tl.where 组合过于复杂 用乘法替代 tl.where:a * mask.to(dtype)
内存越界访问 运行时错误、结果异常、随机崩溃 load/store缺少mask或boundary_check 添加正确的mask或boundary_check保护
Grid超限 编译失败或运行时错误 grid总大小超过65535 使用交错循环for i in range(pid, total, core_num)或连续分块处理
控制流错误 unsupported AST node type: Continue 使用了return/break/continue 改用 if-else 包裹逻辑
while循环错误 编译失败(Ascend后端) 使用了while循环 改用for + if替代:for i in range(MAX): if i < n:
constexpr 索引 ValueError('unsupported tensor index: constexpr[0]') 对 tl.sum 等返回的标量做 [0] 索引 直接使用标量结果,不要索引
切片语法错误 编译失败 使用了b[0]b[i:j]直接切片 使用tl.get_elementtl.extract_slice
tl.arange索引错误 编译失败 tl.arange结果使用get_element 直接计算索引值而非提取
类型转换错误 cast incompatible 隐式 cast 或使用tl.float16(scalar) .to(tl.float16) 或显式指定 acc dtype
constexpr误用 编译失败 在host侧使用tl.constexpr 仅在kernel参数中使用tl.constexpr
Stride设置错误 计算结果错误、数据错位 stride参数计算或传递错误 验证stride设置,检查tensor.stride()
数值不稳定 结果为NaN或Inf softmax/sqrt等操作溢出 减去最大值、检查非负、使用float32
数据竞争 结果不确定、每次运行不同 多program并发写入同一位置 使用tl.atomic_add等原子操作
BLOCK_SIZE过大 编译失败或运行时错误 BLOCK_SIZE超过65536或硬件限制 减小BLOCK_SIZE,使用循环处理
tl.where偏移计算 编译失败(Ascend后端) 在内存偏移中使用tl.where 改用if-else静态分支处理
性能低下 运行缓慢 内存访问不连续、切分不合理 优化内存布局、调整BLOCK_SIZE、使用block_ptr
运行时range边界崩溃 bishengIR crash range()的start/stop混用运行时变量和constexpr 改用全constexpr的range(0, N, BLOCK_K),循环体内用运行时if跳过
Changes

What this file has done since we first saw it

Hashed on every crawl. A supply-chain change to an agent config is a question of when, not whether, so the history is kept rather than the latest state alone.

  1. 9d ago First seen · 85 lines · 69 tokens per session scan A 4b00bc1d6486

Subscribe to this mod's changes

triton-ascend-debugging is a skill published in the GitHub repository mindspore-ai/akg (259 stars, last pushed 1mo ago), licensed Apache-2.0. It adds 69 tokens to every session and 1,583 once invoked, about $0.0003 per session on Opus 5. A static security scan graded it A with 0 findings. No closer match exists in the catalogue, so it is treated as the original; first seen 2026-09-03.

Related

Other skills, from other repositories

spark-optimization

Optimize Apache Spark jobs with partitioning, caching, shuffle optimization, and memory tuning. Use when improving Spark performance, debugging slow jobs, or scaling data processing pipelines.

wshobson/agents · 37 tokens

notebooklm

Install, authenticate, troubleshoot, and operate Gemini Notebook through the notebooklm-py CLI or typed async Python API. Use for notebook and source management, grounded chat and research, and artifact generation or download when the user mentions Gemini Notebook, notebooklm-py, the notebooklm CLI, or its Python API.…

teng-lin/notebooklm-py · 79 tokens

debug-inference

Debug inference clients that use an attached provider and its native endpoint, including hosted APIs and host-local Ollama, vLLM, SGLang, TRT-LLM, LM Studio, or NIM. Use for provider attachment, endpoint policy, credential substitution, topology, and migration from the removed inference.local endpoint. Trigger…

NVIDIA/OpenShell · 119 tokens

oh-my-posh

Install, configure, or troubleshoot Oh My Posh/ohmyposh: shell init, themes, segments, Nerd Font icons, and prompt setup on PowerShell, zsh, bash, or fish.

JanDeDobbeleer/oh-my-posh · 47 tokens

eagle3-triage

Triage a failed EAGLE3 pipeline run. Identifies which step failed (data synthesis, hidden state dump, training, or benchmark), diagnoses root cause from logs, and suggests fixes. Use when user reports an EAGLE3 pipeline failure or asks why a specific step failed. Also helps debug new model support issues.

NVIDIA/Model-Optimizer · 73 tokens

trulens-instrumentation

Instrument LLM apps with TruLens OTEL-based tracing - from setup to debugging and optimization.

truera/trulens · 25 tokens