feat(strata): build and run edge0-145b on AMD HIP (Windows, gfx1100)

edge0-strata/CMakeLists.txt

@@ -2495,6 +2495,16 @@ 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)
+    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 +2522,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 +2624,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,52 @@ 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.
+
 ## Launch contract with edge0-infer
 
 - `edge0/CONTRACT` holds one integer, the contract version (currently 4). `edge0-infer` refuses to start an

edge0-strata/src/core/verify.cpp

@@ -2193,11 +2193,19 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
             // same path the f32 rows take).  Bit-exact to the host's from_float (probe_std); the
             // xlseq doorbell after both copies covers it.
             quantize_q8_0_rows(lat_xl_, (long long) T * lat_G_, (long long) lat_ELL_, lat_act0_dev_, cs);
+#if defined(STRATA_USE_HIP)   // amd-hip port: HIP does not capture a D2H memcpy into the verify window's graph ("end capture: operation
+            // failed due to a previous error during capture"); a copy kernel writes the same bytes into the mapped host rows
+            strata::kernels::copy_bytes_from_mapped(m_latact0_, lat_act0_dev_,
+                                                    (int64_t) ((size_t) T * (size_t) lat_G_ * ((size_t) lat_ELL_ / 32) * 34), cs);
+            strata::kernels::copy_bytes_from_mapped(m_latxlq_, lat_xl_,
+                                                    (int64_t) ((size_t) T * (size_t) lat_G_ * (size_t) lat_ELL_ * 4), cs);
+#else
             cudaMemcpyAsync(m_latact0_, lat_act0_dev_,
                             (size_t) T * (size_t) lat_G_ * ((size_t) lat_ELL_ / 32) * 34,
                             cudaMemcpyDeviceToHost, cs);
             cudaMemcpyAsync(m_latxlq_, lat_xl_, (size_t) T * (size_t) lat_G_ * (size_t) lat_ELL_ * 4,
                             cudaMemcpyDeviceToHost, cs);
+#endif   // STRATA_USE_HIP (amd-hip port)
             if (trace_m_ != nullptr) gpu_stamp(trace_m_, (l * kProfPer + 35) * 2, cs);   // K0-probe: D2H done
             doorbell_publish_value(nullptr, nullptr, nullptr, 0, 0, nullptr, nullptr, nullptr,
                                    (uint32_t*) m_xlseq_, (uint32_t) l + 1, cs);
@@ -2730,6 +2738,7 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
         }
         // #879: whether each row's logits are all finite, beside the pick (read after the window's own sync)
         if (nan_guard_on()) strata::kernels::logits_nonfinite_rows(head_logits_, T, (int) n_vocab_, m_nf_, cs);
+#if !defined(STRATA_USE_HIP)   // amd-hip port: a debug dump that does a synchronous D2H cudaMemcpy inside the window's graph capture; HIP fails the capture on it ("end capture: operation failed due to a previous error during capture")
         if (windows <= 12 && n_vocab_ > 0 && T > 0) {
             static std::vector<float> tmp;
             tmp.resize((size_t) n_vocab_);
@@ -2749,6 +2758,7 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
                          std::sqrt(s2 / (double) n_vocab_));
             std::fflush(stderr);
         }
+#endif   // !STRATA_USE_HIP (amd-hip port)
     }
     stamp(g.n_layers, 1, 0);
     if (tap_on_ && !batch_rec_ && tap_file_ != nullptr) copy_from_mapped(m_tap_, tap_dev_, (int64_t) T * (int64_t) tap_layers_.size() * N, cs);
@@ -2883,6 +2893,7 @@ bool Verifier::capture(int T, std::string& err) {
         err = std::string("verify: end capture: ") + cudaGetErrorString(ce);
         return false;
     }
+#if !defined(STRATA_USE_HIP)   // amd-hip port: this dump uses cudaGraph* names the HIP shim of this file does not map; HIP has its own listing below
     if (graph) {   // K0: dump node order (unconditional — env was not reaching the engine) (D2H / wait_flag / H2D / kernel)
         size_t nn = 0;
         if (cudaGraphGetNodes(graph, nullptr, &nn) == cudaSuccess && nn > 0) {
@@ -2909,6 +2920,7 @@ bool Verifier::capture(int T, std::string& err) {
             }
         }
     }
+#endif   // !STRATA_USE_HIP (amd-hip port)
 #if defined(STRATA_USE_HIP)   // 143-amdk: the same listing on HIP (kernel names by hipKernelNameRefByPtr); STRATA_VERIFY_NODES=2 also lists the first nodes in graph order
     if (std::getenv("STRATA_VERIFY_NODES") != nullptr) {
         size_t nn = 0;

edge0-strata/src/prefill/moe_fused.cu

@@ -1407,11 +1407,13 @@ void experts_latent(const WideBatch& b, const LatentGeom& g, int n_expert, int64
     const int64_t tiles = (n + kTileRows - 1) / kTileRows + (b.e1 - b.e0);
     const unsigned g_gu = (unsigned) std::min<int64_t>(tiles * 5, (int64_t) d.sms * d.occ_gu);
     const unsigned g_d = (unsigned) std::min<int64_t>(tiles * 7, (int64_t) d.sms * d.occ_d);
+#if !defined(__HIPCC__)   // amd-hip port: the latent A kernels are CUDA-only (lat_dev().ok is false on HIP, so this is not reached)
     latent_kernel<true><<<g_gu, LTHREADS, lat_smem_bytes(), s>>>(b, tb, g, (const uint8_t*) xa, src,
                                                                  (uint8_t*) hbuf, nullptr);
     ck(cudaGetLastError(), "experts_latent gu");
     latent_kernel<false><<<g_d, LTHREADS, lat_smem_bytes(), s>>>(b, tb, g, (const uint8_t*) hbuf, src, nullptr, dm);
     ck(cudaGetLastError(), "experts_latent down");
+#endif   // !__HIPCC__ (amd-hip port)
     ck(cudaGetLastError(), "experts_latent");
 }
 

edge0-strata/src/prefill/moe_fused_q3k.cu

@@ -1534,11 +1534,13 @@ void experts_latent_q3k(const WideBatch& b, const LatentGeom& g, int n_expert, i
     const int64_t nfb_d = (int64_t) std::max<int64_t>(1, (int64_t) (g.ell >> 8));
     const unsigned g_gu = (unsigned) std::min<int64_t>(tiles * nfb_gu, (int64_t) d.sms * d.occ_gu);
     const unsigned g_d = (unsigned) std::min<int64_t>(tiles * nfb_d, (int64_t) d.sms * d.occ_d);
+#if !defined(__HIPCC__)   // amd-hip port: the latent A kernels are CUDA-only (lat_dev().ok is false on HIP, so this is not reached)
     latent_kernel_q3k<true><<<g_gu, LTHREADS, lat_smem_bytes(), s>>>(b, tb, g, (const uint8_t*) xa, src,
                                                                  (uint8_t*) hbuf, nullptr);
     ck(cudaGetLastError(), "experts_latent_q3k gu");
     latent_kernel_q3k<false><<<g_d, LTHREADS, lat_smem_bytes(), s>>>(b, tb, g, (const uint8_t*) hbuf, src, nullptr, dm);
     ck(cudaGetLastError(), "experts_latent_q3k down");
+#endif   // !__HIPCC__ (amd-hip port)
     ck(cudaGetLastError(), "experts_latent_q3k");
 }
 

edge0-strata/src/prefill/prefill.cpp

@@ -5338,6 +5338,7 @@ bool Prefill::run_impl(const int64_t* tokens, int64_t n, int64_t pos0, std::stri
     for (size_t q = 0; q <= nb2 && q < 17; ++q)
                                         std::fprintf(stderr, "latmmq-bh[%zu]=%d\n", q, bh[q]);
                                 }
+#if !defined(STRATA_USE_HIP)   // amd-hip port: trace-only block (STRATA_LATENT_MMQ_TRACE), uses cudaPointerGetAttributes which the HIP shim does not map
                                 {   // latmmq-dev: which device each buffer lives on, and the current one
                                     static bool devsaid = false;
                                     if (!devsaid && std::getenv("STRATA_LATENT_MMQ_TRACE") != nullptr) {
@@ -5366,6 +5367,7 @@ bool Prefill::run_impl(const int64_t* tokens, int64_t n, int64_t pos0, std::stri
                                                      devof(m.lat_gdst));
                                     }
                                 }
+#endif   // !STRATA_USE_HIP (amd-hip port)
                                 auto run_p = [&](const mmq::Product& pr) {
                                     static const bool alt = std::getenv("STRATA_LATENT_MMQ_ALTCTX") != nullptr;
                                     static std::unique_ptr<mmq::Context> altctx;