Engage!

1 files changed, 20 insertions, 0 deletions

diff --git a/llama.cpp/ggml/src/ggml-opencl/kernels/clamp.cl b/llama.cpp/ggml/src/ggml-opencl/kernels/clamp.cl
new file mode 100644
index 0000000..ae60324
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-opencl/kernels/clamp.cl
@@ -0,0 +1,20 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+//------------------------------------------------------------------------------
+// clamp
+//------------------------------------------------------------------------------
+kernel void kernel_clamp(
+        global float * src0,
+        ulong offset0,
+        global float * dst,
+        ulong offsetd,
+        float min,
+        float max
+) {
+    src0 = (global float*)((global char*)src0 + offset0);
+    dst = (global float*)((global char*)dst + offsetd);
+
+    dst[get_global_id(0)] = src0[get_global_id(0)] < min ?
+        min :
+        (src0[get_global_id(0)] > max ? max : src0[get_global_id(0)]);
+}