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 pools[GGML_SYCL_MAX_DEVICES]; - std::unique_ptr fattn_bufs[GGML_SYCL_MAX_DEVICES]; - std::unique_ptr host_pools[GGML_SYCL_MAX_DEVICES]; std::vector mmid_row_mapping_host; @@ -418,8 +415,6 @@ struct ggml_backend_sycl_context { static std::unique_ptr new_pool_for_host(queue_ptr qptr, int device); - static std::unique_ptr 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> 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(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 - -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 KV_max(pool); ggml_sycl_pool_alloc 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_backend_sycl_context::new_pool_for_device(q return std::unique_ptr(new ggml_sycl_pool_leg(qptr, device)); } - -std::unique_ptr ggml_backend_sycl_context::new_fattn_kv_buffers(queue_ptr qptr, int device) { - return std::unique_ptr(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,