From 8d51e183b077f84b4c58e66c7d1884946932e7b0 Mon Sep 17 00:00:00 2001 From: shani-f Date: Thu, 23 Oct 2025 12:38:20 +0300 Subject: [PATCH 1/5] =?UTF-8?q?SYCL=20repeat=5Fback=20v1=20=E2=80=94=20add?= =?UTF-8?q?=20core=20op=20+=20switch=20case?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- ggml/src/ggml-sycl/ggml-sycl.cpp | 14 +++++++ ggml/src/ggml-sycl/repeat_back.cpp | 62 ++++++++++++++++++++++++++++++ ggml/src/ggml-sycl/repeat_back.hpp | 9 +++++ 3 files changed, 85 insertions(+) create mode 100644 ggml/src/ggml-sycl/repeat_back.cpp create mode 100644 ggml/src/ggml-sycl/repeat_back.hpp diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index f3407a813d731..b43f37478dc36 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -44,6 +44,7 @@ #include "ggml-sycl/set_rows.hpp" #include "ggml-sycl/sycl_hw.hpp" #include "ggml-sycl/getrows.hpp" +#include "ggml-sycl/repeat_back.hpp" #include "ggml-sycl/quantize.hpp" #include "ggml.h" @@ -2597,6 +2598,10 @@ catch (sycl::exception const &exc) { std::exit(1); } +static void ggml_sycl_repeat_back(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { + scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1); + ggml_sycl_op_repeat_back(ctx, dst); +} static void ggml_sycl_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2); @@ -3616,6 +3621,9 @@ static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct gg case GGML_OP_REPEAT: ggml_sycl_repeat(ctx, dst); break; + case GGML_OP_REPEAT_BACK: + ggml_sycl_repeat_back(ctx, dst); + break; case GGML_OP_GET_ROWS: ggml_sycl_get_rows(ctx, dst); break; @@ -4405,11 +4413,17 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g } return false; } + case GGML_OP_CONCAT: { ggml_type src0_type = op->src[0]->type; return src0_type != GGML_TYPE_I32 && src0_type != GGML_TYPE_I16; } + case GGML_OP_REPEAT_BACK: + { + ggml_type src0_type = op->src[0]->type; + return src0_type == GGML_TYPE_F32; + } case GGML_OP_DUP: case GGML_OP_ARGMAX: case GGML_OP_NONE: diff --git a/ggml/src/ggml-sycl/repeat_back.cpp b/ggml/src/ggml-sycl/repeat_back.cpp new file mode 100644 index 0000000000000..10fef098cdd93 --- /dev/null +++ b/ggml/src/ggml-sycl/repeat_back.cpp @@ -0,0 +1,62 @@ +//בס"ד +#include "repeat_back.hpp" + +#include "common.hpp" + +void ggml_sycl_op_repeat_back(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { + GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32); + GGML_ASSERT(dst->type == GGML_TYPE_F32); + + const float * src0_dd = (const float *) dst->src[0]->data; + float * dst_dd = (float *) dst->data; + + const int64_t ne0 = dst->ne[0]; + const int64_t ne1 = dst->ne[1]; + const int64_t ne2 = dst->ne[2]; + const int64_t ne3 = dst->ne[3]; + const int64_t ne00 = dst->src[0]->ne[0]; + const int64_t ne01 = dst->src[0]->ne[1]; + const int64_t ne02 = dst->src[0]->ne[2]; + const int64_t ne03 = dst->src[0]->ne[3]; + + const int nr0 = (int) (ne00 / ne0); + const int nr1 = (int) (ne01 / ne1); + const int nr2 = (int) (ne02 / ne2); + const int nr3 = (int) (ne03 / ne3); + + const size_t total = ne0 * ne1 * ne2 * ne3; + const int BLOCK_SIZE = 256; + const int num_blocks = (total + BLOCK_SIZE - 1) / BLOCK_SIZE; + + queue_ptr stream = ctx.stream(); + stream->memset(dst_dd, 0, total * sizeof(float)); + + stream->parallel_for( + sycl::nd_range<1>(sycl::range<1>(num_blocks * BLOCK_SIZE), sycl::range<1>(BLOCK_SIZE)), + [=](sycl::nd_item<1> item_ct1) { + const size_t i = item_ct1.get_global_linear_id(); + if (i >= total) { + return; + } + + const int i0 = i % ne0; + const int i1 = (i / ne0) % ne1; + const int i2 = (i / (ne0 * ne1)) % ne2; + const int i3 = i / (ne0 * ne1 * ne2); + + float acc = 0.0f; + + for (int j3 = 0; j3 < nr3; ++j3) { + for (int j2 = 0; j2 < nr2; ++j2) { + for (int j1 = 0; j1 < nr1; ++j1) { + for (int j0 = 0; j0 < nr0; ++j0) { + acc += src0_dd[(i0 + j0 * ne0) + (i1 + j1 * ne1) * ne00 + (i2 + j2 * ne2) * ne00 * ne01 + + (i3 + j3 * ne3) * ne00 * ne01 * ne02]; + } + } + } + } + + dst_dd[i] = acc; + }); +} diff --git a/ggml/src/ggml-sycl/repeat_back.hpp b/ggml/src/ggml-sycl/repeat_back.hpp new file mode 100644 index 0000000000000..6783dead4fe1c --- /dev/null +++ b/ggml/src/ggml-sycl/repeat_back.hpp @@ -0,0 +1,9 @@ +//בס"ד +#ifndef GGML_SYCL_REPEAT_BACK_HPP +#define GGML_SYCL_REPEAT_BACK_HPP + +#include "common.hpp" + +void ggml_sycl_op_repeat_back(ggml_backend_sycl_context & ctx, ggml_tensor * dst); + +#endif // GGML_SYCL_REPEAT_BACK_HPP From d3e88bc5753e1caa4fc0c26697676e9a6a66d60a Mon Sep 17 00:00:00 2001 From: shani-f Date: Thu, 23 Oct 2025 12:40:43 +0300 Subject: [PATCH 2/5] Implement repeat_back SYCL operation and minor fixes --- ggml/src/ggml-sycl/repeat_back.cpp | 13 ++++--------- 1 file changed, 4 insertions(+), 9 deletions(-) diff --git a/ggml/src/ggml-sycl/repeat_back.cpp b/ggml/src/ggml-sycl/repeat_back.cpp index 10fef098cdd93..48ae998a27543 100644 --- a/ggml/src/ggml-sycl/repeat_back.cpp +++ b/ggml/src/ggml-sycl/repeat_back.cpp @@ -4,20 +4,16 @@ #include "common.hpp" void ggml_sycl_op_repeat_back(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { + GGML_ASSERT(dst->src[0]->type == GGML_TYPE_F32); GGML_ASSERT(dst->type == GGML_TYPE_F32); const float * src0_dd = (const float *) dst->src[0]->data; float * dst_dd = (float *) dst->data; - const int64_t ne0 = dst->ne[0]; - const int64_t ne1 = dst->ne[1]; - const int64_t ne2 = dst->ne[2]; - const int64_t ne3 = dst->ne[3]; - const int64_t ne00 = dst->src[0]->ne[0]; - const int64_t ne01 = dst->src[0]->ne[1]; - const int64_t ne02 = dst->src[0]->ne[2]; - const int64_t ne03 = dst->src[0]->ne[3]; + const int64_t ne0 = dst->ne[0], ne1 = dst->ne[1], ne2 = dst->ne[2], ne3 = dst->ne[3]; + const int64_t ne00 = dst->src[0]->ne[0], ne01 = dst->src[0]->ne[1], ne02 = dst->src[0]->ne[2], + ne03 = dst->src[0]->ne[3]; const int nr0 = (int) (ne00 / ne0); const int nr1 = (int) (ne01 / ne1); @@ -29,7 +25,6 @@ void ggml_sycl_op_repeat_back(ggml_backend_sycl_context & ctx, ggml_tensor * dst const int num_blocks = (total + BLOCK_SIZE - 1) / BLOCK_SIZE; queue_ptr stream = ctx.stream(); - stream->memset(dst_dd, 0, total * sizeof(float)); stream->parallel_for( sycl::nd_range<1>(sycl::range<1>(num_blocks * BLOCK_SIZE), sycl::range<1>(BLOCK_SIZE)), From 020dba5b919bd9c51021676391d491bf12b8ebe5 Mon Sep 17 00:00:00 2001 From: shani-f Date: Sun, 26 Oct 2025 12:37:48 +0200 Subject: [PATCH 3/5] Update ggml/src/ggml-sycl/repeat_back.cpp MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-authored-by: Sigbjørn Skjæret --- ggml/src/ggml-sycl/repeat_back.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/ggml/src/ggml-sycl/repeat_back.cpp b/ggml/src/ggml-sycl/repeat_back.cpp index 48ae998a27543..abcd4cee72a48 100644 --- a/ggml/src/ggml-sycl/repeat_back.cpp +++ b/ggml/src/ggml-sycl/repeat_back.cpp @@ -1,4 +1,3 @@ -//בס"ד #include "repeat_back.hpp" #include "common.hpp" From 5c36f73ff2f3e3149b86dfdc078d7c022d77314c Mon Sep 17 00:00:00 2001 From: shani-f Date: Sun, 26 Oct 2025 12:38:13 +0200 Subject: [PATCH 4/5] Update ggml/src/ggml-sycl/repeat_back.hpp MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-authored-by: Sigbjørn Skjæret --- ggml/src/ggml-sycl/repeat_back.hpp | 1 - 1 file changed, 1 deletion(-) diff --git a/ggml/src/ggml-sycl/repeat_back.hpp b/ggml/src/ggml-sycl/repeat_back.hpp index 6783dead4fe1c..17a87f3e159b3 100644 --- a/ggml/src/ggml-sycl/repeat_back.hpp +++ b/ggml/src/ggml-sycl/repeat_back.hpp @@ -1,4 +1,3 @@ -//בס"ד #ifndef GGML_SYCL_REPEAT_BACK_HPP #define GGML_SYCL_REPEAT_BACK_HPP From a18fbd7f0cabfaf1a396dc6a2522decb793d6d2f Mon Sep 17 00:00:00 2001 From: shani-f Date: Sun, 26 Oct 2025 12:38:42 +0200 Subject: [PATCH 5/5] Update ggml/src/ggml-sycl/ggml-sycl.cpp MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-authored-by: Sigbjørn Skjæret --- ggml/src/ggml-sycl/ggml-sycl.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index b43f37478dc36..5eba49ab12484 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -4413,7 +4413,6 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g } return false; } - case GGML_OP_CONCAT: { ggml_type src0_type = op->src[0]->type;