From 92f44ff7f778ef1b94028b2ba6d39943b5ca0ada Mon Sep 17 00:00:00 2001 From: AT Date: Fri, 9 Jun 2023 04:00:51 -0400 Subject: metal : add GELU implementation (#1770) Co-authored-by: Adam Treat --- ggml-metal.metal | 11 +++++++++++ 1 file changed, 11 insertions(+) (limited to 'ggml-metal.metal') 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, -- cgit v1.2.3