Skip to content

Commit 0db418d

Browse files
ggerganovqnixsynapse
ggerganov
authored and
qnixsynapse
committed
metal : add mean kernel (ggml-org#14267)
* metal : add mean kernel ggml-ci * cont : dedup implementation ggml-ci
1 parent 6e9c823 commit 0db418d

File tree

2 files changed

+67
-14
lines changed

2 files changed

+67
-14
lines changed

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

Lines changed: 28 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -498,6 +498,7 @@ static void ggml_backend_metal_device_rel(struct ggml_backend_metal_device_conte
498498
GGML_METAL_KERNEL_TYPE_COS,
499499
GGML_METAL_KERNEL_TYPE_NEG,
500500
GGML_METAL_KERNEL_TYPE_SUM_ROWS,
501+
GGML_METAL_KERNEL_TYPE_MEAN,
501502
GGML_METAL_KERNEL_TYPE_POOL_2D_AVG_F32,
502503
GGML_METAL_KERNEL_TYPE_POOL_2D_MAX_F32,
503504
GGML_METAL_KERNEL_TYPE_ARGMAX,
@@ -1454,6 +1455,7 @@ @implementation GGMLMetalClass
14541455
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_COS, cos, true);
14551456
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_NEG, neg, true);
14561457
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_SUM_ROWS, sum_rows, true);
1458+
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MEAN, mean, true);
14571459
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ARGMAX, argmax, true);
14581460
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_POOL_2D_AVG_F32, pool_2d_avg_f32, true);
14591461
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_POOL_2D_MAX_F32, pool_2d_max_f32, true);
@@ -1653,6 +1655,7 @@ static bool ggml_metal_supports_op(const struct ggml_backend_metal_device_contex
16531655
case GGML_OP_LOG:
16541656
return false; // TODO: implement
16551657
case GGML_OP_SUM_ROWS:
1658+
case GGML_OP_MEAN:
16561659
case GGML_OP_SOFT_MAX:
16571660
case GGML_OP_GROUP_NORM:
16581661
return has_simdgroup_reduction && ggml_is_contiguous(op->src[0]);
@@ -2400,11 +2403,30 @@ static bool ggml_metal_encode_node(
24002403
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
24012404
} break;
24022405
case GGML_OP_SUM_ROWS:
2406+
case GGML_OP_MEAN:
24032407
{
24042408
GGML_ASSERT(src0->nb[0] == ggml_type_size(src0->type));
24052409

2406-
id<MTLComputePipelineState> pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_SUM_ROWS].pipeline;
2410+
id<MTLComputePipelineState> pipeline = nil;
2411+
2412+
switch (dst->op) {
2413+
case GGML_OP_SUM_ROWS:
2414+
pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_SUM_ROWS].pipeline;
2415+
break;
2416+
case GGML_OP_MEAN:
2417+
pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_MEAN].pipeline;
2418+
break;
2419+
default:
2420+
GGML_ABORT("fatal error");
2421+
}
2422+
2423+
int nth = 32; // SIMD width
2424+
2425+
while (nth < ne00 && nth < (int) pipeline.maxTotalThreadsPerThreadgroup) {
2426+
nth *= 2;
2427+
}
24072428

2429+
nth = MIN(nth, ne00);
24082430

24092431
ggml_metal_kargs_sum_rows args = {
24102432
/*.ne00 =*/ ne00,
@@ -2434,11 +2456,12 @@ static bool ggml_metal_encode_node(
24342456
};
24352457

24362458
[encoder setComputePipelineState:pipeline];
2437-
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
2438-
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
2439-
[encoder setBytes:&args length:sizeof(args) atIndex:2];
2459+
[encoder setBytes:&args length:sizeof(args) atIndex:0];
2460+
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:1];
2461+
[encoder setBuffer:id_dst offset:offs_dst atIndex:2];
2462+
[encoder setThreadgroupMemoryLength:32*sizeof(float) atIndex:0];
24402463

2441-
[encoder dispatchThreadgroups:MTLSizeMake(ne01, ne02, ne03) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
2464+
[encoder dispatchThreadgroups:MTLSizeMake(ne01, ne02, ne03) threadsPerThreadgroup:MTLSizeMake(nth, 1, 1)];
24422465
} break;
24432466
case GGML_OP_SOFT_MAX:
24442467
{

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

Lines changed: 39 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -993,31 +993,61 @@ kernel void kernel_neg(
993993
dst[tpig] = -src0[tpig];
994994
}
995995

996+
template <bool norm>
996997
kernel void kernel_sum_rows(
998+
constant ggml_metal_kargs_sum_rows & args,
997999
device const float * src0,
9981000
device float * dst,
999-
constant ggml_metal_kargs_sum_rows & args,
1000-
uint3 tpig[[thread_position_in_grid]]) {
1001-
int64_t i3 = tpig.z;
1002-
int64_t i2 = tpig.y;
1003-
int64_t i1 = tpig.x;
1001+
threadgroup float * shmem_f32 [[threadgroup(0)]],
1002+
uint3 tgpig[[threadgroup_position_in_grid]],
1003+
ushort3 tpitg[[thread_position_in_threadgroup]],
1004+
ushort sgitg[[simdgroup_index_in_threadgroup]],
1005+
ushort tiisg[[thread_index_in_simdgroup]],
1006+
ushort3 ntg[[threads_per_threadgroup]]) {
1007+
int64_t i3 = tgpig.z;
1008+
int64_t i2 = tgpig.y;
1009+
int64_t i1 = tgpig.x;
10041010

10051011
if (i3 >= args.ne03 || i2 >= args.ne02 || i1 >= args.ne01) {
10061012
return;
10071013
}
10081014

1015+
if (sgitg == 0) {
1016+
shmem_f32[tiisg] = 0.0f;
1017+
}
1018+
10091019
device const float * src_row = (device const float *) ((device const char *) src0 + i1*args.nb01 + i2*args.nb02 + i3*args.nb03);
10101020
device float * dst_row = (device float *) ((device char *) dst + i1*args.nb1 + i2*args.nb2 + i3*args.nb3);
10111021

1012-
float row_sum = 0;
1022+
float sumf = 0;
10131023

1014-
for (int64_t i0 = 0; i0 < args.ne00; i0++) {
1015-
row_sum += src_row[i0];
1024+
for (int64_t i0 = tpitg.x; i0 < args.ne00; i0 += ntg.x) {
1025+
sumf += src_row[i0];
10161026
}
10171027

1018-
dst_row[0] = row_sum;
1028+
sumf = simd_sum(sumf);
1029+
1030+
threadgroup_barrier(mem_flags::mem_threadgroup);
1031+
1032+
if (tiisg == 0) {
1033+
shmem_f32[sgitg] = sumf;
1034+
}
1035+
1036+
threadgroup_barrier(mem_flags::mem_threadgroup);
1037+
1038+
sumf = shmem_f32[tiisg];
1039+
sumf = simd_sum(sumf);
1040+
1041+
if (tpitg.x == 0) {
1042+
dst_row[0] = norm ? sumf / args.ne00 : sumf;
1043+
}
10191044
}
10201045

1046+
typedef decltype(kernel_sum_rows<false>) kernel_sum_rows_t;
1047+
1048+
template [[host_name("kernel_sum_rows")]] kernel kernel_sum_rows_t kernel_sum_rows<false>;
1049+
template [[host_name("kernel_mean")]] kernel kernel_sum_rows_t kernel_sum_rows<true>;
1050+
10211051
template<typename T>
10221052
kernel void kernel_soft_max(
10231053
device const char * src0,

0 commit comments

Comments
 (0)