small fixes

Author: rmatif

ggml/src/ggml-opencl/CMakeLists.txt

@@ -98,7 +98,7 @@ set(GGML_OPENCL_KERNELS
     concat
     tsembd
     upscale
-    unary
+    tanh
     pad
     repeat
 )

ggml/src/ggml-opencl/ggml-opencl.cpp

@@ -317,7 +317,7 @@ struct ggml_backend_opencl_context {
     cl_program program_sum_rows_f32;
     cl_program program_repeat;
     cl_program program_pad;
-    cl_program program_unary;
+    cl_program program_tanh;
     cl_program program_upscale;
     cl_program program_concat;
     cl_program program_tsembd;
@@ -1154,24 +1154,24 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx, ggml_cl_ve
         }
     }
 
-    // unary (tanh)
+    // tanh
     {
 #ifdef GGML_OPENCL_EMBED_KERNELS
         const std::string kernel_src {
-            #include "unary.cl.h"
+            #include "tanh.cl.h"
         };
 #else
-        const std::string kernel_src = read_file("unary.cl");
+        const std::string kernel_src = read_file("tanh.cl");
 #endif
         if (!kernel_src.empty()) {
-            backend_ctx->program_unary =
+            backend_ctx->program_tanh =
                 build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
-            CL_CHECK((backend_ctx->kernel_tanh_f32_nd = clCreateKernel(backend_ctx->program_unary, "kernel_tanh_f32_nd", &err), err));
-            CL_CHECK((backend_ctx->kernel_tanh_f16_nd = clCreateKernel(backend_ctx->program_unary, "kernel_tanh_f16_nd", &err), err));
+            CL_CHECK((backend_ctx->kernel_tanh_f32_nd = clCreateKernel(backend_ctx->program_tanh, "kernel_tanh_f32_nd", &err), err));
+            CL_CHECK((backend_ctx->kernel_tanh_f16_nd = clCreateKernel(backend_ctx->program_tanh, "kernel_tanh_f16_nd", &err), err));
             GGML_LOG_CONT(".");
         } else {
-            GGML_LOG_WARN("ggml_opencl: unary kernel source not found or empty. Unary operations like tanh will not be available.\n");
-            backend_ctx->program_unary = nullptr;
+            GGML_LOG_WARN("ggml_opencl: tanh kernel source not found or empty. Tanh operation will not be available.\n");
+            backend_ctx->program_tanh = nullptr;
             backend_ctx->kernel_tanh_f32_nd = nullptr;
             backend_ctx->kernel_tanh_f16_nd = nullptr;
         }
@@ -2150,7 +2150,7 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
         case GGML_OP_NORM:
         case GGML_OP_RMS_NORM:
             return true;
-                case GGML_OP_REPEAT:
+        case GGML_OP_REPEAT:
             return op->src[0]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32; // Assuming F32 for now, can be expanded
         case GGML_OP_PAD:
             return op->src[0]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 &&

ggml/src/ggml-opencl/kernels/concat.cl

@@ -106,4 +106,4 @@ kernel void kernel_concat_f32_non_contiguous(
         y_val_ptr = (global float *)(dst_base + (ulong)current_i3*d_nb3 + (ulong)current_i2*d_nb2 + (ulong)current_i1*d_nb1 + (ulong)current_i0*d_nb0);
         *y_val_ptr = *x_val_ptr;
     }
-}
+}

ggml/src/ggml-opencl/kernels/repeat.cl

@@ -36,4 +36,4 @@ kernel void kernel_repeat(
             current_dst_el_ptr[k] = current_src_el_ptr[k];
         }
     }
-}
+}

ggml/src/ggml-opencl/kernels/unary.cl renamed to ggml/src/ggml-opencl/kernels/tanh.cl

@@ -31,7 +31,7 @@ kernel void kernel_tanh_f32_nd(
 
             ulong dst_offset_in_tensor = (ulong)i0*nb10 + (ulong)i1*nb11 + (ulong)i2*nb12 + (ulong)i3*nb13;
             global float *dst_val_ptr = (global float *)((global char *)p_dst_base + off_dst_abs + dst_offset_in_tensor);
-            
+
             *dst_val_ptr = tanh(*src_val_ptr);
         }
     }
@@ -56,8 +56,8 @@ kernel void kernel_tanh_f16_nd(
 
             ulong dst_offset_in_tensor = (ulong)i0*nb10 + (ulong)i1*nb11 + (ulong)i2*nb12 + (ulong)i3*nb13;
             global half *dst_val_ptr = (global half *)((global char *)p_dst_base + off_dst_abs + dst_offset_in_tensor);
-            
+
             *dst_val_ptr = tanh(*src_val_ptr);
         }
     }
-}
+}

ggml/src/ggml-opencl/kernels/tsembd.cl

@@ -20,8 +20,8 @@ kernel void kernel_timestep_embedding(
     local_timesteps_input_ptr = (global const float *)((global char *)p_timesteps + off_timesteps);
     local_dst_output_base_ptr = (global float *)((global char *)p_dst + off_dst);
 
-    local_i = get_global_id(1); 
-    local_j = get_global_id(0); 
+    local_i = get_global_id(1);
+    local_j = get_global_id(0);
 
     local_half_dim = logical_dim / 2;
     local_embed_data_ptr = (global float *)((global char *)local_dst_output_base_ptr + local_i * dst_nb1_bytes);
@@ -34,15 +34,15 @@ kernel void kernel_timestep_embedding(
         return;
     }
 
-    local_timestep_val = local_timesteps_input_ptr[local_i]; 
-    
-    if (local_half_dim == 0) { 
-        local_freq = 1.0f; 
+    local_timestep_val = local_timesteps_input_ptr[local_i];
+
+    if (local_half_dim == 0) {
+        local_freq = 1.0f;
     } else {
         local_freq = exp(-log((float)max_period) * (float)local_j / (float)local_half_dim);
     }
-    
+
     local_arg = local_timestep_val * local_freq;
     local_embed_data_ptr[local_j] = cos(local_arg);
     local_embed_data_ptr[local_j + local_half_dim] = sin(local_arg);
-}
+}

ggml/src/ggml-opencl/kernels/upscale.cl

@@ -47,16 +47,16 @@ kernel void kernel_upscale_bilinear(
     ulong off_src0,
     global void * p_dst,
     ulong off_dst,
-    ulong nb00, 
-    ulong nb01, 
-    ulong nb02, 
-    ulong nb03, 
-    int ne00_src, 
-    int ne01_src, 
-    int ne10_dst, 
+    ulong nb00,
+    ulong nb01,
+    ulong nb02,
+    ulong nb03,
+    int ne00_src,
+    int ne01_src,
+    int ne10_dst,
     int ne11_dst,
     int ne12_dst,
-    int ne13_dst, 
+    int ne13_dst,
     float sf0,
     float sf1,
     float sf2,
@@ -118,4 +118,4 @@ kernel void kernel_upscale_bilinear(
                    val_d * dx * dy;
 
     dst_base[index] = result;
-}
+}