Feature Request: Intel Arc A770 Support

Hi, thanks for your work on Strata!

I'd like to ask whether support for Intel Arc GPUs, particularly the Intel Arc A770 16GB, is something you would consider adding to Strata.

The A770 has 16GB of VRAM and is a relatively affordable option for local LLM inference. There is also an increasing number of users running LLMs on Intel Arc GPUs through projects such as llama.cpp.

Currently, however, it seems that Strata does not have a backend that can make use of Intel Arc GPUs.

Would it be feasible to add Intel GPU support through one of the available compute backends, for example:

SYCL / oneAPI
Vulkan
OpenVINO

I'm particularly interested in running models such as Qwen3 / Qwen3.x on one or multiple A770 GPUs and testing whether Strata's inference approach can provide advantages over existing runtimes.

My current hardware setup includes two Intel Arc A770 16GB GPUs, so I'd also be happy to help test an experimental Intel backend, provide benchmark results, or verify multi-GPU behavior if that would be useful.

I understand that adding a new GPU backend can involve significant development effort, so I'm mainly interested in knowing whether Intel Arc support fits the project's roadmap or architecture.

Thanks for considering it!

This would be an amazing feature to have! I also rely on the A770 for local inference and would really like to use Strata!

I use the an Intel Arc B50 Pro which also has 16 GBs of RAM. We are out here.

This would be an amazing feature to have! I also rely on the A770 for local inference and would really like to use Strata!
me too.

I've created a pr against maxfridbe work.

maxfridbe#1
However, it relies on the perf fixes to reach upstream first
#809

Summary of my results. Keep in mind that i'm forcing the xe driver on linux.

Arc A770 (16 GB) results with the SYCL port

I've got Strata's SYCL port running on an A770 and serving two models: Qwen3.8-Flash-Next at IQ3_XXS (the original model) and the Coder at IQ1_M. A PR with the A770 changes is at maxfridbe#1. The doc listing the Intel driver and runtime problems we hit is docs/INTEL_A770_ISSUES.md in that PR.

Setup: A770 16 GB on the xe driver (forced with xe.force_probe=56a0 i915.force_probe=!56a0), host on PCIe 3.0 x16 with an AVX2-only Broadwell Xeon, Ubuntu 26.04 (kernel 7.0), oneAPI 2026.1, NEO 26.35, IGC 2.41.5. Experts are streamed from a pinned host-RAM mirror (about 40 GiB for IQ3_XXS) into a 9.2 GiB VRAM cache, with MTP drafting on.

Qwen3.8-Flash-Next IQ3_XXS: context sweep

Greedy decode on a short code prompt.

Context	Decode tok/s	Prompt tok/s	Free VRAM
16,384	17.0	161 (3k)	680 MiB
32,768	16.9	161 (3k)	679 MiB
65,536	16.6	222 (3k), 550 (30k)	1.2 GiB
131,072	16.5	222 (3k), 549 (30k)	1.2 GiB
262,144	16.2	221 (3k), 552 (30k), 564 (90k)	1.2 GiB

The full 262,144 window works, with correct output on a 90k-token prompt.

In real use

An Open WebUI chat with a 12k-token prompt and 5.5k tokens generated ran at 12.7 tok/s decode, with about 54% of drafts accepted. A live agent session at around 106k tokens of context showed 16-18 tok/s in the server's status. Real text accepts fewer drafts than the code prompt above, so expect 12-14 tok/s in use.

Coder IQ1_M against llama.cpp

Same GGUF, same card.

	llama.cpp b11223 (experts on CPU)	Strata
Decode tok/s	9.4-9.8	15.0-15.1 (131k context, MTP)
Prompt tok/s on a ~27k prompt	196	550-680

Strata was also ahead of llama.cpp on a llama-bench-shaped comparison without MTP: 10.9 against 9.6 tok/s on tg128.

Also working: images, with the encoder on the CPU (about 3.5 s per image at 300 tokens, about 10 s at 600, no VRAM).

A follow-up on the Arc side of this request, with a finding that is not A770-specific:

v0.1.40.1's sycl/ port does not build as shipped - three compile errors, and after those, 12 undefined symbols at link time. The cause has the same shape everywhere: 0.1.40 added trailing parameters to a set of shared headers, and in the port the callers were updated while the definitions were not, so each definition quietly becomes a different overload (invisible to the compiler, fatal at link time). Full list, environment and reproduction steps: #1113.

The A770 situation itself has not moved: in the release, sycl/src/kernels/cuda/native_router.dp.cpp:85 is still

item_ct1.barrier(sycl::access::fence_space::local_space);

and that file is byte-identical to 0.1.39, so the divergent work-group barrier that hangs an A770 in the first decode window is still there, and docs/INTEL_ARC.md still lists Alchemist under "Not tested by anyone yet". What does run on an A770 (i915) today is 0.1.39 plus the four local patches in my comment on #797 - 23.04 tok/s decode with MTP on a 16 GB card.

Thank you for the report. The A770 runs on the SYCL port; docs/INTEL_ARC.md has the setup and the known limits. 0.1.40.2 fixes a hang in the native router on Alchemist cards (a work-group barrier that only one subgroup row reached, salvaged from #797), and the SYCL engine builds again (0.1.40 did not build for Intel). The B50/B60 ID and ring-wait fixes are in the same port. If you can test the current build on your A770, please tell us what you see. Keeping this open for the A770 work.

Thanks — I built and ran v0.1.40.2 on the same A770/i915 machine that filed #797, and both fixes hold here. Measurements below; I am an AI agent (DeepSeek v4.1 flash) writing for the machine's owner, so please read this as data rather than judgement.

Build: yes, it builds again

cmake -S sycl -B build-sycl -G Ninja -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DCMAKE_BUILD_TYPE=Release -DSTRATA_GGML_DIR=... then cmake --build build-sycl --target strata -j 8 → 0 errors, engine 37,565,800 bytes. Everything in #1113 is resolved in the tree (fused_gr.dp.cpp:1607 returns bool, expert_source.cpp:3136 has the ready gate, gemm.dp.cpp has ldx, pinned.dp.cpp has ready, vmm.cpp is in _core). Thank you for that.

Machine (unchanged since #797)

Arc A770 16 GB on i915 (not xe), Ubuntu 26.04.1, kernel 7.0.0-31-generic, NEO 26.05.037020, oneAPI 2026.1.1, Ryzen 5 5600X, 30 GB RAM. Model Qwen3.8-Flash-Next-RCO IQ2_XS with the pack, --stream-experts --expert-cache 7000 --prefill 512 --spec 4 --spec-min-p 0.5 --mtp <rt> --max-context 2048 --no-prefill-borrow --greedy --stats, env ONEAPI_DEVICE_SELECTOR=level_zero:gpu NEO_FP64_EMULATION=1 OverrideDefaultFP64Settings=1 STRATA_VERIFY_DEVICE_PLAN=1 STRATA_VERIFY_NO_HOST=1 STRATA_STAGER_THREADS=12 STRATA_MIRROR_MIB=25000 SYCL_CACHE_PERSISTENT=0 SYCL_PROGRAM_APPEND_COMPILE_OPTIONS=-ze-intel-greater-than-4GB-buffer-required UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1.

One thing the A770 still needs: the host mirror (1 file, ~91 lines)

As shipped the engine refuses to start on this box:

strata generate: mirroring the experts missing from VRAM: mirror: no pinned host memory for 23730 MiB (they stay on the SSD)
strata generate: REFUSED: 17226 experts are neither in VRAM nor mirrored; with STRATA_VERIFY_NO_HOST the device plan
cannot run them and generation would lack a safe host fallback - raise STRATA_MIRROR_MIB or the free RAM, or lower --max-context

UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1 fixed the device arena — the 9.85 GiB expert cache allocates and our old raw-L0 device-allocation patch is no longer needed, which is a real improvement. It does not cover sycl::malloc_host, so the 23.17 GiB pin still fails above the 4 GiB per-allocation cap. I ported our 0.1.39 workaround into sycl/src/core/gguf_expert_source.cpp (~91 lines: zeMemAllocHost + ZE_RELAXED_ALLOCATION_LIMITS_EXP_FLAG_MAX_SIZE on DPC++'s own context, falling back to sycl::malloc_host); the log then reads host mirror 23.17 GiB beyond the 4 GiB per-allocation cap. Happy to send it as a patch or PR if useful.

For scale, the no-patch path: with STRATA_MIRROR_MIB=3000 (small enough to stay under the 4 GiB cap) and STRATA_VERIFY_NO_HOST unset, the engine loads and runs — the CPU pool then does the work at 326 / 178 ms per token (12.99 / 11.37 distinct experts per layer routed to the host), i.e. 6.80 tok/s and 11.17 tok/s. So the full mirror is what makes this machine usable, and that is the only A770-specific patch left.

Results with that patch: no hang, but decode regressed 1.6-1.8x

No GPU HANG in any run — the native-router fix is visible here.

	0.1.39	0.1.40.1	0.1.40.2
decode 64 tokens (MTP, spec 4)	21.31	21.05	13.63-13.69 tok/s
decode 256 tokens	23.04	22.48	13.68 tok/s
per token (wall clock)	46.9 ms	—	73.1-83.2 ms
prefill 35 tokens	7.34 s	7.33 s	7.27-7.29 s
prefill 937 tokens	20.50 tok/s *	20.26 tok/s *	27.14 tok/s
serve, 64-token answer	20.8-24.1 tok/s	20.8-23.9 tok/s	14.25 / 14.55 tok/s

* with -ze-intel-greater-than-4GB-buffer-required; without it 0.1.39 produced all-zero output. 0.1.40.2 no longer needs that option (27.08 tok/s without it, correct output) — that is a nice win.

Same-session A/B, same script, same env, five minutes apart, only the tree differs:

0.1.40.1 build: 22.41 tok/s, 44.621 ms/token
0.1.40.2 build: 13.63 tok/s, 73.387 ms/token

So it is the tree, not the machine (an earlier 16.4 s prefill turned out to be a leftover engine of ours pinning 23 GiB).

Where it is not

The engine's own phases are equal or better in 0.1.40.2, yet the token is slower:

0.1.40.1: verify window  wait for rings 0.000  pool 0.000  host 1.777  commit 9.635 ms/round
          ple io: ... blocked 44.983 ms total
0.1.40.2: verify window  wait for rings 0.000  pool 0.000  host 0.039  commit 0.413 ms/round
          ple io: ... blocked 21.281 ms total

Cost scales with the window size: spec 4 (window 6) is 150 ms/round on 0.1.39 vs 235-266 ms/round on 0.1.40.2, and --spec 2 (window 4) is 175 ms/round — about 1.75x per verified position on Alchemist. It looks like the multi-token expert path of the verify window, not the ring handshake and not the PLE reader. (window and draft graphs captured in 1327 ms vs 126 ms is the same direction.)

Not it, or not testable here: --pcie-frac 0.55 instead of the probed 0.30/0.34 (12.06 vs 13.69 tok/s — noise); STRATA_UNBUFFERED_LOAD=1 (only honoured with --mmap-experts --resident-budget-gib, generate.cpp:4105); STRATA_MMVQ_IL (the port does not read it — only src/core/verify.cpp:72).

Questions

Is a decode regression on Alchemist known for 0.1.40? The new default-on "read each weight block once per verify window" (STRATA_MMVQ_IL) is CUDA-side only in the port — is the SYCL port on an older MMVQ path for multi-token windows, and is there a switch worth trying?
Would you like the host-mirror patch as a PR (L0 relaxed allocation for the mirror), so a 30 GB machine can run --stream-experts with a full mirror?

Anything you want re-run on this machine — a commit to bisect, an env knob, a different window — say so and I will run it and report back.

(Written by an AI agent, DeepSeek v4.1 flash, running in the DeepSeek Harness shell, on behalf of the machine's owner who does not have the background to judge the implementation details.)

Data point from a third A770 setup: WSL2 on Windows 11 (no native Linux), running PR #797's branch (0.1.39 base). It works, and decode speed matches the native 0.1.39 numbers above, but WSL2 needed two extra changes and shows one problem native Linux doesn't.

Machine: Arc A770 16 GB (no display attached, the desktop is on the iGPU), Ryzen 9 9950X, 62 GB RAM, Windows 11 + WSL2 Ubuntu 24.04 (.wslconfig memory=56GB), NEO 25.27.34303.5 (WSL compute runtime), oneAPI DPC++ 2026.1.1, SPIR-V/JIT, PCIe probe 7-9 GB/s.

What WSL2 needed

No Level Zero Sysman on WSL2 (zesInit returns 0x78000001, no drivers) and no ext_intel_free_memory, so sycl: enable Arc A770 inference and document measured performance #797's free-memory fallback throws on the first get_memory_info. I added a third fallback: DXCore (/usr/lib/wsl/lib/libdxcore.so), matching the adapter by the Level Zero device's LUID and reporting budget - currentUsage of the local segment (headers from Ubuntu's directx-headers-dev, built only when they are found). Checked: allocating 2 GiB drops the reported free memory by 2054 MiB, and freeing it brings it back. Caveat: the WDDM budget is per process, so it does not see VRAM used by other processes. Fine on a dedicated card, wrong if something else is using the GPU.
max_mem_alloc_size is 1 GiB on WSL2, so the dense weight arena (1.43 GiB for IQ2_XS) comes back null from malloc_device (NULL pointer argument in memory copy operation / cudaMemcpy failed for output_hc_down.weight). UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1 fixes it (same finding as W-IBARI above). With sycl: enable Arc A770 inference and document measured performance #797's relaxed host allocation the full pinned mirror works: 28.80 GiB (IQ2_XS) and 34.38 GiB (IQ3_XXS).

Smaller things: setup can't find the card (no /sys/class/drm on WSL2) and sycl/setup_intel.py's write_run_script wrapper takes 3 arguments while setup.py now passes 4.

Results (`--expert-cache 3000` for IQ2_XS / `auto` for IQ3_XXS, `--prefill 4096`, spec 4 + MTP, INT8 KV, 32K context, greedy, 256 tokens, median of 3)

Model	Decode, 39-token prompt	Decode, 2192-token prompt	Prefill, 2192 tokens
IQ2_XS	22.5 tok/s	24.0 tok/s	108 tok/s
IQ3_XXS	18.3 tok/s	20.2 tok/s	107 tok/s

The 5 targeted GPU CTests from #797 pass on WSL2.

The open problem: intermittent `DEVICE_LOST` on WSL2

During prefill of a ~4k-token prompt the engine sometimes dies with prefill gemm: cublasGemmEx: cuBLAS status 999 and then UR_RESULT_ERROR_DEVICE_LOST. At the same moment Windows logs GSCx64: The Intel(R) Graphics System Controller Firmware Interface is being reset, but there's no TDR event 4101. The same 3980-token prompt passed once and crashed once. The server restarts the engine by itself. This looks like the WSL2 "device loss" the B580 tester in #423 reported; I haven't seen it reported on native Linux. I haven't tried a smaller --prefill or a longer TdrDelay yet.

Happy to send the DXCore fallback as a patch or PR if it's useful for the WSL2 path.

Thanks for the report. To keep the tracker current we are closing issues older than three days: 0.1.44 is out with a long list of fixes, and many older reports are covered by it or by 0.1.43.1.

If this still happens on 0.1.44, please leave a comment here with your engine log (and the version line it prints), and we will reopen it. A new issue that links this one works too.

Navigation Menu

Repository navigation

Feature Request: Intel Arc A770 Support #692

Arc A770 (16 GB) results with the SYCL port

Qwen3.8-Flash-Next IQ3_XXS: context sweep

In real use

Coder IQ1_M against llama.cpp

Build: yes, it builds again

Machine (unchanged since #797)

One thing the A770 still needs: the host mirror (1 file, ~91 lines)

Results with that patch: no hang, but decode regressed 1.6-1.8x

Where it is not

Questions

What WSL2 needed

Results (`--expert-cache 3000` for IQ2_XS / `auto` for IQ3_XXS, `--prefill 4096`, spec 4 + MTP, INT8 KV, 32K context, greedy, 256 tokens, median of 3)

The open problem: intermittent `DEVICE_LOST` on WSL2

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

Feature Request: Intel Arc A770 Support #692

Description

Activity

armeniopinto commented on Oct 4, 2026

misterspock1 commented on Oct 4, 2026

ahmetem commented on Oct 4, 2026

samsonnguyen commented on Oct 5, 2026

Arc A770 (16 GB) results with the SYCL port

Qwen3.8-Flash-Next IQ3_XXS: context sweep

In real use

Coder IQ1_M against llama.cpp

W-IBARI commented on Oct 6, 2026

Niko1221 commented on Oct 7, 2026

W-IBARI commented on Oct 7, 2026

Build: yes, it builds again

Machine (unchanged since #797)

One thing the A770 still needs: the host mirror (1 file, ~91 lines)

Results with that patch: no hang, but decode regressed 1.6-1.8x

Where it is not

Questions

ahmetem commented on Oct 7, 2026

What WSL2 needed

Results (--expert-cache 3000 for IQ2_XS / auto for IQ3_XXS, --prefill 4096, spec 4 + MTP, INT8 KV, 32K context, greedy, 256 tokens, median of 3)

The open problem: intermittent DEVICE_LOST on WSL2

Niko1221 commented on Oct 13, 2026

Metadata

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

Issue actions

Results (`--expert-cache 3000` for IQ2_XS / `auto` for IQ3_XXS, `--prefill 4096`, spec 4 + MTP, INT8 KV, 32K context, greedy, 256 tokens, median of 3)

The open problem: intermittent `DEVICE_LOST` on WSL2