diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index 3e48a9244..b44e860f7 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -4009,10 +4009,7 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_OP_ROPE:
{
const int mode = ((const int32_t *) op->op_params)[2];
- if (mode & GGML_ROPE_TYPE_MROPE) {
- return false;
- }
- if (mode & GGML_ROPE_TYPE_VISION) {
+ if (mode == GGML_ROPE_TYPE_MROPE) {
return false;
}
return ggml_is_contiguous(op->src[0]);
diff --git a/ggml/src/ggml-sycl/rope.cpp b/ggml/src/ggml-sycl/rope.cpp
index bbcb356e9..80e050f24 100644
--- a/ggml/src/ggml-sycl/rope.cpp
+++ b/ggml/src/ggml-sycl/rope.cpp
@@ -1,9 +1,15 @@
#include "rope.hpp"
+#include "ggml-sycl/common.hpp"
+#include "ggml.h"
struct rope_corr_dims {
float v[2];
};
+struct mrope_sections {
+ int v[4];
+};
+
static float rope_yarn_ramp(const float low, const float high, const int i0) {
const float y = (i0 / 2 - low) / sycl::max(0.001f, high - low);
return 1.0f - sycl::min(1.0f, sycl::max(0.0f, y));
@@ -114,6 +120,48 @@ static void rope_neox(
dst[i + n_dims/2] = x0*sin_theta + x1*cos_theta;
}
+template <typename T, bool has_ff>
+static void rope_vision(const T * x, T * dst, const int ne0, const int ne1, const int ne2, const size_t s1,
+ const size_t s2, const int n_dims, const int32_t * pos, 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 sycl::nd_item<3> & item_ct1) {
+ // get index pos
+ const int i0 = 2 * (item_ct1.get_group(1) * item_ct1.get_local_range(1) + item_ct1.get_local_id(1));
+ if (i0 >= ne0) {
+ return;
+ }
+ const int row_dst = (item_ct1.get_group(2) * item_ct1.get_local_range(2)) + item_ct1.get_local_id(2);
+ const int row_x = row_dst % ne1;
+ const int channel_x = row_dst / ne1;
+ const int idst = (row_dst * ne0) + (i0 / 2);
+ const size_t ix = ((size_t) channel_x * s2) + ((size_t) row_x * s1) + (i0 / 2);
+
+ const int sect_dims = sections.v[0] + sections.v[1];
+ const int sector = (i0 / 2) % sect_dims;
+
+ float theta_base = 0.0f;
+ if (sector < sections.v[0]) {
+ const int p = sector;
+ theta_base = pos[channel_x] * sycl::pow(theta_scale, (float) p);
+ } else {
+ // Simplified from CUDA backend code: if (sector >= sections.v[0] && sector < sec_w) which is just sector >= sections.v[0]
+ const int p = sector - sections.v[0];
+ theta_base = pos[channel_x + ne2] * sycl::pow(theta_scale, (float) p);
+ }
+
+ const float freq_factor = has_ff ? freq_factors[i0 / 2] : 1.0f;
+ float cos_theta;
+ float sin_theta;
+ rope_yarn(theta_base / freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor, &cos_theta, &sin_theta);
+ const float x0 = x[ix + 0];
+ const float x1 = x[ix + n_dims];
+
+ // store results in dst
+ dst[idst + 0] = x0 * cos_theta - x1 * sin_theta;
+ dst[idst + n_dims] = x0 * sin_theta + x1 * cos_theta;
+}
+
template <typename T>
static void rope_norm_sycl(
const T *x, T *dst, int ne0, int n_dims, int nr, const int32_t *pos, float freq_scale, int p_delta_rows,
@@ -192,21 +240,58 @@ static void rope_neox_sycl(
}
}
+// rope vision
+template <typename T>
+static void rope_vision_sycl(const T * x, T * dst, const int ne0, const int ne1, const int ne2, const size_t s1,
+ const size_t s2, const int n_dims, const int nr, const int32_t * pos,
+ const float freq_scale, 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, queue_ptr stream) {
+ GGML_ASSERT(ne0 % 2 == 0);
+ const sycl::range<3> block_dims(1, SYCL_ROPE_BLOCK_SIZE, 1);
+ const int n_blocks_y = (ne0 + 2 * SYCL_ROPE_BLOCK_SIZE - 1) / (2 * SYCL_ROPE_BLOCK_SIZE);
+ const sycl::range<3> grid_dims(1, n_blocks_y, nr);
+ const sycl::nd_range<3> nd_range(grid_dims * block_dims, block_dims);
+
+ const float theta_scale = std::pow(freq_base, -2.0f / n_dims);
+ // Add FP16 capability check if T could be sycl::half
+ if constexpr (std::is_same_v<T, sycl::half>) {
+ dpct::has_capability_or_fail(stream->get_device(), { sycl::aspect::fp16 });
+ }
+ // launch kernel
+ if (freq_factors == nullptr) {
+ stream->parallel_for(nd_range, [=](sycl::nd_item<3> item_ct1) {
+ rope_vision<T, false>(x, dst, ne0, ne1, ne2, s1, s2, n_dims, pos, freq_scale, ext_factor, attn_factor,
+ corr_dims, theta_scale, freq_factors, sections, item_ct1);
+ });
+ } else {
+ stream->parallel_for(nd_range, [=](sycl::nd_item<3> item_ct1) {
+ rope_vision<T, true>(x, dst, ne0, ne1, ne2, s1, s2, n_dims, pos, freq_scale, ext_factor, attn_factor,
+ corr_dims, theta_scale, freq_factors, sections, item_ct1);
+ });
+ }
+}
+
void ggml_sycl_op_rope(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32 || dst->src[0]->type == GGML_TYPE_F16);
GGML_ASSERT( dst->type == GGML_TYPE_F32 || dst->type == GGML_TYPE_F16);
GGML_ASSERT(dst->src[0]->type == dst->type);
-
- const int64_t ne00 = dst->src[0]->ne[0];
- const int64_t ne01 = dst->src[0]->ne[1];
+ const int64_t ne00 = dst->src[0]->ne[0]; // head dims
+ const int64_t ne01 = dst->src[0]->ne[1]; // num heads
+ const int64_t ne02 = dst->src[0]->ne[2]; // num heads
const int64_t nr = ggml_nrows(dst->src[0]);
+ const size_t s01 = dst->src[0]->nb[1] / ggml_type_size(dst->src[0]->type);
+ const size_t s02 = dst->src[0]->nb[2] / ggml_type_size(dst->src[0]->type);
+
+
//const int n_past = ((int32_t *) dst->op_params)[0];
const int n_dims = ((int32_t *) dst->op_params)[1];
const int mode = ((int32_t *) dst->op_params)[2];
//const int n_ctx = ((int32_t *) dst->op_params)[3];
const int n_ctx_orig = ((int32_t *) dst->op_params)[4];
+ mrope_sections sections;
// RoPE alteration for extended context
float freq_base;
@@ -222,8 +307,10 @@ void ggml_sycl_op_rope(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
memcpy(&attn_factor, (int32_t *) dst->op_params + 8, sizeof(float));
memcpy(&beta_fast, (int32_t *) dst->op_params + 9, sizeof(float));
memcpy(&beta_slow, (int32_t *) dst->op_params + 10, sizeof(float));
+ memcpy(§ions.v, (int32_t *) dst->op_params + 11, sizeof(int)*4);
const bool is_neox = mode & GGML_ROPE_TYPE_NEOX;
+ const bool is_vision = mode == GGML_ROPE_TYPE_VISION;
const int32_t * pos = (const int32_t *) dst->src[1]->data;
@@ -240,6 +327,7 @@ void ggml_sycl_op_rope(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
// compute
if (is_neox) {
+ GGML_SYCL_DEBUG("%s: neox path\n", __func__);
if (dst->src[0]->type == GGML_TYPE_F32) {
rope_neox_sycl(
(const float *)dst->src[0]->data, (float *)dst->data, ne00, n_dims, nr, pos, freq_scale, ne01, freq_base, ext_factor,
@@ -253,7 +341,19 @@ void ggml_sycl_op_rope(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
} else {
GGML_ABORT("fatal error");
}
+ } else if (is_vision) {
+ GGML_SYCL_DEBUG("%s: vision path\n", __func__);
+ if (dst->src[0]->type == GGML_TYPE_F16) {
+ rope_vision_sycl((const sycl::half *)dst->src[0]->data, (sycl::half *)dst->data, ne00, ne01, ne02, s01, s02, n_dims, nr, pos, freq_scale,
+ freq_base, ext_factor, attn_factor, corr_dims, freq_factors, sections, main_stream);
+ } else if (dst->src[0]->type == GGML_TYPE_F32) {
+ rope_vision_sycl((const float *) dst->src[0]->data, (float *)dst->data, ne00, ne01, ne02, s01, s02, n_dims, nr, pos, freq_scale,
+ freq_base, ext_factor, attn_factor, corr_dims, freq_factors, sections, main_stream);
+ } else {
+ GGML_ABORT("Fatal error: Tensor type unsupported!");
+ }
} else {
+ GGML_SYCL_DEBUG("%s: norm path\n", __func__);
if (dst->src[0]->type == GGML_TYPE_F32) {
rope_norm_sycl(
(const float *)dst->src[0]->data, (float *)dst->data, ne00, n_dims, nr, pos, freq_scale, ne01, freq_base, ext_factor,