Skip to content
Closed
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
19 commits
Select commit Hold shift + click to select a range
6607546
sycl: the 0.1.39 port builds - #626's thread affinity and #559's per-…
Oct 4, 2026
1939dc1
sycl: benchy v1 (the standard Intel bench), the probe programs and th…
Oct 4, 2026
81e0c4f
sycl: prompt-path speedups - the lend mirror, QSA block scores as GEM…
Oct 4, 2026
3b88502
sycl: several Intel cards - setup --check sizes models across every c…
Oct 4, 2026
01efa8d
docs: INTEL.md as the port's reference; every measured speed in INTEL…
Oct 4, 2026
939e3c6
sycl: #559's conversation parking in the port's conversation_state.cp…
Oct 4, 2026
37d03d6
sycl: conversation_snapshot_test frees its pinned buffers with sycl::…
Oct 4, 2026
cb1f5cf
sycl: host-only copies go through the CPU - the Arc's copy engine hun…
Oct 4, 2026
babd71c
sycl: the kernel tests all pass on the B70 - iq/ple_parity registered…
Oct 4, 2026
bdb421d
sycl: run without Docker (a toolbox with oneAPI); benchy v1 on a Gen5…
Oct 5, 2026
99a6dd4
sycl: STRATA_SYCL_RUNNER=podman - the dev image under rootless Podman
Oct 5, 2026
05b5abc
docs: every INTEL_PERFORMANCE.md table re-measured on the Gen5 host; …
Oct 5, 2026
ee282f2
TODO: speedups suggested by another SYCL project on the B70 (oneDNN G…
Oct 5, 2026
eecc387
onednn_gemm_bench: oneMKL fp16 is not the gap on dense shapes; oneDNN…
Oct 5, 2026
70b3361
TODO: review maxious/Strata_SYCL's experiments on top of intel-arc-0.…
Oct 5, 2026
0be097d
TODO: maxious's DPAS / pre-unpack decode matvec re-run on the B65 - 1…
Oct 5, 2026
8ae1d75
sycl: Q6_K / Q5_K decode options from maxious/Strata_SYCL, off by def…
Oct 5, 2026
b3d6bc9
gpustat: count only the reported card's clients (two cards read as 16…
Oct 5, 2026
3f37281
sycl: the port refreshed to upstream 0.1.40 (main 82f46a8); #866 and …
Oct 6, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1,129 changes: 566 additions & 563 deletions docs/INTEL.md

Large diffs are not rendered by default.

903 changes: 903 additions & 0 deletions docs/INTEL_PERFORMANCE.md

Large diffs are not rendered by default.

1 change: 1 addition & 0 deletions sycl/.gitignore
Original file line number Diff line number Diff line change
Expand Up @@ -2,3 +2,4 @@
dpct_report.*
**/*.yaml
migrate.log
benchy-results/
95 changes: 90 additions & 5 deletions sycl/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,7 @@
#
# Device code is JIT-compiled from SPIR-V unless STRATA_SYCL_AOT names a device (bmg-g31 for the Arc Pro B70).
cmake_minimum_required(VERSION 3.24)
project(strata_sycl VERSION 0.1.39 LANGUAGES C CXX) # the upstream engine version this port tracks
project(strata_sycl VERSION 0.1.40 LANGUAGES C CXX) # the upstream engine version this port tracks
set(STRATA_ROOT ${CMAKE_CURRENT_SOURCE_DIR}/..)
set(CMAKE_CXX_STANDARD 20)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
Expand Down Expand Up @@ -81,7 +81,8 @@ target_include_directories(strata_includes INTERFACE
target_compile_options(strata_includes INTERFACE -Wall -Wextra)

# --- CPU-side pieces, identical to the top-level build ------------------------------------------------------------
strata_resolve(_spec src/spec/suffix_drafter.cpp src/spec/controller.cpp src/spec/draft_policy.cpp)
strata_resolve(_spec src/spec/suffix_drafter.cpp src/spec/controller.cpp src/spec/draft_policy.cpp
src/spec/draft_source.cpp) # 0.1.40
add_library(strata_spec STATIC ${_spec})
target_link_libraries(strata_spec PUBLIC strata_includes)

Expand Down Expand Up @@ -121,13 +122,15 @@ if(STRATA_NATIVE_EXPERTS)
add_subdirectory(${strata_llamacpp_SOURCE_DIR}/ggml ${CMAKE_BINARY_DIR}/ggml EXCLUDE_FROM_ALL)
endif()
strata_resolve(_cpu_native src/kernels/cpu/native_expert.cpp src/kernels/cpu/iq_avx512.cpp src/kernels/cpu/iq_avx2.cpp
src/kernels/cpu/kq_avx2.cpp) # 0.1.31: the Q4_K/Q5_K/Q5_1/Q8_0 expert rows
src/kernels/cpu/kq_avx2.cpp # 0.1.31: the Q4_K/Q5_K/Q5_1/Q8_0 expert rows
src/kernels/cpu/kq_avx1.cpp) # 0.1.40: the router lookahead's AVX1 dot
target_sources(strata_kernels_cpu PRIVATE ${_cpu_native})
target_include_directories(strata_kernels_cpu PRIVATE ${STRATA_ROOT}/third_party/ggml)
set_source_files_properties(${STRATA_ROOT}/src/kernels/cpu/iq_avx512.cpp PROPERTIES COMPILE_OPTIONS
"-mavx512f;-mavx512bw;-mavx512vl;-mavx512dq;-mavx512vnni;-mavx512vbmi;-mf16c;-mfma")
set_source_files_properties(${STRATA_ROOT}/src/kernels/cpu/iq_avx2.cpp PROPERTIES COMPILE_OPTIONS "-mavx2;-mfma;-mf16c")
set_source_files_properties(${STRATA_ROOT}/src/kernels/cpu/kq_avx2.cpp PROPERTIES COMPILE_OPTIONS "-mavx2;-mfma;-mf16c")
set_source_files_properties(${STRATA_ROOT}/src/kernels/cpu/kq_avx1.cpp PROPERTIES COMPILE_OPTIONS "-mavx")
target_link_libraries(strata_kernels_cpu PUBLIC ggml-cpu ggml-base)
target_compile_definitions(strata_kernels_cpu PUBLIC STRATA_NATIVE_EXPERTS=1)
endif()
Expand All @@ -154,10 +157,16 @@ target_include_directories(strata_kernels PRIVATE ${STRATA_ROOT}/third_party/ggm
target_link_libraries(strata_kernels PUBLIC strata_core)

strata_resolve(_engine src/core/layer.cpp src/core/session.cpp src/core/expert_source.cpp src/core/remote_experts.cpp
src/core/remote_expert_opt.cu # 0.1.40: the helper GPU's reduced decode rows
src/core/expert_cache.cpp src/core/native_head.cpp src/core/native_dense.cpp src/core/verify.cpp
src/core/mtp.cpp src/core/conversation_snapshot.cpp src/core/conversation_state.cpp
src/core/conversation_memory.cpp # 0.1.31: the conversation snapshots (#216)
src/core/peer_experts.cpp) # 0.1.36+: --peer-device (a stub in the port: no second GPU tier yet)
src/core/peer_experts.cpp # 0.1.36+: --peer-device (a stub in the port: no second GPU tier yet)
src/core/arch_defaults.cpp src/core/conversation_file.cpp) # 0.1.40: arch defaults, session files
# 0.1.40's virtual-memory K/V (--kv-grow, --vram-elastic) is CUDA driver API (cuMemMap): upstream's own no-VMM branch,
# the one HIP builds take, is compiled here - vmm_available() is false and the elastic K/V stays off
list(APPEND _engine ${STRATA_ROOT}/src/core/vmm.cpp)
set_source_files_properties(${STRATA_ROOT}/src/core/vmm.cpp PROPERTIES COMPILE_DEFINITIONS STRATA_USE_HIP=1)
add_library(strata_engine STATIC ${_engine} ${CMAKE_CURRENT_SOURCE_DIR}/src/core/gguf_expert_source.cpp)
target_include_directories(strata_engine PRIVATE ${STRATA_ROOT}/src/core) # conversation_checked.hpp (#include "...")
target_link_libraries(strata_engine PUBLIC strata_core strata_kernels strata_kernels_cpu)
Expand All @@ -169,6 +178,8 @@ add_library(strata_prefill STATIC ${_prefill}
target_compile_options(strata_prefill PUBLIC -qmkl=sequential)
target_link_options(strata_prefill PUBLIC -qmkl=sequential)
target_link_libraries(strata_prefill PUBLIC strata_engine)
# prefill.cpp includes mmq_resident_sort.hpp and wmma_gemm.h from upstream's src/prefill (headers only)
target_include_directories(strata_prefill PRIVATE ${STRATA_ROOT}/src/prefill)

strata_resolve(_main src/program/generate.cpp)
add_executable(strata ${_main})
Expand All @@ -184,8 +195,43 @@ if(STRATA_SYCL_PARITY)
add_executable(${p}_parity ${_src})
target_include_directories(${p}_parity PRIVATE ${STRATA_ROOT}/src/kernels) # the oracle .inc fixtures
target_link_libraries(${p}_parity PRIVATE strata_kernels)
add_test(NAME ${p}_parity COMMAND ${p}_parity --selftest)
if(NOT p STREQUAL "iq" AND NOT p STREQUAL "ple") # those two take inputs: registered below, as upstream does
add_test(NAME ${p}_parity COMMAND ${p}_parity --selftest)
endif()
endforeach()
# iq_parity compares against generated fixtures (tools/iq_fixture.py: numpy + the vendored gguf-py). As in the
# top-level CMakeLists: a setup test writes them, and both are a skip (exit 3) when the generator cannot run; with
# no python3 with numpy neither is registered. (Before, the port ran `iq_parity --selftest`, which read "--selftest"
# as the fixture directory and failed.) By hand: python3 tools/iq_fixture.py --out <dir> && iq_parity <dir>
set(_iq_python "")
foreach(_cand IN ITEMS ${Python3_EXECUTABLE} ${STRATA_ROOT}/.venv/bin/python3 python3)
if(NOT _cand)
continue()
endif()
execute_process(COMMAND ${_cand} -c "import numpy" OUTPUT_QUIET ERROR_QUIET RESULT_VARIABLE _iq_probe)
if(_iq_probe EQUAL 0)
set(_iq_python "${_cand}")
break()
endif()
endforeach()
if(NOT _iq_python STREQUAL "")
add_test(NAME iq_parity_fixtures COMMAND ${_iq_python} ${STRATA_ROOT}/tools/iq_fixture.py --out ${CMAKE_CURRENT_BINARY_DIR}/iq_fixture)
set_tests_properties(iq_parity_fixtures PROPERTIES FIXTURES_SETUP iq_fixture_data SKIP_RETURN_CODE 3)
add_test(NAME iq_parity COMMAND iq_parity ${CMAKE_CURRENT_BINARY_DIR}/iq_fixture)
set_tests_properties(iq_parity PROPERTIES FIXTURES_REQUIRED iq_fixture_data SKIP_RETURN_CODE 3)
else()
message(STATUS "STRATA SYCL tests: no python3 with numpy - iq_parity's fixtures and its test are not registered")
endif()
# ple_parity's oracles are real artifacts (the dense pack, the original GGUF, ggml's captured block in
# ple_in.bin / ple_out.bin): a missing one is a hard error inside the test, so it is registered only where the
# captures exist (STRATA_PLE_FIXTURE_DIR, as at the top level), with the GGUF from $STRATA_PLE_GGUF.
set(STRATA_PLE_FIXTURE_DIR "${STRATA_ROOT}/bench/micro" CACHE PATH "Directory with ple_parity's ple_in.bin / ple_out.bin")
if(EXISTS "${STRATA_PLE_FIXTURE_DIR}/ple_in.bin" AND EXISTS "${STRATA_PLE_FIXTURE_DIR}/ple_out.bin")
add_test(NAME ple_parity COMMAND ple_parity --selftest --in "${STRATA_PLE_FIXTURE_DIR}/ple_in.bin"
--out "${STRATA_PLE_FIXTURE_DIR}/ple_out.bin" WORKING_DIRECTORY ${STRATA_ROOT})
else()
message(STATUS "STRATA SYCL tests: no ple_in.bin / ple_out.bin in ${STRATA_PLE_FIXTURE_DIR} - ple_parity is not registered")
endif()
# 0.1.31: the IQ kernels that decode a weight part once for every column/entry, bitwise vs the per-column ones
# (no --selftest flag: it runs its checks by default), and the rewritten per-hit/grouped expert kernels vs the
# previous ones, bitwise
Expand All @@ -210,6 +256,24 @@ if(STRATA_SYCL_PARITY)
target_link_libraries(conversation_snapshot_test PRIVATE strata_engine)
add_test(NAME conversation_snapshot_test COMMAND conversation_snapshot_test)
set_tests_properties(conversation_snapshot_test PROPERTIES SKIP_RETURN_CODE 77 TIMEOUT 120)
# 0.1.40 (#783's multi-token kernels and the verify window): bitwise / reference checks on synthetic data
foreach(p IN ITEMS mmvq_multi native_multi gr_multi kv_steps resident_plan route_window qsa_topk_active)
strata_resolve(_src src/kernels/${p}_parity.cpp)
add_executable(${p}_parity ${_src})
target_link_libraries(${p}_parity PRIVATE strata_kernels)
add_test(NAME ${p}_parity COMMAND ${p}_parity)
endforeach()
set_tests_properties(qsa_topk_active_parity PROPERTIES SKIP_RETURN_CODE 77)
foreach(p IN ITEMS qsa_topk verify) # these two take --selftest
strata_resolve(_src src/kernels/${p}_parity.cpp)
add_executable(${p}_parity ${_src})
target_link_libraries(${p}_parity PRIVATE strata_kernels)
add_test(NAME ${p}_parity COMMAND ${p}_parity --selftest)
endforeach()
# 0.1.40's gemm_bf16_parity compares against cuBLAS (cublasGemmEx): no SYCL build of it
strata_resolve(_src src/kernels/fused_gr_bench.cpp) # fused GR timing (upstream's bench; takes arguments)
add_executable(fused_gr_bench ${_src})
target_link_libraries(fused_gr_bench PRIVATE strata_kernels)
# the prompt-attention parity takes a context length, not --selftest, and needs the tensor-core kernel (not
# ported yet: it reports "refused the pools" here). Built, not registered.
strata_resolve(_src src/kernels/qsa_prompt_attn_parity.cpp)
Expand All @@ -231,11 +295,32 @@ if(STRATA_SYCL_PARITY)
target_link_libraries(gr_bench PRIVATE strata_kernels)
add_executable(xmx_gemm_bench ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/xmx_gemm_bench.cpp) # XMX quantized GEMM vs dequant+oneMKL
target_link_libraries(xmx_gemm_bench PRIVATE strata_kernels strata_prefill) # -qmkl for the oneMKL reference
add_executable(xmx_int8_bench ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/xmx_int8_bench.cpp) # int8 XMX IQ4_NL GEMM prototype
target_link_libraries(xmx_int8_bench PRIVATE strata_kernels strata_prefill)
# Q6_K / Q5_K decode options from maxious/Strata_SYCL (their experiments 27/28/30/42, opt-in; see sycl/TODO.md 5):
# the one-time pre-unpack to signed bytes (STRATA_MMVQ_PREUNPACK=1) and the DPAS int8 matvec bench
add_executable(q6k_preunpack_parity ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/q6k_preunpack_parity.cpp)
target_link_libraries(q6k_preunpack_parity PRIVATE strata_kernels)
add_test(NAME q6k_preunpack_parity COMMAND q6k_preunpack_parity)
add_executable(q5k_preunpack_parity ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/q5k_preunpack_parity.cpp)
target_link_libraries(q5k_preunpack_parity PRIVATE strata_kernels)
add_test(NAME q5k_preunpack_parity COMMAND q5k_preunpack_parity)
add_executable(q6k_preunpack_bench ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/q6k_preunpack_bench.cpp)
target_link_libraries(q6k_preunpack_bench PRIVATE strata_kernels)
add_executable(q5k_preunpack_bench ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/q5k_preunpack_bench.cpp)
target_link_libraries(q5k_preunpack_bench PRIVATE strata_kernels)
add_executable(xmx_mmvq_bench ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/xmx_mmvq_bench.cpp)
target_link_libraries(xmx_mmvq_bench PRIVATE strata_kernels strata_prefill) # strata_prefill: oneMKL (the sweep)
add_executable(sel_scores_bench ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/sel_scores_bench.cpp) # QSA block scores: kernel vs oneMKL GEMM
target_link_libraries(sel_scores_bench PRIVATE strata_kernels strata_prefill)
add_executable(attn_bench ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/attn_bench.cpp) # prompt QSA attention variants
target_link_libraries(attn_bench PRIVATE strata_kernels)
# one native expert three ways on real GGUF rows (ggml reference, CPU path, GPU native_expert_grouped):
# native_expert_parity <shard1.gguf> [layer ...]
if(STRATA_NATIVE_EXPERTS)
add_executable(native_expert_parity ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/native_expert_parity.cpp)
target_include_directories(native_expert_parity PRIVATE ${STRATA_ROOT}/third_party/ggml)
target_link_libraries(native_expert_parity PRIVATE strata_kernels strata_kernels_cpu)
# 0.1.40's native_expert_bench is raw CUDA (cudaMalloc/cudaMemcpy): not built on SYCL
endif()
endif()
Loading