diff options
| author | AT <manyoso@users.noreply.github.com> | 2023-06-09 04:00:51 -0400 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2023-06-09 11:00:51 +0300 |
| commit | 92f44ff7f778ef1b94028b2ba6d39943b5ca0ada (patch) | |
| tree | 851ddf5dbe9cbd01a44e8d516aac4c11c351e095 /ggml-metal.metal | |
| parent | 245fc3c37da5ac5963f9f11a9f4f2ac08d96afc6 (diff) | |
metal : add GELU implementation (#1770)
Co-authored-by: Adam Treat <adam@nomic.ai>
Diffstat (limited to 'ggml-metal.metal')
| -rw-r--r-- | ggml-metal.metal | 11 |
1 files changed, 11 insertions, 0 deletions
diff --git a/ggml-metal.metal b/ggml-metal.metal index 8e730eb..745fe8a 100644 --- a/ggml-metal.metal +++ b/ggml-metal.metal @@ -81,6 +81,17 @@ kernel void kernel_relu( dst[tpig] = max(0.0f, src0[tpig]); } +constant float GELU_COEF_A = 0.044715f; +constant float SQRT_2_OVER_PI = 0.79788456080286535587989211986876f; + +kernel void kernel_gelu( + device const float * src0, + device float * dst, + uint tpig[[thread_position_in_grid]]) { + float x = src0[tpig]; + dst[tpig] = 0.5f*x*(1.0f + tanhf(SQRT_2_OVER_PI*x*(1.0f + GELU_COEF_A*x*x))); +} + kernel void kernel_soft_max( device const float * src0, device float * dst, |