Skip to content

Commit 34dd506

Browse files
qnixsynapseqnixsynapse
committed
SYCL: Fix im2col
1 parent c94085d commit 34dd506

File tree

2 files changed

+92
-80
lines changed

2 files changed

+92
-80
lines changed

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

+1-2
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

+91-78
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,6 @@
1-
//
21
// MIT license
32
// Copyright (C) 2024 Intel Corporation
43
// SPDX-License-Identifier: MIT
5-
//
6-
74
//
85
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
96
// See https://llvm.org/LICENSE.txt for license information.
@@ -12,110 +9,126 @@
129

1310
#include "im2col.hpp"
1411

15-
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-
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);
12+
#include <sycl/sycl.hpp>
13+
#include <type_traits> // For std::is_same_v
2314

24-
// 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)) {
15+
#include "ggml.h"
2616

17+
template <typename T>
18+
static void im2col_kernel(const float * x, T * dst, int64_t batch_offset, int64_t offset_delta, int64_t IC, int64_t IW,
19+
int64_t IH, int64_t OH, int64_t OW, int64_t KW, int64_t KH, int64_t pelements, int64_t CHW,
20+
int s0, int s1, int p0, int p1, int d0, int d1, const sycl::nd_item<3> & item) {
21+
const int64_t work_group_size_x = item.get_local_range(2);
22+
const int64_t total_threads_x = work_group_size_x * item.get_group_range(2);
23+
const int64_t global_id_x = item.get_global_id(2);
24+
25+
for (int64_t i = global_id_x; i < pelements; i += total_threads_x) {
2726
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];
27+
const int64_t kx = i / ksize;
28+
const int64_t kd = kx * ksize;
29+
const int64_t ky = (i - kd) / OW;
30+
const int64_t ix = i % OW;
31+
32+
const int64_t oh = item.get_group(1);
33+
const int64_t group_z = item.get_group(0);
34+
const int64_t batch = group_z / IC;
35+
const int64_t ic = group_z % 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 = (((batch * OH + oh) * OW + ix) * CHW) + (ic * (KW * KH) + ky * KW + kx);
41+
42+
const int64_t offset_src_base = (ic * offset_delta) + (batch * batch_offset);
43+
const int64_t offset_src = offset_src_base + (iih * IW) + iiw;
44+
45+
const bool out_of_bounds = (iih < 0 || iih >= IH || iiw < 0 || iiw >= IW);
46+
const float src_val = out_of_bounds ? 0.0f : x[offset_src];
47+
48+
if constexpr (std::is_same_v<T, sycl::half>) {
49+
dst[offset_dst] = sycl::half(src_val);
50+
} else if constexpr (std::is_same_v<T, float>) {
51+
dst[offset_dst] = src_val;
5352
}
5453
}
5554
}
5655

5756
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) {
57+
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,
58+
int64_t KH, int64_t IC, int64_t batch, int64_t batch_offset, int64_t offset_delta,
59+
int s0, int s1, int p0, int p1, int d0, int d1, queue_ptr stream) {
6360
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;
65-
66-
// decrease global range when it exceeds the max int
67-
int64_t local_size = downsample_sycl_global_range(batch * IC * OH * num_blocks, SYCL_IM2COL_BLOCK_SIZE);
68-
sycl::range<3> block_nums(batch * IC, OH, num_blocks);
69-
sycl::range<3> local_range(1, 1, local_size);
70-
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-
});
61+
const int64_t block_size_x = SYCL_IM2COL_BLOCK_SIZE;
62+
const int64_t num_groups_x = (parallel_elements + block_size_x - 1) / block_size_x;
63+
64+
sycl::range<3> global_range_groups(batch * IC, OH, num_groups_x);
65+
sycl::range<3> local_range(1, 1, block_size_x);
66+
sycl::range<3> global_range_items = global_range_groups * local_range;
67+
68+
const int64_t CHW = IC * KH * KW;
69+
70+
stream->parallel_for(sycl::nd_range<3>(global_range_items, local_range), [=](sycl::nd_item<3> item) {
71+
im2col_kernel<T>(x, dst, batch_offset, offset_delta, IC, IW, IH, OH, OW, KW, KH, parallel_elements, CHW, s0, s1,
72+
p0, p1, d0, d1, item);
73+
});
74+
}
75+
76+
static void im2col_sycl_f16(const float * x, sycl::half * dst, int64_t IW, int64_t IH, int64_t OW, int64_t OH,
77+
int64_t KW, int64_t KH, int64_t IC, int64_t batch, int64_t batch_offset,
78+
int64_t offset_delta, int s0, int s1, int p0, int p1, int d0, int d1, queue_ptr stream) {
79+
if (!stream->get_device().has(sycl::aspect::fp16)) {
80+
throw sycl::exception(sycl::make_error_code(sycl::errc::kernel_not_supported),
81+
"Device does not support half precision (fp16) operations!");
8282
}
83+
im2col_sycl_internal<sycl::half>(x, dst, IW, IH, OW, OH, KW, KH, IC, batch, batch_offset, offset_delta, s0, s1, p0,
84+
p1, d0, d1, stream);
85+
}
86+
87+
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,
88+
int64_t KH, int64_t IC, int64_t batch, int64_t batch_offset, int64_t offset_delta, int s0,
89+
int s1, int p0, int p1, int d0, int d1, queue_ptr stream) {
90+
im2col_sycl_internal<float>(x, dst, IW, IH, OW, OH, KW, KH, IC, batch, batch_offset, offset_delta, s0, s1, p0, p1,
91+
d0, d1, stream);
8392
}
8493

85-
void ggml_sycl_op_im2col(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
94+
void ggml_sycl_op_im2col(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
8695
const ggml_tensor * src0 = dst->src[0];
8796
const ggml_tensor * src1 = dst->src[1];
8897

89-
GGML_ASSERT(src0->type == GGML_TYPE_F16);
9098
GGML_ASSERT(src1->type == GGML_TYPE_F32);
9199
GGML_ASSERT(dst->type == GGML_TYPE_F16 || dst->type == GGML_TYPE_F32);
92100

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];
101+
const int32_t s0 = ((const int32_t *) (dst->op_params))[0];
102+
const int32_t s1 = ((const int32_t *) (dst->op_params))[1];
103+
const int32_t p0 = ((const int32_t *) (dst->op_params))[2];
104+
const int32_t p1 = ((const int32_t *) (dst->op_params))[3];
105+
const int32_t d0 = ((const int32_t *) (dst->op_params))[4];
106+
const int32_t d1 = ((const int32_t *) (dst->op_params))[5];
99107

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

102110
const int64_t IC = src1->ne[is_2D ? 2 : 1];
103111
const int64_t IH = is_2D ? src1->ne[1] : 1;
104-
const int64_t IW = src1->ne[0];
112+
const int64_t IW = src1->ne[0];
105113

106114
const int64_t KH = is_2D ? src0->ne[1] : 1;
107-
const int64_t KW = src0->ne[0];
115+
const int64_t KW = src0->ne[0];
108116

109117
const int64_t OH = is_2D ? dst->ne[2] : 1;
110-
const int64_t OW = dst->ne[1];
118+
const int64_t OW = dst->ne[1];
111119

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
120+
const size_t delta_offset = src1->nb[is_2D ? 2 : 1] / sizeof(float);
121+
const int64_t batch = src1->ne[is_2D ? 3 : 2];
122+
const size_t batch_offset = src1->nb[is_2D ? 3 : 2] / sizeof(float);
123+
124+
queue_ptr stream = ctx.stream();
115125

116126
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());
127+
im2col_sycl_f16((const float *) src1->data, (sycl::half *) dst->data, IW, IH, OW, OH, KW, KH, IC, batch,
128+
batch_offset, delta_offset, s0, s1, p0, p1, d0, d1, stream);
118129
} 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());
130+
im2col_sycl_f32((const float *) src1->data, (float *) dst->data, IW, IH, OW, OH, KW, KH, IC, batch,
131+
batch_offset, delta_offset, s0, s1, p0, p1, d0, d1, stream);
120132
}
121133
}
134+

0 commit comments

Comments
 (0)