[WIP] support forge agent

.gitignore

@@ -25,6 +25,11 @@ kill.sh
 saved_results
 .mcp.json
 
+# Local forge-agent run scaffolding: contains gateway creds / regenerated per run
+run_aka_forge_example.sh
+config.forge_example.yaml
+run_logs/
+
 # Generated held-out test data (methodology is in held_out/, data is private)
 held_out_tests/
 

agents/forge/__init__.py

@@ -0,0 +1,2 @@
+# Copyright(C) [2026] Advanced Micro Devices, Inc. All rights reserved.
+from .launch_agent import launch_agent  # noqa: F401

agents/forge/agent_config.yaml

@@ -0,0 +1,46 @@
+# Configuration for the KernelForge "forge" agent.
+#
+# This agent bridges AgentKernelArena to KernelForge's `kernel-agents forge-loop`:
+# an autonomous optimization loop (baseline -> agent edit -> 5-stage validate ->
+# bench -> keep/revert) running as a hard-killable subprocess.
+
+# Model passed to KernelForge (exported as KERNEL_AGENTS_MODEL). Must be served
+# by the configured Anthropic gateway (ANTHROPIC_BASE_URL/ANTHROPIC_AUTH_TOKEN).
+model: claude-opus-4-8
+
+# Backend expertise (fellow) is INFERRED by the launcher, per task family:
+#   * repository / image_kernel tasks -> from `repository_language`
+#       hip -> hip-fellow, triton -> triton-fellow
+#       (the source repo like aiter/sglang is NOT the fellow: the aiter fellow
+#        only integrates prebuilt ops, so an aiter-repo HIP kernel uses hip)
+#   * snippet "<src>2<dst>" tasks     -> the target side picks the fellow
+#       triton2triton / instruction2triton -> triton-fellow
+#       hip2hip / cuda2hip / torch2hip      -> hip-fellow
+#       flydsl2flydsl                       -> flydsl-fellow
+# (valid: ck, flydsl, triton, aiter, hip, hipblaslt). To force a specific
+# fellow regardless of the task, uncomment:
+# fellow: triton-fellow
+
+# Loop budget: max_iters is intentionally high so a forge run is bounded by TIME
+# (see timeout_seconds below), not by an iteration count.
+max_iters: 1000
+
+# Correctness gate (dB) used by the 5-stage validation pipeline.
+snr_threshold: 30.0
+
+# The single time budget for a forge run: a hard wall-clock cap (seconds) on the
+# forge-loop subprocess. The loop's own --max-hours is derived from this
+# (timeout_seconds / 3600 minus a ~15min margin) so the loop self-stops just
+# before this hard kill. NOTE: for API-launched runs the bootstrap overwrites
+# this per-run (e.g. 115200 for a 32h run); the value below is only the local
+# fallback default.
+timeout_seconds: 7200
+
+# Claude permission mode. acceptEdits works under root; bypassPermissions is
+# blocked by Claude Code's security policy for root.
+permission_mode: acceptEdits
+
+# NOTE: no shapes_json here on purpose. The generic adapter ignores --shape (the
+# task's own pytest owns its shapes), so passing kernel-specific n= values adds
+# no value and isn't portable across tasks. The forge-loop CLI defaults to "{}".
+# Only set shapes_json if you use a per-kernel driver that parses --shape.

agents/forge/drivers/arena_task_adapter.py

@@ -0,0 +1,114 @@
+# Copyright(C) [2026] Advanced Micro Devices, Inc. All rights reserved.
+"""Generic KernelForge driver that reuses an AgentKernelArena task's own
+correctness and performance measurement.
+
+Instead of re-implementing (or hardcoding) precision/perf parsing per kernel,
+this adapter calls Arena's OWN evaluation functions and translates their results
+into the KernelForge driver contract consumed by
+``kernel_agents.mcp_server.tools.{test,bench}``:
+
+  * correctness mode (default / --mode smoke|stability|determinism):
+        src.evaluator.evaluate_correctness(...) -> "allclose: True/False"
+  * bench mode (--bench-mode):
+        src.performance.measure_performance(...) -> "mean_ms: <mean-of-cases>"
+        (arithmetic mean across test cases, matching Arena's own evaluator
+        aggregation — deliberately a MEAN, not a median, so the label is honest)
+
+Reusing Arena's functions means the adapter is task-type aware and not tied to a
+single report filename: Arena already searches multiple candidate report files
+(``build/performance_report.json``, ``perf/benchmark_results.json``, ...) and
+falls back to stdout parsing, per task type. So this one driver works for any
+task Arena itself can score — no per-kernel driver and no hardcoded filename.
+
+This module is a LIBRARY: it takes the task paths (workspace / task config /
+Arena repo root) as explicit arguments — NO environment variables. The Arena
+forge launcher generates a tiny ``forge_driver.py`` shim with those paths baked
+in that calls :func:`run`.
+"""
+
+import argparse
+import statistics
+import sys
+from pathlib import Path
+
+import yaml
+
+
+def _load_task_config(task_config: str) -> dict:
+    with open(task_config, "r") as f:
+        return yaml.safe_load(f) or {}
+
+
+def do_correctness(workspace: str, task_config: str, arena_root: str = "") -> int:
+    """Reuse Arena's correctness evaluation; emit the forge allclose contract."""
+    try:
+        from src.evaluator import evaluate_correctness
+    except Exception as e:  # noqa: BLE001
+        print("allclose: False")
+        print(f"error: cannot import Arena evaluator (arena_root={arena_root!r}): {e}")
+        return 0
+
+    task_config_data = _load_task_config(task_config)
+    passed, err = evaluate_correctness(Path(workspace), task_config_data)
+    if passed:
+        print("allclose: True")
+    else:
+        print("allclose: False")
+        if err:
+            print(err[-1500:])
+    return 0
+
+
+def do_bench(workspace: str, task_config: str, arena_root: str = "") -> int:
+    """Reuse Arena's performance measurement; emit the forge median_ms contract."""
+    try:
+        from src.performance import measure_performance
+    except Exception as e:  # noqa: BLE001
+        print(f"error: cannot import Arena performance module (arena_root={arena_root!r}): {e}")
+        return 1
+
+    task_config_data = _load_task_config(task_config)
+    # is_baseline=False: bench whatever kernel is currently in the tree (forge
+    # benches the pristine kernel and each edited version the same way).
+    cases = measure_performance(Path(workspace), task_config_data, is_baseline=False)
+    times = [c.execution_time_ms for c in cases
+             if getattr(c, "execution_time_ms", None) and c.execution_time_ms > 0]
+    if not times:
+        print("error: Arena measure_performance returned no usable timing")
+        return 1
+
+    # Aggregate the per-case times into the single wall time the forge-loop uses
+    # for its keep/revert decision. This is the ARITHMETIC MEAN across cases,
+    # matching Arena's own evaluator (sum/len over cases), so forge optimizes the
+    # same quantity Arena scores. Reported as ``mean_ms:`` (not ``median_ms:``)
+    # so the label reflects the statistic actually computed.
+    agg = statistics.mean(times)
+    print(f"mean_ms: {agg:.6f}")
+    return 0
+
+
+def run(workspace: str, task_config: str, arena_root: str = "", argv: list[str] | None = None) -> int:
+    """Driver entry point. All task paths are passed explicitly (no env).
+
+    Args:
+        workspace: dir where the task commands run (the kernel lives here).
+        task_config: path to the task's config.yaml.
+        arena_root: AgentKernelArena repo root, prepended to sys.path so Arena's
+            ``src`` package (evaluator / performance) is importable.
+        argv: driver args from KernelForge (``--shape`` / ``--mode`` /
+            ``--bench-mode`` / ...); defaults to ``sys.argv[1:]``.
+    """
+    if arena_root and arena_root not in sys.path:
+        sys.path.insert(0, arena_root)
+
+    ap = argparse.ArgumentParser()
+    ap.add_argument("--shape", default="default")   # task harness owns its shapes
+    ap.add_argument("--mode", default="full")        # all modes -> task correctness
+    ap.add_argument("--bench-mode", action="store_true")
+    ap.add_argument("--warmup", type=int, default=10)
+    ap.add_argument("--iters", type=int, default=30)
+    args, _unknown = ap.parse_known_args(argv)
+
+    if args.bench_mode:
+        return do_bench(workspace, task_config, arena_root)
+    return do_correctness(workspace, task_config, arena_root)