ROCm/gfx1151: `--metal-kernels` fails 6 assertions in the paired Q8_0 decode case

## Summary

On a Strix Halo (gfx1151) ROCm-only host, `make test-rocm` at `c1d4597` reports
`ds4 tests: 6 failure(s)`. All six are inside `--metal-kernels`, in
`test_metal_q8_0_decode_pair_exact_case`. They are **two independent problems**:

1. **`-ffast-math` breaks the bit-exactness assertion.** The project's default
`CFLAGS` include `-ffast-math`, and `test-rocm` inherits them. Rebuilding the
same test without `-ffast-math` makes the paired kernel bit-identical to the
two standalone dispatches, exactly as the test intends.
2. **`hipMemcpy` H2D fails with `invalid argument`** for the second weight
range, which makes `ds4_gpu_matmul_q8_0_tensor()` return 0 and leaves the
comparison running on uninitialised output. This one is *not* explained by
fast-math and I could not reproduce it standalone.

This is **not a regression from the recent ROCm MXFP4 work** — control below.

## Environment

```
GPU AMD RYZEN AI MAX+ 395 w/ Radeon 8060S (gfx1151), 128 GB unified
OS Ubuntu 26.04 LTS, kernel 7.0.0-30-generic
HIP 7.14.60850-0000000
cc gcc 15.2.0
ds4 c1d4597, built with `make strix-halo` (--offload-arch=gfx1151)
model DeepSeek-V4-Flash-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-chat-v2-imatrix-0731
```

## Reproduce

```sh
export LD_LIBRARY_PATH=/opt/rocm/core-7.14/lib
make strix-halo
make test-rocm # or, much cheaper, just the failing case:
./ds4_test --metal-kernels
```

## Output

```
metal-kernels:
ds4-test: paired Q8_0 exactness mismatches=57/77 max_abs=1.90735e-06, 14/19 max_abs=9.53674e-07
tests/ds4_test.c:807: assertion failed: mismatch0 == 0
tests/ds4_test.c:808: assertion failed: mismatch1 == 0
ds4: ROCm model range copy failed for q8_0 at 0.00/0.32 MiB: invalid argument
tests/ds4_test.c:774: assertion failed: ds4_gpu_matmul_q8_0_tensor(ref1, weights_raw, weight_alloc, weight1_offset, in_dim, out1_dim, x, 1) != 0
ds4: ROCm model range copy failed for q8_0_pair1 at 0.00/0.32 MiB: invalid argument
tests/ds4_test.c:777: assertion failed: ds4_gpu_matmul_q8_0_pair_tensor(pair0, pair1, weights_raw, weight_alloc, 0, weight1_offset, in_dim, out0_dim, out1_dim, x, 1) != 0
ds4-test: paired Q8_0 exactness mismatches=19/19 max_abs=20.6578, 71/77 max_abs=3203.74
tests/ds4_test.c:807: assertion failed: mismatch0 == 0
tests/ds4_test.c:808: assertion failed: mismatch1 == 0
metal-kernels: ERR
```

The `max_abs=20.6578 / 3203.74` line is downstream of the two failed copies —
`ref1` and `pair1` were never written — so it is one bug, not two.

---

## Part 1 — `-ffast-math` vs the bit-exactness assertion

The assertion is a `memcmp` on floats, i.e. strictly bitwise:

```c
if (memcmp(&ref0_host[i], &pair0_host[i], sizeof(float)) != 0) mismatch0++;
```

Rebuilding **the same test at the same commit** with `-ffast-math` removed from
both `CFLAGS` and `ROCM_CFLAGS`:

| build | case `out0=77,out1=19` |
|---|---|
| default flags (`-ffast-math`) | `mismatches=57/77 max_abs=1.90735e-06, 14/19 max_abs=9.53674e-07` |
| same tree, no `-ffast-math` | **no mismatch line — assertion passes** |

So the paired kernel *is* bit-identical to two standalone dispatches; the
divergence is reassociation in the fast-math build.

`test-rocm` passes the project default `CFLAGS` straight through:

```make
test-rocm:
$(MAKE) -B ds4_test ... CFLAGS="$(CFLAGS) $(ROCM_HOST_CFLAGS) -DDS4_ROCM_BUILD" ...
```

and `CFLAGS` contains `-ffast-math`. The Makefile already does
`$(filter-out -ffast-math,$(QUALITY_CFLAGS))` for `q4k-dot-test` /
`mxfp4-dot-test`, so the pattern exists — it just does not reach `ds4_test.o`
or the ROCm device objects.

I did not isolate host vs device: I removed `-ffast-math` from both at once.

## Part 2 — `hipMemcpy` H2D returns `invalid argument`

Call site, `cuda_model_range_acquire()` in `ds4_cuda.cu`:

```c
err = cudaMalloc(&dev, (size_t)bytes); /* succeeds */
const char *src = (const char *)model_map + offset;
err = cudaMemcpy((char *)dev + done, src + done, (size_t)n,
cudaMemcpyHostToDevice); /* invalid argument */
```

Parameters, derived from the failing case
(`test_metal_q8_0_decode_pair_exact_case(out0_dim=19, out1_dim=77)`):

```
in_dim = 4096
row_bytes = (4096/32)*34 = 4352
weight1_off = round_up(19*4352, 4096) = 86016
weight1_len = 77*4352 = 335104 (= 0.32 MiB, matches the message)
weight_alloc = round_up(86016+335104, 4096) (so the range is in bounds)
```

Note it is the **second** range that fails; the range at offset 0 is fine, and
the mirrored case (`out0=77, out1=19`, a 0.08 MiB second range) does not fail.

### What I ruled out

- **Offset / size / alignment.** A standalone HIP program doing exactly
`hipMalloc(335104)` + `hipMemcpy(dev, base+86016, 335104, H2D)` from a
page-aligned `posix_memalign` buffer succeeds. Reproducer attached below.
- **Memory pressure.** Same program with 40, 80 and **95 GiB** of `hipMalloc`
ballast resident: still `no error` in every case. The real failure happens
with ~80 GiB of model resident, so this was the obvious suspect and it is out.
- **Stale range cache.** `ds4_gpu_set_model_map()` calls
`cuda_model_range_release_all()` when the base pointer changes, which clears
`g_model_range_by_offset`, so a stale entry from a previous case is not it.
- **`hipHostRegister` capability.** `DS4_CUDA_WEIGHT_CACHE_VERBOSE=1` prints no
`mapped ... MiB` line, so the mapped path is never taken and every range goes
through `hipMalloc` + `hipMemcpy`. But standalone,
`hipHostRegister(Mapped|ReadOnly)` and `hipHostGetDevicePointer` both succeed
on this GPU — so the capability is there and something sets
`g_model_range_mapping_supported = 0` earlier (that branch swallows the
error, which is also why there is nothing in the log to point at).

I could not get past this without knowing the intended lifetime of the model
map here, so I have stopped and am reporting what is pinned down.

## Control — this is not from the MXFP4 series

`make test-rocm` is new (#526 / `7891f27`), so this is the first time these
assertions could run on a ROCm-only host at all. To check whether the ROCm
MXFP4 kernels caused it:

- `tests/ds4_test.c` is **byte-identical** between `84cc882` and `c1d4597`
(`git diff --quiet` confirms).
- The only `ds4_rocm.cu` change in that range is +14 lines from `39a8f18`: a
`cuda_block_mxfp4` struct, a `static_assert`, and a `__constant__` lookup
table. Nothing on the Q8_0 path.
- Decisive: I built the **current** test harness and linked the **old**
`ds4_rocm.o` from `84cc882`. Byte-identical failures — same counts
(`57/77`, `14/19`, `19/19`, `71/77`) and same `max_abs`
(`1.90735e-06`, `9.53674e-07`, `20.6578`, `3203.74`).

Also worth noting for anyone comparing: #526 reports
`./ds4_test --metal-kernels` → `metal-kernels: OK` on the same GPU on
2026-07-09, but `427e281` (2026-07-18) added this case afterwards. So there is
no contradiction — the case has simply never passed here.

## Attached reproducer

```cpp
// hipcc -O2 --offload-arch=gfx1151 -o repro repro.hip.cpp
// ./repro -> both cases "no error"
// ./repro 95 -> 95 GiB ballast first, still "no error"
#include
#include
#include
#include
#include

static int try_copy(const char *tag, const void *host_base,
unsigned long long off, unsigned long long len) {
void *dev = nullptr;
hipError_t e = hipMalloc(&dev, (size_t)len);
if (e != hipSuccess) { printf(" %-18s hipMalloc -> %s\n", tag, hipGetErrorString(e)); return 1; }
e = hipMemcpy(dev, (const char *)host_base + off, (size_t)len, hipMemcpyHostToDevice);
printf(" %-18s off=%-8llu len=%-8llu -> %s\n", tag, off, len, hipGetErrorString(e));
hipFree(dev);
return e != hipSuccess;
}

int main(int argc, char **argv) {
const unsigned long long page = 4096, row_bytes = (4096ull / 32ull) * 34ull;
std::vector ballast;
if (argc > 1) {
const size_t chunk = 1ull << 30;
for (int i = 0, n = atoi(argv[1]); i < n; i++) {
void *p = nullptr;
if (hipMalloc(&p, chunk) != hipSuccess) break;
hipMemset(p, 0, chunk);
ballast.push_back(p);
}
printf(" ballast: %zu GiB\n", ballast.size());
}
int fails = 0;
const unsigned int dims[2][2] = {{77, 19}, {19, 77}};
const char *tags[2] = {"case A (out1=19)", "case B (out1=77)"};
for (int c = 0; c < 2; c++) {
unsigned long long w0 = dims[c][0] * row_bytes;
unsigned long long off = ((w0 + page - 1) / page) * page;
unsigned long long len = dims[c][1] * row_bytes;
unsigned long long alloc = ((off + len + page - 1) / page) * page;
void *base = nullptr;
if (posix_memalign(&base, (size_t)page, (size_t)alloc)) return 2;
memset(base, 0x5a, (size_t)alloc);
fails += try_copy(tags[c], base, off, len);
free(base);
}
for (void *p : ballast) hipFree(p);
printf(" failures: %d\n", fails);
return fails != 0;
}
```

Output on this machine:

```
case A (out1=19) off=335872 len=82688 -> no error
case B (out1=77) off=86016 len=335104 -> no error
ballast: 95 GiB
case A (out1=19) off=335872 len=82688 -> no error
case B (out1=77) off=86016 len=335104 -> no error
failures: 0
```

## Suggested next steps

- For Part 1: filter `-ffast-math` out of the objects that back exactness
assertions, the way `q4k-dot-test` / `mxfp4-dot-test` already do — or relax
the assertion to a tolerance and say so.
- For Part 2: the `else` branch in `cuda_model_range_acquire()` that sets
`g_model_range_mapping_supported = 0` swallows the `hipHostRegister` error.
Printing it under `DS4_CUDA_WEIGHT_CACHE_VERBOSE` would probably show why the
mapped path is abandoned on this host, which is likely upstream of the copy
failure. Happy to run any instrumented build on this hardware.

Related: #526 (rocm-test target), #525 (GPU kernel tests on ROCm), #16 (Strix
Halo discussion).