From 78fe2d46ef82b8c726ef5c136c4198b14d307a8f Mon Sep 17 00:00:00 2001 From: ppkrol Date: Thu, 12 Nov 2026 19:54:27 +0100 Subject: [PATCH] feat(strata): build and run edge0-145b on AMD HIP (Windows, gfx1100) edge0-strata keeps upstream Strata's HIP backend, but the latent-MoE additions did not build or run with it. Build (HIP only): - CMakeLists.txt: compile lat_q3k_*.cu and moe_fused_q3k.cu in the HIP MMQ branch (their entry points are stubs there) and add latent_dense.cpp and latent_moe.cpp to strata_engine for HIP - moe_fused.cu, moe_fused_q3k.cu: the latent kernel launches were outside the !__HIPCC__ guard that holds the kernels - verify.cpp, prefill.cpp: two debug dumps use cudaGraph* / cudaPointer* names the HIP shim does not map Run (HIP only): - verify.cpp: the D stage's two D2H cudaMemcpyAsync calls and the "verify top3" dump (a synchronous cudaMemcpy) ran inside the verify window's graph capture, which HIP rejects ("end capture: operation failed due to a previous error during capture"); on HIP the copies go through copy_bytes_from_mapped and the dump is skipped Docs: EDGE0.md describes the build and the configuration. The fused latent kernels stay CUDA-only; on HIP the model runs with STRATA_LATENT_FUSED=0, STRATA_LATENT_FUSED_DU=0, STRATA_LAT_FUSED_Q3K=0. Tested on one RX 7900 XTX under Windows 11. The CUDA build was not rebuilt. Co-Authored-By: Claude Fable 5.2 Claude-Session: https://claude.ai/code/session_01VmeXwKqR7k1jsojD2bUxve --- edge0-strata/CMakeLists.txt | 15 ++++++++ edge0-strata/EDGE0.md | 46 +++++++++++++++++++++++ edge0-strata/src/core/verify.cpp | 12 ++++++ edge0-strata/src/prefill/moe_fused.cu | 2 + edge0-strata/src/prefill/moe_fused_q3k.cu | 2 + edge0-strata/src/prefill/prefill.cpp | 2 + 6 files changed, 79 insertions(+) diff --git a/edge0-strata/CMakeLists.txt b/edge0-strata/CMakeLists.txt index afc05c5f..b83eb044 100644 --- a/edge0-strata/CMakeLists.txt +++ b/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}) diff --git a/edge0-strata/EDGE0.md b/edge0-strata/EDGE0.md index 4ec3743e..feb7b430 100644 --- a/edge0-strata/EDGE0.md +++ b/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 --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 diff --git a/edge0-strata/src/core/verify.cpp b/edge0-strata/src/core/verify.cpp index 1750c2eb..c33c21d4 100644 --- a/edge0-strata/src/core/verify.cpp +++ b/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 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; diff --git a/edge0-strata/src/prefill/moe_fused.cu b/edge0-strata/src/prefill/moe_fused.cu index da5d5906..8b38d8d8 100644 --- a/edge0-strata/src/prefill/moe_fused.cu +++ b/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(tiles * 5, (int64_t) d.sms * d.occ_gu); const unsigned g_d = (unsigned) std::min(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<<>>(b, tb, g, (const uint8_t*) xa, src, (uint8_t*) hbuf, nullptr); ck(cudaGetLastError(), "experts_latent gu"); latent_kernel<<>>(b, tb, g, (const uint8_t*) hbuf, src, nullptr, dm); ck(cudaGetLastError(), "experts_latent down"); +#endif // !__HIPCC__ (amd-hip port) ck(cudaGetLastError(), "experts_latent"); } diff --git a/edge0-strata/src/prefill/moe_fused_q3k.cu b/edge0-strata/src/prefill/moe_fused_q3k.cu index 3bcf3356..ff17b0fc 100644 --- a/edge0-strata/src/prefill/moe_fused_q3k.cu +++ b/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(1, (int64_t) (g.ell >> 8)); const unsigned g_gu = (unsigned) std::min(tiles * nfb_gu, (int64_t) d.sms * d.occ_gu); const unsigned g_d = (unsigned) std::min(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<<>>(b, tb, g, (const uint8_t*) xa, src, (uint8_t*) hbuf, nullptr); ck(cudaGetLastError(), "experts_latent_q3k gu"); latent_kernel_q3k<<>>(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"); } diff --git a/edge0-strata/src/prefill/prefill.cpp b/edge0-strata/src/prefill/prefill.cpp index bf1db352..55cfb098 100644 --- a/edge0-strata/src/prefill/prefill.cpp +++ b/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 altctx;