diff --git a/.github/workflows/native_release.yml b/.github/workflows/native_release.yml index 56f6f10..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 @@ -182,7 +184,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 8d7eb1c..3690ced 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-24.04 + 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 d1a0158..bb5f0bd 100644 --- a/docs/platform_backend_strategy.md +++ b/docs/platform_backend_strategy.md @@ -63,6 +63,17 @@ 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 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). - 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..2a50707 100755 --- a/tools/build.py +++ b/tools/build.py @@ -313,6 +313,19 @@ 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) + 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 +344,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(