Pull requests / #1746

#1746 SYCL: the PCIe probe copies from pinned host memory, as the CUDA engine's does (#1745)

open · @sam-fakhreddine · 0 comments · View on GitHub

BenchmarksMulti-GPUAMD / HIPNVIDIA / CUDAModels & quants

Description

## Title
Issue: Resolves #1745

## Summary
The SYCL engine's PCIe probe now copies from pinned host memory (`sycl::malloc_host`), as the CUDA engine's probe does with `cudaMallocHost` and as the function's own comment describes. With `malloc` it timed the driver's pageable staging copy: an Arc Pro B70 on PCIe 5.0 x16 read 7.7 GB/s and got `pcie_frac` 0.21; with the fix it reads 35.6 GB/s and keeps the 0.55 default.

## What changed
- `sycl/src/program/generate.cpp`, `probe_pcie_h2d_gbps()`: the host buffer comes from `sycl::malloc_host` on the same in-order queue instead of `malloc`, the allocation also checks for a null pointer, and both frees are `sycl::free`. A two-line comment says why. Nothing else in the function changes (four 256 MiB bursts, the fastest counts, the `bw / 20` share).

What it leaves alone: the CUDA and HIP engines (already pinned), the share formula, `--pcie-frac` (an explicit value still skips the probe), and every kernel and switch.

## Extra Notes
How it was checked, on an Arc Pro B70 32 GB (PCIe 5.0 x16, Proxmox passthrough VM, Ubuntu 26.04, xe driver, oneAPI 2026.1.1, NEO 26.05.37020.3), IQ2_XS native pack, the documented command (`--spec 4 --spec-min-p 0.5 --kv int8 --prefill 4096 --max-context 32768`, `STRATA_VERIFY_NO_HOST=1 STRATA_VERIFY_DEVICE_PLAN=1`), greedy, 256 new tokens:

- Build: SYCL only (the file is in `sycl/`), AOT `-DSTRATA_SYCL_AOT=bmg-g31`, 0 errors.
- Probe line before: `PCIe probe: 7.7 GB/s host->device (best of 7.5 7.5 7.0 7.7) -> pcie_frac 0.21 (default 0.55)`.
- Probe line after: `PCIe probe: 35.6 GB/s host->device (best of 27.0 35.6 33.0 31.8) -> pcie_frac 0.55 (default 0.55)`.
- Output tokens identical to main on the 20-token and 2,185-token community prompts (`bench/results/2026-10-07-community-arc-b65-v01402/benchy-{short,long}.ids`).
- Decode, main with `--pcie-frac 0.55` (what the fix now picks) against main's probed 0.18, medians of 3 interleaved pairs: 62.6 vs 61.0 tok/s (20-token prompt), 60.3 vs 60.1 (2,185-token prompt). A small difference on IQ2_XS, where 4 to 10% of routed entries come from the mirror; the share matters more for a quant with more experts in the mirror.
- A standalone SYCL copy test on the same card: pinned 34.7 to 35.3 GB/s at 256 MiB to 1 GiB, pageable 8.1 to 14.9 GB/s (table in #1745).

Not tested: other Intel cards, bare metal (only the passthrough VM), and the layer-split path's per-stage probes (`probe_pcie_h2d_gbps` is shared, so they get the same buffer).

Related on strata.com

Editorial links to help you install, pick models, or read release notes — not part of the upstream thread.