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 agentmods add skills/rocm/aiter/flydsl-kernel-authoringnpx skills add ROCm/aiter --skill flydsl-kernel-authoringgit clone --depth 1 https://github.com/ROCm/aiterWhat 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 | $0.00066 | $0.11071 |
| Opus 5 | $0.00033 | $0.05535 |
| Sonnet 5 | $0.00013 | $0.02214 |
| Haiku 4.5 | $0.00007 | $0.01107 |
Grade C, and why
flydsl-kernel-authoring scanned grade C with 1 finding 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 3d 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.
Recursive force deletehighDestructive command
rm -rf with a variable or a broad path is one typo away from removing the wrong tree.
FLYDSL_RUNTIME_ENABLE_CACHE=0 python my_kernel.py # or: rm -rf ~/.flydsl/cache How it starts
The opening of the file, as written. The whole thing — 971 lines — stays where its author put it; the contents beside it link to each section on GitHub.
FlyDSL Kernel Authoring Skill
Overview
FlyDSL is a Python DSL and MLIR-based compiler for writing high-performance GPU kernels on AMD GPUs (MI300X/MI350). It provides explicit layout algebra for controlling data movement, tiling, and memory access patterns. The layout system is the core abstraction that distinguishes FlyDSL from Triton/Gluon.
Repository: /FlyDSL/ (installed in editable mode)
Target GPU: gfx942 (MI300X, CDNA3), gfx950 (MI350, CDNA4)
Python: 3.12, ROCm 7.2
Scope (read this first): This skill is the reference — the full layout-algebra API
surface, per-op tables, MFMA/copy-atom catalogs, environment variables, and an exhaustive
troubleshooting list. Reach for it to look something up while writing or reviewing kernel
code. If instead you want a guided, step-by-step procedure that turns a kernel requirement
into a finished, tested kernel (classify -> skeleton -> compute -> control flow -> test), use
the flydsl-tile-programming skill, which is the wizard companion to this reference. For
diagnosing a kernel that already compiles but produces NaN/inf/wrong results, use the
debug-flydsl-kernel skill. When modernizing or cleaning up an existing kernel —
replacing legacy/deprecated constructs (ArithValue, fx.Index, buffer_ops, raw
MLIR dialects, SmemPtr, per-tile *_atom_call, raw rocdl.mfma_*) with the current
fx.* surface, trimming comments/dead code, or applying the _run_compiled fast launch
path — use the flydsl-kernel-code-cleanup skill.
1. Architecture and Compilation
Pipeline
Python (@flyc.kernel/@flyc.jit)
-> AST Rewriting (for/if -> scf.for/scf.if)
-> MLIR Tracing (generates Fly dialect + gpu/arith/scf/memref/vector ops)
-> MlirCompiler.compile() (Fly -> ROCDL -> LLVM -> HSACO binary)
-> JITCFunction (ExecutionEngine wrapper)
Key Passes
Pipeline is built by RocmBackend._pipeline_parts() and split into three stages — see docs/architecture_guide.md §3 for the per-pass table. Highlights:
fly-rewrite-func-signature- Rewrite DSL types at function / SCF boundaries to packed LLVM structsfly-layout-lowering- Lower layout algebra (fly.crd2idx, partitions, divides) to arithmeticfly-convert-atom-call-to-ssa-form+fly-promote-regmem-to-vectorssa- Lift copy/MMA atom calls and register memory to vector SSAconvert-fly-to-rocdl- Fly ops -> ROCDL intrinsicsgpu-module-to-binary{format=fatbin}- Emit HSACO binary via LLVM AMDGPU backend
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.
- 3d ago First seen · 971 lines · 66 tokens per session scan C 320c7788e8f5
flydsl-kernel-authoring is a skill published in the GitHub repository ROCm/aiter (543 stars, last pushed 3d ago), licensed MIT. It adds 66 tokens to every session and 11,071 once invoked, about $0.0003 per session on Opus 5. A static security scan graded it C with 1 finding (recursive force delete). No closer match exists in the catalogue, so it is treated as the original; first seen 2026-08-30.
Other skills, from other repositories
platform-port
Guide porting FastLED to new MCU platforms, including int.h types, clockless drivers, SPI implementations, and platform detection. Use when adding support for a new microcontroller family or board.
eide
EIDE (Embedded IDE) 工程构建工具，用于扫描 .eide/eide.yml 工程、枚举构建 配置 (ConfigName)、执行 build/rebuild/clean 并解析构建日志，返回可供 jlink/openocd 复用的产物路径。当用户提到 EIDE、Embedded IDE、eide.yml、 unifybuilder、VS Code EIDE 扩展、Cl.eide 时自动触发，也兼容 /eide 显式调用。 即使用户只是说"用 EIDE 编译一下"或"EIDE 烧录到板子上"，只要上下文涉及 EIDE 嵌入式工程就应触发此 skill。.
gcc
GCC 嵌入式工程构建工具（CMake + arm-none-eabi-gcc），用于扫描 CMake 型嵌入式工程、 列出预设、配置、编译、重建、清理和分析 ELF 大小。当用户提到 GCC、arm-none-eabi、 CMake 嵌入式编译、Ninja 构建、ELF 大小分析、arm-gcc、交叉编译、cmake --build、 cmake --preset 时自动触发，也兼容 /gcc 显式调用。即使用户只是说"编译一下"或 "看看固件多大"，只要上下文涉及 CMake 嵌入式 GCC 工程就应触发此 skill。.
tilelang-ascend-tile-api
TileLang-Ascend 新增 Ascend 专属 T.tile.xxx 小 API 的端到端开发流程。用户要求新增、封装、暴露、实现或测试 ascendtile.py 中的 T.tile API / Ascend tile primitive 时必须使用本 skill，尤其适用于需要同时打通 Python 前端、C++ lowering/codegen、Ascend C helper、文档和 CI 测试的任务。.
cpu-optimization-arm
ARM CPU 架构性能优化技巧、NEON SIMD 向量化、数值稳定性和调试策略.
pull_style_cutedsl_megakernel
This tree vendors the kernel team's SM90 FP8 MegaMoE drop — a fork of the same kernel repo that kernelsrc/cutedslmegamoe vendors (Bangyu's SM100 tree). The SM90 work (Vincent's hoppermegamoe branch) moved the shared runtime forward, so this tree duplicates common/, src/, and moenvfp4swapab/ at its own revision instead…