Revert "Fused soft cap and SIMD-ified GeLU #9"

This reverts commit bb500aa.

Author: Nexesenex

ggml/include/ggml.h

@@ -544,7 +544,6 @@ extern "C" {
         GGML_OP_TIMESTEP_EMBEDDING,
         GGML_OP_ARGSORT,
         GGML_OP_LEAKY_RELU,
-        GGML_OP_SOFTCAP,
 
         GGML_OP_FLASH_ATTN_EXT,
         GGML_OP_FLASH_ATTN_BACK,
@@ -1208,19 +1207,6 @@ extern "C" {
             struct ggml_tensor  * a,
             float                 s);
 
-    GGML_API struct ggml_tensor * ggml_softcap(
-            struct ggml_context * ctx,
-            struct ggml_tensor  * a,
-            float                 s_before,
-            float                 s_after);
-
-    // in-place, returns view(a)
-    GGML_API struct ggml_tensor * ggml_softcap_inplace(
-            struct ggml_context * ctx,
-            struct ggml_tensor  * a,
-            float                 s_before,
-            float                 s_after);
-
     // b -> view(a,offset,nb1,nb2,3), return modified a
     GGML_API struct ggml_tensor * ggml_set(
             struct ggml_context * ctx,

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

@@ -1887,15 +1887,15 @@ inline static void ggml_vec_hardswish_f32 (const int n, float * y, const float *
 inline static void ggml_vec_hardsigmoid_f32 (const int n, float * y, const float * x) { for (int i = 0; i < n; ++i) y[i] = fminf(1.0f, fmaxf(0.0f, (x[i] + 3.0f) / 6.0f)); }
 inline static void ggml_vec_exp_f32 (const int n, float * y, const float * x) { for (int i = 0; i < n; ++i) y[i] = expf(x[i]); }
 
-static const float GELU_QUICK_COEF = -1.702f;
 static const float GELU_COEF_A     = 0.044715f;
+static const float GELU_QUICK_COEF = -1.702f;
 static const float SQRT_2_OVER_PI  = 0.79788456080286535587989211986876f;
 
 inline static float ggml_gelu_f32(float x) {
     return 0.5f*x*(1.0f + tanhf(SQRT_2_OVER_PI*x*(1.0f + GELU_COEF_A*x*x)));
 }
 
-/* inline static void ggml_vec_gelu_f16(const int n, ggml_fp16_t * y, const ggml_fp16_t * x) {
+inline static void ggml_vec_gelu_f16(const int n, ggml_fp16_t * y, const ggml_fp16_t * x) {
     const uint16_t * i16 = (const uint16_t *) x;
     for (int i = 0; i < n; ++i) {
         y[i] = ggml_table_gelu_f16[i16[i]];
@@ -1923,7 +1923,7 @@ inline static void ggml_vec_gelu_f32(const int n, float * y, const float * x) {
         y[i] = ggml_gelu_f32(x[i]);
     }
 }
-#endif */
+#endif
 
 inline static float ggml_gelu_quick_f32(float x) {
     return x*(1.0f/(1.0f+expf(GELU_QUICK_COEF*x)));
@@ -1999,34 +1999,7 @@ inline static float32x4_t ggml_v_silu(float32x4_t x) {
     const float32x4_t neg_x = vsubq_f32(zero, x);
     const float32x4_t exp_neg_x = ggml_v_expf(neg_x);
     const float32x4_t one_plus_exp_neg_x = vaddq_f32(one, exp_neg_x);
-    // return vdivq_f32(x, one_plus_exp_neg_x);
-    const uint32x4_t mask = vcgtq_f32(x, vdupq_n_f32(10.f));
-    const float32x4_t res = vdivq_f32(vsubq_f32(exp_two_x, one), vaddq_f32(exp_two_x, one));
-    return vreinterpretq_f32_u32(vorrq_u32(vandq_u32(vreinterpretq_u32_f32(one), mask), vbicq_u32(vreinterpretq_u32_f32(res), mask)));
-    //return vdivq_f32(vsubq_f32(exp_two_x, one), vaddq_f32(exp_two_x, one));
-}
-
-inline static float32x4_t ggml_v_softcap(float32x4_t x, float32x4_t s_before, float32x4_t s_after) {
-    return vmulq_f32(s_after, ggml_v_tanh(vmulq_f32(x, s_before)));
-    //const float32x4_t one = vdupq_n_f32(1.0f);
-    //const float32x4_t two_x = vmulq_f32(x, s_before);
-    //const float32x4_t exp_two_x = ggml_v_expf(two_x);
-    //const float32x4_t th = vdivq_f32(vsubq_f32(exp_two_x, one), vaddq_f32(exp_two_x, one));
-    //return vmulq_f32(th, s_after);
-}
-
-
-// Slower than lookup on my M2-Max
-inline static float32x4_t ggml_v_gelu(float32x4_t x, float32x4_t c1, float32x4_t c2) {
-    const float32x4_t one = vdupq_n_f32(1.0f);
-    //float32x4_t arg = vaddq_f32(one, vmulq_f32(vmulq_f32(x, x), c1));
-    float32x4_t arg = vfmaq_f32(one, c1, vmulq_f32(x, x));
-    arg = vmulq_f32(arg, vmulq_f32(x, c2));
-    float32x4_t exp_arg = ggml_v_expf(arg);
-    float32x4_t gelu = vmulq_f32(x, vdivq_f32(exp_arg, vaddq_f32(exp_arg, one)));
-    uint32x4_t mask = vcgtq_f32(x, vdupq_n_f32(10.f));
-    return vbslq_f32(mask, x, gelu);
-    //return vreinterpretq_f32_u32(vorrq_u32(vandq_u32(vreinterpretq_u32_f32(x), mask), vbicq_u32(vreinterpretq_u32_f32(gelu), mask)));
+    return vdivq_f32(x, one_plus_exp_neg_x);
 }
 
 inline static float32x4_t ggml_v_tanh(float32x4_t x) {
@@ -2082,28 +2055,7 @@ inline static __m512 ggml_v_silu(__m512 x) {
 inline static __m512 ggml_v_tanh(__m512 x) {
     const __m512 one = _mm512_set1_ps(1.0f);
     const __m512 exp_two_x = ggml_v_expf(_mm512_mul_ps(x, _mm512_set1_ps(2.f)));
-    // return _mm512_div_ps(_mm512_sub_ps(exp_two_x, one), _mm512_add_ps(exp_two_x, one));
-    const __mmask16 mask = _mm512_cmp_ps_mask(x, _mm512_set1_ps(10.f), _CMP_GT_OQ);
-    const __m512 res = _mm512_div_ps(_mm512_sub_ps(exp_two_x, one), _mm512_add_ps(exp_two_x, one));
-    return _mm512_mask_blend_ps(mask, res, one);
-}
-
-inline static __m512 ggml_v_softcap(__m512 x, __m512 s_before, __m512 s_after) {
-    const __m512 one = _mm512_set1_ps(1.0f);
-    const __m512 exp_two_x = ggml_v_expf(_mm512_mul_ps(x, s_before));
-    const __m512 th = _mm512_div_ps(_mm512_sub_ps(exp_two_x, one), _mm512_add_ps(exp_two_x, one));
-    return _mm512_mul_ps(th, s_after);
-}
-
-inline static __m512 ggml_v_gelu(__m512 x, __m512 c1, __m512 c2) {
-    const __m512 one = _mm512_set1_ps(1.0f);
-    __m512 arg = _mm512_fmadd_ps(x, _mm512_mul_ps(c1, x), one);
-    //__m512 arg = _mm512_add_ps(one, _mm512_mul_ps(_mm512_mul_ps(x, x), c1));
-    arg = _mm512_mul_ps(arg, _mm512_mul_ps(c2, x));
-    const __mmask16 mask = _mm512_cmp_ps_mask(arg, _mm512_set1_ps(30.f), _CMP_GT_OQ);
-    const __m512 exp_arg = ggml_v_expf(arg);
-    const __m512 ratio = _mm512_div_ps(exp_arg, _mm512_add_ps(exp_arg, one));
-    return _mm512_mul_ps(x, _mm512_mask_blend_ps(mask, ratio, one));
+    return _mm512_div_ps(_mm512_sub_ps(exp_two_x, one), _mm512_add_ps(exp_two_x, one));
 }
 
 #elif defined(__AVX2__) && defined(__FMA__)
@@ -2164,28 +2116,7 @@ inline static __m256 ggml_v_silu(__m256 x) {
 inline static __m256 ggml_v_tanh(__m256 x) {
     const __m256 one = _mm256_set1_ps(1.0f);
     const __m256 exp_two_x = ggml_v_expf(_mm256_mul_ps(x, _mm256_set1_ps(2.f)));
-    // return _mm256_div_ps(_mm256_sub_ps(exp_two_x, one), _mm256_add_ps(exp_two_x, one));
-    const __m256 res = _mm256_div_ps(_mm256_sub_ps(exp_two_x, one), _mm256_add_ps(exp_two_x, one));
-    const __m256 mask = _mm256_cmp_ps(x, _mm256_set1_ps(10.f), _CMP_GT_OQ);
-    return _mm256_or_ps(_mm256_and_ps(mask, one), _mm256_andnot_ps(mask, res));
-}
-
-inline static __m256 ggml_v_softcap(__m256 x, float s_before, float s_after) {
-    return _mm256_mul_ps(_mm256_set1_ps(s_after), ggml_v_tanh(_mm256_mul_ps(x, _mm256_set1_ps(s_before))));
-    //const __m256 one = _mm256_set1_ps(1.0f);
-    //const __m256 exp_two_x = ggml_v_expf(_mm256_mul_ps(x, _mm256_set1_ps(2.f*s_before)));
-    //const __m256 th = _mm256_div_ps(_mm256_sub_ps(exp_two_x, one), _mm256_add_ps(exp_two_x, one));
-    //return _mm256_mul_ps(th, _mm256_set1_ps(s_after));
-}
-
-inline static __m256 ggml_v_gelu(__m256 x, __m256 c1, __m256 c2) {
-    const __m256 one = _mm256_set1_ps(1.0f);
-    const __m256 mask = _mm256_cmp_ps(x, _mm256_set1_ps(10.f), _CMP_GT_OQ);
-    __m256 arg = _mm256_add_ps(one, _mm256_mul_ps(_mm256_mul_ps(x, x), c1));
-    arg = _mm256_mul_ps(arg, _mm256_mul_ps(x, c2));
-    __m256 exp_arg = ggml_v_expf(arg);
-    __m256 gelu = _mm256_mul_ps(x, _mm256_div_ps(exp_arg, _mm256_add_ps(exp_arg, one)));
-    return _mm256_or_ps(_mm256_and_ps(mask, x), _mm256_andnot_ps(mask, gelu));
+    return _mm256_div_ps(_mm256_sub_ps(exp_two_x, one), _mm256_add_ps(exp_two_x, one));
 }
 
 #elif defined(__SSE2__) // __AVX2__ / __ARM_NEON
@@ -2248,13 +2179,6 @@ inline static __m128 ggml_v_tanh(__m128 x) {
     return _mm_div_ps(_mm_sub_ps(exp_two_x, one), _mm_add_ps(exp_two_x, one));
 }
 
-inline static __m128 ggml_v_softcap(__m128 x, float s_before, float s_after) {
-    const __m128 one = _mm_set1_ps(1.0f);
-    const __m128 exp_two_x = ggml_v_expf(_mm_mul_ps(x, _mm_set1_ps(2.f*s_before)));
-    const __m128 th = _mm_div_ps(_mm_sub_ps(exp_two_x, one), _mm_add_ps(exp_two_x, one));
-    return _mm_mul_ps(th, _mm_set1_ps(s_after));
-}
-
 #endif // __ARM_NEON / __AVX2__ / __SSE2__
 
 static void ggml_vec_silu_f32(const int n, float * y, const float * x) {
@@ -2314,108 +2238,6 @@ static void ggml_vec_tanh_f32(const int n, float * y, const float * x) {
     }
 }
 
-static void ggml_vec_softcap_f32(const int n, float * x, float s_before, float s_after) {
-    int i = 0;
-#if defined(__AVX512F__) && defined(__AVX512DQ__)
-    __m512 vs_before = _mm512_set1_ps(2.f*s_before);
-    __m512 vs_after  = _mm512_set1_ps(s_after);
-    //for (; i + 63 < n; i += 64) {
-    //    __m512 x1 = _mm512_loadu_ps(x + i);
-    //    __m512 x2 = _mm512_loadu_ps(x + i + 16);
-    //    __m512 x3 = _mm512_loadu_ps(x + i + 32);
-    //    __m512 x4 = _mm512_loadu_ps(x + i + 48);
-    //    _mm512_storeu_ps(x + i +  0, ggml_v_softcap(x1, vs_before, vs_after));
-    //    _mm512_storeu_ps(x + i + 16, ggml_v_softcap(x2, vs_before, vs_after));
-    //    _mm512_storeu_ps(x + i + 32, ggml_v_softcap(x3, vs_before, vs_after));
-    //    _mm512_storeu_ps(x + i + 48, ggml_v_softcap(x4, vs_before, vs_after));
-    //}
-    for (; i + 15 < n; i += 16) {
-        _mm512_storeu_ps(x + i, ggml_v_softcap(_mm512_loadu_ps(x + i), vs_before, vs_after));
-    }
-#elif defined(__AVX2__) && defined(__FMA__)
-    for (; i + 7 < n; i += 8) {
-        _mm256_storeu_ps(x + i, ggml_v_softcap(_mm256_loadu_ps(x + i), s_before, s_after));
-    }
-#elif defined(__SSE2__)
-    for (; i + 3 < n; i += 4) {
-        _mm_storeu_ps(x + i, ggml_v_softcap(_mm_loadu_ps(x + i), s_before, s_after));
-    }
-#elif defined(__ARM_NEON) && defined(__aarch64__)
-    float32x4_t vs_before = vdupq_n_f32(s_before);
-    float32x4_t vs_after  = vdupq_n_f32(s_after);
-    for (; i + 3 < n; i += 4) {
-        vst1q_f32(x + i, ggml_v_softcap(vld1q_f32(x + i), vs_before, vs_after));
-    }
-#endif
-    for (; i < n; ++i) {
-        x[i] = s_after*tanhf(x[i]*s_before);
-    }
-}
-
-inline static void ggml_vec_gelu_f16(const int n, ggml_fp16_t * y, const ggml_fp16_t * x) {
-    const uint16_t * i16 = (const uint16_t *) x;
-    for (int i = 0; i < n; ++i) {
-        y[i] = ggml_table_gelu_f16[i16[i]];
-    }
-}
-
-//
-// On my AVX512 (Ryzen-7950X) and AVX2 (Ryzen-5975WX) computing gelu directly
-// via SIMD instructions is faster than the fp16-based lookup table.
-// On my M2-Max CPU the lookup table is slightly faster than the SIMD version,
-// hence we use the SIMD version only if GGML_GELU_FP16 is not defined.
-// We do not run into numerical issues for large or small arguments because
-//    0.5f * (1 + tanhf(arg))
-// is computed as
-//    exp(2*arg)/(exp(2*arg) + 1)
-// The ggml_v_expf functions flushes to zero for large enough negative
-// arguments, so the above becomes zero. ggml_v_expf returns INFINITY
-// for large positive arguments, so we would get a NaN if we did nothing. But in the
-// ggml_v_gelu SIMD implementations we override the gelu result with the
-// input argument when the argument is greater than 10, so it is all good.
-//
-inline static void ggml_vec_gelu_f32(const int n, float * y, const float * x) {
-    int i = 0;
-#if defined(__AVX512F__) && defined(__AVX512DQ__)
-    __m512 c1 = _mm512_set1_ps(GELU_COEF_A);
-    __m512 c2 = _mm512_set1_ps(2.f*SQRT_2_OVER_PI);
-    for (; i + 15 < n; i += 16) {
-        _mm512_storeu_ps(y + i, ggml_v_gelu(_mm512_loadu_ps(x + i), c1, c2));
-    }
-#elif defined __AVX2__ && defined __FMA__
-    __m256 c1 = _mm256_set1_ps(GELU_COEF_A);
-    __m256 c2 = _mm256_set1_ps(2.f*SQRT_2_OVER_PI);
-    for (; i + 7 < n; i += 8) {
-        _mm256_storeu_ps(y + i, ggml_v_gelu(_mm256_loadu_ps(x + i), c1, c2));
-    }
-#endif
-#ifdef GGML_GELU_FP16
-    uint16_t t;
-    for (; i < n; ++i) {
-        if (x[i] <= -10.0f) {
-            y[i] = 0.0f;
-        } else if (x[i] >= 10.0f) {
-            y[i] = x[i];
-        } else {
-            ggml_fp16_t fp16 = GGML_FP32_TO_FP16(x[i]);
-            memcpy(&t, &fp16, sizeof(uint16_t));
-            y[i] = GGML_FP16_TO_FP32(ggml_table_gelu_f16[t]);
-        }
-    }
-#else
-#if defined __ARM_NEON
-    float32x4_t c1 = vdupq_n_f32(GELU_COEF_A);
-    float32x4_t c2 = vdupq_n_f32(2.f*SQRT_2_OVER_PI);
-    for (; i + 3 < n; i += 4) {
-        vst1q_f32(y + i, ggml_v_gelu(vld1q_f32(x + i), c1, c2));
-    }
-#endif
-    for (; i < n; ++i) {
-        y[i] = ggml_gelu_f32(x[i]);
-    }
-#endif
-}
-
 static ggml_float ggml_vec_soft_max_f32(const int n, float * y, const float * x, float max) {
     int i = 0;
     ggml_float sum = 0;
@@ -8598,71 +8420,6 @@ static void ggml_compute_forward_scale(
     }
 }
 
-// ggml_compute_forward_softcap
-
-static void ggml_compute_forward_softcap_f32(
-        const struct ggml_compute_params * params,
-        struct ggml_tensor * dst) {
-
-    const struct ggml_tensor * src0 = dst->src[0];
-
-    GGML_ASSERT(ggml_is_contiguous(src0));
-    GGML_ASSERT(ggml_is_contiguous(dst));
-    GGML_ASSERT(ggml_are_same_shape(src0, dst));
-
-    // scale factor
-    float val[2];
-    memcpy(val, dst->op_params, sizeof(val));
-
-    const int ith = params->ith;
-    const int nth = params->nth;
-
-    const int nc = src0->ne[0];
-    const int nr = ggml_nrows(src0);
-
-    // rows per thread
-    const int dr = (nr + nth - 1)/nth;
-
-    // row range for this thread
-    const int ir0 = dr*ith;
-    const int ir1 = MIN(ir0 + dr, nr);
-
-    const size_t nb01 = src0->nb[1];
-
-    const size_t nb1 = dst->nb[1];
-
-    for (int i1 = ir0; i1 < ir1; i1++) {
-        if (dst->data != src0->data) {
-            // src0 is same shape as dst => same indices
-            memcpy((char *)dst->data + i1*nb1, (char *)src0->data + i1*nb01, nc * sizeof(float));
-        }
-        // TODO: better implementation
-        float * row = (float *) ((char *) dst->data + i1*nb1);
-        ggml_vec_softcap_f32(nc, row, val[0], val[1]);
-        //ggml_vec_scale_f32(nc, row, val[0]);
-        //ggml_vec_tanh_f32(nc, row, row);
-        //ggml_vec_scale_f32(nc, row, val[1]);
-    }
-}
-
-static void ggml_compute_forward_softcap(
-        const struct ggml_compute_params * params,
-        struct ggml_tensor * dst) {
-
-    const struct ggml_tensor * src0 = dst->src[0];
-
-    switch (src0->type) {
-        case GGML_TYPE_F32:
-            {
-                ggml_compute_forward_softcap_f32(params, dst);
-            } break;
-        default:
-            {
-                GGML_ASSERT(false);
-            } break;
-    }
-}
-
 // ggml_compute_forward_set
 
 static void ggml_compute_forward_set_f32(
@@ -12961,10 +12718,6 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm
             {
                 ggml_compute_forward_scale(params, tensor);
             } break;
-        case GGML_OP_SOFTCAP:
-            {
-                ggml_compute_forward_softcap(params, tensor);
-            } break;
         case GGML_OP_SET:
             {
                 ggml_compute_forward_set(params, tensor);
@@ -13372,7 +13125,6 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) {
                 n_tasks = 1;
             } break;
         case GGML_OP_SCALE:
-        case GGML_OP_SOFTCAP:
         case GGML_OP_SET:
         case GGML_OP_RESHAPE:
         case GGML_OP_VIEW:

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -40,7 +40,6 @@ bool g_mul_mat_q = false;
 #include "ggml-cuda/unary.cuh"
 #include "ggml-cuda/upscale.cuh"
 #include "ggml-cuda/wkv6.cuh"
-#include "ggml-cuda/softcap.cuh"
 
 #include <algorithm>
 #include <array>
@@ -2313,9 +2312,6 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg
         case GGML_OP_SCALE:
             ggml_cuda_op_scale(ctx, dst);
             break;
-        case GGML_OP_SOFTCAP:
-            ggml_cuda_op_softcap(ctx, dst);
-            break;
         case GGML_OP_SQR:
             ggml_cuda_op_sqr(ctx, dst);
             break;
@@ -3208,7 +3204,6 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
         case GGML_OP_MUL:
         case GGML_OP_DIV:
         case GGML_OP_SCALE:
-        case GGML_OP_SOFTCAP:
         case GGML_OP_SQR:
         case GGML_OP_SQRT:
         case GGML_OP_SIN: