char-fork/llama.cpp

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

Aman Guptaggml: add SWIGLU_CLAMP (#27930)0190529ec

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
4.9 KiB135 linesraw
1#ifndef GGML_SYCL_ELEMENTWISE_HPP
2#define GGML_SYCL_ELEMENTWISE_HPP
3
4#include "common.hpp"
5#include "ggml.h"
6#include <limits> // For std::numeric_limits
7
8#define SYCL_GLU_BLOCK_SIZE 256
9
10template <typename T>
11T neg_infinity() {
12    return -std::numeric_limits<T>::infinity();
13}
14
15template<typename T_Dst, typename T_Src = T_Dst>
16struct typed_data {
17    const T_Src * src;
18    T_Dst * dst;
19};
20
21template<typename T_Dst, typename T_Src = T_Dst>
22typed_data<T_Dst, T_Src> cast_data(ggml_tensor * dst) {
23    return {
24        /* .src = */ static_cast<const T_Src *>(dst->src[0]->data),
25        /* .dst = */ static_cast<T_Dst *>(dst->data)
26    };
27}
28
29const float GELU_QUICK_COEF = -1.702f;
30
31// Single-element activations, shared with the mat-vec kernels that fuse a GLU epilogue
32// (mmvq.cpp), so both apply the same formula.
33template <typename T> static __dpct_inline__ T op_tanh(T x) {
34    if constexpr (std::is_same_v<T, sycl::ext::oneapi::bfloat16>) {
35#if defined(__INTEL_LLVM_COMPILER) && (__INTEL_LLVM_COMPILER >= 20260000)
36        return sycl::ext::oneapi::experimental::tanh(x);
37#else
38        return static_cast<T>(sycl::tanh(static_cast<float>(x)));
39#endif
40    } else {
41        return sycl::tanh(x);
42    }
43}
44
45template <typename T> static __dpct_inline__ T op_gelu(T x) {
46    const T GELU_COEF_A    = static_cast<T>(0.044715f);
47    const T SQRT_2_OVER_PI = static_cast<T>(0.79788456080286535587989211986876f);
48    return static_cast<T>(0.5f) * x *
49           (static_cast<T>(1.0f) +
50            op_tanh(SQRT_2_OVER_PI * x * (static_cast<T>(1.0f) + GELU_COEF_A * x * x)));
51}
52
53template <typename T> static __dpct_inline__ T op_exp(T x) {
54    if constexpr (std::is_same_v<T, sycl::ext::oneapi::bfloat16>) {
55        return sycl::ext::oneapi::experimental::exp(x);
56    } else {
57        return sycl::exp(x);
58    }
59}
60
61template <typename T> static __dpct_inline__ T op_silu(T x) {
62    return x / (static_cast<T>(1.0f) + op_exp(-x));
63}
64
65void ggml_sycl_sqrt(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
66
67void ggml_sycl_sin(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
68
69void ggml_sycl_cos(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
70
71void ggml_sycl_acc(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
72
73void ggml_sycl_gelu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
74
75void ggml_sycl_silu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
76
77void ggml_sycl_gelu_quick(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
78
79void ggml_sycl_swiglu_oai(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
80void ggml_sycl_swiglu_clamp(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
81
82void ggml_sycl_gelu_erf(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
83
84void ggml_sycl_tanh(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
85
86void ggml_sycl_relu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
87
88void ggml_sycl_sigmoid(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
89
90void ggml_sycl_hardsigmoid(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
91
92void ggml_sycl_hardswish(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
93
94void ggml_sycl_exp(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
95
96void ggml_sycl_expm1(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
97
98void ggml_sycl_log(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
99
100void ggml_sycl_softplus(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
101
102void ggml_sycl_neg(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
103
104void ggml_sycl_step(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
105
106void ggml_sycl_leaky_relu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
107
108void ggml_sycl_sqr(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
109
110void ggml_sycl_clamp(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
111
112void ggml_sycl_xielu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
113
114void ggml_sycl_sgn(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
115
116void ggml_sycl_abs(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
117
118void ggml_sycl_elu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
119
120void ggml_sycl_geglu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
121void ggml_sycl_reglu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
122void ggml_sycl_swiglu(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
123void ggml_sycl_geglu_erf(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
124void ggml_sycl_geglu_quick(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
125void ggml_sycl_floor(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
126void ggml_sycl_ceil(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
127void ggml_sycl_round(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
128void ggml_sycl_trunc(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
129
130void ggml_sycl_arange(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
131
132// fused UNARY(silu|sigmoid|softplus) + MUL; see ggml_sycl_can_fuse() for the accepted shapes
133void ggml_sycl_op_unary_mul_fused(ggml_backend_sycl_context & ctx, ggml_tensor * unary_node, ggml_tensor * mul_node);
134
135#endif // GGML_SYCL_ELEMENTWISE_HPP