ROCm: add DeepSeek V4.1 Flash support for Strix Halo (gfx1151)

.gitignore

@@ -84,5 +84,16 @@ __pycache__/
 .*.swp
 .DS_Store
 
+# Linux TP protocol and ROCm ownership harnesses
+tests/test_tp_linux
+ds4-kernel-v41-tp-attention
+ds4-kernel-v41-tp-moe
+ds4-kernel-v41-tp-gates
+
+/ds4-kernel-v41-tp-mmq
+
+/ds4-kernel-v41-tp-down
+
+/ds4-kernel-v41-tp-bind-failure
 /tests/test_qwen4_conv_parallel
 /tests/test_q8_prefill_variants

QA_BEFORE_RELEASES.md

@@ -1102,6 +1102,8 @@ clients.
   and streaming decode. Repeat across OpenAI chat, Responses, Anthropic, and
   completions. Abandoned work must stop at the next backend-safe boundary, and
   a valid request after each cancellation must complete normally.
+- For V4.1 concurrent serving, run `tests/test_server_batching.py --pairs 8 --workers 4 --output batching.json` with the server defaults, then repeat with `--stream --cancel-first 4` and `--same-prompt --case long-greedy`. Cover resident and SSD experts, cold and restored disk checkpoints, and a separately recorded `--mixed-prefill-quantum 2048` control. Preserve full responses and inspect actual slot overlap and chunk boundaries. Fixed-size diagnostic controls do not substitute for testing the defaults. Fixture groups named "batch" in official scoring do not test concurrent requests.
+- Set time limits from observed model loading, prefill, generation and queue duration. If a test is still progressing when its limit expires, preserve the attempt, increase or remove the inadequate limit and complete it; investigate a stalled test. Do not count partial assertions as a completed suite. For the vision-cache test, use `--timeout` and `--thinking-tokens` when the model needs more time or reasoning tokens, recording the settings and checking that answers finish normally.
 - For the repeatable chat-completions cancellation and slot-reuse gate, run
   `python3 tests/test_server_batching.py --url http://127.0.0.1:8000 --pairs 2
   --workers 4 --case short-sampled --max-tokens 12 --cancel-first 4`. Then run
@@ -1664,6 +1666,39 @@ Exclude loading. Only the first row is initial prefill; later rows are appends.
   --step-mul 2 --ctx-alloc 36864 --gen-tokens 128 --show-output --csv RESULT.csv
 ```
 
+## DeepSeek V4.1 Flash (ROCm/gfx1151)
+
+- ROCm 10.0, Strix Halo `gfx1151`, calibrated `DeepSeek-V4.1-Flash-Q2.gguf` (365,713,686,528 bytes; saved SHA-256 `1ce6a8f8806205c13330d7ca287bd198331dc5ca35ccc5d8a9a92a188a6f6f42`). Text/vision, resident experts, SSD streaming and two-machine inference; Engram stays disk-backed.
+- Official scoring: 115 cases / 3,954 teacher-forced targets per ROCm resident and two-machine panel, default arithmetic, no DSpark or `--quality`. CUDA is the recorded `a04f46f` GB10 reference with a 76 GiB SSD cache; it predates later CUDA SSD prefill changes. Metal values are published same-Q2 aggregates. CUDA and Metal figures are recorded references; this ROCm panel did not rerun them.
+
+| Official panel | Targets | ROCm resident NLL | ROCm two-rank NLL | CUDA reference NLL |
+|---|---:|---:|---:|---:|
+| General 100 | 2,994 | 0.368267385 | 0.362317121 | 0.363135483 |
+| Batched 12 | 768 | 0.521719404 | 0.525356473 | 0.521390812 |
+| Selected 8K; 8,197 prompt tokens | 64 | 0.133878180 | 0.138695106 | 0.116477286 |
+| Selected 16K; 16,389 prompt tokens | 64 | 0.488504024 | 0.497056311 | 0.478793408 |
+| Selected 64K; 65,541 prompt tokens | 64 | 0.439256315 | 0.448364408 | 0.437359191 |
+
+| General 100 backend/mode | NLL | Target-logprob MAE | Recorded API top1 agreement |
+|---|---:|---:|---:|
+| ROCm resident | 0.368267385 | 0.230736441 | 2705/2994 |
+| ROCm two-rank | 0.362317121 | 0.223234795 | 2704/2994 |
+| CUDA SSD | 0.363135483 | 0.223383502 | 2696/2994 |
+| Published Metal resident batching | 0.365680596 | — | 2705/2994 |
+| Published Metal scalar/SSD | 0.364576009 | 0.227531809 | 2697/2994 |
+
+- Probability results are mixed. Resident general NLL is 0.71% above published Metal resident and 1.41% above CUDA. The selected long cases each contain only 64 targets; neither these scores nor passing practical checks establish universal generation or distribution equivalence.
+- Fixed-length greedy literary continuations can become repetitive; the native benchmark checks timing and completion, not prose quality. Coding, tool-use and image checks use their own prompts and validators.
+- Practical checks pass in resident and two-machine modes: nine executable coding checks with a 36,992-token prompt, seven real tool-use checks after a 9,509-token prompt and a 24-token cached append, plus 28 sampled checks. Sampled settings: temperature 0.6, top-p 0.95, seed 1729, min-p/top-k 0; thinking and DSpark off.
+- State: 69 exact comparisons and 138 full logit vectors per resident/two-machine panel. Covers interleaved sessions, snapshot restore, rewind, corrupt snapshot rejection and three cancellation/recovery modes, at lengths from 31 through 16,383 tokens. Current transport measurements and complete-output checks are listed in [CLUSTERING_ROCM.md](docs/CLUSTERING_ROCM.md#measured-performance).
+- Vision: six resident and six RoCE image cases pass, plus focused SSD photo/screenshot checks; 168 full vectors and 112 complete states in total. Image facts, image changes and exact replay are checked. Restoring an image snapshot into a new session and every image over every transport were not tested. Image inputs are correctness checks, not timing inputs.
+- SSD: 92 GiB cache, fresh 16K/64K with 69,632 allocated context and 512/128 outputs; complete frontiers and printed continuations match resident execution. Memory/swap results and exact commands are in [STRIX_HALO.md](docs/STRIX_HALO.md#ssd-performance).
+- Kernel checks cover production layouts, full outputs, independent FP64 references, tail rows, canaries, selection ties and nonfinite inputs. The V4.1 split-key decode attention path uses FP32 throughout; matrix instructions remain in the prefill paths. FP16 matrix paths accumulate in FP32; Engram packing checks representability and retains the FP32 fallback. Numerical differences are assessed with the official and practical results above.
+- ROCm/gfx1151/V4.1 guards preserve other model/backend paths. Shared CPU/CUDA/Metal code was checked with preprocessing comparisons; five frontends build/link/help checks and relevant memory/cache/Engram units pass. No fresh physical GLM, CUDA or Metal regression run is claimed.
+- Exactly two machines, each with its assigned experts in RAM; no cluster SSD streaming or DSpark. RoCE uses system-RAM staging. Plain TCP over USB4 Ethernet also works. Long-running production endurance was not tested.
+- Performance, hardware/SSD/profile details and commands: [single-node SSD](docs/STRIX_HALO.md#deepseek-v41-flash), [TCP/RoCE](docs/CLUSTERING_ROCM.md#measured-performance).
+
+
 ### CUDA SSD Streaming
 
 V4.1 CUDA starts with text inference on a single DGX Spark. Test on

README.md

@@ -7,7 +7,7 @@ language models on consumer hardware (that is, hardware that people
 can actually own). To reach this goal, we are building
 a small native inference engine optimized first for
 **DeepSeek V4 Flash** (including the experimental vision model),
-**DeepSeek V4.1 Flash** (Metal, and text inference on CUDA),
+**DeepSeek V4.1 Flash** (Metal, CUDA text inference, and ROCm on Strix Halo),
 and additionally **GLM 5.2 and 5.3**, **GLM 5.3 Flash** and
 **DeepSeek V4 PRO**, and **Qwen3.8 Flash Next** (Metal and CUDA). The code is self-contained and
 deliberately narrow, not a general GGUF runner: you need to use the
@@ -165,10 +165,9 @@ requirements. DeepSeek Vision Experimental uses a different checkpoint from
 Flash 0731; GLM 5.3 Flash and Qwen3.8 Flash Next add vision to the same text
 model through a separate encoder.
 
-DeepSeek V4.1 Flash text and vision run on Metal; text also runs on a DGX Spark.
-Q2 runs with SSD streaming on one 128 GB Mac or Spark, or resident across two
-Macs or two Sparks using RDMA. Q4 needs SSD streaming or a 512 GB Mac.
-Engram tables remain on disk in every mode, so use a fast
+DeepSeek V4.1 Flash text and vision run on Metal and ROCm on Strix Halo (`gfx1151`); CUDA supports text on DGX Spark. ROCm Q2 supports resident weights, SSD streaming and two-machine resident clustering. See [ROCm qualification](QA_BEFORE_RELEASES.md#deepseek-v41-flash-rocmgfx1151) and [clustering setup](docs/CLUSTERING_ROCM.md).
+
+Q2 runs with SSD streaming on one 128 GB Mac or Spark, or resident across two Macs or two Sparks using RDMA. Q4 needs SSD streaming or a 512 GB Mac. Engram tables remain on disk in every mode, so use a fast
 local SSD. See the [model guide](docs/MODELS.md#deepseek-v41-flash) for downloads
 and setup.
 

cuda/mmq/ds4_mmq.cu

@@ -26,6 +26,29 @@
 #include "mmid.cuh"
 #include "ds4_mmq_d2r.cuh"
 
+#if defined(GGML_USE_HIP)
+// One small producer on the existing stream. Thread0 scans384 bounds;128
+// threads fill all575 slots and then their disjoint active expert records.
+static __global__ void ds4_mmq_compact_produce(const int32_t *bounds, ds4_mmq_compact::List *out) {
+    __shared__ uint32_t offsets[ds4_mmq_compact::experts + 1];
+    __shared__ uint32_t status;
+    const uint32_t t = threadIdx.x;
+    for (uint32_t i = t; i < ds4_mmq_compact::capacity; i += blockDim.x)
+        out->entries[i] = {UINT32_MAX, UINT32_MAX};
+    if (t == 0) {
+        status = ds4_mmq_compact::prefix(bounds, offsets);
+        out->status = status;
+        out->count = status ? 0 : offsets[ds4_mmq_compact::experts];
+    }
+    __syncthreads();
+    if (status) return;
+    for (uint32_t e = t; e < ds4_mmq_compact::experts; e += blockDim.x)
+        for (uint32_t i = offsets[e]; i < offsets[e + 1]; i++)
+            out->entries[i] = {e, i - offsets[e]};
+}
+
+#endif
+
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
@@ -1337,6 +1360,11 @@ int ds4_mmq_moe_pair_impl(
     const int64_t s01          = (int64_t)K / blck;
     const int64_t s02          = (int64_t)M * s01;
 
+#if defined(GGML_USE_HIP)
+    // Declared first so the list remains owned through both projections.
+    ggml_cuda_pool_alloc<ds4_mmq_compact::List> compact_alloc;
+    ds4_mmq_compact::List *compact_list = nullptr;
+#endif
     ggml_cuda_pool_alloc<int32_t> ids_src1_alloc;
     ggml_cuda_pool_alloc<int32_t> ids_dst_alloc;
     ggml_cuda_pool_alloc<int32_t> expert_bounds_alloc;
@@ -1472,6 +1500,23 @@ int ds4_mmq_moe_pair_impl(
         ? (int64_t)n_tokens
         : ne_get_rows;
 
+#if defined(GGML_USE_HIP)
+    const bool compact_eligible =
+        type == GGML_TYPE_IQ2_XXS && cc == GGML_CUDA_CC_OFFSET_AMD + 0x1151 &&
+        n_tokens == 2048 && n_experts == 384 && n_expert_used == 6 && K == 5120 && M == 2304 &&
+        !direct_gateup_q8 && !fused_down && !persistent_pair_maps && !xa_soa && !xb_soa &&
+        !use_stream_k && stream == nullptr && routed_ncols_max == 2048 &&
+        get_mmq_y_host(cc) == 64 && get_mmq_x_max_host(cc) == 64 &&
+        ggml_cuda_info().devices[dev].warp_size == 32 && mmq_get_nwarps_host(cc,32) == 4;
+    if (compact_eligible) {
+        compact_list = compact_alloc.alloc(ctx->pool(),1);
+        ds4_mmq_compact_produce<<<1,128,0,stream>>>(expert_bounds,compact_list);
+        err = cudaGetLastError();
+        if (err != cudaSuccess) return -102;
+    }
+
+#endif
+
     /* The materialized path stream-frees gate/up Q8_1 before allocating the
      * down Q8_1. The direct path needs both simultaneously, but writes down
      * Q8_1 into caller-owned gate scratch instead of growing the CUDA pool. */
@@ -1722,6 +1767,10 @@ int ds4_mmq_moe_pair_impl(
         /*x_soa=*/xa_soa,
         /*soa_blocks=*/soa_blocks,
     };
+#if defined(GGML_USE_HIP)
+    args.compact_list = compact_list;
+#endif
+
 
     {
         ds4_mmq_nvtx_scope stage(