fix(kernel): unblock source resolution, the FlyDSL rewrite route, and the KERNEL idle guard - #1335
Open
ZhengGong-amd wants to merge 21 commits into
Open
fix(kernel): unblock source resolution, the FlyDSL rewrite route, and the KERNEL idle guard#1335ZhengGong-amd wants to merge 21 commits into
ZhengGong-amd wants to merge 21 commits into
Conversation
The scratch worktree copytree excluded __pycache__ and build/, but aiter's JIT products sit directly in aiter/jit/ and flydsl_cache/ inside the package tree. Measured after copy: 2.14 GB tracked, of which 2.13 GB (117 .so files) was JIT output. git add -A had to hash that over NFS, blowing the 120 s timeout. Extend _SCRATCH_EXCLUDE_GLOBS to cover __pycache__/, flydsl_cache/, jit/ and *.so, write the full set into .git/info/exclude before the baseline git add, and extend _ignore to skip those names during copytree so the bytes are never copied to begin with. With a source-only index, git add completes in under a second. Restore the scaffold git timeout from 120 s to 60 s: a large value only hides a regression since the operation should now be sub-second. Co-authored-by: Cursor <cursoragent@cursor.com>
scratch_dir derived from USER_DATA_PATH lands on shared NFS storage, where git add -A must hash multi-GB JIT products over the network (even after the exclude list shrinks the index). The existing pattern for this problem is in baseline.py (_ensure_local_inferencex), which already detects network filesystems and mirrors onto local disk. Promote _NETWORK_FS_TYPES / _path_fstype / _is_network_fs to a shared hyperloom/common/fs_utils.py module and re-point baseline.py at it. Add _local_scratch_dir() to forge_submit.py: when output_dir is on a network FS, it places the worktree under $FORGE_LOCAL_SCRATCH_ROOT (default ~/.cache/hyperloom/forge_scratch) while leaving the durable experiment archive on the shared mount. A startup sweep removes orphaned local scratch trees from attempts that crashed before cleanup. On local-disk hosts the function returns the original output_dir/worktree path unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>
…site _candidate_keywords stripped both leading and trailing underscores, so a kernel named _mxfp8_linear_kernel yielded the keyword mxfp8_linear_kernel. _file_defines_symbol then searched for "def mxfp8_linear_kernel", which never appears in source -- the real definition is "def _mxfp8_linear_kernel" -- so the definition-site check failed and _prefer_symbol_definition fell back to generic ranking, picking a re-exporting __init__.py over the file that defines the kernel. Strip only trailing underscores, and let _file_defines_symbol match both the bare keyword and its underscore-prefixed spelling so the check succeeds either way. This branch also carried a second change here, accepting an uncorroborated trace launcher frame as the source_file. It is dropped in favour of main's policy: main removed the per-kernel LLM tier this was defending against and now leaves an unconfirmed launcher blank for the whole-table review to weigh, and the resolution fixes on top of this commit make grep succeed in the cases that used to fall through to the frame.
The scratch copytree filter and git info/exclude glob list were local constants that duplicated knowledge already in KernelForge's path-ownership manifest. Keeping them in sync required manual updates in two codebases. Add _forge_scratch_exclude_globs(), _forge_scratch_exclude_dirs() and _forge_scratch_exclude_suffixes() which import from kernel_agents.loop.path_ownership via the existing _ensure_forge_on_path() seam, using the established try/except ImportError pattern from _git_argv(). When kernel_agents is unavailable (remote node, FORGE_PATH unset) the functions fall back to local constants that match the manifest's current content, so the module is fully usable without FORGE_PATH configured. Wire the helpers into _exclude_bytecode_caches() and the _ignore copytree filter so both derive from a single authoritative source. Co-authored-by: Cursor <cursoragent@cursor.com>
… name A traced Triton kernel is named after its device symbol, which says nothing about the language, and frameworks import Triton through a shim rather than by name. Both signals source_type_for relies on are therefore absent, so a real Triton kernel was reported as `python` -- a language no consumer accepts as portable, which declined every one of them at the rewrite gate. Resolve the traced symbol to a `@triton.jit` def through the AST reader that already backs the bypass route, so the verdict comes from the file itself. The FlyDSL check keeps precedence: such a file may carry a Triton reference path, and its FlyDSL identity is the one consumers act on. Co-authored-by: Cursor <cursoragent@cursor.com>
Pins the four behaviours the definition-based check has to hold: a device symbol whose file imports Triton through a framework shim resolves to `triton`; a FlyDSL source carrying a Triton reference path stays `flydsl`; and a plain `def`, or a Triton def whose name does not match the traced symbol, leaves the kernel unproven rather than claimed. Co-authored-by: Cursor <cursoragent@cursor.com>
…h API A trace reports the launch API it observed around a symbol, so one kernel reached through a graph replay and through a direct module launch arrives as two rows whose names differ only by that prefix. Native sources already stripped it; Python sources did not, so a single Triton kernel became two task groups that would each port the same source file, splitting its share and letting two attempts write the same lines. The launch API, the C return type and the synthetic-op suffix are all trace formatting rather than operator identity, so the cleanup moves ahead of the source-kind split and Itanium demangling stays the one native-only step. Distinct kernels behind one launch API keep distinct keys. Operator identities recorded under the previous unstripped form no longer match, so per-operator attempt and rejection memory restarts for Python sources. Co-authored-by: Cursor <cursoragent@cursor.com>
Asserts both key builders reduce two launch paths of one kernel to the same identity, that distinct kernels behind one launch API stay separate, that native demangling still runs after the shared cleanup, and that the grouping above it now aggregates both rows into one task group. Co-authored-by: Cursor <cursoragent@cursor.com>
Extracting them left three private aliases behind for "backward compatibility with call sites in this module", but the module is the only caller: one alias had a single use and the other two had none. The tests reached for the alias too, so they now patch the imported name. The move also dropped why the type list exists -- these mounts can be revoked mid-run -- and the note that /proc/mounts octal-escapes its mountpoints, which is the only reason that decode is there. Co-authored-by: Cursor <cursoragent@cursor.com>
…py of it The scratch exclusions arrived as three constants and three functions, each wrapping one manifest name in its own import-and-fall-back. The fallback cannot help: a run that cannot import kernel_agents has no forge-loop to invoke, and answering with a stale local copy is how the index silently regains the gigabytes this exclusion exists to remove. Two functions now read the manifest directly, and a missing manifest fails where it is read. The orphan sweep swallowed every OSError twice and re-imported shutil inside its own loop; relocation fell back to the network path it exists to avoid. The copytree filter unioned in a directory the manifest already lists. _exclude_bytecode_caches no longer matched what it excludes. Co-authored-by: Cursor <cursoragent@cursor.com>
The tier ordering, the underscore anchor and the launch-decoration rule each restated in prose what the lines beside them already show. Keep the part a reader cannot derive from the code. Accepting an uncorroborated tracer frame guarded its own frame against None on a branch only a non-None frame can reach. Co-authored-by: Cursor <cursoragent@cursor.com>
locate_source_via_grep's primary pass ranked hits by path shape only, so a package __init__ that re-exports a kernel outranked the module holding the @triton.jit body. The hottest routable MiniMax-M3 kernels were dispatched to backends pointed at a file containing no kernel: _mxfp8_linear_kernel resolved to vllm/model_executor/kernels/linear/__init__.py rather than mxfp8/rocm_native.py, leaving the backend nothing to rewrite. Use _prefer_symbol_definition in the primary pass, the same rule the compound sub-window fallback below it already applied. It degrades to plain ranking when no hit defines the symbol, so mention-only resolutions are unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>
aiter's MoE GEMMs are the two hottest kernels on a MiniMax-M3-MXFP8 trace --
31% of GPU time together -- and both resolved to nothing, so the dispatcher
skipped them and every backend ran against 2-8% kernels instead.
They launch as mfma_moe1_silu_mul_afp8_wfp8_bf16_t32x128x256_... but are
written f"mfma_moe1_silu_mul_a{a_dtype}_w{b_dtype}_{out_s}". Only the head of
the name is literal in source, so the whole-name pass and the trailing
sub-window pass both search text that is never written down.
Add a last pass over leading prefixes -- the mirror of the sub-window pass,
dropping trailing segments rather than leading ones. The budget goes to the
short prefixes, since everything interpolated sits at the tail, and they are
still tried longest-first. Sibling f-strings in sibling files share a short
prefix (mfma_moe2_a{...} vs mfma_moe2_{...}), so hits are ranked by how many
characters of the launched name a file spells out; the file that spells out
more is the one that built it.
Both kernels now resolve to aiter/ops/flydsl/kernels/mixed_moe_gemm_2stage.py.
Co-authored-by: Cursor <cursoragent@cursor.com>
The guard freezes its streak while kernel-lane work is in flight, but it learns what is in flight from the task registry. Kernel requests are awaited straight in the intent router and never become a row there -- a fact that module already documents on _kernel_step_heartbeat, which exists precisely because the task-progress heartbeat cannot see them either. So an integrate re-baseline looked exactly like a dead phase: no registry row, and no ledger write until it returns, so the progress fingerprint held still too. run_optimization is masked from this because its attempt ledger updates as it goes; integrate has no such heartbeat. On a MiniMax-M3-MXFP8 run a nine-minute integrate accrued idle_seconds=648 and wound KERNEL_AGENT down to SWEEP four seconds after it completed, with nine selected candidates never dispatched -- including the two hottest kernels on the trace at 16.4% and 9.3% of GPU time. The run closed at 0% validated gain having used 4.7h of a 24h budget. Stamp kernel_inline_step_seen_unix from the heartbeat that already wraps the inline handler, and treat a fresh stamp as in-flight. It is refreshed rather than set once so a stamp orphaned by a process that died mid-step goes stale on its own instead of muting the guard for good. Co-authored-by: Cursor <cursoragent@cursor.com>
… visible Two constants in this repo could never agree at their defaults. The per-attempt backend budget defaults to 60 minutes; the rewrite route requires 75 (an hour for the producer plus a quarter reserved for apply-back). So every candidate that was otherwise eligible got declined budget_insufficient and quietly fell back to forge-loop -- setting HYPERLOOM_FORGE_REWRITE_BY_FLYDSL bought nothing. Nothing said so. submit() computes the verdict and hangs it on its result as "flydsl_rewrite", but that dict is never persisted: the forge_result in the run log is the forge CLI's own JSON, and none of submit()'s fields appear in any artifact. The operator sees a flydsl-fellow attempt and reasonably concludes the route ran. Two full MiniMax-M3-MXFP8 sessions were read that way, including by me; the route has in fact never executed once. Floor the budget at the route's own minimum when, and only when, the operator opted in -- runs that never asked for the route keep the 60-minute default, and an explicit KERNEL_OPT_BACKEND_BUDGET_MIN still wins in both directions. Then write the verdict to the attempt log, the one channel known to surface. Confirmed against the real k005 candidate: budget_insufficient at 3600s, past that gate at 4500s. Note this only reopens the route for Triton/HIP sources; the hottest MoE kernels are already FlyDSL and are declined earlier by already_flydsl_source, which is correct -- rewrite ports INTO FlyDSL. Co-authored-by: Cursor <cursoragent@cursor.com>
ZhengGong-amd
requested review from
a team,
Ahmedhasssan-aig,
devalshahamd and
tsrikris
as code owners
August 28, 2026 09:26
Hyperloom gates `ruff format --check` in CI; `ruff check` was already clean. Co-authored-by: Cursor <cursoragent@cursor.com>
CI E2E report — ✅ Succeeded
|
The no-git scratch path copytrees an installed package and puts it ahead of site-packages, shadowing the install outright. The copy filter had grown to skip jit/ and dist/ and every .so, which is exactly what that package is imported through: aiter/jit carries core.py and the extension modules `import aiter` loads, aiter/dist carries its distributed sources. Because the copy shadows rather than overlays, there was nothing left to fall back to and the worktree could not import the framework at all. Keeping compiled output out of a git index is a separate, wider question, so the git exclude keeps its own list and still covers *.so. Co-authored-by: Cursor <cursoragent@cursor.com>
The durable archive is .../forge/<session_id>/<attempt>, but the local scratch root was keyed on the attempt alone, which is the kernel's name. Two sessions optimizing one kernel therefore claimed one path, and the second was refused for a retained-workspace collision with the first. The sweep then read that flat layout against the durable tree two levels up, comparing an attempt name against a directory of session ids. Nothing ever matched, so the liveness check was constantly false and every local scratch but the current one was deleted -- including trees another attempt on the same node was still writing into. Mirroring <session_id>/<attempt> locally settles both: the sweep compares session against session, and only reaches a directory shaped like its own output. Co-authored-by: Cursor <cursoragent@cursor.com>
The comment claimed a grep result may annotate but never veto the tracer frame. The code below it does the opposite and always has: source_file is written only from the grep hit, and a frame without one is recorded as trace_launcher_file instead. Leaving the claim in place misdirects whoever edits this next. Co-authored-by: Cursor <cursoragent@cursor.com>
Both scratch defects were silent: one skipped a submit, the other deleted a running attempt's worktree, and neither had a test. Build a durable forge/<session>/<attempt> tree against a local root and assert the live session survives a sweep, a session whose archive is gone does not, and a stranger directory is untouched. Assert the copy filter keeps the names a package is imported through while the git exclude still covers them. Co-authored-by: Cursor <cursoragent@cursor.com>
The lint job runs ruff format --check as well as ruff check. Co-authored-by: Cursor <cursoragent@cursor.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Four independent defects found while driving MiniMax-M3-MXFP8 optimization
sessions. Each one silently cost the run kernels it should have optimized.
__init__that re-exports a kernel outranked the module holding the
@triton.jitbody,so a backend was handed a file with no kernel in it. Kernel names assembled by
an f-string (aiter's MoE GEMMs, 26% of GPU time on this trace) matched nothing
at all, because only the head of the name is literal in source. A definition
check also missed underscored symbols.
shipped budget was 60, so every eligible candidate was declined
budget_insufficientand quietly fell back to forge-loop. The verdict wascomputed and attached to the result, but nothing persisted it, so opting in
looked identical to the route running.
task registry, and
integrate/run_optimizationare awaited directly inthe intent router, so a nine-minute
integratere-baseline read as a deadphase and wound KERNEL down four seconds after it finished, stranding every
remaining candidate.
so
git add -Atimed out on NFS and compiled artefacts reached the patch.Notes for review
main. One commit was dropped as superseded:main'skernel_search_roots()chain replaces the host-discovery fix thisbranch carried, and the remaining resolution fixes were ported onto it.
source_file. That is dropped in favour ofmain's policy of leaving itblank for the whole-table review, which also removed the LLM tier the change
was defending against.
mainraised the per-optimization budget to 90 minutes, which already clearsthe rewrite route's gate. The floor added here stays as a guard so a budget
tuned down for another reason cannot switch the route off again in silence.
forge_submitmust not importhyperloomorkernel_agentsat modulelevel and CI does not install the latter, so both are imported inside their
callers with a local fallback.
Test plan