Skip to content

Claude Code version - #3

Open
bill810975 wants to merge 14 commits into
mainfrom
cc
Open

bill810975 wants to merge 14 commits into
mainfrom
cc

Conversation

@bill810975

Copy link
Copy Markdown
Collaborator

all previous api mode code is in api/

Muheng Li and others added 14 commits May 28, 2026 17:40
Move from the motus API-driven ReActAgent loop (orchestrator.py +
agents.py + main.py) to a Claude-Code subscription path where the
mainthread runs as a `claude` process per workspace and delegates
specialized work to scoped subagents under .claude/agents/.

Architecture:
  - mainthread reads the rewritten CLAUDE.md and operates on one
    workspaceN/ identified by $FERRET_WORKSPACE.
  - .claude/agents/{iterator,planner,profiler,reviewer,codex-dispatcher,
    memory-keeper}.md — each with its own tools allowlist enforcing
    read/write scopes (iterator is Read-only, memory-keeper writes only
    under docs/dev-memory/, reviewer's Task is restricted in prose to
    codex-dispatcher + memory-keeper).
  - docs/dev-memory/ (gitignored) is host-shared knowledge; its tracked
    template lives in docs/dev-memory-seed/ and is bootstrapped by
    cc-init.sh / cc-run.sh on first launch.
  - profile.py — sync `python -m ferret.profile` CLI wrapping
    tools/profiler.py with TMPDIR=/tmp/$USER, pick_gpu.sh,
    .profile_last.json delta — so the profiler subagent doesn't
    hand-craft ncu invocations.
  - scripts/cc-init.sh / scripts/cc-run.sh — new launchers; cc-run.sh
    exports FERRET_WORKSPACE, FERRET_ROOT, PYTHONPATH (so `python -m
    ferret.state` works from inside ferret/), TMPDIR, and picks a GPU
    once per session.

The legacy motus path (orchestrator.py, agents.py, main.py, prompts.py,
cost_tracker.py, scripts/run.sh) is untouched and kept for parity.

Smoke-tested on cc branch: cold-start of workspace1 via headless
`claude -p ...` correctly read CLAUDE.md, dispatched to planner,
which read $MIRAGE_ROOT/include/mirage/persistent_kernel/tasks/
blackwell/mla_mtp_decode_sm100.cuh, extracted the actual
__device__ __noinline__ signatures, and wrote a 6-step Plan +
Untried list into progress.md without touching kernel.cu.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…o Mirage

Chain validation passed end-to-end on workspace1 (planner → 3×
write→compile→run→commit→tag→reviewer→codex-dispatcher; OVERALL_VERDICT
PASS). The codex_timeout graceful-degrade path also fired and was
recorded correctly. Convergence-triggered kernel-extractor was
validated in a follow-up micro-test: extractor read sibling Mirage
task header `mla_mtp_decode_sm100.cuh`, wrote a structurally-correct
`workspace1/kernel.cuh` (pragma once + namespace + __device__
__noinline__ + no host artifacts), and surfaced sanity-compile result
correctly (stub source → flagged FAIL with TODO header).

Concrete changes:

  Deleted (API/motus path, replaced by Claude-Code mainthread):
    orchestrator.py, agents.py, main.py, prompts.py, cost_tracker.py,
    scripts/run.sh, tools/compiler.py, tools/doc_loader.py

  Added:
    .claude/agents/kernel-extractor.md — convergence-only subagent,
      reads kernel.cu + sibling Mirage task header, writes Mirage-ready
      kernel.cuh (device function header). Refuses unless state CLI
      reports advance? True AND every config ✓ AND latest Review's API
      check is not FAIL.
    cc_goal.py — renders concrete /goal text from task.yaml (SOTA name
      + per-config target ratios + scoring policy) so the standing
      goal carries numeric success criteria, not abstract "iterate
      until done".

  Modified:
    .claude/agents/reviewer.md — new step 5 dispatches
      kernel-extractor at convergence; Review block grows a
      `convergence:` line; Task allowlist now includes
      kernel-extractor.
    .claude/agents/codex-dispatcher.md — timeout bumped from ~90s to
      300s (chain test showed 90s clips real verdicts).
    CLAUDE.md — §6.5 documents the dual deliverable (kernel.cu
      standalone + kernel.cuh Mirage-ready); file-ownership table
      shows kernel-extractor owns kernel.cuh; subagent routing table
      lists kernel-extractor as reviewer-invoked at convergence; §10
      no longer references the deleted motus path.
    scripts/cc-run.sh — default goal now generated by `ferret.cc_goal`
      (with hand-written fallback); injection via both
      --append-system-prompt and a leading /goal slash command.
    tools/__init__.py — exports trimmed to profiler symbols only.
    pyproject.toml — drop `lithosai-motus` dep; bump to 0.2.0;
      description updated to reflect the seven subagents.
    README.md — full rewrite for the Claude-Code path; layout diagram
      lists every subagent; usage shows `cc-run.sh` not
      `python -m ferret.main`.

Also added at ~/mirage/.claude/agents/ferret-kernel-agent.md (commit
separately in the mirage repo): a dispatcher subagent that translates
a kernel requirement into a ferret task.yaml, picks a free workspace,
launches cc-run.sh in background, monitors, and reports back the
delivered kernel.cuh path + per-config TFLOPS + Mirage ABI verdict.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Mirage's .gitignore line 85 excludes .claude/, so the dispatcher
subagent can't be shipped inside the Mirage repo. ferret holds the
canonical template at integration/mirage/ferret-kernel-agent.md;
teammates copy it into ~/mirage/.claude/agents/ manually after each
ferret pull. This makes ferret the single source of truth for the
Mirage→ferret CLI contract — when cc-run.sh's surface changes, the
dispatcher description is updated here in the same commit.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…k contract, Codex-MCP dispatcher, clean public seed

- test-mode validation harness: mpk-validator subagent + scripts/mpk_validate.sh +
  templates/ — validate a candidate kernel.cuh through the real MPK compile pipeline
  on ONE exclusive GPU: correctness (PyTorch-ref cos + sentinel guard) AND single-kernel
  performance (profiler_tensor -> scripts/parse_profile.py duration_ns, vs the baseline kernel).
- task_spec.py: structured MPK-format I/O contract (IoTensor) + ReduceHint
  (internal_atomic / tma_reduce[preferred] / external_reduce) + optional target_latency_us.
- Codex cross-check migrated to the MCP protocol (.mcp.json + codex-dispatcher/reviewer):
  read-only, pre-feed, graceful codex_unavailable degradation.
- New ferret task specs: dense-fp8-gemm-decode-splitk, qkva-internal-splitk-fp8-decode,
  fp8-group-gemm-w2-compact-dispatch, dsv3-router-gemm-decode.
- dev-memory seed refreshed with shareable Mirage/CUDA dev knowledge; all host
  paths/usernames/hostnames/UUIDs generalized to \$MIRAGE_ROOT/\$FERRET_ROOT/\$USER/~.
- gitignore the local campaign log (docs/WORKFLOW_AUDIT.md, kept local), calib_scratch/, *.pid.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…edian

The pushed test-mode harness ranked kernels by per-task MEDIAN duration_ns, which on
under-occupied bimodal-CTA decode GEMMs lands on an IDLE CTA and inverts the verdict
(qkv_a split-K measured "0.06x / 16x slower" when it is actually 1.32x FASTER in-MPK).

- mpk_validate.sh: scrape wall-span (splitk_wall_us/mediumm_wall_us/ratio), back-compat
  fallback to legacy fields; verdict relabeled WALL-SPAN.
- mpk-validator.md: contract now requires a WALL-SPAN number (parse_profile.py --stat
  wall) + documents the bimodal-CTA "median is an idle CTA" pitfall; never rank by median.
- templates/README.md: PERFORMANCE section rewritten around wall-span + the pitfall.

Depends on mirage scripts/parse_profile.py --stat wall (max_end-min_begin, 32-bit wrap
corrected). Validated on the qkv_a trace CSVs: split-K 22.27us vs mediumm 29.31us = 1.32x.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
… + auto-update script

Merge the KernelWiki SOTA-kernel knowledge base into ferret as a proper git
submodule (was a loose ~/KernelWiki clone referenced by absolute path), and add
an automated refresh script so the corpus stays current as upstream repos merge
new kernel PRs.

- Submodule: resources/kernelwiki -> https://github.com/mit-han-lab/KernelWiki
  (pinned d1536a0). Populate on fresh clone with
  `git submodule update --init resources/kernelwiki`.
- scripts/update_kernelwiki.sh: two channels, both safe to run from cron and
  graceful when offline/unauthed (the offline read path never breaks):
    (A) upstream sync  — ff-only pull of mit-han-lab master (pushable).
    (B) local refresh  — gh-ingest of newly-merged kernel PRs since cutoff
        (refresh_candidate_ledger -> generate-pr-pages -> fetch_pr_diff ->
        generate-indices), auto-skips with an install hint if `gh` is absent.
  Always finishes with an offline validate + page-count report; flock
  single-flight; --upstream-only / --refresh-only / --repos / --commit-pointer.
- Repointed the skill/planner/doc references from the hardcoded
  /home/muhengl/KernelWiki to the portable ${FERRET_ROOT:-$HOME/ferret}/
  resources/kernelwiki (no machine-path leak in the shared repo).
- .gitignore: track the committed kernelwiki skill (rest of .claude/skills/
  stays local); ignore runtime logs/.

Verified (3-agent adversarial audit): zero stale absolute-path refs, the 5
pipeline CLIs match KernelWiki's actual argv handling, the other 10 submodules
untouched, offline query.py/get_page.py work from the new location, and the
two issues the audit caught (unguarded --repos under set -u; .gitmodules staged
with the local-path url) are fixed.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The planner already queries KernelWiki at cold-start (pick starting template +
anchor target_ratio), but the iterator — the agent that loops on kernel.cu —
never referenced it, so the skill's stated "query by symptom on stall" behavior
was documented but never triggered mid-run. Add the hook to the iterator's
OPTIMIZE-stage tools (ON STALL ONLY, to avoid per-iteration prompt bloat) and to
its stall-handling note, with the M=1/skinny-M + perf-claim-is-a-ceiling caveats.
Now KernelWiki is proactively consulted at BOTH workflow points: planner
cold-start and iterator stall.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…ncy invariants)

- §9 Standing references: add resources/kernelwiki — the SOTA-kernel prior-art
  corpus — and state explicitly that the planner queries it at cold-start and the
  iterator on stall, with the M=1/skinny-M and perf_claim-is-a-ceiling caveats and
  a pointer to scripts/update_kernelwiki.sh. Closes the loop so the mainthread
  knows the resource exists and that its two subagents touch it.
- Also lands the previously-uncommitted "Mirage megakernel co-residency invariants
  (HARD — for any tcgen05 kernel)" §, the same-warp-alloc + minimize-TMEM-cols
  lesson from the FP8 split-K root-cause.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…ream-tracking

Preserve the pre-da135d5 programmatic/API form under a new api/ subpackage so it
isn't lost, kept OUT of the Claude-Code agent's context. The agent (cc-run.sh +
.claude/agents subagents) IGNORES api/.

- api/: restored the 8 deleted files (orchestrator/agents/main/prompts/
  cost_tracker + tools/{compiler,doc_loader} + scripts/run.sh) + __init__/README.
  Relative imports fixed for the new location (root modules task_spec/state/
  tools.profiler via `..`, api-internal via `.`); AGENT_ROOT goes up two levels;
  run.sh launches `python -m ferret.api.main` from ferret's parent; masked the
  hardcoded python path to ${PYTHON:-python3}.
- CLAUDE.md: "IGNORE api/" note for the mainthread; KernelWiki note made explicit
  (submodule tracking upstream mit-han-lab/KernelWiki, content NOT vendored; run
  `git submodule update --init resources/kernelwiki && bash
  scripts/update_kernelwiki.sh` after every clone/pull).
- .gitmodules: kernelwiki branch=master (track upstream).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…ve original API-form README to api/

- README.md (root): reframed around how ferret is actually invoked — dispatched
  as its OWN headless Claude-Code session per workspace, and from Mirage via the
  `ferret-kernel-agent` subagent (linked). Functional content (install/usage/
  cc-run.sh/task authoring/layout/loop/submodule maint) preserved; added the
  `api/` = preserved-API-form-the-agent-ignores note + the kernelwiki
  submodule-update step.
- api/README.md: the ORIGINAL (v0.1 API-form / motus) README, preserved under
  api/ where that code now lives, with launch paths updated to
  `python -m ferret.api.main` / `api/scripts/run.sh` and the layout pointed at api/.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…; keep the api/ design

Resolves the cc->main divergence so the PR merges cleanly:
- main.py (modify/delete) → resolved via rename detection: cc had relocated
  main.py to api/main.py, so main's --remote-host addition auto-merged onto
  api/main.py. The new remote.py was moved to api/remote.py (stdlib-only;
  api/main.py imports `.remote`), keeping the API-mode ssh+rsync routing in its
  natural home (api/, which the Claude-Code agent ignores).
- README.md → combined: main's hook + verified-wins table, plus cc's
  dispatch-model section (the 2 invocation points + the Mirage ferret-kernel-agent
  link), CC-mode install/usage/loop/layout, and the "api/ is preserved + ignored"
  note. Both audiences preserved.
- Brought in main's Qwen3-8B GQA kernel + tasks/gqa-decode-qwen3-8b.yaml +
  docs/assets (hero chart) cleanly.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
… compile+benchmark to a remote GPU

So a ferret Claude-Code session whose box has no/busy GPU runs its own nvcc+./kernel
on a remote host (ssh+rsync) instead of stopping and bouncing back to a Mirage session.

- scripts/remote_run.sh: rsync workspace → remote (same abs path, .git stays local)
  → run the cmd under a REMOTE `timeout` via ssh `bash -s` → rsync back. Transparent
  LOCAL fallback when FERRET_REMOTE_HOST is unset (exec bash -c). Remote cmd stdout
  (KERNEL_RESULT lines) forwarded clean; rsync/diag chatter → stderr; exit = remote rc.
  $CMD's inner $FERRET_WORKSPACE/quotes survive verbatim (single-pass heredoc + quoted
  inner heredoc) and expand on the remote. Codex-audited (rc/quote/stdout) + locally
  simulated (rc + env-expansion + clean stdout verified).
- CLAUDE.md §5: tells the mainthread to wrap compile+benchmark in remote_run.sh when
  FERRET_REMOTE_HOST is set (combine compile+run in one call; remote GPU =
  FERRET_REMOTE_CUDA_DEVICES, not pick_gpu.sh).
- cc-run.sh: passthrough FERRET_REMOTE_HOST / _CUDA_DEVICES / _ENV so a launcher can
  point a session at a remote box.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
@@ -0,0 +1,43 @@
# dev-memory — shared machine knowledge across workspaces

This directory is **gitignored**. It holds operational knowledge that

Copy link
Copy Markdown
Owner

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

is it supposed to be ignored?

Comment thread scripts/cc-init.sh

Copy link
Copy Markdown
Owner

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

do you want to mv the cc related scripts to .claude?

Comment thread scripts/cc-run.sh

Copy link
Copy Markdown
Owner

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

same

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants