Repository navigation
Arc A770 (DG2) support and fixes: 4 GiB handling, host-visible flags, chunked mirror, XMX-off build, A770 docs - #1
Closed
samsonnguyen wants to merge 56 commits into
Conversation
…ix path IGC crashes with a floating point exception (ocloc -device dg2-g10) on the two joint_matrix kernels, xmx_gemm_iq and qsa_prompt_attn_xmx. Both are opt-in and slower than the default paths on the B70, and their callers already fall back when they return false. With -DSTRATA_SYCL_NO_XMX=ON both entry points become `return false` stubs and the kernels are not instantiated (SYCL extracts device code from the template instantiation, so an early return alone is not enough). Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…issing On the xe driver (Arc A-series) the oneAPI Level Zero adapter answers UR_DEVICE_INFO_GLOBAL_MEM_FREE with UR_RESULT_ERROR_UNINITIALIZED even with ZES_ENABLE_SYSMAN=1, so sycl::aspect::ext_intel_free_memory is never set, dpct::get_memory_info reports 0 bytes free and --expert-cache auto sizes the cache to zero slots. Level Zero Sysman itself answers correctly (zesMemoryGetState tracks allocations: 16268 MiB free, 13191 after a 3 GiB allocation, 16263 after freeing it). Fall back to zesMemoryGetState through the device's native Level Zero handle when the aspect is missing. Where the aspect exists the path is unchanged. strata_core now links ze_loader. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
--stream-experts mirrors the experts that did not fit VRAM in one sycl::malloc_host. A single host allocation is capped at about 4 GiB on Arc A-series (max_mem_alloc_size 4095 MiB; UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS lifts the device cap only), so a 14 GiB mirror failed with "no pinned host memory" and every miss then stayed on the SSD. The mirror is now a list of chunks of whole blobs, at most 3 GiB each, and each (layer, expert) keeps a direct pointer instead of an offset. If a later chunk cannot be allocated the earlier ones are kept as a partial mirror. On a 16 GB A770 all 7448 experts missing from VRAM were mirrored (14.2 GiB, 3.4 s). Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…ever reach route<NE> returns early for every work-item with local_id(1) != 0 and then calls a work-group barrier. CUDA lets exited threads satisfy __syncthreads; on Arc A-series (dg2) the barrier never completes, the first decode window hangs with the GPU busy and the xe driver eventually resets the compute engine. The barrier protects nothing (the values live in private arrays), so it is removed. Adds native_router_parity: the native router (256 and 512 experts, one and six tokens, gaussian / negative-offset / all-ties logits) against a host reference. The kernel had no test of its own; router_top10_parity covers the generic kernel only. 0 mismatches on the A770. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
The expert cache is one allocation larger than 4 GiB once the card has the room (9.2 GiB on a 16 GB A770). With UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS=1 the allocation succeeds, but IGC addresses a buffer with 32-bit offsets unless the kernel is compiled with -ze-opt-greater-than-4GB-buffer-required. Kernels reading a slot more than 4 GiB into the cache then misread it without any error: a resident IQ4_NL expert 5.96 GiB into the cache dequantised to NaN (285,388 non-finite values of 1.6M) while an identical expert in a small staging buffer dequantised correctly, which made the prompt path poison the residual stream from layer 1 on and the model emit only token 0. Xe2 addresses buffers statelessly, so the B70 never saw it. With -DSTRATA_SYCL_LARGE_BUFFERS=ON the AOT ocloc options carry the flag. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…eck macro's 0 verify.cpp and session.cpp ported cudaStreamQuery as `q = DPCT_CHECK_ERROR(queue.ext_oneapi_empty())`. The macro evaluates the expression for its error code and returns 0, so `q != 1` was always true and any layer that took longer than one flush interval (2 ms) without ringing was reported as "layer N never rang (graph finished)" while the graph was still running. `q` is now 1 while the queue still has work and 0 once it is empty, which is what the check was written for. On the A770 this was what stopped the host-handshake path (no STRATA_VERIFY_NO_HOST) at layer 1; with it the path runs. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…speed, open issues Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
GgufExpertSource::blob() entered the slot in where_ while holding the lock and then did the three preads outside it, so a second caller asking for the same expert in that window took the "still resident" branch and was handed a pointer to a half-filled buffer. The CPU pool's workers, or the two tokens of one verify window, do ask for the same expert at once. The slot is now reserved under the lock and entered in where_ once the read is complete; a concurrent caller reads the expert into a slot of its own. Found by reading the code while chasing the adaptive-swap corruption on the A770. It is a real race, but fixing it did not by itself change that failure. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
An adaptive swap (the default, --adapt-every 4) evicts a resident expert into the missing set. At startup only the experts that miss VRAM are mirrored in pinned host memory, so an evicted expert is a miss the GPU cannot reach on its own. The device plan then waits for the host to compute it, and where the device-to-host handshake does not hold (the A770: kernel writes to host-mapped memory are not reliably visible while the graph runs) the GPU's bounded wait gives up and the window carries on with garbage. The first window after the first swap produced NaN, which the recurrent GDN state keeps for good: the output was correct for four tokens and token 0 after that. Verified on the A770 that the swap itself was sound (the device residency table matched the host's, and the swapped-in slots were byte-identical to the GGUF), and that mirroring every expert makes the output correct across 16 swap rounds (1,536 experts swapped, tokens identical to the run without swaps). With swaps on, the resident experts are now mirrored after the real misses, within the same RAM cap, so a victim is always reachable. STRATA_MIRROR_ALL=1 forces it and STRATA_MIRROR_ALL=0 turns it off. It costs about 9 GiB more pinned host memory on the 16 GB card. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…ice-to-host wait problem Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
… every card for a split stage_room() never queried free VRAM: dpct dropped the cudaMemGetInfo inside an if-with-initializer, so every later stage of a --layer-split saw 0 bytes and its expert cache had "no room". Restored, plus a fixups.py rule. strata-sycl.sh forwards ONEAPI_DEVICE_SELECTOR and defaults to level_zero:gpu for a --layer-split (the image pins level_zero:0). Reported and verified on 2x B70 in the PR Niko1221#423 review. Also xmx_int8_bench: an int8 XMX IQ4_NL GEMM prototype - correct, but slower than dequant + oneMKL (not ported). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…ndard Intel bench INTEL.md keeps how the port works and how to run it; every measured speed moves to INTEL_PERFORMANCE.md (current numbers, dated history, llama.cpp on the same card), with the results posted on PR Niko1221#423 for the Arc Pro B50, the B580 and 2x B70 copied as their posters reported them. sycl/benchy.sh (v1): every set-up model x 20 / 2,185 / 8K / 40K / 128K / 256K-token prompts (sycl/bench/v1), each with its serve config, from a cold page cache. Reports PP, TTFT, TG, drafts accepted, experts offloaded to the RAM mirror, VRAM, RAM, SSD reads, load time and power, with the system and build. INTEL_PERFORMANCE.md asks submitters for that report plus their configs. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
The reference benchy v1 run: Coder IQ1_M (256K), IQ2_XS and Swift 1.5 (128K) x 20 / 2,185 / 8K / 40K / 128K / 256K-token prompts, cold page cache, every run complete. The 32K test-config numbers stay beside it, labelled. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
The verify window's device waits on the host's flags are bounded (kSpinMax) because an unbounded spin that never sees its flag wedges the card. A wait that gave up simply let the window go on without what the host was to write, and nothing said so: the output was garbage and the recurrent GDN state kept it for good (token 0 forever after one bad window). Each wait that gives up now counts itself, per flag (A: the plan, B: the PCIe copies, M: the CPU's results), in host-visible memory. After the window the verifier reads the counts and, if any wait gave up, prints which, prints the host's own view (how far it got, how often it saw the GPU ring, which flags it raised) and stops the window with an error. STRATA_WAIT_TIMEOUT=warn only reports. The bound can be set with STRATA_SPIN_MAX (default unchanged: 20,000 polls). This is how the adaptive-swap corruption and the missing host-to-GPU stores were pinned down on the A770: the waits gave up although the host reported raising every flag. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…l sees the host's stores The host raises flags A, B and M while a kernel on the GPU polls them. On the Arc A770 (xe driver, oneAPI 2026.1) memory from sycl::malloc_host does not work for that: a kernel polling it did not see a host store made after the kernel started, under any load (system-scope relaxed and acquire atomics, device-scope, volatile, with and without a system fence), after 60 million polls. The host reported raising every flag for all 48 layers while 34-46 GPU waits gave up per window, even with the bound raised 100x, and every layer then paid the whole bound. Memory from zeMemAllocHost with ZE_HOST_MEM_ALLOC_FLAG_BIAS_UNCACHED (or _BIAS_WRITE_COMBINED) was seen within microseconds in every trial. The same call with no flags was not reliable (it failed twice in a row and then worked, in one run; failed in others), which is why the flag is explicit. The flag words are now allocated that way (host_alloc_coherent, HostUse::kControl) on the Level Zero backend. Nothing else needs it: the plan and result rows are read by later kernels, and a kernel launch refreshes the GPU's view; making them uncached cost 42 ms per window. With STRATA_VERIFY_NO_HOST nothing waits on the host and the flags keep the plain allocation: the uncached flags alone cost the same 42 ms there for a reason that was not found. Measured on the A770, Coder IQ1_M, 128 tokens: the host-plan path (CPU computes part of the misses beside the GPU, the design the engine was written for) went from zeros at the first token to correct output at 8.0-8.6 tok/s; the device plan with the host handshake no longer produces garbage when an expert is not mirrored; the GPU-only path is unchanged at 7.14 tok/s. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…ok/s, what stays open Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…nger the limit Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…n GPU-only mode With STRATA_VERIFY_DEVICE_PLAN and STRATA_VERIFY_NO_HOST nothing on the host sees which experts a window routes to (the usage the adaptive tier ranks comes from the host pool), so the swaps never had anything to act on and the static profile's hit rate stayed at about a third: about two thirds of every expert read crossed PCIe from the pinned mirror. unitrace on the A770 showed it: the GPU was 88% busy, and 78 ms of the 141 ms of kernel time per decode round was native_gu_port / native_down_port reading mirrored experts at about 9 GB/s. resident_plan_kernel now counts, per (layer, expert), how often an expert is routed to (a relaxed device atomic per routed id, into a table indexed like the residency table; resident_plan_set_usage). Before each adaptation the engine reads the counters, clears them and adds them to the usage the adaptive tier ranks, in both the serve and the CLI loop. Off unless the run is GPU-only with adaptive swaps on. Measured on the A770, Coder IQ1_M, 128 tokens: 7.14 -> 9.76 and 9.74 tok/s (two runs), 9.77 at 256 tokens; kernel time per round 141 -> 100 ms (the expert kernels 78 -> 31 ms); output identical to the run without swaps. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…7 -> 10.37 tok/s A Q6_K block is 210 bytes, so every 16-byte load in the decode mat-vec starts at a 2-byte-aligned address and the kernel splits it into two aligned loads and a shift (load16_a2). The loaders now pad each block to 224 bytes on the host before the upload (type tag 114, native_mmvq_serving_type / native_mmvq_pack; STRATA_Q6K_STRIDE224=0 keeps the GGUF layout), and a variant of the wide kernel reads the padded blocks with plain aligned loads and the activations as scalar loads, at most four columns per launch (five or six in one launch spill registers: 0.14x). The prefill dequantiser reads tag 114 too; an unknown tag still throws instead of misreading. mmvq_bench (new check against the 210-byte kernel; bit-identical results), same build, two columns: 1.12-1.32x over the model's Q6_K shapes. Whole GPU-only decode on the A770, 128 tokens, identical tokens: 10.37 and 10.36 tok/s against 9.77 and 9.76. Load time is unchanged (the repack of about 1.9 GB is a fraction of a second); the Q6_K weights grow 6.7% (2,019 -> 2,105 MiB). Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
Upstream moved the B70 measurement tables to docs/INTEL_PERFORMANCE.md and rewrote the "Not done" list; the only textual conflict was docs/INTEL.md, resolved by taking upstream's tail and keeping the Arc A770 section after it. The sources auto-merged (generate.cpp, verify.cpp, iq_kernels.dp.cpp, CMakeLists.txt). Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
xmx_int8_bench (new upstream) is a joint_matrix int8 DPAS prototype; IGC for dg2 dies on it with a floating point exception at AOT link, the same as on the fp16 joint_matrix kernels, so the target now sits behind the option that already compiles those out. (Upstream measured it 0.38-0.98x of dequant + oneMKL on the B70: not worth an engine port.) Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…m merge Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…an the runner's CPU Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
….py it cites Fixed: the missing "How to build it" section (docker build + build.sh with AOT/REPO/BUILD_DIR), the strata-sycl.sh path, the llama.cpp-only scope of the persistent-cache rule, the kernel count, where the CUDA headers live (mounted for migration, not in the image), the draft layer's VRAM, stream-all's 1,024-token minimum chunk, the expert XMX switch that is not in the tree, the metrics image that is not in the repo, and the HAND-merge list's reasons. sycl/probe/*.cpp and sycl/tools/rank_kernels.py were cited but lived outside the repo; they are in it now. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
setup_intel.py --check judged each model by the first card only: with two cards the IQ2_XS/IQ3_XXS/IQ3_S fit in VRAM across both (layer split). Also names the Intel path where setup's AMD line still said docs/AMD_HIP.md. INTEL.md: a two-die layer split needs a comma-list AOT build (checked: bmg-g21,bmg-g31 compiles). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…them by DMA With --stream-experts the slots a prompt lends had their experts read from the GGUF for every chunk and again for the refill after the prompt. They are now mirrored once in pinned host memory beside the VRAM misses (the resident RAM mode's idea, upstream), and the stream-all walk DMAs pinned experts. Same bytes into the same slots: outputs identical (Coder 2,185 / 8K / 40K, IQ2_XS 2,185 / 8K, cold cache). Prompt reading on the served configs: Coder 456 -> 610 / 802 -> 946 / 988 -> 1,046 tok/s, IQ2_XS 303 -> 500 / 565 -> 720; refill 720-2,240 -> 200-380 ms. 2.3 GiB more RAM for the Coder, 1.9 GiB for the IQ2_XS; 2.1 s at start (taken out of the one-shot TTFT clock, added to benchy's load time). STRATA_LEND_MIRROR=0: off. Also: host-side timers for the prompt path's per-layer expert grouping (STRATA_PREFILL_TIMING=1). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…ot worth it Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
The warp kernel scored every (query, block) with a sub-group, re-reading the pooled key and spending 20 shuffles per pair (~1.1 TFLOP/s). The scores are now pooled keys x the batch's indexer queries as oneMKL fp32 GEMM tiles (8,192 blocks, 32 MB) plus a relu-sum kernel that also gives block n_bid its dead key. sel_scores_bench: 7.3-8.4x at 40K-256K, relative error 2-4e-7 against FP64. Coder TTFT: 256K 328.6 -> 258.2 s, 128K 138.9 -> 122.9 s. FP32-level, not bitwise (like the CUDA build's 3xTF32 path): outputs part at a near-tie after 13-93 tokens, the same text either way. STRATA_SELECT_GEMM=0: the warp kernel. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…ll chunks) The batched attention's warp kernel split a cell's 256 dims over a sub-group and reduced with a 5-step shuffle tree per query head (60 shuffles a cell). The prompt path now scores with one work-item per cell holding its key row and all 12 heads' dots, the query heads read from local memory (a broadcast), 128-cell chunks; the values pass and the merge are the warp kernel's. Decode keeps the warp kernel (bitwise unchanged). attn_bench (32 queries x 2,048 cells, INT8 KV, 131K context): 0.57 -> 0.37 ms; the K/V gather floor is 0.07 ms, so the kernel is arithmetic-bound, not gather-bound. Coder TTFT 40K 37.2 -> 33.5 s, 128K 123.0 -> 111.3 s, 256K 258.3 -> 235.1 s. FP32 in another summation order (not bitwise): outputs part at a near-tie after 25-44 tokens, the same text. kv_stream_parity and kv_hybrid_parity pass; qsa_prompt_attn_parity passes but for the Q4_0 tensor-core cases the port refuses (as before). All four KV formats. STRATA_ATTN_PERCELL=0: the warp kernel. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…ores, per-cell attention) INTEL.md describes the three changes and their switches, and corrects the old claim that the prompt attention is gather-bound (the K/V fetch is a fifth of its time). INTEL_PERFORMANCE.md adds their A/B measurements and the two new XMX results, and notes that the "host grouping" share of the 40K breakdown was mostly the profiler. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…ze; benchy's mirror column qsa_set_bytes (the count bytes_needed uses) lacked the 32 MB sel_S tile init allocates, so a chunk whose QSA buffer set is the largest (256 tokens: a 20-token prompt) failed with "device buffers ... do not fit". Benchy's "offloaded" column also read the lend mirror's line; it reads the VRAM misses' line only. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
Coder 256K TTFT 328.8 -> 235.1 s, 128K 140.9 -> 111.2 s; IQ2_XS / Swift 128K ~191 -> ~156 s. All 16 runs complete (the 20-token prompts included, after 16e5ae0). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…L >4 GiB GEMM bug and its reproducer Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…nd append support
…versions, PCIe link Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…replaces the machine write-up) Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…d cache safeguards
…rc-0.1.39 A/B on the A770 against 41ee109 (same config, four fixed greedy prompts, two passes): decode and prompt speed within +-3%, outputs differ at the token level but stay correct. The merged tree keeps all A770 fixes (audit counts unchanged). Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…t in INTEL.md Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
…nto the A770 branch strata-sycl.sh: Podman support combined with the explicit environment array; the empty options array no longer trips set -u on older bash. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Building on top of your perf branch - Niko1221#809.
If this isn't merged in, then i can request for a re-pull to Niko1221 upstream without the perf changes merged in.
What this is
The changes needed to run the SYCL port on an Arc A770 (Alchemist / DG2-G10, 16 GB,
xedriver), found while running the Coder IQ1_M and the original model's IQ3_XXS on it. It is based onintel-arc-0.1.39(everything from this branch is already merged in, so the diff is only the A770 work). Most items are bugs that silently produced garbage (token 0 forever, NaNs) rather than crashes, so each has a test or a written root cause.Also added:
docs/INTEL_A770.md(bring-up, what differs from the B70, measurements, the 4 GiB bug write-up) anddocs/INTEL_A770_ISSUES.md(a short catalogue of the Intel driver / Level Zero / IGC / oneMKL problems hit, with workarounds, for anyone who wants to report them upstream to Intel). The A770 material was moved out ofdocs/INTEL.md, which now points to it.Changes, grouped (can be split into separate PRs if you prefer)
Build for DG2
STRATA_SYCL_NO_XMX(IGC for dg2 crashes on thejoint_matrixkernels),STRATA_SYCL_LARGE_BUFFERS(-ze-opt-greater-than-4GB-buffer-required),STRATA_GGML_NATIVE(so a binary built on another CPU does not die with "Illegal instruction");ze_loaderlink; AOT builds needIGC_EnableDPEmulation=1 OverrideDefaultFP64Settings=1(no fp64 on the A770).sycl/include/dpct/device.hpp: Level Zero Sysman fallback for free-memory queries onxe.Host/GPU handshake and waits
zeMemAllocHostand the uncached flag; a kernel polling plainsycl::malloc_hostmemory never saw host stores on the A770. Bounded waits now fail loudly with the flag's name, with per-verifier timeout counters (wait_timeout_registry.hpp).Expert cache, mirror and prompt borrowing
intel-arc-0.1.39's lend mirror; ring slot ownership for the streaming ring (expert_ring_slots.hpp).sycl/tests/repro/mkl_gemm_4gib.cpp). With prompt borrowing and a cache over 4 GiB, prompts of ~1.5k tokens and up decoded to token 0. Fix: the expert cache is laid out in reverse (coldest slots, which are the ones lent to the prompt path, at the lowest addresses) and the loan is capped at 3.9 GiB. It switches on automatically when the arena exceeds 4 GiB;STRATA_CACHE_REVERSE=0/1forces it. Kernels still index slots through the offset table.STRATA_VERIFY_NO_HOSTmode).Kernels
STRATA_Q6K_REPACK=1;mmvq_bench --selftestcompares it against the unpacked kernel and fails on a mismatch.Other:
sycl/tools/Dockerfile.serve(optional serving image).How it was tested
On an A770 (Ubuntu 26.04, kernel 7.0,
xe, oneAPI 2026.1): builds clean; host-only tests (sycl/tests: ring slots, timeout registry, cache layout, mirror) pass;mmvq_bench --selftestpasses (max abs diff 4.9e-4, dequant exact). End to end: the Coder IQ1_M and the original model's IQ3_XXS serve correctly, including a 90k-token prompt at a 262144 context (prompts ~550 tok/s, decode 16-17 tok/s on a code prompt, 12-14 on real chats). A/B of this branch against the previous A770 build on four fixed greedy prompts: speed within +-3%, outputs correct. Numbers are indocs/INTEL_A770.md.Not tested, please look at
🤖 Generated with Claude Code
https://claude.ai/code/session_011CjbiQBP63VFfAvn5SA6tu
yes, AI was used. Combination of claude code and gpt