diff --git a/.github/workflows/ci.yaml b/.github/workflows/ci.yaml index d721d97c..4c67e8e4 100644 --- a/.github/workflows/ci.yaml +++ b/.github/workflows/ci.yaml @@ -90,5 +90,29 @@ jobs: - name: Test with pytest run: | - pytest --timeout=1500 -W ignore::PendingDeprecationWarning tests + pytest -s --timeout=1500 -W ignore::PendingDeprecationWarning tests working-directory: . + + - name: Dump macOS crash reports (if pytest crashed) + if: always() && startsWith(matrix.os, 'macos') + run: | + # give the async ReportCrash daemon a moment to finish writing + sleep 10 + FOUND=0 + for REPORT_DIR in "$HOME/Library/Logs/DiagnosticReports" "/Library/Logs/DiagnosticReports"; do + echo "Looking for crash reports in $REPORT_DIR" + if [ -d "$REPORT_DIR" ]; then + # -mmin, not GNU find's -newermt, for BSD find compatibility on macOS + while IFS= read -r f; do + FOUND=1 + echo "===== $f =====" + cat "$f" + echo "===== end $f =====" + done < <(find "$REPORT_DIR" \( -iname "*python*" -o -iname "*pytest*" \) \( -iname "*.ips" -o -iname "*.crash" \) -mmin -30 -print 2>/dev/null) + fi + done + if [ "$FOUND" -eq 0 ]; then + echo "no matching crash reports found" + fi + working-directory: . + diff --git a/patches/llama.cpp/0002-metal-pin-msl-language-version.patch b/patches/llama.cpp/0002-metal-pin-msl-language-version.patch new file mode 100644 index 00000000..46dbdb3e --- /dev/null +++ b/patches/llama.cpp/0002-metal-pin-msl-language-version.patch @@ -0,0 +1,170 @@ +diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m +index 7d2a68685..0abd203db 100644 +--- a/ggml/src/ggml-metal/ggml-metal-device.m ++++ b/ggml/src/ggml-metal/ggml-metal-device.m +@@ -103,6 +103,76 @@ struct ggml_metal_library { + NSLock * lock; + }; + ++// Minimum Metal Shading Language (MSL) version required by the features enabled ++// on this device, as a raw MTLLanguageVersion value (major << 16 | minor). ++// ++// When MTLCompileOptions.languageVersion is left unset, Metal derives the default ++// from the SDK the *host executable* was linked against - not from the running OS. ++// Hosts built against old SDKs (conda, python.org, ...) then fail in two ways: ++// ++// - bfloat: the __METAL_VERSION__ < 310 guard in ggml-metal.metal drops every ++// bf16 kernel while the device still reports has_bfloat, so the first BF16 op ++// fails with "Function kernel_..._bf16_... was not found in the library" ++// - tensor API: / MetalPerformancePrimitives need MSL 4.0, so ++// the capability probes below fail even on M5+ hardware that supports them ++// ++// Requesting the version the enabled features need fixes both. ++// ++// note: raw values are used because the MTLLanguageVersion3_1 / 4_0 enum ++// constants require a recent macOS SDK to compile against ++// ++// ref: https://github.com/ggml-org/llama.cpp/issues/21381 ++static NSInteger ggml_metal_device_msl_version_min(ggml_metal_device_t dev) { ++ const struct ggml_metal_device_props * props = ggml_metal_device_get_props(dev); ++ ++ if (props->has_tensor) { ++ // + MetalPerformancePrimitives (implies bfloat support too) ++ return (4 << 16 | 0); ++ } ++ ++ if (props->has_bfloat) { ++ // the bf16 kernels need MSL 3.1+ ++ return (3 << 16 | 1); ++ } ++ ++ return 0; ++} ++ ++// compile MSL source into a library, requesting the language version the ++// device's enabled features need (see ggml_metal_device_msl_version_min) ++// ++// note: an MSL version the linked Metal framework does not recognize can raise ++// an uncaught NSException instead of populating `error`, which would ++// abort() the host process - fatal for an embedded library (e.g. a python ++// extension) - so it is caught and treated as an ordinary compile failure ++static id ggml_metal_compile_source(ggml_metal_device_t dev, NSString * src, NSDictionary * prep, NSError ** error) { ++ id device = ggml_metal_device_get_obj(dev); ++ id library = nil; ++ ++ MTLCompileOptions * options = [MTLCompileOptions new]; ++ options.preprocessorMacros = prep; ++ ++ const NSInteger ver_min = ggml_metal_device_msl_version_min(dev); ++ if (ver_min != 0) { ++ options.languageVersion = (MTLLanguageVersion) ver_min; ++ } ++ ++ //[options setFastMathEnabled:false]; ++ ++ @try { ++ library = [device newLibraryWithSource:src options:options error:error]; ++ } @catch (NSException * exception) { ++ GGML_LOG_ERROR("%s: exception while compiling metal source: %s\n", __func__, [[exception reason] UTF8String]); ++ library = nil; ++ } ++ ++#if !__has_feature(objc_arc) ++ [options release]; ++#endif ++ ++ return library; ++} ++ + ggml_metal_library_t ggml_metal_library_init(ggml_metal_device_t dev) { + id library = nil; + id device = ggml_metal_device_get_obj(dev); +@@ -225,20 +295,16 @@ ggml_metal_library_t ggml_metal_library_init(ggml_metal_device_t dev) { + [prep setObject:@"1" forKey:@"GGML_METAL_EMBED_LIBRARY"]; + #endif + +- MTLCompileOptions * options = [MTLCompileOptions new]; +- options.preprocessorMacros = prep; +- +- //[options setFastMathEnabled:false]; +- +- library = [device newLibraryWithSource:src options:options error:&error]; ++ library = ggml_metal_compile_source(dev, src, prep, &error); + if (error) { + GGML_LOG_ERROR("%s: error: %s\n", __func__, [[error description] UTF8String]); + return nil; + } + +-#if !__has_feature(objc_arc) +- [options release]; +-#endif ++ if (!library) { ++ GGML_LOG_ERROR("%s: error: failed to compile the metal library\n", __func__); ++ return nil; ++ } + } + } + +@@ -265,7 +331,6 @@ ggml_metal_library_t ggml_metal_library_init_from_source(ggml_metal_device_t dev + return NULL; + } + +- id device = ggml_metal_device_get_obj(dev); + id library = nil; + NSError * error = nil; + +@@ -282,10 +347,7 @@ ggml_metal_library_t ggml_metal_library_init_from_source(ggml_metal_device_t dev + @autoreleasepool { + NSMutableDictionary * prep = [NSMutableDictionary dictionary]; + +- MTLCompileOptions * options = [MTLCompileOptions new]; +- options.preprocessorMacros = prep; +- +- library = [device newLibraryWithSource:src options:options error:&error]; ++ library = ggml_metal_compile_source(dev, src, prep, &error); + if (error) { + if (verbose) { + GGML_LOG_ERROR("%s: error compiling source: %s\n", __func__, [[error description] UTF8String]); +@@ -294,8 +356,6 @@ ggml_metal_library_t ggml_metal_library_init_from_source(ggml_metal_device_t dev + } + library = nil; + } +- +- [options release]; + } + + [src release]; +@@ -858,6 +918,33 @@ ggml_metal_device_t ggml_metal_device_init(int device) { + } + } + ++ // double-check that bfloat support compiles: the bf16 kernels require ++ // MSL 3.1+ (see ggml_metal_device_msl_version_min) - if the environment ++ // cannot provide it, disable bfloat so BF16 ops fall back to the CPU ++ // backend instead of failing later with "Function kernel_..._bf16_... was ++ // not found in the library" ++ if (dev->props.has_bfloat) { ++ const char * src_bf16 = "\n" ++ "#include \n" ++ "using namespace metal; \n" ++ "kernel void dummy_kernel(device bfloat * dst [[buffer(0)]], device const bfloat * src [[buffer(1)]], uint i [[thread_position_in_grid]]) { dst[i] = src[i]; } \n"; ++ ++ GGML_LOG_INFO("%s: testing bfloat support\n", __func__); ++ ggml_metal_library_t lib = ggml_metal_library_init_from_source(dev, src_bf16, false); ++ if (lib == NULL) { ++ GGML_LOG_WARN("%s: - bfloat is not supported in this environment - disabling bfloat support\n", __func__); ++ dev->props.has_bfloat = false; ++ } else { ++ struct ggml_metal_pipeline_with_params ppl = ggml_metal_library_compile_pipeline(lib, "dummy_kernel", "dummy_kernel", nil); ++ if (!ppl.pipeline) { ++ GGML_LOG_WARN("%s: - bfloat is not supported in this environment - disabling bfloat support\n", __func__); ++ dev->props.has_bfloat = false; ++ } ++ ++ ggml_metal_library_free(lib); ++ } ++ } ++ + dev->props.use_residency_sets = true; + #if defined(GGML_METAL_HAS_RESIDENCY_SETS) + dev->props.use_residency_sets = getenv("GGML_METAL_NO_RESIDENCY") == nil; diff --git a/tests/test_metal_bf16.py b/tests/test_metal_bf16.py new file mode 100644 index 00000000..0185b233 --- /dev/null +++ b/tests/test_metal_bf16.py @@ -0,0 +1,264 @@ +"""Regression tests for Metal bf16 kernel support on Apple Silicon (M1+). + +Background: the wheel embeds the Metal shader source and JIT-compiles it at +runtime. llama.cpp injects GGML_METAL_HAS_BF16 when the GPU supports bfloat +(has_bfloat is true starting with MTLGPUFamilyApple6, i.e. A14/M1, not just +M3+ as earlier drafts of this file assumed), but never pins +MTLCompileOptions.languageVersion, so the Metal compiler derives the default +shading language version from the *host executable's* LC_BUILD_VERSION (the +python interpreter). Interpreters linked against old SDKs (conda, +python.org) default to MSL 2.x, in which case the `__METAL_VERSION__ < 310` +guard in ggml-metal.metal silently strips all bf16 kernels and loading a +model with BF16 tensors fails with: + + Function kernel_mul_mv_ext_bf16_f32_r1_2 was not found in the library + +The fix (patches/llama.cpp/0001-metal-pin-msl-language-version.patch, applied +at build time by scripts/build.py) pins the MSL version to >= 3.1 when the +device reports bfloat support, and probes bf16 compilation at device init: if +the environment cannot compile MSL 3.1, bfloat is disabled (with a warning) so +BF16 ops fall back to CPU instead of failing model loads. + +These tests build a tiny llama GGUF whose weight matrices are BF16 (norm/1D +tensors stay F32, matching how real BF16 GGUF conversions work - see +conversion/base.py in the llama.cpp submodule) and load it: +1. in-process, and +2. under a python host whose LC_BUILD_VERSION is rewritten to an old SDK, + emulating conda/python.org interpreters (the failing scenario). + +On machines without bfloat support llama.cpp falls back to CPU for BF16 ops, +so the tests still pass but do not exercise the Metal bf16 kernels. Where +_EXPECT_BF16 is true (SoC generation >= 3, i.e. M3+ - kept conservative here +even though has_bfloat itself is true starting at M1) the tests additionally +verify that the bf16 kernels were actually built: a silent fallback (bfloat +disabled, model running on CPU) must fail the test, not pass it. +""" + +import os +import platform +import re +import shutil +import subprocess +import sys +import sysconfig +from pathlib import Path + +import pytest + +ROOT = Path(__file__).parent.parent + +pytestmark = pytest.mark.skipif( + platform.system() != "Darwin" or platform.machine() != "arm64", + reason="Metal bf16 kernels are only exercised on Apple Silicon", +) + + +def _apple_silicon_gen() -> int: + """SoC generation (3 for M3, 4 for M4, ...); 0 if unknown.""" + try: + brand = subprocess.run( + ["sysctl", "-n", "machdep.cpu.brand_string"], + capture_output=True, + text=True, + check=True, + ).stdout.strip() + except (subprocess.CalledProcessError, FileNotFoundError): + return 0 + m = re.search(r"Apple M(\d+)", brand) + return int(m.group(1)) if m else 0 + + +# bfloat (Metal3 GPU family) requires M3 or newer +_EXPECT_BF16 = _apple_silicon_gen() >= 3 + + +def _f32_to_bf16(arr): + """Round-to-nearest-even float32 -> bfloat16 bits as uint16.""" + import numpy as np + + bits = arr.astype(np.float32).view(np.uint32) + return ((bits + 0x7FFF + ((bits >> 16) & 1)) >> 16).astype(np.uint16) + + +def _write_bf16_variant(src: Path, dst: Path) -> None: + """Copy a llama GGUF with weight matrices rewritten as BF16 random weights. + + Mirrors llama.cpp's own conversion invariant (conversion/base.py, + ModelBase.write_tensors): 1D tensors and *_norm.weight tensors are always + kept as F32 ("most of the codebase that takes in 1D tensors or norms only + handles F32 tensors") - no real BF16 GGUF file has BF16 norm weights. + Forcing them to BF16 here would create a tensor-type combination + (F32 activation x BF16 norm weight in ggml_mul) that no real model ever + produces and that isn't universally supported by every ggml backend. + """ + gguf = pytest.importorskip("gguf") + import numpy as np + + reader = gguf.GGUFReader(str(src)) + writer = gguf.GGUFWriter(str(dst), arch="llama") + for field in reader.fields.values(): + # GGUF.* and general.architecture are managed by the writer; file_type + # describes the original quantization and no longer applies. + if field.name.startswith("GGUF.") or field.name in ( + "general.architecture", + "general.file_type", + "general.quantization_version", + ): + continue + vtype = field.types[0] + if vtype == gguf.GGUFValueType.ARRAY: + writer.add_array(field.name, field.contents()) + elif vtype == gguf.GGUFValueType.STRING: + writer.add_string(field.name, field.contents()) + else: + writer.add_key_value(field.name, field.contents(), vtype) + + rng = np.random.default_rng(0) + for tensor in reader.tensors: + # reader shapes are in ggml order (ne0 first); the writer expects + # numpy order and reverses them when writing the file. + shape = tuple(int(d) for d in tensor.shape[::-1]) + data = rng.normal(0.0, 0.02, size=shape).astype(np.float32) + if len(shape) <= 1 or tensor.name.endswith("_norm.weight"): + writer.add_tensor(tensor.name, data, raw_dtype=gguf.GGMLQuantizationType.F32) + else: + writer.add_tensor( + tensor.name, + _f32_to_bf16(data).reshape(shape), + raw_dtype=gguf.GGMLQuantizationType.BF16, + ) + + writer.write_header_to_file() + writer.write_kv_data_to_file() + writer.write_tensors_to_file(progress=False) + writer.close() + + +@pytest.fixture(scope="session") +def bf16_model_path(tmp_path_factory): + src = ROOT / "models" / "stories15M-q4_0.gguf" + if not src.exists(): + pytest.skip(f"reference model not found: {src} (run `make download`)") + dst = tmp_path_factory.mktemp("bf16-model") / "tiny-llama-bf16.gguf" + _write_bf16_variant(src, dst) + return str(dst) + + +def _load_model(model_path: str) -> None: + import sys + + import xllamacpp as xlc + + params = xlc.CommonParams() + params.model.path = model_path + params.n_ctx = 512 + params.n_gpu_layers = 99 + # GGML_LOG_INFO lines (e.g. ggml_metal_device_init's "has bfloat"/"has + # tensor"/GPU family diagnostics) are mapped by common_get_verbosity() to + # LOG_LEVEL_TRACE (4), while the default verbosity threshold is + # LOG_LEVEL_INFO (3) - i.e. they are silently dropped by default, + # regardless of any output capturing. Raise it so they are actually + # emitted; this matters most for diagnosing CI-only failures. + params.verbosity = 999 + print("[test_metal_bf16] about to construct xlc.Server(...)", flush=True) + sys.stderr.flush() + # default warmup=True: the warmup run is what compiles the bf16 pipelines + server = xlc.Server(params) + print("[test_metal_bf16] xlc.Server(...) constructed successfully", flush=True) + del server + + +def test_metal_bf16_model_loads(bf16_model_path, capfd): + _load_model(bf16_model_path) + if _EXPECT_BF16: + # a silent fallback (bfloat disabled -> BF16 ops on CPU) would still + # load successfully - make sure that did not happen + assert "disabling bfloat support" not in capfd.readouterr().err + + +_PYHOST_C = "#include \nint main(int argc, char **argv) { return Py_BytesMain(argc, argv); }\n" + + +@pytest.fixture(scope="session") +def old_sdk_python(tmp_path_factory): + """A python executable with an old LC_BUILD_VERSION (like conda/python.org). + + Metal derives the default shading language version from the main + executable's LC_BUILD_VERSION; rewriting it to SDK 11.0 reproduces the + environment in which the bf16 kernels were silently dropped pre-fix. + """ + for tool in ("clang", "vtool", "codesign"): + if not shutil.which(tool): + pytest.skip(f"{tool} is required to build the old-SDK python host") + + tmp = tmp_path_factory.mktemp("pyhost") + src = tmp / "pyhost.c" + src.write_text(_PYHOST_C) + + include = sysconfig.get_paths()["include"] + libdir = sysconfig.get_config_var("LIBDIR") + match = re.fullmatch( + r"lib(.+?)\.(?:dylib|so(?:\..*)?|a)", sysconfig.get_config_var("LDLIBRARY") or "" + ) + lib = match.group(1) if match else f"python{sys.version_info.major}.{sys.version_info.minor}" + + modern = tmp / "pyhost_modern" + subprocess.run( + [ + "clang", + "-mmacosx-version-min=11.0", + str(src), + "-o", + str(modern), + f"-I{include}", + f"-L{libdir}", + f"-l{lib}", + f"-Wl,-rpath,{libdir}", + ], + check=True, + ) + + old = tmp / "python_old" + subprocess.run( + ["vtool", "-set-build-version", "1", "11.0", "11.0", "-output", str(old), str(modern)], + check=True, + ) + subprocess.run( + ["codesign", "--force", "--sign", "-", str(old)], + check=True, + capture_output=True, + ) + subprocess.run([str(old), "--version"], check=True, capture_output=True) + return str(old) + + +def test_metal_bf16_model_loads_with_old_sdk_python(old_sdk_python, bf16_model_path): + code = ( + "import xllamacpp as xlc\n" + "p = xlc.CommonParams()\n" + f"p.model.path = {bf16_model_path!r}\n" + "p.n_ctx = 512\n" + "p.n_gpu_layers = 99\n" + "p.verbosity = 999\n" + "print('about to construct xlc.Server(...)')\n" + "xlc.Server(p)\n" + "print('MODEL_LOADED_OK')\n" + ) + env = dict(os.environ) + env["PYTHONPATH"] = str(ROOT / "src") + proc = subprocess.run( + [old_sdk_python, "-c", code], + capture_output=True, + text=True, + timeout=300, + env=env, + ) + assert proc.returncode == 0, f"Process failed with code {proc.returncode}. Stderr: {proc.stderr}" + assert "was not found in the library" not in proc.stderr + assert "MODEL_LOADED_OK" in proc.stdout + if _EXPECT_BF16: + # the bf16 kernels must have been built for real: bfloat must have + # stayed enabled - a silent disable would still load OK via CPU + # fallback (with bfloat enabled but no bf16 kernels in the library, + # the load above would have failed with "was not found") + assert "disabling bfloat support" not in proc.stderr