#697 Windows HIP / gfx1201: prompt stalls at `GPU rang 0`; post-store `__threadfence_system()` resolves it

Beschreibung

I'm running Strata on Windows with an AMD Radeon AI PRO R9700 (gfx1201).

Prompt processing consistently stalled during the initial prompt. The stall diagnostics showed:

```text
0 layers served
verify window (last window, not the current stage): 0 tokens at position 0, host at layer step 1; the GPU rang 0; flags: served 0, plan (A) 0, copies (B) 0
```

This occurred repeatedly before the change.

### Diagnostic test

As a diagnostic test, I tried adding `__threadfence_system()` immediately after the store to `seq` in `doorbell_publish_kernel`.

This is the only change I made to the kernel:

```diff
 if (threadIdx.x == 0) {
     __threadfence_system();
     *(volatile uint32_t*) seq = *(volatile uint32_t*) seq + 1u;
+    __threadfence_system();
 }
```

After rebuilding Strata with this one-line change, the same configuration successfully completed inference for the first time:

```text
prompt 170 tokens = 0 reused + 170 read in 1871 ms (90.9 tok/s), 235 generated in 3299 ms (71.2 tok/s)
```

I then sent another prompt in the same session, which also completed successfully:

```text
prompt 247 tokens = 165 reused + 82 read in 35556 ms (2.3 tok/s), 239 generated in 6190 ms (38.6 tok/s)
```

The attached log contains both the original stalls and the subsequent successful runs using the patched binary.

### Environment

- OS: Windows
- GPU: AMD Radeon AI PRO R9700
- HIP architecture: gfx1201, wave32
- VRAM: 32 GB
- System RAM: 32 GB
- AMD Display Driver: `32.0.31041.1004` (2026-08-17)
- HIP driver/runtime reported by `strata-device`: `71726391 / 71726391`
- HIP runtime DLL: `amdhip64_7.dll`
- Backend: HIP
- Strata: v0.1.38
- Commit: `99f3dbd0b21d1401b3769e0c0d963913607f380b`
- Model: Qwen3.8-Flash-Next Q2_0
- Context: 32768
- KV cache: int8
- Expert cache: auto
- Resident budget: 8 GiB
- Speculative decoding: 4
- GPU connection: OCuLink / PCIe 4.0 x4

### Note

This change was suggested while troubleshooting the issue with an AI assistant, so I'm not sure whether adding the post-store system fence is the correct upstream fix.

I'm reporting it because the original build consistently stalled before the first layer was served, while the build with this one-line change successfully completed inference and handled a subsequent prompt.

[strata-q2_0-redacted.log](https://github.com/user-attachments/files/33013094/strata-q2_0-redacted.log)