verify: the shared-expert stream fork starts off on HIP (#816)

The fork overlaps the shared expert with the routed experts on a second stream.
It pays off on CUDA, where a36be1d added it, but on HIP the cross-stream event
dependency costs more than the shared expert takes to run, so there the fork
starts off and STRATA_SH_STREAM=1 turns it back on. CUDA's default is unchanged.

Same shape as mtp.cpp's STRATA_MTP_PREFILL_SYNC (#382): the env var always wins,
only the default differs per backend. The profiler takes the fork off to time the
stages, so a profiled run never showed this - which is why STRATA_SH_STREAM=0 was
the only way to see it.

Radeon AI PRO R9700 (gfx1201, ROCm 6.4.3), Qwen3.8-Flash-Next IQ3_XXS, --kv int8
--kv-resident 32768 --resident-experts, temp 0, 3 runs each, same binary and
prompt, 48,067-token prompt so the KV stream path runs:

  stock 0.1.39 (fork on, default)   43.40 t/s   (38.90 - 44.60)
  this change (fork off, default)   62.90 t/s   (55.40 - 63.80)
  STRATA_SH_STREAM=1                43.50 t/s   (39.10 - 45.20)
  STRATA_SH_STREAM=0                63.00 t/s   (55.90 - 64.10)

The fork is a scheduling change only: the same four prompts answer byte for byte
the same with it on and off. #816 reports the same fall on a gfx1030 (RX 6800,
Windows, 42.5 -> 56.4 t/s), and a comment on #646 sees it on a gfx1100.

Short prompts are barely affected here (65.4 against 62.7 t/s at 256 tokens), so
this only shows once the KV stream path is in play.

src/core/verify.cpp

@@ -89,6 +89,21 @@ const bool g_doorbell_store = env_on("STRATA_DOORBELL_STORE");
 #endif
 const bool g_trace = env_on("STRATA_VERIFY_TRACE");
 
+// The shared expert runs on its own stream, forked off and joined back per layer (`sh_fork` in record_window).  The
+// overlap pays off on CUDA, where it was added (cfd3b72).  On HIP a cross-stream event dependency costs more than the
+// shared expert takes to run, so there the fork starts off: same binary, same 48,067-token prompt, three runs each on a
+// Radeon AI PRO R9700 (gfx1201, ROCm 6.4.3), the fork off decodes at 62.9 t/s against 43.4 t/s with it on.  The
+// profiler turns the fork off to time the stages, so a profiled run never showed that.  #816 reports the same fall on a
+// gfx1030 (RX 6800, Windows, 42.5 -> 56.4 t/s).  STRATA_SH_STREAM overrides either default: `=1` forks, `=0` does not.
+const bool g_sh_stream = [] {
+    if (const char* v = std::getenv("STRATA_SH_STREAM")) return v[0] != '0';
+#if defined(STRATA_USE_HIP)
+    return false;
+#else
+    return true;
+#endif
+}();
+
 bool mapped(size_t bytes, void** h, void** d) {
 #if defined(STRATA_USE_HIP)
     if (g_coherent) {
@@ -976,11 +991,7 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
         }
         stamp(l, 16, grp);
         const bool q8_ffn = gr_read_group(1, true, inj_, inj2_);   // true: nat_xq_ holds the MoE input's q8_1 (STRATA_QFUSE)
-        static const bool sh_stream_env = [] {
-            const char* e = std::getenv("STRATA_SH_STREAM");
-            return !e || e[0] != '0';
-        }();
-        const bool sh_fork = sh_stream_env && !prof_on_ && sh_cs_ != nullptr && ev_fork_ != nullptr && ev_join_ != nullptr;
+        const bool sh_fork = g_sh_stream && !prof_on_ && sh_cs_ != nullptr && ev_fork_ != nullptr && ev_join_ != nullptr;
         cudaStream_t sh_stream = sh_fork ? sh_cs_ : cs;
         if (sh_fork) {
             cudaEventRecord(ev_fork_, cs);
@@ -1172,7 +1183,9 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
             }
             moe_hit_add(parts_out, hit_out, p_dst, p_counts + 1, cap, N, cs);
         }
-        if (!prof_on_ && sh_cs_ != nullptr && ev_fork_ != nullptr && ev_join_ != nullptr) {
+        // Same condition as `sh_fork` above, which is per group and out of scope here: wait only when the fork
+        // actually ran (a wait on an event never recorded is a no-op, but saying it outright reads better).
+        if (g_sh_stream && !prof_on_ && sh_cs_ != nullptr && ev_fork_ != nullptr && ev_join_ != nullptr) {
             cudaStreamWaitEvent(cs, ev_join_, 0);
         }
         if (sg_gated_[grp]) {   // STRATA_LFUSE: the shared expert's gate was not applied: the combine applies it (all rows are hits)