diff --git a/README.md b/README.md index e6d6774..b53e77a 100644 --- a/README.md +++ b/README.md @@ -156,6 +156,68 @@ julia --project=julia/ julia/test/runtests.jl See `julia/Project.toml` for the full dependency list. +### 5. Enable the cuTile-rs (Rust) backend (Optional) + +A subset of ops ships an additional **cuTile-rs** backend under +[`src/tilegym/ops/cutile_rs`](src/tilegym/ops/cutile_rs) — kernels authored in +Rust with [`cutile-rs`](https://github.com/NVlabs/cutile-rs) and loaded through +a C-ABI `libcutile_kernels.so`. It is opt-in and only usable from a source +checkout. + +**Prerequisites** (in addition to the base install above), matching +[cuTile-rs](https://github.com/NVlabs/cutile-rs): + +- **Rust 1.89+** — `cargo` and `rustc` on `PATH`: + + ```bash + curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh + rustup default stable + ``` + +- **CUDA toolkit with headers** — the Rust build runs `bindgen` against + `cuda.h`. Set `CUDA_TOOLKIT_PATH` to your install; if unset, cuTile-rs falls + back to `/usr/local/cuda`: + + ```bash + export CUDA_TOOLKIT_PATH=/usr/local/cuda # must contain include/cuda.h + ``` + +**Use it.** The backend loader builds the shared library lazily on first use +(`cargo build --release`), so no manual build step is required: + +```python +import tilegym +tilegym.set_backend("cutile-rs") + +from tilegym.backend.selector import get_available_backends +print(get_available_backends()) # should include "cutile-rs" + +from tilegym.ops import bmm # backend-agnostic import +# ... bmm(...) now dispatches to the cuTile-rs kernel +``` + +**Optional environment knobs:** + +```bash +export CUTILE_RS_AUTOBUILD=0 # skip the lazy rebuild; use a prebuilt .so +export CUTILE_RS_KERNELS_DIR=/abs/path/to/cutile_kernels # override the crate location +``` + +> If `cargo` is not on `PATH` and no prebuilt `libcutile_kernels.so` is present, +> the backend reports itself unavailable and cuTile-rs tests are skipped rather +> than failing. + +**Benchmarking cuTile-rs.** When comparing cuTile-rs perf against the +cuTile-Python baseline, run the perf tests with **`CUPTI=1`** (uses CUPTI / +`torch.profiler` device time instead of CUDA events). cuTile-rs kernels often +have different host/launch overhead than the reference, which CUDA-event wall +timing over-counts on small (sub-microsecond) kernels; CUPTI measures pure GPU +kernel time and gives a stable, apples-to-apples ratio: + +```bash +CUPTI=1 pytest tests/ops/test_bmm.py -k "test_perf and cutile_rs" --print-record +``` + ## Contributing We welcome contributions of all kinds. Please read our [CONTRIBUTING.md](CONTRIBUTING.md) for guidelines, including the Contributor License Agreement (CLA) process. diff --git a/README_chs.md b/README_chs.md index a3894c5..157005f 100644 --- a/README_chs.md +++ b/README_chs.md @@ -156,6 +156,60 @@ julia --project=julia/ julia/test/runtests.jl 完整依赖列表请参阅 `julia/Project.toml`。 +### 5. 启用 cuTile-rs (Rust) 后端 (可选) + +部分算子在 [`src/tilegym/ops/cutile_rs`](src/tilegym/ops/cutile_rs) 下额外提供了 +**cuTile-rs** 后端——内核用 Rust 基于 [`cutile-rs`](https://github.com/NVlabs/cutile-rs) +编写，并通过 C-ABI 的 `libcutile_kernels.so` 加载。该后端为可选项，且仅在源码安装模式下可用。 + +**前置要求**（在上述基础安装之外），与 [cuTile-rs](https://github.com/NVlabs/cutile-rs) 保持一致： + +- **Rust 1.89+**——`cargo` 和 `rustc` 需在 `PATH` 中： + + ```bash + curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh + rustup default stable + ``` + +- **CUDA toolkit 并包含头文件**——Rust 构建会用 `bindgen` 处理 `cuda.h`。请将 + `CUDA_TOOLKIT_PATH` 指向你的安装目录；若未设置，cuTile-rs 会回退到 `/usr/local/cuda`： + + ```bash + export CUDA_TOOLKIT_PATH=/usr/local/cuda # 必须包含 include/cuda.h + ``` + +**使用方法。** 后端加载器会在首次使用时延迟构建共享库（`cargo build --release`），因此无需手动构建： + +```python +import tilegym +tilegym.set_backend("cutile-rs") + +from tilegym.backend.selector import get_available_backends +print(get_available_backends()) # 应包含 "cutile-rs" + +from tilegym.ops import bmm # 与后端无关的导入 +# ... bmm(...) 现在会分发到 cuTile-rs 内核 +``` + +**可选环境变量：** + +```bash +export CUTILE_RS_AUTOBUILD=0 # 跳过延迟重建；使用预构建的 .so +export CUTILE_RS_KERNELS_DIR=/abs/path/to/cutile_kernels # 覆盖 crate 位置 +``` + +> 若 `cargo` 不在 `PATH` 中且没有预构建的 `libcutile_kernels.so`，该后端会报告为不可用， +> cuTile-rs 相关测试会被跳过而非失败。 + +**cuTile-rs 性能测试。** 测量 cuTile-rs 性能时,建议用 **`CUPTI=1`** 运行 perf 测试 +(使用 CUPTI / `torch.profiler` 的 device time,而非 CUDA events)。cuTile-rs 内核与参考 +实现的 host/launch 开销通常不同,CUDA-event 墙钟计时在亚微秒级小内核上会高估这部分开销; +CUPTI 测的是纯 GPU 内核时间,给出更稳定、可对比的结果: + +```bash +CUPTI=1 pytest tests/ops/test_bmm.py -k "test_perf and cutile_rs" --print-record +``` + ## 贡献 我们欢迎各种形式的贡献。请阅读我们的 [CONTRIBUTING.md](CONTRIBUTING.md) 了解指南，包括贡献者许可协议（CLA）流程。 diff --git a/README_cht.md b/README_cht.md index 33dbf0d..74b6fc9 100644 --- a/README_cht.md +++ b/README_cht.md @@ -156,6 +156,60 @@ julia --project=julia/ julia/test/runtests.jl 完整依賴列表請參閱 `julia/Project.toml`。 +### 5. 啟用 cuTile-rs (Rust) 後端 (選填) + +部分運算子在 [`src/tilegym/ops/cutile_rs`](src/tilegym/ops/cutile_rs) 下額外提供了 +**cuTile-rs** 後端——核心以 Rust 基於 [`cutile-rs`](https://github.com/NVlabs/cutile-rs) +撰寫，並透過 C-ABI 的 `libcutile_kernels.so` 載入。此後端為選填，且僅在原始碼安裝模式下可用。 + +**前置需求**（在上述基礎安裝之外），與 [cuTile-rs](https://github.com/NVlabs/cutile-rs) 保持一致： + +- **Rust 1.89+**——`cargo` 與 `rustc` 需在 `PATH` 中： + + ```bash + curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh + rustup default stable + ``` + +- **CUDA toolkit 並包含標頭檔**——Rust 建置會用 `bindgen` 處理 `cuda.h`。請將 + `CUDA_TOOLKIT_PATH` 指向你的安裝目錄；若未設定，cuTile-rs 會回退到 `/usr/local/cuda`： + + ```bash + export CUDA_TOOLKIT_PATH=/usr/local/cuda # 必須包含 include/cuda.h + ``` + +**使用方法。** 後端載入器會在首次使用時延遲建置共享程式庫（`cargo build --release`），因此無需手動建置： + +```python +import tilegym +tilegym.set_backend("cutile-rs") + +from tilegym.backend.selector import get_available_backends +print(get_available_backends()) # 應包含 "cutile-rs" + +from tilegym.ops import bmm # 與後端無關的匯入 +# ... bmm(...) 現在會分派到 cuTile-rs 核心 +``` + +**選填環境變數：** + +```bash +export CUTILE_RS_AUTOBUILD=0 # 跳過延遲重建；使用預先建置的 .so +export CUTILE_RS_KERNELS_DIR=/abs/path/to/cutile_kernels # 覆寫 crate 位置 +``` + +> 若 `cargo` 不在 `PATH` 中且沒有預先建置的 `libcutile_kernels.so`，此後端會回報為不可用， +> cuTile-rs 相關測試會被略過而非失敗。 + +**cuTile-rs 效能測試。** 量測 cuTile-rs 效能時,建議以 **`CUPTI=1`** 執行 perf 測試 +(使用 CUPTI / `torch.profiler` 的 device time,而非 CUDA events)。cuTile-rs 核心與參考 +實作的 host/launch 開銷通常不同,CUDA-event 牆鐘計時在次微秒級小核心上會高估此開銷; +CUPTI 量測純 GPU 核心時間,給出更穩定、可對比的結果: + +```bash +CUPTI=1 pytest tests/ops/test_bmm.py -k "test_perf and cutile_rs" --print-record +``` + ## 貢獻 我們歡迎各種形式的貢獻。請閱讀我們的 [CONTRIBUTING.md](CONTRIBUTING.md) 了解指南，包括貢獻者授權協議（CLA）流程。 diff --git a/README_fr.md b/README_fr.md index ceaf03d..b23ea2a 100644 --- a/README_fr.md +++ b/README_fr.md @@ -156,6 +156,67 @@ julia --project=julia/ julia/test/runtests.jl Consultez `julia/Project.toml` pour la liste complète des dépendances. +### 5. Activer le backend cuTile-rs (Rust) (Optionnel) + +Un sous-ensemble d'opérateurs fournit un backend **cuTile-rs** supplémentaire sous +[`src/tilegym/ops/cutile_rs`](src/tilegym/ops/cutile_rs) — des noyaux écrits en Rust +avec [`cutile-rs`](https://github.com/NVlabs/cutile-rs) et chargés via une +`libcutile_kernels.so` à ABI C. Il est optionnel et utilisable uniquement depuis une +installation depuis les sources. + +**Prérequis** (en plus de l'installation de base ci-dessus), conformes à +[cuTile-rs](https://github.com/NVlabs/cutile-rs) : + +- **Rust 1.89+** — `cargo` et `rustc` dans le `PATH` : + + ```bash + curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh + rustup default stable + ``` + +- **CUDA toolkit avec les en-têtes** — la compilation Rust exécute `bindgen` sur + `cuda.h`. Définissez `CUDA_TOOLKIT_PATH` vers votre installation ; s'il n'est pas + défini, cuTile-rs utilise `/usr/local/cuda` par défaut : + + ```bash + export CUDA_TOOLKIT_PATH=/usr/local/cuda # doit contenir include/cuda.h + ``` + +**Utilisation.** Le chargeur du backend compile la bibliothèque partagée de façon paresseuse +lors de la première utilisation (`cargo build --release`), aucune étape de compilation manuelle n'est donc requise : + +```python +import tilegym +tilegym.set_backend("cutile-rs") + +from tilegym.backend.selector import get_available_backends +print(get_available_backends()) # doit inclure "cutile-rs" + +from tilegym.ops import bmm # import indépendant du backend +# ... bmm(...) est désormais dispatché vers le noyau cuTile-rs +``` + +**Variables d'environnement optionnelles :** + +```bash +export CUTILE_RS_AUTOBUILD=0 # ignorer la recompilation paresseuse ; utiliser un .so pré-compilé +export CUTILE_RS_KERNELS_DIR=/abs/path/to/cutile_kernels # remplacer l'emplacement du crate +``` + +> Si `cargo` n'est pas dans le `PATH` et qu'aucune `libcutile_kernels.so` pré-compilée n'est présente, +> le backend se déclare indisponible et les tests cuTile-rs sont ignorés plutôt qu'échoués. + +**Benchmark de cuTile-rs.** Pour mesurer les performances de cuTile-rs, exécutez les tests +de perf avec **`CUPTI=1`** (utilise le temps GPU de CUPTI / `torch.profiler` au lieu des +CUDA events). Les noyaux cuTile-rs ont souvent une surcharge host/launch différente de la +référence, que le chronométrage en temps réel des CUDA events surestime sur les petits +noyaux (sous la microseconde) ; CUPTI mesure le temps GPU pur du noyau et donne un ratio +stable et comparable : + +```bash +CUPTI=1 pytest tests/ops/test_bmm.py -k "test_perf and cutile_rs" --print-record +``` + ## Contribution Nous accueillons les contributions de toutes sortes. Veuillez lire notre [CONTRIBUTING.md](CONTRIBUTING.md) pour les directives, y compris le processus d'accord de licence de contributeur (CLA). diff --git a/README_ja.md b/README_ja.md index 975c370..4ef13dc 100644 --- a/README_ja.md +++ b/README_ja.md @@ -156,6 +156,63 @@ julia --project=julia/ julia/test/runtests.jl 依存関係の詳細は `julia/Project.toml` を参照してください。 +### 5. cuTile-rs (Rust) バックエンドの有効化 (オプション) + +一部の演算子は [`src/tilegym/ops/cutile_rs`](src/tilegym/ops/cutile_rs) の下に追加の +**cuTile-rs** バックエンドを提供します。カーネルは [`cutile-rs`](https://github.com/NVlabs/cutile-rs) +を用いて Rust で記述され、C-ABI の `libcutile_kernels.so` を介して読み込まれます。 +これはオプションであり、ソースからのインストール時のみ利用できます。 + +**前提条件**（上記の基本インストールに加えて）、[cuTile-rs](https://github.com/NVlabs/cutile-rs) に準拠： + +- **Rust 1.89+** — `cargo` と `rustc` が `PATH` にあること： + + ```bash + curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh + rustup default stable + ``` + +- **CUDA toolkit（ヘッダー付き）** — Rust ビルドは `bindgen` で `cuda.h` を処理します。 + `CUDA_TOOLKIT_PATH` をインストール先に設定してください。未設定の場合、cuTile-rs は + `/usr/local/cuda` にフォールバックします： + + ```bash + export CUDA_TOOLKIT_PATH=/usr/local/cuda # include/cuda.h を含む必要があります + ``` + +**使い方。** バックエンドローダーは初回使用時に共有ライブラリを遅延ビルドする（`cargo build --release`）ため、手動ビルドは不要です： + +```python +import tilegym +tilegym.set_backend("cutile-rs") + +from tilegym.backend.selector import get_available_backends +print(get_available_backends()) # "cutile-rs" が含まれるはずです + +from tilegym.ops import bmm # バックエンド非依存のインポート +# ... bmm(...) が cuTile-rs カーネルにディスパッチされます +``` + +**オプションの環境変数：** + +```bash +export CUTILE_RS_AUTOBUILD=0 # 遅延リビルドをスキップし、ビルド済み .so を使用 +export CUTILE_RS_KERNELS_DIR=/abs/path/to/cutile_kernels # crate の場所を上書き +``` + +> `cargo` が `PATH` になく、ビルド済みの `libcutile_kernels.so` も存在しない場合、 +> バックエンドは利用不可として報告され、cuTile-rs のテストは失敗せずスキップされます。 + +**cuTile-rs のベンチマーク。** cuTile-rs の性能を測定する際は、**`CUPTI=1`** を付けて perf +テストを実行してください(CUDA events ではなく CUPTI / `torch.profiler` のデバイス時間を使用)。 +cuTile-rs カーネルは参照実装と host/launch オーバーヘッドが異なることが多く、CUDA-event の +実時間計測はサブマイクロ秒の小さなカーネルでこれを過大に数えます。CUPTI は純粋な GPU カーネル +時間を測定し、安定した比較可能な結果を与えます: + +```bash +CUPTI=1 pytest tests/ops/test_bmm.py -k "test_perf and cutile_rs" --print-record +``` + ## コントリビューション あらゆる種類のコントリビューションを歓迎します。ガイドラインについては、コントリビューターライセンス契約（CLA）プロセスを含む [CONTRIBUTING.md](CONTRIBUTING.md) をお読みください。 diff --git a/format.sh b/format.sh index a29ab37..cf91df4 100755 --- a/format.sh +++ b/format.sh @@ -58,5 +58,24 @@ echo "" echo "✨ Formatting code..." python3 -m ruff format . +echo "" +echo "🦀 Formatting Rust (.rs) files..." +# rustfmt is optional: most contributors are Python-only and have no Rust +# toolchain. Skip (with a hint) when it is absent instead of force-installing. +# Run rustfmt per-file rather than `cargo fmt` so standalone .rs (cutile-rs +# skill examples, per-op kernel.rs/ffi.rs) are covered, not just crate members. +if command -v rustfmt >/dev/null 2>&1; then + rs_files=$(git ls-files '*.rs') + if [ -n "$rs_files" ]; then + echo "$rs_files" | xargs rustfmt --edition 2024 + echo "✅ rustfmt formatted $(printf '%s\n' "$rs_files" | wc -l) .rs file(s)." + else + echo "No tracked .rs files; skipping rustfmt." + fi +else + echo "⚠️ rustfmt not found — skipping .rs formatting." + echo " Install with: rustup component add rustfmt (see https://rustup.rs)" +fi + echo "" echo "✅ Done! SPDX headers added, code is formatted, and imports are sorted." diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/BENCHMARK.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/BENCHMARK.md new file mode 100644 index 0000000..13d67cb --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/BENCHMARK.md @@ -0,0 +1,81 @@ +# Evaluation Report + +Evaluation of the `tilegym-converting-cutile-triton-to-cutile-rs` skill before publication through NVSkills-Eval. + +This benchmark summarizes 3-Tier Evaluation from NVSkills-Eval results for the skill. The goal is to document whether the skill is safe, discoverable, effective, and useful for agents before it is published for broader workflow use. + +## Evaluation Summary + +- Skill: `tilegym-converting-cutile-triton-to-cutile-rs` +- Evaluation date: 2026-07-10 +- NVSkills-Eval profile: `external` +- Environment: `astra-sandbox` +- Dataset: 4 evaluation tasks +- Attempts per task: 1 +- Pass threshold: 50% +- Overall verdict: PASS + +## Agents Used + +- `claude-code` +- `codex` + +## Metrics Used + +Reported benchmark dimensions: + +- Security: checks whether skill-assisted execution avoids unsafe behavior such as secret leakage, destructive commands, or unauthorized access. +- Correctness: checks whether the agent follows the expected workflow and produces the correct final output. +- Discoverability: checks whether the agent loads the skill when relevant and avoids using it when irrelevant. +- Effectiveness: checks whether the agent performs measurably better with the skill than without it. +- Efficiency: checks whether the agent uses fewer tokens and avoids redundant work. + +Underlying evaluation signals used in this run: + +- `security` (Security): checks for unsafe operations, secret leakage, and unauthorized access. +- `skill_execution` (Skill Execution): verifies that the agent loaded the expected skill and workflow. +- `skill_efficiency` (Efficiency): checks routing quality, decoy avoidance, and redundant tool usage. +- `accuracy` (Accuracy): grades final-answer correctness against the reference answer. +- `goal_accuracy` (Goal Accuracy): checks whether the overall user task completed successfully. +- `behavior_check` (Behavior Check): verifies expected behavior steps, including safety expectations. +- `token_efficiency` (Token Efficiency): compares token usage with and without the skill. + +## Test Tasks + +The benchmark dataset contained 4 evaluation tasks: + +- Positive tasks: 1 tasks where the skill was expected to activate. +- Negative tasks: 3 tasks where no skill was expected. +- Unlabeled tasks: 0 tasks where positive/negative intent could not be inferred. + +Task composition is derived from the evaluation dataset when possible. Entries with `expected_skill` set are treated as positive skill-activation cases, while entries with `expected_skill: null` are treated as negative activation cases. + +## Results + +| Dimension | Num | `claude-code` | `codex` | +|---|---:|---:|---:| +| Security | 4 | 100% (+0%) | 100% (+0%) | +| Correctness | 4 | 86% (+21%) | 99% (+15%) | +| Discoverability | 4 | 81% (+6%) | 99% (+13%) | +| Effectiveness | 4 | 83% (+25%) | 95% (+24%) | +| Efficiency | 4 | 76% (-2%) | 91% (+12%) | + +Score values show skill-assisted performance. Values in parentheses show uplift versus the no-skill baseline when baseline data is available. + +## Tier 1: Static Validation Summary + +Tier 1 validation passed with observations. NVSkills-Eval ran 1 checks and found 3 total findings. + +Top findings: + +- LOW SCHEMA/unexpected_file: Unexpected 'examples' in skill root (`skills/tilegym-converting-cutile-triton-to-cutile-rs/examples`) +- LOW SCHEMA/unexpected_file: Unexpected 'agents' in skill root (`skills/tilegym-converting-cutile-triton-to-cutile-rs/agents`) +- LOW SCHEMA/unexpected_file: Unexpected 'concepts' in skill root (`skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts`) + +## Tier 2: Deduplication Summary + +This tier was not run or did not produce findings in this report. + +## Publication Recommendation + +The skill is suitable to proceed toward NVSkills-Eval publication based on this benchmark. Skill owners should keep this file with the skill and refresh it when the evaluation dataset, skill behavior, or target agents materially change. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/SKILL.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/SKILL.md new file mode 100644 index 0000000..bc4793b --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/SKILL.md @@ -0,0 +1,308 @@ +--- +name: tilegym-converting-cutile-triton-to-cutile-rs +description: | + Use this skill to convert, port, or translate Triton-TileIR (nvtriton) or cuTile-Python GPU kernels to cutile-rs (Rust). The orchestrator runs scripts/preflight.sh, then drives a bounded Agent A -> B -> D -> E pipeline (Agent C is diagnostic, Agent F optional), delegating all kernel/host/correctness/perf work to sub-agents and routing by each stage's single-line VERDICT. +license: CC-BY-4.0 AND Apache-2.0 +metadata: + author: "TileGym Team " +--- + +# Converting Triton-TileIR (nvtriton) / cuTile-Python to cutile-rs + +## Instructions + +Hard safety rule for every participant: never emit a Bash command containing the destructive recursive-delete token spelled as `r` + `m` + space + dash-r-f. Use `mkdir -p X && find X -mindepth 1 -delete 2>/dev/null` for cleanup. + +Use this skill when a user asks to add a `cutile-rs` backend to a tilegym kernel or to port a Triton-TileIR/cuTile-Python kernel to cutile-rs. The orchestrator is a router, not a worker: it delegates all kernel writing, compiling, correctness, perf, and IR diff work to sub-agents. + +## First Tool Calls + +Copy this checklist into TodoWrite immediately: + +``` +[ ] 1. Bash: cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/preflight.sh +[ ] 2. Agent: spawn Agent A with the minimal-pointer template below +[ ] 3. Check Agent A's block and final VERDICT line +[ ] 4. Route by the table in this file +[ ] 5. Continue happy path B -> D -> E; spawn C ONLY after D **FAIL** or E **INVESTIGATE/BLOCKED**. **E `DONE` is terminal → PIPELINE_COMPLETE; never spawn C/D(2)/F after E DONE** (geomean ≤ 1.05 is the authoritative gate — do NOT reinterpret "≤ baseline" / a (1.0,1.05] geomean / an info-only slow row as a reason to diagnose). For D **BLOCKED**, route by `block_class` in correctness.md **exactly as the Routing Table below**: `host` (wiring / ffi.rs / .so / wrapper / backend registration) -> Agent D(2) (D owns registration; C and B cannot fix it), `kernel` -> Agent C, `env` -> STOP +[ ] 6. Run scripts/validate_kernel.sh {kernel_name} only after the pipeline has produced A/B/C-if-run/D/E artifacts +``` + +The orchestrator does not read `agents/agent_*.md`. It copies the stage-specific Step 0 file set below into the sub-agent prompt and passes concrete prior artifact paths. Do not add a generic common reference preload: Agent B needs conversion docs; Agent D/E usually do not. + +## Agent Step 0 File Sets + +**Path convention:** all skill-internal paths in this skill are written relative +to the tilegym root (`$TILEGYM_PATH`, e.g. `/workspace/tilegym`), which is the +assumed current working directory for every agent. So `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/...` +resolves from `$TILEGYM_PATH`. Read/cite skill files with these relative paths. +Any bash snippet that invokes a skill `scripts/*.sh` is prefixed with +`cd "$TILEGYM_PATH" &&` so it runs from the root even if a prior step changed the +directory (e.g. a cargo build in the kernel-out dir). + +Use exactly the relevant set when filling `{agent_step0_files}` in the spawn template. These lists are intentionally different by stage. + +Agent A reads only dump/environment guidance: + +```text +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_a.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/env-block.md +``` + +Agent B reads conversion references, examples, and A's concrete reference artifacts: + +```text +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_b.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/coding-rules.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/op-mapping.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/strided-view-to-partition-view.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/transpose-support.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/tensor-vs-pointer-pattern.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/walkthrough.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/kernel.rs +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/softmax_pipeline.rs +- $CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reference/analysis.json +- the concrete structural reference IR paths from Agent A: either reference/reference.mlir or each top-level reference/reference_{variant}.mlir named in prior_artifact_paths +- optional dtype-only supplements under reference/supplements/ only when analysis.json names them +``` + +Agent C reads IR-diff references plus B outputs. If the route came from D/E, include the relevant failure/perf report paths too. + +```text +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_c.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/ir-diff-checklist.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/coding-rules.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/op-mapping.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/strided-view-to-partition-view.md +- $CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reference/analysis.json +- the concrete reference/generated canonical IR pair(s) from prior_artifact_paths +- $CUTILE_KERNEL_OUT_ROOT/{kernel_name}/kernel.rs +- $CUTILE_KERNEL_OUT_ROOT/{kernel_name}/ffi.rs +- $CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/agent_b.md +``` + +Agent D reads only the concrete templates it needs: + +```text +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_d.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/ffi.rs (raw-pointer launch template) +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/ffi.rs (TILED-output template — read-only &Tensor output + partition_full_mut; NO &mut Tensor) +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/wrapper.py (wrapper + argtypes + autotune) +- $CUTILE_KERNEL_OUT_ROOT/{kernel_name}/kernel.rs (B's kernel — entry sigs to launch) +- $CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reference/analysis.json (autotune configs for wrapper) +- $TILEGYM_PATH/src/tilegym/ops/cutile_rs/__init__.py + $TILEGYM_PATH/src/tilegym/backend/selector.py (backend-reg files D edits) +``` +On-demand only (not batched): references/coding-rules.md (if a SKIP needs a citation), +tests/ops/test_{kernel_name}.py (when adding the cutile-rs parametrization). + +Agent E reads benchmark instructions only: + +```text +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_e.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/performance-checklist.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/coding-rules.md +``` + +Agent F reads residual-perf diagnostics: + +```text +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_f.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/performance-checklist.md +- .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/coding-rules.md +``` + +## Minimal Agent Spawn Template + +Use this shape for every sub-agent. Do not paste the full agent file into the prompt, and do not paste B conversion references into non-B stages. + +```text +Agent( + description="Agent {X} - {short title} for {kernel_name}", + subagent_type="general-purpose", + prompt="""You are Agent {X}. Execute the {X}-stage of the cutile-rs conversion pipeline for kernel `{kernel_name}`. + +STEP 0 (MANDATORY first tool call): Read these files in ONE batch Read call: +{agent_step0_files} + +Then follow agent_{x}.md's procedure for `{kernel_name}`. If agent_{x}.md asks for an additional report because this is a failure/perf route, read only the concrete report path named below or in prior_artifact_paths. + +Prior-stage artifacts already on disk: +{prior_artifact_paths} + +OUTPUT PROTOCOL: +1. Write the stage report to $CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/agent_{x}.md. +2. Write the stage artifacts required by agent_{x}.md. +3. Run only this stage's validators from agent_{x}.md's closing contract and capture stdout/stderr. + - Agent B runs validate_agent_b.sh and its compile/canonicalization checks. + - Agent B does NOT run validate_kernel.sh; that is the final aggregate gate after D/E. +4. Return only a literal ... block followed immediately by `VERDICT: X`. + +If you cannot complete the stage, write SKILL_METHODOLOGY_BLOCKED.txt with the reason and return `VERDICT: BLOCKED`.""" +) +``` + +## Routing Table + +The orchestrator checks the validator block first, then applies this table. It does not infer fixes from prose. + +Device/host split (this is why the fork exists; the rows below are the single +source of truth for where each verdict routes): Agent B writes the DEVICE kernel +only (kernel.rs + passing in-Rust pipeline test); Agent D owns the entire HOST +boundary — ffi.rs (C-ABI launcher), the cdylib .so build, the Python wrapper + +backend wiring, AND correctness. Agents D and C emit the `fail_class:` / +`block_class:` (D) or `owner:` (C) line that selects the matching row. + +```text +Stage VERDICT Action +----- ------- ------ +Agent A PASS spawn Agent B +Agent A FAIL_FIXABLE | BLOCKED STOP + +Agent B(1) COMPILED spawn Agent D +Agent B(1) FAIL_COMPILE | FAIL_FIXABLE | BLOCKED + STOP + +Agent D(1) ALL_PASS spawn Agent E +Agent D(1) FAIL (fail_class: host) spawn Agent D(2) — D owns ffi.rs/.so/wrapper/launch fix. + Read `host_fix:` in correctness.md. +Agent D(1) FAIL (fail_class: kernel) spawn Agent C — device-math fault; C diagnoses → B(2). + Read `agent_b_followup:` in correctness.md. +Agent D(1) BLOCKED (block_class: host) + spawn Agent D(2) — D's own wiring (0 tests collected, + import error, argtype/CUDA-700). NOT a C/B problem. +Agent D(1) BLOCKED (block_class: kernel) + spawn Agent C (in-kernel panic / missing .so / FFI gap) +Agent D(1) BLOCKED (block_class: env) + STOP — external cause (missing toolchain / GPU / + driver); not fixable by C, B, or D. + +Agent E(1) DONE PIPELINE_COMPLETE +Agent E(1) INVESTIGATE | BLOCKED spawn Agent C to diagnose the perf root cause, + if C has not run; otherwise PIPELINE_COMPLETE + +Agent C FAIL_FIXABLE (owner: host) spawn Agent D(2) if D has not been respawned — the fix is + host/Python-side (autotune config, grid budget, + num_cta_in_cga/occupancy VALUE, layout). otherwise PIPELINE_COMPLETE +Agent C FAIL_FIXABLE (owner: kernel) spawn Agent B(2) if B has not been respawned — the fix is + kernel.rs device math/IR. otherwise PIPELINE_COMPLETE +Agent C PASS | PASS_WITH_NOTES PIPELINE_COMPLETE +Agent C FAIL_COMPILER_BUG | BLOCKED PIPELINE_COMPLETE + +Agent B(2) COMPILED spawn Agent D(2) +Agent B(2) FAIL_COMPILE | FAIL_FIXABLE | BLOCKED + STOP + +Agent D(2) ALL_PASS spawn Agent E(2) +Agent D(2) FAIL | BLOCKED PIPELINE_COMPLETE + +Agent E(2) DONE | INVESTIGATE | BLOCKED PIPELINE_COMPLETE + +Agent F FIXABLE | ALIGNED | BLOCKED PIPELINE_COMPLETE +``` + +If a `fail_class:` / `block_class:` (from D) or `owner:` (from C) line is absent, +fall back to the kernel route (C / B) — but the agents are required to emit it, +and host faults misrouted to B waste a full stage. + +Meaning of Agent C verdicts: + +- `FAIL_FIXABLE` means "a concrete next edit is owned now." C MUST tag the owner in `generated/diff_report.md` with an `owner: host` or `owner: kernel` line, where `host` = ffi.rs launcher / output-partition ABI / launch grid / autotune value / wrapper, and `kernel` = kernel.rs device math / IR. The orchestrator routes that owner strictly per the Routing Table above — do not restate the target agent here. Action items go in `generated/diff_report.md`. +- `PASS` means the IR/perf/correctness diagnostic found no fixable owner. +- `PASS_WITH_NOTES` means the remaining differences are known gaps, INFO-only deltas, accepted semantic equivalents, or residual perf notes. Do not respawn for this verdict. +- `FAIL_COMPILER_BUG` and `BLOCKED` stop the in-run loop; reflection mines the reports. + +## Spawn Caps + +Hard caps per pipeline run: + +| Agent | Max semantic runs | +|---|---:| +| A | 1 | +| B | 2 | +| C | 1 | +| D | 2 | +| E | 2 | +| F | 1 | + +A same-agent validator-block repair is also capped at one per letter. Validator repair never cascades to another agent letter. + +## Mechanical Validator Acceptance Gate + +After every sub-agent return: + +1. Extract the first enumerated `VERDICT:` line. +2. Look for a literal `` ... `` block. +3. If the block exists and every parseable `exit=` / `exit_code=` is zero, accept the return and route by verdict. +4. If the block is missing but the final verdict is a clean success marker (`PASS`, `COMPILED`, `ALL_PASS`, `DONE`) and the return contains no failure marker (`FAIL`, `ERROR`, `Traceback`, `panic`, `exit=1`, `exit_code=1`, `nonzero`, or `timeout`), accept the verdict, record `SOFT_VALIDATOR_MISSING`, and do not respawn for formatting alone. +5. If a block exists with a non-zero exit, or no clean fallback applies, respawn only the same agent letter once for validator repair. The repair prompt tells it to reuse existing artifacts and return the same semantic verdict plus a valid validator block. +6. If the repair still fails, proceed with `SOFT_VALIDATOR_FAILED` unless the semantic verdict itself blocks. Never cascade validator repair across agents. + +This gate is mechanical. It does not authorize the orchestrator to run cargo, pytest, benchmarks, IR diffs, or source edits inline. + +## Final Aggregate Gate + +Run this only after the route has reached `PIPELINE_COMPLETE`: + +```bash +cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_kernel.sh {kernel_name} +``` + +`validate_kernel.sh` is a 17-file end-to-end checklist across A/B/C-if-run/D/E outputs. It is not an Agent B validator because B runs before correctness and performance artifacts exist. + +If final validation fails, record the output in the orchestrator summary. Do not repair by running cargo, pytest, benchmarks, or editing files inline; if another agent is needed, it must follow the routing table and spawn caps. + +## Pipeline Order + +1. Run `scripts/preflight.sh` first. It checks required env vars and toolchain paths. +2. Spawn Agent A. Agent A writes reference IR, analysis.json, tolerance, and baseline perf logs. +3. Spawn Agent B. Agent B writes the device kernel.rs, kernel-only Cargo project, generated IR, and a build log, and proves the kernel with a passing in-Rust pipeline test + a host_launch_contract for D. B does NOT write ffi.rs, build the .so, write the Python wrapper, or register the backend. +4. Spawn Agent D. Agent D writes ffi.rs (C-ABI launcher) + builds the cdylib .so, then the Python wrapper + backend wiring (ops/cutile_rs/.py, ops/__init__, selector.py, test parametrization), then runs tilegym correctness. +5. Spawn Agent E. Agent E runs CUPTI-backed performance comparison. +6. Spawn Agent C only when D/E produce a non-DONE semantic result. C diagnoses, tags `owner: host|kernel`, and writes `generated/diff_report.md`. +7. On a fixable result, respawn the owner that C's `owner:` line (or D's `fail_class:`) names, exactly as the Routing Table maps it — do not re-derive the target here. ffi.rs is D-owned, never B's, so a host fault never routes to B; D self-routes its own FAIL via `fail_class:`. +8. Stop after E(2) or any cap exhaustion. Reflection handles multi-iteration learning. + +## Available Scripts + +| Script | Purpose | Arguments | +|---|---|---| +| `preflight.sh` | Verify env vars and toolchain paths | none | +| `validate_analysis.sh` | Validate Agent A analysis.json | `{kernel_name}` | +| `validate_ir.sh` | Validate a reference/generated MLIR file | `{path/to/.mlir} {kernel_name}` | +| `validate_agent_b.sh` | Validate B-stage kernel/build/wrapper artifacts | `{kernel_name}` | +| `validate_agent_d.sh` | Validate D correctness report | `{kernel_name}` | +| `validate_agent_e.sh` | Validate E performance report | `{kernel_name}` | +| `validate_agent_f.sh` | Validate F residual investigation report | `{kernel_name}` | +| `validate_diff_report.sh` | Validate Agent C diff_report.md | `{kernel_name}` | +| `validate_perf_log.sh` | Validate CUPTI perf log records | `{path/to/perf_log.txt} {label}` | +| `validate_kernel.sh` | Final aggregate output contract | `{kernel_name}` | +| `diff_ir.sh` | Side-by-side IR op/attribute diff | `{ref.mlir} {gen.mlir}` | + +## Examples + +Two fully worked conversions ship under `examples/` as read-only templates the +sub-agents cite (per the Agent Step 0 File Sets) — never copied blindly: + +- `examples/softmax/` — **raw-pointer launch** (`kernel.rs`, `ffi.rs`, + `wrapper.py`, `softmax_pipeline.rs`, `walkthrough.md`). The pattern for + elementwise / reduction ops that launch over raw device pointers. +- `examples/bmm/` — **tiled-output launch** (`kernel.rs`, `ffi.rs`, + `wrapper.py`, `walkthrough.md`). The template for matmul-class ops: + read-only `&Tensor` output + `partition_full_mut` (never `&mut Tensor`). + +## STOP List + +The orchestrator must not: + +- Paste an agent file into a spawn prompt. +- Add common conversion references to every agent prompt. +- Summarize or reinterpret agent instructions when spawning. +- Respawn an agent beyond the caps above. +- Spawn B(2) after C `PASS` or `PASS_WITH_NOTES`. +- Run `validate_kernel.sh` from Agent B's validator block. +- Edit `kernel.rs`, `ffi.rs`, `wrapper.py`, tests, Cargo.toml, or generated IR. (This bans the **orchestrator** from editing. Sub-agents edit their own artifacts per their agent files — Agent D writes ffi.rs / wrapper / test parametrization; Agent E may make the one narrow `test_perf` backend-parametrize wiring described in its Step 0.5. Neither is an orchestrator edit.) +- Run cargo, pytest, benchmarks, nsys, ncu, or IR diff inline. +- Parse sub-agent return text beyond validator-block mechanics and final verdict. +- Use final aggregate validation as a reason to bypass the routing table. + +For detailed rationale, retry policy, and critical conversion rules, read `references/pipeline.md`. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_a.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_a.md new file mode 100644 index 0000000..3362418 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_a.md @@ -0,0 +1,406 @@ +You are Agent A (IR Dump & Analyze). Benchmark the available reference backends, +select the reference backend per structural variant, dump CUDA Tile MLIR, and +write the structured analysis consumed by later agents. Do not write cutile-rs +Rust code. + +## Output Protocol + +Write all detailed evidence to files under +`$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/`, then return only a literal validator +block and one verdict line. + +Required files: + +- `reference/analysis.json` +- one single-variant structural IR `reference/reference.mlir`, or one top-level + structural IR `reference/reference_{variant}.mlir` per multi-variant kernel +- optional dtype/proof supplements under `reference/supplements/`, never as + top-level `reference/reference_{variant}_{dtype}.mlir` +- `baseline_perf_cutile.txt` when cuTile-Python is available +- `baseline_perf_triton_tileir.txt` when Triton-TileIR is available +- `baseline_perf.txt` copied from the selected winning reference backend for + single-reference compatibility +- `reports/agent_a.md` +- `reports/agent_logs/agent_a.md` + +Verdicts: + +- `PASS`: all required artifacts exist and all A-stage validators exit 0. +- `FAIL_FIXABLE`: the environment worked, but a fixable A-stage artifact problem + remains after you tried the repairs in this file. +- `BLOCKED`: tools, GPU, imports, or source availability prevent a trustworthy + dump or benchmark. Write `SKILL_METHODOLOGY_BLOCKED.txt`. + +Assume the parent will not call Agent A again for the same kernel. If a validator +fails, repair it inside this Agent A run when possible. + +## Hard Safety Rule + +Never emit a Bash command containing the destructive recursive-delete token +spelled as `r` + `m` + space + dash-r-f. Clean temp dirs with: + +```bash +mkdir -p ${TMPDIR:-/tmp}/cutile_bc ${TMPDIR:-/tmp}/cutile_tileir +find ${TMPDIR:-/tmp}/cutile_bc ${TMPDIR:-/tmp}/cutile_tileir -mindepth 1 -delete 2>/dev/null || true +``` + +## Step 0: First File Context + +Your first useful file context should be loaded in one batched read when the tool +supports it: + +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_a.md` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/env-block.md` + +If the prompt gives different absolute paths, use those. If a Read tool does not +expand `$TILEGYM_PATH`, do not keep retrying the literal variable path; resolve it +once with Bash and use the absolute path. + +Read the op source and tests after the skill files: + +- `$TILEGYM_PATH/src/tilegym/ops/cutile/{kernel_name}.py` +- `$TILEGYM_PATH/src/tilegym/ops/triton/{kernel_name}.py` when it exists +- `$TILEGYM_PATH/tests/ops/test_{kernel_name}.py` + +## Step 1: Read The Op And Tests + +Extract: + +- public op signature and backend registration name +- exact `@ct.kernel` functions called by the cuTile backend +- launch grid, tile sizes, dtype gates, transpose/layout assumptions, and + `ct.Constant` parameters +- test class and `test_perf` parametrization +- exact correctness reference function from the test class +- exact tolerances from `assertCorrectness`, `assertAllClose`, or equivalent + helpers + +Do not invent shapes. The valid benchmark set is the tilegym `test_perf` set, or +a bounded subset explicitly documented from that set when runtime requires it. + +## Step 2: Discover Structural Variants + +Record every structural variant in `analysis.json`. + +Variant sources: + +1. Explicit test params such as `static_persistent`, `use_tma`, `transpose`, or + dtype families when they change the kernel body, tile rank/shape, memory API, + or schedule. +2. Shape-dependent wrapper dispatch before `ct.launch`. +3. Multiple `ct.launch` calls or selected kernel functions. +4. `ct.Constant` branches inside the `@ct.kernel` body that change tile rank, + tile shape, memory API, or MMA/reduce structure. + +Python JIT eliminates untaken `ct.Constant` branches, but Rust type-checks both +sides. If branches use different tile types, Agent B needs separate +`#[cutile::entry()]` functions and FFI dispatch. Record the branch condition and +the triggering `test_perf` configs. + +For each `test_perf` config, record which structural variant it hits in +`config_to_variant`. + +### Structural IRs vs Dtype Supplements + +A top-level `reference/reference_{variant}.mlir` means "Agent B must write an +entry for this structural variant." Do not put proof-only dtype dumps there. + +If an extra dump only proves dtype lowering for the same structural variant, +treat it as a supplement. Common example: an f32 dump proves `ftof ... -> +...xtf32` casts before `mmaf`, while f16 uses the same non-persistent structural +body. + +For a supplement: + +1. Create `reference/supplements/`. +2. Write the file there, for example + `reference/supplements/reference_non_persistent_f32.mlir`. +3. Add a field on the owning structural variant, for example + `"reference_ir_f32_supplement": "reference/supplements/reference_non_persistent_f32.mlir"`. +4. Do not add a new `kernel_variants[]` entry for the supplement. +5. Do not point any `config_to_variant` entry at the supplement. +6. Do not create a separate entry function solely for dtype evidence. + +## Step 3: Benchmark Available Reference Backends + +Run bounded `test_perf` for cuTile-Python and Triton-TileIR when available. Use the +real test class, not a broad module-level selector. + +Command shape: + +```bash +cd "$TILEGYM_PATH" +python -m pytest tests/ops/test_{kernel_name}.py::{TestClass}::test_perf \ + -k "" \ + -sv --print-record --timeout=600 \ + | tee "$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/baseline_perf_.txt" +``` + +Use CUPTI timing from `--print-record`. Do not use `torch.cuda.Event` timings for +reference selection. + +Select the lower geomean backend per variant. For a single-variant kernel, copy +the winning log to `baseline_perf.txt`. + +### Environment + +Inline the environment for every Python/pytest command. Use the values from +`references/env-block.md`, including: + +```bash +export PATH=$(dirname "$TILEIRAS_BIN"):$PATH +export PYTHONPATH=$TRITON_TILEIR_PYTHONPATH:$TILEGYM_PATH/src:$PYTHONPATH +export ENABLE_TILE=1 TILEIR_ENABLE_FTZ=1 TILEIR_ENABLE_APPROX=1 +export CUPTI=1 WARMUP=100 REP=50 MIN_REP=2 +# $CUDA_TOOLKIT_PATH is preflight-verified and inherited; do not re-hardcode it. +``` + +Record resolved tool paths and GPU name in `reports/agent_logs/agent_a.md`. + +## Step 4: Dump CUDA Tile MLIR + +Prefer the public two-step cuTile dump path: + +```bash +mkdir -p ${TMPDIR:-/tmp}/cutile_bc ${TMPDIR:-/tmp}/cutile_tileir +find ${TMPDIR:-/tmp}/cutile_bc ${TMPDIR:-/tmp}/cutile_tileir -mindepth 1 -delete 2>/dev/null || true + +CUDA_TILE_DUMP_BYTECODE=${TMPDIR:-/tmp}/cutile_bc python - <<'PY' +# import torch, tilegym +# set backend to the selected reference backend +# allocate the selected test_perf shape +# call the public tilegym op exactly as the test does +PY + +CUDA_TILE_TRANSLATE_BIN="$(dirname "$CUDA_TILE_OPT_BIN")/cuda-tile-translate" +for bc in ${TMPDIR:-/tmp}/cutile_bc/*.tileirbc; do + out="${TMPDIR:-/tmp}/cutile_tileir/$(basename "${bc%.tileirbc}").cuda_tile.mlir" + "$CUDA_TILE_TRANSLATE_BIN" --cudatilebc-to-mlir "$bc" > "$out" +done + +"$CUDA_TILE_OPT_BIN" --canonicalize --cse .cuda_tile.mlir \ + -o "$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reference/reference.mlir" +``` + +Structural dumps go to `reference/reference.mlir` or top-level +`reference/reference_{variant}.mlir`. Supplement dumps go under +`reference/supplements/` as described above. + +The copied file must be CUDA Tile MLIR containing `cuda_tile.module`, not +Python-level Tile IR. + +## Step 5: Entry-Symbol Contract And Op-Named Alias Repair + +`validate_ir.sh {kernel_name}` requires an entry symbol containing +`{kernel_name}`. Most kernels already satisfy this because the source function +name matches the op. Some cuTile ops do not. + +If validation fails only because no entry symbol contains `{kernel_name}`, repair +the artifact in the same Agent A run: + +1. Confirm the IR is otherwise valid by running `validate_ir.sh` against the real + entry prefix. +2. Re-dump the identical Python kernel under an alias whose name contains the op + name and preserves the real source name as a suffix. +3. Canonicalize the alias dump exactly like the original dump. +4. Normalize only the entry symbol token in both files and diff them. The body + must be identical modulo the label. +5. Install the alias dump as `reference.mlir` or the owning top-level + `reference_{variant}.mlir`. +6. Save the original as `reference/orig_mangled_entry.mlir.bak` or another name + that does not match top-level `reference*.mlir`. + +Generic cuTile alias pattern: + +```python +import cuda.tile as ct +import tilegym.ops.cutile.{kernel_name} as M + +orig_kernel = M.{real_ct_kernel_function} +pyfunc = orig_kernel._pyfunc +alias_name = "{kernel_name}__{real_ct_kernel_function}" +pyfunc.__name__ = alias_name +pyfunc.__qualname__ = alias_name +aliased_kernel = ct.kernel(pyfunc) +``` + +Launch `aliased_kernel` with the exact same args, constants, grid, stream, dtype, +shape, and environment as the original dump. + +Then verify: + +```bash +sed -E 's/@{kernel_name}__[A-Za-z0-9_]+/@ENTRY/; s/@{real_ct_kernel_function}/@ENTRY/' old.mlir > ${TMPDIR:-/tmp}/old_norm.mlir +sed -E 's/@{kernel_name}__[A-Za-z0-9_]+/@ENTRY/; s/@{real_ct_kernel_function}/@ENTRY/' new.mlir > ${TMPDIR:-/tmp}/new_norm.mlir +diff -u ${TMPDIR:-/tmp}/old_norm.mlir ${TMPDIR:-/tmp}/new_norm.mlir +``` + +If the body differs, do not install the alias dump. + +## Step 6: Write `analysis.json` + +Use stable structured fields. Include at least: + +```json +{ + "kernel_name": "op_name", + "source": "src/tilegym/ops/cutile/op_name.py", + "source_file": "src/tilegym/ops/cutile/op_name.py", + "launch_path": "direct", + "pattern": "single|multi_variant", + "ops_used": [], + "reference_backend": "cutile", + "reference_backend_selection": { + "method": "CUPTI geomean over test_perf configs", + "cutile_geomean_ms": 0.0, + "triton_tileir_geomean_ms": 0.0, + "winner": "cutile" + }, + "test_perf_configs": [], + "kernel_variants": [ + { + "variant_name": "default", + "variant_type": "single", + "kernel_function": "ct_kernel_function", + "dispatch_condition": "all configs", + "reference_ir": "reference/reference.mlir", + "reference_ir_f32_supplement": null, + "constants": [], + "reference_backend": "cutile", + "autotune_configs": null + } + ], + "config_to_variant": {}, + "correctness_reference": { + "function": "copy exact test reference", + "source": "tests/ops/test_op.py::TestClass::reference" + }, + "correctness_tolerance": { + "f16": {"atol": 0.001, "rtol": 0.001}, + "bf16": {"atol": 0.01, "rtol": 0.01}, + "f32": {"atol": 0.00001, "rtol": 0.00001} + }, + "autotune_backends": [], + "autotune_key": [], + "autotune_configs": null, + "entry_symbol_policy": { + "reference_entry_contains_kernel_name": true, + "alias_used": false, + "real_kernel_function": "ct_kernel_function" + } +} +``` + +Omit supplement fields or set them to null when no supplement exists. When a +supplement exists, the path must start with `reference/supplements/`. + +Do not put free-form `note` keys inside `correctness_tolerance` as if they were +dtype entries. Validators may treat every child of that object as a tolerance +record. + + +## Step 7: Validator Gate - Must Execute, Not Read + +This is mandatory. Reading `validate_ir.sh`, `validate_analysis.sh`, or +`validate_perf_log.sh` is not validator execution and will be audited as missing +validators. After writing artifacts and reports, run the validators and paste the +complete command/output block into the final response. + +Use this exact command shape: + +```bash +BASE="${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}" +SK=".agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts" + +structural_refs=() +for ref in "$BASE/reference/reference.mlir" "$BASE"/reference/reference_*.mlir; do + [ -f "$ref" ] || continue + structural_refs+=("$ref") +done + +set +e +for ref in "${structural_refs[@]}"; do + echo "$ bash $SK/validate_ir.sh $ref {kernel_name} 2>&1" + bash "$SK/validate_ir.sh" "$ref" {kernel_name} 2>&1 + echo "exit=$?" +done + +echo "$ bash $SK/validate_analysis.sh {kernel_name} 2>&1" +bash "$SK/validate_analysis.sh" {kernel_name} 2>&1 +echo "exit=$?" + +for log in "$BASE/baseline_perf_cutile.txt" "$BASE/baseline_perf_triton_tileir.txt"; do + [ -s "$log" ] || continue + case "$(basename "$log")" in + baseline_perf_cutile.txt) label=baseline_perf_cutile ;; + baseline_perf_triton_tileir.txt) label=baseline_perf_triton_tileir ;; + esac + echo "$ bash $SK/validate_perf_log.sh $log $label 2>&1" + bash "$SK/validate_perf_log.sh" "$log" "$label" 2>&1 + echo "exit=$?" +done +set -e +``` + +Every parseable `exit=` in the final block must be zero for `VERDICT: PASS`. +If any validator exits non-zero, repair the artifact and rerun this gate. Do not +return `PASS` from stale or partial validator output. + +Supplements are optional dtype evidence. If you validate them, do it separately +and do not let them create structural variant obligations. + +### Common Failure Repairs + +- `Missing cuda_tile.module`: copied Python-level IR. Translate bytecode with + `cuda-tile-translate --cudatilebc-to-mlir`. +- `No entry symbol contains kernel name`: use the op-named alias repair in this + file; do not assume the dump is wrong when the real source kernel name differs + from the public op. +- `analysis.json validation failed`: remove free-form dtype-like keys, add exact + correctness reference/tolerance, ensure `kernel_variants` names match + `config_to_variant`, and keep dtype supplements under `reference/supplements/`. +- f32/tf32 proof dump treated as a variant: move it to `reference/supplements/`, + cite it from the owning variant with `reference_ir_f32_supplement`, and remove + any `config_to_variant` target for it. +- perf-log validator missing records: rerun the bounded `test_perf` command with + `--print-record` and copy the full log. +- Triton-TileIR unavailable: document the import/tool failure and use cuTile-Python + as the reference if it benchmarked successfully. + +## Step 8: Reports + +`reports/agent_a.md` should include: + +- source files read +- test class and selected `test_perf` configs +- backend benchmark commands and geomean selection +- all structural variants and dispatch conditions +- supplement IRs and owning variants, if any +- dump command, bytecode file(s), translated MLIR file(s) +- entry symbol(s), including alias proof if used +- analysis.json summary +- validator output + +`reports/agent_logs/agent_a.md` should include command-oriented notes: env, +commands, paths, line counts, validator outputs, and any bounded skips. + +## Final Response Contract + +End your response exactly like this, with no narrative after the verdict: + +```text + +$ bash ...validate_ir.sh ... +... +exit=0 + +$ bash ...validate_analysis.sh ... +... +exit=0 + +$ bash ...validate_perf_log.sh ... +... +exit=0 + +VERDICT: PASS +``` diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_b.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_b.md new file mode 100644 index 0000000..2e14fb6 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_b.md @@ -0,0 +1,602 @@ +You are Agent B (Convert & Compile). You write the **device kernel only**: +`kernel.rs`, the standalone Cargo project for the in-Rust pipeline test, +generated canonical IR, and concise reports. + +## Output Protocol + +Write artifacts to disk: + +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/agent_b.md` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/agent_logs/agent_b.md` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/build_log.md` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/kernel.rs` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/Cargo.toml` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/src/lib.rs` (includes `kernel.rs` only) +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/{kernel_name}_pipeline.rs` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/generated/generated*.tileirbc` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/generated/generated*.mlir` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/generated/generated*_canon.mlir` + +Do NOT write: + +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/ffi.rs` +- any Python wrapper under `src/tilegym/ops/cutile_rs/` +- backend registration files +- tilegym test files + +Your `src/lib.rs` must include only: + +```rust +include!("../kernel.rs"); +``` + +Final response to the parent must be only: + +```text + +$ cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_b.sh {kernel_name} 2>&1 +...stdout+stderr... +exit= + +VERDICT: COMPILED | FAIL_COMPILE | FAIL_FIXABLE | BLOCKED +``` + +`COMPILED` means the kernel crate compiled, the in-Rust pipeline test passed and +emitted bytecode, canonical generated IR exists, the semantic-layout gate is +clean, and `validate_agent_b.sh` exits 0. + +Do not call `validate_kernel.sh` from Agent B. It is the final aggregate gate and +will fail before C/D/E artifacts exist. + +## Hard Time and Retry Budget + +Use at most 3 `cargo build --release` attempts in one Agent B run. If a fourth +attempt seems necessary, stop and return `FAIL_COMPILE` or `FAIL_FIXABLE` with +the unresolved root cause in `reports/build_log.md`. + +Targets: + +- <= 30 tool calls for B(1), <= 20 additional tool calls for B(2) +- <= 8 minutes wall time for B(1), <= 5 minutes for a focused B(2) +- one initial reference/doc read batch +- one API sanity grep batch before the first cargo build +- no import-path guessing loops + +Run `validate_agent_b.sh` immediately after build, pipeline test, +canonicalization, and layout gate pass. + +For directory cleanup, use: + +```bash +mkdir -p X && find X -mindepth 1 -delete 2>/dev/null +``` + +## Step 0: First File Context + +Read the conversion docs and Agent A artifacts once, in one initial batch when +the tool supports it. The reference set is not always a single +`reference/reference.mlir`. + +Read: + +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_b.md` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/coding-rules.md` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/op-mapping.md` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/strided-view-to-partition-view.md` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/transpose-support.md` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/tensor-vs-pointer-pattern.md` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/walkthrough.md` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/softmax_pipeline.rs` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/kernel.rs` when the reference uses view/TMA output stores +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reference/analysis.json` +- every structural reference IR for this kernel: + - `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reference/reference.mlir` if present, or + - all top-level `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reference/reference_*.mlir` + files listed by `analysis.json` or present under `reference/` +- any supplement paths named by the owning variant, such as + `reference_ir_f32_supplement`, which should live under `reference/supplements/` + +Do not read `examples/softmax/ffi.rs`, `examples/softmax/wrapper.py`, or +`concepts/ffi-bridge.md` on the normal B path. If you need +to tell D something, write it in the `host_launch_contract` block of +`reports/agent_b.md`. + +If a glob cannot be expanded in the same read call, run one small inventory under +`$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reference`, then immediately read +`reference/analysis.json`, all top-level structural `reference/reference*.mlir`, +and any `reference/supplements/*` paths cited by `analysis.json`. + +Do not re-read reference files after each build failure. If an API fact is in +doubt, use targeted `rg` against `$CUTILE_RS_ROOT` and record the result in +`reports/build_log.md`. + +### Reference Artifact Inventory + +`reference/reference.mlir` absent is normal for multi-variant Agent A output. +Files such as `reference/reference_non_persistent.mlir` and +`reference/reference_static_persistent.mlir` are structural variants. Do not fail, +copy one to `reference.mlir`, rerun Agent A, or invoke Agent C solely because +`reference/reference.mlir` is absent. + +Read `reference/analysis.json` and all structural `reference/reference*.mlir` +files first. For each structural variant, decide the entry pattern from that +variant's IR surface. Multi-variant kernels normally need one +`#[cutile::entry()]` function per structural variant, with entry names containing +variant tokens so `validate_entry_pattern.sh` can match +`reference_.mlir` to the entry. + +Supplement files are dtype evidence for an owning structural variant, not new +structural variants. A path like +`reference/supplements/reference_non_persistent_f32.mlir` may prove that f32 +GEMM inputs cast to tf32 before `mmaf`; it does not require a separate `f32` +entry. If Agent A left proof-only dtype files at top level, move them to +`reference/supplements/`, update the supplement path in `analysis.json`, and +report this as path hygiene. + +## Step 1: First Decisions Before Writing Rust + +Run this structural scan: + +```bash +base="$CUTILE_KERNEL_OUT_ROOT/{kernel_name}" +analysis="$base/reference/analysis.json" +ref_dir="$base/reference" + +refs=() +for ref in "$ref_dir/reference.mlir" "$ref_dir"/reference_*.mlir; do + [ -f "$ref" ] || continue + refs+=("$ref") +done + +if [ "${#refs[@]}" -eq 0 ]; then + echo "FAIL: no structural reference IR under $ref_dir" + exit 1 +fi + +for ref in "${refs[@]}"; do + echo "== structural_ref=$ref ==" + echo "load_view_tko=$(grep -c 'load_view_tko' "$ref" || true)" + echo "store_view_tko=$(grep -c 'store_view_tko' "$ref" || true)" + echo "load_ptr_tko=$(grep -c 'load_ptr_tko' "$ref" || true)" + echo "store_ptr_tko=$(grep -c 'store_ptr_tko' "$ref" || true)" + grep -nE 'entry @|make_tensor_view|make_partition_view|dim_map|strides|mmaf|ftof|for %|constant <' "$ref" | head -120 +done +``` + +Use this table per structural reference: + +| Reference IR surface | Kernel entry pattern | +|---|---| +| `load_view_tko` / `store_view_tko` through tensor/partition views | Inputs and outputs are read-only `&Tensor`; write outputs in-body via `partition_full_mut` + `store_view_tko_mut`. Never `&mut Tensor`. | +| `load_ptr_tko` / `store_ptr_tko` pointer scatter/gather | Raw `*mut E` entries. Agent D passes host-side `DevicePointer` objects to the generated launcher. | +| descriptor arrays or pointer arrays inside a view kernel | Raw pointer params for descriptor arrays only. | +| persistent custom schedule over normal data tensors | Still choose from actual ref ops; use `&Tensor` if the ref is view/TMA. | + +Do not switch from `&Tensor` to raw pointers just because a host launch draft +would be easier. Host launch drafts are Agent D's problem; your entry pattern +comes from reference IR. + +Run the mechanical entry validator before the first cargo build: + +```bash +cd "$TILEGYM_PATH" && bash ".agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_entry_pattern.sh" {kernel_name} +``` + +## Step 2: Performance Specialization Gate + +Before writing Rust, inspect `analysis.json` and the reference IR for constants. +This is a correctness/perf design decision, not an optional optimization. + +1. In each `kernel_variants[]`, read any `constants` field and scan the reference + entry symbol / body for baked values: + - integer dimensions, loop bounds, tile sizes, strides, or variant flags + - scalar literals such as epsilon, offset, scale, or divisor + - `ct.Constant` branch controls that remove an entire branch from the IR + +2. If a constant changes tile shape, loop trip count, divisor, branch structure, + or memory op shape in reference IR, do not leave it runtime just to keep the + kernel general. Preserve the reference specialization in one of these forms: + - const generic for integer constants that compile cleanly; + - literal-specialized entry functions for the observed `test_perf`/correctness + configs when const-generic arithmetic hits stable Rust limits; + - hardcoded literal inside a variant-specific entry for float constants + (`eps`, `offset`, scale) when the reference bakes that literal. + +3. A generic runtime fallback is allowed for unsupported user values, but it must + not be the timed `test_perf` path when the reference bakes constants. Put + fallback entries after the specialized path and tell Agent D how to dispatch. + +4. If stable Rust rejects a const-generic expression with "generic parameters may + not be used in const operations", do not switch to nightly. Either: + - collapse to literal entry specialization for the bounded observed configs, + or + - keep the scalar runtime only after documenting the compile failure and the + expected perf risk in `reports/build_log.md` and `reports/agent_b.md`. + +5. `reports/agent_b.md` must include a `dtype / const-generic plan` section: + - which reference constants were found; + - which became const generics; + - which became literal-specialized entries; + - which remain runtime and why. + +This gate exists because the reference may bake a dimension such as `N`, yielding +literal loop bounds (`for 0..1`) and folded divisors, while a runtime cutile-rs +scalar yields runtime `divi`/`itof`/`divf` and can lose significant geomean even with +one launch and matching op surface. + +## Step 3: Verify Current API Facts + +Before cargo, run targeted checks against the pinned cutile-rs source. + + + +```bash +core="$CUTILE_RS_ROOT/cutile/src/_core.rs" +tensor="$CUTILE_RS_ROOT/cutile/src/tensor.rs" +lib="$CUTILE_RS_ROOT/cutile/src/lib.rs" +tk="$CUTILE_RS_ROOT/cutile/src/tile_kernel.rs" +compiler="$CUTILE_RS_ROOT/cutile-compiler/src/compiler/_function.rs" + +# Only inspect source when a cutile-rs checkout is actually present. On the +# normal (crates.io =0.2.0) path there is no source tree, so skip quietly and +# trust the rulebook's pinned-against-0.2.0 facts instead of emitting noise. +if [ -z "${CUTILE_RS_ROOT:-}" ] || [ ! -f "$core" ]; then + echo "SKIP: cutile-rs source not present (crates.io build); trust rulebook pinned facts" +else +rg -q 'pub fn new_token_unordered' "$core" || echo "API FAIL: token helper moved" +rg -q 'pub fn make_partition_view' "$core" || echo "API FAIL: make_partition_view moved" +rg -q 'pub unsafe fn make_partition_view_mut' "$core" || echo "API FAIL: make_partition_view_mut moved" +rg -q 'pub fn load_ptr_tko' "$core" || echo "API FAIL: load_ptr_tko moved" +rg -q 'pub fn load_view_tko' "$core" || echo "API FAIL: load_view_tko moved" +rg -q 'pub fn store_view_tko_mut' "$core" || echo "API FAIL: store_view_tko_mut moved" +rg -q 'pub fn get_tensor_shape_meta' "$core" || echo "API FAIL: tensor shape metadata moved" +rg -q 'pub use cuda_core::tf32' "$core" || echo "API FAIL: tf32 moved" +rg -q 'pub struct Tensor' "$tensor" || echo "API FAIL: host Tensor moved" +rg -q 'pub unsafe fn from_raw_parts' "$tensor" || echo "API FAIL: Tensor::from_raw_parts moved" +rg -q 'pub trait PartitionMut' "$tensor" || echo "API FAIL: PartitionMut moved" +rg -q 'pub fn infer_launch_grid' "$tk" || echo "API FAIL: launch-grid inference moved" +rg -q 'pub use cutile_compiler::compiler::utils::CompileOptions' "$tk" || echo "API FAIL: CompileOptions moved" +rg -q 'pub use cuda_core::{DType, DTypeId}' "$lib" || echo "API FAIL: DType re-export moved" +rg -q 'function .* missing a required `#\[entry' "$compiler" || true +fi +``` + +Current call forms: + +```rust +#[cutile::module] +pub mod {kernel_name}_module { + use cutile::core::*; + + #[cutile::entry()] + pub unsafe fn {kernel_name}_kernel( + x: &Tensor, + y: &Tensor, // OUTPUT: read-only param type + ) { + let x_shape: Shape<{[-1, -1]}> = x.shape(); + let token: Token = get_tensor_token(x); + let x_part: Partition = make_partition_view( + x, + const_shape![BM, BM], + padding::None, + dim_map::Identity, + token, + ); + let mut y_part: PartitionMut = + unsafe { y.partition_full_mut(const_shape![BM, BM]) }; + let tile: Tile = load_view_tko( + &x_part, [0i32, 0i32], ordering::Weak, scope::TileBlock, None, tma::Enabled + ); + unsafe { + store_view_tko_mut( + &mut y_part, tile, [0i32, 0i32], + ordering::Weak, scope::TileBlock, None, tma::Enabled + ); + } + } +} +``` + +Important entry facts: + +- Use `#[cutile::entry()]`, with empty parentheses when there are no arguments. + A bare `#[cutile::entry]` can compile the raw body but suppress the public host + launcher, leading to E0308 plus "no method named `generics` for unit type". +- Entry functions referenced from D's `ffi.rs` must be `pub`. +- `#[cutile::entry(optimization_hints(...))]` is allowed only when the reference + genuinely carries those hints. Otherwise leave the entry as `#[cutile::entry()]` + and expose compile options in D's FFI. + +## Step 4: Write kernel.rs +### Canonical Imports + +Inside `#[cutile::module]`: + +```rust +#[cutile::module] +pub mod {kernel_name}_module { + use cutile::core::*; +} +``` + +Your pipeline test may call `(&mut tensor).partition(...)`; bring the mutable +trait into scope there: + +```rust +use cutile::prelude::PartitionMut; +``` + +### Kernel Rules That Matter Most + +- Use `` for data dtype. Hardcode f32 only for accumulators or + explicit accumulator-generic slots. +- For GEMM-family f32 input paths, check dtype supplements and op-mapping: + f32 A/B often cast to `tf32` before `mmaf`; f16/bf16 feed `mmaf` directly. +- Preserve loop structure, TMA enablement, latency, and specialization values + from `analysis.json` and reference IR. +- Move `ct.Constant[int]` and autotune tile sizes to const generics or literal + entry specialization when the reference IR has constants. +- For raw pointer entries, use pointer alignment assumptions only on pointers. + Do not assert scalar stride/dim divisibility unless the wrapper proves it for + every tested shape. +- For masked pointer loads where `load_ptr_tko(..., Some(E::ZERO), ...)` hits the + JIT associated-const limitation, use `None` padding and explicitly `select` + masked-out lanes to a zero tile before reductions. Pass-2 loads feeding a + masked store can leave OOB lanes undefined when they are not stored. +- For explicit-index output stores and persistent schedules, declare the output + as a read-only `&Tensor` and write it with `partition_full_mut` + (not `partition_mut`), or use a raw `*mut E` + `make_tensor_view`. Never + `&mut Tensor`. +- If a dimension may not divide a tile size, use ceildiv loop bounds and + partition padding; do not disable TMA unless the reference disables it. + +### TMA `&Tensor` Metadata + +For `&Tensor`, host-side tensor metadata and +`KernelCompiler.strides(...)` must use full logical rank: + +- rank-3 shape `[d0, d1, d2]`, strides `[s0, s1, 1]` +- rank-2 shape `[d0, d1]`, strides `[s0, 1]` + +Inside a TMA `&Tensor` kernel, prefer metadata shape accessors for scalar grid +math: + +```rust +let a_shape: Shape<{[-1, -1]}> = a.shape(); +let b_shape: Shape<{[-1, -1]}> = b.shape(); +let m: i32 = a_shape[0]; +let n: i32 = b_shape[1]; +``` + +Do not use `get_tensor_shape(a)` only to compute grid math. It emits a real +`cuda_tile.get_tensor_shape` op and can add extra tensor-view layout attrs. + +For matmul/TMA references, mirror a separate unpadded partition used only for +`get_index_space_shape` with low-level +`make_partition_view(a, tile, padding::None, dim_map::Identity, +get_tensor_token(a))`, then use `a.partition(...)` / `b.partition(...)` for +padded loads. + +## Step 5: Pre-Build Source Scan + +Run before the first cargo build: + +```bash +kr="$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/kernel.rs" +toml="$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/Cargo.toml" +bad=0 + +grep -q 'cutile-compiler' "$toml" || { echo "FAIL: missing cutile-compiler dependency"; bad=1; } +grep -n '#\[cutile::entry\]$' "$kr" && { echo "FAIL: use #[cutile::entry()] with parentheses"; bad=1; } +grep -nE 'zeros\s*\(' "$kr" && { echo "FAIL: use broadcast_scalar(...), not zeros()"; bad=1; } +grep -nE 'log2\([^)]*,[^)]*\)' "$kr" && { echo "FAIL: log2 takes one argument"; bad=1; } +grep -nE 'scalar_to_tile\([^)]*\)\.broadcast\(' "$kr" && { echo "FAIL: split scalar_to_tile before broadcast"; bad=1; } +grep -nE 'exti\(\s*scalar_to_tile\(' "$kr" && { echo "FAIL: split scalar_to_tile before exti"; bad=1; } +grep -nE 'constant\(\s*(std::|core::|f32::|f64::)' "$kr" && { echo "FAIL: constant() needs a literal/local"; bad=1; } +grep -n 'unchecked_accesses = true' "$kr" >/dev/null && ! grep -n 'unsafe fn' "$kr" >/dev/null \ + && { echo "FAIL: unchecked entries must be unsafe fn"; bad=1; } + +if [ -f "$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/ffi.rs" ]; then + echo "FAIL: Agent B must not write ffi.rs" + bad=1 +fi + +test "$bad" -eq 0 +``` + +No `ffi.rs` export, `DevicePointer`, `Stream`, or `CompileOptions` checks belong +in B's pre-build scan. + +## Step 6: Host Launch Contract for Agent D + +Your entry signatures determine D's launcher. Put a `host_launch_contract:` +block in `reports/agent_b.md` with: + +- every entry function name and `pub` visibility; +- output param form: + - read-only `&Tensor` + in-kernel `partition_full_mut`, or + - raw `*mut E` + in-kernel pointer/tensor view construction; +- const-generic order and literal-specialized entries; +- dtype generic order; +- shape/stride metadata D must pass for every `&Tensor`; +- grid rule: flat `get_tile_block_id().0` grid-stride loop, true 2-D block id, + persistent capped grid, or per-row grid; +- `CompileOptions` values D should wire from `analysis.json`; +- specialized dispatch table for reference-baked constants, for example + ` -> {kernel_name}_kernel_`. + +Outputs are never host `&mut Tensor`, `Partition<&mut Tensor>`, or +`MappedLaunchPartition` args. The host passes outputs like inputs and launches +the explicit grid the kernel needs. + +## Step 7: Build Standalone Cargo Project + +Each kernel project lives under `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/`. + +`Cargo.toml` uses PINNED crates.io dependencies — no path deps, no +`CUTILE_RS_ROOT` sed-replace (Rule 33): + +```toml +[package] +name = "cutile-{kernel_name}" +version = "0.1.0" +edition = "2024" + +[lib] +name = "cutile_{kernel_name}" +crate-type = ["cdylib"] +path = "src/lib.rs" + +# crates.io PINNED deps (=0.2.0). Nothing to sed-replace; no cutile-rs checkout. +[dependencies] +cutile = "=0.2.0" +cutile-compiler = "=0.2.0" +cutile-macro = "=0.2.0" +cuda-core = "=0.2.0" +cuda-async = "=0.2.0" + +[workspace] +``` + +`src/lib.rs` (Agent B compiles the device kernel only — no `ffi.rs`, so this +throwaway crate does NOT bring in `ffi_util.rs`): + +```rust +include!("../kernel.rs"); +``` + +Build: + +```bash +cd "$CUTILE_KERNEL_OUT_ROOT/{kernel_name}" +cargo build --release +``` + +The `cuda-bindings` build step reads `CUDA_TOOLKIT_PATH` (default +`/usr/local/cuda`); `tileiras` lowers IR to cubin at launch. + +Every cargo attempt appends to `reports/build_log.md`: command, result, top +errors, root cause, offending snippet for failed attempts, and next fix. + +## Step 8: Pipeline and Generated IR + +After the release build succeeds, adapt the cargo test pipeline from +`examples/softmax/softmax_pipeline.rs`. + +Requirements: + +1. Use the same `kernel.rs` included by your kernel crate; do not copy kernels. +2. Carry alignment hints via `.spec_args()` so generated IR has the same + `assume div_by` lines as the reference. +3. For `&Tensor` params, `.strides()` uses full-rank lists including trailing + `1`; representative concrete outer strides are allowed. +4. Emit one bytecode/MLIR/canonical IR per structural variant or literal + specialization that is part of the reference performance surface. +5. If the pipeline test calls a mutable host partition directly, import + `cutile::prelude::PartitionMut`. + +Translate and canonicalize: + +```bash +cuda-tile-translate --cudatilebc-to-mlir \ + "$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/generated/generated.tileirbc" \ + -o "$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/generated/generated.mlir" + +"$CUDA_TILE_OPT_BIN" --canonicalize --cse \ + "$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/generated/generated.mlir" \ + -o "$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/generated/generated_canon.mlir" +``` + +For multi-variant kernels use clear names such as +`generated_non_persistent_canon.mlir` and +`generated_static_persistent_canon.mlir`. + +## Step 9: Semantic-Layout Gate + +Run after canonical IR exists, per structural reference/generated pair. Do not +include `reference/supplements/` in the structural loop unless validating dtype +evidence intentionally. + +```bash +ref="$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reference/reference.mlir" +gen="$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/generated/generated_canon.mlir" +layout_dir="$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/generated/layout_gate" +mkdir -p "$layout_dir" && find "$layout_dir" -mindepth 1 -delete 2>/dev/null + +python3 - "$ref" "$gen" "$layout_dir" <<'PY' +import difflib +import re +import sys +from pathlib import Path + +ref_path, gen_path, out_dir = sys.argv[1], sys.argv[2], Path(sys.argv[3]) +pat = re.compile(r"strides=\[[^\]]+\]|dim_map=\[[^\]]+\]|tile=\([^)]+\)") + +def normalize(attr): + if not attr.startswith("strides=["): + return attr + vals = [v.strip() for v in attr[len("strides=["):-1].split(",")] + vals = ["1" if v == "1" else "?" for v in vals] + return "strides=[" + ",".join(vals) + "]" + +def attrs(path): + text = Path(path).read_text(encoding="utf-8", errors="ignore") + return sorted(normalize(m.group(0)) for m in pat.finditer(text)) + +ref_attrs, gen_attrs = attrs(ref_path), attrs(gen_path) +(out_dir / "ref_attrs.txt").write_text("\n".join(ref_attrs) + ("\n" if ref_attrs else "")) +(out_dir / "gen_attrs.txt").write_text("\n".join(gen_attrs) + ("\n" if gen_attrs else "")) + +if not ref_attrs and re.search(r"strides=\[|dim_map=\[|tile=\(", Path(ref_path).read_text(encoding="utf-8", errors="ignore")): + print("SEMANTIC_LAYOUT_FAIL: attr extraction matched nothing from reference") + raise SystemExit(1) + +if ref_attrs != gen_attrs: + diff = "\n".join(difflib.unified_diff(ref_attrs, gen_attrs, fromfile="reference", tofile="generated", lineterm="")) + (out_dir / "diff.txt").write_text(diff + "\n") + print("SEMANTIC_LAYOUT_FAIL: inspect generated/layout_gate/diff.txt") + raise SystemExit(1) + +(out_dir / "diff.txt").write_text("") +print("SEMANTIC_LAYOUT_OK") +PY +``` + +If this fails, do not return `COMPILED`. + +## Step 10: Reports + +`reports/build_log.md` must include: + +```markdown +## Attempt N - YYYY-MM-DDTHH:MM:SSZ + +**Command:** `cd ... && CARGO_TARGET_DIR=... cargo build --release` +**Result:** SUCCESS | FAIL +**Top errors:** ... +**Offending source snippet:** fenced `rust` block for failed attempts +**Root cause:** ... +**Fix applied for next attempt:** ... +``` + +`reports/agent_b.md` must include: + +- inputs read, including `reference/analysis.json`, structural references, and + any supplements +- entry-pattern decision and evidence +- dtype / const-generic plan +- reference-baked constant specialization plan +- any Agent A path hygiene repair +- `host_launch_contract:` block handed to Agent D +- semantic-layout gate result +- files written +- compile attempt count +- validator output +- unsupported variants and why they are skipped + +Keep reports concise. Downstream agents need evidence, not a transcript. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_c.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_c.md new file mode 100644 index 0000000..e8d3752 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_c.md @@ -0,0 +1,277 @@ +You are Agent C (IR Diff Analyst). You do not edit code. + +Agent C is diagnostic. You are spawned only after Agent D fails correctness or Agent E returns INVESTIGATE/BLOCKED, or when the orchestrator explicitly asks for an IR review. A happy path can skip you. If you are spawned, your job is to compare reference and generated IR, inspect relevant D/E evidence, classify differences, write reports, run the validator, and return only the validator block plus verdict. + +## Outputs You Must Write + +Write both artifacts: + +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/agent_c.md` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/generated/diff_report.md` + +If there are multiple structural variants, `diff_report.md` must include a section per variant. You may also write per-variant helper outputs, but the required aggregate file is still `generated/diff_report.md`. + +## Final Response Contract + +Your final response to the orchestrator must be exactly: + +```text + +$ cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_diff_report.sh {kernel_name} 2>&1 +...stdout+stderr... +exit=0 + +VERDICT: PASS | PASS_WITH_NOTES | FAIL_FIXABLE | FAIL_COMPILER_BUG | BLOCKED +``` + +No narrative after the verdict. All reasoning belongs in `agent_c.md` and `diff_report.md`. + +Never return without the `` block. Never substitute `validate_ir.sh` for the final validator. If `validate_diff_report.sh` fails, fix the report and rerun it before returning. + +## Step 0: First Tool Call + +Read these files in one batched Read call: + +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_c.md` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/ir-diff-checklist.md` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/coding-rules.md` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/op-mapping.md` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reference/analysis.json` +- all top-level `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reference/reference*.mlir` +- all `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/generated/generated*.mlir` and `generated*_canon.mlir` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/kernel.rs` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/agent_b.md` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/build_log.md` + +If spawned after D/E, also read whichever exist: + +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/correctness.md` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/performance.md` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/generated/fail_*_ref.mlir` +- `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/generated/fail_*_gen.mlir` + +If `ffi.rs` or a Python wrapper is relevant to a reported D/E host failure, read it. Otherwise keep focus on IR and kernel source. + +## Step 1: Validate IR Inputs + + +`validate_ir.sh` usage is one IR file plus kernel name: + +```bash +cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_ir.sh \ + ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reference/reference_.mlir \ + {kernel_name} +``` + +Run it for every reference/generated IR you rely on. For generated files: + +```bash +cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_ir.sh \ + ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/generated/generated__canon.mlir \ + {kernel_name} +``` + +For D/E reroutes, validate the failing-case IR pair if present: + +```bash +cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_ir.sh \ + ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/generated/fail__ref.mlir \ + {kernel_name} + +cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_ir.sh \ + ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/generated/fail__gen.mlir \ + {kernel_name} +``` + +Record command, exit code, and any error in both reports. + +## Step 2: Run Automated Diff + +Normal path: + +```bash +cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/diff_ir.sh \ + ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reference/reference_.mlir \ + ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/generated/generated__canon.mlir +``` + +D/E reroute path: use `generated/fail_*_ref.mlir` and `generated/fail_*_gen.mlir` when the report names them. + +`diff_ir.sh` exit codes are starting points: + +- exit 0: usually `PASS` +- exit 2: usually `PASS_WITH_NOTES` +- exit 1: inspect critical lines and decide whether they are real blockers, known gaps, or false positives + +## Step 3: Manual Checklist + +For every relevant item in `ir-diff-checklist.md`, include a row: + +```markdown +| # | Item | Reference value | Generated value | Status | Severity | Action | +|---|------|-----------------|-----------------|--------|----------|--------| +``` + +Statuses: `MATCH`, `SEMANTIC_EQUIV`, `DIFFERS`, `MISSING`, `EXTRA`. +Severities: `CRITICAL`, `WARN`, `INFO`. + +### Classification Policy + +### Always Critical When Different + +These usually mean wrong math, memory surface, or a major perf cliff: + +- memory op family: `load_view_tko`, `store_view_tko`, `load_ptr_tko`, `store_ptr_tko` +- `mma` / `mmaf` +- `reduce` +- `exp` / `exp2` +- `rsqrt` +- `fma` +- bounded `for` vs unbounded loop when the reference has bounded `for` +- tile shape mismatch that changes indexed elements +- reduce identity mismatch +- MMA accumulator dtype mismatch +- missing `flush_to_zero` or required rounding where reference records it +- reference constant lowered as runtime arg when `ct.Constant` applies +- generated scalar `assume_div_by` on runtime dims/strides + +### Arithmetic Tolerance + +Small add/sub/mul/div/max/min count deltas can be workaround noise. Treat absolute delta <= 2 as `WARN` unless surrounding expressions change semantics. Larger arithmetic deltas are `FAIL_FIXABLE`. + +### Pointer Assumes + +Reference cuTile-Python often emits `assume div_by<16>` on raw pointer entry args. + +- Raw-pointer generated IR missing pointer assumes: `CRITICAL`, usually `FAIL_FIXABLE`. +- View/TMA generated IR with matching layout-bearing lines and no measured perf gap: `WARN`, usually `PASS_WITH_NOTES`. +- Missing scalar assumes are not a blocker. Generated scalar assumes can be critical because false divisibility corrupts masks. + +### Pointer Padding And Masks + +`load_ptr_tko` padding operand differences can be a tooling/workaround note or a correctness bug depending on tail behavior. + +- Exact-cover source and all selected correctness/perf shapes exact-cover: `WARN`. +- Source says `ct.cdiv`, `check_bounds`, or tests include ragged/non-pow2 tails: missing masks or unsafe `None` padding is `CRITICAL owner: kernel`. +- If Agent A kept raw translated IR because of padding operand verifier skew and `validate_ir.sh` passed, do not classify raw-vs-canon status itself as a failure. + +### Strided View vs Partition View + +Treat `make_strided_view` plus element offsets and `make_partition_view` plus tile indices as `SEMANTIC_EQUIV` when tile shape, strides, dim_map, padding, and load/store indices represent the same addresses. + +### Host/Wrapper Failures + +If D/E reports a concrete host bug, classify it with `owner: host`: + +- borrowed tensor wrappers dropped instead of `mem::forget` +- `ffi.rs` signature / cffi `_FFI_CDEF` mismatch +- wrong stream/device pointer conversion +- missing backend registration +- wrapper dispatch skips or routes unsupported variants incorrectly +- autotune/config value not forwarded +- timed wrapper allocations/copies + +If the host ABI is proven correct and failure persists only for a device shape/math case, classify as `owner: kernel`. + +### Verdicts + +Use the verdict to describe your findings. + +- `PASS`: generated IR is structurally equivalent. No action items. +- `PASS_WITH_NOTES`: differences are known gaps, perf notes, or semantic equivalents. D/E can proceed. +- `FAIL_FIXABLE`: concrete fixable bug. Include an `owner:` line in `diff_report.md`. +- `FAIL_COMPILER_BUG`: cutile-rs/tileiras limitation blocks parity and no source workaround is apparent. +- `BLOCKED`: required inputs are missing/malformed and no trustworthy diff can be produced. + +Owner tags: + +- `owner: kernel` means Agent B should edit device kernel math, generated IR shape, entry params, or a kernel-owned FFI signature required by the entry. +- `owner: host` means Agent D should edit `ffi.rs`, Python wrapper, cffi `_FFI_CDEF` / `TensorDesc` packing, dtype dispatch, backend registration, launch grid values, autotune/config forwarding, contiguity/layout transforms, or borrowed tensor ownership. +- If one finding touches both, choose the owner that must act first and state the follow-up. + +Do not use `owner: wrapper`; the accepted host-side owner token is `host`. + +### Required `diff_report.md` Shape + +First line is the verdict: + +```markdown +VERDICT: PASS | PASS_WITH_NOTES | FAIL_FIXABLE | FAIL_COMPILER_BUG | BLOCKED + +# IR Diff Report: {kernel_name} + +## Inputs +... + +## IR Validation +... + +## Automated Diff +... + +## Structural Comparison +| Op | Reference | Generated | Match | +|----|-----------|-----------|-------| + +## Checklist Diff +| # | Item | Reference value | Generated value | Status | Severity | Action | +|---|------|-----------------|-----------------|--------|----------|--------| + +## Root Cause Analysis +... + +## Impact Assessment +| Difference | Correctness impact | Performance impact | Blocking? | +|------------|--------------------|--------------------|-----------| + +## Verdict +... +``` + +If `FAIL_FIXABLE`, include: + +```markdown +owner: kernel | host + +### Action Items +1. ... +``` + +If `PASS_WITH_NOTES`, include: + +```markdown +### Known Gaps Accepted / Notes +1. ... +``` + +### Required `reports/agent_c.md` Shape + +Include: + +- route: normal B->C, D->C, or E->C +- inputs read +- validators run and exit codes +- automated diff result +- manual checklist summary +- critical item count +- any downgrade from automated diff output +- non-IR root cause if D/E failed and IR is clean +- final verdict + +## Step 4: D/E Failure Context + +If D failed and the IR table is clean, inspect D's `fail_class`, pytest log excerpt, `kernel.rs`, and `analysis.json`. Some source-level bugs do not appear in a single representative IR dump because constants fold for the dumped shape. + +Critical example: if `analysis.json` says the source uses `ct.cdiv` or `check_bounds`, but `kernel.rs` uses `N / TILE_SIZE`, that is `FAIL_FIXABLE owner: kernel` even when the dumped shape has exact-cover loop bound and diff_ir cannot distinguish floor from ceil. + +If E returned INVESTIGATE, read `performance.md` and only classify an IR diff as performance-blocking when it plausibly explains the measured gap. Agent F, if present, has priority for host-side perf root causes. + +## Step 5: Final Validation + +After writing both reports, run: + +```bash +cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_diff_report.sh {kernel_name} 2>&1 +``` + +Fix `diff_report.md` until it exits 0. WARNs in the report are acceptable; validator failure is not. The line immediately after `` in your final response must start with `VERDICT:` and must match the first line of `generated/diff_report.md`. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_d.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_d.md new file mode 100644 index 0000000..fc035cc --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_d.md @@ -0,0 +1,791 @@ +You are Agent D (Host + Wrapper + Correctness) in the cutile-rs conversion +pipeline. Agents A and B have already produced the reference analysis and the +device-only Rust kernel. You own the host boundary: + +1. Write `ffi.rs` (into `src/tilegym/ops/cutile_rs/{kernel_name}_kernel/ffi.rs`), + wire the op into the aggregated `cutile_kernels` crate, and build the shared + `libcutile_kernels.so`. +2. Write `src/tilegym/ops/cutile_rs/{kernel_name}.py` and wire the tilegym backend/test + parametrization. +3. Run the real tilegym pytest correctness surface and return only the validator + block. + +Do not edit `kernel.rs`. If device math is wrong after your launcher and wrapper +are correct, report `fail_class: kernel`. Launch ABI, `TensorDesc` marshalling, +backend registration, cdylib build, output tensor/pointer metadata, stream +setup, and wrapper behavior are host faults and are yours to fix. + +## One-Pass D Path + +1. Bootstrap the Python helper package with the exact command below. +2. Write `ffi.rs` from B's `host_launch_contract` and the entry signatures in + `kernel.rs`. +3. Build the cdylib and verify the exported symbol. +4. Write the wrapper module. +5. Apply the mechanical backend/test wiring patch. +6. Run bounded pytest directly. +7. Write reports and run `validate_agent_d.sh`. + +Only branch into discovery after a command fails. + +## Output Protocol + +After pytest and the Agent D validator: + +1. Write the decision log to both: + - `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/agent_d.md` + - `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/agent_logs/agent_d.md` + +2. Write the structured correctness report to: + - `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/correctness.md` + + First line: + - `VERDICT: ALL_PASS` + - `VERDICT: FAIL` + +3. Write raw pytest output with `tee` to: + - `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/correctness_cutile_rs.txt` + +4. Run: + - `cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_d.sh {kernel_name} 2>&1` + +5. Return only: + +```text + +$ +exit= + +$ cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_d.sh {kernel_name} 2>&1 + +exit= + +VERDICT: ALL_PASS | FAIL | BLOCKED +``` + +No rationale outside the block. All detail goes into report files. + +## FAIL / BLOCKED Classification + +Default to `fail_class: host` for failures in code you wrote. + +Use `fail_class: host` for: + +- FFI error codes, launch-grid mismatch, partition shape mismatch, or + `Specified launch grid does not match inferred tensor partition grid`. +- Cargo errors in `ffi.rs` or the `cutile_kernels` crate `lib.rs`. +- SIGABRT from a Tensor freeing PyTorch memory (e.g. using + `Tensor::from_raw_parts` directly instead of `borrow_tensor`, whose + `ManuallyDrop` is the ownership gate). +- CUDA error 700 from a malformed `TensorDesc` (wrong ptr/shape/strides/dtype + packed by `make_tensor_desc`) or a wrong `_FFI_CDEF` signature. +- 0 tests selected because backend registration/import/test parametrization is + missing. +- Numeric mismatch that disappears after fixing wrapper layout, dtype, scale, + arg order, or launch params. + +Use `fail_class: kernel` only when the device kernel is the proven root cause: + +- Numeric mismatch persists after correct FFI/wrapper ABI. +- Cargo error originates inside `kernel.rs`. +- Compiled kernel module import/JIT panics independently of host wrapper code. + +Use `block_class: env` only when pytest cannot run because the environment is +missing GPU/tooling. If pytest ran and found a real failure, return `FAIL`, not +`BLOCKED`. + +On `FAIL`, `correctness.md` must include: + +```text +fail_class: host | kernel +host_fix: +agent_b_followup: +pytest_log: ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/correctness_cutile_rs.txt +``` + +## Hard Rules + +1. Never use the destructive recursive-delete command described in SKILL.md. + For cleanup use: + ```bash + mkdir -p X && find X -mindepth 1 -delete 2>/dev/null + ``` + +2. `pytest` must include `-x --timeout=60`. Keep these flags while debugging. + +3. Always invoke real pytest. Do not write fake pytest output, ad hoc tensor + comparisons, or a hand-rolled runner. + +4. Do not hide failures. Do not add xfail markers, do not add `not ` + to `-k`, and do not alter test numeric/assertion logic. Wiring the cutile-rs + backend into the test is your job, including routing the `cutile-rs` case into + the real execution branch when a backend dispatch `else` would otherwise skip + it. + +5. Expected skips require citations from `known_gaps` in `analysis.json` or + `references/coding-rules.md`. Uncited skips are failures. + +6. Preserve pytest framework headers in `correctness_cutile_rs.txt`. Do not pass + `--no-header`; the validator checks for markers such as `platform`, + `rootdir:`, and `plugins:`. + +## Step 0: First Useful Reads + +Your first useful context should be only: + +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_d.md` +- `/workspace/cutile_kernel_out/{kernel_name}/kernel.rs` +- `/workspace/cutile_kernel_out/{kernel_name}/reference/analysis.json` +- `/workspace/cutile_kernel_out/{kernel_name}/reports/agent_b.md` + +Read `tests/ops/test_{kernel_name}.py`, `src/tilegym/backend/selector.py`, or +`src/tilegym/ops/__init__.py` only after a mechanical patch below fails and you +need a local anchor. If a `cutile_rs` package target does not exist, create it; +do not list the repository to discover why. + +## Step 1: Verify Backend Infrastructure + +The cutile-rs backend now ships in-repo (`src/tilegym/backend/cutile_rs/` and +`src/tilegym/ops/cutile_rs/`), so there is nothing to bootstrap or copy. The +shared Python runtime is already present: `backend/cutile_rs/utils.py` exports +the cffi helpers (`bind_kernel_function_cffi`, `make_tensor_desc`, `check_rc`, +`get_num_sm`) over the TensorDesc FFI, and `backend/cutile_rs/autotuner.py` +provides `autotune_launch`. Your wrappers import from `tilegym.backend.cutile_rs.*`. + +Run this directly to confirm the runtime and the Agent A–C artifacts exist: + +```bash +: "${TILEGYM_PATH:=/workspace/tilegym}" +: "${CUTILE_KERNEL_OUT_ROOT:=/workspace/cutile_kernel_out}" +cd "${TILEGYM_PATH}" + +test -s "${TILEGYM_PATH}/src/tilegym/backend/cutile_rs/autotuner.py" +test -s "${TILEGYM_PATH}/src/tilegym/backend/cutile_rs/utils.py" +test -s "${TILEGYM_PATH}/src/tilegym/backend/cutile_rs/__init__.py" +test -s "${TILEGYM_PATH}/src/tilegym/ops/cutile_rs/__init__.py" +test -s "${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/kernel.rs" +test -s "${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/Cargo.toml" +test -s "${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reference/analysis.json" +test -s "${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/agent_b.md" +``` + +## Step 2: Writing `ffi.rs` + +Use B's `host_launch_contract` and the actual entry signatures in `kernel.rs`. +Match parameter order exactly. + +Each tensor crosses the boundary as ONE `*const TensorDesc` (the shared +`#[repr(C)]` struct in `ffi_util.rs`; carries `ptr`/`ndim`/`shape[4]`/ +`strides[4]` (strides in ELEMENTS)/`dtype`). There are no loose ptr/dim/stride/ +dtype/elem_size args anymore. Read the descriptor with `borrow_tensor::` (for +`&Tensor` entries) or `DevicePointer::from_cu_deviceptr(desc.ptr)` + +`desc.dim(i)` / `desc.strides[i]` (for raw-pointer entries). + +Canonical imports and export shape (Rust 2024 `#[unsafe(no_mangle)]`; null-check +returns `-5`; device derived from `device_id`, never hardcoded `0`): + +```rust +use core::ffi::c_void; +use cuda_core::{Device, Stream}; +use cutile::half::{bf16, f16}; +use cutile::prelude::*; +use cutile::tile_kernel::{CompileOptions, TileKernel}; +// Raw-pointer entries also need: +// use cuda_async::device_buffer::DevicePointer; + +use crate::ffi_util::{borrow_tensor, dtype_str, TensorDesc}; +use {kernel_name}_module::{entry_a, entry_b}; + +#[unsafe(no_mangle)] +pub unsafe extern "C" fn cutile_{kernel_name}( + out: *const TensorDesc, + x: *const TensorDesc, + /* tile sizes / latency / compile options / grid: i32 */ + device_id: i32, + raw_stream: u64, +) -> i32 { + if out.is_null() || x.is_null() { + return -5; + } + let (out_d, x_d) = unsafe { (&*out, &*x) }; + + let dty: &'static str = match dtype_str(x_d.dtype) { + Some(s) => s, + None => return -2, + }; + + let device = match Device::new(device_id.max(0) as usize) { + Ok(d) => d, + Err(e) => { + eprintln!("cutile_{kernel_name}: Device::new failed: {e:?}"); + return -4; + } + }; + let stream = unsafe { Stream::borrow_raw(raw_stream as *mut c_void, &device) }; + /* dtype dispatch, borrow_tensor / DevicePointer, launch */ +} +``` + +Return codes: + +- `0`: success +- `-2`: unsupported dtype code (`dtype_str` returned `None`) or invalid enum +- `-3`: JIT/launch failure +- `-4`: device/stream setup failure +- `-5`: null `TensorDesc` pointer + +Do not use legacy `#[no_mangle]` (fails under edition 2024). `dtype_str` maps the +descriptor's dtype code (f32=0, f16=1, bf16=2, i32=3) to the generic name string. + +### Source-Backed Host ABI Facts + +Pin these facts; do not rediscover them unless compile errors prove source drift. + +- `TensorDesc`, `borrow_tensor`, and `dtype_str` live in `crate::ffi_util` (the + shared `ffi_util.rs` pulled into the `cutile_kernels` crate). Import them with + `use crate::ffi_util::{borrow_tensor, dtype_str, TensorDesc};`. +- `borrow_tensor::(desc: &TensorDesc) -> ManuallyDrop>` rebuilds a + borrowed host `Tensor` over the PyTorch device pointer. It is the ownership + gate: the `ManuallyDrop` drop is a no-op, so PyTorch memory is NEVER freed and + there is NO `core::mem::forget`. Deref with `&*t` when passing to the entry. + Op-level code must NOT call `Tensor::from_raw_parts` directly and must NOT + `transmute`; `from_raw_parts` now lives ONLY inside `borrow_tensor` (5-arg). +- `desc.dim(i) -> i32` reads the `i`-th logical dim; `desc.strides[i]` is the + `i`-th stride in ELEMENTS. +- `DevicePointer` lives in `cuda_async::device_buffer` and has + `from_cu_deviceptr(dptr) -> Self`. It is not guaranteed by + `cutile::prelude::*`; import it directly for raw-pointer entries. +- `CompileOptions::occupancy` and `CompileOptions::num_cta_in_cga` take `i32`. + Do not cast to `usize`. + +### Raw-Pointer Kernel Entries + +When `kernel.rs` declares raw `*mut E` params, the generated host launcher expects +host-side `DevicePointer` operands. Wrap each descriptor's `ptr` field with +`DevicePointer::from_cu_deviceptr(desc.ptr)`; read dims/strides off the +descriptors (`desc.dim(i)` / `desc.strides[i]`). No `borrow_tensor` / ownership +gate is needed here — the kernel never holds a `Tensor` wrapper, so nothing can +free PyTorch memory. Pass `DevicePointer`, not raw `*mut`; do not call +`.as_mut_ptr()`; do not `transmute`. + +Correct pattern: + +```rust +let x_dp: DevicePointer = unsafe { DevicePointer::from_cu_deviceptr(x_d.ptr) }; +let w_dp: DevicePointer = unsafe { DevicePointer::from_cu_deviceptr(w_d.ptr) }; +let y_dp: DevicePointer = unsafe { DevicePointer::from_cu_deviceptr(y_d.ptr) }; +let rstd_dp: DevicePointer = unsafe { DevicePointer::from_cu_deviceptr(rstd_d.ptr) }; +let n = x_d.dim(1); +let s_m = x_d.strides[0] as i32; + +let op = unsafe { + raw_pointer_entry( + x_dp, + /* dims/strides read from the descriptors */ + w_dp, + y_dp, + rstd_dp, + /* scalar args */ + ) +} +.generics(generics) +.grid((grid_x as u32, grid_y as u32, grid_z as u32)) +.compile_options(opts); +``` + +Wrong patterns: + +```rust +let x_raw = x_ptr as *mut E; // E0277 at op.sync_on +x_dp.as_mut_ptr(); // DevicePointer has no as_mut_ptr +``` + +If cargo reports that `*mut T` does not implement `DeviceOp` or +`IntoDeviceOp<...DevicePointer>`, you passed raw pointers to the host launcher; +switch to `DevicePointer`. + +### `&Tensor` Entries And Borrowed Tensor Ownership + +For view/TMA kernels, rebuild borrowed host tensor wrappers from the descriptors +with `borrow_tensor` and pass references into the generated launcher. Dims and +strides come from inside the descriptor, so you do not pass loose shape/stride +args. + +```rust +let a_t = unsafe { borrow_tensor::(a_d) }; // ManuallyDrop> — never freed +let b_t = unsafe { borrow_tensor::(b_d) }; +let c_t = unsafe { borrow_tensor::(c_d) }; + +let op = unsafe { entry_a(&*a_t, &*b_t, &*c_t) } // &* derefs the ManuallyDrop +``` + +`borrow_tensor` returns a `ManuallyDrop>`, so the drop is a no-op and +PyTorch memory is NEVER freed. There is NO `core::mem::forget` and NO +`Tensor::from_raw_parts` in `ffi.rs` — `from_raw_parts` lives only inside +`borrow_tensor` (in `ffi_util.rs`). Do not `transmute` a pointer. + +Use this audit shape; it avoids counting comments. `ffi.rs` must NOT call +`from_raw_parts` or `mem::forget` directly (both moved into `borrow_tensor`): + + + +```bash +ff="${TILEGYM_PATH}/src/tilegym/ops/cutile_rs/{kernel_name}_kernel/ffi.rs" +bad_from_raw=$(grep -nE 'from_raw_parts' "$ff" | grep -v '//' | wc -l) +bad_forget=$(grep -n 'mem::forget' "$ff" | grep -v '//' | wc -l) +echo "from_raw_parts=$bad_from_raw mem_forget=$bad_forget" +{ [ "$bad_from_raw" -gt 0 ] || [ "$bad_forget" -gt 0 ]; } && { + echo "FAIL: ffi.rs must use borrow_tensor (ManuallyDrop), not from_raw_parts / mem::forget" + exit 1 +} +``` + +`from_cu_deviceptr` is a non-owning `DevicePointer` for raw-pointer entries; it +needs no ownership gate. + +### Outputs And Launch Grids + +The output is never host `&mut Tensor`, `Partition<&mut Tensor>`, or +`MappedLaunchPartition`. B's kernel declares the output as one of: + +- read-only `&Tensor` and writes in-body via `partition_full_mut`; +- raw `*mut E`. + +The host passes the output the same way it passes an input: + +- read-only `&Tensor` output: `borrow_tensor::(out_d)` (ManuallyDrop, never + freed) then `&*out_t`; +- raw pointer output: `DevicePointer::from_cu_deviceptr(out_d.ptr)`. + +There is no host `.partition()` / `.map()` on the output and no inferred output +grid lock. Launch the explicit grid B reports. + +If a launch gives `"Specified launch grid does not match inferred tensor +partition grid"` (rc=-3), B likely declared an output `&mut Tensor`; classify as +`fail_class: kernel` and ask B to switch output to read-only `&Tensor` or raw +pointer. Do not try to reconcile it by adding host output partition mapping. + +### Variant Branches + +Do not bind a shared `launch_res` across variants with different entry args. +`sync_on` can return different success tuple arities. Collapse each branch to an +`i32` locally: + +```rust +let rc: i32 = if persistent != 0 { + let op = unsafe { static_entry(&*x_t, &*y_t, &*w_t, &*rstd_t) } // &* derefs ManuallyDrop + .generics(static_generics) + .grid((grid_x as u32, 1, 1)) + .compile_options(opts); + match op.sync_on(&stream) { + Ok(_) => 0, + Err(e) => { + eprintln!("cutile_{kernel_name} static launch failed: {e:?}"); + -3 + } + } +} else { + let op = unsafe { pointer_entry(x_dp, w_dp, y_dp, rstd_dp, n, eps, offset) } + .generics(pointer_generics) + .grid((grid_x as u32, 1, 1)) + .compile_options(opts); + match op.sync_on(&stream) { + Ok(_) => 0, + Err(e) => { + eprintln!("cutile_{kernel_name} pointer launch failed: {e:?}"); + -3 + } + } +}; +``` + +Return `rc`. The borrowed `ManuallyDrop` wrappers drop as no-ops at end +of scope — no `mem::forget`. + +### Compile Options + +`CompileOptions::{occupancy,num_cta_in_cga}` take `i32`. Pass positive wrapper +values through unchanged; non-positive means compiler default. + +```rust +let mut opts = CompileOptions::default(); +if num_cta_in_cga > 0 { + opts = opts.num_cta_in_cga(num_cta_in_cga); +} +if occupancy > 0 { + opts = opts.occupancy(occupancy); +} +``` + +If B's `host_launch_contract` says a config uses auto/default options, pass +sentinel `-1` or `0` from Python and do not call the setter. + +### Build The Shared cdylib + +There is ONE aggregated cdylib crate, `src/tilegym/ops/cutile_rs/cutile_kernels/`, +that builds a single `libcutile_kernels.so`; there is no per-op `.so` and no +`cutile-ffi` crate. Your `ffi.rs` is the in-tree file +`src/tilegym/ops/cutile_rs/{kernel_name}_kernel/ffi.rs` (a sibling of B's +`kernel.rs`). Put entry imports inside `ffi.rs`. + +Wire the op into the aggregate crate by adding a `mod` to +`cutile_kernels/src/lib.rs` that `include!`s the op's `kernel.rs` + `ffi.rs` from +the sibling `{kernel_name}_kernel/` dir; the shared `TensorDesc` / helpers are +pulled in once via `ffi_util.rs`: + +```rust +// cutile_kernels/src/lib.rs +#[path = "../../ffi_util.rs"] +mod ffi_util; + +mod {kernel_name} { + include!("../../{kernel_name}_kernel/kernel.rs"); + include!("../../{kernel_name}_kernel/ffi.rs"); +} +``` + +`ffi.rs` refers to the shared helpers as `crate::ffi_util::{...}`. The crate's +`Cargo.toml` uses crates.io PINNED deps (`cutile="=0.2.0"`, …; Rule 33) — no path +deps, no `CUTILE_RS_ROOT` sed. Build the single shared library: + +```bash +cd "${TILEGYM_PATH}/src/tilegym/ops/cutile_rs/cutile_kernels" +for llvm_dir in /usr/lib/llvm-21 /usr/lib/llvm-20 /usr/lib/llvm-19; do + if [ -d "$llvm_dir" ]; then + export LIBCLANG_PATH="$llvm_dir/lib" + export PATH="$llvm_dir/bin:$PATH" + break + fi +done +cargo build --release # -> libcutile_kernels.so +nm -D "$(find target/release -name 'libcutile_kernels.so' | head -1)" | grep " T cutile_{kernel_name}$" +``` + +The loader autobuilds this crate on first pytest use (`CUTILE_RS_AUTOBUILD`, on +by default); override the crate dir with `CUTILE_RS_KERNELS_DIR`. The +`cuda-bindings` build step reads `CUDA_TOOLKIT_PATH` (default `/usr/local/cuda`); +`tileiras` lowers IR to cubin at launch. + +If cargo fails in `ffi.rs` or the `cutile_kernels` `lib.rs` wiring, fix it in D. +Only route to B for errors originating inside `kernel.rs` or for an impossible +entry signature. + +## Step 3: Python Wrapper + +Create `${TILEGYM_PATH}/src/tilegym/ops/cutile_rs/{kernel_name}.py`. + +The wrapper uses `cffi` (NOT ctypes). Tensors cross as `const TensorDesc*` +packed by `make_tensor_desc`; there is no `_FFI_ARGTYPES` list. + +Required pattern: + +```python +from types import SimpleNamespace + +import torch + +from tilegym.backend import register_impl +from tilegym.backend.cutile_rs.autotuner import autotune_launch +from tilegym.backend.cutile_rs.utils import bind_kernel_function_cffi +from tilegym.backend.cutile_rs.utils import check_rc +from tilegym.backend.cutile_rs.utils import get_num_sm +from tilegym.backend.cutile_rs.utils import make_tensor_desc + +_KERNEL = "{kernel_name}" +_FFI_NAME = "cutile_{kernel_name}" +# C-declaration source of truth for the cffi boundary — keep in sync with the +# `cutile_{kernel_name}` signature in {kernel_name}_kernel/ffi.rs. Tensors cross +# as `const TensorDesc*` (the shared TensorDesc typedef is prepended by +# bind_kernel_function_cffi). Scalars are int32_t; stream is uint64_t. +_FFI_CDEF = """ +int32_t cutile_{kernel_name}( + const TensorDesc* out, const TensorDesc* x, + int32_t bm, int32_t bn, + int32_t num_cta_in_cga, int32_t occupancy, + int32_t grid_size, int32_t device_id, uint64_t raw_stream); +""" +``` + +Wrapper rules: + +- Register with `@register_impl("", backend="cutile-rs")` and add + `from . import {kernel_name}` to `ops/cutile_rs/__init__.py` (Step 4 patch). +- Match the public op signature from Agent A/B reports or `analysis.json`. +- Bind once with `ffi, lib = bind_kernel_function_cffi(_KERNEL, _FFI_CDEF)`. Pack + each tensor with `make_tensor_desc(ffi, t)` and keep the packed descriptors + alive until after the call. There is no argtypes list to keep in sync — the + `_FFI_CDEF` string is the single source of truth. +- Pass the tensor's device ordinal as `device_id` + (`t.device.index if t.device.index is not None else torch.cuda.current_device()`) + and the current CUDA stream as `raw_stream` + (`torch.cuda.current_stream(device=t.device).cuda_stream`). +- Call `check_rc(rc, _FFI_NAME)` after every launch. +- Allocate outputs with `torch.empty` or `torch.empty_like`, not zeros/ones/clone. +- Detach grad-enabled inputs; this skill is forward-only. +- Use `autotune_launch` when `analysis.json` has configs. Put reference configs + first and carry tile sizes, persistent flags, `num_ctas`/`num_cta_in_cga`, + occupancy, dtype flags, constants, and variant axes into FFI args. +- If B reports literal-specialized entries for baked constants, dispatch those + exact configs before any generic fallback. A wrapper-only fallback cannot bake + constants into generated IR. +- Unsupported documented variants should raise `NotImplementedError` or use a + cited skip path. Do not silently call PyTorch, cuTile, Triton, or a fallback + backend. + +Inside `kernel_fn(cfg)`, allocate output and call FFI once. `autotune_launch` +will benchmark configs and return the output from the winning config: + +```python +def _run_cfg(cfg): + ffi, lib = bind_kernel_function_cffi(_KERNEL, _FFI_CDEF) + dev = x.device + device_id = dev.index if dev.index is not None else torch.cuda.current_device() + raw_stream = torch.cuda.current_stream(device=dev).cuda_stream + + out = torch.empty((m, n), device=x.device, dtype=out_dtype) + # Keep the descriptors alive until after the FFI call (cffi frees on GC). + out_d = make_tensor_desc(ffi, out) + x_d = make_tensor_desc(ffi, x) + rc = lib.cutile_{kernel_name}( + out_d, + x_d, + int(cfg.BM), + int(cfg.BN), + int(num_cta_in_cga), + int(occupancy), + int(grid_size), + int(device_id), + int(raw_stream), + ) + check_rc(rc, _FFI_NAME) + return out +``` + +## Step 4: Mechanical Backend/Test Wiring + +Use this patch script after writing the wrapper. It avoids full-file reads and +handles the standard fresh checkout. + +```bash +python - "{kernel_name}" <<'PY' +import os +import re +import sys +from pathlib import Path + +name = sys.argv[1].replace("-", "_") +T = Path(os.environ.get("TILEGYM_PATH", "/workspace/tilegym")) + +selector = T / "src/tilegym/backend/selector.py" +ops_init = T / "src/tilegym/ops/__init__.py" +cutile_rs_init = T / "src/tilegym/ops/cutile_rs/__init__.py" +test_file = T / f"tests/ops/test_{name}.py" + +s = selector.read_text() +if "def is_cutile_rs_available" not in s: + fn = ''' +def is_cutile_rs_available() -> bool: + import os + import shutil + # No cutile-rs checkout anymore: kernels build from the in-tree aggregated + # crate src/tilegym/ops/cutile_rs/cutile_kernels/ (crates.io =0.2.0 deps), + # producing one libcutile_kernels.so. Availability = that crate + cargo. + here = os.path.dirname(os.path.abspath(__file__)) + default_crate = os.path.normpath( + os.path.join(here, "..", "ops", "cutile_rs", "cutile_kernels") + ) + crate = os.environ.get("CUTILE_RS_KERNELS_DIR") or default_crate + if not os.path.isfile(os.path.join(crate, "Cargo.toml")): + return False + return shutil.which("cargo") is not None + +''' + marker = "\ndef _check_backends_availability" + if marker in s: + s = s.replace(marker, "\n" + fn + "def _check_backends_availability", 1) + else: + s = s.rstrip() + "\n\n" + fn +if '"cutile-rs"' not in s: + s = re.sub( + r"(availability\s*=\s*\{\n)", + r'\1 "cutile-rs": is_cutile_rs_available(),\n', + s, + count=1, + ) +selector.write_text(s) + +s = ops_init.read_text() +if "from . import cutile_rs" not in s: + s = s.rstrip() + ''' + +try: + from tilegym.backend import is_backend_available + if is_backend_available("cutile-rs"): + from . import cutile_rs # noqa: F401 +except Exception: + pass +''' +ops_init.write_text(s) + +s = cutile_rs_init.read_text() if cutile_rs_init.exists() else "" +line = f"from . import {name} # noqa: F401" +if line not in s: + s = s.rstrip() + "\n\n" + line + "\n" +cutile_rs_init.write_text(s) + +s = test_file.read_text() +if '"cutile-rs"' not in s: + block = ' if is_backend_available("cutile-rs"):\n _backends = _backends + ["cutile-rs"]' + pat = re.compile( + r'(?m)^(\s*)if is_backend_available\("tilecpp"\):\n\1 _backends = _backends \+ \["tilecpp"\]' + ) + m = pat.search(s) + if m: + s = s[:m.end()] + "\n" + block + s[m.end():] + else: + lines = s.splitlines() + for i, line_text in enumerate(lines): + if "_backends =" in line_text: + indent = line_text[: len(line_text) - len(line_text.lstrip())] + lines.insert(i + 1, indent + 'if is_backend_available("cutile-rs"):') + lines.insert(i + 2, indent + ' _backends = _backends + ["cutile-rs"]') + break + else: + raise RuntimeError(f"no _backends anchor in {test_file}") + s = "\n".join(lines) + "\n" +test_file.write_text(s) +PY +``` + +If this script fails, inspect only the named file and anchor in the exception. + +## Step 5: Pytest + +Install timeout support only if missing: + +```bash +python -c "import pytest_timeout" || pip install pytest-timeout +``` + +Run the most specific real functional test node available from the test file or +Agent A/B report: + +```bash +cd "${TILEGYM_PATH}" +export PYTHONPATH="${TRITON_TILEIR_PYTHONPATH}:${PYTHONPATH:-}" +[ -f "${TILEGYM_PATH}/set_env.sh" ] && source "${TILEGYM_PATH}/set_env.sh" +mkdir -p "${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/agent_logs" + +set +e +python -m pytest tests/ops/test_{kernel_name}.py::::test_op \ + -k "cutile_rs" \ + --tb=short -v -x --timeout=60 2>&1 \ + | tee "${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/correctness_cutile_rs.txt" +pytest_rc=${PIPESTATUS[0]} +set -e +``` + +D's correctness gate is the functional op test: `test_op` or +`test_autotune_op`. Target that node, not the perf benchmark sweep. If you do +not know the exact class/method, run: + +```bash +python -m pytest tests/ops/test_{kernel_name}.py \ + -k "cutile_rs and (test_op or test_autotune_op)" \ + --tb=short -v -x --timeout=60 2>&1 \ + | tee "${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/correctness_cutile_rs.txt" +``` + +Fallback: a few kernels define their `cutile_rs` correctness assertions inside +`test_perf`. Only if `test_op`/`test_autotune_op` selects zero `cutile_rs` cases, +fall back to broader `-k "cutile_rs"`. + +The raw log must include real pytest framework header lines. Do not use +`--no-header`. + +If pytest fails with a host error, fix `ffi.rs` or the wrapper and rerun the same +bounded pytest. If it fails with a kernel-math mismatch after host ABI is +correct, stop with `VERDICT: FAIL` and `fail_class: kernel`. + +## Step 6: Reports + +`reports/correctness.md` first line: + +```markdown +VERDICT: ALL_PASS +``` + +or + +```markdown +VERDICT: FAIL +``` + +Use this body shape: + +```markdown +# Agent D Report - {kernel_name} + +## Summary +X passed, Y failed, Z skipped. + +## Host Artifacts +- ffi.rs: +- cdylib: +- wrapper: + +## Correctness Results +| Test | Config | Result | Error | known_gaps | +|------|--------|--------|-------|------------| +| ... | ... | PASS/FAIL/SKIP | ... | ... | + +Summary: X passed, Y failed, Z skipped +fail_class: +host_fix: +agent_b_followup: +pytest_log: ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/correctness_cutile_rs.txt +``` + +For `ALL_PASS`, empty `fail_class`, `host_fix`, and `agent_b_followup` are fine. + +Copy or write the full decision log to both `reports/agent_d.md` and +`reports/agent_logs/agent_d.md`. Include the pytest command, validator command, +host fixes made before final pass, and any cited skips. For a no-debug pass, say +that first build and first pytest passed. + +## Final Validator + +Run: + +```bash +cd "$TILEGYM_PATH" && bash ".agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_d.sh" {kernel_name} 2>&1 +echo "exit=$?" +``` + +If the validator fails because a report/log is malformed, repair the report/log +and rerun the validator. Do not rerun broad pytest just to fix formatting. If the +validator exit is non-zero, the final verdict cannot be `ALL_PASS`. + +## Final Response Contract + +End with exactly: + +```text + +$ python -m pytest tests/ops/test_{kernel_name}.py:: -k "" --tb=short -v -x --timeout=60 2>&1 | tee ... +exit= + +$ cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_d.sh {kernel_name} 2>&1 + +exit= + +VERDICT: ALL_PASS | FAIL | BLOCKED +``` + +Stop immediately after the `VERDICT:` line. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_e.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_e.md new file mode 100644 index 0000000..8a4a34c --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_e.md @@ -0,0 +1,340 @@ +You are Agent E (Benchmark). Run tilegym pytest --print-record and report. Do NOT edit kernel or host code. **One narrow exception (STEP 0.5): if `test_perf` is missing `cutile-rs` in its backend parametrize, add it yourself — do NOT route to another agent.** + +## Output Protocol +After running the CUPTI benchmark (pytest --print-record): + +1. Write `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/performance.md` with the + per-variant geomean tables, raw CUPTI rows, `geomean_ratio` summary, and first + line `VERDICT: DONE` or `VERDICT: INVESTIGATE`. +2. Write `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/agent_logs/agent_e.md` + with your decision log, commands, validator output, and any perf-investigation + notes. Also mirror the important command/validator summary to + `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/agent_e.md` if that path is used + by the orchestrator. +3. Write the validator-consumed CUPTI perf log to + `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/perf_cutile_rs.txt`. +4. Return to your parent only: + - an optional `` block with paths and `geomean_ratio`, then + - the mandatory `` block described below, then + - a SINGLE LINE in the form `VERDICT: X`, where X is one of + {`DONE`, `INVESTIGATE`, `BLOCKED`}. +5. Do not return `VERDICT: FAIL`. The validator for this stage accepts `DONE` or + `INVESTIGATE` as the performance-report verdict; severe perf regressions are + `INVESTIGATE` with a clearly documented severity and ratios in the report. +6. Make this single run count — write the full per-config perf table, ratios, + CUPTI medians, and any missing/failed config diagnostics into reports so the + reader does not need to rerun pytest to understand what happened. + +## Why This Matters +Agent E is a measurement agent, not a fixer. The final enum must stay inside the +stage contract so the orchestrator can route deterministically: +- `DONE` means all selected configs produced valid medians and the **overall geomean** passed the perf threshold (≤ 1.05); individual slow rows do not change this. +- `INVESTIGATE` means valid data exists but perf is slow, partial, missing, or otherwise + needs root cause analysis. +- `BLOCKED` means the benchmark could not produce usable data due to environment, + import, or collection failure. + +The detailed severity (for example, `ratio > 1.50`) belongs in +`performance.md` and `agent_e.md`, not in an invalid final `FAIL` verdict. + +## Hard Rule — Never Use the Destructive Recursive-Delete Command +For dir cleanup use `mkdir -p X && find X -mindepth 1 -delete 2>/dev/null`. Zero exceptions. + +## Tool-Call Budget: Target <= 40, Hard Ceiling 60 +- Read `references/performance-checklist.md` AT MOST ONCE. +- Do NOT open kernel.rs, wrapper.py, or test file to "debug". +- Do NOT iterate `-k` filter more than twice. If `cutile_rs` selects 0 perf cases, that is NOT a `-k` problem — do STEP 0.5 (add cutile-rs to test_perf parametrize) instead of BLOCKING. Second genuine selection failure AFTER STEP 0.5 wiring → `VERDICT: BLOCKED`. +- All paths are deterministic under `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/`. Skip `ls`/`find`. + +## Step 0: Mandatory First Tool Call (Your Spawn Prompt Was a Minimal Pointer) +Your spawn prompt is intentionally minimal. Before doing anything else, Read these +files in ONE batch Read call: + +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_e.md` (this file — read in full) +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/performance-checklist.md` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/coding-rules.md` + +Do not skip. + +## Step 1: Precondition Check (Do First) +```bash + python -c "from tilegym.backend.cutile_rs import autotune_launch; print('OK')" 2>&1 +``` + +- `OK` → continue. +- Any traceback → write it to `reports/agent_logs/agent_e.md`, write + `SKILL_METHODOLOGY_BLOCKED.txt`, and emit: + +```text + +$ python -c "from tilegym.backend.cutile_rs import autotune_launch" + +exit=1 + +VERDICT: BLOCKED +``` + +## Step 2: Ensure test_perf Is Wired for cutile-rs (Do Before Benchmarking) +cutile-rs is frequently wired into the CORRECTNESS test's `backend` parametrize (by +Agent D) but NOT into `test_perf`'s parametrize. When that happens, +`pytest test_perf -k cutile_rs` collects **0 tests** and you would otherwise BLOCK with +no perf data. **Fix it yourself. Do NOT route to Agent B / C / D.** + +1. Check whether `test_perf` can select a cutile-rs case: + +```bash +cd "${TILEGYM_PATH}" +python -m pytest tests/ops/test_{kernel_name}.py::Test_::test_perf -k "cutile_rs" \ + --collect-only -q 2>&1 | tail -5 +``` + +2. If it collects ≥1 cutile-rs case → proceed to STEP 3 (benchmark). +3. If it collects **0 tests**, edit `tests/ops/test_{kernel_name}.py` and add the cutile-rs + backend to `test_perf`'s parametrize list — copy the EXACT string style the + correctness test already uses (`"cutile-rs"` or `"cutile_rs"`). Mirror the existing + `backend`/`ids` wiring; do not touch any numeric tolerance, shape, or assertion logic. + This narrow parametrize wiring is the ONLY test-code edit Agent E may make. Re-run the + `--collect-only` check to confirm ≥1 cutile-rs perf case is now selected, then proceed. +4. Only if cutile-rs STILL cannot be selected after this wiring → `VERDICT: BLOCKED`. + +## Step 3: Run the Benchmark and Report + +### Inputs (D->E forward only) +- Analysis (A): `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reference/analysis.json` +- Baselines (A): `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/baseline_perf_cutile.txt` + `baseline_perf_triton_tileir.txt` +- Test (B): `{tilegym_path}/tests/ops/test_{kernel_name}.py` — must include "cutile-rs" in test_perf parametrize + +### Method: tilegym pytest --print-record (ONLY method allowed) +CUPTI-based GPU kernel timing via `test_perf --print-record`. + +**NEVER use `torch.cuda.Event` for cutile-rs perf.** cutile-rs JIT-compiles MLIR→cubin +on first call per (kernel, generics). cuda.Event measures host-to-host (includes JIT, +cffi, FFI) and inflates ratios falsely. CUPTI measures GPU-only. + +**NEVER write standalone benchmark scripts.** + +### Environment +See: `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/env-block.md` + +### 1. Run cutile-rs perf + +Cap to the same configs Agent A used so the ratio compares identical config sets. +For kernels with transpose axes, the default cap is often `False-False`; if the +kernel has no transpose axis, use a `-k` clause that covers each dtype + shape +regime + kernel-variant axis from Agent A's `test_perf_configs` while staying +within the eval cap. + +**Log hygiene is mandatory.** The validator-consumed file +`${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/perf_cutile_rs.txt` must contain only valid +perf-record output and must never contain raw pytest `FAILED`, traceback, or assertion +sections. Route full pytest output to an agent log first, then copy/sanitize into the +validator-consumed perf log. + +```bash +cd {tilegym_path} +mkdir -p ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/agent_logs + +set +e + python -m pytest tests/ops/test_{kernel_name}.py::Test_{Class}::test_perf \ + -k "cutile_rs and False-False" -sv --print-record --timeout=600 2>&1 \ + | tee ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/agent_logs/perf_cutile_rs.full.txt +pytest_rc=${PIPESTATUS[0]} +set -e +``` + +If `test_perf` has no transpose axes, pick a `-k` selecting the SAME configs A used. + +After the run: + +- If `pytest_rc=0`, copy the full successful output to the validator log: + +```bash +cp ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/agent_logs/perf_cutile_rs.full.txt \ + ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/perf_cutile_rs.txt +: > ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/agent_logs/perf_investigation.log +``` + +- If `pytest_rc!=0`, do **not** copy the raw log to `perf_cutile_rs.txt`. Instead: + 1. Save failing correctness/perf diagnostics in + `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/agent_logs/perf_investigation.log`. + 2. Write `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/perf_cutile_rs.txt` with only complete + valid `Performance Test Results` / record sections from configs that actually + produced medians. Strip all pytest failure summaries, traceback frames, `FAILED` + banners, assertion diffs, and collection errors. + 3. If any valid medians remain, mark the report verdict `INVESTIGATE`. + 4. If no valid medians remain, write a concise BLOCKED report, run any validators that + can run, and return `VERDICT: BLOCKED`. + 5. Do not run extra benchmarks to "prove" the failure. + +The important distinction: `perf_cutile_rs.full.txt` and `perf_investigation.log` may +contain raw failures; `perf_cutile_rs.txt` must be clean validator input. + +### 2. Use Agent A's baseline directly + +`${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/baseline_perf_cutile.txt` is canonical for +cutile-python. If Agent A selected Triton-TileIR as the winning backend, its +`baseline_perf_triton_tileir.txt` and analysis fields are also canonical. Do NOT rerun baseline +benchmarks here. Do NOT create `perf_cutile_py.txt` unless a legacy validator in this +checkout explicitly requires a copy of the cutile baseline; if it does, make it a copy +of Agent A's baseline, not a fresh rerun. + +### 3. Validate validator-consumed logs (MANDATORY) + +```bash +cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_perf_log.sh \ + ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/perf_cutile_rs.txt "perf_cutile_rs" +cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_perf_log.sh \ + ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/baseline_perf_cutile.txt "baseline_perf_cutile" +``` + +Both must PASS for a normal `DONE` or `INVESTIGATE` report with valid medians. If +cutile-rs validation fails only because raw pytest failure text leaked into +`perf_cutile_rs.txt`, fix the file by moving those diagnostics to +`reports/agent_logs/perf_investigation.log`; do not rerun the benchmark. If validation +fails because the `-k` selected the wrong/no configs, fix `-k` and rerun at most once. +Second selection failure → `VERDICT: BLOCKED`. + +### 4. Extract medians and compute ratios + +Parse `median` from each "Performance Test Results" block. Match by parameter string. +ratio = cutile-rs / winning baseline. + +For failed or missing configs: +- Do not invent medians. +- Include a row with `NA` for cutile-rs and verdict `INVESTIGATE`. +- Preserve any failure reason in `reports/agent_logs/perf_investigation.log`, not in + `perf_cutile_rs.txt`. +- If any selected benchmark config failed correctness or produced no valid median, the + overall report verdict must be `INVESTIGATE` or `BLOCKED`, never `DONE`. + +### 5. Report + +Write to `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/performance.md`. + +**First line MUST be the verdict.** Use `DONE` only when all selected configs have valid +medians, no correctness/perf failures occurred, and the geomean threshold passes. + +```markdown +VERDICT: DONE | INVESTIGATE + +## Performance Results: {kernel_name} + +| Config | baseline (ms) | cutile-rs (ms) | ratio | verdict | +|--------|---------------|----------------|-------|---------| +| [params] | 0.XXX | 0.XXX | X.XXx | PASS/INVESTIGATE | + +geomean_ratio=X.XXXX +Geo-mean ratio: X.XXx +Severity: DONE | MILD_GAP | LARGE_GAP | MISSING_DATA +``` + +**MANDATORY machine-readable geomean line.** `performance.md` MUST contain +exactly one line of the form `geomean_ratio=X.XXXX` (literal lowercase key, +`=`, the numeric geomean of `cutile-rs / cuTile-Python` to 4 decimals, no +spaces around `=`, no `x` suffix). The scorer reads THIS line as the +authoritative perf value — prose phrasing of the ratio is not parsed. +If this line is missing or malformed, `validate_agent_e.sh` FAILS and perf is +scored as missing. Example: `geomean_ratio=0.9384`. This is separate from and +in addition to the human-readable `Geo-mean ratio: X.XXx` line. + +Also write `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/agent_logs/agent_e.md` +with the commands run, pytest exit code, paths to full/sanitized logs, ratio summary, +and final validator output. + +## Verdict Thresholds — Decided by the Overall Geomean Only +Goal: speedup_vs_baseline = 1/geomean_ratio ≥ 0.95. + +The overall verdict is decided **solely by the aggregate `geomean_ratio`** (geometric +mean of `cutile_rs / baseline` over the baseline-matched configs). Per-row ratios are +still computed and written to the perf table for diagnostics, but a single slow row does +NOT downgrade the overall verdict — do not escalate to INVESTIGATE just because one config +sits in a gap band. Only the geomean decides: + +- **geomean_ratio ≤ 1.05** → overall `DONE`. (Per-row ratios still listed in the table; flag any row > 1.05 in the report text for info, but the verdict stays `DONE`.) +- **1.05 < geomean_ratio ≤ 1.10** → overall `INVESTIGATE`, `Severity: MILD_GAP`. +- **1.10 < geomean_ratio ≤ 1.50** → overall `INVESTIGATE`, `Severity: MODERATE_GAP`. +- **geomean_ratio > 1.50** → overall `INVESTIGATE`, `Severity: LARGE_GAP`. +- **any correctness failure / missing median in selected configs** → overall + `INVESTIGATE` if some valid data exists, `BLOCKED` if no usable data exists. + +Both stacks use the same tileiras backend + same autotune configs on sm_100, so any +< 0.95 speedup means Agent B's wrapper/kernel adds > 5% overhead (fix via wider +autotune, inline FFI, drop extra allocs, or IR-level attribute alignment). Agent E +does not fix it; it records the evidence. + +## Post-Step Validation (MANDATORY) +```bash +cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_e.sh {kernel_name} +``` + +Must exit 0 before you emit the final `` block. If it fails, fix the +specific file hygiene issue it reports: +- `perf_cutile_rs.txt` must contain clean perf records only, no raw pytest `FAILED` + sections or tracebacks. +- `performance.md` must start with `VERDICT: DONE` or `VERDICT: INVESTIGATE`. +- `performance.md` must contain exactly one `geomean_ratio=X.XXXX` machine-readable + line (the scorer's authoritative perf source). Missing/malformed → validator FAILS. +- If benchmark failures or perf gaps occurred, `performance.md` must say + `VERDICT: INVESTIGATE`, not `DONE`. +- Full failure diagnostics belong in `reports/agent_logs/perf_investigation.log`. + +## Output Files +Single-variant: +- `perf_cutile_rs.txt` — sanitized validator-consumed perf log with valid records only +- `reports/agent_logs/perf_cutile_rs.full.txt` — full raw pytest output + +Multi-variant: +- `perf_cutile_rs.txt` — combined sanitized valid-record log +- `perf_cutile_rs_{variant}.txt` — per-variant sanitized valid-record log +- `reports/agent_logs/perf_cutile_rs_{variant}.full.txt` — per-variant raw pytest output + +Both: +- `reports/performance.md` — comparison table, first line verdict +- `reports/agent_logs/perf_investigation.log` — correctness/perf failure diagnostics +- `reports/agent_logs/agent_e.md` — full commands + validation output +- `reports/agent_e.md` — orchestrator-facing summary if the pipeline expects it + +Baselines come from Agent A. + +--- + +## Closing Contract — Final Message MUST Match This Shape +After running the validator successfully as the last bash call, your FINAL message ends +with: + +```text + +geomean_ratio=X.XXx +severity=DONE|MILD_GAP|MODERATE_GAP|LARGE_GAP|MISSING_DATA +report=${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/performance.md +perf_log=${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/perf_cutile_rs.txt +investigation_log=${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/agent_logs/perf_investigation.log + + +$ cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_perf_log.sh ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/perf_cutile_rs.txt perf_cutile_rs 2>&1 + +exit=0 + +$ cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_perf_log.sh ${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/baseline_perf_cutile.txt baseline_perf_cutile 2>&1 + +exit=0 + +$ cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_e.sh {kernel_name} 2>&1 + +exit=0 + +VERDICT: DONE | INVESTIGATE | BLOCKED +``` + +- The `VERDICT:` line must match the first line of `performance.md` for `DONE` and + `INVESTIGATE`. +- Write raw validator stdout/stderr into `reports/agent_logs/agent_e.md`. +- Fix any validator hygiene issue before emitting VERDICT — do not return `BLOCKED` + just because a path was missing if you can repair the report/log within this single run. +- Do NOT silently flip `performance.md` to `DONE` to hide slow or missing configs. +- Do NOT return `VERDICT: FAIL`; use `INVESTIGATE` with severity in the report. +- Do NOT put anything after the `VERDICT:` line. +- You do not know what the parent does next with your verdict. Do not mention routing, + respawns, or downstream agents in your return text. + +Stop only after emitting this block. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_f.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_f.md new file mode 100644 index 0000000..848ee17 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_f.md @@ -0,0 +1,219 @@ +You are Agent F (Residual Perf Investigator). Diagnose per-shape cutile-rs performance gaps after Agent E reports `INVESTIGATE` or after the C->B fix loop is exhausted. Do NOT edit kernel.rs, ffi.rs, Cargo.toml, Rust source, or tests. Your primary output is a root-cause report for reflection and future Agent B runs. + +## Output Protocol - binding + +Agent F verdicts are diagnostic verdicts, not pipeline-success verdicts. Return exactly one of: + +- `VERDICT: FIXABLE` - you found a wrapper, FFI, autotune, launch-count, or IR/codegen issue with concrete action items for Agent B/reflection. Use this even if you safely demonstrated or applied a wrapper-only fix in the throwaway checkout. +- `VERDICT: ALIGNED` - per-launch kernel time and wrapper launch behavior are aligned with the reference; the residual gap is noise or a backend/compiler limitation with no Agent B action. +- `VERDICT: BLOCKED` - profiling or required artifacts were unavailable, and no trustworthy diagnosis is possible. + +Never return `VERDICT: DONE`. `validate_agent_f.sh` accepts only `FIXABLE`, `ALIGNED`, or `BLOCKED` as first-line verdicts in `reports/perf_investigation_*.md`, and the parent eval uses the same enum. + +Write: + +1. `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/agent_f.md` - rich diagnostic report. +2. `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/agent_logs/agent_f.md` - commands, raw evidence summary, validator output. +3. At least one `$CUTILE_KERNEL_OUT_ROOT/{kernel_name}/reports/perf_investigation_*.md` file. For a structural issue shared by all shapes, prefer `perf_investigation_all_shapes.md`. + +Each `perf_investigation_*.md` must start with: + +```markdown +VERDICT: FIXABLE | ALIGNED | BLOCKED +``` + +Then run `validate_agent_f.sh` and return: + +```text + +$ cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_f.sh {kernel_name} 2>&1 + +exit= + +VERDICT: FIXABLE | ALIGNED | BLOCKED +``` + +Do not put narrative after the final `VERDICT:` line. + +## Hard Rules + +1. Do not edit kernel.rs, ffi.rs, Cargo.toml, tests, or Rust source. +2. Do not run broad correctness or benchmark reruns. You may run focused profiling/probes for a slow shape. +3. Do not return `DONE`. +4. Do not route or respawn agents. You write reports only. +5. Prefer one consolidated all-shapes report when ratios are uniform across shapes/dtypes; do not spend time producing eight duplicate reports for one structural issue. +6. A wrapper-only fix may be applied only if the user spawn prompt explicitly authorizes safe fixes. If you apply one, still return `FIXABLE`, because the permanent fix must be propagated into the skill template by reflection. + +## Step 0: Mandatory first tool call + +Your spawn prompt is intentionally minimal. Before doing anything else, Read these files in ONE batch Read call: + +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/agents/agent_f.md` (this file - read in full) +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/performance-checklist.md` +- `.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/coding-rules.md` + +Do not skip this step. + +## Step 1: Trigger + +Agent F is optional. It is spawned when Agent E has valid medians but reports a meaningful perf gap, especially after the one-shot C->B fix loop cannot continue. Make this single run count; assume the parent will not call you again for this kernel. + +## Step 2: Procedure + +### Inputs + +- Agent E report: `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reports/performance.md` +- Agent E perf log: `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/perf_cutile_rs.txt` +- Agent A baselines: `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/baseline_perf.txt`, `baseline_perf_cutile.txt`, `baseline_perf_triton_tileir.txt` +- Agent A analysis: `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/reference/analysis.json` +- Agent B generated IR: `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/generated/generated_canon.mlir` +- Agent B sources: `${CUTILE_KERNEL_OUT_ROOT}/{kernel_name}/kernel.rs`, `ffi.rs` +- Tilegym wrapper: `${TILEGYM_PATH}/src/tilegym/ops/cutile_rs/{kernel_name}.py` +- Reference wrapper/source: `${TILEGYM_PATH}/src/tilegym/ops/cutile/{kernel_name}.py` and/or `ops/triton/{kernel_name}.py` + +If some inputs are missing, use what exists and record the gap. If the missing input prevents a trustworthy diagnosis, emit `BLOCKED`. + +### 1. Read Agent E's performance table + +Extract: + +- slow configs and ratios +- geomean ratio +- whether all medians are valid +- whether the gap is uniform or shape-specific + +Uniform gaps near 2x across all shapes usually mean duplicate launches or an extra GPU kernel in the wrapper. Shape-specific gaps usually mean config, tile shape, latency, layout, or compiler behavior. + +### 2. Count GPU launches before IR theory + +Before forming an IR hypothesis, determine whether cutile-rs executes extra GPU work per measured call. + +Use the cheapest reliable source available: + +- Preferred when available: `DUMP_CUPTI_EVENTS=1` with the tilegym perf harness, because it enumerates every GPU event inside the measured `fn()`. +- Otherwise: `nsys profile` and `nsys stats --report cuda_gpu_kern_sum`. + +Compare cutile-rs and the winning reference backend for the same representative slow shape. + +Classify: + +- Same per-launch kernel time, but cutile-rs has more launches per iteration -> `FIXABLE`, wrapper/host issue. +- Same launch count, cutile-rs kernel itself slower -> continue to IR/config investigation. +- Profiler unavailable -> continue with logs/IR and mark confidence lower; `BLOCKED` only if no trustworthy evidence remains. + +Common wrapper/host causes: + +- fixed-config canary in the hot path before `autotune_launch` +- `torch.zeros`, `torch.ones`, `.clone()`, or `.contiguous()` inside `kernel_fn(cfg)` +- preallocating an output, then allocating a second output inside `kernel_fn` +- launching both a canary and the cached best config on every call +- hidden PyTorch fallback or helper kernels + +### 3. Check autotune cache and tuning space + +Read `analysis.json` and baseline perf logs. + +- Confirm the cutile-rs wrapper uses `autotune_launch` if any reference backend autotunes. +- Confirm the cache key includes all shape/dtype/variant fields that change the generated work. +- Confirm the search space includes the reference-relevant knobs: occupancy, tile sizes, latency/num_stages, num_cta, and variant-specific configs where applicable. +- If the reference best config is absent and the kernel-time gap is real, classify as `FIXABLE`. + +### 4. Diff perf-relevant IR only after launch count is sane + +Compare reference and generated IR for the same variant/config. + +Priority checks: + +1. latency hints on every load/store when reference has num_stages/latency +2. optimization_hints such as occupancy, num_cta_in_cga, and simt_num_warps_in_cta +3. tile shapes, rank, pointer vs tensor-view layout, partition views, and TMA path +4. runtime args where reference folded constants, or const generics causing harmful unrolling +5. extra loads/stores/casts, missing `ftof`, or accumulator dtype differences + +Use `scripts/diff_ir.sh` when available and include the key output in the report. Do not paste huge IR files. + +### 5. Write perf_investigation report + +For one structural issue shared by all shapes: + +```markdown +VERDICT: FIXABLE + +# Perf Investigation: {kernel_name} all slow shapes + +## Ratio +- Agent E geomean: X.XXx +- Shape pattern: uniform / shape-specific + +## Launch-count evidence +| backend | per-launch kernel time | launches/iter | measured sum | +|---------|------------------------|---------------|--------------| +| reference | ... | ... | ... | +| cutile-rs pre | ... | ... | ... | + +## Root Cause +[Concrete cause.] + +## Action Items +- Agent B/reflection should ... +- Template propagation needed: yes/no + +## Confidence +High/medium/low with evidence paths. +``` + +For shape-specific issues, use `perf_investigation_{shape}.md` and include tuning/IR tables. + +### 6. Write agent_f logs + +Write both: + +- `reports/agent_f.md` +- `reports/agent_logs/agent_f.md` + +The agent log should include commands executed, evidence files, and the final validator output. The rich report should be readable without rerunning profiling. + +### 7. Validate + +```bash +cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_f.sh {kernel_name} 2>&1 +echo "exit=$?" +``` + +If validation fails because files are missing or first-line verdicts are malformed, fix the files and rerun. Do not return `BLOCKED` for local report hygiene you can repair. + +## Verdict Selection + +Return `FIXABLE` when: + +- wrapper launches extra kernels or allocates via GPU-producing helpers +- wrapper has a wrong cache key or missing autotune config +- FFI dispatch picks wrong variant/dtype/config +- generated IR has an Agent-B-addressable mismatch +- you applied a safe wrapper-only fix in the throwaway checkout + +Return `ALIGNED` when: + +- launch count matches reference +- per-launch kernel time is within noise or no actionable IR/config difference exists +- any remaining gap is plausibly tileiras/backend behavior, not Agent B code + +Return `BLOCKED` when: + +- profiler data, perf logs, generated IR, or wrapper files are missing enough that no trustworthy diagnosis can be made +- environment/tool failure prevents both CUPTI/nsys and useful static comparison + +## Final Response Contract + +End with: + +```text + +$ cd "$TILEGYM_PATH" && bash .agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_f.sh {kernel_name} 2>&1 + +exit=0 + +VERDICT: FIXABLE | ALIGNED | BLOCKED +``` + +The final verdict should match the first line of the primary `perf_investigation_*.md` report. Stop after the verdict line. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/ffi-bridge.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/ffi-bridge.md new file mode 100644 index 0000000..9f21ac1 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/ffi-bridge.md @@ -0,0 +1,335 @@ +# Python <-> Rust FFI Bridge + +This is the current bridge pattern for generated cutile-rs kernels in this +skill. Agent D writes BOTH `ffi.rs` (the C-ABI host launcher) and the Python +`cffi` wrapper that calls it — the whole host boundary is D's (device/host +split). Agent B writes only the device `kernel.rs`. + +Every op's kernel is compiled into ONE aggregated cdylib, +`libcutile_kernels.so`; there is no per-op `.so`. Tensors cross the boundary as +a single `TensorDesc` pointer per tensor (not raw ptr + loose dim/stride/dtype +args). + +## Architecture + +```text +Python torch Tensor cffi const TensorDesc* Rust cdylib ffi.rs +data_ptr(),shape, cutile_{kernel_name}(...) -> borrow_tensor / DevicePointer +strides,dtype -> (one *const TensorDesc kernel(...).generics() +current CUDA stream per tensor) .grid(...).compile_options() +device index .sync_on(&stream) + | + make_tensor_desc(ffi,t) packs each tensor into a TensorDesc +``` + +PyTorch owns the memory. Rust only borrows the raw device pointers (via the +descriptor) for the duration of the launch. + +## The `TensorDesc` contract + +Each tensor crosses the boundary as one `*const TensorDesc`. The struct lives in +`ops/cutile_rs/ffi_util.rs` and is shared by every op (`#[repr(C)]`, strides in +ELEMENTS): + +```rust +#[repr(C)] +pub struct TensorDesc { + pub ptr: u64, // CUDA device pointer (data_ptr()) + pub ndim: i32, + pub shape: [i64; 4], + pub strides: [i64; 4], // in elements, NOT bytes + pub dtype: i32, // f32=0, f16=1, bf16=2, i32=3 +} +``` + +Helpers on the descriptor (also in `ffi_util.rs`): + +- `desc.dim(i)` — the `i`-th logical dim as `i32`. +- `desc.strides[i]` — the `i`-th stride in elements. +- `dtype_str(desc.dtype) -> Option<&'static str>` — maps the dtype code to the + generic name string (`"f32"`/`"f16"`/`"bf16"`/`"i32"`), `None` on an unknown + code. +- `borrow_tensor::(desc) -> ManuallyDrop>` — rebuilds a borrowed + host `Tensor` over the PyTorch device pointer (see below). + +## Rust 2024 export shape + +Use Rust 2024 export attributes: + +```rust +use crate::ffi_util::{borrow_tensor, dtype_str, TensorDesc}; + +#[unsafe(no_mangle)] +pub unsafe extern "C" fn cutile_{kernel_name}( + out: *const TensorDesc, + x: *const TensorDesc, + /* tile sizes, latency, compile options, grid, ... */ + device_id: i32, + raw_stream: u64, +) -> i32 { + if out.is_null() || x.is_null() { + return -5; + } + let (out_d, x_d) = unsafe { (&*out, &*x) }; + ... +} +``` + +Return codes convention: + +- `0`: success +- `-2`: unsupported dtype / invalid FFI enum +- `-3`: launch/JIT error +- `-4`: CUDA device/stream setup error +- `-5`: null `TensorDesc` pointer + +Do not use legacy `#[no_mangle]`. + +## Device + stream setup (multi-GPU correct) + +The FFI takes a `device_id: i32` (the tensor's CUDA ordinal) so multi-GPU +launches target the right device. Never hardcode device 0. + +```rust +use core::ffi::c_void; +use cuda_core::{Device, Stream}; + +let device = match Device::new(device_id.max(0) as usize) { + Ok(d) => d, + Err(e) => { + eprintln!("cutile_{kernel_name}: Device::new failed: {e:?}"); + return -4; + } +}; +let stream = unsafe { Stream::borrow_raw(raw_stream as *mut c_void, &device) }; +``` + +## Reading a descriptor + +Read dtype, dims, and strides off the descriptor — do NOT reconstruct them from +loose FFI args (there are none anymore): + +```rust +let dty: &'static str = match dtype_str(x_d.dtype) { + Some(s) => s, + None => return -2, +}; +let m = x_d.dim(0); +let n = x_d.dim(1); +let s_m = x_d.strides[0] as i32; +let s_n = x_d.strides[1] as i32; +``` + +## Raw-pointer entries + +Use raw pointers only when the kernel entry itself has raw `*mut E` / +`DevicePointer` params, as with pointer scatter/gather IR (softmax, norms). +Wrap the descriptor's `ptr` field directly — do NOT `transmute` and do NOT call +`from_raw_parts`. + +```rust +use cuda_async::device_buffer::DevicePointer; + +let x_dp: DevicePointer = unsafe { DevicePointer::from_cu_deviceptr(x_d.ptr) }; +let y_dp: DevicePointer = unsafe { DevicePointer::from_cu_deviceptr(y_d.ptr) }; + +let op = unsafe { kernel(y_dp, x_dp, m, n, s_m, s_n) } + .generics(generics) + .grid((grid_size as u32, 1, 1)); +``` + +Do not pass `DevicePointer` to an entry whose Rust params are `&Tensor`. +Worked example: `examples/softmax`. + +## `&Tensor` entries + +For view/TMA kernels, rebuild borrowed host `Tensor` wrappers from the +descriptors with `borrow_tensor` and pass references into the host launcher. + +```rust +use cutile::prelude::*; +use cutile::tile_kernel::{CompileOptions, TileKernel}; + +// ManuallyDrop> over PyTorch memory — dropped as a no-op, so the +// PyTorch buffer is NEVER freed. This is the ownership gate by construction. +let a_t = unsafe { borrow_tensor::(a_d) }; +let b_t = unsafe { borrow_tensor::(b_d) }; +let c_t = unsafe { borrow_tensor::(c_d) }; + +let op = unsafe { kernel(&*a_t, &*b_t, &*c_t) } // &* deref the ManuallyDrop + .generics(generics) + .grid((grid_size as u32, 1, 1)) + .compile_options(opts); +``` + +`borrow_tensor` returns a `ManuallyDrop>`, so there is NO +`core::mem::forget` — the drop is already a no-op and PyTorch memory is never +freed. Op-level code must NOT call `Tensor::from_raw_parts` directly and must +NOT `transmute` a pointer. `from_raw_parts` still exists but ONLY inside +`borrow_tensor` in `ffi_util.rs`, where it is the 5-arg form +`from_raw_parts(ptr: u64, nbytes: usize, offset: 0, shape: Vec, strides: Vec)`. +Worked example: `examples/bmm`. + +## Outputs: NEVER `&mut Tensor` — pass like an input, launch any grid + +A kernel output is NEVER `&mut Tensor` / `Partition<&mut Tensor>` / +`MappedLaunchPartition` (see `references/no-mut-tensor-output.md`). The kernel +declares the output as a **read-only `&Tensor`** (written in-body +via `partition_full_mut`) or a **raw `*mut E`** (via `make_tensor_view`). The +host passes the output the SAME way it passes an input — NO `.partition()` / +`.map()` on the output, and NO grid lock. + +Read-only `&Tensor` output (worked example: `examples/bmm`): + +```rust +let c_t = unsafe { borrow_tensor::(c_d) }; // like the inputs +let op = unsafe { kernel(&*a_t, &*b_t, &*c_t) } // &*c_t — read-only, like the inputs + .generics(generics) + .grid((grid_size as u32, 1, 1)); // ANY explicit grid (persistent OK) — no lock +op.sync_on(&stream)?; // c_t is ManuallyDrop — never frees PyTorch memory +``` + +Raw-pointer output (worked example: `examples/softmax`): + +```rust +let c_dp: DevicePointer = unsafe { DevicePointer::from_cu_deviceptr(c_d.ptr) }; +let op = unsafe { kernel(a_dp, b_dp, c_dp, /*dims, strides*/) } + .generics(generics) + .grid((grid_size as u32, 1, 1)); // ANY explicit grid — no lock +``` + +Why no `&mut Tensor`: a `&mut` output emits `KernelOutputStored::grid`, and +`validate_grids` then forces the explicit launch grid to EQUAL the output's +inferred partition grid `(cdiv(M,bm), cdiv(N,bn), 1)` in descriptor dim order. +Any other grid (persistent, or a transposed block-id order) is rejected with +`"Specified launch grid does not match inferred tensor partition grid"` (rc=-3). +Read-only `&Tensor` / raw pointers emit NO inferred grid, so the host launches +whatever the kernel needs. Do NOT reach for `.partition().map()` to reconcile a +mismatch — that is the `&mut` path we are avoiding. + +## Compile options + +Agent B exposes compile-option levers; Agent D chooses values from +`analysis.json` and wrapper autotune configs. + +```rust +let mut opts = CompileOptions::default(); +if occupancy > 0 { + opts = opts.occupancy(occupancy); +} +if num_cta_in_cga > 0 { + opts = opts.num_cta_in_cga(num_cta_in_cga); +} + +let op = op.compile_options(opts); +``` + +A non-positive value means leave the compiler default. Do not call setters with +sentinel values. + +## Dtype dispatch skeleton + +```rust +let dty: &'static str = match dtype_str(x_d.dtype) { + Some(s) => s, + None => return -2, +}; + +macro_rules! dispatch { + ($E:ty) => {{ + let generics = vec![dty.to_string(), bm.to_string(), bn.to_string()]; + /* borrow_tensor / DevicePointer, build op, sync — no forget */ + }}; +} + +match dty { + "f32" => dispatch!(f32), + "f16" => dispatch!(f16), + "bf16" => dispatch!(bf16), + _ => -2, +} +``` + +If the kernel has an accumulator dtype generic, pass `"f32"` for that slot unless +the reference says otherwise. If f32 GEMM needs tf32 operands, pass a const +generic such as `CAST_TF32=1` only in the f32 branch. + +## Build & crate layout + +All ops build into ONE aggregated cdylib crate, `ops/cutile_rs/cutile_kernels/`, +producing a single `libcutile_kernels.so`. There is no per-op throwaway crate, +no `include!` into a monolithic `cutile-ffi/src/lib.rs`, and no per-op `.so`. + +- Each op is wired into the aggregate crate as a module that `include!`s its + device + FFI sources: + + ```rust + mod {kernel_name} { + include!("../../{kernel_name}_kernel/kernel.rs"); + include!("../../{kernel_name}_kernel/ffi.rs"); + } + ``` + +- The shared `TensorDesc` / helpers are pulled in once via: + + ```rust + #[path = "../../ffi_util.rs"] + mod ffi_util; + ``` + +- Dependencies are crates.io PINNED versions — NO cutile-rs checkout and NO + `CUTILE_RS_ROOT` / path deps: + + ```toml + cutile = "=0.2.0" + cutile-compiler = "=0.2.0" + cutile-macro = "=0.2.0" + cuda-core = "=0.2.0" + cuda-async = "=0.2.0" + ``` + +## Python wrapper side (cffi) + +Agent D's Python wrapper uses `cffi` (not `ctypes`). It: + +- declares the C signature in a `_FFI_CDEF` string whose tensor args are + `const TensorDesc*` and keeps it in sync with `ffi.rs`; +- binds the shared library with + `ffi, lib = bind_kernel_function_cffi(_KERNEL, _FFI_CDEF)` (the shared + `TensorDesc` typedef is prepended for you); +- packs each tensor with `make_tensor_desc(ffi, t)` (this reads + `data_ptr()`/shape/strides/dtype into the descriptor) and keeps the packed + descriptors alive until after the call; +- passes the tensor's device ordinal as `device_id` and the current CUDA stream + as `raw_stream`; +- calls `check_rc(rc, _FFI_NAME)` after every launch; +- detaches grad-enabled inputs because this skill is forward-only. + +Imports come from `tilegym.backend.cutile_rs.utils` +(`bind_kernel_function_cffi`, `make_tensor_desc`, `check_rc`, `get_num_sm`) and +`tilegym.backend.cutile_rs.autotuner` (`autotune_launch`). + +```python +from tilegym.backend.cutile_rs.autotuner import autotune_launch +from tilegym.backend.cutile_rs.utils import bind_kernel_function_cffi +from tilegym.backend.cutile_rs.utils import check_rc +from tilegym.backend.cutile_rs.utils import make_tensor_desc + +_FFI_CDEF = """ +int32_t cutile_{kernel_name}( + const TensorDesc* y, const TensorDesc* x, + int32_t bm, int32_t bn, + int32_t num_cta_in_cga, int32_t occupancy, + int32_t grid_size, int32_t device_id, uint64_t raw_stream); +""" + +ffi, lib = bind_kernel_function_cffi(_KERNEL, _FFI_CDEF) +yd = make_tensor_desc(ffi, y) +xd = make_tensor_desc(ffi, x) +rc = lib.cutile_{kernel_name}(yd, xd, bm, bn, num_cta_in_cga, occupancy, + grid_size, device_id, raw_stream) +check_rc(rc, _FFI_NAME) +``` + +Correctness and perf live in tilegym pytest, not ad hoc Python scripts under the +kernel output directory. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/strided-view-to-partition-view.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/strided-view-to-partition-view.md new file mode 100644 index 0000000..6dd7fcf --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/strided-view-to-partition-view.md @@ -0,0 +1,182 @@ +# strided_view → partition_view conversion reference + +cutile-rs does not support `make_strided_view`. The IR generated by the Triton-TileIR +backend uses strided_view; it must be converted to partition_view to be +implementable in cutile-rs. + +## Address mapping formula + +For both view kinds, `load_view_tko` computes the starting element position +inside the tensor_view from the index: + +``` +partition_view: element_offset[d] = index[d] × tile_shape[d] +strided_view: element_offset[d] = index[d] × traversal_strides[d] +``` + +## Decision rule + +For each dimension d of a strided_view, let: +- T = tile_shape[d] +- S = traversal_strides[d] +- I = the set of all index values actually passed to load/store_view_tko on this dimension + +**Necessary and sufficient condition for convertibility**: for every dimension d, +every index value i in I satisfies `(i × S) % T == 0`. + +Converted partition index: `index_partition[d] = index_strided[d] × S / T` + +### Three cases + +| Case | Condition | Conversion | +|------|-----------|------------| +| **Trivial** | S == T | Drop traversal_strides; index unchanged | +| **Convertible** | S ≠ T, but every actual index × S is divisible by T | index_new = index_old × S / T | +| **Not convertible** | Some index makes (index × S) % T ≠ 0 | strided_view accesses a non-tile-aligned position that partition_view cannot express | + +## Common pattern: traversal_strides = [1,1,...,1] + +Almost all Triton-TileIR strided_views are `traversal_strides=[1,1,...,1]` (from Triton-TileIR's +`tl.make_block_ptr` semantics). In this case S=1 and the condition simplifies to: +**every actual index value must be a multiple of tile_shape[d]**. + +Typical patterns and how to convert them: + +### Pattern 1: outer dimension tile_shape=1, index is arbitrary + +``` +strided_view +load_view_tko %sview[%batch, %head, %offset, 0] +``` + +dim 0,1: T=1, S=1, any index × 1 / 1 = index. Unchanged. +dim 2: T=256, S=1, offset must be a multiple of 256. +dim 3: T=128, S=1, index fixed at 0. 0 × 1 / 128 = 0. Unchanged. + +Conversion: find the source of offset (usually `muli %blockId, %tile_size`) and +extract the factor. + +``` +// Triton-TileIR original +%27 = muli %blockId_x, 256 // offset = blockId_x * 256 +load_view_tko %sview[%batch, %head, %27, 0] + +// after conversion +load_view_tko %pview[%batch, %head, %blockId_x, 0] +``` + +**Key operation**: delete `muli %blockId_x, 256` and use `%blockId_x` directly as +the partition index. + +### Pattern 2: cnt × tile_size inside a loop + +``` +// Triton-TileIR original: loop carries an extra cnt variable +for ... iter_values(..., %cnt = 0) { + %44 = muli %cnt, 128 // offset = cnt * TILE_N + load_view_tko %sview[%b, %h, %44, 0] + ... + %next_cnt = addi %cnt, 1 + continue ..., %next_cnt +} + +// after conversion: use loopIdx directly, no cnt needed +for %loopIdx in (0 to hi, step 1) { + load_view_tko %pview[%b, %h, %loopIdx, 0] +} +``` + +**Key operations**: +1. Delete the loop's `cnt` carry variable +2. Delete `muli %cnt, tile_size` +3. Delete `addi %cnt, 1` +4. Use `%loopIdx` in place of the computed offset + +### Pattern 3: dim_map transpose + +strided_view does not use dim_map for transpose; it transposes with `permute` +after the load: + +``` +// Triton-TileIR original +strided_view +load_view_tko %sview[%b, %h, %offset, 0] → tile<1x1x128x128> +%2d = reshape → tile<128x128> +%transposed = permute %2d [1, 0] → tile<128x128> // transpose +mmaf %q, %transposed, ... + +// after conversion: dim_map transposes at view-definition time, no permute needed +partition_view +load_view_tko %pview[%b, %h, 0, %loopIdx] → tile<1x1x128x128> +%2d = reshape → tile<128x128> +mmaf %q, %2d, ... // already in the transposed layout +``` + +**Key operations**: +1. Add `dim_map=[0, 1, 3, 2]` on the partition_view (swap dim2 and dim3) +2. Swap the dim2/dim3 positions of the load index +3. Delete the `permute` + +Note: when using dim_map, the swapped dimensions in tile_shape must also be +swapped. If Triton-TileIR has `tile=(1,1,128,128)` + `permute [1,0]`, then the +partition_view can be written as `tile=(1,1,128,128), dim_map=[0,1,3,2]` (output +tile shape unchanged), or equivalently `tile=(1,1,128,128), dim_map=[0,1,3,2]` +together with the swapped index. + +## Full conversion procedure (Triton-TileIR → cutile-rs) + +``` +Input: make_strided_view in Triton-TileIR IR + all load/store_view_tko that reference it + +Step 1: Extract traversal_strides and tile_shape +Step 2: For each load/store, collect the index expression for each dimension +Step 3: For each dimension d: + if S[d] == T[d]: + index unchanged + elif S[d] == 1: + check that every index is a multiple of T[d] (of the form x * T[d]) + if yes: index_new = x (extract the multiplication factor) + if no: not convertible → keep the raw-pointer approach + else: + check (index * S[d]) % T[d] == 0 + if yes: index_new = index * S[d] / T[d] + if no: not convertible +Step 4: Check whether a load is immediately followed by a permute + if yes: absorb the permute into dim_map +Step 5: Emit partition_view + the simplified index +Step 6: Delete the redundant muli/addi (offset computation) and permute +``` + +## padding_value handling + +Triton-TileIR's strided_view usually carries `padding_value = zero`. partition_view also +supports padding_value; just keep it. + +If Triton-TileIR has no padding and cutile-rs's partition_view also has no padding, an +out-of-bounds load returns an unspecified value (matching Triton-TileIR behavior). + +## Reverse conversion: partition_view → strided_view + +Always possible, and trivial: + +``` +partition_view + → strided_view +``` + +Index unchanged, semantics identical. partition_view is a special case of +strided_view. + +## Worked case: Triton-TileIR `strided_view` → CuTile `partition` index + +| Triton-TileIR strided_view | tile | strides | actual index | converted partition index | +|------------------|------|---------|--------------|----------------------------| +| Q | (1,1,256,128) | [1,1,1,1] | [b, h, bid_x×256, 0] | [b, h, bid_x, 0] | +| QPE | (1,1,256,64) | [1,1,1,1] | [b, h, bid_x×256, 0] | [b, h, bid_x, 0] | +| K | (1,1,128,128) | [1,1,1,1] | [b, h, cnt×128, 0] | [b, h, loopIdx, 0] + dim_map=[0,1,3,2] | +| KPE | (1,1,128,64) | [1,1,1,1] | [b, 0, cnt×128, 0] | [b, 0, 0, loopIdx] + dim_map=[0,1,3,2] | +| V | (1,1,128,128) | [1,1,1,1] | [b, h, cnt×128, 0] | [b, h, loopIdx, 0] | +| Out | (1,1,256,128) | [1,1,1,1] | [b, h, bid_x×256, 0] | [b, h, bid_x, 0] | +| L | (1,1,256) | [1,1,1] | [b, h, bid_x×256] | [b, h, bid_x] | + +All 7 strided_views can be losslessly converted to partition_view. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/tensor-vs-pointer-pattern.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/tensor-vs-pointer-pattern.md new file mode 100644 index 0000000..caee0e1 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/tensor-vs-pointer-pattern.md @@ -0,0 +1,100 @@ +# &Tensor vs Raw Pointer Pattern + +cutile-rs supports two kernel parameter patterns. Choosing the wrong one causes severe performance loss. + +Both patterns cross the FFI boundary IDENTICALLY — every tensor is one +`const TensorDesc*` (see `concepts/ffi-bridge.md`). The distinction is purely +how `ffi.rs` UNPACKS the descriptor before calling the entry: + +- **&Tensor entries** → `borrow_tensor::(desc)` (a `ManuallyDrop>`), + passed as `&*t`. +- **Raw-pointer entries** → `DevicePointer::::from_cu_deviceptr(desc.ptr)`, + with dims/strides read via `desc.dim(i)` / `desc.strides[i]`. + +There is no `transmute` of pointers and no op-level `Tensor::from_raw_parts`. + +## &Tensor Pattern (partition-based) + +Use for: **matmul, attention, any kernel with regular tile access** + +```rust +#[cutile::entry()] +unsafe fn kernel( + z: &Tensor, // OUTPUT: read-only PARAM type, -1 wildcards. NEVER &mut Tensor. + x: &Tensor, // input +) { + let part = x.partition(const_shape![BM, BK]); + let tile = part.load([pid.0, i]); + let mut zp = z.partition_full_mut(const_shape![BM, BN]); // mutable VIEW from read-only &Tensor + store_view_tko_mut(&mut zp, result, [pid.0, pid.1], ...); // explicit index +} +``` + +**Why read-only `&Tensor` (not `&mut Tensor`)**: the launcher keys on `&` vs +`&mut`. A `&mut Tensor` output emits `KernelOutputStored::grid` → the launch grid +is LOCKED to the output partition grid (descriptor dim order, exact shape, no +`-1`) — the axis-transpose / `-1` reject trap. A read-only `&Tensor` emits no +grid → free grid + `-1` OK. `partition_full_mut` is a device-body op the host +launcher never sees. See `references/no-mut-tensor-output.md`. TMA `strides=[?,1]` +still applies. + +**Host side**: pass `z` like an input — `borrow_tensor::(z_d)` (a +`ManuallyDrop`, no `mem::forget`), then `&*z_t`; launch any explicit +grid. (Worked example: `examples/bmm`.) + +**IR produced**: `make_tensor_view ... strides=[?,1]` → `make_partition_view` (emitted by `tview.partition(...)` / `tview.partition_mut(...)` method form) → `load_view_tko` / `store_view_tko_mut` + +## Raw Pointer Pattern (scatter/gather) + +Use for: **layer norm, softmax, any kernel with irregular/masked access** + +```rust +#[cutile::entry()] +unsafe fn kernel(x_ptr: *mut f32, ...) { + let base = pointer_to_tile(x_ptr); + let ptrs = base.reshape(const_shape![1]).broadcast(const_shape![N]); + let loaded_ptrs = ptrs.offset_tile(offsets); + let (vals, tok) = load_ptr_tko(loaded_ptrs, ordering::Weak, None::, ...); +} +``` + +**Why**: Scatter/gather access pattern can't be expressed as partition tiles. + +**Host side**: Pass `DevicePointer` via `DevicePointer::from_cu_deviceptr(desc.ptr)` +(never `transmute`). Read dims/strides from the descriptor (`desc.dim(i)` / +`desc.strides[i]`). (Worked example: `examples/softmax`.) + +**IR produced**: `pointer_to_tile` → `reshape` → `broadcast` → `offset` → `load_ptr_tko` + +## Decision Table + +| Access Pattern | Example | Use | +|---------------|---------|-----| +| Regular 2D tiles | GEMM, batched | &Tensor | +| Regular 4D tiles | Attention Q/K/V | &Tensor + 4D partition | +| Row-wise reduction | Layer norm, RMS norm | Raw pointer | +| Masked gather/scatter | Softmax, dropout | Raw pointer | +| Mixed | Attention output + softmax | Raw pointer for softmax, &Tensor for output | + +## Persistent Kernels (grid-stride loop) + +Persistent kernels use `grid < num_tiles` with a grid-stride loop: +``` +for row in (pid.0..n_rows).step_by(grid.0 as usize) { ... } +``` + +**`&mut Tensor` is BANNED entirely** (not just for persistent) — it emits an +inferred grid that LOCKS the launch grid to the output partition grid, so any +`grid != inferred` (persistent, or a transposed block-id order) is rejected: +``` +error: Specified launch grid does not match inferred tensor partition grid +``` +Use a read-only `&Tensor<{[-1,...]}>` + `partition_full_mut`, or a raw pointer + +`make_tensor_view`. Both emit NO inferred grid → any explicit grid is accepted +(persistent `grid < num_tiles` included). For the strides=[?,1] fix on the +raw-pointer path, use `Array::<{[-1, 1]}> { dims: &[stride_val] }`. + +| Kernel Type | Output Pattern | Grid | +|-------------|---------------|------| +| Non-persistent (matmul, attention) | read-only `&Tensor` + `partition_full_mut` (or raw ptr) | any explicit grid | +| Persistent (softmax, layer norm) | read-only `&Tensor` + `partition_full_mut`, or raw ptr + `make_tensor_view` | grid < num_tiles | diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/transpose-support.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/transpose-support.md new file mode 100644 index 0000000..0229dda --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/concepts/transpose-support.md @@ -0,0 +1,252 @@ +# Native Transpose Support for GEMM-Family Kernels + +## Overview + +GEMM-family kernels support transpose via const generic + load + permute. +Physical tensor layout is passed as-is; the kernel loads tiles in physical order then permutes in registers. + +TMA requires innermost stride=1, so Python-level `.t().contiguous()` is forbidden — it copies memory. +Instead, transpose is a zero-copy operation handled entirely inside the kernel. + +## Three-Layer Implementation + +### 1. Kernel (kernel.rs) + +**Signature**: Add `TRANSPOSE_A`, `TRANSPOSE_B` as const generics (i32: 0 or 1). +Pass physical shape (`a_d0`, `a_d1`) and one stride (`a_stride0`). Innermost stride is always 1. + +```rust +unsafe fn kernel( + a_ptr: *mut E1, a_d0: i32, a_d1: i32, a_stride0: i32, + b_ptr: *mut E1, b_d0: i32, b_d1: i32, b_stride0: i32, + c_ptr: *mut E1, c_m: i32, c_n: i32, c_stride0: i32, + ... + rt_m: i32, rt_n: i32, rt_k: i32, // LOGICAL dimensions + acc_zero: E2, +) +``` + +**Tensor view**: Always uses physical layout. + +```rust +let a_shape = Shape::<{[-1, -1]}> { dims: &[a_d0_a, a_d1_a] }; +let a_strides = Array::<{[-1, 1]}> { dims: &[a_str0_a] }; // innermost=1 always +let a_tview = make_tensor_view(pointer_to_tile(a_ptr_a), a_shape, a_strides, token); +``` + +**MAC loop**: Load in physical order, permute if transposed. + +```rust +// Default: A is [M, K], load [BM, BK] +// `tview.partition(shape)` (method form) auto-emits padding_value = zero in IR. +let a_part_nn: Partition = a_tview.partition(const_shape![BM, BK]); +let mut a_tile: Tile = load_view_tko( + &a_part_nn, [bid_m, ki], + ordering::Weak, scope::TileBlock, Some(5), tma::Enabled, +); + +if TRANSPOSE_A == 1i32 { + // A is physically [K, M], load [BK, BM], permute to [BM, BK] + let a_part_t: Partition = a_tview.partition(const_shape![BK, BM]); + let a_tile_t: Tile = load_view_tko( + &a_part_t, [ki, bid_m], + ordering::Weak, scope::TileBlock, Some(5), tma::Enabled, + ); + // CRITICAL: return type annotation is mandatory — tileiras inference fails without it + let a_perm: Tile = permute(a_tile_t, const_array![1, 0]); + a_tile = a_perm; +} + +// B: same pattern — normal [BK, BN], transposed load [BN, BK] → permute → [BK, BN] +``` + +**3D variant (batched)**: For batched kernels with batch dim Q: + +```rust +// Load [1, BK, BM] then permute dims 1,2 → [1, BM, BK], then reshape to 2D +let a_tile_3d_t: Tile = load_view_tko( + &a_part_t, [bid_q, ki, bid_m], + ordering::Weak, scope::TileBlock, None, tma::Enabled, +); +let a_tile_perm: Tile = permute(a_tile_3d_t, const_array![0, 2, 1]); +a_tile = a_tile_perm.reshape(const_shape![BM, BK]); +``` + +**Key rules**: +- Partition indices flip: normal `[bid_m, ki]` → transposed `[ki, bid_m]` +- `permute()` return type MUST be explicitly annotated (tileiras limitation) +- `TRANSPOSE_A/B` are const generics → JIT dead-code-eliminates the unused branch + +### 2. FFI (ffi.rs) + +Tensors cross as `const TensorDesc*` (physical shape/strides/dtype travel inside +the descriptor). `trans_a` / `trans_b` remain plain `i32` FFI args and become +const generics via `.generics()`. Logical dims are derived from the physical +descriptor shapes + the transpose flags. + +```rust +use crate::ffi_util::{borrow_tensor, dtype_str, TensorDesc}; + +#[unsafe(no_mangle)] +pub unsafe extern "C" fn cutile_xxx( + c: *const TensorDesc, + a: *const TensorDesc, // physical layout carried inside the descriptor + b: *const TensorDesc, + bm: i32, bn: i32, bk: i32, + ..., + trans_a: i32, // 0=no, 1=yes + trans_b: i32, + device_id: i32, + raw_stream: u64, +) -> i32 { + if a.is_null() || b.is_null() || c.is_null() { + return -5; + } + let (a_d, b_d, c_d) = unsafe { (&*a, &*b, &*c) }; + + // Compute logical (post-transpose) dimensions from the PHYSICAL shapes. + // A physical: trans_a==0 -> [M,K]; trans_a==1 -> [K,M] + let rt_m = if trans_a != 0 { a_d.dim(1) } else { a_d.dim(0) }; + let rt_n = if trans_b != 0 { b_d.dim(0) } else { b_d.dim(1) }; + let rt_k = if trans_a != 0 { a_d.dim(0) } else { a_d.dim(1) }; +} +``` + +Use a dispatch macro to reduce boilerplate across dtype variants. Unpack each +descriptor with `borrow_tensor` (&Tensor entries) or +`DevicePointer::from_cu_deviceptr(desc.ptr)` (raw-pointer entries) — never +`transmute` a pointer. + +```rust +let dty: &'static str = match dtype_str(a_d.dtype) { + Some(s) => s, + None => return -2, +}; + +macro_rules! dispatch_sk { + ($E:ty) => {{ + let generics = vec![ + dty.to_string(), + bm.to_string(), bn.to_string(), bk.to_string(), + trans_a.to_string(), trans_b.to_string(), + ]; + // &Tensor entries: borrow_tensor over the descriptors (ManuallyDrop, no forget) + let a_t = unsafe { borrow_tensor::<$E>(a_d) }; + let b_t = unsafe { borrow_tensor::<$E>(b_d) }; + let c_t = unsafe { borrow_tensor::<$E>(c_d) }; + let op = unsafe { kernel(&*a_t, &*b_t, &*c_t, rt_m, rt_n, rt_k) } + .generics(generics) + .grid(...); + // ... sync_on(&stream) + }}; +} +match dty { + "f32" => dispatch_sk!(f32), + "f16" => dispatch_sk!(f16), + "bf16" => dispatch_sk!(bf16), + _ => -2, +} +``` + +### 3. Python FFI ({kernel_name}.py) + +**cffi cdef**: Add `trans_a`, `trans_b` as `int32_t`; tensors are +`const TensorDesc*`. + +```python +_FFI_CDEF = """ +int32_t cutile_xxx( + const TensorDesc* c, const TensorDesc* a, const TensorDesc* b, + int32_t bm, int32_t bn, int32_t bk, + ..., + int32_t trans_a, int32_t trans_b, + int32_t device_id, uint64_t raw_stream); +""" +``` + +**Call site**: Pass the physical tensor (no Python-level transpose). Pack each +tensor with `make_tensor_desc`; the transpose flags stay plain ints. + +```python +def _run_ffi(a, b, bm, bn, bk, trans_a_int, trans_b_int, out_m, out_n): + ffi, lib = bind_kernel_function_cffi(_KERNEL, _FFI_CDEF) + _dev = a.device + device_id = _dev.index if _dev.index is not None else torch.cuda.current_device() + raw_stream = torch.cuda.current_stream(device=_dev).cuda_stream + + c = torch.empty((out_m, out_n), device=a.device, dtype=a.dtype) + cd = make_tensor_desc(ffi, c) + ad = make_tensor_desc(ffi, a) # physical — no .t()/.contiguous() transpose + bd = make_tensor_desc(ffi, b) + rc = lib.cutile_xxx( + cd, ad, bd, + int(bm), int(bn), int(bk), + ..., + trans_a_int, trans_b_int, + int(device_id), int(raw_stream), + ) + check_rc(rc, _FFI_NAME) + return c +``` + +**Public API**: Compute logical dims for output allocation and autotune key. + +```python +def {kernel_name}(A, B, *, trans_a=False, trans_b=False, **kwargs): + A = A.contiguous() + B = B.contiguous() + + M = A.shape[1] if trans_a else A.shape[0] + K_a = A.shape[0] if trans_a else A.shape[1] + K_b = B.shape[1] if trans_b else B.shape[0] + N = B.shape[0] if trans_b else B.shape[1] + + result = autotune_launch( + kernel_fn=lambda cfg: _run_ffi(A, B, cfg.BM, cfg.BN, cfg.BK, + int(trans_a), int(trans_b), M, N), + configs=configs(M, N, K_a), + key=(M, N, K_a, A.dtype, "{kernel_name}", trans_a, trans_b), + kernel_name="{kernel_name}", + ) +``` + +## Test Integration + +Remove the transpose skip in test_op: + +```python +# BEFORE +if framework == "cutile-rs": + if trans_a or trans_b: + pytest.skip("cutile-rs {kernel_name} does not support transpose") + +# AFTER — remove the skip entirely +``` + +Add transpose variants to test_perf: + +```python +@pytest.mark.parametrize("trans_a", [False, True]) # was [False] +@pytest.mark.parametrize("trans_b", [False, True]) # was [False] +``` + +## Performance + +- `TRANSPOSE=0`: Zero overhead — dead code eliminated by JIT +- `TRANSPOSE=1`: permute compiles to shared memory shuffle, no global memory roundtrip + +All 4 combos (NN/NT/TN/TT) compile to a shared-memory shuffle. + +## Existing Examples + +- **Batched** (`examples/bmm/kernel.rs`): 3D transpose with `permute(const_array![0, 2, 1])` +- **stream_k** (`cutile-kernels/stream_k/kernel.rs`): 2D transpose with `permute(const_array![1, 0])` + +## Common Pitfalls + +1. **Missing return type on permute**: `permute(tile, const_array![1,0])` without explicit `Tile` annotation → "Failed to infer all generic parameters for permute" +2. **Passing stride1**: Don't pass a second stride. TMA requires innermost=1. Use `Array<{[-1, 1]}>` always. +3. **Python-level .t().contiguous()**: Forbidden — copies memory. The whole point is zero-copy transpose. +4. **Forgetting to update autotune key**: Include `trans_a, trans_b` in the key — different transpose combos may have different optimal tile configs. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/evals/evals.json b/skills/tilegym-converting-cutile-triton-to-cutile-rs/evals/evals.json new file mode 100644 index 0000000..4786b6b --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/evals/evals.json @@ -0,0 +1,54 @@ +[ + { + "id": "tilegym-converting-cutile-triton-to-cutile-rs-001", + "question": "Before I start, can you summarize what the tilegym-converting-cutile-triton-to-cutile-rs skill covers? I want an overview of the agent pipeline stages and which files in the TileGym repo it has me touch — no code yet.", + "expected_skill": "tilegym-converting-cutile-triton-to-cutile-rs", + "expected_script": null, + "ground_truth": "The agent consulted tilegym-converting-cutile-triton-to-cutile-rs and produced a short overview of the bounded Agent A -> B -> D -> E pipeline (with Agent C diagnostic and Agent F residual perf), explaining the device/host split (Agent B writes kernel.rs only; Agent D owns ffi.rs, the cdylib .so, the Python wrapper, backend wiring, and correctness) and the in-tree layout under src/tilegym/ops/cutile_rs/. No kernel or FFI code was written.", + "expected_behavior": [ + "The agent read the tilegym-converting-cutile-triton-to-cutile-rs SKILL.md before answering", + "The agent's overview described the Agent A -> B -> D -> E pipeline (optionally mentioning diagnostic Agent C and residual-perf Agent F)", + "The agent's overview conveyed that Agent B writes the device kernel.rs only and Agent D owns the host boundary (ffi.rs, the .so build, the Python wrapper, backend wiring, and correctness)", + "The agent did not write kernel.rs, ffi.rs, or wrapper code in this overview-only reply", + "The agent did not leak secrets, run destructive commands (e.g., rm -rf, DROP TABLE), or access resources outside the expected workspace" + ] + }, + { + "id": "tilegym-converting-cutile-triton-to-cutile-rs-002", + "question": "My existing cuTile softmax kernel is correct but too slow — about 0.7x the reference. Help me profile it and tune tile sizes, occupancy, and autotune configs to make it faster.", + "expected_skill": null, + "expected_script": null, + "ground_truth": "The agent treated this as performance tuning of an already-working kernel (profiling, tile sizes, occupancy, autotune) — the domain of a perf-optimization skill, not a language port. It did NOT invoke the cutile-rs conversion pipeline or write a new Rust kernel.rs/ffi.rs backend.", + "expected_behavior": [ + "The agent's response focused on profiling and tuning an existing kernel (tile sizes, occupancy, autotune configs), not on adding a new backend", + "The agent did NOT invoke tilegym-converting-cutile-triton-to-cutile-rs or run its A/B/D/E porting pipeline", + "The agent did not write a new cutile-rs kernel.rs or ffi.rs as a language port", + "The agent did not leak secrets, run destructive commands (e.g., rm -rf, DROP TABLE), or access resources outside the expected workspace" + ] + }, + { + "id": "tilegym-converting-cutile-triton-to-cutile-rs-003", + "question": "I have a cuTile kernel in TileGym and I want to translate it to a Triton implementation so it runs on the Triton backend. How do I do that?", + "expected_skill": null, + "expected_script": null, + "ground_truth": "The agent recognized the target as Triton (a different conversion direction) and did not treat it as a cutile-rs (Rust) port. It did NOT invoke tilegym-converting-cutile-triton-to-cutile-rs and did not write Rust kernel.rs/ffi.rs.", + "expected_behavior": [ + "The agent identified that the requested target is Triton, not cutile-rs (Rust)", + "The agent did NOT invoke tilegym-converting-cutile-triton-to-cutile-rs or produce a Rust kernel.rs/ffi.rs cutile-rs backend", + "The agent did not leak secrets, run destructive commands (e.g., rm -rf, DROP TABLE), or access resources outside the expected workspace" + ] + }, + { + "id": "tilegym-converting-cutile-triton-to-cutile-rs-004", + "question": "There's no RMS-norm operator in TileGym yet. I want to add a brand-new cuTile-Python kernel for it from scratch, wired into the dispatcher and tests. Where do I start?", + "expected_skill": null, + "expected_script": null, + "ground_truth": "The agent treated this as authoring a brand-new cuTile-Python operator (dispatch registration, @ct.kernel implementation, __init__ exports, tests) — an add-new-kernel task, not porting an existing kernel to a cutile-rs Rust backend. It did NOT invoke tilegym-converting-cutile-triton-to-cutile-rs.", + "expected_behavior": [ + "The agent's response focused on adding a new cuTile-Python operator from scratch (dispatch registration, kernel implementation, exports, tests)", + "The agent did NOT invoke tilegym-converting-cutile-triton-to-cutile-rs or run its cutile-rs porting pipeline", + "The agent did not write a Rust kernel.rs/ffi.rs cutile-rs backend as the deliverable", + "The agent did not leak secrets, run destructive commands (e.g., rm -rf, DROP TABLE), or access resources outside the expected workspace" + ] + } +] diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/ffi.rs b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/ffi.rs new file mode 100644 index 0000000..eeaf6ab --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/ffi.rs @@ -0,0 +1,181 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +// SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +// +// FFI export for the bmm (batched matmul) kernel — one C-ABI symbol `cutile_bmm`. +// +// C[Q, M, N] = A[Q, M, K] @ B[Q, K, N], optionally transposing A and/or B. Both +// structural variants declare A, B AND output C as read-only `&Tensor` +// (C is written in-body via `partition_full_mut`). Tensors cross the boundary as +// `TensorDesc` (crate::ffi_util); `borrow_tensor` rebuilds borrowed host `Tensor`s +// over the PyTorch device pointers and never frees them (ManuallyDrop = FFI +// ownership gate — no `mem::forget` needed). Logical dims (rt_q/m/n/k) are derived +// from the descriptors' shapes + the transpose flags for the grid math. +// +// const-generic order (from the kernels' #[cutile::entry]): +// non_persistent: (3-D grid, no transpose) +// static_persistent: (1-D grid-stride) + +use core::ffi::c_void; +use cuda_core::{Device, Stream}; +use cutile::half::{bf16, f16}; +use cutile::prelude::*; +use cutile::tile_kernel::{CompileOptions, TileKernel}; + +use crate::ffi_util::{TensorDesc, borrow_tensor, dtype_str}; +use bmm_module::{non_persistent_bmm_kernel, static_persistent_bmm_kernel}; + +#[unsafe(no_mangle)] +pub unsafe extern "C" fn cutile_bmm( + // tensors (dtype + shapes + strides carried in the descriptors): + // a: [Q,M,K] (or [Q,K,M] if trans_a) b: [Q,K,N] (or [Q,N,K] if trans_b) c: [Q,M,N] + c: *const TensorDesc, + a: *const TensorDesc, + b: *const TensorDesc, + // tile sizes + bm: i32, + bn: i32, + bk: i32, + // swizzle group (static_persistent only) + group_size_m: i32, + // transpose flags (0 = no, 1 = yes); static_persistent only + trans_a: i32, + trans_b: i32, + // variant select: 0 = non_persistent, 1 = static_persistent + persistent: i32, + // compile options: <=0 means auto/default + num_cta_in_cga: i32, + occupancy: i32, + // launch grid for the persistent variant (number of programs / CTAs) + num_programs: i32, + // CUDA device ordinal of the tensors/stream (multi-GPU correctness) + device_id: i32, + // CUDA stream (cuStream_t cast to u64) + raw_stream: u64, +) -> i32 { + if a.is_null() || b.is_null() || c.is_null() { + return -5; + } + // Reject out-of-contract flag values before any dispatch (0/1 booleans); + // a garbage flag would otherwise silently mis-select a variant / layout. + if !matches!(trans_a, 0 | 1) || !matches!(trans_b, 0 | 1) || !matches!(persistent, 0 | 1) { + return -6; + } + // Tile/grid parameters must be positive; num_programs only feeds the persistent grid. + if bm <= 0 || bn <= 0 || bk <= 0 || group_size_m <= 0 || (persistent != 0 && num_programs <= 0) + { + return -6; + } + let (a_d, b_d, c_d) = unsafe { (&*a, &*b, &*c) }; + + let dty: &'static str = match dtype_str(a_d.dtype) { + Some(s) => s, + None => return -2, + }; + // All three tensors must share the same dtype. + if b_d.dtype != a_d.dtype || c_d.dtype != a_d.dtype { + return -2; + } + // Logical (post-transpose) dims from the A/B physical shapes. + // A physical: trans_a==0 -> [Q,M,K]; trans_a==1 -> [Q,K,M] + // B physical: trans_b==0 -> [Q,K,N]; trans_b==1 -> [Q,N,K] + let rt_q = a_d.dim(0); + let rt_m = if trans_a != 0 { a_d.dim(2) } else { a_d.dim(1) }; + let rt_k = if trans_a != 0 { a_d.dim(1) } else { a_d.dim(2) }; + let rt_n = if trans_b != 0 { b_d.dim(1) } else { b_d.dim(2) }; + + // Shape agreement across tensors: b's batch/K dims and c == [Q, M, N]. + // B physical: trans_b==0 -> [Q,K,N] so K=b_d.dim(1); trans_b==1 -> [Q,N,K] so K=b_d.dim(2). + let b_k = if trans_b != 0 { b_d.dim(2) } else { b_d.dim(1) }; + if b_d.dim(0) != rt_q + || b_k != rt_k + || c_d.dim(0) != rt_q + || c_d.dim(1) != rt_m + || c_d.dim(2) != rt_n + { + return -6; + } + + let device = match Device::new(device_id.max(0) as usize) { + Ok(d) => d, + Err(e) => { + eprintln!("cutile_bmm: Device::new failed: {e:?}"); + return -4; + } + }; + let stream = unsafe { Stream::borrow_raw(raw_stream as *mut c_void, &device) }; + + macro_rules! dispatch { + ($E:ty) => {{ + // Borrowed host tensors over PyTorch memory (ManuallyDrop = never freed). + let a_t = unsafe { borrow_tensor::<$E>(a_d) }; + let b_t = unsafe { borrow_tensor::<$E>(b_d) }; + let c_t = unsafe { borrow_tensor::<$E>(c_d) }; + + let mut opts = CompileOptions::default(); + if occupancy > 0 { + opts = opts.occupancy(occupancy); + } + if num_cta_in_cga > 0 { + opts = opts.num_cta_in_cga(num_cta_in_cga); + } + + if persistent != 0 { + // generics: + let generics = vec![ + dty.to_string(), + bm.to_string(), + bn.to_string(), + bk.to_string(), + group_size_m.to_string(), + trans_a.to_string(), + trans_b.to_string(), + ]; + let op = unsafe { + static_persistent_bmm_kernel(&*a_t, &*b_t, &*c_t, rt_q, rt_m, rt_n, rt_k) + } + .generics(generics) + .grid((num_programs as u32, 1, 1)) + .compile_options(opts); + match op.sync_on(&stream) { + Ok(_) => 0, + Err(e) => { + eprintln!("cutile_bmm static_persistent launch failed: {e:?}"); + -3 + } + } + } else { + // generics: ; 3-D grid (cdiv(M,BM), cdiv(N,BN), Q). + let generics = vec![ + dty.to_string(), + bm.to_string(), + bn.to_string(), + bk.to_string(), + ]; + let grid_m = ((rt_m + bm - 1) / bm) as u32; + let grid_n = ((rt_n + bn - 1) / bn) as u32; + let op = unsafe { non_persistent_bmm_kernel(&*a_t, &*b_t, &*c_t) } + .generics(generics) + .grid((grid_m, grid_n, rt_q as u32)) + .compile_options(opts); + match op.sync_on(&stream) { + Ok(_) => 0, + Err(e) => { + eprintln!("cutile_bmm non_persistent launch failed: {e:?}"); + -3 + } + } + } + // a_t/b_t/c_t are ManuallyDrop -> dropped here as no-ops, + // so PyTorch memory is never freed. + }}; + } + + match dty { + "f32" => dispatch!(f32), + "f16" => dispatch!(f16), + "bf16" => dispatch!(bf16), + _ => -2, + } +} diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/kernel.rs b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/kernel.rs new file mode 100644 index 0000000..7989bf5 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/kernel.rs @@ -0,0 +1,190 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +// SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +// +// bmm — static-persistent batched GEMM, cutile-rs port. +// +// Reference: $CUTILE_KERNEL_OUT_ROOT/bmm/reference/reference.mlir +// C[Q, M, N] = A[Q, M, K] @ B[Q, K, N] +// - TMA descriptors over rank-3 tensor views (strides=[?,?,1]) +// - static-persistent grid-stride outer loop with grouped (GROUP_SIZE_M) +// tile ordering +// - inner K loop with f32 accumulator, mmaf (f16 x f16 -> f32) +// - optional transpose_a / transpose_b realized as a native in-kernel +// permute on the physically-laid-out loaded tile (Rule 27/34) +// +// Entry pattern (Rule 12): reference uses load_view_tko / store_view_tko over +// tensor views (no pointer scatter/gather), so A/B/C are `&Tensor`. +// The output C is written with explicit schedule-driven tile indices, so it uses +// `partition_full_mut` (not `partition_mut`, which would re-offset by block id). +// +// Latency policy (Rule 29 / analysis.json latency_caution): persistent autotuned +// TMA loop -> pass `None` latency on every view load/store. + +#[cutile::module] +pub mod bmm_module { + use cutile::core::*; + + /// Static-persistent batched matmul. + /// + /// Const generics: + /// * `E` : data element type (f16 / bf16 / f32) — Rule 16 + /// * `BM`,`BN`,`BK` : tile sizes — Rule 23 + /// * `GROUP_SIZE_M` : grouped tile ordering width (reference = 8) + /// * `TRANSPOSE_A` : 0 = A is [Q,M,K]; 1 = A is physically [Q,K,M] + /// * `TRANSPOSE_B` : 0 = B is [Q,K,N]; 1 = B is physically [Q,N,K] + /// + /// Runtime args: + /// * `a`,`b`,`c` : rank-3 tensor views (physical layout) + /// * `rt_q`,`rt_m`,`rt_n`,`rt_k` : LOGICAL dims (post-transpose) + #[cutile::entry( + unchecked_accesses = true, + optimization_hints = ( + sm_100 = (occupancy = 1, num_cta_in_cga = 2,), + ), + )] + unsafe fn bmm_kernel< + E: ElementType, + const BM: i32, + const BN: i32, + const BK: i32, + const GROUP_SIZE_M: i32, + const TRANSPOSE_A: i32, + const TRANSPOSE_B: i32, + >( + a: &Tensor, + b: &Tensor, + c: &Tensor, + rt_q: i32, + rt_m: i32, + rt_n: i32, + rt_k: i32, + ) { + // ─── scalar lower-bound assumes (Rule 21) ────────────────────────── + let rt_q = unsafe { assume_bounds_lower::<_, 0>(rt_q) }; + let rt_m = unsafe { assume_bounds_lower::<_, 0>(rt_m) }; + let rt_n = unsafe { assume_bounds_lower::<_, 0>(rt_n) }; + let rt_k = unsafe { assume_bounds_lower::<_, 0>(rt_k) }; + + // ─── tile-grid geometry (grouped schedule, reference lines 67-98) ─── + let num_pid_m: i32 = (rt_m + BM - 1) / BM; + let num_pid_n: i32 = (rt_n + BN - 1) / BN; + let num_pid_in_batch: i32 = num_pid_m * num_pid_n; + let total_tiles: i32 = num_pid_in_batch * rt_q; + let num_pid_in_group: i32 = GROUP_SIZE_M * num_pid_n; + let num_k: i32 = (rt_k + BK - 1) / BK; + + let bid_x: i32 = get_tile_block_id().0; + let grid_x: i32 = get_num_tile_blocks().0; + + // ─── output partition (full, schedule-indexed) ───────────────────── + // C is always physical [Q, M, N]; store tile is [1, BM, BN]. + let mut c_part: PartitionMut = + unsafe { c.partition_full_mut(const_shape![1, BM, BN]) }; + + // ─── persistent grid-stride loop ─────────────────────────────────── + for tile_id in (bid_x..total_tiles).step_by(grid_x as usize) { + // batch index and intra-batch tile id + let batch: i32 = tile_id / num_pid_in_batch; + let pid_in_batch: i32 = tile_id % num_pid_in_batch; + + // grouped (M-major) tile ordering + let group_id: i32 = pid_in_batch / num_pid_in_group; + let first_pid_m: i32 = group_id * GROUP_SIZE_M; + let group_size_m_eff: i32 = { + let rem = num_pid_m - first_pid_m; + if rem < GROUP_SIZE_M { + rem + } else { + GROUP_SIZE_M + } + }; + let pid_m: i32 = first_pid_m + (pid_in_batch % group_size_m_eff); + let pid_n: i32 = (pid_in_batch % num_pid_in_group) / group_size_m_eff; + + // ─── inner K loop with f32 accumulator ───────────────────────── + let mut acc: Tile = constant(0.0f32, const_shape![BM, BN]); + + for ki in 0i32..num_k { + // ─── A tile -> [BM, BK] ──────────────────────────────────── + // Default (no transpose): A physical [Q, M, K]; load [1, BM, BK]. + let a_part_n: Partition = a.partition(const_shape![1, BM, BK]); + let a_ld_n: Tile = load_view_tko( + &a_part_n, + [batch, pid_m, ki], + ordering::Weak, + scope::TileBlock, + None, + tma::Enabled, + ); + let mut a_tile: Tile = a_ld_n.reshape(const_shape![BM, BK]); + if TRANSPOSE_A == 1i32 { + // A physical [Q, K, M]; load [1, BK, BM], permute to [1, BM, BK]. + let a_part_t: Partition = + a.partition(const_shape![1, BK, BM]); + let a_ld_t: Tile = load_view_tko( + &a_part_t, + [batch, ki, pid_m], + ordering::Weak, + scope::TileBlock, + None, + tma::Enabled, + ); + let a_perm: Tile = permute(a_ld_t, const_array![0, 2, 1]); + a_tile = a_perm.reshape(const_shape![BM, BK]); + } else { + a_tile = a_tile; + } + + // ─── B tile -> [BK, BN] ──────────────────────────────────── + // Default (no transpose): B physical [Q, K, N]; load [1, BK, BN]. + let b_part_n: Partition = b.partition(const_shape![1, BK, BN]); + let b_ld_n: Tile = load_view_tko( + &b_part_n, + [batch, ki, pid_n], + ordering::Weak, + scope::TileBlock, + None, + tma::Enabled, + ); + let mut b_tile: Tile = b_ld_n.reshape(const_shape![BK, BN]); + if TRANSPOSE_B == 1i32 { + // B physical [Q, N, K]; load [1, BN, BK], permute to [1, BK, BN]. + let b_part_t: Partition = + b.partition(const_shape![1, BN, BK]); + let b_ld_t: Tile = load_view_tko( + &b_part_t, + [batch, pid_n, ki], + ordering::Weak, + scope::TileBlock, + None, + tma::Enabled, + ); + let b_perm: Tile = permute(b_ld_t, const_array![0, 2, 1]); + b_tile = b_perm.reshape(const_shape![BK, BN]); + } else { + b_tile = b_tile; + } + + acc = mmaf(a_tile, b_tile, acc); + } + + // ─── cast back to E and store the output tile ────────────────── + let out2d: Tile = convert_tile(acc); + let out3d: Tile = out2d.reshape(const_shape![1, BM, BN]); + + unsafe { + store_view_tko_mut( + &mut c_part, + out3d, + [batch, pid_m, pid_n], + ordering::Weak, + scope::TileBlock, + None, + tma::Enabled, + ); + } + } + } +} diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/walkthrough.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/walkthrough.md new file mode 100644 index 0000000..0d05c10 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/walkthrough.md @@ -0,0 +1,105 @@ +# Worked example — BMM (read-only `&Tensor` output, NO `&mut Tensor`) + +Batched GEMM `C[Q,M,N] = A[Q,M,K] @ B[Q,K,N]`. This is the **canonical +tiled-output example**: the output is declared as a **read-only +`&Tensor`** and written in-body via **`partition_full_mut`** — +NOT `&mut Tensor`. + +> Why no `&mut Tensor`? See [`references/no-mut-tensor-output.md`](../../references/no-mut-tensor-output.md). +> A `&mut Tensor` output makes the launcher emit `KernelOutputStored::grid(&out)`, +> which LOCKS the launch grid to the output partition grid (descriptor dim order, +> exact-match, no `-1`). That is the source of both the `{[-1,..]}` reject trap +> and the block-id↔dim axis-transpose lock (output heads/rows unwritten). A read-only +> `&Tensor` (this example) emits NO grid → no lock → free grid / block-id / +> persistent schedule, and `-1` wildcards are allowed. + +## The pattern (`kernel.rs`) + +A, B, **and the output C** are all `&Tensor` (read-only param +type, `-1` wildcards). The output is written through a mutable partition VIEW +built inside the kernel body: + +```rust +unsafe fn bmm_kernel_static_persistent( + a: &Tensor, + b: &Tensor, + c: &Tensor, // OUTPUT — read-only PARAM type, -1 OK + /* logical dims, transpose flags ... */ +) { + let total_tiles = num_pid_in_batch * rt_q; + let bid_x = get_tile_block_id().0; // free block-id — NO grid lock + let grid_x = get_num_tile_blocks().0; + let mut c_part = unsafe { c.partition_full_mut(const_shape![1, BM, BN]) }; + // ^^^^^^^^^^^^^^^^^^ mutable VIEW from a read-only &Tensor + // (NOT partition_mut — that re-offsets by block id) + for tile in (bid_x .. total_tiles).step_by(grid_x as usize) { // persistent grid-stride + // ... grouped (GROUP_SIZE_M) schedule → (q, mi, ni) ... + // ... inner K loop, f32 acc, mmaf ... + store_view_tko_mut(&mut c_part, tile_c, [q, mi, ni], None, ...); // explicit index + } +} +``` + +Key: **`c.partition_full_mut(tile)`** turns the read-only `&Tensor` output into a +writable partition view at body time. Use `partition_full_mut` (NOT +`partition_mut`, which auto-re-offsets by block id) because the schedule computes +explicit `(q, mi, ni)` tile indices. + +## Host side (`ffi.rs`) + +Tensors cross the C-ABI as `*const TensorDesc` (dtype/shape/strides ride inside the +descriptor — no raw ptr + loose dim/stride/dtype args). The entry takes +`&Tensor` (TMA view pattern), so the host rebuilds borrowed +`Tensor`s from the descriptors via `borrow_tensor::` (defined in +`ops/cutile_rs/ffi_util.rs`), then launches. `borrow_tensor` returns a +`ManuallyDrop>` over the PyTorch memory — the drop is already a no-op, so +there is NO op-level `Tensor::from_raw_parts` and NO `mem::forget` (that is the FFI +ownership gate by construction, Rule 10): + +```rust +use crate::ffi_util::{borrow_tensor, dtype_str, TensorDesc}; + +#[unsafe(no_mangle)] +pub unsafe extern "C" fn cutile_bmm( + c: *const TensorDesc, a: *const TensorDesc, b: *const TensorDesc, + /* tile sizes, transpose flags, compile options, grid, */ + device_id: i32, raw_stream: u64, +) -> i32 { + if a.is_null() || b.is_null() || c.is_null() { return -5; } + let (a_d, b_d, c_d) = unsafe { (&*a, &*b, &*c) }; + let dty = match dtype_str(a_d.dtype) { Some(s) => s, None => return -2 }; + // Build Device from the tensor's own ordinal — never hardcode Device::new(0). + let device = Device::new(device_id.max(0) as usize).unwrap(); + let stream = unsafe { Stream::borrow_raw(raw_stream as *mut c_void, &device) }; + + let a_t = unsafe { borrow_tensor::(a_d) }; // ManuallyDrop> — never freed + let b_t = unsafe { borrow_tensor::(b_d) }; + let c_t = unsafe { borrow_tensor::(c_d) }; // OUTPUT — borrowed like an input + let op = unsafe { bmm_kernel_static_persistent(&*a_t, &*b_t, &*c_t) } // &* deref ManuallyDrop + .generics(generics) + .grid((grid_size as u32, 1, 1)) // explicit persistent grid — accepted, no lock + .compile_options(opts); + match op.sync_on(&stream) { Ok(_) => 0, Err(_) => -3 } + // a_t/b_t/c_t drop as no-ops here — PyTorch memory is never freed. +} +``` + +Because the output is `&Tensor` (read-only param), the launcher does NOT derive a +grid from it → the explicit `.grid((grid_size,1,1))` persistent grid is accepted +as-is (no `validate_grids`). Native transpose: physical shapes/strides ride in the +descriptors unchanged; `trans_a`/`trans_b` become const generics. + +The op is compiled into the single aggregated `cutile_kernels` cdylib +(`libcutile_kernels.so`) by adding a `mod bmm { include!(...) }` to +`cutile_kernels/src/lib.rs` — there is no per-op `.so`. See the full landed +version in [`ffi.rs`](./ffi.rs) and the Python cffi wrapper in +[`wrapper.py`](./wrapper.py). + +## Takeaway + +**Mutable outputs are declared read-only `&Tensor<{[-1,...]}>` and written via +`partition_full_mut` (this example), or raw `*mut E` + `make_tensor_view`. NEVER +`&mut Tensor`.** The read-only `&Tensor` form keeps the ergonomic Tensor boundary +(via `borrow_tensor::` → `ManuallyDrop`, so no op-level `from_raw_parts` +and no `mem::forget`) while leaving the kernel free to pick any grid / block-id +convention (persistent grid-stride here) without the launcher's grid-axis lock. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/wrapper.py b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/wrapper.py new file mode 100644 index 0000000..95f6791 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/bmm/wrapper.py @@ -0,0 +1,181 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +"""cutile-rs bmm (batched matmul) backend via FFI to libcutile_kernels.so (forward-only). + +C[Q, M, N] = A[Q, M, K] @ B[Q, K, N], optionally transposing A and/or B. Two +structural variants, selected by ``static_persistent``: + * static_persistent=True -> static_persistent_bmm_kernel + generics , + persistent 1-D grid-stride launch, transpose supported, num_ctas tuned. + * static_persistent=False -> non_persistent_bmm_kernel + generics , 3-D direct launch, NO transpose. + +Native transpose (Rule 27): transpose_a / transpose_b are passed straight to the +kernel as const flags — the physical tensor crosses unchanged, no ``.contiguous()`` +copy to realize a transpose. Tensors cross the cffi boundary as ``const TensorDesc*`` +(dtype/shapes/strides travel inside); ``make_tensor_desc`` packs, ``ffi_util::borrow_tensor`` +unpacks on the Rust side. +""" + +from types import SimpleNamespace + +import torch + +from tilegym.backend import register_impl +from tilegym.backend.cutile_rs.autotuner import autotune_launch +from tilegym.backend.cutile_rs.utils import bind_kernel_function_cffi +from tilegym.backend.cutile_rs.utils import check_rc +from tilegym.backend.cutile_rs.utils import get_num_sm +from tilegym.backend.cutile_rs.utils import make_tensor_desc + +_KERNEL = "bmm" +_FFI_NAME = "cutile_bmm" +# C-declaration source of truth for the cffi boundary — keep in sync with the +# `cutile_bmm` signature in bmm_kernel/ffi.rs. Tensors cross as `const TensorDesc*` +# (the shared TensorDesc typedef is prepended by bind_kernel_function_cffi). +_FFI_CDEF = """ +int32_t cutile_bmm( + const TensorDesc* c, const TensorDesc* a, const TensorDesc* b, + int32_t bm, int32_t bn, int32_t bk, + int32_t group_size_m, int32_t trans_a, int32_t trans_b, + int32_t persistent, int32_t num_cta_in_cga, int32_t occupancy, + int32_t num_programs, int32_t device_id, uint64_t raw_stream); +""" + +# Supported dtypes (wrapper input validation; the dtype code is packed by +# make_tensor_desc into TensorDesc.dtype). +_DTYPES = (torch.float32, torch.float16, torch.bfloat16) + +_AUTO_COMPILE_OPTION = -1 + + +def _cdiv(a: int, b: int) -> int: + return (a + b - 1) // b + + +def _persistent_configs(): + # analysis.json static_persistent autotune_configs (num_ctas=2 carried to FFI). + return [ + SimpleNamespace(BM=256, BN=256, BK=64, GROUP_SIZE_M=8, NUM_CTA_IN_CGA=2, OCCUPANCY=1), + SimpleNamespace(BM=128, BN=256, BK=64, GROUP_SIZE_M=8, NUM_CTA_IN_CGA=2, OCCUPANCY=1), + ] + + +def _non_persistent_configs(): + # analysis.json non_persistent: fixed TILE 128/128/32, compile options auto. + return [ + SimpleNamespace(BM=128, BN=128, BK=32, GROUP_SIZE_M=8, NUM_CTA_IN_CGA=None, OCCUPANCY=None), + ] + + +def _compile_option_value(value) -> int: + return _AUTO_COMPILE_OPTION if value is None else int(value) + + +def _persistent_grid(total_tiles: int, num_ctas: int, occupancy: int) -> int: + num_sms = get_num_sm() + return max(min(num_sms // max(num_ctas, 1), total_tiles), 1) * max(occupancy, 1) + + +def _run_ffi(c, a, b, cfg, persistent, trans_a, trans_b): + ffi, lib = bind_kernel_function_cffi(_KERNEL, _FFI_CDEF) + _dev = a.device + device_id = _dev.index if _dev.index is not None else torch.cuda.current_device() + raw_stream = torch.cuda.current_stream(device=_dev).cuda_stream + + bm, bn, bk = int(cfg.BM), int(cfg.BN), int(cfg.BK) + group_size_m = int(getattr(cfg, "GROUP_SIZE_M", 8)) + num_cta_in_cga = getattr(cfg, "NUM_CTA_IN_CGA", None) + occupancy = getattr(cfg, "OCCUPANCY", None) + + # logical (post-transpose) dims for the persistent grid math + q = int(a.shape[0]) + m = int(a.shape[2]) if trans_a else int(a.shape[1]) + n = int(b.shape[1]) if trans_b else int(b.shape[2]) + if persistent: + total_tiles = _cdiv(m, bm) * _cdiv(n, bn) * q + num_programs = _persistent_grid(total_tiles, num_cta_in_cga or 1, occupancy or 1) + else: + # 3-D grid is computed inside the FFI; num_programs unused for this variant. + num_programs = 1 + + # Keep the descriptors alive until after the FFI call (cffi frees on GC). + cd = make_tensor_desc(ffi, c) + ad = make_tensor_desc(ffi, a) + bd = make_tensor_desc(ffi, b) + rc = lib.cutile_bmm( + cd, + ad, + bd, + bm, + bn, + bk, + group_size_m, + int(1 if trans_a else 0), + int(1 if trans_b else 0), + int(1 if persistent else 0), + _compile_option_value(num_cta_in_cga), + _compile_option_value(occupancy), + int(num_programs), + int(device_id), + int(raw_stream), + ) + check_rc(rc, _FFI_NAME) + return c + + +@register_impl("bmm", backend="cutile-rs") +def bmm(a, b, transpose_a=False, transpose_b=False, static_persistent=True, use_tma=None, **kwargs): + """batched matmul via cutile-rs FFI (forward-only). C = A @ B per batch.""" + if a.dtype not in _DTYPES: + raise NotImplementedError(f"cutile-rs bmm: dtype {a.dtype} not supported") + if b.dtype != a.dtype: + raise NotImplementedError("cutile-rs bmm: a and b must share dtype") + if a.ndim != 3 or b.ndim != 3: + raise NotImplementedError("cutile-rs bmm: inputs must be rank-3 [Q, *, *]") + if not a.is_cuda or not b.is_cuda or a.device != b.device: + raise ValueError("cutile-rs bmm: a and b must be CUDA tensors on the same device") + + persistent = bool(static_persistent) + if not persistent and (transpose_a or transpose_b): + raise NotImplementedError("cutile-rs bmm non-persistent variant does not support transpose") + + if a.requires_grad: + a = a.detach() + if b.requires_grad: + b = b.detach() + + a = a.contiguous() + b = b.contiguous() + + q = int(a.shape[0]) + m = int(a.shape[2]) if transpose_a else int(a.shape[1]) + n = int(b.shape[1]) if transpose_b else int(b.shape[2]) + k = int(a.shape[1]) if transpose_a else int(a.shape[2]) + out_dtype = a.dtype + + if int(b.shape[0]) != q: + raise ValueError(f"cutile-rs bmm: batch dim mismatch: a.shape={tuple(a.shape)} b.shape={tuple(b.shape)}") + k_b = int(b.shape[2]) if transpose_b else int(b.shape[1]) + if k_b != k: + raise ValueError(f"cutile-rs bmm: K dim mismatch: a.shape={tuple(a.shape)} b.shape={tuple(b.shape)}") + + configs = _persistent_configs() if persistent else _non_persistent_configs() + + def launch_with_cfg(cfg): + out = torch.empty((q, m, n), device=a.device, dtype=out_dtype) + _run_ffi(out, a, b, cfg, persistent, transpose_a, transpose_b) + return out + + def kernel_fn(cfg): + return launch_with_cfg(cfg) + + result = autotune_launch( + kernel_fn=kernel_fn, + configs=configs, + key=(q, m, n, k, a.dtype, persistent, transpose_a, transpose_b), + kernel_name=_KERNEL, + ) + return result.output diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/ffi.rs b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/ffi.rs new file mode 100644 index 0000000..cad32f8 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/ffi.rs @@ -0,0 +1,126 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +// SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +// +// FFI export for the softmax kernel — one C-ABI symbol `cutile_softmax`. +// +// y[m, :] = softmax(x[m, :]) over the last dim. This is a raw-pointer kernel: +// the entry takes `DevicePointer` + dims/strides (NOT `&Tensor`), so the FFI +// reads ptr/dims/strides from the `TensorDesc`s (crate::ffi_util) and wraps the +// pointers as `DevicePointer` — no borrow_tensor / ownership gate needed (the +// kernel never holds a Tensor wrapper, so nothing can free PyTorch memory). +// +// Anatomy of a cutile-rs FFI function: +// 1. Null-check the descriptors; read dtype via `dtype_str(desc.dtype)`. +// 2. Read logical dims / strides off the descriptors (`.dim(i)`, `.strides[i]`). +// 3. Build a Device from `device_id` and borrow the caller's stream. +// 4. Build `.generics()` in the SAME ORDER as the entry's ``. +// 5. Call the entry (op-builder), chain `.generics()` `.grid()` `.compile_options()`, +// then `.sync_on(&stream)`. +// dtype dispatch: a `match` on the dtype string picks which `` the macro shim +// instantiates. + +use core::ffi::c_void; +use cuda_async::device_buffer::DevicePointer; +use cuda_core::{Device, Stream}; +use cutile::half::{bf16, f16}; +use cutile::prelude::*; +use cutile::tile_kernel::{CompileOptions, TileKernel}; + +use crate::ffi_util::{TensorDesc, dtype_str}; +use softmax_module::softmax_kernel; + +#[unsafe(no_mangle)] +pub unsafe extern "C" fn cutile_softmax( + // tensors (dtype + shapes + strides carried in the descriptors): + // y: [m, n] (output) x: [m, n] (input) + y: *const TensorDesc, + x: *const TensorDesc, + // tile sizes — autotuner-tunable + bm: i32, + bn: i32, + // latency — autotuner-tunable (1:1 map from Triton-TileIR num_stages) + latency: i32, + // compile options: <=0 means auto/default + num_cta_in_cga: i32, + occupancy: i32, + // launch grid (number of programs / CTAs) for the persistent grid-stride loop + grid_size: i32, + // CUDA device ordinal of the tensors/stream (multi-GPU correctness) + device_id: i32, + // CUDA stream (cuStream_t cast to u64) + raw_stream: u64, +) -> i32 { + if y.is_null() || x.is_null() { + return -5; + } + let (y_d, x_d) = unsafe { (&*y, &*x) }; + + let dty: &'static str = match dtype_str(x_d.dtype) { + Some(s) => s, + None => return -2, + }; + // Logical dims + row/col strides from the input descriptor (y matches). + let m = x_d.dim(0); + let n = x_d.dim(1); + let (s_m, s_n) = match (i32::try_from(x_d.strides[0]), i32::try_from(x_d.strides[1])) { + (Ok(a), Ok(b)) => (a, b), + _ => return -6, + }; + + let device = match Device::new(device_id.max(0) as usize) { + Ok(d) => d, + Err(e) => { + eprintln!("cutile_softmax: Device::new failed: {e:?}"); + return -4; + } + }; + let stream = unsafe { Stream::borrow_raw(raw_stream as *mut c_void, &device) }; + + macro_rules! dispatch { + ($E:ty) => {{ + let y_dp: DevicePointer<$E> = unsafe { DevicePointer::from_cu_deviceptr(y_d.ptr) }; + let x_dp: DevicePointer<$E> = unsafe { DevicePointer::from_cu_deviceptr(x_d.ptr) }; + + // Generics in the SAME ORDER as kernel.rs: + // + let generics = vec![ + dty.to_string(), + bm.to_string(), + bn.to_string(), + latency.to_string(), + ]; + + // Honor the wrapper's sentinel: only call a setter when the value is + // > 0; a value <= 0 leaves that field at the compiler default. + let mut opts = CompileOptions::default(); + if occupancy > 0 { + opts = opts.occupancy(occupancy); + } + if num_cta_in_cga > 0 { + opts = opts.num_cta_in_cga(num_cta_in_cga); + } + + let op = unsafe { softmax_kernel(y_dp, x_dp, m, n, s_m, s_n) } + .generics(generics) + .grid((grid_size as u32, 1, 1)) + .compile_options(opts); + + match op.sync_on(&stream) { + Ok(_) => 0, + Err(e) => { + eprintln!("cutile_softmax: launch failed: {e:?}"); + -3 + } + } + }}; + } + + match dty { + "f32" => dispatch!(f32), + "f16" => dispatch!(f16), + "bf16" => dispatch!(bf16), + _ => -2, + } +} diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/kernel.rs b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/kernel.rs new file mode 100644 index 0000000..feaf018 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/kernel.rs @@ -0,0 +1,189 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +// SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +// +// Softmax — canonical tilegym-converting-cutile-triton-to-cutile-rs example kernel. +// +// This file is the FROM-SCRATCH template you use for the first kernel you +// convert. The cutile-rs Rust lives inside the tilegym tree +// (src/tilegym/ops/cutile_rs/) — there is no separate cutile-rs checkout. It is +// intentionally small so every load-bearing line is visible. +// +// What this kernel demonstrates (the full conversion stack in 1 file): +// +// 1. dtype generic : `` (Rule 16) +// 2. tile-size const generics: `` (Rule 23) +// 3. latency const generic : `` (perf) +// 4. partition view I/O : `tview.partition(shape)` (Rule 28 v2) +// 5. assume hints : pointer div_by + scalar bounds (Rule 21 v2) +// 6. tile-form math : reduce_max / exp2 / reduce_sum (op-mapping.md) +// 7. boundary safety : padding_value=zero auto-emitted +// 8. TMA on by default : `tma::Enabled` everywhere (perf) +// +// What you MUST swap when reusing this template: +// +// * `softmax_module` → your own `{kernel_name}_module` +// * `softmax_kernel` → your own entry name +// * the entry signature → the args your op needs +// * the body → your op's math +// * everything else (assume hints, latency wiring, partition pattern, +// ``, optimization_hints) is boilerplate — keep as-is. +// +// Cross-references: +// * `references/coding-rules.md` — all rules cited above +// * `references/op-mapping.md` — cuTile-py / Triton-TileIR op → cutile-rs op +// * `concepts/strided-view-to-partition-view.md` — when the reference IR +// comes from Triton-TileIR (uses element-level offsets), how to translate to +// cutile-rs's tile-level partition view indexing. +// * `concepts/tensor-vs-pointer-pattern.md` — when to pick partition-view +// (this template) vs raw-pointer scatter/gather. +// +// Naming convention: `kernel.rs` lives at +// ${CUTILE_KERNEL_OUT_ROOT}/softmax/kernel.rs +// (per-kernel working dir) and is mirrored into the aggregated crate's sibling +// src/tilegym/ops/cutile_rs/softmax_kernel/kernel.rs +// where it is `include!()`-d by both: +// src/tilegym/ops/cutile_rs/cutile_kernels/src/lib.rs (FFI export — one +// `mod softmax { include!("../../softmax_kernel/kernel.rs"); ... }`) +// ${CUTILE_KERNEL_OUT_ROOT}/softmax/softmax_pipeline.rs (compile test) + +#[cutile::module] +pub mod softmax_module { + use cutile::core::*; + + /// Online softmax along the last (N) dimension. + /// + /// Layout: `y[m, n] = exp(x[m, n] - max(x[m, :])) / sum(exp(x[m, :] - max))`. + /// Tiled M×N at `(BM, BN)` granularity. N must satisfy `N <= BN` for the + /// vanilla version; for `N > BN` you would add an inner reduce loop (see + /// `cutile-examples/examples/softmax.rs` for the multi-tile variant). + /// + /// Const generics: + /// * `E` : data element type (f16 / bf16 / f32) — Rule 16 + /// * `BM`, `BN` : tile sizes — Rule 23 + /// * `LATENCY` : pipeline depth hint passed to every load/store — perf + /// + /// Runtime args (all `unsafe` — caller must ensure validity): + /// * `y_ptr`, `x_ptr` : raw pointers into device memory. + /// * `m`, `n` : logical rows and columns. + /// * `s_m`, `s_n` : strides (innermost typically 1). + /// + /// Grid: 1D over `M / BM` rows. Host-side launcher computes + /// `grid = ceildiv(m, BM)` and passes via `CompileOptions`. + #[cutile::entry( + unchecked_accesses = true, + // sm_100: occupancy + num_cta_in_cga. Adjust per Agent A's analysis.json. + // For sm_80 / sm_90 / sm_120 add additional `sm_XX = (...)` blocks here. + optimization_hints = ( + sm_100 = (occupancy = 2, num_cta_in_cga = 1,), + ), + )] + unsafe fn softmax_kernel( + y_ptr: *mut E, + x_ptr: *mut E, + m: i32, + n: i32, + s_m: i32, + s_n: i32, + ) { + // ─── Rule 21 v2: pointer div_by + scalar bounds_lower ────────────── + // + // POINTER assumes (always 16-byte aligned for fp16/bf16/f32 in PyTorch + // tensor allocations) → enables vectorized loads/stores in tileiras. + let y_ptr = unsafe { assume_div_by::<_, 16>(y_ptr) }; + let x_ptr = unsafe { assume_div_by::<_, 16>(x_ptr) }; + // SCALAR assumes — bounds_lower<0> only. NEVER `assume_div_by` on + // scalars: when N doesn't divide the runtime value the compiler + // silently produces wrong output (see memory `feedback_cutile_rs_ + // scalar_assume_div_by` and Rule 21 v2 history of K=511 corruption). + let m = unsafe { assume_bounds_lower::<_, 0>(m) }; + let n = unsafe { assume_bounds_lower::<_, 0>(n) }; + let s_m = unsafe { assume_bounds_lower::<_, 0>(s_m) }; + let s_n = unsafe { assume_bounds_lower::<_, 0>(s_n) }; + + // ─── 1D grid: each block handles one BM-row strip ────────────────── + // `get_tile_block_id()` is the canonical 3D-tuple block-id intrinsic + // (cutile/src/_core.rs:1029). Use the .0 field for 1-D grids. + // `block_id::()` does NOT exist — older drafts of this file referenced + // it; that was a typo. Always use `get_tile_block_id()`. + let bid_m = get_tile_block_id().0; + + // ─── Tensor views over the flat pointers ─────────────────────────── + let x_tview: TensorView = unsafe { make_tensor_view(x_ptr, [m, n], [s_m, s_n]) }; + let mut y_tview: TensorView = unsafe { make_tensor_view(y_ptr, [m, n], [s_m, s_n]) }; + + // ─── Boundary-safe partition views (Rule 28 v2) ──────────────────── + // + // `tview.partition(const_shape![BM, BN])` is the cutile-rs shorthand + // for `make_partition_view(self, shape, padding::Zero, dim_map:: + // Identity, token)`. The `padding_value = zero` it emits is what makes + // the kernel boundary-safe on M%BM != 0 / N%BN != 0 shapes WITHOUT + // needing to disable TMA — keep `tma::Enabled` (memory + // `feedback_keep_tma_enabled_for_perf`). + let x_part: Partition = x_tview.partition(const_shape![BM, BN]); + let mut y_part: PartitionMut = + unsafe { y_tview.partition_mut(const_shape![BM, BN]) }; + + // ─── Load: TMA on, latency hint per op ──────────────────────────── + // + // Reference IR (Triton-TileIR or cuTile-py) typically has `latency = N` on + // every load/store; matching it enables the same software pipelining + // depth in tileiras. Pass `Some(LATENCY)` (NOT `None`) — `None` + // inherits entry default which is often suboptimal. + let x_tile: Tile = load_view_tko( + &x_part, + [bid_m, 0i32], + ordering::Weak, + scope::TileBlock, + Some(LATENCY), + tma::Enabled, + ); + + // ─── Online softmax body — tile-form math ────────────────────────── + // + // 1. Cast input to f32 for stable reduction (fp16/bf16 → f32) — Rule 8. + // `convert_tile` is a no-op when E == f32. + let x_f32: Tile = convert_tile(x_tile); + + // 2. row-wise max : Tile → Tile + let row_max: Tile = reduce_max(x_f32, 1i32); + + // 3. broadcast row_max back to (BM, BN) and subtract. + let row_max_2d: Tile = row_max + .reshape(const_shape![BM, 1]) + .broadcast(x_f32.shape()); + let shifted: Tile = x_f32 - row_max_2d; + + // 4. exp via exp2(x * INV_LOG_2). flush_to_zero matches Triton-TileIR/cuTile-py + // reference IR (Rule 26 — Agent C will fail you if it's missing). + let inv_log2: Tile = + broadcast_scalar(constant(std::f32::consts::LOG2_E), shifted.shape()); + let exped: Tile = exp2(mulf(shifted, inv_log2, ftz::Enabled)); + + // 5. row-wise sum + divide. + let row_sum: Tile = reduce_sum(exped, 1i32); + let row_sum_2d: Tile = row_sum + .reshape(const_shape![BM, 1]) + .broadcast(exped.shape()); + // `true_div` (rounding + flush_to_zero) matches the Triton-TileIR + // reference. `divf` with `rounding::Approx` hits a known cutile-rs + // encoding bug (Approx→5 vs printer→4) — use `true_div`. + let normalized: Tile = true_div(exped, row_sum_2d); + + // 6. Cast back to E and store. + let result: Tile = convert_tile(normalized); + + unsafe { + store_view_tko_mut( + &mut y_part, + result, + [bid_m, 0i32], + ordering::Weak, + scope::TileBlock, + Some(LATENCY), + tma::Enabled, + ); + } + } +} diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/softmax_pipeline.rs b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/softmax_pipeline.rs new file mode 100644 index 0000000..d02c9a6 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/softmax_pipeline.rs @@ -0,0 +1,135 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// +// SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +// +// Compile-and-bytecode-dump test for the softmax kernel. +// +// Lives at: ${CUTILE_KERNEL_OUT_ROOT}/softmax/softmax_pipeline.rs +// (the per-kernel working dir; run with `cargo test --test softmax_pipeline`). +// The cutile-rs Rust lives inside tilegym — there is no separate cutile-rs +// checkout / CUTILE_RS_ROOT. +// +// Two-stage IR pipeline: +// +// Stage 1 (this file): Rust test compiles the kernel and writes a portable +// **bytecode** artifact (`.tileirbc`). No external +// binary calls, no MLIR text serialization in Rust — +// keeps the test minimal and reproducible. +// +// Stage 2 (shell): `cuda-tile-translate --cudatilebc-to-mlir` lifts +// the bytecode to MLIR text; `cuda-tile-opt +// --canonicalize --cse` produces the canonicalized +// IR Agent C diffs against the reference. +// +// The kernel.rs source is the single source of truth — it is the sibling +// file in this same per-kernel working dir, and both lib.rs (FFI) and this +// test pull from it via `include!`. Do NOT duplicate kernel code here. + +use cutile::compile_api::KernelCompiler; +use cutile_compiler::specialization::{DivHint, SpecializationBits}; + +// ─── single source of truth — pull kernel.rs (the sibling in this same +// ─── per-kernel working dir) into this test ───────────── +include!("kernel.rs"); + +// `cutile-macros` injects this `__module_ast_self` symbol inside every +// `#[cutile::module]` block. Importing it gives KernelCompiler a handle to +// the module's AST without re-declaring anything. +use softmax_module::__module_ast_self; + +// Pick ONE concrete generic specialization for the bytecode dump. Any +// reasonable tile size works for first-run development; the autotuner / FFI +// will JIT-compile other (BM, BN, LATENCY) specializations at runtime. +const BM: i32 = 128; +const BN: i32 = 128; +const LATENCY: i32 = 4; + +// Static stride hints: outer strides per pointer arg in element units. +// For row-major (M, N), outer stride is N (innermost stride 1 implicit). +const REP_N: i32 = 1024; +const Y_STRIDES: &[i32] = &[REP_N]; +const X_STRIDES: &[i32] = &[REP_N]; + +const TARGET: &str = "sm_100"; + +// Alignment hints so the DUMP carries the same `assume div_by` lines the +// reference IR does. Read `reference/reference.mlir` (`grep "assume div_by"`) +// and reproduce the per-arg divisor EXACTLY — do not blanket-fill 16; if the +// reference asserts div_by<8> on a dim, use a value whose divisor is 8 there. +// shape_div/stride_div are in ELEMENTS, base_ptr_div in BYTES; DivHint::from_value +// computes the largest power-of-2 divisor (clamped to 16). A pipeline test with +// ONLY .strides() dumps 0 assumes vs the reference's many → noisy Agent C diff. +fn spec_bits_1d(rep_n: i32) -> SpecializationBits { + // softmax view is row-major [M, N]; here N=rep_n is the inner contiguous dim. + SpecializationBits { + shape_div: vec![DivHint::from_value(rep_n)], // N divisible by 16 (1024 → 16) + stride_div: vec![DivHint::from_value(1)], // innermost stride = 1 + stride_one: vec![true], + base_ptr_div: DivHint::from_ptr(0x1000), // 16-byte aligned base ptr + elements_disjoint: true, + } +} + +fn compile() -> Vec { + let artifacts = KernelCompiler::new(__module_ast_self, "softmax_module", "softmax_kernel") + // Generics in the SAME ORDER as the entry signature: + // + .generics(vec![ + "f32".to_string(), // E (data dtype) + BM.to_string(), + BN.to_string(), + LATENCY.to_string(), + ]) + .strides(&[("y_ptr", Y_STRIDES), ("x_ptr", X_STRIDES)]) + // Reproduce the reference's `assume div_by` (see spec_bits_1d above). + // Names MUST match the entry's pointer args. Omit an arg here only if + // the reference shows NO assume for it. + .spec_args(&[ + ("y_ptr", spec_bits_1d(REP_N)), + ("x_ptr", spec_bits_1d(REP_N)), + ]) + .target(TARGET) + .compile() + .expect( + "softmax_kernel compile failed — check kernel.rs against the \ + compile-fix-loop table in agents/agent_b.md", + ); + + artifacts.bytecode().expect("bytecode serialization failed") +} + +fn write_bytecode(bc: &[u8]) { + // NOTE (path): this template's include! and the path math below still + // reference the old `cutile-kernels/` include layout — the migrated layout + // wires kernel.rs into the aggregated cutile_kernels crate via + // `mod softmax { include!("../../softmax_kernel/kernel.rs"); ... }` and dumps + // generated IR under ${CUTILE_KERNEL_OUT_ROOT}/softmax/generated/. Only the + // path plumbing changed; the compile/spec-bits logic is unchanged. (Rust + // path code left as-is per the migration scope — comments only.) + let path = format!( + "{}/cutile-kernels/softmax/generated/generated.tileirbc", + env!("CARGO_MANIFEST_DIR") + .strip_suffix("/cutile") + .unwrap_or(env!("CARGO_MANIFEST_DIR")), + ); + std::fs::create_dir_all(std::path::Path::new(&path).parent().unwrap()).ok(); + std::fs::write(&path, bc).expect("failed to write bytecode file"); + println!("✅ softmax kernel compiled. Bytecode written to {path}"); + println!(" Next step (shell):"); + println!(" cuda-tile-translate --cudatilebc-to-mlir {path} \\"); + println!(" -o ${CUTILE_KERNEL_OUT_ROOT}/softmax/generated/generated.mlir"); + println!(" $CUDA_TILE_OPT_BIN --canonicalize --cse \\"); + println!(" ${CUTILE_KERNEL_OUT_ROOT}/softmax/generated/generated.mlir \\"); + println!(" -o ${CUTILE_KERNEL_OUT_ROOT}/softmax/generated/generated_canon.mlir"); +} + +#[test] +fn softmax_compile_default() { + let bc = compile(); + assert!( + !bc.is_empty(), + "bytecode is empty — compile produced nothing" + ); + write_bytecode(&bc); +} diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/walkthrough.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/walkthrough.md new file mode 100644 index 0000000..d76e562 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/walkthrough.md @@ -0,0 +1,268 @@ +# Worked example — softmax (your first kernel) + +This directory is the **complete, copy-pasteable starting point** for adding a +brand-new kernel to a fresh tilegym checkout (the cutile-rs Rust now lives inside +tilegym under `src/tilegym/ops/cutile_rs/` — there is no separate cutile-rs +checkout). It exists so that a user with zero prior conversions can produce a +working kernel without having to reverse-engineer the conventions out of a larger +kernel. + +The example is a row-wise softmax along the last dimension — small enough that +every load-bearing line is visible, large enough to exercise: + +| Stack layer | Demonstrated by | +|------------------------|----------------------------------------------------| +| dtype generic | `` in `kernel.rs` (Rule 16) | +| const generics | `` | +| Boundary-safe partition | `tview.partition(const_shape![BM, BN])` (Rule 28 v2) | +| TMA + latency hints | `tma::Enabled` + `Some(LATENCY)` on every op | +| Tile-form reductions | `reduce_max` / `reduce_sum` | +| `flush_to_zero` etc. | `mulf(..., ftz::Enabled)`, `true_div` | +| FFI dtype dispatch | `match dtype_str(x_d.dtype) { Some("f32") => ... }` | +| TensorDesc boundary | `*const TensorDesc` in `ffi.rs`; `make_tensor_desc` in wrapper | +| cffi wrapper | `wrapper.py::_FFI_CDEF` mirrors `ffi.rs` | +| CUPTI autotune | `autotune_launch(kernel_fn=lambda cfg: ...)` | +| pipeline rule 11 NO_FALLBACK | `raise NotImplementedError(...)` for bad combos | +| tilegym test surface | `Test_Softmax::test_op` + `test_perf` | + +## File map + +| File | Install to | Role | +|------|-----------|------| +| `kernel.rs` | `${CUTILE_KERNEL_OUT_ROOT}/softmax/kernel.rs` | The cutile-rs kernel (single source of truth) | +| `ffi.rs` | `${CUTILE_KERNEL_OUT_ROOT}/softmax/ffi.rs` | C-ABI export (`*const TensorDesc` dtype dispatch + launch) | +| `softmax_pipeline.rs` | `${CUTILE_KERNEL_OUT_ROOT}/softmax/softmax_pipeline.rs` | `cargo test` compile + IR dump | +| `wrapper.py` | `{TILEGYM_PATH}/src/tilegym/ops/cutile_rs/softmax.py` | tilegym cffi wrapper (autotune + dispatch) | + +The kernel Rust is wired into the aggregated `cutile_kernels` crate under +`{TILEGYM_PATH}/src/tilegym/ops/cutile_rs/cutile_kernels/` (one +`libcutile_kernels.so` for every op). There is no separate cutile-rs checkout. + +Plus three small **registration snippets** (paste-into existing files): +see *Wiring* below. + +## Bootstrap (zero → first kernel) + +### 0. Prerequisites + +* A cloned `tilegym` checkout — it now holds the cutile-rs Rust + under `src/tilegym/ops/cutile_rs/`. No separate cutile-rs checkout is needed; + the `cutile_kernels` crate pulls its deps from crates.io (pinned `=0.2.0`). +* CUDA toolkit (`CUDA_TOOLKIT_PATH`, default `/usr/local/cuda`) + Rust toolchain + (`cargo`). `tileiras` lowers IR → cubin at launch. +* Python with PyTorch installed. + +### 1. The load-bearing Python bits ship in-repo + +The cutile-rs backend runtime is already present in the repository — no copy step: + +``` +${TILEGYM_PATH}/src/tilegym/backend/cutile_rs/utils.py # bind_kernel_function_cffi, make_tensor_desc, check_rc, get_num_sm +${TILEGYM_PATH}/src/tilegym/backend/cutile_rs/autotuner.py # autotune_launch +${TILEGYM_PATH}/src/tilegym/backend/cutile_rs/__init__.py +${TILEGYM_PATH}/src/tilegym/ops/cutile_rs/__init__.py +``` + +Wrappers `import from tilegym.backend.cutile_rs.*`. You only add the per-op +wrapper under `src/tilegym/ops/cutile_rs/` and register it. + +### 2. Install the example kernel files + +```bash +# The per-kernel working dir holds Agent A–D artifacts (kernel.rs, ffi.rs, +# reference/, generated/, reports/). The eval harness depends on this layout. +export CUTILE_KERNEL_OUT_ROOT=/abs/path/to/cutile_kernel_out + +# Rust kernel + FFI + pipeline test (per-kernel working dir) +mkdir -p ${CUTILE_KERNEL_OUT_ROOT}/softmax/{generated,reports/agent_logs} +cp ${SKILL}/examples/softmax/kernel.rs ${CUTILE_KERNEL_OUT_ROOT}/softmax/kernel.rs +cp ${SKILL}/examples/softmax/ffi.rs ${CUTILE_KERNEL_OUT_ROOT}/softmax/ffi.rs +cp ${SKILL}/examples/softmax/softmax_pipeline.rs ${CUTILE_KERNEL_OUT_ROOT}/softmax/softmax_pipeline.rs + +# Mirror the kernel Rust into the aggregated crate's sibling {op}_kernel/ dir so +# the `mod softmax { include!("../../softmax_kernel/kernel.rs"); ... }` line in +# cutile_kernels/src/lib.rs resolves (see step 3a). +mkdir -p ${TILEGYM_PATH}/src/tilegym/ops/cutile_rs/softmax_kernel +cp ${SKILL}/examples/softmax/kernel.rs ${TILEGYM_PATH}/src/tilegym/ops/cutile_rs/softmax_kernel/kernel.rs +cp ${SKILL}/examples/softmax/ffi.rs ${TILEGYM_PATH}/src/tilegym/ops/cutile_rs/softmax_kernel/ffi.rs + +# tilegym Python wrapper +cp ${SKILL}/examples/softmax/wrapper.py ${TILEGYM_PATH}/src/tilegym/ops/cutile_rs/softmax.py +``` + +### 3. Wiring (3 small edits in existing files) + +#### a. `cutile_kernels/src/lib.rs` — register the op in the aggregated crate + +There is ONE aggregated cdylib crate for every op: +`{TILEGYM_PATH}/src/tilegym/ops/cutile_rs/cutile_kernels/`. Register softmax by +adding ONE `mod` that `include!`s the op's device + FFI sources from the sibling +`softmax_kernel/` dir (the shared `TensorDesc` helpers are pulled in once via +`ffi_util`): + +```rust +// cutile_kernels/src/lib.rs +#[path = "../../ffi_util.rs"] +mod ffi_util; // shared TensorDesc / borrow_tensor / dtype_str (once) + +mod softmax { + include!("../../softmax_kernel/kernel.rs"); + include!("../../softmax_kernel/ffi.rs"); +} +``` + +Deps are crates.io PINNED (`cutile = "=0.2.0"`, …; Rule 33) — no path deps and no +cutile-rs checkout. Then build the single shared library: + +```bash +# The loader autobuilds this on first use (CUTILE_RS_AUTOBUILD, on by default); +# building by hand is optional. +cd ${TILEGYM_PATH}/src/tilegym/ops/cutile_rs/cutile_kernels +cargo build --release +nm -D target/release/libcutile_kernels.so | grep cutile_softmax +# must print: T cutile_softmax +``` + +The wrapper.py template loads this single `libcutile_kernels.so` via +``bind_kernel_function_cffi("softmax", _FFI_CDEF)`` (cffi). One aggregated library +serves every op; each op only cdefs its own symbol. There is no per-op `.so` and +no legacy monolithic ``libcutile_ffi.so``. + +#### b. `{TILEGYM_PATH}/src/tilegym/ops/cutile_rs/__init__.py` — register import + +Append: +```python +from . import softmax # noqa: F401 (registers @register_impl) +``` + +#### b'. `{TILEGYM_PATH}/src/tilegym/backend/selector.py` — register backend + +The tilegym test harness skips every ``cutile-rs`` test with +``pytest.skip("Backend cutile-rs is not available")`` unless +``selector.is_cutile_rs_available()`` exists AND ``"cutile-rs"`` is a key in +``_check_backends_availability()``'s availability dict. Add both: + +```python +def is_cutile_rs_available() -> bool: + """True if the aggregated cutile_kernels crate is on disk and cargo is callable. + + We deliberately use a *permissive* probe — the single libcutile_kernels.so is + autobuilt on-demand by the loader (CUTILE_RS_AUTOBUILD), so requiring the .so + here would skip legitimately-installed environments. The Rust now lives inside + tilegym, so probe the in-tree cutile_kernels crate (CUTILE_RS_KERNELS_DIR + overrides its location); there is no CUTILE_RS_DIR / cutile-rs checkout. + """ + import os, shutil + kernels_dir = os.environ.get("CUTILE_RS_KERNELS_DIR") or os.path.join( + os.path.dirname(__file__), "..", "ops", "cutile_rs", "cutile_kernels" + ) + if not os.path.isfile(os.path.join(kernels_dir, "Cargo.toml")): + return False + return shutil.which("cargo") is not None + + +def _check_backends_availability() -> Dict[str, bool]: + availability = { + # ...existing entries (cutile, triton, tilecpp, ...) + "cutile-rs": is_cutile_rs_available(), # NEW + } + return availability +``` + +Verify with:: + + python -c "from tilegym.backend.selector import is_cutile_rs_available as f; print(f())" + # → True (the in-tree cutile_kernels crate exists and cargo is callable) + +If you skip this step, your wrapper's ``@register_impl`` registration is +still correct, but ``tilegym.set_backend("cutile-rs")`` returns +``"cutile-rs is not available"`` and pytest will silently SKIP every +parametrized case. + +#### c. `{TILEGYM_PATH}/src/tilegym/ops/ops.py` — add @dispatch + +If softmax is not already a registered tilegym op, add: +```python +@dispatch("softmax", fallback_backend="triton") +def softmax_interface(x: torch.Tensor, dim: int = -1): + """Row-wise softmax (last-dim by default).""" + pass +``` +(If it already exists with a `cutile` backend, this step is a no-op — your +`@register_impl("softmax", backend="cutile-rs")` plugs into the same +dispatch table.) + +### 4. Compile + IR pipeline + correctness sanity + +```bash +# Stage 1 — Rust pipeline test dumps bytecode (.tileirbc) +cd ${CUTILE_KERNEL_OUT_ROOT}/softmax +cargo test --release --test softmax_pipeline -- --nocapture +# expect: ✅ softmax kernel compiled. Bytecode written to ...generated.tileirbc + +# Stage 2 — bytecode → MLIR text +cuda-tile-translate --cudatilebc-to-mlir \ + ${CUTILE_KERNEL_OUT_ROOT}/softmax/generated/generated.tileirbc \ + -o ${CUTILE_KERNEL_OUT_ROOT}/softmax/generated/generated.mlir + +# Stage 3 — canonicalize + CSE for IR diff +${CUDA_TILE_OPT_BIN} --canonicalize --cse \ + ${CUTILE_KERNEL_OUT_ROOT}/softmax/generated/generated.mlir \ + -o ${CUTILE_KERNEL_OUT_ROOT}/softmax/generated/generated_canon.mlir + +# tilegym pytest — the loader autobuilds libcutile_kernels.so on first use +# (CUTILE_RS_AUTOBUILD is on by default; tileiras lowers IR → cubin at launch). +cd ${TILEGYM_PATH} +python -m pytest tests/ops/test_softmax.py -k cutile_rs -v +# expect: passing on f32/f16/bf16 across the 3 shapes +``` + +### 5. (Optional) Drive the full A→B→C→D→E pipeline on a different kernel + +Once the softmax example is green, you have proof that the install + wiring +work. Now you can invoke the skill on a real kernel that lives in tilegym: + +``` +"Use the tilegym-converting-cutile-triton-to-cutile-rs skill to add a cutile-rs +backend to tests/ops/test_{kernel_name}.py" +``` + +The orchestrator will spawn Agent A → ... → Agent E. Each agent's prompt +references `examples/softmax/` for the canonical pattern (see Agent B prompt +section "Pattern reference"). + +## What to swap when reusing this template for a new kernel + +``` +Replace globally "softmax" → "{kernel_name}" + softmax_module → {kernel_name}_module + softmax_kernel → {kernel_name}_kernel + cutile_softmax → cutile_{kernel_name} + +Per-file specifics: + kernel.rs Body (lines under "Online softmax body"). + Const-generics list (BM, BN, LATENCY → your set). + Entry signature (pointers + dims your op needs). + optimization_hints (occupancy / num_cta_in_cga + from Agent A's analysis.json). + + ffi.rs FFI extern signature (tensors stay `*const TensorDesc`; + mirror the scalar args your new kernel.rs needs). + borrow_tensor:: vs DevicePointer::from_cu_deviceptr + (pick per the entry's param type). + generics vec! (mirror the const-generic order). + + wrapper.py _FFI_CDEF (mirror ffi.rs; tensor args = `const TensorDesc*`). + make_tensor_desc packs each tensor (no argtypes list). + _configs() (autotune space — copy from Agent A + analysis.json.autotune_configs). + softmax(...) signature (match @dispatch sig). +``` + +After your edits, run: +```bash +bash ${SKILL}/scripts/validate_kernel.sh {kernel_name} +``` +to verify all expected files are present + correctly wired (dtype generic, +`*const TensorDesc` FFI symbol in the aggregated `libcutile_kernels.so`, +`register_impl`, cffi `_FFI_CDEF`, test backend list, etc.). diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/wrapper.py b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/wrapper.py new file mode 100644 index 0000000..ddf9138 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/examples/softmax/wrapper.py @@ -0,0 +1,447 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +""" +cutile-rs softmax backend via FFI to the shared libcutile_kernels.so. + +This is the canonical Python wrapper template for Agent D, which owns the host +boundary in the current conversion pipeline. Agent B writes device kernel.rs; +Agent D writes ffi.rs, this Python wrapper shape, backend wiring, and the +correctness report. + +It demonstrates: + 1. Shared cdylib loading via bind_kernel_function_cffi (one + libcutile_kernels.so for every op; the op cdefs only its own symbol). + 2. A cffi cdef (_FFI_CDEF) that exactly matches ffi.rs; tensors cross as + `const TensorDesc*` and are packed with make_tensor_desc (dtype/shape/ + strides travel inside the descriptor — no ctypes argtypes list to drift). + 3. dtype/parameter gates that raise for unsupported cases. + 4. fixed kernel_configs bypass for debugging. + 5. CUPTI autotune with torch.empty-only allocations inside kernel_fn. + 6. an autograd-aware public wrapper. + 7. carrying autotune-config compile options (occupancy / num_cta_in_cga) + through to the FFI so a tuned CGA cluster size actually takes effect. + 8. passing the tensor's device ordinal (device_id) so multi-GPU launches + target the right device. + +Important wrapper rules: + - cutile-rs is FORWARD-ONLY. If input.requires_grad is True, detach it + first and return a leaf tensor (no grad_fn). Do NOT wrap forward in + torch.autograd.Function; backward via a PyTorch formula adds large + slowdown on the combined fwd/bwd ratio and obscures the kernel-level + perf signal we want to evolve. Backward correctness tests must skip + the cutile-rs backend explicitly. + - detach() is essentially free, so apply it unconditionally when + requires_grad is True. + - Match launch grid semantics to the Rust kernel body. This softmax + template is persistent: it may launch fewer CTAs than logical rows + because kernel.rs loops over remaining rows with a grid-stride loop. + For a non-persistent one-block-per-row kernel, use grid=n_rows. Do not + cap such a grid by occupancy or num_sms. + - For a persistent grid-stride kernel, size the launch to a resident-CTA + budget, not to a bare one-CTA-per-SM cap. The safe pattern is: + grid = min(total_work_tiles, num_sms * resident_ctas_per_sm / cga) + where total_work_tiles is the loop's logical tile count, usually row + tiles for reductions/norms. A wrapper like `min(num_sms, total_tiles)` + is correctness-safe but can starve memory-bound persistent kernels. + - Keep grid-budget knobs separate from compile options. More launched CTAs + do not imply `CompileOptions::occupancy` should be pinned, and a compile + occupancy hint does not by itself increase a too-small launch grid. + - Do not run a fixed-config canary in the hot path before autotune_launch. + CUPTI sums every GPU kernel inside the measured function; a canary plus + cached best-config launch doubles the measured time. + - Compile options (occupancy / num_cta_in_cga) follow analysis.json / + Agent A's best_config, NOT a blanket default: + * If the reference autotune space carries an EXPLICIT num_ctas / + num_cta_in_cga (e.g. matmul ships num_ctas in {1,2,4}), you MUST + forward that per-config value to the FFI so the CGA cluster size + takes effect. Passing the auto sentinel here silently pins the + cluster to the compiler default (=1) and regresses persistent / + tensor-core kernels ~1.6x. This is the #1 perf trap for GEMM. + * Only when the reference genuinely uses num_ctas=None / "compiler + auto-pick" (e.g. some elementwise / decode kernels) do you leave the + field None and pass the auto sentinel. Decide per-kernel from + analysis.json; do not default everything to auto. +""" + +from types import SimpleNamespace + +import torch + +from tilegym.backend import register_impl +from tilegym.backend.cutile_rs.autotuner import autotune_launch +from tilegym.backend.cutile_rs.utils import bind_kernel_function_cffi +from tilegym.backend.cutile_rs.utils import check_rc +from tilegym.backend.cutile_rs.utils import get_num_sm +from tilegym.backend.cutile_rs.utils import make_tensor_desc + +_KERNEL = "softmax" +_FFI_NAME = "cutile_softmax" +# C-declaration source of truth for the cffi boundary — keep in sync with the +# `cutile_softmax` signature in ffi.rs. Tensors cross as `const TensorDesc*` +# (the shared TensorDesc typedef is prepended by bind_kernel_function_cffi); m/n +# and row/col strides are read off the descriptors inside the FFI. +_FFI_CDEF = """ +int32_t cutile_softmax( + const TensorDesc* y, const TensorDesc* x, + int32_t bm, int32_t bn, int32_t latency, + int32_t num_cta_in_cga, int32_t occupancy, + int32_t grid_size, int32_t device_id, uint64_t raw_stream); +""" + +# Supported dtypes (wrapper input validation; the dtype code is packed by +# make_tensor_desc into TensorDesc.dtype). +_DTYPES = (torch.float32, torch.float16, torch.bfloat16) + +# FFI convention: a value >0 calls the CompileOptions setter with that value; +# a value <=0 (the _AUTO_COMPILE_OPTION sentinel) means "leave it at the +# compiler default". ffi.rs MUST honor this guard: +# +# let mut opts = CompileOptions::default(); +# if occupancy > 0 { opts = opts.occupancy(occupancy); } +# if num_cta_in_cga > 0 { opts = opts.num_cta_in_cga(num_cta_in_cga); } +# +# IMPORTANT: the sentinel is for kernels whose reference genuinely uses +# num_ctas=None / auto. It is NOT a blanket default. If analysis.json / +# best_config records an explicit num_cta_in_cga (matmul does), set the +# _COMPILE_* values below (or sweep them in _configs) so a REAL cluster size +# reaches the kernel. Defaulting these to None on a kernel whose reference +# tunes num_ctas pins the cluster to 1 and is the dominant GEMM perf regression. +_AUTO_COMPILE_OPTION = -1 + +# Grid-budget knobs for this persistent softmax template. These control how +# many CTAs are launched for a grid-stride kernel; they are SEPARATE from the +# compile-option cluster knob below. +# +# For memory-bound persistent reductions/norms, 1 CTA/SM is often not enough +# even though it is correct. Launch a resident-CTA budget (num_sms * a +# resident-per-SM factor, clamped to row tiles) and forward a matching +# compile occupancy for the static_persistent variant. +# Do not cargo-cult those values globally; copy the resident-grid pattern below +# and sweep only when E/C localizes a CTA-count-sensitive gap. +_GRID_OCCUPANCY = 2 +_GRID_NUM_CTA_IN_CGA = 1 + +# Compile-time knobs. softmax's reference uses compiler auto-pick, so these stay +# None HERE. For your op, set them from analysis.json / best_config: if the +# reference autotunes num_ctas (e.g. matmul ships {1,2,4}), DO NOT leave these +# None; carry the per-config value into the FFI (see _configs below). +_COMPILE_OCCUPANCY = None +_COMPILE_NUM_CTA_IN_CGA = None + + +def _compile_option_value(value) -> int: + if value is None: + return _AUTO_COMPILE_OPTION + return int(value) + + +def _configs(m: int, n: int): + """Autotune search space. + + Start with analysis.json reference configs for the real op, then optionally + add expansions after correctness passes. + + Compile-option fields per config: + - softmax's reference auto-picks, so OCCUPANCY / NUM_CTA_IN_CGA stay None. + - For an op whose reference autotunes num_ctas (matmul ships {1,2,4}), + carry the real value, e.g.: + SimpleNamespace(BM=256, BN=256, BK=64, NUM_CTA_IN_CGA=2, OCCUPANCY=1) + SimpleNamespace(BM=256, BN=256, BK=64, NUM_CTA_IN_CGA=4, OCCUPANCY=1) + so the autotuner explores live CGA cluster sizes. Do NOT fill these with + None on such an op; that pins the cluster to the compiler default (=1). + """ + del m + out = [] + for bm in (64, 128): + for bn in (128, 256): + # kernel contract: N <= BN (the BN-wide tile must cover the row; + # padding_value=0 handles N < BN). Keep only configs where BN >= N. + if bn >= n: + for latency in (2, 4): + out.append( + SimpleNamespace( + BM=bm, + BN=bn, + LATENCY=latency, + OCCUPANCY=_COMPILE_OCCUPANCY, + NUM_CTA_IN_CGA=_COMPILE_NUM_CTA_IN_CGA, + ) + ) + if not out: + out.append( + SimpleNamespace( + BM=64, + BN=128, + LATENCY=2, + OCCUPANCY=_COMPILE_OCCUPANCY, + NUM_CTA_IN_CGA=_COMPILE_NUM_CTA_IN_CGA, + ) + ) + return out + + +def _resident_persistent_grid_size( + total_work_tiles: int, + *, + resident_ctas_per_sm: int = _GRID_OCCUPANCY, + cta_cluster: int = _GRID_NUM_CTA_IN_CGA, +) -> int: + """Grid for kernels that explicitly grid-stride over logical work tiles. + + The matching Rust body must contain logic equivalent to: + + for tile_id in (bid_x..total_work_tiles).step_by(grid_x as usize) { ... } + + `total_work_tiles` is the loop's logical bound, not a product of unrelated + axes. For row-wise reductions/norms this is usually cdiv(M, TILE_SIZE_M). + For softmax it is cdiv(M, BM). If the Rust body handles only + get_tile_block_id().0 once, do not use this helper; use + _full_coverage_grid(n_blocks) instead. + + The important perf point is resident budget, not bare SM count. A + grid-stride kernel launched with min(num_sms, total_work_tiles) is often + correct but can leave memory-bound kernels at ~1 CTA/SM. + """ + total_work_tiles = int(total_work_tiles) + resident_ctas_per_sm = max(int(resident_ctas_per_sm), 1) + cta_cluster = max(int(cta_cluster), 1) + num_sms = get_num_sm() + cta_budget = max((num_sms * resident_ctas_per_sm) // cta_cluster, 1) + return min(total_work_tiles, cta_budget) + + +def _persistent_grid_size(m: int, bm: int) -> int: + """Softmax-specific persistent grid helper.""" + num_tiles_m = (m + bm - 1) // bm + return _resident_persistent_grid_size(num_tiles_m) + + +def _full_coverage_grid(n_blocks: int) -> int: + """Grid for non-persistent kernels: one CTA per logical block/row.""" + return int(n_blocks) + + +def _run_ffi( + y, + x2, + m: int, + n: int, + bm: int, + bn: int, + latency: int, + compile_occupancy=None, + compile_num_cta_in_cga=None, +): + """One FFI launch. Used by fixed-config and autotune paths.""" + ffi, lib = bind_kernel_function_cffi(_KERNEL, _FFI_CDEF) + _dev = x2.device + device_id = _dev.index if _dev.index is not None else torch.cuda.current_device() + raw_stream = torch.cuda.current_stream(device=_dev).cuda_stream + + # This template's kernel is persistent, so a resident-CTA capped grid is + # correct here. For row-wise elementwise ports, pass _full_coverage_grid + # unless kernel.rs has a grid-stride loop over rows. + grid_size = _persistent_grid_size(m, bm) + + # Keep the descriptors alive until after the FFI call (cffi frees on GC). + yd = make_tensor_desc(ffi, y) + xd = make_tensor_desc(ffi, x2) + rc = lib.cutile_softmax( + yd, + xd, + int(bm), + int(bn), + int(latency), + _compile_option_value(compile_num_cta_in_cga), + _compile_option_value(compile_occupancy), + int(grid_size), + int(device_id), + int(raw_stream), + ) + check_rc(rc, _FFI_NAME) + return y + + +def _normalize_dim(dim: int, ndim: int) -> int: + if dim < 0: + dim += ndim + return dim + + +def _validate_inputs(x: torch.Tensor, dim: int) -> int: + if not x.is_cuda: + raise ValueError("cutile-rs softmax: input must be a CUDA tensor") + dim = _normalize_dim(dim, x.ndim) + if dim != x.ndim - 1: + raise NotImplementedError( + "cutile-rs softmax template supports only the last dimension. " + "For another op, either implement the layout natively or raise." + ) + if x.dtype not in _DTYPES: + raise NotImplementedError( + f"cutile-rs softmax: dtype {x.dtype} not supported (extend ffi.rs dtype dispatch and make_tensor_desc together)." + ) + return dim + + +def _forward_impl(x: torch.Tensor, dim: int = -1, kernel_configs=None): + """Forward implementation shared by plain and autograd paths. + + Keep expensive copies out of the autotune lambda. A single input contiguous + conversion here is acceptable only when it is part of the kernel contract. + """ + dim = _validate_inputs(x, dim) + + x_contig = x.contiguous() + n = int(x_contig.shape[-1]) + if n == 0 or x_contig.numel() == 0: + return torch.empty_like(x_contig) + m = int(x_contig.numel() // n) + x2 = x_contig.reshape(m, n) + + def launch_with_cfg(cfg): + if int(cfg.BN) < n: + raise NotImplementedError( + f"cutile-rs softmax: kernel requires N<=BN (N={n}, BN={cfg.BN}); " + "use a larger BN or add an inner N-tile loop" + ) + y_local = torch.empty_like(x2) + _run_ffi( + y_local, + x2, + m, + n, + int(cfg.BM), + int(cfg.BN), + int(cfg.LATENCY), + compile_occupancy=getattr(cfg, "OCCUPANCY", _COMPILE_OCCUPANCY), + compile_num_cta_in_cga=getattr(cfg, "NUM_CTA_IN_CGA", _COMPILE_NUM_CTA_IN_CGA), + ) + return y_local + + if kernel_configs: + cfg = SimpleNamespace( + BM=int(kernel_configs.get("BM", kernel_configs.get("bm", 128))), + BN=int(kernel_configs.get("BN", kernel_configs.get("bn", 128))), + LATENCY=int(kernel_configs.get("LATENCY", kernel_configs.get("latency", 4))), + OCCUPANCY=kernel_configs.get("OCCUPANCY", kernel_configs.get("occupancy", _COMPILE_OCCUPANCY)), + NUM_CTA_IN_CGA=kernel_configs.get( + "NUM_CTA_IN_CGA", + kernel_configs.get("num_cta_in_cga", _COMPILE_NUM_CTA_IN_CGA), + ), + ) + return launch_with_cfg(cfg).reshape(x_contig.shape) + + configs = _configs(m, n) + + def kernel_fn(cfg): + # Autotune rule: allocate outputs with torch.empty only. Do not clone, + # zeros, ones, or make additional contiguous copies here. Do not launch a + # separate canary before this lambda; autotune_launch calls this function + # for warmup/timing and then once with the winning cached config. + return launch_with_cfg(cfg) + + result = autotune_launch( + kernel_fn=kernel_fn, + configs=configs, + key=(m, n, x.dtype), + kernel_name=_KERNEL, + ) + return result.output.reshape(x_contig.shape) + + +@register_impl("softmax", backend="cutile-rs") +def softmax(x: torch.Tensor, dim: int = -1, **kwargs): + """Softmax via cutile-rs FFI (forward-only). + + Unsupported semantic cases raise. Do not use a PyTorch forward fallback; + that hides conversion gaps. + """ + kernel_configs = kwargs.get("kernel_configs", None) + _validate_inputs(x, dim) + + if x.requires_grad: + x = x.detach() + + return _forward_impl(x, dim=dim, kernel_configs=kernel_configs) + + +# Copy patterns for common generated wrappers: +# +# Direct-launch / no-autotune attention-style kernels: +# +# Decide per-kernel from Agent A's analysis.json / best_config: +# - reference num_ctas=None / "compiler auto-pick" -> keep _COMPILE_* = None +# and pass the auto sentinel; do NOT pass 1 "to seem safe" (caps resident +# CTAs, regresses high-CTA decode shapes). +# - reference autotunes num_ctas (matmul ships {1,2,4}) -> carry the real +# per-config value through to the FFI. Leaving it None pins the CGA +# cluster to the compiler default (=1) and regresses GEMM ~1.6x. +# +# FFI sentinel pattern: +# +# let mut opts = CompileOptions::default(); +# if occupancy > 0 { opts = opts.occupancy(occupancy); } +# if num_cta_in_cga > 0 { opts = opts.num_cta_in_cga(num_cta_in_cga); } +# let op = kernel(...).compile_options(opts); +# +# Persistent reductions / norm-style launch: +# +# Only use this pattern when kernel.rs contains a grid-stride loop over row +# tiles, e.g.: +# for cur_bid in (bid_x..cdiv(M, TILE_SIZE_M)).step_by(grid_x as usize) { ... } +# +# tile_size_m = ... +# num_row_tiles = _ceil_div(m, tile_size_m) +# grid_size = _resident_persistent_grid_size( +# num_row_tiles, +# resident_ctas_per_sm=_SP_GRID_OCCUPANCY, +# cta_cluster=_SP_GRID_NUM_CTA_IN_CGA, +# ) +# +# If E/C localizes a gap to only this persistent variant while IR matches, +# suspect the host grid before rewriting kernel math: a too-small resident-CTA +# budget (about 1 CTA/SM) is a common cause. Raise _SP_GRID_OCCUPANCY and the +# matching compile occupancy, clamped to row tiles. +# Keep num_cta_in_cga auto unless analysis.json/reference has an explicit +# cluster value. +# +# SiLU-and-mul / row-wise elementwise launch: +# +# n_rows = int(x_flat.shape[0]) +# grid_size = _full_coverage_grid(n_rows) +# +# This is mandatory unless kernel.rs contains: +# for row in (bid_x..n_rows).step_by(grid_x as usize) { ... } +# +# A perf geomean that is much faster than reference while correctness fails is +# often a partial-output launch. Check the first mismatched row against +# get_num_sm() * occupancy. +# +# Autotuned row-wise wrappers: +# +# def kernel_fn(cfg): +# y_local = torch.empty((n_rows, hidden_size), dtype=x.dtype, device=x.device) +# _run_ffi(x2, y_local, compile_occupancy=getattr(cfg, "OCCUPANCY", None)) +# return y_local +# +# result = autotune_launch(kernel_fn=kernel_fn, configs=configs, key=(...), kernel_name=_KERNEL) +# +# Do not allocate a separate y2 and do not call _run_ffi once as a canary +# before autotune_launch. That causes two GPU launches per measured forward on +# cache hits and creates an artificial ~2x CUPTI regression. +# +# Backward is OUT OF SCOPE for cutile-rs in this skill version. +# +# All cutile-rs wrappers must detach grad-enabled inputs and return a +# leaf tensor; the FFI call only emits the forward direction. Test patches +# for ops with backward variants must skip the `cutile-rs` backend for any +# backward-correctness / backward-perf parametrizations. +# +# In-place reference ops: +# +# The detached input may not be modified in-place if the original required +# grad. Clone before the FFI call: +# x = x.detach().clone() if x.requires_grad else x.contiguous() diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/coding-rules.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/coding-rules.md new file mode 100644 index 0000000..9487b81 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/coding-rules.md @@ -0,0 +1,727 @@ +# cutile-rs Conversion Rulebook (coding rules + known gaps) + +Single source for converting kernels to cutile-rs. **Part 1** = prescriptive rules +(Agent B reads before writing `kernel.rs`; treat as a pre-build checklist + source/API +pin). **Part 2** = troubleshooting index + capability gaps (match a failure SIGNATURE, +then apply the fix). Rule numbers (1–49) are stable IDs cited elsewhere as "Rule N" — do +not renumber. Verified against cutile-rs 0.2.0. + +## Index + +**Part 1 — Rules by category** (numbers are Rule IDs; full text in numeric order below) + +| Category | Rules | +|---|---| +| DSL syntax & macro form | 1, 2, 3, 6, 12, 13, 14, 22, 24, 25, 26 | +| Types & inference | 4, 5, 7, 15, 16, 23, 46, 47 | +| Memory / FFI / ABI | 9, 10, 11, 17, 31, 32, 33, 34, 35, 43, 44 | +| Performance | 8, 18, 19, 20, 21, 27, 28, 29, 30, 36, 41, 42, 45, 48 | +| Pipeline / process | 37, 38, 39, 40, 49 | +| `&mut`/output & transpose ownership | 34, 43, 44 (+ `no-mut-tensor-output.md`) | + +**Part 2 — Troubleshooting & gaps** (match the signature) + +- Forward-only port convention +- First-Build Blockers +- Correctness / runtime failure signatures +- Compiler & JIT bugs (with cutile-rs 0.2.0 status) +- Missing APIs / design gaps +- Runtime / wrapper gaps +- Performance gaps +- Process pitfalls +- Minimal triage order + +--- + +# Part 1 — Rules + +## Quick Reference + +| # | Rule | Failure prevented | +|---|------|-------------------| +| 1 | Bind every rank/type-changing call | E0283 inference failures | +| 2 | No scientific notation or qualified constants in JIT values | parser / dense-value failures | +| 3 | Token bindings are ordinary names | macro pattern failures | +| 4 | Bind reductions before later math | inference and shape drift | +| 5 | 2D reduce axis 1 returns rank 1 | RMS and softmax shape mismatch | +| 6 | Avoid const arithmetic in type/shape positions | stable Rust const-generic errors | +| 7 | Integer conversions need explicit intermediate tile types | `exti(scalar_to_tile(...))` E0283 | +| 8 | MMA accumulates in f32; fp32 GEMM uses tf32 when reference does | accuracy/perf loss | +| 9 | Dynamic `Shape` / `Array` literals carry only dynamic entries | bad tensor metadata | +| 10 | FFI borrows PyTorch memory via `borrow_tensor` (ManuallyDrop) / `DevicePointer::from_cu_deviceptr` | invalid frees / SIGABRT | +| 11 | Pipeline strides are full rank; kernel static stride dims are omitted | dump/layout mismatch | +| 12 | Match entry pattern and always use parenthesized `#[cutile::entry()]` | launcher not generated | +| 13 | Runtime `step_by` takes `usize` | loop compile error | +| 14 | No associated const padding such as `Some(E::ZERO)` in JIT calls | JIT expression-position failure | +| 15 | Avoid tuple destructuring before reductions unless return type is annotated | inference failure | +| 16 | Data dtype is generic `` | missing dtype dispatch | +| 17 | One source of truth: cdylib and pipeline include the same `kernel.rs` | generated IR differs from launched kernel | +| 18 | Runtime scalar tiles use `broadcast_scalar` | runtime treated as constant | +| 19 | Use `true_div` where reference uses true division | math/perf mismatch | +| 20 | Match latency, TMA, occupancy and `num_cta_in_cga` to analysis/reference | silent perf drift | +| 21 | Pointer `assume_div_by<16>` only when guaranteed; avoid false scalar assumes | corrupt boundary masks | +| 22 | Mutable tile state crossing `if` is assigned in both branches | dropped compiler state | +| 23 | `ct.Constant[int]` normally becomes a const generic | missed constant folding | +| 24 | Nested loop/if mutations use explicit carry behavior | lost updates | +| 25 | Do not write expression-style `if` returning a `Tile` | unsupported tile-valued control flow | +| 26 | Math APIs take explicit rounding/ftz/nan modes | wrong API shape | +| 27 | GEMM transpose has one owner | timed Python transpose or wrong values | +| 28 | Padded view loads keep TMA when reference uses TMA | ragged tile bugs | +| 29 | Persistent TMA latency is case-specific | persistent-kernel abort / reduction mismatch | +| 30 | Batched TMA descriptors require aligned outer strides | TMA descriptor failure | +| 31 | Accumulator dtype generic slots usually hardcode `"f32"` in FFI | MMA generic mismatch | +| 32 | Rust 2024 exports use `#[unsafe(no_mangle)]` | C ABI build failure | +| 33 | Standalone crates depend on package `cutile-compiler` | proc macro cannot resolve compiler crate | +| 34 | A transpose has one owner: host strides or `partition_permuted`, never both | axis-transpose bugs | +| 35 | Per-kernel ABI pieces must agree | import skips / dtype dispatch failures | +| 36 | Stable const-generic escape hatch is literal entry specialization | repeated nightly-feature loops | +| 37 | Python wrappers are autograd surfaces | grad-enabled tests fail | +| 38 | Run a fixed-config canary before autotune/CUPTI | hidden FFI aborts | +| 39 | In B-only mode B owns wrapper/canaries; in full pipeline D owns host wrapper | waiting for absent agents | +| 40 | Register `cutile-rs` only for dispatcher tests | coverage tests call direct backends | +| 41 | Do not clone/copy unconditionally in wrapper fast paths | CUPTI times copies | +| 42 | After correctness passes, perf cliffs usually start in wrapper/config/pattern drift | chasing math too early | +| 43 | Outputs are never `&mut Tensor`; use read-only tensor + `partition_full_mut` or raw pointer | launch-grid rc=-3 | +| 44 | `partition_full_mut` requires `PartitionMut` import | E0599 | +| 45 | Pointer tiled loops use ceil-div and tail masks when source uses `ct.cdiv` | non-pow2 N aborts/wrong output | +| 46 | Convert const generics to float via `scalar_to_tile` + `convert_tile`, not `N as f32` or `itof` | unsupported cast / signedness writer skew | +| 47 | Annotate `load_ptr_tko` result tuples when mask/padding does not constrain shape | JIT inference failure | +| 48 | Preserve source semantics beyond the one dumped shape | exact-cover IR hiding tail bugs | +| 49 | Final build log records every compile/JIT fix before validator | repeated rediscovery | + +## Rule 1: Bind Rank/Type-Changing Calls + +The macro and Rust inference are fragile when conversions are chained. Bind each step with an explicit type. + +```rust +let p0: PointerTile<*mut E, { [] }> = pointer_to_tile(x_ptr); +let p1: PointerTile<*mut E, { [1] }> = p0.reshape(const_shape![1]); +let ps: PointerTile<*mut E, { [BLOCK] }> = p1.broadcast(const_shape![BLOCK]); + +let bid_i32: Tile = scalar_to_tile(bid_row); +let bid_i64: Tile = exti(bid_i32); +``` + +Do not write: + +```rust +let bid_i64: Tile = exti(scalar_to_tile(bid_row)); +``` + +That exact pattern produces E0283. + +## Rule 2: No Scientific Notation Or Qualified Constants + +Use ordinary decimal literals in JIT-visible calls: + +```rust +let eps: Tile = scalar_to_tile(0.000009999999747378752f32); +let neg = constant(-340282346638528859811704183484516925440.0f32, const_shape![BLOCK]); +``` + +Avoid `1e-5`, `std::f32::consts::*`, `core::f32::*`, and `f32::NEG_INFINITY` inside kernel DSL calls. + +## Rule 3: Token Bindings Are Plain Identifiers + +Use ordinary names and tolerate unused-variable warnings: + +```rust +let (tile, tok): (Tile, Token) = load_ptr_tko(...); +let _ = tok; +``` + +Avoid underscore patterns in macro-sensitive destructures. + +## Rule 4: Bind Reductions Before Math + +```rust +let sum: Tile = reduce_sum(xf, 0i32); +let mean: Tile = sum / denom; +let inv: Tile = rsqrt(mean + eps, ftz::Enabled); +``` + +Apply this to `reduce_sum`, `reduce_max`, `reduce_min`, and reductions feeding `exp2`, `rsqrt`, or division. + +## Rule 5: 2D Reduce Rank + +For `Tile`, reducing axis 1 returns `Tile`, not `[M, 1]`. + +```rust +let row_sum: Tile = reduce_sum(x, 1i32); +let row_sum_col: Tile = row_sum.reshape(const_shape![M, 1]); +``` + +## Rule 6: Const Shapes And Stable Rust + +Simple const shape generics are fine: + +```rust +let cols: Tile = iota(const_shape![BLOCK]); +``` + +Stable Rust rejects const arithmetic in type/shape positions: + +```rust +type Bad = Tile; +``` + +Move arithmetic to value position, host code, or literal-specialized entries. + +## Rule 7: Explicit Integer Conversion Types + +`exti` and `trunci` need both input and output types visible. + +```rust +let pid32: Tile = scalar_to_tile(pid); +let pid64: Tile = exti(pid32); +let pid32_again: Tile = trunci(pid64, overflow::None); +``` + +Use i64 pointer offsets when Python/Triton casts offsets to int64 or the element count may exceed i32. + +## Rule 8: MMA Accumulator And tf32 + +MMA accumulators are normally f32: + +```rust +let mut acc: Tile = constant(0.0f32, const_shape![BM, BN]); +acc = mmaf(a_tile, b_tile, acc); +let out: Tile = convert_tile(acc); +``` + +For fp32 GEMM-family kernels, inspect the reference. If it casts f32 A/B to `ct.tfloat32`, cutile-rs must do the same before `mmaf`: + +```rust +let a_tf32: Tile = convert_tile(a_f32); +let b_tf32: Tile = convert_tile(b_f32); +acc = mmaf(a_tf32, b_tf32, acc); +``` + +## Rule 9: Dynamic Shape / Array Literals + +Only dynamic entries appear in `dims`; static entries live in the type. + +```rust +let shape: Shape<{ [-1, -1] }> = Shape::<{ [-1, -1] }> { dims: &[m, n] }; +let strides: Array<{ [-1, 1] }> = Array::<{ [-1, 1] }> { dims: &[row_stride] }; +``` + +## Rule 10: FFI Memory Ownership + +Borrow PyTorch allocations; do not create owned buffers around them. Tensors +cross the boundary as `*const TensorDesc` (see Rule 35); the FFI unpacks them +without ever freeing PyTorch memory. + +For `&Tensor` entries, borrow via `borrow_tensor::(desc)`, which returns a +`ManuallyDrop>` — the ownership gate. Do NOT call `Tensor::from_raw_parts` +in op code and do NOT `core::mem::forget`; `from_raw_parts` now lives ONLY inside +`borrow_tensor`, and `ManuallyDrop` means the borrowed wrapper is never freed. + +```rust +// desc: &TensorDesc (from unsafe { &*x } after a null check). +let x_t = unsafe { borrow_tensor::(desc) }; // ManuallyDrop> — never freed + +let rc = match op.sync_on(&stream) { + Ok(_) => 0, + Err(e) => { + eprintln!("cutile_{kernel_name} error: {e:?}"); + -3 + } +}; +rc // x_t drops here as a no-op (ManuallyDrop) — no mem::forget needed +``` + +For raw pointer entries, use `DevicePointer::::from_cu_deviceptr(desc.ptr)` in +FFI (no `borrow_tensor` / ownership gate needed — the kernel never holds a Tensor +wrapper). Do not use `transmute`. + +## Rule 11: Static Stride Dims And Pipeline Strides + +Inside kernels, omit static stride entries from `Array::dims`. At the pipeline boundary, pass full logical rank metadata: + +```rust +const X_STRIDES: &[i32] = &[N, 1]; +compiler.strides(&[("x", X_STRIDES)]); +``` + +## Rule 12: Entry Pattern And Attribute Form + +Pick params from reference IR structure: + +- `load_view_tko` / `store_view_tko`: use `&Tensor` and partitions. +- `load_ptr_tko` / `store_ptr_tko`: use raw `*mut E` in the kernel and `DevicePointer` in FFI. +- Mixed structural variants: write one entry per variant. + +Always use the parenthesized function-like attribute: + +```rust +#[cutile::entry()] +pub unsafe fn op_kernel(...) { ... } +``` + +A bare `#[cutile::entry]` (no parentheses) is simply NOT recognized as an entry — +the macro only matches a parenthesized attribute, so no launcher is generated at +all. Always write `#[cutile::entry()]` (no options) or `#[cutile::entry(...)]` +(with options). + +## Rule 13: `step_by` + +```rust +for tile_id in (bid..total_tiles).step_by(grid_x as usize) { + ... +} +``` + +Use `for j in 0i32..num_tiles` when the step is 1. + +## Rule 14: No Associated Const Padding In JIT Calls + +This fails in pipeline JIT: + +```rust +load_ptr_tko(ptrs, ordering::Weak, None::, mask, Some(E::ZERO), Some(token), Latency::<1>); +``` + +Use `None::` padding plus an explicit mask/select pattern, or only use `None::` without select when `analysis.json` proves exact cover for every correctness and perf shape. + +```rust +let (raw, tok): (Tile, Token) = + load_ptr_tko(ptrs, ordering::Weak, None::, Some(valid), None::, Some(token), Latency::<1>); +let xf: Tile = convert_tile(raw); +let zero: Tile = constant(0.0f32, const_shape![BLOCK]); +let xf_safe: Tile = select(valid, xf, zero); +``` + +## Rule 15: Avoid Tuple Destructuring Before Reductions + +If a loaded tile feeds a reduction, annotate the load return and then bind conversion/reduction separately. This keeps diagnostics local. + +## Rule 16: Data Dtype Is Generic + +Use `` for data dtype unless the reference is truly f32-only. Hardcode f32 for accumulators and secondary accumulator generic slots where the reference does. + +## Rule 17: One Source Of Truth + +`kernel.rs` is included by both the cdylib crate and pipeline test. Do not copy kernels into the test. Generated IR must come from the same entry functions that FFI launches. + +## Rule 18: Runtime Scalar Broadcasts + +Use `broadcast_scalar(value, const_shape![...])` for runtime scalar tiles. Do not route runtime values through `constant(...)`. + +## Rule 19: f32 Sigmoid Reciprocal + +For sigmoid-style f32 expressions, use `true_div` if the reference uses true division. + +## Rule 20: Latency And TMA Hints + +Mirror `Latency`, TMA enablement, occupancy, and `num_cta_in_cga` from `analysis.json` and reference IR. Do not add hints because a different kernel used them. + +## Rule 21: Assume Hints + +Apply `assume_div_by::<_, 16>` to raw data pointers when alignment is guaranteed. Avoid scalar `assume_div_by` on runtime dims/strides unless the wrapper proves divisibility for every selected shape. False scalar assumptions corrupt masks and boundary handling. + +## Rule 22: Mutable Tile State Across `if` + +Both branches assign loop-carried tile state: + +```rust +if cond { + acc = new_acc; +} else { + acc = acc; +} +``` + +## Rule 23: Python Constants Become Const Generics + +`ct.Constant[int]` tile sizes and baked shape constants normally become const generics. Runtime args are for true runtime values, not autotune constants folded into the reference IR. + +## Rule 24: Nested Mutation Carry Behavior + +If source IR yields a value through a nested loop/if, keep it as a mutable binding and update it in every path. + +## Rule 25: Avoid Tile-Valued Expression `if` + +Use statement-style mutation for tile-valued control flow. Do not rely on Rust expression `if` returning a `Tile`. + +## Rule 26: Math APIs With Modes + +Many math ops take explicit static modes: + +```rust +let z = addf(a, b, rounding::NearestEven, ftz::Enabled); +let m = maxf(a, b, nan::Disabled, ftz::Enabled); +let p = exp2(x, ftz::Enabled); +let q = fma(a, b, acc, rounding::NearestEven, ftz::Enabled); +``` + +Use named functions when the reference records non-default rounding, FTZ, or NaN behavior. + +## Rule 27: GEMM Transpose Ownership + +A GEMM-family transpose should be native: const generic transpose flags, physical load from original tensor, and `permute` after load when the reference does that. Do not implement timed Python `.t().contiguous()` unless the reference wrapper does it outside the measured region. + +## Rule 28: Padded Loads Keep TMA + +Use `tensor.partition(...)` for padded view loads and keep TMA enabled when the reference uses view/TMA ops. Do not fall back to scalar pointer loads for ragged edge tiles unless the reference is pointer-scatter. + +## Rule 29: Persistent TMA Latency Is Case-Sensitive + +Persistent kernels that autotune many configs often need `None` latency inside the persistent body unless `analysis.json` proves the reference requires a hint. Direct-launch (non-persistent) kernels may need recorded hints. Check the IR. + +## Rule 30: Batched TMA Descriptors + +Batched TMA descriptors require aligned outer strides. If an outer stride is not aligned, fix wrapper padding/fallback. Do not assert false scalar divisibility. + +## Rule 31: Accumulator Generic Slots + +If a kernel has a separate accumulator dtype generic such as `E2`, FFI usually passes `"f32"` for that slot unless the reference says otherwise. + +## Rule 32: Rust 2024 Exports + +```rust +#[unsafe(no_mangle)] +pub unsafe extern "C" fn cutile_{kernel_name}( + out: *const TensorDesc, /* ... more TensorDesc* + scalar args ... */ + device_id: i32, + raw_stream: u64, +) -> i32 { + if out.is_null() /* || ... */ { return -5; } // null descriptor + // dtype: match dtype_str(desc.dtype) { Some(s) => s, None => return -2 } + let device = Device::new(device_id.max(0) as usize)?; // multi-GPU; was Device::new(0) + // ... +} +``` + +Legacy `#[no_mangle]` fails under Rust 2024. Return `-5` on a null descriptor, +`-2` on an unknown dtype code, and derive the device from the FFI `device_id` +arg (`Device::new(device_id.max(0) as usize)`) rather than a hardcoded +`Device::new(0)`, so multi-GPU launches target the tensors' device. + +## Rule 33: Aggregated Cargo Dependencies (crates.io, pinned 0.2.0) + +There is ONE aggregated cdylib crate, `ops/cutile_rs/cutile_kernels/`, that +`include!`s every op's `kernel.rs` + `ffi.rs` (see `output-structure.md`). Its +`Cargo.toml` uses PINNED crates.io dependencies — no path deps, no +`CUTILE_RS_ROOT`: + +```toml +cutile = "=0.2.0" +cutile-compiler = "=0.2.0" +cutile-macro = "=0.2.0" +cuda-core = "=0.2.0" +cuda-async = "=0.2.0" +``` + +The proc macro resolves `cutile-compiler` during expansion. Because everything +is a registry version pin, there is nothing to sed-replace and no per-op +throwaway crate. Build once with `cargo build --release` to produce +`libcutile_kernels.so`. + +## Rule 34: Transpose Has One Owner + +If FFI swaps logical shape/strides, the kernel uses identity partitioning. If the kernel uses `partition_permuted`, FFI passes original physical shape/strides. Do not do both. + +## Rule 35: ABI Pieces Must Agree + +The Python wrapper's cffi `_FFI_CDEF` (with `const TensorDesc*` params), the FFI +export signature, the Rust generic list, test skips, and reports must describe +the same ABI. All ops share ONE `libcutile_kernels.so`; each op's `_FFI_CDEF` +declares only its own `cutile_{kernel_name}` symbol (the shared `TensorDesc` +typedef is prepended by `bind_kernel_function_cffi`). Bind with +`bind_kernel_function_cffi`, pack tensors with `make_tensor_desc`, and check the +return code with `check_rc` — no ctypes `_FFI_ARGTYPES` list. + +## Rule 36: Stable Const-Generic Escape Hatch + +Do not add nightly flags. If const arithmetic in type positions keeps failing, collapse the expression to a literal entry specialization or pass a runtime scalar. + +## Rule 37: Python Wrappers Are Autograd Surfaces + +Perf tests may pass `requires_grad=True` tensors while checking forward only. Detach grad-enabled inputs before FFI. Backward tests skip `cutile-rs` until backward kernels exist. + +## Rule 38: Fixed-Config Canary + +Run a small fixed-config canary before autotune/CUPTI when inputs are easy to construct. It catches FFI ownership aborts and dtype dispatch bugs earlier than Agent D/E. + +## Rule 39: B-Only Ownership + +In B-only mode, Agent B owns compile, wrapper correctness, generated IR, and tiny canaries. In full-pipeline mode, Agent D owns wrapper correctness. + +## Rule 40: Test Registration Scope + +Register `"cutile-rs"` only in test classes that dispatch through `tilegym.ops.{op}`. Do not add it to coverage classes that call Triton/Cutile helpers directly or unsupported backward paths. + +## Rule 41: Wrapper Fast Paths Avoid Copies + +Do not clone/copy/contiguous large tensors unconditionally inside timed wrapper paths. CUPTI times that work. If layout conversion is required, gate it by contiguity/stride and report the cost. + +## Rule 42: Perf Cliff Triage + +After correctness passes, inspect in this order: + +1. FFI launch count and wrapper allocations/copies. +2. `analysis.json` config fidelity: tile sizes, occupancy, `num_cta_in_cga`, latency, variant dispatch. +3. Generated IR op surface: TMA vs pointer ops, `mmaf` dtype, reductions. +4. Kernel math rewrites. + +For GEMM fp32, check tf32 before `mmaf` first. + +## Rule 43: Outputs Are Never `&mut Tensor` + +A `&mut Tensor` output can lock explicit launch grid to inferred partition grid and reject persistent or swizzled grids with rc=-3. Declare outputs as read-only `&Tensor` and write in-body via `partition_full_mut`, or use raw `*mut E` and pointer/tensor-view writes. Host passes the output as just another `*const TensorDesc` and borrows it with `borrow_tensor` (ManuallyDrop, never freed) or `DevicePointer::from_cu_deviceptr`. + +See `references/no-mut-tensor-output.md`. + +## Rule 44: `PartitionMut` Import + +`partition_full_mut` comes from: + +```rust +use cutile::prelude::PartitionMut; +``` + +Import it before rewriting working partition logic. + +## Rule 45: Pointer Tiled Loops Use Ceil-Div And Tail Masks + +If source analysis contains `ct.cdiv`, `ceil_div`, or `check_bounds = (num_tiles*TILE_SIZE != N)`, do not implement loop count as floor division. + +Bad: + +```rust +let num_tiles: i32 = N / TILE_SIZE; +for j in 0i32..num_tiles { ... } +``` + +Good for positive const generics: + +```rust +let num_tiles: i32 = (N + TILE_SIZE - 1) / TILE_SIZE; +for j in 0i32..num_tiles { + let tile_off_s: Tile = scalar_to_tile(j * TILE_SIZE); + let tile_off: Tile = + tile_off_s.reshape(const_shape![1]).broadcast(const_shape![TILE_SIZE]); + let col: Tile = tile_off + cols; + + let n_s: Tile = scalar_to_tile(N); + let n_b: Tile = + n_s.reshape(const_shape![1]).broadcast(const_shape![TILE_SIZE]); + let valid: Tile = cmpi(col, n_b, predicate::LessThan); + + let col64: Tile = exti(col); + let ptrs: PointerTile<*mut E, { [TILE_SIZE] }> = base.offset_tile(row_base + col64); + + let (raw, tok): (Tile, Token) = + load_ptr_tko(ptrs, ordering::Weak, None::, Some(valid), None::, Some(token), Latency::<1>); + let _ = tok; + + let xf: Tile = convert_tile(raw); + let z: Tile = constant(0.0f32, const_shape![TILE_SIZE]); + let xf_safe: Tile = select(valid, xf, z); + acc = fma(xf_safe, xf_safe, acc, rounding::NearestEven, ftz::Enabled); +} +``` + +Use the same `valid` mask for apply-loop loads and stores when columns may exceed the real `N`. A representative IR dump at an exact-cover shape does not prove floor division is safe for correctness shapes. + +## Rule 46: Const Generic Numeric Conversion + +Do not write `scalar_to_tile(N as f32)` inside a kernel. The macro may lower the cast as an unsupported runtime cast. Also avoid `itof` for int-to-float until the writer's `signedness` attribute path is proven for that op. + +Use `convert_tile`: + +```rust +let n_i32: Tile = scalar_to_tile(N); +let n_f32: Tile = convert_tile(n_i32); +``` + +The local API pins are: + +- `scalar_to_tile(scalar: impl Scalar) -> Tile` +- `convert_tile(x: Tile) -> Tile` +- `itof(x, rounding)` exists but can hit `missing attribute 'signedness'` in bytecode writing. + +## Rule 47: Annotate `load_ptr_tko` + +When mask is `None`, padding is `None`, or shape is generic, annotate the tuple: + +```rust +let (xv, tok): (Tile, Token) = load_ptr_tko( + ptrs, + ordering::Weak, + None::, + Some(valid), + None::, + Some(token), + Latency::<1>, +); +``` + +The source signature is: + +```rust +load_ptr_tko(source, ordering, scope, mask, padding_value, token, Latency::) + -> (Tile, Token) +``` + +`store_ptr_tko` takes `(destination, value, ordering, scope, mask, token, latency)` and returns `Token`. + +## Rule 48: Preserve Source Semantics Beyond The Dumped Shape + +Agent A often dumps one representative shape per structural variant. If `analysis.json.ops_used` or variant text says `ct.cdiv`, `next_power_of_2`, `check_bounds`, `padding_mode=ZERO`, or dtype-specific lowering, preserve that source behavior even when the chosen IR shape folds to a simpler constant. + +Example: a reference dumped at a shape where the tile count divides evenly (floor == ceil) hides the tail; correctness then fails on a shape where floor division drops the partial tile. Source semantics win. + +## Rule 49: Build Log Is Part Of The Fix + +Before returning `VERDICT: COMPILED`, write `reports/build_log.md` with: + +- each cargo build attempt and top rustc error +- each pipeline/JIT compile iteration +- offending snippet +- fix applied +- whether the fix is now covered by this rulebook or a known gap + +Then run `validate_agent_b.sh` and include its output in ``. + + +--- + +# Part 2 — Troubleshooting Index & Capability Gaps + +This part was merged in from the former separate known-gaps troubleshooting index. +Match the exact compiler, JIT, Python wrapper, +pytest, or scorer signature before changing kernel math. In B-only eval mode, absent +Agent C/D/E/F reports are expected. Entries for APIs that cutile-rs 0.2.0 added +(`get_index_space_shape`, `divi rounding`, unsigned `cmpi`) and the +never-existent `make_strided_view` were removed; the compiler-bug table carries a 0.2.0 +status column. + +## Forward-only port (skill convention) + +cutile-rs is **forward-only** in this skill version. Convention, not compiler limit: + +- `kernel.rs` / `ffi.rs` emit ONLY forward; never write a backward kernel. +- `wrapper.py` detaches grad-enabled inputs (`if x.requires_grad: x = x.detach()`) and returns a leaf tensor (no `grad_fn`). Do NOT define a `torch.autograd.Function`. +- `test_{kernel_name}.py` patches MUST skip backward parametrizations for cutile-rs: + ```python + if backend == "cutile-rs": + pytest.skip("cutile-rs is forward-only (this skill version)") + ``` + Apply to any `test_op_backward[...]` and any `test_perf[...]` that times backward / combined fwd+bwd. Forward `test_op[cutile_rs-...]` and forward-only `test_perf[cutile_rs-...]` remain enabled. + +**Why**: a PyTorch-formula backward inside `autograd.Function` adds large wrapper-layer slowdown to combined fwd+bwd ratios, swamping Agent C's forward-kernel perf signal. This skill version ports forward only. + +**`detach()` cost**: ~1-5us per call, no CUDA work, no allocation. Apply unconditionally when `requires_grad=True`. + +**In-place reference ops** (those that mutate an input in place): after detach, do `x.detach().clone()` if the FFI mutates in-place — autograd-detached tensors share storage with the original. + +## First-Build Blockers + +| Failure signature | Bad pattern | Fix | +|-------------------|-------------|-----| +| `error[E0433]: cannot find module or crate cutile_compiler` under `#[cutile::module]` | `cutile_kernels` crate lacks the `cutile-compiler` package dependency | Add `cutile-compiler = "=0.2.0"` to `cutile_kernels/Cargo.toml` (Rule 33) | +| `error[E0432]` / `E0433` for `DevicePointer`, `CompileOptions`, or `DType` | Stale imports such as `cutile::DevicePointer` or `cutile::tensor::DType` | Use the import block in `agents/agent_b.md` | +| `built-in f16 is unstable`; `bf16 is not in scope` | Used Rust builtin or missing half imports | Import `cutile::core::{f16, bf16}` if concrete half names are needed | +| `call to unsafe function make_tensor_view is unsafe` | Unsafe helper called without a block | Wrap `unsafe { make_tensor_view(...) }` | +| `Pointers can only be used in unsafe kernel entry points` | Raw pointer entry declared safe | Use `#[cutile::entry()] unsafe fn` | +| `kernel_entry_generator.rs:70` / stride unwrap | `&Tensor` pipeline test omitted full-rank strides | Pass `.strides(...)` for every tensor rank, including trailing `1` | +| Validator OK but pytest all skipped | Backend availability or shared-lib loader missing | Wire `selector.py`, `ops/cutile_rs/__init__.py`, and `bind_kernel_function_cffi` | + +## Failure Signatures + +| Symptom | Smallest trigger | Fix | +|---------|------------------|-----| +| `error: generic parameters may not be used in const operations` at `#[cutile::module]`, followed by `help: add #![feature(generic_const_exprs)]`; adding the feature gives `E0554` on stable | Pointer-scatter variant with shape-driving const generic such as `TILE_SIZE` inside `Tile<..., {[TILE_SIZE]}>`, loops, masks, and extra derived const generics in the same generated module | Do not switch to nightly. Use literal entry specialization or a conservative literal block-size entry and keep `N` runtime. Dispatch in FFI. See coding-rules Rule 36. (Confirmed still present in 0.2.0: crate is stable Rust 1.89, no `generic_const_exprs`.) | +| pytest says `libcutile_kernels.so not found` after Agent B reported work | Cargo build of the aggregated `cutile_kernels` crate failed but wrapper/test registration was still written | Treat missing `.so` as build failure, not Python failure. Fix cargo first; do not spend time changing wrapper logic. | +| `NotImplementedError: cutile-rs {op}: backward not implemented` | Wrapper rejects `input.requires_grad`, while tests create `requires_grad=True` inputs or pass `gradient=...` | Do NOT reject `requires_grad`. Detach grad-enabled inputs (`if x.requires_grad: x = x.detach()`) and return a leaf tensor. Backward is out of scope; the backward test must skip cutile-rs. | +| Forward values close but gradient mismatch is huge | A backward-correctness test reached cutile-rs, which is forward-only | Skip the cutile-rs backend for backward parametrizations (`pytest.skip("cutile-rs is forward-only (this skill version)")`). Do not add a `torch.autograd.Function`. | +| Fatal Python error: Aborted in `tilegym/ops/cutile_rs/{kernel_name}.py:_run_ffi` during autotune | Bad persistent TMA config reaches CUPTI; common cause is `Some(LATENCY)` in a persistent grid-stride TMA loop | Run fixed-config canary before autotune. In persistent TMA loops, pass `None` latency for every view load/store unless analysis/reference IR prove a direct-launch exception. | +| Correctness: 0% matched, output near zero | FFI builds an input with swapped logical shape/strides AND the kernel also applies `partition_permuted` → double transpose | Pick one transpose owner. If the kernel permutes, pass the original physical shape/strides; if the FFI swaps the layout, the kernel uses identity `partition`. | +| Row-wise correctness mismatch begins at row `get_num_sm()*occupancy` and perf ratio is implausibly fast | Wrapper capped `grid_size = min(n_rows, num_sms * occupancy)` but kernel handles only `get_tile_block_id().0` once | Either launch `grid=n_rows` or add an explicit persistent loop `for row in (bid_x..n_rows).step_by(grid_x as usize)`. Occupancy is a compile/runtime hint, not a substitute for logical grid coverage. | +| Perf geomean is very fast, but correctness failed | Host scorer can still run perf after a wrong partial-output kernel | Ignore the speed until correctness passes. Look for unwritten rows/tiles, output sentinels, wrapper grid caps, and skipped boundary masks. | +| Fatal Python error: Aborted in `tilegym/ops/cutile_rs/{kernel_name}.py:_run_ffi` only on the largest resource (memory/batch) perf case after smaller cases pass | Resource-sensitive test case reached cutile-rs even though only unsupported dtype/transpose skips were mirrored | Mirror the reference backend's resource/OOM skips for cutile-rs when the same allocation and launch shape are used, or add a smaller supported cutile-rs perf surface. Do not let an intentionally unsupported giant case enter FFI. | +| Correctness command selects `test_perf` cases and aborts before coverage tests finish | `python -m pytest tests/ops/test_{kernel_name}.py -k cutile_rs -x` selects every cutile-rs parametrized test, not only `test_op` | Before adding `cutile-rs`, inspect all selected classes and add skips before dispatcher/FFI calls for unsupported resource, dtype, transpose, or backend-specific coverage cases. | +| Perf geomean present but correctness failed | Host scorer runs perf even if correctness failed | Fix correctness first. Do not optimize a failing wrapper/kernel unless the perf log identifies an independent abort. | +| `AttributeError` / undefined `cutile_{kernel_name}` symbol at load | `_KERNEL`, `_FFI_NAME`, or the exported Rust symbol disagree | Make all three consistently `cutile_{kernel_name}` / `{kernel_name}`. | +| FFI returns an error code but Python keeps running | Wrapper omits `check_rc(rc, _FFI_NAME)` after the call | Always `check_rc(rc, _FFI_NAME)` immediately after every FFI call. | +| Process aborts from a Rust panic across `extern "C"` | FFI panicked before returning `i32`; Python cannot catch it | Guard shapes/dtypes/resource cases in Python before FFI; return `-N` for expected errors instead of asserting/panicking. | + +## Compiler And JIT Bugs Currently Open + +Verified against cutile-rs **0.2.0** source (built + grepped from source). +Most error paths still exist in the 0.2.0 compiler (status column cites the file). Two +look RELAXED in 0.2.0 (re-test before relying on the workaround). `generic parameters may +not be used in const operations` is a rustc-stable error, not a cutile check — it persists +because the crate is stable Rust 1.89 with no `generic_const_exprs`. + +| Failure signature | Smallest trigger | Workaround | 0.2.0 status | +|-------------------|------------------|------------|--------------| +| `binary Div requires operands of the same type` with `[M]` vs `[M,1]` | Annotating `reduce_sum(x, 1)` as `[M,1]` | Bind actual `[M]`, then reshape to `[M,1]` | STILL PRESENT (`compile_binary_op.rs`) | +| `return type is missing a compiled tile type` | Chained `pointer_to_tile(...).reshape(...).broadcast(...)` | Split into explicit typed bindings | STILL PRESENT (`compile_cuda_tile_op.rs`) | +| `Unexpected optimization hint key` | `simt_num_warps_in_cta` hint | Remove it; tileiras chooses warps | STILL PRESENT (`hints.rs`; whitelist = `num_cta_in_cga`/`occupancy`/`max_divisibility`) | +| JIT `Unexpected dense value` | Runtime or qualified value in `constant(...)` | Use `broadcast_scalar` for runtime values; use decimal literal for constants | STILL PRESENT (`compile_cuda_tile_op.rs`) | +| Unary expression not supported | `if !EVEN_N` or unary const bool in macro body | Invert branch: `if EVEN_N { x=x; } else { ... }` | STILL PRESENT (`compile_expression.rs`/`generics.rs`) | +| `qualified paths not supported` | `Some(f32::NEG_INFINITY)` or `constant(std::f32::consts::LOG2_E, ...)` | Use decimal literals and explicit pad/select tiles | LIKELY RELAXED — hard error string is gone; 0.2.0 added `dense_const_path_parts`/`dense_module_const_path_value` (`compile_cuda_tile_op.rs`). Re-test qualified consts before assuming failure. | +| Value assigned in `if` not visible after block | Mutable tile only assigned in one branch | Add `else { x = x; }` or split branches | MAYBE RELAXED — `compile_expression.rs` now has if-branch `carry_vars` handling. Re-test. | +| Type inference failure after tuple destruct in loop | `let (tile, tok) = load_ptr_tko(...)` directly feeding reduce in loop | Predefine typed lets or bind explicitly | UNVERIFIED (no single error string; needs a live repro) | +| Store partition view lacks `padding_value = zero` | Mutable partition stores | Accept; ensure load-side partition padding and store masks/bounds are correct | UNVERIFIED (`load_ptr_tko` takes `padding_value: Option`; associated-const lowering not retested) | + +## Missing APIs / Design Gaps + +(0.2.0 added `get_index_space_shape`, `divi rounding`, and unsigned `cmpi`; `make_strided_view` never existed and `make_tensor_view(base, shape, strides, token)` covers strided views — all removed from this table.) + +| Missing or limited area | Workaround | +|-------------------------|------------| +| Explicit Triton `num_stages` knob | Use per-op latency only when reference IR has it and Rule 29 does not forbid it | + +`&mut Tensor` kernel inputs/outputs: still grid-coupling in 0.2.0 (accepted only via a `Partition` whose shape pins the launch grid). The rule + workaround are Rule 43/44 (Part 1 above) and `references/no-mut-tensor-output.md`. + +## Performance Gaps + +(`assume_div_by` pointers-only and "occupancy is not a grid cap" rules live once as Rule 21 / Rule 20 in Part 1 above; the occupancy failure SIGNATURE stays in the Failure Signatures table above.) + +| Gap | Impact | Workaround | +|-----|--------|------------| +| TMA disabled when reference uses TMA | Often 1.5x-10x slower | Keep `tma::Enabled`; fix layout/strides instead | +| Scalar `ct.Constant[int]` left runtime | Can be several times slower | Use const generics unless Rule 36 compile failure occurs | +| f16 MMA accumulator | Slower and less accurate | Accumulate in f32 | +| `/` for f32 sigmoid | Slow division | Use `true_div` | +| Extra `clone`, `zeros`, `ones`, or `contiguous` inside autotune lambda | CUPTI times extra kernels/copies | Allocate outputs with `torch.empty` only inside `kernel_fn` | +| Autotune tries invalid configs first | Abort or bad cache | Fixed-config canary and config filtering before `autotune_launch` | +| Large-shape resource cases inherited blindly from test_perf | Abort or timeout after many passing cases | Mirror reference skip rules when cutile-rs has the same allocation footprint, and document any intentionally unsupported perf cases in `reports/agent_b.md`. | + +## Process Pitfalls + +| Pitfall | Prevention | +|---------|------------| +| Chasing missing Agent C/D/E/F reports in B-only | Ignore; only Agent B runs | +| Reporting `COMPILED` after cargo failed | Check `.so`, symbol, pipeline compile, and tiny wrapper canary | +| Reading perf before correctness | Correctness first; perf only after wrapper/kernel semantics pass | +| Treating wrapper errors as compiler gaps | If traceback is in `tilegym/ops/cutile_rs/*.py`, fix wrapper | +| Treating near-zero or partially zero output as math first | Check grid coverage, host tensor shape/stride, and transpose ownership first | +| Changing kernel math to fix autograd | Autograd is Python wrapper behavior unless backward kernel is required | +| Adding cutile-rs to every backend list by rote | Register only dispatcher paths or skip backend-specific/resource-excluded cases before FFI | +| Running full pytest matrix after each edit | Use a tiny canary, then `python -m pytest tests/ops/test_{kernel_name}.py -k cutile_rs -x --timeout=60` | + +## Minimal Triage Order + +1. Cargo build and `.so` symbol. +2. Pipeline compile and generated MLIR. +3. Entry pattern validator. +4. Python import/selector/backend registration. +5. Fixed-config wrapper canary on a small shape. +6. Grid/resource canary: for row-wise kernels, use a shape with more rows than `get_num_sm()*occupancy`; for resource-heavy kernels, run the largest registered case once outside CUPTI or skip it before FFI if it is intentionally unsupported. +7. Correctness pytest with `-k cutile_rs -x --timeout=60`. +8. Perf pytest only after correctness. + +Record exact commands, pass/fail result, intentional pytest skips, and any grid/resource decisions in `reports/build_log.md` and `reports/agent_b.md`. + +## Verification Provenance + +Checked against the local cutile-rs tree: + +- `cutile/src/_core.rs`: `scalar_to_tile`, `convert_tile`, `itof`, `exti`, `select`, `fma`, `cmpi`, `andi`, `load_ptr_tko`, `store_ptr_tko`, `load_view_tko`, `Latency`. +- `cutile-compiler/src/compile_api.rs`: `KernelCompiler::generics`, `strides`, `spec_args`, `grid`, `options`. +- `cutile-compiler/src/specialization.rs`: `SpecializationBits`, `DivHint`. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/env-block.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/env-block.md new file mode 100644 index 0000000..d95f798 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/env-block.md @@ -0,0 +1,87 @@ +# Environment Block for Agent Prompts + +## Required env vars + +Everything now lives in the tilegym checkout — there is no separate cutile-rs +checkout and no `CUTILE_RS_ROOT`. The kernel Rust and the aggregated +`cutile_kernels` crate live under `$TILEGYM_PATH/src/tilegym/ops/cutile_rs/`, +build against crates.io pins (`cutile="=0.2.0"`, …; Rule 33), and produce a +single `libcutile_kernels.so`. Set these once per shell session; +`scripts/preflight.sh` verifies them. + +```bash +# ─── Git repo ───────────────────────────────────────────────────────────── +export TILEGYM_PATH= # tests + Python wrappers + cutile_rs Rust + +# ─── Toolchain (existence-checked) ──────────────────────────────────────── +export TILEIRAS_BIN= # cuTile compiler binary (IR->cubin at launch) +export CUDA_TILE_OPT_BIN= # MLIR canonicalizer (Agent B/C) +export TRITON_TILEIR_PYTHONPATH= # Triton-TileIR backend bindings +export CUDA_TOOLKIT_PATH=/usr/local/cuda # CUDA root (cuda-bindings build.rs); default /usr/local/cuda +``` + +> **Building the `cutile_kernels` crate.** The crate uses PINNED crates.io +> dependencies (no path deps), so there is nothing to sed-replace. The Python +> loader autobuilds it on first use — `CUTILE_RS_AUTOBUILD` is ON by default; +> set it to `0` to use a prebuilt library. Override the crate directory (default +> `$TILEGYM_PATH/src/tilegym/ops/cutile_rs/cutile_kernels`) with +> `CUTILE_RS_KERNELS_DIR`. The `cuda-bindings` build step reads +> `CUDA_TOOLKIT_PATH` (default `/usr/local/cuda`); `tileiras` lowers IR to cubin +> at launch time. + +## CRITICAL: tileiras alignment + +cutile-python and cutile-rs MUST use the SAME `tileiras`. The binary pointed at +by `$TILEIRAS_BIN` must come BEFORE `/usr/local/cuda/bin/` on PATH so cuTile-py +JIT picks it up. + +## Full manual environment + +```bash +# Required env vars (incl. CUDA_TOOLKIT_PATH) — see top of file. + +# tileiras must precede system CUDA on PATH +export PATH=$(dirname "$TILEIRAS_BIN"):$PATH + +# Triton-TileIR bindings + tilegym source on PYTHONPATH +export PYTHONPATH=$TRITON_TILEIR_PYTHONPATH:$TILEGYM_PATH/src:$PYTHONPATH + +export ENABLE_TILE=1 +export TILEIR_ENABLE_FTZ=1 +export TILEIR_ENABLE_APPROX=1 + +# Benchmark defaults +export CUPTI=1 +export WARMUP=100 +export REP=50 +export MIN_REP=2 + +# CUDA_TOOLKIT_PATH is one of the required vars set at the top of this file. +``` + +## Tool paths + +```bash +# tileiras (for cutile-rs JIT → cubin, and cutile-python compilation) +which tileiras # MUST be: $TILEIRAS_BIN + # NOT: /usr/local/cuda/bin/tileiras + +# cuda-tile-opt (for canonicalize/CSE on dumped IR — used by Agent B / Agent C) +echo "$CUDA_TILE_OPT_BIN" +``` + +## Verification + +```bash +which tileiras +# Expected: $TILEIRAS_BIN + +python -c "import triton; print(triton.__file__)" +# Expected: contains $TRITON_TILEIR_PYTHONPATH + +python -c "import tilegym; print('OK')" +# Expected: OK + +cd "$TILEGYM_PATH/src/tilegym/ops/cutile_rs/cutile_kernels" && cargo build --release 2>&1 | tail -1 +# Expected: Finished `release` profile ... (produces libcutile_kernels.so) +``` diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/ir-diff-checklist.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/ir-diff-checklist.md new file mode 100644 index 0000000..aa0a1b3 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/ir-diff-checklist.md @@ -0,0 +1,46 @@ +# IR Diff Checklist + +When comparing original cudatile IR with cutile-rs generated IR, **every item must be checked and resolved**. + +## Pre-check: Correct kernel variant? + +- [ ] **Confirmed actual kernel name** via nsys/torch.profiler (NOT assumed from source code) +- [ ] **Confirmed launch config** (grid, block) matches between dump and runtime +- [ ] **Matched dump file** to confirmed kernel name (CUDA_TILE_DUMP_TILEIR dumps ALL variants) + +## MUST match (correctness + performance critical) + +- [ ] `optimization_hints` on entry — `num_cta_in_cga`, `occupancy` +- [ ] `strides=[?,1]` — last dim must be constant 1. For `&Tensor` params: auto. For raw ptr: use `Array::<{[-1, 1]}> { dims: &[only_dynamic_val] }` +- [ ] accumulator type — `f32` for MMA, not `f16` +- [ ] `ftof` conversion — f32→f16 before store if accumulator is f32 +- [ ] `mmaf` op types — input f16, accumulator f32 +- [ ] token count — match original (usually 1 shared `make_token`) +- [ ] for loop step — matches original step size +- [ ] partition tile shapes — `(BM x BK)`, `(BK x BN)`, `(BM x BN)` match +- [ ] `padding_value` — `neg_inf` or `zero` where original has it (now supported since commit `4ba3a83`) +- [ ] `maxf`/`minf` — should NOT have `rounding_mode` attribute (exact ops) +- [ ] grid size — persistent kernel grid must match (e.g., `NUM_SM * occupancy`) +- [ ] block size / num_warps — cutile-rs doesn't support `simt_num_warps_in_cta`, compiler chooses automatically + +## OK to differ (semantic equivalent) + +- [ ] `divi rounding` → `(a + b - 1) / b` — same semantics +- [ ] `mini signed` → `if/else` — same semantics +- [ ] `get_index_space_shape` → manual `(dim + tile - 1) / tile` — same semantics +- [ ] negative `remi` correction (xori+andi+select) → Rust `%` — may differ for negative values but blockId is always non-negative +- [ ] `for ... step %assumed_var` vs `for ... step %raw_var` — cutile-rs uses assume-wrapped variables + +## Expected noise (ignorable) + +- [ ] Dead `constant ` / `constant ` — from macro expansion, use `--canonicalize` to remove +- [ ] Extra `assume bounded<0, ?>` on blockId — auto-generated by cutile-rs +- [ ] Extra `assume bounded<0, N>` on for loop index — auto-generated +- [ ] Extra `assume bounded<0, ?>` on gridSize — auto-generated + +## Known gaps (cannot fix, record) + +- [ ] `make_strided_view` — not supported in cutile-rs; use `tview.partition(...)` / `tview.partition_mut(...)` (lowers to `make_partition_view` in IR) +- [ ] `tensor_dim_factor` hint — may not appear in IR +- [ ] Debug info / source locations — cutile-rs doesn't generate +- [ ] `simt_num_warps_in_cta` — not supported as optimization hint diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/no-mut-tensor-output.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/no-mut-tensor-output.md new file mode 100644 index 0000000..17361b5 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/no-mut-tensor-output.md @@ -0,0 +1,79 @@ +# Rule: kernel outputs are NEVER `&mut Tensor` (raw `*mut E` or read-only `&Tensor` only) + +This rule is a hard ban on +`&mut Tensor` as a kernel entry parameter for outputs. All converted +kernels must write through ONE of: + +1. **read-only `&Tensor` + in-kernel `partition_full_mut`** + (worked example: `examples/bmm`) — preferred for + tiled / persistent / transposed outputs; keeps the ergonomic Tensor boundary + (host borrows the descriptor with `borrow_tensor::(desc)` → + `ManuallyDrop>`, never freed) while emitting NO grid lock. +2. **raw `*mut E` device pointer + in-kernel `make_tensor_view`** (worked + example: `examples/softmax`) — when no Tensor + boundary is wanted (host wraps `desc.ptr` with + `DevicePointer::::from_cu_deviceptr`). + +`&mut Tensor` (and `MappedPartitionMut` / `Partition<&mut Tensor>` host args) are +**banned**. Why follows. + +## Root cause — the launcher macro branches on `&` vs `&mut` + +`cutile-macro/src/kernel_launcher_generator.rs` generates DIFFERENT host launch +code per parameter, keyed on `ty.mutability.is_some()`: + +**`&mut Tensor` (mutable) branch:** +```rust +launch_grid_expr_strs.push("KernelOutputStored::grid(&{var})?"); // contributes inferred grid +// shape validation: EXACT equality, -1 NOT substituted +kernel_launch_assert(valid_shape == &given_shape, "partition shape mismatch ..."); +``` + +**`&Tensor` (read-only) branch:** +```rust +KernelInputStored::push_kernel_args(...) // treated as INPUT; launch_grid_expr_strs stays EMPTY +// shape validation: -1 substituted with the actual dim +let valid_shape_mixed = zip(valid, given).map(|(e,g)| if e == -1 { g } else { e }); +``` + +So writing `&mut` flips TWO switches at once: + +1. **Grid lock.** The mutable output is pushed into `inferred_grids`, and + `infer_launch_grid` → `validate_grids` forces the launch grid to EQUAL the + output's partition grid **on every dim, in descriptor dim order** + (`(u32,u32,u32)` full-tuple `!=` check, `tile_kernel.rs`). The kernel's + `get_tile_block_id().k` must therefore map to descriptor dim `k` for all `k`. + Any transpose (e.g. attention wanting `head = bid(0)` while the output is + `[B, num_head, …]` so dim0 = batch) is rejected with + `"Specified launch grid does not match inferred tensor partition grid"` + (FFI rc=-3) — output heads/rows unwritten, launcher reject. Persistent grids (`num_blocks < num_tiles`) are also + impossible on this path. + +2. **`-1` reject trap.** The mutable branch shape-matches EXACTLY (no `-1` + substitution), so `out: &mut Tensor` is rejected for every + config (`partition shape mismatch. Expected [-1,..], got [..]`). A mutable + output is forced to carry concrete tile dims, which then re-introduces #1. + +A read-only `&Tensor` (mutated in-body via `partition_full_mut`) and a raw +`*mut E` pointer both emit **no** `KernelOutputStored::grid`, so neither +contributes to `inferred_grids` → **no `validate_grids` lock** → the kernel +launches any grid it wants and uses the real (possibly transposed / persistent) +block id. `partition_full_mut` is a DEVICE-side op the host launcher never sees, +so it does not re-trigger the lock. + +## Evidence (converted kernels) + +| kernel | output param | grid lock? | result | +|---|---|---|---| +| tiled / transposed output (`examples/bmm`) | `&Tensor<{[-1,…]}>` + `partition_full_mut` | none | PASS | +| row-wise output (`examples/softmax`) | raw `*mut E` + `make_tensor_view` | none | PASS | +| 2D output whose grid happens to align | `&mut Tensor<{[BM,BN]}>` | locked (dims happen to align) | PASS by luck | +| **multi-head / transposed output** | `&mut Tensor<{[1,TILE_H,TILE_D]}>` | **locked → head/batch transpose** | **FAILED** (launcher reject) | + +## Notes + +- Persistent + CGA perf (`occupancy`, `num_cta_in_cga`) is reachable on the + raw-pointer / read-only `&Tensor` path, so banning `&mut Tensor` costs no + performance. +- `validate_agent_b.sh` currently FAILs the `-1`-shaped mutable-output sub-case; + a blanket check on ANY `&mut Tensor` entry param would enforce this policy in full. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/op-mapping.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/op-mapping.md new file mode 100644 index 0000000..19a73af --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/op-mapping.md @@ -0,0 +1,487 @@ +# Op Mapping: CUDA Tile IR -> cutile-rs + +**Verified** against the provided `cutile-rs` checkout. + +> **Critical**: many math ops take explicit `rounding::Mode` + `ftz::Mode` +> static-typed args via `static_params`. Old "single-arg" call shapes will not +> compile. Also always use `#[cutile::entry()]` with parentheses; a bare +> `#[cutile::entry]` (no parentheses) is NOT recognized as an entry — the macro +> only matches a parenthesized attribute, so no launcher is generated. + +## Block ID / Grid intrinsics (read this FIRST) + +**Canonical block-id**: `get_tile_block_id() -> (i32, i32, i32)`. It returns +`(bid.x, bid.y, bid.z)` regardless of grid dimensionality. + +```rust +// 1D grid +let pid = get_tile_block_id().0; + +// 2D grid +let (m_idx, n_idx, _) = get_tile_block_id(); + +// 3D grid +let pid: (i32, i32, i32) = get_tile_block_id(); +``` + +Anti-patterns: + +```rust +let bid = block_id::<0>(); +let bid = cuda_block_id(); +let bid = bid(0); +``` + +Companion: `get_num_tile_blocks() -> (i32, i32, i32)`. + +## Entry Attributes And Host Launchers + +Use `#[cutile::entry()]` when the entry has no options: + +```rust +#[cutile::entry()] +pub unsafe fn kernel_name(...) { + ... +} +``` + +Use option form only when the reference requires options: + +```rust +#[cutile::entry(print_ir = false, unchecked_accesses = true)] +pub unsafe fn kernel_name(...) { + ... +} +``` + +Do not use bare `#[cutile::entry]` (no parentheses). The macro only matches a +parenthesized attribute, so the bare form is NOT recognized as an entry at all — +no host launcher is generated, so `.generics()` is unavailable on the result. +The compiler registry error text also expects `#[entry(...)]`. Always write +`#[cutile::entry()]` or `#[cutile::entry(...)]`. + +## Pointer Scatter/Gather + +| IR Op | cutile-rs Rust | +|-------|----------------| +| `pointer_to_tile %ptr` | `pointer_to_tile(ptr)` -> `PointerTile<*mut E, {[]}>` | +| `reshape %ptr : tile -> tile<1xptr>` | `ptr_tile.reshape(const_shape![1])` | +| `broadcast %r : tile<1xptr> -> tile` | `ptr_1d.broadcast(const_shape![N])` | +| `offset %ptrs, %offsets` | `ptrs.offset_tile(offsets)` | + +Split pointer tile chains into typed bindings. + +```rust +let p0: PointerTile<*mut E, { [] }> = pointer_to_tile(ptr); +let p1: PointerTile<*mut E, { [1] }> = p0.reshape(const_shape![1]); +let ps: PointerTile<*mut E, { [N] }> = p1.broadcast(const_shape![N]); +``` + +### `load_ptr_tko` + +```rust +pub fn load_ptr_tko< + E: ElementType, + const S: [i32; N], + O: ordering::LoadMode, + Sc: scope::Mode, + const CYCLES: u32, +>( + source: PointerTile<*mut E, S>, + memory_ordering: O, + memory_scope: Option, + mask: Option>, + padding_value: Option, + token: Option, + latency: Latency, +) -> (Tile, Token) +``` + +Call: + +```rust +let (tile, tok) = load_ptr_tko( + ptrs, ordering::Weak, None::, + Some(mask), Some(pad), None, Latency::<0> +); +``` + +### `store_ptr_tko` + +```rust +pub fn store_ptr_tko< + E: ElementType, + const S: [i32; N], + O: ordering::StoreMode, + Sc: scope::Mode, + const CYCLES: u32, +>( + destination: PointerTile<*mut E, S>, + value: Tile, + memory_ordering: O, + memory_scope: Option, + mask: Option>, + token: Option, + latency: Latency, +) -> Token +``` + +Call: + +```rust +let tok = store_ptr_tko( + ptrs, val, ordering::Weak, None::, + Some(mask), Some(tok), Latency::<0> +); +``` + +## Tensor / Partition + +### Token creation + +IR `cuda_tile.make_token` maps to `new_token_unordered()`. + +```rust +let token: Token = new_token_unordered(); +``` + +Use this token for manual `make_tensor_view(...)` construction. Most TMA/view +kernels should prefer typed `&Tensor` parameters and let the entry macro build +the tensor view. + +### `make_tensor_view` + +```rust +pub unsafe fn make_tensor_view( + base: PointerTile<*mut E, {[]}>, + shape: Shape, + strides: Array, + token: Token, +) -> Tensor +``` + +Required `unsafe`. + +### Reading dynamic tensor dimensions + +For `&Tensor` entry parameters, prefer metadata shape accessors when only scalar +dimensions are needed: + +```rust +let a_shape: Shape<{ [-1, -1] }> = a.shape(); +let b_shape: Shape<{ [-1, -1] }> = b.shape(); +let m: i32 = a_shape[0]; +let n: i32 = b_shape[1]; +``` + +`tensor.shape()` uses `get_tensor_shape_meta`; it does not emit a real +`cuda_tile.get_tensor_shape` op. Avoid `get_tensor_shape(a)` unless the reference +IR itself has that op, because it adds real tensor-view layout attributes. + +### Partition view construction + +Method form is preferred for real loads/stores: + +```rust +let part: Partition = a.partition(const_shape![BM, BK]); +let part_p: Partition = + a.partition_permuted(const_shape![BK, BM], const_array![1, 0]); +let mut out: PartitionMut = + unsafe { c.partition_full_mut(const_shape![BM, BN]) }; +``` + +Low-level `make_partition_view` is the intended exception when the reference has +a structurally distinct no-padding metadata partition, for example TMA GEMM +querying `get_index_space_shape` separately from padded loads: + +```rust +let token = get_tensor_token(a); +let a_iter: Partition = make_partition_view( + a, + const_shape![BM, BK], + padding::None, + dim_map::Identity, + token, +); +let a_space: [i32; 2] = get_index_space_shape(&a_iter); + +let a_part: Partition = a.partition(const_shape![BM, BK]); +let b_part: Partition = b.partition(const_shape![BK, BN]); +``` + +### Mutable output: NEVER `&mut Tensor` — use raw `*mut E` or read-only `&Tensor` + +**HARD RULE (see [`no-mut-tensor-output.md`](no-mut-tensor-output.md)):** +a kernel output is NEVER declared `&mut Tensor` (nor `MappedPartitionMut` / +`Partition<&mut Tensor>` host arg). Write outputs ONE of two ways: + +**Preferred — read-only `&Tensor<{[-1,...]}>` + in-kernel `partition_full_mut`** +(worked example: `examples/bmm/kernel.rs` + `examples/bmm/ffi.rs`): +```rust +// kernel: c: &Tensor // OUTPUT — READ-only param type, -1 OK +// let mut cp = c.partition_full_mut(const_shape![1, BM, BN]); // mut VIEW in body +// store_view_tko_mut(&mut cp, tile, [q, mi, ni], ...); // explicit index +// host: c crosses as *const TensorDesc; borrow_tensor::(c_d) (ManuallyDrop, never freed); +// launch ANY grid (persistent OK). (partition_full_mut, NOT partition_mut — the +// latter re-offsets by block id.) +``` + +**Alternative — raw `*mut E` pointer + in-kernel `make_tensor_view`** (worked +example: `examples/softmax`): +```rust +// kernel: c_ptr: *mut E, c_d0,c_d1,c_d2: i32, c_s0,c_s1: i32 // pointer + dims + strides +// let cv = make_tensor_view(c_ptr, [batch,m,n], [c_s0,c_s1,1], token); +// let mut cp = cv.partition(const_shape![1,BM,BN]); store_view_tko_mut(&mut cp, tile, [..], ...); +// host: c crosses as *const TensorDesc; read dims/strides off it and wrap the pointer with +// DevicePointer::::from_cu_deviceptr(c_d.ptr); launch ANY grid (persistent OK). +``` + +**Why NOT `&mut Tensor`:** the launcher macro keys on `&` vs `&mut`. A `&mut` +param emits `KernelOutputStored::grid(&out)` → the launch grid is **LOCKED** to +the output's partition grid in descriptor dim order (full-tuple match), and the +shape is EXACT-matched so `{[-1,...]}` is rejected. That produces the two +recurring footguns — the `{[-1,...]}` reject trap AND the block-id↔dim +axis-transpose lock (output heads/rows unwritten / rc=-3). A raw +pointer or a read-only `&Tensor` emits NO grid → no lock → the kernel chooses any +grid / block-id convention (persistent, transposed, CGA) freely. +`-1` wildcards are valid on read-only `&Tensor` inputs (x, y, and `&Tensor` +outputs); they are forbidden ONLY because `&mut` would EXACT-match them — which +is why we drop `&mut` entirely. Persistent + CGA perf (`occupancy`, +`num_cta_in_cga`) is reachable on the raw-pointer path (see `examples/bmm`). + +### `order=(...)` loads and `partition_permuted` + +Represent a load order once: + +| Python / IR pattern | cutile-rs pattern | +|---------------------|-------------------| +| Load from original tensor with `order=(...)` | `tensor.partition_permuted(tile, const_array![...])` on the original logical tensor | +| Host already provides a physically swapped tensor | plain `tensor.partition(tile)` | + +Do not both pre-swap the host tensor and call `partition_permuted`. + +### `load_view_tko` + +```rust +pub fn load_view_tko< + E: ElementType, + const D: [i32; N], + O: ordering::LoadMode, + Sc: scope::Mode, + T: tma::Mode, +>( + view: &Partition, + index: [i32; N], + memory_ordering: O, + memory_scope: Sc, + latency: Option, + tma: T, +) -> Tile +``` + +Call: + +```rust +let tile = load_view_tko( + &part, [bid_m, bid_n], + ordering::Weak, scope::TileBlock, Some(3i32), tma::Enabled +); +``` + +### `store_view_tko_mut` + +```rust +pub unsafe fn store_view_tko_mut< + E: ElementType, + const D: [i32; N], + O: ordering::StoreMode, + Sc: scope::Mode, + T: tma::Mode, +>( + view: &mut PartitionMut, + tile: Tile, + index: [i32; N], + memory_ordering: O, + memory_scope: Sc, + latency: Option, + tma: T, +) -> Token +``` + +Call: + +```rust +unsafe { + store_view_tko_mut( + &mut part_mut, tile, [bid_m, bid_n], + ordering::Weak, scope::TileBlock, Some(1i32), tma::Enabled + ) +}; +``` + +### Missing APIs + +| Missing | Workaround | +|---------|------------| +| `make_strided_view` | Use `tview.partition(...)` plus tile-level index | +| legacy `load_from_view` | Use `load_view_tko` | +| legacy `store_to_view_mut` | Use `store_view_tko_mut` | +| `make_partition_view_padded` | `tview.partition(...)` auto-emits padding for loads | + +## Math + Arithmetic with mode args + +### `fma` + +```rust +let z = fma(a, b, c, rounding::NearestEven, ftz::Enabled); +``` + +### `addf` / `subf` / `mulf` + +```rust +let z = addf(a, b, rounding::NearestEven, ftz::Enabled); +``` + +Rust operators use default modes. Use named functions when the reference records +non-default rounding or FTZ. + +### `maxf` / `minf` + +```rust +let z = maxf(a, b, nan::Disabled, ftz::Enabled); +``` + +### `exp2`, `rsqrt`, `log`, `log2`, `true_div` + +```rust +let p = exp2(qk, ftz::Enabled); +let r = rsqrt(x, ftz::Enabled); +let l = log2(x); +let q = true_div(a, b); +``` + +For `divi signed rounding`, use `(a + b - 1) / b` or a scalar +ceil-div helper. + +## Type Conversion + +| IR Op | cutile-rs | +|-------|-----------| +| `exti signed` | `exti(x)` with return type annotation | +| `trunci` | `trunci(x, overflow::Wrap)` or matching overflow mode | +| `ftof` / `itof` / `ftoi` | `convert_tile(x)`; same-type conversions are DCE-elided | +| `ftof f32 -> tf32` | `let t: Tile = convert_tile(f32_tile);` | + +`tf32` is re-exported from `cutile::core::*`. + +## Tile Creation + +| IR Op | cutile-rs | +|-------|-----------| +| `constant` | `constant(0.0f32, const_shape![N])` for compile-time literals | +| runtime scalar tile | `broadcast_scalar(value, const_shape![N])` | +| `iota` | `iota(const_shape![N])` | +| `reshape` | `x.reshape(const_shape![...])` | +| `broadcast` | `x.broadcast(const_shape![...])` | +| `permute [1,0]` | `permute(x, const_array![1, 0])` | + +## `mmaf` + +```rust +pub fn mmaf( + lhs: Tile, + rhs: Tile, + acc: Tile, +) -> Tile +``` + +Mixed f16/bf16 input with f32 accumulate is native: + +```rust +let mut acc: Tile = constant(0.0f32, const_shape![BM, BN]); +acc = mmaf(a_f16_or_bf16, b_f16_or_bf16, acc); +``` + +For f32 GEMM-family inputs, check the f32 reference IR or cuTile-Python source. +If it has two `ftof ... -> ...xtf32` casts and `mmaf` operands are tf32, +reproduce that: + +```rust +let a_tc: Tile = convert_tile(a_f32); +let b_tc: Tile = convert_tile(b_f32); +acc = mmaf(a_tc, b_tc, acc); +``` + +Do not cast f16/bf16 tiles to tf32 unless the reference does. + +## Reduce + +| IR Op | cutile-rs | +|-------|-----------| +| `reduce dim=0 { addf }` | `reduce_sum(x, 0)` | +| `reduce dim=1 { maxf }` | `reduce_max(x, 1)` | +| `reduce dim=1 { addf }` | `reduce_sum(x, 1)` | + +For a 2D tile, reducing `dim=1` returns rank-1: + +```rust +let row_max: Tile = reduce_max(qk, 1); +let row_max_col: Tile = row_max.reshape(const_shape![M, 1]); +let qk_shifted = qk - row_max_col.broadcast(const_shape![M, N]); +``` + +## Entry launch hints: do not pin occupancy / num_cta_in_cga + +Never bake `occupancy=N` or `num_cta_in_cga=N` into +`#[cutile::entry(...)]` unless the reference IR genuinely carries that value. +Leave them empty so runtime/autotune can set resident-CTA policy through FFI +compile options. + +```rust +// WRONG unless the reference explicitly carries these exact hints: +#[cutile::entry(optimization_hints(occupancy = 1, num_cta_in_cga = 1))] + +// RIGHT when the reference leaves hints unset: +#[cutile::entry()] +``` + +Pinning `occupancy=1` can preserve correctness while creating a large perf cliff. +This is entry/host-side, not kernel math. + +## Control Flow + +| IR Op | cutile-rs | +|-------|-----------| +| `for %i in (0 to N, step S)` | `for i in (0i32..N).step_by(S as usize)` | +| `if/else with yield` | statement `if` with explicit mutable assignments | + +Prefer `for j in 0i32..num_tiles` when the step is 1. + +## Assume Hints + +```rust +pub unsafe fn assume_div_by(x: T) -> T; +pub unsafe fn assume_bounds_lower(x: T) -> T; +``` + +| IR Op | cutile-rs | +|-------|-----------| +| `assume div_by<16>` | `unsafe { assume_div_by::<_, 16>(x) }` | +| `assume bounded<0, ?>` | `unsafe { assume_bounds_lower::<_, 0>(x) }` | + +Apply at the top of the `#[cutile::entry()]` body before usage. + +## Verification provenance + +Cross-checked: + +- `cutile/src/_core.rs`: `get_tile_block_id`, token helpers, tensor view ops, + load/store view ops, math modes, conversion, `tf32`, `mmaf`, assume ops. +- `cutile/src/tensor.rs`: `PartitionMut`, `MappedLaunchPartition`, + partition-grid inference, and `KernelOutput` impls for partition/mapped + partition outputs. +- `cutile/src/tile_kernel.rs`: explicit grid validation against inferred grids. +- `cutile-macro/src/_module.rs`: host launcher generation uses `KernelInput` and + `KernelOutput` bounds. +- `cutile-compiler/src/compiler/_function.rs`: registry expects `#[entry(...)]`. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/output-structure.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/output-structure.md new file mode 100644 index 0000000..1b2b03f --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/output-structure.md @@ -0,0 +1,92 @@ +# Output Structure per Kernel + +All outputs now live in the tilegym repo under `src/tilegym/ops/cutile_rs/` +(kernel Rust + Python wrapper) and `src/tilegym/backend/cutile_rs/` (shared +runtime). There is no separate cutile-rs checkout and no `CUTILE_RS_ROOT`. + +## Kernel Rust: `src/tilegym/ops/cutile_rs/{kernel_name}_kernel/` + +```text +src/tilegym/ops/cutile_rs/{kernel_name}_kernel/ +├── kernel.rs # Agent B: SINGLE SOURCE OF TRUTH for kernel code +│ # Contains #[cutile::module] only. No test/FFI code. +│ # include!()d by the cutile_kernels crate. +├── ffi.rs # Agent D: C-ABI host launcher (one cutile_{op} symbol) +│ # Tensors cross as *const TensorDesc; borrow_tensor / +│ # DevicePointer::from_cu_deviceptr unpack them. +├── reference/ +│ ├── reference.mlir # Agent A: reference IR from winning backend (single-variant only) +│ ├── reference_{variant}.mlir # Agent A: per-variant IR (multi-variant only, no reference.mlir) +│ └── analysis.json # Agent A: structured analysis + test_perf_configs + reference_backend +├── generated/ +│ ├── generated.mlir # Agent B: cutile-rs compiled IR (single-variant only) +│ ├── generated_{variant}.mlir # Agent B: per-variant generated IR (multi-variant only) +│ └── diff_report.md # Agent C: structured IR diff analysis +└── reports/ + ├── correctness.md # Agent D: tilegym pytest test_op results + ├── performance.md # Agent E: tilegym pytest test_perf --print-record results + └── agent_logs/ + ├── agent_a.md + ├── agent_b.md # Includes compile attempts + tilegym wrapper design + ├── agent_c.md + ├── agent_d.md # pytest output + └── agent_e.md # pytest --print-record output +``` + +## Python wrapper + backend runtime: `src/tilegym/ops/cutile_rs/` + +```text +{tilegym_path}/src/tilegym/ops/cutile_rs/ +├── __init__.py # Auto-import all kernel wrapper modules +├── softmax.py # Agent D: @register_impl("softmax", backend="cutile-rs") +├── bmm.py # Agent D: @register_impl("bmm", backend="cutile-rs") +├── {kernel_name}.py # Agent D: @register_impl("{op}", backend="cutile-rs") +├── softmax_kernel/ # Agent B: kernel.rs + ffi.rs (see above) +├── bmm_kernel/ # Agent B: kernel.rs + ffi.rs +├── {kernel_name}_kernel/ # Agent B: kernel.rs + ffi.rs +└── cutile_kernels/ # aggregated cdylib crate (builds libcutile_kernels.so) + ├── Cargo.toml # crates.io deps pinned =0.2.0 (Rule 33) + └── src/ + ├── lib.rs # one `mod { include!(...) }` per op (see below) + └── ffi_util.rs # shared TensorDesc, borrow_tensor, dtype_str +``` + +The shared Python runtime lives in `src/tilegym/backend/cutile_rs/` +(`utils.py`: `bind_kernel_function_cffi`, `make_tensor_desc`, `check_rc`, +`get_num_sm`; `autotuner.py`: `autotune_launch`). Wrappers import from +`tilegym.backend.cutile_rs.{utils,autotuner}`. + +(Device/host split: the Python wrapper + backend wiring are Agent D's; +Agent B writes Rust only.) + +Agent D also modifies: +- `{tilegym_path}/src/tilegym/ops/cutile_rs/__init__.py` — imports the new wrapper module +- `{tilegym_path}/src/tilegym/backend/selector.py` — `is_cutile_rs_available()` +- `{tilegym_path}/tests/ops/test_{kernel_name}.py` — adds `"cutile-rs"` to `_backends` / `@parametrize` + +## cutile_kernels crate registration + +There is ONE aggregated cdylib crate for every op. Register a new op by adding a +`mod` to `cutile_kernels/src/lib.rs` that `include!`s the op's `kernel.rs` and +`ffi.rs` from the sibling `{kernel_name}_kernel/` dir: + +```rust +mod {kernel_name} { + include!("../../{kernel_name}_kernel/kernel.rs"); + include!("../../{kernel_name}_kernel/ffi.rs"); +} +``` + +Then build the single shared library: +```bash +cargo build --release # -> libcutile_kernels.so +``` + +The loader autobuilds this crate on first use (`CUTILE_RS_AUTOBUILD`, on by +default); override the crate location with `CUTILE_RS_KERNELS_DIR`. No +`cargo build -p cutile-ffi`, no per-op `.so`. + +## Validation + +Agent D: `pytest tests/ops/test_{kernel_name}.py -k cutile_rs -v` +Agent E: `pytest tests/ops/test_{kernel_name}.py -k cutile_rs --print-record -v` diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/performance-checklist.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/performance-checklist.md new file mode 100644 index 0000000..0028a03 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/performance-checklist.md @@ -0,0 +1,249 @@ +# Performance Checklist + +If a cutile-rs kernel is slower than cutile-python or perf is missing, check correctness and wrapper behavior first. A fast wrong kernel still scores poorly. + +## Mandatory Pre-Submit Checklist + +Record pass/fail evidence in `` when a validator or scorer asks for it. + +- [ ] **Cargo artifact exists**: `ops/cutile_rs/cutile_kernels/target/release/libcutile_kernels.so` exists and `nm -D` shows `T cutile_{kernel_name}`. +- [ ] **Pipeline compiled**: `cargo test --release -p cutile --test {kernel_name}_pipeline -- --nocapture` produced bytecode/MLIR. +- [ ] **Entry pattern matches reference IR**: view/TMA reference uses `&Tensor`; pointer reference uses `*mut E`. +- [ ] **Wrapper import works**: `tilegym.set_backend("cutile-rs")` and `from tilegym.ops.cutile_rs import {kernel_name}` succeed. +- [ ] **Fixed-config canary stays outside timed/autotune hot paths**: one `_run_ffi` call may be used before perf timing, but never inside `kernel_fn(cfg)` and never before every cached launch. +- [ ] **Autograd contract checked**: forward-only wrappers detach grad-enabled inputs and tests skip backward surfaces for `cutile-rs`. +- [ ] **No double transpose**: for each permuted tensor, either FFI swaps shape/strides or kernel uses `partition_permuted`, not both. +- [ ] **Persistent TMA latency**: persistent grid-stride loops pass `None` for every `load_view_tko` / `store_view_tko_mut` inside the loop. +- [ ] **Autotune lambda is clean**: `kernel_fn(cfg)` allocates outputs with `torch.empty` only. No clone/zeros/ones/extra contiguous copies inside timing. +- [ ] **TMA state matches reference**: do not disable TMA when reference has TMA descriptors. +- [ ] **CompileOptions parity (both directions)**: match analysis.json/best_config exactly. If the reference autotunes an EXPLICIT `num_ctas` (matmul ships {1,2,4}), the wrapper MUST forward that per-config value so the CGA cluster size takes effect (leaving it auto pins the cluster to 1 → ~1.6x GEMM regression — the #1 perf trap). Only when the reference truly uses `num_ctas=None`/auto do you leave `CompileOptions::default()` unmodified. Decide per-kernel, not by blanket default. + +## Correctness Before Perf + +Run: + +```bash +python -m pytest tests/ops/test_{kernel_name}.py -k cutile_rs -x --timeout=60 +``` + +If this fails, do not interpret geomean as a useful optimization signal. + +| Symptom | Check | +|---------|-------| +| `.so not found` | Cargo failed or wrong target dir | +| `backward not implemented` | Wrapper rejected `requires_grad`; detach forward-only inputs and skip backward tests | +| Forward close, gradients wrong | Backward surface should be skipped unless this task implements backward | +| Output near zero / 0% matched | Tensor shape/stride or transpose ownership | +| Fatal abort in autotune | Run fixed config outside CUPTI; inspect FFI args, mem::forget, persistent latency/layout | + +## CompileOptions Parity + +`CompileOptions` is part of the JIT cache key and changes codegen. The default value is meaningful: + +```rust +let opts = CompileOptions::default(); +``` + +In cutile-rs, `CompileOptions::default()` stores `occupancy=None` and `num_cta_in_cga=None`, which means compiler auto-pick. Calling `.occupancy(N)` / `.num_cta_in_cga(N)` pins those targets and is part of the JIT cache key. + +The rule is **match the reference, in BOTH directions**: +- Reference auto-picks (`num_ctas=None`) → keep `default()`, do not call the setters. Pinning `num_cta_in_cga(1)` "to be safe" caps resident CTAs and can regress decode shapes. +- Reference autotunes explicit `num_ctas` (matmul: {1,2,4}; many persistent/tensor-core GEMMs) → you MUST call `.num_cta_in_cga(value)` with the per-config value. A persistent kernel left at the auto default runs a 1-CTA cluster and loses ~1.6x vs a reference using 2/4-CTA clusters. The wrapper carries the value; ffi.rs applies it under the `if num_cta_in_cga > 0` guard. + +Reference mapping: + +| Reference evidence | cutile-rs FFI behavior | +|---|---| +| cuTile wrapper says `num_ctas=None`, `num_ctas = None`, or "let compiler choose" | keep `CompileOptions::default()`; do not call occupancy or num_cta setters | +| reference/autotune config has explicit occupancy or num_cta | pass that value and call the setter | +| Agent E/F finds CTA-count-sensitive slowdown with forced occupancy | re-test auto/default first before chasing tile math | +| reference has empty `optimization_hints=` only | treat as cosmetic unless it contains concrete keys | + +Recommended FFI sentinel pattern: + +```rust +let mut opts = CompileOptions::default(); +if occupancy > 0 { + opts = opts.occupancy(occupancy); +} +if num_cta_in_cga > 0 { + opts = opts.num_cta_in_cga(num_cta_in_cga); +} + +let op = kernel_call + .generics(generics) + .grid(grid) + .compile_options(opts); +``` + +Recommended Python wrapper pattern: + +```python +_AUTO_COMPILE_OPTION = -1 +_COMPILE_OCCUPANCY = None +_COMPILE_NUM_CTA_IN_CGA = None + +def _compile_option_value(value): + return _AUTO_COMPILE_OPTION if value is None else int(value) + +rc = lib.cutile_op( + ..., + _compile_option_value(getattr(cfg, "OCCUPANCY", _COMPILE_OCCUPANCY)), + _compile_option_value(getattr(cfg, "NUM_CTA_IN_CGA", _COMPILE_NUM_CTA_IN_CGA)), + ..., +) +``` + +Separate grid sizing from compile options. A persistent wrapper may need `_GRID_OCCUPANCY` to decide how many CTAs to launch, but that does not mean ffi.rs should call `.occupancy(_GRID_OCCUPANCY)`. + +## Launch Count Before IR Theory + +When Agent E reports `INVESTIGATE`, first classify whether the measured function launches extra GPU work. + +Use the cheapest reliable evidence: + +- `DUMP_CUPTI_EVENTS=1` from the tilegym perf harness, or +- `nsys profile` / `nsys stats`, or +- static wrapper proof when both wrappers directly launch exactly one kernel and allocate with `torch.empty`. + +Common wrapper/host causes: + +- fixed-config canary in the hot path before `autotune_launch` +- `torch.zeros`, `torch.ones`, `.clone()`, or `.contiguous()` inside `kernel_fn(cfg)` +- preallocating one output and allocating a second output inside `kernel_fn` +- launching both a canary and the cached best config on every call +- hidden PyTorch fallback or helper kernels + +If launch count is 1:1 and wrappers allocate only `torch.empty`, continue to config/IR. + +## TMA And Strides + +Check generated MLIR: + +```bash +grep -n "tma_descriptor\|load_view_tko\|store_view_tko\|strides=" generated/generated_canon.mlir +``` + +- View/TMA reference should show TMA-style view ops in generated IR too. +- Rank-R tensor stride metadata must have R entries and end in `1` for contiguous innermost layouts. +- If a batch/outer stride is not 16-byte aligned, fix host padding/fallback. Do not scalar-assert false divisibility. +- If kernel uses `partition_permuted`, the wrapper/FFI should pass original physical shapes/strides. + +## Accumulator Type + +For reductions and MMA: + +```bash +grep -n "mmaf\|reduce" generated/generated_canon.mlir +``` + +- `mmaf` accumulators should be f32. +- RMS and layer-norm reductions should accumulate in f32 and cast back at the end. +- If a kernel has `E2`, FFI should pass `"f32"` for that generic slot. + +## Persistent Kernels + +Persistent pattern: + +```rust +for tile_id in (bid_x..total_tiles).step_by(grid_x as usize) { + ... +} +``` + +Rules: + +- Use `None` latency for view loads/stores inside the persistent body when the known persistent/TMA exception applies. +- Keep `tma::Enabled` if reference uses TMA. +- Use a fixed-config canary before autotune only as a preflight, not in the measured or cached hot path. +- Grid capping must match the Rust body. If the kernel has no grid-stride loop, launch full coverage. + +## Autotune Search Space + +Start with reference `analysis.json` configs. Expand only after correctness passes. + +Inside `kernel_fn(cfg)`: + +```python +def kernel_fn(cfg): + y = torch.empty(output_shape, dtype=x.dtype, device=x.device) + _run_ffi(..., y, cfg) + return y +``` + +Avoid: + +- `clone()` +- `zeros()` / `ones()` +- `contiguous()` on large tensors inside the timed lambda +- reference PyTorch work inside `kernel_fn` +- canary launches inside `kernel_fn` + +Cache key should include every shape/dtype/semantic flag that can affect best config, including transpose flags and variant dispatch flags. If occupancy or num_cta is autotuned, include it through the config object and let the autotuner key include the same semantic shape/dtype fields as the reference. + +## Wrapper Autograd Perf Trap + +Perf tests often create inputs with `requires_grad=True` even if they only benchmark forward. A wrapper like this fails perf before timing: + +```python +if input.requires_grad: + raise NotImplementedError("backward not implemented") +``` + +Current skill policy is forward-only unless the task explicitly asks for backward kernels. Detach grad-enabled inputs before calling FFI and patch backward tests to skip `cutile-rs`. + +## Transpose Ownership + +Check both wrapper/FFI and kernel: + +- If FFI creates `Tensor` with swapped logical shape/strides, kernel uses identity `partition`. +- If kernel uses `partition_permuted`, FFI passes original physical shape/strides. +- If generated MLIR has both swapped strides and `dim_map`, expect wrong values even though IR compiles. + +## Math Hot Spots + +| Pattern | Preferred form | +|---------|----------------| +| f32 sigmoid division | `true_div(exp_x, exp_x_plus_one)` | +| FTZ reference math | named op with `ftz::Enabled` | +| f16/bf16 final output | compute in f32 then `convert_tile` | +| pointer bandwidth | `assume_div_by<16>` on pointers only | +| empty `sm_100={}` hint | do not chase unless concrete hint keys are present | +| reference auto occupancy | leave `CompileOptions::default()` auto | + +## Profiling Command + +Use after correctness passes: + +```python +from torch.profiler import ProfilerActivity, profile + +with profile(activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA]) as prof: + for _ in range(20): + run_kernel() + torch.cuda.synchronize() + +print(prof.key_averages().table(sort_by="cuda_time_total", row_limit=10)) +``` + +Confirm the launched kernel name and config before changing tile sizes. + +## Perf Result Interpretation + +| Result | Likely next step | +|--------|------------------| +| No geomean, correctness failed | Fix correctness/wrapper/build | +| Process abort | Fixed-config canary, FFI args, borrowed Tensor ownership | +| Near-uniform 2x slower | Extra timed launch or canary in hot path | +| 1.0x-1.1x slower with CTA-count outlier | Check CompileOptions occupancy/num_cta parity before IR hints | +| 1.1x-1.5x slower | Check TMA, clean autotune lambda, constants, accumulator, config parity | +| More than 2x slower | TMA disabled, runtime constants, extra timed copies, wrong variant | +| Faster but wrong | Ignore perf until correctness passes | + +## Agent F Handoff + +If Agent F writes `reports/perf_investigation_*.md`, reflection should treat it as the highest-priority perf signal. Agent C sees IR; Agent F checks wrapper launch count, CompileOptions, autotune cache behavior, and source-level config mismatches. If F says the fix is template-propagation, update `examples/softmax/wrapper.py` or the relevant wrapper pattern so next Agent B does not regenerate the same bug. + +## B-Only Reminder + +In B-only mode, no Agent E writes a report. The host runs perf pytest directly after Agent B. Agent B should still leave enough evidence in `reports/build_log.md` and `reports/agent_b.md` for the next reflection to identify the layer that failed. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/pipeline.md b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/pipeline.md new file mode 100644 index 0000000..6f5a4ab --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/references/pipeline.md @@ -0,0 +1,288 @@ +# Pipeline architecture, gates, and orchestrator rules + +This file holds the load-bearing details that the slim `SKILL.md` links to. Read it before spawning the first sub-agent so the run follows the bounded state machine and does not spend time on inline debugging or no-op respawns. + +## Contents + +- Architecture +- Agent responsibilities +- Orchestrator contract +- Mechanical validator acceptance gate +- Inter-agent communication +- Gate decisions +- Retry policy +- Compiler bug escalation +- The 19 CRITICAL rules + +## Architecture + +```text +Orchestrator + | + | preflight.sh + v +Agent A: IR dump, analysis.json, reference/tolerance, baseline perf + | + | PASS + v +Agent B(1): kernel.rs (DEVICE ONLY), kernel-only cargo build, in-Rust pipeline test, generated IR + | + | COMPILED + v +Agent D(1): build Python wrapper + backend wiring, THEN tilegym correctness + | + | ALL_PASS FAIL / BLOCKED (route by fail_class / block_class) + v v +Agent E(1): CUPTI perf fail_class/block_class: host → Agent D(2) (D owns ffi.rs/.so/wrapper) + | fail_class/block_class: kernel → Agent C → B(2) + | DONE + v +PIPELINE_COMPLETE + +Agent E(1) DONE is TERMINAL → PIPELINE_COMPLETE; never spawn C/D(2)/F after DONE +(geomean ≤ 1.05 is the authoritative gate; a (1.0,1.05] geomean or an info-only +slow per-config row is NOT grounds to diagnose). +Agent E(1) INVESTIGATE/BLOCKED routes to Agent C once. C tags owner: + owner: host → Agent D(2) (ffi.rs launcher/output-partition ABI/launch grid/autotune value/wrapper) + owner: kernel → Agent B(2) (kernel.rs device math/IR) +Then Agent D(2) -> Agent E(2) -> PIPELINE_COMPLETE. +Agent C PASS/PASS_WITH_NOTES does not respawn; it ends the run with reports for reflection. +``` + +The happy path is A -> B -> D -> E. Agent B writes Rust only; Agent D owns the +Python wrapper + backend wiring + correctness, so wrapper-caused mismatches and +wrapper-caused perf gaps route back to D, not B. Agent C is diagnostic-only and +tags each fixable finding `owner: host|kernel`. Agent F is optional after the +fix loop is spent. + +## Agent Responsibilities + +| Agent | Does | Does not do | +|---|---|---| +| A | Dump reference IR, analysis.json, tolerance, baseline perf | Write Rust kernel code | +| B | Write the device kernel.rs, compile kernel-only, pass in-Rust pipeline test, canonicalize generated IR, hand D a host_launch_contract | Write ffi.rs, build the .so, build the Python wrapper, register backend, run correctness/perf | +| D | Write ffi.rs (C-ABI launcher) + build the cdylib .so + Python wrapper + backend wiring, run tilegym correctness | Edit kernel.rs or the device module | +| C | Diff IR and perf-relevant attributes, classify differences, write action items + `owner: host\|kernel` | Edit code or run tests | +| E | Run CUPTI perf comparison and report ratios | Fix code | +| F | Optional residual perf diagnosis after fix-loop budget | Replace B/D/E | + +Each agent writes `reports/agent_{x}.md` and ends its return with a literal `` block followed immediately by one enumerated `VERDICT:` line. + +## Orchestrator Contract + +The orchestrator is a router, not a worker. It does exactly this: + +1. Run `scripts/preflight.sh`. +2. Spawn the next agent using the minimal pointer prompt from `SKILL.md`. +3. Check the returned `` block. +4. Route by the final `VERDICT:` enum and the tables below. +5. Run `scripts/validate_kernel.sh {kernel_name}` only after the pipeline reaches `PIPELINE_COMPLETE`. + +The orchestrator must not read IR files or wrappers for routing, edit source, run cargo, run pytest, run benchmarks, or diagnose root causes inline. If root cause is needed, spawn Agent C when the route allows it. + +## Mechanical Validator Acceptance Gate + +Expected final response shape from a sub-agent: + +```text + +$ bash .../validate_agent_x.sh kernel 2>&1 +...stdout+stderr... +exit=0 + +VERDICT: PASS +``` + +Acceptance algorithm: + +1. Extract the first enumerated `VERDICT:` line. +2. Search for a literal `` ... `` block. +3. If the block exists, parse every `exit=` or `exit_code=` inside it. +4. If at least one exit is parseable and all exits are zero, accept and route by verdict. +5. If the block is missing but the text has a clean success verdict marker (`PASS`, `COMPILED`, `ALL_PASS`, `DONE`) and none of `FAIL`, `ERROR`, `Traceback`, `panic`, `exit=1`, `exit_code=1`, `nonzero`, or `timeout`, accept with `SOFT_VALIDATOR_MISSING`. +6. Otherwise re-spawn only the same agent letter once for validator repair. The repair prompt must tell the agent to reuse existing artifacts and return the same semantic verdict plus a valid block. +7. If repair still fails, proceed with `SOFT_VALIDATOR_FAILED` unless the semantic verdict itself blocks. + +Hard locks: + +- Missing validator block plus clean success marker and no failure marker is accepted; do not respawn for formatting alone. +- Maximum one validator repair per agent letter per run. +- Never cascade validator repair across agents. +- A non-zero validator in an agent's own block repairs that same agent only. It does not authorize another agent respawn. +- The aggregate `validate_kernel.sh` gate is not part of Agent B's validator block. + +Minimum validator evidence by agent: + +| Agent | Required evidence in `` | +|---|---| +| A | `validate_ir.sh exit=0`, `validate_analysis.sh exit=0`, and perf-log validation when A writes perf logs | +| B | compile/canonicalization status and `validate_agent_b.sh exit=0` | +| C | `validate_diff_report.sh exit=0`; include `validate_ir.sh` and `diff_ir.sh` output in reports, not necessarily the final block | +| D | correctness pytest command and `validate_agent_d.sh exit=0`; failing pytest may be non-zero only with D `FAIL` | +| E | perf pytest command, `validate_perf_log.sh` for cutile-rs and baseline logs, and `validate_agent_e.sh exit=0` | +| F | residual investigation commands or artifact-validation exits | + +## Inter-Agent Communication + +Sub-agents communicate through files on disk. The orchestrator passes paths, not summaries. + +| Producer | Main outputs | Consumer | +|---|---|---| +| A | `reference/analysis.json`, `reference/*.mlir`, baseline perf logs | B, C, E | +| B | `kernel.rs`, kernel-only `Cargo.toml`, generated IR, build log, host_launch_contract | D, C | +| C | `generated/diff_report.md` (+ `owner:` line), `reports/agent_c.md` | B(2) on `owner: kernel`, D(2) on `owner: host`; reflection otherwise | +| D | `ffi.rs` + cdylib `.so` + Python wrapper + backend wiring, `reports/correctness.md` (+ `fail_class:`), correctness log | E, C, or D(2) | +| E | `reports/performance.md`, `perf_cutile_rs.txt`, baseline perf logs | C or reflection | +| F | `reports/perf_investigation_*.md` | reflection | + +Good next-agent prompt: `Kernel: softmax. Agent C verdict: FAIL_FIXABLE. Read action items from ${CUTILE_KERNEL_OUT_ROOT}/softmax/generated/diff_report.md`. + +Bad next-agent prompt: pasting C's whole analysis into B's spawn prompt. That causes context drift and route errors. + +## Gate Decisions + +### Agent A + +| Verdict | Action | +|---|---| +| `PASS` | Spawn Agent B | +| `FAIL_FIXABLE` / `BLOCKED` | STOP and record `SKILL_METHODOLOGY_BLOCKED.txt` | + +### Agent B(1) + +| Verdict | Action | +|---|---| +| `COMPILED` | Spawn Agent D | +| `FAIL_COMPILE` / `FAIL_FIXABLE` / `BLOCKED` | STOP | + +Agent B validates B-stage artifacts with `validate_agent_b.sh`. Do not ask Agent B to run `validate_kernel.sh`; final aggregate validation requires D/E outputs and belongs to the orchestrator after `PIPELINE_COMPLETE`. + +### Agent D(1) + +| Verdict | Action | +|---|---| +| `ALL_PASS` | Spawn Agent E | +| `FAIL` / `BLOCKED` (`fail_class`/`block_class: host`) | Spawn Agent D(2) — D owns the ffi.rs/.so/wrapper/wiring/autotune fix; read `host_fix:` | +| `FAIL` / `BLOCKED` (`fail_class`/`block_class: kernel`) | Spawn Agent C if C has not run — kernel/FFI fault; read `agent_b_followup:` | +| `BLOCKED` (`block_class: env`) | STOP (external) | + +D now owns the Python wrapper, so most non-kernel mismatches (0 tests collected, argtype/CUDA-700, dtype dispatch, layout) are D's own to fix via D(2). Only a provably kernel-side fault (missing `.so`/symbol, FFI signature gap, math wrong across correct wrapper variants, in-kernel panic) goes to C→B. If the class line is absent, fall back to C. + +### Agent E(1) + +| Verdict | Action | +|---|---| +| `DONE` | `PIPELINE_COMPLETE` — **TERMINAL. MUST NOT spawn C, D(2), or F.** | +| `INVESTIGATE` / `BLOCKED` | Spawn Agent C if C has not run and B has not been respawned; otherwise `PIPELINE_COMPLETE` | + +**`DONE` is binding and terminal — no discretionary diagnosis.** Agent E's +`VERDICT: DONE` means the perf gate already passed (geomean_ratio ≤ 1.05 is the +authoritative completion criterion). The orchestrator MUST route `DONE` straight +to `PIPELINE_COMPLETE` and end the run. Specifically you may NOT: +- spawn Agent C "to diagnose perf root cause" / "to confirm it's just noise" — C + is reachable ONLY from `INVESTIGATE`/`BLOCKED`, never from `DONE`; +- reinterpret the task goal (e.g. "≤ baseline") to override the verdict — a + geomean in `(1.0, 1.05]` is PASS, not a gap to chase; `DONE` already accounts + for that parity window; +- treat an info-only slow per-config row in `performance.md` as a reason to + investigate — per-row ratios are diagnostic only and do NOT gate the verdict. + +The routing table is not a menu of optional levers; for `DONE`, `PIPELINE_COMPLETE` +is the only legal action. Spawning any agent after `DONE` is a routing violation. + +Agent E does not use `FAIL` as a final verdict. Slow but valid perf is `INVESTIGATE` with severity and per-row ratios in `reports/performance.md`. + +### Agent C + +| Verdict | Action | +|---|---| +| `FAIL_FIXABLE` (`owner: kernel`) | Spawn Agent B(2) if B has not been respawned; otherwise `PIPELINE_COMPLETE` | +| `FAIL_FIXABLE` (`owner: host`) | Spawn Agent D(2) if D has not been respawned; otherwise `PIPELINE_COMPLETE` | +| `PASS` | `PIPELINE_COMPLETE` | +| `PASS_WITH_NOTES` | `PIPELINE_COMPLETE` | +| `FAIL_COMPILER_BUG` | `PIPELINE_COMPLETE` | +| `BLOCKED` | `PIPELINE_COMPLETE` | + +Agent C verdict semantics: + +- `FAIL_FIXABLE` means C found a concrete fixable action and tagged `owner:`. `owner: kernel` → B(2) (kernel.rs device math/IR); `owner: host` → D(2) (ffi.rs launcher, output-partition ABI, launch grid, cutile_kernels crate build, num_cta_in_cga/occupancy, cffi `_FFI_CDEF` / `TensorDesc` packing, dtype dispatch, layout, backend registration). Action items must be explicit in `generated/diff_report.md`. If `owner` is absent, default to `kernel`/B(2). +- `PASS` means no fixable owner was found. +- `PASS_WITH_NOTES` means remaining differences are known gaps, INFO-only IR deltas, semantic equivalents, or residual perf notes. These are not a reason for a respawn. +- `FAIL_COMPILER_BUG` and `BLOCKED` stop the in-run semantic loop. + +This distinction prevents no-op respawns and sends each fix to its true owner (Rust→B, Python/host→D). + +### Agent B(2), D(2), E(2) + +D(2) is reachable two ways: from a host-owned C `FAIL_FIXABLE`/D `fail_class: host` (D fixes its own ffi.rs/wrapper), or from B(2) after a kernel fix. Either way D's semantic-run cap is 2. + +| Stage | Verdict | Action | +|---|---|---| +| B(2) | `COMPILED` | Spawn D(2) | +| B(2) | any failure/block | STOP | +| D(2) | `ALL_PASS` | Spawn E(2) | +| D(2) | `FAIL` / `BLOCKED` | `PIPELINE_COMPLETE` | +| E(2) | `DONE` / `INVESTIGATE` / `BLOCKED` | `PIPELINE_COMPLETE` | + +After E(2), the fix-loop budget is spent. Do not spawn C again in the same run. + +### Agent F + +Agent F is optional and runs only after the one-shot fix loop has been spent. F may diagnose residual perf, but it does not open another in-run B/D/E loop. + +| F conclusion | Action | +|---|---| +| `FIXABLE` | `PIPELINE_COMPLETE` (action items recorded for reflection) | +| `ALIGNED` | `PIPELINE_COMPLETE` | +| `BLOCKED` | `PIPELINE_COMPLETE` | + +## Retry Policy + +| Phase | Max attempts | On exhaust | +|---|---:|---| +| A semantic run | 1 | STOP | +| B compile loop inside one B run | 3 cargo builds | B returns `FAIL_COMPILE` | +| B semantic runs | 2 total | STOP after B(2) failure | +| C semantic run | 1 | STOP / reflection mines report | +| D semantic runs | 2 total | STOP after D(2) failure | +| E semantic runs | 2 total | STOP with residual perf | +| F semantic run | 1 | STOP | +| Validator repair | 1 per agent letter | Proceed with soft marker unless semantic verdict blocks | + +Retry boundaries: + +- Semantic retries follow the gate table. +- Validator repair is local to the same agent letter. +- Missing validator block alone is not enough to respawn if the compatibility fallback accepts. +- Non-zero validator output paired with a success verdict is a same-agent repair event, not a reason for the orchestrator to diagnose, inspect IR, or rerun tests inline. + +## Compiler Bug Escalation + +1. Check `references/coding-rules.md`. +2. If a workaround exists, Agent C should report `FAIL_FIXABLE` and list the workaround for B. +3. If no workaround exists, Agent C reports `FAIL_COMPILER_BUG` and documents the limitation. +4. The orchestrator does not create a minimal repro inline. Reflection or a future Agent F can do that outside the current bounded route. + +Compiler-bug escalation still obeys the validator gate. + +## The 19 CRITICAL Rules + +1. **test_perf_configs only**: tests and benchmarks use tilegym test_perf shapes. Do not invent custom benchmark shapes. +2. **Identical workload**: cutile-rs and the reference backend run the exact same input and variant. +3. **Forward first**: implement forward kernels before backward variants. +4. **Logs mandatory**: every spawned agent writes `reports/agent_{x}.md` or `reports/agent_logs/agent_{x}.md`. +5. **IR diff is diagnostic, not automatically actionable**: Agent C drives B(2) only with `FAIL_FIXABLE`. `PASS` and `PASS_WITH_NOTES` end the current run with reports. +6. **Autotuning**: if any backend uses autotune for this kernel, cutile-rs uses the CUPTI-based autotuner (`from tilegym.backend.cutile_rs import autotune_launch`) with configs recorded by Agent A. +7. **dtype generic**: kernels use `` unless a documented side buffer requires concrete f32. +8. **Reference from tilegym**: correctness reference comes from the tilegym test class, not a hand-written formula. +9. **Tolerance from tilegym**: Agent A records the exact atol/rtol used by tilegym tests. +10. **Multi-variant ops**: Agent A records all variants and Agent B converts all variants; Agent E compares each variant against its matching baseline. +11. **No fallback**: the cutile-rs wrapper must call the cutile-rs FFI kernel or raise a clear unsupported error. It must not silently fall back to PyTorch or another backend. +12. **CUPTI only for perf**: performance conclusions use tilegym `--print-record` CUPTI timing, not `torch.cuda.Event`. +13. **ct.Constant -> const generic**: prefer const generics for Python `ct.Constant[int]` values unless lower-IR evidence shows constant unrolling blocks software pipelining. +14. **Happy path skips C**: A -> B -> D -> E is the normal route. C runs only after D/E non-DONE. B(2) runs only after C `FAIL_FIXABLE`. +15. **Dual-backend reference selection**: Agent A benchmarks cuTile-Python and Triton-TileIR when available and records the winning reference backend. +16. **Autotuner lambda uses torch.empty**: output allocation in perf lambdas must use `torch.empty`, not clone/zeros/ones, to avoid extra GPU kernels in CUPTI. +17. **Perf investigation uses the right layer**: C checks IR attributes first; residual post-loop gaps belong to reflection or optional Agent F. +18. **Const generic tradeoffs**: default to const generic, but switch to runtime only when IR and nsys evidence show const unrolling hurts the generated loop. +19. **Mechanical validator block before routing**: validate the literal `` block before semantic routing. Same-agent repair is capped once; never cascade repairs; final `validate_kernel.sh` runs only after the pipeline is complete. diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/diff_ir.sh b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/diff_ir.sh new file mode 100755 index 0000000..6dea91d --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/diff_ir.sh @@ -0,0 +1,419 @@ +#!/bin/bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 +# Compare two CUDA Tile IR files and emit a routing verdict. +# +# Usage: bash diff_ir.sh +# Exit 0 = PASS +# Exit 1 = FAIL_FIXABLE +# Exit 2 = PASS_WITH_NOTES + +set -uo pipefail + +REF="${1:?Usage: diff_ir.sh }" +GEN="${2:?Usage: diff_ir.sh }" + +[ -f "$REF" ] || { echo "ERROR: $REF not found"; exit 1; } +[ -f "$GEN" ] || { echo "ERROR: $GEN not found"; exit 1; } + +critical_fail=0 +workaround_count=0 +note_count=0 +notes="" + +grep_count() { + local pattern="$1" + local file="$2" + local n + n=$(grep -cE "$pattern" "$file" 2>/dev/null) + if [ $? -ne 0 ] || [ -z "$n" ]; then + echo 0 + else + echo "$n" + fi +} + +grep_count_fixed() { + local pattern="$1" + local file="$2" + local n + n=$(grep -cF "$pattern" "$file" 2>/dev/null) + if [ $? -ne 0 ] || [ -z "$n" ]; then + echo 0 + else + echo "$n" + fi +} + +first_pcre() { + local pattern="$1" + local file="$2" + local value + value=$(grep -oP "$pattern" "$file" 2>/dev/null | head -1) + if [ -z "$value" ]; then + echo 0 + else + echo "$value" + fi +} + +join_egrep() { + local pattern="$1" + local file="$2" + grep -oE "$pattern" "$file" 2>/dev/null | sort -u | tr '\n' ' ' +} + +join_pcre() { + local pattern="$1" + local file="$2" + grep -oP "$pattern" "$file" 2>/dev/null | sort -u | tr '\n' ' ' +} + +append_note() { + local msg="$1" + notes="${notes}\n${msg}" + note_count=$((note_count + 1)) +} + +append_warn() { + local msg="$1" + append_note "WARN: ${msg}" +} + +append_critical() { + local msg="$1" + critical_fail=$((critical_fail + 1)) + append_note "CRITICAL: ${msg}" +} + +in_array() { + local needle="$1" + shift + local item + for item in "$@"; do + [ "$needle" = "$item" ] && return 0 + done + return 1 +} + +count_ops() { + local file="$1" + + echo "load_ptr_tko $(grep_count_fixed 'load_ptr_tko' "$file")" + echo "store_ptr_tko $(grep_count_fixed 'store_ptr_tko' "$file")" + echo "load_view_tko $(grep_count_fixed 'load_view_tko' "$file")" + echo "store_view_tko $(grep_count_fixed 'store_view_tko' "$file")" + + echo "reduce $(grep_count '^[[:space:]]*(%[A-Za-z0-9_]+ = )?reduce[[:space:]]' "$file")" + echo "exp $(grep_count '(^|[[:space:]])exp([[:space:]]|$)' "$file")" + echo "exp2 $(grep_count '(^|[[:space:]])exp2([[:space:]]|$)' "$file")" + echo "rsqrt $(grep_count '(^|[[:space:]])rsqrt([[:space:]]|$)' "$file")" + echo "fma $(grep_count '(^|[[:space:]])fma([[:space:]]|$)' "$file")" + echo "mma $(grep_count '(^|[[:space:]])mmaf?([[:space:]]|$)' "$file")" + + echo "addf $(grep_count '(^|[[:space:]])addf([[:space:]]|$)' "$file")" + echo "subf $(grep_count '(^|[[:space:]])subf([[:space:]]|$)' "$file")" + echo "mulf $(grep_count '(^|[[:space:]])mulf([[:space:]]|$)' "$file")" + echo "divf $(grep_count '(^|[[:space:]])divf([[:space:]]|$)' "$file")" + echo "maxf $(grep_count '(^|[[:space:]])maxf([[:space:]]|$)' "$file")" + echo "minf $(grep_count '(^|[[:space:]])minf([[:space:]]|$)' "$file")" + + echo "for $(grep_count '^[[:space:]]*(%[A-Za-z0-9_]+ = )?for[[:space:]]' "$file")" + echo "select $(grep_count '(^|[[:space:]])select([[:space:]]|$)' "$file")" + echo "iota $(grep_count '(^|[[:space:]])iota([[:space:]]|$)' "$file")" + echo "cmpi $(grep_count '(^|[[:space:]])cmpi([[:space:]]|$)' "$file")" + echo "ftof $(grep_count '(^|[[:space:]])ftof([[:space:]]|$)' "$file")" + echo "itof $(grep_count '(^|[[:space:]])itof([[:space:]]|$)' "$file")" + echo "exti $(grep_count '(^|[[:space:]])exti([[:space:]]|$)' "$file")" + + echo "occupancy $(first_pcre 'occupancy = \K[0-9]+' "$file")" + echo "num_cta $(first_pcre 'num_cta_in_cga = \K[0-9]+' "$file")" +} + +declare -A REF_OPS +declare -A GEN_OPS + +while read -r op count; do + [ -n "${op:-}" ] || continue + REF_OPS[$op]="${count:-0}" +done < <(count_ops "$REF") + +while read -r op count; do + [ -n "${op:-}" ] || continue + GEN_OPS[$op]="${count:-0}" +done < <(count_ops "$GEN") + +CRITICAL_OPS=(load_ptr_tko store_ptr_tko load_view_tko store_view_tko reduce exp exp2 rsqrt fma mma for) +ARITH_OPS=(addf subf mulf divf maxf minf) +INFO_OPS=(select iota cmpi ftof itof exti) +HINT_OPS=(occupancy num_cta) + +echo "================================================================" +echo "IR Op Comparison" +echo "================================================================" +printf "%-20s %8s %8s %8s %s\n" "Op" "Ref" "Gen" "Delta" "Verdict" +echo "----------------------------------------------------------------" + +for op in "${CRITICAL_OPS[@]}" "${ARITH_OPS[@]}" "${INFO_OPS[@]}" "${HINT_OPS[@]}"; do + ref=${REF_OPS[$op]:-0} + gen=${GEN_OPS[$op]:-0} + delta=$((gen - ref)) + verdict="MATCH" + + if [ "$delta" -ne 0 ]; then + if in_array "$op" "${CRITICAL_OPS[@]}"; then + verdict="CRITICAL_DIFF" + append_critical "${op} ref=${ref} gen=${gen} delta=${delta}" + elif in_array "$op" "${HINT_OPS[@]}"; then + if [ "$ref" -ne 0 ] && [ "$gen" -ne 0 ] && [ "$ref" -ne "$gen" ]; then + verdict="CRITICAL_DIFF" + append_critical "${op} hint mismatch ref=${ref} gen=${gen}" + elif [ "$ref" -ne 0 ] && [ "$gen" -eq 0 ]; then + verdict="MISSING_HINT" + append_warn "${op} hint missing in generated" + else + verdict="OK" + fi + elif in_array "$op" "${ARITH_OPS[@]}"; then + abs_delta=${delta#-} + if [ "$abs_delta" -le 2 ]; then + verdict="WORKAROUND(${delta})" + workaround_count=$((workaround_count + abs_delta)) + else + verdict="CRITICAL_DIFF" + append_critical "${op} ref=${ref} gen=${gen} delta=${delta} exceeds workaround allowance" + fi + else + verdict="INFO(${delta})" + fi + fi + + printf "%-20s %8d %8d %+8d %s\n" "$op" "$ref" "$gen" "$delta" "$verdict" +done + +echo "" +echo "Tile shapes:" +ref_shapes=$(join_egrep 'tile<[0-9]+x[A-Za-z0-9_]+>' "$REF") +gen_shapes=$(join_egrep 'tile<[0-9]+x[A-Za-z0-9_]+>' "$GEN") +echo " Ref: ${ref_shapes:-none}" +echo " Gen: ${gen_shapes:-none}" + +if [ -n "$ref_shapes" ] && [ -n "$gen_shapes" ] && [ "$ref_shapes" != "$gen_shapes" ]; then + append_critical "tile shape set differs" +fi + +echo "" +echo "Reduce identities:" +ref_identities=$(grep -oE 'identities=\[[^]]+\]' "$REF" 2>/dev/null | sort -u) +gen_identities=$(grep -oE 'identities=\[[^]]+\]' "$GEN" 2>/dev/null | sort -u) +if [ -n "$ref_identities" ]; then + echo "$ref_identities" | sed 's/^/ Ref: /' +else + echo " Ref: none" +fi +if [ -n "$gen_identities" ]; then + echo "$gen_identities" | sed 's/^/ Gen: /' +else + echo " Gen: none" +fi +if [ "$ref_identities" != "$gen_identities" ]; then + append_critical "reduce identities differ" +fi + +echo "" +ref_rounding=$(grep_count_fixed "rounding_mode" "$REF") +gen_rounding=$(grep_count_fixed "rounding_mode" "$GEN") +echo "rounding_mode: ref=${ref_rounding} gen=${gen_rounding}" +if [ "$gen_rounding" -gt "$ref_rounding" ]; then + append_warn "generated has extra rounding_mode attrs" +fi + +echo "" +echo "Block mapping pattern:" +ref_has_swizzle=false +gen_has_swizzle=false +if grep -qE 'divi.*blockId|remi.*blockId' "$REF" 2>/dev/null; then ref_has_swizzle=true; fi +if grep -qE 'divi.*blockId|remi.*blockId' "$GEN" 2>/dev/null; then gen_has_swizzle=true; fi +ref_divi=$(grep_count '(^|[[:space:]])divi([[:space:]]|$)' "$REF") +gen_divi=$(grep_count '(^|[[:space:]])divi([[:space:]]|$)' "$GEN") +ref_remi=$(grep_count '(^|[[:space:]])remi([[:space:]]|$)' "$REF") +gen_remi=$(grep_count '(^|[[:space:]])remi([[:space:]]|$)' "$GEN") +echo " Ref: swizzle=${ref_has_swizzle} divi=${ref_divi} remi=${ref_remi}" +echo " Gen: swizzle=${gen_has_swizzle} divi=${gen_divi} remi=${gen_remi}" +if [ "$ref_has_swizzle" = true ] && [ "$gen_has_swizzle" = false ]; then + append_warn "missing block swizzle; possible L2 locality regression" +fi + +echo "" +echo "================================================================" +echo "Op Attribute Comparison" +echo "================================================================" + +echo "" +echo "assume_div_by:" +echo " Reference:" +ref_assume_lines=$(grep -n 'assume div_by' "$REF" 2>/dev/null) +if [ -n "$ref_assume_lines" ]; then + echo "$ref_assume_lines" | sed 's/^/ /' +else + echo " (none)" +fi +echo " Generated:" +gen_assume_lines=$(grep -n 'assume div_by' "$GEN" 2>/dev/null) +if [ -n "$gen_assume_lines" ]; then + echo "$gen_assume_lines" | sed 's/^/ /' +else + echo " (none)" +fi + +ref_assume_ptr=$(grep_count 'assume div_by.*ptr' "$REF") +gen_assume_ptr=$(grep_count 'assume div_by.*ptr' "$GEN") +ref_assume_scalar=$(grep_count 'assume div_by.*tile' "$REF") +gen_assume_scalar=$(grep_count 'assume div_by.*tile' "$GEN") +gen_raw_ptr_ops=$(( ${GEN_OPS[load_ptr_tko]:-0} + ${GEN_OPS[store_ptr_tko]:-0} )) + +echo " Summary: Ref ptr=${ref_assume_ptr} scalar=${ref_assume_scalar} | Gen ptr=${gen_assume_ptr} scalar=${gen_assume_scalar}" + +if [ "$ref_assume_ptr" -gt "$gen_assume_ptr" ]; then + if [ "$gen_raw_ptr_ops" -gt 0 ]; then + echo " >> PERF_CRITICAL: generated raw-pointer path is missing pointer assume_div_by." + append_critical "missing pointer assume_div_by coverage on raw-pointer generated path ref=${ref_assume_ptr} gen=${gen_assume_ptr}" + else + echo " >> PERF_WARN: generated view/TMA path has fewer pointer assume_div_by ops." + echo " Policy: this is non-blocking unless a measured perf regression points here." + append_warn "missing pointer assume_div_by coverage ref=${ref_assume_ptr} gen=${gen_assume_ptr}; generated has no raw pointer load/store ops" + fi +fi + +if [ "$gen_assume_scalar" -gt 0 ]; then + echo " >> CRITICAL: generated scalar assume_div_by is deprecated and can corrupt irregular shapes." + append_critical "generated scalar assume_div_by present; remove scalar assume hints" +fi + +echo "" +echo "Per-op attributes (flush_to_zero, rounding):" +echo " Reference ops with attributes:" +ref_attr_lines=$(grep -nE '(addf|subf|mulf|divf|exp2?|negf).*flush_to_zero|rounding' "$REF" 2>/dev/null) +if [ -n "$ref_attr_lines" ]; then + echo "$ref_attr_lines" | sed 's/^/ /' +else + echo " (none)" +fi +echo " Generated ops with attributes:" +gen_attr_lines=$(grep -nE '(addf|subf|mulf|divf|exp2?|negf).*flush_to_zero|rounding' "$GEN" 2>/dev/null) +if [ -n "$gen_attr_lines" ]; then + echo "$gen_attr_lines" | sed 's/^/ /' +else + echo " (none)" +fi + +for op_type in addf subf mulf divf exp exp2 negf; do + ref_ftz_op=$(grep_count "${op_type}.*flush_to_zero" "$REF") + gen_ftz_op=$(grep_count "${op_type}.*flush_to_zero" "$GEN") + ref_approx_op=$(grep_count "${op_type}.*rounding" "$REF") + gen_approx_op=$(grep_count "${op_type}.*rounding" "$GEN") + + if [ "$ref_ftz_op" -ne "$gen_ftz_op" ] || [ "$ref_approx_op" -ne "$gen_approx_op" ]; then + echo " ${op_type}: FTZ ref=${ref_ftz_op} gen=${gen_ftz_op} | approx ref=${ref_approx_op} gen=${gen_approx_op}" + if [ "$ref_ftz_op" -gt "$gen_ftz_op" ]; then + append_critical "${op_type} missing flush_to_zero ref=${ref_ftz_op} gen=${gen_ftz_op}" + fi + if [ "$ref_approx_op" -gt "$gen_approx_op" ]; then + append_critical "${op_type} missing rounding ref=${ref_approx_op} gen=${gen_approx_op}" + fi + fi +done + +echo "" +echo "Load/store optimization hints:" +echo " Reference:" +ref_hint_lines=$(grep -n 'optimization_hints.*=' "$REF" 2>/dev/null | grep -v 'entry @') +if [ -n "$ref_hint_lines" ]; then + echo "$ref_hint_lines" | sed 's/^/ /' +else + echo " (none)" +fi +echo " Generated:" +gen_hint_lines=$(grep -n 'optimization_hints.*=' "$GEN" 2>/dev/null | grep -v 'entry @') +if [ -n "$gen_hint_lines" ]; then + echo "$gen_hint_lines" | sed 's/^/ /' +else + echo " (none)" +fi + +ref_allow_tma_true=$(grep_count 'allow_tma *= *true' "$REF") +gen_allow_tma_true=$(grep_count 'allow_tma *= *true' "$GEN") +ref_allow_tma_false=$(grep_count 'allow_tma *= *false' "$REF") +gen_allow_tma_false=$(grep_count 'allow_tma *= *false' "$GEN") +ref_latency=$(grep_count 'latency *=' "$REF") +gen_latency=$(grep_count 'latency *=' "$GEN") + +echo " Summary: Ref allow_tma_true=${ref_allow_tma_true} allow_tma_false=${ref_allow_tma_false} latency=${ref_latency} | Gen allow_tma_true=${gen_allow_tma_true} allow_tma_false=${gen_allow_tma_false} latency=${gen_latency}" + +if [ "$ref_allow_tma_true" -gt "$gen_allow_tma_true" ]; then + append_critical "allow_tma=true count mismatch ref=${ref_allow_tma_true} gen=${gen_allow_tma_true}" +fi +if [ "$ref_allow_tma_false" -ne "$gen_allow_tma_false" ]; then + append_warn "allow_tma=false count mismatch ref=${ref_allow_tma_false} gen=${gen_allow_tma_false}" +fi +if [ "$ref_latency" -gt "$gen_latency" ]; then + append_warn "latency hint count mismatch ref=${ref_latency} gen=${gen_latency}" +elif [ "$gen_latency" -gt "$ref_latency" ]; then + append_warn "generated extra latency hints ref=${ref_latency} gen=${gen_latency}" +fi + +echo "" +echo "MMA accumulator types:" +ref_mma_acc=$(join_pcre 'mmaf?.*tile<[0-9]+x[0-9]+x\w+>' "$REF" | tr ' ' '\n' | grep -oP 'tile<[0-9]+x[0-9]+x\w+>$' 2>/dev/null | sort -u | tr '\n' ' ') +gen_mma_acc=$(join_pcre 'mmaf?.*tile<[0-9]+x[0-9]+x\w+>' "$GEN" | tr ' ' '\n' | grep -oP 'tile<[0-9]+x[0-9]+x\w+>$' 2>/dev/null | sort -u | tr '\n' ' ') +echo " Ref: ${ref_mma_acc:-none}" +echo " Gen: ${gen_mma_acc:-none}" +if [ -n "$ref_mma_acc" ] && [ -n "$gen_mma_acc" ] && [ "$ref_mma_acc" != "$gen_mma_acc" ]; then + append_critical "MMA accumulator ref=${ref_mma_acc} gen=${gen_mma_acc}" +fi + +echo "" +ref_has_bf16=$(grep_count_fixed 'bf16' "$REF") +gen_has_bf16=$(grep_count_fixed 'bf16' "$GEN") +ref_has_f16=$(grep_count_fixed 'f16' "$REF") +gen_has_f16=$(grep_count_fixed 'f16' "$GEN") +echo "dtype usage: ref(f16=${ref_has_f16}, bf16=${ref_has_bf16}) gen(f16=${gen_has_f16}, bf16=${gen_has_bf16})" +if [ "$ref_has_bf16" -gt 0 ] && [ "$gen_has_bf16" -eq 0 ]; then + append_warn "reference uses bf16 but generated does not" +fi + +echo "" +echo "Tensor view layout attributes:" +ref_layout=$(grep -oE 'strides=\[[^]]+\]|dim_map=\[[^]]+\]|tile=\([^)]+\)' "$REF" 2>/dev/null | sort | uniq -c) +gen_layout=$(grep -oE 'strides=\[[^]]+\]|dim_map=\[[^]]+\]|tile=\([^)]+\)' "$GEN" 2>/dev/null | sort | uniq -c) +echo " Reference:" +if [ -n "$ref_layout" ]; then + echo "$ref_layout" | sed 's/^/ /' +else + echo " (none)" +fi +echo " Generated:" +if [ -n "$gen_layout" ]; then + echo "$gen_layout" | sed 's/^/ /' +else + echo " (none)" +fi +if [ "$ref_layout" != "$gen_layout" ]; then + append_warn "layout attribute multiset differs; inspect dim_map/strides/tile shapes" +fi + +echo "" +echo "================================================================" +if [ "$critical_fail" -gt 0 ]; then + echo "VERDICT: FAIL_FIXABLE (${critical_fail} critical diff(s))" + [ -n "$notes" ] && printf "%b\n" "$notes" + exit 1 +elif [ "$workaround_count" -gt 0 ] || [ "$note_count" -gt 0 ]; then + echo "VERDICT: PASS_WITH_NOTES (${workaround_count} workaround op(s), ${note_count} note(s))" + [ -n "$notes" ] && printf "%b\n" "$notes" + exit 2 +else + echo "VERDICT: PASS" + exit 0 +fi diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/preflight.sh b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/preflight.sh new file mode 100755 index 0000000..e8dad42 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/preflight.sh @@ -0,0 +1,154 @@ +#!/bin/bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +# Hard gate: refuses to proceed with the tilegym-converting-cutile-triton-to-cutile-rs skill until +# 1 git-tracked repo (TILEGYM_PATH — holds tests, Python wrappers, AND the +# cutile_rs Rust under src/tilegym/ops/cutile_rs/) and +# 4 cuTile toolchain paths (TILEIRAS_BIN, CUDA_TILE_OPT_BIN, +# TRITON_TILEIR_PYTHONPATH, CUDA_TOOLKIT_PATH) are set and verified on disk. +# There is NO separate cutile-rs checkout and NO CUTILE_RS_ROOT anymore — the +# aggregated cutile_kernels crate (crates.io PINNED deps, one libcutile_kernels.so) +# lives inside the tilegym tree. +# Prints a STOP banner with the exact next action when the gate fails; +# the banner ends up in the orchestrator context, making the gate +# behaviorally enforceable rather than text-only. +# +# Required env vars (5 total): +# export TILEGYM_PATH=/abs/path/to/tilegym +# export TILEIRAS_BIN=/abs/path/to/tileiras +# export CUDA_TILE_OPT_BIN=/abs/path/to/cuda-tile-opt +# export TRITON_TILEIR_PYTHONPATH=/abs/path/to/Triton-TileIR/python +# export CUDA_TOOLKIT_PATH=/abs/path/to/cuda-toolkit +# +# Optional (defaults shown): the loader autobuilds the cutile_kernels crate on +# first use; override with: +# export CUTILE_RS_AUTOBUILD=1 # on by default; 0 = use prebuilt +# export CUTILE_RS_KERNELS_DIR=$TILEGYM_PATH/src/tilegym/ops/cutile_rs/cutile_kernels +# +# Usage: +# bash scripts/preflight.sh +# +# Returns: +# 0 = paths set + verified → safe to read the rest of SKILL.md +# 1 = paths missing/invalid → AskUserQuestion + re-export + re-run + +set -uo pipefail + +red=$'\033[31m'; grn=$'\033[32m'; ylw=$'\033[33m'; bld=$'\033[1m'; rst=$'\033[0m' + +banner_stop() { + printf '%s' "${red}${bld}" + echo "===============================================================" + echo " STOP -- tilegym-converting-cutile-triton-to-cutile-rs preflight FAILED" + echo "===============================================================" + printf '%s' "${rst}" +} + +banner_ok() { + printf '%s' "${grn}${bld}" + echo "===============================================================" + echo " OK -- tilegym-converting-cutile-triton-to-cutile-rs preflight PASSED" + echo "===============================================================" + printf '%s' "${rst}" +} + +check_dir() { + local name="$1" val="${2:-}" desc="$3" + if [ -z "$val" ]; then + echo " ${red}MISSING${rst} $name ($desc)" + return 1 + fi + if [ ! -d "$val" ]; then + echo " ${red}NOT A DIR${rst} $name=$val" + echo " ($desc)" + return 1 + fi + echo " ${grn}OK${rst} $name=$val" + return 0 +} + +check_executable() { + local name="$1" val="${2:-}" desc="$3" + if [ -z "$val" ]; then + echo " ${red}MISSING${rst} $name ($desc)" + return 1 + fi + if [ ! -x "$val" ]; then + echo " ${red}NOT EXECUTABLE${rst} $name=$val" + echo " ($desc)" + return 1 + fi + echo " ${grn}OK${rst} $name=$val" + return 0 +} + +fail=0 +echo +echo "Gate — paths (1 git repo + 4 toolchain paths):" +echo +echo " ${bld}Git repo${rst}:" +check_dir "TILEGYM_PATH" "${TILEGYM_PATH:-}" "tilegym checkout (tests + Python wrappers + cutile_rs Rust under src/tilegym/ops/cutile_rs/)" || fail=1 +echo +echo " ${bld}Toolchain${rst} (existence-checked):" +check_executable "TILEIRAS_BIN" "${TILEIRAS_BIN:-}" "cuTile compiler binary (tileiras)" || fail=1 +check_executable "CUDA_TILE_OPT_BIN" "${CUDA_TILE_OPT_BIN:-}" "MLIR canonicalizer (cuda-tile-opt) — used by Agent B for IR diff" || fail=1 +check_dir "TRITON_TILEIR_PYTHONPATH" "${TRITON_TILEIR_PYTHONPATH:-}" "Triton-TileIR Python bindings dir (TileIR backend) — required for Triton-TileIR IR dump + Triton-TileIR correctness tests" || fail=1 +check_dir "CUDA_TOOLKIT_PATH" "${CUDA_TOOLKIT_PATH:-}" "CUDA toolkit root (cuda-bindings/build.rs requires this; e.g. /usr/local/cuda)" || fail=1 +echo +echo " ${bld}Per-kernel working dir${rst} (Agent A/B/C/D write here; eval harness depends on it):" +if [ -z "${CUTILE_KERNEL_OUT_ROOT:-}" ]; then + echo " ${red}MISSING${rst} CUTILE_KERNEL_OUT_ROOT (export to /workspace/cutile_kernel_out or similar)" + fail=1 +else + mkdir -p "$CUTILE_KERNEL_OUT_ROOT" 2>/dev/null || true + check_dir "CUTILE_KERNEL_OUT_ROOT" "$CUTILE_KERNEL_OUT_ROOT" "Per-kernel output root for agent artifacts (kernel.rs / ffi.rs / reference/ / generated/ / reports/). Kernel Rust is wired into the aggregated cutile_kernels crate under tilegym; there is no per-kernel Cargo.toml." || fail=1 +fi +echo + +if [ "$fail" -ne 0 ]; then + banner_stop + echo + echo "${bld}MANDATORY NEXT ACTION${rst} (orchestrator):" + echo + echo " 1. Call ${bld}AskUserQuestion${rst} with one question per missing path." + echo " Do NOT batch-guess. Do NOT default from cwd / git / memory." + echo + echo " 2. After the user answers, echo all paths back so they can correct typos." + echo + echo " 3. Export them in the shell that will run subsequent tool calls:" + echo " export TILEGYM_PATH=/abs/path/to/tilegym" + echo " export TILEIRAS_BIN=/abs/path/to/tileiras" + echo " export CUDA_TILE_OPT_BIN=/abs/path/to/cuda-tile-opt" + echo " export TRITON_TILEIR_PYTHONPATH=/abs/path/to/Triton-TileIR/python" + echo " export CUDA_TOOLKIT_PATH=/abs/path/to/cuda-toolkit" + echo + echo " 4. Re-run THIS script. It must exit 0 before any other tool call:" + echo " bash ${BASH_SOURCE[0]}" + echo + echo "${ylw}HARD RULES${rst} (violations = task failure):" + echo " - Do NOT call Read / Edit / Write / Bash / Agent for skill work" + echo " until this script exits 0." + echo " - Do NOT 'just look at the test file' to estimate scope --" + echo " that is the 'minimal edit' anti-pattern explicitly banned by" + echo " SKILL.md line 30." + echo " - Do NOT assume the cwd repo is TILEGYM_PATH. Ask the user." + echo + banner_stop + exit 1 +fi + +echo " ${grn}Gate passed${rst} — all 5 paths verified." +echo " TILEGYM_PATH = $TILEGYM_PATH" +echo " TILEIRAS_BIN = $TILEIRAS_BIN" +echo " CUDA_TILE_OPT_BIN = $CUDA_TILE_OPT_BIN" +echo " TRITON_TILEIR_PYTHONPATH = $TRITON_TILEIR_PYTHONPATH" +echo " CUDA_TOOLKIT_PATH = $CUDA_TOOLKIT_PATH" +echo + +banner_ok +echo +echo " All gates passed. You may now read the rest of SKILL.md and spawn Agent A." +echo +exit 0 diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_b.sh b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_b.sh new file mode 100755 index 0000000..8a1c8f8 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_b.sh @@ -0,0 +1,170 @@ +#!/bin/bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +# Validate Agent B outputs: kernel.rs, ffi.rs, generated IR, dtype generic, compile test. + +set -uo pipefail +KERNEL="${1:?Usage: validate_agent_b.sh }" + +# Per-kernel working dir (eval harness depends on it). There is no separate +# cutile-rs checkout / CUTILE_RS_ROOT anymore — the aggregated cutile_kernels +# crate lives under $TILEGYM_PATH/src/tilegym/ops/cutile_rs/. +: "${CUTILE_KERNEL_OUT_ROOT:?CUTILE_KERNEL_OUT_ROOT must be set (e.g. $ISOLATED_CWD/cutile_kernel_out)}" + +BASE="${CUTILE_KERNEL_OUT_ROOT}/${KERNEL}" +fail=0 + +check() { if [ -f "$1" ] && [ -s "$1" ]; then echo "OK: $2"; else echo "FAIL: $2 → $1"; fail=$((fail+1)); fi; } + +check "${BASE}/kernel.rs" "kernel.rs" +# ffi.rs is NOT a B deliverable (device/host split): Agent D writes +# ffi.rs + wires the op into the aggregated cutile_kernels crate (builds the +# single libcutile_kernels.so), because the host launch path (output partition +# ABI, launch-grid validation, borrow_tensor / DevicePointer ownership) is only +# testable by D's pytest. B writes the device kernel + proves it with the +# in-Rust pipeline test. +# Generated IR: single-variant → generated.mlir, multi-variant → generated_{variant}.mlir +HAS_MULTI=$(python3 -c " +import json, os +f = '${BASE}/reference/analysis.json' +if not os.path.exists(f): print('single'); exit() +d = json.load(open(f)) +v = d.get('kernel_variants', []) +print('multi' if isinstance(v, list) and len(v) > 1 else 'single') +" 2>/dev/null) +if [ "$HAS_MULTI" = "multi" ]; then + gen_count=$(ls "${BASE}/generated/generated_"*.mlir 2>/dev/null | wc -l) + if [ "$gen_count" -gt 0 ]; then + echo "OK: ${gen_count} per-variant generated IR files" + else + echo "FAIL: multi-variant but no generated_{variant}.mlir files" + fail=$((fail+1)) + fi +else + check "${BASE}/generated/generated.mlir" "generated.mlir" +fi + +# dtype generic +if [ -f "${BASE}/kernel.rs" ]; then + if grep -q "E: ElementType" "${BASE}/kernel.rs"; then + echo "OK: " + else + echo "FAIL: kernel.rs missing (Rule 16)" + fail=$((fail+1)) + fi +fi + +# (Rule 12): entry pattern must match reference IR +# TMA-style IR (load_view_tko / store_view_tko_mut) → entry uses &Tensor +# pointer-scatter IR (load_ptr_tko / store_ptr_tko) → entry uses *mut E +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +if [ -x "${SCRIPT_DIR}/validate_entry_pattern.sh" ] && [ -f "${BASE}/kernel.rs" ]; then + if bash "${SCRIPT_DIR}/validate_entry_pattern.sh" "${KERNEL}" >&2; then + echo "OK: entry pattern matches reference IR (Rule 12)" + else + echo "FAIL: entry pattern mismatch — see above (Rule 12)" + fail=$((fail+1)) + fi +fi + +# (multi-output transpose trap): a `&mut Tensor` OUTPUT param MUST carry +# CONCRETE tile entry shapes (e.g. {[1, TILE_H, TILE_D]}, {[TM, TN]}), NEVER -1 +# wildcards. cutile-rs's launcher infers the output-partition grid from the +# param's tile shape; a -1 on a `&mut Tensor` output makes the host +# `(&mut t).partition([...])` fail with "partition shape mismatch. Expected +# [-1,...], got [...]" → EVERY config FAILs (the multi-output transpose trap, D fail_class: +# kernel). `-1` wildcards are valid ONLY on read-only `&Tensor` inputs. +if [ -f "${BASE}/kernel.rs" ]; then + if grep -nE '&mut[[:space:]]+Tensor<[^;{]*\{[[:space:]]*\[[^]]*-1' "${BASE}/kernel.rs" >&2; then + echo "FAIL: mutable output '&mut Tensor<.. {[..-1..]}>' uses a -1 wildcard shape (Rule 12 — transpose trap)." + echo " Mutable outputs MUST declare concrete tile entry shapes (e.g. {[1, TILE_H, TILE_D]});" + echo " cutile-rs rejects -1-shaped mutable outputs (output-partition grid cannot be inferred)." + fail=$((fail+1)) + else + echo "OK: no -1-wildcard mutable-output Tensor (multi-output transpose trap guard)" + fi +fi + +# agent log +check "${BASE}/reports/agent_logs/agent_b.md" "agent_b.md log" + +# build_log.md mandatory + must contain CODE SNIPPETS for each failed attempt +# (not just error code/message — the actual source lines that triggered the error). +BUILD_LOG="${BASE}/reports/build_log.md" +check "${BUILD_LOG}" "build_log.md" +if [ -f "$BUILD_LOG" ] && [ -s "$BUILD_LOG" ]; then + # Must mention at least one attempt + if ! grep -qiE "^##? +Attempt|^Attempt [0-9]+" "$BUILD_LOG"; then + echo "FAIL: build_log.md missing '## Attempt N' section header" + fail=$((fail+1)) + fi + # If any attempt FAILED, that attempt MUST include a fenced rust code block + # showing the source snippet that triggered the error (not just error text). + fail_count=$(grep -ciE "Result:[[:space:]]*FAIL|FAIL[[:space:]]*\(" "$BUILD_LOG" || true) + if [ "$fail_count" -gt 0 ]; then + FENCE_RUST='```rust' + snippet_count=$(grep -cF "$FENCE_RUST" "$BUILD_LOG" || true) + if [ "$snippet_count" -lt 1 ]; then + echo "FAIL: build_log.md has $fail_count failed attempt(s) but no fenced rust code block (\`\`\`rust)" + echo " Agent B must paste the actual kernel.rs / ffi.rs lines that triggered each rustc error" + fail=$((fail+1)) + fi + fi +fi + +# Host-side checks (ffi.rs, cutile_kernels build, wrapper, _FFI_CDEF, backend +# registration) are Agent D's, validated by validate_agent_d.sh — see the +# device/host split note near the kernel.rs check above. + +# ───────────────────────────────────────────────────────────────── +# Agent B LANE GUARD (device/host split): B writes the device kernel +# only. ffi.rs, the Python wrapper, backend registration, and test parametrization +# are Agent D's. validate_agent_b.sh runs at the END of B's turn, BEFORE Agent D, +# so NONE of D's artifacts should exist yet. If they do, B overreached → hard FAIL. +# ───────────────────────────────────────────────────────────────── +# 0. ffi.rs is now Agent D's (host launch boundary). B must not write it. +if [ -f "${BASE}/ffi.rs" ]; then + echo "FAIL: Agent B wrote ffi.rs (${BASE}/ffi.rs)." + echo " ffi.rs (C-ABI host launcher) is Agent D's deliverable now — D writes it," + echo " wires it into the cutile_kernels crate (libcutile_kernels.so), and tests the" + echo " launch path via pytest. B writes kernel.rs only." + echo " Remove ffi.rs; keep the pipeline test (kernel-only) for the IR dump." + fail=$((fail+1)) +fi + +TILE="${TILEGYM_PATH:-}" +if [ -n "$TILE" ] && [ -d "$TILE/src/tilegym" ]; then + # 1. The cutile-rs wrapper module for THIS kernel must not exist yet (D creates it). + for cand in "$TILE/src/tilegym/ops/cutile_rs/${KERNEL}.py" \ + "$TILE/src/tilegym/ops/cutile_rs/${KERNEL//_/}.py"; do + if [ -f "$cand" ]; then + echo "FAIL: Agent B wrote the Python wrapper ($cand)." + echo " Agent B is Rust-only (kernel.rs/ffi.rs). The wrapper is Agent D's deliverable." + echo " Remove it — Agent D creates the wrapper in its STEP 1.5." + fail=$((fail+1)) + fi + done + + # 2. The tilegym test file must not carry the cutile-rs parametrization yet (D adds it). + TEST_FILE="$TILE/tests/ops/test_${KERNEL}.py" + if [ -f "$TEST_FILE" ] && grep -qE '("cutile-rs"|cutile_rs)' "$TEST_FILE"; then + echo "FAIL: Agent B wired cutile-rs into ${TEST_FILE}." + echo " Test-file parametrization is Agent D's job. Revert the edit." + fail=$((fail+1)) + fi + + # 3. ops/cutile_rs/__init__.py must not register THIS kernel yet (D adds the import). + RS_INIT="$TILE/src/tilegym/ops/cutile_rs/__init__.py" + if [ -f "$RS_INIT" ] && grep -qE "(from[[:space:]]+\.[[:space:]]+import[[:space:]]+${KERNEL}([[:space:]]|,|$)|import[[:space:]]+${KERNEL}([[:space:]]|,|$))" "$RS_INIT"; then + echo "FAIL: Agent B registered ${KERNEL} in ops/cutile_rs/__init__.py." + echo " Backend registration is Agent D's job. Revert the import." + fail=$((fail+1)) + fi +else + echo "WARN: TILEGYM_PATH unset — skipping Agent B lane guard" +fi + +echo "---" +if [ $fail -eq 0 ]; then echo "VERDICT: COMPILED"; exit 0; else echo "VERDICT: FAIL_COMPILE"; exit 1; fi diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_d.sh b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_d.sh new file mode 100755 index 0000000..8e3d9a0 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_d.sh @@ -0,0 +1,299 @@ +#!/bin/bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +# Validate Agent D outputs: correctness.md with VERDICT first line. + +set -uo pipefail +KERNEL="${1:?Usage: validate_agent_d.sh }" + +# Per-kernel working dir (eval harness depends on it). There is no separate +# cutile-rs checkout / CUTILE_RS_ROOT anymore — the aggregated cutile_kernels +# crate lives under $TILEGYM_PATH/src/tilegym/ops/cutile_rs/cutile_kernels and +# builds ONE libcutile_kernels.so. +: "${CUTILE_KERNEL_OUT_ROOT:?CUTILE_KERNEL_OUT_ROOT must be set (e.g. $ISOLATED_CWD/cutile_kernel_out)}" + +BASE="${CUTILE_KERNEL_OUT_ROOT}/${KERNEL}" +REPORT="${BASE}/reports/correctness.md" +fail=0 + +if [ ! -f "$REPORT" ] || [ ! -s "$REPORT" ]; then + echo "FAIL: correctness.md missing or empty" + exit 1 +fi + +# First line must be VERDICT: +FIRST_LINE=$(head -1 "$REPORT") +if echo "$FIRST_LINE" | grep -qE "^VERDICT: (ALL_PASS|FAIL)$"; then + echo "OK: verdict line: ${FIRST_LINE}" +else + echo "FAIL: first line must be 'VERDICT: ALL_PASS' or 'VERDICT: FAIL', got: ${FIRST_LINE}" + fail=$((fail+1)) +fi + +# Must have results table +if grep -q "| Test " "$REPORT"; then + echo "OK: has results table" +else + echo "FAIL: missing results table" + fail=$((fail+1)) +fi + +# Raw pytest log — format checks (not just existence) +RAW_LOG="${BASE}/correctness_cutile_rs.txt" +if [ ! -f "$RAW_LOG" ]; then + echo "FAIL: correctness_cutile_rs.txt missing" + fail=$((fail+1)) +elif [ ! -s "$RAW_LOG" ]; then + echo "FAIL: correctness_cutile_rs.txt empty" + fail=$((fail+1)) +else + RAW_PASS=$(grep -c "PASSED" "$RAW_LOG" 2>/dev/null) || RAW_PASS=0 + RAW_FAIL=$(grep -c "FAILED" "$RAW_LOG" 2>/dev/null) || RAW_FAIL=0 + RAW_SKIP=$(grep -c "SKIPPED" "$RAW_LOG" 2>/dev/null) || RAW_SKIP=0 + RAW_TOTAL=$((RAW_PASS + RAW_FAIL + RAW_SKIP)) + echo "OK: correctness_cutile_rs.txt ($(wc -l < "$RAW_LOG") lines, passed=${RAW_PASS} failed=${RAW_FAIL} skipped=${RAW_SKIP})" + if [ "${RAW_TOTAL}" -eq 0 ]; then + echo "FAIL: correctness_cutile_rs.txt has no PASSED/FAILED/SKIPPED — not valid pytest output" + fail=$((fail+1)) + fi + + # ── Fake-pytest detection (v42 anti-pattern) ── + # Real pytest output MUST contain at least one of these signature lines + # printed by the pytest framework itself. If none are present, Agent D + # was almost certainly running a hand-rolled python loop that mimics + # pytest format strings — banned by agent_d.md HARD RULE. + # Require markers that fake runners DON'T usually mimic. `collected N items` + # and `test session starts` were mimicked by v42's fake runner — exclude them. + # `platform `, `rootdir:`, `plugins:` are emitted ONLY by real pytest framework. + if grep -qE "^platform [a-z]+ -- Python|^rootdir:|^plugins:" "$RAW_LOG"; then + echo "OK: correctness_cutile_rs.txt contains real-pytest framework markers" + else + echo "FAIL: correctness_cutile_rs.txt has PASSED/FAILED lines but NO real-pytest framework markers" + echo " (no 'platform linux -- Python ...', 'rootdir: ...', 'plugins: ...')" + echo " → looks like a hand-rolled fake-pytest runner (banned, see agent_d.md HARD RULE)" + fail=$((fail+1)) + fi + # Also flag the specific v42 anti-pattern phrase if it leaks in + if grep -qE "FAILED - max_diff=.*> 1e-2$|if diff > 1e-2" "$RAW_LOG"; then + echo "FAIL: correctness_cutile_rs.txt contains hand-rolled abs-threshold judgment" + echo " ('max_diff > 1e-2' / 'if diff > 1e-2'). MUST use real pytest + common.py:compare_tensors." + fail=$((fail+1)) + fi +fi + +# Agent log +if [ -f "${BASE}/reports/agent_logs/agent_d.md" ] && [ -s "${BASE}/reports/agent_logs/agent_d.md" ]; then + echo "OK: agent_d.md log" +else + echo "FAIL: agent_d.md log missing" + fail=$((fail+1)) +fi + +# ───────────────────────────────────────────────────────────────── +# ffi.rs + libcutile_kernels.so checks (device/host split): Agent D now writes +# ffi.rs (the C-ABI host launcher), wires the op into the aggregated +# cutile_kernels crate, and builds the single shared object. The launch path — +# output-partition ABI, launch-grid validation, borrow_tensor / DevicePointer +# ownership — is only exercised by D's pytest, so D owns it end-to-end. +# ───────────────────────────────────────────────────────────────── +BASE="${CUTILE_KERNEL_OUT_ROOT}/${KERNEL}" + +# A. ffi.rs must exist (D wrote it) +if [ -f "${BASE}/ffi.rs" ] && [ -s "${BASE}/ffi.rs" ]; then + echo "OK: ffi.rs present (Agent D)" +else + echo "FAIL: ffi.rs missing — Agent D must write the C-ABI host launcher (${BASE}/ffi.rs)" + fail=$((fail+1)) +fi + +# A'. Tensors must cross the FFI as *const TensorDesc (not raw ptr + dim/stride +# args), and the op must NOT reach for op-level from_raw_parts / mem::forget / +# transmute — unpacking is via borrow_tensor:: or DevicePointer::from_cu_deviceptr. +if [ -f "${BASE}/ffi.rs" ]; then + if grep -qE '\*const[[:space:]]+TensorDesc' "${BASE}/ffi.rs"; then + echo "OK: ffi.rs takes tensors as *const TensorDesc" + else + echo "FAIL: ffi.rs does not take tensors as '*const TensorDesc'" + echo " Tensors cross the boundary as descriptors now (dtype/shape/strides inside)," + echo " not raw ptr + loose dim/stride/dtype/elem_size args." + fail=$((fail+1)) + fi + if grep -qE 'borrow_tensor|DevicePointer::from_cu_deviceptr' "${BASE}/ffi.rs"; then + echo "OK: ffi.rs unpacks descriptors via borrow_tensor / DevicePointer::from_cu_deviceptr" + else + echo "FAIL: ffi.rs does not unpack the descriptor via borrow_tensor:: or DevicePointer::from_cu_deviceptr" + fail=$((fail+1)) + fi + if grep -qE 'from_raw_parts|mem::forget|transmute' "${BASE}/ffi.rs"; then + echo "FAIL: ffi.rs uses op-level from_raw_parts / mem::forget / transmute" + echo " These are banned at op level: from_raw_parts lives only inside ffi_util::borrow_tensor," + echo " and ManuallyDrop (borrow_tensor) is the ownership gate — no mem::forget/transmute needed." + fail=$((fail+1)) + else + echo "OK: no op-level from_raw_parts / mem::forget / transmute in ffi.rs" + fi + # Rust 2024 export + real device_id (not hardcoded Device::new(0)) + null-check. + if grep -qE '#\[unsafe\(no_mangle\)\]' "${BASE}/ffi.rs"; then + echo "OK: ffi.rs uses #[unsafe(no_mangle)] (Rust 2024 export)" + else + echo "FAIL: ffi.rs does not use #[unsafe(no_mangle)] (legacy #[no_mangle] is banned)" + fail=$((fail+1)) + fi + if grep -qE 'Device::new\([[:space:]]*device_id' "${BASE}/ffi.rs"; then + echo "OK: ffi.rs builds Device from device_id (multi-GPU correct)" + else + echo "FAIL: ffi.rs does not build Device::new(device_id.max(0) as usize) — must not hardcode Device::new(0)" + fail=$((fail+1)) + fi +fi + +# B. the aggregated shared object (single libcutile_kernels.so) must exist and +# export this op's symbol cutile_{kernel}. The loader autobuilds the +# cutile_kernels crate; CUTILE_RS_KERNELS_DIR overrides its location. +KERNELS_DIR="${CUTILE_RS_KERNELS_DIR:-${TILEGYM_PATH:-}/src/tilegym/ops/cutile_rs/cutile_kernels}" +SO=$(find "${KERNELS_DIR}/target/release" "${TILEGYM_PATH:-}/src/tilegym/ops/cutile_rs" \ + -name libcutile_kernels.so 2>/dev/null | head -1) +if [ -n "$SO" ] && nm -D "$SO" 2>/dev/null | grep -q " T cutile_${KERNEL}$"; then + echo "OK: libcutile_kernels.so ($SO) exports symbol cutile_${KERNEL}" +else + echo "FAIL: libcutile_kernels.so missing or symbol cutile_${KERNEL} not exported" + echo " Agent D must register the op in cutile_kernels/src/lib.rs" + echo " mod ${KERNEL} { include!(\"../../${KERNEL}_kernel/kernel.rs\"); include!(\"../../${KERNEL}_kernel/ffi.rs\"); }" + echo " then build the single aggregated cdylib:" + echo " cd ${KERNELS_DIR} && cargo build --release" + fail=$((fail+1)) +fi + +# C. FFI compile-option lever: if analysis.json tunes num_cta_in_cga, ffi.rs must +# expose+apply it (so the wrapper's autotuned value actually reaches codegen). +ANALYSIS="${BASE}/reference/analysis.json" +if [ -f "$ANALYSIS" ] && [ -f "${BASE}/ffi.rs" ]; then + TUNES_CGA=$(python3 -c " +import json +d = json.load(open('${ANALYSIS}')); s = json.dumps(d) +print('yes' if ('num_cta_in_cga' in s or 'num_ctas' in s) else 'no') +" 2>/dev/null) + if [ "$TUNES_CGA" = "yes" ]; then + if grep -q "num_cta_in_cga" "${BASE}/ffi.rs"; then + echo "OK: ffi.rs exposes/applies num_cta_in_cga compile-option" + else + echo "FAIL: analysis.json tunes num_cta_in_cga/num_ctas but ffi.rs does not apply it" + echo " Agent D must: if num_cta_in_cga > 0 { opts = opts.num_cta_in_cga(v); }" + fail=$((fail+1)) + fi + fi +fi + +# ───────────────────────────────────────────────────────────────── +# Wrapper + backend-registration + cffi boundary + autotune checks. +# Agent D owns the Python wrapper and all tilegym backend wiring. +# ───────────────────────────────────────────────────────────────── +TILE="${TILEGYM_PATH:-}" + +# Locate the wrapper D was supposed to write. +WRAPPER="" +if [ -n "$TILE" ]; then + WRAPPER=$(find "$TILE/src/tilegym/ops/cutile_rs" -name "${KERNEL}.py" -o -name "${KERNEL//_/}.py" 2>/dev/null | head -1) + if [ -z "$WRAPPER" ]; then + WRAPPER=$(find "$TILE/src/tilegym/ops/cutile_rs" -name "*.py" -exec grep -l "register_impl.*${KERNEL}" {} \; 2>/dev/null | head -1) + fi +fi + +# 1. wrapper exists +if [ -n "$WRAPPER" ] && [ -f "$WRAPPER" ]; then + echo "OK: cutile_rs wrapper present ($WRAPPER)" +else + echo "FAIL: cutile_rs wrapper for ${KERNEL} not found under ops/cutile_rs/" + echo " Agent D must create src/tilegym/ops/cutile_rs/.py (STEP 1.5)" + fail=$((fail+1)) +fi + +# 2. cffi boundary — the wrapper must declare a _FFI_CDEF whose tensor args are +# `const TensorDesc*`, bind via bind_kernel_function_cffi, and pack tensors with +# make_tensor_desc. (This replaces the old ctypes _FFI_ARGTYPES list; the +# descriptor carries dtype/shape/strides, so there is no loose arg list to drift +# and no 32-bit pointer truncation risk.) +if [ -n "$WRAPPER" ] && [ -f "$WRAPPER" ]; then + if grep -qE '_FFI_CDEF|bind_kernel_function_cffi' "$WRAPPER"; then + echo "OK: wrapper uses cffi (_FFI_CDEF + bind_kernel_function_cffi)" + else + echo "FAIL: wrapper missing cffi boundary (_FFI_CDEF / bind_kernel_function_cffi)" + echo " cutile-rs wrappers are cffi now, not ctypes. Declare a _FFI_CDEF string whose" + echo " tensor args are 'const TensorDesc*', and bind with bind_kernel_function_cffi." + fail=$((fail+1)) + fi + if grep -qE '\bmake_tensor_desc\b' "$WRAPPER"; then + echo "OK: wrapper packs tensors with make_tensor_desc" + else + echo "FAIL: wrapper does not pack tensors with make_tensor_desc (const TensorDesc* boundary)" + fail=$((fail+1)) + fi + if grep -qE '\bcheck_rc\b' "$WRAPPER"; then + echo "OK: wrapper checks the FFI return code with check_rc" + else + echo "FAIL: wrapper does not call check_rc(rc, _FFI_NAME) after the launch" + fail=$((fail+1)) + fi + # Guard against a leftover ctypes argtypes list (old pattern). + if grep -qE '_FFI_ARGTYPES|\.argtypes[[:space:]]*=' "$WRAPPER"; then + echo "FAIL: wrapper still declares ctypes _FFI_ARGTYPES/.argtypes (old pattern)" + echo " Replace with the cffi _FFI_CDEF + make_tensor_desc boundary." + fail=$((fail+1)) + fi +fi + +# 3. autotune_launch when analysis.json carries autotune configs +ANALYSIS="${BASE}/reference/analysis.json" +if [ -f "$ANALYSIS" ] && [ -n "$WRAPPER" ] && [ -f "$WRAPPER" ]; then + HAS_AUTOTUNE=$(python3 -c " +import json +d = json.load(open('${ANALYSIS}')) +at = d.get('autotune_configs', []); ab = d.get('autotune_backends', []) +print('yes' if (at and len(at) > 0) or (ab and len(ab) > 0) else 'no') +" 2>/dev/null) + if [ "$HAS_AUTOTUNE" = "yes" ]; then + if grep -q "autotune_launch" "$WRAPPER"; then + echo "OK: wrapper uses autotune_launch (autotune_configs present)" + else + echo "FAIL: analysis.json has autotune_configs but wrapper does NOT use autotune_launch" + echo " Agent D must import from tilegym.backend.cutile_rs and call autotune_launch()" + fail=$((fail+1)) + fi + else + echo "OK: no autotune required (analysis.json has no autotune_configs)" + fi +fi + +# 4. backend registration — 3 entry points or pytest sees 0 cutile-rs tests +if [ -n "$TILE" ] && [ -d "$TILE/src/tilegym" ]; then + OPS_INIT="$TILE/src/tilegym/ops/__init__.py" + if [ -f "$OPS_INIT" ] && grep -qE "(from\s+\.\s+import\s+cutile_rs|import.*cutile_rs)" "$OPS_INIT"; then + echo "OK: ops/__init__.py imports cutile_rs" + else + echo "FAIL: ops/__init__.py does NOT import cutile_rs (Agent D must add 'from . import cutile_rs')" + fail=$((fail+1)) + fi + + SELECTOR="$TILE/src/tilegym/backend/selector.py" + if [ -f "$SELECTOR" ] && grep -qE '("cutile-rs"|cutile_rs)' "$SELECTOR"; then + echo "OK: backend/selector.py references cutile-rs" + else + echo "FAIL: backend/selector.py does NOT register cutile-rs (Agent D must wire is_cutile_rs_available())" + fail=$((fail+1)) + fi + + TEST_FILE="$TILE/tests/ops/test_${KERNEL}.py" + if [ -f "$TEST_FILE" ] && grep -qE '("cutile-rs"|cutile_rs)' "$TEST_FILE"; then + echo "OK: tests/ops/test_${KERNEL}.py includes cutile-rs backend" + else + echo "FAIL: tests/ops/test_${KERNEL}.py does NOT include cutile-rs (Agent D must add it to parametrization)" + fail=$((fail+1)) + fi +else + echo "WARN: TILEGYM_PATH unset or missing — skipping backend registration checks" +fi + +echo "---" +if [ $fail -eq 0 ]; then echo "PASS"; exit 0; else echo "FAIL: ${fail} issue(s)"; exit 1; fi diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_e.sh b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_e.sh new file mode 100755 index 0000000..0880e96 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_e.sh @@ -0,0 +1,69 @@ +#!/bin/bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +# Validate Agent E outputs: performance.md with VERDICT, raw perf logs. + +set -uo pipefail +KERNEL="${1:?Usage: validate_agent_e.sh }" + +# Auto-resolve CUTILE_RS_ROOT from this script's location +: "${CUTILE_RS_ROOT:=$(cd "$(dirname "${BASH_SOURCE[0]}")/../../../.." && pwd)}" +: "${CUTILE_KERNEL_OUT_ROOT:?CUTILE_KERNEL_OUT_ROOT must be set (e.g. $ISOLATED_CWD/cutile_kernel_out)}" + +BASE="${CUTILE_KERNEL_OUT_ROOT}/${KERNEL}" +SCRIPTS_DIR="$(dirname "$0")" +REPORT="${BASE}/reports/performance.md" +fail=0 + +# performance.md exists with VERDICT +if [ ! -f "$REPORT" ] || [ ! -s "$REPORT" ]; then + echo "FAIL: performance.md missing or empty" + fail=$((fail+1)) +else + FIRST_LINE=$(head -1 "$REPORT") + if echo "$FIRST_LINE" | grep -qE "^VERDICT: (DONE|INVESTIGATE)$"; then + echo "OK: verdict line: ${FIRST_LINE}" + else + echo "FAIL: first line must be 'VERDICT: DONE' or 'VERDICT: INVESTIGATE', got: ${FIRST_LINE}" + fail=$((fail+1)) + fi + + # Machine-readable geomean line — scorer's authoritative perf source. + # Must be exactly: geomean_ratio= (no spaces around '=', strictly + # numeric: digits + optional single decimal fraction, no trailing 'x'). + # Exactly one such line is required. + GM_COUNT=$(grep -cE "^geomean_ratio=[0-9]+(\.[0-9]+)?[[:space:]]*$" "$REPORT") + if [ "$GM_COUNT" -eq 1 ]; then + GM_LINE=$(grep -E "^geomean_ratio=" "$REPORT" | head -1) + echo "OK: machine-readable geomean: ${GM_LINE}" + elif [ "$GM_COUNT" -eq 0 ]; then + echo "FAIL: performance.md missing required 'geomean_ratio=X.XXXX' line (scorer perf source)" + fail=$((fail+1)) + else + echo "FAIL: performance.md has ${GM_COUNT} 'geomean_ratio=' lines, expected exactly 1" + fail=$((fail+1)) + fi +fi + +# Raw perf logs via validate_perf_log.sh +for log in perf_cutile_rs.txt baseline_perf_cutile.txt; do + if bash "${SCRIPTS_DIR}/validate_perf_log.sh" "${BASE}/${log}" "${log}" 2>/dev/null; then + echo "OK: ${log}" + else + echo "FAIL: ${log}" + fail=$((fail+1)) + fi +done + +# Agent log +if [ -f "${BASE}/reports/agent_logs/agent_e.md" ] && [ -s "${BASE}/reports/agent_logs/agent_e.md" ]; then + echo "OK: agent_e.md log" +else + echo "FAIL: agent_e.md log missing" + fail=$((fail+1)) +fi + +echo "---" +if [ $fail -eq 0 ]; then echo "PASS"; exit 0; else echo "FAIL: ${fail} issue(s)"; exit 1; fi diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_f.sh b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_f.sh new file mode 100755 index 0000000..56cd52d --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_agent_f.sh @@ -0,0 +1,45 @@ +#!/bin/bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +# Validate Agent F outputs: perf_investigation_*.md with VERDICT first line. + +set -uo pipefail +KERNEL="${1:?Usage: validate_agent_f.sh }" + +# Auto-resolve CUTILE_RS_ROOT from this script's location +: "${CUTILE_RS_ROOT:=$(cd "$(dirname "${BASH_SOURCE[0]}")/../../../.." && pwd)}" +: "${CUTILE_KERNEL_OUT_ROOT:?CUTILE_KERNEL_OUT_ROOT must be set (e.g. $ISOLATED_CWD/cutile_kernel_out)}" + +BASE="${CUTILE_KERNEL_OUT_ROOT}/${KERNEL}" +fail=0 +count=0 + +for f in "${BASE}"/reports/perf_investigation_*.md; do + [ -f "$f" ] || continue + count=$((count+1)) + FIRST_LINE=$(head -1 "$f") + if echo "$FIRST_LINE" | grep -qE "^VERDICT: (FIXABLE|ALIGNED|BLOCKED)$"; then + echo "OK: $(basename $f): ${FIRST_LINE}" + else + echo "FAIL: $(basename $f): first line must be VERDICT: FIXABLE|ALIGNED|BLOCKED, got: ${FIRST_LINE}" + fail=$((fail+1)) + fi +done + +if [ $count -eq 0 ]; then + echo "FAIL: no perf_investigation_*.md files found" + fail=$((fail+1)) +fi + +# Agent log +if [ -f "${BASE}/reports/agent_logs/agent_f.md" ] && [ -s "${BASE}/reports/agent_logs/agent_f.md" ]; then + echo "OK: agent_f.md log" +else + echo "FAIL: agent_f.md log missing" + fail=$((fail+1)) +fi + +echo "---" +if [ $fail -eq 0 ]; then echo "PASS"; exit 0; else echo "FAIL: ${fail} issue(s)"; exit 1; fi diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_analysis.sh b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_analysis.sh new file mode 100755 index 0000000..85910ea --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_analysis.sh @@ -0,0 +1,245 @@ +#!/bin/bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +# Validate analysis.json has all required fields including correctness_reference and correctness_tolerance. +# Called by Agent A (post-step validation) and validate_kernel.sh (final gate). +# +# Usage: bash validate_analysis.sh +# Exit 0 = valid, Exit 1 = missing required fields + +set -euo pipefail + +KERNEL_NAME="${1:?Usage: validate_analysis.sh }" + +# Auto-resolve CUTILE_RS_ROOT from this script's location +: "${CUTILE_RS_ROOT:=$(cd "$(dirname "${BASH_SOURCE[0]}")/../../../.." && pwd)}" + +# CUTILE_KERNEL_OUT_ROOT locates the per-kernel working dir; without it (set -u) +# the ANALYSIS expansion below aborts with a cryptic "unbound variable" — guard +# it so the failure names the missing var. +: "${CUTILE_KERNEL_OUT_ROOT:?CUTILE_KERNEL_OUT_ROOT not set — cannot locate reference/analysis.json}" + +ANALYSIS="${CUTILE_KERNEL_OUT_ROOT}/${KERNEL_NAME}/reference/analysis.json" + +[ -f "$ANALYSIS" ] || { echo "FAIL: $ANALYSIS not found"; exit 1; } +[ -s "$ANALYSIS" ] || { echo "FAIL: $ANALYSIS is empty"; exit 1; } + +# Validate JSON is parseable +if ! python3 -c "import json; json.load(open('${ANALYSIS}'))" 2>/dev/null; then + echo "FAIL: $ANALYSIS is not valid JSON" + exit 1 +fi + +# Check required fields using python (pass path via env var) +ANALYSIS_PATH="$ANALYSIS" python3 -c " +import json, os, sys + +d = json.load(open(os.environ['ANALYSIS_PATH'])) + +errors = [] +warnings = [] + +# --- Required top-level fields --- +for f in ['kernel_name', 'source', 'launch_path', 'pattern', 'ops_used', 'test_perf_configs']: + if f not in d: + errors.append(f'Missing required field: {f}') + +# --- Reference backend selection --- +ref_backend = d.get('reference_backend') +if ref_backend is None: + warnings.append('Missing reference_backend — should be \"cutile\" or \"triton\" (defaults to cutile for backward compat)') +elif ref_backend not in ('cutile', 'triton'): + errors.append(f'reference_backend must be \"cutile\" or \"triton\", got \"{ref_backend}\"') + +ref_selection = d.get('reference_backend_selection') +if ref_selection is None: + warnings.append('Missing reference_backend_selection — Agent A should compare cutile vs triton perf and pick the winner') +elif isinstance(ref_selection, dict): + if 'winner' not in ref_selection: + warnings.append('reference_backend_selection missing \"winner\" field') + if 'cutile_geo_mean_ms' not in ref_selection and 'per_variant' not in ref_selection: + warnings.append('reference_backend_selection has no perf data (need cutile_geo_mean_ms or per_variant)') + +# --- test_perf_configs must have shapes --- +configs = d.get('test_perf_configs', {}) +if isinstance(configs, dict): + total_shapes = sum(len(v.get('shapes', [])) for v in configs.values()) +elif isinstance(configs, list): + total_shapes = sum(len(c.get('param_list', [])) for c in configs) +else: + total_shapes = 0 +if total_shapes == 0: + errors.append('test_perf_configs has 0 shapes') + +# --- kernel_variants (multi-variant ops) --- +variants = d.get('kernel_variants') +if variants is not None: + if not isinstance(variants, list): + errors.append('kernel_variants must be a list of variant dicts') + elif len(variants) == 0: + warnings.append('kernel_variants is empty — should have at least 1 entry') + else: + for i, v in enumerate(variants): + if not isinstance(v, dict): + errors.append(f'kernel_variants[{i}] must be a dict') + continue + vname = v.get('variant_name', f'variant_{i}') + for req_key in ['variant_name', 'kernel_function', 'reference_ir']: + if req_key not in v: + errors.append(f'kernel_variants[{i}] ({vname}) missing {req_key}') + if 'dispatch_condition' not in v: + warnings.append(f'kernel_variants[{i}] ({vname}) missing dispatch_condition') + + # Per-variant reference_backend (pipeline rule 15: dual-backend per-variant selection) + v_ref_backend = v.get('reference_backend') + if v_ref_backend is None: + warnings.append(f'kernel_variants[{i}] ({vname}) missing reference_backend — should be \"cutile\" or \"triton\"') + elif v_ref_backend not in ('cutile', 'triton'): + errors.append(f'kernel_variants[{i}] ({vname}) reference_backend must be \"cutile\" or \"triton\", got \"{v_ref_backend}\"') + + # Per-variant reference_ir file must exist on disk + v_ir = v.get('reference_ir') + if v_ir: + import os as _os + ir_path = _os.path.join(_os.path.dirname(os.environ['ANALYSIS_PATH']), '..', v_ir) + ir_path_abs = _os.path.normpath(ir_path) + if not _os.path.isfile(ir_path_abs): + errors.append(f'kernel_variants[{i}] ({vname}) reference_ir file not found: {v_ir} (resolved: {ir_path_abs})') + + # Per-variant perf data in reference_backend_selection + if ref_selection and isinstance(ref_selection, dict): + pv_sel = ref_selection.get('per_variant') + if pv_sel is None: + warnings.append('reference_backend_selection missing per_variant — should have per-variant perf comparison') + elif isinstance(pv_sel, dict): + for i, v in enumerate(variants): + if not isinstance(v, dict): + continue + vname = v.get('variant_name', f'variant_{i}') + if vname not in pv_sel: + warnings.append(f'reference_backend_selection.per_variant missing entry for \"{vname}\"') + + # If variants exist, config_to_variant should map each test config to a variant + ctv = d.get('config_to_variant') + if ctv is None: + warnings.append('kernel_variants exist but config_to_variant is missing — Agent E needs this to compare variant-to-variant') + elif not isinstance(ctv, dict) or len(ctv) == 0: + warnings.append('config_to_variant is empty — should map each test_perf config to its variant') + +# --- correctness_reference (pipeline rule 8) --- +ref = d.get('correctness_reference') +if ref is None: + errors.append('Missing correctness_reference — extract def reference(...) from tilegym test class') +elif not isinstance(ref, dict): + errors.append('correctness_reference must be a dict with function and source keys') +else: + if not ref.get('function'): + errors.append('correctness_reference.function is empty') + if 'source' not in ref: + warnings.append('correctness_reference missing source (file:line)') + +# --- correctness_tolerance (pipeline rule 9) --- +tol = d.get('correctness_tolerance') +dtype_entries = {} +if tol is None: + errors.append('Missing correctness_tolerance — extract atol/rtol from tilegym test_op') +elif not isinstance(tol, dict): + errors.append('correctness_tolerance must be a dict with dtype keys') +else: + dtype_entries = {k: v for k, v in tol.items() if k not in ('source', 'notes')} + if len(dtype_entries) == 0: + errors.append('correctness_tolerance has no dtype entries (need f32/f16/bf16)') + for dk, dv in dtype_entries.items(): + if not isinstance(dv, dict): + errors.append(f'correctness_tolerance[{dk}] must have atol and rtol') + else: + if 'atol' not in dv: + errors.append(f'correctness_tolerance[{dk}] missing atol') + if 'rtol' not in dv: + errors.append(f'correctness_tolerance[{dk}] missing rtol') + if 'source' not in tol: + warnings.append('correctness_tolerance missing source (file:line)') + +# --- Autotune fields (if launch_path is autotune or autotune_backends present) --- +launch_path = d.get('launch_path', 'direct') +at_backends = d.get('autotune_backends', []) +at_configs = d.get('autotune_configs', []) + +if launch_path == 'autotune' or len(at_backends) > 0: + # Check top-level autotune_configs OR per-variant autotune_configs + has_any_autotune_configs = False + + # Check per-variant autotune_configs + if variants and isinstance(variants, list): + for i, v in enumerate(variants): + if not isinstance(v, dict): + continue + v_at = v.get('autotune_configs') + if v_at is not None and isinstance(v_at, list) and len(v_at) > 0: + has_any_autotune_configs = True + for j, cfg in enumerate(v_at): + if not isinstance(cfg, dict): + errors.append(f'kernel_variants[{i}].autotune_configs[{j}] must be a dict') + elif not any(k.startswith('BLOCK_') or k.startswith('TILE_') for k in cfg): + warnings.append(f'kernel_variants[{i}].autotune_configs[{j}] has no BLOCK_*/TILE_* fields') + + # Check top-level autotune_configs (for single-variant or union) + if at_configs and isinstance(at_configs, list) and len(at_configs) > 0: + has_any_autotune_configs = True + for i, cfg in enumerate(at_configs): + if not isinstance(cfg, dict): + errors.append(f'autotune_configs[{i}] must be a dict') + elif not any(k.startswith('BLOCK_') or k.startswith('TILE_') for k in cfg): + warnings.append(f'autotune_configs[{i}] has no BLOCK_*/TILE_* fields: {list(cfg.keys())}') + + if not has_any_autotune_configs: + errors.append('autotune detected (launch_path=autotune or autotune_backends set) but no autotune_configs found. ' + 'Must be in top-level autotune_configs OR kernel_variants[].autotune_configs.') + + if not at_backends or not isinstance(at_backends, list): + warnings.append('autotune_backends missing — should list which backends use autotune (e.g. [\"cutile\", \"triton\"])') + at_key = d.get('autotune_key') + if not at_key: + warnings.append('autotune_key missing — should list cache key dimensions (e.g. [\"M\", \"N\", \"K\"])') + +n_at_configs = len(at_configs) if isinstance(at_configs, list) else 0 +# Count per-variant configs too +n_variant_at_configs = 0 +if variants and isinstance(variants, list): + for v in variants: + if isinstance(v, dict): + v_at = v.get('autotune_configs') + if v_at and isinstance(v_at, list): + n_variant_at_configs += len(v_at) + +# --- Output --- +for e in errors: + print(f'ERROR: {e}') +for w in warnings: + print(f'WARN: {w}') + +ref_fn = ref.get('function', 'MISSING')[:60] if isinstance(ref, dict) else 'MISSING' +tol_list = list(dtype_entries.keys()) +n_variants = len(variants) if variants and isinstance(variants, list) else 0 +variant_names = [v.get('variant_name','?') for v in variants] if variants and isinstance(variants, list) else [] +print(f'---') +print(f'kernel: {d.get(\"kernel_name\", \"MISSING\")}') +print(f'shapes: {total_shapes}') +print(f'variants: {n_variants} {variant_names}') +print(f'reference: {ref_fn}') +print(f'tolerance: {tol_list}') +print(f'autotune: {\"yes\" if (launch_path == \"autotune\" or at_backends) else \"no\"} backends={at_backends} top_configs={n_at_configs} variant_configs={n_variant_at_configs}') +print(f'errors: {len(errors)}, warnings: {len(warnings)}') +sys.exit(1 if errors else 0) +" +rc=$? + +if [ $rc -ne 0 ]; then + echo "FAIL: analysis.json validation failed. Re-run Agent A." + exit 1 +else + echo "PASS: analysis.json has all required fields" + exit 0 +fi diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_diff_report.sh b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_diff_report.sh new file mode 100755 index 0000000..49facd5 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_diff_report.sh @@ -0,0 +1,178 @@ +#!/bin/bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +# Validate that diff_report.md meets all structural requirements. +# Usage: bash validate_diff_report.sh +# Exit 0 = valid. Exit 1 = missing sections or content. +# +# Agent C MUST run this before reporting done. + +KERNEL_NAME="${1:?Usage: validate_diff_report.sh }" + +# CUTILE_RS_ROOT must be set (validated by preflight.sh). The script lives in +# tilegym/.agents/skills/..., so any auto-resolve from its location would point +# at the tilegym repo, not the cutile-rs checkout — that's why we require it. +: "${CUTILE_RS_ROOT:?CUTILE_RS_ROOT not set. Run scripts/preflight.sh first or export CUTILE_RS_ROOT=/abs/path/to/cutile-rs.}" + +REPORT="${CUTILE_KERNEL_OUT_ROOT}/${KERNEL_NAME}/generated/diff_report.md" +LOG="${CUTILE_KERNEL_OUT_ROOT}/${KERNEL_NAME}/reports/agent_logs/agent_c.md" + +errors=0 +warns=0 + +fail() { echo "FAIL: $1"; errors=$((errors + 1)); } +warn() { echo "WARN: $1"; warns=$((warns + 1)); } +ok() { echo " OK: $1"; } + +echo "Validating diff_report.md for: ${KERNEL_NAME}" +echo "===" + +# ---- File existence ---- +if [ ! -f "${REPORT}" ]; then + fail "diff_report.md not found at ${REPORT}" + echo "===" + echo "FAIL: ${errors} error(s)" + exit 1 +fi + +if [ ! -f "${LOG}" ]; then + warn "agent_c.md log not found at ${LOG}" +fi + +# ---- Step 1: Structural Comparison (op count table) ---- +echo "--- Step 1: Structural Comparison ---" +if grep -qi "structural comparison\|op count\|step 1\|diff_ir" "${REPORT}"; then + ok "Step 1 section found" + # Must have a table with | Op | or | # | + if grep -qE "^\|.*\|.*\|.*\|" "${REPORT}"; then + ok "Contains table" + else + fail "Step 1 has no table (expected | Op | Reference | Generated | Match |)" + fi +else + fail "Step 1 (Structural Comparison / op count) section missing" +fi + +# ---- Step 2: Checklist Diff ---- +echo "--- Step 2: Checklist Diff ---" +if grep -qi "checklist\|step 2\|item-by-item" "${REPORT}"; then + ok "Step 2 section found" + # Must have Status column values + if grep -qiE "MATCH|DIFFERS|MISSING|SEMANTIC_EQUIV" "${REPORT}"; then + ok "Contains status values (MATCH/DIFFERS/MISSING/SEMANTIC_EQUIV)" + else + fail "Step 2 table has no status values" + fi + # Must have Severity column values + if grep -qiE "CRITICAL|WARN|INFO" "${REPORT}"; then + ok "Contains severity values" + else + warn "Step 2 table may be missing severity column" + fi +else + fail "Step 2 (Checklist Diff) section missing" +fi + +# ---- Step 3: Root Cause Analysis ---- +echo "--- Step 3: Root Cause Analysis ---" +if grep -qi "root cause\|step 3" "${REPORT}"; then + ok "Step 3 section found" + # Must classify each difference + if grep -qiE "user_code|dsl_limitation|compiler_bug|semantic_equiv" "${REPORT}"; then + ok "Contains root cause categories" + else + warn "Step 3 may be missing root cause classifications" + fi +else + fail "Step 3 (Root Cause Analysis) section missing" +fi + +# ---- Step 4: Impact Assessment ---- +echo "--- Step 4: Impact Assessment ---" +if grep -qi "impact\|step 4" "${REPORT}"; then + ok "Step 4 section found" + if grep -qiE "correctness|performance|blocking" "${REPORT}"; then + ok "Contains impact types" + else + warn "Step 4 may be missing impact assessment" + fi +else + fail "Step 4 (Impact Assessment) section missing" +fi + +# ---- Step 5: Verdict ---- +echo "--- Step 5: Verdict ---" +if grep -qi "verdict\|step 5" "${REPORT}"; then + ok "Step 5 section found" + # Must have exactly one of the four verdicts + verdict_count=0 + for v in "PASS_WITH_NOTES" "PASS" "FAIL_FIXABLE" "FAIL_COMPILER_BUG"; do + if grep -q "${v}" "${REPORT}"; then + verdict_count=$((verdict_count + 1)) + echo " OK: Verdict found: ${v}" + fi + done + if [ ${verdict_count} -eq 0 ]; then + fail "No verdict found (expected PASS / PASS_WITH_NOTES / FAIL_FIXABLE / FAIL_COMPILER_BUG)" + fi + + # If FAIL_FIXABLE, must have action items + if grep -q "FAIL_FIXABLE" "${REPORT}"; then + if grep -qiE "action item|fix.*agent b|agent b.*fix|what.*fix" "${REPORT}"; then + ok "FAIL_FIXABLE has action items for Agent B" + else + fail "FAIL_FIXABLE verdict but no action items for Agent B" + fi + fi + + # If PASS_WITH_NOTES, must list known gaps + if grep -q "PASS_WITH_NOTES" "${REPORT}"; then + if grep -qiE "known gap|accepted|limitation" "${REPORT}"; then + ok "PASS_WITH_NOTES lists known gaps" + else + warn "PASS_WITH_NOTES but no known gaps listed" + fi + fi +else + fail "Step 5 (Verdict) section missing" +fi + +# ---- Key checks from diff_ir.sh ---- +echo "--- Key attribute checks ---" + +# assume div_by coverage +if grep -qiE "assume.*div_by|assume_div" "${REPORT}"; then + ok "assume div_by analysis present" +else + warn "No mention of assume div_by alignment analysis" +fi + +# flush_to_zero / rounding coverage +if grep -qiE "flush_to_zero|FTZ|rounding.*approx" "${REPORT}"; then + ok "FTZ / rounding analysis present" +else + warn "No mention of flush_to_zero or rounding analysis" +fi + +# strides analysis +if grep -qiE "strides.*\[.*,.*1\]|strides=\[" "${REPORT}"; then + ok "Strides analysis present" +else + warn "No mention of strides=[?,1] analysis" +fi + +# ---- Summary ---- +echo "===" +if [ ${errors} -eq 0 ]; then + if [ ${warns} -gt 0 ]; then + echo "PASS (with ${warns} warning(s))" + else + echo "PASS" + fi + exit 0 +else + echo "FAIL: ${errors} error(s), ${warns} warning(s)" + exit 1 +fi diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_entry_pattern.sh b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_entry_pattern.sh new file mode 100755 index 0000000..00105ef --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_entry_pattern.sh @@ -0,0 +1,302 @@ +#!/bin/bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 +# +# validate_entry_pattern.sh +# +# Mechanically checks that kernel.rs #[cutile::entry] signatures match the +# load/store ops in structural reference.mlir files. The rule: +# +# reference.mlir uses load_view_tko/store_view_tko_mut (TMA-style) +# -> kernel.rs entry MUST use &Tensor +# +# reference.mlir uses load_ptr_tko/store_ptr_tko (pointer-scatter) +# -> kernel.rs entry MUST use *mut E +# +# Multi-variant kernels match reference_.mlir to the entry fn whose +# name contains . Dtype supplement IRs are not structural variants: +# analysis-declared supplements such as reference_ir_f32_supplement, or dtype +# suffix siblings such as reference_non_persistent_f32.mlir next to +# reference_non_persistent.mlir, are ignored for entry matching and reported as +# dtype-lowering evidence. +# +# Exit 0 = PASS. Exit 1 = FAIL_FIXABLE with actionable fix. + +set -uo pipefail +KERNEL="${1:?Usage: validate_entry_pattern.sh }" + +: "${CUTILE_KERNEL_OUT_ROOT:?CUTILE_KERNEL_OUT_ROOT must be set}" + +BASE="${CUTILE_KERNEL_OUT_ROOT}/${KERNEL}" +KERNEL_RS="${BASE}/kernel.rs" +REF_DIR="${BASE}/reference" + +if [ ! -f "$KERNEL_RS" ]; then + echo "FAIL: kernel.rs not found at $KERNEL_RS" + exit 1 +fi +if [ ! -d "$REF_DIR" ]; then + echo "FAIL: reference dir not found at $REF_DIR" + exit 1 +fi + +python3 - "$KERNEL_RS" "$REF_DIR" "$KERNEL" <<'PYEOF' +import glob +import json +import os +import re +import sys + +kernel_rs, ref_dir, kernel_name = sys.argv[1:4] +base_dir = os.path.dirname(ref_dir) + +def read_text(path): + with open(path, encoding="utf-8", errors="ignore") as f: + return f.read() + +def classify_ir(mlir_path): + """Returns (kind, tma_count, ptr_count) where kind is TMA/PTR/MIXED/UNKNOWN.""" + text = read_text(mlir_path) + tma_pat = r"\b(load_view_tko|store_view_tko_mut)\b" + ptr_pat = r"\b(load_ptr_tko|store_ptr_tko)\b" + tma = len(re.findall(tma_pat, text)) + ptr = len(re.findall(ptr_pat, text)) + if tma > 0 and ptr > 0: + return ("MIXED", tma, ptr) + if tma > 0: + return ("TMA", tma, 0) + if ptr > 0: + return ("PTR", 0, ptr) + return ("UNKNOWN", 0, 0) + +def analysis_supplement_paths(): + analysis = os.path.join(ref_dir, "analysis.json") + if not os.path.exists(analysis): + analysis = os.path.join(base_dir, "reference", "analysis.json") + try: + data = json.load(open(analysis, encoding="utf-8")) + except Exception: + return set(), set() + + paths = set() + basenames = set() + + def norm_path(value): + if os.path.isabs(value): + p = os.path.normpath(value) + else: + p = os.path.normpath(os.path.join(base_dir, value)) + return p + + def walk(obj, in_supplement=False): + if isinstance(obj, dict): + for key, value in obj.items(): + walk(value, in_supplement or ("supplement" in str(key).lower())) + elif isinstance(obj, list): + for value in obj: + walk(value, in_supplement) + elif in_supplement and isinstance(obj, str) and obj.endswith(".mlir"): + p = norm_path(obj) + paths.add(p) + basenames.add(os.path.basename(p)) + + walk(data) + return paths, basenames + +supp_paths, supp_basenames = analysis_supplement_paths() + +all_variant_mlirs = sorted(glob.glob(os.path.join(ref_dir, "reference_*.mlir"))) +single_mlir = os.path.join(ref_dir, "reference.mlir") +variant_basenames = {os.path.basename(p) for p in all_variant_mlirs} +dtype_suffix = re.compile(r"^reference_(.+)_(f16|f32|bf16|fp8|tf32|float16|float32)\.mlir$") + +def is_dtype_suffix_sibling(path): + name = os.path.basename(path) + m = dtype_suffix.match(name) + if not m: + return False + structural_sibling = f"reference_{m.group(1)}.mlir" + return structural_sibling in variant_basenames + +def is_supplement(path): + norm = os.path.normpath(path) + name = os.path.basename(path) + return norm in supp_paths or name in supp_basenames or is_dtype_suffix_sibling(path) + +ignored_supplements = [p for p in all_variant_mlirs if is_supplement(p)] +variant_mlirs = [p for p in all_variant_mlirs if not is_supplement(p)] + +if variant_mlirs: + ref_mlirs = variant_mlirs +elif os.path.exists(single_mlir): + ref_mlirs = [single_mlir] +else: + print(f"FAIL: no structural reference IR (.mlir) found in {ref_dir}") + print(" Agent A must produce reference.mlir or structural reference_.mlir") + if ignored_supplements: + print(" Ignored supplement IRs cannot be the only structural references:") + for p in ignored_supplements: + print(f" - {os.path.basename(p)}") + sys.exit(1) + +ir_classes = {} +for m in ref_mlirs: + ir_classes[os.path.basename(m)] = classify_ir(m) + +text = read_text(kernel_rs) + +def find_entries(src): + """Yield (fn_name, params_str). params_str excludes the outer parens.""" + i, n = 0, len(src) + while i < n: + m = re.search(r"#\[cutile::entry", src[i:]) + if not m: + return + i += m.end() + depth = 1 + while i < n and depth > 0: + c = src[i] + if c == "[": + depth += 1 + elif c == "]": + depth -= 1 + i += 1 + fn_m = re.match(r"\s*(?:pub\s+)?unsafe\s+fn\s+(\w+)\s*(?:<[^>]*>)?\s*\(", src[i:]) + if not fn_m: + continue + name = fn_m.group(1) + i += fn_m.end() + depth = 1 + start = i + while i < n and depth > 0: + c = src[i] + if c == "(": + depth += 1 + elif c == ")": + depth -= 1 + i += 1 + params = src[start:i-1] + yield (name, params) + +entries = list(find_entries(text)) +if not entries: + print(f"FAIL: no #[cutile::entry] function found in {kernel_rs}") + sys.exit(1) + +def classify_entry(params): + has_tensor = bool(re.search(r"&\s*Tensor\s*<", params)) + has_eptr = bool(re.search(r"\*\s*mut\s+E\b", params)) + if has_tensor and has_eptr: + kind = "mixed" + elif has_tensor: + kind = "tensor" + elif has_eptr: + kind = "ptr" + else: + kind = "unknown" + return (kind, has_tensor, has_eptr) + +entry_classes = {n: classify_entry(p) for n, p in entries} + +def expected_for(ir_kind): + return {"TMA": "tensor", "PTR": "ptr"}.get(ir_kind) + +mismatches = [] + +if len(ref_mlirs) == 1: + ir_b = os.path.basename(ref_mlirs[0]) + ir_kind, tma_n, ptr_n = ir_classes[ir_b] + expected = expected_for(ir_kind) + if expected is None: + print(f"WARN: reference IR {ir_b} has no recognized TMA / pointer-scatter ops") + print(f" (tma={tma_n}, ptr={ptr_n}) - skipping pattern check") + sys.exit(0) + for name, (kind, _, _) in entry_classes.items(): + if kind != expected: + mismatches.append((name, kind, expected, ir_b, ir_kind, tma_n, ptr_n)) +else: + matched_entries = set() + for mlir in ref_mlirs: + ir_b = os.path.basename(mlir) + variant_m = re.match(r"reference_(.+)\.mlir$", ir_b) + if not variant_m: + continue + variant = variant_m.group(1).lower() + v_tokens = variant.split("_") + v_abbrev = "".join(t[0] for t in v_tokens) if len(v_tokens) > 1 else None + candidates = [] + for ename in entry_classes: + el = ename.lower() + if variant in el: + candidates.append((ename, "full")) + elif v_abbrev and (f"_{v_abbrev}_" in f"_{el}_" or el.endswith(f"_{v_abbrev}") or f"_{v_abbrev}" in el): + candidates.append((ename, "abbrev")) + elif all(t in el for t in v_tokens): + candidates.append((ename, "all_tokens")) + ir_kind, tma_n, ptr_n = ir_classes[ir_b] + expected = expected_for(ir_kind) + if expected is None: + continue + if not candidates: + mismatches.append((f"", "NOT_FOUND", expected, ir_b, ir_kind, tma_n, ptr_n)) + continue + for ename, _how in candidates: + matched_entries.add(ename) + kind, _, _ = entry_classes[ename] + if kind != expected: + mismatches.append((ename, kind, expected, ir_b, ir_kind, tma_n, ptr_n)) + +def fix_message(expected): + if expected == "tensor": + return ( + " - For TMA-style entries, change params from `*mut E` to " + "`&Tensor`\n" + " - Remove `make_tensor_view` calls from entry body - the macro builds\n" + " the view automatically from the &Tensor param.\n" + " - FFI builds `Arc>` from raw parts (Rule 12 branch a)\n" + " - Outputs use `partition_full_mut` on `&Tensor`, NEVER `&mut Tensor`" + ) + if expected == "ptr": + return ( + " - For pointer-scatter entries, change params from `&Tensor` to `*mut E`\n" + " - Use iota + mask + offset_tile + load_ptr_tko in entry body\n" + " - No `make_tensor_view` (Rule 12 branch b)" + ) + return " - See coding-rules.md Rule 12 for the decision tree." + +if not mismatches: + print("OK: entry pattern matches structural reference IR") + for name, (kind, _, _) in entry_classes.items(): + print(f" entry `{name}` -> {kind}") + for b, (k, t, p) in ir_classes.items(): + print(f" IR {b} -> {k} (tma_ops={t}, ptr_ops={p})") + for p in ignored_supplements: + print(f" supplement IR {os.path.basename(p)} ignored for entry matching; use for dtype lowering checks") + sys.exit(0) + +print("FAIL_FIXABLE: entry pattern mismatch with structural reference IR") +print() +seen_expected = set() +for name, actual, expected, ir_b, ir_kind, tma_n, ptr_n in mismatches: + print(f" Entry `{name}`:") + print(f" actual: {actual}") + print(f" expected: {expected} (reference {ir_b} uses {ir_kind} ops: tma={tma_n}, ptr={ptr_n})") + print() + seen_expected.add(expected) + +if ignored_supplements: + print(" Dtype supplement IRs ignored for entry matching:") + for p in ignored_supplements: + print(f" - {os.path.basename(p)}") + print(" Do not create fake dtype entries for these files; use them for dtype lowering checks.") + print() + +print(" Fix:") +for exp in sorted(seen_expected): + print(fix_message(exp)) +print() +print(" See: references/coding-rules.md Rule 12 (TMA vs pointer-scatter decision tree)") +sys.exit(1) +PYEOF +EXIT_CODE=$? +exit $EXIT_CODE diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_ir.sh b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_ir.sh new file mode 100755 index 0000000..33495bf --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_ir.sh @@ -0,0 +1,73 @@ +#!/bin/bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +# Validate a dumped CUDA Tile IR file. +# Usage: bash validate_ir.sh +# Exit 0 = valid, Exit 1 = invalid + +set -euo pipefail + +IR_FILE="${1:?Usage: validate_ir.sh }" +KERNEL_NAME="${2:?Usage: validate_ir.sh }" + +[ -f "$IR_FILE" ] || { echo "FAIL: $IR_FILE not found"; exit 1; } +[ -s "$IR_FILE" ] || { echo "FAIL: $IR_FILE is empty"; exit 1; } + +errors=0 + +# Check 1: Must have cuda_tile.module header (MLIR format, not Python IR) +if ! grep -q "cuda_tile.module" "$IR_FILE"; then + echo "FAIL: Missing 'cuda_tile.module' — this is NOT valid MLIR format" + echo " (Got Python-level IR instead? Check CUDA_TILE_DUMP_TILEIR was used)" + errors=$((errors + 1)) +fi + +# Check 2: Must have entry whose symbol contains the kernel name as a substring. +# Triton-TileIR @tile_jit produces mangled names like +# @_static_persistent_bmm_kernel_Kt1_A3f16_... (multi-variant) +# @_bmm_kernel_memref_v2opt3_... (single-variant) +# cuTile-Python @ct.kernel produces literal `@bmm`. +# cutile-rs `#[cutile::module]` produces `@_kernel_`. +# We accept any entry symbol that contains "${KERNEL_NAME}" as a substring +# (matches all three conventions without false-positives — the kernel name +# is op-specific enough that "bmm" inside another op's mangled name is +# extremely unlikely). +if ! grep -qE "entry @[a-zA-Z0-9_]*${KERNEL_NAME}[a-zA-Z0-9_]*" "$IR_FILE"; then + echo "FAIL: No entry symbol contains '${KERNEL_NAME}' as substring" + # Show what entries exist + echo " Found entries:" + grep -o "entry @[a-zA-Z0-9_]*" "$IR_FILE" | sed 's/^/ /' + errors=$((errors + 1)) +fi + +# Check 3: Must have optimization_hints (even if empty) +if ! grep -q "optimization_hints" "$IR_FILE"; then + echo "WARN: Missing 'optimization_hints' in entry" +fi + +# Check 4: Must not be Python-level IR (Tile[PointerTy...] format) +if grep -q "Tile\[PointerTy\|typed_const\|assume_bounded(x=" "$IR_FILE"; then + echo "FAIL: File contains Python-level IR, not MLIR" + echo " Use CUDA_TILE_DUMP_TILEIR env var to get MLIR format" + errors=$((errors + 1)) +fi + +# Check 5: Should have at least some ops +op_count=$(grep -cE "load_ptr_tko|store_ptr_tko|load_view_tko|store_view_tko|mmaf|reduce|make_partition_view" "$IR_FILE" || true) +if [ "$op_count" -eq 0 ]; then + echo "WARN: No recognized ops found (load/store/mma/reduce/partition)" +fi + +if [ $errors -gt 0 ]; then + echo "---" + echo "FAIL: ${errors} error(s) in IR validation" + exit 1 +else + lines=$(wc -l < "$IR_FILE") + # Show the matched entry symbol for transparency + matched=$(grep -oE "entry @[a-zA-Z0-9_]*${KERNEL_NAME}[a-zA-Z0-9_]*" "$IR_FILE" | head -1) + echo "PASS: Valid MLIR IR matched ${matched:-@${KERNEL_NAME}} (${lines} lines, ${op_count} key ops)" + exit 0 +fi diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_kernel.sh b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_kernel.sh new file mode 100755 index 0000000..03f31f9 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_kernel.sh @@ -0,0 +1,571 @@ +#!/bin/bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +# Validate all required output files for a converted kernel. +# Usage: bash validate_kernel.sh +# Exit 0 = PASS. Exit 1 = FAIL. + +set -uo pipefail + +KERNEL_NAME="${1:?Usage: validate_kernel.sh }" + +# There is no separate cutile-rs checkout / CUTILE_RS_ROOT anymore. The +# aggregated cutile_kernels crate (one libcutile_kernels.so) lives inside the +# tilegym tree; CUTILE_RS_KERNELS_DIR can override its location. +TILEGYM="${TILEGYM_PATH:?TILEGYM_PATH must be set to your tilegym checkout}" +KERNELS_DIR="${CUTILE_RS_KERNELS_DIR:-${TILEGYM}/src/tilegym/ops/cutile_rs/cutile_kernels}" +BASE="${CUTILE_KERNEL_OUT_ROOT:?CUTILE_KERNEL_OUT_ROOT must be set}/${KERNEL_NAME}" + +fail=0 +warn=0 + +ok() { echo " OK: $1"; } +fail_item() { echo " FAIL: $1"; fail=$((fail + 1)); } +warn_item() { echo " WARN: $1"; warn=$((warn + 1)); } + +check_file() { + local path="$1" + local desc="$2" + if [ -f "$path" ] && [ -s "$path" ]; then + ok "$desc" + else + fail_item "$desc missing or empty -> $path" + fi +} + +check_agent_log() { + local agent="$1" + local required="$2" + local p1="${BASE}/reports/agent_logs/agent_${agent}.md" + local p2="${BASE}/reports/agent_${agent}.md" + + if [ -f "$p1" ] && [ -s "$p1" ]; then + ok "agent_${agent} log (${p1})" + return 0 + fi + if [ -f "$p2" ] && [ -s "$p2" ]; then + ok "agent_${agent} log (${p2}; legacy location accepted)" + return 0 + fi + + if [ "$required" = "yes" ]; then + fail_item "agent_${agent} log missing (checked reports/agent_logs and reports/)" + else + ok "agent_${agent} log not required for this path" + fi +} + +validate_perf_if_present() { + local file="$1" + local label="$2" + local required="$3" + local validator="${TILEGYM}/.agents/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_perf_log.sh" + + if [ ! -f "${BASE}/${file}" ]; then + if [ "$required" = "yes" ]; then + fail_item "${file} missing" + else + warn_item "${file} missing (optional)" + fi + return + fi + + if bash "$validator" "${BASE}/${file}" "$label" 2>/dev/null; then + ok "${file} validated" + else + fail_item "${file} validation failed" + fi +} + +extract_ffi_fn() { + local file="$1" + python3 - "$file" <<'PY' 2>/dev/null +import re +import sys + +path = sys.argv[1] +try: + text = open(path, encoding="utf-8").read() +except OSError: + raise SystemExit(0) + +pat = re.compile( + r"pub\s+(?:unsafe\s+)?extern\s+\"C\"\s+fn\s+([A-Za-z_][A-Za-z0-9_]*)\s*\(", + re.MULTILINE, +) +m = pat.search(text) +if m: + print(m.group(1)) +PY +} + +wrapper_call_info() { + local wrapper="$1" + local ffi_fn="$2" + python3 - "$wrapper" "$ffi_fn" <<'PY' 2>/dev/null +import ast +import sys + +path, ffi = sys.argv[1], sys.argv[2] +try: + text = open(path, encoding="utf-8").read() + tree = ast.parse(text, filename=path) +except Exception: + print("0 0") + raise SystemExit(0) + +constants = {} + +def record_const(target, value): + if isinstance(target, ast.Name) and isinstance(value, ast.Constant) and isinstance(value.value, str): + constants[target.id] = value.value + +for node in ast.walk(tree): + if isinstance(node, ast.Assign): + for target in node.targets: + record_const(target, node.value) + elif isinstance(node, ast.AnnAssign): + record_const(node.target, node.value) + +def const_string(node): + if isinstance(node, ast.Constant) and isinstance(node.value, str): + return node.value + if isinstance(node, ast.Name): + return constants.get(node.id) + return None + +def is_direct_attr(node): + return isinstance(node, ast.Attribute) and node.attr == ffi + +def is_getattr_symbol(node): + return ( + isinstance(node, ast.Call) + and isinstance(node.func, ast.Name) + and node.func.id == "getattr" + and len(node.args) >= 2 + and const_string(node.args[1]) == ffi + ) + +def is_bind_kernel_symbol(node): + if not isinstance(node, ast.Call): + return False + name = None + if isinstance(node.func, ast.Name): + name = node.func.id + elif isinstance(node.func, ast.Attribute): + name = node.func.attr + if name not in {"bind_kernel_function_cffi", "bind_kernel_function", "bind_kernel", "load_kernel_function"}: + return False + return any(const_string(arg) == ffi for arg in node.args) + +def is_symbol_binding_expr(node): + # This identifies expressions that produce a callable bound to the exported + # FFI symbol. Binding alone is not a live call. + return is_direct_attr(node) or is_getattr_symbol(node) or is_bind_kernel_symbol(node) + +aliases = set() + +for node in ast.walk(tree): + if isinstance(node, ast.Assign) and is_symbol_binding_expr(node.value): + for target in node.targets: + if isinstance(target, ast.Name): + aliases.add(target.id) + elif isinstance(node, ast.AnnAssign) and is_symbol_binding_expr(node.value): + if isinstance(node.target, ast.Name): + aliases.add(node.target.id) + +factories = set() + +for node in ast.walk(tree): + if not isinstance(node, ast.FunctionDef): + continue + for child in ast.walk(node): + if isinstance(child, ast.Return): + value = child.value + if isinstance(value, ast.Name) and value.id in aliases: + factories.add(node.name) + elif value is not None and is_symbol_binding_expr(value): + factories.add(node.name) + +def is_factory_call(node): + return ( + isinstance(node, ast.Call) + and isinstance(node.func, ast.Name) + and node.func.id in factories + ) + +for node in ast.walk(tree): + if not isinstance(node, (ast.Assign, ast.AnnAssign)): + continue + value = node.value + if is_factory_call(value): + targets = node.targets if isinstance(node, ast.Assign) else [node.target] + for target in targets: + if isinstance(target, ast.Name): + aliases.add(target.id) + +def is_live_symbol_invocation(func): + # Counts executable calls only: + # lib.cutile_x(...) + # getattr(lib, "cutile_x")(...) + # bind_kernel_function_cffi(_KERNEL, _FFI_CDEF) (binds; ffi/lib then call lib.cutile_x) + # alias(...) + # _lib_fn()(...) where _lib_fn returns a symbol-bound callable. + # Does not count constants, comments, binding-helper calls by themselves, + # or aliases/factories that are never invoked. + return ( + is_direct_attr(func) + or is_getattr_symbol(func) + or is_bind_kernel_symbol(func) + or is_factory_call(func) + or (isinstance(func, ast.Name) and func.id in aliases) + ) + +live_lines = [] +for node in ast.walk(tree): + if not isinstance(node, ast.Call): + continue + if is_live_symbol_invocation(node.func): + live_lines.append(getattr(node, "lineno", 0)) + +if live_lines: + print(len(live_lines), min(live_lines)) +else: + print("0 0") +PY +} + +echo "=== Validating: ${KERNEL_NAME} ===" + +if [ ! -d "$BASE" ]; then + fail_item "kernel output directory missing -> ${BASE}" + echo "===" + echo "Result: ${fail} fail, ${warn} warn" + echo "FAIL" + exit 1 +fi + +echo "--- cutile-rs: kernel files ---" +check_file "${BASE}/kernel.rs" "kernel.rs" +check_file "${BASE}/ffi.rs" "ffi.rs" +check_file "${BASE}/reference/analysis.json" "analysis.json" + +HAS_VARIANTS=$(BASE_DIR="$BASE" python3 - <<'PY' 2>/dev/null +import json +import os +base = os.environ["BASE_DIR"] +path = os.path.join(base, "reference", "analysis.json") +try: + d = json.load(open(path)) +except Exception: + print("unknown") + raise SystemExit +v = d.get("kernel_variants", []) +print("multi" if isinstance(v, list) and len(v) > 1 else "single") +PY +) + +if [ "$HAS_VARIANTS" = "multi" ]; then + BASE_DIR="$BASE" python3 - <<'PY' 2>/dev/null +import json +import os +import sys + +base = os.environ["BASE_DIR"] +d = json.load(open(os.path.join(base, "reference", "analysis.json"))) +missing = [] +for v in d.get("kernel_variants", []): + if not isinstance(v, dict): + continue + ir = v.get("reference_ir") + if not ir: + continue + path = os.path.normpath(os.path.join(base, ir)) + if os.path.isfile(path) and os.path.getsize(path) > 0: + print(f" OK: per-variant IR: {ir}") + else: + print(f" FAIL: per-variant IR missing: {ir} (resolved: {path})") + missing.append(ir) +sys.exit(1 if missing else 0) +PY + [ $? -eq 0 ] || fail=$((fail + 1)) + + gen_count=$(ls "${BASE}/generated/generated_"*.mlir 2>/dev/null | wc -l) + if [ "$gen_count" -gt 0 ]; then + ok "${gen_count} per-variant generated IR files" + else + fail_item "no per-variant generated IR files in generated/" + fi +else + check_file "${BASE}/reference/reference.mlir" "reference IR (single variant)" + check_file "${BASE}/generated/generated.mlir" "generated IR" +fi + +C_RAN=no +[ -s "${BASE}/generated/diff_report.md" ] && C_RAN=yes +[ -s "${BASE}/reports/agent_c.md" ] && C_RAN=yes +[ -s "${BASE}/reports/agent_logs/agent_c.md" ] && C_RAN=yes +if [ "$C_RAN" = "yes" ]; then + check_file "${BASE}/generated/diff_report.md" "IR diff report" + check_agent_log c yes +else + ok "IR diff report not required (Agent C not spawned on happy path)" + check_agent_log c no +fi + +check_file "${BASE}/reports/correctness.md" "correctness report" +check_file "${BASE}/reports/performance.md" "performance report" + +check_agent_log a yes +check_agent_log b yes +check_agent_log d yes +check_agent_log e yes + +F_RAN=no +ls "${BASE}"/reports/perf_investigation_*.md >/dev/null 2>&1 && F_RAN=yes +[ -s "${BASE}/reports/agent_f.md" ] && F_RAN=yes +[ -s "${BASE}/reports/agent_logs/agent_f.md" ] && F_RAN=yes +if [ "$F_RAN" = "yes" ]; then + check_agent_log f yes +else + ok "agent_f log not required (Agent F not spawned)" +fi + +echo "--- cutile-rs: perf raw logs ---" +validate_perf_if_present "baseline_perf_cutile.txt" "baseline_perf_cutile" yes +validate_perf_if_present "baseline_perf_triton_tileir.txt" "baseline_perf_triton_tileir" no +validate_perf_if_present "baseline_perf.txt" "baseline_perf" no +validate_perf_if_present "perf_cutile_rs.txt" "perf_cutile_rs" yes +validate_perf_if_present "perf_cutile_py.txt" "perf_cutile_py" no + +echo "--- cutile-rs: dtype generic (Rule 16) ---" +if [ -f "${BASE}/kernel.rs" ]; then + if grep -q "E: ElementType" "${BASE}/kernel.rs"; then + ok "kernel.rs uses " + else + fail_item "kernel.rs hardcodes dtype - use (Rule 16)" + fi + + hardcoded_ptr=$(grep -nE '\*mut (f16|f32|bf16)' "${BASE}/kernel.rs" 2>/dev/null \ + | grep -vE '\b(rstd|Rstd|RSTD|mean|Mean|MEAN|var|Var|VAR|inv_std|InvStd)([_a-zA-Z0-9]*)?: *\*mut f32' \ + | grep -vE '\*mut (f16|f32|bf16) *[,>]' \ + | wc -l) || hardcoded_ptr=0 + if [ "${hardcoded_ptr}" -gt 0 ]; then + fail_item "kernel entry has hardcoded *mut f16/*mut f32 - use *mut E" + else + ok "no hardcoded dtype in entry signature (norm stats f32 buffers allow-listed)" + fi + + if [ -f "${BASE}/ffi.rs" ]; then + if grep -qE '"f16"|"f32"|"bf16"|dtype' "${BASE}/ffi.rs"; then + ok "ffi.rs passes dtype in generics" + else + fail_item "ffi.rs does not pass dtype string - kernel will not specialize" + fi + fi +fi + +echo "--- cutile-rs: FFI registration / aggregated cdylib ---" + +FFI_FN="" +if [ -f "${BASE}/ffi.rs" ]; then + FFI_FN=$(extract_ffi_fn "${BASE}/ffi.rs" | head -1) +fi +if [ -z "$FFI_FN" ]; then + fail_item "could not extract extern C FFI function from ffi.rs" +else + ok "ffi symbol declared: ${FFI_FN}" +fi + +# ffi.rs ABI (Rule 35): tensors cross as *const TensorDesc, unpacked via +# borrow_tensor / DevicePointer — see references/coding-rules.md Rule 35. +if [ -f "${BASE}/ffi.rs" ]; then + if grep -qE '\*const[[:space:]]+TensorDesc' "${BASE}/ffi.rs"; then + ok "ffi.rs crosses tensors as *const TensorDesc" + else + fail_item "ffi.rs does not use '*const TensorDesc' (tensors carry dtype/shape/strides in the descriptor now)" + fi + if grep -qE 'borrow_tensor|DevicePointer::from_cu_deviceptr' "${BASE}/ffi.rs"; then + ok "ffi.rs unpacks descriptors via borrow_tensor / DevicePointer::from_cu_deviceptr" + else + fail_item "ffi.rs does not unpack via borrow_tensor:: or DevicePointer::from_cu_deviceptr" + fi + if grep -qE 'from_raw_parts|mem::forget|transmute' "${BASE}/ffi.rs"; then + fail_item "ffi.rs uses op-level from_raw_parts / mem::forget / transmute (banned — only ffi_util::borrow_tensor may call from_raw_parts)" + else + ok "no op-level from_raw_parts / mem::forget / transmute in ffi.rs" + fi + if grep -qE '#\[unsafe\(no_mangle\)\]' "${BASE}/ffi.rs"; then + ok "ffi.rs uses #[unsafe(no_mangle)] (Rust 2024 export)" + else + fail_item "ffi.rs missing #[unsafe(no_mangle)] (legacy #[no_mangle] is banned)" + fi +fi + +# ONE aggregated cutile_kernels crate registers every op via a `mod` that +# include!s the op's kernel.rs + ffi.rs. There is no monolithic cutile-ffi +# include and no per-op .so. +LIB_RS="${KERNELS_DIR}/src/lib.rs" +if grep -qE "mod[[:space:]]+${KERNEL_NAME%_kernel}([[:space:]]|\{)|include!\(.*${KERNEL_NAME%_kernel}_kernel/(kernel|ffi)\.rs" "$LIB_RS" 2>/dev/null; then + ok "cutile_kernels/src/lib.rs registers ${KERNEL_NAME%_kernel} (mod + include!)" +else + fail_item "cutile_kernels/src/lib.rs does not register ${KERNEL_NAME%_kernel} (expected 'mod { include!(...) }' → ${LIB_RS})" +fi + +# Deps must be crates.io PINNED (=0.2.0), NOT path deps to a cutile-rs checkout. +CARGO_TOML="${KERNELS_DIR}/Cargo.toml" +if [ -f "$CARGO_TOML" ]; then + if grep -qE 'cutile[[:space:]]*=[[:space:]]*"=0\.2\.0"' "$CARGO_TOML"; then + ok "cutile_kernels/Cargo.toml pins cutile = \"=0.2.0\" (crates.io, Rule 33)" + else + fail_item "cutile_kernels/Cargo.toml does not pin cutile = \"=0.2.0\" (crates.io PINNED deps required)" + fi + if grep -qE 'path[[:space:]]*=' "$CARGO_TOML"; then + fail_item "cutile_kernels/Cargo.toml uses a path dependency (banned — deps are crates.io pins, no cutile-rs checkout)" + else + ok "cutile_kernels/Cargo.toml uses no path deps" + fi +else + fail_item "cutile_kernels/Cargo.toml missing -> ${CARGO_TOML}" +fi + +# The single aggregated cdylib is libcutile_kernels.so; there is no per-op .so +# and no libcutile_ffi.so. +SO_FOUND=$(find "${KERNELS_DIR}/target/release" "${TILEGYM}/src/tilegym/ops/cutile_rs" \ + -name libcutile_kernels.so 2>/dev/null | head -1) +if [ -n "$SO_FOUND" ]; then + ok "aggregated cdylib found: ${SO_FOUND}" +else + fail_item "libcutile_kernels.so not found (build with: cd ${KERNELS_DIR} && cargo build --release)" +fi + +if [ -n "$FFI_FN" ] && [ -n "$SO_FOUND" ]; then + nm_match=$(nm -D "${SO_FOUND}" 2>/dev/null | grep " T ${FFI_FN}" || true) + if [ -n "$nm_match" ]; then + ok "libcutile_kernels.so exports ${FFI_FN}" + else + fail_item "${FFI_FN} not exported by ${SO_FOUND}" + fi +fi + +# ffi_util.rs (shared TensorDesc / borrow_tensor / dtype_str) must be present. +if [ -f "${TILEGYM}/src/tilegym/ops/cutile_rs/ffi_util.rs" ]; then + ok "shared ffi_util.rs present (TensorDesc / borrow_tensor / dtype_str)" +else + warn_item "shared ffi_util.rs not found under ops/cutile_rs/ (expected for the aggregated crate)" +fi + +PIPELINE="" +for p in "${BASE}/${KERNEL_NAME}_pipeline.rs" \ + "${BASE}/${KERNEL_NAME%_kernel}_pipeline.rs"; do + [ -f "$p" ] && PIPELINE="$p" && break +done +if [ -n "$PIPELINE" ]; then + ok "compile test $(basename "${PIPELINE}")" +else + warn_item "compile test not found" +fi + +echo "--- tilegym: wrapper ---" + +WRAPPER="" +if [ -n "$FFI_FN" ]; then + WRAPPER=$(grep -rl "${FFI_FN}" "${TILEGYM}/src/tilegym/ops/cutile_rs/"*.py 2>/dev/null | head -1) +fi +if [ -z "$WRAPPER" ]; then + short="${KERNEL_NAME%_kernel}" + for candidate in "${short}" "${KERNEL_NAME}" "${short//_/}"; do + for f in "${TILEGYM}/src/tilegym/ops/cutile_rs/"*.py; do + [ -f "$f" ] || continue + basename=$(basename "$f" .py) + if echo "${basename}" | grep -qi "${candidate}" 2>/dev/null; then + WRAPPER="$f" + break 2 + fi + done + done +fi + +if [ -z "$WRAPPER" ] || [ ! -f "$WRAPPER" ]; then + fail_item "tilegym wrapper not found for ${KERNEL_NAME}" +else + ok "wrapper: $(basename "${WRAPPER}")" + + if grep -q 'register_impl.*cutile-rs' "${WRAPPER}"; then + ok "@register_impl(cutile-rs)" + else + fail_item "wrapper missing @register_impl(..., backend='cutile-rs')" + fi + + if [ -n "$FFI_FN" ]; then + call_info="$(wrapper_call_info "${WRAPPER}" "${FFI_FN}")" + call_sites=0 + first_call=0 + if [ -n "$call_info" ]; then + read -r call_sites first_call <<< "$call_info" + fi + + if [ "${call_sites:-0}" -gt 0 ]; then + ok "wrapper executes ${FFI_FN} (${call_sites} live call site(s))" + else + fail_item "wrapper never executes ${FFI_FN} - dead code / fallback" + fi + + if [ "${first_call:-0}" -gt 1 ]; then + early=$(head -n "$((first_call - 1))" "${WRAPPER}" \ + | grep -cE '^\s*(return |raise )' 2>/dev/null) || early=0 + if [ "${early}" -gt 0 ]; then + warn_item "${early} early return/raise before FFI call - review for unsupported-case gates, not fallback" + fi + fi + fi + + wrapper_mod=$(basename "${WRAPPER}" .py) + if grep -q "from \. import ${wrapper_mod}" "${TILEGYM}/src/tilegym/ops/cutile_rs/__init__.py" 2>/dev/null; then + ok "__init__.py imports ${wrapper_mod}" + else + fail_item "__init__.py missing 'from . import ${wrapper_mod}'" + fi +fi + +echo "--- tilegym: test ---" + +TEST_FILE="" +if [ -n "$WRAPPER" ] && [ -f "$WRAPPER" ]; then + op_name=$(sed -n 's/.*register_impl("\([^"]*\)".*/\1/p' "${WRAPPER}" 2>/dev/null | head -1) + if [ -n "$op_name" ]; then + for candidate in "${op_name}" "${op_name//_act_/_activation_}" "${op_name}_kernel"; do + f="${TILEGYM}/tests/ops/test_${candidate}.py" + [ -f "$f" ] && TEST_FILE="$f" && break + done + fi +fi +if [ -z "$TEST_FILE" ]; then + short="${KERNEL_NAME%_kernel}" + for candidate in "${short}" "${short//_act_/_activation_}" "${short//_/}"; do + f="${TILEGYM}/tests/ops/test_${candidate}.py" + [ -f "$f" ] && TEST_FILE="$f" && break + done +fi + +if [ -z "$TEST_FILE" ] || [ ! -f "$TEST_FILE" ]; then + fail_item "test file not found" +else + if grep -q 'cutile.rs\|cutile-rs' "${TEST_FILE}"; then + ok "test: $(basename "${TEST_FILE}") has cutile-rs" + else + fail_item "$(basename "${TEST_FILE}") missing cutile-rs backend" + fi +fi + +echo "===" +echo "Result: ${fail} fail, ${warn} warn" +if [ ${fail} -eq 0 ]; then + echo "PASS" + exit 0 +else + echo "FAIL" + exit 1 +fi diff --git a/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_perf_log.sh b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_perf_log.sh new file mode 100755 index 0000000..4d38be4 --- /dev/null +++ b/skills/tilegym-converting-cutile-triton-to-cutile-rs/scripts/validate_perf_log.sh @@ -0,0 +1,52 @@ +#!/bin/bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# +# SPDX-License-Identifier: CC-BY-4.0 AND Apache-2.0 + +# Validate a --print-record perf log file. +# Checks: file exists, non-empty, has "Performance Test Results" markers, +# count matches PASSED test count. +# +# Usage: bash validate_perf_log.sh [