From b333b06772c89d96aacb5490d6a219fba7c09cc6 Mon Sep 17 00:00:00 2001
From: Mitja Felicijan <mitja.felicijan@gmail.com>
Date: Thu, 12 Feb 2026 20:57:17 +0100
Subject: Engage!

---
 llama.cpp/ggml/src/ggml-opencl/kernels/tri.cl | 32 +++++++++++++++++++++++++++
 1 file changed, 32 insertions(+)
 create mode 100644 llama.cpp/ggml/src/ggml-opencl/kernels/tri.cl

(limited to 'llama.cpp/ggml/src/ggml-opencl/kernels/tri.cl')

diff --git a/llama.cpp/ggml/src/ggml-opencl/kernels/tri.cl b/llama.cpp/ggml/src/ggml-opencl/kernels/tri.cl
new file mode 100644
index 0000000..35cdd54
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-opencl/kernels/tri.cl
@@ -0,0 +1,32 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+//------------------------------------------------------------------------------
+// tri
+//------------------------------------------------------------------------------
+__kernel void kernel_tri_f32(
+        global float * src0,
+        ulong offset0,
+        global float * dst,
+        ulong offsetd,
+        int n,
+        int ne0,
+        int ne1,
+        int tri_type
+) {
+    src0 = (global float*)((global char*)src0 + offset0);
+    dst = (global float*)((global char*)dst + offsetd);
+
+    int idx = get_global_id(0);
+    if (idx >= n) return;
+
+    int i0 = idx % ne0;
+    int i1 = (idx / ne0) % ne1;
+
+    int keep = 0;
+    if (tri_type == 0) keep = (i0 >= i1);
+    else if (tri_type == 1) keep = (i0 >  i1);
+    else if (tri_type == 2) keep = (i0 <= i1);
+    else                    keep = (i0 <  i1);
+
+    dst[idx] = keep ? src0[idx] : 0.0f;
+}
-- 
cgit v1.2.3