Skip to content

Commit ebc9ed3

Browse files
committed
ggml: metal unary exp & neg
There isn't much peformance gain though. Just for more op coverage Signed-off-by: Molly Sophia <[email protected]>
1 parent d672ecc commit ebc9ed3

File tree

2 files changed

+46
-0
lines changed

2 files changed

+46
-0
lines changed

ggml/src/ggml-metal/ggml-metal.m

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -134,6 +134,8 @@ static void ggml_backend_metal_device_rel(struct ggml_backend_metal_device_conte
134134
GGML_METAL_KERNEL_TYPE_SCALE_4,
135135
GGML_METAL_KERNEL_TYPE_CLAMP,
136136
GGML_METAL_KERNEL_TYPE_TANH,
137+
GGML_METAL_KERNEL_TYPE_EXP,
138+
GGML_METAL_KERNEL_TYPE_NEG,
137139
GGML_METAL_KERNEL_TYPE_RELU,
138140
GGML_METAL_KERNEL_TYPE_SIGMOID,
139141
GGML_METAL_KERNEL_TYPE_GELU,
@@ -741,6 +743,8 @@ @implementation GGMLMetalClass
741743
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_SCALE_4, scale_4, true);
742744
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_CLAMP, clamp, true);
743745
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_TANH, tanh, true);
746+
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_EXP, exp, true);
747+
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_NEG, neg, true);
744748
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_RELU, relu, true);
745749
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_SIGMOID, sigmoid, true);
746750
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_GELU, gelu, true);
@@ -1180,6 +1184,8 @@ static bool ggml_metal_supports_op(const struct ggml_backend_metal_device_contex
11801184
case GGML_UNARY_OP_GELU_QUICK:
11811185
case GGML_UNARY_OP_SILU:
11821186
case GGML_UNARY_OP_ELU:
1187+
case GGML_UNARY_OP_EXP:
1188+
case GGML_UNARY_OP_NEG:
11831189
return ggml_is_contiguous(op->src[0]);
11841190
default:
11851191
return false;
@@ -1747,6 +1753,30 @@ static void ggml_metal_encode_node(
17471753

17481754
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
17491755
} break;
1756+
case GGML_UNARY_OP_EXP:
1757+
{
1758+
id<MTLComputePipelineState> pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_EXP].pipeline;
1759+
1760+
[encoder setComputePipelineState:pipeline];
1761+
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
1762+
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
1763+
1764+
const int64_t n = ggml_nelements(dst);
1765+
1766+
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
1767+
} break;
1768+
case GGML_UNARY_OP_NEG:
1769+
{
1770+
id<MTLComputePipelineState> pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_NEG].pipeline;
1771+
1772+
[encoder setComputePipelineState:pipeline];
1773+
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
1774+
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
1775+
1776+
const int64_t n = ggml_nelements(dst);
1777+
1778+
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
1779+
} break;
17501780
case GGML_UNARY_OP_RELU:
17511781
{
17521782
id<MTLComputePipelineState> pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_RELU].pipeline;

ggml/src/ggml-metal/ggml-metal.metal

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -840,6 +840,22 @@ kernel void kernel_tanh(
840840
dst[tpig] = precise::tanh(x);
841841
}
842842

843+
kernel void kernel_exp(
844+
device const float * src0,
845+
device float * dst,
846+
uint tpig[[thread_position_in_grid]]) {
847+
device const float & x = src0[tpig];
848+
dst[tpig] = precise::exp(x);
849+
}
850+
851+
kernel void kernel_neg(
852+
device const float * src0,
853+
device float * dst,
854+
uint tpig[[thread_position_in_grid]]) {
855+
device const float & x = src0[tpig];
856+
dst[tpig] = -x;
857+
}
858+
843859
constant float GELU_COEF_A = 0.044715f;
844860
constant float GELU_QUICK_COEF = -1.702f;
845861
constant float SQRT_2_OVER_PI = 0.79788456080286535587989211986876f;

0 commit comments

Comments
 (0)