#1208 HIP layer split / gfx1030: batched prompt path produces NaN logits (output "!!!!") unless STRATA_HIP_PROMPT_F16=1 and the gfx1030 card is the first stage

Descrição

> Strata 0.1.40.1 (`82f46a8`), Fedora 44, ROCm 7.1.1 system packages (rocBLAS 7.1.1, hipBLAS 7.1.0, no hipBLASLt),
> build `gfx1030;gfx1201`, `STRATA_PREFILL_MMQ=ON`. RX 6800 (gfx1030) + RX 9070 (gfx1201), Qwen3.8-Flash-Next
> GSQ-RCO IQ3_S, `--kv int8 --max-context 131072 --spec 4 --mtp`. The same bug was seen first on a single RX 6800 with
> setup's TheRock 7.10 wheels.
>
> - Any prompt part > 64 tokens (`short_read`, batched path) gives 248320/248320 non-finite logits → token 0 (`!`)
>   repeated. ≤ 64-token parts (verify windows) are fine.
> - `STRATA_DBG_NAN=1`: the first non-finite MoE layer is the first layer that runs on the gfx1030 card (layer 12 in a
>   `[9070, 6800]` split; layer 0-3 on the 6800 alone). The RX 9070 alone is fine.
> - `STRATA_PREFILL_MMQ=0`, `STRATA_PROMPT_ATTN_OLD=1`, `--no-prefill-borrow`, `STRATA_HC_SPLIT=0`,
>   `STRATA_DENSE_MMQ=1`, `STRATA_GR_UNFUSED=1`: still NaN. A standalone `hipblasGemmEx` BF16→FP32 (OP_T/OP_N, same
>   shapes, 32 MiB workspace) is correct on gfx1030 with the same rocBLAS.
> - `STRATA_HIP_PROMPT_F16=1` fixes the 6800 alone, and fixes the split **only when the gfx1030 card is stage 0**
>   (`gpu [0,1]`): 69 to 20183-token prompts OK, 0 FP16 overflows. With gfx1030 as stage 1, `STRATA_DBG_NAN` counts
>   non-finite FP16 GEMM outputs from the first GEMM of the stage (hc up, N=10240 K=320). The engine also exits
>   sometimes with `prefill gr_silu: no kernel image is available` / `prefill swiglu_interleaved: no kernel image…`
>   in that order (cf. #690).
> - Suspects: a BF16-activation prompt kernel on gfx1030, and the per-device state of the second stage
>   (`set_act_f16` symbol / code object selection, #690).