Gitiles
Code Review Sign In
review.mlplatform.org / ml / ComputeLibrary / da37e2fa661f85090bd4a3b93c07178df268bd8c / . / src / core / CL / cl_kernels / gemm.cl
blob: db15720ad0d06e8c3eab59b6625df4115cfb96a3 [file] [log] [blame]
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1/*
2 * Copyright (c) 2017 ARM Limited.
3 *
4 * SPDX-License-Identifier: MIT
5 *
6 * Permission is hereby granted, free of charge, to any person obtaining a copy
7 * of this software and associated documentation files (the "Software"), to
8 * deal in the Software without restriction, including without limitation the
9 * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
10 * sell copies of the Software, and to permit persons to whom the Software is
11 * furnished to do so, subject to the following conditions:
12 *
13 * The above copyright notice and this permission notice shall be included in all
14 * copies or substantial portions of the Software.
15 *
16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
19 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
21 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
22 * SOFTWARE.
23 */
24#include "helpers.h"
25
Gian Marco Iodice368da832017-07-03 12:33:49 +0100[diff] [blame]26#ifdef FIXED_POINT_POSITION
27#include "fixed_point.h"
28#endif // FIXED_POINT_POSITION
29
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]30/** This OpenCL kernel computes the "vector" 1x4 transposition of input matrix
31 *
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]32 * @param[in] src_ptr Pointer to the source matrix. Supported data types: U32/S32/F32
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]33 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes)
34 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
35 * @param[in] src_stride_y Stride of the source matrix in Y dimension (in bytes)
36 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
37 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]38 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]39 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
40 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
41 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
42 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
43 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
44 */
Gian Marco Iodice9f89bae2017-06-22 12:09:49 +0100[diff] [blame]45__kernel void gemm_transpose1x4(IMAGE_DECLARATION(src),
46 IMAGE_DECLARATION(dst))
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]47{
48 uint x = get_global_id(0);
49 uint y = get_global_id(1);
50
51 /* Compute address for Matrix B - source */
52 Image src = CONVERT_TO_IMAGE_STRUCT(src);
53
54 /* Compute address for Matrix B transposed - destination. X and Y are swapped */
55 uint dst_addr_in_bytes = y * 16 + ((x * dst_stride_y + dst_offset_first_element_in_bytes));
56
Gian Marco Iodice9f89bae2017-06-22 12:09:49 +0100[diff] [blame]57 uint4 b0 = vload4(0, (__global uint *)src.ptr);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]58
Gian Marco Iodice9f89bae2017-06-22 12:09:49 +0100[diff] [blame]59 vstore4(b0, 0, (__global uint *)(dst_ptr + dst_addr_in_bytes));
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]60}
61
62/** This OpenCL kernel computes the "vector" 1x8 transposition of input matrix
63 *
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]64 * @param[in] src_ptr Pointer to the source matrix. Supported data types: U16/S16/QS16/F16
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]65 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes)
66 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
67 * @param[in] src_stride_y Stride of the source matrix in Y dimension (in bytes)
68 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
69 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]70 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]71 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
72 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
73 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
74 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
75 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
76 */
Gian Marco Iodice9f89bae2017-06-22 12:09:49 +0100[diff] [blame]77__kernel void gemm_transpose1x8(IMAGE_DECLARATION(src),
78 IMAGE_DECLARATION(dst))
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]79{
80 uint x = get_global_id(0);
81 uint y = get_global_id(1);
82
83 /* Compute address for Matrix B - source */
84 Image src = CONVERT_TO_IMAGE_STRUCT(src);
85
86 /* Compute address for Matrix B transposed - destination. X and Y are swapped */
87 uint dst_addr_in_bytes = y * 16 + ((x * dst_stride_y + dst_offset_first_element_in_bytes));
88
Gian Marco Iodice9f89bae2017-06-22 12:09:49 +0100[diff] [blame]89 ushort8 b0 = vload8(0, (__global ushort *)src.ptr);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]90
Gian Marco Iodice9f89bae2017-06-22 12:09:49 +0100[diff] [blame]91 vstore8(b0, 0, (__global ushort *)(dst_ptr + dst_addr_in_bytes));
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]92}
93
94/** This OpenCL kernel computes the "vector" 1x16 transposition of input matrix
95 *
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]96 * @param[in] src_ptr Pointer to the source matrix. Supported data types: U8/S8/QS8
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]97 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes)
98 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
99 * @param[in] src_stride_y Stride of the source matrix in Y dimension (in bytes)
100 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
101 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]102 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]103 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
104 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
105 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
106 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
107 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
108 */
Gian Marco Iodice9f89bae2017-06-22 12:09:49 +0100[diff] [blame]109__kernel void gemm_transpose1x16(IMAGE_DECLARATION(src),
110 IMAGE_DECLARATION(dst))
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]111{
112 uint x = get_global_id(0);
113 uint y = get_global_id(1);
114
115 /* Compute address for Matrix B - source */
116 Image src = CONVERT_TO_IMAGE_STRUCT(src);
117
118 /* Compute address for Matrix B transposed - destination. X and Y are swapped */
119 uint dst_addr_in_bytes = y * 16 + ((x * dst_stride_y + dst_offset_first_element_in_bytes));
120
121 uchar16 b0 = vload16(0, (__global uchar *)src.ptr);
122
123 vstore16(b0, 0, (__global uchar *)(dst_ptr + dst_addr_in_bytes));
124}
125
126/** This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values
127 *
128 * @param[in] src_ptr Pointer to the source matrix. Supported data types: U32/S32/F32
129 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes)
130 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
131 * @param[in] src_stride_y Stride of the source matrix in Y dimension (in bytes)
132 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
133 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]134 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]135 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
136 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
137 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
138 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
139 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
140 */
141__kernel void gemm_interleave4x4_32bit(IMAGE_DECLARATION(src),
142 IMAGE_DECLARATION(dst))
143{
144 /* Compute source and destination addresses */
145 Image src = CONVERT_TO_IMAGE_STRUCT(src);
146 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
147
148 /* Load values from Matrix A */
Gian Marco Iodiceb93f5de2017-07-05 15:48:39 +0100[diff] [blame]149 uint4 a0 = vload4(0, (__global uint *)(offset(&src, 0, 0)));
150 uint4 a1 = vload4(0, (__global uint *)(offset(&src, 0, 1)));
151 uint4 a2 = vload4(0, (__global uint *)(offset(&src, 0, 2)));
152 uint4 a3 = vload4(0, (__global uint *)(offset(&src, 0, 3)));
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]153
Gian Marco Iodiceb93f5de2017-07-05 15:48:39 +0100[diff] [blame]154 uint4 val0 = (uint4)(a0.s0, a1.s0, a2.s0, a3.s0);
155 vstore4(val0, 0, ((__global uint *)dst.ptr) + 0);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]156
Gian Marco Iodiceb93f5de2017-07-05 15:48:39 +0100[diff] [blame]157 val0 = (uint4)(a0.s1, a1.s1, a2.s1, a3.s1);
158 vstore4(val0, 0, ((__global uint *)dst.ptr) + 4);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]159
Gian Marco Iodiceb93f5de2017-07-05 15:48:39 +0100[diff] [blame]160 val0 = (uint4)(a0.s2, a1.s2, a2.s2, a3.s2);
161 vstore4(val0, 0, ((__global uint *)dst.ptr) + 8);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]162
Gian Marco Iodiceb93f5de2017-07-05 15:48:39 +0100[diff] [blame]163 val0 = (uint4)(a0.s3, a1.s3, a2.s3, a3.s3);
164 vstore4(val0, 0, ((__global uint *)dst.ptr) + 12);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]165}
166
167/** This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values
168 *
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]169 * @param[in] src_ptr Pointer to the source matrix. Supported data types: U16/S16/QS16/F16
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]170 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes)
171 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
172 * @param[in] src_stride_y Stride of the source matrix in Y dimension (in bytes)
173 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
174 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]175 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]176 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
177 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
178 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
179 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
180 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
181 */
182__kernel void gemm_interleave4x4_16bit(IMAGE_DECLARATION(src),
183 IMAGE_DECLARATION(dst))
184{
185 /* Compute source and destination addresses */
186 Image src = CONVERT_TO_IMAGE_STRUCT(src);
187 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
188
189 /* Load values from Matrix A */
Gian Marco Iodiceb93f5de2017-07-05 15:48:39 +0100[diff] [blame]190 ushort8 a0 = vload8(0, (__global ushort *)(offset(&src, 0, 0)));
191 ushort8 a1 = vload8(0, (__global ushort *)(offset(&src, 0, 1)));
192 ushort8 a2 = vload8(0, (__global ushort *)(offset(&src, 0, 2)));
193 ushort8 a3 = vload8(0, (__global ushort *)(offset(&src, 0, 3)));
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]194
Gian Marco Iodiceb93f5de2017-07-05 15:48:39 +0100[diff] [blame]195 ushort8 val0 = (ushort8)((ushort4)(a0.s0, a1.s0, a2.s0, a3.s0), (ushort4)(a0.s1, a1.s1, a2.s1, a3.s1));
196 vstore8(val0, 0, ((__global ushort *)dst.ptr) + 0);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]197
Gian Marco Iodiceb93f5de2017-07-05 15:48:39 +0100[diff] [blame]198 val0 = (ushort8)((ushort4)(a0.s2, a1.s2, a2.s2, a3.s2), (ushort4)(a0.s3, a1.s3, a2.s3, a3.s3));
199 vstore8(val0, 0, ((__global ushort *)dst.ptr) + 8);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]200
Gian Marco Iodiceb93f5de2017-07-05 15:48:39 +0100[diff] [blame]201 val0 = (ushort8)((ushort4)(a0.s4, a1.s4, a2.s4, a3.s4), (ushort4)(a0.s5, a1.s5, a2.s5, a3.s5));
202 vstore8(val0, 0, ((__global ushort *)dst.ptr) + 16);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]203
Gian Marco Iodiceb93f5de2017-07-05 15:48:39 +0100[diff] [blame]204 val0 = (ushort8)((ushort4)(a0.s6, a1.s6, a2.s6, a3.s6), (ushort4)(a0.s7, a1.s7, a2.s7, a3.s7));
205 vstore8(val0, 0, ((__global ushort *)dst.ptr) + 24);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]206}
207
208/** This OpenCL kernel reshapes the input matrix transposing each 4x4 block and interleaving the values
209 *
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]210 * @param[in] src_ptr Pointer to the source matrix. Supported data types: U8/S8/QS8
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]211 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes)
212 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
213 * @param[in] src_stride_y Stride of the source matrix in Y dimension (in bytes)
214 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
215 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]216 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]217 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
218 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
219 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
220 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
221 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
222 */
223__kernel void gemm_interleave4x4_8bit(IMAGE_DECLARATION(src),
224 IMAGE_DECLARATION(dst))
225{
226 /* Compute source and destination addresses */
227 Image src = CONVERT_TO_IMAGE_STRUCT(src);
228 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
229
230 /* Load values from Matrix A */
231 uchar16 a0 = vload16(0, (__global uchar *)(offset(&src, 0, 0)));
232 uchar16 a1 = vload16(0, (__global uchar *)(offset(&src, 0, 1)));
233 uchar16 a2 = vload16(0, (__global uchar *)(offset(&src, 0, 2)));
234 uchar16 a3 = vload16(0, (__global uchar *)(offset(&src, 0, 3)));
235
236 uchar16 val0 = (uchar16)((uchar4)(a0.s0, a1.s0, a2.s0, a3.s0), (uchar4)(a0.s1, a1.s1, a2.s1, a3.s1),
237 (uchar4)(a0.s2, a1.s2, a2.s2, a3.s2), (uchar4)(a0.s3, a1.s3, a2.s3, a3.s3));
238 vstore16(val0, 0, ((__global uchar *)dst.ptr) + 0);
239
240 val0 = (uchar16)((uchar4)(a0.s4, a1.s4, a2.s4, a3.s4), (uchar4)(a0.s5, a1.s5, a2.s5, a3.s5),
241 (uchar4)(a0.s6, a1.s6, a2.s6, a3.s6), (uchar4)(a0.s7, a1.s7, a2.s7, a3.s7));
242 vstore16(val0, 0, ((__global uchar *)dst.ptr) + 16);
243
244 val0 = (uchar16)((uchar4)(a0.s8, a1.s8, a2.s8, a3.s8), (uchar4)(a0.s9, a1.s9, a2.s9, a3.s9),
245 (uchar4)(a0.sA, a1.sA, a2.sA, a3.sA), (uchar4)(a0.sB, a1.sB, a2.sB, a3.sB));
246 vstore16(val0, 0, ((__global uchar *)dst.ptr) + 32);
247
248 val0 = (uchar16)((uchar4)(a0.sC, a1.sC, a2.sC, a3.sC), (uchar4)(a0.sD, a1.sD, a2.sD, a3.sD),
249 (uchar4)(a0.sE, a1.sE, a2.sE, a3.sE), (uchar4)(a0.sF, a1.sF, a2.sF, a3.sF));
250 vstore16(val0, 0, ((__global uchar *)dst.ptr) + 48);
251}
252
253/** This kernel accumulates each row with the biases vector
254 *
Gian Marco Iodice578ab612017-06-23 09:34:33 +0100[diff] [blame]255 * @note The data type must be passed at compile time -DDATA_TYPE=type. e.g. -DDATA_TYPE=short
256 *
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]257 * @param[in, out] accum_ptr Pointer to the accumulate tensor. Supported data type: U8/S8/QS8/U16/S16/F16/U32/S32/F32
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]258 * @param[in] accum_stride_x Stride of the accmulate tensor in X dimension (in bytes)
259 * @param[in] accum_step_x accum_stride_x * number of elements along X processed per workitem(in bytes)
260 * @param[in] accum_stride_y Stride of the accumlulate tensor in Y dimension (in bytes)
261 * @param[in] accum_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
262 * @param[in] accum_offset_first_element_in_bytes The offset of the first element in the accumulate tensor
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]263 * @param[in] biases_ptr Pointer to the biases vector. Same as @p accum_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]264 * @param[in] biases_stride_x Stride of the destination tensor in X dimension (in bytes)
265 * @param[in] biases_step_x dst_stride_x * number of elements along X processed per workitem(in bytes)
266 * @param[in] biases_offset_first_element_in_bytes The offset of the first element in the destination tensor
267 */
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]268#ifdef DATA_TYPE
Gian Marco Iodice578ab612017-06-23 09:34:33 +0100[diff] [blame]269__kernel void gemm_accumulate_biases(
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]270 IMAGE_DECLARATION(accum),
271 VECTOR_DECLARATION(biases))
272{
273 Image accum = CONVERT_TO_IMAGE_STRUCT(accum);
274 Vector biases = CONVERT_TO_VECTOR_STRUCT(biases);
275
Gian Marco Iodice578ab612017-06-23 09:34:33 +0100[diff] [blame]276 VEC_DATA_TYPE(DATA_TYPE, 16)
277 accum_value = vload16(0, (__global DATA_TYPE *)accum.ptr);
278 VEC_DATA_TYPE(DATA_TYPE, 16)
279 biases_value = vload16(0, (__global DATA_TYPE *)biases.ptr);
Gian Marco Iodice368da832017-07-03 12:33:49 +0100[diff] [blame]280#ifdef FIXED_POINT_POSITION
281 accum_value = ADD_SAT_OP_EXPAND(biases_value, accum_value, DATA_TYPE, 16);
282#else // FIXED_POINT_POSITION
283 accum_value = biases_value + accum_value;
284#endif // FIXED_POINT_POSITION
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]285
286 // Store result in the accummulate buffer
Gian Marco Iodice578ab612017-06-23 09:34:33 +0100[diff] [blame]287 vstore16(accum_value, 0, (__global DATA_TYPE *)accum.ptr);
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]288}
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]289#endif /* DATA_TYPE */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]290
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]291#ifdef WIDTH_MATRIX_B
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]292/** This OpenCL kernel computes the matrix multiplication between matrix A (src0) and matrix B (src1)
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]293 * Matrix A and matrix B must be reshaped respectively with @ref gemm_interleave4x4_8bit and @ref gemm_transpose1x16 before running the matrix multiplication
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]294 *
295 * @attention The width of matrix B and the alpha's value need to be passed at compile time using -DWIDTH_MATRIX_B
296 *
297 * @param[in] src0_ptr Pointer to the source matrix. Supported formats: U8
298 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes)
299 * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
300 * @param[in] src0_stride_y Stride of the source matrix in Y dimension (in bytes)
301 * @param[in] src0_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
302 * @param[in] src0_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]303 * @param[in] src1_ptr Pointer to the source matrix. Supported formats: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]304 * @param[in] src1_stride_x Stride of the source matrix in X dimension (in bytes)
305 * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
306 * @param[in] src1_stride_y Stride of the source matrix in Y dimension (in bytes)
307 * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
308 * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]309 * @param[out] dst_ptr Pointer to the destination matrix Supported formats: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]310 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
311 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
312 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
313 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
314 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
315 * @param[in] a_offset Offset to be added to each element of the matrix A
316 * @param[in] b_offset Offset to be added to each element of the matrix B.
317 * @param[in] c_offset Offset to be added to each element of the matrix C.
318 * @param[in] c_mult_int Multiplied with each element of the matrix C.
319 * @param[in] shift Number of bits to shift right the result.
320 */
321__kernel void gemm_mm_u8(IMAGE_DECLARATION(src0),
322 IMAGE_DECLARATION(src1),
323 IMAGE_DECLARATION(dst),
324 int a_offset,
325 int b_offset,
326 int c_offset,
327 int c_mult_int,
328 int shift)
329{
330 /* src_addr.s0 = address of matrix A */
331 /* src_addr.s1 = address of matrix B */
332
333 /* Compute address for matrix A and B */
334 int2 src_addr = (int2)(get_global_id(1), get_global_id(0)) * (int2)((src0_stride_y),
335 (src1_stride_y));
336
337 /* Add offset_first_element_in_bytes */
338 src_addr = src_addr + ((int2)(src0_offset_first_element_in_bytes, src1_offset_first_element_in_bytes));
339
340 /* Compute end row address for matrix B */
341 int end_row_mtx_b = src_addr.s1 + WIDTH_MATRIX_B;
342
343 /* Reset accumulators */
344 int16 c00 = 0.0f;
345 int16 c10 = 0.0f;
346 int16 c20 = 0.0f;
347 int16 c30 = 0.0f;
348
349 for(; src_addr.s1 <= (end_row_mtx_b - 8); src_addr += (int2)(8, 32))
350 {
351 /* Load values from matrix A (interleaved) and matrix B (transposed) */
352 int8 a0 = (int8)a_offset + convert_int8(vload8(0, ((__global uchar *)src0_ptr) + src_addr.s0));
353 int16 b0 = (int16)b_offset + convert_int16(vload16(0, ((__global uchar *)src1_ptr) + src_addr.s1));
354
355 c00 += (int16)a0.s0 * b0;
356 c10 += (int16)a0.s1 * b0;
357 c20 += (int16)a0.s2 * b0;
358 c30 += (int16)a0.s3 * b0;
359
360 int16 b1 = (int16)b_offset + convert_int16(vload16(0, ((__global uchar *)src1_ptr) + src_addr.s1 + 16));
361
362 c00 += (int16)a0.s4 * b1;
363 c10 += (int16)a0.s5 * b1;
364 c20 += (int16)a0.s6 * b1;
365 c30 += (int16)a0.s7 * b1;
366 }
367
368 for(; src_addr.s1 < end_row_mtx_b; src_addr += (int2)(4, 16))
369 {
370 /* Load values from matrix A (interleaved) and matrix B (transposed) */
371 int4 a0 = (int4)a_offset + convert_int4(vload4(0, ((__global uchar *)src0_ptr) + src_addr.s0));
372 int16 b0 = (int16)b_offset + convert_int16(vload16(0, ((__global uchar *)src1_ptr) + src_addr.s1));
373
374 c00 += (int16)a0.s0 * b0;
375 c10 += (int16)a0.s1 * b0;
376 c20 += (int16)a0.s2 * b0;
377 c30 += (int16)a0.s3 * b0;
378 }
379
380 /* Compute destination address */
381 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
382
383 /* Multiply by the weight of matrix product */
384 c00 = (((int16)c_offset + c00) * (int16)c_mult_int) >> shift;
385 c10 = (((int16)c_offset + c10) * (int16)c_mult_int) >> shift;
386 c20 = (((int16)c_offset + c20) * (int16)c_mult_int) >> shift;
387 c30 = (((int16)c_offset + c30) * (int16)c_mult_int) >> shift;
388
389 /* Store 4x16 block */
390 vstore16(convert_uchar16_sat(c00), 0, (__global uchar *)(offset(&dst, 0, 0)));
391 vstore16(convert_uchar16_sat(c10), 0, (__global uchar *)(offset(&dst, 0, 1)));
392 vstore16(convert_uchar16_sat(c20), 0, (__global uchar *)(offset(&dst, 0, 2)));
393 vstore16(convert_uchar16_sat(c30), 0, (__global uchar *)(offset(&dst, 0, 3)));
394}
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]395#endif /* WIDTH_MATRIX_B */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]396
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]397#if defined(WIDTH_MATRIX_B) && defined(ALPHA)
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]398/** This OpenCL kernel is optimised for Midgard. It computes the matrix multiplication between matrix A (src0) and matrix B (src1)
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]399 * Matrix A and matrix B must be reshaped respectively with @ref gemm_interleave4x4_32bit and @ref gemm_transpose1x4 before running the matrix multiplication
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]400 *
401 * @attention The width of matrix B and the alpha's value need to be passed at compile time using -DWIDTH_MATRIX_B and -DALPHA
402 *
403 * @param[in] src0_ptr Pointer to the source matrix. Supported data types: F32
404 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes)
405 * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
406 * @param[in] src0_stride_y Stride of the source matrix in Y dimension (in bytes)
407 * @param[in] src0_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
408 * @param[in] src0_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]409 * @param[in] src1_ptr Pointer to the source matrix. Supported data types: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]410 * @param[in] src1_stride_x Stride of the source matrix in X dimension (in bytes)
411 * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
412 * @param[in] src1_stride_y Stride of the source matrix in Y dimension (in bytes)
413 * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
414 * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]415 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]416 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
417 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
418 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
419 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
420 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
421 */
422__kernel void gemm_mm_f32_midgard(IMAGE_DECLARATION(src0),
423 IMAGE_DECLARATION(src1),
424 IMAGE_DECLARATION(dst))
425{
426 /* src_addr.s0 = address of matrix A */
427 /* src_addr.s1 = address of matrix B */
428
429 /* Compute address for matrix A and B */
430 int2 src_addr = (int2)(get_global_id(1), get_global_id(0)) * (int2)((src0_stride_y),
431 (src1_stride_y));
432
433 /* Add offset_first_element_in_bytes */
434 src_addr = src_addr + ((int2)(src0_offset_first_element_in_bytes, src1_offset_first_element_in_bytes));
435
436 /* Divide by 4 in order to get the src_addr in unit of float */
437 src_addr = src_addr >> 2;
438
439 /* Compute end row address for matrix B */
440 int end_row_mtx_b = src_addr.s1 + WIDTH_MATRIX_B;
441
442 /* Reset accumulators */
443 float4 c00 = 0.0f;
444 float4 c10 = 0.0f;
445 float4 c20 = 0.0f;
446 float4 c30 = 0.0f;
447
448 for(; src_addr.s1 <= (end_row_mtx_b - 8); src_addr += (int2)(8, 8))
449 {
450 /* Load values from matrix A (interleaved) and matrix B (transposed) */
451 float4 a0 = vload4(0, ((__global float *)src0_ptr) + src_addr.s0);
452 float4 b0 = vload4(0, ((__global float *)src1_ptr) + src_addr.s1);
453
454 c00 += (float4)a0.s0 * b0;
455 c10 += (float4)a0.s1 * b0;
456 c20 += (float4)a0.s2 * b0;
457 c30 += (float4)a0.s3 * b0;
458
459 /* Load values from matrix A (interleaved) and matrix B (transposed) */
460 a0 = vload4(0, ((__global float *)src0_ptr) + src_addr.s0 + 4);
461 b0 = vload4(0, ((__global float *)src1_ptr) + src_addr.s1 + 4);
462
463 c00 += (float4)a0.s0 * b0;
464 c10 += (float4)a0.s1 * b0;
465 c20 += (float4)a0.s2 * b0;
466 c30 += (float4)a0.s3 * b0;
467 }
468
469 for(; src_addr.s1 < end_row_mtx_b; src_addr += (int2)(4, 4))
470 {
471 /* Load values from matrix A (interleaved) and matrix B (transposed) */
472 float4 a0 = vload4(0, ((__global float *)src0_ptr) + src_addr.s0);
473 float4 b0 = vload4(0, ((__global float *)src1_ptr) + src_addr.s1);
474
475 c00 += (float4)a0.s0 * b0;
476 c10 += (float4)a0.s1 * b0;
477 c20 += (float4)a0.s2 * b0;
478 c30 += (float4)a0.s3 * b0;
479 }
480
481 /* Compute destination address */
482 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
483
484 /* Multiply by the weight of matrix product */
485 c00 = c00 * (float4)ALPHA;
486 c10 = c10 * (float4)ALPHA;
487 c20 = c20 * (float4)ALPHA;
488 c30 = c30 * (float4)ALPHA;
489
490 /* Store 4x4 block */
491 vstore4(c00, 0, (__global float *)(offset(&dst, 0, 0)));
492 vstore4(c10, 0, (__global float *)(offset(&dst, 0, 1)));
493 vstore4(c20, 0, (__global float *)(offset(&dst, 0, 2)));
494 vstore4(c30, 0, (__global float *)(offset(&dst, 0, 3)));
495}
496
497/** This OpenCL kernel is optimised for Bifrost. It computes the matrix multiplication between matrix A (src0) and matrix B (src1)
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]498 * Matrix A and matrix B must be reshaped respectively with @ref gemm_interleave4x4_32bit and @ref gemm_transpose1x4 before running the matrix multiplication
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]499 *
500 * @attention The width of matrix B and the alpha's value need to be passed at compile time using -DWIDTH_MATRIX_B and -DALPHA
501 *
502 * @param[in] src0_ptr Pointer to the source matrix. Supported data types: F32
503 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes)
504 * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
505 * @param[in] src0_stride_y Stride of the source matrix in Y dimension (in bytes)
506 * @param[in] src0_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
507 * @param[in] src0_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]508 * @param[in] src1_ptr Pointer to the source matrix. Supported data types: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]509 * @param[in] src1_stride_x Stride of the source matrix in X dimension (in bytes)
510 * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
511 * @param[in] src1_stride_y Stride of the source matrix in Y dimension (in bytes)
512 * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
513 * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]514 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]515 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
516 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
517 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
518 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
519 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
520 */
521__kernel void gemm_mm_f32_bifrost(IMAGE_DECLARATION(src0),
522 IMAGE_DECLARATION(src1),
523 IMAGE_DECLARATION(dst))
524{
525 // src_addr_a = address of matrix A
526 // src_addr_b = address of matrix B
527 __global float *src_addr_a = (__global float *)(src0_ptr + get_global_id(1) * src0_stride_y + src0_offset_first_element_in_bytes);
528 __global float *src_addr_b = (__global float *)(src1_ptr + get_global_id(0) * src1_stride_y + src1_offset_first_element_in_bytes);
529
530 // Compute end row address for matrix B
531 __global float *src_end_addr_b = src_addr_b + WIDTH_MATRIX_B;
532
533 // Reset accumulators
534 float c00 = 0.0f;
535 float c01 = 0.0f;
536 float c02 = 0.0f;
537 float c03 = 0.0f;
538 float c10 = 0.0f;
539 float c11 = 0.0f;
540 float c12 = 0.0f;
541 float c13 = 0.0f;
542 float c20 = 0.0f;
543 float c21 = 0.0f;
544 float c22 = 0.0f;
545 float c23 = 0.0f;
546 float c30 = 0.0f;
547 float c31 = 0.0f;
548 float c32 = 0.0f;
549 float c33 = 0.0f;
550
551 for(; src_addr_b <= (src_end_addr_b - 16); src_addr_a += 16, src_addr_b += 16)
552 {
553 // Load values from matrix A (interleaved) and matrix B (transposed)
554 float4 a0 = vload4(0, src_addr_a);
555 float4 b0 = vload4(0, src_addr_b);
556
557 c00 = fma(a0.s0, b0.s0, c00);
558 c01 = fma(a0.s0, b0.s1, c01);
559 c02 = fma(a0.s0, b0.s2, c02);
560 c03 = fma(a0.s0, b0.s3, c03);
561
562 c10 = fma(a0.s1, b0.s0, c10);
563 c11 = fma(a0.s1, b0.s1, c11);
564 c12 = fma(a0.s1, b0.s2, c12);
565 c13 = fma(a0.s1, b0.s3, c13);
566
567 c20 = fma(a0.s2, b0.s0, c20);
568 c21 = fma(a0.s2, b0.s1, c21);
569 c22 = fma(a0.s2, b0.s2, c22);
570 c23 = fma(a0.s2, b0.s3, c23);
571
572 c30 = fma(a0.s3, b0.s0, c30);
573 c31 = fma(a0.s3, b0.s1, c31);
574 c32 = fma(a0.s3, b0.s2, c32);
575 c33 = fma(a0.s3, b0.s3, c33);
576
577 // Load values from matrix A (interleaved) and matrix B (transposed)
578 a0 = vload4(0, src_addr_a + 4);
579 b0 = vload4(0, src_addr_b + 4);
580
581 c00 = fma(a0.s0, b0.s0, c00);
582 c01 = fma(a0.s0, b0.s1, c01);
583 c02 = fma(a0.s0, b0.s2, c02);
584 c03 = fma(a0.s0, b0.s3, c03);
585
586 c10 = fma(a0.s1, b0.s0, c10);
587 c11 = fma(a0.s1, b0.s1, c11);
588 c12 = fma(a0.s1, b0.s2, c12);
589 c13 = fma(a0.s1, b0.s3, c13);
590
591 c20 = fma(a0.s2, b0.s0, c20);
592 c21 = fma(a0.s2, b0.s1, c21);
593 c22 = fma(a0.s2, b0.s2, c22);
594 c23 = fma(a0.s2, b0.s3, c23);
595
596 c30 = fma(a0.s3, b0.s0, c30);
597 c31 = fma(a0.s3, b0.s1, c31);
598 c32 = fma(a0.s3, b0.s2, c32);
599 c33 = fma(a0.s3, b0.s3, c33);
600
601 // Load values from matrix A (interleaved) and matrix B (transposed)
602 a0 = vload4(0, src_addr_a + 8);
603 b0 = vload4(0, src_addr_b + 8);
604
605 c00 = fma(a0.s0, b0.s0, c00);
606 c01 = fma(a0.s0, b0.s1, c01);
607 c02 = fma(a0.s0, b0.s2, c02);
608 c03 = fma(a0.s0, b0.s3, c03);
609
610 c10 = fma(a0.s1, b0.s0, c10);
611 c11 = fma(a0.s1, b0.s1, c11);
612 c12 = fma(a0.s1, b0.s2, c12);
613 c13 = fma(a0.s1, b0.s3, c13);
614
615 c20 = fma(a0.s2, b0.s0, c20);
616 c21 = fma(a0.s2, b0.s1, c21);
617 c22 = fma(a0.s2, b0.s2, c22);
618 c23 = fma(a0.s2, b0.s3, c23);
619
620 c30 = fma(a0.s3, b0.s0, c30);
621 c31 = fma(a0.s3, b0.s1, c31);
622 c32 = fma(a0.s3, b0.s2, c32);
623 c33 = fma(a0.s3, b0.s3, c33);
624
625 // Load values from matrix A (interleaved) and matrix B (transposed)
626 a0 = vload4(0, src_addr_a + 12);
627 b0 = vload4(0, src_addr_b + 12);
628
629 c00 = fma(a0.s0, b0.s0, c00);
630 c01 = fma(a0.s0, b0.s1, c01);
631 c02 = fma(a0.s0, b0.s2, c02);
632 c03 = fma(a0.s0, b0.s3, c03);
633
634 c10 = fma(a0.s1, b0.s0, c10);
635 c11 = fma(a0.s1, b0.s1, c11);
636 c12 = fma(a0.s1, b0.s2, c12);
637 c13 = fma(a0.s1, b0.s3, c13);
638
639 c20 = fma(a0.s2, b0.s0, c20);
640 c21 = fma(a0.s2, b0.s1, c21);
641 c22 = fma(a0.s2, b0.s2, c22);
642 c23 = fma(a0.s2, b0.s3, c23);
643
644 c30 = fma(a0.s3, b0.s0, c30);
645 c31 = fma(a0.s3, b0.s1, c31);
646 c32 = fma(a0.s3, b0.s2, c32);
647 c33 = fma(a0.s3, b0.s3, c33);
648 }
649
650 for(; src_addr_b < src_end_addr_b; src_addr_a += 4, src_addr_b += 4)
651 {
652 // Load values from matrix A (interleaved) and matrix B (transposed)
653 float4 a0 = vload4(0, src_addr_a);
654 float4 b0 = vload4(0, src_addr_b);
655
656 c00 = fma(a0.s0, b0.s0, c00);
657 c01 = fma(a0.s0, b0.s1, c01);
658 c02 = fma(a0.s0, b0.s2, c02);
659 c03 = fma(a0.s0, b0.s3, c03);
660
661 c10 = fma(a0.s1, b0.s0, c10);
662 c11 = fma(a0.s1, b0.s1, c11);
663 c12 = fma(a0.s1, b0.s2, c12);
664 c13 = fma(a0.s1, b0.s3, c13);
665
666 c20 = fma(a0.s2, b0.s0, c20);
667 c21 = fma(a0.s2, b0.s1, c21);
668 c22 = fma(a0.s2, b0.s2, c22);
669 c23 = fma(a0.s2, b0.s3, c23);
670
671 c30 = fma(a0.s3, b0.s0, c30);
672 c31 = fma(a0.s3, b0.s1, c31);
673 c32 = fma(a0.s3, b0.s2, c32);
674 c33 = fma(a0.s3, b0.s3, c33);
675 }
676
677 // Compute destination address
678 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
679
680 // Multiply by the weight of matrix product
681 c00 = c00 * ALPHA;
682 c01 = c01 * ALPHA;
683 c02 = c02 * ALPHA;
684 c03 = c03 * ALPHA;
685 c10 = c10 * ALPHA;
686 c11 = c11 * ALPHA;
687 c12 = c12 * ALPHA;
688 c13 = c13 * ALPHA;
689 c20 = c20 * ALPHA;
690 c21 = c21 * ALPHA;
691 c22 = c22 * ALPHA;
692 c23 = c23 * ALPHA;
693 c30 = c30 * ALPHA;
694 c31 = c31 * ALPHA;
695 c32 = c32 * ALPHA;
696 c33 = c33 * ALPHA;
697
698 barrier(CLK_GLOBAL_MEM_FENCE);
699
700 // Store 4x4 block
701 vstore4((float4)(c00, c01, c02, c03), 0, (__global float *)(offset(&dst, 0, 0)));
702 vstore4((float4)(c10, c11, c12, c13), 0, (__global float *)(offset(&dst, 0, 1)));
703 vstore4((float4)(c20, c21, c22, c23), 0, (__global float *)(offset(&dst, 0, 2)));
704 vstore4((float4)(c30, c31, c32, c33), 0, (__global float *)(offset(&dst, 0, 3)));
705}
706
707/** This OpenCL kernel computes the matrix multiplication between matrix A (src0) and matrix B (src1)
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]708 * Matrix A and matrix B must be reshaped respectively with @ref gemm_interleave4x4_16bit and @ref gemm_transpose1x8 before running the matrix multiplication
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]709 *
710 * @attention The width of matrix B and the alpha's value need to be passed at compile time using -DWIDTH_MATRIX_B and -DALPHA
711 *
712 * @param[in] src0_ptr Pointer to the source matrix. Supported data types: F16
713 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes)
714 * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
715 * @param[in] src0_stride_y Stride of the source matrix in Y dimension (in bytes)
716 * @param[in] src0_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
717 * @param[in] src0_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]718 * @param[in] src1_ptr Pointer to the source matrix. Supported data types: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]719 * @param[in] src1_stride_x Stride of the source matrix in X dimension (in bytes)
720 * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
721 * @param[in] src1_stride_y Stride of the source matrix in Y dimension (in bytes)
722 * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
723 * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]724 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]725 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
726 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
727 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
728 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
729 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
730 */
731__kernel void gemm_mm_f16(IMAGE_DECLARATION(src0),
732 IMAGE_DECLARATION(src1),
733 IMAGE_DECLARATION(dst))
734{
735 /* src_addr.s0 = address of matrix A */
736 /* src_addr.s1 = address of matrix B */
737
738 /* Compute address for matrix A and B */
739 int2 src_addr = (int2)(get_global_id(1), get_global_id(0)) * (int2)((src0_stride_y),
740 (src1_stride_y));
741
742 /* Add offset_first_element_in_bytes */
743 src_addr = src_addr + ((int2)(src0_offset_first_element_in_bytes, src1_offset_first_element_in_bytes));
744
745 /* Divide by 2 in order to get the src_addr in unit of half */
746 src_addr = src_addr >> 1;
747
748 /* Compute end row address for matrix B */
749 int end_row_mtx_b = src_addr.s1 + WIDTH_MATRIX_B;
750
751 /* Reset accumulators */
752 half8 c00 = 0.0f;
753 half8 c10 = 0.0f;
754 half8 c20 = 0.0f;
755 half8 c30 = 0.0f;
756
757 for(; src_addr.s1 <= (end_row_mtx_b - 8); src_addr += (int2)(8, 16))
758 {
759 /* Load values from matrix A (interleaved) and matrix B (transposed) */
760 half4 a0 = vload4(0, ((__global half *)src0_ptr) + src_addr.s0);
761 half8 b0 = vload8(0, ((__global half *)src1_ptr) + src_addr.s1);
762
763 c00 += (half8)a0.s0 * b0;
764 c10 += (half8)a0.s1 * b0;
765 c20 += (half8)a0.s2 * b0;
766 c30 += (half8)a0.s3 * b0;
767
768 /* Load values from matrix A (interleaved) and matrix B (transposed) */
769 a0 = vload4(0, ((__global half *)src0_ptr) + src_addr.s0 + 4);
770 b0 = vload8(0, ((__global half *)src1_ptr) + src_addr.s1 + 8);
771
772 c00 += (half8)a0.s0 * b0;
773 c10 += (half8)a0.s1 * b0;
774 c20 += (half8)a0.s2 * b0;
775 c30 += (half8)a0.s3 * b0;
776 }
777
778 for(; src_addr.s1 < end_row_mtx_b; src_addr += (int2)(4, 8))
779 {
780 /* Load values from matrix A (interleaved) and matrix B (transposed) */
781 half4 a0 = vload4(0, ((__global half *)src0_ptr) + src_addr.s0);
782 half8 b0 = vload8(0, ((__global half *)src1_ptr) + src_addr.s1);
783
784 c00 += (half8)a0.s0 * b0;
785 c10 += (half8)a0.s1 * b0;
786 c20 += (half8)a0.s2 * b0;
787 c30 += (half8)a0.s3 * b0;
788 }
789
790 /* Compute destination address */
791 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
792
793 /* Multiply by the weight of matrix product */
794 c00 = c00 * (half8)ALPHA;
795 c10 = c10 * (half8)ALPHA;
796 c20 = c20 * (half8)ALPHA;
797 c30 = c30 * (half8)ALPHA;
798
799 /* Store 4x8 block */
800 vstore8(c00, 0, (__global half *)(offset(&dst, 0, 0)));
801 vstore8(c10, 0, (__global half *)(offset(&dst, 0, 1)));
802 vstore8(c20, 0, (__global half *)(offset(&dst, 0, 2)));
803 vstore8(c30, 0, (__global half *)(offset(&dst, 0, 3)));
804}
805
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]806#ifdef FIXED_POINT_POSITION
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]807/** This OpenCL kernel computes the matrix multiplication between matrix A (src0) and matrix B (src1) in 8 bit fixed point precision
808 * Matrix A and matrix B must be reshaped respectively with @ref gemm_interleave4x4_8bit and @ref gemm_transpose1x16 before running the matrix multiplication
809 *
810 * @attention The width of matrix B, the alpha's value and fixed point position need to be passed at compile time using -DWIDTH_MATRIX_B -DALPHA and -DFIXED_POINT_POSITION
811 *
812 * @note: ALPHA must be passed in 8 bit fixed point format
813 *
814 * @param[in] src0_ptr Pointer to the source matrix. Supported data types: QS8
815 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes)
816 * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
817 * @param[in] src0_stride_y Stride of the source matrix in Y dimension (in bytes)
818 * @param[in] src0_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
819 * @param[in] src0_offset_first_element_in_bytes The offset of the first element in the source matrix
820 * @param[in] src1_ptr Pointer to the source matrix. Supported data types: same as @p src0_ptr
821 * @param[in] src1_stride_x Stride of the source matrix in X dimension (in bytes)
822 * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
823 * @param[in] src1_stride_y Stride of the source matrix in Y dimension (in bytes)
824 * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
825 * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source matrix
826 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src0_ptr
827 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
828 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
829 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
830 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
831 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
832 */
833__kernel void gemm_mm_qs8(IMAGE_DECLARATION(src0),
834 IMAGE_DECLARATION(src1),
835 IMAGE_DECLARATION(dst))
836{
837 /* src_addr.s0 = address of matrix A */
838 /* src_addr.s1 = address of matrix B */
839
840 /* Compute address for matrix A and B */
841 int2 src_addr = (int2)(get_global_id(1), get_global_id(0)) * (int2)((src0_stride_y),
842 (src1_stride_y));
843
844 /* Add offset_first_element_in_bytes */
845 src_addr = src_addr + ((int2)(src0_offset_first_element_in_bytes, src1_offset_first_element_in_bytes));
846
847 /* Compute end row address for matrix B */
848 int end_row_mtx_b = src_addr.s1 + WIDTH_MATRIX_B;
849
850 /* Reset accumulators */
851 short8 c00 = 0.0f;
852 short8 c10 = 0.0f;
853 short8 c20 = 0.0f;
854 short8 c30 = 0.0f;
855 short8 c01 = 0.0f;
856 short8 c11 = 0.0f;
857 short8 c21 = 0.0f;
858 short8 c31 = 0.0f;
859
860 /* This for loop performs 1 accumulation for each iteration */
861 for(; src_addr.s1 <= (end_row_mtx_b - 16); src_addr += (int2)(4, 16))
862 {
863 /* Load values from matrix A (interleaved) and matrix B (transposed) */
864 char4 a0 = vload4(0, ((__global char *)src0_ptr) + src_addr.s0);
865 char16 b0 = vload16(0, ((__global char *)src1_ptr) + src_addr.s1);
866
867 c00 = mlal_sat_qs8x8(c00, (char8)a0.s0, b0.s01234567, FIXED_POINT_POSITION);
868 c10 = mlal_sat_qs8x8(c10, (char8)a0.s1, b0.s01234567, FIXED_POINT_POSITION);
869 c20 = mlal_sat_qs8x8(c20, (char8)a0.s2, b0.s01234567, FIXED_POINT_POSITION);
870 c30 = mlal_sat_qs8x8(c30, (char8)a0.s3, b0.s01234567, FIXED_POINT_POSITION);
871
872 c01 = mlal_sat_qs8x8(c01, (char8)a0.s0, b0.s89ABCDEF, FIXED_POINT_POSITION);
873 c11 = mlal_sat_qs8x8(c11, (char8)a0.s1, b0.s89ABCDEF, FIXED_POINT_POSITION);
874 c21 = mlal_sat_qs8x8(c21, (char8)a0.s2, b0.s89ABCDEF, FIXED_POINT_POSITION);
875 c31 = mlal_sat_qs8x8(c31, (char8)a0.s3, b0.s89ABCDEF, FIXED_POINT_POSITION);
876 }
877
878 /* Compute destination address */
879 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
880
881 /* Multiply by the weight of matrix product */
882 char16 c00_qs8 = convert_char16_sat((short16)(c00, c01));
883 char16 c10_qs8 = convert_char16_sat((short16)(c10, c11));
884 char16 c20_qs8 = convert_char16_sat((short16)(c20, c21));
885 char16 c30_qs8 = convert_char16_sat((short16)(c30, c31));
886
887 c00_qs8 = mul_sat_qs8x16(c00_qs8, (char16)ALPHA, FIXED_POINT_POSITION);
888 c10_qs8 = mul_sat_qs8x16(c10_qs8, (char16)ALPHA, FIXED_POINT_POSITION);
889 c20_qs8 = mul_sat_qs8x16(c20_qs8, (char16)ALPHA, FIXED_POINT_POSITION);
890 c30_qs8 = mul_sat_qs8x16(c30_qs8, (char16)ALPHA, FIXED_POINT_POSITION);
891
892 /* Store 16x4 block */
893 vstore16(c00_qs8, 0, (__global char *)(offset(&dst, 0, 0)));
894 vstore16(c10_qs8, 0, (__global char *)(offset(&dst, 0, 1)));
895 vstore16(c20_qs8, 0, (__global char *)(offset(&dst, 0, 2)));
896 vstore16(c30_qs8, 0, (__global char *)(offset(&dst, 0, 3)));
897}
Gian Marco Iodice8a383692017-07-03 17:41:47 +0100[diff] [blame]898
899/** This OpenCL kernel computes the matrix multiplication between matrix A (src0) and matrix B (src1) in 16 bit fixed point precision
900 * Matrix A and matrix B must be reshaped respectively with @ref gemm_interleave4x4_16bit and @ref gemm_transpose1x8 before running the matrix multiplication
901 *
902 * @attention The width of matrix B, the alpha's value and fixed point position need to be passed at compile time using -DWIDTH_MATRIX_B -DALPHA and -DFIXED_POINT_POSITION
903 *
904 * @note: ALPHA must be passed in 16 bit fixed point format
905 *
906 * @param[in] src0_ptr Pointer to the source matrix. Supported data types: QS16
907 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes)
908 * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
909 * @param[in] src0_stride_y Stride of the source matrix in Y dimension (in bytes)
910 * @param[in] src0_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
911 * @param[in] src0_offset_first_element_in_bytes The offset of the first element in the source matrix
912 * @param[in] src1_ptr Pointer to the source matrix. Supported data types: same as @p src0_ptr
913 * @param[in] src1_stride_x Stride of the source matrix in X dimension (in bytes)
914 * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
915 * @param[in] src1_stride_y Stride of the source matrix in Y dimension (in bytes)
916 * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
917 * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source matrix
918 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src0_ptr
919 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
920 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
921 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
922 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
923 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
924 */
925__kernel void gemm_mm_qs16(IMAGE_DECLARATION(src0),
926 IMAGE_DECLARATION(src1),
927 IMAGE_DECLARATION(dst))
928{
929 /* src_addr.s0 = address of matrix A */
930 /* src_addr.s1 = address of matrix B */
931
932 /* Compute address for matrix A and B */
933 int2 src_addr = (int2)(get_global_id(1), get_global_id(0)) * (int2)((src0_stride_y),
934 (src1_stride_y));
935
936 /* Add offset_first_element_in_bytes */
937 src_addr = src_addr + ((int2)(src0_offset_first_element_in_bytes, src1_offset_first_element_in_bytes));
938
939 /* Divide by 2 in order to get the src_addr in unit of short */
940 src_addr = src_addr >> 1;
941
942 /* Compute end row address for matrix B */
943 int end_row_mtx_b = src_addr.s1 + WIDTH_MATRIX_B;
944
945 /* Reset accumulators */
946 int8 c00 = 0.0f;
947 int8 c10 = 0.0f;
948 int8 c20 = 0.0f;
949 int8 c30 = 0.0f;
950
951 /* This for loop performs 1 accumulation for each iteration */
952 for(; src_addr.s1 <= (end_row_mtx_b - 8); src_addr += (int2)(4, 8))
953 {
954 /* Load values from matrix A (interleaved) and matrix B (transposed) */
955 short4 a0 = vload4(0, ((__global short *)src0_ptr) + src_addr.s0);
956 short8 b0 = vload8(0, ((__global short *)src1_ptr) + src_addr.s1);
957
958 c00 = mlal_sat_qs16x8(c00, (short8)a0.s0, b0, FIXED_POINT_POSITION);
959 c10 = mlal_sat_qs16x8(c10, (short8)a0.s1, b0, FIXED_POINT_POSITION);
960 c20 = mlal_sat_qs16x8(c20, (short8)a0.s2, b0, FIXED_POINT_POSITION);
961 c30 = mlal_sat_qs16x8(c30, (short8)a0.s3, b0, FIXED_POINT_POSITION);
962 }
963
964 /* Compute destination address */
965 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
966
967 /* Multiply by the weight of matrix product */
968 short8 c00_qs16 = convert_short8_sat(c00);
969 short8 c10_qs16 = convert_short8_sat(c10);
970 short8 c20_qs16 = convert_short8_sat(c20);
971 short8 c30_qs16 = convert_short8_sat(c30);
972
973 c00_qs16 = mul_sat_qs16x8(c00_qs16, (short8)ALPHA, FIXED_POINT_POSITION);
974 c10_qs16 = mul_sat_qs16x8(c10_qs16, (short8)ALPHA, FIXED_POINT_POSITION);
975 c20_qs16 = mul_sat_qs16x8(c20_qs16, (short8)ALPHA, FIXED_POINT_POSITION);
976 c30_qs16 = mul_sat_qs16x8(c30_qs16, (short8)ALPHA, FIXED_POINT_POSITION);
977
978 /* Store 8x4 block */
979 vstore8(c00_qs16, 0, (__global short *)(offset(&dst, 0, 0)));
980 vstore8(c10_qs16, 0, (__global short *)(offset(&dst, 0, 1)));
981 vstore8(c20_qs16, 0, (__global short *)(offset(&dst, 0, 2)));
982 vstore8(c30_qs16, 0, (__global short *)(offset(&dst, 0, 3)));
983}
984#endif // defined(FIXED_POINT_POSITION)
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]985
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]986#ifdef WIDTH_VECTOR_A
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]987/** This OpenCL kernel computes the vector by matrix multiplication between the vector A (src0) and matrix B (src1)
988 *
989 * @attention The width of vector A, the width of matrix B and the alpha's value need to be passed at compile time using -DWIDTH_VECTOR_A -DWIDTH_MATRIX_B and -DALPHA
990 *
991 * @attention The input vector A and matrix B must not be reshaped
992 *
993 * @param[in] src0_ptr Pointer to the source matrix. Supported data types: F32
994 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes)
995 * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
996 * @param[in] src0_stride_y Stride of the source matrix in Y dimension (in bytes)
997 * @param[in] src0_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
998 * @param[in] src0_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]999 * @param[in] src1_ptr Pointer to the source matrix. Supported data types: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1000 * @param[in] src1_stride_x Stride of the source matrix in X dimension (in bytes)
1001 * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
1002 * @param[in] src1_stride_y Stride of the source matrix in Y dimension (in bytes)
1003 * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
1004 * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]1005 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1006 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
1007 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
1008 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
1009 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
1010 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
1011 */
1012__kernel void gemm_vm_f32(IMAGE_DECLARATION(src0),
1013 IMAGE_DECLARATION(src1),
1014 IMAGE_DECLARATION(dst))
1015{
1016 int idx = get_global_id(0) * 4;
1017
1018 /* Compute the address for the vector A and matrix B */
1019 int2 src_addr = ((int2)(src0_offset_first_element_in_bytes, src1_offset_first_element_in_bytes));
1020 src_addr.s1 += idx * sizeof(float);
1021
1022 int end_row_vec_a = src_addr.s0 + (WIDTH_VECTOR_A * sizeof(float));
1023
1024 float4 acc = 0.0f;
1025
1026 for(; src_addr.s0 <= (end_row_vec_a - 2 * sizeof(float)); src_addr += (int2)(2 * sizeof(float), 2 * src1_stride_y))
1027 {
1028 float2 a0 = vload2(0, (__global float *)(src0_ptr + src_addr.s0));
1029 float4 b0 = vload4(0, (__global float *)(src1_ptr + src_addr.s1));
1030 float4 b1 = vload4(0, (__global float *)(src1_ptr + src_addr.s1 + src1_stride_y));
1031
1032 acc += b0 * (float4)a0.s0;
1033 acc += b1 * (float4)a0.s1;
1034 }
1035
1036 for(; src_addr.s0 < end_row_vec_a; src_addr += (int2)(sizeof(float), src1_stride_y))
1037 {
1038 float a0 = *((__global float *)(src0_ptr + src_addr.s0));
1039 float4 b0 = vload4(0, (__global float *)(src1_ptr + src_addr.s1));
1040
1041 acc += b0 * (float4)a0;
1042 }
1043
1044 /* Compute destination address */
1045 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
1046
1047 /* Multiply by the weight of vector-matrix product */
1048 acc = acc * (float4)ALPHA;
1049
1050 vstore4(acc, 0, (__global float *)(offset(&dst, 0, 0)));
1051}
1052
1053/** This OpenCL kernel computes the vector by matrix multiplication between the vector A (src0) and matrix B (src1)
1054 *
1055 * @attention The width of vector A, the width of matrix B and the alpha's value need to be passed at compile time using -DWIDTH_VECTOR_A -DWIDTH_MATRIX_B and -DALPHA
1056 *
1057 * @attention The input vector A and matrix B must not be reshaped
1058 *
1059 * @param[in] src0_ptr Pointer to the source matrix. Supported data types: F16
1060 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes)
1061 * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
1062 * @param[in] src0_stride_y Stride of the source matrix in Y dimension (in bytes)
1063 * @param[in] src0_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
1064 * @param[in] src0_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]1065 * @param[in] src1_ptr Pointer to the source matrix. Supported data types: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1066 * @param[in] src1_stride_x Stride of the source matrix in X dimension (in bytes)
1067 * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
1068 * @param[in] src1_stride_y Stride of the source matrix in Y dimension (in bytes)
1069 * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
1070 * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]1071 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1072 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
1073 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
1074 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
1075 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
1076 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
1077 */
1078__kernel void gemm_vm_f16(IMAGE_DECLARATION(src0),
1079 IMAGE_DECLARATION(src1),
1080 IMAGE_DECLARATION(dst))
1081{
1082 int idx = get_global_id(0) * 8;
1083
1084 /* Compute the address for the vector A and matrix B */
1085 int2 src_addr = ((int2)(src0_offset_first_element_in_bytes, src1_offset_first_element_in_bytes));
1086 src_addr.s1 += idx * sizeof(half);
1087
1088 int end_row_vec_a = src_addr.s0 + (WIDTH_VECTOR_A * sizeof(half));
1089
1090 half8 acc = 0.0f;
1091
1092 for(; src_addr.s0 <= (end_row_vec_a - 4 * sizeof(half)); src_addr += (int2)(4 * sizeof(half), 4 * src1_stride_y))
1093 {
1094 half4 a0 = vload4(0, (__global half *)(src0_ptr + src_addr.s0));
1095 half8 b0 = vload8(0, (__global half *)(src1_ptr + src_addr.s1 + 0 * src1_stride_y));
1096 half8 b1 = vload8(0, (__global half *)(src1_ptr + src_addr.s1 + 1 * src1_stride_y));
1097 half8 b2 = vload8(0, (__global half *)(src1_ptr + src_addr.s1 + 2 * src1_stride_y));
1098 half8 b3 = vload8(0, (__global half *)(src1_ptr + src_addr.s1 + 3 * src1_stride_y));
1099
1100 acc += b0 * (half8)a0.s0;
1101 acc += b1 * (half8)a0.s1;
1102 acc += b2 * (half8)a0.s2;
1103 acc += b3 * (half8)a0.s3;
1104 }
1105
1106 for(; src_addr.s0 < end_row_vec_a; src_addr += (int2)(sizeof(half), src1_stride_y))
1107 {
1108 half a0 = *((__global half *)(src0_ptr + src_addr.s0));
1109 half8 b0 = vload8(0, (__global half *)(src1_ptr + src_addr.s1));
1110
1111 acc += b0 * (half8)a0;
1112 }
1113
1114 /* Compute destination address */
1115 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
1116
1117 /* Multiply by the weight of vector-matrix product */
1118 acc = acc * (half8)ALPHA;
1119
1120 vstore8(acc, 0, (__global half *)(offset(&dst, 0, 0)));
1121}
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]1122
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]1123#ifdef FIXED_POINT_POSITION
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]1124/** This OpenCL kernel computes the vector by matrix multiplication between the vector A (src0) and matrix B (src1) in 8 bit fixed point
1125 *
1126 * @attention The width of vector A, the width of matrix B, the alpha's value and the fixed point position need to be passed at compile time using -DWIDTH_VECTOR_A -DWIDTH_MATRIX_B, -DALPHA and -DFIXED_POINT_POSITION
1127 *
1128 * @attention The input vector A and matrix B must not be reshaped
1129 *
1130 * @note: ALPHA must be passed in 8 bit fixed point format
1131 *
1132 * @param[in] src0_ptr Pointer to the source matrix. Supported data types: QS8
1133 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes)
1134 * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
1135 * @param[in] src0_stride_y Stride of the source matrix in Y dimension (in bytes)
1136 * @param[in] src0_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
1137 * @param[in] src0_offset_first_element_in_bytes The offset of the first element in the source matrix
1138 * @param[in] src1_ptr Pointer to the source matrix. Supported data types: same as @p src0_ptr
1139 * @param[in] src1_stride_x Stride of the source matrix in X dimension (in bytes)
1140 * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
1141 * @param[in] src1_stride_y Stride of the source matrix in Y dimension (in bytes)
1142 * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
1143 * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source matrix
1144 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src0_ptr
1145 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
1146 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
1147 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
1148 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
1149 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
1150 */
1151__kernel void gemm_vm_qs8(IMAGE_DECLARATION(src0),
1152 IMAGE_DECLARATION(src1),
1153 IMAGE_DECLARATION(dst))
1154{
1155 int idx = get_global_id(0) * 16;
1156
1157 /* Compute the address for the vector A and matrix B */
1158 int2 src_addr = ((int2)(src0_offset_first_element_in_bytes, src1_offset_first_element_in_bytes));
1159 src_addr.s1 += idx;
1160
1161 int end_row_vec_a = src_addr.s0 + WIDTH_VECTOR_A;
1162
1163 short8 acc0 = 0;
1164 short8 acc1 = 0;
1165
1166 /* This for loop performs 4 accumulations per iteration */
1167 for(; src_addr.s0 <= (end_row_vec_a - 4); src_addr += (int2)(4, 4 * src1_stride_y))
1168 {
1169 char4 a0 = vload4(0, (__global char *)(src0_ptr + src_addr.s0));
1170 char16 b0 = vload16(0, (__global char *)(src1_ptr + src_addr.s1 + 0 * src1_stride_y));
1171 char16 b1 = vload16(0, (__global char *)(src1_ptr + src_addr.s1 + 1 * src1_stride_y));
1172 char16 b2 = vload16(0, (__global char *)(src1_ptr + src_addr.s1 + 2 * src1_stride_y));
1173 char16 b3 = vload16(0, (__global char *)(src1_ptr + src_addr.s1 + 3 * src1_stride_y));
1174
1175 acc0 = mlal_sat_qs8x8(acc0, (char8)a0.s0, b0.s01234567, FIXED_POINT_POSITION);
1176 acc0 = mlal_sat_qs8x8(acc0, (char8)a0.s1, b1.s01234567, FIXED_POINT_POSITION);
1177 acc0 = mlal_sat_qs8x8(acc0, (char8)a0.s2, b2.s01234567, FIXED_POINT_POSITION);
1178 acc0 = mlal_sat_qs8x8(acc0, (char8)a0.s3, b3.s01234567, FIXED_POINT_POSITION);
1179
1180 acc1 = mlal_sat_qs8x8(acc1, (char8)a0.s0, b0.s89ABCDEF, FIXED_POINT_POSITION);
1181 acc1 = mlal_sat_qs8x8(acc1, (char8)a0.s1, b1.s89ABCDEF, FIXED_POINT_POSITION);
1182 acc1 = mlal_sat_qs8x8(acc1, (char8)a0.s2, b2.s89ABCDEF, FIXED_POINT_POSITION);
1183 acc1 = mlal_sat_qs8x8(acc1, (char8)a0.s3, b3.s89ABCDEF, FIXED_POINT_POSITION);
1184 }
1185
1186 /* Left-over accumulations */
1187 for(; src_addr.s0 < end_row_vec_a; src_addr += (int2)(1, src1_stride_y))
1188 {
1189 char a0 = *((__global char *)(src0_ptr + src_addr.s0));
1190 char16 b0 = vload16(0, (__global char *)(src1_ptr + src_addr.s1));
1191
1192 acc0 = mlal_sat_qs8x8(acc0, (char8)a0, b0.s01234567, FIXED_POINT_POSITION);
1193 acc1 = mlal_sat_qs8x8(acc1, (char8)a0, b0.s89ABCDEF, FIXED_POINT_POSITION);
1194 }
1195
1196 /* Compute destination address */
1197 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
1198
1199 /* Multiply by the weight of matrix product */
1200 char16 acc_qs8 = convert_char16_sat((short16)(acc0, acc1));
1201
1202 acc_qs8 = mul_sat_qs8x16(acc_qs8, (char16)ALPHA, FIXED_POINT_POSITION);
1203
1204 /* Store 16 values */
1205 vstore16(acc_qs8, 0, (__global char *)(offset(&dst, 0, 0)));
1206}
Gian Marco Iodice8a383692017-07-03 17:41:47 +0100[diff] [blame]1207
1208/** This OpenCL kernel computes the vector by matrix multiplication between the vector A (src0) and matrix B (src1) in 16 bit fixed point
1209 *
1210 * @attention The width of vector A, the width of matrix B, the alpha's value and the fixed point position need to be passed at compile time using -DWIDTH_VECTOR_A -DWIDTH_MATRIX_B, -DALPHA and -DFIXED_POINT_POSITION
1211 *
1212 * @attention The input vector A and matrix B must not be reshaped
1213 *
1214 * @note: ALPHA must be passed in 16 bit fixed point format
1215 *
1216 * @param[in] src0_ptr Pointer to the source matrix. Supported data types: QS16
1217 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes)
1218 * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
1219 * @param[in] src0_stride_y Stride of the source matrix in Y dimension (in bytes)
1220 * @param[in] src0_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
1221 * @param[in] src0_offset_first_element_in_bytes The offset of the first element in the source matrix
1222 * @param[in] src1_ptr Pointer to the source matrix. Supported data types: same as @p src0_ptr
1223 * @param[in] src1_stride_x Stride of the source matrix in X dimension (in bytes)
1224 * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
1225 * @param[in] src1_stride_y Stride of the source matrix in Y dimension (in bytes)
1226 * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
1227 * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source matrix
1228 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src0_ptr
1229 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
1230 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
1231 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
1232 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
1233 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
1234 */
1235__kernel void gemm_vm_qs16(IMAGE_DECLARATION(src0),
1236 IMAGE_DECLARATION(src1),
1237 IMAGE_DECLARATION(dst))
1238{
1239 int idx = get_global_id(0) * 8;
1240
1241 /* Compute the address for the vector A and matrix B */
1242 int2 src_addr = ((int2)(src0_offset_first_element_in_bytes, src1_offset_first_element_in_bytes));
1243 src_addr.s1 += idx * sizeof(short);
1244
1245 int end_row_vec_a = src_addr.s0 + (WIDTH_VECTOR_A * sizeof(short));
1246
1247 /* Reset accumulator */
1248 int8 acc0 = 0;
1249
1250 /* This for loop performs 4 accumulations per iteration */
1251 for(; src_addr.s0 <= (end_row_vec_a - 4 * sizeof(short)); src_addr += (int2)(4 * sizeof(short), 4 * src1_stride_y))
1252 {
1253 short4 a0 = vload4(0, (__global short *)(src0_ptr + src_addr.s0));
1254 short8 b0 = vload8(0, (__global short *)(src1_ptr + src_addr.s1 + 0 * src1_stride_y));
1255 short8 b1 = vload8(0, (__global short *)(src1_ptr + src_addr.s1 + 1 * src1_stride_y));
1256 short8 b2 = vload8(0, (__global short *)(src1_ptr + src_addr.s1 + 2 * src1_stride_y));
1257 short8 b3 = vload8(0, (__global short *)(src1_ptr + src_addr.s1 + 3 * src1_stride_y));
1258
1259 acc0 = mlal_sat_qs16x8(acc0, (short8)a0.s0, b0, FIXED_POINT_POSITION);
1260 acc0 = mlal_sat_qs16x8(acc0, (short8)a0.s1, b1, FIXED_POINT_POSITION);
1261 acc0 = mlal_sat_qs16x8(acc0, (short8)a0.s2, b2, FIXED_POINT_POSITION);
1262 acc0 = mlal_sat_qs16x8(acc0, (short8)a0.s3, b3, FIXED_POINT_POSITION);
1263 }
1264
1265 /* Left-over accumulations */
1266 for(; src_addr.s0 < end_row_vec_a; src_addr += (int2)(sizeof(short), src1_stride_y))
1267 {
1268 short a0 = *((__global short *)(src0_ptr + src_addr.s0));
1269 short8 b0 = vload8(0, (__global short *)(src1_ptr + src_addr.s1));
1270
1271 acc0 = mlal_sat_qs16x8(acc0, (short8)a0, b0, FIXED_POINT_POSITION);
1272 }
1273
1274 /* Compute destination address */
1275 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
1276
1277 /* Multiply by the weight of matrix product */
1278 short8 acc_qs16 = convert_short8_sat(acc0);
1279
1280 acc_qs16 = mul_sat_qs16x8(acc_qs16, (short8)ALPHA, FIXED_POINT_POSITION);
1281
1282 /* Store 8 values */
1283 vstore8(acc_qs16, 0, (__global short *)(offset(&dst, 0, 0)));
1284}
1285#endif /* defined(FIXED_POINT_POSITION) */
1286#endif /* defined(WIDTH_VECTOR_A) */
1287#endif /* defined(WIDTH_MATRIX_B) && defined(ALPHA) */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1288
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]1289#ifdef BETA
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1290/** This OpenCL kernel performs the in-place matrix addition between 2 matrices taking into account that the second matrix might be weighted by a scalar value beta:
1291 *
1292 * @attention The beta's value need to be passed at compile time using -DBETA
1293 *
1294 * @param[in] src_ptr Pointer to the source matrix. Supported data types: F32
1295 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes)
1296 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
1297 * @param[in] src_stride_y Stride of the source matrix in Y dimension (in bytes)
1298 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
1299 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]1300 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1301 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
1302 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
1303 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
1304 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
1305 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
1306 */
1307__kernel void gemm_ma_f32(IMAGE_DECLARATION(src),
1308 IMAGE_DECLARATION(dst))
1309{
1310 /* Compute source and destination addresses */
1311 Image src = CONVERT_TO_IMAGE_STRUCT(src);
1312 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
1313
1314 /* Load values from A x B */
1315 float4 alpha_ab = vload4(0, (__global float *)dst.ptr);
1316
1317 /* Load values from Matrix C */
1318 float4 c = vload4(0, (__global float *)src.ptr);
1319
1320 /* Computes alpha * axb + beta * c */
1321 float4 out = alpha_ab + (float4)BETA * c;
1322
1323 /* Store final result in axb matrix */
1324 vstore4(out, 0, (__global float *)dst.ptr);
1325}
1326
1327/** This OpenCL kernel performs the in-place matrix addition between 2 matrices taking into account that the second matrix might be weighted by a scalar value beta:
1328 *
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]1329 * @attention The beta's value need to be passed at compile time using -DBETA
1330 *
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1331 * @param[in] src_ptr Pointer to the source matrix. Supported data types: F16
1332 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes)
1333 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
1334 * @param[in] src_stride_y Stride of the source matrix in Y dimension (in bytes)
1335 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
1336 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]1337 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1338 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
1339 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
1340 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
1341 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
1342 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
1343 */
1344__kernel void gemm_ma_f16(IMAGE_DECLARATION(src),
1345 IMAGE_DECLARATION(dst))
1346{
1347 /* Compute source and destination addresses */
1348 Image src = CONVERT_TO_IMAGE_STRUCT(src);
1349 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
1350
1351 /* Load values from A x B */
1352 half8 alpha_ab = vload8(0, (__global half *)dst.ptr);
1353
1354 /* Load values from Matrix C */
1355 half8 c = vload8(0, (__global half *)src.ptr);
1356
1357 /* Computes alpha * axb + beta * c */
1358 half8 out = alpha_ab + (half8)BETA * c;
1359
1360 /* Store final result in axb matrix */
1361 vstore8(out, 0, (__global half *)dst.ptr);
1362}
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]1363
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]1364#ifdef FIXED_POINT_POSITION
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]1365/** This OpenCL kernel performs the in-place matrix addition between 2 matrices in 8 bit fixed point taking into account that the second matrix might be weighted by a scalar value beta:
1366 *
1367 * @attention The beta's value and the fixed point position need to be passed at compile time using -DBETA and -DFIXED_POINT_POSITION
1368 *
1369 * @note: BETA must be passed in 8 bit fixed point format
1370 *
1371 * @param[in] src_ptr Pointer to the source matrix. Supported data types: QS8
1372 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes)
1373 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
1374 * @param[in] src_stride_y Stride of the source matrix in Y dimension (in bytes)
1375 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
1376 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source matrix
1377 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src_ptr
1378 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
1379 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
1380 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
1381 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
1382 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
1383 */
1384__kernel void gemm_ma_qs8(IMAGE_DECLARATION(src),
1385 IMAGE_DECLARATION(dst))
1386{
1387 /* Compute source and destination addresses */
1388 Image src = CONVERT_TO_IMAGE_STRUCT(src);
1389 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
1390
1391 /* Load values from A x B */
1392 char16 alpha_ab = vload16(0, (__global char *)dst.ptr);
1393
1394 /* Load values from Matrix C */
1395 char16 c = vload16(0, (__global char *)src.ptr);
1396
1397 /* Computes alpha * axb + beta * c */
1398 char16 out = mla_sat_qs8x16(alpha_ab, (char16)BETA, c, FIXED_POINT_POSITION);
1399
1400 /* Store final result in axb matrix */
1401 vstore16(out, 0, (__global char *)dst.ptr);
1402}
Gian Marco Iodice8a383692017-07-03 17:41:47 +0100[diff] [blame]1403
1404/** This OpenCL kernel performs the in-place matrix addition between 2 matrices in 16 bit fixed point taking into account that the second matrix might be weighted by a scalar value beta:
1405 *
1406 * @attention The beta's value and the fixed point position need to be passed at compile time using -DBETA and -DFIXED_POINT_POSITION
1407 *
1408 * @note: BETA must be passed in 16 bit fixed point format
1409 *
1410 * @param[in] src_ptr Pointer to the source matrix. Supported data types: QS16
1411 * @param[in] src_stride_x Stride of the source matrix in X dimension (in bytes)
1412 * @param[in] src_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
1413 * @param[in] src_stride_y Stride of the source matrix in Y dimension (in bytes)
1414 * @param[in] src_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
1415 * @param[in] src_offset_first_element_in_bytes The offset of the first element in the source matrix
1416 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src_ptr
1417 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
1418 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
1419 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
1420 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
1421 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
1422 */
1423__kernel void gemm_ma_qs16(IMAGE_DECLARATION(src),
1424 IMAGE_DECLARATION(dst))
1425{
1426 /* Compute source and destination addresses */
1427 Image src = CONVERT_TO_IMAGE_STRUCT(src);
1428 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
1429
1430 /* Load values from A x B */
1431 short8 alpha_ab = vload8(0, (__global short *)dst.ptr);
1432
1433 /* Load values from Matrix C */
1434 short8 c = vload8(0, (__global short *)src.ptr);
1435
1436 /* Computes alpha * axb + beta * c */
1437 short8 out = mla_sat_qs16x8(alpha_ab, (short8)BETA, c, FIXED_POINT_POSITION);
1438
1439 /* Store final result in axb matrix */
1440 vstore8(out, 0, (__global short *)dst.ptr);
1441}
1442#endif /* defined(FIXED_POINT_POSITION) */
1443#endif /* defined(BETA) */
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1444
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]1445#ifdef WIDTH_VECTOR_A
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1446/** This OpenCL kernel computes the vector by matrix multiplication between each row of A (src0) and matrix B (src1) used for locally connected layer
1447 *
1448 * @attention The width of A need to be passed at compile time using -DWIDTH_VECTOR_A
1449 *
1450 * @attention The input A and matrix B must not be reshaped
1451 *
1452 * @param[in] src0_ptr Pointer to the source matrix. Supported data types: F32
1453 * @param[in] src0_stride_x Stride of the source matrix in X dimension (in bytes)
1454 * @param[in] src0_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
1455 * @param[in] src0_stride_y Stride of the source matrix in Y dimension (in bytes)
1456 * @param[in] src0_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
1457 * @param[in] src0_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]1458 * @param[in] src1_ptr Pointer to the source matrix. Supported data types: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1459 * @param[in] src1_stride_x Stride of the source matrix in X dimension (in bytes)
1460 * @param[in] src1_step_x src_stride_x * number of elements along X processed per workitem(in bytes)
1461 * @param[in] src1_stride_y Stride of the source matrix in Y dimension (in bytes)
1462 * @param[in] src1_step_y src_stride_y * number of elements along Y processed per workitem(in bytes)
1463 * @param[in] src1_stride_z Stride of the source matrix in Z dimension (in bytes)
1464 * @param[in] src1_step_z src_stride_z * number of elements along Z processed per workitem(in bytes)
1465 * @param[in] src1_offset_first_element_in_bytes The offset of the first element in the source matrix
Gian Marco Iodice3a3066b2017-06-23 13:38:14 +0100[diff] [blame]1466 * @param[out] dst_ptr Pointer to the destination matrix Supported data types: same as @p src0_ptr
Anthony Barbier6ff3b192017-09-04 18:44:23 +0100[diff] [blame]1467 * @param[in] dst_stride_x Stride of the destination matrix in X dimension (in bytes)
1468 * @param[in] dst_step_x dst_gx_stride_x * number of elements along X processed per workitem(in bytes)
1469 * @param[in] dst_stride_y Stride of the destination matrix in Y dimension (in bytes)
1470 * @param[in] dst_step_y dst_gx_stride_y * number of elements along Y processed per workitem(in bytes)
1471 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination matrix
1472 */
1473__kernel void gemm_lc_vm_f32(IMAGE_DECLARATION(src0),
1474 TENSOR3D_DECLARATION(src1),
1475 IMAGE_DECLARATION(dst))
1476{
1477 int idx = get_global_id(0) * 4;
1478 int idy = get_global_id(1);
1479
1480 /* Compute the address for the vector A and matrix B */
1481 int2 src_addr = ((int2)(src0_offset_first_element_in_bytes + src0_stride_y * idy, src1_offset_first_element_in_bytes + src1_stride_z * idy));
1482 src_addr.s1 += idx * sizeof(float);
1483
1484 int end_row_vec_a = src_addr.s0 + (WIDTH_VECTOR_A * sizeof(float));
1485
1486 float4 acc = 0.0f;
1487
1488 for(; src_addr.s0 <= (end_row_vec_a - 2 * sizeof(float)); src_addr += (int2)(2 * sizeof(float), 2 * src1_stride_y))
1489 {
1490 float2 a0 = vload2(0, (__global float *)(src0_ptr + src_addr.s0));
1491 float4 b0 = vload4(0, (__global float *)(src1_ptr + src_addr.s1));
1492 float4 b1 = vload4(0, (__global float *)(src1_ptr + src_addr.s1 + src1_stride_y));
1493
1494 acc += b0 * (float4)a0.s0;
1495 acc += b1 * (float4)a0.s1;
1496 }
1497
1498 for(; src_addr.s0 < end_row_vec_a; src_addr += (int2)(sizeof(float), src1_stride_y))
1499 {
1500 float a0 = *((__global float *)(src0_ptr + src_addr.s0));
1501 float4 b0 = vload4(0, (__global float *)(src1_ptr + src_addr.s1));
1502
1503 acc += b0 * (float4)a0;
1504 }
1505
1506 /* Compute destination address */
1507 Image dst = CONVERT_TO_IMAGE_STRUCT(dst);
1508
1509 vstore4(acc, 0, (__global float *)(offset(&dst, 0, 0)));
1510}
Anthony Barbierac69aa12017-07-03 17:39:37 +0100[diff] [blame]1511#endif /* WIDTH_VECTOR_A */