1/* 2 * Copyright (c) 2022 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 "activation_float_helpers.h" 25#include "helpers.h" 26#include "tile_helpers.h" 27#if defined(GEMMLOWP_MM_RESHAPED_ONLY_RHS_MMUL) 28/** This OpenCL kernel computes the matrix multiplication between 2 matrices using the MMUL extension: 29 * 30 * The LHS matrix is NOT reshaped 31 * The RHS is reshaped with @ref ClGemmMatrixMultiplyReshapedOnlyRhsKernel and the block K0xN0 is transposed 32 * 33 * @note The block's dimensions used for reshaping the RHS matrix (N0 and K0) must be passed at compile time using -DN0 and -DK0 (e.g. -DN0=1, -DK0=1). 34 * @note The number of M0 rows to process must be passed at compile time using -DM0 (e.g. -DM0=1) 35 * @note The number of output columns processed by the the cooperative mmul extension must be passed at compile time using -DMMUL_N0 (e.g., -DMMUL_N0=4) 36 * @note The number of output rows processed by the the cooperative mmul extension must be passed at compile time using -DMMUL_M0 (e.g., -DMMUL_M0=4) 37 * @note The number of lhs columns (or rhs rows) processed by the the cooperative mmul extension must be passed at compile time using -DMMUL_K0 (e.g., -DMMUL_K0=16) 38 * @note Only the following configurations of M0, N0 and K0 are currently supported: 39 * - M0 = 1, 2, 4 40 * - N0 = 1, 4, 8 41 * - K0 = 4 42 * 43 * @note If the activation type were passed at compile time through -DACTIVATION_TYPE (e.g. -DACTIVATION_TYPE=RELU), A, B variables, required by some activation functions, should be passed at compile time as well using -DA_VAL= and -DB_VAL= respectively. 44 * The activation function is performed after the bias addition 45 * 46 * @param[in] lhs_ptr Pointer to the LHS tensor. Supported data types: QASYMM8/QASYMM8_SIGNED 47 * @param[in] lhs_stride_y Stride of the LHS tensor in Y dimension (in bytes) 48 * @param[in] lhs_stride_z Stride of the LHS tensor in Z dimension (in bytes) 49 * @param[in] lhs_w The size of the width dimension of the LHS tensor 50 * @param[in] lhs_h The size of the height dimension of the LHS tensor 51 * @param[in] lhs_n The size of the depth dimension of the LHS tensor 52 * @param[in] lhs_offset_first_element_in_bytes The offset of the first element in the LHS tensor 53 * @param[in] rhs_ptr Pointer to the RHS reshaped tensor. Supported data type: same as @p lhs_ptr 54 * @param[in] rhs_stride_y Stride of the RHS tensor in Y dimension (in bytes) 55 * @param[in] rhs_stride_z Stride of the RHS tensor in Z dimension (in bytes) 56 * @param[in] rhs_w The size of the width dimension of the RHS tensor 57 * @param[in] rhs_h The size of the height dimension of the RHS tensor 58 * @param[in] rhs_n The size of the depth dimension of the RHS tensor 59 * @param[in] rhs_offset_first_element_in_bytes The offset of the first element in the RHS tensor 60 * @param[in] bia_ptr (Optional) Pointer to the bias tensor. Supported data type: S32 61 * @param[in] bia_stride_y (Optional) Stride of the bias tensor in Y dimension (in bytes) 62 * @param[in] bia_stride_z (Optional) Stride of the bias tensor in Z dimension (in bytes) 63 * @param[in] bia_w (Optional) The size of the width dimension of the bias tensor 64 * @param[in] bia_h (Optional) The size of the height dimension of the bias tensor 65 * @param[in] bia_n (Optional) The size of the depth dimension of the bias tensor 66 * @param[in] bia_offset_first_element_in_bytes (Optional) The offset of the first element in the bias tensor 67 * @param[out] dst_ptr Pointer to the destination tensor. Supported data type: same as @p lhs_ptr or S32 68 * @param[in] dst_stride_y Stride of the destination tensor in Y dimension (in bytes) 69 * @param[in] dst_stride_z Stride of the destination tensor in Z dimension (in bytes) 70 * @param[in] dst_w The size of the width dimension of the destination tensor 71 * @param[in] dst_h The size of the height dimension of the destination tensor 72 * @param[in] dst_n The size of the depth dimension of the destination tensor 73 * @param[in] dst_offset_first_element_in_bytes The offset of the first element in the destination tensor 74 * @param[in] M Number of rows in LHS matrix not reshaped 75 * @param[in] N Number of columns in RHS matrix not reshaped 76 * @param[in] K Number of columns in LHS matrix and rows in RHS matrix not reshaped 77 * @param[in] sum_col_ptr (Optional) Pointer to the source tensor. Supported data type: S32 78 * @param[in] sum_col_stride_x (Optional) Stride of the source tensor in X dimension (in bytes) 79 * @param[in] sum_col_step_x (Optional) sum_col_stride_x * number of elements along X processed per workitem(in bytes) 80 * @param[in] sum_col_stride_y (Optional) Stride of the source tensor in Y dimension (in bytes) 81 * @param[in] sum_col_step_y (Optional) sum_col_stride_y * number of elements along Y processed per workitem(in bytes) 82 * @param[in] sum_col_offset_first_element_in_bytes (Optional) The offset of the first element in the source tensor 83 * @param[in] sum_row_ptr (Optional) Pointer to the source tensor. Supported data type: S32 84 * @param[in] sum_row_stride_x (Optional) Stride of the source tensor in X dimension (in bytes) 85 * @param[in] sum_row_step_x (Optional) sum_row_stride_x * number of elements along X processed per workitem(in bytes) 86 * @param[in] sum_row_stride_y (Optional) Stride of the source tensor in Y dimension (in bytes) 87 * @param[in] sum_row_step_y (Optional) sum_row_stride_y * number of elements along Y processed per workitem(in bytes) 88 * @param[in] sum_row_offset_first_element_in_bytes (Optional) The offset of the first element in the source tensor 89 */ 90__kernel void gemmlowp_mm_reshaped_only_rhs_mmul( 91 TENSOR3D_T(lhs, BUFFER), 92 TENSOR3D_T(rhs, BUFFER), 93#if defined(ADD_BIAS) 94 TENSOR3D_T(bia, BUFFER), 95#endif // defined(ADD_BIAS) 96 TENSOR3D_T(dst, BUFFER), 97 const int M, 98 const int N, 99 const int K 100#if defined(A_OFFSET) 101 , 102 TENSOR3D_T(sum_col, BUFFER) 103#endif // defined(A_OFFSET) 104#if defined(B_OFFSET) 105 , 106 TENSOR3D_T(sum_row, BUFFER) 107#endif // defined(B_OFFSET) 108) 109{ 110#define MMUL_BLOCK_SIZE (MMUL_N0 * MMUL_M0) 111#define VEC_SIZE 4 // For int8 types input to mmul instruction is a length 4 vector 112 113 uint x0 = get_global_id(0); 114 uint y0 = get_global_id(1); 115 uint z = get_global_id(2); 116 117 // Get block ID and thread ID within the block 118 uint block_id = (x0 / MMUL_BLOCK_SIZE); 119 uint thread_id = (x0 % MMUL_BLOCK_SIZE); 120 121 // Coordinate within a block 122 uint block_x = thread_id % MMUL_N0; 123 uint block_y = (thread_id / MMUL_M0); 124 125 // Starting destination coordinates 126 uint dst_x = min(block_x * N0 + block_id * MMUL_N0 * N0, (uint)(N - 1)); 127 uint dst_y = min(block_y * M0 + y0 * M0 * MMUL_M0, (uint)(M - M0)); 128 129 uint lhs_x = VEC_SIZE * block_x; 130 uint lhs_y = dst_y; 131 132 uint rhs_x = VEC_SIZE * N0 * block_y; 133 uint rhs_y = 4 * block_id + block_x; 134 135 // Compute LHS/RHS/DST matrix address 136 lhs_offset_first_element_in_bytes += lhs_x * sizeof(DATA_TYPE) + lhs_y * lhs_stride_y + z * lhs_stride_z; 137 rhs_offset_first_element_in_bytes += rhs_x * sizeof(DATA_TYPE) + rhs_y * rhs_stride_y + z * rhs_stride_z; 138 dst_offset_first_element_in_bytes += dst_x * sizeof(OUT_DATA_TYPE) + dst_y * dst_stride_y + z * dst_stride_z; 139 140 TILE(ACC_DATA_TYPE, M0, N0, c); 141 LOOP_UNROLLING(int, i, 0, 1, M0, 142 { 143 c[i].v = 0; 144 }) 145 146 for(int k = 0; k <= K - MMUL_K0; k += MMUL_K0) 147 { 148 TILE(DATA_TYPE, M0, VEC_SIZE, a); 149 T_LOAD(DATA_TYPE, M0, VEC_SIZE, BUFFER, lhs, 0, 0, 1, lhs_stride_y, a); 150 151 TILE(DATA_TYPE, N0, VEC_SIZE, b); 152 T_LOAD(DATA_TYPE, N0, VEC_SIZE, BUFFER, rhs, 0, 0, 1, VEC_SIZE, b); 153 154 LOOP_UNROLLING(int, m0, 0, 1, M0, 155 { 156 LOOP_UNROLLING(int, n0, 0, 1, N0, 157 { 158 VEC_TYPE vec_a = (VEC_TYPE)(a[m0].s[0], a[m0].s[1], a[m0].s[2], a[m0].s[3]); 159 VEC_TYPE vec_b = (VEC_TYPE)(b[n0].s[0], b[n0].s[1], b[n0].s[2], b[n0].s[3]); 160 c[m0].s[n0] = arm_matrix_multiply(vec_a, vec_b, c[m0].s[n0]); 161 }) 162 }) 163 164 lhs_offset_first_element_in_bytes += MMUL_K0 * sizeof(DATA_TYPE); 165 rhs_offset_first_element_in_bytes += MMUL_K0 * N0 * sizeof(DATA_TYPE); 166 } 167 168 if(block_x * N0 + block_id * MMUL_N0 * N0 >= N) 169 { 170 return; 171 } 172 173 if(block_y * M0 + y0 * M0 * MMUL_M0 >= M) 174 { 175 return; 176 } 177 178#if defined(FUSED_OUTPUT_STAGE_FIXED_POINT) 179 180 TILE(int, M0, N0, offset_s32); 181 LOOP_UNROLLING(int, i, 0, 1, M0, 182 { 183 offset_s32[i].v = (VEC_DATA_TYPE(int, N0))K_OFFSET; 184 }) 185 186#if defined(A_OFFSET) 187 188 TILE(int, 1, N0, a_offset_s32); 189 190 T_LOAD(int, 1, N0, BUFFER, sum_col, dst_x, z, 1, sum_col_stride_z, a_offset_s32); 191 192 a_offset_s32[0].v *= A_OFFSET; 193 194 T_ELTWISE_BROADCAST_ADD_X(int, M0, N0, offset_s32, a_offset_s32, offset_s32); 195#endif // defined(A_OFFSET) 196 197#if defined(B_OFFSET) 198 199 TILE(int, M0, 1, b_offset_s32); 200 201 T_LOAD(int, M0, 1, BUFFER, sum_row, dst_y, z * M, 1, 4, b_offset_s32); 202 203 LOOP_UNROLLING(int, m0, 0, 1, M0, 204 { 205 offset_s32[m0].v += b_offset_s32[m0].v *B_OFFSET; 206 }) 207 208#endif // defined(B_OFFSET) 209 210#if defined(ADD_BIAS) 211#if defined(BROADCAST_BIAS) 212 bia_offset_first_element_in_bytes += dst_x * sizeof(ACC_DATA_TYPE) + z * bia_stride_y; 213 214 TILE(int, M0, N0, bias); 215 216 T_LOAD(int, M0, N0, BUFFER, bia, dst_x, dst_y, 1, 1, bias); 217 218 T_ADD(ACC_DATA_TYPE, M0, N0, offset_s32, bias, offset_s32); 219 220#else // defined(BROADCAST_BIAS) 221 bia_offset_first_element_in_bytes += dst_x * sizeof(ACC_DATA_TYPE); 222 223 TILE(int, 1, N0, bias); 224 225 if(dst_x + N0 <= N || N0_LEFTOVER == 0) 226 { 227 bias[0].v = VLOAD(N0)(0, (ACC_DATA_TYPE *)(bia_ptr + bia_offset_first_element_in_bytes)); 228 } 229 else 230 { 231 VLOAD_PARTIAL(N0, N0_LEFTOVER) 232 (bias[0].v, 0, (ACC_DATA_TYPE *)(bia_ptr + bia_offset_first_element_in_bytes)); 233 } 234 235 T_ELTWISE_BROADCAST_ADD_X(int, M0, N0, offset_s32, bias, offset_s32); 236 237#endif // defined(BROADCAST_BIAS) 238#endif // defined(ADD_BIAS) 239 240 T_ADD(ACC_DATA_TYPE, M0, N0, c, offset_s32, c); 241 TILE(OUT_DATA_TYPE, M0, N0, c_lp); 242 T_QUANTIZE8(ACC_DATA_TYPE, OUT_DATA_TYPE, PER_TENSOR, M0, N0, RESULT_OFFSET, RESULT_SHIFT, RESULT_MULTIPLIER, c, 0, 0, c_lp); 243 244#if defined(MIN_BOUND) 245 LOOP_UNROLLING(int, i, 0, 1, M0, 246 { 247 c_lp[i].v = max(c_lp[i].v, (VEC_DATA_TYPE(OUT_DATA_TYPE, N0))MIN_BOUND); 248 }) 249#endif // defined(MIN_BOUND) 250#if defined(MAX_BOUND) 251 LOOP_UNROLLING(int, i, 0, 1, M0, 252 { 253 c_lp[i].v = min(c_lp[i].v, (VEC_DATA_TYPE(OUT_DATA_TYPE, N0))MAX_BOUND); 254 }) 255#endif // defined(MAX_BOUND) 256 257 T_ACTIVATION(DATA_TYPE, M0, N0, ACTIVATION_TYPE, A_VAL, B_VAL, c, c); 258 259 if(dst_x + N0 <= N || N0_LEFTOVER == 0) 260 { 261 LOOP_UNROLLING(int, m0, 0, 1, M0, 262 { 263 if(dst_y + m0 < M || M0_LEFTOVER == 0) 264 { 265 VSTORE(N0) 266 (c_lp[m0].v, 0, (__global OUT_DATA_TYPE *)(dst_ptr + dst_offset_first_element_in_bytes + m0 * dst_stride_y)); 267 } 268 }) 269 } 270 else 271 { 272 LOOP_UNROLLING(int, m0, 0, 1, M0, 273 { 274 if(dst_y + m0 < M || M0_LEFTOVER == 0) 275 { 276 VSTORE_PARTIAL(N0, N0_LEFTOVER) 277 (c_lp[m0].v, 0, (__global OUT_DATA_TYPE *)(dst_ptr + dst_offset_first_element_in_bytes + m0 * dst_stride_y)); 278 } 279 }) 280 } 281 282#else // FUSED_OUTPUT_STAGE_FIXED_POINT 283 // Store 284 if(dst_x + N0 <= N || N0_LEFTOVER == 0) 285 { 286 LOOP_UNROLLING(int, m0, 0, 1, M0, 287 { 288 if(dst_y + m0 < M || M0_LEFTOVER == 0) 289 { 290 VSTORE(N0) 291 (c[m0].v, 0, (__global OUT_DATA_TYPE *)(dst_ptr + dst_offset_first_element_in_bytes + m0 * dst_stride_y)); 292 } 293 }) 294 } 295 else 296 { 297 LOOP_UNROLLING(int, m0, 0, 1, M0, 298 { 299 if(dst_y + m0 < M || M0_LEFTOVER == 0) 300 { 301 VSTORE_PARTIAL(N0, N0_LEFTOVER) 302 (c[m0].v, 0, (__global OUT_DATA_TYPE *)(dst_ptr + dst_offset_first_element_in_bytes + m0 * dst_stride_y)); 303 } 304 }) 305 } 306#endif // FUSED_OUTPUT_STAGE_FIXED_POINT 307} 308 309#endif // defined(GEMMLOWP_MM_RESHAPED_ONLY_RHS_MMUL) 310