# Gpu Kernel Episode Loop > Run the evidence loop of one long-horizon GPU kernel optimization episode. Use this skill to reconstruct the incumbent, profile and localize a bottleneck, research progressively, plan one coherent direction, implement and repair, validate development correctness and performance, and record every decisive experiment in the episode journal. - Skill: `alibaba/gpu-kernel-episode-loop` (Agent Skill) - Install (CLI): `npx skillmds@latest add alibaba/gpu-kernel-episode-loop` - Raw SKILL.md: https://api.skillmd.com/api/skills/alibaba/gpu-kernel-episode-loop/raw - Safety review: pending - Works with: Claude Code, Claude.ai, OpenAI Codex - Category: Research & Search - Author: alibaba (https://skillmd.com/u/alibaba) - Updated: 2026-09-17 - Page: https://skillmd.com/skills/alibaba/gpu-kernel-episode-loop --- # GPU Kernel Episode Loop ## When to Use Use this skill when an orchestrator episode prompt hands you one optimization episode in an isolated Git worktree and points here for the evidence loop. It does not apply to the V0 baseline session (`gpu-kernel-baseline`) or to a workspace without an episode journal. ## Episode bindings The episode prompt supplies a concrete value for every ALL-CAPS bracketed name below. Substitute them before running any command; never invent a path. Lowercase bracketed names in the command examples are values you fill in from the campaign, or a choice among the listed alternatives, not bindings. | Placeholder | Meaning | | --- | --- | | `` | profile output directory for this episode | | `` | evidence draft for the canonical version | | `` | generated plan for the canonical version | | `` | backend-native plan generator invocation | | `` | episode journal command prefix | | `` | episode journal path, already shell-quoted | The episode prompt's ownership rules, execution boundary, mode policy, and framework-escalation directive outrank this skill. Where they conflict, follow the prompt. ## Telemetry Telemetry is best-effort and must not block engineering work. Mark phase boundaries with standalone commands and keep at most one phase active: ```bash python3 tools/iteration_trace.py phase-start python3 tools/iteration_trace.py phase-end python3 tools/iteration_trace.py source-read ``` Never put credentials, private URL parameters, absolute user paths, raw tool output, or transcript text into telemetry. ## Loop Repeat this evidence loop until the direction yields a mature candidate or is exhausted. The numbered steps map onto the telemetry phases above: `profile`, `research`, `planning`, `implementation`, `correctness`/`benchmark`, and `recording`. ### 1. Reconstruct the incumbent and choose a hypothesis Read the workspace goal, unmasked `memory/v*.json`, and prior plans/profiles. Prior-episode summaries are carried only by canonical memory and are not injected into the episode prompt. Identify attempted dead ends and open directions from those records, including each record's compact `experience.experiments`. For PPU, also inspect `profile_evidence.accepted_ppu_diagnostics` and reuse a conclusion only when its recorded identities remain comparable and none of its `invalidation_conditions` holds. Start with one falsifiable hypothesis tied to the current bottleneck. ### 2. Profile and localize For a PPU target, do not apply the NVIDIA/AMD default below. Read `skills/ppu-acu-joint-profile/SKILL.md` first and let its PPU-specific per-iteration rule decide whether new PPU profiler evidence is needed. That route may use source or compiler inspection, the probe-free benchmark, or still-valid PPU evidence instead of collecting a new profile. Reuse a profile only when it matches the current committed kernel. Otherwise profile through the sandbox using the vendor-appropriate tooling. Both wrappers run `python `, so the profiled file is the immutable `profile_driver.py` seeded next to `kernel.py` — never `kernel.py` itself, which the evaluator only ever imports: ```bash # NVIDIA python tools/sandbox.py --kind profile --sync -- \ bash tools/profile_nvidia.sh profile_driver.py --output-dir --source # AMD python tools/sandbox.py --kind profile --sync -- \ bash tools/profile_kernel.sh profile_driver.py --output-dir ``` `profile_driver.py` imports the current `kernel.py`, builds real inputs from the campaign contract (`definition.json` + `workload.jsonl`, a privately injected generalized Atrex-Bench real shape, or legacy `shapes.json` + `input.py`), warms up, and then invokes the candidate repeatedly. Select what it drives with environment variables rather than editing it — it is a protected path and a candidate that modifies it is rejected: ```bash PROFILE_ITERS=30 PROFILE_WORKLOAD_IDX=2 python tools/sandbox.py ... # SOL: one workload PROFILE_ITERS=30 PROFILE_SHAPE_ID=3 python tools/sandbox.py ... # generalized or legacy Atrex-Bench ``` When several shapes or workloads need profiling, run one sandbox command per id in waves of at most four concurrent jobs. Give every job its own `/shape-` sync/output directory and wait for the whole wave before starting the next one. For generalized Atrex-Bench tasks, choose `PROFILE_SHAPE_ID` from the previous canonical memory's complete opaque-id `performance.latency_us_by_shape` map. The sandbox privately resolves that id and injects only its real input case into the ephemeral remote profile job; the driver deletes the case JSON before importing candidate code. Profile the highest-cost ids and additional ids representing distinct latency regimes, but do not infer or reconstruct the complete hidden input table. Extract a concrete bottleneck and source-level target. Use PTX/SASS/TTGIR inspection when compiler lowering or instruction selection is part of the hypothesis. Do not make speculative optimization changes before obtaining usable evidence. When ordinary profiling has isolated one kernel but cannot distinguish a specific in-kernel timing hypothesis, read `skills/autonomous-gpu-kernel-timeline/SKILL.md` and run its autonomous loop. Use standalone CUDA/inline PTX through its CUDA backend and CuTe DSL through IKeT. Keep every attempt under `/timeline/attempt-N`; when the remote command reads backend files, pass that specific skill path with sandbox `--input` and sync only the attempt output directory. For PPU, use the routing and capture contracts in the PPU skill linked above. Timeline instrumentation is a temporary working snapshot on this episode's single HEAD line, not a candidate. Preserve the clean source and each useful instrumented source or reversible patch before replacing it. After the evidence answers the question, restore or rewrite a probe-free `kernel.py` before correctness/performance validation, commit, journal finalization, and handoff. Never submit a profiling snapshot as `candidate_commit`; its latency and failures do not count as promotion attempts or framework-stall events. Escalate through the typed profile funnel instead of collecting everything at once: `--profile-level survey` to enumerate kernels, `sol` (the default) for the bottleneck class, and `deep --kernel-regex '^$'` for one named kernel, especially a Triton `@triton.jit` entry. Take that name verbatim from the survey/SOL result; never guess a substring. Raw `.ncu-rep`/ATT artifacts stay remote unless `--include-raw-profile` is justified. On NVIDIA, `summary.txt` carries a `LOCALIZE` line naming the analysis files that pin a symptom to source lines. Those files exist only on a `--source` run: never pin a source-level claim to a profile collected without it. #### When the seeded driver cannot represent the work Build a local fallback driver at `/harness/profile_driver.py` and profile that file instead when the seeded driver cannot express the case — a multi-kernel sequence, a new synthetic case inside the public domain, or a driver that needs sibling helper modules. It must import `kernel.py` plus the immutable input module, select a representative workload, warm up, invoke the entry point repeatedly, and never write memory files. Because it lives below `/harness/`, Python does not put the workspace root on its import path; add it before importing anything local: ```python import sys from pathlib import Path WORKSPACE_ROOT = Path(__file__).resolve().parents[3] if str(WORKSPACE_ROOT) not in sys.path: sys.path.insert(0, str(WORKSPACE_ROOT)) ``` The sandbox uploads that whole `harness/` directory automatically, so sibling helpers need no extra flags. Only a file opened dynamically by command code needs a repeatable `--input ` option before `--`, which routes the job through the dev interface. ### 3. Research progressively Search in this order and stop when one actionable direction is supported: 1. **GPU Wiki through the natural-language front door.** Profile first, then describe the measured problem rather than trying to guess query flags. PPU is the exception: start from the decisive evidence selected by `ppu-acu-joint-profile`, whether or not it required a new profile: ```bash python3 gpu-wiki/tools/query_nl.py "" --brief python3 gpu-wiki/tools/query_nl.py --file research_request.txt ``` Include the true target product and authoritative runtime architecture exactly as supplied. Ask for the full product specification and relevant architecture/ISA facts so the response contains isolated `hardware_wiki` and `kernel_wiki` records. Also include the operator, framework, shapes, dtypes, profile numbers, what was already tried, exact failures, competing hypotheses, and the fact that would end this line of work. For a PPU iteration that skipped profiling, identify the source, compiler, clean-benchmark, or prior PPU evidence used instead. Do not translate the hardware identity or pre-compress the prose into keywords. Read the compact response before acting: records are keyed by stable id, every `payload` is isolated, `store` distinguishes `gpu_wiki` from namespaced `internal_gpu_wiki` records, `match.arch` states its reach, and `notes` reports deterministic normalization, widening, truncation, or store gaps. Pass `--exclude ` on later queries and use `--max-bytes` for a hard context bound. The structured `query_wiki.py` and `query_hardware.py` tools remain available when the exact address is already known; never drop architecture scope to manufacture a match. Copy the response's top-level `query_id` and each used record's own emitted canonical `wiki_id` exactly so the decisive experiment can declare its Wiki attribution in the native journal. Never reconstruct them from response mapping keys or prose. 2. `reference-projects/` only when the local wiki is insufficient. 3. Public primary sources only when local sources do not answer the question. After repeated rejected episodes, expand across DSLs targeting the same architecture instead of repeating local parameter tweaks. Record stable Wiki ids and the evidence-to-action chain. ### 4. Plan a coherent direction Write or update `` with profile evidence, research findings, concrete edits, risks, rollback points, and measurable acceptance criteria. For a PPU iteration that did not need a new profile, record the decisive PPU evidence selected by its routing skill instead. Then produce `` with the backend-native plan generator ``. The episode may contain multiple related experiments, but they must advance one coherent engineering direction. Checkpoint useful intermediate states so failed sub-steps can be reverted without losing the whole direction. ### 5. Implement and repair Modify only candidate source/metadata files allowed by policy. Compile and probe through the sandbox. On compile or correctness failure, diagnose and repair while the direction remains viable. Do not publish an intermediate checkpoint as a candidate. Land one optimization category per edit — vectorized load, swizzle, double buffering, tiling change, and so on — and attribute each edit as `evidence -> inference -> action`. Do not mix unrelated refactors, formatting, or cleanup into the same change: a bundled edit makes a regression unattributable. When the evidence localizes a symptom to specific lines, change those lines only. ### 6. Development correctness and performance Use the immutable evaluator for development measurements: ```bash python tools/sandbox.py --kind run --no-sync -- \ python test_kernel.py --version vlong --no-memory python tools/sandbox.py --kind run --no-sync -- \ python test_kernel.py --version vlong --multi-seed 5 --no-memory ``` All workloads and all additional seeds must pass. Never depend on tensor values, pointer identity, cached outputs, evaluator ordering, or hidden workload IDs. Shape/dtype/layout dispatch is allowed, and pre-converting stable weights (transpose, contiguous) is allowed because weights do not change during evaluation. Before trusting a large delta — especially a regression beyond roughly 30% — re-run the same command on the same sandbox hardware and compare. GPU selection belongs to the gateway; never set a local `CUDA_VISIBLE_DEVICES` to steer it. Repeated development measurements are not promotion authority; the supervisor reruns incumbent and candidate in one ABBA allocation. ### 7. Record every decisive experiment immediately Immediately after each decisive experiment, append it to the single episode journal. Do not batch these writes at the end of the episode: every append refreshes the non-canonical `memory/live.json` progress view in the incumbent workspace. ```bash append --path \ --experiment-json '{"name":"...","hypothesis":"...","change":"...","evidence":"...","result":"...","evaluation":{"correctness":"pass|fail|unknown","performance":"improved|not_improved|unknown","latency_us":null,"kernel_hash":""},"decision":"keep_as_best|promote|reject_and_continue|revert|pivot|blocked","wiki_usage_status":"declared","wiki_query_ids":[""],"wiki_usage":[{"query_id":"","wiki_id":"","disposition":"applied|partially_applied|reference_only|rejected","use":"...","evidence":"..."}]}' ``` Use `declared` only with non-empty `wiki_usage`. Include `wiki_query_ids` for both `declared` and `no_material_use`; omit both arrays for `not_queried`. Do not invent ids and do not collapse repeated use across experiments. Invalid Wiki rows are omitted with `wiki_usage_errors`; this diagnostic field never blocks the experiment or handoff. ## Leaving the loop Leave the loop as soon as one coherent candidate passes the full development correctness check and has credible performance evidence, or as soon as the direction is exhausted or blocked. Then follow the episode prompt's terminal contract for finalizing the journal and publishing the handoff. For a PPU full episode, include `outcome.accepted_ppu_diagnostics` using the schema in `skills/ppu-acu-joint-profile/SKILL.md`; retain only evidence that still applies to the terminal probe-free kernel. Each retained row must bind an accepted decision-grade artifact by path, SHA-256, schema, and evidence id. This is optional when no reusable PPU profiler evidence exists.