[misc]: H3 MK tier: drop narration comments, move tests under the collected transformers lane

[bugfix]: H3 MK kernels: int64 QKV batch offset, broadcast gated-residual index over batch
[perf] MiniMax H3: opt-in faster kernel tier for the DiT fused ops
2026-10-30 13:26:54 -07:00 · 2026-10-30 13:23:26 -07:00 · 2026-10-30 13:04:59 -07:00
59 changed files with 1648 additions and 1302 deletions
@@ -49,6 +49,7 @@ export default function HardwarePicker({
  value,
  onChange,
  gpuCount = 1,
+  purpose = "target",
  reload,
 }: {
  backend: string;
@@ -58,6 +59,7 @@ export default function HardwarePicker({
  value: HardwareSelection;
  onChange: (selection: HardwareSelection) => void;
  gpuCount?: number;
+  purpose?: "target" | "filter";
  reload?: () => void;
 }) {
  const id = useId();
@@ -69,9 +71,9 @@ export default function HardwarePicker({
  const [adding, setAdding] = useState(false);
  const [draft, setDraft] = useState<Draft>(() => draftFrom(value.custom, backend));
  const record = resolveHardware(value, catalog);
-  const unknown = status === "ready" && !!value.id && value.id !== "custom" && !record;
-  const loading = status === "loading" && !!value.id && value.id !== "custom";
-  const candidates = catalog.filter((item) => item.backend === (backend === "mlx" ? "mlx" : "cuda"));
+  const unknown = !!value.id && value.id !== "custom" && !record;
+  const candidates =
+    purpose === "filter" ? catalog : catalog.filter((item) => item.backend === (backend === "mlx" ? "mlx" : "cuda"));
  const matches = searchHardware(candidates, query).slice(0, 20);
  const draftMemory = amount(draft.memory);
  const draftBudget = amount(draft.budget);
@@ -142,7 +144,8 @@ export default function HardwarePicker({
            </span>
            {record.usable_budget_gb ? (
              <span className="hardware-selected-meta">
-                Usable memory: {record.usable_budget_gb} {record.memory_unit}
+                {value.id === "custom" ? "Your budget" : "Usable budget"}: {record.usable_budget_gb}{" "}
+                {record.memory_unit}
              </span>
            ) : null}
            {record.host_ram_gb && record.memory_type !== "unified" ? (
@@ -152,24 +155,20 @@ export default function HardwarePicker({
            ) : null}
          </div>
          <button type="button" className="hardware-clear" onClick={() => onChange({ id: "" })}>
-            Clear hardware
+            Clear hardware {purpose}
          </button>
        </div>
      ) : value.id === "custom" ? (
        <p className="hardware-current" role="status">
          Add this device's memory to finish selecting it.
        </p>
-      ) : loading ? (
-        <p className="hardware-current" role="status">
-          Loading hardware catalog…
-        </p>
      ) : unknown ? (
        <p className="hardware-current" role="status">
-          The saved device ({value.id}) is not in this catalog.
+          Saved hardware target ({value.id}) is not in this catalog.
        </p>
      ) : (
        <p className="hardware-current" role="status">
-          No hardware selected.
+          No hardware {purpose} selected.
        </p>
      )}
      {record ? (
@@ -185,7 +184,7 @@ export default function HardwarePicker({
      <div className="hardware-picker-actions">
        {!record && value.id ? (
          <button type="button" className="hardware-clear" onClick={() => onChange({ id: "" })}>
-            Clear hardware
+            Clear hardware {purpose}
          </button>
        ) : null}
        <button
@@ -199,11 +198,21 @@ export default function HardwarePicker({
            setAdding(false);
          }}
        >
-          {record ? "Change hardware" : backend === "mlx" ? "+ Find or add a Mac" : "+ Find or add a GPU"}
+          {record
+            ? "Change hardware"
+            : purpose === "filter"
+              ? "Choose hardware"
+              : backend === "mlx"
+                ? "+ Find or add a Mac"
+                : "+ Find or add a GPU"}
        </button>
      </div>
      {record || open ? (
-        <p className="hardware-picker-caption">Used for the setup link and, on single-unit systems, the GPU count.</p>
+        <p className="hardware-picker-caption">
+          {purpose === "filter"
+            ? "Filters by runtime. Memory fit is not estimated."
+            : "Single-GPU devices set GPU count to 1. Nothing else in the configuration changes."}
+        </p>
      ) : null}
      {open ? (
        <div className="hardware-lookup" id={panelId} onKeyDown={(event) => event.key === "Escape" && close()}>
@@ -287,13 +296,13 @@ export default function HardwarePicker({
                    value={draft.memory_type}
                    onChange={(event) => setDraft({ ...draft, memory_type: event.target.value as MemoryType })}
                  >
-                    <option value="dedicated">GPU memory (VRAM)</option>
-                    <option value="unified">Unified memory</option>
-                    <option value="system">System memory (RAM)</option>
+                    <option value="dedicated">Dedicated GPU memory (VRAM)</option>
+                    <option value="unified">Unified memory (CPU + GPU)</option>
+                    <option value="system">System memory (host RAM)</option>
                  </select>
                </label>
                <label className="hardware-form-field">
-                  <span>Total memory</span>
+                  <span>Nominal memory</span>
                  <Input
                    type="number"
                    min="1"
@@ -314,7 +323,7 @@ export default function HardwarePicker({
                  </select>
                </label>
                <label className="hardware-form-field">
-                  <span>Usable memory (optional)</span>
+                  <span>Usable budget (optional)</span>
                  <Input
                    type="number"
                    min="1"
@@ -339,14 +348,16 @@ export default function HardwarePicker({
                  </label>
                ) : null}
              </div>
-              {budgetTooLarge ? <p className="hardware-form-error">Usable memory cannot exceed total memory.</p> : null}
+              {budgetTooLarge ? (
+                <p className="hardware-form-error">The usable budget cannot exceed the nominal memory.</p>
+              ) : null}
              <p className="hardware-form-note">
-                Total memory is the device's capacity. Usable memory is the part FastVideo may use, such as a MIG slice
-                or a limit you set. Unified memory is shared with the CPU.
+                Nominal memory is device capacity, not what FastVideo can use. A budget is your own limit and is not
+                verified. Unified memory is shared with the CPU; it is not dedicated VRAM.
              </p>
              <div className="hardware-form-actions">
                <button type="submit" disabled={!validDraft}>
-                  Save device
+                  Use this device
                </button>
                <button
                  type="button"
@@ -1,8 +1,8 @@
-import { useId, useState } from "react";
-import type { MemoryAssessment } from "@/lib/capacity";
+import { useState } from "react";

 export type Validity = "valid" | "invalid" | "loading";
 export type Fit = "fits" | "tight" | "exceeds" | "unsupported" | "unknown";
+export type Evidence = "recorded" | "none" | "stale";

 const validityLabel: Record<Validity, string> = {
  valid: "Configuration valid",
@@ -12,51 +12,48 @@ const validityLabel: Record<Validity, string> = {
 const fitLabel: Record<Fit, string> = {
  fits: "Fits estimate",
  tight: "Tight fit",
-  exceeds: "Exceeds memory",
-  unsupported: "Unsupported",
-  unknown: "Memory not estimated",
+  exceeds: "Likely exceeds memory",
+  unsupported: "Unsupported combination",
+  unknown: "Fit unknown",
 };
-const weightsLabel: Record<Fit, string> = {
-  fits: "Weights fit",
-  tight: "Weights fit, little headroom",
-  exceeds: "Weights exceed memory",
-  unsupported: "Weights unsupported",
-  unknown: "Weights not estimated",
+const evidenceLabel: Record<Evidence, string> = {
+  recorded: "Run evidence available",
+  none: "No recorded run",
+  stale: "Run evidence does not match these settings",
 };

 /**
- * Two facts, each with its own evidence: the recipe schema accepts the values,
- * and the analytical estimate covers the selected device. Neither is a promise
- * that the recipe runs.
+ * Three separate facts: the recipe schema accepts the values, a memory estimate
+ * covers the configuration, and a recorded run matches it. No positive claim is
+ * made until the shared API supplies the corresponding fact.
 */
 export default function RecipeFacts({
  validity,
-  memory,
+  fit = "unknown",
+  evidence = "none",
  subject = "configuration",
+  detail,
  validityText,
  validityDetail,
-  detail,
 }: {
  validity: Validity;
-  memory?: MemoryAssessment;
+  fit?: Fit;
+  evidence?: Evidence;
  subject?: "configuration" | "example";
+  detail?: string;
  validityText?: string;
  validityDetail?: string;
-  detail?: string;
 }) {
  const [open, setOpen] = useState(false);
-  const id = useId();
-  const fit: Fit = memory?.status ?? "unknown";
-  const estimated = fit === "fits" || fit === "tight" || fit === "exceeds";
-  const weightsBasis = memory?.basis === "weights";
-  const label = weightsBasis ? weightsLabel[fit] : fitLabel[fit];
+  const id = `recipe-facts-${validity}-${fit}-${evidence}`;
  return (
    <div className="recipe-facts">
-      <div className="recipe-fact-list" aria-label="Recipe checks">
+      <div className="recipe-fact-list" aria-label="Recipe facts">
        <span data-state={validity}>
          {validity === "valid" ? (validityText ?? `Valid ${subject}`) : validityLabel[validity]}
        </span>
-        <span data-state={fit}>{label}</span>
+        <span data-state={fit}>{fitLabel[fit]}</span>
+        <span data-state={evidence}>{evidenceLabel[evidence]}</span>
      </div>
      <button
        type="button"
@@ -73,19 +70,24 @@ export default function RecipeFacts({
            <strong>{validityLabel[validity]}.</strong>{" "}
            {validityDetail ??
              (validity === "valid"
-                ? "These values pass the recipe schema and runtime topology checks."
+                ? "Selected values pass the recipe schema and runtime topology checks."
                : validity === "invalid"
-                  ? "Output is hidden until the highlighted fields are fixed."
+                  ? "Output is withheld until the highlighted fields are corrected."
                  : "Values are still being checked.")}
          </p>
          <p>
-            <strong>{label}.</strong>{" "}
-            {weightsBasis
-              ? `${memory?.reason ?? ""} Resident weights need ${memory?.requiredGiB?.toFixed(1)} GiB of ${memory?.capacityGiB?.toFixed(1)} GiB available.`
-              : estimated
-                ? `The analytical estimate needs ${memory?.requiredGiB?.toFixed(1)} GiB of ${memory?.capacityGiB?.toFixed(1)} GiB available. It models weights, encoders, and the configured workload. It is not a measured guarantee.`
-                : (memory?.reason ??
-                  "FastVideo has no capacity model for this checkpoint, so it does not estimate memory use here.")}
+            <strong>{fitLabel[fit]}.</strong>{" "}
+            {fit === "unknown"
+              ? "No memory estimate covers this deployment, device, and request yet. This is not a fit claim."
+              : "The shared capacity API supplied this estimate; it is not a measured guarantee."}
+          </p>
+          <p>
+            <strong>{evidenceLabel[evidence]}.</strong>{" "}
+            {evidence === "none"
+              ? "Successful runs on other checkpoints, settings, software revisions, or GPU layouts do not establish this configuration."
+              : evidence === "stale"
+                ? "A recorded run exists, but edited settings fall outside what it covers."
+                : "A recorded run matches the deployment and settings shown."}
          </p>
          {detail && <p>{detail}</p>}
        </div>
@@ -4,17 +4,12 @@ description: Recipes for video, image, audio, and world generation.
 mode: wide
 ---

-Pick a model, choose your hardware, and copy the command. Every recipe starts from checked-in YAML, so the settings you see are the settings that run.
-
-Before you pick a device, check [weights and quantization](weights/) for measured
-checkpoint sizes and the FastH3 precision variants.
-
 <a id="inference-cookbook"></a>
+<a id="planned-families-heading"></a>
 <a id="video-models-heading"></a>
 <a id="image-models-heading"></a>
 <a id="audio-models-heading"></a>
 <a id="world-models-heading"></a>
-<a id="planned-families-heading"></a>

 <Cookbook />

@@ -1,14 +1,7 @@
 ---
-title: FastH3 Omni v1
-description: Configure FastH3 Omni v1 reference-to-video serving and copy inference or serving commands.
+title: FastH3 OmniRef
+description: Configure FastH3 OmniRef reference-to-video serving on an RTX 5090, RTX PRO 6000, or DGX Spark.
 mode: frame
 ---

 <Cookbook modelId="fasth3-omniref" />
-
-## Apple silicon
-
-The MLX path composes the BF16 Omni release with the NVFP4 text encoder, converts
-the DiT to INT6, and runs the same image and audio references on Metal. Follow
-[FastH3 Omni on MLX](/inference/fasth3_omniref_mlx).
-For copyable local commands, open the [MLX recipe](/cookbook/minimax-h3/models/fasth3-omniref-mlx).
@@ -1,112 +0,0 @@
---
-title: Weights and quantization
-description: Measured checkpoint sizes for every cookbook model, plus the FastH3 quantized variants and which one fits a given GPU.
-mode: wide
---
-
-Every number here is the size of the safetensors files in the model repository,
-measured on 2026-10-27 with the Hugging Face tree API. Treat it as the weight
-floor. A run also needs activations, VSA workspace, and decode buffers, so
-weights that fit do not guarantee a run that fits.
-
-The cookbook editor starts from each recipe's maintained deployment. Changing
-a quantization control does not select one of the published exports below.
-
-## Catalog weights
-
-Sizes are GiB at the checkpoint's runtime precision. Repositories that store
-F32 weights are converted to BF16 here, which is what the loader uses.
-
-| Model                   | Weights | Transformer | Text encoder | VAE  |
-| ----------------------- | ------- | ----------- | ------------ | ---- |
-| FastH3 8-Step V2        | 137.7   | 65.3        | 62.1         | 10.3 |
-| FastH3 4-Step V1        | 137.7   | 65.3        | 62.1         | 10.3 |
-| FastH3 Trim             | 56.8    | 34.9        | 15.3         | 6.5  |
-| FastH3 Omni v1          | 92.5    | 65.3        | 15.3         | 11.9 |
-| Pika H3 Flex            | 136.7   | 64.3        | 62.1         | 10.3 |
-| FastWan2.1 1.3B         | 13.6    | 2.8         | 10.6         | 0.3  |
-| Wan2.1 T2V 14B          | 37.4    | 26.6        | 10.6         | 0.3  |
-| Wan2.1 I2V 14B          | 41.9    | 30.5        | 10.6         | 0.3  |
-| Wan2.2 I2V A14B         | 58.8    | 53.2        | 5.3          | 0.3  |
-| Hunyuan 1.5 I2V 480p    | 32.2    | 15.5        | 13.6         | 2.3  |
-| Hunyuan 1.5 T2V 480p    | 24.9    | 15.5        | 7.0          | 2.4  |
-| Kandinsky 5 Lite        | 11.1    | 2.2         | 8.5          | 0.5  |
-| Kandinsky 6 Pro         | 41.9    | 32.5        | 8.0          | 1.4  |
-| Kandinsky 6 Pro distill | 37.5    | 28.1        | 8.0          | 1.4  |
-| LTX2 Distilled          | 66.9    | 35.2        | 25.4         | 3.6  |
-
-Pika H3 Flex's transformer is the 61.7 GiB H3 trunk plus its 2.6 GiB task adapter.
-Ref2VA swaps in `transformer_ref` and the `ref2va` adapter, which are the same size.
-
-The FastVideo organization publishes the same FastH3 model at several
-precisions. The transformer shrinks with precision; the text encoder and VAE
-only shrink when the repository ships them quantized too.
-
-| Checkpoint        | Variant        | Weights | Transformer | Text encoder | VAE  |
-| ----------------- | -------------- | ------- | ----------- | ------------ | ---- |
-| 8-Step V2         | BF16           | 137.7   | 65.3        | 62.1         | 10.3 |
-| 8-Step V2         | FP8            | 69.2    | 47.4        | 15.3         | 6.5  |
-| 8-Step V2         | NVFP4          | 114.5   | 49.8        | 62.1         | 2.5  |
-| 8-Step V2         | NVFP4 Consumer | 58.7    | 36.9        | 15.3         | 6.5  |
-| Trim              | BF16           | 56.8    | 34.9        | 15.3         | 6.5  |
-| Trim              | FP8            | 41.7    | 19.9        | 15.3         | 6.5  |
-| Trim              | NVFP4          | 32.9    | 11.1        | 15.3         | 6.5  |
-| Trim              | MLX INT6       | 33.0    | 14.3        | 14.2         | 4.5  |
-| 4-Step V1         | BF16           | 137.7   | 65.3        | 62.1         | 10.3 |
-| 4-Step V1         | NVFP4          | 114.5   | 49.8        | 62.1         | 2.5  |
-| Omni v1           | BF16 DiT       | 65.3    | 65.3        | n/a          | n/a  |
-| Omni v1           | FP8            | 67.2    | 47.4        | 15.3         | 4.5  |
-| Omni v1           | NVFP4 QAD      | 63.3    | 37.0        | 15.3         | 11.0 |
-| NVFP4 VAE decoder | quality        | 5.8     | n/a         | n/a          | 5.8  |
-| NVFP4 VAE decoder | fast           | 1.3     | n/a         | n/a          | 1.3  |
-
-The NVFP4 checkpoint keeps the BF16 text encoder, which is why it is larger
-than FP8. The Consumer variant pairs the smaller NVFP4 transformer with the
-FP8 text encoder and VAE, which brings the listed weights to 58.7 GiB. These
-export sizes do not establish device fit. Loading, offload, quantization
-profiles, and temporary buffers affect runtime residency. The decoder-only
-rows exclude the VAE encoder and can contain master weights that the loader
-packs before inference.
-
-## Pick by GPU
-
-| GPU family       | Examples                            | Use                                             | Why                                                               |
-| ---------------- | ----------------------------------- | ----------------------------------------------- | ----------------------------------------------------------------- |
-| Blackwell        | RTX 5090, RTX PRO 6000, B200, GB200 | NVFP4 or FP8 with a matching maintained profile | Native FP4 tensor cores cut the transformer to about half of BF16 |
-| Hopper           | H100, H200                          | FP8                                             | FP8 tensor cores; the NVFP4 path requires Blackwell               |
-| Ada              | RTX 4090, L40S                      | FP8                                             | FP8 works, NVFP4 does not                                         |
-| Ampere and older | A100, RTX 3090                      | BF16                                            | No FP8 or FP4 hardware                                            |
-| Apple silicon    | M-series                            | MLX INT6 or INT8                                | The MLX conversions are INT6 and INT8, not FP4                    |
-| DGX Spark        | GB10                                | NVFP4 with sequential load                      | Unified memory, so the host RAM counts toward the budget          |
-
-## What the cookbook does now
-
-For configurations without transformer quantization or a decoder swap, FastH3
-pages compare checkpoint weights with the selected device. This is a weights
-check, not a request peak. Host-only RAM cannot establish GPU fit. Quantized transformers, TAEH3, and NVFP4 decoder swaps stay
-unestimated until metadata covers the effective layer profile, retained
-components, and packed runtime storage. The editor explains the missing
-coverage instead of substituting an unrelated export or decoder file size.
-
-The transformer quantization control writes
-`generator.engine.quantization.transformer_quant`, so the copied YAML names the
-precision. Encoder and decoder choices use component paths and stay visible in
-the YAML too.
-
-## Automatic selection
-
-The editor does not pick a precision from the device yet. When it does, the
-rule is:
-
-1. Read the device architecture and memory from the hardware catalog.
-2. Prefer the smallest variant the architecture accelerates: NVFP4 on
-   Blackwell, FP8 on Hopper and Ada, INT6 on Apple, BF16 everywhere else.
-3. Keep the BF16 text encoder when the quantized checkpoint ships one, and
-   switch the decoder to NVFP4 only on Blackwell.
-4. Expose the choice as a deployment setting, not a hidden default, so the
-   copied command still shows which checkpoint runs.
-
-That needs one deployment per variant in
-[`docs/cookbook/recipes`](https://github.com/hao-ai-lab/FastVideo/tree/main/docs/cookbook/recipes)
-and a match rule in the editor. Until then, select the precision by hand and
-read the weights fact.
@@ -2,7 +2,7 @@ import { useEffect, useId, useRef, useState, type ReactNode } from "react";
 import { MotionConfig } from "motion/react";
 import { ArrowUpRight } from "lucide-react";
 import { Tabs } from "radix-ui";
-import { modelDirectory, recipePages, recentlyAdded, sitePath, type DirectoryModel } from "@/lib/cookbook-pages";
+import { modelDirectory, recipePages, sitePath, type DirectoryModel } from "@/lib/cookbook-pages";
 import { Badge } from "@/components/ui/badge";
 import { Button } from "@/components/ui/button";
 import { Input } from "@/components/ui/input";
@@ -19,7 +19,6 @@ import {
  CardCurtainRevealTitle,
 } from "@/components/ui/card-curtain-reveal";
 import RecipeFacts from "@/components/cookbook/RecipeFacts";
-import { assessRecipe, assessWeightsRecipe, useCapacityModel, weightsFor } from "@/lib/capacity";
 import { configAndLaunch, inferenceOutput, sdkCommand } from "@/lib/cookbook-commands";
 import {
  gpuCountFrom,
@@ -131,7 +130,7 @@ function FieldControl({
      </div>
      {choices ? (
        <div role="group" aria-label={label} className="option-chips">
-          {value === undefined && !choices.some((choice) => choice === null) && (
+          {value === undefined && (
            <button type="button" aria-pressed={!changed} onClick={() => update(field.path, undefined, field)}>
              Default
            </button>
@@ -143,9 +142,18 @@ function FieldControl({
              aria-pressed={value === choice}
              onClick={() => update(field.path, JSON.stringify(choice), field)}
            >
-              {choice === null ? "Default" : choiceLabel(field.path, choice)}
+              {choiceLabel(field.path, choice)}
            </button>
          ))}
+          {nullable && !choices.includes(null) && (
+            <button
+              type="button"
+              aria-pressed={value === null}
+              onClick={() => update(field.path, "null", { ...field, schema: { ...field.schema, enum: [null] } })}
+            >
+              null
+            </button>
+          )}
        </div>
      ) : kind.includes("object") || kind.includes("array") ? (
        <Textarea
@@ -305,41 +313,6 @@ function Editor({
  const hardwareRecord = resolveHardware(hardware, hardwareCatalog.records);
  const gpuCount = gpuCountFrom(result?.config);
  const model = modelDirectory.find((model) => model.id === catalog.id.split("/")[0])!;
-  const capacityModel = useCapacityModel(model.id);
-  const weights = weightsFor(model.id);
-  const memory =
-    result && !failure
-      ? hardwareRecord
-        ? capacityModel.model
-          ? assessRecipe({
-              model: capacityModel.model,
-              example: { workload: catalog.workload, env: catalog.env },
-              hardware: hardwareRecord,
-              config: result.config,
-              budgetGb: hardwareRecord.usable_budget_gb,
-              hostGb: hardwareRecord.host_ram_gb,
-            })
-          : catalog.runtime.backend === "mlx" && capacityModel.status !== "loading"
-            ? {
-                status: "unknown" as const,
-                reason: "MLX checkpoints use a different precision and load pattern; memory is not estimated.",
-              }
-            : capacityModel.status !== "loading" && weights
-              ? assessWeightsRecipe({
-                  weights,
-                  config: result.config,
-                  hardware: hardwareRecord,
-                  budgetGb: hardwareRecord.usable_budget_gb,
-                })
-              : {
-                  status: "unknown" as const,
-                  reason:
-                    capacityModel.status === "loading"
-                      ? "Loading the memory estimate."
-                      : "No capacity model ships for this checkpoint.",
-                }
-        : { status: "unknown" as const, reason: "Select a device to estimate memory for this recipe." }
-      : undefined;
  const renderField = (field: Field) => (
    <FieldControl
      key={`${reset}-${field.path}`}
@@ -587,7 +560,7 @@ function Editor({
              </a>
            </div>
          </div>
-          <RecipeFacts validity={failure ? "invalid" : result ? "valid" : "loading"} memory={memory} />
+          <RecipeFacts validity={failure ? "invalid" : result ? "valid" : "loading"} fit="unknown" evidence="none" />
        </MatrixRow>
        {catalog.workload === "i2v" && runMode !== "serving" && (
          <MatrixRow title="Inputs">
@@ -608,7 +581,7 @@ function Editor({
                }}
              />
              <p id={`${sourceImageId}-hint`} className="matrix-note">
-                Local path to the first frame. FastVideo reads it on the machine that runs the command.
+                Path to an image on the machine running FastVideo. Included as request.inputs.image_path.
              </p>
            </div>
          </MatrixRow>
@@ -698,11 +671,6 @@ function MatrixRow({ title, children }: { title: string; children: ReactNode })
  );
 }

-function deploymentFromUrl() {
-  const requested = new URL(window.location.href).searchParams.get("deployment");
-  return requested === "fasth3-omniref/cuda-rest" ? "fasth3-omniref/rtx5090-rest" : requested;
-}
-
 export default function Cookbook({ modelId }: { modelId?: string }) {
  const [mounted, setMounted] = useState(false);
  useEffect(() => setMounted(true), []);
@@ -725,7 +693,7 @@ export default function Cookbook({ modelId }: { modelId?: string }) {
          throw new Error("No cookbook deployments are available. Check the catalog and retry.");
        }
        setIndex(data);
-        const requested = deploymentFromUrl();
+        const requested = new URL(window.location.href).searchParams.get("deployment");
        const requestedModel = data.models.find((model) =>
          model.deployments.some((deployment) => deployment.id === requested),
        );
@@ -771,7 +739,7 @@ export default function Cookbook({ modelId }: { modelId?: string }) {
  }
  useEffect(() => {
    const restore = () => {
-      const requested = deploymentFromUrl();
+      const requested = new URL(window.location.href).searchParams.get("deployment");
      const available = index?.models.find((item) => item.id === modelId)?.deployments ?? [];
      const next = available.find((deployment) => deployment.id === requested)?.id ?? available[0]?.id;
      if (!next) return;
@@ -838,8 +806,11 @@ export default function Cookbook({ modelId }: { modelId?: string }) {
 }

 const GPU_COUNT_PATH = "generator.engine.num_gpus";
-function singleGpuDevice(record: { form_factor?: string; kind?: string }) {
-  return record.kind === "system" || record.form_factor === "soc" || record.form_factor === "laptop";
+const SINGLE_GPU_FORM_FACTOR = /desktop|laptop|workstation|system|soc/;
+
+/** Desktop and workstation cards, DGX Spark, and Macs are one-GPU machines; datacenter parts keep the recipe default. */
+function singleGpuDevice(record: { form_factor?: string }) {
+  return SINGLE_GPU_FORM_FACTOR.test(record.form_factor ?? "");
 }

 /** Returns new edits where every fixed parallel size that no longer divides `gpus` becomes automatic. */
@@ -862,13 +833,6 @@ function deploymentLabel(deployment: DeploymentChoice, deployments: DeploymentCh
 }

 /** Mirrors the editor: MLX and WebSocket deployments offer serving only. */
-function modelModes(index: Index | undefined, model: DirectoryModel): string[] | undefined {
-  const modes = model.configurable ? configurableModes(index, model.id) : exampleModes(model);
-  const hasMlx = model.recipes.some((recipe) => recipe.gpu_types.includes("Apple Silicon"));
-  if (hasMlx && modes && !modes.some((mode) => /mlx/i.test(mode))) return [...modes, "MLX"];
-  return modes;
-}
-
 function configurableModes(index: Index | undefined, modelId: string): string[] | undefined {
  const deployments = index?.models.find((model) => model.id === modelId)?.deployments;
  if (!deployments?.length) return undefined;
@@ -883,20 +847,8 @@ function exampleModes(model: DirectoryModel): string[] {
  return ["Python SDK", "CLI", ...(model.recipes.some((recipe) => recipe.serving) ? ["Serving"] : [])];
 }

-const FASTH3_SECTION = new Set(["fasth3-omniref", "fasth3-8step", "fasth3-4step", "fasth3-trim"]);
-
 /** Rest state shows identity only; hover or keyboard focus opens a curtain over the details. */
-function RecipeCard({
-  model,
-  href,
-  modes,
-  badge,
-}: {
-  model: DirectoryModel;
-  href: string;
-  modes?: string[];
-  badge?: string;
-}) {
+function RecipeCard({ model, href, modes }: { model: DirectoryModel; href: string; modes?: string[] }) {
  const [focused, setFocused] = useState(false);
  const details = useRef<HTMLDivElement>(null);
  const [restOffset, setRestOffset] = useState(72);
@@ -917,7 +869,6 @@ function RecipeCard({
            <img src={sitePath(`/assets/logos/${model.logo}.webp`)} alt="" width="40" height="40" />
            <span className="recipe-card-kicker">{model.provider}</span>
            {model.featured ? <span className="recipe-featured">Featured</span> : null}
-            {badge ? <span className="recipe-featured recipe-new">{badge}</span> : null}
          </div>
          <CardCurtainRevealTitle
            as={model.configurable ? "h3" : "h4"}
@@ -981,7 +932,6 @@ function RecipeDirectory({ index, error, retry }: { index?: Index; error: string
  };
  const visible = modelDirectory.filter(
    (model) =>
-      !model.hidden &&
      (category === "All models" || model.category === category) &&
      (family === "All families" || model.family === family) &&
      (task === "All tasks" || model.task === task) &&
@@ -997,10 +947,9 @@ function RecipeDirectory({ index, error, retry }: { index?: Index; error: string
    setRuntime("all");
    setTask("All tasks");
  }
-  const fastH3 = visible.filter((model) => FASTH3_SECTION.has(model.id));
-  const recent = visible.filter((model) => !FASTH3_SECTION.has(model.id) && recentlyAdded(model.model));
-  const rest = visible.filter((model) => !FASTH3_SECTION.has(model.id) && !recentlyAdded(model.model));
-  const families = [...new Set(rest.map((model) => model.family))];
+  const configurable = visible.filter((model) => model.configurable);
+  const examples = visible.filter((model) => !model.configurable);
+  const exampleFamilies = [...new Set(examples.map((model) => model.family))];
  const filtered =
    !!query.trim() ||
    category !== "All models" ||
@@ -1008,7 +957,7 @@ function RecipeDirectory({ index, error, retry }: { index?: Index; error: string
    runtime !== "all" ||
    task !== "All tasks";
  const categoryCount = (value: string) =>
-    modelDirectory.filter((model) => !model.hidden && (value === "All models" || model.category === value)).length;
+    modelDirectory.filter((model) => value === "All models" || model.category === value).length;
  return (
    <MotionConfig reducedMotion="user">
      <div className="cookbook not-prose recipe-directory">
@@ -1042,6 +991,7 @@ function RecipeDirectory({ index, error, retry }: { index?: Index; error: string
        <div className="recipe-directory-subbar">
          <p role="status">
            {visible.length} {visible.length === 1 ? "model" : "models"}
+            {configurable.length ? ` · ${configurable.length} configurable` : ""}
            {runtime !== "all" ? ` · ${runtime === "mlx" ? "MLX" : "CUDA"} only` : ""}
            {task !== "All tasks" ? ` · ${task}` : ""}
            {filtered && visible.length ? (
@@ -1093,62 +1043,59 @@ function RecipeDirectory({ index, error, retry }: { index?: Index; error: string
            </Button>
          </div>
        ) : null}
-        {fastH3.length ? (
-          <section className="recipe-section" aria-labelledby="recipe-fasth3-heading">
+        {configurable.length ? (
+          <section className="recipe-section" aria-labelledby="recipe-configurable-heading">
            <header className="recipe-section-head">
-              <h2 id="recipe-fasth3-heading">FastVideo FastH3</h2>
+              <h2 id="recipe-configurable-heading">Configurable recipes</h2>
+              <p>Edit settings, then copy one config for the Python SDK, the CLI, or an OpenAI-compatible server.</p>
            </header>
            <div className="recipe-card-grid recipe-card-grid-configurable">
-              {fastH3.map((model) => (
-                <RecipeCard key={model.id} model={model} href={sitePath(model.path)} modes={modelModes(index, model)} />
-              ))}
-            </div>
-          </section>
-        ) : null}
-        {recent.length ? (
-          <section className="recipe-section" aria-labelledby="recipe-new-heading">
-            <header className="recipe-section-head">
-              <h2 id="recipe-new-heading">New</h2>
-            </header>
-            <div className="recipe-card-grid recipe-card-grid-configurable">
-              {recent.map((model) => (
+              {configurable.map((model) => (
                <RecipeCard
                  key={model.id}
                  model={model}
                  href={sitePath(model.path)}
-                  modes={modelModes(index, model)}
-                  badge="New"
+                  modes={configurableModes(index, model.id)}
                />
              ))}
            </div>
          </section>
        ) : null}
-        {families.map((name) => {
-          const members = rest.filter((model) => model.family === name);
-          return (
-            <section className="recipe-section" key={name} aria-label={name}>
-              <div className="recipe-family-group">
-                <div className="recipe-family-head">
-                  <h3>{name}</h3>
-                  <span>
-                    {members.length} {members.length === 1 ? "recipe" : "recipes"}
-                  </span>
-                  <a href={sitePath(members[0].familyPath)}>Family guide →</a>
+        {examples.length ? (
+          <section className="recipe-section" aria-labelledby="recipe-examples-heading">
+            <header className="recipe-section-head">
+              <h2 id="recipe-examples-heading">Example recipes</h2>
+              <p>
+                Maintained example scripts, grouped by model family. Adjust their documented settings and run them as
+                written.
+              </p>
+            </header>
+            {exampleFamilies.map((name) => {
+              const members = examples.filter((model) => model.family === name);
+              return (
+                <div className="recipe-family-group" key={name}>
+                  <div className="recipe-family-head">
+                    <h3>{name}</h3>
+                    <span>
+                      {members.length} {members.length === 1 ? "recipe" : "recipes"}
+                    </span>
+                    <a href={sitePath(members[0].familyPath)}>Family guide →</a>
+                  </div>
+                  <div className="recipe-card-grid">
+                    {members.map((model) => (
+                      <RecipeCard
+                        key={model.id}
+                        model={model}
+                        href={sitePath(model.path)}
+                        modes={exampleModes(model)}
+                      />
+                    ))}
+                  </div>
                </div>
-                <div className="recipe-card-grid">
-                  {members.map((model) => (
-                    <RecipeCard
-                      key={model.id}
-                      model={model}
-                      href={sitePath(model.path)}
-                      modes={modelModes(index, model)}
-                    />
-                  ))}
-                </div>
-              </div>
-            </section>
-          );
-        })}
+              );
+            })}
+          </section>
+        ) : null}
        {!visible.length ? (
          <div className="recipe-empty">
            <h2>No matching recipes</h2>
@@ -14,7 +14,7 @@ export default function FastModels() {
      models: group.models.filter(
        (model) =>
          (family === "All models" || family === group.name) &&
-          `${group.name} ${model.id}`.toLowerCase().includes(query.trim().toLowerCase()),
+          `${group.name} ${group.summary} ${model.id}`.toLowerCase().includes(query.trim().toLowerCase()),
      ),
    }))
    .filter((group) => group.models.length);
@@ -48,7 +48,7 @@ export default function FastModels() {
        <header className="model-catalog-heading">
          <div>
            <h2 id="checkpoint-catalog">Browse checkpoints</h2>
-            <p>Pick the checkpoint that matches your runtime.</p>
+            <p>Find the export your runtime needs.</p>
          </div>
          <a href="https://huggingface.co/collections/FastVideo/fastvideo-fasth3">
            Earlier FastH3 releases <span aria-hidden="true">↗</span>
@@ -67,7 +67,7 @@ export default function FastModels() {
            <input
              id={id}
              type="search"
-              placeholder="Name, family, or precision…"
+              placeholder="Name, precision, or runtime…"
              value={query}
              onChange={(event) => setQuery(event.target.value)}
            />
@@ -4,18 +4,17 @@ import { modelDirectory, sitePath } from "@/lib/cookbook-pages";
 export default function RecipeFamily({ family }: { family: string }) {
  const models = modelDirectory.filter(
    (model) =>
-      !model.hidden &&
-      (model.recipes.some((recipe) => recipe.family === family) ||
-        (family === "minimax_h3" && model.family === "FastH3")),
+      model.recipes.some((recipe) => recipe.family === family) ||
+      (family === "minimax_h3" && model.family === "FastH3"),
  );
  return (
    <div className="recipe-directory not-prose">
      <a className="recipe-back" href={sitePath("/cookbook")}>
-        ← All recipes
+        ← All models
      </a>
      <div className="recipe-card-grid">
        {models.map((model) => (
-          <a className="recipe-card recipe-card-static" key={model.id} href={sitePath(model.path)}>
+          <a className="recipe-card" key={model.id} href={sitePath(model.path)}>
            <div className="recipe-card-top">
              <img src={sitePath(`/assets/logos/${model.logo}.webp`)} alt="" width="40" height="40" />
              <span className="recipe-card-kicker">{model.provider}</span>
@@ -484,7 +484,9 @@ export default function SourceRecipe({ model }: { model: DirectoryModel }) {
            <RecipeFacts
              validity={failure ? "invalid" : example ? "valid" : "loading"}
              validityText={yamlSource ? "YAML parses" : "Example settings applied"}
-              validityDetail="Checks that the YAML parses and that edited numbers are valid. It does not run the recipe."
+              validityDetail="This page checks that the YAML parses and that edited numeric fields are finite and supported by the example. It does not run the full recipe schema or runtime topology validation."
+              fit="unknown"
+              evidence="none"
              subject={yamlSource ? "configuration" : "example"}
            />
          </div>
@@ -1,226 +0,0 @@
-import { useEffect, useState } from "react";
-// @ts-expect-error The shared ES module has no TypeScript declarations.
-import { assessCapacity } from "../../docs/js/cookbook-memory.mjs";
-import weightsData from "../../docs/assets/cookbook-weights.json";
-
-export type MemoryStatus = "fits" | "tight" | "exceeds" | "unsupported" | "unknown";
-
-/** The subset of an `assessCapacity` result the UI shows. */
-export interface MemoryAssessment {
-  status: MemoryStatus;
-  reason?: string;
-  requiredGiB?: number;
-  capacityGiB?: number;
-  basis?: string;
-  parts?: Record<string, number>;
-}
-
-export interface WeightsModel {
-  parts: Record<string, number>;
-  alternates?: Record<string, number>;
-  sources?: string[];
-}
-
-const WEIGHTS = weightsData as { models: Record<string, WeightsModel> };
-
-/** Measured checkpoint sizes for recipes that have no full capacity model yet. */
-export function weightsFor(modelId: string | undefined): WeightsModel | undefined {
-  return modelId ? WEIGHTS.models[modelId] : undefined;
-}
-
-/**
- * Compare resident checkpoint weights with one device. This answers whether the
- * weights load, not whether a request peaks below the device limit: activations,
- * VSA workspace, and decode buffers are not modeled.
- */
-export function assessWeightsRecipe({
-  weights,
-  config,
-  hardware,
-  budgetGb,
-}: {
-  weights: WeightsModel;
-  config: Record<string, unknown>;
-  hardware: { memory_gb?: number; usable_budget_gb?: number; memory_unit?: string; memory_type?: string };
-  budgetGb?: number;
-}): MemoryAssessment | undefined {
-  if (hardware.memory_type !== "dedicated" && hardware.memory_type !== "unified") {
-    return {
-      status: "unknown",
-      reason:
-        "A weights fit check requires dedicated GPU memory or unified memory. Host-only system RAM is not GPU capacity.",
-    };
-  }
-  const generator = (config.generator ?? {}) as {
-    engine?: { offload?: Record<string, boolean> };
-    pipeline?: {
-      components?: Record<string, string | null>;
-      experimental?: Record<string, unknown>;
-    };
-  };
-  const parts = { ...weights.parts };
-  const components = generator.pipeline?.components ?? {};
-  const experimental = generator.pipeline?.experimental ?? {};
-  if (experimental.video_decode_backend === "taeh3") {
-    return {
-      status: "unknown",
-      reason: "TAEH3 decoder weights are not modeled. This backend does not load the bundled H3 video VAE.",
-    };
-  }
-  if (experimental.h3_vae_nvfp4_decoder) {
-    return {
-      status: "unknown",
-      reason:
-        "NVFP4 decoder residency is not modeled: the VAE retains its encoder and packs decoder weights after loading. Decoder file size is not the resident VAE size.",
-    };
-  }
-  if (components.text_encoder_weights && weights.alternates?.text_encoder_quantized !== undefined) {
-    parts.text_encoder = weights.alternates.text_encoder_quantized;
-  }
-  const quant = (config.generator as { engine?: { quantization?: { transformer_quant?: string | null } } })?.engine
-    ?.quantization?.transformer_quant;
-  if (quant) {
-    return {
-      status: "unknown",
-      reason:
-        "Transformer quantization residency is not modeled for this checkpoint and layer profile. Published export sizes do not describe on-the-fly quantization of this configuration.",
-    };
-  }
-  const offload = generator.engine?.offload ?? {};
-  if (offload.text_encoder) delete parts.text_encoder;
-  const resident = Object.values(parts).reduce((sum, value) => sum + value, 0);
-  const capacityGb = budgetGb ?? hardware.usable_budget_gb ?? hardware.memory_gb;
-  if (!capacityGb) return undefined;
-  const capacityGiB = hardware.memory_unit === "GiB" ? capacityGb : (capacityGb * 1e9) / GIB;
-  const label: Record<string, string> = {
-    transformer: "transformer",
-    text_encoder: "text encoder",
-    vae: "VAE",
-    audio_vae: "audio VAE",
-  };
-  const split = Object.entries(parts)
-    .map(([role, value]) => `${label[role] ?? role} ${value.toFixed(1)} GiB`)
-    .join(", ");
-  const status: MemoryStatus = resident > capacityGiB ? "exceeds" : resident > capacityGiB * 0.85 ? "tight" : "fits";
-  const offloaded = offload.text_encoder ? " The text encoder is offloaded." : "";
-  return {
-    status,
-    basis: "weights",
-    requiredGiB: resident,
-    capacityGiB,
-    parts,
-    reason: `Weights only: ${split}.${offloaded} Activations, VSA workspace, and decode buffers are not modeled.`,
-  };
-}
-
-interface CapacityModel {
-  id: string;
-  name?: string;
-  examples?: { deployment_id?: string }[];
-  reference_runs?: unknown[];
-  [key: string]: unknown;
-}
-
-interface CapacityState {
-  status: "loading" | "ready" | "unavailable";
-  model?: CapacityModel;
-}
-
-const CAPACITY_INDEX = `${import.meta.env.BASE_URL.replace(/\/$/, "")}/assets/cookbook-config/capacity/index.json`;
-const GIB = 1024 ** 3;
-
-let indexPromise: Promise<{ models: { id: string; catalog_url: string }[] }> | null = null;
-
-function loadIndex() {
-  indexPromise ??= fetch(CAPACITY_INDEX)
-    .then((response) => {
-      if (!response.ok) throw new Error(`Capacity index: HTTP ${response.status}`);
-      return response.json() as Promise<{ models: { id: string; catalog_url: string }[] }>;
-    })
-    .catch((error) => {
-      indexPromise = null;
-      throw error;
-    });
-  return indexPromise;
-}
-
-/**
- * Load the capacity model for one cookbook entry. Most recipes do not have one;
- * the hook stays `unavailable` and the UI says so instead of guessing.
- */
-export function useCapacityModel(modelId: string | undefined): CapacityState {
-  const [state, setState] = useState<CapacityState>({ status: "loading" });
-  useEffect(() => {
-    if (!modelId) {
-      setState({ status: "unavailable" });
-      return;
-    }
-    let live = true;
-    setState({ status: "loading" });
-    loadIndex()
-      .then(async (index) => {
-        const entry = index.models.find((model) => model.id === modelId);
-        if (!entry) throw new Error("No capacity model ships for this checkpoint.");
-        const url = new URL(entry.catalog_url, new URL(CAPACITY_INDEX, window.location.origin));
-        const response = await fetch(url);
-        if (!response.ok) throw new Error(`Capacity model: HTTP ${response.status}`);
-        const model = (await response.json()) as CapacityModel;
-        if (live) setState({ status: "ready", model });
-      })
-      .catch(() => {
-        if (live) setState({ status: "unavailable" });
-      });
-    return () => {
-      live = false;
-    };
-  }, [modelId]);
-  return state;
-}
-
-/**
- * Run the analytical estimate for the edited configuration. Returns `undefined`
- * when no capacity model covers the checkpoint, so the caller can say so.
- */
-export function assessRecipe({
-  model,
-  example,
-  hardware,
-  config,
-  budgetGb,
-  hostGb,
-}: {
-  model: CapacityModel;
-  example: { workload?: string; env?: Record<string, string> };
-  hardware?: {
-    memory_gb?: number;
-    memory_unit?: string;
-    memory_type?: string;
-    usable_budget_gb?: number;
-    host_ram_gb?: number;
-    id?: string;
-  };
-  config?: Record<string, unknown>;
-  budgetGb?: number;
-  hostGb?: number;
-}): MemoryAssessment | undefined {
-  if (!config || !hardware?.memory_gb) return undefined;
-  try {
-    const result = assessCapacity({
-      model,
-      example,
-      hardware,
-      config,
-      memoryBudgetGiB: budgetGb ? (hardware.memory_unit === "GiB" ? budgetGb : (budgetGb * 1e9) / GIB) : undefined,
-      hostMemoryGb: hostGb,
-    }) as MemoryAssessment;
-    return {
-      status: (result.status ?? "unknown") as MemoryStatus,
-      reason: result.reason,
-      requiredGiB: result.requiredGiB,
-      capacityGiB: result.capacityGiB,
-      basis: result.basis,
-    };
-  } catch {
-    return undefined;
-  }
-}
@@ -2,15 +2,6 @@ import sourceCatalog from "../../docs/assets/cookbook-recipes.json" with { type:

 /** Presentation metadata only. Deployments and editable fields come from the cookbook backend. */
 export const recipePages = [
-  {
-    id: "fasth3-preview",
-    family: "MiniMax H3",
-    title: "FastH3 Preview V0.2",
-    mark: "H3",
-    path: "/cookbook/minimax-h3/fasth3-preview",
-    description: "The few-step FastH3 preview checkpoint on CUDA or native Apple Silicon MLX.",
-    tags: ["Text to video", "Audio", "4-step"],
-  },
  {
    id: "fasth3-8step",
    family: "MiniMax H3",
@@ -29,6 +20,15 @@ export const recipePages = [
    description: "The smaller FastH3 model for eight-step video generation with synchronized audio.",
    tags: ["Text to video", "Audio", "NVFP4", "8-step"],
  },
+  {
+    id: "fasth3-preview",
+    family: "MiniMax H3",
+    title: "FastH3 Preview V0.2",
+    mark: "H3",
+    path: "/cookbook/minimax-h3/fasth3-preview",
+    description: "The few-step FastH3 preview checkpoint on CUDA or native Apple Silicon MLX.",
+    tags: ["Text to video", "Audio", "4-step"],
+  },
  {
    id: "fasth3-4step",
    family: "MiniMax H3",
@@ -41,7 +41,7 @@ export const recipePages = [
  {
    id: "fasth3-omniref",
    family: "MiniMax H3",
-    title: "FastH3 Omni v1",
+    title: "FastH3 OmniRef",
    mark: "H3",
    path: "/cookbook/minimax-h3/fasth3-omniref",
    description: "Video and audio from a prompt plus ordered image, video, and audio references.",
@@ -128,16 +128,6 @@ const backendModels: Record<string, string> = {
  "Wan-AI/Wan2.1-I2V-14B-480P-Diffusers": "wan21-i2v",
  "FastVideo/LTX2-Distilled-Diffusers": "ltx2-distilled",
 };
-const hiddenBackends = new Set(["fasth3-preview"]);
-
-const modelAddedOn: Record<string, string> = {
-  "Qwen/Qwen-Image-2.1": "2026-10-25",
-};
-export function recentlyAdded(model: string) {
-  const added = modelAddedOn[model];
-  return !!added && Date.now() - Date.parse(added) < 30 * 86_400_000;
-}
-
 const TASK_GROUPS: Array<[RegExp, string]> = [
  [/world/, "Text to world"],
  [/video\/text to audio|video.*to audio/, "Video to audio"],
@@ -157,69 +147,57 @@ export function taskLabel(value: string): string {
  return TASK_GROUPS.find(([pattern]) => pattern.test(normalized))?.[1] ?? "Text to video";
 }

-const mappedModels = [...new Set(sourceRecipes.map((recipe) => recipe.model))].map((model) => {
-  const recipes = sourceRecipes.filter((recipe) => recipe.model === model);
-  const first = recipes[0];
-  const family = familyInfo[first.family];
-  const backend = recipePages.find((page) => page.id === backendModels[model]);
-  const ownModel = model.startsWith("FastVideo/Fast");
-  const logo = ownModel ? "fastvideo" : family.logo;
-  const providers: Record<string, string> = {
-    fastvideo: "FastVideo",
-    minimax: "MiniMax",
-    "wan-ai": "Wan AI",
-    ltx: "Lightricks",
-    "tencent-hunyuan": "Tencent",
-    kandinsky: "Kandinsky Lab",
-    "black-forest-labs": "Black Forest Labs",
-    qwen: "Qwen",
-    stabilityai: "Stability AI",
-    nvidia: "NVIDIA",
-  };
-  return {
-    id: backend?.id ?? first.id,
-    model,
-    recipes,
-    title:
-      backend?.title ??
-      (model.includes("FastH3-4-step")
-        ? "FastH3 4-Step V1"
-        : model.includes("FastH3-Preview-v0.2")
-          ? "FastH3 Preview V0.2"
-          : model === "MiniMaxAI/MiniMax-H3"
-            ? "MiniMax H3"
-            : first.label.split(" on ")[0]),
-    family: backend?.family ?? family.title,
-    familyPath: `/cookbook/${family.path}`,
-    path: backend?.path ?? `/cookbook/${family.path}/models/${first.id}`,
-    logo,
-    provider: providers[logo],
-    category: family.category,
-    description: backend?.description ?? first.summary,
-    tags: backend?.tags ?? [...new Set(recipes.map((recipe) => recipe.task))],
-    task: taskLabel(backend?.tags[0] ?? first.task),
-    platforms: [...new Set(recipes.flatMap((recipe) => recipe.gpu_types))],
-    configurable: !!backend,
-    hidden: hiddenBackends.has(backend?.id ?? ""),
-    featured: backend?.featured ?? false,
-  };
-});
-
-const byPage = new Map<string, (typeof mappedModels)[number]>();
-for (const entry of mappedModels) {
-  const existing = byPage.get(entry.id);
-  if (!existing) {
-    byPage.set(entry.id, entry);
-    continue;
-  }
-  existing.recipes.push(...entry.recipes.filter((recipe) => !existing.recipes.some((seen) => seen.id === recipe.id)));
-  existing.platforms = [...new Set([...existing.platforms, ...entry.platforms])];
-}
-
-export const modelDirectory = [...byPage.values()].sort(
-  (a, b) =>
-    Number(b.featured) - Number(a.featured) ||
-    Number(b.family === "MiniMax H3") - Number(a.family === "MiniMax H3") ||
-    Number(b.configurable) - Number(a.configurable),
-);
+export const modelDirectory = [...new Set(sourceRecipes.map((recipe) => recipe.model))]
+  .map((model) => {
+    const recipes = sourceRecipes.filter((recipe) => recipe.model === model);
+    const first = recipes[0];
+    const family = familyInfo[first.family];
+    const backend = recipePages.find((page) => page.id === backendModels[model]);
+    const ownModel = model.startsWith("FastVideo/Fast");
+    const logo = ownModel ? "fastvideo" : family.logo;
+    const providers: Record<string, string> = {
+      fastvideo: "FastVideo",
+      minimax: "MiniMax",
+      "wan-ai": "Wan AI",
+      ltx: "Lightricks",
+      "tencent-hunyuan": "Tencent",
+      kandinsky: "Kandinsky Lab",
+      "black-forest-labs": "Black Forest Labs",
+      qwen: "Qwen",
+      stabilityai: "Stability AI",
+      nvidia: "NVIDIA",
+    };
+    return {
+      id: backend?.id ?? first.id,
+      model,
+      recipes,
+      title:
+        backend?.title ??
+        (model.includes("FastH3-4-step")
+          ? "FastH3 4-Step V1"
+          : model.includes("FastH3-Preview-v0.2")
+            ? "FastH3 Preview V0.2"
+            : model === "MiniMaxAI/MiniMax-H3"
+              ? "MiniMax H3"
+              : first.label.split(" on ")[0]),
+      family: backend?.family ?? family.title,
+      familyPath: `/cookbook/${family.path}`,
+      path: backend?.path ?? `/cookbook/${family.path}/models/${first.id}`,
+      logo,
+      provider: providers[logo],
+      category: family.category,
+      description: backend?.description ?? first.summary,
+      tags: backend?.tags ?? [...new Set(recipes.map((recipe) => recipe.task))],
+      task: taskLabel(backend?.tags[0] ?? first.task),
+      platforms: [...new Set(recipes.flatMap((recipe) => recipe.gpu_types))],
+      configurable: !!backend,
+      featured: backend?.featured ?? false,
+    };
+  })
+  .sort(
+    (a, b) =>
+      Number(b.featured) - Number(a.featured) ||
+      Number(b.family === "MiniMax H3") - Number(a.family === "MiniMax H3") ||
+      Number(b.configurable) - Number(a.configurable),
+  );
 export type DirectoryModel = (typeof modelDirectory)[number];
@@ -1,45 +0,0 @@
-const NEWS_HEADINGS = ["news", "recent updates", "what's new", "latest news", "announcements"];
-
-const PROJECT_HEADINGS = [
-  "awesome work using fastvideo or our research projects",
-  "projects using fastvideo",
-  "built on fastvideo",
-  "awesome work",
-];
-
-function normalizeHeading(line) {
-  return line
-    .replace(/^#{2,3}\s+/, "")
-    .replace(/[`*_]/g, "")
-    .replace(/[^\p{L}\p{N}\s'’-]/gu, "")
-    .replace(/\s+/g, " ")
-    .trim()
-    .toLowerCase();
-}
-
-function sectionBody(readme, names) {
-  const lines = readme.split("\n");
-  const start = lines.findIndex((line) => /^##\s+/.test(line) && names.includes(normalizeHeading(line)));
-  if (start < 0) return null;
-  const rest = lines.slice(start + 1);
-  const end = rest.findIndex((line) => /^##\s+/.test(line));
-  return (end < 0 ? rest : rest.slice(0, end)).join("\n").trim();
-}
-
-const BULLET = /^-\s+(?:`(\d{4})[/-](\d{2})[/-](\d{2})`:|\*\*(\d{4})[/-](\d{2})[/-](\d{2}):?\*\*:?)\s*(.+)$/;
-
-export function readReadmeSections(readme) {
-  const news = sectionBody(readme, NEWS_HEADINGS);
-  const projects = sectionBody(readme, PROJECT_HEADINGS);
-  if (news === null || projects === null) return { news: null, projects: null };
-  const items = [];
-  for (const line of news.split("\n")) {
-    const match = BULLET.exec(line.trim());
-    if (!match) continue;
-    const [, y1, m1, d1, y2, m2, d2, text] = match;
-    items.push({ date: `${y1 ?? y2}/${m1 ?? m2}/${d1 ?? d2}`, text: text.trim() });
-  }
-  return { news: items, projects };
-}
-
-export const README_HEADINGS = { news: NEWS_HEADINGS, projects: PROJECT_HEADINGS };
@@ -141,17 +141,21 @@
        "label": "All recipes",
        "href": "/cookbook"
      },
+      {
+        "label": "H3 server and playground",
+        "href": "/cookbook/openai-api"
+      },
      {
        "label": "MiniMax H3",
        "display": "group",
-        "collapsed": true,
+        "collapsed": false,
        "items": [
          {
            "label": "MiniMax H3",
            "href": "/cookbook/minimax-h3/models/minimax-h3-t2v"
          },
          {
-            "label": "FastH3 Omni v1",
+            "label": "FastH3 OmniRef",
            "href": "/cookbook/minimax-h3/fasth3-omniref"
          },
          {
@@ -162,6 +166,10 @@
            "label": "FastH3 Trim",
            "href": "/cookbook/minimax-h3/fasth3-trim"
          },
+          {
+            "label": "FastH3 Preview V0.2",
+            "href": "/cookbook/minimax-h3/fasth3-preview"
+          },
          {
            "label": "FastH3 4-Step V1",
            "href": "/cookbook/minimax-h3/fasth3-4step"
@@ -179,7 +187,7 @@
      {
        "label": "Wan",
        "display": "group",
-        "collapsed": true,
+        "collapsed": false,
        "items": [
          {
            "label": "FastWan2.1 1.3B",
@@ -15,7 +15,7 @@
    "test": "playwright test",
    "format": "prettier --write .",
    "format:check": "prettier --check .",
-    "test:content": "node --test tests/content.test.mjs tests/commands.test.mjs tests/readme-sections.test.mjs"
+    "test:content": "node --test tests/content.test.mjs tests/commands.test.mjs"
  },
  "dependencies": {
    "blume": "2.2.3",
@@ -6,7 +6,6 @@ import { cp, mkdir, readdir, readFile, rm, stat, writeFile } from "node:fs/promi
 import { dirname, join, relative, resolve, sep } from "node:path";
 import { fileURLToPath } from "node:url";
 import { parse, stringify } from "yaml";
-import { readReadmeSections, README_HEADINGS } from "../lib/readme.mjs";

 const site = fileURLToPath(new URL("../", import.meta.url));
 const repo = resolve(site, "..");
@@ -177,11 +176,12 @@ export async function transferDocs() {
      return (await stat(path)).isDirectory() || /\.mdx?$/.test(path);
    },
  });
+  // Native callout directives and islands compile as MDX, while source links keep .md filenames.
  for (const file of await files(content)) {
    if (!file.endsWith(".md")) continue;
    const text = await readFile(file, "utf8");
    if (/^:::(?:note|tip|warning|danger|info|success)\b/m.test(text) || text.includes("<RecipeFamily ")) {
-      await writeFile(file.replace(/\.md$/, ".mdx"), text.replace(/\{#([^}]+)\}/g, "[#$1]"));
+      await writeFile(file.replace(/\.md$/, ".mdx"), text);
      await rm(file);
    }
  }
@@ -251,15 +251,10 @@ export async function transferDocs() {

  // The README owns announcements and the community list; the welcome page renders a short news digest.
  const readme = await readFile(join(repo, "README.md"), "utf8");
-  const { news, projects } = readReadmeSections(readme);
-  if (!news || !projects) {
-    const accepted = (names) => names.map((name) => `"## ${name.replace(/\b\w/g, (c) => c.toUpperCase())}"`).join(", ");
-    throw new Error(
-      `README section not found. Accepted news headings: ${accepted(README_HEADINGS.news)}. ` +
-        `Accepted project headings: ${accepted(README_HEADINGS.projects)}.`,
-    );
-  }
-  const recent = news.slice(0, 2).map(({ date, text }) => {
+  const news = readme.match(/^## NEWS\n([\s\S]*?)(?=^##|^###)/m)?.[1];
+  const projects = readme.match(/^## Awesome work using FastVideo or our research projects\n([\s\S]*?)(?=^## )/m)?.[1];
+  if (!news || !projects) throw new Error("The README news or project section boundary needs updating.");
+  const recent = [...news.matchAll(/^- `([\d/]+)`: (.+)$/gm)].slice(0, 2).map(([, date, text]) => {
    const links = [...text.matchAll(/\[([^\]]+)\]\(([^)]+)\)/g)];
    const headline = text
      .replace(/\[([^\]]+)\]\([^)]+\)/g, "$1")
@@ -31,7 +31,7 @@ test("converted guides render admonitions, preserve content, and publish local i
  const home = await readFile(new URL("dist/index.html", site), "utf8");
  assert.ok(home.includes("Welcome to FastVideo"));
  assert.ok(home.includes("/FastVideo/assets/logos/logo.svg"));
-  assert.ok(home.includes("FastVideo is an inference and post-training framework for video diffusion models."));
+  assert.ok(home.includes("FastVideo is an inference and post-training framework for diffusion models."));
  assert.ok(!home.includes("About FastVideo and performance overview"));
  await readFile(new URL("dist/assets/logos/logo.svg", site));
 });
@@ -223,18 +223,18 @@ test("inherited values, explicit false, zero and null retain their meanings", as
 test("catalog discovery opens dedicated model recipes and filters by family and task", async ({ page }) => {
  await page.goto("/FastVideo/cookbook");
  const cards = page.locator(".recipe-card");
-  await expect(cards).toHaveCount(36);
-  await expect(cards.first()).toContainText("FastH3 Omni v1");
+  await expect(cards).toHaveCount(37);
+  await expect(cards.first()).toContainText("FastH3 OmniRef");
  await page.getByLabel("Model family", { exact: true }).selectOption("Wan");
  await page.getByLabel("Search recipes").fill("Wan2.1 I2V");
  await expect(cards).toHaveCount(1);
  await cards.click();
-  await expect(page).toHaveURL(/cookbook\/wan\/wan21-i2v\/?$/);
+  await expect(page).toHaveURL(/cookbook\/wan\/wan21-i2v$/);
  await expect(page.getByLabel("Base model")).toHaveCount(0);
  await page.getByRole("link", { name: "← All recipes", exact: true }).click();
  await page.getByLabel("Search recipes").fill("FastWan2.1");
  await page.locator(".recipe-card").click();
-  await expect(page).toHaveURL(/cookbook\/wan\/fastwan21\/?$/);
+  await expect(page).toHaveURL(/cookbook\/wan\/fastwan21$/);
  await expect(page.getByRole("heading", { level: 1 })).toHaveText("FastWan2.1 1.3B");
 });

@@ -253,7 +253,8 @@ test("hardware preferences stay distinct from schema-backed GPU count and topolo
  await expect(yaml).toContainText('"num_gpus": 1');
  await expect(yaml).not.toContainText('"hardware"');
  await expect(page.getByText("Valid configuration", { exact: true })).toBeVisible();
-  await expect(page.getByText("Weights fit", { exact: true })).toBeVisible();
+  await expect(page.getByText("Fit unknown", { exact: true })).toBeVisible();
+  await expect(page.getByText("No recorded run", { exact: true })).toBeVisible();
  await page.getByLabel("GPU count", { exact: true }).fill("0");
  await expect(yaml).toHaveCount(0);
  await expect(page.getByLabel("GPU count", { exact: true })).toHaveAttribute("aria-invalid", "true");
@@ -276,7 +277,7 @@ test("output tabs support keyboard navigation and precision chips use the backen
  await page.getByLabel("Search recipes").fill("does not exist");
  await expect(page.getByRole("heading", { name: "No matching recipes" })).toBeVisible();
  await page.getByRole("button", { name: "Clear filters" }).click();
-  await expect(page.locator(".recipe-card")).toHaveCount(36);
+  await expect(page.locator(".recipe-card")).toHaveCount(37);
 });

 test("cookbook cards and controls remain readable in dark mode and reflow at 320px", async ({ page }) => {
@@ -461,16 +462,16 @@ test("custom devices persist and no preview control can block generated YAML", a
  await page.getByRole("button", { name: "Can't find it? Add a device", exact: true }).click();
  await page.getByLabel("Device name", { exact: true }).fill("RTX 4090");
  await page.getByRole("combobox", { name: "Memory type", exact: true }).selectOption("dedicated");
-  await page.getByRole("spinbutton", { name: "Total memory", exact: true }).fill("24");
-  await page.getByRole("spinbutton", { name: "Usable memory (optional)", exact: true }).fill("32");
-  await expect(page.getByText("Usable memory cannot exceed total memory.", { exact: true })).toBeVisible();
-  await expect(page.getByRole("button", { name: "Save device", exact: true })).toBeDisabled();
-  await page.getByRole("spinbutton", { name: "Usable memory (optional)", exact: true }).fill("22");
+  await page.getByRole("spinbutton", { name: "Nominal memory", exact: true }).fill("24");
+  await page.getByRole("spinbutton", { name: "Usable budget (optional)", exact: true }).fill("32");
+  await expect(page.getByText("The usable budget cannot exceed the nominal memory.", { exact: true })).toBeVisible();
+  await expect(page.getByRole("button", { name: "Use this device", exact: true })).toBeDisabled();
+  await page.getByRole("spinbutton", { name: "Usable budget (optional)", exact: true }).fill("22");
  await page.getByRole("spinbutton", { name: "Host RAM (optional)", exact: true }).fill("64");
-  await page.getByRole("button", { name: "Save device", exact: true }).click();
+  await page.getByRole("button", { name: "Use this device", exact: true }).click();
  await expect(page.getByText("RTX 4090", { exact: true })).toBeVisible();
  await expect(page.locator(".hardware-selected")).toContainText("24 GB VRAM");
-  await expect(page.locator(".hardware-selected")).toContainText("Usable memory: 22 GB");
+  await expect(page.locator(".hardware-selected")).toContainText("Your budget: 22 GB");
  await expect(page.locator(".hardware-selected")).toContainText("Host RAM: 64 GB");
  await expect(page).toHaveURL(/hardware=custom/);
  await expect(page).toHaveURL(/hwName=RTX\+4090/);
@@ -479,26 +480,28 @@ test("custom devices persist and no preview control can block generated YAML", a
  await expect(page.locator(".hardware-selected")).toContainText("24 GB VRAM");
  await expect(page.locator(".component-plan")).toHaveCount(0);
  await expect(page.getByRole("button", { name: "Download YAML", exact: true })).toBeEnabled();
-  await page.getByRole("button", { name: "Clear hardware", exact: true }).click();
+  await page.getByRole("button", { name: "Clear hardware target", exact: true }).click();
  await expect(page.getByText("RTX 4090", { exact: true })).toHaveCount(0);
-  await expect(page.getByText("No hardware selected.", { exact: true })).toBeVisible();
+  await expect(page.getByText("No hardware target selected.", { exact: true })).toBeVisible();
 });

-test("hardware is never preselected and a desktop card keeps the recipe GPU count", async ({ page }) => {
+test("hardware is never preselected; a single-GPU device sets one GPU with valid parallelism", async ({ page }) => {
  await page.goto(route);
  await expect(page.locator(".hardware-picker")).toBeVisible({ timeout: 30000 });
-  await expect(page.getByText("No hardware selected.", { exact: true })).toBeVisible();
+  await expect(page.getByText("No hardware target selected.", { exact: true })).toBeVisible();
  await output(page, "config.yaml");
  const yaml = page.getByLabel("Generated YAML", { exact: true });
  await expect(yaml).toContainText('"num_gpus": 4');
  await chooseHardware(page, "RTX 5090");
  await expect(page).toHaveURL(/hardware=geforce-rtx-5090-32gb/);
  await expect(page.getByText(/GeForce RTX 5090 .*32 GB VRAM/)).toBeVisible();
-  await expect(yaml).toContainText('"num_gpus": 4');
+  await expect(page.getByText(/Set to 1 GPU for .*RTX 5090/)).toBeVisible();
+  await expect(yaml).toContainText('"num_gpus": 1');
+  await expect(yaml).toContainText('"sp_size": -1');
  await expect(yaml).not.toContainText('"hardware"');
  await expect(page.getByText("Valid configuration", { exact: true })).toBeVisible();
-  await page.getByRole("button", { name: "Clear hardware", exact: true }).click();
-  await expect(page.getByText("No hardware selected.", { exact: true })).toBeVisible();
+  await page.getByRole("button", { name: "Clear hardware target", exact: true }).click();
+  await expect(page.getByText("No hardware target selected.", { exact: true })).toBeVisible();
 });

 test("lowering the GPU count resets parallelism that no longer fits", async ({ page }) => {
@@ -519,7 +522,7 @@ test("browser back after a hardware-only change keeps the editor loaded", async
  await expect(page).toHaveURL(/hardware=geforce-rtx-5090-32gb/);
  await page.goBack();
  await expect(page).not.toHaveURL(/hardware=/);
-  await expect(page.getByText("No hardware selected.", { exact: true })).toBeVisible();
+  await expect(page.getByText("No hardware target selected.", { exact: true })).toBeVisible();
  await expect(page.getByRole("button", { name: "Download YAML", exact: true })).toBeEnabled();
  await expect(page.getByText("Loading deployment settings…")).toHaveCount(0);
 });
@@ -532,135 +535,15 @@ test("hardware links resolve catalog aliases and non-CUDA devices are not CUDA t
  await expect(page.getByRole("radio", { name: /Arc B580/ })).toHaveCount(0);
 });

-test("recipe facts state validity and why memory has no estimate", async ({ page }) => {
+test("recipe facts separate validity, fit, and run evidence with inspectable assumptions", async ({ page }) => {
  await page.goto(route);
-  const facts = page.locator(".recipe-facts");
+  const facts = page.getByLabel("Recipe facts");
  await expect(facts.getByText("Valid configuration", { exact: true })).toBeVisible();
-  await expect(facts.getByText("Memory not estimated", { exact: true })).toBeVisible();
+  await expect(facts.getByText("Fit unknown", { exact: true })).toBeVisible();
+  await expect(facts.getByText("No recorded run", { exact: true })).toBeVisible();
  await page.getByRole("button", { name: "Details", exact: true }).click();
-  await expect(page.getByText(/Select a device to estimate memory/)).toBeVisible();
-});
-
-test("host-only system RAM cannot produce a GPU weights fit verdict", async ({ page }) => {
-  await page.goto(`${route}?hardware=custom&hwName=Host&hwType=system&hwMemory=128&hwUnit=GB&hwBackend=cuda`);
-  const facts = page.locator(".recipe-facts");
-  await expect(facts.getByText("Memory not estimated", { exact: true })).toBeVisible();
-  await page.getByRole("button", { name: "Details", exact: true }).click();
-  await expect(facts.getByText(/Host-only system RAM is not GPU capacity/)).toBeVisible();
-  await expect(facts.getByText(/Resident weights need/)).toHaveCount(0);
-  await expect(page.getByRole("button", { name: "Download YAML", exact: false })).toBeEnabled();
-});
-
-test("TAEH3 stops using the bundled VAE weights in memory facts", async ({ page }) => {
-  await page.goto(`${route}?hardware=custom&hwName=GPU&hwType=dedicated&hwMemory=80&hwUnit=GB&hwBackend=cuda`);
-  const facts = page.locator(".recipe-facts");
-  await expect(facts.getByText("Weights exceed memory", { exact: true })).toBeVisible();
-  await chip(page, "Video decoder", "taeh3");
-  await expect(facts.getByText("Memory not estimated", { exact: true })).toBeVisible();
-  await page.getByRole("button", { name: "Details", exact: true }).click();
-  await expect(facts.getByText(/TAEH3 decoder weights are not modeled/)).toBeVisible();
-  await expect(facts.getByText(/Resident weights need/)).toHaveCount(0);
-  await output(page, "config.yaml");
-  await expect(page.getByLabel("Generated YAML")).toContainText('"video_decode_backend": "taeh3"');
-  await chip(page, "Video decoder", "h3-vae");
-  await expect(facts.getByText("Weights exceed memory", { exact: true })).toBeVisible();
-});
-
-test("FastH3 keeps unmodeled quantization and decoder residency unknown", async ({ page }) => {
-  await page.goto(route);
-  await chooseHardware(page, "RTX 5090");
-  const facts = page.locator(".recipe-facts");
-  await expect(facts.getByText("Weights exceed memory", { exact: true })).toBeVisible();
-  await page.getByRole("button", { name: "Details", exact: true }).click();
-  await expect(page.getByText(/Resident weights need 75\.6 GiB of 29\.8 GiB/)).toBeVisible();
-  await page
-    .getByRole("group", { name: "Transformer quantization", exact: true })
-    .getByRole("button", { name: "NVFP4", exact: true })
-    .click();
-  await expect(facts.getByText("Memory not estimated", { exact: true })).toBeVisible();
-  await expect(facts.getByText(/checkpoint and layer profile/)).toBeVisible();
-  await expect(facts.getByText(/Resident weights need/)).toHaveCount(0);
-  await page
-    .getByRole("group", { name: "NVFP4 video decoder", exact: true })
-    .getByRole("button", { name: /8L · fast/ })
-    .click();
-  await expect(facts.getByText("Memory not estimated", { exact: true })).toBeVisible();
-  await expect(facts.getByText(/VAE retains its encoder and packs decoder weights/)).toBeVisible();
-  await expect(facts.getByText(/Resident weights need/)).toHaveCount(0);
-  await output(page, "config.yaml");
-  await expect(page.getByLabel("Generated YAML")).toContainText('"transformer_quant": "NVFP4"');
-  await expect(page.getByLabel("Generated YAML")).toContainText("fast-8L/decoder-nvfp4.safetensors");
-});
-
-test("4-Step quantization and full-depth decoder swaps also stay unestimated", async ({ page }) => {
-  await page.goto("/FastVideo/cookbook/minimax-h3/fasth3-4step");
-  await chooseHardware(page, "RTX 5090");
-  const facts = page.locator(".recipe-facts");
-  await page.getByRole("button", { name: "Details", exact: true }).click();
-  await chip(page, "Transformer quantization", "FP8");
-  await expect(facts.getByText(/checkpoint and layer profile/)).toBeVisible();
-  await chip(page, "Transformer quantization", "Default");
-  await page
-    .getByRole("group", { name: "NVFP4 video decoder", exact: true })
-    .getByRole("button", { name: /36L/ })
-    .click();
-  await expect(facts.getByText(/VAE retains its encoder and packs decoder weights/)).toBeVisible();
-  await expect(facts.getByText(/Resident weights need/)).toHaveCount(0);
-});
-
-test("the retired Preview route and deployment links keep their checkpoint", async ({ page }) => {
-  await page.goto("/FastVideo/cookbook/minimax-h3/fasth3-preview");
-  await output(page, "config.yaml");
-  await expect(page.getByLabel("Generated YAML")).toContainText("FastVideo-Minimax-FastH3-Preview-v0.2");
-  await page.goto("/FastVideo/cookbook?deployment=fasth3-preview%2Fcuda-rest");
-  await expect(page).toHaveURL(/minimax-h3\/fasth3-preview/);
-  await output(page, "config.yaml");
-  await expect(page.getByLabel("Generated YAML")).toContainText("FastVideo-Minimax-FastH3-Preview-v0.2");
-});
-
-for (const [deployment, label, sequential, tileBatch] of [
-  ["rtx5090-rest", "RTX 5090", true, "4"],
-  ["rtx-pro6000-rest", "RTX PRO 6000", false, "8"],
-  ["spark-rest", "DGX Spark", false, "12"],
-  ["cuda-rest", "RTX 5090", true, "4"],
-] as const) {
-  test(`Omni deployment link ${deployment} preserves its hardware baseline`, async ({ page }) => {
-    await page.goto(`/FastVideo/cookbook?deployment=fasth3-omniref%2F${deployment}`);
-    await expect(page).toHaveURL(/minimax-h3\/fasth3-omniref/);
-    await expect(page.getByRole("button", { name: label, exact: true })).toHaveAttribute("aria-pressed", "true");
-    await page.getByRole("button", { name: "Serving", exact: true }).click();
-    await output(page, "config.yaml");
-    const yaml = page.getByLabel("Generated YAML");
-    await expect(yaml).toContainText(`"h3_sequential_load": ${sequential}`);
-    await expect(yaml).toContainText('"layer_profile": "h3_dit_vsa"');
-    await output(page, "Commands");
-    await expect(page.getByLabel("Server command")).toContainText(`FASTVIDEO_H3_VAE_TILE_BATCH=${tileBatch}`);
-    if (deployment === "spark-rest")
-      await expect(page.getByLabel("Server command")).toContainText("FASTVIDEO_NVFP4_MM_BACKEND=cudnn");
-  });
-}
-
-test("packed FastH3 checkpoints keep their fixed precision", async ({ page }) => {
-  await page.goto("/FastVideo/cookbook/minimax-h3/fasth3-trim");
-  await page.getByRole("tab", { name: "config.yaml" }).click();
-  await expect(page.getByLabel("Generated YAML")).toContainText('"transformer_quant": "NVFP4"');
-  await expect(page.getByLabel("Generated YAML")).toContainText('"layer_profile": "h3_dit_vsa"');
-  await expect(page.getByText("Valid configuration", { exact: true })).toBeVisible();
-  await expect(page.getByRole("group", { name: "Transformer quantization", exact: true })).toHaveCount(0);
-});
-
-test("MLX FastH3 stays unestimated and a GiB device keeps its unit", async ({ page }) => {
-  await page.goto("/FastVideo/cookbook/minimax-h3/fasth3-8step?deployment=fasth3-8step%2Fmlx-rest");
-  await page.getByRole("button", { name: "+ Find or add a Mac", exact: true }).click();
-  await page.getByLabel("Search hardware", { exact: true }).fill("M4 Max");
-  await page.getByRole("radio", { name: /Apple M4 Max 128 GB/ }).click();
-  await expect(page.getByText("Memory not estimated", { exact: true })).toBeVisible();
-  await page.getByRole("button", { name: "Details", exact: true }).click();
-  await expect(page.getByText(/MLX checkpoints use a different precision/)).toBeVisible();
-  await page.goto(
-    "/FastVideo/cookbook/minimax-h3/fasth3-8step?hardware=custom&hwName=Big&hwType=dedicated&hwMemory=80&hwUnit=GiB&hwBackend=cuda",
-  );
-  await expect(page.getByText("Weights fit, little headroom", { exact: true })).toBeVisible();
+  await expect(page.getByText(/This is not a fit claim/)).toBeVisible();
+  await expect(page.getByText(/do not establish this configuration/)).toBeVisible();
 });

 test("catalog lookup resolves aliases, variants, and unified memory without pooling", async ({ page }) => {
@@ -669,7 +552,7 @@ test("catalog lookup resolves aliases, variants, and unified memory without pool
  await expect(page).toHaveURL(/hardware=rtx-pro-6000-blackwell-workstation-96gb/);
  await expect(page.locator(".hardware-selected")).toContainText("96 GB VRAM");
  await expect(page.locator(".hardware-address-note")).toContainText("Per device, not a combined pool");
-  await page.getByRole("button", { name: "Clear hardware", exact: true }).click();
+  await page.getByRole("button", { name: "Clear hardware target", exact: true }).click();
  await page.getByRole("button", { name: "+ Find or add a GPU", exact: true }).click();
  await page.getByLabel("Search hardware", { exact: true }).fill("GB10");
  await page.getByRole("radio", { name: /DGX Spark/ }).click();
@@ -689,9 +572,9 @@ test("MLX lookup uses unified memory and a custom Mac stays reusable", async ({
  await page.getByRole("radio", { name: /Apple M4 Max/ }).click();
  await expect(page.locator(".hardware-selected")).toContainText("128 GB unified memory");
  await expect(page.locator(".hardware-selected")).toContainText("Apple MLX");
-  await expect(page.locator(".hardware-selected")).toContainText("Usable memory: 96 GB");
+  await expect(page.locator(".hardware-selected")).toContainText("Usable budget: 96 GB");
  await expect(page.locator(".hardware-address-note")).toContainText("not dedicated VRAM");
-  await expect(page.getByRole("button", { name: "Clear hardware", exact: true })).toBeVisible();
+  await expect(page.getByRole("button", { name: "Clear hardware target", exact: true })).toBeVisible();
 });

 test("the directory filters by runtime without a hardware picker", async ({ page }) => {
@@ -705,33 +588,6 @@ test("the directory filters by runtime without a hardware picker", async ({ page
  await expect(page.locator(".recipe-card").filter({ hasText: "FastWan2.1 1.3B" })).toHaveCount(0);
 });

-test("MLX Omni discovery and its original route expose the MLX launch command", async ({ page }) => {
-  await page.goto("/FastVideo/cookbook");
-  await page.getByLabel("Runtime", { exact: true }).selectOption("mlx");
-  await page.getByLabel("Search recipes").fill("FastH3 Omni");
-  const cards = page.locator(".recipe-card");
-  await expect(cards).toHaveCount(1);
-  await expect(cards).toHaveAttribute("href", "/FastVideo/cookbook/minimax-h3/models/fasth3-omniref-mlx");
-  await cards.click();
-  await expect(page.getByLabel("Inference command")).toContainText("convert_minimax_h3_mlx.py");
-  await expect(page.getByLabel("Inference command")).toContainText("--mlx-checkpoint ./FastH3-OmniRef-MLX/int6");
-  await expect(page.getByLabel("Inference command")).not.toContainText("fastvideo serve");
-  await page.reload();
-  await expect(page.getByLabel("Inference command")).toContainText("--mlx-checkpoint ./FastH3-OmniRef-MLX/int6");
-});
-
-test("category counts match visible cards and exclude retired Preview", async ({ page }) => {
-  await page.goto("/FastVideo/cookbook");
-  await expect(page.locator(".recipe-card")).toHaveCount(36);
-  for (const category of ["All models", "Video", "Image", "Audio", "World"]) {
-    const button = page.getByRole("button", { name: category, exact: true });
-    const count = Number(await button.locator(".recipe-filter-count").innerText());
-    await button.click();
-    await expect(page.locator(".recipe-card")).toHaveCount(count);
-    await expect(page.locator(".recipe-card").filter({ hasText: "FastH3 Preview V0.2" })).toHaveCount(0);
-  }
-});
-
 test("task discovery filters the directory and round-trips through the URL", async ({ page }) => {
  await page.goto("/FastVideo/cookbook");
  await page.getByLabel("Task", { exact: true }).selectOption("Image to video");
@@ -757,10 +613,10 @@ test("directory filters round-trip through the URL", async ({ page }) => {

 test("a saved target missing from the catalog stays visible and clearable", async ({ page }) => {
  await page.goto(`${route}?hardware=retired-gpu`);
-  await expect(page.getByText(/The saved device \(retired-gpu\) is not in this catalog/)).toBeVisible();
-  await page.getByRole("button", { name: "Clear hardware", exact: true }).click();
+  await expect(page.getByText(/Saved hardware target \(retired-gpu\) is not in this catalog/)).toBeVisible();
+  await page.getByRole("button", { name: "Clear hardware target", exact: true }).click();
  await expect(page).not.toHaveURL(/hardware=/);
-  await expect(page.getByText("No hardware selected.", { exact: true })).toBeVisible();
+  await expect(page.getByText("No hardware target selected.", { exact: true })).toBeVisible();
 });

 test("catalog search renders one plain no-results message", async ({ page }) => {
@@ -1,19 +1,5 @@
 import { expect, test } from "@playwright/test";

-test("Omni Apple silicon links open the published guide and MLX recipe", async ({ page }) => {
-  for (const slash of ["", "/"]) {
-    await page.goto(`/FastVideo/cookbook/minimax-h3/fasth3-omniref${slash}`);
-    const guide = page.getByRole("link", { name: "FastH3 Omni on MLX", exact: true });
-    await expect(guide).toHaveAttribute("href", "/FastVideo/inference/fasth3_omniref_mlx");
-    await guide.click();
-    await expect(page).toHaveURL(/\/FastVideo\/inference\/fasth3_omniref_mlx\/?$/);
-    await expect(page.getByRole("heading", { level: 1 })).toBeVisible();
-  }
-  await page.goto("/FastVideo/cookbook/minimax-h3/fasth3-omniref");
-  await page.getByRole("link", { name: "MLX recipe", exact: true }).click();
-  await expect(page.getByLabel("Inference command")).toContainText("--mlx-checkpoint");
-});
-
 test("GitHub stars refresh automatically and survive a failed refresh", async ({ page }) => {
  await page.clock.install();
  let requests = 0;
@@ -173,12 +159,6 @@ test("the model catalog filters by family and checkpoint name", async ({ page })
  await expect(catalog.getByRole("status")).toHaveText("3 checkpoints");
  await page.getByRole("searchbox", { name: "Search checkpoints", exact: true }).fill("14B");
  await expect(catalog.getByRole("status")).toHaveText("1 checkpoint");
-  await page.getByRole("button", { name: "All models", exact: true }).click();
-  await page.getByRole("searchbox").fill("NVFP4");
-  const codes = await catalog.locator(".model-checkpoints code").allInnerTexts();
-  expect(codes.length).toBeGreaterThan(0);
-  expect(codes.every((text) => text.includes("NVFP4"))).toBe(true);
-  await expect(catalog.getByText("FastVideo/FastVideo-FastH3-8-Step-V2-FP8", { exact: true })).toHaveCount(0);
  await page.getByRole("searchbox").fill("no-such-model");
  await expect(catalog.getByRole("status")).toHaveText("0 checkpoints");
  await page.getByRole("button", { name: "Clear filters", exact: true }).click();
@@ -198,7 +178,7 @@ test("family runbooks use React model cards without the old cookbook runtime", a
  await page.goto("/FastVideo/cookbook/wan/");
  await expect(page.locator(".recipe-card").first()).toBeVisible();
  await expect(page.getByRole("button", { name: "Search", exact: true })).toBeVisible();
-  await expect(page.getByRole("link", { name: "← All recipes", exact: true })).toHaveAttribute(
+  await expect(page.getByRole("link", { name: "← All models", exact: true })).toHaveAttribute(
    "href",
    /\/FastVideo\/cookbook$/,
  );
@@ -1,64 +0,0 @@
-import assert from "node:assert/strict";
-import { test } from "node:test";
-import { readReadmeSections } from "../lib/readme.mjs";
-
-const legacy = `# FastVideo
-
-## NEWS
- \`2026/10/06\`: FastH3 V2 runs on a single consumer machine.
- \`2026/09/15\`: Release FastH3 8-Step V2.
-
-## Awesome work using FastVideo or our research projects
- [Some project](https://example.com)
-
-## Community
- Slack
-`;
-
-const rewritten = `# FastVideo
-
-## Overview
-FastVideo is a framework.
-
-## Recent updates
- **2026-10-06:** FastH3 V2 runs on a single consumer machine.
- **2026-09-15:** Release FastH3 8-Step V2.
-
-## Projects using FastVideo
- [Some project](https://example.com)
-
-## Community and contributing
- Slack
-`;
-
-test("reads the legacy README headings and dates", () => {
-  const { news, projects } = readReadmeSections(legacy);
-  assert.equal(news.length, 2);
-  assert.deepEqual(
-    news.map((item) => item.date),
-    ["2026/10/06", "2026/09/15"],
-  );
-  assert.match(news[0].text, /^FastH3 V2/);
-  assert.match(projects, /Some project/);
-});
-
-test("reads a rewritten README with sentence case headings and bold dates", () => {
-  const { news, projects } = readReadmeSections(rewritten);
-  assert.equal(news?.length, 2);
-  assert.deepEqual(
-    news.map((item) => item.date),
-    ["2026/10/06", "2026/09/15"],
-  );
-  assert.doesNotMatch(news[0].text, /\*\*/);
-  assert.match(projects, /Some project/);
-  assert.doesNotMatch(projects, /Slack/);
-});
-
-test("stops at the next heading and ignores other sections", () => {
-  const { news } = readReadmeSections(legacy);
-  assert.doesNotMatch(news[1].text, /Awesome work/);
-});
-
-test("returns null when a section is missing", () => {
-  assert.deepEqual(readReadmeSections("# FastVideo\n\n## Overview\n"), { news: null, projects: null });
-});
@@ -798,39 +798,6 @@
 .recipe-card-grid-configurable .recipe-card {
  border-block-start: 2px solid var(--blume-accent);
 }
-/* Family runbook cards are static: no curtain, so they need their own rhythm. */
-.recipe-card-static {
-  display: flex;
-  flex-direction: column;
-  gap: 0.5rem;
-  min-block-size: 13.5rem;
-  padding: 1.25rem 1.25rem 1rem;
-}
-.recipe-card-static .recipe-card-top {
-  margin-block-end: 0;
-  padding-block-end: 0;
-}
-.recipe-card-static h2 {
-  margin: 0;
-  font-size: 1rem;
-  line-height: 1.35;
-  font-weight: 600;
-  letter-spacing: -0.01em;
-}
-.recipe-card-static footer {
-  display: flex;
-  align-items: center;
-  justify-content: space-between;
-  gap: 0.75rem;
-  margin-block-start: auto;
-  margin-inline: 0;
-  padding-block: 0.75rem 0;
-  border-block-start: 1px solid var(--recipe-line);
-  font-size: 0.6875rem;
-}
-.recipe-card-static footer span:first-child {
-  color: var(--blume-muted-foreground);
-}
 .recipe-empty {
  padding-block: 2rem;
  text-align: center;
@@ -1655,25 +1622,3 @@
 .recipe-request-code summary {
  padding-block: 0.25rem;
 }
-
-/* Article rhythm.
-   Blume's prose gives <details> no vertical margin, so consecutive cookbook
-   sections touch. Tables inherit the body line height, which makes dense cells
-   taller than the prose around them. Both are set here, once, instead of in
-   each page's markup. */
-.prose details {
-  margin-block: 1rem;
-}
-.prose details + details {
-  margin-top: 1rem;
-}
-.prose summary {
-  cursor: pointer;
-}
-.prose table {
-  line-height: 1.5;
-}
-.prose td,
-.prose th {
-  vertical-align: top;
-}
@@ -1,92 +0,0 @@
-{
-  "generated": "2026-10-28",
-  "note": "Measured from Hugging Face safetensors files. Sizes are GiB at the checkpoint's own precision.",
-  "models": {
-    "fasth3-8step": {
-      "parts": {
-        "transformer": 65.29,
-        "text_encoder": 62.13,
-        "vae": 9.7,
-        "audio_vae": 0.56
-      },
-      "alternates": {
-        "transformer_nvfp4": 36.91,
-        "transformer_fp8": 47.35,
-        "vae_nvfp4_quality": 5.82,
-        "vae_nvfp4_fast": 1.33,
-        "text_encoder_quantized": 15.33
-      },
-      "sources": [
-        "https://huggingface.co/FastVideo/FastVideo-FastH3-8-Step-V2/tree/main"
-      ]
-    },
-    "fasth3-4step": {
-      "parts": {
-        "transformer": 65.29,
-        "text_encoder": 62.13,
-        "vae": 9.7,
-        "audio_vae": 0.56
-      },
-      "alternates": {
-        "transformer_nvfp4": 36.91,
-        "transformer_fp8": 47.35,
-        "vae_nvfp4_quality": 5.82,
-        "vae_nvfp4_fast": 1.33,
-        "text_encoder_quantized": 15.33
-      },
-      "sources": [
-        "https://huggingface.co/FastVideo/FastVideo-FastH3-4-step-Preview-v1-VSA-DataFree/tree/main"
-      ]
-    },
-    "fasth3-trim": {
-      "parts": {
-        "transformer": 34.94,
-        "text_encoder": 15.33,
-        "vae": 5.93,
-        "audio_vae": 0.56
-      },
-      "alternates": {
-        "transformer_nvfp4": 11.1,
-        "transformer_fp8": 19.88,
-        "vae_nvfp4_quality": 5.82,
-        "vae_nvfp4_fast": 1.33
-      },
-      "sources": [
-        "https://huggingface.co/FastVideo/FastVideo-FastH3-Trim-8-Step/tree/main"
-      ]
-    },
-    "fasth3-omniref": {
-      "parts": {
-        "transformer": 65.33,
-        "text_encoder": 15.33,
-        "vae": 9.7,
-        "audio_vae": 0.56
-      },
-      "alternates": {
-        "transformer_nvfp4": 36.95,
-        "transformer_fp8": 47.39,
-        "vae_nvfp4_quality": 5.82,
-        "vae_nvfp4_fast": 1.33
-      },
-      "sources": [
-        "https://huggingface.co/FastVideo/FastVideo-FastH3-Omni-8-Step-V1/tree/main",
-        "https://huggingface.co/FastVideo/FastH3-Omni-NVFP4-QAD/tree/main"
-      ]
-    },
-    "pika-h3-flex-t2va": {
-      "parts": {
-        "transformer": 61.73,
-        "adapter": 2.58,
-        "text_encoder": 62.13,
-        "vae": 9.7,
-        "audio_vae": 0.56
-      },
-      "alternates": {
-        "vae_nvfp4_quality": 5.82,
-        "vae_nvfp4_fast": 1.33,
-        "text_encoder_quantized": 15.33
-      },
-      "sources": ["https://huggingface.co/Pika-Labs/Pika-H3-Flex/tree/main"]
-    }
-  }
-}
@@ -157,22 +157,15 @@ __all__ = [

 ## 4. Register the Backend

-Registration is two edits:
+Update `fastvideo/attention/backends/__init__.py` to export your new class.

-1. Add a member to `AttentionBackendEnum` in `fastvideo/platforms/interface.py`.
-2. Add the selection branch in `fastvideo/platforms/cuda.py`, next to the
-   existing backends, so `--attention-backend` maps the enum to your class.
+```python
+from .mynew_attn import MyNewAttnBackend
+```

-`fastvideo/attention/backends/__init__.py` is empty and does not register
-anything.
+## 5. Platform Integration

-## 5. Platform integration
-
-If your backend needs platform checks, such as an H100 guard, put them in
-`fastvideo/platforms/cuda.py` or in the backend constructor.
-
-The maintained version of this guide is
-[Add an attention backend](../../contributing/attention_backend.md).
+If your backend requires specific platform checks (e.g., checking for H100 support), handle that in `fastvideo/platforms/cuda.py` or within your backend's `__init__`.

 ## 6. Add Documentation

@@ -85,10 +85,10 @@ from fastvideo_kernel import video_sparse_attn
 # topk: Number of blocks to attend

 output = video_sparse_attn(
-    q, k, v,
-    variable_block_sizes,
-    q_variable_block_sizes,
-    topk=32,
+    q, k, v, 
+    block_sizes,
+    block_sizes,
+    topk=32
 )
 ```

@@ -49,7 +49,7 @@ node docs/js/build-cookbook-capacity.mjs --check
 node docs/js/build-cookbook-capacity.mjs --preview fastwan21
 ```

-The check currently prints `Validated 1 capacity model(s) and 52 hardware variant(s).`
+The check currently prints `Validated 1 capacity model(s) and 50 hardware variant(s).`
 The preview prints the model with component references expanded and each
 example's baseline under `base_config`. Inspect its identity, counts, geometry,
 and evidence. Both commands validate deployments and capacity metadata without
@@ -404,41 +404,6 @@ matching metric, scope, phases, and execution context. One point can calibrate a
 existing scaling assumption, not establish accuracy across modes. Missing startup
 evidence leaves startup coverage unknown. Keep observations intact when deriving coefficients.

-## FastH3 coverage
-
-No FastH3 checkpoint has a full capacity model yet. A weights-only preflight
-is live on the FastH3 pages: it compares measured resident checkpoint weights
-with the selected device and labels the result as weights, not a request peak.
-Host-only system RAM is not accepted as GPU capacity. Transformer quantization,
-TAEH3, and NVFP4 decoder swaps return unknown until their
-effective layer profiles, retained components, and packed storage are modeled.
-A full estimate still stops early: `checkCoverage()` in `docs/js/cookbook-memory.mjs` still stops before the
-estimate for every FastH3 serving profile.
-
-| Profile | First stop | Why |
-| --- | --- | --- |
-| OmniRef | Workload | Runs image-to-video; the estimator models text-to-video only |
-| 8-Step CUDA | Topology | Defaults to four GPUs and Ulysses sequence parallelism |
-| Trim Spark | Attention | `VIDEO_SPARSE_ATTN_H3` is not a dense backend, so VSA workspace is unmodeled |
-| 4-Step Spark | Execution | `h3_sequential_load` loads components on demand |
-| Omni NVFP4 decoder | Precision | The decoder swap changes decoder precision |
-
-To make a FastH3 estimate possible, in order:
-
-1. Add capacity entries for each checkpoint: transformer parameter counts and
-   geometry, the text encoder, the VAE, and the NVFP4 decoder variants.
-2. Model an image-to-video request: reference latents and audio tokens change
-   the working set, and OmniRef passes ordered image, video, and audio inputs.
-3. Model per-rank sharding for four-GPU Ulysses, including replicated weights
-   and per-rank sequence slices, so the estimate compares against one device.
-4. Add a VSA-H3 workspace term, or a dense upper bound that is labeled as one.
-5. Support quantized decoder precision instead of rejecting the configuration.
-6. Record successful runs on the RTX 5090, RTX PRO 6000, and DGX Spark profiles
-   so the analytical estimate can be calibrated rather than trusted blindly.
-
-Until those land, the cookbook reports a weights fact for FastH3 instead of a
-peak estimate. That is the correct result, not a bug.
-
 ## Static consumers

 [`cookbook-capacity.mjs`](../js/cookbook-capacity.mjs) exposes these async loaders,
@@ -208,7 +208,7 @@ longer exists also fails the test, so the fixing pull request deletes its entry.
 | `FASTVIDEO_VAE_PARALLEL_DECODE_STRATEGY`                   | str                                                                | unset                                                 | performance | Collective that moves chunks in parallel VAE decode: gather (used when unset) or all_gather.                                                                                                                                                                                                                                                                                                                                                |
 | `FASTVIDEO_H3_ENCODER_SPLIT`                               | bool                                                               | `0`                                                   | distributed | MiniMax-H3 runs the Qwen3-VL text encoder on a dedicated group of nodes, so the denoising ranks never load it. Same as FastVideoArgs.h3_encoder_split=True.                                                                                                                                                                                                                                                                                 |
 | `FASTVIDEO_H3_ENCODER_NODES`                               | int                                                                | `1`                                                   | distributed | Number of leading nodes that FASTVIDEO_H3_ENCODER_SPLIT dedicates to the MiniMax-H3 text encoder. Values below 1 count as 1.                                                                                                                                                                                                                                                                                                                |
-| `FASTVIDEO_MINIMAX_H3_FUSIONS`                             | str                                                                | `""`                                                  | performance | MiniMax-H3 inference-only Triton fusions: all, 1, or a comma-separated subset of modulate,qknorm_rope,swiglu. Empty, 0, or none keeps the eager implementation.                                                                                                                                                                                                                                                                             |
+| `FASTVIDEO_MINIMAX_H3_FUSIONS`                             | str                                                                | `""`                                                  | performance | MiniMax-H3 inference-only Triton fusions: all, 1, or a comma-separated subset of modulate,qknorm_rope,swiglu. The whole-tier value mk selects the MotionKernel kernels instead of the Sol-Engine ones. Empty, 0, or none keeps the eager implementation.                                                                                                                                                                                    |
 | `FASTVIDEO_MINIMAX_H3_EXACT_KERNELS`                       | str                                                                | `""`                                                  | performance | MiniMax-H3 inference-only Triton kernels that reproduce the eager output bit for bit: all, 1, or a comma-separated subset of rope,modulate,swiglu. An op whose FASTVIDEO_MINIMAX_H3_FUSIONS fusion is on keeps the fusion. Empty, 0, or none keeps the eager implementation.                                                                                                                                                                |
 | `FASTVIDEO_FSDP2_AUTOWRAP`                                 | bool                                                               | `0`                                                   | performance | FSDP2 shards modules by parameter count instead of the model's shard conditions. Not supported by self-forcing distillation.                                                                                                                                                                                                                                                                                                                |
 | `FASTVIDEO_FSDP2_MIN_PARAMS`                               | int                                                                | `10000000`                                            | performance | Minimum parameter count of a module that FASTVIDEO_FSDP2_AUTOWRAP shards.                                                                                                                                                                                                                                                                                                                                                                   |
@@ -59,7 +59,7 @@
          </tr>
          <tr>
            <td>DGX Spark</td>
-            <td>FastH3 V1 on one GB10. Two Sparks run one clip faster over QSFP RoCE (<code>sp_size=2</code>), but this page configures one Spark.</td>
+            <td>FastH3 V1 on one GB10, or two Sparks with Ray sequence parallel (<code>sp_size=2</code>) over QSFP RoCE. Select NVIDIA DGX Spark, then 1 Spark or 2 Sparks.</td>
            <td>Not wired</td>
          </tr>
          <tr>
@@ -100,7 +100,7 @@ cd FastVideo</code></pre>
        <li>The FastH3 CUDA performance profile was measured on four GB200 GPUs. Use its strict profile when exact operation order matters more than the measured performance configuration.</li>
        <li>The MLX source runtime supports T2VA, optional temporal <code>--fast</code>, optional spatial <code>--fast-spatial</code>, and opt-in VSA on <code>--include-vsa</code> checkpoints. FastH3 V2 MLX converts with <code>--include-vsa</code> and runs eight forwards. FL2VA, Ref2VA, and two-pass refinement are not wired.</li>
        <li>GPU count and VAE decode backend are configurable in the builder above for FastH3 CUDA recipes. Only the value shown by default has a recorded run; other supported values are unmeasured here.</li>
-        <li>DGX Spark is a runtime on FastH3 V1, not a separate family card. This page configures one Spark. For two Sparks, run <code>examples/inference/basic/basic_fasth3_spark_pair.yaml</code>; there is no serving baseline. The CUDA GPU-count knob does not apply to Spark.</li>
+        <li>DGX Spark is a runtime on FastH3 V1, not a separate family card. Select NVIDIA DGX Spark, then 1 Spark or 2 Sparks. The CUDA GPU-count knob does not apply to Spark.</li>
        <li>CompactH3 NVFP4 is one Blackwell GPU. RTX 5090 (32 GB) parks the encoder in pinned host RAM. RTX PRO 6000 Blackwell (96 GB) keeps encoder, DiT, and VAE resident. Keep <code>FASTVIDEO_ATTENTION_BACKEND=ATTN_QAT_INFER</code>, <code>FASTVIDEO_FA4=0</code>, <code>FASTVIDEO_VSA_PTX_BLACKWELL=0</code>, and <code>FLASHINFER_CUDA_ARCH_LIST=12.0a</code>. On PRO 6000 enable VAE compile and leave DiT <code>inference_torch_compile</code> off. CompactH3 is a dense prune; do not enable <code>VIDEO_SPARSE_ATTN_H3</code> until a VSA-trained student exists.</li>
        <li>GB10 has no FA4 / ptx_blackwell VSA kernel. Keep <code>FASTVIDEO_FA4=0</code> and <code>FASTVIDEO_VSA_PTX_BLACKWELL=0</code>. Legal <code>num_frames</code> values are <code>17n+5</code>, capped at 362 (15.08 s). A 345-frame request on one Spark can OOM. Native 16:9 sizes include 832×480 and 1344×768.</li>
        <li>Gated or missing checkpoints: run <code>huggingface-cli login</code> and confirm you accepted the model's license on Hugging Face.</li>
@@ -111,6 +111,6 @@ cd FastVideo</code></pre>
 <details  id="cookbook-evidence">
  <summary>Evidence status</summary>
  <div >
-      <p>Every command, model ID, and flag on this page maps to a checked-in FastVideo source. No recipe here carries a recorded run for a specific device, so the editor does not claim one. The full H3 CUDA examples remain <strong>Source-backed</strong> where the source records a GPU count but no GPU model or memory requirement. Unlisted hardware is unknown, not unsupported.</p>
+      <p>Every command, model ID, and flag on this page maps to a checked-in FastVideo source. Recipes marked <strong>Verified</strong> also have a recorded hardware path in linked FastVideo evidence. The full H3 CUDA examples remain <strong>Source-backed</strong> where the source records a GPU count but no GPU model or memory requirement. Unlisted hardware is unknown, not unsupported.</p>
  </div>
 </details>
@@ -1,27 +1,30 @@
 # Run H3 with a server and playground

-Start FastH3 on CUDA, one DGX Spark, or Apple Silicon MLX, then send requests from a
-browser, cURL, or your app. Several clients can share one server process.
+Start FastH3 on CUDA, one DGX Spark, or Apple Silicon MLX, then iterate on prompts in a browser,
+with cURL, or from your app. The server and clients can run on the same machine.

-CUDA requests reuse a loaded `VideoGenerator`. MLX keeps one pipeline object and
-a prompt-embedding cache, but loads and releases heavyweight components between
-phases to fit unified memory. MLX serving still pays that component-loading cost
-on subsequent requests.
+The server supports the OpenAI-compatible video-job API. The Python and
+JavaScript client examples use that interface, but requests go to FastVideo.
+You do not need an OpenAI account or cloud key.

-The server implements the OpenAI-compatible video-job API. The Python and
-JavaScript clients use that API, but requests go to FastVideo. You do not need an
-OpenAI account or cloud key.
+The [H3 recipe selector](minimax-h3.md) provides the same workflow with runtime
+selection. This guide covers FastH3 V1 (four forwards) and FastH3 V2
+(eight forwards). Start a server, then iterate in the playground or with the
+OpenAI Python client. Other H3 recipes keep their direct Python commands.

-This guide covers FastH3 V1 (four forwards) and FastH3 V2 (eight forwards). For
-direct Python generation, see the [H3 cookbook](minimax-h3.md).
+CUDA requests reuse one loaded `VideoGenerator`. The Python SDK can do the same
+when you reuse the generator across `generate()` calls. MLX keeps one
+`MiniMaxH3MLXPipeline` and a prompt-embedding cache, but still loads and releases
+components between phases to fit unified memory. MLX serving does not keep all
+weights resident or remove phase-loading time. Use a server when you want
+separate clients to share the process without restarting scripts.

 ## Install and start the server

-All runtimes need a FastVideo clone and an activated Python environment.
-
 ### CUDA

-Complete the [CUDA installation requirements](../getting_started/installation/gpu.md)
+Use a FastVideo clone and an activated Python environment. Complete the
+[CUDA installation requirements](../getting_started/installation/gpu.md)
 before running these commands:

 ```bash
@@ -30,9 +33,9 @@ fastvideo serve --config examples/serving/openai_fasth3.yaml --server.host 127.0
 ```

 The configuration loads `FastVideo/FastVideo-Minimax-FastH3-Preview-v0.2` and
-advertises it as `fasth3`. It configures four CUDA GPUs and does not record a GPU
-model or VRAM requirement, so treat it as a starting point, not a measured
-profile. Compilation is disabled.
+advertises it as `fasth3`. It configures four CUDA GPUs but does not record
+a GPU model or VRAM requirement. This is a source-backed server profile, not
+the measured GB200 Python performance profile. Compilation is disabled.

 For FastH3 V2, use the same install and the 8-step config (nine sigma
 points, eight DiT forwards, VSA sparsity 0.8):
@@ -97,14 +97,6 @@ runtimes:
        title: Compile transformer
        type: boolean
        default: false
-      generator.engine.quantization.transformer_quant:
-        title: Transformer quantization
-        description: >-
-          Quantizes the DiT at load time. NVFP4 needs a Blackwell GPU. FP8 runs on Hopper and Ada.
-          Leave empty to use the checkpoint's own precision.
-        type: [string, "null"]
-        enum: [null, NVFP4, FP8]
-        default: null
      # MiniMax-H3 only (fastvideo_args.video_decode_backend). TAEH3 fetches its own SHA-pinned weights.
      generator.pipeline.experimental.video_decode_backend:
        title: Video decoder
@@ -24,7 +24,6 @@ deployments:
      - generator.engine.parallelism.tp_size
      - generator.engine.parallelism.sp_size
      - generator.engine.offload.text_encoder
-      - generator.engine.quantization.transformer_quant
      - generator.engine.offload.vae
      - generator.pipeline.experimental.video_decode_backend
      - generator.pipeline.experimental.h3_vae_nvfp4_decoder
@@ -21,7 +21,6 @@ deployments:
      - generator.engine.parallelism.*
      - "!generator.engine.parallelism.hsdp_*"
      - generator.engine.offload.text_encoder
-      - generator.engine.quantization.transformer_quant
      - generator.engine.offload.vae
      - generator.pipeline.experimental.video_decode_backend
      - generator.pipeline.experimental.h3_vae_nvfp4_decoder
@@ -1,4 +1,4 @@
-title: FastH3 Omni v1
+title: FastH3 OmniRef
 summary: Video and synchronized audio from a prompt plus ordered image, video, and audio references, in eight forwards.
 guide: docs/inference/fasth3_omniref.md
 deployments:
@@ -39,12 +39,10 @@ deployments:
        directories in server.media_roots.
      - Based on the maintained inference recipe; this serving adaptation has not been GPU-validated.
      - "Optional NVFP4 decoder (Blackwell only): hf download FastVideo/FastH3-NVFP4-VAE-Decoder --include 'quality-36L/decoder-nvfp4.safetensors' 'fast-8L/decoder-nvfp4.safetensors' --local-dir fasth3-nvfp4-vae"
-      - "Optional NVFP4 text encoder (Blackwell only): hf download FastVideo/FastVideo-FastH3-8-Step-V2-NVFP4-Consumer --include 'text_encoder/*' --local-dir fasth3-nvfp4-encoder"
    options: &omniref_options
      - server.*
      - generator.engine.offload.text_encoder
      - generator.pipeline.experimental.h3_vae_nvfp4_decoder
-      - generator.pipeline.components.text_encoder_weights
      - default_request.sampling.num_frames
      - default_request.sampling.height
      - default_request.sampling.width
@@ -94,7 +92,6 @@ deployments:
        directories in server.media_roots.
      - Based on the maintained inference recipe; this serving adaptation has not been GPU-validated.
      - "Optional NVFP4 decoder (Blackwell only): hf download FastVideo/FastH3-NVFP4-VAE-Decoder --include 'quality-36L/decoder-nvfp4.safetensors' 'fast-8L/decoder-nvfp4.safetensors' --local-dir fasth3-nvfp4-vae"
-      - "Optional NVFP4 text encoder (Blackwell only): hf download FastVideo/FastVideo-FastH3-8-Step-V2-NVFP4-Consumer --include 'text_encoder/*' --local-dir fasth3-nvfp4-encoder"
    options: *omniref_options
    overrides:
      default_request.sampling.num_frames:
@@ -145,6 +142,5 @@ deployments:
        directories in server.media_roots.
      - Based on the maintained inference recipe; this serving adaptation has not been GPU-validated.
      - "Optional NVFP4 decoder (Blackwell only): hf download FastVideo/FastH3-NVFP4-VAE-Decoder --include 'quality-36L/decoder-nvfp4.safetensors' 'fast-8L/decoder-nvfp4.safetensors' --local-dir fasth3-nvfp4-vae"
-      - "Optional NVFP4 text encoder (Blackwell only): hf download FastVideo/FastVideo-FastH3-8-Step-V2-NVFP4-Consumer --include 'text_encoder/*' --local-dir fasth3-nvfp4-encoder"
    options: *omniref_options
    overrides: *omniref_overrides
@@ -21,7 +21,6 @@ deployments:
      - Based on the released one-Spark inference recipe. This serving adaptation and edited hardware combinations have not been GPU-validated.
      - The checkpoint includes NVFP4 weights and encoder, a lightweight video VAE, and its trained eight-forward schedule. No local CompactH3 assembly is needed.
      - "Optional NVFP4 decoder (Blackwell only): hf download FastVideo/FastH3-NVFP4-VAE-Decoder --include 'quality-36L/decoder-nvfp4.safetensors' 'fast-8L/decoder-nvfp4.safetensors' --local-dir fasth3-nvfp4-vae"
-      - "Optional NVFP4 text encoder (Blackwell only): hf download FastVideo/FastVideo-FastH3-8-Step-V2-NVFP4-Consumer --include 'text_encoder/*' --local-dir fasth3-nvfp4-encoder"
    options:
      - server.*
      - generator.engine.num_gpus
@@ -31,7 +30,6 @@ deployments:
      - generator.engine.offload.vae
      - generator.pipeline.experimental.video_decode_backend
      - generator.pipeline.experimental.h3_vae_nvfp4_decoder
-      - generator.pipeline.components.text_encoder_weights
      - default_request.prompt
      - default_request.sampling.num_frames
      - default_request.sampling.height
@@ -35,7 +35,7 @@ The preprocessing steps are identical to training. Run the appropriate preproces

 ```bash
 # For Wan2.1 T2V
-bash scripts/preprocess/v1_preprocess_wan_data_t2v.sh
+bash scripts/preprocess/v1_preprocess_wan_data_t2v

 # For Wan2.2 TI2V  
 bash examples/distill/Wan2.2-TI2V-5B-Diffusers/crush_smol/preprocess_wan_data_ti2v_5b.sh
@@ -11,16 +11,16 @@ Attention. See [Attn-QAT Training](../training/attn_qat.md).

 ## 📊 Model Overview

-We provide three distilled models:
+We provide two distilled models:

 - **[FastWan2.1-T2V-1.3B-Diffusers](https://huggingface.co/FastVideo/FastWan2.1-T2V-1.3B-Diffusers)**: 3-step inference, up to **16 FPS** on H100 GPU
 - **[FastWan2.1-T2V-14B-480P-Diffusers](https://huggingface.co/FastVideo/FastWan2.1-T2V-14B-480P-Diffusers)**: 3-step inference, up to **60x speed up** at 480P, **90x speed up** at 720P for denoising loop
 - **[FastWan2.2-TI2V-5B-FullAttn-Diffusers](https://huggingface.co/FastVideo/FastWan2.2-TI2V-5B-FullAttn-Diffusers)**: 3-step inference, up to **50x speed up** at 720P for denoising loop

-All three models are trained on 61 frames at 448×832 but support generating videos with **any resolution** (1.3B  model mainly support 480P, 14B model support 480P and 720P, quality may degrade for different resolutions).
+Both models are trained on **61×448×832** resolution but support generating videos with **any resolution** (1.3B  model mainly support 480P, 14B model support 480P and 720P, quality may degrade for different resolutions).

 ## ⚙️ Inference
-First install [VSA](../attention/vsa/index.md), then run the 1.3B config. The config pins the checkpoint; edit `model_path` there to use a local copy.
+First install [VSA](../attention/vsa/index.md). Set `MODEL_BASE` to your own model path and run:

 ```bash
 FASTVIDEO_ATTENTION_BACKEND=VIDEO_SPARSE_ATTN \
@@ -45,7 +45,7 @@ python scripts/huggingface/download_hf.py \

 ### Wan2.1 1.3B Model Sparse-Distill

-For the 1.3B model, we use **8 nodes with 64 H200 GPUs** (8 GPUs per node):
+For the 1.3B model, we use **4 nodes with 32 H200 GPUs** (8 GPUs per node):

 ```bash
 # Multi-node training (8 nodes, 64 GPUs total)
@@ -6,9 +6,9 @@ load stands down when lazy owns deferral). Two boxes connected
 by the QSFP ConnectX-7 cables can run **one clip faster** and can hold a
 **longer clip** (up to the FastH3 15 s cap).

-The [MiniMax H3 cookbook](../../cookbook/minimax-h3.md) covers one Spark. For
-two, run `examples/inference/basic/basic_fasth3_spark_pair.yaml`; it has no
-serving baseline.
+Copy-paste commands for both counts are on the
+[MiniMax H3 cookbook](../../cookbook/minimax-h3.md): FastH3 V1 → NVIDIA DGX
+Spark → 1 Spark or 2 Sparks.

 This is FastVideo sequence parallel (`sp_size=2`) over Ray, not a third-party
 xDiT vendor. Do not install xDiT for this path.
@@ -156,7 +156,7 @@ lightweight video VAE. Use the environment above plus the release kernel
 settings, then run one of these configs from the head:

 ```bash
-export FASTVIDEO_MINIMAX_H3_FUSIONS=all FASTVIDEO_NVFP4_MM_BACKEND=cudnn
+export FASTVIDEO_MINIMAX_H3_FUSIONS=all FASTVIDEO_NVFP4_MM_BACKEND=cutlass
 export FASTVIDEO_H3_VAE_TILE_BATCH=1 FASTVIDEO_VSA_TRITON=1
 export FASTVIDEO_VSA_PTX_BLACKWELL=0 FASTVIDEO_FA4=0
 export FASTVIDEO_ATTENTION_BACKEND=VIDEO_SPARSE_ATTN_H3 FASTVIDEO_STAGE_LOGGING=1
@@ -1,44 +1,28 @@
-# Quick start
+# 🚀 Quick Start

-Install FastVideo, generate one video, then start from a maintained recipe.
+Get up and running with FastVideo in minutes!

-## Install
+## Installation

-Requires Python 3.10 to 3.12 and CUDA 12.6 or newer.
+First, install FastVideo:

 ```bash
+# If you previously used Conda, use uv instead for a faster, more stable setup
 uv venv --python 3.12 --seed
 source .venv/bin/activate

-# CUDA 12. Use cu130 on CUDA 13.
+# Install FastVideo on NVIDIA CUDA 12
 UV_TORCH_BACKEND=cu126 uv pip install fastvideo
 ```

-For the fastest attention kernels, install flash-attn:
+Use `UV_TORCH_BACKEND=cu130` instead on CUDA 13.
+
+Also optionally install flash-attn:

 ```bash
 uv pip install flash-attn --no-build-isolation -v
 ```

-## Generate a video
-
-```python
-from fastvideo import VideoGenerator
-
-generator = VideoGenerator.from_pretrained(
-    "Wan-AI/Wan2.1-T2V-1.3B-Diffusers",
-    num_gpus=1,
-    text_encoder_cpu_offload=True,
-)
-generator.generate_video(
-    "A raccoon peers through a field of sunflowers.",
-    output_path="videos/",
-    save_video=True,
-)
-```
-
-The first run downloads the Wan2.1 1.3B checkpoint. The call writes an MP4 under `videos/`. The full example is `examples/inference/basic/basic.py`.
-
 ## Choose a maintained recipe

 The cookbook selects complete, checked-in recipes instead of mixing model,
@@ -1,6 +1,6 @@
 # Welcome to FastVideo

-FastVideo is an inference and post-training framework for video diffusion models. One pipeline covers data preprocessing, fine-tuning, distillation, and inference. Attention backends, parallelism, and quantization are separate modules, so you can add one without changing model code.
+FastVideo is an inference and post-training framework for diffusion models. It features an end-to-end unified pipeline for accelerating diffusion models, starting from data preprocessing to model training, finetuning, distillation, and inference. FastVideo is designed to be modular and extensible, allowing users to easily add new optimizations and techniques. Whether it is training-free optimizations or post-training optimizations, FastVideo has you covered.

 ## Key Features

@@ -90,7 +90,7 @@ PipelineConfig                    (fastvideo/configs/pipelines/base.py)
 │   ├── WanT2V720PConfig
 │   └── WanI2V480PConfig
 ├── HunyuanConfig                 (fastvideo/configs/pipelines/hunyuan.py)
-├── LTX2T2VConfig                 (fastvideo/pipelines/basic/ltx2/pipeline_configs.py)
+├── LTX2T2VConfig                 (fastvideo/configs/pipelines/ltx2.py)
 ├── CosmosConfig                  (fastvideo/configs/pipelines/cosmos.py)
 └── ... (15+ model families)
 ```
@@ -136,7 +136,7 @@ Concrete hierarchy: `DiTConfig` → `DiTArchConfig`, `VAEConfig` →
 ### SamplingParam (`fastvideo/api/sampling_param.py`)

 Generation parameters separate from pipeline config. Each model family
-provides defaults via a profile (see `fastvideo/pipelines/basic/<family>/presets.py`):
+provides defaults via a profile (see `fastvideo/pipelines/basic/<family>/profiles.py`):

 ```python
 sp = SamplingParam.from_pretrained("Wan-AI/Wan2.1-T2V-1.3B-Diffusers")
@@ -505,7 +505,7 @@ User: generator.generate_video(prompt, ...)
   flow_shift, precision defaults.

 2. **Sampling param profile** — Create a profile in
-   `fastvideo/pipelines/basic/<model>/presets.py` with default height,
+   `fastvideo/pipelines/basic/<model>/profiles.py` with default height,
   width, num_frames, guidance_scale, num_inference_steps.

 3. **Register configs** — In `fastvideo/registry.py`, add a
@@ -72,7 +72,7 @@ def main():
    sampling_param.guidance_scale = 7.5
    
    # Random seed for reproducibility
-    sampling_param.seed = 42  # Optional; the default is 1024
+    sampling_param.seed = 42  # Optional, leave unset for random results

    # Generate video with custom parameters
    prompt = "A beautiful sunset over a calm ocean, with gentle waves."
@@ -0,0 +1 @@
+/Users/aryank/os/FastVideo-review-int/docs/node_modules
@@ -12,25 +12,6 @@ function at(value, path) {
  return path.split(".").reduce((current, key) => current?.[key], value);
 }

-test("every local NVFP4 encoder choice includes its download requirement", () => {
-  const path = "generator.pipeline.components.text_encoder_weights";
-  const catalogs = buildCatalogs().filter((catalog) => catalog.controls.some((control) => control.path === path));
-  assert.ok(catalogs.length > 0);
-  for (const catalog of catalogs) {
-    const option = catalog.controls.find((control) => control.path === path);
-    if (!option.schema.enum?.includes("./fasth3-nvfp4-encoder/text_encoder")) continue;
-    assert.ok(
-      catalog.requirements.some(
-        (requirement) =>
-          requirement.includes("hf download ") &&
-          requirement.includes("--include 'text_encoder/*'") &&
-          requirement.includes("--local-dir fasth3-nvfp4-encoder"),
-      ),
-      `${catalog.id} must explain how to create the selected encoder directory`,
-    );
-  }
-});
-
 function fixture({ interfaceName = "rest", backend = "cuda", workload = "t2v" } = {}) {
  return {
    id: "example/deployment",
@@ -56,7 +56,7 @@ torchrun --nproc_per_node=$GPU_NUM \
 | `--preprocess.dataset_path` | Path to dataset (HF repo ID or local folder) |
 | `--preprocess.dataset_output_dir` | Output directory for Parquet files |
 | `--preprocess.video_loader_type` | Video decoder: `torchcodec` or `torchvision` |
-| `--preprocess.max_height` / `--preprocess.max_width` | Target resolution for videos |
+| `--preprocess.max_height` / `max_width` | Target resolution for videos |
 | `--preprocess.num_frames` | Number of frames to extract per video |
 | `--preprocess.train_fps` | Target FPS for frame extraction |

@@ -1,12 +1,12 @@
-# Fine-tuning
+# 🧠 Finetuning

 This guide covers finetuning video diffusion models with FastVideo, including full finetuning and LoRA.

-## Training arguments
+## Training Arguments

 FastVideo training scripts use several argument groups:

-### Loop and model arguments
+### Training Arguments

 | Argument | Description |
 |----------|-------------|
@@ -12,7 +12,7 @@ parquet that step produces.

 ## What you need

-| Requirement | Value |
+| | |
 |---|---|
 | GPU | 48GB or more for the LoRA overfit; multi-GPU for full-parameter runs |
 | Disk | ~40GB for the checkpoint (DiT 17GB + Qwen2.5-VL 16GB + ByT5/VAE) |
@@ -145,10 +145,16 @@ a real problem — stop and investigate rather than burning GPU hours.
 ## 4. Scale out

 `examples/train/configs/fine_tuning/hunyuan15/t2v.yaml` runs
-full-parameter training across 8 GPUs with FSDP. One 80 GB card cannot hold it:
+full-parameter training across 8 GPUs with FSDP. Full-parameter training
+does not fit one 80GB card:

-| Item | Per rank (bf16) |
-|---|---:|
+It reads `data/hunyuan15_preprocessed` (and its `validation_prompts.json`),
+not the overfit directory above — either preprocess into that path (set
+`FASTVIDEO_TEST_HUNYUAN15_OVERFIT_OUTPUT_DIR=data/hunyuan15_preprocessed`) or
+pass `--training.data.data_path` and `--callbacks.validation.dataset_file`.
+
+| | bf16 |
+|---|---|
 | Weights | 16.7 GB |
 | Gradients | 16.7 GB |
 | AdamW `exp_avg` + `exp_avg_sq` | 33.4 GB |
@@ -157,27 +163,21 @@ full-parameter training across 8 GPUs with FSDP. One 80 GB card cannot hold it:
 | **Total** | **~80 GB** |

 That subtotal is independent of sequence length, so shortening the clip
-(`--training.data.num_latent_t`) barely helps: measured peak memory moved by
-0.07 GB going from 21 latent frames to 9. Options that do move the needle:
+(`--training.data.num_latent_t`) barely helps — measured peak memory moved
+by 0.07GB going from 21 latent frames down to 9. Options that do move the
+needle:

- **LoRA** (what the overfit config uses): gradients and optimizer state
-  shrink to the adapter parameters, about 30 GB total.
- **FSDP sharding across GPUs**: the 66.8 GB subtotal divides by the number of
-  shards.
- **A memory-light optimizer**: SGD with momentum keeps a single state tensor
-  instead of two. `torch.optim.Adafactor` is smaller still, but is not
-  DTensor-safe. Its in-place `.square_()` fails under FSDP2 with
-  `aten.pow_.Scalar: in-place operations that require placement changes are
-  not supported`.
+- **LoRA** (what the overfit config uses) — gradients and optimizer state
+  shrink to the adapter parameters, ~30GB total.
+- **FSDP sharding across GPUs** — the 66.8GB subtotal divides by the number
+  of shards.
+- **A memory-light optimizer** — SGD with momentum keeps a single state
+  tensor instead of two. `torch.optim.Adafactor` is smaller still, but is
+  not DTensor-safe: its in-place `.square_()` fails under FSDP2 with
+  `aten.pow_.Scalar: in-place operations that require placement changes
+  are not supported`.

-The config reads `data/hunyuan15_preprocessed` and its
-`validation_prompts.json`, not the overfit directory above. Either preprocess
-into that path (set
-`FASTVIDEO_TEST_HUNYUAN15_OVERFIT_OUTPUT_DIR=data/hunyuan15_preprocessed`) or
-pass `--training.data.data_path` and `--callbacks.validation.dataset_file`.
-
-The shipped overfit config already sets `training.dit_precision: bf16`. If you
-write your own config, set it for any single-device run. The default
+Set `training.dit_precision: bf16` for any single-device run. The default
 is fp32, which doubles weights, gradients, and optimizer state.

 ## Troubleshooting
@@ -18,7 +18,7 @@ should use the config-driven system documented here.
 ### Launch with the helper script

 ```bash
-bash examples/train/run.sh examples/train/configs/distribution_matching/wan/dmd2_t2v.yaml
+bash examples/train/run.sh examples/train/distill_wan2.1_t2v_1.3B_dmd2.yaml
 ```

 The script auto-detects available GPUs and sets up `torchrun`. Override with
@@ -35,7 +35,7 @@ NUM_GPUS=4 NNODES=2 NODE_RANK=0 \
 ```bash
 torchrun --nproc_per_node=8 \
    fastvideo/train/entrypoint/train.py \
-    --config examples/train/configs/distribution_matching/wan/dmd2_t2v.yaml
+    --config examples/train/distill_wan2.1_t2v_1.3B_dmd2.yaml
 ```

 ### CLI flags
@@ -136,16 +136,44 @@ class DistributedAttention(nn.Module):
        """
        # Check input shapes
        assert q.dim() == 4 and k.dim() == 4 and v.dim() == 4, "Expected 4D tensors"
-        batch_size, _, num_heads, _ = q.shape
+        # Stack QKV
+        qkv = torch.cat([q, k, v], dim=0)  # [3*batch, seq_len, num_heads, head_dim]
+        return self.forward_stacked(
+            qkv,
+            original_seq_len=original_seq_len,
+            replicated_q=replicated_q,
+            replicated_k=replicated_k,
+            replicated_v=replicated_v,
+            freqs_cis=freqs_cis,
+            attn_metadata=attn_metadata,
+        )
+
+    def forward_stacked(
+        self,
+        qkv: torch.Tensor,
+        original_seq_len: int | None = None,
+        replicated_q: torch.Tensor | None = None,
+        replicated_k: torch.Tensor | None = None,
+        replicated_v: torch.Tensor | None = None,
+        freqs_cis: tuple[torch.Tensor, torch.Tensor] | None = None,
+        attn_metadata: AttentionMetadata | None = None,
+    ) -> tuple[torch.Tensor, torch.Tensor | None]:
+        """Attention over a pre-stacked ``[3*batch, seq_len, num_heads, head_dim]`` QKV buffer.
+
+        Identical to :meth:`forward` after its ``torch.cat``, so a fused QKV
+        prologue that already wrote the stack (e.g. MiniMax H3's MotionKernel
+        prologue) can skip that copy. Sequence-parallel redistribution,
+        ``freqs_cis``, replicated tokens and backend pre/post-processing are
+        handled exactly as in :meth:`forward`.
+        """
+        stacked, _, num_heads, _ = qkv.shape
+        batch_size = stacked // 3
        local_rank = get_sp_parallel_rank()
        world_size = get_sp_world_size()

        forward_context: ForwardContext = get_forward_context()
        ctx_attn_metadata = attn_metadata if attn_metadata is not None else forward_context.attn_metadata

-        # Stack QKV
-        qkv = torch.cat([q, k, v], dim=0)  # [3*batch, seq_len, num_heads, head_dim]
-
        # Redistribute heads across sequence dimension
        qkv = sequence_model_parallel_all_to_all_4D(qkv, scatter_dim=2, gather_dim=1)

@@ -427,7 +427,8 @@ FASTVIDEO_MINIMAX_H3_FUSIONS = EnvStr(
    "",
    category="performance",
    doc="MiniMax-H3 inference-only Triton fusions: all, 1, or a comma-separated subset of "
-    "modulate,qknorm_rope,swiglu. Empty, 0, or none keeps the eager implementation.")
+    "modulate,qknorm_rope,swiglu. The whole-tier value mk selects the MotionKernel kernels instead of "
+    "the Sol-Engine ones. Empty, 0, or none keeps the eager implementation.")
 FASTVIDEO_MINIMAX_H3_EXACT_KERNELS = EnvStr(
    "",
    category="performance",
@@ -38,6 +38,7 @@ from fastvideo.models.dits.base import BaseDiT
 from fastvideo.models.dits.minimax_h3_vsa_fp4 import (STAGES, vsa_fp4_attention, vsa_fp4_attention_sp,
                                                       vsa_fp4_requested, vsa_tile_first_attention)
 from fastvideo.models.dits.minimax_h3_fusions import exact as exact_kernels
+from fastvideo.models.dits.minimax_h3_fusions import mk as mk_kernels
 from fastvideo.models.dits.minimax_h3_fusions import (
    HAVE_TRITON,
    fused_qknorm_rope,
@@ -52,9 +53,13 @@ from fastvideo.utils import get_compute_dtype
 logger = init_logger(__name__)

 MINIMAX_H3_MODALITY_NUM = 3
+#: eps baked into the MotionKernel H3 kernels; their call sites require it.
+_MK_NORM_EPS = 1e-5
 _CFG = MiniMaxH3Config()
 _MINIMAX_H3_FUSION_NAMES = frozenset({"modulate", "qknorm_rope", "swiglu"})
 _MINIMAX_H3_EXACT_KERNEL_NAMES = frozenset({"rope", "modulate", "swiglu"})
+#: Whole-tier opt-in that replaces the Sol kernels with the MotionKernel ones.
+_MINIMAX_H3_MK_MODE = "mk"


 def _enabled_minimax_h3_fusions(value: str | None = None) -> frozenset[str]:
@@ -65,7 +70,14 @@ def _enabled_minimax_h3_fusions(value: str | None = None) -> frozenset[str]:
        return frozenset()
    if normalized in {"1", "all"}:
        return _MINIMAX_H3_FUSION_NAMES
+    if normalized == _MINIMAX_H3_MK_MODE:
+        # ``mk`` is a whole-tier mode: every fusion is on, and each fused call
+        # site dispatches to the MotionKernel kernel instead of the Sol one.
+        return _MINIMAX_H3_FUSION_NAMES | {_MINIMAX_H3_MK_MODE}
    enabled = frozenset(item.strip() for item in normalized.split(",") if item.strip())
+    if _MINIMAX_H3_MK_MODE in enabled:
+        raise ValueError(f"'{_MINIMAX_H3_MK_MODE}' is a whole-tier mode and cannot be combined with individual "
+                         f"fusions: {sorted(enabled - {_MINIMAX_H3_MK_MODE})}")
    unknown = enabled - _MINIMAX_H3_FUSION_NAMES
    if unknown:
        supported = ",".join(sorted(_MINIMAX_H3_FUSION_NAMES))
@@ -139,6 +151,7 @@ class MiniMaxH3FeedForward(nn.Module):
        prefix: str = "",
        fuse_swiglu: bool = False,
        exact_swiglu: bool = False,
+        mk_swiglu: bool = False,
    ) -> None:
        super().__init__()
        self.fc_in = ReplicatedLinear(
@@ -157,6 +170,7 @@ class MiniMaxH3FeedForward(nn.Module):
        )
        self.fuse_swiglu = fuse_swiglu
        self.exact_swiglu = exact_swiglu
+        self.mk_swiglu = mk_swiglu
        self.use_mxfp8 = isinstance(self.fc_in.quant_method, MXFP8QuantizeMethod) and isinstance(
            self.fc_out.quant_method, MXFP8QuantizeMethod)
        # Inference-only token chunking: the 2 * ffn_dim intermediate is ~5.3x the block input
@@ -181,7 +195,10 @@ class MiniMaxH3FeedForward(nn.Module):
            return mxfp8_swiglu_feed_forward(hidden_states, self.fc_in, self.fc_out)
        hidden_states, _ = self.fc_in(hidden_states)
        if self.fuse_swiglu and _can_run_minimax_h3_fusion(hidden_states):
-            hidden_states = minimax_h3_swiglu(hidden_states)
+            if self.mk_swiglu and mk_kernels.HAVE_TRITON and not torch.compiler.is_compiling():
+                hidden_states = mk_kernels.swiglu(hidden_states)
+            else:
+                hidden_states = minimax_h3_swiglu(hidden_states)
        elif self.exact_swiglu and exact_kernels.supports_rowwise(hidden_states):
            hidden_states = exact_kernels.swiglu(hidden_states)
        else:
@@ -206,6 +223,7 @@ class MiniMaxH3Attention(nn.Module):
        fuse_qknorm_rope: bool = False,
        fa4_packed_varlen: bool = False,
        exact_rope: bool = False,
+        mk_qkv: bool = False,
        vsa_fp4: dict[str, Any] | None = None,
    ) -> None:
        super().__init__()
@@ -244,6 +262,7 @@ class MiniMaxH3Attention(nn.Module):
        )
        self.fuse_qknorm_rope = fuse_qknorm_rope
        self.exact_rope = exact_rope
+        self.mk_qkv = mk_qkv
        # VSA carries a learned gate on its pooled-compression branch. The H3
        # checkpoint has no such weight, so the loader zero-initializes it
        # (ALLOWED_NEW_PARAM_PATTERNS) and the branch is exactly disabled
@@ -401,6 +420,22 @@ class MiniMaxH3Attention(nn.Module):
            cos, sin = rotary_emb
            cos = cos.to(query.dtype)
            sin = sin.to(query.dtype)
+            if (self.mk_qkv and mk_kernels.HAVE_TRITON and not torch.compiler.is_compiling()
+                    and self.norm_q.eps == _MK_NORM_EPS
+                    and self.norm_k.eps == _MK_NORM_EPS and type(self.distributed_attention) is DistributedAttention
+                    and not (self.to_gate_compress is not None and self._gate_active())):
+                # One launch replaces the two Sol QK-norm/RoPE kernels and the
+                # ``torch.cat`` stack: the prologue writes a contiguous
+                # ``[3*B, S, H, D]`` buffer that ``forward_stacked`` feeds to the
+                # backend directly, so q/k/v are never copied again.
+                qkv = mk_kernels.qkv_prologue(query, key, value, self.norm_q.weight, self.norm_k.weight, cos, sin)
+                with STAGES.span("attention"):
+                    hidden_states, _ = self.distributed_attention.forward_stacked(
+                        qkv, original_seq_len=original_seq_len, attn_metadata=attn_metadata)
+                hidden_states = hidden_states.flatten(2, 3).type_as(qkv)
+                with STAGES.span("out_proj"):
+                    hidden_states, _ = self.to_out(hidden_states)
+                return hidden_states
            query = fused_qknorm_rope(query, self.norm_q.weight, cos, sin, self.norm_q.eps)
            key = fused_qknorm_rope(key, self.norm_k.weight, cos, sin, self.norm_k.eps)
        else:
@@ -658,6 +693,7 @@ class MiniMaxH3TransformerBlock(nn.Module):
        fa4_packed_varlen: bool = False,
        exact_kernels: frozenset[str] = frozenset(),
        vsa_fp4: dict[str, Any] | None = None,
+        mk_tier: bool = False,
    ) -> None:
        super().__init__()
        self.norm1 = nn.RMSNorm(hidden_size, eps=norm_eps)
@@ -673,6 +709,7 @@ class MiniMaxH3TransformerBlock(nn.Module):
            fa4_packed_varlen=fa4_packed_varlen,
            exact_rope="rope" in exact_kernels,
            vsa_fp4=vsa_fp4,
+            mk_qkv=mk_tier,
        )
        self.norm2 = nn.RMSNorm(hidden_size, eps=norm_eps)
        self.ff = MiniMaxH3FeedForward(
@@ -682,6 +719,7 @@ class MiniMaxH3TransformerBlock(nn.Module):
            prefix=f"{prefix}.ff",
            fuse_swiglu=fuse_swiglu,
            exact_swiglu="swiglu" in exact_kernels,
+            mk_swiglu=mk_tier,
        )
        self.adaln_proj = MiniMaxH3AdaLayerNormModulation(
            time_embed_dim,
@@ -692,6 +730,7 @@ class MiniMaxH3TransformerBlock(nn.Module):
        )
        self.fuse_modulate = fuse_modulate
        self.exact_modulate = "modulate" in exact_kernels
+        self.mk_tier = mk_tier

    def forward(
        self,
@@ -711,15 +750,20 @@ class MiniMaxH3TransformerBlock(nn.Module):
                              and exact_kernels.supports_rowwise(hidden_states, shift_msa, scale_msa, gate_msa,
                                                                 shift_mlp, scale_mlp, gate_mlp))
        if use_modulate_fusion:
-            with nvtx_range("minimax_h3.transformer_block.modulate_fusion"):
-                norm_hidden_states = fused_rmsnorm_modulate(
-                    hidden_states,
-                    self.norm1.weight,
-                    scale_msa,
-                    shift_msa,
-                    adaln_indices,
-                    self.norm1.eps,
-                )
+            if (self.mk_tier and mk_kernels.HAVE_TRITON and not torch.compiler.is_compiling()
+                    and self.norm1.eps == _MK_NORM_EPS):
+                norm_hidden_states = mk_kernels.rmsnorm_modulate(hidden_states, self.norm1.weight, scale_msa,
+                                                                 shift_msa, adaln_indices)
+            else:
+                with nvtx_range("minimax_h3.transformer_block.modulate_fusion"):
+                    norm_hidden_states = fused_rmsnorm_modulate(
+                        hidden_states,
+                        self.norm1.weight,
+                        scale_msa,
+                        shift_msa,
+                        adaln_indices,
+                        self.norm1.eps,
+                    )
        elif use_exact_modulate:
            with nvtx_range("minimax_h3.transformer_block.exact_modulate"):
                norm_hidden_states = exact_kernels.modulate(self.norm1(hidden_states), scale_msa, shift_msa,
@@ -735,17 +779,23 @@ class MiniMaxH3TransformerBlock(nn.Module):
        # overlaps three full-width activations during the residual fusion.
        del norm_hidden_states
        if use_modulate_fusion:
-            with nvtx_range("minimax_h3.transformer_block.modulate_fusion"):
-                hidden_states, norm_hidden_states = fused_residual_gate_rmsnorm_modulate(
-                    hidden_states,
-                    attention_output,
-                    gate_msa,
-                    self.norm2.weight,
-                    scale_mlp,
-                    shift_mlp,
-                    adaln_indices,
-                    self.norm2.eps,
-                )
+            if (self.mk_tier and mk_kernels.HAVE_TRITON and not torch.compiler.is_compiling()
+                    and self.norm2.eps == _MK_NORM_EPS):
+                hidden_states, norm_hidden_states = mk_kernels.residual_gate_rmsnorm_modulate(
+                    hidden_states, attention_output, gate_msa, self.norm2.weight, scale_mlp, shift_mlp,
+                    adaln_indices)
+            else:
+                with nvtx_range("minimax_h3.transformer_block.modulate_fusion"):
+                    hidden_states, norm_hidden_states = fused_residual_gate_rmsnorm_modulate(
+                        hidden_states,
+                        attention_output,
+                        gate_msa,
+                        self.norm2.weight,
+                        scale_mlp,
+                        shift_mlp,
+                        adaln_indices,
+                        self.norm2.eps,
+                    )
        elif use_exact_modulate:
            with nvtx_range("minimax_h3.transformer_block.exact_modulate"):
                hidden_states = exact_kernels.gate_residual(hidden_states, gate_msa, attention_output, adaln_indices)
@@ -761,6 +811,8 @@ class MiniMaxH3TransformerBlock(nn.Module):
        with nvtx_range("minimax_h3.transformer_block.feed_forward"), STAGES.span("feed_forward"):
            feed_forward_output = self.ff(norm_hidden_states)
        if use_modulate_fusion and not torch.compiler.is_compiling():
+            if self.mk_tier and mk_kernels.HAVE_TRITON and not torch.compiler.is_compiling():
+                return mk_kernels.gated_residual(hidden_states, gate_mlp, feed_forward_output, adaln_indices)
            return _gated_residual(hidden_states, gate_mlp, adaln_indices, feed_forward_output)
        if use_exact_modulate:
            return exact_kernels.gate_residual(hidden_states, gate_mlp, feed_forward_output, adaln_indices)
@@ -821,6 +873,10 @@ class MiniMaxH3Transformer3DModel(BaseDiT):
                logger.warning(
                    "FASTVIDEO_MINIMAX_H3_FUSIONS requested %s but Triton is unavailable; "
                    "every forward stays on the eager path.", ",".join(sorted(self.enabled_fusions)))
+        self.use_mk_kernels = _MINIMAX_H3_MK_MODE in self.enabled_fusions
+        if self.use_mk_kernels and not mk_kernels.HAVE_TRITON:
+            logger.warning("FASTVIDEO_MINIMAX_H3_FUSIONS=mk requested but the MotionKernel kernels are "
+                           "unavailable (Triton missing); every forward stays on the eager path.")
        self.enabled_exact_kernels = _enabled_minimax_h3_exact_kernels()
        if self.enabled_exact_kernels:
            if exact_kernels.HAVE_TRITON:
@@ -935,6 +991,7 @@ class MiniMaxH3Transformer3DModel(BaseDiT):
                fa4_packed_varlen=envs.FASTVIDEO_MINIMAX_H3_FA4_PACKED_VARLEN.get(),
                exact_kernels=self.enabled_exact_kernels,
                vsa_fp4=config.vsa_fp4,
+                mk_tier=self.use_mk_kernels,
            ) for index in range(arch.num_layers)
        ])
        self.norm_out = MiniMaxH3AdaLayerNormOut(
@@ -0,0 +1,49 @@
+# SPDX-License-Identifier: Apache-2.0
+"""MotionKernel H3 kernels (opt-in through ``FASTVIDEO_MINIMAX_H3_FUSIONS=mk``).
+
+These are the MotionKernel agent-discovered Triton kernels for the MiniMax-H3
+inference blocks, kept separate from the Sol-Engine kernels in the parent
+package so the default path stays byte-for-byte identical. Each kernel is a
+drop-in replacement for one Sol fusion or eager expression:
+
+============================================  ======================================
+MotionKernel                                   replaces
+============================================  ======================================
+``qkv_prologue``                              2x ``fused_qknorm_rope`` + the
+                                              ``torch.cat`` QKV stack
+``rmsnorm_modulate``                          ``fused_rmsnorm_modulate``
+``residual_gate_rmsnorm_modulate``            ``fused_residual_gate_rmsnorm_modulate``
+``swiglu``                                    ``minimax_h3_swiglu``
+``gated_residual``                            the compiled ``_gated_residual``
+============================================  ======================================
+
+All five run in fp32 internally with a single rounding at the bf16/fp16 store
+and require the production H3 norm eps of ``1e-5`` (enforced at the call sites).
+"""
+
+try:
+    import triton  # noqa: F401  (each module imports it directly)
+
+    from .gated_residual import gated_residual
+    from .qkv_prologue import qkv_prologue
+    from .residual_gate_rmsnorm_modulate import residual_gate_rmsnorm_modulate
+    from .rmsnorm_modulate import rmsnorm_modulate
+    from .swiglu import swiglu
+
+    HAVE_TRITON = True
+except ImportError:  # pragma: no cover - environments without Triton
+    gated_residual = None
+    qkv_prologue = None
+    residual_gate_rmsnorm_modulate = None
+    rmsnorm_modulate = None
+    swiglu = None
+    HAVE_TRITON = False
+
+__all__ = [
+    "HAVE_TRITON",
+    "gated_residual",
+    "qkv_prologue",
+    "residual_gate_rmsnorm_modulate",
+    "rmsnorm_modulate",
+    "swiglu",
+]
@@ -0,0 +1,151 @@
+# SPDX-License-Identifier: Apache-2.0
+"""MotionKernel H3 kernel: Fused H3 row-indexed gated residual (hidden + gate[index] * branch).
+
+Found with MotionKernel (agent-driven kernel search). Tuned on GB200 / sm_100 for the
+MiniMax-H3 inference shapes; opt-in through ``FASTVIDEO_MINIMAX_H3_FUSIONS=mk``.
+
+Triton kernel for MiniMax H3's post-FFN row-indexed gated residual.
+
+``out = hidden + gate[index] * branch`` with the math in fp32 and a single
+bf16 rounding at the store.
+
+Design notes (GB200, sm_100a):
+  * The op moves 3 activation-sized streams (2 reads + 1 write); it is purely
+    bandwidth bound.  A flat 1-D grid over ``tokens * hidden`` gives fully
+    coalesced, 16-byte-vectorized accesses on all three streams with no lane
+    padding (the 5376-wide rows are not padded to a power of two).
+  * Split full-block / tail code paths.  Every non-tail block runs completely
+    unmasked, so the compiler emits unpredicated 128-bit loads/stores; only the
+    final block pays for predicates.  Measured 175.6 -> 173.6 us weighted over
+    the five production shapes (~1.1%), consistently on every shape.
+  * The gathered AdaLN gate row is tiny (<= 12 rows) so it stays resident in
+    L2; the per-element row/column split needed to address it costs only ALU,
+    of which this kernel uses ~1% of peak.
+  * 32-bit offsets (all production shapes fit) halve index register pressure.
+  * ``evict_last`` on the three loads is worth ~2%: it keeps the streamed
+    lines in L2 long enough for the DRAM scheduler to coalesce the
+    read/write turnaround (measured 269 -> 266 us at 59k tokens).
+  * Roofline: a minimal unmasked 2-read/1-write Triton ``c = a + b`` kernel
+    (same bytes, no gather) runs at 267.7 us / 7.11 TB/s at the largest shape;
+    this kernel measures 263.2 us / 7.24 TB/s, i.e. it is at (slightly above)
+    the measured 3-stream streaming roofline, so the body is not tunable
+    further.
+"""
+
+import torch
+import triton
+import triton.language as tl
+
+
+@triton.jit
+def _h3_gated_residual_kernel(
+    out_ptr,
+    hidden_ptr,
+    branch_ptr,
+    gate_ptr,
+    index_ptr,
+    n_elems,
+    n_index,
+    D: tl.constexpr,
+    GS: tl.constexpr,
+    BLOCK: tl.constexpr,
+    NFULL: tl.constexpr,
+    WIDE: tl.constexpr,
+):
+    pid = tl.program_id(0)
+    if WIDE:
+        off = pid.to(tl.int64) * BLOCK + tl.arange(0, BLOCK)
+        mask = off < n_elems
+        row = off // D
+        col = off - row * D
+        table_row = tl.load(index_ptr + row % n_index, mask=mask, other=0).to(tl.int32)
+        hidden = tl.load(hidden_ptr + off, mask=mask, eviction_policy="evict_last").to(tl.float32)
+        branch = tl.load(branch_ptr + off, mask=mask, eviction_policy="evict_last").to(tl.float32)
+        gate = tl.load(gate_ptr + table_row * GS + col, mask=mask, eviction_policy="evict_last").to(tl.float32)
+        tl.store(
+            out_ptr + off,
+            (hidden + gate * branch).to(out_ptr.dtype.element_ty),
+            mask=mask,
+        )
+    else:
+        base = pid * BLOCK
+        if base + BLOCK <= NFULL:
+            off = base + tl.arange(0, BLOCK)
+            row = off // D
+            col = off - row * D
+            table_row = tl.load(index_ptr + row % n_index).to(tl.int32)
+            hidden = tl.load(hidden_ptr + off, eviction_policy="evict_last").to(tl.float32)
+            branch = tl.load(branch_ptr + off, eviction_policy="evict_last").to(tl.float32)
+            gate = tl.load(gate_ptr + table_row * GS + col, eviction_policy="evict_last").to(tl.float32)
+            tl.store(out_ptr + off, (hidden + gate * branch).to(out_ptr.dtype.element_ty))
+        else:
+            off = base + tl.arange(0, BLOCK)
+            mask = off < n_elems
+            row = off // D
+            col = off - row * D
+            table_row = tl.load(index_ptr + row % n_index, mask=mask, other=0).to(tl.int32)
+            hidden = tl.load(hidden_ptr + off, mask=mask, eviction_policy="evict_last").to(tl.float32)
+            branch = tl.load(branch_ptr + off, mask=mask, eviction_policy="evict_last").to(tl.float32)
+            gate = tl.load(
+                gate_ptr + table_row * GS + col,
+                mask=mask,
+                eviction_policy="evict_last",
+            ).to(tl.float32)
+            tl.store(
+                out_ptr + off,
+                (hidden + gate * branch).to(out_ptr.dtype.element_ty),
+                mask=mask,
+            )
+
+
+def gated_residual(
+    hidden: torch.Tensor,
+    gate: torch.Tensor,
+    branch: torch.Tensor,
+    index: torch.Tensor,
+) -> torch.Tensor:
+    """``hidden + gate[index] * branch`` in one pass."""
+    if not hidden.is_cuda:
+        raise ValueError("h3_gated_residual requires CUDA tensors")
+    if branch.shape != hidden.shape:
+        raise ValueError("branch must match hidden")
+    hidden_size = hidden.shape[-1]
+    if index.ndim != 1 or index.numel() != hidden.shape[-2]:
+        raise ValueError("index must hold one table row per token")
+    if gate.shape[-1] != hidden_size:
+        raise ValueError("gate table must match the hidden width")
+    flat_hidden = hidden.reshape(-1, hidden_size)
+    flat_branch = branch.reshape(-1, hidden_size)
+    if not flat_hidden.is_contiguous():
+        flat_hidden = flat_hidden.contiguous()
+    if not flat_branch.is_contiguous():
+        flat_branch = flat_branch.contiguous()
+    if gate.stride(-1) != 1:
+        gate = gate.contiguous()
+    gate_row_stride = gate.stride(0)
+    index = index.contiguous()
+    out = torch.empty_like(flat_hidden)
+
+    n_elems = flat_hidden.numel()
+    if n_elems == 0:
+        return out.view_as(hidden)
+    block = 4096
+    wide = n_elems >= 2**31
+    nfull = (n_elems // block) * block
+    grid = (triton.cdiv(n_elems, block), )
+    _h3_gated_residual_kernel[grid](
+        out,
+        flat_hidden,
+        flat_branch,
+        gate,
+        index,
+        n_elems,
+        index.numel(),
+        D=hidden_size,
+        GS=gate_row_stride,
+        BLOCK=block,
+        NFULL=nfull,
+        WIDE=wide,
+        num_warps=8,
+    )
+    return out.view_as(hidden)
@@ -0,0 +1,190 @@
+# SPDX-License-Identifier: Apache-2.0
+"""MotionKernel H3 kernel: Fused H3 QKV prologue: per-head QK RMSNorm + partial RoPE + stacked QKV output.
+
+Found with MotionKernel (agent-driven kernel search). Tuned on GB200 / sm_100 for the
+MiniMax-H3 inference shapes; opt-in through ``FASTVIDEO_MINIMAX_H3_FUSIONS=mk``.
+
+Fused QKV prologue: per-head QK RMSNorm + partial RoPE on Q and K plus a straight V copy.
+
+Production runs two Sol-Engine ``_qknorm_partial_rope_kernel`` launches (one Triton program per
+128-wide (token, head) row of Q and of K) plus ``torch.cat([q, k, v])``. Here a single launch
+produces Q, K and V together, so every byte is read and written exactly once.
+
+Each program owns ``HB`` consecutive head rows of one token, i.e. a *contiguous flat* chunk of
+``HB * head_dim`` elements. Because the tile is flat and contiguous, each tensor costs exactly
+one wide (128-bit) load and one wide store per thread instead of a pile of per-sub-block
+addressing. The grid is ordered so that the head-group index is the *fastest* varying dimension:
+consecutive CTAs then cover one token's whole ``H * D`` slab in address order, which turns the
+six DRAM streams (q, k, v in; q, k, v out) into pure sequential scans (~7.1 vs ~6.6 TB/s).
+
+All three input tiles are issued as loads *before* any store (q, k and v in flight together,
+then the stores), so the six streams overlap instead of being serialised by the store that
+follows each load. With ``HB=4`` (512 elements, one warp) this reaches the measured streaming
+roofline of this GPU: 7.08 TB/s versus 7.10 TB/s for a pure three-stream copy of the same bytes
+(the previous load/store/load/store schedule measured 7.00 TB/s).
+
+The rotate-half partner ``x[i +/- rotary_dim/2]`` is materialised with an in-register
+``tl.gather`` along the head dim (warp-shuffle lowering for a single warp), so the whole row is
+loaded once: no gather loads from memory, no lane padding. All arithmetic is fp32 with one
+rounding at the bf16 store; V is copied verbatim.
+"""
+
+import torch
+import triton
+import triton.language as tl
+
+H3_NORM_EPS = 1e-5
+
+
+@triton.jit
+def _qkv_norm_rope_kernel(
+    q_ptr,
+    k_ptr,
+    v_ptr,
+    qw_ptr,
+    kw_ptr,
+    cos_ptr,
+    sin_ptr,
+    out_ptr,
+    num_heads,
+    seq_len,
+    plane,  # elements per output plane = B*S*H*D
+    eps,
+    HD: tl.constexpr,
+    RD: tl.constexpr,
+    HB: tl.constexpr,
+    NHG: tl.constexpr,
+    I64: tl.constexpr,
+):
+    # 1-D grid: gridDim.y is capped at 65535, below the 73k-token shape.
+    pid = tl.program_id(0)
+    h0 = (pid % NHG) * HB
+    bt = pid // NHG
+    tok = bt % seq_len
+    batch = bt // seq_len
+
+    NB: tl.constexpr = HB * HD
+    off = tl.arange(0, NB)
+    hrow = h0 + off // HD
+    mask = hrow < num_heads
+
+    if I64:
+        base = (batch.to(tl.int64) * seq_len + tok) * num_heads * HD + h0 * HD
+    else:
+        base = (batch * seq_len + tok) * num_heads * HD + h0 * HD
+
+    idx = tl.arange(0, HD)
+    pidx = tl.where(idx < RD, tl.where(idx < RD // 2, idx + RD // 2, idx - RD // 2), idx)
+    gidx = tl.broadcast_to(pidx[None, :], (HB, HD))
+    c = tl.load(cos_ptr + tok * RD + idx, mask=idx < RD, other=1.0).to(tl.float32)[None, :]
+    s = tl.load(sin_ptr + tok * RD + idx, mask=idx < RD, other=0.0).to(tl.float32)
+    s = tl.where(idx < RD // 2, -s, s)[None, :]
+
+    wq = tl.load(qw_ptr + idx).to(tl.float32)[None, :]
+    wk = tl.load(kw_ptr + idx).to(tl.float32)[None, :]
+
+    xq = tl.load(q_ptr + base + off, mask=mask, other=0.0).to(tl.float32)
+    xk = tl.load(k_ptr + base + off, mask=mask, other=0.0).to(tl.float32)
+    xv = tl.load(v_ptr + base + off, mask=mask, other=0.0)
+
+    tl.store(out_ptr + 2 * plane + base + off, xv, mask=mask)
+
+    y = tl.reshape(xq, (HB, HD))
+    inv = tl.math.rsqrt(tl.sum(y * y, 1) / HD + eps)[:, None]
+    y = y * inv * wq
+    p = tl.gather(y, gidx, 1)
+    y = y * c + p * s
+    tl.store(out_ptr + base + off, tl.reshape(y, (NB, )).to(out_ptr.dtype.element_ty), mask=mask)
+
+    y = tl.reshape(xk, (HB, HD))
+    inv = tl.math.rsqrt(tl.sum(y * y, 1) / HD + eps)[:, None]
+    y = y * inv * wk
+    p = tl.gather(y, gidx, 1)
+    y = y * c + p * s
+    tl.store(out_ptr + plane + base + off, tl.reshape(y, (NB, )).to(out_ptr.dtype.element_ty), mask=mask)
+
+
+def _pick_config(num_heads: int):
+    """(HB, num_warps) -- consecutive head rows per program.
+
+    One warp per program with 4 rows gives 16 elements (two 128-bit accesses) per thread per
+    tensor, the measured streaming optimum for this GPU: every row stays inside one warp so the
+    RMSNorm reduction and the rotate-half gather are pure warp shuffles, and 32 such CTAs per SM
+    keep enough loads in flight to run at the copy roofline.
+    """
+    if num_heads >= 4:
+        return 4, 1
+    hb = max(triton.next_power_of_2(max(num_heads, 1)), 1)
+    return hb, 1
+
+
+def _launch(
+    query: torch.Tensor,
+    key: torch.Tensor,
+    value: torch.Tensor,
+    q_weight: torch.Tensor,
+    k_weight: torch.Tensor,
+    cos: torch.Tensor,
+    sin: torch.Tensor,
+) -> torch.Tensor:
+    b, s, h, d = query.shape
+    out = torch.empty((3 * b, s, h, d), device=query.device, dtype=query.dtype)
+    hb, warps = _pick_config(h)
+    nhg = triton.cdiv(h, hb)
+    _qkv_norm_rope_kernel[(b * s * nhg, )](
+        query,
+        key,
+        value,
+        q_weight,
+        k_weight,
+        cos,
+        sin,
+        out,
+        h,
+        s,
+        b * s * h * d,
+        H3_NORM_EPS,
+        HD=d,
+        RD=cos.shape[1],
+        HB=hb,
+        NHG=nhg,
+        I64=out.numel() >= 2**31,
+        num_warps=warps,
+        num_stages=1,
+    )
+    return out
+
+
+def qkv_prologue(
+    query: torch.Tensor,
+    key: torch.Tensor,
+    value: torch.Tensor,
+    q_weight: torch.Tensor,
+    k_weight: torch.Tensor,
+    cos: torch.Tensor,
+    sin: torch.Tensor,
+) -> torch.Tensor:
+    """Returns the [3*B, S, H, D] stack (rotated Q, rotated K, V), as FastVideo's attention path builds it."""
+    if query.ndim != 4 or key.shape != query.shape or value.shape != query.shape:
+        raise ValueError("query/key/value must share a [B, S, H, D] shape")
+    if cos.shape != sin.shape or cos.shape[0] != query.shape[1]:
+        raise ValueError("cos/sin must be [S, rotary_dim] tables")
+    if not query.is_contiguous():
+        query = query.contiguous()
+    if not key.is_contiguous():
+        key = key.contiguous()
+    if not value.is_contiguous():
+        value = value.contiguous()
+    if not q_weight.is_contiguous():
+        q_weight = q_weight.contiguous()
+    if not k_weight.is_contiguous():
+        k_weight = k_weight.contiguous()
+    if cos.dtype != query.dtype:
+        cos = cos.to(query.dtype)
+    if sin.dtype != query.dtype:
+        sin = sin.to(query.dtype)
+    if not cos.is_contiguous():
+        cos = cos.contiguous()
+    if not sin.is_contiguous():
+        sin = sin.contiguous()
+    return _launch(query, key, value, q_weight, k_weight, cos, sin)
@@ -0,0 +1,137 @@
+# SPDX-License-Identifier: Apache-2.0
+"""MotionKernel H3 kernel: Fused H3 residual gate + RMSNorm + AdaLN modulation.
+
+Found with MotionKernel (agent-driven kernel search). Tuned on GB200 / sm_100 for the
+MiniMax-H3 inference shapes; opt-in through ``FASTVIDEO_MINIMAX_H3_FUSIONS=mk``.
+
+Chunked gated residual + RMSNorm + AdaLN modulation (H3 post-attention).
+"""
+
+import torch
+import triton
+import triton.language as tl
+
+H3_NORM_EPS = 1e-5
+
+_BLOCK = 1024
+_NUM_WARPS = 4
+_NUM_STAGES = 3
+
+
+@triton.jit
+def _residual_gate_rmsnorm_modulate_kernel(
+    hidden_out_ptr,
+    normed_out_ptr,
+    residual_ptr,
+    branch_ptr,
+    gate_ptr,
+    weight_ptr,
+    scale_ptr,
+    shift_ptr,
+    index_ptr,
+    n_cols,
+    n_index,
+    eps,
+    stride_gate,
+    stride_scale,
+    stride_shift,
+    BLOCK: tl.constexpr,
+    IDENTITY_INDEX: tl.constexpr,
+):
+    row = tl.program_id(0).to(tl.int64)
+    cols = tl.arange(0, BLOCK)
+    base = row * n_cols
+    if IDENTITY_INDEX:
+        table_row = tl.load(index_ptr + row).to(tl.int64)
+    else:
+        table_row = tl.load(index_ptr + (row % n_index)).to(tl.int64)
+    gbase = table_row * stride_gate
+    sbase = table_row * stride_scale
+    hbase = table_row * stride_shift
+
+    acc = tl.zeros((BLOCK, ), tl.float32)
+    for c in range(0, n_cols, BLOCK):
+        cur = c + cols
+        m = cur < n_cols
+        off = base + cur
+        residual = tl.load(residual_ptr + off, mask=m, other=0.0).to(tl.float32)
+        branch = tl.load(branch_ptr + off, mask=m, other=0.0).to(tl.float32)
+        gate = tl.load(gate_ptr + gbase + cur, mask=m, other=0.0).to(tl.float32)
+        hidden = residual + gate * branch
+        tl.store(
+            hidden_out_ptr + off,
+            hidden.to(hidden_out_ptr.dtype.element_ty),
+            mask=m,
+        )
+        acc += hidden * hidden
+
+    rstd = tl.math.rsqrt(tl.sum(acc, axis=0) / n_cols + eps)
+
+    for c in range(0, n_cols, BLOCK):
+        cur = c + cols
+        m = cur < n_cols
+        off = base + cur
+        hidden = tl.load(hidden_out_ptr + off, mask=m, other=0.0, cache_modifier=".cg").to(tl.float32)
+        weight = tl.load(weight_ptr + cur, mask=m, other=0.0).to(tl.float32)
+        scale = tl.load(scale_ptr + sbase + cur, mask=m, other=0.0).to(tl.float32)
+        shift = tl.load(shift_ptr + hbase + cur, mask=m, other=0.0).to(tl.float32)
+        output = hidden * rstd * weight * (1.0 + scale) + shift
+        tl.store(
+            normed_out_ptr + off,
+            output.to(normed_out_ptr.dtype.element_ty),
+            mask=m,
+        )
+
+
+def _row_addressable(table: torch.Tensor) -> torch.Tensor:
+    return table if table.stride(-1) == 1 else table.contiguous()
+
+
+def residual_gate_rmsnorm_modulate(
+    residual: torch.Tensor,
+    branch: torch.Tensor,
+    gate: torch.Tensor,
+    weight: torch.Tensor,
+    scale: torch.Tensor,
+    shift: torch.Tensor,
+    index: torch.Tensor,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    """Return ``(hidden, normed)`` exactly as FastVideo's optimized profile does."""
+    if not residual.is_cuda:
+        raise ValueError("h3_residual_gate_rmsnorm_modulate requires CUDA tensors")
+    if branch.shape != residual.shape or branch.dtype != residual.dtype:
+        raise ValueError("branch must match residual shape and dtype")
+    hidden_size = residual.shape[-1]
+    if index.ndim != 1 or index.numel() != residual.shape[-2]:
+        raise ValueError("index must hold one table row per token")
+    flat_residual = residual.reshape(-1, hidden_size).contiguous()
+    flat_branch = branch.reshape(-1, hidden_size).contiguous()
+    weight = weight.contiguous()
+    gate = _row_addressable(gate)
+    scale = _row_addressable(scale)
+    shift = _row_addressable(shift)
+    index = index.contiguous()
+    hidden = torch.empty_like(flat_residual)
+    modulated = torch.empty_like(flat_residual)
+    _residual_gate_rmsnorm_modulate_kernel[(flat_residual.shape[0], )](
+        hidden,
+        modulated,
+        flat_residual,
+        flat_branch,
+        gate,
+        weight,
+        scale,
+        shift,
+        index,
+        hidden_size,
+        index.numel(),
+        H3_NORM_EPS,
+        gate.stride(0),
+        scale.stride(0),
+        shift.stride(0),
+        IDENTITY_INDEX=index.numel() == flat_residual.shape[0],
+        BLOCK=_BLOCK,
+        num_warps=_NUM_WARPS,
+        num_stages=_NUM_STAGES,
+    )
+    return hidden.view_as(residual), modulated.view_as(residual)
@@ -0,0 +1,238 @@
+# SPDX-License-Identifier: Apache-2.0
+"""MotionKernel H3 kernel: Fused H3 RMSNorm + row-indexed AdaLN modulation.
+
+Found with MotionKernel (agent-driven kernel search). Tuned on GB200 / sm_100 for the
+MiniMax-H3 inference shapes; opt-in through ``FASTVIDEO_MINIMAX_H3_FUSIONS=mk``.
+
+MiniMax H3 RMSNorm + row-indexed AdaLN modulation (two-kernel fast path).
+
+Math (FastVideo ``minimax_h3.py`` lines 729-731), FP32 internally with one
+rounding at the store::
+
+    y = rms_norm(x, weight, eps) * (1 + scale[idx]) + shift[idx]
+
+Shape notes that drive the design (GB200, hidden D = 5376 = 2**8 * 21, so D is
+*not* a power of two):
+
+* The vendored Sol-Engine starter runs one program per row with ``BLOCK =
+  next_pow2(D) = 8192`` and ``num_warps=16``; 34% of its lanes are masked off
+  and the whole 8192-lane tile is live in registers.
+* The AdaLN tables are row-strided views (row stride ``6 * D``) with only
+  ``table_rows`` (<= 12) distinct rows, so every token re-reads the same
+  ``weight``, ``scale`` and ``shift`` vectors: 3 * D * 2 bytes of gather per
+  row on top of the D * 2 bytes of activation.
+
+This candidate keeps the public interface and splits the work:
+
+1. ``_fuse_adaln_kernel`` (tiny: ``table_rows`` rows) folds the per-table-row
+   modulation into ``a[t] = weight * (1 + scale[t])`` and ``b[t] = shift[t]``,
+   stored in the activation dtype. The main pass then gathers two vectors per
+   token instead of three, and they are read at the activation width.
+2. ``_rmsnorm_modulate_kernel`` is a chunked single-row program with no lane
+   padding: ``ceil(D / BLOCK)`` chunks of ``BLOCK`` lanes (BLOCK = 2048 for
+   D = 5376, i.e. three chunks with only the last one partially masked). It
+   makes a first pass over the row to accumulate the sum of squares, then a
+   second pass that re-reads the row (an L1/L2 hit, the row is 10.5 KiB) and
+   applies the store. That keeps the live register tile at one chunk instead
+   of a whole padded row, which is what the starter loses to.
+
+Performance notes (GB200, measured against a minimal Triton streaming copy of
+the same byte count on the same shapes, i.e. the true roofline):
+
+* copy roofline per production shape: 58.9 / 98.6 / 129.1 / 140.8 / 183.9 us
+  (6.2-6.9 TB/s), mean 122.3 us.
+* this kernel: mean of five shapes ~126.3 us, i.e. ~94.5% of the copy
+  roofline re-measured for this attempt (55.6 / 95.1 / 124.7 / 138.0 / 184.6 us,
+  mean 119.6 us); on the 59k shape 188.2 us vs 184.6 us (98%).
+
+The step that mattered most over the starter (134.8 us -> 128.5 us) was making
+the hidden width a ``tl.constexpr`` and unrolling the chunk loops with
+``tl.static_range``: the lane masks, the row offset and the reduction become
+static, so every activation/table/store access compiles to an aligned 16-byte
+vector access instead of scalar traffic on an unknown-alignment base pointer.
+"""
+
+import torch
+import triton
+import triton.language as tl
+
+H3_NORM_EPS = 1e-5
+
+_MAIN_BLOCK = 2048
+_TABLE_BLOCK = 1024
+
+
+@triton.jit
+def _fuse_adaln_kernel(
+    a_ptr,
+    b_ptr,
+    weight_ptr,
+    scale_ptr,
+    shift_ptr,
+    stride_scale_row,
+    stride_shift_row,
+    NC: tl.constexpr,
+    BLOCK: tl.constexpr,
+):
+    """One program per (table row, column chunk): a = w * (1 + scale), b = shift."""
+    row = tl.program_id(0).to(tl.int64)
+    cols = tl.program_id(1) * BLOCK + tl.arange(0, BLOCK)
+    mask = cols < NC
+    w = tl.load(weight_ptr + cols, mask=mask, other=0.0).to(tl.float32)
+    s = tl.load(scale_ptr + row * stride_scale_row + cols, mask=mask, other=0.0).to(tl.float32)
+    h = tl.load(shift_ptr + row * stride_shift_row + cols, mask=mask, other=0.0).to(tl.float32)
+    out = row * NC + cols
+    tl.store(a_ptr + out, (w * (1.0 + s)).to(a_ptr.dtype.element_ty), mask=mask)
+    tl.store(b_ptr + out, h.to(b_ptr.dtype.element_ty), mask=mask)
+
+
+@triton.jit
+def _rmsnorm_modulate_kernel(
+    out_ptr,
+    x_ptr,
+    a_ptr,
+    b_ptr,
+    index_ptr,
+    n_index,
+    eps,
+    stride_a_row,
+    stride_b_row,
+    NC: tl.constexpr,
+    BLOCK: tl.constexpr,
+    NCHUNK: tl.constexpr,
+):
+    """One program per row: sumsq pass over the row, then an apply pass.
+
+    ``NC`` (the hidden width) is a ``constexpr`` so that lane masks, row
+    offsets and the reduction are fully static: every activation, table and
+    store access is then a 16-byte vector access on an aligned base, which is
+    what lifts this pass to the streaming-copy roofline. The row is 10.5 KiB,
+    so the second pass re-reads it from L1 instead of DRAM.
+    """
+    row = tl.program_id(0).to(tl.int64)
+    cols = tl.arange(0, BLOCK)
+    table_row = tl.load(index_ptr + row % n_index).to(tl.int64)
+    base = row * NC
+
+    acc = tl.zeros([BLOCK], dtype=tl.float32)
+    for chunk in tl.static_range(NCHUNK):
+        off = chunk * BLOCK + cols
+        x = tl.load(x_ptr + base + off, mask=off < NC, other=0.0, eviction_policy="evict_last").to(tl.float32)
+        acc += x * x
+    rstd = 1.0 / tl.sqrt(tl.sum(acc, axis=0) / NC + eps)
+
+    chunk = NCHUNK - 1
+    off = chunk * BLOCK + cols
+    mask = off < NC
+    x = tl.load(x_ptr + base + off, mask=mask, other=0.0, eviction_policy="evict_last").to(tl.float32)
+    a = tl.load(
+        a_ptr + table_row * stride_a_row + off,
+        mask=mask,
+        other=0.0,
+        eviction_policy="evict_last",
+    ).to(tl.float32)
+    b = tl.load(
+        b_ptr + table_row * stride_b_row + off,
+        mask=mask,
+        other=0.0,
+        eviction_policy="evict_last",
+    ).to(tl.float32)
+    y = x * rstd * a + b
+    tl.store(
+        out_ptr + base + off,
+        y.to(out_ptr.dtype.element_ty),
+        mask=mask,
+        cache_modifier=".cs",
+    )
+
+    for chunk in tl.static_range(NCHUNK - 1):
+        off = chunk * BLOCK + cols
+        mask = off < NC
+        x = tl.load(x_ptr + base + off, mask=mask, other=0.0).to(tl.float32)
+        a = tl.load(a_ptr + table_row * stride_a_row + off, mask=mask, other=0.0).to(tl.float32)
+        b = tl.load(b_ptr + table_row * stride_b_row + off, mask=mask, other=0.0).to(tl.float32)
+        y = x * rstd * a + b
+        tl.store(
+            out_ptr + base + off,
+            y.to(out_ptr.dtype.element_ty),
+            mask=mask,
+            cache_modifier=".cs",
+        )
+
+
+def _next_power_of_two(value: int) -> int:
+    return 1 << (value - 1).bit_length()
+
+
+def _main_block(hidden: int) -> int:
+    return min(_MAIN_BLOCK, _next_power_of_two(hidden))
+
+
+def _num_warps(block: int) -> int:
+    if block >= 512:
+        return 4
+    return 2
+
+
+def _row_addressable(table: torch.Tensor) -> torch.Tensor:
+    return table if table.stride(-1) == 1 else table.contiguous()
+
+
+def rmsnorm_modulate(
+    x: torch.Tensor,
+    weight: torch.Tensor,
+    scale: torch.Tensor,
+    shift: torch.Tensor,
+    index: torch.Tensor,
+) -> torch.Tensor:
+    """RMSNorm + per-token AdaLN modulation with one activation read and write."""
+    if not x.is_cuda:
+        raise ValueError("h3_rmsnorm_modulate requires CUDA tensors")
+    hidden = x.shape[-1]
+    if index.ndim != 1 or index.numel() != x.shape[-2]:
+        raise ValueError("index must hold one table row per token")
+    if scale.shape[-1] != hidden or shift.shape[-1] != hidden:
+        raise ValueError("modulation tables must match the hidden width")
+    flat_x = x.reshape(-1, hidden).contiguous()
+    weight = weight.contiguous()
+    scale = _row_addressable(scale)
+    shift = _row_addressable(shift)
+    index = index.contiguous()
+    output = torch.empty_like(flat_x)
+
+    table_rows = scale.shape[0]
+    tab_dtype = flat_x.dtype if flat_x.dtype in (torch.bfloat16, torch.float16) else torch.float32
+    fused_a = torch.empty((table_rows, hidden), dtype=tab_dtype, device=flat_x.device)
+    fused_b = torch.empty((table_rows, hidden), dtype=tab_dtype, device=flat_x.device)
+
+    table_block = min(_TABLE_BLOCK, _next_power_of_two(hidden))
+    _fuse_adaln_kernel[(table_rows, triton.cdiv(hidden, table_block))](
+        fused_a,
+        fused_b,
+        weight,
+        scale,
+        shift,
+        scale.stride(0),
+        shift.stride(0),
+        NC=hidden,
+        BLOCK=table_block,
+        num_warps=_num_warps(table_block),
+    )
+
+    block = _main_block(hidden)
+    _rmsnorm_modulate_kernel[(flat_x.shape[0], )](
+        output,
+        flat_x,
+        fused_a,
+        fused_b,
+        index,
+        index.numel(),
+        H3_NORM_EPS,
+        fused_a.stride(0),
+        fused_b.stride(0),
+        NC=hidden,
+        BLOCK=block,
+        NCHUNK=triton.cdiv(hidden, block),
+        num_warps=_num_warps(block),
+    )
+    return output.view_as(x)
@@ -0,0 +1,123 @@
+# SPDX-License-Identifier: Apache-2.0
+"""MotionKernel H3 kernel: Fused H3 SwiGLU over the packed gate/value halves.
+
+Found with MotionKernel (agent-driven kernel search). Tuned on GB200 / sm_100 for the
+MiniMax-H3 inference shapes; opt-in through ``FASTVIDEO_MINIMAX_H3_FUSIONS=mk``.
+
+MiniMax H3 value-first packed SwiGLU -- chunked, unpadded, streaming.
+
+Math (FastVideo ``minimax_h3.py`` @ 7203b4976, lines 188-189)::
+
+    value, gate = hidden_states.chunk(2, dim=-1)   # value first, gate second
+    out = value * F.silu(gate)
+
+Kept in FP32 throughout with one rounding at the store (matches Sol-Engine).
+
+Why this beats the vendored Sol-Engine kernel: that kernel runs one program per
+row with ``BLOCK = next_pow2(ffn) = 16384`` lanes, so 12.5% of every lane group
+is idle and each program keeps two 16384-wide FP32 tiles live (heavy register
+pressure, low occupancy). F=14336 = 14 x 1024, so a chunked grid needs no
+padding at all. Each program handles one 1024-element column chunk of one row:
+one 128-bit vector load per thread per half, FP32 math, one store.
+
+Launch-shape notes measured on GB200 (Pika shapes):
+  * A flat 1D grid whose column chunk is ``pid % NTILES`` (so concurrently
+    resident CTAs stream adjacent 2 KiB blocks of the same row) is the fastest
+    enumeration: ~443 us at 37718 tokens = 7.33 TB/s, versus a measured
+    ~7.4 TB/s pure-read / 7.8 TB/s pure-write streaming roofline. Enumerating
+    rows fastest costs ~8% (478 us); BLOCK=512 or BLOCK=2048 with more warps
+    cost 2-28%.
+  * The flat grid also removes the ``grid.y <= 65535`` limit that the 768p
+    243-frame shape (73394 rows) exceeds: adding a row-group loop there cost
+    4% (888 us vs 852 us).
+"""
+
+import torch
+import triton
+import triton.language as tl
+
+_BLOCK = 1024
+_NUM_WARPS = 4
+_I32_LIMIT = 0x7FFFFFFF
+
+
+@triton.jit
+def _h3_swiglu_kernel(
+    out_ptr,
+    x_ptr,
+    ffn_dim,
+    n_rows,
+    stride_in_row,
+    stride_out_row,
+    NTILES: tl.constexpr,
+    BLOCK: tl.constexpr,
+    EXACT: tl.constexpr,
+    GUARD: tl.constexpr,
+    I64: tl.constexpr,
+):
+    pid = tl.program_id(0)
+    col = pid % NTILES
+    row = pid // NTILES
+    cols = col * BLOCK + tl.arange(0, BLOCK)
+    if I64:
+        row = row.to(tl.int64)
+    if GUARD:
+        if row < n_rows:
+            _body(out_ptr, x_ptr, ffn_dim, row, cols, stride_in_row, stride_out_row, BLOCK, EXACT)
+    else:
+        _body(out_ptr, x_ptr, ffn_dim, row, cols, stride_in_row, stride_out_row, BLOCK, EXACT)
+
+
+@triton.jit
+def _body(out_ptr, x_ptr, ffn_dim, row, cols, stride_in_row, stride_out_row, BLOCK: tl.constexpr, EXACT: tl.constexpr):
+    in_off = row * stride_in_row + cols
+    if EXACT:
+        value = tl.load(x_ptr + in_off).to(tl.float32)
+        gate = tl.load(x_ptr + in_off + ffn_dim).to(tl.float32)
+        out = value * (gate * tl.sigmoid(gate))
+        tl.store(out_ptr + row * stride_out_row + cols, out.to(out_ptr.dtype.element_ty))
+    else:
+        mask = cols < ffn_dim
+        value = tl.load(x_ptr + in_off, mask=mask, other=0.0).to(tl.float32)
+        gate = tl.load(x_ptr + in_off + ffn_dim, mask=mask, other=0.0).to(tl.float32)
+        out = value * (gate * tl.sigmoid(gate))
+        tl.store(out_ptr + row * stride_out_row + cols, out.to(out_ptr.dtype.element_ty), mask=mask)
+
+
+def swiglu(x: torch.Tensor) -> torch.Tensor:
+    """Launch the chunked unpadded SwiGLU."""
+    if not x.is_cuda:
+        raise ValueError("h3_swiglu requires a CUDA tensor")
+    packed_width = x.shape[-1]
+    if packed_width == 0 or packed_width % 2:
+        raise ValueError("h3_swiglu needs a positive even last dimension")
+    ffn_dim = packed_width // 2
+    flat = x.reshape(-1, packed_width)
+    if not flat.is_contiguous():
+        flat = flat.contiguous()
+    output_shape = (*x.shape[:-1], ffn_dim)
+    rows = flat.shape[0]
+    if rows == 0:
+        return torch.empty(output_shape, dtype=x.dtype, device=x.device)
+    out = torch.empty((rows, ffn_dim), dtype=x.dtype, device=x.device)
+
+    block = min(_BLOCK, triton.next_power_of_2(ffn_dim))
+    exact = (ffn_dim % block) == 0
+    ntiles = (ffn_dim + block - 1) // block
+    i64 = rows * packed_width > _I32_LIMIT
+    num_warps = _NUM_WARPS if block >= 1024 else 2
+    _h3_swiglu_kernel[(ntiles * rows, )](
+        out,
+        flat,
+        ffn_dim,
+        rows,
+        flat.stride(0),
+        out.stride(0),
+        NTILES=ntiles,
+        BLOCK=block,
+        EXACT=exact,
+        GUARD=False,
+        I64=i64,
+        num_warps=num_warps,
+    )
+    return out.view(output_shape)
@@ -0,0 +1,40 @@
+# SPDX-License-Identifier: Apache-2.0
+"""``DistributedAttention.forward`` is the stacked path plus a ``torch.cat``."""
+import torch
+
+from fastvideo.attention.layer import DistributedAttention
+
+
+class _Recorder:
+
+    def __init__(self):
+        self.qkv = None
+        self.kwargs = None
+
+
+def test_forward_cats_then_delegates_to_forward_stacked():
+    rec = _Recorder()
+
+    def fake_forward_stacked(self, qkv, **kwargs):
+        rec.qkv = qkv
+        rec.kwargs = kwargs
+        return torch.zeros(2, 3, 4, 5), None
+
+    original = DistributedAttention.forward_stacked
+    DistributedAttention.forward_stacked = fake_forward_stacked
+    try:
+        q = torch.arange(2 * 3 * 4 * 5, dtype=torch.float32).reshape(2, 3, 4, 5)
+        k = q + 1
+        v = q + 2
+        # Bypass __init__ (it builds a real backend); ``forward`` only needs the
+        # compiler-policy flag and the patched ``forward_stacked``.
+        dummy = object.__new__(DistributedAttention)
+        dummy._compile_forward_enabled = False
+        out, replicated = DistributedAttention.forward(dummy, q, k, v, original_seq_len=3)
+    finally:
+        DistributedAttention.forward_stacked = original
+
+    assert replicated is None
+    assert rec.qkv.shape == (6, 3, 4, 5)
+    assert torch.equal(rec.qkv[:2], q) and torch.equal(rec.qkv[2:4], k) and torch.equal(rec.qkv[4:], v)
+    assert rec.kwargs["original_seq_len"] == 3
@@ -0,0 +1,35 @@
+# SPDX-License-Identifier: Apache-2.0
+"""The MotionKernel ``mk`` tier is opt-in and never changes the default fusion set."""
+import pytest
+
+from fastvideo.models.dits.minimax_h3 import (_MINIMAX_H3_FUSION_NAMES, _MINIMAX_H3_MK_MODE,
+                                              _enabled_minimax_h3_fusions)
+
+
+def test_default_is_unchanged():
+    for value in ("", "0", "none", None):
+        assert _enabled_minimax_h3_fusions(value) == frozenset()
+
+
+def test_all_is_unchanged():
+    assert _enabled_minimax_h3_fusions("all") == _MINIMAX_H3_FUSION_NAMES
+    assert _enabled_minimax_h3_fusions("1") == _MINIMAX_H3_FUSION_NAMES
+
+
+def test_subset_is_unchanged():
+    assert _enabled_minimax_h3_fusions("modulate,swiglu") == frozenset({"modulate", "swiglu"})
+
+
+def test_mk_enables_every_fusion_plus_the_tier_marker():
+    assert _enabled_minimax_h3_fusions("mk") == _MINIMAX_H3_FUSION_NAMES | {_MINIMAX_H3_MK_MODE}
+    assert _enabled_minimax_h3_fusions(" MK ") == _MINIMAX_H3_FUSION_NAMES | {_MINIMAX_H3_MK_MODE}
+
+
+def test_mk_cannot_be_combined_with_individual_fusions():
+    with pytest.raises(ValueError, match="whole-tier mode"):
+        _enabled_minimax_h3_fusions("mk,swiglu")
+
+
+def test_unknown_fusion_still_raises():
+    with pytest.raises(ValueError, match="Unknown MiniMax H3 fusion"):
+        _enabled_minimax_h3_fusions("modulate,nope")
@@ -0,0 +1,120 @@
+# SPDX-License-Identifier: Apache-2.0
+"""A real H3 transformer block runs the MotionKernel tier and matches the Sol tier.
+
+Small random-initialised block, SDPA backend, no checkpoint: this exercises the
+``mk_tier`` dispatch of every fused call site (including the stacked QKV
+prologue) on the actual module wiring rather than on standalone kernels.
+"""
+import pytest
+import torch
+
+from fastvideo.forward_context import set_forward_context
+from fastvideo.models.dits.minimax_h3 import MiniMaxH3TransformerBlock
+from fastvideo.models.dits.minimax_h3_fusions import mk
+from fastvideo.platforms import AttentionBackendEnum
+
+pytestmark = pytest.mark.skipif(not torch.cuda.is_available() or not mk.HAVE_TRITON,
+                                reason="MotionKernel H3 kernels need CUDA + Triton")
+
+HIDDEN, HEADS, HEAD_DIM, ROTARY_DIM, FFN, TABLE_ROWS, TOKENS = 512, 4, 128, 96, 1024, 6, 256
+
+
+def _block(mk_tier: bool, fused: bool = True) -> MiniMaxH3TransformerBlock:
+    torch.manual_seed(0)
+    torch.set_default_dtype(torch.bfloat16)
+    try:
+        return MiniMaxH3TransformerBlock(
+            hidden_size=HIDDEN,
+            num_attention_heads=HEADS,
+            attention_head_dim=HEAD_DIM,
+            ffn_dim=FFN,
+            time_embed_dim=128,
+            norm_eps=1e-5,
+            qk_norm_eps=1e-5,
+            supported_attention_backends=(AttentionBackendEnum.TORCH_SDPA, ),
+            quant_config=None,
+            prefix="test",
+            fuse_modulate=fused,
+            fuse_qknorm_rope=fused,
+            fuse_swiglu=fused,
+            mk_tier=mk_tier,
+        ).cuda()
+    finally:
+        torch.set_default_dtype(torch.float32)
+
+
+def _inputs():
+    torch.manual_seed(1)
+    return (torch.randn(1, TOKENS, HIDDEN, device="cuda",
+                        dtype=torch.bfloat16), torch.randn(TABLE_ROWS, 128, device="cuda", dtype=torch.bfloat16),
+            torch.randint(0, TABLE_ROWS, (TOKENS, ),
+                          device="cuda", dtype=torch.int64), (torch.randn(TOKENS,
+                                                                          ROTARY_DIM,
+                                                                          device="cuda",
+                                                                          dtype=torch.bfloat16),
+                                                              torch.randn(TOKENS,
+                                                                          ROTARY_DIM,
+                                                                          device="cuda",
+                                                                          dtype=torch.bfloat16)), TOKENS)
+
+
+def test_mk_tier_dispatches_and_matches_sol(monkeypatch, distributed_setup):
+    calls = {
+        name: 0
+        for name in ("qkv_prologue", "rmsnorm_modulate", "residual_gate_rmsnorm_modulate", "swiglu", "gated_residual")
+    }
+    for name in calls:
+        target = getattr(mk, name)
+
+        def counted(*args, _target=target, _name=name, **kwargs):
+            calls[_name] += 1
+            return _target(*args, **kwargs)
+
+        monkeypatch.setattr(mk, name, counted)
+
+    sol, mk_block, eager = _block(False), _block(True), _block(False, fused=False)
+    # The loader zero-initialises ReplicatedLinear (weights arrive from the
+    # checkpoint) and that would make the block an identity, hiding every kernel
+    # difference. Give all three copies the same non-trivial weights instead.
+    torch.manual_seed(2)
+    for name, param in sol.named_parameters():
+        if name.endswith(".weight") and param.dim() == 2 or name.endswith(".bias"):
+            param.data.normal_(0.0, 0.02)
+        elif name.endswith(".weight"):
+            param.data.fill_(1.0)
+    mk_block.load_state_dict(sol.state_dict())
+    eager.load_state_dict(sol.state_dict())
+    inputs = _inputs()
+    with torch.inference_mode(), set_forward_context(current_timestep=0, attn_metadata=None):
+        expected = sol(*inputs)
+        actual = mk_block(*inputs)
+        unfused = eager(*inputs)
+
+    assert all(count > 0 for count in calls.values()), calls
+    base = torch.linalg.vector_norm(expected.float())
+
+    def rel(other):
+        return (torch.linalg.vector_norm(other.float() - expected.float()) / base).item()
+
+    floor, rel = rel(unfused), rel(actual)
+    print(f"eager-vs-sol block rel-L2 = {floor:.3e}; mk-vs-sol block rel-L2 = {rel:.3e}")
+    assert rel <= 3 * floor, (rel, floor)
+
+
+def test_mk_tier_defers_to_sol_while_compiling(monkeypatch, distributed_setup):
+    """A traced block must keep using the compile-safe Sol custom ops."""
+    calls = []
+    for name in ("qkv_prologue", "rmsnorm_modulate", "residual_gate_rmsnorm_modulate", "swiglu", "gated_residual"):
+        target = getattr(mk, name)
+
+        def counted(*args, _target=target, _name=name, **kwargs):
+            calls.append(_name)
+            return _target(*args, **kwargs)
+
+        monkeypatch.setattr(mk, name, counted)
+    monkeypatch.setattr(torch.compiler, "is_compiling", lambda: True)
+
+    block = _block(True)
+    with torch.inference_mode(), set_forward_context(current_timestep=0, attn_metadata=None):
+        block(*_inputs())
+    assert calls == []
@@ -0,0 +1,94 @@
+# SPDX-License-Identifier: Apache-2.0
+"""MotionKernel H3 kernels match the in-tree Sol/eager paths on the production shapes.
+
+GPU-only (Triton on CUDA); the shapes use H3's real hidden width, head layout and
+AdaLN table stride, only the token count is reduced to keep the test quick.
+"""
+import pytest
+import torch
+
+from fastvideo.models.dits.minimax_h3_fusions import (fused_qknorm_rope, fused_residual_gate_rmsnorm_modulate,
+                                                      fused_rmsnorm_modulate, minimax_h3_swiglu)
+from fastvideo.models.dits.minimax_h3_fusions import mk
+
+pytestmark = pytest.mark.skipif(not torch.cuda.is_available() or not mk.HAVE_TRITON,
+                                reason="MotionKernel H3 kernels need CUDA + Triton")
+
+HIDDEN, HEADS, HEAD_DIM, ROTARY_DIM, FFN_DIM, EPS = 5376, 56, 128, 96, 14336, 1e-5
+TOKENS = 512
+
+
+def _rel_l2(actual, expected):
+    return (torch.linalg.vector_norm(actual.float() - expected.float()) /
+            torch.linalg.vector_norm(expected.float())).item()
+
+
+@pytest.fixture(scope="module")
+def production_inputs():
+    torch.manual_seed(0)
+    device = "cuda"
+    mod = torch.randn(6, 6, HIDDEN, device=device, dtype=torch.bfloat16)
+    # The production AdaLN tables are row-strided views of this [rows, 6*D] buffer.
+    shift_msa, scale_msa, gate_msa, shift_mlp, scale_mlp, gate_mlp = (mod[:, i, :] for i in range(6))
+    return {
+        "x": torch.randn(TOKENS, HIDDEN, device=device, dtype=torch.bfloat16),
+        "branch": torch.randn(TOKENS, HIDDEN, device=device, dtype=torch.bfloat16),
+        "weight": torch.randn(HIDDEN, device=device, dtype=torch.bfloat16),
+        "qk_weight": torch.randn(HEAD_DIM, device=device, dtype=torch.bfloat16),
+        "scale_msa": scale_msa,
+        "shift_msa": shift_msa,
+        "gate_msa": gate_msa,
+        "scale_mlp": scale_mlp,
+        "shift_mlp": shift_mlp,
+        "gate_mlp": gate_mlp,
+        "index": torch.randint(0, 6, (TOKENS, ), device=device, dtype=torch.int64),
+    }
+
+
+def test_rmsnorm_modulate_matches_sol(production_inputs):
+    i = production_inputs
+    expected = fused_rmsnorm_modulate(i["x"], i["weight"], i["scale_msa"], i["shift_msa"], i["index"], EPS)
+    actual = mk.rmsnorm_modulate(i["x"], i["weight"], i["scale_msa"], i["shift_msa"], i["index"])
+    assert actual.shape == expected.shape and actual.dtype == expected.dtype
+    assert _rel_l2(actual, expected) < 5e-3
+
+
+def test_residual_gate_rmsnorm_modulate_matches_sol(production_inputs):
+    i = production_inputs
+    args = (i["branch"], i["gate_msa"], i["weight"], i["scale_mlp"], i["shift_mlp"], i["index"])
+    exp_hidden, exp_normed = fused_residual_gate_rmsnorm_modulate(i["x"], *args, EPS)
+    act_hidden, act_normed = mk.residual_gate_rmsnorm_modulate(i["x"], *args)
+    assert _rel_l2(act_hidden, exp_hidden) < 1e-3
+    assert _rel_l2(act_normed, exp_normed) < 5e-3
+
+
+def test_swiglu_matches_sol(production_inputs):
+    packed = torch.randn(TOKENS, 2 * FFN_DIM, device="cuda", dtype=torch.bfloat16)
+    assert _rel_l2(mk.swiglu(packed), minimax_h3_swiglu(packed)) < 1e-3
+
+
+def test_gated_residual_matches_eager(production_inputs):
+    i = production_inputs
+    expected = i["x"] + i["gate_mlp"].index_select(0, i["index"]) * i["branch"]
+    assert _rel_l2(mk.gated_residual(i["x"], i["gate_mlp"], i["branch"], i["index"]), expected) < 5e-3
+
+
+def test_qkv_prologue_matches_two_sol_launches(production_inputs):
+    i = production_inputs
+    q = torch.randn(1, TOKENS, HEADS, HEAD_DIM, device="cuda", dtype=torch.bfloat16)
+    k = torch.randn_like(q)
+    v = torch.randn_like(q)
+    cos = torch.randn(TOKENS, ROTARY_DIM, device="cuda", dtype=torch.bfloat16)
+    sin = torch.randn(TOKENS, ROTARY_DIM, device="cuda", dtype=torch.bfloat16)
+    expected = torch.cat(
+        [fused_qknorm_rope(q, i["qk_weight"], cos, sin, EPS),
+         fused_qknorm_rope(k, i["qk_weight"], cos, sin, EPS), v],
+        dim=0)
+    actual = mk.qkv_prologue(q, k, v, i["qk_weight"], i["qk_weight"], cos, sin)
+    assert actual.shape == (3, TOKENS, HEADS, HEAD_DIM)
+    assert _rel_l2(actual, expected) < 1e-4
+
+
+def test_kernels_reject_cpu_tensors():
+    with pytest.raises(ValueError):
+        mk.swiglu(torch.randn(4, 8))
Author	SHA1	Message	Date
Aryan Kumar	bfc4577034	[misc]: H3 MK tier: drop narration comments, move tests under the collected transformers lane	2026-10-30 13:26:54 -07:00
Aryan Kumar	5e1e29196a	[bugfix]: H3 MK kernels: int64 QKV batch offset, broadcast gated-residual index over batch	2026-10-30 13:23:26 -07:00
Aryan Kumar	cb6d12b55a	[perf] MiniMax H3: opt-in faster kernel tier for the DiT fused ops	2026-10-30 13:04:59 -07:00
				`@@ -0,0 +1 @@`
				`/Users/aryank/os/FastVideo-review-int/docs/node_modules`