Add gfx1100 HIP support with tuned prefill and bounded expert uploads

README.md

@@ -81,6 +81,9 @@ one later with `SETUP.bat` (the same as `START-HERE.bat --setup`; on Linux `./se
 For **OrcaRouter's Flash-Next Uncensored IQ3_XXS**, see the [manual compatibility setup](docs/ORCA.md).
 It needs an explicit packing conversion and is not an installer menu option.
 
+For an **experimental RX 7900 XTX (gfx1100) Linux source build**, see the
+[AMD HIP instructions](docs/AMD_HIP.md). The automatic installer below remains NVIDIA-only.
+
 ## Install
 
 **You need:** an NVIDIA RTX 30, 40 or 50 card with 12 GB of VRAM or more, enough RAM for the size you pick (above),

cmake/hip_backend.cmake

@@ -0,0 +1,58 @@
+# Opt-in HIP configuration. Strata's CUDA-shaped kernels currently target RDNA3 gfx1100 wave32.
+# CMake/compiler discovery stays machine-independent; pass CMAKE_HIP_COMPILER when it is not on PATH.
+if(NOT DEFINED CMAKE_HIP_ARCHITECTURES OR CMAKE_HIP_ARCHITECTURES STREQUAL "")
+  set(CMAKE_HIP_ARCHITECTURES gfx1100 CACHE STRING "Strata HIP target architecture")
+endif()
+if(NOT CMAKE_HIP_ARCHITECTURES STREQUAL "gfx1100")
+  message(FATAL_ERROR
+    "Strata HIP currently supports only gfx1100 wave32; CMAKE_HIP_ARCHITECTURES is '${CMAKE_HIP_ARCHITECTURES}'")
+endif()
+
+enable_language(HIP)
+find_package(hip CONFIG REQUIRED)
+find_package(hipblas CONFIG REQUIRED)
+find_package(hipblaslt CONFIG QUIET)
+
+if(NOT TARGET hip::host)
+  message(FATAL_ERROR "The ROCm hip CMake package did not provide hip::host")
+endif()
+if(NOT TARGET roc::hipblas)
+  message(FATAL_ERROR "The ROCm hipblas CMake package did not provide roc::hipblas")
+endif()
+if(TARGET roc::hipblaslt)
+  set(STRATA_HIPBLASLT_AVAILABLE ON)
+else()
+  set(STRATA_HIPBLASLT_AVAILABLE OFF)
+  message(STATUS "Strata: hipBLASLt not found; solution-table dispatch is unavailable")
+endif()
+
+# HIP's link step produces a PIE; make Strata and ggml objects PIC for the ROCm linker.
+set(CMAKE_POSITION_INDEPENDENT_CODE ON)
+
+set(STRATA_HIP_COMPAT_INCLUDE_DIR "${CMAKE_CURRENT_SOURCE_DIR}/include/strata/hip_compat")
+add_library(strata_hip_runtime INTERFACE)
+target_include_directories(strata_hip_runtime BEFORE INTERFACE
+  "${STRATA_HIP_COMPAT_INCLUDE_DIR}" "${CMAKE_CURRENT_SOURCE_DIR}/include")
+target_compile_definitions(strata_hip_runtime INTERFACE STRATA_USE_HIP=1)
+target_link_libraries(strata_hip_runtime INTERFACE hip::host)
+foreach(_language IN ITEMS CXX HIP)
+  target_compile_options(strata_hip_runtime INTERFACE
+    "$<$<COMPILE_LANGUAGE:${_language}>:-include>"
+    "$<$<COMPILE_LANGUAGE:${_language}>:${STRATA_HIP_COMPAT_INCLUDE_DIR}/cuda_runtime.h>")
+endforeach()
+
+# CMake does not infer HIP from Strata's existing CUDA-shaped .cu suffixes.
+file(GLOB_RECURSE _strata_hip_sources CONFIGURE_DEPENDS
+  "${CMAKE_CURRENT_SOURCE_DIR}/src/*.cu"
+  "${CMAKE_CURRENT_SOURCE_DIR}/bench/*.cu"
+  "${CMAKE_CURRENT_SOURCE_DIR}/tests/*.cu")
+if(_strata_hip_sources)
+  set_source_files_properties(${_strata_hip_sources} PROPERTIES LANGUAGE HIP)
+endif()
+foreach(_source IN ITEMS tests/hip/intrinsics.cpp tests/hip/native_qsa_score.cpp)
+  if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${_source}")
+    set_source_files_properties("${_source}" PROPERTIES LANGUAGE HIP)
+  endif()
+endforeach()
+
+message(STATUS "Strata: HIP enabled, arch ${CMAKE_HIP_ARCHITECTURES}")

docs/AMD_HIP.md

@@ -0,0 +1,138 @@
+# Experimental AMD HIP backend (gfx1100)
+
+This is a manual Linux source build for the RX 7900 XTX. It is opt-in; the
+NVIDIA installer and CUDA build remain the default. Other AMD architectures,
+wave64, Windows HIP, and mixed AMD/NVIDIA execution are outside this contribution.
+
+The backend maps the CUDA-shaped runtime and BLAS calls to HIP/hipBLAS, uses
+RDNA3's signed integer dot instruction for quantized kernels, and supplies
+wave32 shuffle/packed-byte operations. CUDA-only QSA matrix instructions have
+an ordered FP32 fallback. Prefill supports both dequantization plus hipBLAS GEMM and opt-in HIP ggml MMQ.
+An optional, calibrated hipBLASLt path accelerates dense projections on gfx1100.
+This does not claim bit-identical model answers across backends. See
+[performance settings and evidence](AMD_HIP_PERFORMANCE.md).
+
+## Build
+
+Requirements: a working ROCm driver/runtime, HIP development headers and
+compiler, hipBLAS and (for tuned dense prefill) hipBLASLt development files, CMake 3.24+, a C++20 host compiler, and Git.
+The model still needs sufficient system RAM and fast SSD storage for its PLE
+table. VRAM occupancy alone is not a throughput measurement.
+
+```sh
+cmake -S . -B build-hip \
+  -DCMAKE_BUILD_TYPE=Release \
+  -DSTRATA_ENABLE_HIP=ON -DSTRATA_ENABLE_CUDA=OFF \
+  -DCMAKE_HIP_ARCHITECTURES=gfx1100
+cmake --build build-hip --target strata -j2
+```
+
+If CMake cannot find the HIP compiler, add
+`-DCMAKE_HIP_COMPILER=/path/to/rocm/llvm/bin/clang++`. On the Fedora-family test
+host this was `/usr/lib64/rocm/llvm/bin/clang++`. Use the compiler and libraries
+from the same ROCm installation. Do not enable both GPU backends in one build.
+
+Native IQ experts are enabled by default. CMake fetches the llama.cpp revision
+pinned by this repository. For an offline build, point `STRATA_GGML_DIR` at a
+checkout of that exact revision; using an arbitrary newer checkout changes the
+dependency being tested.
+
+## Model and serving configuration
+
+Prepare a supported model pack and its MTP runtime using the existing tools.
+For OrcaRouter IQ3_XXS, follow [the explicit compatibility conversion](ORCA.md);
+this backend does not change quantization, model licenses, or tokenizers.
+Add `--experts-bin` to the `tools/iq_pack.py` command: mmap requires the pack's
+`experts.bin`, which the default native packing command does not emit. This
+consumes additional disk space. Original GSQ-RCO IQ3_XXS uses shard 2 for PLE;
+Orca uses shard 1. Keep each model's own pack and tokenizer together.
+
+Use `build-hip/strata` as the executable in the server JSON. Select the AMD
+device with `ROCR_VISIBLE_DEVICES`/`HIP_VISIBLE_DEVICES` if necessary. Start with
+a modest context and prefill chunk size before measuring larger workloads.
+Keep the PLE file on an SSD. Do not assume a configured context length proves
+successful full-window inference.
+
+Large pinned host allocations can fail on ROCm even when ordinary RAM is
+available. The existing `--mmap-experts` path avoids allocating the full pinned
+expert arena; it still depends on OS file-cache residency and may stall on
+storage reads. It does not make SSD access equivalent to RAM.
+
+Starting args for the original GSQ-RCO IQ3_XXS model (replace the paths):
+
+```sh
+build-hip/strata --serve \
+  --pack packs/iq3xxs \
+  --native /path/to/Qwen3.8-Flash-Next-GSQ-RCO-IQ3_XXS-00001-of-00002.gguf \
+  --ple-gguf /path/to/Qwen3.8-Flash-Next-GSQ-RCO-IQ3_XXS-00002-of-00002.gguf \
+  --mmap-experts --expert-profile data/expert-profile.bin --expert-cache auto \
+  --prefill 512 --spec 4 --spec-min-p 0.5 --mtp mtp/rt \
+  --max-context 4096 --kv int8 --pool-workers 15 \
+  --adapt-every 0 --pcie-frac 0 --vram-reserve-mib 1024
+```
+
+`--serve` is the engine's internal token protocol. For a browser or OpenAI API,
+put these args in the `args` array of the [server JSON example](ORCA.md#local-server),
+set `exe` to `build-hip/strata`, and use the matching pack's `tokenizer` directory.
+Launch with `.venv/bin/python -m serve.server --engine strata --config strata-hip.json --port 8080`.
+The worker count above was used on a 16-core CPU; measure it for your CPU.
+The 4K context is a smoke-test starting point, not a model limit. The expert cache
+sizes itself automatically and leaves 1 GiB of VRAM headroom.
+
+The automatic installer,
+vision helper, local gateway integrations, and multi-GPU experiments are not
+included in this backend patch.
+
+## Original backend validation (PR #94)
+
+The following is historical validation of the original backend, not a fresh
+test count for this replacement. Current build/test evidence and performance
+limits are recorded in [AMD_HIP_PERFORMANCE.md](AMD_HIP_PERFORMANCE.md).
+
+Tested against upstream `c1e903310f211e6630780c3bd2038778c071c68d` (0.1.20),
+with pinned llama.cpp `3cf03257f219afbe7334045ff7c6a06ac68c627d`.
+The host has an RX 7900 XTX (24 GiB), Ryzen 9 7950X3D, 64 GiB system RAM,
+and Fedora-family Linux with ROCm 7.1.52802 / Clang 20. The original model used
+for the smoke check was on mechanical RAID0; cold page faults were slow. Keep
+latency-sensitive data on SSD and measure warm residency separately.
+
+- Complete HIP Release build, including the executable and parity targets.
+- All 25 tests selected by the command below passed on 2026-09-29.
+- Intrinsic parity includes signed packed-byte dot products and overflow.
+- CPU/GPU mapped-memory handoff covers delayed publication, changing payloads,
+  and repeated graph replay in both directions.
+- GR coverage includes the maximum eight-token fused batch and graph replay
+  after inputs change. Native QSA checks the scalar HIP fallback.
+- Q8_K quantization is byte-exact against the existing reference. HIP's
+  `__fmul_rn` can become ordinary multiplication; disabling contraction in
+  `quantize_act.cu` preserves the separately rounded product before the magic
+  rounding bias is added. CUDA's compile flags remain unchanged.
+- The mmap regression covers canonical and differing native layer sizes,
+  truncated/oversized files, lookup bounds, and reopening after errors.
+- CUDA 13 / GCC 15 / sm_89 executable build passed; GR, sampler, quantization,
+  and mmap tests passed on the same tree. This is a regression check, not a
+  claim of complete CUDA inference validation.
+- CPU-only `strata-plan` build passed without either GPU backend.
+- Real GSQ-RCO IQ3_XXS server smoke with mmap, the shipped expert profile,
+  automatic GPU cache, and MTP passed arithmetic, executable Python addition,
+  system-marker recall, and 1,170-token batched-prefill recall. All four requests
+  ended normally with correct answers; one engine process served them all.
+  The cache held 11,142 experts (18.08 GiB), with 874 MiB of VRAM free after
+  startup. These short checks do not establish a throughput or quality benchmark.
+
+Build all targets before running the registered focused checks:
+
+```sh
+cmake --build build-hip -j2
+ctest --test-dir build-hip --output-on-failure --timeout 60 \
+  -E '^(ple_parity|platform_memory_test)$'
+```
+
+`ple_parity` requires an external model fixture. `platform_memory_test` requests
+256 MiB of locked memory; the test shell's 8 MiB memlock limit prevented that
+test from passing. These are explicit exclusions, not skipped tests counted as
+passes. Additional unpublished upstream fixture suites are not covered.
+
+The published performance table explicitly identifies the measured development
+snapshot; it must not be read as a benchmark of every subsequent rebase. Vision, long-context stress, broad answer-quality
+equivalence, other AMD cards, and mixed-vendor inference are not validated here.

docs/AMD_HIP_PERFORMANCE.md

@@ -0,0 +1,144 @@
+# RX 7900 XTX support and performance evidence
+
+This opt-in Linux `gfx1100` backend supersedes the initial support in
+[PR #94](https://github.com/Niko1221/Strata/pull/94). It retains HIP runtime and
+wave32 integer-dot compatibility, native mmap layout validation, and MTP, then
+adds HIP MMQ, optional calibrated dense hipBLASLt GEMM, native prefill batching,
+and the host-memory / SSD paths used by the measured configuration. HIP blocking
+expert uploads use one reusable pinned staging buffer, including prompt-cache
+refills, to avoid repeated pageable-source registrations.
+
+## Reproduce the configuration
+
+Build instructions are in [AMD_HIP.md](AMD_HIP.md). Use the repository-pinned
+llama.cpp dependency; do not silently substitute another revision. Build with
+`STRATA_PREFILL_MMQ=ON` and use these runtime variables for the measured arm:
+
+```sh
+export STRATA_PREFILL_MMQ=1
+export STRATA_HIPBLASLT_TUNING="$PWD/tools/hip/gfx1100-hipblaslt-100100.txt"
+export STRATA_PF_IDX_BATCH=1
+export STRATA_PF_EMBED_BATCH=1
+export STRATA_PREFILL_RING=96
+export STRATA_PLE_IO_THREADS=32
+```
+
+The supplied table is calibrated for gfx1100 and hipBLASLt version 100100.
+It is not a universal ROCm tuning table. The runtime guards architecture,
+library version and actual shape/stride/workspace requirements, falling back
+when a table entry is unavailable or incompatible. Batched embedding/indexer
+paths and HIP MMQ are opt-in at runtime. Default CUDA selection is preserved.
+
+Measured engine configuration: Orca Flash Next IQ3_XXS, native pack plus matching
+GGUF/tokenizer/template, `--mmap-experts --resident-cpu-experts`, fixed ranked
+`--expert-profile`, `--expert-cache auto`, `--prefill 8192`, `--spec 4`,
+`--spec-min-p 0.5`, matching MTP runtime, `--max-context 262144`, `--kv int8`,
+`--kv-resident 32768`, `--pool-workers 15`, `--adapt-every 0`, `--pcie-frac 0`,
+and `--vram-reserve-mib 1024`. Keep PLE on SSD.
+
+`--resident-cpu-experts` copies the complement of the static GPU cache to ordinary
+RAM. It requires mmap, a fixed expert profile, and disabled adaptation. It does
+not lock all data into RAM: temporary prefill borrowing and cache refill may
+still use the mapped fallback. Allocation needs sufficient available RAM;
+on Linux this option requires readable standard cgroup-v2 mounts. For each
+finite cgroup ancestor, the guard credits only `inactive_file` after subtracting
+`file_dirty` and `file_writeback`, capped by current usage; it remains bounded by
+the ancestor limit and host `MemAvailable`, with 8 GiB headroom. This accounts
+for reclaimable clean file cache but cannot reserve memory against concurrent
+system or process allocations.
+The POSIX PLE path issues direct reads through a configurable worker pool.
+
+Use a dedicated idle server, restart it between arms, and capture its engine log:
+
+```sh
+python3 tools/hip/bench_prefill.py \
+  --model MODEL_NAME --url http://127.0.0.1:8080 \
+  --engine-log /path/to/engine.log --label candidate --output candidate.json
+```
+
+The script uses the same synthetic source prompts and request order as the
+reported measurements: a small warmup, then 140/280/140/280 functions, with one
+follow-up after each fresh prompt. It enforces zero reused tokens for the fresh
+prompts and records prefill, decode, wall time and stop reason. It requires the
+unbuffered log to contain exactly one completed timing record per request. The
+128-token output cap is intentional for throughput measurement, not task success.
+Never use cancelled-request timing lines as throughput evidence.
+
+## Final revision benchmark
+
+Measured on 2026-09-29, source revision `9568e78da663e6b224ab96d57baf450717b36e73`.
+HIP executable SHA-256: `705f5925a788bc227bad72abfb5466032e0879d752fea9afc3eb32051b35182b`.
+The control and candidate below were both measured anew in this session; no
+historical benchmark values are mixed into the table. The control binary SHA-256
+is `4522d4937ca4a41ca294d31b60df8c2bc4df5b7a6caf65eaad10c036e08491f7`;
+its source was an existing local runtime snapshot rather than a clean upstream commit.
+Both arms use an RX 7900 XTX 24 GiB / gfx1100, Ryzen 9 7950X3D,
+64 GiB installed RAM, Fedora-family Linux, ROCm 7.1 and a 272 W GPU cap.
+Other model services are stopped for measurement and restored afterward.
+Sampling: temperature 0, top-k 1, top-p 1, min-p 0, seed 42, reasoning disabled.
+
+| Fresh request, execution order | Control prefill t/s | Candidate prefill t/s | Candidate output t/s | Candidate request wall time |
+| --- | ---: | ---: | ---: | ---: |
+| 4,210 tokens, first use | 240.0 | 447.5 | 55.5 | 11.729 s |
+| 8,830 tokens, first use at this size | 484.8 | 873.7 | 55.7 | 12.420 s |
+| 4,210 tokens, warmed | 457.2 | 901.1 | 56.7 | 6.937 s |
+| 8,830 tokens, warmed | 478.9 | 966.5 | 59.2 | 11.314 s |
+
+Each request generated exactly 128 tokens and finished at the intentional length
+cap. All four fresh prompts had zero reused KV tokens. Fresh-prompt prefill was
+1.80–2.02x the freshly measured existing AMD runtime; fresh-request wall time was
+40–46% lower. This is a comparison with our existing HIP runtime, not unmodified
+upstream, which does not provide this backend. The candidate also completed four
+cached follow-ups in 2.95–3.26 seconds. Those follow-ups are not apples-to-apples
+prefill comparisons: cache reuse differs with generated text and checkpoint
+selection. Full sanitized measurements are in
+[the fresh-run JSON](benchmarks/2026-09-29-gfx1100.json).
+
+Before this revision, an intermediate package missing the HIP upload staging
+buffer stalled on the 8,830-token request and was rejected. Restoring staging
+allowed the complete sequence above to finish. That failed attempt is not
+included in the throughput table. No watchdog limit was raised to obtain these
+results.
+
+Zero KV reuse is not equivalent to a cold filesystem cache. First-use speed
+and warmed speed are reported separately. Individual observations do not
+establish confidence intervals or a general rate at every context length.
+
+## Correctness and quality limits
+
+Both HIP and CUDA executables built from the measured source revision. CUDA was
+compile-checked, not performance-tested by this run. The selected HIP CTest suite
+passed **29/29** with the tuning table enabled, including expert-upload readback,
+asynchronous handoff, QSA, MMQ, Lt GEMM, KV streaming and PLE reading. The excluded
+`ple_parity` requires an external fixture; `platform_memory_test` requires a
+larger locked-memory limit than the test account provides. Separately, three
+real IQ3_XXS PLE matrix graph replays passed, and the POSIX direct-file test passed
+queued reads, short EOF, wake, close/drain and reopen checks.
+
+Development checks passed eight MMQ numerical comparisons, four actual Lt GEMM
+comparisons, and native QSA/indexer/embedding parity including tail states,
+chunk continuation and image overrides. Real quantized MMQ relative L2 error
+was approximately 0.0026 against raw-FP32/dequantized-weight reference, reflecting
+Q8 activation arithmetic; this is not bitwise numerical equivalence.
+
+End-to-end coding quality is not established by these numerical checks or capped
+throughput requests. Validate completed tasks with independent runtime tests; this
+contribution makes no broad answer-quality or agentic-reliability claim.
+
+The changes do not claim better model reasoning, verified full-context behavior,
+end-to-end vision validation, Windows HIP, other AMD architectures, or mixed
+AMD/NVIDIA execution. Existing CUDA multi-GPU code remains present but is not
+validation of HIP multi-GPU support.
+
+## Attribution and rejected experiments
+
+Native embedding gather and QSA batching are selectively adapted from
+[PR #108](https://github.com/Niko1221/Strata/pull/108), commit
+`acd487233c0bbe2217a6881c5bb43f8a283b0de5`. Its optional GDN parallel/split path is
+not included. Existing upstream MMQ orchestration is retained and enabled for
+HIP with AMD architecture identification and the correct backend compilation.
+
+Static 16K chunks, expanded FP16 expert tuning, and a 192-slot ring did not offer
+a consistent end-to-end win in the evaluated workload. The table includes only
+the selected 30 dense-shape rows. Microkernel speedups alone were not sufficient
+to select a configuration. No increase in GPU power cap was used.