pass src tensor instead

Author: CISC

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

@@ -2848,15 +2848,8 @@ static void evaluate_and_capture_cuda_graph(ggml_backend_cuda_context * cuda_ctx
                     }
 
                     if (ggml_cuda_can_fuse(cgraph, i, { GGML_OP_SCALE, GGML_OP_UNARY, GGML_OP_SCALE }, { GGML_UNARY_OP_TANH })) {
-                        ggml_tensor * src0 = node->src[0];
-                        float scale = ggml_get_op_params_f32(node, 0);
-
-                        i += 2; node = cgraph->nodes[i];
-
-                        ggml_set_op_params_f32(node, 1, scale);
-                        node->src[0] = src0;
-
-                        ggml_cuda_op_softcap(*cuda_ctx, node);
+                        i += 2;
+                        ggml_cuda_op_softcap(*cuda_ctx, cgraph->nodes[i], node);
                         continue;
                     }
                 }

ggml/src/ggml-cuda/softcap.cu

@@ -15,8 +15,8 @@ static void softcap_f32_cuda(const float * x, float * dst, const float scale, co
     softcap_f32<<<num_blocks, CUDA_SOFTCAP_BLOCK_SIZE, 0, stream>>>(x, dst, scale, softcap, k);
 }
 
-void ggml_cuda_op_softcap(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
-    const ggml_tensor * src0 = dst->src[0];
+void ggml_cuda_op_softcap(ggml_backend_cuda_context & ctx, ggml_tensor * dst, ggml_tensor * src) {
+    const ggml_tensor * src0 = src->src[0];
     const float * src0_d = (const float *)src0->data;
     float * dst_d = (float *)dst->data;
     cudaStream_t stream = ctx.stream();
@@ -26,8 +26,8 @@ void ggml_cuda_op_softcap(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
 
     float scale;
     float softcap;
+    memcpy(&scale,   (float *) src->op_params + 0, sizeof(float));
     memcpy(&softcap, (float *) dst->op_params + 0, sizeof(float));
-    memcpy(&scale,   (float *) dst->op_params + 1, sizeof(float));
 
     softcap_f32_cuda(src0_d, dst_d, scale, softcap, ggml_nelements(src0), stream);
 }

ggml/src/ggml-cuda/softcap.cuh

@@ -2,4 +2,4 @@
 
 #define CUDA_SOFTCAP_BLOCK_SIZE 256
 
-void ggml_cuda_op_softcap(ggml_backend_cuda_context & ctx, ggml_tensor * dst);
+void ggml_cuda_op_softcap(ggml_backend_cuda_context & ctx, ggml_tensor * dst, ggml_tensor * src);