From 20c31a5bab309f5d334dde16232fb792100b734a Mon Sep 17 00:00:00 2001 From: Graffioh Date: Thu, 27 Aug 2026 06:16:44 +0000 Subject: [PATCH 1/3] fix(pflash): use HIP-specific sparse alpha default --- server/src/server/server_main.cpp | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/server/src/server/server_main.cpp b/server/src/server/server_main.cpp index efae13bbe..b425d015e 100644 --- a/server/src/server/server_main.cpp +++ b/server/src/server/server_main.cpp @@ -888,8 +888,15 @@ int main(int argc, char ** argv) { bool pflash_enabled = (sconfig.pflash_mode != ServerConfig::PflashMode::OFF); if (pflash_enabled) { set_environment_variable("DFLASH_FP_USE_BSA", "1", false); - set_environment_variable("DFLASH_FP_ALPHA", "0.85", false); +#if defined(DFLASH27B_BACKEND_HIP) + constexpr const char * kPflashAlphaDefault = "0.95"; +#else + constexpr const char * kPflashAlphaDefault = "0.85"; +#endif + set_environment_variable("DFLASH_FP_ALPHA", kPflashAlphaDefault, false); set_environment_variable("DFLASH27B_FA_WINDOW", "0", false); + std::fprintf(stderr, "[server] PFlash sparse alpha: %s\n", + std::getenv("DFLASH_FP_ALPHA")); } if (sconfig.draft_residency == DraftResidencyPolicy::RequestScoped && From c3ab61a636046c41512808f1d752e887c6344b94 Mon Sep 17 00:00:00 2001 From: Graffioh Date: Thu, 27 Aug 2026 06:18:49 +0000 Subject: [PATCH 2/3] fix(pflash): seed session keep ratio from config --- server/src/server/adaptive_keep_ratio.h | 18 ++++++++++++++++++ server/src/server/http_server.cpp | 4 ++-- server/src/server/http_server.h | 2 +- server/test/test_adaptive_keep_ratio.cpp | 16 ++++++++++++++++ server/test/test_server_unit.cpp | 12 ++++++++++++ 5 files changed, 49 insertions(+), 3 deletions(-) diff --git a/server/src/server/adaptive_keep_ratio.h b/server/src/server/adaptive_keep_ratio.h index 959b87bce..57ca1c47c 100644 --- a/server/src/server/adaptive_keep_ratio.h +++ b/server/src/server/adaptive_keep_ratio.h @@ -66,6 +66,24 @@ class HttpServerSessions { } } + float get_or_create_keep_ratio( + const std::string& session_id, float initial_keep) { + std::lock_guard lock(mu_); + auto it = map_.find(session_id); + if (it == map_.end()) { + evict_if_full_locked(); + AdaptiveKeepRatioState state; + state.last_keep = std::clamp( + initial_keep, kBanditKeepMin, kBanditKeepMax); + lru_.push_front(session_id); + it = map_.emplace( + session_id, Entry{state, lru_.begin()}).first; + } else { + lru_.splice(lru_.begin(), lru_, it->second.lru_it); + } + return it->second.state.last_keep; + } + float get_keep_ratio(const std::string& session_id) const { std::lock_guard lock(mu_); auto it = map_.find(session_id); diff --git a/server/src/server/http_server.cpp b/server/src/server/http_server.cpp index a909344e8..4ba44a880 100644 --- a/server/src/server/http_server.cpp +++ b/server/src/server/http_server.cpp @@ -242,10 +242,10 @@ bool flowkv_should_activate(const ServerConfig & config, float resolve_pflash_keep_ratio(float configured_ratio, const std::string & session_id, - const HttpServerSessions & sessions) { + HttpServerSessions & sessions) { return session_id.empty() ? configured_ratio - : sessions.get_keep_ratio(session_id); + : sessions.get_or_create_keep_ratio(session_id, configured_ratio); } bool should_clamp_flowkv_disk_cache( diff --git a/server/src/server/http_server.h b/server/src/server/http_server.h index 52c36473b..d6072e0f0 100644 --- a/server/src/server/http_server.h +++ b/server/src/server/http_server.h @@ -248,7 +248,7 @@ bool flowkv_should_activate(const ServerConfig & config, int aged_token_estimate); float resolve_pflash_keep_ratio(float configured_ratio, const std::string & session_id, - const HttpServerSessions & sessions); + HttpServerSessions & sessions); bool should_clamp_flowkv_disk_cache( bool flowkv, const DiskPrefixCachePolicy & policy); bool canonical_turn_matches_checkpoint( diff --git a/server/test/test_adaptive_keep_ratio.cpp b/server/test/test_adaptive_keep_ratio.cpp index 495fc28ee..02fda47a0 100644 --- a/server/test/test_adaptive_keep_ratio.cpp +++ b/server/test/test_adaptive_keep_ratio.cpp @@ -120,6 +120,22 @@ TEST_CASE(AdaptiveKeepRatioFixture, unknown_session_returns_default) { CHECK(mgr.turn_count("no-such-session") == 0); } +TEST_CASE(AdaptiveKeepRatioFixture, session_seed_is_clamped_and_sticky) { + HttpServerSessions mgr; + + CHECK(approx_eq( + mgr.get_or_create_keep_ratio("low", 0.0f), kBanditKeepMin)); + CHECK(approx_eq( + mgr.get_or_create_keep_ratio("high", 1.0f), kBanditKeepMax)); + CHECK(approx_eq( + mgr.get_or_create_keep_ratio("configured", 0.05f), 0.05f)); + + // Once created, a later request's curve/config value must not reset the bandit. + CHECK(approx_eq( + mgr.get_or_create_keep_ratio("configured", 0.15f), 0.05f)); + CHECK(mgr.size() == 3); +} + TEST_CASE(AdaptiveKeepRatioFixture, get_ema_reflects_post_update_value) { HttpServerSessions mgr; CHECK(approx_eq(mgr.get_ema("s1"), 0.0f)); diff --git a/server/test/test_server_unit.cpp b/server/test/test_server_unit.cpp index 01dad8dcc..494acdb71 100644 --- a/server/test/test_server_unit.cpp +++ b/server/test/test_server_unit.cpp @@ -5971,6 +5971,18 @@ TEST_CASE(ServerUnitFixture, test_flowkv_session_keep_ratio_override) { TEST_ASSERT(std::fabs(adaptive_ratio - 0.09f) < 1e-6f); } +TEST_CASE(ServerUnitFixture, test_fresh_session_uses_configured_keep_ratio) { + HttpServerSessions sessions; + + const float configured_ratio = 0.05f; + const float resolved_ratio = http_detail::resolve_pflash_keep_ratio( + configured_ratio, "fresh", sessions); + + TEST_ASSERT(std::fabs(resolved_ratio - configured_ratio) < 1e-6f); + TEST_ASSERT(std::fabs(sessions.get_keep_ratio("fresh") - configured_ratio) < 1e-6f); + TEST_ASSERT(sessions.size() == 1); +} + // ═══════════════════════════════════════════════════════════════════════ // Qwen3-0.6B drafter loader: truncated GGUF guard (bug #438) // ═══════════════════════════════════════════════════════════════════════ From c594b8792e0c250cf6cc837ecc60f77bbeb59fce Mon Sep 17 00:00:00 2001 From: Graffioh Date: Thu, 27 Aug 2026 06:23:50 +0000 Subject: [PATCH 3/3] feat(pflash): auto-enable rocWMMA kernels on HIP --- server/CMakeLists.txt | 32 ++++++++++++++++++++++++-------- 1 file changed, 24 insertions(+), 8 deletions(-) diff --git a/server/CMakeLists.txt b/server/CMakeLists.txt index 925d5839d..cd7689cb5 100644 --- a/server/CMakeLists.txt +++ b/server/CMakeLists.txt @@ -187,13 +187,32 @@ if(DFLASH27B_ROCMFP2_AFFINE) add_compile_definitions(ROCMFP2_AFFINE=1) endif() -# HIP Phase 2 — opt-in rocWMMA flashprefill kernels (Strix Halo / gfx1151). -# Default OFF (Phase 1 = ggml q8 fallback). Set ON for the 1.7-2.7× compress -# speedup at 8K-32K context. Requires rocwmma headers installed under -# ${ROCM_PATH}/include (e.g. `sudo apt install rocwmma`). +# HIP Phase 2 — rocWMMA flashprefill kernels (Strix Halo / gfx1151). +# Fresh HIP builds enable these automatically when the headers are installed. +# Preserve an explicit -DDFLASH27B_HIP_SM80_EQUIV=OFF opt-out. +set(_dflash_hip_sm80_equiv_explicit OFF) +if(DEFINED CACHE{DFLASH27B_HIP_SM80_EQUIV}) + set(_dflash_hip_sm80_equiv_explicit ON) +endif() + +set(_dflash_hip_sm80_equiv_default OFF) +if(DFLASH27B_GPU_BACKEND STREQUAL "hip") + find_path(DFLASH27B_ROCWMMA_INCLUDE_DIR rocwmma/rocwmma.hpp + HINTS "${_dflash_rocm_root}/include" /opt/rocm/include + NO_DEFAULT_PATH) + if(DFLASH27B_ROCWMMA_INCLUDE_DIR) + set(_dflash_hip_sm80_equiv_default ON) + elseif(NOT _dflash_hip_sm80_equiv_explicit) + message(WARNING + "rocWMMA headers were not found; PFlash drafter will use the slow q8 " + "fallback. Install rocwmma to enable the HIP FlashPrefill kernels.") + endif() +endif() option(DFLASH27B_HIP_SM80_EQUIV "HIP: build the rocWMMA flashprefill kernels (Phase 2). Requires rocwmma." - OFF) + "${_dflash_hip_sm80_equiv_default}") +unset(_dflash_hip_sm80_equiv_default) +unset(_dflash_hip_sm80_equiv_explicit) # Resolve the CUDA architecture list up-front so downstream logic (notably # the consumer-Blackwell ggml workaround below) can inspect the actual arches @@ -616,9 +635,6 @@ if(DFLASH27B_GPU_BACKEND STREQUAL "hip") # GPU draft top-K path instead of the CPU fallback. target_compile_definitions(dflash_common PUBLIC DFLASH27B_HAVE_DRAFT_TOPK=1) if(DFLASH27B_HIP_SM80_EQUIV) - find_path(DFLASH27B_ROCWMMA_INCLUDE_DIR rocwmma/rocwmma.hpp - HINTS "${_dflash_rocm_root}/include" /opt/rocm/include - NO_DEFAULT_PATH) if(NOT DFLASH27B_ROCWMMA_INCLUDE_DIR) message(FATAL_ERROR "DFLASH27B_HIP_SM80_EQUIV=ON but rocwmma/rocwmma.hpp not found. "