Models continue to improve in quality for a similar memory size. Sometimes, though, one box just isn’t enough room for those that want to run higher quality/intelligent models. DeepSeek V4 Flash 0731 at Q4_K is a 153 GiB single-file model – and without SSD streaming, that 128 GB Strix Halo won’t fit it – even assuming it somehow fits, the Strix Halo just doesn’t have the memory bandwidth to generate tokens very quickly. Naturally, this begs the question – What happens if you add one more Strix Halo for 256 GB URAM to run higher-end models?
For starters, we have shared a few guides on running pipeline parallelism, but having spent the past few weeks tinkering with VLLM and other methods, we believe Tensor Parallelism is the sequential step to squeeze more juice out of the same box(es).
What is Tensor Parallelism? Simply put, think of it as two Strix Halo can proceed layers at the same time to give you closer to double the output on prefill (prompt processing) and decode (tokens / second). As a usual note – good things come in pairs (2,4,8…etc. TP doesn’t really support odd number machines for now).
The human parts end here, but we did manage to achieve: ~280 t/s prefill, ~19-21 t/s decode, and 800,000 tokens of context. Granted, it’s not great but this is certainly usable territory considering you are running a very high quality version of Deepseek V4 Flash 0731 (likewise, we suspect with more tweaking you can run the near lossless Q8 as well!)
A big favorite for the local hosting community’s engine remains ds4, antirez’s purpose-built DeepSeek V4 Flash inference engine, and what we’re running here is the wkljohn/ds4-strix-halo-tp-odinlink fork, which adds two-node tensor parallelism over RoCE v2 (or OdinLink, if you’d rather use USB4 – we didn’t). This is the complete, reproducible recipe: every step, every gotcha, and the benchmarks we measured. (or if you want to DIY, go ahead – you may need some technical knowledge) – Enjoy!
What you need
- 2× Strix Halo mini-PCs (MS-S1 MAX class) with the Ryzen AI MAX+ 395 and 128 GiB RAM each.
- 2× Mellanox ConnectX-4 Lx 25 GbE NICs plus a direct-attach cable. (We got them from Amazon, but any NIC (10Gbe, or 50, or even 100Gbe! will likely work as long as they support RDMA)
- Ubuntu (any recent release; the two boxes don’t have to run the same kernel – that matters later).
- podman and the container image
kyuz0/strix-halo-ds4-toolbox:multi-node-rocm-7.14(ROCm 7.14 for the gfx1151 iGPU). - The model: DeepSeek V4 Flash 0731, Q4_K – a 153.32 GiB single-file GGUF from antirez’s
deepseek-v4-gguf.
Two networks are used for our case – one is the LAN for network serving, and the second is for the 25Gbe for TP=2: 192.x.x.x is our existing 10 GbE LAN for management, SSH, and internet; 10.0.1.0/24 is the point-to-point 25 GbE data path (10.0.1.1 = coordinator, 10.0.1.2 = worker). If you’ve followed our single-box ds4 recipe, the BIOS and kernel work will look familiar – but read Step 1 anyway, because the two-node build really should use as much URAM as available to fit models and maximum context.
The recipe
Six steps, both nodes, in order. The two roles in this dance are the coordinator (rank 0 – it runs the API server) and the worker (rank 1 – it holds the other half of the model).
Step 1 – Fix the BIOS carve-out
Many vendors ship Strix Halo with a huge static GPU carve-out. One of our boxes had UMA Frame Buffer Size = 96 GB, which meant the OS could see only 30 GiB of RAM. In BIOS (Advanced -> AMD CBS -> NBIO -> GFX Configuration):
- UMA Frame Buffer Size -> Auto (or <= 1 GB – with Auto, the GPU dynamically claims GTT memory instead)
- IOMMU -> Off
After boot, verify the OS sees ~122 GiB:
free -g # want ~122 total
cat /sys/class/drm/card*/device/mem_info_vram_totalDo not skip this. With a 96 GiB carve-out the OS can’t run an 80 GiB model shard; with Auto and a large GTT it can. The same total RAM, two completely different outcomes.
Step 2 – Kernel parameters (GTT)
Append to GRUB_CMDLINE_LINUX_DEFAULT in /etc/default/grub:
amd_iommu=off amdgpu.gttsize=126976 ttm.pages_limit=32505856 ttm.page_pool_size=32505856Then reboot and verify:
cat /proc/cmdline
cat /sys/module/ttm/parameters/pages_limit # 32505856This lets the GPU claim up to ~124 GB of system RAM as GTT. The model shard (80.8 GiB per rank) plus KV cache has to fit in it.
Step 3 – RDMA over 25 GbE
Give each 25G port a static IP on the private subnet (netplan, one file per node):
# /etc/netplan/99-25g-rdma.yaml (coordinator; use .2 on the worker)
network:
version: 2
ethernets:
enp197s0f1np1: # your ConnectX-4 Lx port
addresses: [10.0.1.1/24]sudo netplan applyInstall the userspace RDMA tools:
sudo apt-get install -y ibverbs-providers ibverbs-utils rdma-core sshpassThen verify the link at line rate on both nodes:
ibv_devices # note the device name (this matters - see the gotcha section)
ibv_devinfo -d <device> # port ACTIVE, link layer Ethernet
ib_write_bw -d <device> --report_gbits -x 3 # GID index 3 = RoCE v2We measured 23.15 Gb/s with ib_write_bw – close enough to the 25G nominal rate. Now the classic RDMA killer: memlock. Model registration pins gigabytes of memory, and the default limit is 8192 KB. Create /etc/security/limits.d/99-rdma-unlimited.conf:
* soft memlock unlimited
* hard memlock unlimitedThe trap: this file only applies to fresh PAM login sessions. Your current desktop session keeps the old limit no matter what the file says. When launching from a script or desktop app, go through a fresh session and check:
echo '<sudo-pass>' | sudo -S -u "$USER" bash -c 'ulimit -l' # -> unlimitedStep 4 – Containers and build
Everything runs inside containers so the host stays clean. Pull the base toolchain image, then derive a build image with the compiler and RDMA headers:
podman pull docker.io/kyuz0/strix-halo-ds4-toolbox:multi-node-rocm-7.14
podman build -t ds4-build2:rocm714 - <<'EOF'
FROM kyuz0/strix-halo-ds4-toolbox:multi-node-rocm-7.14
RUN microdnf -y --nodocs --setopt=install_weak_deps=0 install \
amdrocm-core-devel7.14-gfx1151 amdrocm-core7.14 rdma-core-devel \
make diffutils && microdnf clean all
EOFClone the fork and stage the model at identical absolute paths on both nodes – the two filesystems are independent, and the engine expects everything in the same place:
sudo mkdir -p /srv/ds4-tp /srv/ds4-models
sudo git clone https://github.com/wkljohn/ds4-strix-halo-tp-odinlink /srv/ds4-tp
# copy the 153 GiB GGUF over - at ~1.37 GB/s across the 25G link, that's about 2 minutesBuild inside the container with /srv/ds4-tp mounted at /work. The container’s ld.lld links PIE by default, so the objects need -fPIC:
podman run --rm -v /srv/ds4-tp:/work -w /work localhost/ds4-build2:rocm714 bash -c '
make -B ds4 ds4-server \
CORE_OBJS="ds4.o ds4_distributed.o ds4_tp.o ds4_ssd.o ds4_rocm.o ds4_rocm_compat.o ds4_rocm_unavailable.o ds4_layer_pack.o ds4_glm5_kda.o ds4_glm5_next_runtime.o ds4_glm5_next_state.o ds4_glm5_next_exec.o" \
CFLAGS="-O3 -ffast-math -fPIC -Wall -Wextra -std=c99 -D_GNU_SOURCE -fno-finite-math-only -DDS4_ROCM_BUILD -DDS4_ROCM_TP_READY=1" \
DS4_LINK="/opt/rocm/bin/hipcc -O3 -ffast-math -g -fno-finite-math-only -pthread -D__HIP_PLATFORM_AMD__ -Wno-unused-command-line-argument --offload-arch=gfx1151 -mno-wavefrontsize64 -DDS4_GFX1151_WAVE32=1 " \
DS4_LINK_LIBS="-L/opt/rocm/lib -Wl,-rpath,/opt/rocm/lib -lm -pthread -lhipblas -lhipblaslt"'Then compare binary hashes across nodes before every run – the fork refuses to benchmark on mismatched builds:
sha256sum /srv/ds4-tp/ds4 # must match on both nodesStep 5 – Launch the cluster
GPU access needs the render and video groups (sudo usermod -aG render,video $USER, then re-login). Launch the worker first (it retries until the coordinator appears), then the coordinator – both from sessions with unlimited memlock:
# ---- worker (node2) ----
podman run --rm --network=host --device=/dev/kfd --device=/dev/dri \
--device=/dev/infiniband -v /srv/ds4-tp:/work -v /srv/ds4-models:/srv/ds4-models:ro \
-w /work localhost/ds4-build2:rocm714 ./ds4 \
-m /srv/ds4-models/ds4-0731.gguf --rocm --tensor-parallel \
--role worker --coordinator 10.0.1.1 5599 --transport rdma \
--rdma-device rocep101s0f0 --rdma-gid-index 3 -c 800000# ---- coordinator (node1) - ds4-server for the OpenAI-compatible API ----
podman run --rm --network=host --device=/dev/kfd --device=/dev/dri \
--device=/dev/infiniband -v /srv/ds4-tp:/work -v /srv/ds4-models:/srv/ds4-models:ro \
-w /work localhost/ds4-build2:rocm714 ./ds4-server \
-m /srv/ds4-models/ds4-0731.gguf --rocm --tensor-parallel \
--role coordinator --listen 10.0.1.1 5599 --transport rdma \
--rdma-device mlx5_1 --rdma-gid-index 3 -c 800000 \
--host 0.0.0.0 --port 8000Note the --rdma-device values differ between the two boxes (rocep101s0f0 vs mlx5_1) – that’s not a typo, and it’s the seeds of the story in the next section. Each rank loads its 80.8 GiB shard in about 40-60 seconds. When the coordinator prints listening on http://0.0.0.0:8000, the API is live.
Step 6 – The API
| Base URL | http://<node1-LAN-IP>:8000 |
| Endpoints | /v1/chat/completions – /v1/completions – /v1/responses – /v1/messages (Anthropic) – /v1/models |
| Model id | deepseek-v4-flash |
| Context | 800,000 tokens |
| Extras | streaming, tools, temperature/top_p/seed, reasoning_effort |
One thing to know before you plug it into anything: this is a reasoning model. The chain-of-thought arrives in reasoning_content and the answer in content – budget your max_tokens accordingly, or you’ll get answers that stop mid-thought.
The gotcha that cost us a day
Here is the debugging story we promised. Two identical mini-PCs, identical NICs, and yet the RDMA path only worked on one of them.
The cause was almost too boring to believe: our two boxes run different kernel versions, and the kernel names the same NIC differently. On the coordinator (kernel 7.0.0-30), the ConnectX-4 Lx shows up as mlx5_0 and mlx5_1. On the worker (kernel 7.0.0-22), the same hardware is rocep101s0f0 and rocep101s0f1.
That would be a harmless cosmetic difference – except ds4 identified Mellanox NICs by name prefix (mlx5_*). On the rocep* node, the mlx5 code path silently disabled, and three things went wrong at once: the slab was allocated as GPU device memory (unregistrable, so reg_mr: Bad address), and the full 270+ MiB slab got registered as one MR instead of three segments. No crash, no clear error – just a worker that couldn’t register its memory and a coordinator wondering where its peer went.
The fix is one paragraph of C, now in our tree of the fork (ds4_tp.c): identify Mellanox hardware by PCI vendor ID, not by name:
r->is_mlx5 = strncmp(name, "mlx5_", 5) == 0;
if (!r->is_mlx5 && !r->is_odinlink) {
char p[192]; snprintf(p, sizeof p,
"/sys/class/infiniband/%s/device/vendor", name);
FILE *vf = fopen(p, "r");
if (vf) { unsigned v; if (fscanf(vf, "%i", &v) == 1 && v == 0x15b3)
r->is_mlx5 = true; fclose(vf); }
}So if your cluster “works” on one node and throws reg_mr: Bad address on the other: check ibv_devices on both, and don’t trust the names.
Where the pieces come from
This cluster stands on three layers of other people’s work, and they all deserve credit. ds4 itself is the latest from antirez – the creator of Redis – who builds his Dwarfstar inference engines the way he built Redis: small, fast, few dependencies (source on GitHub, antirez’s blog). The two-node tensor-parallel work over RoCE v2 and OdinLink is wkljohn’s fork, ds4-strix-halo-tp-odinlink – including the benchmark harness behind the calibration table below, and the vendor-ID fix we’ve fed back into our tree of it. And the container packaging that keeps the ROCm 7.14 toolchain off your host is kyuz0’s strix-halo-ds4-toolbox line, which our single-box recipe also runs on. One paragraph of C from us; everything else was already standing.
Hooking it up to your tools
Because ds4-server speaks both OpenAI-compatible and Anthropic-compatible APIs on port 8000, it drops into agentic harnesses with no template wrangling – the same dual compat that sold us on the single-box build:
- Claude Code – point it at
http://<node1-LAN-IP>:8000and it talks to V4 Flash through the Anthropic-compatible/v1/messagesroute. - Codex CLI and other OpenAI-native clients – the
/v1/responsesendpoint covers the Responses API, and/v1/chat/completionscovers everything else. - Any client that speaks either API – Aider, Cline, Roo Code, your own curl script – works out of the box.
One workflow note for agentic use: because the reasoning arrives separately in reasoning_content, harnesses that count it against max_tokens will burn budget on invisible thinking. If your agent seems to stop mid-thought, that’s the knob to turn, not the model.
The numbers
Our measured numbers – this exact cluster, Q4_K, balanced 50/50 expert split:
| RDMA link (ib_write_bw, RoCE v2) | 23.15 Gb/s |
| Model file copy node->node | ~1.37 GB/s (~2 min for 153 GiB) |
| Smoke test, ctx 4096, greedy | 31.9 t/s prefill / 19.6 t/s decode |
| Server, ctx 800k, ~71k-token prompt (chunked prefill) | ~195-218 t/s prefill |
| Server, ctx 800k, decode | 18.5-19 t/s sustained |
| Per-node memory @ 800k ctx | ~106 GiB (80.8 shard + 10.6 KV + 16.7 buffers) |
| Per-node memory @ ctx 4096 | ~81 GiB |
For calibration, the fork’s published reference numbers (TP=2, 2048-token prefill chunk, both nodes at 128 GB):
| Configuration | Prefill | Decode |
|---|---|---|
| Q4_K over RoCE v2 (ours) | 280.58 t/s | 21.30 t/s |
| Q4_K over OdinLink (USB4/TB5) | 233.04 t/s | 19.17 t/s |
| Q2_K over RoCE v2 | 209.60 t/s | 20.09 t/s |
| Original Q4_K baseline (pre-acceleration) | 34.11 t/s | 9.96 t/s |
Two caveats, stated plainly: the fork’s numbers come from its benchmark harness (long-form prompts with exact output-fingerprint checks), while our server numbers come from live API traffic at 800k context. The two aren’t directly comparable – and both are honest.
The memory line is the quiet star of the table. Context memory scales cheaply here: the raw KV window clamps at 8192 rows and the compressed cache is what grows, so 800k tokens cost only ~10.6 GiB per rank. That’s the economics that let a 153 GiB model serve 800k context on 128 GiB nodes – the same DeepSeek KV-cache architecture we unpacked in our single-box guide, and this cluster is that math working in your favor.
Running a single box? Start with our DeepSeek V4 Flash Q2 on Strix Halo 128GB recipe for the ds4 basics, and the two-node cluster guide for the hardware and network setup underneath this build.
