Skip to content

Commit 65730bc

Browse files
ngxsonngxson
committed
rename na --> erf
1 parent bddf57b commit 65730bc

File tree

7 files changed

+34
-33
lines changed

7 files changed

+34
-33
lines changed

‎ggml/include/ggml.h‎

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -536,7 +536,7 @@ extern "C" {
536536
GGML_UNARY_OP_HARDSWISH,
537537
GGML_UNARY_OP_HARDSIGMOID,
538538
GGML_UNARY_OP_EXP,
539-
GGML_UNARY_OP_GELU_NA,
539+
GGML_UNARY_OP_GELU_ERF,
540540

541541
GGML_UNARY_OP_COUNT,
542542
};
@@ -1025,12 +1025,13 @@ extern "C" {
10251025
struct ggml_context * ctx,
10261026
struct ggml_tensor * a);
10271027

1028-
// GELU without approximation (na = not approximated)
1029-
GGML_API struct ggml_tensor * ggml_gelu_na(
1028+
// GELU using erf (error function) when possible
1029+
// some backends may fallback to approximation based on Abramowitz and Stegun formula
1030+
GGML_API struct ggml_tensor * ggml_gelu_erf(
10301031
struct ggml_context * ctx,
10311032
struct ggml_tensor * a);
10321033

1033-
GGML_API struct ggml_tensor * ggml_gelu_na_inplace(
1034+
GGML_API struct ggml_tensor * ggml_gelu_erf_inplace(
10341035
struct ggml_context * ctx,
10351036
struct ggml_tensor * a);
10361037

‎ggml/src/ggml-cpu/ggml-cpu.c‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2197,7 +2197,7 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) {
21972197
case GGML_UNARY_OP_HARDSWISH:
21982198
case GGML_UNARY_OP_HARDSIGMOID:
21992199
case GGML_UNARY_OP_EXP:
2200-
case GGML_UNARY_OP_GELU_NA:
2200+
case GGML_UNARY_OP_GELU_ERF:
22012201
{
22022202
n_tasks = 1;
22032203
} break;

‎ggml/src/ggml-cpu/ops.cpp‎

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -2691,9 +2691,9 @@ static void ggml_compute_forward_gelu(
26912691
}
26922692
}
26932693

2694-
// ggml_compute_forward_gelu_na
2694+
// ggml_compute_forward_gelu_erf
26952695

2696-
static void ggml_compute_forward_gelu_na_f32(
2696+
static void ggml_compute_forward_gelu_erf_f32(
26972697
const ggml_compute_params * params,
26982698
ggml_tensor * dst) {
26992699

@@ -2717,7 +2717,7 @@ static void ggml_compute_forward_gelu_na_f32(
27172717
const int ir1 = MIN(ir0 + dr, nr);
27182718

27192719
for (int i1 = ir0; i1 < ir1; i1++) {
2720-
ggml_vec_gelu_na_f32(nc,
2720+
ggml_vec_gelu_erf_f32(nc,
27212721
(float *) ((char *) dst->data + i1*( dst->nb[1])),
27222722
(float *) ((char *) src0->data + i1*(src0->nb[1])));
27232723

@@ -2732,7 +2732,7 @@ static void ggml_compute_forward_gelu_na_f32(
27322732
}
27332733
}
27342734

2735-
static void ggml_compute_forward_gelu_na_f16(
2735+
static void ggml_compute_forward_gelu_erf_f16(
27362736
const ggml_compute_params * params,
27372737
ggml_tensor * dst) {
27382738

@@ -2756,7 +2756,7 @@ static void ggml_compute_forward_gelu_na_f16(
27562756
const int ir1 = MIN(ir0 + dr, nr);
27572757

27582758
for (int i1 = ir0; i1 < ir1; i1++) {
2759-
ggml_vec_gelu_na_f16(nc,
2759+
ggml_vec_gelu_erf_f16(nc,
27602760
(ggml_fp16_t *) ((char *) dst->data + i1*( dst->nb[1])),
27612761
(ggml_fp16_t *) ((char *) src0->data + i1*(src0->nb[1])));
27622762

@@ -2772,7 +2772,7 @@ static void ggml_compute_forward_gelu_na_f16(
27722772
}
27732773
}
27742774

2775-
static void ggml_compute_forward_gelu_na(
2775+
static void ggml_compute_forward_gelu_erf(
27762776
const ggml_compute_params * params,
27772777
ggml_tensor * dst) {
27782778

@@ -2781,11 +2781,11 @@ static void ggml_compute_forward_gelu_na(
27812781
switch (src0->type) {
27822782
case GGML_TYPE_F32:
27832783
{
2784-
ggml_compute_forward_gelu_na_f32(params, dst);
2784+
ggml_compute_forward_gelu_erf_f32(params, dst);
27852785
} break;
27862786
case GGML_TYPE_F16:
27872787
{
2788-
ggml_compute_forward_gelu_na_f16(params, dst);
2788+
ggml_compute_forward_gelu_erf_f16(params, dst);
27892789
} break;
27902790
default:
27912791
{
@@ -7852,9 +7852,9 @@ void ggml_compute_forward_unary(
78527852
{
78537853
ggml_compute_forward_gelu(params, dst);
78547854
} break;
7855-
case GGML_UNARY_OP_GELU_NA:
7855+
case GGML_UNARY_OP_GELU_ERF:
78567856
{
7857-
ggml_compute_forward_gelu_na(params, dst);
7857+
ggml_compute_forward_gelu_erf(params, dst);
78587858
} break;
78597859
case GGML_UNARY_OP_GELU_QUICK:
78607860
{

‎ggml/src/ggml-cpu/vec.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -441,7 +441,7 @@ inline static void ggml_vec_gelu_f16(const int n, ggml_fp16_t * y, const ggml_fp
441441
}
442442
}
443443

444-
inline static void ggml_vec_gelu_na_f16(const int n, ggml_fp16_t * y, const ggml_fp16_t * x) {
444+
inline static void ggml_vec_gelu_erf_f16(const int n, ggml_fp16_t * y, const ggml_fp16_t * x) {
445445
for (int i = 0; i < n; ++i) {
446446
float xi = GGML_FP16_TO_FP32(x[i]);
447447
float res = 0.5f*xi*(1.0f + erff(xi*SQRT_2_INV));
@@ -472,7 +472,7 @@ inline static void ggml_vec_gelu_f32(const int n, float * y, const float * x) {
472472
}
473473
#endif
474474

475-
inline static void ggml_vec_gelu_na_f32(const int n, float * y, const float * x) {
475+
inline static void ggml_vec_gelu_erf_f32(const int n, float * y, const float * x) {
476476
for (int i = 0; i < n; ++i) {
477477
float xi = x[i];
478478
y[i] = 0.5f*xi*(1.0f + erff(xi*SQRT_2_INV));

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

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -149,8 +149,8 @@ static void ggml_backend_metal_device_rel(struct ggml_backend_metal_device_conte
149149
GGML_METAL_KERNEL_TYPE_SIGMOID,
150150
GGML_METAL_KERNEL_TYPE_GELU,
151151
GGML_METAL_KERNEL_TYPE_GELU_4,
152-
GGML_METAL_KERNEL_TYPE_GELU_NA,
153-
GGML_METAL_KERNEL_TYPE_GELU_NA_4,
152+
GGML_METAL_KERNEL_TYPE_GELU_ERF,
153+
GGML_METAL_KERNEL_TYPE_GELU_ERF_4,
154154
GGML_METAL_KERNEL_TYPE_GELU_QUICK,
155155
GGML_METAL_KERNEL_TYPE_GELU_QUICK_4,
156156
GGML_METAL_KERNEL_TYPE_SILU,
@@ -1105,8 +1105,8 @@ @implementation GGMLMetalClass
11051105
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_SIGMOID, sigmoid, true);
11061106
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_GELU, gelu, true);
11071107
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_GELU_4, gelu_4, true);
1108-
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_GELU_NA, gelu_na, true);
1109-
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_GELU_NA_4, gelu_na_4, true);
1108+
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_GELU_ERF, gelu_erf, true);
1109+
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_GELU_ERF_4, gelu_erf_4, true);
11101110
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_GELU_QUICK, gelu_quick, true);
11111111
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_GELU_QUICK_4, gelu_quick_4, true);
11121112
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_SILU, silu, true);
@@ -1617,7 +1617,7 @@ static bool ggml_metal_supports_op(const struct ggml_backend_metal_device_contex
16171617
case GGML_UNARY_OP_RELU:
16181618
case GGML_UNARY_OP_SIGMOID:
16191619
case GGML_UNARY_OP_GELU:
1620-
case GGML_UNARY_OP_GELU_NA:
1620+
case GGML_UNARY_OP_GELU_ERF:
16211621
case GGML_UNARY_OP_GELU_QUICK:
16221622
case GGML_UNARY_OP_SILU:
16231623
case GGML_UNARY_OP_ELU:
@@ -2256,17 +2256,17 @@ static bool ggml_metal_encode_node(
22562256

22572257
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
22582258
} break;
2259-
case GGML_UNARY_OP_GELU_NA:
2259+
case GGML_UNARY_OP_GELU_ERF:
22602260
{
22612261
int64_t n = ggml_nelements(dst);
22622262

22632263
id<MTLComputePipelineState> pipeline = nil;
22642264

22652265
if (n % 4 == 0) {
2266-
pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_GELU_NA_4].pipeline;
2266+
pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_GELU_ERF_4].pipeline;
22672267
n /= 4;
22682268
} else {
2269-
pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_GELU_NA].pipeline;
2269+
pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_GELU_ERF].pipeline;
22702270
}
22712271

22722272
[encoder setComputePipelineState:pipeline];

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -916,7 +916,7 @@ T erf_approx(T x) {
916916
return sign_x * y;
917917
}
918918

919-
kernel void kernel_gelu_na(
919+
kernel void kernel_gelu_erf(
920920
device const float * src0,
921921
device float * dst,
922922
uint tpig[[thread_position_in_grid]]) {
@@ -925,7 +925,7 @@ kernel void kernel_gelu_na(
925925
dst[tpig] = 0.5f*x*(1.0f+erf_approx<float>(x*SQRT_2_INV));
926926
}
927927

928-
kernel void kernel_gelu_na_4(
928+
kernel void kernel_gelu_erf_4(
929929
device const float4 * src0,
930930
device float4 * dst,
931931
uint tpig[[thread_position_in_grid]]) {

‎ggml/src/ggml.c‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1099,7 +1099,7 @@ static const char * GGML_UNARY_OP_NAME[GGML_UNARY_OP_COUNT] = {
10991099
"HARDSWISH",
11001100
"HARDSIGMOID",
11011101
"EXP",
1102-
"GELU_NA",
1102+
"GELU_ERF",
11031103
};
11041104

11051105
static_assert(GGML_UNARY_OP_COUNT == 15, "GGML_UNARY_OP_COUNT != 15");
@@ -2502,18 +2502,18 @@ struct ggml_tensor * ggml_gelu_inplace(
25022502
return ggml_unary_inplace(ctx, a, GGML_UNARY_OP_GELU);
25032503
}
25042504

2505-
// ggml_gelu_na
2505+
// ggml_gelu_erf
25062506

2507-
struct ggml_tensor * ggml_gelu_na(
2507+
struct ggml_tensor * ggml_gelu_erf(
25082508
struct ggml_context * ctx,
25092509
struct ggml_tensor * a) {
2510-
return ggml_unary(ctx, a, GGML_UNARY_OP_GELU_NA);
2510+
return ggml_unary(ctx, a, GGML_UNARY_OP_GELU_ERF);
25112511
}
25122512

2513-
struct ggml_tensor * ggml_gelu_na_inplace(
2513+
struct ggml_tensor * ggml_gelu_erf_inplace(
25142514
struct ggml_context * ctx,
25152515
struct ggml_tensor * a) {
2516-
return ggml_unary_inplace(ctx, a, GGML_UNARY_OP_GELU_NA);
2516+
return ggml_unary_inplace(ctx, a, GGML_UNARY_OP_GELU_ERF);
25172517
}
25182518

25192519
// ggml_gelu_quick

0 commit comments

Comments
 (0)