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.
npx skills add mindspore-ai/akg --skill triton-cuda-debugginggit clone --depth 1 https://github.com/mindspore-ai/akgWrote 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.
[](https://agentmods.dev/skills/mindspore-ai/akg/triton-cuda-debugging)<a href="https://agentmods.dev/skills/mindspore-ai/akg/triton-cuda-debugging"><img src="https://agentmods.dev/badge/skills/mindspore-ai/akg/triton-cuda-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.
<a href="https://agentmods.dev/skills/mindspore-ai/akg/triton-cuda-debugging"><img src="https://agentmods.dev/badge/skills/mindspore-ai/akg/triton-cuda-debugging.svg" alt="Reviewed on agentmods" width="80" height="20"></a>- NVIDIA SkillSpector pass
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.
| Model | Per session | Once invoked |
|---|---|---|
| Fable 5.1 | $0.00066 | $0.02279 |
| Opus 5 | $0.00033 | $0.01140 |
| Sonnet 5 | $0.00013 | $0.00456 |
| Haiku 4.5 | $0.00007 | $0.00228 |
Grade A, and why
triton-cuda-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 7d 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.
How it starts
The opening of the file, as written. The whole thing — 264 lines — stays where its author put it; the contents beside it link to each section on GitHub.
调试与排查清单
完整调试清单
内存访问问题
- 所有 load/store 是否都有 mask 或 boundary_check？
- stride 参数设置是否正确？
- 数组索引是否越界？
- 是否使用了
.contiguous()确保内存连续？ - 2D 数据是否使用了
tl.make_block_ptr？ - 内存访问是否合并（coalesced）？
控制流检查
- 是否误用了 return/break/continue？
- 复杂条件是否用 mask 组合实现？
-
tl.constexpr是否只在内核参数中使用？ - 是否有 lambda 表达式（不支持）？
Grid 与 Block 配置检查
- BLOCK_SIZE 是否为 2 的幂？
- num_warps 是否合理（2-8）？
- num_stages 是否合理（2-5）？
- Grid 总大小是否不超过硬件限制？
并发与原子操作检查
- 并发写入是否使用了原子操作（
tl.atomic_add等）？ - 原子操作是否必要（能否避免）？
- 是否有数据竞争（多个程序写同一位置）？
性能优化检查
- 是否使用了 autotune？
- MatMul 是否使用了 Grouped Ordering？
- 是否使用 float32 进行中间累加？
- Reduce 操作是否有数值稳定性处理？
常见错误速查表
编译错误
| 错误类型 | 典型症状 | 常见原因 | 解决方案 |
|---|---|---|---|
| Return 语句 | 编译失败 | Kernel 中使用 return | 移除 return，使用 mask 代替 |
| Break/Continue | 编译失败 | 不支持控制流跳转 | 用 mask 或重构逻辑 |
| Lambda 表达式 | 编译失败 | 不支持 lambda | 改用普通函数或内联 |
| 类型错误 | 编译失败 | constexpr 类型不匹配 | 检查 tl.constexpr 声明 |
运行时错误
| 错误类型 | 典型症状 | 常见原因 | 解决方案 |
|---|---|---|---|
| 内存越界 | CUDA error | 缺少 mask | 添加 mask 或 boundary_check |
| 形状不匹配 | 维度错误 | stride 计算错误 | 检查 stride 参数 |
| 非法内存访问 | Segfault | 指针计算错误 | 验证偏移计算 |
| 共享内存溢出 | Launch failed | num_stages 过大 | 减少 num_stages |
数值错误
| 错误类型 | 典型症状 | 常见原因 | 解决方案 |
|---|---|---|---|
| NaN/Inf | 结果异常 | Softmax 溢出 | 减去最大值 |
| 精度损失 | 结果不准确 | 全程使用 fp16 累加 | 使用 float32 累加 |
| 除零错误 | NaN | 方差或和为零 | 添加 eps |
| 负数开方 | NaN | 方差为负 | tl.maximum(var, 0.0) |
性能问题
| 问题类型 | 典型症状 | 常见原因 | 解决方案 |
|---|---|---|---|
| 性能差 | 比 PyTorch 慢 | 未使用 autotune | 添加 autotune |
| 带宽低 | 内存受限 | 非合并访问 | 确保合并访问 |
| Occupancy 低 | GPU 利用率低 | 寄存器/共享内存超限 | 减小 BLOCK_SIZE |
| L2 缓存差 | MatMul 性能低 | 未使用 Grouped Ordering | 添加 L2 缓存优化 |
分类调试流程
1. 编译失败
步骤:
- 检查错误信息中的关键词（return, break, lambda）
- 查看是否使用了不支持的语法
- 参考"API 使用限制"部分修改代码
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.
- 7d ago First seen · 264 lines · 66 tokens per session scan A 570e05200cc7
triton-cuda-debugging is a skill published in the GitHub repository mindspore-ai/akg (259 stars, last pushed 1mo ago), licensed Apache-2.0. It adds 66 tokens to every session and 2,279 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.
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.
debug-inference
Debug why inference.local, direct external inference, or supervisor-only system inference is failing. Use when the user cannot reach a local model server, has provider base URL issues, sees inference verification failures, hits protocol mismatches, or needs to diagnose inference on local vs remote gateways. Trigger…
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.…
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.
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.
trulens-instrumentation
Instrument LLM apps with TruLens OTEL-based tracing - from setup to debugging and optimization.