mirror of
https://github.com/ggml-org/llama.cpp.git
synced 2026-07-08 12:10:25 +00:00
[SYCL] support op col2im_1d (#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
This commit is contained in:
Neo Zhang
+1
-1
@@ -24,7 +24,7 @@ Legend:
|
||||
| ARGSORT | ❌ | ✅ | ✅ | ✅ | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
|
||||
| CEIL | ❌ | ❌ | ✅ | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
|
||||
| CLAMP | ❌ | ✅ | ✅ | ✅ | ✅ | 🟡 | ✅ | 🟡 | ✅ | ❌ | ❌ |
|
||||
| COL2IM_1D | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
|
||||
| COL2IM_1D | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
|
||||
| CONCAT | ❌ | ✅ | ✅ | 🟡 | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
|
||||
| CONT | ❌ | 🟡 | ✅ | ✅ | ✅ | 🟡 | 🟡 | ✅ | 🟡 | ❌ | ❌ |
|
||||
| CONV_2D | ❌ | ❌ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
|
||||
|
||||
+33
-33
@@ -5053,39 +5053,39 @@
|
||||
"SYCL0","CONV_TRANSPOSE_1D","ne_input=[3,2,1,1],ne_kernel=[3,2,2,1],s0=1,p0=0,d0=1","support","1","yes","SYCL"
|
||||
"SYCL0","CONV_TRANSPOSE_1D","ne_input=[3,2,1,1],ne_kernel=[3,1,2,1],s0=1,p0=0,d0=1","support","1","yes","SYCL"
|
||||
"SYCL0","CONV_TRANSPOSE_1D","ne_input=[2,1,1,1],ne_kernel=[3,1,1,1],s0=1,p0=0,d0=1","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=16,OC=32,T_in=197,s0=8,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=4,OC=3,T_in=7,s0=2,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=1,OC=5,T_in=13,s0=1,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=6,OC=4,T_in=11,s0=3,p0=1","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=2,OC=3,T_in=9,s0=3,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=5,OC=4,T_in=11,s0=2,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=8,OC=4,T_in=13,s0=4,p0=2","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=4,OC=3,T_in=1,s0=2,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=16,OC=1,T_in=197,s0=8,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=1,OC=5,T_in=13,s0=3,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=8,OC=2,T_in=3,s0=2,p0=5","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=16,OC=32,T_in=197,s0=8,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=4,OC=3,T_in=7,s0=2,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=1,OC=5,T_in=13,s0=1,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=6,OC=4,T_in=11,s0=3,p0=1","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=2,OC=3,T_in=9,s0=3,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=5,OC=4,T_in=11,s0=2,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=8,OC=4,T_in=13,s0=4,p0=2","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=4,OC=3,T_in=1,s0=2,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=16,OC=1,T_in=197,s0=8,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=1,OC=5,T_in=13,s0=3,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=8,OC=2,T_in=3,s0=2,p0=5","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=16,OC=32,T_in=197,s0=8,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=4,OC=3,T_in=7,s0=2,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=1,OC=5,T_in=13,s0=1,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=6,OC=4,T_in=11,s0=3,p0=1","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=2,OC=3,T_in=9,s0=3,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=5,OC=4,T_in=11,s0=2,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=8,OC=4,T_in=13,s0=4,p0=2","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=4,OC=3,T_in=1,s0=2,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=16,OC=1,T_in=197,s0=8,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=1,OC=5,T_in=13,s0=3,p0=0","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=8,OC=2,T_in=3,s0=2,p0=5","support","0","no","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=16,OC=32,T_in=197,s0=8,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=4,OC=3,T_in=7,s0=2,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=1,OC=5,T_in=13,s0=1,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=6,OC=4,T_in=11,s0=3,p0=1","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=2,OC=3,T_in=9,s0=3,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=5,OC=4,T_in=11,s0=2,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=8,OC=4,T_in=13,s0=4,p0=2","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=4,OC=3,T_in=1,s0=2,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=16,OC=1,T_in=197,s0=8,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=1,OC=5,T_in=13,s0=3,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f32,K=8,OC=2,T_in=3,s0=2,p0=5","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=16,OC=32,T_in=197,s0=8,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=4,OC=3,T_in=7,s0=2,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=1,OC=5,T_in=13,s0=1,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=6,OC=4,T_in=11,s0=3,p0=1","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=2,OC=3,T_in=9,s0=3,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=5,OC=4,T_in=11,s0=2,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=8,OC=4,T_in=13,s0=4,p0=2","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=4,OC=3,T_in=1,s0=2,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=16,OC=1,T_in=197,s0=8,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=1,OC=5,T_in=13,s0=3,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=f16,K=8,OC=2,T_in=3,s0=2,p0=5","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=16,OC=32,T_in=197,s0=8,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=4,OC=3,T_in=7,s0=2,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=1,OC=5,T_in=13,s0=1,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=6,OC=4,T_in=11,s0=3,p0=1","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=2,OC=3,T_in=9,s0=3,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=5,OC=4,T_in=11,s0=2,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=8,OC=4,T_in=13,s0=4,p0=2","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=4,OC=3,T_in=1,s0=2,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=16,OC=1,T_in=197,s0=8,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=1,OC=5,T_in=13,s0=3,p0=0","support","1","yes","SYCL"
|
||||
"SYCL0","COL2IM_1D","type=bf16,K=8,OC=2,T_in=3,s0=2,p0=5","support","1","yes","SYCL"
|
||||
"SYCL0","CONV_TRANSPOSE_2D","kernel_type=f32,ne_input=[3,2,3,1],ne_kernel=[2,2,1,3],stride=1","support","1","yes","SYCL"
|
||||
"SYCL0","CONV_TRANSPOSE_2D","kernel_type=f32,ne_input=[10,10,9,1],ne_kernel=[3,3,1,9],stride=2","support","1","yes","SYCL"
|
||||
"SYCL0","CONV_TRANSPOSE_2D","kernel_type=f32,ne_input=[129,63,35,1],ne_kernel=[3,3,48,35],stride=1","support","1","yes","SYCL"
|
||||
|
||||
|
Can't render this file because it is too large.
|
@@ -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"
|
||||
|
||||
@@ -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);
|
||||
}
|
||||
}
|
||||
@@ -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
|
||||
@@ -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) &&
|
||||
|
||||
@@ -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
|
||||
|
||||
Reference in New Issue
Block a user