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.

Episode mode

Use the Python-supplied ATREX_EPISODE_MODE (default: full), as defined in skills/gen-plan/SKILL.md. Fast/full episodes advance one coherent engineering direction. Goal episodes can repeat this loop across a roadmap of directions, preserving the best validated checkpoint until the roadmap is complete or exhausted.

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 (fast/full) or roadmap (goal) 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.