From 0853cea20a5e85705f69884feee405a9510bd1ff Mon Sep 17 00:00:00 2001 From: Jhin Lee Date: Thu, 24 Sep 2026 08:07:06 -0400 Subject: [PATCH 1/4] ci: build Windows ARM64 with Visual Studio 2026 GitHub is migrating the windows-11-arm label to the Windows 11 Arm64 with Visual Studio 2026 image between 2026-09-21 and 2026-09-30 (actions/runner-images#14602). That image has no VS 2022 instance, so the windows-arm64-full preset's hardcoded "Visual Studio 17 2022" generator fails at configure. Native release run 35976011426 lost both arm64 lanes (blas, vulkan) this way; the same lanes passed on the VS 2022 image in run 35850331498 on the same commit. Switch the preset to "Visual Studio 18 2026" and pin the arm64 jobs to windows-11-vs2026-arm so the image cannot flip mid-rollout. ClangCL, the ARM64 and x64 MSVC tools, and CMake 4.4 are all present on that image. --- .github/workflows/native_release.yml | 4 ++-- .github/workflows/validate_wrapper.yml | 2 +- CMakePresets.json | 2 +- docs/platform_backend_strategy.md | 6 ++++++ 4 files changed, 10 insertions(+), 4 deletions(-) diff --git a/.github/workflows/native_release.yml b/.github/workflows/native_release.yml index b76e433..56f6f10 100644 --- a/.github/workflows/native_release.yml +++ b/.github/workflows/native_release.yml @@ -697,7 +697,7 @@ jobs: backend_glob: "*ggml-blas*.dll" - arch: arm64 vcpkg_triplet: arm64-windows - runs_on: windows-11-arm + runs_on: windows-11-vs2026-arm backend: vulkan include_core: true needs_cuda: false @@ -706,7 +706,7 @@ jobs: backend_glob: "*ggml-vulkan*.dll" - arch: arm64 vcpkg_triplet: arm64-windows - runs_on: windows-11-arm + runs_on: windows-11-vs2026-arm backend: blas include_core: false needs_cuda: false diff --git a/.github/workflows/validate_wrapper.yml b/.github/workflows/validate_wrapper.yml index f63edae..8d7eb1c 100644 --- a/.github/workflows/validate_wrapper.yml +++ b/.github/workflows/validate_wrapper.yml @@ -134,7 +134,7 @@ jobs: windows-arm64-kleidiai: needs: changes if: needs.changes.outputs.native == 'true' - runs-on: windows-11-arm + runs-on: windows-11-vs2026-arm timeout-minutes: 30 strategy: fail-fast: false diff --git a/CMakePresets.json b/CMakePresets.json index 2ac4822..9fb51bd 100644 --- a/CMakePresets.json +++ b/CMakePresets.json @@ -175,7 +175,7 @@ "name": "windows-arm64-full", "inherits": "windows-base", "binaryDir": "${sourceDir}/build/wa64", - "generator": "Visual Studio 17 2022", + "generator": "Visual Studio 18 2026", "architecture": { "value": "ARM64", "strategy": "set" diff --git a/docs/platform_backend_strategy.md b/docs/platform_backend_strategy.md index 4047e42..bdde67f 100644 --- a/docs/platform_backend_strategy.md +++ b/docs/platform_backend_strategy.md @@ -69,6 +69,12 @@ The Android artifact check is `tools/validate_android_cpu_isa.py --help`. - `third_party/opencl-stubs/` - auto-built OpenCL ICD loader from `third_party/OpenCL-ICD-Loader` + `third_party/OpenCL-Headers` - Linux arm64 builds on x64 runners require `aarch64-linux-gnu-gcc/g++`, `libopenblas-dev:arm64`, and `libvulkan-dev:arm64`. +- Windows ARM64 uses the `Visual Studio 18 2026` generator (CMake >= 4.2) with + the ClangCL toolset. CI pins the `windows-11-vs2026-arm` runner because the + `windows-11-arm` label moved to that image in September 2026 + (actions/runner-images#14602), and that image has no VS 2022 instance. With + only VS 2022 installed, pass `-G "Visual Studio 17 2022"` after + `--preset windows-arm64-full`. - Windows MSVC builds disable IPO/LTCG for `llama-common` and `mtmd` by default. Current MSVC `link.exe` can access-violate when linking the large `llama-common` utility DLL with `/LTCG`. CMake's automatic Windows export From ddc4f847ddde074e4ad43030b5a8e8c27b4a203c Mon Sep 17 00:00:00 2001 From: Jhin Lee Date: Thu, 24 Sep 2026 08:10:54 -0400 Subject: [PATCH 2/4] fix: build Android Vulkan shaders with a coopmat-capable host glslc llama.cpp v0.5.0 (4ceb17191, #24406) compiles the new fa_decode_ph1/ph2 shaders as coopmat variants unconditionally, and ggml-vulkan.cpp references them without the GGML_VULKAN_COOPMAT_GLSLC_SUPPORT guard every other _cm1 shader has. NDK glslc (shaderc v2022.3, through r29) lacks GL_KHR_cooperative_matrix, so both Android Vulkan lanes of native release run 35976011426 failed in vulkan-shaders-gen. Reported upstream as ggml-org/llama.cpp#29373; no upstream source is patched here. tools/build.py now accepts ANDROID_VULKAN_GLSLC to replace the NDK glslc, and the release Android Vulkan lanes set it to the Ubuntu glslc package the Linux Vulkan lanes already use. That glslc adds exactly one detected shader feature over the NDK one (KHR coopmat), so Android now ships the same shader feature set as Linux. validate_wrapper gains android-vulkan-shaders, which builds the release arm64 ggml-vulkan target with that glslc on the pinned upstream and on v0.5.0. --- .github/workflows/native_release.yml | 10 ++++- .github/workflows/validate_wrapper.yml | 58 +++++++++++++++++++++++++- docs/platform_backend_strategy.md | 8 ++++ tests/test_android_vulkan_glslc.py | 41 ++++++++++++++++++ tests/test_ci_scope.py | 2 +- tools/build.py | 14 ++++++- tools/ci_scope.py | 2 +- 7 files changed, 130 insertions(+), 5 deletions(-) create mode 100644 tests/test_android_vulkan_glslc.py diff --git a/.github/workflows/native_release.yml b/.github/workflows/native_release.yml index b76e433..50f7dbb 100644 --- a/.github/workflows/native_release.yml +++ b/.github/workflows/native_release.yml @@ -182,7 +182,15 @@ jobs: echo "ANDROID_NDK_HOME=$NDK_DIR" >> "$GITHUB_ENV" echo "Using pinned Android NDK: $NDK_DIR" - name: Install build deps - run: sudo apt-get update && sudo apt-get install -y ninja-build ccache + run: | + sudo apt-get update && sudo apt-get install -y ninja-build ccache + if [ "${{ matrix.backend }}" = "vulkan" ]; then + # Host glslc, same package as the Linux Vulkan lanes; the NDK's lacks + # GL_KHR_cooperative_matrix (ggml-org/llama.cpp#29373). + sudo apt-get install -y glslc + glslc --version + echo "ANDROID_VULKAN_GLSLC=$(command -v glslc)" >> "$GITHUB_ENV" + fi - name: Set up ccache uses: ./.github/actions/setup-ccache with: diff --git a/.github/workflows/validate_wrapper.yml b/.github/workflows/validate_wrapper.yml index f63edae..5377f04 100644 --- a/.github/workflows/validate_wrapper.yml +++ b/.github/workflows/validate_wrapper.yml @@ -16,6 +16,7 @@ env: # Qualification only: never moves the production submodule or release channel. LLAMADART_QUALIFICATION_SHA: 73ab7599b553c03f6f5d2db24a18ad76f2eb36a3 LLAMADART_V041_QUALIFICATION_SHA: b29c606e28a01b1bc8c1351026a0fa6e616bf6c4 + LLAMADART_V050_QUALIFICATION_SHA: 7fe450e19305b828c199d602c23a8337aaa1f03b jobs: changes: @@ -84,6 +85,61 @@ jobs: --kleidiai-source build/android-isa/_deps/kleidiai-src \ build/android-isa/bin/libggml-cpu.so + android-vulkan-shaders: + needs: changes + if: needs.changes.outputs.native == 'true' + runs-on: ubuntu-latest + timeout-minutes: 45 + strategy: + fail-fast: false + matrix: + upstream: [pinned, v0.5.0] + steps: + - uses: actions/checkout@v7 + with: + persist-credentials: false + submodules: recursive + - name: Select candidate upstream + if: matrix.upstream != 'pinned' + run: | + git -C third_party/llama.cpp fetch --depth=1 origin "$LLAMADART_V050_QUALIFICATION_SHA" + git -C third_party/llama.cpp checkout --detach "$LLAMADART_V050_QUALIFICATION_SHA" + test "$(git -C third_party/llama.cpp rev-parse HEAD)" = "$LLAMADART_V050_QUALIFICATION_SHA" + - uses: actions/setup-java@v5 + with: + distribution: temurin + java-version: '17' + - uses: android-actions/setup-android@v4 + with: + packages: platform-tools ndk;28.2.13676358 + - name: Install build dependencies and the release glslc + run: | + sudo apt-get update && sudo apt-get install -y ninja-build glslc + glslc --version + - name: Build the release Android Vulkan backend + run: | + set -euo pipefail + export ANDROID_NDK_HOME="${ANDROID_SDK_ROOT}/ndk/28.2.13676358" + ANDROID_VULKAN_GLSLC="$(command -v glslc)" + export ANDROID_VULKAN_GLSLC + python3 - <<'PY' 2>&1 | tee configure-build.log + import os, sys + from pathlib import Path + sys.path.insert(0, 'tools') + import build + abi = 'arm64-v8a' + build_dir = Path('build/android-vulkan').resolve() + env = dict(os.environ) + args = build.android_configure_args( + abi, build_dir=build_dir, ndk=Path(env['ANDROID_NDK_HOME']), env=env, + cache_vars=build.android_backend_cache_vars(abi, 'vulkan'), jobs=4) + build.configure_and_build_dir( + build_dir, jobs=4, cmake_args=args, build_target='ggml-vulkan', env=env) + PY + grep -F -- "-DVulkan_GLSLC_EXECUTABLE=$ANDROID_VULKAN_GLSLC" configure-build.log + grep -Fx -- '-- GL_KHR_cooperative_matrix supported by glslc' configure-build.log + test -f build/android-vulkan/bin/libggml-vulkan.so + kleidiai-dispatch-emulated: needs: changes if: needs.changes.outputs.native == 'true' @@ -324,7 +380,7 @@ jobs: wrapper-validation: if: always() - needs: [changes, android-arm64-isa, kleidiai-dispatch-emulated, windows-arm64-kleidiai, linux-artifact-contract, wrapper-contract, msvc-mtmd-link-contract] + needs: [changes, android-arm64-isa, android-vulkan-shaders, kleidiai-dispatch-emulated, windows-arm64-kleidiai, linux-artifact-contract, wrapper-contract, msvc-mtmd-link-contract] runs-on: ubuntu-latest steps: - uses: actions/checkout@v7 diff --git a/docs/platform_backend_strategy.md b/docs/platform_backend_strategy.md index 4047e42..e2912d1 100644 --- a/docs/platform_backend_strategy.md +++ b/docs/platform_backend_strategy.md @@ -63,6 +63,14 @@ The Android artifact check is `tools/validate_android_cpu_isa.py --help`. - CUDA lanes require `nvcc` availability. - HIP/ROCm lanes require `hipcc`, `rocblas-dev`, and `hipblas-dev` (Linux x64 only, built in a separate release job). - Android Vulkan lanes require NDK-provided `libvulkan.so`. +- Android Vulkan shaders need a glslc with `GL_KHR_cooperative_matrix` from + llama.cpp v0.5.0 on: upstream compiles the `fa_decode` shaders as coopmat + unconditionally (ggml-org/llama.cpp#29373), and NDK glslc (shaderc v2022.3 + through r29) lacks the extension. Set `ANDROID_VULKAN_GLSLC` to a host glslc; + release CI uses the Ubuntu `glslc` package, like the Linux Vulkan lanes. + Upstream enables shader features from glslc, so this also compiles the KHR + coopmat `mul_mm`/flash-attention variants, used at runtime only on devices + reporting `VK_KHR_cooperative_matrix` (`GGML_VK_DISABLE_COOPMAT=1` opts out). - Vulkan lanes use vendored `third_party/SPIRV-Headers` for SPIR-V registry headers required by upstream `llama.cpp`. - Android OpenCL lanes require `CL/cl.h` and `libOpenCL.so` from one of: - env overrides (`OPENCL_INCLUDE_DIR`, `OPENCL_LIBRARY_ANDROID_`) diff --git a/tests/test_android_vulkan_glslc.py b/tests/test_android_vulkan_glslc.py new file mode 100644 index 0000000..3a6346b --- /dev/null +++ b/tests/test_android_vulkan_glslc.py @@ -0,0 +1,41 @@ +from __future__ import annotations + +from pathlib import Path +import sys +import tempfile +import unittest + +ROOT = Path(__file__).resolve().parents[1] +sys.path.insert(0, str(ROOT / "tools")) + +import build # noqa: E402 + + +class AndroidVulkanGlslcTest(unittest.TestCase): + def test_override_replaces_ndk_glslc_and_missing_override_fails(self) -> None: + with tempfile.TemporaryDirectory() as temp: + ndk = Path(temp) / "ndk" + ndk_glslc = ndk / "shader-tools/linux-x86_64/glslc" + ndk_glslc.parent.mkdir(parents=True) + ndk_glslc.write_text("") + host_glslc = Path(temp) / "glslc" + host_glslc.write_text("") + + self.assertEqual(ndk_glslc, build.android_vulkan_glslc(ndk, {})) + self.assertEqual( + host_glslc, + build.android_vulkan_glslc(ndk, {"ANDROID_VULKAN_GLSLC": str(host_glslc)}), + ) + with self.assertRaises(SystemExit): + build.android_vulkan_glslc(ndk, {"ANDROID_VULKAN_GLSLC": str(Path(temp) / "missing")}) + + def test_release_android_vulkan_lanes_use_host_glslc(self) -> None: + workflow = (ROOT / ".github/workflows/native_release.yml").read_text(encoding="utf-8") + android_job = workflow.split("\n build-android:\n", 1)[1].split("\n build-apple:\n", 1)[0] + self.assertIn('if [ "${{ matrix.backend }}" = "vulkan" ]; then', android_job) + self.assertIn("sudo apt-get install -y glslc", android_job) + self.assertIn('echo "ANDROID_VULKAN_GLSLC=$(command -v glslc)" >> "$GITHUB_ENV"', android_job) + + +if __name__ == "__main__": + unittest.main() diff --git a/tests/test_ci_scope.py b/tests/test_ci_scope.py index 15521bc..b5dff71 100644 --- a/tests/test_ci_scope.py +++ b/tests/test_ci_scope.py @@ -12,7 +12,7 @@ sys.path.insert(0, str(ROOT / 'scripts')) from verify_release_provenance import workflow_job -JOBS = ('android-arm64-isa', 'kleidiai-dispatch-emulated', +JOBS = ('android-arm64-isa', 'android-vulkan-shaders', 'kleidiai-dispatch-emulated', 'windows-arm64-kleidiai', 'linux-artifact-contract', 'wrapper-contract', 'msvc-mtmd-link-contract') diff --git a/tools/build.py b/tools/build.py index a6c996f..26433f4 100755 --- a/tools/build.py +++ b/tools/build.py @@ -313,6 +313,18 @@ def android_base_cmake_args(abi: str, ndk: Path) -> list[str]: ] +def android_vulkan_glslc(ndk: Path, env: dict[str, str]) -> Path | None: + # The NDK's glslc (shaderc v2022.3 through r29) lacks GL_KHR_cooperative_matrix, + # which llama.cpp v0.5.0 needs unconditionally (ggml-org/llama.cpp#29373). + override = env.get("ANDROID_VULKAN_GLSLC") + if override: + glslc = Path(override) + if not glslc.is_file(): + fail(f"ANDROID_VULKAN_GLSLC does not name a file: {glslc}") + return glslc + return find_file_with_suffix(ndk, "glslc") or find_file_with_suffix(ndk, "glslc.exe") + + def android_configure_args( abi: str, *, @@ -331,7 +343,7 @@ def android_configure_args( write_android_host_toolchain(toolchain) cmake_args.append(f"-DGGML_VULKAN_SHADERS_GEN_TOOLCHAIN={toolchain}") - glslc = find_file_with_suffix(ndk, "glslc") or find_file_with_suffix(ndk, "glslc.exe") + glslc = android_vulkan_glslc(ndk, env) if glslc: cmake_args.append(f"-DVulkan_GLSLC_EXECUTABLE={glslc}") diff --git a/tools/ci_scope.py b/tools/ci_scope.py index 70b35c2..4ec9c64 100644 --- a/tools/ci_scope.py +++ b/tools/ci_scope.py @@ -48,7 +48,7 @@ def validate_results(needs): if selected not in ('true', 'false'): return False expected = 'success' if selected == 'true' else 'skipped' - jobs = {'android-arm64-isa', 'kleidiai-dispatch-emulated', + jobs = {'android-arm64-isa', 'android-vulkan-shaders', 'kleidiai-dispatch-emulated', 'windows-arm64-kleidiai', 'linux-artifact-contract', 'wrapper-contract', 'msvc-mtmd-link-contract'} return set(needs) == jobs | {'changes'} and all( From 4a88e863340ecedc122513987fdc891a01911d29 Mon Sep 17 00:00:00 2001 From: Jhin Lee Date: Thu, 24 Sep 2026 08:19:45 -0400 Subject: [PATCH 3/4] docs: scope the VS 2022 fallback to direct cmake configures --- docs/platform_backend_strategy.md | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/docs/platform_backend_strategy.md b/docs/platform_backend_strategy.md index bdde67f..d1a0158 100644 --- a/docs/platform_backend_strategy.md +++ b/docs/platform_backend_strategy.md @@ -72,9 +72,11 @@ The Android artifact check is `tools/validate_android_cpu_isa.py --help`. - Windows ARM64 uses the `Visual Studio 18 2026` generator (CMake >= 4.2) with the ClangCL toolset. CI pins the `windows-11-vs2026-arm` runner because the `windows-11-arm` label moved to that image in September 2026 - (actions/runner-images#14602), and that image has no VS 2022 instance. With - only VS 2022 installed, pass `-G "Visual Studio 17 2022"` after - `--preset windows-arm64-full`. + (actions/runner-images#14602), and that image has no VS 2022 instance. + `tools/build.py windows --arch arm64` therefore needs VS 2026. With only + VS 2022 installed, configure directly with + `cmake --preset windows-arm64-full -G "Visual Studio 17 2022"` in a fresh + build directory. - Windows MSVC builds disable IPO/LTCG for `llama-common` and `mtmd` by default. Current MSVC `link.exe` can access-violate when linking the large `llama-common` utility DLL with `/LTCG`. CMake's automatic Windows export From 95e85f4d361e1f22caacab818e7de671d04f889a Mon Sep 17 00:00:00 2001 From: Jhin Lee Date: Thu, 24 Sep 2026 08:20:41 -0400 Subject: [PATCH 4/4] ci: pin the Android runner so the host glslc cannot drift --- .github/workflows/native_release.yml | 4 +++- .github/workflows/validate_wrapper.yml | 2 +- docs/platform_backend_strategy.md | 5 ++++- tools/build.py | 1 + 4 files changed, 9 insertions(+), 3 deletions(-) diff --git a/.github/workflows/native_release.yml b/.github/workflows/native_release.yml index 99c86ec..c4507da 100644 --- a/.github/workflows/native_release.yml +++ b/.github/workflows/native_release.yml @@ -124,7 +124,9 @@ jobs: build-android: needs: resolve-tag - runs-on: ubuntu-latest + # Pinned: the Vulkan lanes' glslc comes from this image's apt archive, and a + # newer shaderc would silently enable more Android shader features. + runs-on: ubuntu-24.04 env: CCACHE_DIR: ${{ github.workspace }}/.ccache CCACHE_MAXSIZE: 2G diff --git a/.github/workflows/validate_wrapper.yml b/.github/workflows/validate_wrapper.yml index d4a9f6e..3690ced 100644 --- a/.github/workflows/validate_wrapper.yml +++ b/.github/workflows/validate_wrapper.yml @@ -88,7 +88,7 @@ jobs: android-vulkan-shaders: needs: changes if: needs.changes.outputs.native == 'true' - runs-on: ubuntu-latest + runs-on: ubuntu-24.04 timeout-minutes: 45 strategy: fail-fast: false diff --git a/docs/platform_backend_strategy.md b/docs/platform_backend_strategy.md index 3faf5b4..7ee8eb2 100644 --- a/docs/platform_backend_strategy.md +++ b/docs/platform_backend_strategy.md @@ -67,7 +67,10 @@ The Android artifact check is `tools/validate_android_cpu_isa.py --help`. llama.cpp v0.5.0 on: upstream compiles the `fa_decode` shaders as coopmat unconditionally (ggml-org/llama.cpp#29373), and NDK glslc (shaderc v2022.3 through r29) lacks the extension. Set `ANDROID_VULKAN_GLSLC` to a host glslc; - release CI uses the Ubuntu `glslc` package, like the Linux Vulkan lanes. + release CI uses the Ubuntu 24.04 `glslc` package (shaderc 2023.8), like the + Linux Vulkan lanes, on a pinned `ubuntu-24.04` runner so the shader feature set + cannot drift with `ubuntu-latest`. This applies to every release built with + the override, including rebuilds of the current pin. Upstream enables shader features from glslc, so this also compiles the KHR coopmat `mul_mm`/flash-attention variants, used at runtime only on devices reporting `VK_KHR_cooperative_matrix` (`GGML_VK_DISABLE_COOPMAT=1` opts out). diff --git a/tools/build.py b/tools/build.py index 26433f4..2a50707 100755 --- a/tools/build.py +++ b/tools/build.py @@ -316,6 +316,7 @@ def android_base_cmake_args(abi: str, ndk: Path) -> list[str]: def android_vulkan_glslc(ndk: Path, env: dict[str, str]) -> Path | None: # The NDK's glslc (shaderc v2022.3 through r29) lacks GL_KHR_cooperative_matrix, # which llama.cpp v0.5.0 needs unconditionally (ggml-org/llama.cpp#29373). + # Drop the release override once a pinned upstream gates those shaders. override = env.get("ANDROID_VULKAN_GLSLC") if override: glslc = Path(override)