Commit dcbff2166 for llama.cpp
commit dcbff21661e5f5f1833bdd1b3a2517d7885f29da
Author: Neo Zhang <zhang.jianyu@outlook.com>
Date: Thu Nov 12 14:54:58 2026 +0800
sycl : fix AOT for more than 4GB (#31873)
* fix aot for >4g garbage output issue; add --aot in build script
* correct the aot flag
* support no-xmx build, update the code for no-xmx building, rename nossl to no-ssl
diff --git a/examples/sycl/build.sh b/examples/sycl/build.sh
index dc9e44a8e..c3017b6c5 100755
--- a/examples/sycl/build.sh
+++ b/examples/sycl/build.sh
@@ -4,55 +4,87 @@
# SPDX-License-Identifier: MIT
print_usage() {
- echo "Usage: ./build.sh [fp32|fp16] [ssl|nossl] [--help]"
+ echo "Usage: ./build.sh [fp32|fp16] [ssl|no-ssl] [no-xmx] [--aot aot_flags] [--help]"
echo ""
echo "Options:"
echo " fp32 Build with FP32 precision (default)"
echo " fp16 Build with FP16 precision (faster for long-prompt inference)"
echo " ssl Build with OpenSSL (default)"
- echo " nossl Build without OpenSSL"
+ echo " no-ssl Build without OpenSSL"
+ echo " no-xmx Disable XMX (DPAS) support"
+ echo " --aot Set SYCL AOT device architecture, e.g. arl-h,acm-g10,bmg-g21 or bmg-g21"
echo " --help Print this help message"
}
PRECISION=fp32
SSL=ON
+AOT_FLAGS=""
+SYCL_SUPPORT_XMX=ON
-for arg in "$@"; do
- case "$arg" in
+while [[ $# -gt 0 ]]; do
+ case "$1" in
--help)
print_usage
exit 0
;;
fp32|fp16)
- PRECISION="$arg"
+ PRECISION="$1"
+ shift
;;
ssl)
SSL=ON
+ shift
;;
- nossl)
+ no-ssl)
SSL=OFF
+ shift
+ ;;
+ no-xmx)
+ SYCL_SUPPORT_XMX=OFF
+ shift
+ ;;
+ --aot)
+ if [[ $# -lt 2 || "$2" == --* ]]; then
+ echo "Error: --aot requires aot_flags"
+ print_usage
+ exit 1
+ fi
+ AOT_FLAGS="$2"
+ shift 2
;;
*)
- echo "Error: unknown option '$arg'"
+ echo "Error: unknown option '$1'"
print_usage
exit 1
;;
esac
done
+AOT_OPTION=()
+if [[ -n "$AOT_FLAGS" ]]; then
+ AOT_OPTION=("-DGGML_SYCL_DEVICE_ARCH=$AOT_FLAGS")
+fi
+
mkdir -p build
cd build
source /opt/intel/oneapi/setvars.sh
if [ "$PRECISION" = "fp16" ]; then
#for FP16
- echo "cmake .. -DGGML_SYCL=ON -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DGGML_SYCL_F16=ON -DLLAMA_OPENSSL=$SSL # faster for long-prompt inference
-"
- cmake .. -DGGML_SYCL=ON -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DGGML_SYCL_F16=ON -DLLAMA_OPENSSL="$SSL" # faster for long-prompt inference
+ printf 'cmake .. -DGGML_SYCL=ON -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DGGML_SYCL_F16=ON -DGGML_SYCL_SUPPORT_XMX=%s -DLLAMA_OPENSSL=%s' "$SYCL_SUPPORT_XMX" "$SSL"
+ if [[ ${#AOT_OPTION[@]} -gt 0 ]]; then
+ printf ' %s' "${AOT_OPTION[@]}"
+ fi
+ echo " # faster for long-prompt inference"
+ cmake .. -DGGML_SYCL=ON -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DGGML_SYCL_F16=ON -DGGML_SYCL_SUPPORT_XMX="$SYCL_SUPPORT_XMX" -DLLAMA_OPENSSL="$SSL" "${AOT_OPTION[@]}" # faster for long-prompt inference
else
#for FP32
- echo "cmake .. -DGGML_SYCL=ON -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DLLAMA_OPENSSL=$SSL"
- cmake .. -DGGML_SYCL=ON -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DLLAMA_OPENSSL="$SSL"
+ printf 'cmake .. -DGGML_SYCL=ON -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DGGML_SYCL_SUPPORT_XMX=%s -DLLAMA_OPENSSL=%s' "$SYCL_SUPPORT_XMX" "$SSL"
+ if [[ ${#AOT_OPTION[@]} -gt 0 ]]; then
+ printf ' %s' "${AOT_OPTION[@]}"
+ fi
+ echo
+ cmake .. -DGGML_SYCL=ON -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DGGML_SYCL_SUPPORT_XMX="$SYCL_SUPPORT_XMX" -DLLAMA_OPENSSL="$SSL" "${AOT_OPTION[@]}"
fi
#build example/main
@@ -62,5 +94,6 @@ fi
#cmake --build . --config Release --target llama-bench
#build all binary
-echo "cmake --build . --config Release -j$((($(nproc)+1)/2)) -v"
-cmake --build . --config Release -j$((($(nproc)+1)/2)) -v
+BUILD_JOBS=$((($(nproc) + 1) / 2))
+echo "cmake --build . --config Release -j$BUILD_JOBS -v"
+cmake --build . --config Release -j"$BUILD_JOBS" -v
diff --git a/examples/sycl/win-build-sycl.bat b/examples/sycl/win-build-sycl.bat
index 4b3906a63..eac95ece0 100644
--- a/examples/sycl/win-build-sycl.bat
+++ b/examples/sycl/win-build-sycl.bat
@@ -6,6 +6,8 @@
SET "PRECISION=fp32"
SET "SSL=ON"
SET "OPENSSL_OPTION="
+SET "AOT_OPTION="
+SET "XMX_OPTION="
:PARSE_ARGS
IF "%~1"=="" GOTO ARGS_DONE
@@ -13,7 +15,9 @@ IF /I "%~1"=="--help" GOTO HELP
IF /I "%~1"=="fp32" GOTO ARG_FP32
IF /I "%~1"=="fp16" GOTO ARG_FP16
IF /I "%~1"=="ssl" GOTO ARG_SSL
-IF /I "%~1"=="nossl" GOTO ARG_NOSSL
+IF /I "%~1"=="no-ssl" GOTO ARG_NO_SSL
+IF /I "%~1"=="no-xmx" GOTO ARG_NO_XMX
+IF /I "%~1"=="--aot" GOTO ARG_AOT
GOTO INVALID_ARG
:ARG_FP32
@@ -29,11 +33,23 @@ SET "SSL=ON"
SET "OPENSSL_OPTION="
GOTO NEXT_ARG
-:ARG_NOSSL
+:ARG_NO_SSL
SET "SSL=OFF"
SET "OPENSSL_OPTION=-DLLAMA_OPENSSL=OFF"
GOTO NEXT_ARG
+:ARG_NO_XMX
+SET "XMX_OPTION=-DGGML_SYCL_SUPPORT_XMX=OFF"
+GOTO NEXT_ARG
+
+:ARG_AOT
+IF "%~2"=="" GOTO MISSING_AOT
+SET "AOT_VALUE=%~2"
+IF "%AOT_VALUE:~0,2%"=="--" GOTO MISSING_AOT
+SET "AOT_OPTION=-DGGML_SYCL_DEVICE_ARCH=%~2"
+SHIFT
+GOTO NEXT_ARG
+
:NEXT_ARG
SHIFT
GOTO PARSE_ARGS
@@ -45,15 +61,22 @@ GOTO PRINT_USAGE
:INVALID_ARG
SET "ARG_ERROR=1"
echo Error: unknown option '%~1'
+GOTO PRINT_USAGE
+
+:MISSING_AOT
+SET "ARG_ERROR=1"
+echo Error: --aot requires aot_flags
:PRINT_USAGE
-echo Usage: win-build-sycl.bat [fp32^|fp16] [ssl^|nossl] [--help]
+echo Usage: win-build-sycl.bat [fp32^|fp16] [ssl^|no-ssl] [no-xmx] [--aot aot_flags] [--help]
echo.
echo Options:
echo fp32 Build with FP32 precision ^(default^)
echo fp16 Build with FP16 precision ^(faster for long-prompt inference^)
echo ssl Build with BoringSSL ^(default^)
-echo nossl Build without BoringSSL
+echo no-ssl Build without BoringSSL
+echo no-xmx Disable XMX ^(DPAS^) support
+echo --aot Set SYCL AOT device architecture, e.g. arl-h,acm-g10,bmg-g21 or bmg-g21
echo --help Print this help message
IF "%ARG_ERROR%"=="1" exit /B 1
exit /B 0
@@ -70,10 +93,10 @@ if %errorlevel% neq 0 goto ERROR
IF /I "%PRECISION%"=="fp16" (
:: for FP16
:: faster for long-prompt inference
- cmake -G "MinGW Makefiles" .. -DLLAMA_BUILD_BORINGSSL=%SSL% %OPENSSL_OPTION% -DGGML_SYCL=ON -DCMAKE_CXX_COMPILER=icx -DBUILD_SHARED_LIBS=ON -DCMAKE_BUILD_TYPE=Release -DGGML_SYCL_F16=ON
+ cmake -G "MinGW Makefiles" .. -DLLAMA_BUILD_BORINGSSL=%SSL% %OPENSSL_OPTION% -DGGML_SYCL=ON %XMX_OPTION% %AOT_OPTION% -DCMAKE_CXX_COMPILER=icx -DBUILD_SHARED_LIBS=ON -DCMAKE_BUILD_TYPE=Release -DGGML_SYCL_F16=ON
) ELSE (
:: for FP32
- cmake -G "Ninja" .. -DLLAMA_BUILD_BORINGSSL=%SSL% %OPENSSL_OPTION% -DGGML_SYCL=ON -DCMAKE_C_COMPILER=cl -DCMAKE_CXX_COMPILER=icx -DBUILD_SHARED_LIBS=ON -DCMAKE_BUILD_TYPE=Release
+ cmake -G "Ninja" .. -DLLAMA_BUILD_BORINGSSL=%SSL% %OPENSSL_OPTION% -DGGML_SYCL=ON %XMX_OPTION% %AOT_OPTION% -DCMAKE_C_COMPILER=cl -DCMAKE_CXX_COMPILER=icx -DBUILD_SHARED_LIBS=ON -DCMAKE_BUILD_TYPE=Release
)
if %errorlevel% neq 0 goto ERROR
diff --git a/ggml/src/ggml-sycl/CMakeLists.txt b/ggml/src/ggml-sycl/CMakeLists.txt
index 35d36e583..8312a8305 100644
--- a/ggml/src/ggml-sycl/CMakeLists.txt
+++ b/ggml/src/ggml-sycl/CMakeLists.txt
@@ -221,7 +221,7 @@ if (GGML_SYCL_TARGET STREQUAL "INTEL")
if (NOT GGML_SYCL_DEVICE_ARCH)
target_link_options(ggml-sycl PRIVATE -Xs -ze-intel-greater-than-4GB-buffer-required)
else()
- message(STATUS "Skipping -ze-intel-greater-than-4GB-buffer-required for spir64_gen AOT")
+ message(STATUS "spir64_gen AOT: -ze-intel-greater-than-4GB-buffer-required passed to ocloc")
endif()
# Link against Intel oneMKL
@@ -251,7 +251,7 @@ if (GGML_SYCL_DEVICE_ARCH)
target_compile_options(
ggml-sycl PRIVATE
-fsycl-targets=spir64_gen
- "SHELL:-Xsycl-target-backend=spir64_gen \"-device ${GGML_SYCL_DEVICE_ARCH}\""
+ "SHELL:-Xsycl-target-backend=spir64_gen \"-device ${GGML_SYCL_DEVICE_ARCH} -options -ze-intel-greater-than-4GB-buffer-required\""
)
# Pass through parallel job (process) count for parallelising the
@@ -266,7 +266,7 @@ if (GGML_SYCL_DEVICE_ARCH)
target_link_options(
ggml-sycl PRIVATE
-fsycl-targets=spir64_gen
- "SHELL:-Xsycl-target-backend=spir64_gen \"-device ${GGML_SYCL_DEVICE_ARCH}\""
+ "SHELL:-Xsycl-target-backend=spir64_gen \"-device ${GGML_SYCL_DEVICE_ARCH} -options -ze-intel-greater-than-4GB-buffer-required\""
-fsycl-max-parallel-link-jobs=${GGML_SYCL_MAX_PARALLEL_LINK_JOBS}
)
diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp
index 99da9713e..f46e79784 100644
--- a/ggml/src/ggml-sycl/common.hpp
+++ b/ggml/src/ggml-sycl/common.hpp
@@ -237,9 +237,9 @@ typedef sycl::float2 dfloat2;
#define MMVQ_MAX_BATCH_SIZE 8
-#if defined(__INTEL_LLVM_COMPILER) && __has_include(<sycl/ext/intel/esimd/xmx/dpas.hpp>)
+#if defined(GGML_SYCL_SUPPORT_XMX) && defined(__INTEL_LLVM_COMPILER) && __has_include(<sycl/ext/intel/esimd/xmx/dpas.hpp>)
# define GGML_SYCL_MMVQ_HAS_XMX
-#endif // __INTEL_LLVM_COMPILER
+#endif // GGML_SYCL_SUPPORT_XMX && __INTEL_LLVM_COMPILER
// default of the most columns the XMX mul_mat_vec_q handles (GGML_SYCL_XMX_MAX_COLS), wider batches use other kernels
#define GGML_SYCL_XMX_MAX_COLS_DEFAULT 80
diff --git a/ggml/src/ggml-sycl/esimd.hpp b/ggml/src/ggml-sycl/esimd.hpp
index 85cc93e21..568173e7a 100644
--- a/ggml/src/ggml-sycl/esimd.hpp
+++ b/ggml/src/ggml-sycl/esimd.hpp
@@ -2,11 +2,14 @@
#define GGML_SYCL_ESIMD_HPP
#include <sycl/ext/intel/esimd.hpp>
-#include <sycl/ext/intel/esimd/xmx/dpas.hpp>
#include <sycl/ext/intel/experimental/grf_size_properties.hpp>
#include "common.hpp"
+#ifdef GGML_SYCL_MMVQ_HAS_XMX
+#include <sycl/ext/intel/esimd/xmx/dpas.hpp>
+#endif
+
namespace ggml_sycl_esimd {
constexpr int GGML_SYCL_DMMV_ESIMD_WG_SIZE = 4;
@@ -586,6 +589,7 @@ template <> struct esimd_reorder_q_traits<GGML_TYPE_Q6_K> {
}
};
+
//
// Multi-column mul_mat_vec for reordered K-quant weights on the XMX engines (int8 DPAS).
// A thread owns 16 weight rows (B operand, one dword is 4 weights of a row), the q8_1 activation
@@ -1024,6 +1028,7 @@ template <typename T, int NC, bool FUSED> constexpr int xmx_grf() {
return FUSED || NC >= T::nc_wide ? 256 : 128;
}
+#ifdef GGML_SYCL_MMVQ_HAS_XMX
// One DPAS of a block of RC columns: the int32 products of the weights B and the activations A, as float and scaled
// per weight row. Every quant type sums one or two of these per group of 32.
template <int RC, int NT>
@@ -1394,6 +1399,8 @@ static void xmx_mul_mat_ncols(const void * vx,
}
}
+#endif // GGML_SYCL_MMVQ_HAS_XMX
+
} // namespace ggml_sycl_esimd
#endif // GGML_SYCL_ESIMD_HPP