|
1 | 1 | [ |
2 | 2 | { |
3 | | - "id": "01-convert-matmul-cutile-to-triton", |
4 | | - "question": "Convert the following cuTile matmul kernel (which uses `ct.mma` and `ct.load`/`ct.store` with TMA-style tile loads) to Triton (`@triton.jit`). Place the Triton implementation in `triton_matmul.py` alongside the cuTile original. After conversion, run the relevant tests to confirm correctness on the current GPU.\n\nKey requirements:\n- Use `tl.make_tensor_descriptor` for 2D+ block loads (do NOT use raw `tl.load(ptr+offs, mask=...)` for full tiles — this is the most common source of large performance regressions).\n- Use Triton's bracket launch syntax `kernel[grid](args)` (with normal ASCII brackets in code).\n- If there is a `transpose` flag in the cuTile source, follow the dual-kernel + META-grid pattern.\n- Don't modify the original cuTile kernel.", |
| 3 | + "id": "converting-cutile-to-triton-001-conversion-checklist", |
| 4 | + "question": "I'm converting a CuTile `@ct.kernel` matmul to Triton (`@triton.jit`). List the conversion checkpoints I need to handle — name each cuTile → Triton mapping and one-line what changes. Don't write a full runnable kernel.", |
5 | 5 | "expected_skill": "converting-cutile-to-triton", |
6 | 6 | "expected_script": null, |
7 | | - "ground_truth": "Agent follows the phase-gated workflow (analyze → convert → validate → test → benchmark) from `translations/workflow.md`. Agent uses `tl.make_tensor_descriptor` for 2D+ tile loads (TMA), maps `ct.mma` to `tl.dot`, uses bracket-launch syntax, and runs `pytest tests/ops/test_matmul.py -k triton -vs` to verify correctness before declaring done. If a `transpose` flag exists, agent emits two kernels with META-grid selection.", |
| 7 | + "ground_truth": "The agent follows the phase-gated workflow (analyze → convert → validate → test → benchmark) and lists the conversion checkpoints without a full kernel. The list covers: (1) Pre-flight: count `@ct.kernel` definitions and find `ct.load`/`ct.store`, `ct.launch`, `ct.Constant`, layout flags. (2) `ct.load(ptr, index=(i,j), shape=(M,N))` for 2D+ block-shaped tile loads → MUST use `tl.make_tensor_descriptor` (TMA), NOT raw `tl.load(ptr+offs, mask=...)` — skipping TMA is the most common source of 10-50× regressions. (3) `ct.mma(acc, a, b)` → `tl.dot(a, b, acc)` with explicit accumulator. (4) `ct.Constant[int]` kernel args → Triton `tl.constexpr` parameters. (5) `ct.launch(stream, grid, kernel, args)` → Triton bracket-launch syntax `kernel[grid](args)` (ASCII brackets in code). (6) Layout flags (`transpose=True/False`) → dual-kernel + META-grid pattern in `translations/advanced-patterns.md`, NOT one kernel + `tl.trans`. (7) For attention / FMHA / Gemma / GQA / soft-cap / sliding-window kernels: read `references/optimization-strategy.md` §4 and apply the Gemma FMHA checklist before treating conversion as optimized. (8) Validate with `pytest tests/ops/test_<op>.py -k triton -vs` before benchmarking. The agent does NOT modify the original cuTile kernel.", |
8 | 8 | "expected_behavior": [ |
9 | | - "Agent reads the converting-cutile-to-triton SKILL.md and the `references/api-mapping.md`", |
10 | | - "Agent runs the Pre-flight Analysis grep commands on the cuTile source (count `@ct.kernel`, find `ct.load`/`ct.store`, layout flags)", |
11 | | - "Agent uses `tl.make_tensor_descriptor` for 2D+ block-shaped tile loads (TMA), not raw pointer+mask loads", |
12 | | - "Agent uses Triton's bracket launch syntax `kernel[grid](args)` with normal ASCII brackets", |
13 | | - "Agent handles layout flags (e.g. `transpose=True/False`) via the dual-kernel + META-grid pattern documented in `translations/advanced-patterns.md` (if present in source)", |
14 | | - "Agent does NOT modify the original cuTile kernel file", |
15 | | - "Agent runs `pytest tests/ops/test_*.py -k triton -vs` (or equivalent) to validate before declaring done", |
16 | | - "If benchmarking shows 10-50× slowdown, agent reads the CRITICAL PERFORMANCE PATTERNS section before declaring done" |
17 | | - ] |
18 | | - }, |
19 | | - { |
20 | | - "id": "02-react-hooks-negative", |
21 | | - "question": "What is the difference between `useEffect` and `useLayoutEffect` in React? When would I use one over the other?", |
22 | | - "expected_skill": null, |
23 | | - "expected_script": null, |
24 | | - "should_trigger": false, |
25 | | - "ground_truth": "Agent explains that `useEffect` runs asynchronously after the browser paints, while `useLayoutEffect` runs synchronously before paint. Use `useLayoutEffect` only when you need to read layout from the DOM and synchronously re-render (e.g., measuring tooltip positioning). The converting-cutile-to-triton skill is NOT activated.", |
26 | | - "expected_behavior": [ |
27 | | - "The converting-cutile-to-triton skill is NOT loaded", |
28 | | - "Agent explains the React `useEffect` vs `useLayoutEffect` distinction", |
29 | | - "Agent does not mention cuTile, Triton, `@ct.kernel`, or `@triton.jit`", |
30 | | - "Agent does not run destructive commands" |
| 9 | + "Lists the conversion checkpoints in order following analyze → convert → validate → test → benchmark", |
| 10 | + "Names `tl.make_tensor_descriptor` (TMA) for 2D+ block-shaped tile loads — NOT raw `tl.load` with masks", |
| 11 | + "Maps `ct.mma` → `tl.dot` with explicit accumulator argument", |
| 12 | + "Maps `ct.Constant[int]` → `tl.constexpr` for kernel parameters", |
| 13 | + "Names Triton's bracket-launch syntax `kernel[grid](args)` with ASCII brackets", |
| 14 | + "Mentions the dual-kernel + META-grid pattern for `transpose` layout flags (per advanced-patterns.md)", |
| 15 | + "Does NOT modify the original cuTile kernel source", |
| 16 | + "Validates via `pytest tests/ops/test_<op>.py -k triton -vs` before benchmarking" |
31 | 17 | ] |
32 | 18 | } |
33 | 19 | ] |
0 commit comments