Commit 31424b7cf for llama.cpp
commit 31424b7cf14604ca65a9195123da1413f31c4472
Author: shaofeiqi <shaoqi@qti.qualcomm.com>
Date: Fri Nov 6 21:44:05 2026 -0800
opencl: add arange (#31829)
Assisted-by: Claude Opus 5.5
Co-authored-by: Li He <lih@qti.qualcomm.com>
diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt
index 58a6d438e..0ef7f8fe6 100644
--- a/ggml/src/ggml-opencl/CMakeLists.txt
+++ b/ggml/src/ggml-opencl/CMakeLists.txt
@@ -71,6 +71,7 @@ set(GGML_OPENCL_KERNELS
lightning_indexer
tri
fill
+ arange
fwht
clamp
cpy
diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index e0ee5518b..1a4c03519 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -856,6 +856,7 @@ struct ggml_backend_opencl_context {
cl_kernel kernel_sigmoid_f32_nc, kernel_sigmoid_f16_nc;
cl_kernel kernel_tri;
cl_kernel kernel_fill;
+ cl_kernel kernel_arange;
cl_kernel kernel_clamp, kernel_clamp_nc;
cl_kernel kernel_geglu, kernel_reglu, kernel_swiglu, kernel_swiglu_oai, kernel_swiglu_clamp, kernel_geglu_erf,
kernel_geglu_quick, kernel_geglu_f16, kernel_reglu_f16, kernel_swiglu_f16, kernel_swiglu_clamp_f16,
@@ -2048,6 +2049,24 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
CL_CHECK(clReleaseProgram(prog));
}
+ // arange
+ {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+ const std::string kernel_src {
+ #include "arange.cl.h"
+ };
+#else
+ const std::string kernel_src = read_file("arange.cl");
+#endif
+ cl_program prog =
+ build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts);
+
+ CL_CHECK((backend_ctx->kernel_arange = clCreateKernel(prog, "kernel_arange_f32", &err), err));
+ GGML_LOG_CONT(".");
+
+ CL_CHECK(clReleaseProgram(prog));
+ }
+
// clamp
{
#ifdef GGML_OPENCL_EMBED_KERNELS
@@ -11792,6 +11811,8 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
return op->type == GGML_TYPE_F32 && ggml_is_contiguous(op);
case GGML_OP_FILL:
return op->type == GGML_TYPE_F32 && ggml_is_contiguous(op);
+ case GGML_OP_ARANGE:
+ return op->type == GGML_TYPE_F32;
case GGML_OP_CLAMP:
return op->src[0]->type == GGML_TYPE_F32;
case GGML_OP_SOFT_MAX:
@@ -19593,6 +19614,40 @@ static void ggml_cl_fill(ggml_backend_t backend, const ggml_tensor * src0, const
backend_ctx->enqueue_ndrange_kernel(kernel, 1, global_work_size, local_work_size, dst);
}
+static void ggml_cl_arange(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+ GGML_ASSERT(dst);
+ GGML_ASSERT(dst->extra);
+ GGML_ASSERT(dst->type == GGML_TYPE_F32);
+
+ UNUSED(src0);
+ UNUSED(src1);
+
+ ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context;
+
+ ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra;
+ cl_ulong offsetd = extrad->offset + dst->view_offs;
+
+ float start;
+ float step;
+ memcpy(&start, (float *) dst->op_params + 0, sizeof(float));
+ memcpy(&step, (float *) dst->op_params + 2, sizeof(float));
+
+ const int n = dst->ne[0];
+
+ cl_kernel kernel = backend_ctx->kernel_arange;
+
+ CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extrad->data_device));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offsetd));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), &n));
+ CL_CHECK(clSetKernelArg(kernel, 3, sizeof(float), &start));
+ CL_CHECK(clSetKernelArg(kernel, 4, sizeof(float), &step));
+
+ size_t local_work_size[1] = { 64 };
+ size_t global_work_size[1] = { ((size_t)n + local_work_size[0] - 1) / local_work_size[0] * local_work_size[0] };
+
+ backend_ctx->enqueue_ndrange_kernel(kernel, 1, global_work_size, local_work_size, dst);
+}
+
static void ggml_cl_clamp(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
GGML_ASSERT(src0);
GGML_ASSERT(src0->extra);
@@ -38649,6 +38704,12 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor
}
func = ggml_cl_fill;
break;
+ case GGML_OP_ARANGE:
+ if (!any_on_device) {
+ return false;
+ }
+ func = ggml_cl_arange;
+ break;
case GGML_OP_CLAMP:
if (!any_on_device) {
return false;
diff --git a/ggml/src/ggml-opencl/kernels/arange.cl b/ggml/src/ggml-opencl/kernels/arange.cl
new file mode 100644
index 000000000..53dc383e6
--- /dev/null
+++ b/ggml/src/ggml-opencl/kernels/arange.cl
@@ -0,0 +1,14 @@
+__kernel void kernel_arange_f32(
+ __global float * dst,
+ ulong offsetd,
+ int n,
+ float start,
+ float step
+) {
+ dst = (global float*)((global char*)dst + offsetd);
+
+ const int i = get_global_id(0);
+ if (i < n) {
+ dst[i] = start + step * i;
+ }
+}