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.

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:

python3 tools/iteration_trace.py phase-start <profile|research|planning|implementation|correctness|benchmark|recording>
python3 tools/iteration_trace.py phase-end <profile|research|planning|implementation|correctness|benchmark|recording>
python3 tools/iteration_trace.py source-read <gpu_wiki|reference_projects|workspace|public_web> <safe-relative-reference>

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 <file>, 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:

# NVIDIA
python tools/sandbox.py --kind profile --sync <PROFILE_DIR> -- \
  bash tools/profile_nvidia.sh profile_driver.py --output-dir <PROFILE_DIR> --source

# AMD
python tools/sandbox.py --kind profile --sync <PROFILE_DIR> -- \
  bash tools/profile_kernel.sh profile_driver.py --output-dir <PROFILE_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:

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 <PROFILE_DIR>/shape-<index> 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 <PROFILE_DIR>/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 '^<exact_base_function_name>$' 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 <PROFILE_DIR>/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 <PROFILE_DIR>/harness/, Python does not put the workspace root on its import path; add it before importing anything local:

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 <relative-path> 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:

   python3 gpu-wiki/tools/query_nl.py "<your description>" --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 <ids-already-read> 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 <PLAN_DRAFT> 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 <PLAN_FILE> with the backend-native plan generator <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:

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.

<JOURNAL_CLI> append --path <JOURNAL_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":["<emitted-query-id>"],"wiki_usage":[{"query_id":"<emitted-query-id>","wiki_id":"<emitted-canonical-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.