Lumpiasty/llama.cpp
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

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

Browse Source 
* ggml: support ggml_rope_set_offset on opencl, sycl, wgpu, hexagon

* rm inplace optimization
This commit is contained in:
Xuan-Son Nguyen
2026-08-21 00:36:57 +02:00
committed by GitHub
parent 0e1d9185c5
commit 749f688fca
8 changed files with 152 additions and 112 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-hexagon/ggml-hexagon.cpp
+3 -2
View File
@@ -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) { 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]; const int32_t * op_params = &op->op_params[0];
if (op_params[15] != 0) { // ggml_rope_set_offset: HVX kernels need a VLEN-aligned window start (32 f32 elems)
return false; // FIXME: support ggml_rope_set_offset if (op_params[15] % 32 != 0) {
return false;
} }
int mode = op_params[2]; int mode = op_params[2];
ggml/src/ggml-hexagon/htp/rope-ops.c
+16 -6
View File
@@ -53,6 +53,7 @@
struct htp_rope_context { struct htp_rope_context {
int32_t n_dims; int32_t n_dims;
int32_t n_offs;
int32_t mode; int32_t mode;
int32_t n_ctx_orig; int32_t n_ctx_orig;
int32_t sections[4]; int32_t sections[4];
@@ -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, 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) { 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) #pragma unroll(4)
for (uint32_t i = 0; i < nr; i++) { for (uint32_t i = 0; i < nr; i++) {
float * d = (float *) (dst + i * rctx->dst_row_size_aligned); float * d = (float *) (dst + i * rctx->dst_row_size_aligned);
float * s = (float *) (src + i * rctx->src0_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 // fill the remain channels with data from src tensor
if (rctx->n_dims < ne0) { if (n_offs > 0) {
hvx_copy_f32_uu((uint8_t *)(d + rctx->n_dims), (uint8_t *)(s + rctx->n_dims), ne0 - rctx->n_dims); 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, 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) { 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) #pragma unroll(4)
for (uint32_t i = 0; i < nr; i++) { for (uint32_t i = 0; i < nr; i++) {
float * d = (float *) (dst + i * rctx->dst_row_size_aligned); float * d = (float *) (dst + i * rctx->dst_row_size_aligned);
float * s = (float *) (src + i * rctx->src0_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 // fill the remain channels with data from src tensor
if (rctx->n_dims < ne0) { if (n_offs > 0) {
hvx_copy_f32_uu((uint8_t *)(d + rctx->n_dims), (uint8_t *)(s + rctx->n_dims), ne0 - rctx->n_dims); 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);
} }
} }
} }
@@ -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.n_dims = ((const int32_t *) op_params)[1];
rctx.mode = ((const int32_t *) op_params)[2]; rctx.mode = ((const int32_t *) op_params)[2];
rctx.n_ctx_orig = ((const int32_t *) op_params)[4]; 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_base, (int32_t *) op_params + 5, sizeof(float));
memcpy(&rctx.freq_scale, (int32_t *) op_params + 6, sizeof(float)); memcpy(&rctx.freq_scale, (int32_t *) op_params + 6, sizeof(float));
ggml/src/ggml-opencl/ggml-opencl.cpp
+8 -3
View File
@@ -7434,9 +7434,6 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
case GGML_OP_DIAG_MASK_INF: case GGML_OP_DIAG_MASK_INF:
return op->ne[3] == 1; return op->ne[3] == 1;
case GGML_OP_ROPE: { 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 int mode = ((const int32_t *) op->op_params)[2];
const bool is_mrope = mode & GGML_ROPE_TYPE_MROPE; const bool is_mrope = mode & GGML_ROPE_TYPE_MROPE;
const bool is_vision = mode == GGML_ROPE_TYPE_VISION; const bool is_vision = mode == GGML_ROPE_TYPE_VISION;
@@ -23910,6 +23907,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 n_dims = ((int *) dst->op_params)[1];
const int mode = ((int *) dst->op_params)[2]; const int mode = ((int *) dst->op_params)[2];
const int n_ctx_orig = ((int32_t *) dst->op_params)[4]; 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_base;
float freq_scale; float freq_scale;
@@ -23938,6 +23936,7 @@ static void ggml_cl_rope(ggml_backend_t backend, const ggml_tensor * src0, const
if (is_vision) { if (is_vision) {
GGML_ASSERT(n_dims == ne00/2); 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; cl_kernel kernel;
@@ -24029,6 +24028,12 @@ static void ggml_cl_rope(ggml_backend_t backend, const ggml_tensor * src0, const
if (is_mrope && !is_vision) { if (is_mrope && !is_vision) {
CL_CHECK(clSetKernelArg(kernel, 34, sizeof(int), &is_imrope)); 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 global_work_size[] = {(size_t)ne01*nth, (size_t)ne02, (size_t)ne03};
size_t local_work_size[] = {(size_t)nth, 1, 1}; size_t local_work_size[] = {(size_t)nth, 1, 1};
ggml/src/ggml-opencl/kernels/rope.cl
+52 -40
View File
@@ -75,7 +75,8 @@ kernel void kernel_rope_norm_f32(
float ext_factor, float ext_factor,
float attn_factor, float attn_factor,
float beta_fast, float beta_fast,
float beta_slow float beta_slow,
int n_offs
) { ) {
src0 = (global void*)((global char*)src0 + offset0); src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1); src1 = (global int*)((global char*)src1 + offset1);
@@ -94,14 +95,15 @@ kernel void kernel_rope_norm_f32(
float inv_ndims = -1.f/n_dims; float inv_ndims = -1.f/n_dims;
for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) { for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) {
if (i0 < n_dims) { if (i0 >= n_offs && i0 < n_offs + n_dims) {
int ic = i0/2; 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; 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 * 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); global float * dst_data = (global float *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
@@ -154,7 +156,8 @@ kernel void kernel_rope_norm_f16(
float ext_factor, float ext_factor,
float attn_factor, float attn_factor,
float beta_fast, float beta_fast,
float beta_slow float beta_slow,
int n_offs
) { ) {
src0 = (global void*)((global char*)src0 + offset0); src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1); src1 = (global int*)((global char*)src1 + offset1);
@@ -173,14 +176,15 @@ kernel void kernel_rope_norm_f16(
float inv_ndims = -1.f/n_dims; float inv_ndims = -1.f/n_dims;
for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) { for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) {
if (i0 < n_dims) { if (i0 >= n_offs && i0 < n_offs + n_dims) {
int ic = i0/2; 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; 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 * 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); global half * dst_data = (global half *)((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
@@ -233,7 +237,8 @@ kernel void kernel_rope_neox_f32(
float ext_factor, float ext_factor,
float attn_factor, float attn_factor,
float beta_fast, float beta_fast,
float beta_slow float beta_slow,
int n_offs
) { ) {
src0 = (global void*)((global char*)src0 + offset0); src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1); src1 = (global int*)((global char*)src1 + offset1);
@@ -252,17 +257,18 @@ kernel void kernel_rope_neox_f32(
float inv_ndims = -1.f/n_dims; float inv_ndims = -1.f/n_dims;
for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) { for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) {
if (i0 < n_dims) { if (i0 >= n_offs && i0 < n_offs + n_dims) {
int ic = i0/2; 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; 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 * 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 + ic*nb0); 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 x0 = src[0];
const float x1 = src[n_dims/2]; const float x1 = src[n_dims/2];
@@ -312,7 +318,8 @@ kernel void kernel_rope_neox_f16(
float ext_factor, float ext_factor,
float attn_factor, float attn_factor,
float beta_fast, float beta_fast,
float beta_slow float beta_slow,
int n_offs
) { ) {
src0 = (global void*)((global char*)src0 + offset0); src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1); src1 = (global int*)((global char*)src1 + offset1);
@@ -331,17 +338,18 @@ kernel void kernel_rope_neox_f16(
float inv_ndims = -1.f/n_dims; float inv_ndims = -1.f/n_dims;
for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) { for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) {
if (i0 < n_dims) { if (i0 >= n_offs && i0 < n_offs + n_dims) {
int ic = i0/2; 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; 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 * 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 + ic*nb0); 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 x0 = src[0];
const float x1 = src[n_dims/2]; const float x1 = src[n_dims/2];
@@ -393,7 +401,8 @@ kernel void kernel_rope_multi_f32(
float beta_fast, float beta_fast,
float beta_slow, float beta_slow,
int4 sections, int4 sections,
int is_imrope int is_imrope,
int n_offs
) { ) {
src0 = (global void*)((global char*)src0 + offset0); src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1); src1 = (global int*)((global char*)src1 + offset1);
@@ -414,10 +423,11 @@ kernel void kernel_rope_multi_f32(
float inv_ndims = -1.f/n_dims; float inv_ndims = -1.f/n_dims;
for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) { for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) {
if (i0 < n_dims) { if (i0 >= n_offs && i0 < n_offs + n_dims) {
int ic = i0/2; 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; float theta_base = 0.0f;
if (is_imrope) { if (is_imrope) {
@@ -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; 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 * 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 + ic*nb0); 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 x0 = src[0];
const float x1 = src[n_dims/2]; const float x1 = src[n_dims/2];
@@ -504,7 +514,8 @@ kernel void kernel_rope_multi_f16(
float beta_fast, float beta_fast,
float beta_slow, float beta_slow,
int4 sections, int4 sections,
int is_imrope int is_imrope,
int n_offs
) { ) {
src0 = (global void*)((global char*)src0 + offset0); src0 = (global void*)((global char*)src0 + offset0);
src1 = (global int*)((global char*)src1 + offset1); src1 = (global int*)((global char*)src1 + offset1);
@@ -525,10 +536,11 @@ kernel void kernel_rope_multi_f16(
float inv_ndims = -1.f/n_dims; float inv_ndims = -1.f/n_dims;
for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) { for (int i0 = 2*get_local_id(0); i0 < ne0; i0 += 2*get_local_size(0)) {
if (i0 < n_dims) { if (i0 >= n_offs && i0 < n_offs + n_dims) {
int ic = i0/2; 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; float theta_base = 0.0f;
if (is_imrope) { if (is_imrope) {
@@ -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; 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 * 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 + ic*nb0); 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 x0 = src[0];
const float x1 = src[n_dims/2]; const float x1 = src[n_dims/2];
ggml/src/ggml-sycl/ggml-sycl.cpp
-2
View File
@@ -6242,8 +6242,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
} }
case GGML_OP_ROPE: case GGML_OP_ROPE:
case GGML_OP_ROPE_BACK: 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:
case GGML_OP_IM2COL_3D: case GGML_OP_IM2COL_3D:
case GGML_OP_UPSCALE: case GGML_OP_UPSCALE:
ggml/src/ggml-sycl/rope.cpp
+58 -48
View File
@@ -41,7 +41,7 @@ template <bool forward, bool has_ff, typename T, typename D>
static void rope_norm(const T *x, D *dst, const int ne00, const int ne01, static void rope_norm(const T *x, D *dst, const int ne00, const int ne01,
const int ne02, const int s01, const int s02, const int ne02, const int s01, const int s02,
const int s03, const int s1, const int s2, const int s3, const int s03, const int s1, const int s2, const int s3,
const int n_dims, const int32_t *pos, const int n_dims, const int n_offs, const int32_t *pos,
const float freq_scale, const float ext_factor, const float freq_scale, const float ext_factor,
const float attn_factor, const rope_corr_dims corr_dims, const float attn_factor, const rope_corr_dims corr_dims,
const float theta_scale, const float *freq_factors, const float theta_scale, const float *freq_factors,
@@ -78,19 +78,21 @@ static void rope_norm(const T *x, D *dst, const int ne00, const int ne01,
ggml_sycl_memcpy_1<4>(dst + idst, &v); ggml_sycl_memcpy_1<4>(dst + idst, &v);
} }
}; };
if (i0 >= n_dims) { if (i0 < n_offs || i0 >= n_offs + n_dims) {
store_coaelsced(x[ix + 0], x[ix + 1]); store_coaelsced(x[ix + 0], x[ix + 1]);
return; return;
} }
const float theta_base = pos[i2] * dpct::pow(theta_scale, i0 / 2.0f); const int iw = i0 - n_offs; // relative idx
const float freq_factor = has_ff ? freq_factors[i0 / 2] : 1.0f; const float theta_base = pos[i2] * dpct::pow(theta_scale, iw / 2.0f);
const float freq_factor = has_ff ? freq_factors[iw / 2] : 1.0f;
float cos_theta; float cos_theta;
float sin_theta; float sin_theta;
rope_yarn<forward>(theta_base / freq_factor, freq_scale, corr_dims, i0, rope_yarn<forward>(theta_base / freq_factor, freq_scale, corr_dims, iw,
ext_factor, attn_factor, cos_theta, sin_theta); ext_factor, attn_factor, cos_theta, sin_theta);
const float x0 = x[ix + 0]; const float x0 = x[ix + 0];
@@ -104,7 +106,7 @@ template <bool forward, bool has_ff, typename T, typename D>
static void rope_neox(const T *x, D *dst, const int ne00, const int ne01, static void rope_neox(const T *x, D *dst, const int ne00, const int ne01,
const int ne02, const int s01, const int s02, const int ne02, const int s01, const int s02,
const int s03, const int s1, const int s2, const int s3, const int s03, const int s1, const int s2, const int s3,
const int n_dims, const int32_t *pos, const int n_dims, const int n_offs, const int32_t *pos,
const float freq_scale, const float ext_factor, const float freq_scale, const float ext_factor,
const float attn_factor, const rope_corr_dims corr_dims, const float attn_factor, const rope_corr_dims corr_dims,
const float theta_scale, const float *freq_factors, const float theta_scale, const float *freq_factors,
@@ -132,35 +134,38 @@ static void rope_neox(const T *x, D *dst, const int ne00, const int ne01,
idst += row_indices[i2] * set_rows_stride; idst += row_indices[i2] * set_rows_stride;
} }
if (i0 >= n_dims) { if (i0 < n_offs || i0 >= n_offs + n_dims) {
dst[idst + i0 / 2 + 0] = ggml_sycl_cast<D>(x[ix + i0 / 2 + 0]); 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]); dst[idst + i0 / 2 + 1] = ggml_sycl_cast<D>(x[ix + i0 / 2 + 1]);
return; return;
} }
const float theta_base = pos[i2] * dpct::pow(theta_scale, i0 / 2.0f); const int iw = i0 - n_offs; // relative idx
const float freq_factor = has_ff ? freq_factors[i0 / 2] : 1.0f; const float theta_base = pos[i2] * dpct::pow(theta_scale, iw / 2.0f);
const float freq_factor = has_ff ? freq_factors[iw / 2] : 1.0f;
float cos_theta; float cos_theta;
float sin_theta; float sin_theta;
rope_yarn<forward>(theta_base / freq_factor, freq_scale, corr_dims, i0, rope_yarn<forward>(theta_base / freq_factor, freq_scale, corr_dims, iw,
ext_factor, attn_factor, cos_theta, sin_theta); ext_factor, attn_factor, cos_theta, sin_theta);
const float x0 = x[ix + 0]; // idst/ix point at channel i0/2; the first channel of the rotated pair is n_offs + iw/2 = i0/2 + n_offs/2
const float x1 = x[ix + n_dims / 2]; const float x0 = x[ix + n_offs / 2 + 0];
const float x1 = x[ix + n_offs / 2 + n_dims / 2];
dst[idst + 0] = ggml_sycl_cast<D>(x0 * cos_theta - x1 * sin_theta); dst[idst + n_offs / 2 + 0] = ggml_sycl_cast<D>(x0 * cos_theta - x1 * sin_theta);
dst[idst + n_dims / 2] = ggml_sycl_cast<D>(x0 * sin_theta + x1 * cos_theta); dst[idst + n_offs / 2 + n_dims / 2] = ggml_sycl_cast<D>(x0 * sin_theta + x1 * cos_theta);
} }
template <bool forward, bool has_ff, typename T> template <bool forward, bool has_ff, typename T>
static void rope_multi(const T *x, T *dst, const int ne00, const int ne01, static void rope_multi(const T *x, T *dst, const int ne00, const int ne01,
const int ne02, const int s01, const int s02, const int ne02, const int s01, const int s02,
const int s03, const int s1, const int s2, const int s3, const int s03, const int s1, const int s2, const int s3,
const int n_dims, const int32_t *pos, const int n_dims, const int n_offs, const int32_t *pos,
const float freq_scale, const float ext_factor, const float freq_scale, const float ext_factor,
const float attn_factor, const rope_corr_dims corr_dims, const float attn_factor, const rope_corr_dims corr_dims,
const float theta_scale, const float *freq_factors, const float theta_scale, const float *freq_factors,
@@ -183,54 +188,57 @@ static void rope_multi(const T *x, T *dst, const int ne00, const int ne01,
int idst = i0 / 2 + i1 * s1 + i2 * s2 + i3 * s3; int idst = i0 / 2 + i1 * s1 + i2 * s2 + i3 * s3;
const int ix = i0 / 2 + i1 * s01 + i2 * s02 + i3 * s03; const int ix = i0 / 2 + i1 * s01 + i2 * s02 + i3 * s03;
if (i0 >= n_dims) { if (i0 < n_offs || i0 >= n_offs + n_dims) {
dst[idst + i0 / 2 + 0] = x[ix + i0 / 2 + 0]; dst[idst + i0 / 2 + 0] = x[ix + i0 / 2 + 0];
dst[idst + i0 / 2 + 1] = x[ix + i0 / 2 + 1]; dst[idst + i0 / 2 + 1] = x[ix + i0 / 2 + 1];
return; return;
} }
const int iw = i0 - n_offs; // relative idx
const int sect_dims = const int sect_dims =
sections.v[0] + sections.v[1] + sections.v[2] + sections.v[3]; sections.v[0] + sections.v[1] + sections.v[2] + sections.v[3];
const int sec_w = sections.v[1] + sections.v[0]; const int sec_w = sections.v[1] + sections.v[0];
const int sector = (i0 / 2) % sect_dims; const int sector = (iw / 2) % sect_dims;
float theta_base = 0.0; float theta_base = 0.0;
if (is_imrope) { if (is_imrope) {
if (sector % 3 == 1 && sector < 3 * sections.v[1]) { // h if (sector % 3 == 1 && sector < 3 * sections.v[1]) { // h
theta_base = pos[i2 + ne02 * 1] * dpct::pow(theta_scale, i0 / 2.0f); theta_base = pos[i2 + ne02 * 1] * dpct::pow(theta_scale, iw / 2.0f);
} else if (sector % 3 == 2 && sector < 3 * sections.v[2]) { // w } else if (sector % 3 == 2 && sector < 3 * sections.v[2]) { // w
theta_base = pos[i2 + ne02 * 2] * dpct::pow(theta_scale, i0 / 2.0f); theta_base = pos[i2 + ne02 * 2] * dpct::pow(theta_scale, iw / 2.0f);
} else if (sector % 3 == 0 && sector < 3 * sections.v[0]) { // t } else if (sector % 3 == 0 && sector < 3 * sections.v[0]) { // t
theta_base = pos[i2] * dpct::pow(theta_scale, i0 / 2.0f); theta_base = pos[i2] * dpct::pow(theta_scale, iw / 2.0f);
} else { } else {
theta_base = pos[i2 + ne02 * 3] * dpct::pow(theta_scale, i0 / 2.0f); theta_base = pos[i2 + ne02 * 3] * dpct::pow(theta_scale, iw / 2.0f);
} }
} else { } else {
if (sector < sections.v[0]) { if (sector < sections.v[0]) {
theta_base = pos[i2] * dpct::pow(theta_scale, i0 / 2.0f); theta_base = pos[i2] * dpct::pow(theta_scale, iw / 2.0f);
} else if (sector >= sections.v[0] && sector < sec_w) { } else if (sector >= sections.v[0] && sector < sec_w) {
theta_base = pos[i2 + ne02 * 1] * dpct::pow(theta_scale, i0 / 2.0f); theta_base = pos[i2 + ne02 * 1] * dpct::pow(theta_scale, iw / 2.0f);
} else if (sector >= sec_w && sector < sec_w + sections.v[2]) { } else if (sector >= sec_w && sector < sec_w + sections.v[2]) {
theta_base = pos[i2 + ne02 * 2] * dpct::pow(theta_scale, i0 / 2.0f); theta_base = pos[i2 + ne02 * 2] * dpct::pow(theta_scale, iw / 2.0f);
} else if (sector >= sec_w + sections.v[2]) { } else if (sector >= sec_w + sections.v[2]) {
theta_base = pos[i2 + ne02 * 3] * dpct::pow(theta_scale, i0 / 2.0f); theta_base = pos[i2 + ne02 * 3] * dpct::pow(theta_scale, iw / 2.0f);
} }
} }
const float freq_factor = has_ff ? freq_factors[i0 / 2] : 1.0f; const float freq_factor = has_ff ? freq_factors[iw / 2] : 1.0f;
float cos_theta; float cos_theta;
float sin_theta; float sin_theta;
rope_yarn<forward>(theta_base / freq_factor, freq_scale, corr_dims, i0, rope_yarn<forward>(theta_base / freq_factor, freq_scale, corr_dims, iw,
ext_factor, attn_factor, cos_theta, sin_theta); ext_factor, attn_factor, cos_theta, sin_theta);
const float x0 = x[ix + 0]; // idst/ix point at channel i0/2; the first channel of the rotated pair is n_offs + iw/2 = i0/2 + n_offs/2
const float x1 = x[ix + n_dims / 2]; const float x0 = x[ix + n_offs / 2 + 0];
const float x1 = x[ix + n_offs / 2 + n_dims / 2];
dst[idst + 0] = x0 * cos_theta - x1 * sin_theta; dst[idst + n_offs / 2 + 0] = x0 * cos_theta - x1 * sin_theta;
dst[idst + n_dims / 2] = x0 * sin_theta + x1 * cos_theta; dst[idst + n_offs / 2 + n_dims / 2] = x0 * sin_theta + x1 * cos_theta;
} }
template <bool forward, bool has_ff, typename T> template <bool forward, bool has_ff, typename T>
@@ -293,7 +301,7 @@ static void
rope_norm_sycl(const T *x, D *dst, const int ne00, const int ne01, rope_norm_sycl(const T *x, D *dst, const int ne00, const int ne01,
const int ne02, const int s01, const int s02, const int s03, const int ne02, const int s01, const int s02, const int s03,
const int s1, const int s2, const int s3, const int n_dims, const int s1, const int s2, const int s3, const int n_dims,
const int nr, const int32_t *pos, const float freq_scale, const int n_offs, const int nr, const int32_t *pos, const float freq_scale,
const float freq_base, const float ext_factor, const float freq_base, const float ext_factor,
const float attn_factor, const rope_corr_dims corr_dims, const float attn_factor, const rope_corr_dims corr_dims,
const float *freq_factors, const int64_t *row_indices, const float *freq_factors, const int64_t *row_indices,
@@ -313,7 +321,7 @@ rope_norm_sycl(const T *x, D *dst, const int ne00, const int ne01,
GGML_UNUSED(item_ct1); GGML_UNUSED(item_ct1);
rope_norm<forward, false>( rope_norm<forward, false>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
pos, freq_scale, ext_factor, attn_factor, corr_dims, n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
theta_scale, freq_factors, row_indices, set_rows_stride); theta_scale, freq_factors, row_indices, set_rows_stride);
}); });
} else { } else {
@@ -323,7 +331,7 @@ rope_norm_sycl(const T *x, D *dst, const int ne00, const int ne01,
GGML_UNUSED(item_ct1); GGML_UNUSED(item_ct1);
rope_norm<forward, true>( rope_norm<forward, true>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
pos, freq_scale, ext_factor, attn_factor, corr_dims, n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
theta_scale, freq_factors, row_indices, set_rows_stride); theta_scale, freq_factors, row_indices, set_rows_stride);
}); });
} }
@@ -334,7 +342,7 @@ static void
rope_neox_sycl(const T *x, D *dst, const int ne00, const int ne01, rope_neox_sycl(const T *x, D *dst, const int ne00, const int ne01,
const int ne02, const int s01, const int s02, const int s03, const int ne02, const int s01, const int s02, const int s03,
const int s1, const int s2, const int s3, const int n_dims, const int s1, const int s2, const int s3, const int n_dims,
const int nr, const int32_t *pos, const float freq_scale, const int n_offs, const int nr, const int32_t *pos, const float freq_scale,
const float freq_base, const float ext_factor, const float freq_base, const float ext_factor,
const float attn_factor, const rope_corr_dims corr_dims, const float attn_factor, const rope_corr_dims corr_dims,
const float *freq_factors, const int64_t *row_indices, const float *freq_factors, const int64_t *row_indices,
@@ -354,7 +362,7 @@ rope_neox_sycl(const T *x, D *dst, const int ne00, const int ne01,
GGML_UNUSED(item_ct1); GGML_UNUSED(item_ct1);
rope_neox<forward, false>( rope_neox<forward, false>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
pos, freq_scale, ext_factor, attn_factor, corr_dims, n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
theta_scale, freq_factors, row_indices, set_rows_stride); theta_scale, freq_factors, row_indices, set_rows_stride);
}); });
} else { } else {
@@ -364,7 +372,7 @@ rope_neox_sycl(const T *x, D *dst, const int ne00, const int ne01,
GGML_UNUSED(item_ct1); GGML_UNUSED(item_ct1);
rope_neox<forward, true>( rope_neox<forward, true>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
pos, freq_scale, ext_factor, attn_factor, corr_dims, n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
theta_scale, freq_factors, row_indices, set_rows_stride); theta_scale, freq_factors, row_indices, set_rows_stride);
}); });
} }
@@ -375,7 +383,7 @@ static void
rope_multi_sycl(const T *x, T *dst, const int ne00, const int ne01, rope_multi_sycl(const T *x, T *dst, const int ne00, const int ne01,
const int ne02, const int s01, const int s02, const int s03, const int ne02, const int s01, const int s02, const int s03,
const int s1, const int s2, const int s3, const int n_dims, const int s1, const int s2, const int s3, const int n_dims,
const int nr, const int32_t *pos, const float freq_scale, const int n_offs, const int nr, const int32_t *pos, const float freq_scale,
const float freq_base, const float ext_factor, const float freq_base, const float ext_factor,
const float attn_factor, const rope_corr_dims corr_dims, const float attn_factor, const rope_corr_dims corr_dims,
const float *freq_factors, const mrope_sections sections, const float *freq_factors, const mrope_sections sections,
@@ -395,7 +403,7 @@ rope_multi_sycl(const T *x, T *dst, const int ne00, const int ne01,
GGML_UNUSED(item_ct1); GGML_UNUSED(item_ct1);
rope_multi<forward, false, T>( rope_multi<forward, false, T>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
pos, freq_scale, ext_factor, attn_factor, corr_dims, n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
theta_scale, freq_factors, sections, is_imrope); theta_scale, freq_factors, sections, is_imrope);
}); });
} else { } else {
@@ -405,7 +413,7 @@ rope_multi_sycl(const T *x, T *dst, const int ne00, const int ne01,
GGML_UNUSED(item_ct1); GGML_UNUSED(item_ct1);
rope_multi<forward, true, T>( rope_multi<forward, true, T>(
x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims, x, dst, ne00, ne01, ne02, s01, s02, s03, s1, s2, s3, n_dims,
pos, freq_scale, ext_factor, attn_factor, corr_dims, n_offs, pos, freq_scale, ext_factor, attn_factor, corr_dims,
theta_scale, freq_factors, sections, is_imrope); theta_scale, freq_factors, sections, is_imrope);
}); });
} }
@@ -497,6 +505,7 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst,
const int n_dims = ((int32_t *)dst->op_params)[1]; const int n_dims = ((int32_t *)dst->op_params)[1];
const int mode = ((int32_t *)dst->op_params)[2]; const int mode = ((int32_t *)dst->op_params)[2];
const int n_ctx_orig = ((int32_t *)dst->op_params)[4]; const int n_ctx_orig = ((int32_t *)dst->op_params)[4];
const int n_offs = ((int32_t *)dst->op_params)[15];
mrope_sections sections; mrope_sections sections;
float freq_base; float freq_base;
@@ -526,6 +535,7 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst,
if (is_vision) { if (is_vision) {
GGML_ASSERT(n_dims == ne00 / 2); GGML_ASSERT(n_dims == ne00 / 2);
GGML_ASSERT(n_offs == 0); // offset not supported for vision, as the rotated pairs span the whole row
} }
const int32_t *pos = (const int32_t *)src1_d; const int32_t *pos = (const int32_t *)src1_d;
@@ -545,19 +555,19 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst,
if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F32) { if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F32) {
rope_neox_sycl<forward, float, float>( rope_neox_sycl<forward, float, float>(
(const float *)src0_d, (float *)dst_d, ne00, ne01, ne02, s01, (const float *)src0_d, (float *)dst_d, ne00, ne01, ne02, s01,
s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base, 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, ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
set_rows_stride, stream); set_rows_stride, stream);
} else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) { } else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) {
rope_neox_sycl<forward, float, sycl::half>( rope_neox_sycl<forward, float, sycl::half>(
(const float *)src0_d, (sycl::half *)dst_d, ne00, ne01, ne02, (const float *)src0_d, (sycl::half *)dst_d, ne00, ne01, ne02,
s01, s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale,
freq_base, ext_factor, attn_factor, corr_dims, freq_factors, freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
row_indices, set_rows_stride, stream); row_indices, set_rows_stride, stream);
} else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) { } else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) {
rope_neox_sycl<forward, sycl::half, sycl::half>( rope_neox_sycl<forward, sycl::half, sycl::half>(
(const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01, (const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01,
ne02, s01, s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, 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, freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
row_indices, set_rows_stride, stream); row_indices, set_rows_stride, stream);
} else { } else {
@@ -568,13 +578,13 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst,
if (src0->type == GGML_TYPE_F32) { if (src0->type == GGML_TYPE_F32) {
rope_multi_sycl<forward>((const float *)src0_d, (float *)dst_d, rope_multi_sycl<forward>((const float *)src0_d, (float *)dst_d,
ne00, ne01, ne02, s01, s02, s03, s1, s2, ne00, ne01, ne02, s01, s02, s03, s1, s2,
s3, n_dims, nr, pos, freq_scale, freq_base, s3, n_dims, n_offs, nr, pos, freq_scale, freq_base,
ext_factor, attn_factor, corr_dims, ext_factor, attn_factor, corr_dims,
freq_factors, sections, is_imrope, stream); freq_factors, sections, is_imrope, stream);
} else if (src0->type == GGML_TYPE_F16) { } else if (src0->type == GGML_TYPE_F16) {
rope_multi_sycl<forward>( rope_multi_sycl<forward>(
(const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01, (const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01,
ne02, s01, s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, 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, freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
sections, is_imrope, stream); sections, is_imrope, stream);
} else { } else {
@@ -602,19 +612,19 @@ void ggml_sycl_op_rope_impl(ggml_backend_sycl_context &ctx, ggml_tensor *dst,
if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F32) { if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F32) {
rope_norm_sycl<forward, float, float>( rope_norm_sycl<forward, float, float>(
(const float *)src0_d, (float *)dst_d, ne00, ne01, ne02, s01, (const float *)src0_d, (float *)dst_d, ne00, ne01, ne02, s01,
s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base, 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, ext_factor, attn_factor, corr_dims, freq_factors, row_indices,
set_rows_stride, stream); set_rows_stride, stream);
} else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) { } else if (src0->type == GGML_TYPE_F32 && dst_type == GGML_TYPE_F16) {
rope_norm_sycl<forward, float, sycl::half>( rope_norm_sycl<forward, float, sycl::half>(
(const float *)src0_d, (sycl::half *)dst_d, ne00, ne01, ne02, (const float *)src0_d, (sycl::half *)dst_d, ne00, ne01, ne02,
s01, s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, s01, s02, s03, s1, s2, s3, n_dims, n_offs, nr, pos, freq_scale,
freq_base, ext_factor, attn_factor, corr_dims, freq_factors, freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
row_indices, set_rows_stride, stream); row_indices, set_rows_stride, stream);
} else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) { } else if (src0->type == GGML_TYPE_F16 && dst_type == GGML_TYPE_F16) {
rope_norm_sycl<forward, sycl::half, sycl::half>( rope_norm_sycl<forward, sycl::half, sycl::half>(
(const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01, (const sycl::half *)src0_d, (sycl::half *)dst_d, ne00, ne01,
ne02, s01, s02, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, 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, freq_base, ext_factor, attn_factor, corr_dims, freq_factors,
row_indices, set_rows_stride, stream); row_indices, set_rows_stride, stream);
} else { } else {
ggml/src/ggml-webgpu/ggml-webgpu.cpp
+4 -4
View File
@@ -2714,6 +2714,7 @@ static webgpu_encoded_op ggml_webgpu_rope(webgpu_context & ctx,
const int n_dims = ((int32_t *) dst->op_params)[1]; const int n_dims = ((int32_t *) dst->op_params)[1];
const int mode = ((int32_t *) dst->op_params)[2]; const int mode = ((int32_t *) dst->op_params)[2];
const int n_ctx_orig = ((int32_t *) dst->op_params)[4]; 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_base;
float freq_scale; float freq_scale;
@@ -2762,7 +2763,8 @@ static webgpu_encoded_op ggml_webgpu_rope(webgpu_context & ctx,
(uint32_t) sections[0], (uint32_t) sections[0],
(uint32_t) sections[1], (uint32_t) sections[1],
(uint32_t) sections[2], (uint32_t) sections[2],
(uint32_t) sections[3] (uint32_t) sections[3],
(uint32_t) n_offs
}; };
std::vector<wgpu::BindGroupEntry> entries = { ggml_webgpu_make_tensor_bind_group_entry(ctx, 0, src0), std::vector<wgpu::BindGroupEntry> entries = { ggml_webgpu_make_tensor_bind_group_entry(ctx, 0, src0),
@@ -4472,9 +4474,7 @@ static bool ggml_backend_webgpu_device_supports_op(ggml_backend_dev_t dev, const
supports_op = (op->type == GGML_TYPE_F32 && src0->type == GGML_TYPE_F32) && ggml_is_contiguous_rows(src0); supports_op = (op->type == GGML_TYPE_F32 && src0->type == GGML_TYPE_F32) && ggml_is_contiguous_rows(src0);
break; break;
case GGML_OP_ROPE: case GGML_OP_ROPE:
// FIXME: support ggml_rope_set_offset supports_op = op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16;
supports_op =
(op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16) && ((const int32_t *) op->op_params)[15] == 0;
break; break;
case GGML_OP_GLU: case GGML_OP_GLU:
switch (ggml_get_glu_op(op)) { switch (ggml_get_glu_op(op)) {
ggml/src/ggml-webgpu/wgsl-shaders/rope.wgsl
+11 -7
View File
@@ -38,7 +38,8 @@ struct Params {
sections0: u32, sections0: u32,
sections1: u32, sections1: u32,
sections2: u32, sections2: u32,
sections3: u32 sections3: u32,
n_offs: u32
}; };
@group(0) @binding(0) @group(0) @binding(0)
@@ -126,7 +127,8 @@ fn rope_yarn(theta_extrap: f32, i: u32) -> vec2<f32> {
fn pair_base(i0: u32, div_2: bool) -> u32 { fn pair_base(i0: u32, div_2: bool) -> u32 {
if (div_2) { if (div_2) {
return i0 / 2; // first channel of the rotated pair: n_offs + (i0 - n_offs)/2
return i0 / 2 + params.n_offs / 2;
} else { } else {
return i0; return i0;
} }
@@ -165,20 +167,22 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
let i_src_row = params.offset_src0 + i3 * params.stride_src03 + i2 * params.stride_src02 + i1 * params.stride_src01; let i_src_row = params.offset_src0 + i3 * params.stride_src03 + i2 * params.stride_src02 + i1 * params.stride_src01;
let i_dst_row = params.offset_dst + i3 * params.stride_dst3 + i2 * params.stride_dst2 + i1 * params.stride_dst1; let i_dst_row = params.offset_dst + i3 * params.stride_dst3 + i2 * params.stride_dst2 + i1 * params.stride_dst1;
if (i0 >= params.n_dims && !is_vision) { if ((i0 < params.n_offs || i0 >= params.n_offs + params.n_dims) && !is_vision) {
let i_src = i_src_row + i0; let i_src = i_src_row + i0;
let i_dst = i_dst_row + i0; let i_dst = i_dst_row + i0;
rotate(i_dst, i_dst + 1, f32(src0[i_src]), f32(src0[i_src + 1])); rotate(i_dst, i_dst + 1, f32(src0[i_src]), f32(src0[i_src + 1]));
return; return;
} }
let iw = i0 - params.n_offs; // relative idx
var theta_base_mult: u32 = 0; var theta_base_mult: u32 = 0;
var theta_scale_pwr: u32 = i0 / 2; var theta_scale_pwr: u32 = iw / 2;
if (is_mrope) { if (is_mrope) {
let sect_dims = params.sections0 + params.sections1 + params.sections2 + params.sections3; let sect_dims = params.sections0 + params.sections1 + params.sections2 + params.sections3;
let sec_w = params.sections1 + params.sections0; let sec_w = params.sections1 + params.sections0;
let sec_e = params.sections2 + sec_w; let sec_e = params.sections2 + sec_w;
let sector = (i0 / 2) % sect_dims; let sector = (iw / 2) % sect_dims;
if (is_imrope) { if (is_imrope) {
if (sector % 3 == 1 && sector < 3 * params.sections1) { if (sector % 3 == 1 && sector < 3 * params.sections1) {
theta_base_mult = 1; theta_base_mult = 1;
@@ -203,7 +207,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
} else if (sector >= sec_e) { } else if (sector >= sec_e) {
if (is_vision) { if (is_vision) {
theta_scale_pwr = sector - sec_e; theta_scale_pwr = sector - sec_e;
theta_scale_pwr = (i0 / 2) % sec_e; theta_scale_pwr = (iw / 2) % sec_e;
} }
theta_base_mult = 3; theta_base_mult = 3;
} else if (is_vision) { } else if (is_vision) {
@@ -212,7 +216,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
} }
} }
let theta_base = f32(src1[params.offset_src1 + i2 + params.ne2 * theta_base_mult]) * pow(params.theta_scale, f32(theta_scale_pwr)); let theta_base = f32(src1[params.offset_src1 + i2 + params.ne2 * theta_base_mult]) * pow(params.theta_scale, f32(theta_scale_pwr));
let thetas = rope_yarn(theta_base/freq_factor(i0), i0); let thetas = rope_yarn(theta_base/freq_factor(iw), iw);
let i_src = i_src_row + pair_base(i0, is_neox || is_mrope || is_vision); let i_src = i_src_row + pair_base(i0, is_neox || is_mrope || is_vision);
let i_dst = i_dst_row + pair_base(i0, is_neox || is_mrope || is_vision); let i_dst = i_dst_row + pair_base(i0, is_neox || is_mrope || is_vision);