char-fork/llama.cpp

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

Akarshan Biswassycl: Improve mul_mat_id memory efficiency and add BF16 fast path (#22119)225088ea7

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
2.2 KiB64 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_CONVERT_HPP
14#define GGML_SYCL_CONVERT_HPP
15
16#include "common.hpp"
17
18template <typename T>
19using to_t_sycl_t = void (*)(const void * __restrict__ x, T * __restrict__ y, int64_t k, dpct::queue_ptr stream);
20typedef to_t_sycl_t<float>      to_fp32_sycl_t;
21typedef to_t_sycl_t<sycl::half> to_fp16_sycl_t;
22
23to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst);
24to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor * dst);
25
26#ifdef GGML_SYCL_HAS_BF16
27typedef to_t_sycl_t<sycl::ext::oneapi::bfloat16> to_bf16_sycl_t;
28to_bf16_sycl_t ggml_get_to_bf16_sycl(ggml_type type, ggml_tensor * dst);
29#endif
30
31// Nc = Non-contiguous
32template <typename T>
33using to_t_nc_sycl_t = void (*)(const void * x, T * y, int64_t ne00, int64_t ne01, int64_t ne02, int64_t ne03,
34                                   int64_t s01, int64_t s02, int64_t s03, dpct::queue_ptr queue);
35
36typedef to_t_nc_sycl_t<sycl::half> to_fp16_nc_sycl_t;
37to_fp16_nc_sycl_t ggml_get_to_fp16_nc_sycl(ggml_type type);
38
39template<typename dst_t, typename src_t>
40 inline dst_t ggml_sycl_cast(src_t x) {
41    if constexpr (std::is_same_v<dst_t, src_t>) {
42        return x;
43#ifdef GGML_SYCL_HAS_BF16
44    } else if constexpr (std::is_same_v<dst_t, sycl::ext::oneapi::bfloat16>) {
45        return sycl::ext::oneapi::bfloat16(float(x));
46    } else if constexpr (std::is_same_v<src_t, sycl::ext::oneapi::bfloat16>) {
47        return static_cast<float>(x);
48#endif
49    } else if constexpr (std::is_same_v<src_t, sycl::float2> && std::is_same_v<dst_t, sycl::half2>) {
50        return x.template convert<sycl::half, sycl::rounding_mode::rte>();
51#ifdef GGML_SYCL_HAS_BF16
52    } else if constexpr (std::is_same_v<src_t, sycl::float2> &&
53                         std::is_same_v<dst_t, sycl::vec<sycl::ext::oneapi::bfloat16, 2>>) {
54        return {x.x, x.y};
55#endif
56    } else if constexpr(std::is_same_v<dst_t, int32_t>) {
57        return int32_t(x);
58    } else {
59        return float(x);
60    }
61}
62
63
64#endif  // GGML_SYCL_CONVERT_HPP