public_git/whisper.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/whisper.cpp.git synced 2026-10-20 08:52:24 +02:00
Code Issues Packages Projects Releases Wiki Activity

sycl : fattn_kv_buffers cleanup (llama/27689)

Browse Source
This commit is contained in:
Pratyush Kumar
2026-10-19 09:47:10 +02:00
committed by Daniel Bevenius
parent 372eda18ab
commit 7eeb1caf65
8 changed files with 4 additions and 155 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-sycl/common.hpp
-16
View File
@@ -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
ggml/src/ggml-sycl/fattn-buffers.cpp
-60
View File
@@ -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)));
}
}
ggml/src/ggml-sycl/fattn-buffers.hpp
-63
View File
@@ -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
ggml/src/ggml-sycl/fattn-common.hpp
+4 -8
View File
@@ -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);
ggml/src/ggml-sycl/fattn-mkl.cpp
-1
View File
@@ -8,7 +8,6 @@
#include "common.hpp"
#include "fattn-common.hpp"
#include "fattn-buffers.hpp"
#include "convert.hpp"
#include "fattn.hpp"
ggml/src/ggml-sycl/ggml-sycl.cpp
-5
View File
@@ -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,
ggml/src/ggml-sycl/memtrace.cpp
-1
View File
@@ -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);
}
ggml/src/ggml-sycl/memtrace.hpp
-1
View File
@@ -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,