Revert "Reapply "IK Softcap op""

Author: Nexesenex

ggml/include/ggml.h

@@ -634,7 +634,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,
@@ -1464,19 +1463,6 @@ extern "C" {
         float                 s,
         float                 b);
 
-    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-alloc.c

@@ -45,7 +45,6 @@ static bool ggml_op_can_inplace(enum ggml_op op) {
         case GGML_OP_RMS_NORM_BACK:
         case GGML_OP_SOFT_MAX:
         case GGML_OP_SOFT_MAX_BACK:
-        case GGML_OP_SOFTCAP:	
             return true;
 
         default:

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

@@ -2542,10 +2542,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);
@@ -2998,7 +2994,6 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) {
             {
                 n_tasks = 1; //TODO
             } break;
-        case GGML_OP_SOFTCAP:
         case GGML_OP_SOFT_MAX:
             {
                 n_tasks = MIN(n_threads, ggml_nrows(node->src[0]));

ggml/src/ggml-cpu/ops.cpp

@@ -5486,79 +5486,6 @@ void ggml_compute_forward_scale(
     }
 }
 
-// ggml_compute_forward_softcap
-
-static void ggml_compute_forward_softcap_f32(
-        const ggml_compute_params * params,
-        ggml_tensor * dst) {
-
-    const 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];
-
-    //if (ith == 0) printf("%s: nc = %d, nr = %d, nth = %d, params = %g, %g,  %d\n", __func__, nc, nr, nth, val[0], val[1], dst->data == src0->data ? 1 : 0);
-
-    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)); */
-        float * dst_row = (float *) ((char *) dst->data + i1*nb1);
-        if (dst->data == src0->data) {
-            ggml_vec_softcap_f32(nc, dst_row, val[0], val[1]);
-        } else {
-            const float * src_row = (const float *)((const char *)src0->data + i1*nb01);
-            ggml_vec_cpy_softcap_f32(nc, src_row, dst_row, val[0], val[1]);
-        // 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]);
-        }
-    }
-}
-
-void ggml_compute_forward_softcap(
-        const ggml_compute_params * params,
-        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);
-            }
-    }
-}
-
 // ggml_compute_forward_set
 
 static void ggml_compute_forward_set_f32(

ggml/src/ggml-cpu/ops.h

@@ -66,7 +66,6 @@ void ggml_compute_forward_diag_mask_inf(const struct ggml_compute_params * param
 void ggml_compute_forward_diag_mask_zero(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_soft_max(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_soft_max_ext_back(const struct ggml_compute_params * params, struct ggml_tensor * dst);
-void ggml_compute_forward_softcap(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_rope(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_rope_back(const struct ggml_compute_params * params, struct ggml_tensor * dst);
 void ggml_compute_forward_clamp(const struct ggml_compute_params * params, struct ggml_tensor * dst);

ggml/src/ggml-cpu/vec.h

@@ -771,15 +771,6 @@ inline static float32x4_t ggml_v_tanh(float32x4_t x) {
     //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);
@@ -845,13 +836,6 @@ inline static __m512 ggml_v_tanh(__m512 x) {
     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);
@@ -927,14 +911,6 @@ inline static __m256 ggml_v_tanh(__m256 x) {
     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);
@@ -1005,13 +981,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__
 
 inline static void ggml_vec_silu_f16(const int n, ggml_fp16_t * y, const ggml_fp16_t * x) {
@@ -1140,72 +1109,6 @@ static void ggml_vec_tanh_f32(const int n, float * y, const float * x) {
     }
 }
 
-static void ggml_vec_cpy_softcap_f32(const int n, const float * x, float * y, 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 + 15 < n; i += 16) {
-        _mm512_storeu_ps(y + 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(y + 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(y + 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(y + i, ggml_v_softcap(vld1q_f32(x + i), vs_before, vs_after));
-    }
-#endif
-    for (; i < n; ++i) {
-        y[i] = s_after*tanhf(x[i]*s_before);
-    }
-}
-
-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);
-    }
-}
-
 //
 // On my AVX512 (Ryzen-7950X) and AVX2 (Ryzen-5975WX) computing gelu directly
 // via SIMD instructions is faster than the fp16-based lookup table.

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

@@ -50,8 +50,6 @@ bool g_mul_mat_q = true;
 #include "ggml-cuda/set-rows.cuh"
 #include "ggml.h"
 
-#include "ggml-cuda/softcap.cuh"
-
 // #include "ggml-cuda/iqk_mmvq.cuh"
 
 #include <algorithm>
@@ -2625,9 +2623,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;
@@ -3673,7 +3668,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: