llama.cpp-Strix-Halo-DFlash2-Windows-HIP

llama.cpp DFlash2 for Strix Halo — Windows HIP community build

This repository provides an unofficial, reproducible Windows x64 HIP build of llama-server.exe for AMD Strix Halo (gfx1151) with upstream DFlash/DFlash2 support and a one-file correction for silent long-prompt corruption.

It is not an official build from AMD, the llama.cpp maintainers, Hugging Face, or RUBI. It was built and validated on one Ryzen AI MAX+ 395 system. Read the validation boundary before treating it as production-ready.

Why this build exists

Current llama.cpp supports draft-dflash, but its HIP integrated-GPU host-buffer path can produce plausible-looking wrong logits on gfx1151 when a prompt is split across micro-batches. The defect is documented in ggml-org/llama.cpp#28211.

This build applies the narrowly scoped fix already present in public llama.cpp commit 865374b: preserve integrated-device classification while refusing direct HIP host-buffer compute specifically on gfx1151.

Download and identity

• Binary: bin/llama-server.exe
• Size: 86,799,360 bytes
• SHA-256: 6587aecb1f636954a6e1d787a896cc461fac343797a1fe9ab07a098ca011b3aa
• Binary version: 0.4.0-dev (build 1, commit 465e49b)
• Source: llama.cpp commit 465e49b9cea78a68b9c244ffb48d0ee24a82873d, tag b10830
• Patch: patches/gfx1151-disable-direct-host-buft.patch

Verify before running:

Get-FileHash .\bin\llama-server.exe -Algorithm SHA256

The executable is not code-signed. A matching checksum proves it matches this upload; it does not substitute for reviewing and rebuilding the source.

Runtime requirements

• Windows 11 x64
• AMD Strix Halo / Radeon 8060S (gfx1151)
• A compatible AMD display driver
• TheRock/ROCm Windows runtime on PATH
• Microsoft Visual C++ runtime

The validated machine used:

• Ryzen AI MAX+ 395 with Radeon 8060S
• Windows 11 Pro build 26200
• AMD driver 32.0.31041.1004
• TheRock gfx1151-7.13.0
• HIP 7.13.99004-3309c6114a

ROCm DLLs are deliberately not included. The executable dynamically imports amdhip64_7.dll and hipblas.dll; install them through an appropriate TheRock distribution and keep one coherent runtime version. See the official llama.cpp Windows HIP guide.

Running DFlash2

The target and draft GGUF files are not included. DFlash drafts are trained for a specific target family; do not assume an unrelated draft is compatible merely because it loads.

.\scripts\run_dflash2_example.ps1 `
  -ServerExe .\bin\llama-server.exe `
  -TargetModel D:\Models\Target.gguf `
  -DraftModel D:\Models\Target-DFlash2-Q8_0.gguf `
  -RocmBin C:\TheRock\bin `
  -Port 8080

The example uses p_min=0.90 and n_max=7. For a DFlash2 checkpoint with a trained block size of eight, requesting n_max=8 is clamped by llama.cpp to seven draft tokens, so seven states the effective setting directly.

Verify the live surface:

Invoke-RestMethod http://127.0.0.1:8080/health
Invoke-RestMethod http://127.0.0.1:8080/v1/models
(Invoke-WebRequest http://127.0.0.1:8080/metrics).Content

Measured results

All rows below use the same Qwen3.8-27B UD-Q4_K_XL target weights. The old HIP and fixed HIP rows used the locally available Q8 DFlash2 draft where identified.

Additional fixed-build evidence:

• On the actual 27B target, three 6,974-token positional probes—early, middle, and late—were all exact with the fixed build.
• In a matched A/B using Qwen3.5-0.8B Q4 weights, the first-pass HIP binary failed the same early and middle probes and passed only the late probe; the fixed HIP binary passed all three.
• WikiText sanity test, changing only micro-batch size: ubatch=512 → PPL 17.1071; ubatch=2048 → PPL 17.1097 (0.015% difference).
• Aggregate DFlash validation: 41,905 drafted tokens, 40,906 accepted (97.616% acceptance), across 6,969 verification steps.
• Aggregate generation: 48,736 predicted tokens in 2,038.06 generation seconds (23.913 tok/s).
• Aggregate prompt processing: 61,099 tokens in 194.812 seconds (313.63 tok/s).

The public evidence/ directory contains the exact first-pass source shims and a machine-readable, sanitized result summary. Raw private battery outputs are deliberately excluded because publishing answer text would contaminate future evaluation. The complete two-pass investigation is in docs/DEBUGGING_WORKUP.md, and test interpretation is in docs/VALIDATION.md.

Rebuilding

The public build script takes explicit paths and refuses the wrong source commit:

git clone https://github.com/ggml-org/llama.cpp.git
git -C .\llama.cpp checkout 465e49b9cea78a68b9c244ffb48d0ee24a82873d

.\scripts\build_llama_windows_gfx1151fix.ps1 `
  -SourceDir .\llama.cpp `
  -ToolchainDir C:\TheRock `
  -CMakeExe C:\Tools\CMake\bin\cmake.exe `
  -NinjaExe C:\Tools\Ninja\ninja.exe `
  -GitExe "C:\Program Files\Git\cmd\git.exe" `
  -RcExe "C:\Program Files (x86)\Windows Kits\10\bin\10.0.26100.0\x64\rc.exe"

The validated build used CMake 4.4.3, Ninja 1.13.2, Git 2.55.0.windows.3, Clang 23.0.0, and Windows SDK resource compiler 10.0.26100.0. Build flags are documented in docs/BUILD_AND_RUNTIME.md.

Validation boundary

Verified:

• Windows 11 on one gfx1151 Radeon 8060S machine
• Full-GPU Qwen3.8-27B target plus Q8 DFlash2 draft
• 32,768 configured context
• Long prompts beyond the 512-token micro-batch boundary
• Consumer-surface OpenAI-compatible requests and DFlash metrics

Not established:

• Other AMD architectures or driver/runtime combinations
• Multi-user or --parallel > 1 production soak behavior
• 262,144-token context stability
• A complete rerun of the 38-question main battery on the fixed build
• General ladder improvement: rung 9 is one seed/configuration, not proof of a causal quality increase from the backend patch

The correction fixes a demonstrated HIP data-placement defect. It should not be described as improving model intelligence, DFlash acceptance, or every form of long-context behavior.

Attribution and disclosure

llama.cpp is MIT licensed; its license is included. The gfx1151 patch is derived from the public llama.cpp commit linked above. The build, investigation, testing, and documentation were produced collaboratively by Shane (shanecol), the RUBI fleet, and OpenAI Codex. Public claims are summarized under evidence/; the raw private battery records were retained by the project and are not redistributed.