#2788 hip: route the asynchronous adaptive tier through adapt_copy_h2d (#884)

本文

## Title
<!-- If Applicable, reference the GitHub issue -->
Issue: relates to #884 — it makes that switch effective on a path it never reached. It does not claim to fix the ROCm barrier itself.

## Summary
`STRATA_HIP_ADAPT_KERNEL_COPY=1` (#884) is dead code whenever `--adapt-async 1` is on. The asynchronous adaptive tier copies its swaps in with a raw `cudaMemcpyAsync` and never calls `adapt_copy_h2d()`, so every swap goes back over SDMA — the combination #884 exists to avoid. On gfx1030 the queue then dies with `HSA_STATUS_ERROR_INVALID_PACKET_FORMAT` (the malformed barrier packet), which the server reports as `verify: layer 0 never rang (unspecified launch failure)`.

## What changed
`src/program/generate.cpp`, `a_in_one` (the async tier's step-2 copy in): `cudaMemcpyAsync(...)` → `adapt_copy_h2d(...)`, the helper the blocking tier already calls. Two lines.

The stream, the asynchrony and the fallback are unchanged: `adapt_copy_h2d` ends in the identical `cudaMemcpyAsync` when the switch is off or the source is not mapped, so with the switch off this is behaviourally the old code.

Leaves alone: the blocking tier, `a_back` (the D2H side has no kernel-copy helper), and `sycl/src/program/generate.cpp` (its own `adapt_copy_h2d`, same gap).

## Evidence
gfx1030 (RX 6800), ROCm 7.13.0a, 0.1.44.1 source, `--adapt-async 1 --adapt-swaps 160 --kv-resident 32768 --cache-suffix-reuse 64`, arena mode, 250 MiB VRAM free, same config and box:

| binary | outcome | elapsed | AQL aborts |
| --- | --- | ---: | ---: |
| unpatched | aborted | 72.7 s | 2 |
| patched | completed | 185.2 s | 0 |

The switch's own line (`adaptive swaps copy with a kernel, not SDMA (#884)`) appears once with the patch and never without it. It is printed on the helper's first call, and `a_in_one` is the only GPU copy this config makes, so the line is the proof the helper is now reachable.

Cost — 3 alternating rounds, 512 greedy tokens each, medians:

| prompt | kernel copy | SDMA |
| --- | ---: | ---: |
| essay | decode 39.5 tok/s | 40.4 tok/s |
| code | decode 36.7 tok/s | 36.6 tok/s |

−2.2% / +0.3%: no measurable cost, as the #884 reporter also measured.

## Extra Notes
- **Not built or run: CUDA and SYCL.** Only HIP/gfx1030 could be built here (no CUDA toolkit, no oneAPI installed). `src/program/generate.cpp` is one target for HIP and CUDA; the SYCL engine builds a separate file, `sycl/src/program/generate.cpp`, whose `adapt_copy_h2d` has the same gap — left alone deliberately rather than hand-edited dpct output I cannot compile.
- The switch stays opt-in: with it off, the async tier keeps using SDMA and the abort returns.
- Reported from hardware the maintainers may not have (single RX 6800, not the reporter's 2x RX 6900 XT), so this is one machine's result.