#889 sycl: read doorbell flags with an uncached L1+L3 hint (the GPU never saw the host's store on an Arc Pro B60)

Description

This fix was found, tested and written by Claude on my machine, at my request. I'm not a developer; I'm submitting it because it may help other Arc owners.

## What

`strata::sys_load` (the device side of every doorbell wait: `wait_flag_ge_kernel`, `wait_flag_ge_or_kernel`, the elementwise spins) now reads the flag through an `annotated_ptr` with an uncached L1+L3 `read_hint`, after an acquire fence, instead of a system-scope atomic load. One file, `sycl/include/strata/sycl_doorbell.hpp`. The old load is kept behind `-DSTRATA_DOORBELL_ATOMIC_LOAD`.

## Why

This is the "GPU not seeing the host's flag store" cost noted in #866 (and the visibility question in #667). On an Arc Pro B60, a system-scope `atomic_ref` load of `malloc_host` memory is still served from the GPU's cache, so a device spin does not see the host's store and runs to `kSpinMax` every time. That bound is what made the per-layer waits slow, and why a smaller `kSpinMax` looked faster but produced degenerate output (comment on #866).

A standalone, bounded host/GPU ping-pong (GPU waits for `flag == i`, answers `ack = i`; host times each round; 200 rounds, every device spin bounded) shows it directly:

| device read of the host flag | round trip (median) | device waits that hit the bound |
|---|---|---|
| `malloc_host`, system-scope atomic load (current) | 263 ms | 168 / 200 |
| `malloc_host`, atomic `fetch_add(0)` | 273 ms | 125 / 200 |
| `zeMemAllocHost` with `BIAS_UNCACHED`, atomic load | 263 ms | 192 / 200 |
| **`malloc_host`, `read_hint` uncached L1+L3** | **2.4 us** | **0 / 200** |

(The host side always saw the GPU's `ack` stores, so only the device-side load needed changing.)

## Effect in Strata

Same machine as my #866 comment: Arc Pro B60 24 GB (`8086:e211`) over OCuLink PCIe 4.0 x4, NEO 26.31.39395.14, oneAPI 2026.1.1, port commit `7ba023e` + #866 + this change, AOT `bmg-g21`. Flash-Next GSQ-RCO IQ2_XS, 32K context, int8 KV, `--spec 4` with MTP, 10,813 experts resident and 13,763 in the pinned host mirror, per-layer waits (no `STRATA_VERIFY_NO_HOST`):

| | before | after |
|---|---|---|
| decode, 300-token answer | 8.7 - 10.5 tok/s | **19.3 tok/s** |
| decode, RAM-arena mode (no `--stream-experts`) | 7.6 tok/s | **16.0 tok/s** |
| `NO_HOST=1` (all-device plan, no waits) | 8.1 tok/s | 8.2 tok/s (unchanged, as expected) |
| 21.5K-token prompt | ~59 s | ~58 s (unchanged) |

Output is unchanged in quality: on a 16-task tool-calling test set run twice, the same model scored 28/30 before and 28/29 after (the one miss is the same kind of model slip in both runs), so the data the host hands back is read correctly too, not only the flag.

## Not tested

- Any card other than one Arc Pro B60 (no B70 / B580 / multi-GPU runs).
- Windows, WSL2, JIT (non-AOT) builds.
- Whether `memory_scope::system` atomic loads behave differently on other driver versions; this change does not depend on that.

<details><summary>The ping-pong test (standalone, ~100 lines)</summary>

```cpp
// Doorbell fix candidates on the Arc B60 (05/10). Same bounded ping-pong as pingpong.cpp.
#include <sycl/sycl.hpp>
#include <sycl/ext/oneapi/backend/level_zero.hpp>
#include <sycl/ext/intel/experimental/cache_control_properties.hpp>
#include <level_zero/ze_api.h>
#include <atomic>
#include <chrono>
#include <cstdio>
#include <vector>
#include <algorithm>
namespace syclex = sycl::ext::oneapi::experimental;
namespace intelex = sycl::ext::intel::experimental;
using sys_u32 = sycl::atomic_ref<uint32_t, sycl::memory_order::relaxed, sycl::memory_scope::system>;
constexpr uint32_t kRounds = 200, kSpin = 200000;
using UCread = decltype(syclex::properties(intelex::read_hint<
    intelex::cache_control<intelex::cache_mode::uncached, syclex::cache_level::L1, syclex::cache_level::L3>>));
enum Mode { ATOMIC_LOAD = 0, UNCACHED_HINT = 1, RMW = 2 };

void run(sycl::queue& q, const char* name, uint32_t* mem, Mode mode) {
    uint32_t* flag = mem; uint32_t* ack = mem + 16;
    *flag = 0; *ack = 0; std::atomic_thread_fence(std::memory_order_seq_cst);
    uint32_t* timeouts = sycl::malloc_shared<uint32_t>(1, q); *timeouts = 0;
    auto ev = q.single_task([=] {
        syclex::annotated_ptr<uint32_t, UCread> uf(flag);
        for (uint32_t i = 1; i <= kRounds; ++i) {
            uint32_t s = 0, v = 0;
            for (; s < kSpin; ++s) {
                if (mode == ATOMIC_LOAD) v = sys_u32(*flag).load();
                else if (mode == UNCACHED_HINT) { sycl::atomic_fence(sycl::memory_order::acquire, sycl::memory_scope::system); v = uf[0]; }
                else v = sys_u32(*flag).fetch_add(0u);
                if (v >= i) break;
            }
            if (s == kSpin) sys_u32(*timeouts).fetch_add(1);
            sys_u32(*ack).store(i);
            sycl::atomic_fence(sycl::memory_order::release, sycl::memory_scope::system);
        }
    });
    std::vector<double> us; volatile uint32_t* vack = ack; volatile uint32_t* vflag = flag; uint32_t lost = 0;
    for (uint32_t i = 1; i <= kRounds; ++i) {
        auto t0 = std::chrono::steady_clock::now();
        *vflag = i; std::atomic_thread_fence(std::memory_order_seq_cst);
        auto lim = t0 + std::chrono::milliseconds(300);
        while (*vack < i && std::chrono::steady_clock::now() < lim) {}
        if (*vack < i) ++lost;
        us.push_back(std::chrono::duration<double, std::micro>(std::chrono::steady_clock::now() - t0).count());
    }
    ev.wait(); std::sort(us.begin(), us.end());
    std::printf("%-44s median %9.1f us  p90 %9.1f us  max %9.1f us  GPU timeouts %3u/%u  host misses %u\n", name,
                us[us.size()/2], us[us.size()*9/10], us.back(), *timeouts, kRounds, lost);
    sycl::free(timeouts, q);
}
uint32_t* ze_host(sycl::queue& q, ze_host_mem_alloc_flags_t flags) {
    auto zc = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(q.get_context());
    ze_host_mem_alloc_desc_t hd{ZE_STRUCTURE_TYPE_HOST_MEM_ALLOC_DESC, nullptr, flags};
    void* p = nullptr;
    if (zeMemAllocHost(zc, &hd, 4096, 4096, &p) != ZE_RESULT_SUCCESS) { std::printf("zeMemAllocHost failed\n"); return nullptr; }
    return (uint32_t*) p;
}
int main() {
    setvbuf(stdout, nullptr, _IONBF, 0);
    sycl::queue q{sycl::gpu_selector_v};
    auto zc = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(q.get_context());
    std::printf("device: %s\n", q.get_device().get_info<sycl::info::device::name>().c_str());
    uint32_t* h = sycl::malloc_host<uint32_t>(1024, q);
    run(q, "malloc_host, atomic load (Strata today)", h, ATOMIC_LOAD);
    run(q, "malloc_host, uncached L1+L3 read hint", h, UNCACHED_HINT);
    run(q, "malloc_host, atomic fetch_add(0)", h, RMW);
    sycl::free(h, q);
    if (uint32_t* u = ze_host(q, ZE_HOST_MEM_ALLOC_FLAG_BIAS_UNCACHED)) {
        run(q, "ze host BIAS_UNCACHED, atomic load", u, ATOMIC_LOAD);
        run(q, "ze host BIAS_UNCACHED, uncached hint", u, UNCACHED_HINT);
        run(q, "ze host BIAS_UNCACHED, fetch_add(0)", u, RMW);
        zeMemFree(zc, u);
    }
}
```
Build: `icpx -fsycl -O2 -fsycl-targets=spir64_gen -Xs "-device bmg-g21" pingpong2.cpp -o pingpong2 -lze_loader`
</details>

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