Commit 5e5b628eb for llama.cpp
commit 5e5b628eb5515d50de5c7f98a01937fc25b844b9
Author: Pratyush Kumar <praty.bachchas@gmail.com>
Date: Wed Oct 7 12:53:28 2026 +0530
sycl : fattn_kv_buffers cleanup (#27689)
diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp
index 11fe86cc0..5e904bd2a 100644
--- a/ggml/src/ggml-sycl/common.hpp
+++ b/ggml/src/ggml-sycl/common.hpp
@@ -26,7 +26,6 @@
#include "presets.hpp"
#include "type.hpp"
#include "sycl_hw.hpp"
-#include "fattn-buffers.hpp"
#include "memtrace.hpp"
namespace syclexp = sycl::ext::oneapi::experimental;
@@ -408,8 +407,6 @@ struct ggml_backend_sycl_context {
// pool
std::unique_ptr<ggml_sycl_pool> pools[GGML_SYCL_MAX_DEVICES];
- std::unique_ptr<ggml_sycl_fattn_kv_buffers> fattn_bufs[GGML_SYCL_MAX_DEVICES];
-
std::unique_ptr<ggml_sycl_pool> host_pools[GGML_SYCL_MAX_DEVICES];
std::vector<mmid_row_mapping> mmid_row_mapping_host;
@@ -418,8 +415,6 @@ struct ggml_backend_sycl_context {
static std::unique_ptr<ggml_sycl_pool> new_pool_for_host(queue_ptr qptr, int device);
- static std::unique_ptr<ggml_sycl_fattn_kv_buffers> new_fattn_kv_buffers(queue_ptr qptr, int device);
-
ggml_sycl_pool & pool(int device) {
if (pools[device] == nullptr) {
pools[device] = new_pool_for_device(stream(device,0), device);
@@ -431,17 +426,6 @@ struct ggml_backend_sycl_context {
return pool(device);
}
- ggml_sycl_fattn_kv_buffers & fattn_buffers(int device) {
- if (fattn_bufs[device] == nullptr) {
- fattn_bufs[device] = new_fattn_kv_buffers(stream(device, 0), device);
- }
- return *fattn_bufs[device];
- }
-
- ggml_sycl_fattn_kv_buffers & fattn_buffers() {
- return fattn_buffers(device);
- }
-
#ifdef GGML_SYCL_GRAPH
std::unique_ptr<sycl_ex::command_graph<sycl_ex::graph_state::executable>> exec_graph = nullptr;
#endif
diff --git a/ggml/src/ggml-sycl/fattn-buffers.cpp b/ggml/src/ggml-sycl/fattn-buffers.cpp
deleted file mode 100644
index 78a52d2ab..000000000
--- a/ggml/src/ggml-sycl/fattn-buffers.cpp
+++ /dev/null
@@ -1,60 +0,0 @@
-//
-// MIT license
-// Copyright (C) 2025 Intel Corporation
-// SPDX-License-Identifier: MIT
-//
-
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-
-#include "common.hpp"
-
-sycl::half * ggml_sycl_fattn_kv_buffers::kv_buffer::ensure_half(size_t n_elems) {
- const size_t need_bytes = n_elems * sizeof(sycl::half);
-
- if (capacity >= need_bytes) {
- return ptr;
- }
-
- if (ptr) {
- SYCL_CHECK(CHECK_TRY_ERROR(qptr->wait()));
- ggml_sycl_memtrace_del(ptr);
- SYCL_CHECK(CHECK_TRY_ERROR(sycl::free(ptr, *qptr)));
- ptr = nullptr;
- capacity = 0;
- }
-
- size_t cap = 0;
- while (cap < need_bytes) {
- cap += CHUNK_SIZE;
- }
-
- void * dev_ptr;
- SYCL_CHECK(
- CHECK_TRY_ERROR(dev_ptr = sycl::malloc_device(
- cap, *qptr)));
-
- if (!dev_ptr) {
- GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on device\n", __func__, cap);
- ggml_sycl_memtrace_fail(GGML_SYCL_MEM_FATTN_KV, cap);
- GGML_ABORT("fattn buffer alloc failed");
- }
-
- ptr = static_cast<sycl::half *>(dev_ptr);
- capacity = cap;
- ggml_sycl_memtrace_add(GGML_SYCL_MEM_FATTN_KV, ptr, cap);
- return ptr;
-}
-
-ggml_sycl_fattn_kv_buffers::kv_buffer::~kv_buffer() {
-#ifdef DEBUG_SYCL_POOL
- GGML_LOG_INFO("ggml_sycl_fattn_kv_buffer[%d]: %.2f MiB\n", device, capacity / 1024.0 / 1024.0);
-#endif
- if (ptr) {
- ggml_sycl_memtrace_del(ptr);
- SYCL_CHECK(CHECK_TRY_ERROR(sycl::free(ptr, *qptr)));
- }
-}
diff --git a/ggml/src/ggml-sycl/fattn-buffers.hpp b/ggml/src/ggml-sycl/fattn-buffers.hpp
deleted file mode 100644
index c00461de6..000000000
--- a/ggml/src/ggml-sycl/fattn-buffers.hpp
+++ /dev/null
@@ -1,63 +0,0 @@
-//
-// MIT license
-// Copyright (C) 2025 Intel Corporation
-// SPDX-License-Identifier: MIT
-//
-
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-
-#ifndef GGML_SYCL_FATTN_BUFFERS_HPP
-#define GGML_SYCL_FATTN_BUFFERS_HPP
-
-#include <sycl/sycl.hpp>
-
-typedef sycl::queue *queue_ptr;
-
-struct ggml_sycl_fattn_kv_buffers {
- // buffers grow in chunks of this size
- static constexpr size_t CHUNK_SIZE = 16ull << 20; // 16 MiB
-
- struct kv_buffer {
- kv_buffer(queue_ptr qptr_, int device_) : qptr(qptr_), device(device_) {}
- ~kv_buffer();
-
- kv_buffer(const kv_buffer &) = delete;
- kv_buffer & operator=(const kv_buffer &) = delete;
-
- sycl::half * ensure_half(size_t n_elems);
-
- private:
- sycl::half * ptr = nullptr;
- size_t capacity = 0;
- queue_ptr qptr = nullptr;
- [[maybe_unused]] int device = 0;
- };
-
- kv_buffer K;
- kv_buffer V;
-
- ggml_sycl_fattn_kv_buffers(queue_ptr qptr, int device) : K(qptr, device), V(qptr, device) {}
-
- ggml_sycl_fattn_kv_buffers(const ggml_sycl_fattn_kv_buffers &) = delete;
- ggml_sycl_fattn_kv_buffers & operator=(const ggml_sycl_fattn_kv_buffers &) = delete;
-};
-
-/**
- * Imitates `ggml_sycl_pool_alloc` to keep the code calling alloc unchanged.
- */
-struct ggml_sycl_fattn_alloc {
- ggml_sycl_fattn_kv_buffers::kv_buffer & buf;
- sycl::half * ptr = nullptr;
-
- explicit ggml_sycl_fattn_alloc(ggml_sycl_fattn_kv_buffers::kv_buffer & buf_) : buf(buf_) {}
-
- sycl::half * alloc(size_t n_elems) {
- ptr = buf.ensure_half(n_elems);
- return ptr;
- }
-};
-#endif
diff --git a/ggml/src/ggml-sycl/fattn-common.hpp b/ggml/src/ggml-sycl/fattn-common.hpp
index 3c2d1a776..2e8ac6af3 100644
--- a/ggml/src/ggml-sycl/fattn-common.hpp
+++ b/ggml/src/ggml-sycl/fattn-common.hpp
@@ -6,7 +6,6 @@
#include "common.hpp"
#include "convert.hpp"
#include "vecdotq.hpp"
-#include "fattn-buffers.hpp"
#include "fattn.hpp"
#include "ggml.h"
@@ -933,13 +932,10 @@ void launch_fattn(
GGML_ASSERT(!mask || mask->type == GGML_TYPE_F16);
ggml_sycl_pool & pool = ctx.pool();
- ggml_sycl_fattn_kv_buffers & fbuf = ctx.fattn_buffers();
dpct::queue_ptr main_stream = ctx.stream();
const int id = ggml_sycl_get_device();
const int nsm = ggml_sycl_info().devices[id].nsm;
- ggml_sycl_fattn_alloc K_f16(fbuf.K);
- ggml_sycl_fattn_alloc V_f16(fbuf.V);
const ggml_sycl_fattn_extra extra = ggml_sycl_fattn_get_extra(dst);
ggml_sycl_pool_alloc<int> KV_max(pool);
ggml_sycl_pool_alloc<float> dst_tmp(pool);
@@ -959,8 +955,8 @@ void launch_fattn(
const size_t bs = ggml_blck_size(K->type);
const size_t ts = ggml_type_size(K->type);
- sycl::half * K_f16_ptr = extra.K_buffer_ptr ? (sycl::half *) extra.K_buffer_ptr
- : K_f16.alloc(ggml_nelements(K));
+ GGML_ASSERT(extra.K_buffer_ptr);
+ sycl::half * K_f16_ptr = (sycl::half *) extra.K_buffer_ptr;
if (ggml_is_contiguously_allocated(K)) {
to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(K->type, dst);
to_fp16(K_data, K_f16_ptr, ggml_nelements(K), main_stream);
@@ -993,8 +989,8 @@ void launch_fattn(
const size_t bs = ggml_blck_size(V->type);
const size_t ts = ggml_type_size(V->type);
- sycl::half * V_f16_ptr = extra.V_buffer_ptr ? (sycl::half *) extra.V_buffer_ptr
- : V_f16.alloc(ggml_nelements(V));
+ GGML_ASSERT(extra.V_buffer_ptr);
+ sycl::half * V_f16_ptr = (sycl::half *) extra.V_buffer_ptr;
if (ggml_is_contiguously_allocated(V)) {
to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(V->type, dst);
to_fp16(V_data, V_f16_ptr, ggml_nelements(V), main_stream);
diff --git a/ggml/src/ggml-sycl/fattn-mkl.cpp b/ggml/src/ggml-sycl/fattn-mkl.cpp
index 30947b17b..27a8bce8b 100644
--- a/ggml/src/ggml-sycl/fattn-mkl.cpp
+++ b/ggml/src/ggml-sycl/fattn-mkl.cpp
@@ -8,7 +8,6 @@
#include "common.hpp"
#include "fattn-common.hpp"
-#include "fattn-buffers.hpp"
#include "convert.hpp"
#include "fattn.hpp"
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index 49148d7c5..c6aff3955 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -2064,11 +2064,6 @@ std::unique_ptr<ggml_sycl_pool> ggml_backend_sycl_context::new_pool_for_device(q
return std::unique_ptr<ggml_sycl_pool>(new ggml_sycl_pool_leg(qptr, device));
}
-
-std::unique_ptr<ggml_sycl_fattn_kv_buffers> ggml_backend_sycl_context::new_fattn_kv_buffers(queue_ptr qptr, int device) {
- return std::unique_ptr<ggml_sycl_fattn_kv_buffers>(new ggml_sycl_fattn_kv_buffers(qptr, device));
-}
-
/// kernels
typedef void (*ggml_sycl_op_mul_mat_t)(
ggml_backend_sycl_context & ctx,
diff --git a/ggml/src/ggml-sycl/memtrace.cpp b/ggml/src/ggml-sycl/memtrace.cpp
index 9c4f88539..ecd82ff6d 100644
--- a/ggml/src/ggml-sycl/memtrace.cpp
+++ b/ggml/src/ggml-sycl/memtrace.cpp
@@ -15,7 +15,6 @@ static const char * mem_type_name(ggml_sycl_mem_type type) {
case GGML_SYCL_MEM_POOL_LEG: return "pool_leg";
case GGML_SYCL_MEM_POOL_VMM: return "pool_vmm";
case GGML_SYCL_MEM_ASYNC: return "async";
- case GGML_SYCL_MEM_FATTN_KV: return "fattn_kv";
case GGML_SYCL_MEM_DIRECT: return "direct";
default: GGML_ABORT("[%s] The type value %d is not supported\n", __func__, (int) type);
}
diff --git a/ggml/src/ggml-sycl/memtrace.hpp b/ggml/src/ggml-sycl/memtrace.hpp
index 426d90963..c7da47f31 100644
--- a/ggml/src/ggml-sycl/memtrace.hpp
+++ b/ggml/src/ggml-sycl/memtrace.hpp
@@ -10,7 +10,6 @@ enum ggml_sycl_mem_type {
GGML_SYCL_MEM_POOL_LEG,
GGML_SYCL_MEM_POOL_VMM,
GGML_SYCL_MEM_ASYNC,
- GGML_SYCL_MEM_FATTN_KV,
GGML_SYCL_MEM_DIRECT,
GGML_SYCL_MEM_TYPE_COUNT,