b9088

The llama.cpp project, maintained by ggml-org, released version b9088 with a critical fix for the SYCL backend. The update adds BF16 (bfloat16) support to the GET_ROWS operation, which previously only handled F16 and F32 types. This was causing a severe performance regression: models like Gemma4 that use BF16 for embedding tensors (e.g., per_layer_token_embd.weight) would fall back to CPU execution for the GET_ROWS op, forcing a full GPU-to-CPU tensor transfer on every token generation step. The fix patches both the supports_op check and the kernel dispatch to correctly handle BF16, reusing the existing get_rows_sycl_float template with sycl::ext::oneapi::bfloat16 – the same pattern already used for half-precision (F16) and single-precision (F32) types.

The broader impact is significant for LLM inference on Intel GPUs and other SYCL-compatible hardware. BF16 embeddings are becoming more common in modern models like Gemma4 due to their memory efficiency and training advantages. By eliminating the CPU fallback, this update unlocks native GPU throughput for these architectures without costly data transfers. The release also underscores llama.cpp's commitment to supporting diverse hardware backends (CUDA, Vulkan, ROCm, SYCL, etc.) and ensures that cutting-edge models run efficiently across platforms. For developers deploying Gemma4 or similar models on Intel GPUs via SYCL, this fix is essential for maintaining low latency and high throughput.