fix: build Android Vulkan shaders with a coopmat-capable host glslc

.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:

.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

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_<ABI>`)

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()

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')
 

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}")
 

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(