d3 v3.0.4

MODEL_MANIFEST.json

@@ -564,7 +564,9 @@
  },
  "runtime": {
   "files_sha256": {
-   "d3_runtime.py": "a4b577dedcc946ad749589ad47e8314a3db7490245c89f896109a32ec4a54f53",
+   "d3_runtime.py": "88fcaf3451edcc2139ffcb0db53ea4f9d03152e2736e0ffbb47386b40f618d52",
+   "d3_fast.py": "c8f6e74968e2835732e0865814b98a21e4b604c80ecb0416c270219eb8e64c05",
+   "d3_kernels.py": "8f7534ca1e1846ebf7c4c73058002f3be4dbf0d5d426f0815150ae0cbe7677bd",
    "modeling_d3.py": "55314f36206484a505e9525db373621d672df7447ddc07b71329a1f2857e8227",
    "pipeline_d3.py": "c0975a91316bcda9c09c5f14bf634ceec94e6980ed7e76a29b718fb451fe4091",
    "d3_server.py": "434d6aa5e8572b649e1068588bf1811ff970412e409684b80832bd5544080bca",
@@ -586,11 +588,11 @@
    }
   ]
  },
- "built_utc": "2026-10-10T20:41:39+00:00",
+ "built_utc": "2026-10-10T21:44:47+00:00",
  "files_sha256": {
   "LICENSE": "c71d239df91726fc519c6eb72d318ec65820627232b2f796219e87dcf35d0ab4",
   "NOTICE": "61f4ea488e1fed0c3fce97171506a23152c2591116442f78f15827fb9e6af469",
-  "README.md": "089caef8180dfbc3a7725e1e37331f4d71509557d87e5c36cec050b8f56dce1b",
+  "README.md": "8cde7822a0840c7c641df439f0f31d64e94216f2bd5df15d6b55eb0ba6aec264",
   "assets/banner.png": "eb7ecc949232e1de544247899dd55ee629a36cfe276042d5795f65178b9c64c5",
   "assets/example-receipt.png": "b9dbb8b103bacd45f1de824d844a3c07b81d57cb1a7466712ef060ad7dcb3926",
   "assets/index-areas.png": "49ecc07e5c271a69a42572240cb48afb648b8f9b99cd6bf934c9bf7ff79e4a77",
@@ -598,8 +600,10 @@
   "chat_template.jinja": "c3cf9e34abf4f9e36c2d72165aa9c132d3e2a725b6c2586aaa3a8af9d7a81041",
   "config.json": "7e16284fafd2d54b73c10073c0391bfd6804886ceaf2cbe37fdb62d4c6813dea",
   "d3_engine.py": "ab638d785325d944d3dac891f77448dfc91cad1b8d55a7302bd527c5f1e0e251",
+  "d3_fast.py": "c8f6e74968e2835732e0865814b98a21e4b604c80ecb0416c270219eb8e64c05",
   "d3_format.py": "e5036154d2e54793f59b320c8726632957bc343c382bac26d825b263a6e9ce62",
-  "d3_runtime.py": "a4b577dedcc946ad749589ad47e8314a3db7490245c89f896109a32ec4a54f53",
+  "d3_kernels.py": "8f7534ca1e1846ebf7c4c73058002f3be4dbf0d5d426f0815150ae0cbe7677bd",
+  "d3_runtime.py": "88fcaf3451edcc2139ffcb0db53ea4f9d03152e2736e0ffbb47386b40f618d52",
   "d3_server.py": "434d6aa5e8572b649e1068588bf1811ff970412e409684b80832bd5544080bca",
   "decision_config.json": "6b80ca11bd6ba481df4b3c187db8a5d1983786563ca05e2edccb0bdac0d2fc4f",
   "merges.txt": "a9d356d7bdf1ef4949e3e748e95b8e10ad9d4e2e838eddc38a0a7b6b94d1db8d",
@@ -636,8 +640,10 @@
   "chat_template.jinja": 8952,
   "config.json": 3949,
   "d3_engine.py": 3726,
+  "d3_fast.py": 21284,
   "d3_format.py": 4882,
-  "d3_runtime.py": 49044,
+  "d3_kernels.py": 16129,
+  "d3_runtime.py": 57795,
   "d3_server.py": 8295,
   "decision_config.json": 5503,
   "merges.txt": 3353259,

README.md

@@ -28,10 +28,10 @@ tags:
 
 ## Highlights
 
-- **Jev Decision Index 0.3, public suite: 65.15**, measured with the official 0.3 kit on the released weights: all 140,178 public requests answered, none unsupported.
+- **Jev Decision Index 0.3, public suite: 65.17**, measured with the official 0.3 kit on the released weights: all 140,178 public requests answered, none unsupported.
 - **+8.2 on the public suite over Decision 2.0** (its 27B model: 56.97 on the board), ahead in all five areas.
 - **Reads images:** multiple images per request (PNG, JPEG or WebP), given as paths, URLs, PIL images or base64 data URLs; every question of the request sees all of them.
-- **Speed:** text requests take a median of 84 ms, and requests with an image a median of 342 ms, on one AMD Instinct MI325X, one request at a time.
+- **Speed:** text requests take a median of 54 ms, and requests with an image a median of 277 ms, on one AMD Instinct MI325X, one request at a time.
 - **Many questions, one call:** Choice, Yes / No and Score questions about the same input are answered together, each from its own forward pass over the input, with a probability for every option.
 
 ## Quickstart

d3_fast.py

@@ -0,0 +1,512 @@
+"""ROCm fast path of the d3 runtime: same function, same weights, less host work.
+
+``d3_runtime.D3`` installs it on AMD GPUs (``torch.version.hip``) for Qwen3.5 backbones with noncausal full
+attention; every other device and model keeps the plain path, and ``D3_FAST=0`` turns it off.
+
+- Text forward passes replay HIP graphs captured at warm-up. A pass (one batch of questions, left-padded to
+  its longest prompt exactly as in the plain path) is extended on the right with masked padding to a length
+  bucket and replays the graph of its (questions, bucket) shape. The real tokens keep their positions,
+  chunk boundaries and convolution taps; the right padding is masked out of the full-attention keys and
+  comes after every real token in the causal Gated DeltaNet layers; the readout reads each prompt's last
+  real token. Passes over the graph budget (``D3_GRAPH_TOKENS`` rows x bucket tokens) run eagerly at their
+  own shape.
+- Eager passes build both attention masks from the host-known prompt lengths, so nothing inside the forward
+  pass waits for the GPU.
+"""
+
+from __future__ import annotations
+
+import os
+import time
+from typing import Any
+
+GRAPH_TOKENS = 0
+BUCKET = 64
+CAPTURE_WARM_RUNS = 2
+
+
+def enabled(model: Any) -> str | None:
+    """None when the fast path applies to this loaded model, else the reason it does not."""
+    torch = model.torch
+    if os.environ.get("D3_FAST", "").strip().lower() in ("0", "false", "no", "off"):
+        return "D3_FAST=0"
+    if model.device.type != "cuda" or not getattr(torch.version, "hip", None):
+        return "not a ROCm GPU"
+    if model.attention_mode != "noncausal_full_attention":
+        return f"attention mode {model.attention_mode}"
+    if type(model.backbone).__name__ != "Qwen3_5Model":
+        return f"backbone {type(model.backbone).__name__}"
+    config = model.backbone.language_model.config
+    if set(config.layer_types[: config.num_hidden_layers]) - {
+        "full_attention",
+        "linear_attention",
+    }:
+        return "layer types other than full / linear attention"
+    return None
+
+
+def fusable(model: Any) -> str | None:
+    """None when the fused decoder-layer kernels (``d3_kernels.py``) apply to this model, else the reason."""
+    torch = model.torch
+    arch = torch.cuda.get_device_properties(model.device).gcnArchName.split(":")[0]
+    if arch != "gfx942":
+        return f"fused kernels verified on gfx942 only, not {arch}"
+    try:
+        import triton  # noqa: F401
+    except ImportError:
+        return "triton is not installed"
+    config = model.backbone.language_model.config
+    chunk = model.kernels.get("torch_chunk_gated_delta_rule", "")
+    checks = {
+        "hidden size a multiple of 256": config.hidden_size % 256 == 0,
+        "128-wide gated-delta heads": config.linear_key_head_dim == 128
+        and config.linear_value_head_dim == 128,
+        "256-wide attention heads": getattr(config, "head_dim", None) == 256,
+        "4-tap convolution": config.linear_conv_kernel_dim == 4,
+        "SiLU MLP": config.hidden_act == "silu",
+        "FLA chunk kernel": chunk.startswith("fla"),
+        "SDPA attention": config._attn_implementation == "sdpa",
+    }
+    missing = [name for name, ok in checks.items() if not ok]
+    return "needs " + ", ".join(missing) if missing else None
+
+
+class FusedLayers:
+    """The decoder layers and final norm of a Qwen3.5 text model through ``d3_kernels`` (eager values)."""
+
+    def __init__(self, lm: Any, torch: Any):
+        import importlib
+
+        from transformers.models.qwen3_5 import modeling_qwen3_5 as modeling
+
+        try:
+            from . import d3_kernels as kernels
+        except ImportError:
+            kernels = importlib.import_module("d3_kernels")
+        self.k = kernels
+        self.torch = torch
+        self.chunk = modeling.torch_chunk_gated_delta_rule
+        self.repeat_kv = modeling.repeat_kv
+        config = lm.config
+        self.layers = list(lm.layers[: config.num_hidden_layers])
+        self.params = []
+        with torch.no_grad():
+            for layer in self.layers:
+                p = {
+                    "linear": layer.block_type == "linear_attention",
+                    "eps": layer.input_layernorm.eps,
+                    "w1_in": (1.0 + layer.input_layernorm.weight.float()).contiguous(),
+                    "w1_post": (
+                        1.0 + layer.post_attention_layernorm.weight.float()
+                    ).contiguous(),
+                }
+                if p["linear"]:
+                    m = layer.linear_attn
+                    p.update(
+                        conv_w=m.conv1d.weight.squeeze(1).contiguous(),
+                        A_log=m.A_log.float().contiguous(),
+                        dt_bias=m.dt_bias.float().contiguous(),
+                    )
+                else:
+                    m = layer.self_attn
+                    p.update(
+                        qw1=(1.0 + m.q_norm.weight.float()).contiguous(),
+                        kw1=(1.0 + m.k_norm.weight.float()).contiguous(),
+                    )
+                self.params.append(p)
+            self.final_eps = lm.norm.eps
+            self.final_w1 = (1.0 + lm.norm.weight.float()).contiguous()
+
+    def gated_delta(self, m: Any, p: dict[str, Any], x: Any) -> Any:
+        k = self.k
+        rows, length, _ = x.shape
+        q, key, v, g, beta = k.gdn_prep(
+            m.in_proj_qkv(x),
+            m.in_proj_b(x),
+            m.in_proj_a(x),
+            p["conv_w"],
+            p["A_log"],
+            p["dt_bias"],
+            m.num_k_heads,
+            m.head_k_dim,
+        )
+        core, _ = self.chunk(
+            q,
+            key,
+            v,
+            g=g,
+            beta=beta,
+            initial_state=None,
+            output_final_state=False,
+            use_qk_l2norm_in_kernel=True,
+            cu_seqlens=None,
+            use_cache=False,
+        )
+        out = k.gated_rmsnorm(
+            core.reshape(-1, m.head_v_dim),
+            m.in_proj_z(x).reshape(-1, m.head_v_dim),
+            m.norm.weight,
+            m.norm.variance_epsilon,
+        )
+        return m.out_proj(out.reshape(rows, length, -1))
+
+    def attention(self, m: Any, p: dict[str, Any], x: Any, rotary, full_mask) -> Any:
+        rows, length, _ = x.shape
+        hd = m.head_dim
+        qp = m.q_proj(x)
+        kp = m.k_proj(x)
+        heads = qp.shape[-1] // (2 * hd)
+        kv_heads = kp.shape[-1] // hd
+        cos, sin = rotary
+        q, key = self.k.attn_prep(
+            qp, kp, p["qw1"], p["kw1"], cos, sin, heads, kv_heads, hd, m.q_norm.eps
+        )
+        v = self.repeat_kv(
+            m.v_proj(x).view(rows, length, kv_heads, hd).transpose(1, 2),
+            heads // kv_heads,
+        )
+        attn = self.torch.nn.functional.scaled_dot_product_attention(
+            q,
+            key,
+            v,
+            attn_mask=full_mask,
+            dropout_p=0.0,
+            scale=m.scaling,
+            is_causal=False,
+        )
+        gate = qp.view(rows, length, heads, 2 * hd)[..., hd:]
+        return m.o_proj(self.k.sigmoid_gate(attn.transpose(1, 2), gate))
+
+    def forward(self, embeds: Any, full_mask: Any, rowmask: Any | None, rotary) -> Any:
+        """Final-norm hidden states [B, L, D]; ``rowmask`` marks real tokens when the batch is padded."""
+        k = self.k
+        hidden, delta = embeds, None
+        for layer, p in zip(self.layers, self.params):
+            hidden, x = k.add_rmsnorm(
+                hidden, delta, p["w1_in"], p["eps"], rowmask if p["linear"] else None
+            )
+            if p["linear"]:
+                delta = self.gated_delta(layer.linear_attn, p, x)
+            else:
+                delta = self.attention(layer.self_attn, p, x, rotary, full_mask)
+            hidden, x = k.add_rmsnorm(hidden, delta, p["w1_post"], p["eps"])
+            mlp = layer.mlp
+            delta = mlp.down_proj(k.silu_mul(mlp.gate_proj(x), mlp.up_proj(x)))
+        return k.add_rmsnorm(hidden, delta, self.final_w1, self.final_eps)[1]
+
+
+class FastText:
+    """Text passes of one loaded model: HIP graphs per (rows, bucket) plus a synchronization-free eager pass."""
+
+    def __init__(
+        self,
+        model: Any,
+        *,
+        graph_tokens: int,
+        bucket: int,
+        max_rows: int,
+        fused: FusedLayers | None = None,
+        fused_skipped: str | None = None,
+    ):
+        torch = model.torch
+        self.model = model
+        self.torch = torch
+        self.fused = fused
+        self.fused_skipped = fused_skipped
+        self.cpu_threads: dict[str, int] | None = None
+        self.lm = model.backbone.language_model
+        config = self.lm.config
+        self.layers = list(self.lm.layers[: config.num_hidden_layers])
+        self.kinds = list(config.layer_types[: config.num_hidden_layers])
+        self.graph_tokens = graph_tokens
+        self.bucket = bucket
+        self.max_rows = max_rows
+        self.codes = torch.arange(model.readout.shape[0], device=model.device)
+        self.graphs: dict[tuple[int, int], dict[str, Any]] = {}
+        self.failed: dict[tuple[int, int], str] = {}
+        self.pool = None
+        self.stats = {"replays": 0, "eager": 0, "captures": 0, "capture_seconds": 0.0}
+
+    # ------------------------------------------------------------------ forward pieces
+
+    def hidden(self, ids, mask, linear_mask):
+        """Final-norm hidden states [B, L, D]: ``Qwen3_5Model.forward`` for text with the d3 masks given."""
+        torch = self.torch
+        lm = self.lm
+        embeds = lm.embed_tokens(ids)
+        rows, length = ids.shape
+        positions = torch.arange(length, device=ids.device).view(1, 1, -1)
+        positions = positions.expand(4, rows, -1)
+        text_positions, positions = positions[0], positions[1:]
+        rotary = lm.rotary_emb(embeds, positions)
+        full = mask.bool()[:, None, None, :]
+        if self.fused is not None:
+            rowmask = None if linear_mask is None else linear_mask.reshape(-1)
+            return self.fused.forward(embeds, full, rowmask, rotary)
+        hidden = embeds
+        for layer, kind in zip(self.layers, self.kinds):
+            hidden = layer(
+                hidden,
+                position_embeddings=rotary,
+                attention_mask=full if kind == "full_attention" else linear_mask,
+                position_ids=text_positions,
+                past_key_values=None,
+                use_cache=False,
+            )
+        return lm.norm(hidden)
+
+    def image_hidden(self, inputs: dict[str, Any], padded: bool):
+        """Final-norm hidden states of an image batch: ``Qwen3_5Model.forward`` with the fused text layers."""
+        torch = self.torch
+        backbone = self.model.backbone
+        ids = inputs["input_ids"]
+        mask = inputs["attention_mask"]
+        embeds = backbone.get_input_embeddings()(ids)
+        features = backbone.get_image_features(
+            inputs["pixel_values"], inputs["image_grid_thw"], return_dict=True
+        ).pooler_output
+        features = torch.cat(features, dim=0).to(embeds.device, embeds.dtype)
+        image_mask, _ = backbone.get_placeholder_mask(
+            ids, inputs_embeds=embeds, image_features=features
+        )
+        embeds = embeds.masked_scatter(image_mask, features)
+        positions = backbone.compute_3d_position_ids(
+            input_ids=ids,
+            image_grid_thw=inputs["image_grid_thw"],
+            inputs_embeds=embeds,
+            attention_mask=mask,
+            past_key_values=None,
+            mm_token_type_ids=inputs["mm_token_type_ids"],
+        )
+        rotary = self.lm.rotary_emb(embeds, positions)
+        full = mask.bool()[:, None, None, :]
+        return self.fused.forward(
+            embeds, full, mask.reshape(-1) if padded else None, rotary
+        )
+
+    def probabilities_of(self, last, counts):
+        """``D3.logits`` -> temperature -> softmax on the last-token hidden states [B, D]."""
+        model = self.model
+        if model.readout_dtype == "float32":
+            logits = last.float() @ model.readout.T
+        else:
+            logits = self.torch.nn.functional.linear(last, model.readout).float()
+        invalid = self.codes[None] >= counts[:, None]
+        return (logits.masked_fill(invalid, float("-inf")) / model.temperature).softmax(
+            -1
+        )
+
+    # ------------------------------------------------------------------ passes
+
+    def bucket_of(self, rows: int, width: int) -> int | None:
+        size = -(-width // self.bucket) * self.bucket
+        if rows > self.max_rows or rows * size > self.graph_tokens:
+            return None
+        return size
+
+    def probabilities(self, sequences, counts) -> list[list[float]]:
+        torch = self.torch
+        rows = len(sequences)
+        width = max(len(s) for s in sequences)
+        size = self.bucket_of(rows, width)
+        entry = self.graphs.get((rows, size)) if size is not None else None
+        length = size if entry is not None else width
+        ids = torch.full((rows, length), self.model.pad_id, dtype=torch.long)
+        mask = torch.zeros((rows, length), dtype=torch.long)
+        for i, sequence in enumerate(sequences):
+            ids[i, width - len(sequence) : width] = torch.as_tensor(
+                sequence, dtype=torch.long
+            )
+            mask[i, width - len(sequence) : width] = 1
+        counts_cpu = torch.as_tensor(list(counts), dtype=torch.long)
+        if entry is not None:
+            entry["ids"].copy_(ids, non_blocking=True)
+            entry["mask"].copy_(mask, non_blocking=True)
+            entry["counts"].copy_(counts_cpu, non_blocking=True)
+            entry["last"].fill_(width - 1)
+            entry["graph"].replay()
+            self.stats["replays"] += 1
+            probs = entry["probs"].cpu().tolist()
+        else:
+            device = self.model.device
+            ids = ids.to(device, non_blocking=True)
+            mask = mask.to(device, non_blocking=True)
+            padded = any(len(s) != width for s in sequences)
+            with torch.inference_mode():
+                last = self.hidden(ids, mask, mask if padded else None)[:, -1]
+                probs = (
+                    self.probabilities_of(last, counts_cpu.to(device)).cpu().tolist()
+                )
+            self.stats["eager"] += 1
+        return [p[:c] for p, c in zip(probs, counts)]
+
+    # ------------------------------------------------------------------ graphs
+
+    def shapes(self) -> list[tuple[int, int]]:
+        out = []
+        for rows in range(1, self.max_rows + 1):
+            size = self.bucket
+            while rows * size <= self.graph_tokens:
+                out.append((rows, size))
+                size += self.bucket
+        return out
+
+    def capture(self, rows: int, size: int) -> bool:
+        torch = self.torch
+        key = (rows, size)
+        if key in self.graphs:
+            return True
+        device = self.model.device
+        token = self.model.token_ids[0]
+        static = {
+            "ids": torch.full((rows, size), token, dtype=torch.long, device=device),
+            "mask": torch.ones((rows, size), dtype=torch.long, device=device),
+            "counts": torch.full((rows,), 2, dtype=torch.long, device=device),
+            "last": torch.full((1,), size - 1, dtype=torch.long, device=device),
+        }
+
+        def body():
+            hidden = self.hidden(static["ids"], static["mask"], static["mask"])
+            last = hidden.index_select(1, static["last"]).squeeze(1)
+            return self.probabilities_of(last, static["counts"])
+
+        started = time.perf_counter()
+        try:
+            with torch.inference_mode():
+                if self.pool is None:
+                    self.pool = torch.cuda.graph_pool_handle()
+                stream = torch.cuda.Stream(device=device)
+                stream.wait_stream(torch.cuda.current_stream(device))
+                with torch.cuda.stream(stream):
+                    for _ in range(CAPTURE_WARM_RUNS):
+                        body()
+                torch.cuda.current_stream(device).wait_stream(stream)
+                torch.cuda.synchronize(device)
+                graph = torch.cuda.CUDAGraph()
+                with torch.cuda.graph(graph, pool=self.pool):
+                    probs = body()
+            torch.cuda.synchronize(device)
+        except Exception as exc:  # noqa: BLE001 - the shape then runs eagerly
+            torch.cuda.synchronize(device)
+            self.failed[key] = f"{type(exc).__name__}: {str(exc)[:200]}"
+            return False
+        self.graphs[key] = {**static, "graph": graph, "probs": probs}
+        self.stats["captures"] += 1
+        self.stats["capture_seconds"] += time.perf_counter() - started
+        return True
+
+    def capture_all(self) -> float:
+        started = time.perf_counter()
+        for rows, size in self.shapes():
+            self.capture(rows, size)
+        return time.perf_counter() - started
+
+    def report(self) -> dict[str, Any]:
+        return {
+            "fused_layers": len(self.fused.layers) if self.fused is not None else 0,
+            "fused_skipped": self.fused_skipped,
+            "cpu_threads": self.cpu_threads,
+            "graph_tokens": self.graph_tokens,
+            "bucket": self.bucket,
+            "max_rows": self.max_rows,
+            "graphs": len(self.graphs),
+            "failed": dict(list(self.failed.items())[:5]),
+            "failed_count": len(self.failed),
+            **{
+                k: (round(v, 1) if isinstance(v, float) else v)
+                for k, v in self.stats.items()
+            },
+        }
+
+
+def dedupe_image_processing(processor: Any, torch: Any) -> None:
+    """Process each distinct image of one processor call once and repeat its rows for the other copies.
+
+    The runtime hands the processor one copy of the request's images per question; every copy gives the same
+    pixel rows and grid, so the outputs are identical, only the repeated resizing is skipped.
+    """
+    original = processor._process_images
+
+    def process_images(images, **kwargs):
+        flat = list(images) if isinstance(images, (list, tuple)) else [images]
+        index: dict[int, int] = {}
+        unique, order = [], []
+        for image in flat:
+            order.append(index.setdefault(id(image), len(unique)))
+            if len(unique) < len(index):
+                unique.append(image)
+        if len(unique) == len(flat):
+            return original(images, **kwargs)
+        processed = processor.image_processor(unique, **kwargs)
+        if set(processed.keys()) != {"pixel_values", "image_grid_thw"}:
+            return original(images, **kwargs)
+        grid = processed["image_grid_thw"]
+        rows = torch.split(processed["pixel_values"], grid.prod(-1).tolist())
+        processed["pixel_values"] = torch.cat([rows[i] for i in order])
+        processed["image_grid_thw"] = grid[order]
+        replacements = [
+            processor.replace_image_token(processed, image_idx=i, **kwargs)
+            for i in range(len(flat))
+        ]
+        return processed, replacements
+
+    processor._process_images = process_images
+
+
+def cpu_quota() -> int | None:
+    """CPUs this process may use: the cgroup CPU quota (containers), else the affinity mask."""
+    try:
+        with open("/sys/fs/cgroup/cpu.max", encoding="utf-8") as stream:
+            quota, period = stream.read().split()[:2]
+        if quota != "max":
+            return max(1, int(quota) // int(period))
+    except (OSError, ValueError):
+        pass
+    try:
+        return len(os.sched_getaffinity(0))
+    except (AttributeError, OSError):
+        return None
+
+
+def cap_cpu_threads(torch: Any) -> dict[str, int] | None:
+    """Keep torch's CPU threads within the container's CPU quota when OMP_NUM_THREADS is not set.
+
+    torch sizes its pool by the host's CPU count; in a container with a smaller quota the image
+    preprocessing then oversubscribes it and the idle-spinning workers starve the thread that feeds the GPU.
+    """
+    if "OMP_NUM_THREADS" in os.environ:
+        return None
+    limit, current = cpu_quota(), torch.get_num_threads()
+    if limit is None or current <= limit:
+        return None
+    torch.set_num_threads(limit)
+    return {"from": current, "to": limit}
+
+
+def install(model: Any) -> FastText | None:
+    """The fast path for a loaded ``D3`` model, or None (the reason is in ``model.fast_skipped``)."""
+    reason = enabled(model)
+    if reason is not None:
+        model.fast_skipped = reason
+        return None
+    graph_tokens = int(os.environ.get("D3_GRAPH_TOKENS", GRAPH_TOKENS))
+    bucket = int(os.environ.get("D3_GRAPH_BUCKET", BUCKET))
+    if model.processor is not None and hasattr(model.processor, "_process_images"):
+        dedupe_image_processing(model.processor, model.torch)
+    fused, skipped = None, None
+    if os.environ.get("D3_FUSED", "").strip().lower() in ("0", "false", "no", "off"):
+        skipped = "D3_FUSED=0"
+    else:
+        skipped = fusable(model)
+    if skipped is None:
+        fused = FusedLayers(model.backbone.language_model, model.torch)
+    fast = FastText(
+        model,
+        graph_tokens=graph_tokens,
+        bucket=bucket,
+        max_rows=model.batch_size,
+        fused=fused,
+        fused_skipped=skipped,
+    )
+    fast.cpu_threads = cap_cpu_threads(model.torch)
+    return fast

d3_kernels.py

@@ -0,0 +1,523 @@
+"""Fused Triton kernels for the element-wise ops of the d3 Qwen3.5 decoder layers (ROCm gfx942).
+
+Each kernel computes what the BF16 eager forward computes, rounding where the eager ops round: the BF16
+residual adds, the norms' FP32 math rounded to BF16 at the end, SiLU / sigmoid outputs in BF16, the depthwise
+convolution summed in FP64 and rounded to BF16 like MIOpen's naive kernel, RoPE products and sums each rounded
+to BF16. Transcendentals call the same OCML functions as the ATen kernels, floating-point contraction is off,
+the RMSNorm means are summed in the order of ATen's ROCm row reduction (one 64-lane wavefront per row, four
+accumulators per lane, then a lane tree; rows of 128 use 32 lanes) and ``torch.rsqrt``, which is correctly
+rounded on ROCm, is reproduced through FP64. GEMMs, attention and the gated-delta chunk kernel are unchanged.
+
+Kernels:
+    add_rmsnorm     BF16 residual add + zero-centred RMSNorm (optionally zeroing padding rows)
+    silu_mul        SiLU(gate) * up
+    gdn_prep        causal conv + SiLU + q / k (repeated to the value heads) / v split + beta + g
+    gated_rmsnorm   Gated DeltaNet output norm with the SiLU(z) gate
+    attn_prep       q / gate split, q / k RMSNorm, partial RoPE, [B, H, T, D] layout (k repeated)
+    sigmoid_gate    attention output * sigmoid(gate)
+"""
+
+from __future__ import annotations
+
+from typing import Any
+
+import numpy as np
+import torch
+import triton
+import triton.language as tl
+from triton.language.extra import libdevice
+
+EXACT = {"enable_fp_fusion": False}
+
+
+@triton.jit
+def rsqrt_rn(x):
+    """``torch.rsqrt`` on ROCm is correctly rounded; OCML's FP32 rsqrt is not."""
+    return libdevice.rsqrt(x.to(tl.float64)).to(tl.float32)
+
+
+@triton.jit
+def bf16(x):
+    return x.to(tl.bfloat16).to(tl.float32)
+
+
+@triton.jit
+def lane_tree64(v, R: tl.constexpr):
+    """[R, 64] -> [R]: the shfl_down tree of offsets 1, 2, ..., 32."""
+    a, b = tl.split(tl.reshape(v, [R, 32, 2]))
+    v = a + b
+    a, b = tl.split(tl.reshape(v, [R, 16, 2]))
+    v = a + b
+    a, b = tl.split(tl.reshape(v, [R, 8, 2]))
+    v = a + b
+    a, b = tl.split(tl.reshape(v, [R, 4, 2]))
+    v = a + b
+    a, b = tl.split(tl.reshape(v, [R, 2, 2]))
+    v = a + b
+    a, b = tl.split(tl.reshape(v, [R, 1, 2]))
+    v = a + b
+    return tl.reshape(v, [R])
+
+
+@triton.jit
+def combine_vec4(acc, R: tl.constexpr):
+    """[R, 64, 4] accumulators -> [R, 64]: ((a0 + a1) + a2) + a3."""
+    even, odd = tl.split(tl.reshape(acc, [R, 64, 2, 2]))
+    a0, a2 = tl.split(even)
+    a1, a3 = tl.split(odd)
+    return ((a0 + a1) + a2) + a3
+
+
+@triton.jit
+def sumsq_128(x, R: tl.constexpr):
+    """[R, 128] -> [R]: 32 lanes of four accumulators, then a five-level lane tree."""
+    even, odd = tl.split(tl.reshape(x * x, [R, 32, 2, 2]))
+    a0, a2 = tl.split(even)
+    a1, a3 = tl.split(odd)
+    v = ((a0 + a1) + a2) + a3
+    a, b = tl.split(tl.reshape(v, [R, 16, 2]))
+    v = a + b
+    a, b = tl.split(tl.reshape(v, [R, 8, 2]))
+    v = a + b
+    a, b = tl.split(tl.reshape(v, [R, 4, 2]))
+    v = a + b
+    a, b = tl.split(tl.reshape(v, [R, 2, 2]))
+    v = a + b
+    a, b = tl.split(tl.reshape(v, [R, 1, 2]))
+    return tl.reshape(a + b, [R])
+
+
+@triton.jit
+def sumsq_256(x, R: tl.constexpr):
+    """[R, 256] -> [R]: one four-wide load per lane, then the lane tree."""
+    return lane_tree64(combine_vec4(tl.reshape(x * x, [R, 64, 4]), R), R)
+
+
+@triton.jit
+def _add_rmsnorm_kernel(
+    res_ptr,
+    delta_ptr,
+    w1_ptr,
+    rowmask_ptr,
+    hidden_ptr,
+    out_ptr,
+    M,
+    H,
+    inv_h,
+    eps,
+    HAS_DELTA: tl.constexpr,
+    HAS_MASK: tl.constexpr,
+    R: tl.constexpr,
+):
+    rows = tl.program_id(0) * R + tl.arange(0, R)
+    rmask = (rows < M)[:, None]
+    base = rows[:, None].to(tl.int64) * H
+    cols = tl.arange(0, 256)[None, :]
+    acc = tl.zeros([R, 64, 4], dtype=tl.float32)
+    for c in range(0, H // 256):
+        offs = base + c * 256 + cols
+        x = tl.load(res_ptr + offs, mask=rmask, other=0.0).to(tl.float32)
+        if HAS_DELTA:
+            x = bf16(
+                x + tl.load(delta_ptr + offs, mask=rmask, other=0.0).to(tl.float32)
+            )
+            tl.store(hidden_ptr + offs, x.to(tl.bfloat16), mask=rmask)
+        acc = acc + tl.reshape(x * x, [R, 64, 4])
+    var = lane_tree64(combine_vec4(acc, R), R) * inv_h
+    rstd = rsqrt_rn(var + eps)[:, None]
+    if HAS_MASK:
+        keep = tl.load(rowmask_ptr + rows, mask=rows < M, other=0).to(tl.float32)[
+            :, None
+        ]
+    for c in range(0, H // 256):
+        offs = base + c * 256 + cols
+        x = tl.load(res_ptr + offs, mask=rmask, other=0.0).to(tl.float32)
+        if HAS_DELTA:
+            x = bf16(
+                x + tl.load(delta_ptr + offs, mask=rmask, other=0.0).to(tl.float32)
+            )
+        y = bf16((x * rstd) * tl.load(w1_ptr + c * 256 + cols))
+        if HAS_MASK:
+            y = y * keep
+        tl.store(out_ptr + offs, y.to(tl.bfloat16), mask=rmask)
+
+
+def add_rmsnorm(
+    residual: Any,
+    delta: Any | None,
+    weight_plus_one: Any,
+    eps: float,
+    rowmask: Any | None = None,
+) -> tuple[Any, Any]:
+    """(hidden, normed) BF16: ``hidden = residual + delta`` (or ``residual``), then ``Qwen3_5RMSNorm``.
+
+    ``weight_plus_one`` is the FP32 ``1 + w``; ``rowmask`` (one integer per row) zeroes the normed rows of
+    padding as the Gated DeltaNet layer's padding multiply does. Rows of a multiple of 256.
+    """
+    H = residual.shape[-1]
+    rows = residual.numel() // H
+    hidden = residual if delta is None else torch.empty_like(residual)
+    out = torch.empty_like(residual)
+    _add_rmsnorm_kernel[(triton.cdiv(rows, 2),)](
+        residual,
+        delta if delta is not None else residual,
+        weight_plus_one,
+        rowmask if rowmask is not None else residual,
+        hidden,
+        out,
+        rows,
+        H,
+        float(np.float32(1.0) / np.float32(H)),
+        eps,
+        HAS_DELTA=delta is not None,
+        HAS_MASK=rowmask is not None,
+        R=2,
+        num_warps=4,
+        **EXACT,
+    )
+    return hidden, out
+
+
+@triton.jit
+def _silu_mul_kernel(g_ptr, u_ptr, out_ptr, n_cols, BLOCK: tl.constexpr):
+    row = tl.program_id(0).to(tl.int64)
+    cols = tl.program_id(1) * BLOCK + tl.arange(0, BLOCK)
+    mask = cols < n_cols
+    g = tl.load(g_ptr + row * n_cols + cols, mask=mask, other=0.0).to(tl.float32)
+    u = tl.load(u_ptr + row * n_cols + cols, mask=mask, other=0.0).to(tl.float32)
+    s = bf16(g / (1.0 + libdevice.exp(-g)))
+    tl.store(out_ptr + row * n_cols + cols, (s * u).to(tl.bfloat16), mask=mask)
+
+
+def silu_mul(gate: Any, up: Any) -> Any:
+    """``silu(gate) * up`` of two contiguous BF16 tensors of one shape."""
+    n = gate.shape[-1]
+    rows = gate.numel() // n
+    out = torch.empty_like(gate)
+    _silu_mul_kernel[(rows, triton.cdiv(n, 1024))](
+        gate, up, out, n, BLOCK=1024, num_warps=4, **EXACT
+    )
+    return out
+
+
+@triton.jit
+def _gdn_prep_kernel(
+    x_ptr,
+    w_ptr,
+    b_ptr,
+    a_ptr,
+    alog_ptr,
+    dtb_ptr,
+    q_ptr,
+    k_ptr,
+    v_ptr,
+    g_ptr,
+    beta_ptr,
+    T,
+    C,
+    NK,
+    NV,
+    DK: tl.constexpr,
+    REP: tl.constexpr,
+    BT: tl.constexpr,
+    KW: tl.constexpr,
+):
+    pid_t = tl.program_id(0)
+    j = tl.program_id(1)
+    nblk = tl.cdiv(T, BT)
+    bidx = (pid_t // nblk).to(tl.int64)
+    t = (pid_t % nblk) * BT + tl.arange(0, BT)
+    tmask = t < T
+    d = tl.arange(0, DK)
+    c = j * DK + d
+    row = bidx * T + t
+    acc = tl.zeros([BT, DK], dtype=tl.float64)
+    for w in tl.static_range(KW):
+        tt = t - (KW - 1) + w
+        m = (tt >= 0) & tmask
+        xv = tl.load(
+            x_ptr + (bidx * T + tt)[:, None] * C + c[None, :],
+            mask=m[:, None],
+            other=0.0,
+        )
+        wv = tl.load(w_ptr + c * KW + w)
+        acc = acc + wv.to(tl.float64)[None, :] * xv.to(tl.float64)
+    conv = bf16(acc.to(tl.float32))
+    y = (conv / (1.0 + libdevice.exp(-conv))).to(tl.bfloat16)
+    if j < NK:
+        for r in tl.static_range(REP):
+            tl.store(
+                q_ptr + row[:, None] * (NV * DK) + ((j * REP + r) * DK + d)[None, :],
+                y,
+                mask=tmask[:, None],
+            )
+    elif j < 2 * NK:
+        for r in tl.static_range(REP):
+            tl.store(
+                k_ptr
+                + row[:, None] * (NV * DK)
+                + (((j - NK) * REP + r) * DK + d)[None, :],
+                y,
+                mask=tmask[:, None],
+            )
+    else:
+        hv = j - 2 * NK
+        tl.store(
+            v_ptr + row[:, None] * (NV * DK) + (hv * DK + d)[None, :],
+            y,
+            mask=tmask[:, None],
+        )
+        bb = tl.load(b_ptr + row * NV + hv, mask=tmask, other=0.0).to(tl.float32)
+        tl.store(
+            beta_ptr + row * NV + hv,
+            (1.0 / (1.0 + libdevice.exp(-bb))).to(tl.bfloat16),
+            mask=tmask,
+        )
+        s = tl.load(a_ptr + row * NV + hv, mask=tmask, other=0.0).to(
+            tl.float32
+        ) + tl.load(dtb_ptr + hv)
+        sp = tl.where(s > 20.0, s, libdevice.log1p(libdevice.exp(s)))
+        g = -libdevice.exp(tl.load(alog_ptr + hv)) * sp
+        tl.store(g_ptr + row * NV + hv, g, mask=tmask)
+
+
+def gdn_prep(
+    mixed_qkv: Any,
+    b: Any,
+    a: Any,
+    conv_weight: Any,
+    A_log: Any,
+    dt_bias: Any,
+    k_heads: int,
+    head_dim: int,
+) -> tuple[Any, Any, Any, Any, Any]:
+    """(q, k, v, g, beta) as the eager layer hands them to the chunk kernel (q / k repeated to the value heads).
+
+    ``mixed_qkv`` is the contiguous [B, T, C] BF16 ``in_proj_qkv`` output, ``conv_weight`` the [C, KW] BF16
+    depthwise filter, ``A_log`` / ``dt_bias`` FP32 copies of the parameters; q / k / v are SiLU(conv) in BF16,
+    beta = sigmoid(b) in BF16 and g = -exp(A_log) * softplus(a + dt_bias) in FP32.
+    """
+    B, T, C = mixed_qkv.shape
+    nv = (C - 2 * k_heads * head_dim) // head_dim
+    dev = mixed_qkv.device
+    q = torch.empty(B, T, nv, head_dim, dtype=torch.bfloat16, device=dev)
+    k = torch.empty_like(q)
+    v = torch.empty_like(q)
+    g = torch.empty(B, T, nv, dtype=torch.float32, device=dev)
+    beta = torch.empty(B, T, nv, dtype=torch.bfloat16, device=dev)
+    _gdn_prep_kernel[(B * triton.cdiv(T, 16), 2 * k_heads + nv)](
+        mixed_qkv,
+        conv_weight,
+        b,
+        a,
+        A_log,
+        dt_bias,
+        q,
+        k,
+        v,
+        g,
+        beta,
+        T,
+        C,
+        k_heads,
+        nv,
+        DK=head_dim,
+        REP=nv // k_heads,
+        BT=16,
+        KW=conv_weight.shape[-1],
+        num_warps=4,
+        **EXACT,
+    )
+    return q, k, v, g, beta
+
+
+@triton.jit
+def _gated_rmsnorm_kernel(
+    x_ptr, z_ptr, w_ptr, out_ptr, N, eps, inv_d, D: tl.constexpr, BR: tl.constexpr
+):
+    r = (tl.program_id(0) * BR + tl.arange(0, BR)).to(tl.int64)
+    d = tl.arange(0, D)
+    m = (r < N)[:, None]
+    offs = r[:, None] * D + d[None, :]
+    x = tl.load(x_ptr + offs, mask=m, other=0.0).to(tl.float32)
+    rstd = rsqrt_rn(sumsq_128(x, BR) * inv_d + eps)
+    xn = bf16(x * rstd[:, None])
+    y = bf16(tl.load(w_ptr + d).to(tl.float32)[None, :] * xn)
+    z = tl.load(z_ptr + offs, mask=m, other=0.0).to(tl.float32)
+    y = y * (z / (1.0 + libdevice.exp(-z)))
+    tl.store(out_ptr + offs, y.to(tl.bfloat16), mask=m)
+
+
+def gated_rmsnorm(core: Any, z: Any, weight: Any, eps: float) -> Any:
+    """``Qwen3_5RMSNormGated`` (BF16 weight) on contiguous [N, 128] BF16 rows ``core`` and gates ``z``."""
+    D = core.shape[-1]
+    n = core.numel() // D
+    out = torch.empty_like(core)
+    _gated_rmsnorm_kernel[(triton.cdiv(n, 16),)](
+        core,
+        z,
+        weight,
+        out,
+        n,
+        eps,
+        float(np.float32(1.0) / np.float32(D)),
+        D=D,
+        BR=16,
+        num_warps=4,
+        **EXACT,
+    )
+    return out
+
+
+@triton.jit
+def _attn_prep_kernel(
+    x_ptr,
+    w1_ptr,
+    cos_ptr,
+    sin_ptr,
+    o_ptr,
+    T,
+    NH,
+    eps,
+    inv_d,
+    cs_bstride,
+    D: tl.constexpr,
+    ROT: tl.constexpr,
+    QW: tl.constexpr,
+    REP: tl.constexpr,
+):
+    bt = tl.program_id(0).to(tl.int64)
+    hid = tl.program_id(1)
+    b = bt // T
+    t = bt % T
+    d = tl.arange(0, D)
+    half: tl.constexpr = ROT // 2
+    partner = tl.where(d < half, d + half, tl.where(d < ROT, d - half, d))
+    src = x_ptr + bt * (NH * QW) + hid * QW
+    x = tl.load(src + d).to(tl.float32)
+    sumsq = sumsq_256(tl.reshape(x, [1, D]), 1)
+    rstd = rsqrt_rn(tl.reshape(sumsq, []) * inv_d + eps)
+    xp = tl.load(src + partner).to(tl.float32)
+    y = bf16((x * rstd) * tl.load(w1_ptr + d))
+    yp = bf16((xp * rstd) * tl.load(w1_ptr + partner))
+    yp = tl.where(d < half, -yp, yp)
+    rot = d < ROT
+    cs = b * cs_bstride + t * ROT + d
+    c = tl.load(cos_ptr + cs, mask=rot, other=1.0).to(tl.float32)
+    s = tl.load(sin_ptr + cs, mask=rot, other=0.0).to(tl.float32)
+    out = tl.where(rot, bf16(bf16(y * c) + bf16(yp * s)), y).to(tl.bfloat16)
+    for r in tl.static_range(REP):
+        tl.store(o_ptr + ((b * NH * REP + hid * REP + r) * T + t) * D + d, out)
+
+
+def attn_prep(
+    q_proj_out: Any,
+    k_proj_out: Any,
+    q_norm_w1: Any,
+    k_norm_w1: Any,
+    cos: Any,
+    sin: Any,
+    heads: int,
+    kv_heads: int,
+    head_dim: int,
+    eps: float,
+) -> tuple[Any, Any]:
+    """(q [B, H, T, D], k [B, H, T, D]) in BF16: head RMSNorms, then the partial RoPE, k repeated to H heads.
+
+    ``q_proj_out`` [B, T, H * 2D] holds query and gate per head, ``k_proj_out`` [B, T, Hkv * D] (both
+    contiguous BF16); the norms multiply by ``1 + w`` (FP32) and round to BF16 before the RoPE, whose products
+    and sum round to BF16 like the eager BF16 ops; ``cos`` / ``sin`` are the contiguous [B or 1, T, rotary dim]
+    BF16 rotary tables; the head dim is 256.
+    """
+    if head_dim != 256:
+        raise ValueError("attn_prep reduces rows of 256")
+    B, T = q_proj_out.shape[:2]
+    dev = q_proj_out.device
+    q = torch.empty(B, heads, T, head_dim, dtype=torch.bfloat16, device=dev)
+    k = torch.empty_like(q)
+    rot = cos.shape[-1]
+    # One launch per tensor: Triton's AMD pointer canonicalization fails on a runtime branch between two
+    # pointers when only one of their tensors fits the 2 GiB buffer range.
+    for x, w1, out, n, width, rep in (
+        (q_proj_out, q_norm_w1, q, heads, 2 * head_dim, 1),
+        (k_proj_out, k_norm_w1, k, kv_heads, head_dim, heads // kv_heads),
+    ):
+        _attn_prep_kernel[(B * T, n)](
+            x,
+            w1,
+            cos,
+            sin,
+            out,
+            T,
+            n,
+            eps,
+            float(np.float32(1.0) / np.float32(head_dim)),
+            0 if cos.shape[0] == 1 else T * rot,
+            D=head_dim,
+            ROT=rot,
+            QW=width,
+            REP=rep,
+            num_warps=2,
+            **EXACT,
+        )
+    return q, k
+
+
+@triton.jit
+def _sigmoid_gate_kernel(
+    a_ptr,
+    g_ptr,
+    out_ptr,
+    T,
+    H,
+    sab,
+    sat,
+    sah,
+    sgb,
+    sgt,
+    sgh,
+    D: tl.constexpr,
+    HB: tl.constexpr,
+):
+    bt = tl.program_id(0).to(tl.int64)
+    h = tl.program_id(1) * HB + tl.arange(0, HB)
+    b = bt // T
+    t = bt % T
+    d = tl.arange(0, D)
+    m = (h < H)[:, None]
+    a = tl.load(
+        a_ptr + b * sab + t * sat + h[:, None] * sah + d[None, :], mask=m, other=0.0
+    ).to(tl.float32)
+    gt = tl.load(
+        g_ptr + b * sgb + t * sgt + h[:, None] * sgh + d[None, :], mask=m, other=0.0
+    ).to(tl.float32)
+    s = bf16(1.0 / (1.0 + libdevice.exp(-gt)))
+    tl.store(
+        out_ptr + bt * (H * D) + h[:, None] * D + d[None, :],
+        (a * s).to(tl.bfloat16),
+        mask=m,
+    )
+
+
+def sigmoid_gate(attn_out: Any, gate: Any) -> Any:
+    """``attn_out * sigmoid(gate)`` -> [B, T, H * D] BF16 from [B, T, H, D] views with unit last stride."""
+    B, T, H, D = attn_out.shape
+    out = torch.empty(B, T, H * D, dtype=torch.bfloat16, device=attn_out.device)
+    _sigmoid_gate_kernel[(B * T, triton.cdiv(H, 4))](
+        attn_out,
+        gate,
+        out,
+        T,
+        H,
+        attn_out.stride(0),
+        attn_out.stride(1),
+        attn_out.stride(2),
+        gate.stride(0),
+        gate.stride(1),
+        gate.stride(2),
+        D=D,
+        HB=4,
+        num_warps=4,
+        **EXACT,
+    )
+    return out

d3_runtime.py

@@ -13,7 +13,8 @@ probability for every option. No text is generated and no input is truncated.
     model.system_one(state="...", questions={...}, images=["photo.png", "label.jpg"])
 
 The checkpoint directory holds ``config.json`` + ``model*.safetensors`` (a transformers
-``Qwen3_5Model``), ``readout.safetensors`` (``{"weight": [255, hidden]}``), ``decision_config.json``
+``Qwen3_5Model``, or a ``Qwen3VLModel`` when ``config.json`` says ``model_type: qwen3_vl``),
+``readout.safetensors`` (``{"weight": [255, hidden]}``), ``decision_config.json``
 (prompt family, answer codes, attention mode, pooling, temperature, input limit) and the tokenizer.
 ``d3_format.py`` next to this file is the prompt and answer-code contract of the model.
 
@@ -27,10 +28,17 @@ with the vision tower (``visual.*`` weights).
 
 Numerics: BF16 backbone with SDPA attention, FP32 readout and softmax (unless the checkpoint says
 otherwise). A request's questions run in request order, ``batch_size`` per forward pass, each batch
-left-padded to its longest prompt. ``noncausal_full_attention`` lets the full-attention layers see the
-whole prompt while the Gated DeltaNet layers stay causal. The Gated DeltaNet kernels are the ones
-transformers binds at import: flash-linear-attention (and causal-conv1d) when installed, its PyTorch
-reference implementation otherwise.
+left-padded to its longest prompt. ``noncausal_full_attention`` (Qwen3.5 backbones only) lets the
+full-attention layers see the whole prompt while the Gated DeltaNet layers stay causal. The Gated
+DeltaNet kernels are the ones transformers binds at import: flash-linear-attention (and causal-conv1d)
+when installed, its PyTorch reference implementation otherwise. Qwen3-VL backbones have no linear-attention
+layers and run causal attention only.
+
+``permutation_average=True`` (off by default) also scores every choice question with two or more options
+with its options in reversed order, in the same forward passes as the original order, and answers with
+the per-option mean of the two distributions. Noul and score questions are scored once. A pass carrying both
+orders that would exceed ``MERGE_TOKENS`` padded tokens runs the reversed prompts in a pass of their own, so
+peak memory stays that of the original order.
 
 The noncausal attention mask hook is adapted from perplexity-ai/pplx-decider-v1.1-27b, Copyright
 Perplexity AI, Apache License 2.0.
@@ -73,11 +81,19 @@ except ImportError:
         to_answer,
         user_prompt,
     )
+try:
+    from . import d3_fast
+except ImportError:
+    try:
+        import d3_fast
+    except ImportError:
+        d3_fast = None
 
 RUNTIME = "d3-runtime/1"
 FORMAT_VERSION = 1
 PROMPTS = ("d3",)
 ATTENTION_MODES = ("causal", "noncausal_full_attention")
+BACKBONES = ("qwen3_5", "qwen3_vl")
 DEFAULT_BATCH_SIZE = 8
 SCORE_LEVELS = (2, 10)
 MANIFEST = "MODEL_MANIFEST.json"
@@ -97,6 +113,15 @@ MAX_IMAGE_SOURCE_PIXELS = 16_000_000
 IMAGE_FORMATS = ("PNG", "JPEG", "WEBP")
 DOWNLOAD_TIMEOUT_SECONDS = 30
 MAX_DOWNLOAD_BYTES = 64 << 20
+MERGE_TOKENS = 8192
+# What the processor returns for an image batch; the fast path runs exactly these through the fused layers.
+IMAGE_INPUTS = (
+    "input_ids",
+    "attention_mask",
+    "mm_token_type_ids",
+    "pixel_values",
+    "image_grid_thw",
+)
 
 
 class MaxLengthExceeded(ValueError):
@@ -114,9 +139,7 @@ class Question:
 
     kind: str  # choice | noul | score
     original: Mapping[str, Any]
-    rendered: dict[
-        str, Any
-    ]  # a choice or noul question in the d3_format contract
+    rendered: dict[str, Any]  # a choice or noul question in the d3_format contract
     keys: list[str]
     descriptions: list[Any]
 
@@ -216,6 +239,21 @@ def image_messages(
     raise ValueError(f"unknown prompt family {prompt!r}")
 
 
+def reversed_question(question: Question) -> dict[str, Any] | None:
+    """The rendered choice question with its options in reverse order (None when there is nothing to permute)."""
+    if question.kind != "choice" or len(question.keys) < 2:
+        return None
+    criteria = question.rendered["criteria"]
+    return dict(
+        question.rendered, criteria={key: criteria[key] for key in reversed(criteria)}
+    )
+
+
+def average_orders(forward: Sequence[float], backward: Sequence[float]) -> list[float]:
+    """Per-option mean of the original-order and reversed-order distributions, in original option order."""
+    return [(a + b) / 2 for a, b in zip(forward, reversed(backward))]
+
+
 def _canonical(value: Any) -> str:
     return (
         value
@@ -272,10 +310,19 @@ def product_answer(
 def _data_url_payload(value: str, strict: bool) -> bytes:
     header, separator, encoded = value.partition(",")
     kind = header.strip().lower()
-    if not separator or not kind.startswith("data:image/") or not kind.endswith(";base64"):
+    if (
+        not separator
+        or not kind.startswith("data:image/")
+        or not kind.endswith(";base64")
+    ):
         raise ValueError("a data URL image is data:image/<format>;base64,<data>")
     if strict:
-        if kind[len("data:image/") : -len(";base64")] not in ("png", "jpeg", "jpg", "webp"):
+        if kind[len("data:image/") : -len(";base64")] not in (
+            "png",
+            "jpeg",
+            "jpg",
+            "webp",
+        ):
             raise ValueError("images must be base64 PNG, JPEG or WebP data URLs")
         if len(encoded) > 4 * -(-MAX_IMAGE_BYTES // 3):
             raise ValueError(f"each image must be at most {MAX_IMAGE_BYTES:,} bytes")
@@ -292,7 +339,9 @@ def _download(url: str) -> bytes:
     with urllib.request.urlopen(request, timeout=DOWNLOAD_TIMEOUT_SECONDS) as response:
         payload = response.read(MAX_DOWNLOAD_BYTES + 1)
     if len(payload) > MAX_DOWNLOAD_BYTES:
-        raise ValueError(f"the image at {url} is larger than {MAX_DOWNLOAD_BYTES:,} bytes")
+        raise ValueError(
+            f"the image at {url} is larger than {MAX_DOWNLOAD_BYTES:,} bytes"
+        )
     return payload
 
 
@@ -339,7 +388,12 @@ def load_image(value: Any, *, strict: bool = False):
                 image.verify()
         with Image.open(io.BytesIO(payload)) as image:
             return image.convert("RGB")
-    except (OSError, SyntaxError, UnidentifiedImageError, Image.DecompressionBombError) as exc:
+    except (
+        OSError,
+        SyntaxError,
+        UnidentifiedImageError,
+        Image.DecompressionBombError,
+    ) as exc:
         raise ValueError(f"invalid image data ({type(exc).__name__})") from exc
 
 
@@ -473,6 +527,27 @@ def vision_weights_present(root: Path) -> bool:
 # ---------------------------------------------------------------------------------------------
 
 
+def backbone_type(root: Path) -> str:
+    """``model_type`` of the checkpoint's ``config.json``: ``qwen3_5`` or ``qwen3_vl``."""
+    kind = json.loads((Path(root) / "config.json").read_text(encoding="utf-8")).get(
+        "model_type"
+    )
+    if kind not in BACKBONES:
+        raise ValueError(f"unsupported backbone model_type {kind!r}")
+    return kind
+
+
+def backbone_class(kind: str):
+    """The transformers backbone class of a ``model_type``."""
+    if kind == "qwen3_vl":
+        from transformers.models.qwen3_vl.modeling_qwen3_vl import Qwen3VLModel
+
+        return Qwen3VLModel
+    from transformers.models.qwen3_5.modeling_qwen3_5 import Qwen3_5Model
+
+    return Qwen3_5Model
+
+
 def enable_noncausal_full_attention(text_model) -> None:
     """Let the softmax-attention layers see future tokens; keep padding and the causal recurrence.
 
@@ -586,6 +661,10 @@ class Prepared:
     images: list[Any] = field(default_factory=list)
     texts: dict[str, str] = field(default_factory=dict)
     lengths: dict[str, int] = field(default_factory=dict)
+    # Reversed-option prompts of the choice questions (permutation_average only).
+    reversed_sequences: dict[str, list[int]] = field(default_factory=dict)
+    reversed_texts: dict[str, str] = field(default_factory=dict)
+    reversed_lengths: dict[str, int] = field(default_factory=dict)
 
     @property
     def runnable(self) -> list[str]:
@@ -605,11 +684,11 @@ class D3:
         max_length: int | None = None,
         readout_dtype: str | None = None,
         model_name: str | None = None,
+        permutation_average: bool = False,
     ):
         import torch
         from safetensors.torch import load_file
         from transformers import AutoTokenizer
-        from transformers.models.qwen3_5.modeling_qwen3_5 import Qwen3_5Model
 
         self.torch = torch
         self.root = Path(root)
@@ -621,12 +700,20 @@ class D3:
         if config.get("format_version") != FORMAT_VERSION:
             raise ValueError("unsupported decision_config.json format_version")
         self.config = config
+        self.backbone_type = backbone_type(self.root)
         self.prompt = config.get("prompt", "d3")
         if self.prompt not in PROMPTS:
             raise ValueError(f"unknown prompt family {self.prompt!r}")
         self.attention_mode = config.get("attention_mode", "causal")
         if self.attention_mode not in ATTENTION_MODES:
             raise ValueError(f"unknown attention mode {self.attention_mode!r}")
+        if (
+            self.attention_mode == "noncausal_full_attention"
+            and self.backbone_type != "qwen3_5"
+        ):
+            raise ValueError(
+                "noncausal_full_attention is defined for Qwen3.5 backbones only"
+            )
         if config.get("pooling", "last") != "last":
             raise ValueError(f"unsupported pooling {config.get('pooling')!r}")
         self.temperature = float(config.get("temperature", 1.0))
@@ -643,6 +730,7 @@ class D3:
         self.model_name = (
             model_name or (manifest or {}).get("model_name") or self.root.name
         )
+        self.permutation_average = bool(permutation_average)
 
         self.tokenizer = AutoTokenizer.from_pretrained(str(self.root))
         self.tokenizer.padding_side = "left"
@@ -667,7 +755,7 @@ class D3:
         self.device = torch.device(device)
         if self.device.type == "cuda":
             torch.cuda.set_device(self.device)
-        self.kernels = kernel_report()
+        self.kernels = kernel_report() if self.backbone_type == "qwen3_5" else {}
         if self.device.type == "cpu" and any(
             v.startswith(("fla", "causal_conv1d"))
             for k, v in self.kernels.items()
@@ -678,7 +766,7 @@ class D3:
                 "use an environment without them for CPU inference"
             )
         torch.manual_seed(20260919)
-        self.backbone = Qwen3_5Model.from_pretrained(
+        self.backbone = backbone_class(self.backbone_type).from_pretrained(
             str(self.root),
             dtype=torch.bfloat16,
             attn_implementation="sdpa",
@@ -697,16 +785,22 @@ class D3:
         self.processor = None
         self.image_unavailable: str | None = None
         if not vision_weights_present(self.root):
-            self.image_unavailable = "the checkpoint has no vision tower (visual.* weights)"
+            self.image_unavailable = (
+                "the checkpoint has no vision tower (visual.* weights)"
+            )
         elif not linearize_patch_embed(self.backbone):
             self.image_unavailable = "the vision patch embedding was not found"
         else:
             try:
                 self.processor = load_processor(self.root)
-            except Exception as exc:  # noqa: BLE001 - text requests do not use the processor
+            except (
+                Exception
+            ) as exc:  # noqa: BLE001 - text requests do not use the processor
                 self.image_unavailable = (
                     f"the image processor failed to load ({type(exc).__name__}: {exc})"
                 )
+        self.fast_skipped = None if d3_fast else "d3_fast.py is not present"
+        self.fast = d3_fast.install(self) if d3_fast else None
         self.loaded_seconds = time.perf_counter() - started
 
     @classmethod
@@ -725,12 +819,14 @@ class D3:
         max_length: int | None = None,
         readout_dtype: str | None = None,
         model_name: str | None = None,
+        permutation_average: bool = False,
     ) -> D3:
         """Load a package directory or Hub repository.
 
         ``verify``: ``fast`` (default) hashes every file of ``MODEL_MANIFEST.json`` up to 64 MiB and checks
         the size of the weight shards; ``full`` hashes every file; ``none`` skips the check. A checkpoint
-        without a manifest (a plain code-readout export) loads unverified.
+        without a manifest (a plain code-readout export) loads unverified. ``permutation_average``: see the
+        module docstring.
         """
         root = resolve_dir(
             name_or_path,
@@ -749,6 +845,7 @@ class D3:
             max_length=max_length,
             readout_dtype=readout_dtype,
             model_name=model_name,
+            permutation_average=permutation_average,
         )
 
     # ------------------------------------------------------------------ requests
@@ -769,7 +866,9 @@ class D3:
         if not images:
             return []
         if self.image_unavailable is not None:
-            raise ValueError(f"image inputs are not available: {self.image_unavailable}")
+            raise ValueError(
+                f"image inputs are not available: {self.image_unavailable}"
+            )
         decoded = []
         for number, value in enumerate(images):
             try:
@@ -829,8 +928,53 @@ class D3:
                     }
                     continue
                 prepared.sequences[key] = sequence
+        if self.permutation_average:
+            self._prepare_reversed(state, prepared)
         return prepared
 
+    def _prepare_reversed(
+        self, state: Any, prepared: Prepared, visual: int = 0
+    ) -> None:
+        """Render and tokenize the reversed-option prompt of every runnable choice question.
+
+        ``visual``: the image tokens of the request (image path), counted like ``_prepare_images`` counts them.
+        """
+        n_images = len(prepared.images)
+        texts = []
+        for key in prepared.runnable:
+            flipped = reversed_question(prepared.questions[key])
+            if flipped is not None:
+                texts.append(
+                    (
+                        key,
+                        (
+                            self.image_text(state, flipped, n_images)
+                            if n_images
+                            else self.text(state, flipped)
+                        ),
+                    )
+                )
+        if not texts:
+            return
+        tokenizer = self.processor.tokenizer if n_images else self.tokenizer
+        ids = tokenizer([t for _, t in texts], add_special_tokens=False)["input_ids"]
+        for (key, text), sequence in zip(texts, ids):
+            length = len(sequence) - n_images + visual
+            if self.max_length is not None and length > self.max_length:
+                for planned in (prepared.sequences, prepared.texts, prepared.lengths):
+                    planned.pop(key, None)
+                prepared.errors[key] = {
+                    "type": prepared.questions[key].kind,
+                    "error": "max_length_exceeded",
+                    "message": f"the reversed-option prompt has {length} tokens, over the maximum context "
+                    f"length of {self.max_length} tokens; nothing was truncated",
+                }
+            elif n_images:
+                prepared.reversed_texts[key] = text
+                prepared.reversed_lengths[key] = length
+            else:
+                prepared.reversed_sequences[key] = sequence
+
     def _prepare_images(
         self, state: Any, questions: Mapping[str, Any], images: list[Any]
     ) -> Prepared:
@@ -882,6 +1026,8 @@ class D3:
                     continue
                 prepared.texts[key] = text
                 prepared.lengths[key] = length
+        if self.permutation_average:
+            self._prepare_reversed(state, prepared, sum(visual))
         return prepared
 
     def logits(self, sequences: Sequence[Sequence[int]], counts: Sequence[int]):
@@ -914,6 +1060,8 @@ class D3:
         self, sequences: Sequence[Sequence[int]], counts: Sequence[int]
     ) -> list[list[float]]:
         """Softmax over each prompt's own codes, in option order."""
+        if self.fast is not None:
+            return self.fast.probabilities(sequences, counts)
         probs = (
             (self.logits(sequences, counts) / self.temperature)
             .softmax(-1)
@@ -946,8 +1094,15 @@ class D3:
                 f"planned {width} input tokens, the processor produced {encoded['input_ids'].shape[1]}"
             )
         inputs = {name: value.to(self.device) for name, value in encoded.items()}
+        fused = self.fast is not None and self.fast.fused is not None
         with torch.inference_mode(), sdpa_backends(self.device):
-            hidden = self.backbone(**inputs, use_cache=False).last_hidden_state[:, -1]
+            if fused and set(inputs) == set(IMAGE_INPUTS):
+                padded = not bool(encoded["attention_mask"].all())
+                hidden = self.fast.image_hidden(inputs, padded)[:, -1]
+            else:
+                hidden = self.backbone(**inputs, use_cache=False).last_hidden_state[
+                    :, -1
+                ]
             if self.readout_dtype == "float32":
                 logits = hidden.float() @ self.readout.T
             else:
@@ -976,14 +1131,31 @@ class D3:
         out: dict[str, list[float]] = {}
         for start in range(0, len(keys), self.batch_size):
             chunk = keys[start : start + self.batch_size]
-            probs = self.image_probabilities(
-                [prepared.texts[k] for k in chunk],
-                prepared.images,
-                [len(prepared.questions[k].keys) for k in chunk],
-                max(prepared.lengths[k] for k in chunk),
-            )
+            extra = [k for k in chunk if k in prepared.reversed_texts]
+            texts = [prepared.texts[k] for k in chunk] + [
+                prepared.reversed_texts[k] for k in extra
+            ]
+            counts = [len(prepared.questions[k].keys) for k in chunk + extra]
+            widths = [prepared.lengths[k] for k in chunk] + [
+                prepared.reversed_lengths[k] for k in extra
+            ]
+            n = len(chunk)
+            if extra and len(widths) * max(widths) > MERGE_TOKENS:
+                probs = self.image_probabilities(
+                    texts[:n], prepared.images, counts[:n], max(widths[:n])
+                ) + self.image_probabilities(
+                    texts[n:], prepared.images, counts[n:], max(widths[n:])
+                )
+            else:
+                probs = self.image_probabilities(
+                    texts, prepared.images, counts, max(widths)
+                )
             out.update(zip(chunk, probs))
-        return out, sum(prepared.lengths[k] for k in keys)
+            for key, backward in zip(extra, probs[n:]):
+                out[key] = average_orders(out[key], backward)
+        return out, sum(prepared.lengths[k] for k in keys) + sum(
+            prepared.reversed_lengths.values()
+        )
 
     def run(self, prepared: Prepared) -> tuple[dict[str, list[float]], int]:
         """Probabilities per runnable question (request order, ``batch_size`` per pass) and the input tokens."""
@@ -993,12 +1165,24 @@ class D3:
         out: dict[str, list[float]] = {}
         for start in range(0, len(keys), self.batch_size):
             chunk = keys[start : start + self.batch_size]
-            probs = self.probabilities(
-                [prepared.sequences[k] for k in chunk],
-                [len(prepared.questions[k].keys) for k in chunk],
-            )
+            extra = [k for k in chunk if k in prepared.reversed_sequences]
+            sequences = [prepared.sequences[k] for k in chunk] + [
+                prepared.reversed_sequences[k] for k in extra
+            ]
+            counts = [len(prepared.questions[k].keys) for k in chunk + extra]
+            n = len(chunk)
+            if extra and len(sequences) * max(map(len, sequences)) > MERGE_TOKENS:
+                probs = self.probabilities(
+                    sequences[:n], counts[:n]
+                ) + self.probabilities(sequences[n:], counts[n:])
+            else:
+                probs = self.probabilities(sequences, counts)
             out.update(zip(chunk, probs))
-        return out, sum(len(prepared.sequences[k]) for k in keys)
+            for key, backward in zip(extra, probs[n:]):
+                out[key] = average_orders(out[key], backward)
+        return out, sum(len(prepared.sequences[k]) for k in keys) + sum(
+            len(s) for s in prepared.reversed_sequences.values()
+        )
 
     def respond(
         self,
@@ -1049,7 +1233,8 @@ class D3:
         lengths: Sequence[int] = (37, 64, 320, 333, 1000, 1024),
         images: bool = True,
     ) -> float:
-        """Compile and autotune the kernels for every batch size up to ``batch_size`` before serving.
+        """Compile and autotune the kernels for every batch size up to ``batch_size`` (twice that with
+        ``permutation_average``) before serving.
 
         The Gated DeltaNet kernels take the batch size as a compile-time constant, and Triton specializes their
         length and chunk-count arguments on being 1 or a multiple of 16; these lengths cover every combination
@@ -1058,7 +1243,10 @@ class D3:
         four 1.6 MP images) also warm the vision tower. Answers are unchanged.
         """
         started = time.perf_counter()
-        for size in range(1, self.batch_size + 1):
+        if self.fast is not None:
+            self.fast.capture_all()
+        widest = self.batch_size * (2 if self.permutation_average else 1)
+        for size in range(1, widest + 1):
             for length in lengths:
                 sequences = [
                     [
@@ -1091,6 +1279,8 @@ class D3:
         self.backbone.to(target)
         self.readout = self.readout.to(target)
         self.device = target
+        if self.fast is not None:
+            self.fast, self.fast_skipped = None, "moved after loading"
         return self
 
     def parameter_count(self) -> int:
@@ -1108,7 +1298,7 @@ class D3:
             if (self.root / name).is_file()
         }
         identity = (self.manifest or {}).get("identity", {})
-        return {
+        record = {
             "kind": "d3-code-readout",
             "runtime": RUNTIME,
             "model_name": self.model_name,
@@ -1131,6 +1321,13 @@ class D3:
             "refused); no option filtering; one fixed prompt for every request.",
             "images": self.image_contract(),
         }
+        if self.permutation_average:
+            record["permutation_average"] = True
+            record["policy"] += (
+                " Choice questions with two or more options are also scored with their options in reversed "
+                "order, in the same forward passes, and answered with the per-option mean of both distributions."
+            )
+        return record
 
     def image_contract(self) -> dict[str, Any]:
         """How image inputs are read (or why they are not available)."""
@@ -1167,4 +1364,11 @@ class D3:
         }
         if self.device.type == "cuda":
             info["gpu"] = torch.cuda.get_device_name(self.device)
+        info["fast_path"] = self.fast_report()
         return info
+
+    def fast_report(self) -> dict[str, Any]:
+        """What the ROCm fast path (``d3_fast.py``) does in this process, or why it is off."""
+        if self.fast is None:
+            return {"active": False, "reason": self.fast_skipped}
+        return {"active": True, **self.fast.report()}