Skip to content

Run-level kernel knowledge base for GEAK; production-size variants of four launch-floor tasks - #135

Closed
zili-amd wants to merge 4 commits into
AMD-AGI:mainfrom
zili-amd:knowledge-base-and-large-tasks
Closed

zili-amd wants to merge 4 commits into
AMD-AGI:mainfrom
zili-amd:knowledge-base-and-large-tasks

Conversation

@zili-amd

Copy link
Copy Markdown

Stacked on #134: the variants' own-kernels rule relies on its harness guard, and both PRs change the Docker runner. Review only the last two commits until #134 lands.

Summary

1. A run-level kernel knowledge base for GEAK

  • A run enables knowledge_wiki: {enabled: true} and names a checkout that follows the KernelWiki layout (for example ROCmKernelWiki) with AKA_KNOWLEDGE_WIKI_ROOT.
  • The Docker runner mounts the checkout read-only, and forwards the variable, only for runs that enable the key. The arm of an A/B pair without it therefore cannot read it. The runner reads the key through a new src.serving_runtime config-knowledge-wiki mode, so the host needs no PyYAML, and only for configs that mention the key.
  • src/knowledge_wiki.py validates the checkout and appends the same instructions to GEAK's Arena contract on every role and Workflow call. It is also the command line the agents query (search, page, grep).
  • Every call is logged with its arguments, exit code, duration and the pages returned. The adapter status records the checkout's commit, whether it was dirty, and the version and SHA-256 of the delivered instructions, so two arms can be checked for identical delivery before their results are compared.
  • A run that enables the key for an agent without support fails before its first task. Docs: docs/how-to/knowledge-wiki.md.

Forge support needs the engine to serve the knowledge base to its planning roles, which have no shell. That arrives with AMD-AGI/Hyperloom#2034 and will follow separately.

2. Production-size variants of four tasks whose cases sit at the launch floor

On MI300X, every case of these four originals finishes within the roughly 5 us launch-and-timing floor, so no candidate can show a speedup and the tasks cannot rank agents. Each variant keeps its original's kernel, references, tolerances, timing and editable scope, and replaces only what its README states.

Variant Original's cases (MI300X baseline per case) Variant's cases (MI300X baseline per case) Replaced
hip2hip/SimpleMatmulModule_large 11 cases, matrices up to 128x128 (1.6-5.5 us) 8 FP32 GEMMs, 1024^3 to 4096^3 and batched (166 us-11.4 ms) case list
hip2hip/SoftmaxModule_large 5 cases up to 64x64x64 (3.3-5.6 us) 8 inputs of 4M-131M elements (30-721 us) case list
triton2triton/vllm_large/triton_awq_gemm 5 cases, N and K at most 512 (7.9-13.6 us) 5 Llama-3.1-70B layers at M 1 and 32 (125-778 us) case table, CPU reference, weights, one gate
triton2triton/vllm_large/triton_w8a8_block_scaled_mm 5 cases (5.0-9.3 us) 12 DeepSeek-V3 layers at M 128 and 4096 (20 us-2.3 ms) case table, CPU reference

The baselines are medians of 5 rounds; the largest run-to-run spread is 5% on any variant. The two vLLM kernels are the ones vLLM runs on ROCm by default for AWQ int4 and FP8 block-scaled linear layers. The CPU references are the original loops, vectorized, because the loops take hours per call at these sizes.

tests/test_production_size_variants.py holds each variant to exactly these differences from its current original. A variant must have the same file set; only the stated files may differ; and within them, only the stated regions. The HIP migration suite excludes the two hip2hip variants, as it already excludes paged_attention_large and wvsplitk. Each variant's README names its source task and what changed. The variants declare required_arch: gfx942, the only architecture they are validated on.

Review points

  • The AWQ variant changes what counts as correct. With the original's uniform random nibbles and roughly 0.09 scales, outputs at LLM K reach 40-70, and the production kernel's fp16 accumulation alone exceeds the original atol. The variant therefore:
    • quantizes a Gaussian weight the way AWQ does;
    • scales the replay perturbation instead of adding a constant sized for the old magnitudes;
    • adds a 5e-3 normwise relative error bound, which an FP8-activation path would otherwise pass at M=1.
  • FP32 GEMM tolerance. SimpleMatmulModule_large keeps the original's rtol=1e-4, atol=1e-5. By estimate (not measured), BF16x3 emulation of the FP32 products may pass this at K up to 8192, while XF32 would not. Should the task require FP32 arithmetic explicitly, or tighten the tolerance?
  • Inherited whitespace. Copied files keep their originals' bytes, trailing whitespace included, so git diff --check reports lines the originals already have.

Not included

  • Forge wiring for the knowledge base, which waits for Hyperloom#2034.
  • Open-launcher variants of four Triton tasks, which make the launcher editable. Upstream is narrowing edit scopes for integrity, so I'd like to discuss that separately first.

Validation

CPU/mock suite: CPython 3.12, CPU PyTorch wheel, Linux.

  • make check-perf-helpers: stubs in sync.
  • bash tests/test_docker_benchmark.sh: PASS.
  • The knowledge-base, GEAK, docs, layout, perf-helper and HIP/Triton migration tests, and the new tests/test_production_size_variants.py: pass.
  • Full suite: see this PR's CI.

GPU: MI300X (gfx942), ROCm 7.2.0, image lmsysorg/sglang:v0.5.12-rocm720-mi30x.

  • The baselines above are from screening before rebasing onto current main. The hip2hip variants are byte-identical to the screened versions. The vLLM variants differ only in upstream's narrower editable scope and README text.
  • Forge campaigns on the earlier versions show the tasks can be solved and leave room between agents. Final speedups (geometric mean over cases), all runs passing correctness:
    • SimpleMatmulModule_large: 6.3-7.3x (3 runs);
    • SoftmaxModule_large: 3.1-3.2x (7 runs);
    • triton_awq_gemm: 2.8-10.1x (12 runs);
    • triton_w8a8_block_scaled_mm: 1.8-2.1x (12 runs).
  • The vLLM runs had the whole source file editable, before upstream narrowed the originals to the kernel symbols.
  • Still to do before review: initial task validation and 5-round baselines of these exact task packages on MI300X.

zili-amd and others added 4 commits October 21, 2026 11:38
…n it

A forge run config now gets what a GEAK run gets: its agent backend's CLI
(Claude unless the run sets agent_backend: codex), the pinned Hyperloom
checkout named by HYPERLOOM_ROOT mounted read-only, and Claude credentials
forwarded by name. `make docker-setup-forge` installs Hyperloom[forge] from
an exported copy of that checkout's committed HEAD into the host-mounted user
base, records the installed revision, and skips an unchanged one; the
container bootstrap prepends that engine only for Forge runs. Other agents
never receive the checkout, and setup-forge receives no credentials.

`docker_benchmark.sh exec <command>` runs one command non-interactively in the
runtime container, and ANTHROPIC_CUSTOM_HEADERS is forwarded with the other
Anthropic variables to GEAK and Forge containers.

Co-authored-by: Cursor <cursoragent@cursor.com>
… operator calls

Every evaluation path verifies the workspace harness before it trusts a
measurement. That check now also refuses a candidate file that adds, relative
to the task's starting file:

- runtime patching in Python (exec/eval/setattr/globals, frame and code
  object access, importlib reloads), which can make a protected launcher,
  reference or timing loop run differently from how it reads;
- library operator calls in C++/HIP sources (hipBLAS, rocBLAS, Tensile,
  MIOpen, and the ATen/torch products, softmax, norms and reductions a task
  asks the candidate to implement);
- PyTorch or library operators in the host code of a Triton candidate
  (bodies of @triton.jit functions are skipped).

Sites the task's own starting file already has stay allowed. The per-file
starting counts are recorded with the task session, so a resumed invocation
compares against the same state, and GEAK's bridge carries them through its
saved snapshot.

Co-authored-by: Cursor <cursoragent@cursor.com>

Two existing tests change with it. An actual-gate payload such as
`torch.allclose = lambda *args, **kwargs: True` was accepted under a
whole-file editable scope and refused only under the symbol scope; it is now
refused under both. In the softmax builder test, `vars(sys)` and
`sys.__dict__` reads are refused by the guard before the task's own
dependency check, which still rejects every listed statement.

Co-authored-by: Cursor <cursoragent@cursor.com>
A run enables `knowledge_wiki: {enabled: true}` and names a checkout that
follows the KernelWiki layout with AKA_KNOWLEDGE_WIKI_ROOT. The Docker runner
mounts that checkout read-only, and forwards the variable, only for such runs,
so the arm of an A/B pair without it cannot read it. It reads the key through
src.serving_runtime, which needs no host PyYAML, and only for configs that
name it.

src/knowledge_wiki.py resolves and validates the checkout, appends the same
instructions to GEAK's Arena contract on every role and Workflow call, and is
the command line the agents query (search, page, grep). Every call is logged
with its arguments, exit code, duration and returned pages; the adapter
status records the checkout's commit, whether it was dirty, and the version
and SHA-256 of the delivered instructions. A run that enables the key for an
agent without support fails before its first task.

Co-authored-by: Cursor <cursoragent@cursor.com>
…nch floor

On MI300X every case of these originals finishes in 1.6-13.6 us, at the
~5 us launch and timing floor, so no candidate can show a speedup. Each
variant keeps its original's kernel, references, tolerances, timing and
editable scope and replaces only what its README states:

- hip2hip/SimpleMatmulModule_large and SoftmaxModule_large: the case list
  in get_inputs() of both reference modules and workload.json. Eight FP32
  GEMM cases from 1024^3 to 4096^3 plus batched products, and eight softmax
  inputs of 4M-131M elements.
- triton2triton/vllm_large/triton_awq_gemm: decode-time Llama-3.1-70B
  layers at M 1 and 32. The CPU reference is the original loops,
  vectorized; the weights are an AWQ-quantized Gaussian, so outputs have a
  real layer's magnitudes; the replay perturbation scales instead of adding a
  constant sized for the old magnitudes; and a 5e-3 normwise error bound is
  added, which an FP8-activation path would otherwise pass at M=1.
- triton2triton/vllm_large/triton_w8a8_block_scaled_mm: DeepSeek-V3 layers
  with 128x128 FP8 blocks at M 128 and 4096, with the original loops
  vectorized as the CPU reference.

Each README states the operator must be computed in the candidate's own
kernels, which the harness guard enforces. The variants declare
required_arch gfx942, the only architecture they are validated on.
tests/test_production_size_variants.py holds each variant to exactly these
differences from its current original; the HIP migration suite excludes the
two hip2hip variants as it excludes the other non-migrated tasks. Copied
files keep their originals' bytes, trailing whitespace included.

Co-authored-by: Cursor <cursoragent@cursor.com>
@zili-amd
zili-amd force-pushed the knowledge-base-and-large-tasks branch from fb46fde to 63cd8af Compare October 21, 2026 20:59
@irvineoy

Copy link
Copy Markdown
Collaborator

Thanks for your PRs. Please let me know when they are ready to review

@zili-amd

Copy link
Copy Markdown
Author

Closing this PR for now.

@zili-amd zili-amd closed this Oct 26, 2026
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