Skip to content

Commit 75afa0a

Browse files
qnixsynapseqnixsynapse
authored
SYCL: Fix im2col (#12910)
* SYCL: Fix im2col * restore local workgroup size adjustments for large inputs * restore format
1 parent c772d54 commit 75afa0a

File tree

2 files changed

+83
-69
lines changed

2 files changed

+83
-69
lines changed

‎ggml/src/ggml-sycl/ggml-sycl.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4018,8 +4018,7 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g
40184018
return ggml_is_contiguous(op->src[0]);
40194019
}
40204020
case GGML_OP_IM2COL:
4021-
// TODO: add support for the new F32 operations
4022-
return op->src[0]->type == GGML_TYPE_F16;
4021+
return true;
40234022
case GGML_OP_UPSCALE:
40244023
return op->src[0]->type == GGML_TYPE_F32 && op->op_params[0] == GGML_SCALE_MODE_NEAREST;
40254024
case GGML_OP_POOL_2D:

‎ggml/src/ggml-sycl/im2col.cpp

Lines changed: 82 additions & 67 deletions
Original file line numberDiff line numberDiff line change
@@ -12,110 +12,125 @@
1212

1313
#include "im2col.hpp"
1414

15+
#include <sycl/sycl.hpp>
16+
#include <type_traits> // For std::is_same_v
17+
18+
#include "ggml.h"
19+
1520
template <typename T>
16-
static void im2col_kernel(
17-
const float *x, T *dst, int64_t batch_offset, int64_t offset_delta,
18-
int64_t IC, int64_t IW, int64_t IH, int64_t OH, int64_t OW, int64_t KW, int64_t KH,
19-
int64_t pelements, int64_t CHW, int s0, int s1, int p0, int p1, int d0, int d1,
20-
const sycl::nd_item<3> &item_ct1) {
21+
static void im2col_kernel(const float * x, T * dst, int64_t batch_offset, int64_t offset_delta, int64_t IC, int64_t IW,
22+
int64_t IH, int64_t OH, int64_t OW, int64_t KW, int64_t KH, int64_t pelements, int64_t CHW,
23+
int s0, int s1, int p0, int p1, int d0, int d1, const sycl::nd_item<3> & item_ct1) {
2124
const int64_t work_group_size = item_ct1.get_local_range(2);
22-
const int64_t global_id = item_ct1.get_local_id(2) + work_group_size * item_ct1.get_group(2);
25+
const int64_t global_id = item_ct1.get_local_id(2) + (work_group_size * item_ct1.get_group(2));
2326

2427
// make each work-item deal with more elements since sycl global range can not exceed max int
25-
for (int64_t i = global_id; i < pelements; i += work_group_size * item_ct1.get_group_range(2)) {
26-
28+
for (int64_t i = global_id; i < pelements; i += (work_group_size * item_ct1.get_group_range(2))) {
2729
const int64_t ksize = OW * (KH > 1 ? KW : 1);
28-
const int64_t kx = i / ksize;
29-
const int64_t kd = kx * ksize;
30-
const int64_t ky = (i - kd) / OW;
31-
const int64_t ix = i % OW;
32-
33-
const int64_t oh = item_ct1.get_group(1);
34-
const int64_t batch = item_ct1.get_group(0) / IC;
35-
const int64_t ic = item_ct1.get_group(0) % IC;
36-
37-
const int64_t iiw = ix * s0 + kx * d0 - p0;
38-
const int64_t iih = oh * s1 + ky * d1 - p1;
39-
40-
const int64_t offset_dst =
41-
((batch * OH + oh) * OW + ix) * CHW +
42-
(ic * (KW * KH) + ky * KW + kx);
43-
44-
if (iih < 0 || iih >= IH || iiw < 0 || iiw >= IW) {
45-
dst[offset_dst] =
46-
sycl::vec<float, 1>(0.0f)
47-
.convert<sycl::half, sycl::rounding_mode::automatic>()[0];
48-
} else {
49-
const int64_t offset_src = ic * offset_delta + batch * batch_offset;
50-
dst[offset_dst] =
51-
sycl::vec<float, 1>(x[offset_src + iih * IW + iiw])
52-
.convert<sycl::half, sycl::rounding_mode::automatic>()[0];
30+
const int64_t kx = i / ksize;
31+
const int64_t kd = kx * ksize;
32+
const int64_t ky = (i - kd) / OW;
33+
const int64_t ix = i % OW;
34+
35+
const int64_t oh = item_ct1.get_group(1);
36+
const int64_t batch = item_ct1.get_group(0) / IC;
37+
const int64_t ic = item_ct1.get_group(0) % IC;
38+
39+
const int64_t iiw = (ix * s0) + (kx * d0) - p0;
40+
const int64_t iih = (oh * s1) + (ky * d1) - p1;
41+
42+
const int64_t offset_dst = (((batch * OH + oh) * OW + ix) * CHW) + (ic * (KW * KH) + ky * KW + kx);
43+
44+
const int64_t offset_src_base = (ic * offset_delta) + (batch * batch_offset);
45+
const int64_t offset_src = offset_src_base + (iih * IW) + iiw;
46+
47+
const bool out_of_bounds = (iih < 0 || iih >= IH || iiw < 0 || iiw >= IW);
48+
const float src_val = out_of_bounds ? 0.0f : x[offset_src];
49+
50+
if constexpr (std::is_same_v<T, sycl::half>) {
51+
dst[offset_dst] = sycl::half(src_val);
52+
} else if constexpr (std::is_same_v<T, float>) {
53+
dst[offset_dst] = src_val;
5354
}
5455
}
5556
}
5657

5758
template <typename T>
58-
static void im2col_sycl(
59-
const float *x, T *dst, int64_t IW, int64_t IH, int64_t OW, int64_t OH, int64_t KW,
60-
int64_t KH, int64_t IC, int64_t batch, int64_t batch_offset, int64_t offset_delta,
61-
int s0, int s1, int p0, int p1, int d0, int d1,
62-
queue_ptr stream) {
59+
static void im2col_sycl_internal(const float * x, T * dst, int64_t IW, int64_t IH, int64_t OW, int64_t OH, int64_t KW,
60+
int64_t KH, int64_t IC, int64_t batch, int64_t batch_offset, int64_t offset_delta,
61+
int s0, int s1, int p0, int p1, int d0, int d1, queue_ptr stream) {
6362
const int64_t parallel_elements = OW * KW * KH;
64-
const int64_t num_blocks = (parallel_elements + SYCL_IM2COL_BLOCK_SIZE - 1) / SYCL_IM2COL_BLOCK_SIZE;
63+
const int64_t num_blocks = (parallel_elements + SYCL_IM2COL_BLOCK_SIZE - 1) / SYCL_IM2COL_BLOCK_SIZE;
6564

6665
// decrease global range when it exceeds the max int
6766
int64_t local_size = downsample_sycl_global_range(batch * IC * OH * num_blocks, SYCL_IM2COL_BLOCK_SIZE);
67+
6868
sycl::range<3> block_nums(batch * IC, OH, num_blocks);
6969
sycl::range<3> local_range(1, 1, local_size);
7070

71-
{
72-
dpct::has_capability_or_fail(stream->get_device(),
73-
{sycl::aspect::fp16});
74-
75-
stream->parallel_for(
76-
sycl::nd_range<3>(block_nums * local_range, local_range),
77-
[=](sycl::nd_item<3> item_ct1) {
78-
im2col_kernel(x, dst, batch_offset, offset_delta, IC, IW, IH, OH, OW, KW, KH,
79-
parallel_elements, (IC * KH * KW), s0, s1, p0,
80-
p1, d0, d1, item_ct1);
81-
});
71+
const int64_t CHW = IC * KH * KW;
72+
73+
stream->parallel_for(sycl::nd_range<3>(block_nums * local_range, local_range), [=](sycl::nd_item<3> item_ct1) {
74+
im2col_kernel<T>(x, dst, batch_offset, offset_delta, IC, IW, IH, OH, OW, KW, KH, parallel_elements, CHW, s0, s1,
75+
p0, p1, d0, d1, item_ct1);
76+
});
77+
}
78+
79+
static void im2col_sycl_f16(const float * x, sycl::half * dst, int64_t IW, int64_t IH, int64_t OW, int64_t OH,
80+
int64_t KW, int64_t KH, int64_t IC, int64_t batch, int64_t batch_offset,
81+
int64_t offset_delta, int s0, int s1, int p0, int p1, int d0, int d1, queue_ptr stream) {
82+
if (!stream->get_device().has(sycl::aspect::fp16)) {
83+
throw sycl::exception(sycl::make_error_code(sycl::errc::kernel_not_supported),
84+
"Device does not support half precision (fp16) operations!");
8285
}
86+
im2col_sycl_internal<sycl::half>(x, dst, IW, IH, OW, OH, KW, KH, IC, batch, batch_offset, offset_delta, s0, s1, p0,
87+
p1, d0, d1, stream);
88+
}
89+
90+
static void im2col_sycl_f32(const float * x, float * dst, int64_t IW, int64_t IH, int64_t OW, int64_t OH, int64_t KW,
91+
int64_t KH, int64_t IC, int64_t batch, int64_t batch_offset, int64_t offset_delta, int s0,
92+
int s1, int p0, int p1, int d0, int d1, queue_ptr stream) {
93+
im2col_sycl_internal<float>(x, dst, IW, IH, OW, OH, KW, KH, IC, batch, batch_offset, offset_delta, s0, s1, p0, p1,
94+
d0, d1, stream);
8395
}
8496

85-
void ggml_sycl_op_im2col(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
97+
void ggml_sycl_op_im2col(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
8698
const ggml_tensor * src0 = dst->src[0];
8799
const ggml_tensor * src1 = dst->src[1];
88100

89-
GGML_ASSERT(src0->type == GGML_TYPE_F16);
90101
GGML_ASSERT(src1->type == GGML_TYPE_F32);
91102
GGML_ASSERT(dst->type == GGML_TYPE_F16 || dst->type == GGML_TYPE_F32);
92103

93-
const int32_t s0 = ((const int32_t*)(dst->op_params))[0];
94-
const int32_t s1 = ((const int32_t*)(dst->op_params))[1];
95-
const int32_t p0 = ((const int32_t*)(dst->op_params))[2];
96-
const int32_t p1 = ((const int32_t*)(dst->op_params))[3];
97-
const int32_t d0 = ((const int32_t*)(dst->op_params))[4];
98-
const int32_t d1 = ((const int32_t*)(dst->op_params))[5];
104+
const int32_t s0 = ((const int32_t *) (dst->op_params))[0];
105+
const int32_t s1 = ((const int32_t *) (dst->op_params))[1];
106+
const int32_t p0 = ((const int32_t *) (dst->op_params))[2];
107+
const int32_t p1 = ((const int32_t *) (dst->op_params))[3];
108+
const int32_t d0 = ((const int32_t *) (dst->op_params))[4];
109+
const int32_t d1 = ((const int32_t *) (dst->op_params))[5];
99110

100-
const bool is_2D = ((const int32_t*)(dst->op_params))[6] == 1;
111+
const bool is_2D = ((const int32_t *) (dst->op_params))[6] == 1;
101112

102113
const int64_t IC = src1->ne[is_2D ? 2 : 1];
103114
const int64_t IH = is_2D ? src1->ne[1] : 1;
104-
const int64_t IW = src1->ne[0];
115+
const int64_t IW = src1->ne[0];
105116

106117
const int64_t KH = is_2D ? src0->ne[1] : 1;
107-
const int64_t KW = src0->ne[0];
118+
const int64_t KW = src0->ne[0];
108119

109120
const int64_t OH = is_2D ? dst->ne[2] : 1;
110-
const int64_t OW = dst->ne[1];
121+
const int64_t OW = dst->ne[1];
122+
123+
const size_t delta_offset = src1->nb[is_2D ? 2 : 1] / sizeof(float);
124+
const int64_t batch = src1->ne[is_2D ? 3 : 2];
125+
const size_t batch_offset = src1->nb[is_2D ? 3 : 2] / sizeof(float);
111126

112-
const size_t delta_offset = src1->nb[is_2D ? 2 : 1] / 4; // nb is byte offset, src is type float32
113-
const int64_t batch = src1->ne[3];
114-
const size_t batch_offset = src1->nb[3] / 4; // nb is byte offset, src is type float32
127+
queue_ptr stream = ctx.stream();
115128

116129
if (dst->type == GGML_TYPE_F16) {
117-
im2col_sycl((const float *) src1->data, (sycl::half *)dst->data, IW, IH, OW, OH, KW, KH, IC, batch, batch_offset, delta_offset, s0, s1, p0, p1, d0, d1, ctx.stream());
130+
im2col_sycl_f16((const float *) src1->data, (sycl::half *) dst->data, IW, IH, OW, OH, KW, KH, IC, batch,
131+
batch_offset, delta_offset, s0, s1, p0, p1, d0, d1, stream);
118132
} else {
119-
im2col_sycl((const float *) src1->data, (float *)dst->data, IW, IH, OW, OH, KW, KH, IC, batch, batch_offset, delta_offset, s0, s1, p0, p1, d0, d1, ctx.stream());
133+
im2col_sycl_f32((const float *) src1->data, (float *) dst->data, IW, IH, OW, OH, KW, KH, IC, batch,
134+
batch_offset, delta_offset, s0, s1, p0, p1, d0, d1, stream);
120135
}
121136
}

0 commit comments

Comments
 (0)