From b2d79a7e7fd6101e2695b225a2903c2887152eca Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Tue, 28 Jul 2026 19:02:22 +0300 Subject: [PATCH] metal: fuse snake activation (mul, sin, sqr, mul, add) (#25459) --- ggml/src/ggml-metal/kernels/misc.metal | 28 ++++++++++++++++++++++++++ 1 file changed, 28 insertions(+) diff --git a/ggml/src/ggml-metal/kernels/misc.metal b/ggml/src/ggml-metal/kernels/misc.metal index fe5daa92a1..3c3b4b8587 100644 --- a/ggml/src/ggml-metal/kernels/misc.metal +++ b/ggml/src/ggml-metal/kernels/misc.metal @@ -345,3 +345,31 @@ kernel void kernel_count_equal( typedef decltype(kernel_count_equal) kernel_count_equal_t; template [[host_name("kernel_count_equal_i32")]] kernel kernel_count_equal_t kernel_count_equal; + +template +kernel void kernel_snake( + constant ggml_metal_kargs_snake & args, + device const T * x, + device const float * a, + device const float * inv_b, + device T * dst, + uint tgpig [[threadgroup_position_in_grid]], + uint tpitg [[thread_position_in_threadgroup]], + uint ntg [[threads_per_threadgroup]]) { + + const int idx = tgpig * ntg + tpitg; + if (idx >= args.T * args.C) { + return; + } + + const int c = idx / args.T; // x is [T, C], a / inv_b collapse to [1, C] + const float xi = float(x[idx]); + const float si = sin(a[c] * xi); + dst[idx] = T(xi + si * si * inv_b[c]); +} + +template [[host_name("kernel_snake_f32")]] kernel void kernel_snake(constant ggml_metal_kargs_snake &, device const float *, device const float *, device const float *, device float *, uint, uint, uint); +template [[host_name("kernel_snake_f16")]] kernel void kernel_snake(constant ggml_metal_kargs_snake &, device const half *, device const float *, device const float *, device half *, uint, uint, uint); +#if defined(GGML_METAL_HAS_BF16) +template [[host_name("kernel_snake_bf16")]] kernel void kernel_snake(constant ggml_metal_kargs_snake &, device const bfloat *, device const float *, device const float *, device bfloat *, uint, uint, uint); +#endif