feat(pflash): improve AMD defaults and enable rocWMMA

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. "

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<std::mutex> 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<std::mutex> lock(mu_);
         auto it = map_.find(session_id);

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(

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(

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 &&

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));

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)
 // ═══════════════════════════════════════════════════════════════════════