Superminaren/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-09-04 02:37:27 +02:00
Code Issues Packages Projects Releases Wiki Activity

rm inplace optimization

Browse Source
This commit is contained in:
Xuan Son Nguyen
2026-08-19 14:07:39 +02:00
parent efee18d262
commit 2dbbb34514
4 changed files with 27 additions and 79 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-opencl/ggml-opencl.cpp
+1 -6
View File
@@ -23705,9 +23705,6 @@ static void ggml_cl_rope(ggml_backend_t backend, const ggml_tensor * src0, const
const int n_ctx_orig = ((int32_t *) dst->op_params)[4];
const int n_offs = ((int32_t *) dst->op_params)[15];
// when dst aliases src0, the channels outside the rotated window already hold the correct data
const int inplace = extrad->data_device == extra0->data_device && offsetd == offset0;
float freq_base;
float freq_scale;
float ext_factor;
@@ -23827,13 +23824,11 @@ static void ggml_cl_rope(ggml_backend_t backend, const ggml_tensor * src0, const
if (is_mrope && !is_vision) {
CL_CHECK(clSetKernelArg(kernel, 34, sizeof(int), &is_imrope));
}
// norm and neox have n_offs and inplace after beta_slow, mrope has them after is_imrope
// norm and neox have n_offs after beta_slow, mrope has it after is_imrope
if (!is_mrope && !is_vision) {
CL_CHECK(clSetKernelArg(kernel, 33, sizeof(int), &n_offs));
CL_CHECK(clSetKernelArg(kernel, 34, sizeof(int), &inplace));
} else if (is_mrope && !is_vision) {
CL_CHECK(clSetKernelArg(kernel, 35, sizeof(int), &n_offs));
CL_CHECK(clSetKernelArg(kernel, 36, sizeof(int), &inplace));
}
size_t global_work_size[] = {(size_t)ne01*nth, (size_t)ne02, (size_t)ne03};
ggml/src/ggml-opencl/kernels/rope.cl
+6 -30
View File
@@ -76,8 +76,7 @@ kernel void kernel_rope_norm_f32(
float attn_factor,
float beta_fast,
float beta_slow,
int n_offs,
int inplace
int n_offs
) {
src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1);
@@ -115,9 +114,6 @@ kernel void kernel_rope_norm_f32(
dst_data[0] = x0*cos_sin_theta.s0 - x1*cos_sin_theta.s1;
dst_data[1] = x0*cos_sin_theta.s1 + x1*cos_sin_theta.s0;
} else {
if (inplace) {
continue;
}
global float * src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
global float * dst_data = (global float *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
@@ -161,8 +157,7 @@ kernel void kernel_rope_norm_f16(
float attn_factor,
float beta_fast,
float beta_slow,
int n_offs,
int inplace
int n_offs
) {
src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1);
@@ -200,9 +195,6 @@ kernel void kernel_rope_norm_f16(
dst_data[0] = x0*cos_sin_theta.s0 - x1*cos_sin_theta.s1;
dst_data[1] = x0*cos_sin_theta.s1 + x1*cos_sin_theta.s0;
} else {
if (inplace) {
continue;
}
global half * src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
global half * dst_data = (global half *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
@@ -246,8 +238,7 @@ kernel void kernel_rope_neox_f32(
float attn_factor,
float beta_fast,
float beta_slow,
int n_offs,
int inplace
int n_offs
) {
src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1);
@@ -285,9 +276,6 @@ kernel void kernel_rope_neox_f32(
dst_data[0] = x0*cos_sin_theta.s0 - x1*cos_sin_theta.s1;
dst_data[n_dims/2] = x0*cos_sin_theta.s1 + x1*cos_sin_theta.s0;
} else {
if (inplace) {
continue;
}
global float * const src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
global float * dst_data = (global float *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
@@ -331,8 +319,7 @@ kernel void kernel_rope_neox_f16(
float attn_factor,
float beta_fast,
float beta_slow,
int n_offs,
int inplace
int n_offs
) {
src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1);
@@ -370,9 +357,6 @@ kernel void kernel_rope_neox_f16(
dst_data[0] = x0*cos_sin_theta.s0 - x1*cos_sin_theta.s1;
dst_data[n_dims/2] = x0*cos_sin_theta.s1 + x1*cos_sin_theta.s0;
} else {
if (inplace) {
continue;
}
global half * const src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
global half * dst_data = (global half *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
@@ -418,8 +402,7 @@ kernel void kernel_rope_multi_f32(
float beta_slow,
int4 sections,
int is_imrope,
int n_offs,
int inplace
int n_offs
) {
src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1);
@@ -487,9 +470,6 @@ kernel void kernel_rope_multi_f32(
dst_data[0] = x0*cos_sin_theta.s0 - x1*cos_sin_theta.s1;
dst_data[n_dims/2] = x0*cos_sin_theta.s1 + x1*cos_sin_theta.s0;
} else {
if (inplace) {
continue;
}
global float * const src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
global float * dst_data = (global float *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
@@ -535,8 +515,7 @@ kernel void kernel_rope_multi_f16(
float beta_slow,
int4 sections,
int is_imrope,
int n_offs,
int inplace
int n_offs
) {
src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1);
@@ -604,9 +583,6 @@ kernel void kernel_rope_multi_f16(
dst_data[0] = x0*cos_sin_theta.s0 - x1*cos_sin_theta.s1;
dst_data[n_dims/2] = x0*cos_sin_theta.s1 + x1*cos_sin_theta.s0;
} else {
if (inplace) {
continue;
}
global half * const src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
global half * dst_data = (global half *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
ggml/src/ggml-sycl/rope.cpp
+20 -39
View File
@@ -45,8 +45,7 @@ static void rope_norm(const T *x, D *dst, const int ne00, const int ne01,
const float freq_scale, const float ext_factor,
const float attn_factor, const rope_corr_dims corr_dims,
const float theta_scale, const float *freq_factors,
const int64_t *row_indices, const int set_rows_stride,
const bool inplace) {
const int64_t *row_indices, const int set_rows_stride) {
auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
const int i0 = 2 * (item_ct1.get_local_range(1) * item_ct1.get_group(1) +
item_ct1.get_local_id(1));
@@ -80,9 +79,6 @@ static void rope_norm(const T *x, D *dst, const int ne00, const int ne01,
}
};
if (i0 < n_offs || i0 >= n_offs + n_dims) {
if (inplace) {
return;
}
store_coaelsced(x[ix + 0], x[ix + 1]);
return;
}
@@ -114,8 +110,7 @@ static void rope_neox(const T *x, D *dst, const int ne00, const int ne01,
const float freq_scale, const float ext_factor,
const float attn_factor, const rope_corr_dims corr_dims,
const float theta_scale, const float *freq_factors,
const int64_t *row_indices, const int set_rows_stride,
const bool inplace) {
const int64_t *row_indices, const int set_rows_stride) {
auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
const int i0 = 2 * (item_ct1.get_local_range(1) * item_ct1.get_group(1) +
item_ct1.get_local_id(1));
@@ -140,9 +135,6 @@ static void rope_neox(const T *x, D *dst, const int ne00, const int ne01,
}
if (i0 < n_offs || i0 >= n_offs + n_dims) {
if (inplace) {
return;
}
dst[idst + i0 / 2 + 0] = ggml_sycl_cast<D>(x[ix + i0 / 2 + 0]);
dst[idst + i0 / 2 + 1] = ggml_sycl_cast<D>(x[ix + i0 / 2 + 1]);
@@ -177,8 +169,7 @@ static void rope_multi(const T *x, T *dst, const int ne00, const int ne01,
const float freq_scale, const float ext_factor,
const float attn_factor, const rope_corr_dims corr_dims,
const float theta_scale, const float *freq_factors,
const mrope_sections sections, const bool is_imrope,
const bool inplace) {
const mrope_sections sections, const bool is_imrope) {
auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
const int i0 = 2 * (item_ct1.get_local_range(1) * item_ct1.get_group(1) +
item_ct1.get_local_id(1));
@@ -198,9 +189,6 @@ static void rope_multi(const T *x, T *dst, const int ne00, const int ne01,
const int ix = i0 / 2 + i1 * s01 + i2 * s02 + i3 * s03;
if (i0 < n_offs || i0 >= n_offs + n_dims) {
if (inplace) {
return;
}
dst[idst + i0 / 2 + 0] = x[ix + i0 / 2 + 0];
dst[idst + i0 / 2 + 1] = x[ix + i0 / 2 + 1];
@@ -317,7 +305,7 @@ rope_norm_sycl(const T *x, D *dst, const int ne00, const int ne01,
const float freq_base, const float ext_factor,
const float attn_factor, const rope_corr_dims corr_dims,
const float *freq_factors, const int64_t *row_indices,
const int set_rows_stride, const bool inplace, dpct::queue_ptr stream) {
const int set_rows_stride, dpct::queue_ptr stream) {
GGML_ASSERT(ne00 % 2 == 0);
const dpct::dim3 block_dims(1, SYCL_ROPE_BLOCK_SIZE, 1);
const int n_blocks_x =
@@ -334,8 +322,7 @@ rope_norm_sycl(const T *x, D *dst, const int ne00, const int ne01,
rope_norm<forward, false>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
theta_scale, freq_factors, row_indices, set_rows_stride,
inplace);
theta_scale, freq_factors, row_indices, set_rows_stride);
});
} else {
stream->parallel_for(
@@ -345,8 +332,7 @@ rope_norm_sycl(const T *x, D *dst, const int ne00, const int ne01,
rope_norm<forward, true>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
theta_scale, freq_factors, row_indices, set_rows_stride,
inplace);
theta_scale, freq_factors, row_indices, set_rows_stride);
});
}
}
@@ -360,7 +346,7 @@ rope_neox_sycl(const T *x, D *dst, const int ne00, const int ne01,
const float freq_base, const float ext_factor,
const float attn_factor, const rope_corr_dims corr_dims,
const float *freq_factors, const int64_t *row_indices,
const int set_rows_stride, const bool inplace, dpct::queue_ptr stream) {
const int set_rows_stride, dpct::queue_ptr stream) {
GGML_ASSERT(ne00 % 2 == 0);
const dpct::dim3 block_dims(1, SYCL_ROPE_BLOCK_SIZE, 1);
const int n_blocks_x =
@@ -377,8 +363,7 @@ rope_neox_sycl(const T *x, D *dst, const int ne00, const int ne01,
rope_neox<forward, false>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
theta_scale, freq_factors, row_indices, set_rows_stride,
inplace);
theta_scale, freq_factors, row_indices, set_rows_stride);
});
} else {
stream->parallel_for(
@@ -388,8 +373,7 @@ rope_neox_sycl(const T *x, D *dst, const int ne00, const int ne01,
rope_neox<forward, true>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
theta_scale, freq_factors, row_indices, set_rows_stride,
inplace);
theta_scale, freq_factors, row_indices, set_rows_stride);
});
}
}
@@ -403,7 +387,7 @@ rope_multi_sycl(const T *x, T *dst, const int ne00, const int ne01,
const float freq_base, const float ext_factor,
const float attn_factor, const rope_corr_dims corr_dims,
const float *freq_factors, const mrope_sections sections,
const bool is_imrope, const bool inplace, dpct::queue_ptr stream) {
const bool is_imrope, dpct::queue_ptr stream) {
GGML_ASSERT(ne00 % 2 == 0);
const dpct::dim3 block_dims(1, SYCL_ROPE_BLOCK_SIZE, 1);
const int n_blocks_x =
@@ -420,7 +404,7 @@ rope_multi_sycl(const T *x, T *dst, const int ne00, const int ne01,
rope_multi<forward, false, T>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
theta_scale, freq_factors, sections, is_imrope, inplace);
theta_scale, freq_factors, sections, is_imrope);
});
} else {
stream->parallel_for(
@@ -430,7 +414,7 @@ rope_multi_sycl(const T *x, T *dst, const int ne00, const int ne01,
rope_multi<forward, true, T>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
theta_scale, freq_factors, sections, is_imrope, inplace);
theta_scale, freq_factors, sections, is_imrope);
});
}
}
@@ -524,9 +508,6 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst,
const int n_offs = ((int32_t *)dst->op_params)[15];
mrope_sections sections;
// when dst aliases src0, the channels outside the rotated window already hold the correct data
const bool inplace = dst_d == src0->data;
float freq_base;
float freq_scale;
float ext_factor;
@@ -576,19 +557,19 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst,
(const float *)src0_d, (float *)dst_d, ne00, ne01, ne02, s01,
s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
set_rows_stride, inplace, stream);
set_rows_stride, stream);
} else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) {
rope_neox_sycl<forward, float, sycl::half>(
(const float *)src0_d, (sycl::half *)dst_d, ne00, ne01, ne02,
s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale,
freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
row_indices, set_rows_stride, inplace, stream);
row_indices, set_rows_stride, stream);
} else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) {
rope_neox_sycl<forward, sycl::half, sycl::half>(
(const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01,
ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale,
freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
row_indices, set_rows_stride, inplace, stream);
row_indices, set_rows_stride, stream);
} else {
GGML_ABORT("Fatal error: Tensor type unsupported!");
}
@@ -599,13 +580,13 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst,
ne00, ne01, ne02, s01, s02, s03, s1, s2,
s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
ext_factor, attn_factor, corr_dims,
freq_factors, sections, is_imrope, inplace, stream);
freq_factors, sections, is_imrope, stream);
} else if (src0->type == GGML_TYPE_F16) {
rope_multi_sycl<forward>(
(const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01,
ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale,
freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
sections, is_imrope, inplace, stream);
sections, is_imrope, stream);
} else {
GGML_ABORT("Fatal error: Tensor type unsupported!");
}
@@ -633,19 +614,19 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst,
(const float *)src0_d, (float *)dst_d, ne00, ne01, ne02, s01,
s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
set_rows_stride, inplace, stream);
set_rows_stride, stream);
} else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) {
rope_norm_sycl<forward, float, sycl::half>(
(const float *)src0_d, (sycl::half *)dst_d, ne00, ne01, ne02,
s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale,
freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
row_indices, set_rows_stride, inplace, stream);
row_indices, set_rows_stride, stream);
} else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) {
rope_norm_sycl<forward, sycl::half, sycl::half>(
(const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01,
ne02, s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale,
freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
row_indices, set_rows_stride, inplace, stream);
row_indices, set_rows_stride, stream);
} else {
GGML_ABORT("Fatal error: Tensor type unsupported!");
}
ggml/src/ggml-webgpu/wgsl-shaders/rope.wgsl
-4
View File
@@ -168,14 +168,10 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
let i_dst_row = params.offset_dst + i3 * params.stride_dst3 + i2 * params.stride_dst2 + i1 * params.stride_dst1;
if ((i0 < params.n_offs || i0 >= params.n_offs + params.n_dims) && !is_vision) {
#ifdef INPLACE
return;
#else
let i_src = i_src_row + i0;
let i_dst = i_dst_row + i0;
rotate(i_dst, i_dst + 1, f32(src0[i_src]), f32(src0[i_src + 1]));
return;
#endif
}
let iw = i0 - params.n_offs; // relative idx