Skip to content

Commit 98bab63

Browse files
ngxsonngxson
authored
ggml : add ggml_scale_bias (#14417)
* ggml : add ggml_scale_bias * ggml_vec_mad1_f32 * add more simd * add CUDA * sycl * vulkan * cann (placeholder) * opencl * will this fix cpu? * fix cuda * suggestions from coderabbit * fix cann compile error * vDSP_vsmsa * rm __ARM_FEATURE_SVE * use memcpy for op params * make code looks more consistent * use scalar for __ARM_FEATURE_SVE * add x param to ggml_vec_mad1_f32
1 parent 26a48ad commit 98bab63

File tree

14 files changed

+139
-38
lines changed

14 files changed

+139
-38
lines changed

‎ggml/include/ggml.h

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1297,6 +1297,19 @@ extern "C" {
12971297
struct ggml_tensor * a,
12981298
float s);
12991299

1300+
// x = s * a + b
1301+
GGML_API struct ggml_tensor * ggml_scale_bias(
1302+
struct ggml_context * ctx,
1303+
struct ggml_tensor * a,
1304+
float s,
1305+
float b);
1306+
1307+
GGML_API struct ggml_tensor * ggml_scale_bias_inplace(
1308+
struct ggml_context * ctx,
1309+
struct ggml_tensor * a,
1310+
float s,
1311+
float b);
1312+
13001313
// b -> view(a,offset,nb1,nb2,3), return modified a
13011314
GGML_API struct ggml_tensor * ggml_set(
13021315
struct ggml_context * ctx,

‎ggml/src/ggml-cann/ggml-cann.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2188,7 +2188,6 @@ static bool ggml_backend_cann_supports_op(ggml_backend_dev_t dev,
21882188
case GGML_OP_MUL:
21892189
case GGML_OP_DIV:
21902190
case GGML_OP_RMS_NORM:
2191-
case GGML_OP_SCALE:
21922191
case GGML_OP_SQR:
21932192
case GGML_OP_SQRT:
21942193
case GGML_OP_CLAMP:
@@ -2210,6 +2209,10 @@ static bool ggml_backend_cann_supports_op(ggml_backend_dev_t dev,
22102209
case GGML_OP_PAD_REFLECT_1D:
22112210
case GGML_OP_COUNT_EQUAL:
22122211
return true;
2212+
case GGML_OP_SCALE:
2213+
float bias;
2214+
memcpy(&bias, (float*)op->op_params + 1, sizeof(float));
2215+
return bias == 0.0f; // TODO: support bias != 0.0f
22132216
case GGML_OP_SOFT_MAX:
22142217
// TODO: support broadcast
22152218
// ref: https://github.com/ggml-org/llama.cpp/pull/14435

‎ggml/src/ggml-cpu/ops.cpp

Lines changed: 20 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -4643,9 +4643,11 @@ static void ggml_compute_forward_scale_f32(
46434643
GGML_ASSERT(ggml_is_contiguous(dst));
46444644
GGML_ASSERT(ggml_are_same_shape(src0, dst));
46454645

4646-
// scale factor
4647-
float v;
4648-
memcpy(&v, dst->op_params, sizeof(float));
4646+
float s; // scale factor
4647+
float b; // bias
4648+
4649+
memcpy(&s, (float *) dst->op_params + 0, sizeof(float));
4650+
memcpy(&b, (float *) dst->op_params + 1, sizeof(float));
46494651

46504652
const int ith = params->ith;
46514653
const int nth = params->nth;
@@ -4664,12 +4666,22 @@ static void ggml_compute_forward_scale_f32(
46644666

46654667
const size_t nb1 = dst->nb[1];
46664668

4667-
for (int i1 = ir0; i1 < ir1; i1++) {
4668-
if (dst->data != src0->data) {
4669-
// src0 is same shape as dst => same indices
4670-
memcpy((char *)dst->data + i1*nb1, (char *)src0->data + i1*nb01, nc * sizeof(float));
4669+
if (b == 0.0f) {
4670+
for (int i1 = ir0; i1 < ir1; i1++) {
4671+
if (dst->data != src0->data) {
4672+
// src0 is same shape as dst => same indices
4673+
// TODO: add x parameter to ggml_vec_scale_f32 and remove this memcpy
4674+
memcpy((char *)dst->data + i1*nb1, (char *)src0->data + i1*nb01, nc * sizeof(float));
4675+
}
4676+
ggml_vec_scale_f32(nc, (float *) ((char *) dst->data + i1*nb1), s);
4677+
}
4678+
} else {
4679+
for (int i1 = ir0; i1 < ir1; i1++) {
4680+
ggml_vec_mad1_f32(nc,
4681+
(float *) ((char *) dst->data + i1*nb1),
4682+
(float *) ((char *) src0->data + i1*nb1),
4683+
s, b);
46714684
}
4672-
ggml_vec_scale_f32(nc, (float *) ((char *) dst->data + i1*nb1), v);
46734685
}
46744686
}
46754687

‎ggml/src/ggml-cpu/vec.h

Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -351,6 +351,45 @@ inline static void ggml_vec_mad_f32_unroll(const int n, const int xs, const int
351351
#endif
352352
}
353353

354+
inline static void ggml_vec_mad1_f32(const int n, float * y, const float * x, const float s, const float b) {
355+
#if defined(GGML_USE_ACCELERATE)
356+
vDSP_vsmsa(x, 1, &s, &b, y, 1, n);
357+
#elif defined(GGML_SIMD)
358+
#if defined(__ARM_FEATURE_SVE)
359+
// scalar ; TODO: Write SVE code
360+
for (int i = 0; i < n; ++i) {
361+
y[i] = x[i]*s + b;
362+
}
363+
#else
364+
const int np = (n & ~(GGML_F32_STEP - 1));
365+
366+
GGML_F32_VEC vs = GGML_F32_VEC_SET1(s);
367+
GGML_F32_VEC vb = GGML_F32_VEC_SET1(b);
368+
369+
GGML_F32_VEC ay[GGML_F32_ARR];
370+
371+
for (int i = 0; i < np; i += GGML_F32_STEP) {
372+
for (int j = 0; j < GGML_F32_ARR; j++) {
373+
ay[j] = GGML_F32_VEC_LOAD(x + i + j*GGML_F32_EPR);
374+
ay[j] = GGML_F32_VEC_FMA(ay[j], vs, vb);
375+
376+
GGML_F32_VEC_STORE(y + i + j*GGML_F32_EPR, ay[j]);
377+
}
378+
}
379+
380+
// leftovers
381+
for (int i = np; i < n; ++i) {
382+
y[i] = x[i]*s + b;
383+
}
384+
#endif
385+
#else
386+
// scalar
387+
for (int i = 0; i < n; ++i) {
388+
y[i] = x[i]*s + b;
389+
}
390+
#endif
391+
}
392+
354393
//inline static void ggml_vec_scale_f32(const int n, float * y, const float v) { for (int i = 0; i < n; ++i) y[i] *= v; }
355394
inline static void ggml_vec_scale_f32(const int n, float * y, const float v) {
356395
#if defined(GGML_USE_ACCELERATE)

‎ggml/src/ggml-cuda/scale.cu

Lines changed: 8 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,18 +1,18 @@
11
#include "scale.cuh"
22

3-
static __global__ void scale_f32(const float * x, float * dst, const float scale, const int k) {
3+
static __global__ void scale_f32(const float * x, float * dst, const float scale, const float bias, const int k) {
44
const int i = blockDim.x*blockIdx.x + threadIdx.x;
55

66
if (i >= k) {
77
return;
88
}
99

10-
dst[i] = scale * x[i];
10+
dst[i] = scale * x[i] + bias;
1111
}
1212

13-
static void scale_f32_cuda(const float * x, float * dst, const float scale, const int k, cudaStream_t stream) {
13+
static void scale_f32_cuda(const float * x, float * dst, const float scale, const float bias, const int k, cudaStream_t stream) {
1414
const int num_blocks = (k + CUDA_SCALE_BLOCK_SIZE - 1) / CUDA_SCALE_BLOCK_SIZE;
15-
scale_f32<<<num_blocks, CUDA_SCALE_BLOCK_SIZE, 0, stream>>>(x, dst, scale, k);
15+
scale_f32<<<num_blocks, CUDA_SCALE_BLOCK_SIZE, 0, stream>>>(x, dst, scale, bias, k);
1616
}
1717

1818
void ggml_cuda_op_scale(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
@@ -25,7 +25,9 @@ void ggml_cuda_op_scale(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
2525
GGML_ASSERT( dst->type == GGML_TYPE_F32);
2626

2727
float scale;
28-
memcpy(&scale, dst->op_params, sizeof(float));
28+
float bias;
29+
memcpy(&scale, (float *) dst->op_params + 0, sizeof(float));
30+
memcpy(&bias, (float *) dst->op_params + 1, sizeof(float));
2931

30-
scale_f32_cuda(src0_d, dst_d, scale, ggml_nelements(src0), stream);
32+
scale_f32_cuda(src0_d, dst_d, scale, bias, ggml_nelements(src0), stream);
3133
}

‎ggml/src/ggml-metal/ggml-metal.m

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2256,7 +2256,9 @@ static bool ggml_metal_encode_node(
22562256
GGML_ASSERT(ggml_is_contiguous(src0));
22572257

22582258
float scale;
2259-
memcpy(&scale, dst->op_params, sizeof(scale));
2259+
float bias;
2260+
memcpy(&scale, ((const int32_t *) dst->op_params) + 0, sizeof(float));
2261+
memcpy(&bias, ((const int32_t *) dst->op_params) + 1, sizeof(float));
22602262

22612263
int64_t n = ggml_nelements(dst);
22622264

@@ -2273,6 +2275,7 @@ static bool ggml_metal_encode_node(
22732275
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:0];
22742276
[encoder setBuffer:id_dst offset:offs_dst atIndex:1];
22752277
[encoder setBytes:&scale length:sizeof(scale) atIndex:2];
2278+
[encoder setBytes:&bias length:sizeof(bias) atIndex:3];
22762279

22772280
[encoder dispatchThreadgroups:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
22782281
} break;

‎ggml/src/ggml-metal/ggml-metal.metal

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1014,16 +1014,18 @@ kernel void kernel_scale(
10141014
device const float * src0,
10151015
device float * dst,
10161016
constant float & scale,
1017+
constant float & bias,
10171018
uint tpig[[thread_position_in_grid]]) {
1018-
dst[tpig] = src0[tpig] * scale;
1019+
dst[tpig] = src0[tpig] * scale + bias;
10191020
}
10201021

10211022
kernel void kernel_scale_4(
10221023
device const float4 * src0,
10231024
device float4 * dst,
10241025
constant float & scale,
1026+
constant float & bias,
10251027
uint tpig[[thread_position_in_grid]]) {
1026-
dst[tpig] = src0[tpig] * scale;
1028+
dst[tpig] = src0[tpig] * scale + bias;
10271029
}
10281030

10291031
kernel void kernel_clamp(

‎ggml/src/ggml-opencl/ggml-opencl.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -5587,7 +5587,9 @@ static void ggml_cl_scale(ggml_backend_t backend, const ggml_tensor * src0, cons
55875587
ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context;
55885588

55895589
float scale;
5590-
memcpy(&scale, dst->op_params, sizeof(scale));
5590+
float bias;
5591+
memcpy(&scale, ((int32_t *) dst->op_params) + 0, sizeof(float));
5592+
memcpy(&bias, ((int32_t *) dst->op_params) + 1, sizeof(float));
55915593

55925594
ggml_tensor_extra_cl * extra0 = (ggml_tensor_extra_cl *)src0->extra;
55935595
ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra;
@@ -5602,6 +5604,7 @@ static void ggml_cl_scale(ggml_backend_t backend, const ggml_tensor * src0, cons
56025604
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extrad->data_device));
56035605
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offsetd));
56045606
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(float), &scale));
5607+
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(float), &bias));
56055608

56065609
int n = ggml_nelements(dst)/4;
56075610

‎ggml/src/ggml-opencl/kernels/scale.cl

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -8,9 +8,10 @@ kernel void kernel_scale(
88
ulong offset0,
99
global float4 * dst,
1010
ulong offsetd,
11-
float scale
11+
float scale,
12+
float bias
1213
) {
1314
src0 = (global float4*)((global char*)src0 + offset0);
1415
dst = (global float4*)((global char*)dst + offsetd);
15-
dst[get_global_id(0)] = src0[get_global_id(0)] * scale;
16+
dst[get_global_id(0)] = src0[get_global_id(0)] * scale + bias;
1617
}

‎ggml/src/ggml-sycl/ggml-sycl.cpp

Lines changed: 8 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1695,7 +1695,7 @@ static void diag_mask_inf_f32(const float * x, float * dst, const int ncols, con
16951695
dst[i] = x[i] - (col > n_past + row % rows_per_channel) * FLT_MAX;
16961696
}
16971697

1698-
static void scale_f32(const float * x, float * dst, const float scale, const int k,
1698+
static void scale_f32(const float * x, float * dst, const float scale, const float bias, const int k,
16991699
const sycl::nd_item<3> &item_ct1) {
17001700
const int i = item_ct1.get_local_range(2) * item_ct1.get_group(2) +
17011701
item_ct1.get_local_id(2);
@@ -1704,7 +1704,7 @@ static void scale_f32(const float * x, float * dst, const float scale, const int
17041704
return;
17051705
}
17061706

1707-
dst[i] = scale * x[i];
1707+
dst[i] = scale * x[i] + bias;
17081708
}
17091709

17101710

@@ -1842,15 +1842,15 @@ static void ggml_mul_mat_vec_nc_f16_f32_sycl(
18421842

18431843

18441844

1845-
static void scale_f32_sycl(const float *x, float *dst, const float scale,
1845+
static void scale_f32_sycl(const float *x, float *dst, const float scale, const float bias,
18461846
const int k, queue_ptr stream) {
18471847
const int num_blocks = (k + SYCL_SCALE_BLOCK_SIZE - 1) / SYCL_SCALE_BLOCK_SIZE;
18481848
stream->parallel_for(
18491849
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) *
18501850
sycl::range<3>(1, 1, SYCL_SCALE_BLOCK_SIZE),
18511851
sycl::range<3>(1, 1, SYCL_SCALE_BLOCK_SIZE)),
18521852
[=](sycl::nd_item<3> item_ct1) {
1853-
scale_f32(x, dst, scale, k, item_ct1);
1853+
scale_f32(x, dst, scale, bias, k, item_ct1);
18541854
});
18551855
}
18561856

@@ -2319,9 +2319,11 @@ inline void ggml_sycl_op_scale(ggml_backend_sycl_context & ctx, ggml_tensor * ds
23192319
float * dst_dd = static_cast<float *>(dst->data);
23202320

23212321
float scale;
2322-
memcpy(&scale, dst->op_params, sizeof(float));
2322+
float bias;
2323+
memcpy(&scale, (float *) dst->op_params + 0, sizeof(float));
2324+
memcpy(&bias, (float *) dst->op_params + 1, sizeof(float));
23232325

2324-
scale_f32_sycl(src0_dd, dst_dd, scale, ggml_nelements(dst->src[0]), main_stream);
2326+
scale_f32_sycl(src0_dd, dst_dd, scale, bias, ggml_nelements(dst->src[0]), main_stream);
23252327
/*
23262328
DPCT1010:87: SYCL uses exceptions to report errors and does not use the
23272329
error codes. The call was replaced with 0. You need to rewrite this code.

0 commit comments

Comments
 (0)