CS348Project/llama.cpp
5
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2025-11-14 11:07:10 +00:00
Code Issues Packages Projects Releases Wiki Activity

opencl

Browse Source
This commit is contained in:
Xuan Son Nguyen
2025-07-08 23:36:47 +02:00
parent 477a97ad87
commit 0e51a0a8b0
2 changed files with 6 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
ggml/src/ggml-opencl/ggml-opencl.cpp
View File

@@ -5586,8 +5586,8 @@ static void ggml_cl_scale(ggml_backend_t backend, const ggml_tensor * src0, cons
ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context; ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context;
float scale; float scale = ((const float *)(dst->op_params))[0];
memcpy(&scale, dst->op_params, sizeof(scale)); float bias = ((const float *)(dst->op_params))[1];
ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *)src0->extra; ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *)src0->extra;
ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra; ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra;
@@ -5602,6 +5602,7 @@ static void ggml_cl_scale(ggml_backend_t backend, const ggml_tensor * src0, cons
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extrad->data_device)); CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extrad->data_device));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offsetd)); CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offsetd));
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(float), &scale)); CL_CHECK(clSetKernelArg(kernel, 4, sizeof(float), &scale));
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(float), &bias));
int n = ggml_nelements(dst)/4; int n = ggml_nelements(dst)/4;

5
ggml/src/ggml-opencl/kernels/scale.cl
View File

@@ -8,9 +8,10 @@ kernel void kernel_scale(
ulong offset0, ulong offset0,
global float4 * dst, global float4 * dst,
ulong offsetd, ulong offsetd,
float scale float scale,
float bias
) { ) {
src0 = (global float4*)((global char*)src0 + offset0); src0 = (global float4*)((global char*)src0 + offset0);
dst = (global float4*)((global char*)dst + offsetd); dst = (global float4*)((global char*)dst + offsetd);
dst[get_global_id(0)] = src0[get_global_id(0)] * scale; dst[get_global_id(0)] = src0[get_global_id(0)] * scale + bias;
} }