1//
2// MIT license
3// Copyright (C) 2024 Intel Corporation
4// SPDX-License-Identifier: MIT
5//
6
7//
8// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
9// See https://llvm.org/LICENSE.txt for license information.
10// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
11//
12
13#include "im2col.hpp"
14
15#include <sycl/sycl.hpp>
16#include <type_traits> // For std::is_same_v
17
18#include "ggml.h"
19
20template <typename T>
21static 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) {
24 const int64_t work_group_size = item_ct1.get_local_range(2);
25 const int64_t global_id = item_ct1.get_local_id(2) + (work_group_size * item_ct1.get_group(2));
26
27 // make each work-item deal with more elements since sycl global range can not exceed max int
28 for (int64_t i = global_id; i < pelements; i += (work_group_size * item_ct1.get_group_range(2))) {
29 const int64_t ksize = OW * KH;
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;
54 }
55 }
56}
57
58template <typename T>
59static 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) {
62 const int64_t parallel_elements = OW * KW * KH;
63 const int64_t num_blocks = (parallel_elements + SYCL_IM2COL_BLOCK_SIZE - 1) / SYCL_IM2COL_BLOCK_SIZE;
64
65 // decrease global range when it exceeds the max int
66 int64_t local_size = downsample_sycl_global_range(batch * IC * OH * num_blocks, SYCL_IM2COL_BLOCK_SIZE);
67
68 sycl::range<3> block_nums(batch * IC, OH, num_blocks);
69 sycl::range<3> local_range(1, 1, local_size);
70
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
79static 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!");
85 }
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
90static 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);
95}
96
97void ggml_sycl_op_im2col(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
98 const ggml_tensor * src0 = dst->src[0];
99 const ggml_tensor * src1 = dst->src[1];
100
101 GGML_ASSERT(src1->type == GGML_TYPE_F32);
102 GGML_ASSERT(dst->type == GGML_TYPE_F16 || dst->type == GGML_TYPE_F32);
103
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];
110
111 const bool is_2D = ((const int32_t *) (dst->op_params))[6] == 1;
112
113 const int64_t IC = src1->ne[is_2D ? 2 : 1];
114 const int64_t IH = is_2D ? src1->ne[1] : 1;
115 const int64_t IW = src1->ne[0];
116
117 const int64_t KH = is_2D ? src0->ne[1] : 1;
118 const int64_t KW = src0->ne[0];
119
120 const int64_t OH = is_2D ? dst->ne[2] : 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);
126
127 queue_ptr stream = ctx.stream();
128
129 if (dst->type == GGML_TYPE_F16) {
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);
132 } else {
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);
135 }
136}