support op col2im_1d (llama/25264)

* support op col2im_1d

* update ops.md

* rm unused words

* update for bf16

* optimize 1%-11% as the review comments

* fix the format issue

* update as the review comments

diff --git a/ggml/src/ggml-sycl/backend.hpp b/ggml/src/ggml-sycl/backend.hpp
index 1f5a91272663182102033b7d3ac55ea551def4fc..2d92a95661e854ccc29f8e6e295656fedb412124 100644 (file)
--- a/ggml/src/ggml-sycl/backend.hpp
+++ b/ggml/src/ggml-sycl/backend.hpp
@@ -14,6 +14,7 @@
 #define GGML_SYCL_BACKEND_HPP
 
 #include "binbcast.hpp"
+#include "col2im-1d.hpp"
 #include "common.hpp"
 #include "concat.hpp"
 #include "conv.hpp"

diff --git a/ggml/src/ggml-sycl/col2im-1d.cpp b/ggml/src/ggml-sycl/col2im-1d.cpp
new file mode 100644 (file)
index 0000000..c4f0900
--- /dev/null
+++ b/ggml/src/ggml-sycl/col2im-1d.cpp
@@ -0,0 +1,102 @@
+#include "col2im-1d.hpp"
+
+template <typename T>
+static void col2im_1d_sycl(
+        const T * col,
+        T * dst,
+        const int T_in,
+        const sycl::uint3 T_out_fd,
+        const int K,
+        const int K_OC,
+        const int32_t s0,
+        const int32_t p0,
+        const int total,
+        dpct::queue_ptr stream) {
+
+    const uint32_t block_size = SYCL_COL2IM_1D_BLOCK_SIZE;
+    const uint32_t num_blocks = (uint32_t) ((total + block_size - 1) / block_size);
+
+    stream->parallel_for(
+        sycl::nd_range<3>(
+            sycl::range<3>(1, 1, num_blocks * block_size),
+            sycl::range<3>(1, 1, block_size)),
+        [=](sycl::nd_item<3> item_ct1) {
+            const int idx = (int) item_ct1.get_global_id(2);
+            if (idx >= total) {
+                return;
+            }
+
+            const sycl::uint2 qr = fast_div_modulo((uint32_t) idx, T_out_fd);
+            const int oc    = (int) qr.x();
+            const int t_out = (int) qr.y();
+            const int t_abs = t_out + p0;
+
+            int t_in_min = (t_abs - K + s0) / s0;
+            if (t_in_min < 0) {
+                t_in_min = 0;
+            }
+            int t_in_max = t_abs / s0;
+            if (t_in_max >= T_in) {
+                t_in_max = T_in - 1;
+            }
+
+            float sum = 0.0f;
+            for (int t_in = t_in_min; t_in <= t_in_max; ++t_in) {
+                const int k = t_abs - t_in * s0;
+                sum += static_cast<float>(col[(oc * K + k) + t_in * K_OC]);
+            }
+
+            dst[idx] = static_cast<T>(sum);
+        });
+}
+
+void ggml_sycl_op_col2im_1d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    const ggml_tensor * src0 = dst->src[0];
+
+    GGML_ASSERT(src0 != nullptr);
+    GGML_ASSERT(ggml_is_contiguous(src0));
+    GGML_ASSERT(src0->type == dst->type);
+
+    const int32_t s0 = ((const int32_t *) dst->op_params)[0];
+    const int32_t OC = ((const int32_t *) dst->op_params)[1];
+    const int32_t p0 = ((const int32_t *) dst->op_params)[2];
+
+    const int K_OC  = (int) src0->ne[0];
+    const int T_in  = (int) src0->ne[1];
+    const int K     = K_OC / OC;
+    const int T_out = (int) dst->ne[0];
+
+    GGML_ASSERT(OC > 0);
+    GGML_ASSERT(K_OC % OC == 0);
+
+    const sycl::uint3 T_out_fd = init_fastdiv_values((uint32_t) T_out);
+
+    const int total = T_out * OC;
+
+    dpct::queue_ptr stream = ctx.stream();
+
+    switch (src0->type) {
+        case GGML_TYPE_F32:
+            col2im_1d_sycl<float>(
+                (const float *) src0->data,
+                (float *) dst->data,
+                T_in, T_out_fd, K, K_OC, s0, p0, total, stream);
+            break;
+        case GGML_TYPE_F16:
+            col2im_1d_sycl<sycl::half>(
+                (const sycl::half *) src0->data,
+                (sycl::half *) dst->data,
+                T_in, T_out_fd, K, K_OC, s0, p0, total, stream);
+            break;
+#ifdef GGML_SYCL_HAS_BF16
+        case GGML_TYPE_BF16:
+            col2im_1d_sycl<sycl::ext::oneapi::bfloat16>(
+                (const sycl::ext::oneapi::bfloat16 *) src0->data,
+                (sycl::ext::oneapi::bfloat16 *) dst->data,
+                T_in, T_out_fd, K, K_OC, s0, p0, total, stream);
+            break;
+#endif
+        default:
+            GGML_ABORT("col2im_1d: unsupported type %d", src0->type);
+    }
+}

diff --git a/ggml/src/ggml-sycl/col2im-1d.hpp b/ggml/src/ggml-sycl/col2im-1d.hpp
new file mode 100644 (file)
index 0000000..cfb53da
--- /dev/null
+++ b/ggml/src/ggml-sycl/col2im-1d.hpp
@@ -0,0 +1,8 @@
+#ifndef GGML_SYCL_COL2IM_1D_HPP
+#define GGML_SYCL_COL2IM_1D_HPP
+
+#include "common.hpp"
+
+void ggml_sycl_op_col2im_1d(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
+
+#endif // GGML_SYCL_COL2IM_1D_HPP

diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index c1d2eb97eceee71f212eac9f841c18d7322e9acf..74831ce8884aec062eab8e5a9cfe1773ed021227 100644 (file)
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -4771,6 +4771,11 @@ static void ggml_sycl_im2col_3d(ggml_backend_sycl_context & ctx, ggml_tensor * d
     ggml_sycl_op_im2col_3d(ctx, dst);
 }
 
+static void ggml_sycl_col2im_1d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+    scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/1);
+    ggml_sycl_op_col2im_1d(ctx, dst);
+}
+
 static void ggml_sycl_conv_3d(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
     scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/2);
     ggml_sycl_op_conv_3d(ctx, dst);
@@ -5097,6 +5102,9 @@ static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct gg
         case GGML_OP_IM2COL_3D:
             ggml_sycl_im2col_3d(ctx, dst);
             break;
+        case GGML_OP_COL2IM_1D:
+            ggml_sycl_col2im_1d(ctx, dst);
+            break;
         case GGML_OP_POOL_2D:
             ggml_sycl_pool2d(ctx, dst);
             break;
@@ -5646,7 +5654,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
                 // TODO: This specific configuration can fail with oneDNN and needs more debugging
                 if (!ggml_is_permuted(a) && ggml_is_permuted(b) && b->ne[2] > 1 && b->ne[3] > 1 &&
                     a->ne[0] > 128 && a->ne[2] == 1 && src0_type == GGML_TYPE_F16) {
-                        printf("zjy 2\n");
                     return false;
                 }
                 return true;
@@ -5842,6 +5849,14 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
         case GGML_OP_IM2COL_3D:
         case GGML_OP_UPSCALE:
             return true;
+        case GGML_OP_COL2IM_1D:
+            return ggml_is_contiguous(op->src[0]) &&
+                   (op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16
+#ifdef GGML_SYCL_HAS_BF16
+                    || op->type == GGML_TYPE_BF16
+#endif
+                   ) &&
+                   op->src[0]->type == op->type;
         case GGML_OP_CONV_3D:
             return op->type == GGML_TYPE_F32 &&
                    (op->src[0]->type == GGML_TYPE_F32 || op->src[0]->type == GGML_TYPE_F16) &&

diff --git a/ggml/src/ggml-sycl/presets.hpp b/ggml/src/ggml-sycl/presets.hpp
index 3bd1c86eae81aa3dfeb9cd6f98db6194fb8969b4..502e3b6105061ebe4186b4885c1478463855b5dd 100644 (file)
--- a/ggml/src/ggml-sycl/presets.hpp
+++ b/ggml/src/ggml-sycl/presets.hpp
@@ -19,6 +19,7 @@
 #define WARP_SIZE GGML_SYCL_WARP_SIZE
 #define MATRIX_ROW_PADDING 512 // last row of quant. matrices is a multiple of this to avoid out-of-bounds memory accesses
 
+#define SYCL_COL2IM_1D_BLOCK_SIZE 256
 #define SYCL_GELU_BLOCK_SIZE 256
 #define SYCL_SILU_BLOCK_SIZE 256
 #define SYCL_TANH_BLOCK_SIZE 256