summaryrefslogtreecommitdiff
path: root/llama.cpp/ggml/src/ggml-opencl/kernels/sigmoid.cl
diff options
context:
space:
mode:
Diffstat (limited to 'llama.cpp/ggml/src/ggml-opencl/kernels/sigmoid.cl')
-rw-r--r--llama.cpp/ggml/src/ggml-opencl/kernels/sigmoid.cl29
1 files changed, 29 insertions, 0 deletions
diff --git a/llama.cpp/ggml/src/ggml-opencl/kernels/sigmoid.cl b/llama.cpp/ggml/src/ggml-opencl/kernels/sigmoid.cl
new file mode 100644
index 0000000..e3f669d
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-opencl/kernels/sigmoid.cl
@@ -0,0 +1,29 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+//------------------------------------------------------------------------------
+// sigmoid
+//------------------------------------------------------------------------------
+
+kernel void kernel_sigmoid_f32(
+ global 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)] = 1.0f / (1.0f + exp(-src0[get_global_id(0)]));
+}
+
+kernel void kernel_sigmoid_f16(
+ global 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)] = 1.0f / (1.0f + exp(-src0[get_global_id(0)]));
+}