ggml-opencl-chatterbox-ops.patch

diff --git a/include/ggml-opencl.h b/include/ggml-opencl.h
index 6b617713..3f8312cf 100644
--- a/include/ggml-opencl.h
+++ b/include/ggml-opencl.h
@@ -11,6 +11,7 @@ extern "C" {
 //
 // backend API
 //
+// May return NULL when no OpenCL device is available or the driver/context failed to init.
 GGML_BACKEND_API ggml_backend_t ggml_backend_opencl_init(void);
 GGML_BACKEND_API bool ggml_backend_is_opencl(ggml_backend_t backend);
 
diff --git a/src/ggml-opencl/CMakeLists.txt b/src/ggml-opencl/CMakeLists.txt
index 540942b1..d9f93004 100644
--- a/src/ggml-opencl/CMakeLists.txt
+++ b/src/ggml-opencl/CMakeLists.txt
@@ -144,15 +144,22 @@ set(GGML_OPENCL_KERNELS
     tsembd
     upscale
     tanh
+    sin
+    cos
+    abs
+    elu
     exp
     expm1
     softplus
+    mish
     pad
     repeat
     mul_mat_f16_f32
     mul_mm_f16_f32_kq_kqv
     conv2d
     conv2d_f16_f32
+    conv_transpose_1d_f32
+    conv_transpose_1d_f16_f32
     flash_attn_f32_f16
     flash_attn_f16
     flash_attn_f32
diff --git a/src/ggml-opencl/ggml-opencl.cpp b/src/ggml-opencl/ggml-opencl.cpp
index 6f3fc588..80aacdbf 100644
--- a/src/ggml-opencl/ggml-opencl.cpp
+++ b/src/ggml-opencl/ggml-opencl.cpp
@@ -28,6 +28,7 @@
 #include <memory>
 #include <charconv>
 #include <mutex>
+#include <cstdlib>
 
 #undef MIN
 #undef MAX
@@ -468,6 +469,8 @@ struct ggml_backend_opencl_context {
     cl_program program_conv_2d_f16;
     cl_program program_conv_2d_f32;
     cl_program program_conv_2d_f16_f32;
+    cl_program program_conv_transpose_1d_f32;
+    cl_program program_conv_transpose_1d_f16_f32;
     cl_program program_tsembd;
     cl_program program_gemv_moe_mxfp4_f32, program_gemm_moe_mxfp4_f32;
     cl_program program_mul_mv_id_q4_0_f32_8x_flat;
@@ -478,7 +481,7 @@ struct ggml_backend_opencl_context {
     cl_program program_mul_mm_f16_f32_l4_lm;
     cl_program program_mul_mm_q8_0_f32_l4_lm;
 
-    cl_kernel kernel_add, kernel_add_row, kernel_add_f16, kernel_add_row_f16;
+    cl_kernel kernel_add, kernel_add_row, kernel_add_row_residual, kernel_add_f16, kernel_add_row_f16;
     cl_kernel kernel_mul, kernel_mul_row, kernel_mul_f16, kernel_mul_row_f16;
     cl_kernel kernel_div, kernel_div_row, kernel_div_f16, kernel_div_row_f16;
     cl_kernel kernel_sub, kernel_sub_row, kernel_sub_f16, kernel_sub_row_f16;
@@ -492,6 +495,7 @@ struct ggml_backend_opencl_context {
     cl_kernel kernel_gelu_erf, kernel_gelu_erf_4;
     cl_kernel kernel_gelu_quick, kernel_gelu_quick_4;
     cl_kernel kernel_relu;
+    cl_kernel kernel_leaky_relu;
     cl_kernel kernel_sigmoid_f32, kernel_sigmoid_f16;
     cl_kernel kernel_tri;
     cl_kernel kernel_fill;
@@ -560,6 +564,12 @@ struct ggml_backend_opencl_context {
     cl_kernel kernel_pad;
     cl_kernel kernel_tanh_f32, kernel_tanh_f32_4, kernel_tanh_f32_nc;
     cl_kernel kernel_tanh_f16, kernel_tanh_f16_4, kernel_tanh_f16_nc;
+    cl_kernel kernel_sin_f32, kernel_sin_f32_4, kernel_sin_f32_nc;
+    cl_kernel kernel_sin_f16, kernel_sin_f16_4, kernel_sin_f16_nc;
+    cl_kernel kernel_cos_f32, kernel_cos_f32_4, kernel_cos_f32_nc;
+    cl_kernel kernel_cos_f16, kernel_cos_f16_4, kernel_cos_f16_nc;
+    cl_kernel kernel_abs_f32, kernel_abs_f32_4, kernel_abs_f32_nc;
+    cl_kernel kernel_elu_f32, kernel_elu_f32_4, kernel_elu_f32_nc;
     cl_kernel kernel_neg_f32, kernel_neg_f32_4, kernel_neg_f32_nc;
     cl_kernel kernel_neg_f16, kernel_neg_f16_4, kernel_neg_f16_nc;
     cl_kernel kernel_exp_f32, kernel_exp_f32_4, kernel_exp_f32_nc;
@@ -568,12 +578,15 @@ struct ggml_backend_opencl_context {
     cl_kernel kernel_expm1_f16, kernel_expm1_f16_4, kernel_expm1_f16_nc;
     cl_kernel kernel_softplus_f32, kernel_softplus_f32_4, kernel_softplus_f32_nc;
     cl_kernel kernel_softplus_f16, kernel_softplus_f16_4, kernel_softplus_f16_nc;
+    cl_kernel kernel_mish_f32, kernel_mish_f32_4;
     cl_kernel kernel_upscale;
     cl_kernel kernel_upscale_bilinear;
     cl_kernel kernel_concat_f32;
     cl_kernel kernel_conv_2d_f16;
     cl_kernel kernel_conv_2d_f32;
     cl_kernel kernel_conv_2d_f16_f32;
+    cl_kernel kernel_conv_transpose_1d_f32;
+    cl_kernel kernel_conv_transpose_1d_f16_f32;
     cl_kernel kernel_ssm_conv_f32_f32, kernel_ssm_conv_f32_f32_4;
     cl_kernel kernel_timestep_embedding;
     cl_kernel kernel_gemv_moe_mxfp4_f32, kernel_gemm_moe_mxfp4_f32;
@@ -591,10 +604,10 @@ struct ggml_backend_opencl_context {
 
     std::vector<ProfilingInfo> profiling_info;
 
-    void write_profiling_info() {
-        FILE * fperf = fopen("cl_profiling.csv", "w");
+    void write_profiling_info(const char * csv_path = "cl_profiling.csv", const char * trace_path = "cl_trace.json") {
+        FILE * fperf = fopen(csv_path, "w");
         if (!fperf) {
-            GGML_LOG_ERROR("Failed to open cl_profiling.csv\n");
+            GGML_LOG_ERROR("Failed to open %s\n", csv_path);
             return;
         }
 
@@ -649,9 +662,9 @@ struct ggml_backend_opencl_context {
         fclose(fperf);
 
         // Dump a simple chrome trace
-        FILE* ftrace = fopen("cl_trace.json", "w");
+        FILE* ftrace = fopen(trace_path, "w");
         if (!ftrace) {
-            GGML_LOG_ERROR("Failed to open cl_trace.json\n");
+            GGML_LOG_ERROR("Failed to open %s\n", trace_path);
             return;
         }
 
@@ -814,6 +827,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve
 
         CL_CHECK((backend_ctx->kernel_add         = clCreateKernel(backend_ctx->program_add, "kernel_add", &err), err));
         CL_CHECK((backend_ctx->kernel_add_row     = clCreateKernel(backend_ctx->program_add, "kernel_add_row", &err), err));
+        CL_CHECK((backend_ctx->kernel_add_row_residual = clCreateKernel(backend_ctx->program_add, "kernel_add_row_residual", &err), err));
         CL_CHECK((backend_ctx->kernel_add_f16     = clCreateKernel(backend_ctx->program_add, "kernel_add_f16", &err), err));
         CL_CHECK((backend_ctx->kernel_add_row_f16 = clCreateKernel(backend_ctx->program_add, "kernel_add_row_f16", &err), err));
         GGML_LOG_CONT(".");
@@ -1623,7 +1637,8 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve
         backend_ctx->program_relu =
             build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
 
-        CL_CHECK((backend_ctx->kernel_relu = clCreateKernel(backend_ctx->program_relu, "kernel_relu", &err), err));
+        CL_CHECK((backend_ctx->kernel_relu       = clCreateKernel(backend_ctx->program_relu, "kernel_relu",       &err), err));
+        CL_CHECK((backend_ctx->kernel_leaky_relu = clCreateKernel(backend_ctx->program_relu, "kernel_leaky_relu", &err), err));
         GGML_LOG_CONT(".");
     }
 
@@ -2095,6 +2110,84 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve
         GGML_LOG_CONT(".");
     }
 
+    // sin (GGML_OP_SIN)
+    {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+        const std::string kernel_src {
+            #include "sin.cl.h"
+        };
+#else
+        const std::string kernel_src = read_file("sin.cl");
+#endif
+        cl_program prog =
+            build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
+        CL_CHECK((backend_ctx->kernel_sin_f32    = clCreateKernel(prog, "kernel_sin_f32", &err), err));
+        CL_CHECK((backend_ctx->kernel_sin_f32_4  = clCreateKernel(prog, "kernel_sin_f32_4", &err), err));
+        CL_CHECK((backend_ctx->kernel_sin_f32_nc = clCreateKernel(prog, "kernel_sin_f32_nc", &err), err));
+        CL_CHECK((backend_ctx->kernel_sin_f16    = clCreateKernel(prog, "kernel_sin_f16", &err), err));
+        CL_CHECK((backend_ctx->kernel_sin_f16_4  = clCreateKernel(prog, "kernel_sin_f16_4", &err), err));
+        CL_CHECK((backend_ctx->kernel_sin_f16_nc = clCreateKernel(prog, "kernel_sin_f16_nc", &err), err));
+        CL_CHECK(clReleaseProgram(prog));
+        GGML_LOG_CONT(".");
+    }
+
+    // cos (GGML_OP_COS)
+    {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+        const std::string kernel_src {
+            #include "cos.cl.h"
+        };
+#else
+        const std::string kernel_src = read_file("cos.cl");
+#endif
+        cl_program prog =
+            build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
+        CL_CHECK((backend_ctx->kernel_cos_f32    = clCreateKernel(prog, "kernel_cos_f32", &err), err));
+        CL_CHECK((backend_ctx->kernel_cos_f32_4  = clCreateKernel(prog, "kernel_cos_f32_4", &err), err));
+        CL_CHECK((backend_ctx->kernel_cos_f32_nc = clCreateKernel(prog, "kernel_cos_f32_nc", &err), err));
+        CL_CHECK((backend_ctx->kernel_cos_f16    = clCreateKernel(prog, "kernel_cos_f16", &err), err));
+        CL_CHECK((backend_ctx->kernel_cos_f16_4  = clCreateKernel(prog, "kernel_cos_f16_4", &err), err));
+        CL_CHECK((backend_ctx->kernel_cos_f16_nc = clCreateKernel(prog, "kernel_cos_f16_nc", &err), err));
+        CL_CHECK(clReleaseProgram(prog));
+        GGML_LOG_CONT(".");
+    }
+
+    // abs
+    {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+        const std::string kernel_src {
+            #include "abs.cl.h"
+        };
+#else
+        const std::string kernel_src = read_file("abs.cl");
+#endif
+        cl_program prog =
+            build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
+        CL_CHECK((backend_ctx->kernel_abs_f32    = clCreateKernel(prog, "kernel_abs_f32", &err), err));
+        CL_CHECK((backend_ctx->kernel_abs_f32_4  = clCreateKernel(prog, "kernel_abs_f32_4", &err), err));
+        CL_CHECK((backend_ctx->kernel_abs_f32_nc = clCreateKernel(prog, "kernel_abs_f32_nc", &err), err));
+        CL_CHECK(clReleaseProgram(prog));
+        GGML_LOG_CONT(".");
+    }
+
+    // elu
+    {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+        const std::string kernel_src {
+            #include "elu.cl.h"
+        };
+#else
+        const std::string kernel_src = read_file("elu.cl");
+#endif
+        cl_program prog =
+            build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
+        CL_CHECK((backend_ctx->kernel_elu_f32    = clCreateKernel(prog, "kernel_elu_f32", &err), err));
+        CL_CHECK((backend_ctx->kernel_elu_f32_4  = clCreateKernel(prog, "kernel_elu_f32_4", &err), err));
+        CL_CHECK((backend_ctx->kernel_elu_f32_nc = clCreateKernel(prog, "kernel_elu_f32_nc", &err), err));
+        CL_CHECK(clReleaseProgram(prog));
+        GGML_LOG_CONT(".");
+    }
+
     // neg
     {
 #ifdef GGML_OPENCL_EMBED_KERNELS
@@ -2116,6 +2209,23 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve
         GGML_LOG_CONT(".");
     }
 
+    // mish fused pattern (softplus -> tanh -> mul)
+    {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+        const std::string kernel_src {
+            #include "mish.cl.h"
+        };
+#else
+        const std::string kernel_src = read_file("mish.cl");
+#endif
+        cl_program prog =
+            build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
+        CL_CHECK((backend_ctx->kernel_mish_f32   = clCreateKernel(prog, "kernel_mish_f32", &err), err));
+        CL_CHECK((backend_ctx->kernel_mish_f32_4 = clCreateKernel(prog, "kernel_mish_f32_4", &err), err));
+        CL_CHECK(clReleaseProgram(prog));
+        GGML_LOG_CONT(".");
+    }
+
     // exp
     {
 #ifdef GGML_OPENCL_EMBED_KERNELS
@@ -2309,6 +2419,45 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve
                 }
     }
 
+    // conv_transpose_1d
+    {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+        const std::string kernel_src_f32 {
+            #include "conv_transpose_1d_f32.cl.h"
+        };
+        const std::string kernel_src_f16 {
+            #include "conv_transpose_1d_f16_f32.cl.h"
+        };
+#else
+        const std::string kernel_src_f32  = read_file("conv_transpose_1d_f32.cl");
+        const std::string kernel_src_f16 = read_file("conv_transpose_1d_f16_f32.cl");
+#endif
+        if (!kernel_src_f32.empty()) {
+            backend_ctx->program_conv_transpose_1d_f32 =
+                build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src_f32.c_str(), compile_opts);
+            CL_CHECK((backend_ctx->kernel_conv_transpose_1d_f32 = clCreateKernel(
+                          backend_ctx->program_conv_transpose_1d_f32, "kernel_conv_transpose_1d_f32", &err),
+                err));
+            GGML_LOG_CONT(".");
+        } else {
+            GGML_LOG_WARN("ggml_opencl: conv_transpose_1d_f32 kernel source not found. CONV_TRANSPOSE_1D (f32) will be unavailable.\n");
+            backend_ctx->program_conv_transpose_1d_f32 = nullptr;
+            backend_ctx->kernel_conv_transpose_1d_f32  = nullptr;
+        }
+        if (!kernel_src_f16.empty()) {
+            backend_ctx->program_conv_transpose_1d_f16_f32 = build_program_from_source(
+                backend_ctx->context, backend_ctx->device, kernel_src_f16.c_str(), compile_opts);
+            CL_CHECK((backend_ctx->kernel_conv_transpose_1d_f16_f32 = clCreateKernel(
+                          backend_ctx->program_conv_transpose_1d_f16_f32, "kernel_conv_transpose_1d_f16_f32", &err),
+                err));
+            GGML_LOG_CONT(".");
+        } else {
+            GGML_LOG_WARN("ggml_opencl: conv_transpose_1d_f16_f32 kernel source not found. CONV_TRANSPOSE_1D (f16) will be unavailable.\n");
+            backend_ctx->program_conv_transpose_1d_f16_f32 = nullptr;
+            backend_ctx->kernel_conv_transpose_1d_f16_f32  = nullptr;
+        }
+    }
+
     // ssm_conv
     {
 #ifdef GGML_OPENCL_EMBED_KERNELS
@@ -3716,6 +3865,89 @@ static bool ggml_opencl_can_fuse(const struct ggml_cgraph * cgraph, int node_idx
         if (!ggml_is_contiguous(gn->src[0]) || !ggml_is_contiguous(w) || !ggml_is_contiguous(b)) {
             return false;
         }
+    } else if (ops.size() == 2 && ops.begin()[0] == GGML_OP_ADD && ops.begin()[1] == GGML_OP_ADD) {
+        const ggml_tensor * add0 = cgraph->nodes[node_idx];
+        const ggml_tensor * add1 = cgraph->nodes[node_idx+1];
+
+        if (add0->type != GGML_TYPE_F32 || add1->type != GGML_TYPE_F32) {
+            return false;
+        }
+
+        const ggml_tensor * a = add0->src[0];
+        const ggml_tensor * b = add0->src[1];
+        const ggml_tensor * row = nullptr;
+        const ggml_tensor * full = nullptr;
+        if (ggml_nelements(a) == a->ne[0] && ggml_is_contiguous(a)) {
+            row = a;
+            full = b;
+        } else if (ggml_nelements(b) == b->ne[0] && ggml_is_contiguous(b)) {
+            row = b;
+            full = a;
+        } else {
+            return false;
+        }
+
+        const ggml_tensor * residual = add1->src[0] == add0 ? add1->src[1] : add1->src[0];
+
+        if (full->type != GGML_TYPE_F32 || row->type != GGML_TYPE_F32 || residual->type != GGML_TYPE_F32) {
+            return false;
+        }
+        if (!ggml_are_same_shape(full, add0) || !ggml_are_same_shape(residual, add0)) {
+            return false;
+        }
+        if (!ggml_is_contiguous(full) || !ggml_is_contiguous(residual) || !ggml_is_contiguous(add1)) {
+            return false;
+        }
+        if (add0->ne[0] % 4 != 0 || row->ne[0] != add0->ne[0]) {
+            return false;
+        }
+    } else if (ops.size() == 3 && ops.begin()[0] == GGML_OP_UNARY && ops.begin()[1] == GGML_OP_UNARY && ops.begin()[2] == GGML_OP_MUL) {
+        const ggml_tensor * softplus = cgraph->nodes[node_idx];
+        const ggml_tensor * tanh     = cgraph->nodes[node_idx+1];
+        const ggml_tensor * mul      = cgraph->nodes[node_idx+2];
+        const ggml_tensor * x        = softplus->src[0];
+
+        if (ggml_get_unary_op(softplus) != GGML_UNARY_OP_SOFTPLUS ||
+            ggml_get_unary_op(tanh)     != GGML_UNARY_OP_TANH) {
+            return false;
+        }
+        if (tanh->src[0] != softplus) {
+            return false;
+        }
+        if (mul->src[0] != x && mul->src[1] != x) {
+            return false;
+        }
+        if (x->type != GGML_TYPE_F32 || mul->type != GGML_TYPE_F32) {
+            return false;
+        }
+        if (!ggml_is_contiguous(x) || !ggml_is_contiguous(mul)) {
+            return false;
+        }
+    } else if ((ops.size() == 2 && ops.begin()[0] == GGML_OP_MUL_MAT && ops.begin()[1] == GGML_OP_ADD) ||
+               (ops.size() == 3 && ops.begin()[0] == GGML_OP_MUL_MAT && ops.begin()[1] == GGML_OP_ADD && ops.begin()[2] == GGML_OP_ADD)) {
+        const ggml_tensor * mul = cgraph->nodes[node_idx];
+        const ggml_tensor * add0 = cgraph->nodes[node_idx + 1];
+        const ggml_tensor * bias = add0->src[0] == mul ? add0->src[1] : add0->src[0];
+
+        if (mul->src[0]->type != GGML_TYPE_Q4_0 || mul->src[1]->type != GGML_TYPE_F32 ||
+            mul->type != GGML_TYPE_F32 || add0->type != GGML_TYPE_F32) {
+            return false;
+        }
+        if (mul->src[1]->ne[1] != 1 || mul->src[1]->ne[2] != 1 || mul->src[1]->ne[3] != 1) {
+            return false;
+        }
+        const bool bias_is_row = ggml_nelements(bias) == bias->ne[0] && bias->ne[0] == mul->ne[0];
+        if (bias->type != GGML_TYPE_F32 || !(bias_is_row || ggml_are_same_shape(bias, mul)) || !ggml_is_contiguous(bias)) {
+            return false;
+        }
+        if (ops.size() == 3) {
+            const ggml_tensor * add1 = cgraph->nodes[node_idx + 2];
+            const ggml_tensor * residual = add1->src[0] == add0 ? add1->src[1] : add1->src[0];
+            if (add1->type != GGML_TYPE_F32 || residual->type != GGML_TYPE_F32 ||
+                !ggml_are_same_shape(residual, mul) || !ggml_is_contiguous(residual)) {
+                return false;
+            }
+        }
     }
 
     return true;
@@ -3724,6 +3956,9 @@ static bool ggml_opencl_can_fuse(const struct ggml_cgraph * cgraph, int node_idx
 static void ggml_opencl_op_rms_norm_fused(ggml_backend_t backend, ggml_tensor * rms_norm_tensor, ggml_tensor * mul_tensor);
 static void ggml_opencl_op_norm_fused(ggml_backend_t backend, ggml_tensor * norm_tensor, ggml_tensor * mul_tensor, ggml_tensor * add_tensor);
 static void ggml_opencl_op_group_norm_fused(ggml_backend_t backend, ggml_tensor * gn_tensor, ggml_tensor * mul_tensor, ggml_tensor * add_tensor);
+static void ggml_opencl_op_add_row_residual_fused(ggml_backend_t backend, ggml_tensor * add0, ggml_tensor * add1);
+static void ggml_opencl_op_mish_fused(ggml_backend_t backend, ggml_tensor * softplus, ggml_tensor * tanh, ggml_tensor * mul);
+static void ggml_opencl_op_mul_mat_add_fused(ggml_backend_t backend, ggml_tensor * mul, ggml_tensor * add0, ggml_tensor * add1);
 
 static ggml_status ggml_backend_opencl_graph_compute(ggml_backend_t backend, ggml_cgraph * cgraph) {
     ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context;
@@ -3759,6 +3994,26 @@ static ggml_status ggml_backend_opencl_graph_compute(ggml_backend_t backend, ggm
             i++;
             continue;
         }
+        if (!backend_ctx->disable_fusion && backend_ctx->gpu_family == ADRENO && ggml_opencl_can_fuse(cgraph, i, { GGML_OP_MUL_MAT, GGML_OP_ADD, GGML_OP_ADD })) {
+            ggml_opencl_op_mul_mat_add_fused(backend, node, cgraph->nodes[i+1], cgraph->nodes[i+2]);
+            i += 2;
+            continue;
+        }
+        if (!backend_ctx->disable_fusion && backend_ctx->gpu_family == ADRENO && ggml_opencl_can_fuse(cgraph, i, { GGML_OP_MUL_MAT, GGML_OP_ADD })) {
+            ggml_opencl_op_mul_mat_add_fused(backend, node, cgraph->nodes[i+1], nullptr);
+            i++;
+            continue;
+        }
+        if (!backend_ctx->disable_fusion && ggml_opencl_can_fuse(cgraph, i, { GGML_OP_ADD, GGML_OP_ADD })) {
+            ggml_opencl_op_add_row_residual_fused(backend, node, cgraph->nodes[i+1]);
+            i++;
+            continue;
+        }
+        if (!backend_ctx->disable_fusion && ggml_opencl_can_fuse(cgraph, i, { GGML_OP_UNARY, GGML_OP_UNARY, GGML_OP_MUL })) {
+            ggml_opencl_op_mish_fused(backend, node, cgraph->nodes[i+1], cgraph->nodes[i+2]);
+            i += 2;
+            continue;
+        }
 
         bool ok = ggml_cl_compute_forward(backend, node);
         if (!ok) {
@@ -3767,6 +4022,19 @@ static ggml_status ggml_backend_opencl_graph_compute(ggml_backend_t backend, ggm
         GGML_ASSERT(ok);
     }
 
+#ifdef GGML_OPENCL_PROFILING
+    if (std::getenv("GGML_OPENCL_PROFILE_EACH_GRAPH")) {
+        static int graph_profile_idx = 0;
+        char csv_path[64];
+        char trace_path[64];
+        snprintf(csv_path, sizeof(csv_path), "cl_profiling_%04d.csv", graph_profile_idx);
+        snprintf(trace_path, sizeof(trace_path), "cl_trace_%04d.json", graph_profile_idx);
+        backend_ctx->write_profiling_info(csv_path, trace_path);
+        backend_ctx->profiling_info.clear();
+        graph_profile_idx++;
+    }
+#endif
+
     return GGML_STATUS_SUCCESS;
 }
 
@@ -3865,8 +4133,15 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
         case GGML_OP_SQRT:
             return (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16) &&
                     ggml_is_contiguous(op->src[0]);
+        case GGML_OP_SIN:
+        case GGML_OP_COS:
+            return (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16) &&
+                   (op->type == op->src[0]->type);
         case GGML_OP_UNARY:
             switch (ggml_get_unary_op(op)) {
+                case GGML_UNARY_OP_ABS:
+                case GGML_UNARY_OP_ELU:
+                   return op->src[0]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32;
                 case GGML_UNARY_OP_GELU:
                 case GGML_UNARY_OP_SILU:
                 case GGML_UNARY_OP_RELU:
@@ -3902,6 +4177,9 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
             return op->type == GGML_TYPE_F32 && ggml_is_contiguous(op);
         case GGML_OP_FILL:
             return op->type == GGML_TYPE_F32 && ggml_is_contiguous(op);
+        case GGML_OP_LEAKY_RELU:
+            return op->src[0]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 &&
+                   ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op);
         case GGML_OP_CLAMP:
             return op->src[0]->type == GGML_TYPE_F32;
         case GGML_OP_SOFT_MAX:
@@ -3929,6 +4207,20 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
             return (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F16 && op->type == GGML_TYPE_F16) ||
                    (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32) ||
                    (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32);
+        case GGML_OP_CONV_TRANSPOSE_1D:
+            if (ggml_get_op_params_i32(op, 1) != 0 || ggml_get_op_params_i32(op, 2) != 1) {
+                return false;
+            }
+            if (op->type != GGML_TYPE_F32 || op->src[1]->type != GGML_TYPE_F32) {
+                return false;
+            }
+            if (op->src[0]->type == GGML_TYPE_F32) {
+                return backend_ctx->kernel_conv_transpose_1d_f32 != nullptr;
+            }
+            if (op->src[0]->type == GGML_TYPE_F16) {
+                return backend_ctx->kernel_conv_transpose_1d_f16_f32 != nullptr;
+            }
+            return false;
         case GGML_OP_SSM_CONV:
             return (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32);
         case GGML_OP_CONCAT:
@@ -4075,8 +4367,15 @@ static ggml_backend_i ggml_backend_opencl_i = {
 };
 
 ggml_backend_t ggml_backend_opencl_init(void) {
-    ggml_backend_dev_t dev = ggml_backend_reg_dev_get(ggml_backend_opencl_reg(), 0);
+    ggml_backend_reg_t reg = ggml_backend_opencl_reg();
+    if (ggml_backend_reg_dev_count(reg) == 0) {
+        return nullptr;
+    }
+    ggml_backend_dev_t dev = ggml_backend_reg_dev_get(reg, 0);
     ggml_backend_opencl_context *backend_ctx = ggml_cl2_init(dev);
+    if (!backend_ctx) {
+        return nullptr;
+    }
 
     ggml_backend_t backend = new ggml_backend {
         /* .guid    = */ ggml_backend_opencl_guid(),
@@ -5987,8 +6286,8 @@ static ggml_backend_dev_t ggml_backend_opencl_reg_device_get(ggml_backend_reg_t
 
 static struct ggml_backend_reg_i ggml_backend_opencl_reg_i = {
     /* .get_name         = */ ggml_backend_opencl_reg_get_name,
-    /* .device_count     = */ ggml_backend_opencl_reg_device_count,
-    /* .device_get       = */ ggml_backend_opencl_reg_device_get,
+    /* .get_device_count = */ ggml_backend_opencl_reg_device_count,
+    /* .get_device       = */ ggml_backend_opencl_reg_device_get,
     /* .get_proc_address = */ NULL,
 };
 
@@ -6631,6 +6930,90 @@ static void ggml_cl_add(ggml_backend_t backend, const ggml_tensor * src0, const
     }
 }
 
+static void ggml_opencl_op_add_row_residual_fused(ggml_backend_t backend, ggml_tensor * add0, ggml_tensor * add1) {
+    ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *) backend->context;
+
+    const ggml_tensor * a = add0->src[0];
+    const ggml_tensor * b = add0->src[1];
+    const ggml_tensor * row = nullptr;
+    const ggml_tensor * full = nullptr;
+
+    if (ggml_nelements(a) == a->ne[0]) {
+        row = a;
+        full = b;
+    } else {
+        row = b;
+        full = a;
+    }
+
+    const ggml_tensor * residual = add1->src[0] == add0 ? add1->src[1] : add1->src[0];
+
+    ggml_tensor_extra_cl * extraf = (ggml_tensor_extra_cl *) full->extra;
+    ggml_tensor_extra_cl * extrar = (ggml_tensor_extra_cl *) row->extra;
+    ggml_tensor_extra_cl * extrares = (ggml_tensor_extra_cl *) residual->extra;
+    ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *) add1->extra;
+
+    cl_ulong offsetf = extraf->offset + full->view_offs;
+    cl_ulong offsetr = extrar->offset + row->view_offs;
+    cl_ulong offsetres = extrares->offset + residual->view_offs;
+    cl_ulong offsetd = extrad->offset + add1->view_offs;
+
+    cl_kernel kernel = backend_ctx->kernel_add_row_residual;
+    const int ne = add1->ne[0] / 4;
+
+    CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem),   &extraf->data_device));
+    CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offsetf));
+    CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem),   &extrar->data_device));
+    CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offsetr));
+    CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem),   &extrares->data_device));
+    CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offsetres));
+    CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem),   &extrad->data_device));
+    CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_ulong), &offsetd));
+    CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int),      &ne));
+
+    int n = ggml_nelements(add1) / 4;
+    size_t global_work_size[] = {(size_t)n, 1, 1};
+    size_t local_work_size[] = {64, 1, 1};
+
+    backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, add1);
+}
+
+static void ggml_opencl_op_mish_fused(ggml_backend_t backend, ggml_tensor * softplus, ggml_tensor * tanh, ggml_tensor * mul) {
+    UNUSED(tanh);
+
+    ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *) backend->context;
+
+    const ggml_tensor * src0 = softplus->src[0];
+    ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *) src0->extra;
+    ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *) mul->extra;
+
+    cl_ulong offset0 = extra0->offset + src0->view_offs;
+    cl_ulong offsetd = extrad->offset + mul->view_offs;
+
+    int n = ggml_nelements(mul);
+    cl_kernel kernel = nullptr;
+    if (n % 4 == 0) {
+        kernel = backend_ctx->kernel_mish_f32_4;
+        n /= 4;
+    } else {
+        kernel = backend_ctx->kernel_mish_f32;
+    }
+
+    CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem),   &extra0->data_device));
+    CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0));
+    CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem),   &extrad->data_device));
+    CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offsetd));
+
+    size_t global_work_size[] = {(size_t)n, 1, 1};
+    size_t local_work_size[]  = {64, 1, 1};
+    size_t * local_work_size_ptr = local_work_size;
+    if (n % 64 != 0 && !backend_ctx->non_uniform_workgroups) {
+        local_work_size_ptr = nullptr;
+    }
+
+    backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size_ptr, mul);
+}
+
 static void ggml_cl_add_id(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
     GGML_ASSERT(src0);
     GGML_ASSERT(src0->extra);
@@ -7495,6 +7878,44 @@ static void ggml_cl_relu(ggml_backend_t backend, const ggml_tensor * src0, const
     backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size_ptr, dst);
 }
 
+static void ggml_cl_leaky_relu(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    GGML_ASSERT(src0);
+    GGML_ASSERT(src0->extra);
+    GGML_ASSERT(dst);
+    GGML_ASSERT(dst->extra);
+
+    UNUSED(src1);
+
+    ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context;
+
+    ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *)src0->extra;
+    ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra;
+
+    cl_ulong offset0 = extra0->offset + src0->view_offs;
+    cl_ulong offsetd = extrad->offset + dst->view_offs;
+
+    cl_kernel kernel = backend_ctx->kernel_leaky_relu;
+    const float negative_slope = ggml_get_op_params_f32(dst, 0);
+
+    CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem),   &extra0->data_device));
+    CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0));
+    CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem),   &extrad->data_device));
+    CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offsetd));
+    CL_CHECK(clSetKernelArg(kernel, 4, sizeof(float),    &negative_slope));
+
+    const int64_t n = ggml_nelements(dst);
+
+    size_t global_work_size[] = {(size_t)n, 1, 1};
+    size_t local_work_size[] = {64, 1, 1};
+
+    size_t * local_work_size_ptr = local_work_size;
+    if (n % 64 != 0 && !backend_ctx->non_uniform_workgroups) {
+        local_work_size_ptr = nullptr;  // Let driver choose the work-group sizes.
+    }
+
+    backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size_ptr, dst);
+}
+
 static void ggml_cl_sigmoid(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
     GGML_ASSERT(src0);
     GGML_ASSERT(src0->extra);
@@ -8225,6 +8646,295 @@ static void ggml_cl_tanh(ggml_backend_t backend, const ggml_tensor * src0, const
     }
 }
 
+static void ggml_cl_sin(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    GGML_ASSERT(src0);
+    GGML_ASSERT(src0->extra);
+    GGML_ASSERT(dst);
+    GGML_ASSERT(dst->extra);
+
+    UNUSED(src1);
+
+    ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *) backend->context;
+
+    ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *) src0->extra;
+    ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *) dst->extra;
+
+    cl_ulong offset0 = extra0->offset + src0->view_offs;
+    cl_ulong offsetd = extrad->offset + dst->view_offs;
+
+    const int ne00 = src0->ne[0];
+    const int ne01 = src0->ne[1];
+    const int ne02 = src0->ne[2];
+    const int ne03 = src0->ne[3];
+
+    const cl_ulong nb00 = src0->nb[0];
+    const cl_ulong nb01 = src0->nb[1];
+    const cl_ulong nb02 = src0->nb[2];
+    const cl_ulong nb03 = src0->nb[3];
+
+    const cl_ulong nb0 = dst->nb[0];
+    const cl_ulong nb1 = dst->nb[1];
+    const cl_ulong nb2 = dst->nb[2];
+    const cl_ulong nb3 = dst->nb[3];
+
+    cl_kernel kernel;
+
+    if (ggml_is_contiguous(src0)) {
+        int n = ggml_nelements(dst);
+        if (n % 4 == 0) {
+            if (src0->type == GGML_TYPE_F32) {
+                kernel = backend_ctx->kernel_sin_f32_4;
+            } else {
+                kernel = backend_ctx->kernel_sin_f16_4;
+            }
+            n /= 4;
+        } else {
+            if (src0->type == GGML_TYPE_F32) {
+                kernel = backend_ctx->kernel_sin_f32;
+            } else {
+                kernel = backend_ctx->kernel_sin_f16;
+            }
+        }
+
+        CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem),   &extra0->data_device));
+        CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0));
+        CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem),   &extrad->data_device));
+        CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offsetd));
+
+        size_t global_work_size[] = {(size_t) n, 1, 1};
+        size_t local_work_size[]  = {64, 1, 1};
+
+        size_t * local_work_size_ptr = local_work_size;
+        if (n % 64 != 0 && !backend_ctx->non_uniform_workgroups) {
+            local_work_size_ptr = nullptr;
+        }
+
+        backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size_ptr, dst);
+    } else {
+        if (src0->type == GGML_TYPE_F32) {
+            kernel = backend_ctx->kernel_sin_f32_nc;
+        } else {
+            kernel = backend_ctx->kernel_sin_f16_nc;
+        }
+
+        CL_CHECK(clSetKernelArg(kernel,  0, sizeof(cl_mem),   &extra0->data_device));
+        CL_CHECK(clSetKernelArg(kernel,  1, sizeof(cl_ulong), &offset0));
+        CL_CHECK(clSetKernelArg(kernel,  2, sizeof(cl_mem),   &extrad->data_device));
+        CL_CHECK(clSetKernelArg(kernel,  3, sizeof(cl_ulong), &offsetd));
+        CL_CHECK(clSetKernelArg(kernel,  4, sizeof(int),      &ne00));
+        CL_CHECK(clSetKernelArg(kernel,  5, sizeof(cl_ulong), &nb00));
+        CL_CHECK(clSetKernelArg(kernel,  6, sizeof(cl_ulong), &nb01));
+        CL_CHECK(clSetKernelArg(kernel,  7, sizeof(cl_ulong), &nb02));
+        CL_CHECK(clSetKernelArg(kernel,  8, sizeof(cl_ulong), &nb03));
+        CL_CHECK(clSetKernelArg(kernel,  9, sizeof(cl_ulong), &nb0));
+        CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_ulong), &nb1));
+        CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_ulong), &nb2));
+        CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_ulong), &nb3));
+
+        int nth = 64;
+
+        size_t global_work_size[] = {(size_t) ne01 * nth, (size_t) ne02, (size_t) ne03};
+        size_t local_work_size[]  = {(size_t) nth, 1, 1};
+
+        backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
+    }
+}
+
+static void ggml_cl_cos(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    GGML_ASSERT(src0);
+    GGML_ASSERT(src0->extra);
+    GGML_ASSERT(dst);
+    GGML_ASSERT(dst->extra);
+
+    UNUSED(src1);
+
+    ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *) backend->context;
+
+    ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *) src0->extra;
+    ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *) dst->extra;
+
+    cl_ulong offset0 = extra0->offset + src0->view_offs;
+    cl_ulong offsetd = extrad->offset + dst->view_offs;
+
+    const int ne00 = src0->ne[0];
+    const int ne01 = src0->ne[1];
+    const int ne02 = src0->ne[2];
+    const int ne03 = src0->ne[3];
+
+    const cl_ulong nb00 = src0->nb[0];
+    const cl_ulong nb01 = src0->nb[1];
+    const cl_ulong nb02 = src0->nb[2];
+    const cl_ulong nb03 = src0->nb[3];
+
+    const cl_ulong nb0 = dst->nb[0];
+    const cl_ulong nb1 = dst->nb[1];
+    const cl_ulong nb2 = dst->nb[2];
+    const cl_ulong nb3 = dst->nb[3];
+
+    cl_kernel kernel;
+
+    if (ggml_is_contiguous(src0)) {
+        int n = ggml_nelements(dst);
+        if (n % 4 == 0) {
+            if (src0->type == GGML_TYPE_F32) {
+                kernel = backend_ctx->kernel_cos_f32_4;
+            } else {
+                kernel = backend_ctx->kernel_cos_f16_4;
+            }
+            n /= 4;
+        } else {
+            if (src0->type == GGML_TYPE_F32) {
+                kernel = backend_ctx->kernel_cos_f32;
+            } else {
+                kernel = backend_ctx->kernel_cos_f16;
+            }
+        }
+
+        CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem),   &extra0->data_device));
+        CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0));
+        CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem),   &extrad->data_device));
+        CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offsetd));
+
+        size_t global_work_size[] = {(size_t) n, 1, 1};
+        size_t local_work_size[]  = {64, 1, 1};
+
+        size_t * local_work_size_ptr = local_work_size;
+        if (n % 64 != 0 && !backend_ctx->non_uniform_workgroups) {
+            local_work_size_ptr = nullptr;
+        }
+
+        backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size_ptr, dst);
+    } else {
+        if (src0->type == GGML_TYPE_F32) {
+            kernel = backend_ctx->kernel_cos_f32_nc;
+        } else {
+            kernel = backend_ctx->kernel_cos_f16_nc;
+        }
+
+        CL_CHECK(clSetKernelArg(kernel,  0, sizeof(cl_mem),   &extra0->data_device));
+        CL_CHECK(clSetKernelArg(kernel,  1, sizeof(cl_ulong), &offset0));
+        CL_CHECK(clSetKernelArg(kernel,  2, sizeof(cl_mem),   &extrad->data_device));
+        CL_CHECK(clSetKernelArg(kernel,  3, sizeof(cl_ulong), &offsetd));
+        CL_CHECK(clSetKernelArg(kernel,  4, sizeof(int),      &ne00));
+        CL_CHECK(clSetKernelArg(kernel,  5, sizeof(cl_ulong), &nb00));
+        CL_CHECK(clSetKernelArg(kernel,  6, sizeof(cl_ulong), &nb01));
+        CL_CHECK(clSetKernelArg(kernel,  7, sizeof(cl_ulong), &nb02));
+        CL_CHECK(clSetKernelArg(kernel,  8, sizeof(cl_ulong), &nb03));
+        CL_CHECK(clSetKernelArg(kernel,  9, sizeof(cl_ulong), &nb0));
+        CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_ulong), &nb1));
+        CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_ulong), &nb2));
+        CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_ulong), &nb3));
+
+        int nth = 64;
+
+        size_t global_work_size[] = {(size_t) ne01 * nth, (size_t) ne02, (size_t) ne03};
+        size_t local_work_size[]  = {(size_t) nth, 1, 1};
+
+        backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
+    }
+}
+
+static void ggml_cl_unary_f32(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
+        cl_kernel kernel_f32, cl_kernel kernel_f32_4, cl_kernel kernel_f32_nc) {
+    GGML_ASSERT(src0);
+    GGML_ASSERT(src0->extra);
+    GGML_ASSERT(dst);
+    GGML_ASSERT(dst->extra);
+    GGML_ASSERT(src0->type == GGML_TYPE_F32);
+    GGML_ASSERT(dst->type == GGML_TYPE_F32);
+
+    UNUSED(src1);
+
+    ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *) backend->context;
+
+    ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *) src0->extra;
+    ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *) dst->extra;
+
+    cl_ulong offset0 = extra0->offset + src0->view_offs;
+    cl_ulong offsetd = extrad->offset + dst->view_offs;
+
+    const int ne00 = src0->ne[0];
+    const int ne01 = src0->ne[1];
+    const int ne02 = src0->ne[2];
+    const int ne03 = src0->ne[3];
+
+    const cl_ulong nb00 = src0->nb[0];
+    const cl_ulong nb01 = src0->nb[1];
+    const cl_ulong nb02 = src0->nb[2];
+    const cl_ulong nb03 = src0->nb[3];
+
+    const cl_ulong nb0 = dst->nb[0];
+    const cl_ulong nb1 = dst->nb[1];
+    const cl_ulong nb2 = dst->nb[2];
+    const cl_ulong nb3 = dst->nb[3];
+
+    cl_kernel kernel;
+
+    if (ggml_is_contiguous(src0) && ggml_is_contiguous(dst)) {
+        int n = ggml_nelements(dst);
+        if (n % 4 == 0) {
+            kernel = kernel_f32_4;
+            n /= 4;
+        } else {
+            kernel = kernel_f32;
+        }
+
+        CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem),   &extra0->data_device));
+        CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0));
+        CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem),   &extrad->data_device));
+        CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offsetd));
+
+        size_t global_work_size[] = {(size_t) n, 1, 1};
+        size_t local_work_size[]  = {64, 1, 1};
+
+        size_t * local_work_size_ptr = local_work_size;
+        if (n % 64 != 0 && !backend_ctx->non_uniform_workgroups) {
+            local_work_size_ptr = nullptr;
+        }
+
+        backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size_ptr, dst);
+    } else {
+        kernel = kernel_f32_nc;
+
+        CL_CHECK(clSetKernelArg(kernel,  0, sizeof(cl_mem),   &extra0->data_device));
+        CL_CHECK(clSetKernelArg(kernel,  1, sizeof(cl_ulong), &offset0));
+        CL_CHECK(clSetKernelArg(kernel,  2, sizeof(cl_mem),   &extrad->data_device));
+        CL_CHECK(clSetKernelArg(kernel,  3, sizeof(cl_ulong), &offsetd));
+        CL_CHECK(clSetKernelArg(kernel,  4, sizeof(int),      &ne00));
+        CL_CHECK(clSetKernelArg(kernel,  5, sizeof(cl_ulong), &nb00));
+        CL_CHECK(clSetKernelArg(kernel,  6, sizeof(cl_ulong), &nb01));
+        CL_CHECK(clSetKernelArg(kernel,  7, sizeof(cl_ulong), &nb02));
+        CL_CHECK(clSetKernelArg(kernel,  8, sizeof(cl_ulong), &nb03));
+        CL_CHECK(clSetKernelArg(kernel,  9, sizeof(cl_ulong), &nb0));
+        CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_ulong), &nb1));
+        CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_ulong), &nb2));
+        CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_ulong), &nb3));
+
+        int nth = 64;
+
+        size_t global_work_size[] = {(size_t) ne01 * nth, (size_t) ne02, (size_t) ne03};
+        size_t local_work_size[]  = {(size_t) nth, 1, 1};
+
+        backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
+    }
+}
+
+static void ggml_cl_abs(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *) backend->context;
+    ggml_cl_unary_f32(backend, src0, src1, dst,
+            backend_ctx->kernel_abs_f32,
+            backend_ctx->kernel_abs_f32_4,
+            backend_ctx->kernel_abs_f32_nc);
+}
+
+static void ggml_cl_elu(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *) backend->context;
+    ggml_cl_unary_f32(backend, src0, src1, dst,
+            backend_ctx->kernel_elu_f32,
+            backend_ctx->kernel_elu_f32_4,
+            backend_ctx->kernel_elu_f32_nc);
+}
+
 static void ggml_cl_neg(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
     GGML_ASSERT(src0);
     GGML_ASSERT(src0->extra);
@@ -9247,6 +9957,83 @@ static void ggml_cl_conv_2d(ggml_backend_t backend, const ggml_tensor * src0, co
     backend_ctx->enqueue_ndrange_kernel(kernel, 2, global_work_size, local_work_size, dst);
 }
 
+static void ggml_cl_conv_transpose_1d(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *) backend->context;
+
+    GGML_ASSERT(dst->type == GGML_TYPE_F32);
+    GGML_ASSERT(src1->type == GGML_TYPE_F32);
+    GGML_ASSERT(ggml_get_op_params_i32(dst, 1) == 0 && ggml_get_op_params_i32(dst, 2) == 1);
+
+    const int32_t s0 = ggml_get_op_params_i32(dst, 0);
+
+    ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *) src0->extra;
+    ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *) src1->extra;
+    ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *) dst->extra;
+
+    cl_ulong offset0 = extra0->offset + src0->view_offs;
+    cl_ulong offset1 = extra1->offset + src1->view_offs;
+    cl_ulong offsetd = extrad->offset + dst->view_offs;
+
+    const cl_uint K     = (cl_uint) src0->ne[0];
+    const cl_uint Cout  = (cl_uint) src0->ne[1];
+    const cl_uint Cin   = (cl_uint) src0->ne[2];
+    const cl_uint L     = (cl_uint) src1->ne[0];
+    const cl_uint KL    = (cl_uint) dst->ne[0];
+    const cl_uint nbatch = (cl_uint) MAX((int64_t) 1, dst->ne[2]);
+
+    const cl_uint ts0 = (cl_uint) ggml_type_size(src0->type);
+    const cl_uint ts1 = (cl_uint) ggml_type_size(src1->type);
+    const cl_uint tsd = (cl_uint) ggml_type_size(dst->type);
+
+    const cl_uint k_nb0  = (cl_uint) (src0->nb[0] / ts0);
+    const cl_uint k_nb01 = (cl_uint) (src0->nb[1] / ts0);
+    const cl_uint k_nb02 = (cl_uint) (src0->nb[2] / ts0);
+    const cl_uint in_nb0 = (cl_uint) (src1->nb[0] / ts1);
+    const cl_uint in_nb1 = (cl_uint) (src1->nb[1] / ts1);
+    const cl_uint in_nb2 = (cl_uint) (src1->nb[2] / ts1);
+    const cl_uint d_nb0  = (cl_uint) (dst->nb[0] / tsd);
+    const cl_uint d_nb1  = (cl_uint) (dst->nb[1] / tsd);
+    const cl_uint d_nb2  = (cl_uint) (dst->nb[2] / tsd);
+
+    cl_kernel kernel = nullptr;
+    if (src0->type == GGML_TYPE_F32) {
+        kernel = backend_ctx->kernel_conv_transpose_1d_f32;
+    } else if (src0->type == GGML_TYPE_F16) {
+        kernel = backend_ctx->kernel_conv_transpose_1d_f16_f32;
+    } else {
+        GGML_ABORT("ggml_cl_conv_transpose_1d: unsupported src0 type");
+    }
+    GGML_ASSERT(kernel != nullptr);
+
+    cl_uint idx = 0;
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem),   &extra0->data_device));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_ulong), &offset0));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem),   &extra1->data_device));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_ulong), &offset1));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_mem),   &extrad->data_device));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_ulong), &offsetd));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &k_nb0));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &k_nb01));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &k_nb02));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &in_nb0));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &in_nb1));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &in_nb2));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &d_nb0));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &d_nb1));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &d_nb2));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &K));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &Cout));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &Cin));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &L));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &KL));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(cl_uint),  &nbatch));
+    CL_CHECK(clSetKernelArg(kernel, idx++, sizeof(int32_t),  &s0));
+
+    const size_t total     = (size_t) KL * (size_t) Cout * (size_t) nbatch;
+    size_t       global_work_size[] = { total, 1, 1 };
+    backend_ctx->enqueue_ndrange_kernel(kernel, 1, global_work_size, nullptr, dst);
+}
+
 static void ggml_cl_mul_mat_kq_kqv_adreno(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
     ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context;
 
@@ -10190,7 +10977,8 @@ static void ggml_cl_mul_mat_q6_K_f32_adreno(ggml_backend_t backend, const ggml_t
 #endif
 }
 
-static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+static void ggml_cl_mul_mat_impl(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
+        const ggml_tensor * bias = nullptr, const ggml_tensor * residual = nullptr) {
     GGML_ASSERT(src0);
     GGML_ASSERT(src0->extra);
     GGML_ASSERT(src1);
@@ -10206,10 +10994,14 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
     ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *)src0->extra;
     ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *)src1->extra;
     ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra;
+    ggml_tensor_extra_cl * extrabias = bias ? (ggml_tensor_extra_cl *) bias->extra : nullptr;
+    ggml_tensor_extra_cl * extraresidual = residual ? (ggml_tensor_extra_cl *) residual->extra : nullptr;
 
     cl_ulong offset0 = extra0->offset + src0->view_offs;
     cl_ulong offset1 = extra1->offset + src1->view_offs;
     cl_ulong offsetd = extrad->offset + dst->view_offs;
+    cl_ulong offsetbias = extrabias ? extrabias->offset + bias->view_offs : 0;
+    cl_ulong offsetresidual = extraresidual ? extraresidual->offset + residual->view_offs : 0;
 
 #ifdef GGML_OPENCL_SOA_Q
     ggml_tensor_extra_cl_q4_0 * extra0_q4_0 = (ggml_tensor_extra_cl_q4_0 *)src0->extra;
@@ -10527,6 +11319,16 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
             CL_CHECK(clSetKernelArg(kernel,  k_arg++, sizeof(int),      &ne1));
             CL_CHECK(clSetKernelArg(kernel,  k_arg++, sizeof(int),      &r2));
             CL_CHECK(clSetKernelArg(kernel,  k_arg++, sizeof(int),      &r3));
+            cl_mem bias_mem = extrabias ? extrabias->data_device : extra1->data_device;
+            cl_mem residual_mem = extraresidual ? extraresidual->data_device : extra1->data_device;
+            const int has_bias = bias != nullptr;
+            const int has_residual = residual != nullptr;
+            CL_CHECK(clSetKernelArg(kernel,  k_arg++, sizeof(cl_mem),   &bias_mem));
+            CL_CHECK(clSetKernelArg(kernel,  k_arg++, sizeof(cl_ulong), &offsetbias));
+            CL_CHECK(clSetKernelArg(kernel,  k_arg++, sizeof(cl_mem),   &residual_mem));
+            CL_CHECK(clSetKernelArg(kernel,  k_arg++, sizeof(cl_ulong), &offsetresidual));
+            CL_CHECK(clSetKernelArg(kernel,  k_arg++, sizeof(int),      &has_bias));
+            CL_CHECK(clSetKernelArg(kernel,  k_arg++, sizeof(int),      &has_residual));
         } else {
             region.origin = extrad->offset; // Specify the starting offset (in bytes)
             region.size = M * N * sizeof(float); // Specify the size of the sub-buffer
@@ -10675,7 +11477,6 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
                 CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int),      &batch_stride_d));
                 CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int),      &r2));
                 CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int),      &r3));
-
                 // 64 is block tile size BM and BN - change here when BM and BN in the kernel are changed.
                 size_t global_work_size[] = {(size_t)(CEIL_DIV(ne01, 64)*nth0), (size_t)(CEIL_DIV(ne11, 64)), (size_t)ne12*ne13};
                 size_t local_work_size[] = {(size_t)nth0, 1, 1};
@@ -10742,7 +11543,6 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
                 CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int),      &batch_stride_d));
                 CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int),      &r2));
                 CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int),      &r3));
-
                 // 64 is block tile size BM and BN - change here when BM and BN in the kernel are changed.
                 size_t global_work_size[] = {(size_t)(CEIL_DIV(ne01, 64)*nth0), (size_t)(CEIL_DIV(ne11, 64)), (size_t)ne12*ne13};
                 size_t local_work_size[] = {(size_t)nth0, 1, 1};
@@ -10784,6 +11584,18 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
                 CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int),      &batch_stride_d));
                 CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int),      &r2));
                 CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int),      &r3));
+                {
+                    cl_mem bias_mem = extrabias ? extrabias->data_device : extra1->data_device;
+                    cl_mem residual_mem = extraresidual ? extraresidual->data_device : extra1->data_device;
+                    const int has_bias = bias != nullptr;
+                    const int has_residual = residual != nullptr;
+                    CL_CHECK(clSetKernelArg(kernel, 19, sizeof(cl_mem),   &bias_mem));
+                    CL_CHECK(clSetKernelArg(kernel, 20, sizeof(cl_ulong), &offsetbias));
+                    CL_CHECK(clSetKernelArg(kernel, 21, sizeof(cl_mem),   &residual_mem));
+                    CL_CHECK(clSetKernelArg(kernel, 22, sizeof(cl_ulong), &offsetresidual));
+                    CL_CHECK(clSetKernelArg(kernel, 23, sizeof(int),      &has_bias));
+                    CL_CHECK(clSetKernelArg(kernel, 24, sizeof(int),      &has_residual));
+                }
 
                 // 64 is block tile size BM and BN - change here when BM and BN in the kernel are changed.
                 size_t global_work_size[] = {(size_t)(CEIL_DIV(ne01, 64)*nth0), (size_t)(CEIL_DIV(ne11, 64)), (size_t)ne12*ne13};
@@ -11622,6 +12434,23 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
     }
 }
 
+static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    ggml_cl_mul_mat_impl(backend, src0, src1, dst);
+}
+
+static void ggml_opencl_op_mul_mat_add_fused(ggml_backend_t backend, ggml_tensor * mul, ggml_tensor * add0, ggml_tensor * add1) {
+    const ggml_tensor * bias = add0->src[0] == mul ? add0->src[1] : add0->src[0];
+    const ggml_tensor * residual = nullptr;
+    ggml_tensor * dst = add0;
+
+    if (add1) {
+        residual = add1->src[0] == add0 ? add1->src[1] : add1->src[0];
+        dst = add1;
+    }
+
+    ggml_cl_mul_mat_impl(backend, mul->src[0], mul->src[1], dst, bias, residual);
+}
+
 static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
     GGML_ASSERT(src0);
     GGML_ASSERT(src0->extra);
@@ -13235,6 +14064,24 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor
             }
             func = ggml_cl_sqrt;
             break;
+        case GGML_OP_LEAKY_RELU:
+            if (!any_on_device) {
+                return false;
+            }
+            func = ggml_cl_leaky_relu;
+            break;
+        case GGML_OP_SIN:
+            if (!any_on_device) {
+                return false;
+            }
+            func = ggml_cl_sin;
+            break;
+        case GGML_OP_COS:
+            if (!any_on_device) {
+                return false;
+            }
+            func = ggml_cl_cos;
+            break;
         case GGML_OP_MEAN:
             if (!any_on_device) {
                 return false;
@@ -13243,6 +14090,18 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor
             break;
         case GGML_OP_UNARY:
             switch (ggml_get_unary_op(tensor)) {
+                case GGML_UNARY_OP_ABS:
+                    if (!any_on_device) {
+                        return false;
+                    }
+                    func = ggml_cl_abs;
+                    break;
+                case GGML_UNARY_OP_ELU:
+                    if (!any_on_device) {
+                        return false;
+                    }
+                    func = ggml_cl_elu;
+                    break;
                 case GGML_UNARY_OP_GELU:
                     if (!any_on_device) {
                         return false;
@@ -13384,6 +14243,12 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor
             }
             func = ggml_cl_conv_2d;
             break;
+        case GGML_OP_CONV_TRANSPOSE_1D:
+            if (!any_on_device) {
+                return false;
+            }
+            func = ggml_cl_conv_transpose_1d;
+            break;
         case GGML_OP_SSM_CONV:
             if (!any_on_device) {
                 return false;
diff --git a/src/ggml-opencl/kernels/abs.cl b/src/ggml-opencl/kernels/abs.cl
new file mode 100644
index 00000000..5b38fd57
--- /dev/null
+++ b/src/ggml-opencl/kernels/abs.cl
@@ -0,0 +1,55 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+kernel void kernel_abs_f32(
+        global const float * src0,
+        ulong                offset0,
+        global       float * dst,
+        ulong                offsetd
+) {
+    src0 = (global float*)((global char*)src0 + offset0);
+    dst  = (global float*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = fabs(src0[get_global_id(0)]);
+}
+
+kernel void kernel_abs_f32_4(
+        global const float4 * src0,
+        ulong                 offset0,
+        global       float4 * dst,
+        ulong                 offsetd
+) {
+    src0 = (global float4*)((global char*)src0 + offset0);
+    dst  = (global float4*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = fabs(src0[get_global_id(0)]);
+}
+
+kernel void kernel_abs_f32_nc(
+        global const char * src0,
+        ulong               offset0,
+        global       char * dst,
+        ulong               offsetd,
+        int   ne00,
+        ulong nb00,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        ulong nb0,
+        ulong nb1,
+        ulong nb2,
+        ulong nb3
+) {
+    src0 = src0 + offset0;
+    dst  = dst + offsetd;
+
+    const int i3 = get_group_id(2);
+    const int i2 = get_group_id(1);
+    const int i1 = get_group_id(0);
+
+    for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {
+        global const float * x = (global const float *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
+        global       float * y = (global       float *)(dst  + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
+
+        *y = fabs(*x);
+    }
+}
diff --git a/src/ggml-opencl/kernels/add.cl b/src/ggml-opencl/kernels/add.cl
index 509bf173..0495fe2b 100644
--- a/src/ggml-opencl/kernels/add.cl
+++ b/src/ggml-opencl/kernels/add.cl
@@ -82,6 +82,29 @@ kernel void kernel_add_row(
     dst[gid] = src0[gid] + src1[idx1];
 }
 
+// Fused pattern: dst = src0 + row + residual.
+// src0 and residual have dst shape; row is broadcast along rows.
+kernel void kernel_add_row_residual(
+        global float4 * src0,
+        ulong  offset0,
+        global float4 * row,
+        ulong  offset_row,
+        global float4 * residual,
+        ulong  offset_residual,
+        global float4 * dst,
+        ulong  offsetd,
+        int ne
+) {
+    src0     = (global float4*)((global char*)src0     + offset0);
+    row      = (global float4*)((global char*)row      + offset_row);
+    residual = (global float4*)((global char*)residual + offset_residual);
+    dst      = (global float4*)((global char*)dst      + offsetd);
+
+    uint gid = get_global_id(0);
+    uint idx_row = gid - (gid/ne)*ne; // get_global_id(0) % ne
+    dst[gid] = src0[gid] + row[idx_row] + residual[gid];
+}
+
 kernel void kernel_add_f16(
         global char * src0,
         ulong  offset0,
diff --git a/src/ggml-opencl/kernels/conv_transpose_1d_f16_f32.cl b/src/ggml-opencl/kernels/conv_transpose_1d_f16_f32.cl
new file mode 100644
index 00000000..1bf11ea0
--- /dev/null
+++ b/src/ggml-opencl/kernels/conv_transpose_1d_f16_f32.cl
@@ -0,0 +1,57 @@
+// CONV_TRANSPOSE_1D — f16 kernel / f32 input / f32 out
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#define CTD_IDX(ox, i1, ib, dnb0, dnb1, dnb2) \
+    ((ox) * (dnb0) + (i1) * (dnb1) + (ib) * (dnb2))
+
+kernel void kernel_conv_transpose_1d_f16_f32(
+    global void * p_k, ulong off_k,
+    global void * p_in, ulong off_in,
+    global void * p_d, ulong off_d,
+    uint k_nb0, uint k_nb01, uint k_nb02,
+    uint in_nb0, uint in_nb1, uint in_nb2,
+    uint d_nb0, uint d_nb1, uint d_nb2,
+    uint K, uint Cout, uint Cin, uint L, uint KL,
+    uint nbatch,
+    int s0
+) {
+    global const half * kA = (global const half *) ((global const char *) p_k + off_k);
+    global const float * bB = (global const float *) ((global const char *) p_in + off_in);
+    global float * dD = (global float *) ((global char *) p_d + off_d);
+
+    const uint total = KL * Cout * nbatch;
+    const uint gid = get_global_id(0);
+    if (gid >= total) {
+        return;
+    }
+
+    const uint ox  = gid % KL;
+    const uint t1  = gid / KL;
+    const uint i1  = t1 % Cout;
+    const uint ib  = t1 / Cout;
+
+    const int ox_i = (int) ox;
+    const int K_i  = (int) K;
+    const int L_i  = (int) L;
+
+    // Only input positions with 0 <= ox - s0*i10 < K contribute.
+    int i10_min = (ox_i - K_i + s0) / s0; // ceil((ox - K + 1) / s0)
+    int i10_max = ox_i / s0;              // floor(ox / s0)
+    i10_min = max(i10_min, 0);
+    i10_max = min(i10_max, L_i - 1);
+
+    float acc = 0.0f;
+    for (int i10 = i10_min; i10 <= i10_max; i10++) {
+        const int i00 = (int) ox - s0 * (int) i10;
+        const uint ka0 = (uint) i00 * k_nb0 + i1 * k_nb01;
+        const uint ba0 = (uint) i10 * in_nb0 + ib * in_nb2;
+        for (uint ic = 0; ic < Cin; ic++) {
+            float kv = (float) kA[ka0 + ic * k_nb02];
+            float iv = bB[ba0 + ic * in_nb1];
+            acc = fma(kv, iv, acc);
+        }
+    }
+
+    dD[CTD_IDX(ox, i1, ib, d_nb0, d_nb1, d_nb2)] = acc;
+}
diff --git a/src/ggml-opencl/kernels/conv_transpose_1d_f32.cl b/src/ggml-opencl/kernels/conv_transpose_1d_f32.cl
new file mode 100644
index 00000000..e7d8abeb
--- /dev/null
+++ b/src/ggml-opencl/kernels/conv_transpose_1d_f32.cl
@@ -0,0 +1,56 @@
+// CONV_TRANSPOSE_1D — f32 / f32 / f32
+// src0: [K, Cout, Cin], src1: [L, Cin, ...], dst: [KL, Cout, ...]
+
+#define CTD_IDX(ox, i1, ib, dnb0, dnb1, dnb2) \
+    ((ox) * (dnb0) + (i1) * (dnb1) + (ib) * (dnb2))
+
+kernel void kernel_conv_transpose_1d_f32(
+    global void * p_k, ulong off_k,
+    global void * p_in, ulong off_in,
+    global void * p_d, ulong off_d,
+    uint k_nb0, uint k_nb01, uint k_nb02,
+    uint in_nb0, uint in_nb1, uint in_nb2,
+    uint d_nb0, uint d_nb1, uint d_nb2,
+    uint K, uint Cout, uint Cin, uint L, uint KL,
+    uint nbatch,
+    int s0
+) {
+    global const float * kA = (global const float *) ((global const char *) p_k + off_k);
+    global const float * bB = (global const float *) ((global const char *) p_in + off_in);
+    global float * dD = (global float *) ((global char *) p_d + off_d);
+
+    const uint total = KL * Cout * nbatch;
+    const uint gid = get_global_id(0);
+    if (gid >= total) {
+        return;
+    }
+
+    const uint ox  = gid % KL;
+    const uint t1  = gid / KL;
+    const uint i1  = t1 % Cout;
+    const uint ib  = t1 / Cout;
+
+    const int ox_i = (int) ox;
+    const int K_i  = (int) K;
+    const int L_i  = (int) L;
+
+    // Only input positions with 0 <= ox - s0*i10 < K contribute.
+    int i10_min = (ox_i - K_i + s0) / s0; // ceil((ox - K + 1) / s0)
+    int i10_max = ox_i / s0;              // floor(ox / s0)
+    i10_min = max(i10_min, 0);
+    i10_max = min(i10_max, L_i - 1);
+
+    float acc = 0.0f;
+    for (int i10 = i10_min; i10 <= i10_max; i10++) {
+        const int i00 = (int) ox - s0 * (int) i10;
+        const uint ka0 = (uint) i00 * k_nb0 + i1 * k_nb01;
+        const uint ba0 = (uint) i10 * in_nb0 + ib * in_nb2;
+        for (uint ic = 0; ic < Cin; ic++) {
+            float kv = kA[ka0 + ic * k_nb02];
+            float iv = bB[ba0 + ic * in_nb1];
+            acc = fma(kv, iv, acc);
+        }
+    }
+
+    dD[CTD_IDX(ox, i1, ib, d_nb0, d_nb1, d_nb2)] = acc;
+}
diff --git a/src/ggml-opencl/kernels/cos.cl b/src/ggml-opencl/kernels/cos.cl
new file mode 100644
index 00000000..8c340356
--- /dev/null
+++ b/src/ggml-opencl/kernels/cos.cl
@@ -0,0 +1,109 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+kernel void kernel_cos_f32(
+        global const float * src0,
+        ulong                offset0,
+        global       float * dst,
+        ulong                offsetd
+) {
+    src0 = (global float*)((global char*)src0 + offset0);
+    dst  = (global float*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = cos(src0[get_global_id(0)]);
+}
+
+kernel void kernel_cos_f32_4(
+        global const float4 * src0,
+        ulong                 offset0,
+        global       float4 * dst,
+        ulong                 offsetd
+) {
+    src0 = (global float4*)((global char*)src0 + offset0);
+    dst  = (global float4*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = cos(src0[get_global_id(0)]);
+}
+
+kernel void kernel_cos_f16(
+        global const half * src0,
+        ulong               offset0,
+        global       half * dst,
+        ulong               offsetd
+) {
+    src0 = (global half*)((global char*)src0 + offset0);
+    dst  = (global half*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = cos(src0[get_global_id(0)]);
+}
+
+kernel void kernel_cos_f16_4(
+        global const half4 * src0,
+        ulong                offset0,
+        global       half4 * dst,
+        ulong                offsetd
+) {
+    src0 = (global half4*)((global char*)src0 + offset0);
+    dst  = (global half4*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = cos(src0[get_global_id(0)]);
+}
+
+kernel void kernel_cos_f32_nc(
+        global const char * src0,
+        ulong               offset0,
+        global       char * dst,
+        ulong               offsetd,
+        int   ne00,
+        ulong nb00,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        ulong nb0,
+        ulong nb1,
+        ulong nb2,
+        ulong nb3
+) {
+    src0 = src0 + offset0;
+    dst  = dst + offsetd;
+
+    const int i3 = get_group_id(2);
+    const int i2 = get_group_id(1);
+    const int i1 = get_group_id(0);
+
+    for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {
+        global const float * x = (global const float *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
+        global       float * y = (global       float *)(dst  + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
+
+        *y = cos(*x);
+    }
+}
+
+kernel void kernel_cos_f16_nc(
+        global const char * src0,
+        ulong               offset0,
+        global       char * dst,
+        ulong               offsetd,
+        int   ne00,
+        ulong nb00,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        ulong nb0,
+        ulong nb1,
+        ulong nb2,
+        ulong nb3
+) {
+    src0 = src0 + offset0;
+    dst  = dst + offsetd;
+
+    const int i3 = get_group_id(2);
+    const int i2 = get_group_id(1);
+    const int i1 = get_group_id(0);
+
+    for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {
+        global const half * x = (global const half *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
+        global       half * y = (global       half *)(dst  + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
+
+        *y = cos(*x);
+    }
+}
diff --git a/src/ggml-opencl/kernels/elu.cl b/src/ggml-opencl/kernels/elu.cl
new file mode 100644
index 00000000..af113dc8
--- /dev/null
+++ b/src/ggml-opencl/kernels/elu.cl
@@ -0,0 +1,63 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+static inline float elu_f32(float x) {
+    return fmax(x, 0.0f) + fmin(expm1(x), 0.0f);
+}
+
+static inline float4 elu_f32_4(float4 x) {
+    return fmax(x, (float4)(0.0f)) + fmin(expm1(x), (float4)(0.0f));
+}
+
+kernel void kernel_elu_f32(
+        global const float * src0,
+        ulong                offset0,
+        global       float * dst,
+        ulong                offsetd
+) {
+    src0 = (global float*)((global char*)src0 + offset0);
+    dst  = (global float*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = elu_f32(src0[get_global_id(0)]);
+}
+
+kernel void kernel_elu_f32_4(
+        global const float4 * src0,
+        ulong                 offset0,
+        global       float4 * dst,
+        ulong                 offsetd
+) {
+    src0 = (global float4*)((global char*)src0 + offset0);
+    dst  = (global float4*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = elu_f32_4(src0[get_global_id(0)]);
+}
+
+kernel void kernel_elu_f32_nc(
+        global const char * src0,
+        ulong               offset0,
+        global       char * dst,
+        ulong               offsetd,
+        int   ne00,
+        ulong nb00,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        ulong nb0,
+        ulong nb1,
+        ulong nb2,
+        ulong nb3
+) {
+    src0 = src0 + offset0;
+    dst  = dst + offsetd;
+
+    const int i3 = get_group_id(2);
+    const int i2 = get_group_id(1);
+    const int i1 = get_group_id(0);
+
+    for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {
+        global const float * x = (global const float *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
+        global       float * y = (global       float *)(dst  + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
+
+        *y = elu_f32(*x);
+    }
+}
diff --git a/src/ggml-opencl/kernels/gemv_noshuffle.cl b/src/ggml-opencl/kernels/gemv_noshuffle.cl
index ee5c79f0..e2eb5f96 100644
--- a/src/ggml-opencl/kernels/gemv_noshuffle.cl
+++ b/src/ggml-opencl/kernels/gemv_noshuffle.cl
@@ -206,7 +206,13 @@ __kernel void kernel_gemv_noshuffle(
         int ne0,                // M
         int ne1,                // N
         int r2,                 // 1
-        int r3)
+        int r3,
+        global float * bias,
+        ulong offset_bias,
+        global float * residual,
+        ulong offset_residual,
+        int has_bias,
+        int has_residual)
 {
     uint groupId = get_local_id(1);
     uint gid     = get_global_id(0);
@@ -262,6 +268,14 @@ __kernel void kernel_gemv_noshuffle(
     // 2 outputs per fiber in wave 0
     if (groupId == 0) {
         dst = (global float*)((global char*)dst + offsetd);
+        if (has_bias) {
+            bias = (global float*)((global char*)bias + offset_bias);
+            totalSum += vload2(0, &(bias[gid * 2]));
+        }
+        if (has_residual) {
+            residual = (global float*)((global char*)residual + offset_residual);
+            totalSum += vload2(0, &(residual[gid * 2]));
+        }
         vstore2(totalSum, 0, &(dst[gid * 2]));
     }
 
diff --git a/src/ggml-opencl/kernels/gemv_noshuffle_general.cl b/src/ggml-opencl/kernels/gemv_noshuffle_general.cl
index 469d3ede..8a77e27f 100644
--- a/src/ggml-opencl/kernels/gemv_noshuffle_general.cl
+++ b/src/ggml-opencl/kernels/gemv_noshuffle_general.cl
@@ -206,7 +206,13 @@ __kernel void kernel_gemv_noshuffle(
         int ne0,                // M
         int ne1,                // N
         int r2,                 // 1
-        int r3)
+        int r3,
+        global float * bias,
+        ulong offset_bias,
+        global float * residual,
+        ulong offset_residual,
+        int has_bias,
+        int has_residual)
 {
     uint groupId = get_local_id(1);
     uint gid     = get_global_id(0);
@@ -268,6 +274,14 @@ __kernel void kernel_gemv_noshuffle(
     // 2 outputs per fiber in wave 0
     if (groupId == 0) {
         dst = (global float*)((global char*)dst + offsetd);
+        if (has_bias) {
+            bias = (global float*)((global char*)bias + offset_bias);
+            totalSum += vload2(0, &(bias[gid * 2]));
+        }
+        if (has_residual) {
+            residual = (global float*)((global char*)residual + offset_residual);
+            totalSum += vload2(0, &(residual[gid * 2]));
+        }
         vstore2(totalSum, 0, &(dst[gid * 2]));
     }
 
diff --git a/src/ggml-opencl/kernels/mish.cl b/src/ggml-opencl/kernels/mish.cl
new file mode 100644
index 00000000..a1ebf57f
--- /dev/null
+++ b/src/ggml-opencl/kernels/mish.cl
@@ -0,0 +1,35 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+static inline float mish_f32(float x) {
+    const float sp = (x > 20.0f) ? x : log(1.0f + exp(x));
+    return x * tanh(sp);
+}
+
+static inline float4 mish_f32_4(float4 x) {
+    const float4 sp = (x > (float4)(20.0f)) ? x : log((float4)(1.0f) + exp(x));
+    return x * tanh(sp);
+}
+
+kernel void kernel_mish_f32(
+        global const float * src0,
+        ulong                offset0,
+        global       float * dst,
+        ulong                offsetd
+) {
+    src0 = (global float*)((global char*)src0 + offset0);
+    dst  = (global float*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = mish_f32(src0[get_global_id(0)]);
+}
+
+kernel void kernel_mish_f32_4(
+        global const float4 * src0,
+        ulong                 offset0,
+        global       float4 * dst,
+        ulong                 offsetd
+) {
+    src0 = (global float4*)((global char*)src0 + offset0);
+    dst  = (global float4*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = mish_f32_4(src0[get_global_id(0)]);
+}
diff --git a/src/ggml-opencl/kernels/mul_mm_q4_0_f32_l4_lm.cl b/src/ggml-opencl/kernels/mul_mm_q4_0_f32_l4_lm.cl
index 4100e308..394fc1e0 100644
--- a/src/ggml-opencl/kernels/mul_mm_q4_0_f32_l4_lm.cl
+++ b/src/ggml-opencl/kernels/mul_mm_q4_0_f32_l4_lm.cl
@@ -32,7 +32,13 @@ kernel void kernel_mul_mm_q4_0_f32_l4_lm(
     int batch_stride_d,
 
     int r2,
-    int r3
+    int r3,
+    global float * bias,
+    ulong offset_bias,
+    global float * residual,
+    ulong offset_residual,
+    int has_bias,
+    int has_residual
 ) {
     src1 = (global float4*)((global char*)src1 + offset1);
     dst  = (global float *)((global char*)dst  + offsetd);
@@ -152,11 +158,21 @@ kernel void kernel_mul_mm_q4_0_f32_l4_lm(
     const int dc = ic * BN + th_c * TN;
 
     const int offsets = batch_idx * batch_stride_d;
+    bias = (global float*)((global char*)bias + offset_bias);
+    residual = (global float*)((global char*)residual + offset_residual);
 
     for (int cc = 0; cc < TN; cc++) {
         for (int cr = 0; cr < TM; cr++) {
             if (dr + cr < ne01 && dc + cc < ne11) {
-                dst[offsets + (dc + cc) * stride_d + dr + cr] = sums[cc * TM + cr];
+                const int idx = offsets + (dc + cc) * stride_d + dr + cr;
+                float v = sums[cc * TM + cr];
+                if (has_bias) {
+                    v += bias[dr + cr];
+                }
+                if (has_residual) {
+                    v += residual[idx];
+                }
+                dst[idx] = v;
             }
         }
     }
diff --git a/src/ggml-opencl/kernels/relu.cl b/src/ggml-opencl/kernels/relu.cl
index 60ff28a6..249f5e99 100644
--- a/src/ggml-opencl/kernels/relu.cl
+++ b/src/ggml-opencl/kernels/relu.cl
@@ -14,3 +14,17 @@ kernel void kernel_relu(
 
     dst[get_global_id(0)] = fmax(0.0f, src0[get_global_id(0)]);
 }
+
+kernel void kernel_leaky_relu(
+        global float * src0,
+        ulong offset0,
+        global float * dst,
+        ulong offsetd,
+        float negative_slope
+) {
+    src0 = (global float*)((global char*)src0 + offset0);
+    dst = (global float*)((global char*)dst + offsetd);
+
+    const float x = src0[get_global_id(0)];
+    dst[get_global_id(0)] = x >= 0.0f ? x : x * negative_slope;
+}
diff --git a/src/ggml-opencl/kernels/sin.cl b/src/ggml-opencl/kernels/sin.cl
new file mode 100644
index 00000000..4ccbf5af
--- /dev/null
+++ b/src/ggml-opencl/kernels/sin.cl
@@ -0,0 +1,109 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+kernel void kernel_sin_f32(
+        global const float * src0,
+        ulong                offset0,
+        global       float * dst,
+        ulong                offsetd
+) {
+    src0 = (global float*)((global char*)src0 + offset0);
+    dst  = (global float*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = sin(src0[get_global_id(0)]);
+}
+
+kernel void kernel_sin_f32_4(
+        global const float4 * src0,
+        ulong                 offset0,
+        global       float4 * dst,
+        ulong                 offsetd
+) {
+    src0 = (global float4*)((global char*)src0 + offset0);
+    dst  = (global float4*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = sin(src0[get_global_id(0)]);
+}
+
+kernel void kernel_sin_f16(
+        global const half * src0,
+        ulong               offset0,
+        global       half * dst,
+        ulong               offsetd
+) {
+    src0 = (global half*)((global char*)src0 + offset0);
+    dst  = (global half*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = sin(src0[get_global_id(0)]);
+}
+
+kernel void kernel_sin_f16_4(
+        global const half4 * src0,
+        ulong                offset0,
+        global       half4 * dst,
+        ulong                offsetd
+) {
+    src0 = (global half4*)((global char*)src0 + offset0);
+    dst  = (global half4*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = sin(src0[get_global_id(0)]);
+}
+
+kernel void kernel_sin_f32_nc(
+        global const char * src0,
+        ulong               offset0,
+        global       char * dst,
+        ulong               offsetd,
+        int   ne00,
+        ulong nb00,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        ulong nb0,
+        ulong nb1,
+        ulong nb2,
+        ulong nb3
+) {
+    src0 = src0 + offset0;
+    dst  = dst + offsetd;
+
+    const int i3 = get_group_id(2);
+    const int i2 = get_group_id(1);
+    const int i1 = get_group_id(0);
+
+    for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {
+        global const float * x = (global const float *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
+        global       float * y = (global       float *)(dst  + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
+
+        *y = sin(*x);
+    }
+}
+
+kernel void kernel_sin_f16_nc(
+        global const char * src0,
+        ulong               offset0,
+        global       char * dst,
+        ulong               offsetd,
+        int   ne00,
+        ulong nb00,
+        ulong nb01,
+        ulong nb02,
+        ulong nb03,
+        ulong nb0,
+        ulong nb1,
+        ulong nb2,
+        ulong nb3
+) {
+    src0 = src0 + offset0;
+    dst  = dst + offsetd;
+
+    const int i3 = get_group_id(2);
+    const int i2 = get_group_id(1);
+    const int i1 = get_group_id(0);
+
+    for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {
+        global const half * x = (global const half *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
+        global       half * y = (global       half *)(dst  + i3*nb3  + i2*nb2  + i1*nb1  + i0*nb0);
+
+        *y = sin(*x);
+    }
+}