Run Edge0 model folders with edge0-strata, as a llama-swap model

containers/llama-swap/Containerfile.hip

@@ -36,6 +36,60 @@ RUN HIPCXX="$(hipconfig -l)/clang" HIP_PATH="$(hipconfig -R)" \
 COPY collect-rocm-libs.sh /usr/local/bin/collect-rocm-libs.sh
 RUN /usr/local/bin/collect-rocm-libs.sh /out "${GPU_TARGETS}" build/bin/llama-server build/bin/llama-bench build/bin/llama-batched-bench
 
+# Edge0 (edge0-strata) for Radeon RX 7900 cards. The engine is not llama.cpp:
+# it runs an Edge0 latent MoE model from a model folder, and its Python server
+# gives the OpenAI API. llama-swap starts it through edge0-launch.py.
+#
+# EDGE0_REF is a commit of github.com/diffsec/Edge0, the fork that carries the
+# gfx11 kernels. EDGE0_GPU_TARGETS must stay gfx1100: the kernels and
+# the hipBLASLt tuning table are for that architecture only.
+FROM docker.io/rocm/dev-ubuntu-24.04:7.2.4-complete@sha256:92f309c51b52cef8762867848f1529dee821624f23cd5df38455e819538f762f AS edge0-build
+
+ARG EDGE0_REF=49d5def3bf41807acc7eb5b0b07319f75bbcff88
+# The llama.cpp commit that edge0-strata/EDGE0.md names. It gives ggml to the build.
+ARG LLAMA_CPP_REF=3cf03257f219afbe7334045ff7c6a06ac68c627d
+ARG EDGE0_GPU_TARGETS="gfx1100"
+# ROCm 7.2.4 has hipBLASLt 1.2.2, which is tuning table 100202.
+ARG HIPBLASLT_TABLE=gfx1100-hipblaslt-100202.txt
+
+RUN apt-get update \
+    && apt-get install -y --no-install-recommends build-essential cmake ninja-build git ca-certificates curl \
+    && rm -rf /var/lib/apt/lists/*
+
+RUN git init -q /edge0 \
+    && git -C /edge0 fetch -q --depth 1 https://github.com/diffsec/Edge0 "${EDGE0_REF}" \
+    && git -C /edge0 checkout -q FETCH_HEAD \
+    && mkdir -p /edge0/edge0-strata/third_party/llama.cpp \
+    && curl -fsSL "https://github.com/ggml-org/llama.cpp/archive/${LLAMA_CPP_REF}.tar.gz" \
+        | tar -xz --strip-components=1 -C /edge0/edge0-strata/third_party/llama.cpp \
+    && ln -s /edge0/edge0-strata /src
+
+# STRATA_PREFILL_MMQ=ON is required: without it the link stops with undefined
+# symbols such as strata::prefill::fused_q3k::experts_latent_q3k.
+WORKDIR /src
+RUN HIP_PATH="$(hipconfig -R)" \
+    cmake -S . -B build -G Ninja \
+        -DCMAKE_BUILD_TYPE=Release \
+        -DSTRATA_ENABLE_HIP=ON \
+        -DSTRATA_ENABLE_CUDA=OFF \
+        -DCMAKE_HIP_ARCHITECTURES="${EDGE0_GPU_TARGETS}" \
+        -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang++" \
+        -DCMAKE_PREFIX_PATH=/opt/rocm \
+        -DSTRATA_PREFILL_MMQ=ON \
+        -DSTRATA_NATIVE_EXPERTS=ON \
+        -DSTRATA_GGML_DIR=/src/third_party/llama.cpp \
+    && cmake --build build --target strata -j"$(nproc)"
+
+# The engine, its shared libraries, and only the rocBLAS and hipBLASLt kernel
+# files for GPU_TARGETS. The engine reads the tuning table from tuning/ beside
+# its binary. Without the table a 5.3K prompt ran at 725 tok/s, not 1150.
+COPY collect-rocm-libs.sh /usr/local/bin/collect-rocm-libs.sh
+RUN /usr/local/bin/collect-rocm-libs.sh /out "${EDGE0_GPU_TARGETS}" build/strata \
+    && mkdir -p /out/opt/edge0 /out/opt/llama/bin/tuning \
+    && cp "tools/hip/${HIPBLASLT_TABLE}" /out/opt/llama/bin/tuning/ \
+    && cp -r serve strata_control strata_cli tools requirements.txt /out/opt/edge0/ \
+    && find /out/opt/edge0 -name '__pycache__' -prune -exec rm -rf {} +
+
 FROM docker.io/ubuntu:24.04 AS llama-swap-dl
 
 ARG LLAMA_SWAP_VERSION=266
@@ -50,13 +104,19 @@ FROM docker.io/ubuntu:24.04
 
 RUN apt-get update \
     && apt-get install -y --no-install-recommends ca-certificates libgomp1 libssl3t64 curl \
-        libdrm2 libdrm-amdgpu1 libelf1t64 libnuma1 \
+        libdrm2 libdrm-amdgpu1 libelf1t64 libnuma1 python3 python3-venv \
     && rm -rf /var/lib/apt/lists/* \
     && groupadd --system --gid 10001 app \
     && useradd --system --uid 10001 --gid 10001 --home /app --no-create-home app
 
 COPY --from=build /out/ /
 COPY --from=llama-swap-dl /out/llama-swap /app/llama-swap
+# The Edge0 engine, its server, and the server's Python packages at the
+# versions that the engine pins.
+COPY --from=edge0-build /out/ /
+RUN python3 -m venv /opt/venv \
+    && /opt/venv/bin/pip install --no-cache-dir -r /opt/edge0/requirements.txt
+COPY edge0-launch.py /opt/edge0/launch.py
 COPY config.yaml /app/config/config.yaml
 COPY check-libs.sh /usr/local/bin/check-libs.sh
 # Fail the build if any binary or library has an unresolved dependency.
@@ -70,7 +130,9 @@ ENV PATH=/app:/opt/llama/bin:${PATH} \
     ROCBLAS_TENSILE_LIBPATH=/opt/rocm/lib/rocblas/library \
     HIPBLASLT_TENSILE_LIBPATH=/opt/rocm/lib/hipblaslt/library \
     GPU_MAX_HW_QUEUES=1 \
-    SIBYL_GPU_BACKEND=ROCm
+    SIBYL_GPU_BACKEND=ROCm \
+    PYTHONDONTWRITEBYTECODE=1 \
+    PYTHONUNBUFFERED=1
 
 USER 10001:10001
 WORKDIR /app

containers/llama-swap/edge0-launch.py

@@ -0,0 +1,137 @@
+#!/usr/bin/env python3
+"""Start edge0-strata's server for one Edge0 model folder.
+
+llama-swap runs this as the command of an Edge0 model. The folder's
+model.json is a complete server configuration for an NVIDIA start from inside
+the folder. This script changes what differs here, writes the result to a
+work directory, and replaces itself with the server.
+"""
+import argparse
+import json
+import os
+import sys
+
+ENGINE = "/opt/llama/bin/strata"
+SERVER = "/opt/edge0/serve/server.py"
+ROCM_LIB = "/opt/rocm/lib"
+# Options whose value differs here, with the number of values each one takes.
+DROP = {"--pool-workers": 1, "--max-context": 1, "--vram-reserve-later-mib": 1}
+# A card that another model also uses can hold less than the engine planned
+# for. On the test host a later card with 240 MiB free stalled the first
+# request, and 1400 MiB of reserve did not.
+RESERVE_LATER_MIB = "1400"
+
+
+def absolute(value, model_dir):
+    """Resolve a "./" path of the folder's configuration against the folder."""
+    if value == ".":
+        return model_dir
+    if isinstance(value, str) and value.startswith("./"):
+        return os.path.join(model_dir, value[2:])
+    return value
+
+
+# On one GPU the engine keeps the expert file in system RAM, plus about
+# RAM_OVERHEAD. On a layer split it keeps only the experts that no GPU holds
+# (--resident-experts). On two 24 GB cards that was 21.4 GiB of a 37.3 GiB
+# file, which is 57%, so SPAN_SHARE leaves a margin. The same figures are in
+# internal/models (Edge0RAM).
+RAM_OVERHEAD = 6 << 30
+SPAN_SHARE = 0.65
+
+
+def ram_short(experts_path, span=False, meminfo="/proc/meminfo"):
+    """Return (needed, available) in bytes when RAM is short, else None."""
+    try:
+        experts = os.path.getsize(experts_path)
+        need = (int(experts * SPAN_SHARE) if span else experts) + RAM_OVERHEAD
+        with open(meminfo, encoding="ascii") as f:
+            for line in f:
+                if line.startswith("MemAvailable:"):
+                    avail = int(line.split()[1]) << 10
+                    return (need, avail) if avail < need else None
+    except (OSError, ValueError, IndexError):
+        pass
+    return None  # the size or the memory figure cannot be read
+
+
+def configure(manifest, model_dir, alias, gpus, context, work):
+    """Return the server configuration for this container."""
+    args, src, i = [], manifest["args"], 0
+    while i < len(src):
+        arg = src[i]
+        # The engine picks the worker count from the CPU when the option is absent.
+        if arg in DROP:
+            i += 1 + DROP[arg]
+            continue
+        # --pcie-mode dma is not tested on a layer split with the engine's
+        # split fix. Before the fix it stopped the first request.
+        if arg == "--pcie-mode" and len(gpus) > 1:
+            i += 2
+            continue
+        args.append(absolute(arg, model_dir))
+        i += 1
+    args += ["--max-context", str(context)]
+    if len(gpus) > 1:
+        args += ["--vram-reserve-later-mib", RESERVE_LATER_MIB]
+        # Two cards hold most of the experts, so RAM keeps only the rest.
+        # Without this the engine keeps the whole expert file in RAM too.
+        if "--resident-experts" not in args:
+            args.append("--resident-experts")
+
+    env = {k: absolute(v, model_dir) for k, v in manifest.get("env", {}).items()}
+    return {
+        "model_name": alias,
+        "backend": "hip",
+        "exe": ENGINE,
+        "args": args,
+        "cwd": work,
+        "log": os.path.join(work, "engine.log"),
+        "tokenizer": absolute(manifest.get("tokenizer", "./tokenizer"), model_dir),
+        "gpus_asked": True,
+        "gpu": gpus if len(gpus) > 1 else gpus[0],
+        "lib_dirs": [ROCM_LIB],
+        "env": env,
+    }
+
+
+def main():
+    ap = argparse.ArgumentParser(description=__doc__.splitlines()[0])
+    ap.add_argument("--model-dir", required=True)
+    ap.add_argument("--alias", required=True, help="the model name that the server answers to")
+    ap.add_argument("--gpus", required=True, help="GPU numbers as HIP numbers them here, such as 0 or 0,1")
+    ap.add_argument("--context", required=True, type=int)
+    ap.add_argument("--port", required=True, type=int)
+    ap.add_argument("--work", default="/tmp/edge0", help="a writable directory for the configuration and the engine log")
+    ap.add_argument("--print-config", action="store_true", help="print the configuration and stop")
+    opt = ap.parse_args()
+
+    gpus = [int(g) for g in opt.gpus.split(",")]
+    with open(os.path.join(opt.model_dir, "model.json"), encoding="utf-8") as f:
+        manifest = json.load(f)
+    work = os.path.join(opt.work, opt.alias)
+    cfg = configure(manifest, opt.model_dir, opt.alias, gpus, opt.context, work)
+    if opt.print_config:
+        json.dump(cfg, sys.stdout, indent=1)
+        return
+    # Without the RAM, the engine pushes the host into swap or the kernel
+    # stops a process. Refuse the load with a message that says why.
+    short = ram_short(cfg["env"].get("STRATA_LATENT_A", os.path.join(opt.model_dir, "experts.bin")), span=len(gpus) > 1)
+    if short:
+        sys.exit("edge0: %s keeps %.1f GiB of system RAM, and %.1f GiB is available"
+                 % (opt.alias, short[0] / (1 << 30), short[1] / (1 << 30)))
+    # The image sets GPU_MAX_HW_QUEUES=1 for idle llama-server processes. This
+    # engine waits on the GPU in one stream for a copy in another stream. On
+    # one hardware queue the copy never starts, and the first request stops
+    # with "verify: timed out at layer N".
+    os.environ.pop("GPU_MAX_HW_QUEUES", None)
+    os.makedirs(work, exist_ok=True)
+    path = os.path.join(work, "config.json")
+    with open(path, "w", encoding="utf-8") as f:
+        json.dump(cfg, f, indent=1)
+    os.execv(sys.executable, [sys.executable, SERVER, "--engine", "strata", "--config", path,
+                              "--host", "127.0.0.1", "--port", str(opt.port), "--no-api-key"])
+
+
+if __name__ == "__main__":
+    main()

docs/design/overview.md

@@ -9,7 +9,7 @@ Sibyl is an all-in-one inference server. It is a hardened, immutable Linux host
 - Each GPU vendor has its own host image variant: `-amd`, `-nvidia` and `-intel`.
 - The AMD variant is tested on a host with two Radeon RX 7900 XTX cards (the "test host" in this document).
 - The NVIDIA and Intel variants build and pass the same checks, but no one has run them on hardware yet.
-- Models are GGUF files only, quantized to fit consumer GPUs. Sibyl does not use vLLM.
+- Models are GGUF files, quantized to fit consumer GPUs. The one exception is an Edge0 model folder on the AMD variant. Sibyl does not use vLLM.
 
 ## Decisions
 
@@ -19,8 +19,9 @@ Sibyl is an all-in-one inference server. It is a hardened, immutable Linux host
 | Container runtime | Podman with Quadlet, rootless | Fedora default. Quadlet units start at boot without a daemon. |
 | Container UI | Arcane (BSD-3-Clause) | Compose editor, Git Sync, native Podman and Quadlet pages. Dockhand is source-available (BSL 1.1), so it is out. |
 | Inference engine | llama.cpp `llama-server` behind llama-swap | One engine for every GPU vendor, with fast single-stream GGUF inference. llama-swap swaps models and places them on GPUs. |
+| Second engine | edge0-strata for Edge0 model folders, HIP image only, as a llama-swap model | llama.cpp cannot run the Edge0 latent MoE models. llama-swap starts the engine like any other model, so the swap rules, the router and the load routes do not change. |
 | GPU backends | HIP (ROCm) and Vulkan for AMD, CUDA and Vulkan for NVIDIA, Vulkan for Intel | The fastest backend differs by GPU, quant and context size. The benchmark compares the variant's backends. |
-| Model format | GGUF only | Quantized models fit consumer GPUs. No vLLM. |
+| Model format | GGUF, and Edge0 model folders | Quantized models fit consumer GPUs. No vLLM. |
 | Multi-GPU | Per-model layout through the llama-swap matrix router | One model can span both cards, or two models can run on one card each. |
 | GPU vendors | Separate image variants: `-amd`, `-nvidia`, `-intel` | Only the NVIDIA variant carries out-of-tree kernel modules, signed with the project MOK key. NVIDIA and Intel are untested on hardware. |
 | Base for service images | Wolfi where possible | Smaller images with fewer CVEs. The inference images use the vendors' bases, because Wolfi does not package ROCm or CUDA. |
@@ -153,6 +154,31 @@ The HIP image sets `GPU_MAX_HW_QUEUES=1`. With the default of 4 hardware queues
 | `LLAMA_SWAP_VERSION` | pinned release | llama-swap release |
 | `GPU_TARGETS` | `gfx1030;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1200;gfx1201` | RDNA2 (RX 6800 and 6900, PRO W6800), RDNA3 (RX 7000, PRO W7000), RDNA3.5 (Ryzen AI 300, Strix Halo), RDNA4 (RX 9000, AI PRO R9700) |
 | `GGML_HIP_RCCL` | `OFF` | RCCL helps only `--split-mode tensor`. Upstream disables it by default. |
+| `EDGE0_REF` | pinned commit of `diffsec/Edge0` | The Edge0 engine, HIP image only |
+
+### Edge0 models
+
+An Edge0 model is a folder, not a GGUF file. The folder holds a `model.json` with `"family": "Edge0"`, an expert file, and the rest of the pack. The scanner lists the folder as one model with `engine` set to `edge0` (`models.Scan`). llama.cpp cannot run it. The engine is edge0-strata, from the fork `github.com/diffsec/Edge0`, which carries the kernels for the Radeon RX 7900 matrix cores.
+
+The HIP image builds the engine in a stage of its own, for `gfx1100` only, and adds its Python server. `EDGE0_REF` pins the fork's commit. The model's llama-swap command is `containers/llama-swap/edge0-launch.py`. It reads the folder's `model.json`, which is a server configuration for a start from inside the folder, and changes these parts:
+
+- The paths become absolute, and the work directory and the engine log go to `/tmp/edge0/<model>`. The container's root is read-only.
+- `--max-context` comes from the model's `ctx_size`. With no setting it is 131072, the value in `model.json` and the value that the engine was measured with.
+- `--pool-workers` is removed. The engine then takes the count from the CPU.
+- It removes `GPU_MAX_HW_QUEUES` from the engine's environment. The engine waits on the GPU in one stream for a copy in another stream. On one hardware queue the copy never starts, and the first request stops with `verify: timed out at layer 1`.
+- On more than one GPU it adds `--resident-experts`, the engine's low-RAM mode: RAM keeps only the experts that no GPU holds. It also adds `--vram-reserve-later-mib 1400` and removes `--pcie-mode dma`, which is not tested on a layer split.
+
+The model's layout works as for a GGUF model. A model that spans the selected GPUs runs as a layer split, in a matrix set of its own. A model pinned to one GPU shares the `pinned` set, so a GGUF model on another GPU runs beside it. The other llama-server settings do not apply, and the TUI and the console show only the layout, the context size and the hidden flag. The benchmark and the optimizer refuse an Edge0 model, because both measure llama.cpp.
+
+On one GPU the engine keeps the whole expert file in system RAM and copies experts from there to the GPU. Sibyl counts the size of `experts.bin` plus 6 GiB as the model's RAM (`models.Edge0RAM`), which is 43.3 GiB for Edge0 145B. On the test host that model took 41 GiB of RAM, beside 22.6 GB of VRAM on one card. A model that spans GPUs keeps less, because the cards hold more of the experts: on two 24 GB cards, 27.0 GiB of the 37.3 GiB of experts were in VRAM, 21.4 GiB stayed in RAM, and the process took 32.5 GB. For a span Sibyl counts 65% of the expert file plus 6 GiB (`Model.Edge0RAMOn`), which is 30.3 GiB for Edge0 145B. `POST /models/{id}/load` refuses the load with `422` when `MemAvailable` in `/proc/meminfo` is lower, unless the model already runs. The launcher makes the same check, because a chat request also starts a load, and llama-swap then answers the request with the launcher's message.
+
+On the test host, Edge0 145B (Q3_K) on one RX 7900 XTX in this image, with the Quadlet's restrictions, decoded chat answers at 44.1 to 56.1 tokens per second. It read a 5.3K-token prompt at 1029 to 1201 tokens per second and a 21K-token prompt at 1291. The idle card drew 17 W with the model loaded. A layer split across both cards decoded at 32.2 to 51.2 tokens per second, and it read a 5.3K-token prompt at 273 to 285 tokens per second and a 21K-token prompt at 413. One of the test host's cards is on a PCIe 3.0 x1 link, and the prompt path streams experts to it.
+
+The span needs the fork at `49d5def` or later. Before that commit, the second card of a split was given pointers to the first card's latent weights and token embedding table. A card cannot read another card's VRAM, so the first request never finished.
+
+The download picker offers an Edge0 repository as one variant, the whole folder (`hf.Client.Edge0Variant`). A repository counts as Edge0 when `model.json` and `experts.bin` are at its root. The variant leaves out the `assets` folder and hidden files, and its fit is `split` when the host's total RAM holds the model, or `no`. The download keeps the subfolders. Hugging Face gives a SHA-256 for LFS files only, so the download checks each small file against its Git blob ID, which is the SHA-1 of `blob <size>`, a zero byte, and the content. The commit pins both kinds of file.
+
+Only the Q3_K folders run. The NVFP4 folders need kernels that the fork has for CUDA only, and Sibyl does not stop a download of one.
 
 The base images are plain `FROM` lines, so Renovate can update them. The HIP build uses `docker.io/rocm/dev-ubuntu-24.04:7.2.4-complete`. Upstream llama.cpp tests ROCm 7.2.x, and ROCm 7.14 is a benchmark candidate. Renovate leaves the ROCm and Ubuntu releases alone, because a new release changes the llama.cpp build. A maintainer moves them by hand after a benchmark.
 

docs/design/web-console.md

@@ -41,7 +41,7 @@ The console uses Go `html/template` pages on the server with htmx for partial up
 
 The pages use one stylesheet with a light and a dark theme from `prefers-color-scheme`. The layout works at phone width. Live views, such as download progress and the router's last decisions, refresh every 5 seconds with `hx-trigger="every 5s"`, as the TUI does.
 
-The model settings form of the console and the TUI's form share `internal/modelform`, so both check a value the same way.
+The model settings form of the console and the TUI's form share `internal/modelform`, so both check a value the same way. For an Edge0 model, both forms show only the layout, the context size and the hidden flag, because the other settings are llama-server flags.
 
 ## Components