Run Edge0-145B Q3_K on AMD HIP (Linux, gfx1100)

edge0-strata/CMakeLists.txt

@@ -2495,6 +2495,18 @@ if(STRATA_NATIVE_EXPERTS)
       src/prefill/moe_f16_kq.cu
       src/prefill/ggml_cuda_host.cu
       "${_strata_quantize_cu}")
+    # amd-hip port: the latent A stage sources of the CUDA branch.  On HIP the lat_q3k_* entry points are stubs that
+    # exit ("CUDA-only"), so run with STRATA_LATENT_FUSED=0 STRATA_LATENT_FUSED_DU=0 STRATA_LAT_FUSED_Q3K=0; they are
+    # compiled here so prefill.cpp links.
+    set(_strata_lat_hip_sources
+      src/prefill/lat_q3k_down.cu
+      src/prefill/lat_q3k_gu.cu
+      src/prefill/lat_q3k_swiglu.cu
+      src/prefill/lat_q3k_fused.cu
+      src/prefill/moe_fused_q3k.cu
+      src/prefill/lat_dense_w11.cu    # the latent D / U projections on the gfx11 matrix cores
+      src/prefill/lat_q3k_w11.cu)     # the latent A stage (Q3_K gate/up, IQ4_NL down) on the gfx11 matrix cores
+    list(APPEND _strata_mmq_sources ${_strata_lat_hip_sources})
     if(STRATA_MMQ_KQUANTS)
       # the same K-quant instances as the CUDA branch: on HIP they serve the dense GGUF projections
       # of Gemm::native (STRATA_DENSE_MMQ=1), the mixed-quant packs' attention/PLE tensors
@@ -2512,6 +2524,8 @@ if(STRATA_NATIVE_EXPERTS)
     # moe_fused.cu is Strata's own CUDA-shaped source: it takes the HIP compat headers, ggml's MMQ sources do not
     set_source_files_properties(src/prefill/moe_fused.cu src/prefill/moe_fused_iq.cu src/prefill/moe_f16_kq.cu
                                 PROPERTIES INCLUDE_DIRECTORIES "${STRATA_HIP_COMPAT_INCLUDE_DIR}")
+    set_source_files_properties(${_strata_lat_hip_sources}
+                                PROPERTIES INCLUDE_DIRECTORIES "${STRATA_HIP_COMPAT_INCLUDE_DIR}")   # amd-hip port
     add_library(strata_mmq STATIC ${_strata_mmq_sources})
     target_include_directories(strata_mmq
       PRIVATE "${_strata_gg}/include" "${_strata_gg}/src" "${_strata_mmq}"
@@ -2612,6 +2626,9 @@ if(STRATA_NATIVE_EXPERTS)
     target_include_directories(latent_cpu_verify PRIVATE ${CMAKE_SOURCE_DIR}/include)
     target_link_libraries(latent_cpu_verify PRIVATE strata_engine strata_kernels_cpu CUDA::cudart)
   endif()
+  if(STRATA_ENABLE_HIP)   # amd-hip port: the CUDA block above adds these to the engine; HIP needs them too
+    target_sources(strata_engine PRIVATE src/core/latent_dense.cpp src/latent/latent_moe.cpp)
+  endif()
   if(STRATA_ENABLE_CUDA OR STRATA_ENABLE_HIP)
     add_executable(native_expert_parity src/kernels/native_expert_parity.cpp)
     target_include_directories(native_expert_parity PRIVATE ${_strata_gpu_include_directories})

edge0-strata/EDGE0.md

@@ -193,6 +193,132 @@ interpreter. To run `build\strata.exe` by hand, put the CUDA DLLs on `PATH` firs
 (`set "PATH=%CUDA_PATH%\bin;%PATH%"`). The engine writes temporary files under `%TMP%`, which must be
 writable for the user who runs `edge0-infer`.
 
+## Build (Windows, AMD, experimental)
+
+Edge0 145B (Q3_K) has been run on one AMD card with the HIP backend this fork inherits from Strata: a Radeon RX 7900
+XTX (gfx1100, 24 GB) under Windows 11, ROCm 10.2.0a20260930. Other Radeon cards, Linux, Edge0 120B and the NVFP4
+folders have not been tried. `edge0-infer` does not start this build (it expects CUDA), so the server is started by
+hand.
+
+Requirements: Visual Studio 2022 Build Tools with the C++ workload, Python 3.10 or newer, Git, and a current AMD
+display driver. No ROCm or HIP SDK installation: the build script installs ROCm from AMD's wheels into `.rocm-win`
+(about 7 GB).
+
+1. Build and package the engine for the card's architecture, in this directory:
+
+   ```bat
+   set STRATA_HIP_ARCHS=gfx1100
+   set STRATA_NO_VISION=1
+   tools\hip\build_windows.bat
+   ```
+
+   The result is `dist\strata-windows-x64-hip.zip` (`strata.exe`, the ROCm DLLs in `rocm\bin`, the hipBLASLt tuning
+   table in `tuning\`). Unpack it into a folder, for example `engine-hip\`.
+2. Create the server's Python environment as in step 3 of [Build (Windows, NVIDIA)](#build-windows-nvidia).
+3. Write a server configuration from `edge0/models/edge0-145b.json`: replace the placeholders by hand, set `exe` to
+   the unpacked `strata.exe`, add `"backend": "hip"`, and set `lib_dirs` to the unpacked `rocm\bin`. In `env`:
+
+   | Variable | Value | Why |
+   | --- | --- | --- |
+   | `STRATA_LATENT_FUSED`, `STRATA_LATENT_FUSED_DU`, `STRATA_LAT_FUSED_Q3K` | `0` | the fused latent kernels are CUDA-only; the engine exits if they are selected on HIP |
+   | `STRATA_LATENT_MMQ` | `1` (optional) | the latent A stage of the prompt on MMQ; about a third faster than the default path on this card |
+   | `STRATA_LATENT_DENSE_FMT` | `8` or `4` | the engine reads the number, not `q8` / `q4` |
+   | `STRATA_LATENT_A`, `STRATA_LATENT_DENSE` | paths with `/` | `labels.bin` is looked up beside `STRATA_LATENT_A` by its last `/`, so a path with only `\` is not found |
+   | `STRATA_HIPBLASLT_TUNING` | the table in `tuning\` | as in upstream Strata's AMD configurations |
+
+4. Start it: `.venv\Scripts\python.exe serve\server.py --engine strata --config <file> --port 8000`.
+
+Measured with `tools\strata_bench.py` on the card above (Ryzen 7 5800X3D, 64 GB RAM, `--vram-reserve-mib 3072`, about
+4100 experts resident; medians):
+
+| A stage of the prompt | Chat decode tok/s | 4K prompt tok/s | 16K prompt tok/s | Agent decode tok/s |
+| --- | ---: | ---: | ---: | ---: |
+| default (`native_expert_grouped`) | 37.9 | 243.6 | 241.3 | 43.6 |
+| `STRATA_LATENT_MMQ=1` | 40.4 | 335.0 | 319.3 | 47.2 |
+
+The prompt speed is far below the NVIDIA figures because the fused latent kernels have no HIP version yet. Answer
+quality on this build has only been checked by hand.
+
+## Build (Linux, AMD, experimental)
+
+Edge0 145B (Q3_K) runs on a Radeon RX 7900 XTX (gfx1100, 24 GB) under Linux with ROCm 7.2.4. The expert stage and
+the two dense projections of a latent layer have gfx11 kernels (`src/prefill/lat_q3k_w11.cu`,
+`src/prefill/lat_dense_w11.cu`). `edge0-infer` does not start this build, so start the server by hand.
+
+A layer split across two cards works (`"gpu": [0, 1]`). Each card holds the latent D and U weights of its own
+layers, and on a split the token embedding table stays in mapped host memory whatever `STRATA_EMBED_VRAM` says. A
+card cannot read another card's VRAM: HIP answers "an illegal memory access was encountered", and the window or the
+draft layer then never finishes. Add `--resident-experts` to keep in RAM only the experts that no card holds.
+
+Do not set `--pcie-mode dma` on a split. It was not tested again after the fix above. Before the fix, the first
+expert that a card did not hold was never copied, and the request stopped with `verify: timed out at layer N`.
+
+Requirements: ROCm 7 with the HIP compiler, hipBLAS and hipBLASLt (the commands below were run in the
+`docker.io/rocm/dev-ubuntu-24.04:7.2.4-complete` image), CMake 3.24 or newer, Ninja, and Python 3.10 or newer.
+
+1. Unpack llama.cpp into `third_party/llama.cpp` as in step 1 of [Build (Linux, NVIDIA)](#build-linux-nvidia).
+2. Build the engine. `STRATA_PREFILL_MMQ=ON` is required: without it the link stops with undefined symbols such as
+   `strata::prefill::fused_q3k::experts_latent_q3k`.
+
+   ```
+   export HIP_PATH="$(hipconfig -R)"
+   cmake -S . -B build-hip -G Ninja -DCMAKE_BUILD_TYPE=Release \
+     -DSTRATA_ENABLE_HIP=ON -DSTRATA_ENABLE_CUDA=OFF -DCMAKE_HIP_ARCHITECTURES=gfx1100 \
+     -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang++" -DCMAKE_PREFIX_PATH=/opt/rocm \
+     -DSTRATA_PREFILL_MMQ=ON -DSTRATA_NATIVE_EXPERTS=ON \
+     -DSTRATA_GGML_DIR="$PWD/third_party/llama.cpp"
+   cmake --build build-hip --target strata
+   ```
+
+3. Copy the hipBLASLt tuning table for the card and the hipBLASLt version into `tuning/` beside the engine. The
+   engine reads it from there and prints `hipBLASLt tuning enabled`. Without it the dense prompt products use plain
+   hipBLAS, and a 5.3K prompt ran at 725 tok/s where it runs at 1150 tok/s with the table. ROCm 7.2.4 has hipBLASLt
+   1.2.2, which is table `100202`:
+
+   ```
+   mkdir -p build-hip/tuning
+   cp tools/hip/gfx1100-hipblaslt-100202.txt build-hip/tuning/
+   ```
+
+4. Create the server's Python environment as in step 3 of [Build (Linux, NVIDIA)](#build-linux-nvidia).
+5. Write a server configuration from `edge0/models/edge0-145b.json`. Replace the placeholders, set `exe` to
+   `build-hip/strata`, add `"backend": "hip"`, and set `lib_dirs` to the ROCm library directory (`/opt/rocm/lib`).
+   Set `cwd` and `log` to a directory that the server can write to. In `env`, set `STRATA_LATENT_FUSED` and
+   `STRATA_LAT_FUSED_Q3K` to `1`, and `STRATA_LATENT_DENSE_FMT` to `8`. Do not set `--pool-workers` above the number
+   of physical cores; without the option the engine chooses.
+6. Start it: `.venv/bin/python serve/server.py --engine strata --config <file> --port 8000`. The process needs
+   `/dev/kfd` and `/dev/dri`. With more than one card, select one with `HIP_VISIBLE_DEVICES`.
+
+Use a card on a full-width PCIe link. The engine moves experts between host memory and the card while it decodes,
+and it prints each card's link speed at the start of a layer split (`link 0.8 GB/s`). `current_link_width` of each
+bridge above the card in `/sys/bus/pci/devices` shows a narrow link.
+
+Measured on a Ryzen 5 5600X with 64 GB RAM and two RX 7900 XTX cards, one card in use, `--vram-reserve-mib 700`,
+6652 experts resident (the card also held 1.8 GiB of another process). The client sends three chat questions of 150
+tokens and one fresh prompt of each size, thinking off, one run each; it is not `tools/strata_bench.py`.
+
+| Card link | Configuration | Chat decode tok/s | 5.3K prompt tok/s | 21K prompt tok/s |
+| --- | --- | ---: | ---: | ---: |
+| PCIe 4.0 x16 | all of the above | not measured | 1181.8 | 1302.1 |
+| PCIe 4.0 x16 | expert stage on ggml MMQ: `STRATA_LATENT_FUSED=0`, `STRATA_LAT_FUSED_Q3K=0`, `STRATA_LATENT_MMQ=1` | 46.2 to 51.1 | 1157.2 to 1202.0 | 1285.3 to 1316.1 |
+| PCIe 4.0 x16 | MMQ expert stage, no tuning table | not measured | 724.9 | not measured |
+| PCIe 4.0 x16 | MMQ expert stage, no tuning table, `STRATA_LATENT_FUSED_DU=0` | 46.4 to 55.1 | 269.1 | 280.2 |
+| PCIe 4.0 x16 | native expert stage, no tuning table, `STRATA_LATENT_FUSED_DU=0` | 47.7 to 57.1 | 214.3 | 216.7 |
+| PCIe 3.0 x1 | native expert stage, no tuning table, `STRATA_LATENT_FUSED_DU=0` | 17.8 to 23.3 | 108.3 | 129.1 |
+
+The gfx11 expert stage and the MMQ expert stage ran in different sessions, and the same MMQ configuration measured
+1157.2 and 1202.0 tok/s on the 5.3K prompt, so the two stages are level within the run-to-run spread.
+
+Two cards in a layer split with `--resident-experts`, layers 0 to 21 on the PCIe 3.0 x1 card and 22 to 47 on the x16
+card: the cards held 27.0 GiB of the 37.3 GiB of experts, and the engine kept 21.4 GiB of experts in RAM, where one
+card keeps all 37.3 GiB. The process took 32.5 GB of RAM. Chat answers decoded at 32.2 to 51.2 tok/s. A 5.3K prompt
+ran at 272.9 to 284.6 tok/s and a 21K prompt at 413.2 tok/s, because the prompt path streams experts to the x1 card.
+
+Answer quality on this build has only been checked by hand and by the first token's log-probabilities over four
+prompts. The gfx11 expert stage moves the top-8 log-probabilities by at most 0.48 against the MMQ stage. The gfx11
+projections move them by at most 0.57 against `STRATA_LATENT_FUSED_DU=0`. The two upstream expert stages (native and
+MMQ) differ from each other by up to 2.41 on the same prompts.
+
 ## Launch contract with edge0-infer
 
 - `edge0/CONTRACT` holds one integer, the contract version (currently 4). `edge0-infer` refuses to start an

edge0-strata/include/strata/core/latent_share.hpp

@@ -1,4 +1,4 @@
-// Latent MoE: the one session-level D/U device pair, stored by the engine's LatentTab and pulled by the
+// Latent MoE: the session-level D/U device pair of each device, stored by the engine's LatentTab and pulled by the
 // prefill and verify arms inside their own init (no per-object external wiring to miss).  A non-latent pack
 // never stores a pair, so both accessors stay null.
 #pragma once
@@ -7,8 +7,18 @@
 namespace strata {
 namespace core {
 
-inline uint8_t*& latent_shared_d() { static uint8_t* p = nullptr; return p; }
-inline uint8_t*& latent_shared_u() { static uint8_t* p = nullptr; return p; }
+// The pair is device memory, so there is one per device: a kernel cannot read another device's buffer (HIP answers
+// "an illegal memory access was encountered").  A layer split's later stage pulls the pair of its own device, or
+// finds none and uploads its own copy.  `device` is the CUDA / HIP device number.
+struct LatentSharedPair {
+    uint8_t *d = nullptr, *u = nullptr;
+};
+inline LatentSharedPair& latent_shared(int device) {
+    static LatentSharedPair pairs[32];
+    return pairs[device & 31];
+}
+inline uint8_t*& latent_shared_d(int device) { return latent_shared(device).d; }
+inline uint8_t*& latent_shared_u(int device) { return latent_shared(device).u; }
 
 inline const uint8_t*& latent_xl_current() { static const uint8_t* p = nullptr; return p; }
 inline const uint32_t*& latent_xlseq_current() { static const uint32_t* p = nullptr; return p; }

edge0-strata/include/strata/core/native_head.hpp

@@ -51,7 +51,10 @@ class NativeEmbed {
     ~NativeEmbed();
     NativeEmbed(const NativeEmbed&) = delete;
     NativeEmbed& operator=(const NativeEmbed&) = delete;
-    bool load(const std::vector<std::string>& shards, int64_t n_embd, int64_t n_vocab, std::string& err);
+    /// `one_device` = false (a layer split): the table stays in mapped host memory whatever STRATA_EMBED_VRAM
+    /// says, because a later stage's device cannot read another device's VRAM.
+    bool load(const std::vector<std::string>& shards, int64_t n_embd, int64_t n_vocab, std::string& err,
+              bool one_device = true);
     /// Rows for device token ids.
     void gather_dev(const int32_t* tokens, int64_t n_tok, float* out, void* stream) const;
     /// One row for a host token id.

edge0-strata/include/strata/hip_compat/cuda_runtime.h

@@ -11,6 +11,9 @@
 #endif
 #include <hip/hip_runtime.h>
 #include <hip/hip_version.h>
+#include <atomic>
+#include <cstdio>
+#include <cstdlib>
 // Do not let HIP's legacy macro corrupt libstdc++ attribute names.
 #ifdef __noinline__
 #undef __noinline__
@@ -94,17 +97,13 @@ hipError_t mem_get_info(size_t* free_bytes, size_t* total_bytes);
 #else
 #define cudaMemGetInfo hipMemGetInfo
 #endif
-#define cudaMemcpy hipMemcpy
 #define cudaMemcpy2DAsync hipMemcpy2DAsync
-#define cudaMemcpyAsync hipMemcpyAsync
 #define cudaMemcpyPeerAsync hipMemcpyPeerAsync
 #define cudaMemcpyPeer hipMemcpyPeer
 #define cudaMemcpyDefault hipMemcpyDefault
 #define cudaMemcpyDeviceToDevice hipMemcpyDeviceToDevice
 #define cudaMemcpyDeviceToHost hipMemcpyDeviceToHost
 #define cudaMemcpyHostToDevice hipMemcpyHostToDevice
-#define cudaMemset hipMemset
-#define cudaMemsetAsync hipMemsetAsync
 #define cudaPeekAtLastError hipPeekAtLastError
 #define cudaRuntimeGetVersion hipRuntimeGetVersion
 #define cudaSetDevice hipSetDevice
@@ -131,6 +130,86 @@ inline hipError_t malloc_host(T** pointer, size_t bytes) {
 }
 }  // namespace strata::hip_compat
 
+// ---- copies and fills on the buffer's own device.
+// CUDA takes a memcpy or a memset of any device's buffer whatever the current device is, and the engine's layer
+// split relies on that.  HIP runs such a call on the CURRENT device's queue.  A blocking hipMemset of another
+// device's buffer then returns before that device sees the bytes (measured on two gfx1100 cards, ROCm 7.2.4: a
+// kernel launched on the buffer's device right after the call still read the old bytes), so the fill lands late,
+// over whatever that device wrote in between.  With more than one device, these wrappers make the buffer's device
+// current for the call.  A call on a stream of its own is left alone: the stream names the device.
+// STRATA_HIP_XDEV_LOG=1 prints the first calls that needed the switch.
+namespace strata::hip_compat {
+inline bool multi_device() {
+    static const bool multi = [] {
+        int n = 0;
+        return hipGetDeviceCount(&n) == hipSuccess && n > 1;
+    }();
+    return multi;
+}
+// the device that owns `p`, or -1 for host memory and for a pointer HIP does not know
+inline int device_of(const void* p) {
+    hipPointerAttribute_t a{};
+    if (p == nullptr) return -1;
+    if (hipPointerGetAttributes(&a, p) != hipSuccess) {
+        (void) hipGetLastError();   // a plain host pointer: not an error of the caller's
+        return -1;
+    }
+    return a.type == hipMemoryTypeDevice ? a.device : -1;
+}
+struct DeviceScope {
+    int previous = -1;
+    DeviceScope(int device, const char* what, const char* file, int line, size_t bytes) {
+        int cur = 0;
+        if (device < 0 || hipGetDevice(&cur) != hipSuccess || cur == device || hipSetDevice(device) != hipSuccess) return;
+        previous = cur;
+        static const bool log = [] {
+            const char* v = std::getenv("STRATA_HIP_XDEV_LOG");
+            return v != nullptr && v[0] == '1';
+        }();
+        static std::atomic<int> shown{0};
+        if (log && shown.fetch_add(1) < 200)
+            std::fprintf(stderr, "strata hip: %s of %zu bytes on device %d while device %d was current (%s:%d)\n", what,
+                         bytes, device, cur, file, line);
+    }
+    ~DeviceScope() {
+        if (previous >= 0) (void) hipSetDevice(previous);
+    }
+    DeviceScope(const DeviceScope&) = delete;
+    DeviceScope& operator=(const DeviceScope&) = delete;
+};
+// the device side of a copy: the source of a device-to-host copy, else the destination
+inline int copy_device(const void* dst, const void* src, hipMemcpyKind kind) {
+    if (kind == hipMemcpyDeviceToHost) return device_of(src);
+    const int d = device_of(dst);
+    return d >= 0 ? d : device_of(src);
+}
+inline hipError_t memcpy_on(const char* file, int line, void* dst, const void* src, size_t n, hipMemcpyKind kind) {
+    if (!multi_device()) return hipMemcpy(dst, src, n, kind);
+    const DeviceScope scope(copy_device(dst, src, kind), "memcpy", file, line, n);
+    return hipMemcpy(dst, src, n, kind);
+}
+inline hipError_t memcpy_async_on(const char* file, int line, void* dst, const void* src, size_t n, hipMemcpyKind kind,
+                                  hipStream_t stream = nullptr) {
+    if (stream != nullptr || !multi_device()) return hipMemcpyAsync(dst, src, n, kind, stream);
+    const DeviceScope scope(copy_device(dst, src, kind), "memcpy on the null stream", file, line, n);
+    return hipMemcpyAsync(dst, src, n, kind, stream);
+}
+inline hipError_t memset_on(const char* file, int line, void* dst, int value, size_t n) {
+    if (!multi_device()) return hipMemset(dst, value, n);
+    const DeviceScope scope(device_of(dst), "memset", file, line, n);
+    return hipMemset(dst, value, n);
+}
+inline hipError_t memset_async_on(const char* file, int line, void* dst, int value, size_t n, hipStream_t stream = nullptr) {
+    if (stream != nullptr || !multi_device()) return hipMemsetAsync(dst, value, n, stream);
+    const DeviceScope scope(device_of(dst), "memset on the null stream", file, line, n);
+    return hipMemsetAsync(dst, value, n, stream);
+}
+}  // namespace strata::hip_compat
+#define cudaMemcpy(...) strata::hip_compat::memcpy_on(__FILE__, __LINE__, __VA_ARGS__)
+#define cudaMemcpyAsync(...) strata::hip_compat::memcpy_async_on(__FILE__, __LINE__, __VA_ARGS__)
+#define cudaMemset(...) strata::hip_compat::memset_on(__FILE__, __LINE__, __VA_ARGS__)
+#define cudaMemsetAsync(...) strata::hip_compat::memset_async_on(__FILE__, __LINE__, __VA_ARGS__)
+
 template <typename Kernel>
 inline hipError_t cudaFuncSetAttribute(Kernel kernel, hipFuncAttribute attribute, int value) {
     return hipFuncSetAttribute(reinterpret_cast<const void*>(kernel), attribute, value);