Skip to content

ggml: support ggml_rope_set_offset on opencl, sycl, wgpu, hexagon #27345

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
ngxson merged 2 commits into from
Aug 20, 2026
Merged

ggml: support ggml_rope_set_offset on opencl, sycl, wgpu, hexagon #27345

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
efee18d
ggml: support ggml_rope_set_offset on opencl, sycl, wgpu, hexagon
ngxson Aug 18, 2026
2dbbb34
rm inplace optimization
ngxson Aug 18, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 3 additions & 2 deletions ggml/src/ggml-hexagon/ggml-hexagon.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3180,8 +3180,9 @@ static bool ggml_hexagon_supported_argsort(const struct ggml_hexagon_session * s
static bool ggml_hexagon_supported_rope(const struct ggml_hexagon_session * sess, const struct ggml_tensor * op) {
const int32_t * op_params = &op->op_params[0];

if (op_params[15] != 0) {
return false; // FIXME: support ggml_rope_set_offset
// ggml_rope_set_offset: HVX kernels need a VLEN-aligned window start (32 f32 elems)
if (op_params[15] % 32 != 0) {
return false;
}

int mode = op_params[2];
Expand Down
22 changes: 16 additions & 6 deletions ggml/src/ggml-hexagon/htp/rope-ops.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,6 +53,7 @@

struct htp_rope_context {
int32_t n_dims;
int32_t n_offs;
int32_t mode;
int32_t n_ctx_orig;
int32_t sections[4];
Expand Down Expand Up @@ -405,32 +406,40 @@ static inline void hvx_rope_f32_aa(float * restrict dst, const float * restrict

static void inline rope_basic_f32(struct htp_rope_context * rctx, uint8_t * restrict dst, uint8_t * restrict src,
uint32_t nr, uint32_t ne0, const float * restrict theta_cache) {
const uint32_t n_offs = rctx->n_offs; // VLEN-aligned (enforced by supports_op)
#pragma unroll(4)
for (uint32_t i = 0; i < nr; i++) {
float * d = (float *) (dst + i * rctx->dst_row_size_aligned);
float * s = (float *) (src + i * rctx->src0_row_size_aligned);

hvx_rope_f32_aa(d, s, rctx->n_dims, theta_cache);
hvx_rope_f32_aa(d + n_offs, s + n_offs, rctx->n_dims, theta_cache);

// fill the remain channels with data from src tensor
if (rctx->n_dims < ne0) {
hvx_copy_f32_uu((uint8_t *)(d + rctx->n_dims), (uint8_t *)(s + rctx->n_dims), ne0 - rctx->n_dims);
if (n_offs > 0) {
hvx_copy_f32_uu((uint8_t *) d, (uint8_t *) s, n_offs);
}
if (n_offs + rctx->n_dims < ne0) {
hvx_copy_f32_uu((uint8_t *)(d + n_offs + rctx->n_dims), (uint8_t *)(s + n_offs + rctx->n_dims), ne0 - n_offs - rctx->n_dims);
}
}
}

static void inline rope_neox_f32(struct htp_rope_context * rctx, uint8_t * restrict dst, uint8_t * restrict src,
uint32_t nr, uint32_t ne0, const float * restrict theta_cache) {
const uint32_t n_offs = rctx->n_offs; // VLEN-aligned (enforced by supports_op)
#pragma unroll(4)
for (uint32_t i = 0; i < nr; i++) {
float * d = (float *) (dst + i * rctx->dst_row_size_aligned);
float * s = (float *) (src + i * rctx->src0_row_size_aligned);

hvx_rope_neox_f32_aa(d, s, rctx->n_dims, theta_cache);
hvx_rope_neox_f32_aa(d + n_offs, s + n_offs, rctx->n_dims, theta_cache);

// fill the remain channels with data from src tensor
if (rctx->n_dims < ne0) {
hvx_copy_f32_uu((uint8_t *)(d + rctx->n_dims), (uint8_t *)(s + rctx->n_dims), ne0 - rctx->n_dims);
if (n_offs > 0) {
hvx_copy_f32_uu((uint8_t *) d, (uint8_t *) s, n_offs);
}
if (n_offs + rctx->n_dims < ne0) {
hvx_copy_f32_uu((uint8_t *)(d + n_offs + rctx->n_dims), (uint8_t *)(s + n_offs + rctx->n_dims), ne0 - n_offs - rctx->n_dims);
}
}
}
Expand Down Expand Up @@ -673,6 +682,7 @@ static int execute_op_rope_f32(struct htp_ops_context * octx) {
rctx.n_dims = ((const int32_t *) op_params)[1];
rctx.mode = ((const int32_t *) op_params)[2];
rctx.n_ctx_orig = ((const int32_t *) op_params)[4];
rctx.n_offs = ((const int32_t *) op_params)[15];

memcpy(&rctx.freq_base, (int32_t *) op_params + 5, sizeof(float));
memcpy(&rctx.freq_scale, (int32_t *) op_params + 6, sizeof(float));
Expand Down
11 changes: 8 additions & 3 deletions ggml/src/ggml-opencl/ggml-opencl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7376,9 +7376,6 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
case GGML_OP_DIAG_MASK_INF:
return op->ne[3] == 1;
case GGML_OP_ROPE: {
if (((const int32_t *) op->op_params)[15] != 0) {
return false; // FIXME: support ggml_rope_set_offset
}
const int mode = ((const int32_t *) op->op_params)[2];
const bool is_mrope = mode & GGML_ROPE_TYPE_MROPE;
const bool is_vision = mode == GGML_ROPE_TYPE_VISION;
Expand Down Expand Up @@ -23706,6 +23703,7 @@ static void ggml_cl_rope(ggml_backend_t backend, const ggml_tensor * src0, const
const int n_dims = ((int *) dst->op_params)[1];
const int mode = ((int *) dst->op_params)[2];
const int n_ctx_orig = ((int32_t *) dst->op_params)[4];
const int n_offs = ((int32_t *) dst->op_params)[15];

float freq_base;
float freq_scale;
Expand Down Expand Up @@ -23734,6 +23732,7 @@ static void ggml_cl_rope(ggml_backend_t backend, const ggml_tensor * src0, const

if (is_vision) {
GGML_ASSERT(n_dims == ne00/2);
GGML_ASSERT(n_offs == 0); // offset not supported for vision, as the rotated pairs span the whole row
}

cl_kernel kernel;
Expand Down Expand Up @@ -23825,6 +23824,12 @@ 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 after beta_slow, mrope has it after is_imrope
if (!is_mrope && !is_vision) {
CL_CHECK(clSetKernelArg(kernel, 33, sizeof(int), &n_offs));
} else if (is_mrope && !is_vision) {
CL_CHECK(clSetKernelArg(kernel, 35, sizeof(int), &n_offs));
}

size_t global_work_size[] = {(size_t)ne01*nth, (size_t)ne02, (size_t)ne03};
size_t local_work_size[] = {(size_t)nth, 1, 1};
Expand Down
92 changes: 52 additions & 40 deletions ggml/src/ggml-opencl/kernels/rope.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -75,7 +75,8 @@ kernel void kernel_rope_norm_f32(
float ext_factor,
float attn_factor,
float beta_fast,
float beta_slow
float beta_slow,
int n_offs
) {
src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1);
Expand All @@ -94,14 +95,15 @@ kernel void kernel_rope_norm_f32(
float inv_ndims = -1.f/n_dims;

for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) {
if (i0 < n_dims) {
int ic = i0/2;
if (i0 >= n_offs && i0 < n_offs + n_dims) {
int iw = i0 - n_offs; // relative idx
int ic = iw/2;

float theta = theta_base * pow(freq_base, inv_ndims*i0);
float theta = theta_base * pow(freq_base, inv_ndims*iw);

float freq_factor = src2 != src0 ? src2[ic] : 1.0f;

float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor);
float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor);

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);
Expand Down Expand Up @@ -154,7 +156,8 @@ kernel void kernel_rope_norm_f16(
float ext_factor,
float attn_factor,
float beta_fast,
float beta_slow
float beta_slow,
int n_offs
) {
src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1);
Expand All @@ -173,14 +176,15 @@ kernel void kernel_rope_norm_f16(
float inv_ndims = -1.f/n_dims;

for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) {
if (i0 < n_dims) {
int ic = i0/2;
if (i0 >= n_offs && i0 < n_offs + n_dims) {
int iw = i0 - n_offs; // relative idx
int ic = iw/2;

float theta = theta_base * pow(freq_base, inv_ndims*i0);
float theta = theta_base * pow(freq_base, inv_ndims*iw);

float freq_factor = src2 != src0 ? src2[ic] : 1.0f;

float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor);
float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor);

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);
Expand Down Expand Up @@ -233,7 +237,8 @@ kernel void kernel_rope_neox_f32(
float ext_factor,
float attn_factor,
float beta_fast,
float beta_slow
float beta_slow,
int n_offs
) {
src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1);
Expand All @@ -252,17 +257,18 @@ kernel void kernel_rope_neox_f32(
float inv_ndims = -1.f/n_dims;

for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) {
if (i0 < n_dims) {
int ic = i0/2;
if (i0 >= n_offs && i0 < n_offs + n_dims) {
int iw = i0 - n_offs; // relative idx
int ic = iw/2;

const float theta = theta_base * pow(freq_base, inv_ndims*i0);
const float theta = theta_base * pow(freq_base, inv_ndims*iw);

const float freq_factor = src2 != src0 ? src2[ic] : 1.0f;

float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor);
float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor);

global float * src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00);
global float * dst_data = (global float *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0);
global float * src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + (n_offs + ic)*nb00);
global float * dst_data = (global float *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + (n_offs + ic)*nb0);

const float x0 = src[0];
const float x1 = src[n_dims/2];
Expand Down Expand Up @@ -312,7 +318,8 @@ kernel void kernel_rope_neox_f16(
float ext_factor,
float attn_factor,
float beta_fast,
float beta_slow
float beta_slow,
int n_offs
) {
src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1);
Expand All @@ -331,17 +338,18 @@ kernel void kernel_rope_neox_f16(
float inv_ndims = -1.f/n_dims;

for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) {
if (i0 < n_dims) {
int ic = i0/2;
if (i0 >= n_offs && i0 < n_offs + n_dims) {
int iw = i0 - n_offs; // relative idx
int ic = iw/2;

const float theta = theta_base * pow(freq_base, inv_ndims*i0);
const float theta = theta_base * pow(freq_base, inv_ndims*iw);

const float freq_factor = src2 != src0 ? src2[ic] : 1.0f;

float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor);
float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor);

global half * src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00);
global half * dst_data = (global half *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0);
global half * src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + (n_offs + ic)*nb00);
global half * dst_data = (global half *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + (n_offs + ic)*nb0);

const float x0 = src[0];
const float x1 = src[n_dims/2];
Expand Down Expand Up @@ -393,7 +401,8 @@ kernel void kernel_rope_multi_f32(
float beta_fast,
float beta_slow,
int4 sections,
int is_imrope
int is_imrope,
int n_offs
) {
src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1);
Expand All @@ -414,10 +423,11 @@ kernel void kernel_rope_multi_f32(
float inv_ndims = -1.f/n_dims;

for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) {
if (i0 < n_dims) {
int ic = i0/2;
if (i0 >= n_offs && i0 < n_offs + n_dims) {
int iw = i0 - n_offs; // relative idx
int ic = iw/2;

const int sector = (i0 / 2) % sect_dims;
const int sector = ic % sect_dims;
float theta_base = 0.0f;

if (is_imrope) {
Expand Down Expand Up @@ -445,14 +455,14 @@ kernel void kernel_rope_multi_f32(
}
}

const float theta = theta_base * pow(freq_base, inv_ndims*i0);
const float theta = theta_base * pow(freq_base, inv_ndims*iw);

const float freq_factor = src2 != src0 ? src2[ic] : 1.0f;

float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor);
float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor);

global float * src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00);
global float * dst_data = (global float *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0);
global float * src = (global float *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + (n_offs + ic)*nb00);
global float * dst_data = (global float *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + (n_offs + ic)*nb0);

const float x0 = src[0];
const float x1 = src[n_dims/2];
Expand Down Expand Up @@ -504,7 +514,8 @@ kernel void kernel_rope_multi_f16(
float beta_fast,
float beta_slow,
int4 sections,
int is_imrope
int is_imrope,
int n_offs
) {
src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1);
Expand All @@ -525,10 +536,11 @@ kernel void kernel_rope_multi_f16(
float inv_ndims = -1.f/n_dims;

for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) {
if (i0 < n_dims) {
int ic = i0/2;
if (i0 >= n_offs && i0 < n_offs + n_dims) {
int iw = i0 - n_offs; // relative idx
int ic = iw/2;

const int sector = (i0 / 2) % sect_dims;
const int sector = ic % sect_dims;
float theta_base = 0.0f;

if (is_imrope) {
Expand Down Expand Up @@ -556,14 +568,14 @@ kernel void kernel_rope_multi_f16(
}
}

const float theta = theta_base * pow(freq_base, inv_ndims*i0);
const float theta = theta_base * pow(freq_base, inv_ndims*iw);

const float freq_factor = src2 != src0 ? src2[ic] : 1.0f;

float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor);
float2 cos_sin_theta = rope_yarn(theta/freq_factor, freq_scale, corr_dims, iw, ext_factor, attn_factor);

global half * src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + ic*nb00);
global half * dst_data = (global half *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + ic*nb0);
global half * src = (global half *)((global char *) src0 + i3*nb03 + i2*nb02 + i1*nb01 + (n_offs + ic)*nb00);
global half * dst_data = (global half *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + (n_offs + ic)*nb0);

const float x0 = src[0];
const float x1 = src[n_dims/2];
Expand Down
2 changes: 0 additions & 2 deletions ggml/src/ggml-sycl/ggml-sycl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6235,8 +6235,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
}
case GGML_OP_ROPE:
case GGML_OP_ROPE_BACK:
// FIXME: support ggml_rope_set_offset
return ((const int32_t *) op->op_params)[15] == 0;
case GGML_OP_IM2COL:
case GGML_OP_IM2COL_3D:
case GGML_OP_UPSCALE:
Expand Down
Loading
Loading