char-fork/llama.cpp

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

Ozymandias_EBONSYCL: Refactor GGML_SYCL_ENABLE_MKL_FA to global var (#26863)f9f09f02c

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
18.1 KiB452 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
14#include <sycl/sycl.hpp>
15#include "dpct/helper.hpp"
16#include "common.hpp"
17#include "fattn-common.hpp"
18#include "fattn-tile.hpp"
19#include "fattn-vec.hpp"
20#include "fattn.hpp"
21#include "fattn-onednn.hpp"
22
23
24#define FATTN_VEC_CASE(D, type_K, type_V)                                                                        \
25    {                                                                                                            \
26        const bool type_K_okay = K->type == (type_K) || (K->type == GGML_TYPE_F32 && (type_K) == GGML_TYPE_F16); \
27        const bool type_V_okay = V->type == (type_V) || (V->type == GGML_TYPE_F32 && (type_V) == GGML_TYPE_F16); \
28        if (Q->ne[0] == (D) && type_K_okay && type_V_okay) {                                                     \
29            ggml_sycl_flash_attn_ext_vec_case<D, type_K, type_V>(ctx, dst);                                      \
30            return;                                                                                              \
31        }                                                                                                        \
32    }                                                                    \
33
34#define FATTN_VEC_CASES_ALL_D(type_K, type_V) \
35    FATTN_VEC_CASE( 64, type_K, type_V)       \
36    FATTN_VEC_CASE(128, type_K, type_V)       \
37    FATTN_VEC_CASE(256, type_K, type_V)       \
38    FATTN_VEC_CASE(512, type_K, type_V)       \
39
40static void ggml_sycl_flash_attn_ext_vec(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
41    ggml_tensor * Q = dst->src[0];
42    ggml_tensor * K = dst->src[1];
43    ggml_tensor * V = dst->src[2];
44
45#ifdef GGML_SYCL_FA_ALL_QUANTS
46    FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16,  GGML_TYPE_F16)
47    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_F16)
48    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_F16)
49    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_F16)
50    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_F16)
51    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_F16)
52
53    FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16,  GGML_TYPE_Q4_0)
54    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q4_0)
55    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_Q4_0)
56    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_Q4_0)
57    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q4_0)
58    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q4_0)
59
60    FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16,  GGML_TYPE_Q4_1)
61    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q4_1)
62    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_Q4_1)
63    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_Q4_1)
64    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q4_1)
65    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q4_1)
66
67    FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16,  GGML_TYPE_Q5_0)
68    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q5_0)
69    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_Q5_0)
70    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_Q5_0)
71    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q5_0)
72    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q5_0)
73
74    FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16,  GGML_TYPE_Q5_1)
75    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q5_1)
76    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_Q5_1)
77    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_Q5_1)
78    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q5_1)
79    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q5_1)
80
81    FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16,  GGML_TYPE_Q8_0)
82    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q8_0)
83    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_Q8_0)
84    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_Q8_0)
85    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q8_0)
86    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q8_0)
87#else
88    FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16,  GGML_TYPE_F16)
89    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q4_0)
90    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q8_0)
91#endif // GGML_SYCL_FA_ALL_QUANTS
92
93    GGML_ABORT("Not match KV type in vec");
94}
95
96// Best FlashAttention kernel for a specific GPU:
97enum best_fattn_kernel {
98    BEST_FATTN_KERNEL_NONE     =   0,
99    BEST_FATTN_KERNEL_VEC      = 100,
100    BEST_FATTN_KERNEL_ONEDNN   = 150, // oneDNN SDPA: native F16 (PR #25222)
101    BEST_FATTN_KERNEL_TILE     = 200,
102    BEST_FATTN_KERNEL_MKL      = 300,
103};
104
105
106static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const ggml_tensor * dst) {
107#ifndef SYCL_FLASH_ATTN
108    GGML_UNUSED(dst);
109    return BEST_FATTN_KERNEL_NONE;
110#endif// SYCL_FLASH_ATTN
111
112    if(!g_ggml_sycl_enable_flash_attention) return BEST_FATTN_KERNEL_NONE;
113
114    const ggml_tensor * KQV   = dst;
115    const ggml_tensor * Q     = dst->src[0];
116    const ggml_tensor * K     = dst->src[1];
117    const ggml_tensor * V     = dst->src[2];
118    const ggml_tensor * mask  = dst->src[3];
119    const ggml_tensor * sinks = dst->src[4];
120
121    const int gqa_ratio = Q->ne[2] / K->ne[2];
122    GGML_ASSERT(Q->ne[2] % K->ne[2] == 0);
123
124    float max_bias = 0.0f;
125    memcpy(&max_bias, (const float *) KQV->op_params + 1, sizeof(float));
126
127    float logit_softcap = 0.0f;
128    memcpy(&logit_softcap, (const float *) KQV->op_params + 2, sizeof(float));
129
130    bool gqa_opt_applies = gqa_ratio >= 2 && mask && max_bias == 0.0f && K->ne[1] % FATTN_KQ_STRIDE == 0;
131
132    // XMX-accelerated path: oneDNN SDPA (native F16 and dequant+non-F16).
133    // ONEDNN requires min 32 query tokens — short-circuit decode to avoid
134    // calling _supported() on every decode FA call.
135    if (Q->ne[1] >= 32
136        && ggml_sycl_flash_attn_ext_onednn_supported(dst)) {
137        return BEST_FATTN_KERNEL_ONEDNN;
138    }
139
140    // MKL path: XMX-accelerated GEMM for prompt processing (all KV cache types).
141    // The MKL kernel converts non-F16 K/V to F16 via to_fp16_sycl before GEMM,
142    // so quantized, F16, BF16, and F32 caches all benefit from XMX acceleration.
143    // Activates automatically when flash-attn is enabled (--flash-attn on or -fa)
144    // and n_kv >= 1024. Falls through to TILE/VEC for ALiBi, logit softcap,
145    // and mismatched batch dimensions (unsupported by the MKL kernel).
146    // Set GGML_SYCL_ENABLE_MKL_FA=0 to force TILE/VEC path for A/B testing.
147    // Example: GGML_SYCL_ENABLE_MKL_FA=0 llama-cli -m model.gguf -fa -ngl 99 ...
148    // Note: MKL GEMM calls are incompatible with SYCL graph capture replay.
149    // MKL is validated for the mainstream GQA envelope: grouped-query
150    // (gqa_ratio >= 2), head_dim a multiple of 64 in [64,512] with matching
151    // K/V head size, mask, no sinks/ALiBi/softcap. Gemma's global layers use
152    // head_dim 512, so the cap must include it. Head sizes not a multiple of
153    // 64 (72/80/96), MHA (gqa_ratio == 1), and MLA (DKQ != DV, e.g. 576/512)
154    // fall through to TILE/VEC; see follow-up work.
155    if (g_ggml_sycl_enable_mkl_fa == 1 && mask && !sinks && gqa_ratio >= 2 &&
156        Q->ne[0] >= 64 && Q->ne[0] <= 512 && Q->ne[0] % 64 == 0 &&
157        Q->ne[0] == V->ne[0] &&
158        Q->ne[1] >= 32 && K->ne[1] >= 1024 &&
159        max_bias == 0.0f && logit_softcap == 0.0f &&
160        (Q->ne[3] == K->ne[3] || K->ne[3] == 1)) {
161        // F16 K/V strides must be a multiple of ne[0]*2 (the natural row size
162        // in bytes). This passes both dense (nb1 == ne0*2) and interleaved
163        // (nb1 == H * ne0*2). Only pathological test strides like nb1=32 or
164        // nb1=75 for ne0=40 fall through to TILE.
165        bool kv_strides_ok = true;
166        for (const ggml_tensor * t : {K, V}) {
167            if (t->type == GGML_TYPE_F16 && t->nb[1] % (t->ne[0] * 2) != 0) {
168                kv_strides_ok = false;
169                break;
170            }
171        }
172        if (kv_strides_ok) {
173            return BEST_FATTN_KERNEL_MKL;
174        }
175    }
176    for (const ggml_tensor * t : {Q, K, V, mask}) {
177        if (t == nullptr || ggml_is_quantized(t->type)) {
178            continue;
179        }
180        for (size_t i = 1; i < GGML_MAX_DIMS; ++i) {
181            if (t->nb[i] % 16 != 0) {
182                gqa_opt_applies = false;
183                break;
184            }
185        }
186    }
187
188    switch (K->ne[0]) {
189        case  40:
190        case  64:
191        case  72:
192        case  80:
193        case  96:
194        case 128:
195        case 112:
196        case 256:
197        case 512:
198            if (V->ne[0] != K->ne[0]) {
199                return BEST_FATTN_KERNEL_NONE;
200            }
201            break;
202        case 576:
203            if (V->ne[0] != 512) {
204                return BEST_FATTN_KERNEL_NONE;
205            }
206            if (!gqa_opt_applies) {
207                return BEST_FATTN_KERNEL_NONE;
208            }
209            break;
210        default:
211            return BEST_FATTN_KERNEL_NONE;
212    }
213
214#ifndef GGML_SYCL_FA_ALL_QUANTS
215    if (K->type != V->type) {
216        return BEST_FATTN_KERNEL_NONE;
217    }
218#endif // GGML_SYCL_FA_ALL_QUANTS
219
220    switch (K->type) {
221        case GGML_TYPE_F32:
222        case GGML_TYPE_F16:
223        case GGML_TYPE_BF16:
224            break;
225        case GGML_TYPE_Q4_1:
226        case GGML_TYPE_Q5_0:
227        case GGML_TYPE_Q5_1:
228#ifndef GGML_SYCL_FA_ALL_QUANTS
229            return BEST_FATTN_KERNEL_NONE;
230#endif // GGML_SYCL_FA_ALL_QUANTS
231        case GGML_TYPE_Q4_0:
232        case GGML_TYPE_Q8_0:
233            break;
234        default:
235            return BEST_FATTN_KERNEL_NONE;
236    }
237
238    if (mask && mask->ne[2] != 1) {
239        return BEST_FATTN_KERNEL_NONE;
240    }
241
242    // For small batch sizes the vector kernel may be preferable over the kernels optimized for large batch sizes.
243    // BF16 is excluded: the VEC kernel has no BF16 template (it needs GGML_SYCL_FA_ALL_QUANTS for non-F16/Q4_0/Q8_0).
244    const bool has_bf16 = (K->type == GGML_TYPE_BF16 || V->type == GGML_TYPE_BF16);
245    const bool can_use_vector_kernel = Q->ne[0] <= 512 && Q->ne[0] % 64 == 0 && K->ne[1] % FATTN_KQ_STRIDE == 0
246        && !has_bf16;
247
248    // Fused-XMX path: oneDNN Graph SDPA (flash attention). Strictly
249    // additive -- taken only when statically supported, otherwise falls through to VEC/TILE below.
250    if (ggml_sycl_flash_attn_ext_onednn_supported(dst)) {
251        return BEST_FATTN_KERNEL_ONEDNN;
252    }
253
254    // If there are no tensor cores available, use the generic tile kernel:
255    if (can_use_vector_kernel) {
256        if (!ggml_is_quantized(K->type) && !ggml_is_quantized(V->type)) {
257            if (Q->ne[1] == 1) {
258                if (!gqa_opt_applies) {
259                    return BEST_FATTN_KERNEL_VEC;
260                }
261            }
262        } else {
263            if (Q->ne[1] <= 2) {
264                // TILE is faster for quantized KV decode on Xe2 (BMG); keep VEC on untested archs
265                const gpu_arch arch = ggml_sycl_info().devices[device].hw_info.arch;
266                if (arch == gpu_arch::intel_gpu_bmg_g21 || arch == gpu_arch::intel_gpu_bmg_g31) {
267                    return BEST_FATTN_KERNEL_TILE;
268                }
269                return BEST_FATTN_KERNEL_VEC;
270            }
271        }
272    }
273    return BEST_FATTN_KERNEL_TILE;
274}
275
276void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
277    ggml_sycl_set_device(ctx.device);
278
279    // n_kv watchdog: log when n_kv differs from the last FA call with
280    // the same D — helps detect cache-truncation issues.
281    static int nkv_debug = ggml_sycl_get_env("GGML_SYCL_MKL_FA_DEBUG", 0);
282    if (nkv_debug == 1) {
283        const ggml_tensor * K_dbg = dst->src[1];
284        const ggml_tensor * V_dbg = dst->src[2];
285        static int64_t last_nkv_d256 = 0, last_nkv_d512 = 0;
286        static int fa_call_seq = 0;
287        fa_call_seq++;
288        int64_t cur_nkv = K_dbg->ne[1];
289        int Dk = (int)K_dbg->ne[0];
290        const char * kname = "TILE";
291        best_fattn_kernel k = ggml_sycl_get_best_fattn_kernel(ctx.device, dst);
292        if (k == BEST_FATTN_KERNEL_MKL)  kname = "MKL";
293        if (k == BEST_FATTN_KERNEL_ONEDNN)  kname = "ONEDNN";
294        if (k == BEST_FATTN_KERNEL_VEC)  kname = "VEC";
295        int64_t delta = 0;
296        if (Dk == 256) {
297            delta = cur_nkv - last_nkv_d256;
298            last_nkv_d256 = cur_nkv;
299        } else if (Dk == 512) {
300            delta = cur_nkv - last_nkv_d512;
301            last_nkv_d512 = cur_nkv;
302        }
303        GGML_LOG_INFO("[FA-DISP] #%d %s D=%d n_kv=%lld delta=%lld "
304                "V_ne1=%lld\n",
305                fa_call_seq, kname, Dk,
306                (long long)cur_nkv, (long long)delta,
307                (long long)V_dbg->ne[1]);
308    }
309
310    const best_fattn_kernel fk = ggml_sycl_get_best_fattn_kernel(ggml_sycl_get_device(), dst);
311    switch (fk) {
312        case BEST_FATTN_KERNEL_NONE:
313            GGML_ABORT("Not support Flash-Attention");
314        case BEST_FATTN_KERNEL_ONEDNN:
315            // guarded: ggml_sycl_flash_attn_ext_onednn() is only defined under GGML_SYCL_DNNL;
316            // the reference must be compiled out here or the GGML_SYCL_DNNL=0 build fails to link.
317#if GGML_SYCL_DNNL
318            ggml_sycl_flash_attn_ext_onednn(ctx, dst);
319#endif
320            break;
321        case BEST_FATTN_KERNEL_TILE:
322            ggml_sycl_flash_attn_ext_tile(ctx, dst);
323            break;
324        case BEST_FATTN_KERNEL_VEC:
325            ggml_sycl_flash_attn_ext_vec(ctx, dst);
326            break;
327        case BEST_FATTN_KERNEL_MKL:
328            ggml_sycl_flash_attn_ext_mkl(ctx, dst);
329            break;
330    }
331
332    // --- Output fingerprint (GGML_SYCL_MKL_FA_DIAG=1) ---
333    // Copy first 64 float output values to host for fingerprinting.
334    // Compare MKL vs TILE (GGML_SYCL_ENABLE_MKL_FA=0) to detect divergence.
335    // Only fingerprints the first 6 FA calls with n_kv >= 1024.
336    static int fa_diag = ggml_sycl_get_env("GGML_SYCL_MKL_FA_DIAG", 0);
337    static int fa_diag_count = 0;
338    if (fa_diag == 1 && fa_diag_count < 6) {
339        const ggml_tensor * K_diag = dst->src[1];
340        const ggml_tensor * V_diag = dst->src[2];
341        const ggml_tensor * Q_diag = dst->src[0];
342        if (K_diag->ne[1] >= 1024) {
343            fa_diag_count++;
344            float diag_buf[64];
345            dpct::queue_ptr q = ctx.stream();
346            q->memcpy(diag_buf, dst->data, 64 * sizeof(float));
347            q->wait();
348            const char * kname = "???";
349            best_fattn_kernel kb = ggml_sycl_get_best_fattn_kernel(ctx.device, dst);
350            if (kb == BEST_FATTN_KERNEL_ONEDNN) kname = "ONEDNN";
351            if (kb == BEST_FATTN_KERNEL_MKL) kname = "MKL";
352            if (kb == BEST_FATTN_KERNEL_TILE) kname = "TILE";
353            if (kb == BEST_FATTN_KERNEL_VEC) kname = "VEC";
354            GGML_LOG_INFO("[FA-DIAG] #%d %s D=%d n_kv=%lld n_q=%lld "
355                    "n_qh=%lld n_kvh=%lld K=%s V=%s "
356                    "nb1=%zu nb2=%zu first 64 floats:\n",
357                    fa_diag_count, kname,
358                    (int)K_diag->ne[0], (long long)K_diag->ne[1],
359                    (long long)Q_diag->ne[1],
360                    (long long)Q_diag->ne[2], (long long)K_diag->ne[2],
361                    ggml_type_name(K_diag->type),
362                    ggml_type_name(V_diag->type),
363                    K_diag->nb[1], K_diag->nb[2]);
364            for (int i = 0; i < 64; i += 8) {
365                GGML_LOG_INFO("  [%2d] %08x %08x %08x %08x %08x %08x %08x %08x\n",
366                        i,
367                        *(unsigned *)&diag_buf[i+0], *(unsigned *)&diag_buf[i+1],
368                        *(unsigned *)&diag_buf[i+2], *(unsigned *)&diag_buf[i+3],
369                        *(unsigned *)&diag_buf[i+4], *(unsigned *)&diag_buf[i+5],
370                        *(unsigned *)&diag_buf[i+6], *(unsigned *)&diag_buf[i+7]);
371            }
372        }
373    }
374
375}
376
377bool ggml_sycl_flash_attn_ext_supported(int device, const ggml_tensor * dst) {
378    return ggml_sycl_get_best_fattn_kernel(device, dst) != BEST_FATTN_KERNEL_NONE;
379}
380
381static uintptr_t ggml_sycl_fattn_reserve_halves(ggml_sycl_fattn_extra & extra, size_t n_halves) {
382    if (n_halves == 0) {
383        return 0;
384    }
385    extra.end = GGML_PAD(extra.end, SYCL_BUFFER_ALIGNMENT);
386    const uintptr_t block = extra.end;
387    extra.end += n_halves * sizeof(sycl::half);
388    return block;
389}
390
391ggml_sycl_fattn_extra ggml_sycl_fattn_get_extra(const ggml_tensor * dst) {
392    ggml_sycl_fattn_extra extra;
393
394    extra.end = (uintptr_t) dst->data + ggml_nbytes(dst);
395
396    if (dst->op != GGML_OP_FLASH_ATTN_EXT) {
397        return extra;
398    }
399
400    const ggml_tensor * Q = dst->src[0];
401    const ggml_tensor * K = dst->src[1];
402    const ggml_tensor * V = dst->src[2];
403    if (!Q || !K || !V) {
404        return extra;
405    }
406
407    const int64_t d = K->ne[0];
408    const int64_t H = Q->ne[2];
409    const int64_t q = Q->ne[1];
410
411    // calculate the worst-case memory consumption across all kernels
412    const bool onednn_supported = ggml_sycl_flash_attn_ext_onednn_supported(dst, /* use_shape_limit */ false);
413
414    const bool tile_needs_K = K->type != GGML_TYPE_F16;
415    const bool tile_needs_V = V->type != GGML_TYPE_F16;
416
417    const bool V_is_K_view = V->view_src &&
418        (V->view_src == K || (V->view_src == K->view_src && V->view_offs == K->view_offs));
419
420    size_t need_K = 0, need_V = 0, need_Q = 0, need_out = 0, need_scale = 0;
421    if (onednn_supported) {
422        need_Q     = (size_t) H * q * d;
423        need_out   = (size_t) H * q * d;
424        need_scale = 1;
425        // an f16 cache is bound in place, so it needs no staging copy
426        if (!ggml_sycl_fattn_onednn_binds_kv(K, V)) {
427            need_K = (size_t) ggml_nelements(K);
428            need_V = (size_t) ggml_nelements(V);
429        }
430    }
431    if (tile_needs_K) {
432        need_K = std::max(need_K, (size_t) ggml_nelements(K));
433    }
434    if (tile_needs_V) {
435        need_V = std::max(need_V, (size_t) ggml_nelements(V));
436    }
437
438    extra.Q_buffer_ptr = ggml_sycl_fattn_reserve_halves(extra, need_Q);
439    extra.K_buffer_ptr = ggml_sycl_fattn_reserve_halves(extra, need_K);
440    extra.V_buffer_ptr = (V_is_K_view && !onednn_supported && need_V)
441                       ? extra.K_buffer_ptr
442                       : ggml_sycl_fattn_reserve_halves(extra, need_V);
443    extra.scale_buffer_ptr = ggml_sycl_fattn_reserve_halves(extra, need_scale);
444    extra.out_buffer_ptr   = ggml_sycl_fattn_reserve_halves(extra, need_out);
445
446    return extra;
447}
448
449size_t ggml_sycl_flash_attn_ext_get_alloc_size(const ggml_tensor * dst) {
450    const ggml_sycl_fattn_extra extra = ggml_sycl_fattn_get_extra(dst);
451    return (size_t) (extra.end - (uintptr_t) dst->data);
452}