# Building llama.cpp with ROCm/HIP on Windows — Technical Reference > **For AI agents and advanced users.** The human-followable download/install steps live in > [SETUP_GUIDE.md](SETUP_GUIDE.md). **This document assumes those steps are done** and covers > checking the setup, configuring, building, verifying, running, and fixing — in detail. > > Every command is inline PowerShell / CMake; nothing depends on helper scripts. --- ## Workflow 1. **Inventory** — find ROCm, identify every GPU and its architecture, verify the toolchain. 2. **Configure** — one CMake command. 3. **Build** — with Ninja. 4. **Verify** — confirm the GPU is actually used. 5. **Handle multi-GPU systems** — if an iGPU is present, fix device visibility *before* running. 6. **Run the server** — and smoke-test it. 7. **Fix** — exact error signatures → cause → fix (troubleshooting table). Write down the findings from step 1; every later step depends on them. --- ## 1. Inventory (check before building) ### 1.1 Locate ROCm ```powershell Get-ChildItem "C:\Program Files\AMD\ROCm" -Directory | Sort-Object Name -Descending ``` Use the **newest** version present and record it: ```powershell $env:ROCM_ROOT = (Get-ChildItem "C:\Program Files\AMD\ROCm" -Directory | Sort-Object Name -Descending | Select-Object -First 1).FullName Write-Host "ROCM_ROOT = $env:ROCM_ROOT" ``` If multiple versions exist, note them all. Mixed versions are a common source of weird failures (stale `HIP_PATH*` env vars, the wrong runtime DLL winning on `PATH`, …). ### 1.2 Identify every GPU — the single most important check ```powershell & "$env:ROCM_ROOT\bin\hipInfo.exe" ``` For **each** listed device, record: `device#`, `Name:`, `gcnArchName:`, `isIntegrated:`. | device# | Name | gcnArchName | isIntegrated | | --- | --- | --- | --- | | 0 | e.g. `AMD Radeon(TM) Graphics` (iGPU) | e.g. `gfx1036` | 1 | | 1 | e.g. `AMD Radeon RX 9070 XT` (discrete) | e.g. `gfx1201` | 0 | Rules: - The build target is the **exact `gcnArchName` string** of the GPU you will run on (e.g. `gfx1201`). Never derive it from the GPU's marketing name, and don't use the `major:`/`minor:` numbers — `major: 12, minor: 0` does *not* mean `gfx1200`. - **More than one device with different architectures ⇒ high risk of the multi-GPU crash** (step 5). `isIntegrated: 1` = the CPU's iGPU. - If `hipInfo` lists nothing or errors, the AMD driver/ROCm isn't working — check the driver install and the antivirus exclusion from SETUP_GUIDE.md. ### 1.3 Verify the toolchain ```powershell git --version cmake --version ninja --version clang --version # must be ROCm's clang — look for "AMD" in the version banner ``` If `ninja` is missing, install it or fall back to `-G "Visual Studio 17 2022"` in step 2. ### 1.4 Check for stale environment variables ```powershell Get-ChildItem Env: | Where-Object Name -match "HIP|ROCM|HSA" ``` `HIP_PATH*` / `ROCM_PATH` values pointing at uninstalled ROCm versions can confuse auto-detection. Remove stale ones. --- ## 2. Configure Run in the **same PowerShell session** where `$env:ROCM_ROOT` was set in 1.1, and where `$env:GPU_TARGET` is set to the discrete GPU's `gcnArchName` from 1.2: ```powershell $env:GPU_TARGET = "" $rocmShort = "$env:ROCM_ROOT".Replace("C:\Program Files", "C:/PROGRA~1") cmake -S . -B build -G Ninja ` -DGGML_HIP=ON ` "-DGPU_TARGETS=$env:GPU_TARGET" ` -DCMAKE_BUILD_TYPE=Release ` "-DCMAKE_PREFIX_PATH=$env:ROCM_ROOT" ` "-DCMAKE_C_COMPILER=$env:ROCM_ROOT\bin\clang.exe" ` "-DCMAKE_CXX_COMPILER=$env:ROCM_ROOT\bin\clang++.exe" ` "-DCMAKE_CXX_FLAGS=--rocm-device-lib-path=$rocmShort/amdgcn/bitcode" ``` Flag reference: - `-DGGML_HIP=ON` — enable the ROCm backend. Without it: CPU-only build. - `-DGPU_TARGETS=` — compile kernels for the given architecture(s); ggml maps it to HIP's architecture list. **Wrong value = runtime "no binary for GPU" crashes.** - `-DCMAKE_C_COMPILER` / `-DCMAKE_CXX_COMPILER` — on Windows, HIP code must be compiled with ROCm's bundled `clang`/`clang++`, not MSVC. - `-DCMAKE_CXX_FLAGS=--rocm-device-lib-path=...` — tells clang where the AMD device bitcode lives; the `C:/PROGRA~1/...` short path avoids Windows splitting the argument at the space in "Program Files". **Changing `-DGPU_TARGETS` later?** Re-run cmake with the new value, or delete `build/` first — the old value is cached and silently reused otherwise. --- ## 3. Build ```powershell cmake --build build -j %NUMBER_OF_PROCESSORS% ``` - Warnings about `__declspec(dllimport)` / `GNU extension` are harmless. - If **every** `.cu` file fails with the math-function error (`error: __device__ function 'isgreater' cannot overload ...`), see [step 7](#7-rocm-wrapper-bug-fix-build-failure-isgreater). --- ## 4. Verify the GPU is actually used ```powershell build\bin\llama-cli.exe -m -ngl 99 -p "Hello" --no-display-prompt ``` Expected: `llama_model_load: offloaded NN/NN layers to GPU`. If output shows `Backend: CPU` or no "offloaded" line, re-check step 2 (`-DGGML_HIP=ON`, `-DGPU_TARGETS`). --- ## 5. Multi-GPU systems (iGPU + discrete GPU) ⚠️ **This is the #1 cause of "build succeeded but the server crashes instantly".** ### Symptom The server (or `llama-cli`) starts, then dies within a second or two: ``` I cmn init: llama threadpool init, n_threads = 8 E ROCm error: device kernel image is invalid E current device: 1, in function ggml_cuda_kernel_launch at ...\common.cuh:1667 ``` …or the variant `E ROCm error: invalid kernel file`. ### Cause The CPU's iGPU (reported by Windows as "AMD Radeon(TM) Graphics") and the discrete AMD GPU are **both visible to HIP** and have **different architectures** (e.g. `gfx1036` iGPU next to `gfx1201` discrete). The build contains kernels for **one** architecture only, and on some driver/runtime combinations the discrete GPU cannot even launch its own kernels while the iGPU is present. **Confirmation:** `hipInfo.exe` (1.2) lists more than one AMD GPU ⇒ this section applies. A healthy single-GPU machine lists exactly one device and never sees this error. ### Options (pick one) #### Option A — Hide everything except the real GPU (recommended; no rebuild) Set `HIP_VISIBLE_DEVICES` to the discrete GPU's index from 1.2: ```powershell # per session (before starting the server) $env:HIP_VISIBLE_DEVICES = "" # OR permanent (applies to NEW terminals only) setx HIP_VISIBLE_DEVICES ``` Verify: `hipInfo.exe` now lists only the discrete GPU. ⚠️ If the iGPU is later disabled in the BIOS (Option B), the discrete GPU usually becomes index `0` — update or remove the variable, or HIP will see no devices. #### Option B — Disable the iGPU in the BIOS (cleanest long-term) Restart → enter BIOS/UEFI (`Del`/`F2`) → find the integrated-graphics setting (iGPU / Integrated Graphics / Onboard Graphics / IGD) → set to **Disabled** → save & exit. The iGPU disappears from Windows entirely, so no env var is needed and the enumeration can't change. Bonus: frees a slice of reserved system RAM. #### Option C — Build kernels for every GPU If both GPUs must be usable, pass multiple targets (quote the value in PowerShell): ```powershell "-DGPU_TARGETS=;" # e.g. gfx1201;gfx1036 ``` ⚠️ **This does not always fix the crash.** In the interfering-devices scenario the discrete GPU fails even on its own architecture's kernels while the iGPU is present. If it still crashes, use Option A or B. Costs longer builds and bigger binaries. #### Option D — Unify the AMD driver versions The iGPU and discrete GPU can end up on different driver generations (compare in Task Manager → Performance → GPU). Installing the newest Adrenalin driver, which covers both, *sometimes* resolves the multi-GPU launch failure. Not guaranteed alone; combine with A or B. ### After any option: verify ```powershell build\bin\llama-cli.exe --list-devices ``` Exactly one `ROCm` device (the discrete GPU) should be listed. Then re-run the step-4 test. --- ## 6. Run the server ```powershell build\bin\llama-server.exe -m -ngl 99 -c --port 8080 ``` - `-ngl 99` offloads all layers to the GPU. If VRAM is insufficient it falls back to CPU (slow) — use a smaller `-c` or a smaller model. - Ready when the log shows `offloaded NN/NN layers to GPU` and `listening on http://127.0.0.1:`. Smoke test (OpenAI-compatible API): ```powershell curl http://localhost:8080/v1/chat/completions ` -H "Content-Type: application/json" ` -d "{\"model\":\"test\",\"messages\":[{\"role\":\"user\",\"content\":\"Hi\"}]}" ``` --- ## 7. ROCm wrapper bug fix (build failure: `isgreater`) The **only** case where the ROCm installation itself must be modified. A bug in ROCm's clang wrapper header only shows on Windows with the MSVC standard library: **every** `.cu` file fails identically with ``` error: __device__ function 'isgreater' cannot overload __host__ __device__ function 'isgreater' ... 20 errors generated when compiling for . ``` **Cause:** `\lib\clang\\include\__clang_hip_runtime_wrapper.h` includes `` **before** `__clang_cuda_math_forward_declares.h`. MSVC's `` defines the comparison functions as implicitly `__host__ __device__`, which then conflicts with HIP's `__device__`-only declarations. Linux (libstdc++/libc++) is unaffected. **Fix (once per ROCm install; edit one file, as Administrator):** 1. Find the header (path contains the clang version bundled with ROCm): `C:\Program Files\AMD\ROCm\\lib\clang\\include\__clang_hip_runtime_wrapper.h` 2. Back it up: copy to `__clang_hip_runtime_wrapper.h.orig` next to it. 3. Open it in an **elevated** editor (the file is under `C:\Program Files`). Inside the `#if !defined(__HIPCC_RTC__)` block, insert one include line **before** `#include `: ``` #if !defined(__HIPCC_RTC__) +#include <__clang_cuda_math_forward_declares.h> #include ``` (The later existing include of the same header becomes a no-op thanks to its include guard.) If the file doesn't match this layout, stop — newer ROCm versions may differ. 4. Save and rebuild (step 3). > Each ROCm version ships its own copy of the header — repeat after any ROCm upgrade. --- ## Troubleshooting | Symptom | Cause / fix | | --- | --- | | `ROCm error: device kernel image is invalid` or `invalid kernel file` at startup | Multi-GPU interference — [step 5](#5-multi-gpu-systems-igpu--discrete-gpu): set `HIP_VISIBLE_DEVICES` to the discrete GPU's index, or disable the iGPU in BIOS | | `hipErrorNoBinaryForGpu` at runtime | Built for the wrong architecture — set `-DGPU_TARGETS` to the exact `gcnArchName` from 1.2 | | Build OK but CPU-only output (no "offloaded … to GPU") | `GGML_HIP` off or wrong `GPU_TARGETS` — reconfigure with `-DGGML_HIP=ON -DGPU_TARGETS=` | | `__device__ function 'isgreater' cannot overload …` on every `.cu` file | ROCm wrapper include-order bug — [step 7](#7-rocm-wrapper-bug-fix-build-failure-isgreater) | | `ninja: error: … clang++.exe … FAILED` on the first CUDA file | ROCm/compiler integration issue — re-check the step-2 flags | | `argument unused during compilation: '--rocm-device-lib-path=…'` | Harmless warning on non-HIP files; keep the flag | | Server loads the model but answers very slowly | Model doesn't fit in VRAM → CPU fallback. Use a smaller model/quantization or smaller `-c` | | `hipInfo` still shows the iGPU after Option A/B | Variable not visible in the current terminal (`echo $env:HIP_VISIBLE_DEVICES`), or reboot after a BIOS change | | Weird behavior with multiple ROCm versions installed | Stale `HIP_PATH*`/`ROCM_PATH` env vars or wrong DLL on `PATH` — check 1.4 and `PATH` ordering | | `llama_server: invalid value for main_gpu: N (available devices: M)` | `--main-gpu` indexes llama.cpp's device list (discrete GPUs only), not HIP's raw indices | --- ## Appendix A — Environment variables | Variable | Purpose | Example | | --- | --- | --- | | `HIP_VISIBLE_DEVICES` | Restrict which GPUs HIP sees — **the step-5 fix** | `1` (discrete GPU only) | | `ROCM_PATH` | Points CMake/HIP tools at the ROCm root automatically | `C:\Program Files\AMD\ROCm\7.2` | | `HSA_OVERRIDE_GFX_VERSION` | Force a gfx version at runtime (rare, advanced) | `12.0.1` | | `HIP_LOG_LEVEL` | Verbose HIP runtime logging for debugging | `debug` | | `GGML_LOG_DISABLE` | Disable ggml logging | `0` | ## Appendix B — Building for multiple GPU architectures Pass multiple targets (Option C of step 5): ```powershell "-DGPU_TARGETS=;" ``` Longer builds, bigger binaries, but one build runs on any of those GPUs. Pick only what you need. ## Appendix C — Debug builds & runtime diagnostics Swap the build type for `RelWithDebInfo` (keeps debug symbols, still optimized): ```powershell cmake -S . -B build -G Ninja ` -DGGML_HIP=ON ` "-DGPU_TARGETS=$env:GPU_TARGET" ` -DCMAKE_BUILD_TYPE=RelWithDebInfo ` ... (rest of the step-2 flags) cmake --build build -j 16 ``` Check which backend is active: ```powershell build\bin\llama-cli.exe -m -ngl 99 -v -p "test" --no-display-prompt 2>&1 | findstr /i "backend gpu hip rocm" ``` Expect `Backend: ROCm` or `offloaded NN/NN layers to GPU`. ## Appendix D — How the pieces fit 1. **CMake** turns llama.cpp + the ROCm SDK into a build with a HIP backend. 2. `-DGPU_TARGETS` decides which GPU architectures the kernels can run on. 3. At runtime the **HIP driver** decides which GPUs exist; **llama.cpp** picks which to use. A visible GPU with no matching kernels = crash. 4. Most post-build crashes come from a GPU being *visible* but *not built for* — fix the visibility (step 5), not the build.