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/scale.cl | 27 +++++++++++++++++++++++++ 1 file changed, 27 insertions(+) create mode 100644 llama.cpp/ggml/src/ggml-opencl/kernels/scale.cl (limited to 'llama.cpp/ggml/src/ggml-opencl/kernels/scale.cl') diff --git a/llama.cpp/ggml/src/ggml-opencl/kernels/scale.cl b/llama.cpp/ggml/src/ggml-opencl/kernels/scale.cl new file mode 100644 index 0000000..17ed97f --- /dev/null +++ b/llama.cpp/ggml/src/ggml-opencl/kernels/scale.cl @@ -0,0 +1,27 @@ +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +kernel void kernel_scale_f32( + global float * src0, + ulong offset0, + global float * dst, + ulong offsetd, + float scale, + float bias +) { + src0 = (global float*)((global char*)src0 + offset0); + dst = (global float*)((global char*)dst + offsetd); + dst[get_global_id(0)] = src0[get_global_id(0)] * scale + bias; +} + +kernel void kernel_scale_f32_4( + global float4 * src0, + ulong offset0, + global float4 * dst, + ulong offsetd, + float scale, + float bias +) { + src0 = (global float4*)((global char*)src0 + offset0); + dst = (global float4*)((global char*)dst + offsetd); + dst[get_global_id(0)] = src0[get_global_id(0)] * scale + bias; +} -- cgit v1.2.3