From b333b06772c89d96aacb5490d6a219fba7c09cc6 Mon Sep 17 00:00:00 2001 From: Mitja Felicijan Date: Thu, 12 Feb 2026 20:57:17 +0100 Subject: Engage! --- llama.cpp/ggml/src/ggml-opencl/kernels/silu.cl | 30 ++++++++++++++++++++++++++ 1 file changed, 30 insertions(+) create mode 100644 llama.cpp/ggml/src/ggml-opencl/kernels/silu.cl (limited to 'llama.cpp/ggml/src/ggml-opencl/kernels/silu.cl') diff --git a/llama.cpp/ggml/src/ggml-opencl/kernels/silu.cl b/llama.cpp/ggml/src/ggml-opencl/kernels/silu.cl new file mode 100644 index 0000000..1d95e1b --- /dev/null +++ b/llama.cpp/ggml/src/ggml-opencl/kernels/silu.cl @@ -0,0 +1,30 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +//------------------------------------------------------------------------------ +// silu +//------------------------------------------------------------------------------ +kernel void kernel_silu( + global float * src0, + ulong offset0, + global float * dst, + ulong offsetd +) { + src0 = (global float*)((global char*)src0 + offset0); + dst = (global float*)((global char*)dst + offsetd); + + float x = src0[get_global_id(0)]; + dst[get_global_id(0)] = x / (1.0f + exp(-x)); +} + +kernel void kernel_silu_4( + global float4 * src0, + ulong offset0, + global float4 * dst, + ulong offsetd +) { + src0 = (global float4*)((global char*)src0 + offset0); + dst = (global float4*)((global char*)dst + offsetd); + + float4 x = src0[get_global_id(0)]; + dst[get_global_id(0)] = x / (1.0f + exp(-x)); +} -- cgit v1.2.3