#2021 [Bug]: gfx1151: a 4K prompt stays under STRATA_PF_SWITCH_MIN_T=4096 (its chunk is 4,095 tokens) and is ~16% slower

说明

## What happened

`docs/STRIX_HALO.md` section 5 sets `STRATA_PF_SWITCH_MIN_T=4096`, so that prompts shorter than 4,096 tokens keep the default numerics. A prompt of 4,096 tokens (what the benchmarks call "4K") is read as a chunk of 4,095 tokens, so it is below the threshold. The hyper-connection and padding switches, and the dense GEMM that goes with them, then stay off for it, and the dense projections run hipBLASLt solutions that are not in `tools/hip/gfx1151-hipblaslt-100401.txt` (`Cijk_..._HSS_..._MT96x96x32`; 26% of the GPU time of a 4K prompt in a rocprofv3 trace, against `pfg::kernel_bk64` when the switches are on).

Measured on one server, 4 requests of the same 4K prompt in a row (UD-Q4_K_XL, `STRATA_PF_FUSED=1 STRATA_PF_FUSED_KQ=1`, the section 5 switches, only `STRATA_PF_SWITCH_MIN_T` changed), prompt tok/s:

| `STRATA_PF_SWITCH_MIN_T` | range over 4 requests | per request |
|---|---:|---:|
| 4096 | 1,052-1,106 | |
| 2048 | 1,262-1,298 | about 0.58 s shorter |
| 1024 | 1,246-1,295 | about 0.58 s shorter |

I know `4096` was chosen to keep the 1.0-1.9K prompts of the test set on the default output (b6b7a87d); the code default is 0 (every chunk). So this is not a request to change the number. The question is which of these you would want:

1. a lower default in section 5 (for example 2048; chunks of 2,048-4,095 tokens then differ from the default by rounding, which I can measure with first-token KL the way I did for #2020);
2. a note in section 5 that a "4K" prompt is below 4096, with the number to use if you want it above;
3. the threshold compared with the prompt length instead of the chunk length (a 4K prompt as a whole is long enough to benefit, even when the chunk is 4,095 tokens).

I have no preference between them, and I can make the change in whichever form you pick.

## Strata version, GPU, OS

v0.1.43 · Radeon 8060S (Ryzen AI Max+ 395, gfx1151) · Linux 7.0, ROCm 7.14.1

## Engine log

Not attached: it is the same on every start. The rocprofv3 traces (4096 against 2048) are available on request.

🤖 Generated with [Claude Code](https://claude.com/claude-code)