[SYCL] fix unsupport ACC UT cases for noncontiguous (#25124)

* fix unsupport ACC UT cases for noncontiguous

* update ops.md

diff --git a/ggml/src/ggml-sycl/element_wise.cpp b/ggml/src/ggml-sycl/element_wise.cpp
index 0c82ceb969f45c31d76d6fa95c9e51899d2cd592..bae157a487a1f20c4d6cdf7f2a5911dca09ccd8e 100644 (file)
--- a/ggml/src/ggml-sycl/element_wise.cpp
+++ b/ggml/src/ggml-sycl/element_wise.cpp
@@ -9,9 +9,12 @@
 #define SYCL_LOCAL_ID_CALC(ITEM, IDX) \
     (ITEM.get_local_range(IDX) * ITEM.get_group(IDX) + ITEM.get_local_id(IDX))
 
-static void acc_f32(const float * x, const float * y, float * dst, const int64_t ne,
-        const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
-        const int64_t s11, const int64_t s12, const int64_t s13, const int64_t offset) {
+static void acc_f32(const char * x, const char * y, float * dst, const int64_t ne,
+    const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3,
+    const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03,
+    const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
+    const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
+    const int64_t s11, const int64_t s12, const int64_t s13, const int64_t offset) {
     auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
     const int64_t i = SYCL_LOCAL_ID_CALC(item_ct1, 2);
 
@@ -30,9 +33,18 @@ static void acc_f32(const float * x, const float * y, float * dst, const int64_t
     tmp -= i11 * s11;
     const int64_t i10 = tmp;
 
-    float val = x[i];
+    int64_t tmp_dst = i;
+    const int64_t i3 = tmp_dst / (ne2*ne1*ne0);
+    tmp_dst -= i3 * (ne2*ne1*ne0);
+    const int64_t i2 = tmp_dst / (ne1*ne0);
+    tmp_dst -= i2 * (ne1*ne0);
+    const int64_t i1 = tmp_dst / ne0;
+    tmp_dst -= i1 * ne0;
+    const int64_t i0 = tmp_dst;
+
+    float val = *(const float *) (x + i0*nb00 + i1*nb01 + i2*nb02 + i3*nb03);
     if (src1_idx >= 0 && i10 < ne10 && i11 < ne11 && i12 < ne12 && i13 < ne13) {
-        val += y[((i13*ne12 + i12) * ne11 + i11) * ne10 + i10];
+        val += *(const float *) (y + i10*nb10 + i11*nb11 + i12*nb12 + i13*nb13);
     }
     dst[i] = val;
 }
@@ -422,15 +434,24 @@ static void gated_op_fused_geglu_quick(const T * x, const T * g, T * dst, const
 }
 
 namespace ggml_sycl_detail {
-static void acc_f32_sycl(const float *x, const float *y, float *dst,
-                         const int64_t n_elements, const int64_t ne10, const int64_t ne11,
-                         const int64_t ne12, const int64_t ne13, const int64_t s1, const int64_t s2, const int64_t s3,
+static void acc_f32_sycl(const char *x, const char *y, float *dst,
+                         const int64_t n_elements,
+                         const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3,
+                         const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03,
+                         const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
+                         const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
+                         const int64_t s1, const int64_t s2, const int64_t s3,
                          const int64_t offset, queue_ptr stream) {
     const int num_blocks = (n_elements + SYCL_ACC_BLOCK_SIZE - 1) / SYCL_ACC_BLOCK_SIZE;
     stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_ACC_BLOCK_SIZE),
                                            sycl::range<3>(1, 1, SYCL_ACC_BLOCK_SIZE)),
                          [=](sycl::nd_item<3> /*item_ct1*/) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                             acc_f32(x, y, dst, n_elements, ne10, ne11, ne12, ne13, s1, s2, s3, offset);
+                            acc_f32(x, y, dst, n_elements,
+                                ne0, ne1, ne2, ne3,
+                                nb00, nb01, nb02, nb03,
+                                ne10, ne11, ne12, ne13,
+                                nb10, nb11, nb12, nb13,
+                                s1, s2, s3, offset);
                          });
 }
 
@@ -843,8 +864,8 @@ static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor
     const ggml_tensor * src0 = dst->src[0];
     const ggml_tensor * src1 = dst->src[1];
 
-    const float * src0_d = (const float *) src0->data;
-    const float * src1_d = (const float *) src1->data;
+    const char  * src0_d = (const char  *) src0->data;
+    const char  * src1_d = (const char  *) src1->data;
     float       * dst_d  = (float       *)  dst->data;
 
     dpct::queue_ptr stream = ctx.stream();
@@ -853,17 +874,20 @@ static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor
     GGML_ASSERT(src1->type == GGML_TYPE_F32);
     GGML_ASSERT( dst->type == GGML_TYPE_F32);
 
-    GGML_ASSERT(ggml_is_contiguous(src1));
     GGML_ASSERT(dst->nb[0] == ggml_element_size(dst));
     GGML_ASSERT(ggml_is_contiguously_allocated(dst));
+    GGML_ASSERT(ggml_are_same_shape(src0, dst));
 
-    const int64_t s1     = dst->op_params[0] / sizeof(float);
-    const int64_t s2     = dst->op_params[1] / sizeof(float);
-    const int64_t s3     = dst->op_params[2] / sizeof(float);
-    const int64_t offset = dst->op_params[3] / sizeof(float);
+    const int64_t s1     = (int64_t) ((const int32_t *) dst->op_params)[0] / (int64_t) sizeof(float);
+    const int64_t s2     = (int64_t) ((const int32_t *) dst->op_params)[1] / (int64_t) sizeof(float);
+    const int64_t s3     = (int64_t) ((const int32_t *) dst->op_params)[2] / (int64_t) sizeof(float);
+    const int64_t offset = (int64_t) ((const int32_t *) dst->op_params)[3] / (int64_t) sizeof(float);
 
     ggml_sycl_detail::acc_f32_sycl(src0_d, src1_d, dst_d, ggml_nelements(dst),
+        dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3],
+        src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
         src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3],
+        src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],
         s1, s2, s3, offset, stream);
 }
 

diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index 4cb24a7227761bba92218391a74b8294a4460c8a..0b7576ca3f9506118395b10475229ddedbe8fa9a 100644 (file)
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -5852,9 +5852,8 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
         }
         case GGML_OP_POOL_2D:
         case GGML_OP_POOL_1D:
-            return true;
         case GGML_OP_ACC:
-            return ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op->src[1]);
+            return true;
         case GGML_OP_PAD:
             if (ggml_get_op_params_i32(op, 8) != 0) {
                 return false;