char-fork/llama.cpp

git clone https://git.t4t.associates/char-fork/llama.cpp

Alexey KopytkoSYCL: reduce allocation overhead during flash attention (#22732)e20b83930

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
1.7 KiB63 linesraw
1//
2// MIT license
3// Copyright (C) 2025 Intel Corporation
4// SPDX-License-Identifier: MIT
5//
6
7//
8// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
9// See https://llvm.org/LICENSE.txt for license information.
10// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
11//
12
13#ifndef GGML_SYCL_FATTN_BUFFERS_HPP
14#define GGML_SYCL_FATTN_BUFFERS_HPP
15
16#include <sycl/sycl.hpp>
17
18typedef sycl::queue *queue_ptr;
19
20struct ggml_sycl_fattn_kv_buffers {
21    // buffers grow in chunks of this size
22    static constexpr size_t CHUNK_SIZE = 16ull << 20; // 16 MiB
23
24    struct kv_buffer {
25        kv_buffer(queue_ptr qptr_, int device_) : qptr(qptr_), device(device_) {}
26        ~kv_buffer();
27
28        kv_buffer(const kv_buffer &) = delete;
29        kv_buffer & operator=(const kv_buffer &) = delete;
30
31        sycl::half * ensure_half(size_t n_elems);
32
33    private:
34        sycl::half * ptr      = nullptr;
35        size_t       capacity = 0;
36        queue_ptr    qptr     = nullptr;
37        [[maybe_unused]] int device = 0;
38    };
39
40    kv_buffer K;
41    kv_buffer V;
42
43    ggml_sycl_fattn_kv_buffers(queue_ptr qptr, int device) : K(qptr, device), V(qptr, device) {}
44
45    ggml_sycl_fattn_kv_buffers(const ggml_sycl_fattn_kv_buffers &) = delete;
46    ggml_sycl_fattn_kv_buffers & operator=(const ggml_sycl_fattn_kv_buffers &) = delete;
47};
48
49/**
50 * Imitates `ggml_sycl_pool_alloc` to keep the code calling alloc unchanged.
51 */
52struct ggml_sycl_fattn_alloc {
53    ggml_sycl_fattn_kv_buffers::kv_buffer & buf;
54    sycl::half *                         ptr = nullptr;
55
56    explicit ggml_sycl_fattn_alloc(ggml_sycl_fattn_kv_buffers::kv_buffer & buf_) : buf(buf_) {}
57
58    sycl::half * alloc(size_t n_elems) {
59        ptr = buf.ensure_half(n_elems);
60        return ptr;
61    }
62};
63#endif