diff options
Diffstat (limited to 'src/core/CL')
9 files changed, 6 insertions, 3893 deletions
diff --git a/src/core/CL/CLUtils.cpp b/src/core/CL/CLUtils.cpp index 03f78697bc..7e56a3ba18 100644 --- a/src/core/CL/CLUtils.cpp +++ b/src/core/CL/CLUtils.cpp @@ -23,16 +23,14 @@ */ #include "src/core/CL/CLUtils.h" -#include "arm_compute/core/utils/ActivationFunctionUtils.h" #include "arm_compute/core/CL/CLCompileContext.h" #include "arm_compute/core/CL/CLKernelLibrary.h" #include "arm_compute/core/CL/ICLTensor.h" #include "arm_compute/core/Validate.h" +#include "arm_compute/core/utils/ActivationFunctionUtils.h" #include "arm_compute/core/utils/StringUtils.h" #include "support/StringSupport.h" -#include "src/core/experimental/PostOpUtils.h" - namespace arm_compute { cl::Image2D create_image2d_from_tensor(const ICLTensor *tensor, CLImage2DType image_type) @@ -40,7 +38,7 @@ cl::Image2D create_image2d_from_tensor(const ICLTensor *tensor, CLImage2DType im ARM_COMPUTE_ERROR_ON_NULLPTR(tensor); const cl::Context &ctx = CLKernelLibrary::get().context(); - const cl::Buffer &buffer = tensor->cl_buffer(); + const cl::Buffer &buffer = tensor->cl_buffer(); const ITensorInfo *info = tensor->info(); ARM_COMPUTE_ERROR_ON_MSG(info->lock_paddings(), "Tensor paddings must not be locked to allow extending paddings to satisfy cl_image pitch alignment requirement"); @@ -113,112 +111,4 @@ cl::Image2D create_image2d_from_buffer(const cl::Context &ctx, const cl::Buffer return cl::Image2D(cl_image); } - -namespace experimental -{ -PostOpCLKernelUtils::PostOpCLKernelUtils(const Config &supported_config) - : _supported_config(supported_config) -{ - ARM_COMPUTE_ERROR_ON_MSG(supported_config.empty(), "Empty PostOp CL kernel support configuration is not allowed"); - for(auto it = _supported_config.begin(); it != _supported_config.end(); ++it) - { - auto post_op_sequence = it->first; - auto post_op_slots = std::get<1>(it->second); - ARM_COMPUTE_ERROR_ON_MSG(post_op_sequence.size() != post_op_slots.size(), "The number of PostOps must be the same as that of the assigned slots"); - } -} - -bool PostOpCLKernelUtils::are_post_op_shapes_compliant(const ITensorInfo *dst, const experimental::PostOpList<ITensorInfo *> &post_ops) -{ - for(const auto &op : post_ops.get_list()) - { - for(const auto &tensor : op->arguments()) - { - const TensorShape &out_shape = TensorShape::broadcast_shape(dst->tensor_shape(), (*tensor)->tensor_shape()); - // All post ops must be elementwise and must not alter the shape of the original dst tensor after broadcasting - if(detail::have_different_dimensions(out_shape, dst->tensor_shape(), 0)) - { - return false; - } - // NOTE: Kernel limitation: currently only the following broadcasting types are supported: - // 1. Post op arg is scalar, broadcast in both first and second dims - // 2. Post op arg is of shape: second dim=1, first dim=N, broadcast only in second dim - // This means this case: Post op arg is of shape: second dim=M, first dim=1, broadcast only in first dim, is NOT supported - if(dst->dimension(0) > 1 && dst->dimension(1) > 1 && (*tensor)->dimension(0) == 1 && (*tensor)->dimension(1) > 1) - { - return false; - } - } - } - return true; -} - -bool PostOpCLKernelUtils::is_post_op_sequence_supported(const PostOpList<ITensorInfo *> &post_ops) const -{ - if(post_ops.size() == 0) - { - return true; // Always support cases where no post op is specified - } - const auto post_op_sequence = get_post_op_sequence(post_ops); - - return _supported_config.find(post_op_sequence) != _supported_config.end(); -} - -void PostOpCLKernelUtils::set_post_ops_cl_build_options(CLBuildOptions &build_opts, const PostOpList<ITensorInfo *> &post_ops) const -{ - const auto post_op_sequence = get_post_op_sequence(post_ops); - const auto slots = std::get<1>(_supported_config.at(post_op_sequence)); - for(size_t post_op_id = 0; post_op_id < post_ops.size(); ++post_op_id) - { - const auto &post_op = post_ops.get_list().at(post_op_id); - const auto slot_prefix = "-DP" + support::cpp11::to_string(slots[post_op_id]); - if(post_op->type() == experimental::PostOpType::Activation) - { - const auto _post_op = utils::cast::polymorphic_downcast<const experimental::PostOpAct<ITensorInfo *> *>(post_op.get()); - const auto act_type = slot_prefix + "_ACTIVATION_TYPE=" + lower_string(string_from_activation_func(_post_op->_act_info.activation())); - const auto act_a_val = slot_prefix + "_ACTIVATION_A_VAL=" + float_to_string_with_full_precision(_post_op->_act_info.a()); - const auto act_b_val = slot_prefix + "_ACTIVATION_B_VAL=" + float_to_string_with_full_precision(_post_op->_act_info.b()); - build_opts.add_option(act_type); - build_opts.add_option(act_a_val); - build_opts.add_option(act_b_val); - } - else if(post_op->type() == experimental::PostOpType::Eltwise_Add) - { - size_t arg_id = 1; - const auto eltwise_op = slot_prefix + "_ELTWISE_OP=ADD" + "_X_POS_" + support::cpp11::to_string(post_op->prev_dst_pos()); - build_opts.add_option(eltwise_op); - for(const auto &tensor : post_op->arguments()) - { - const auto height = slot_prefix + "_ELTWISE_ARG" + support::cpp11::to_string(arg_id) + "_HEIGHT=" + support::cpp11::to_string((*tensor)->dimension(1)); - const auto width = slot_prefix + "_ELTWISE_ARG" + support::cpp11::to_string(arg_id) + "_WIDTH=" + support::cpp11::to_string((*tensor)->dimension(0)); - build_opts.add_option(height); - build_opts.add_option(width); - ++arg_id; - } - } - else if(post_op->type() == experimental::PostOpType::Eltwise_PRelu) - { - size_t arg_id = 1; - const auto eltwise_op = slot_prefix + "_ELTWISE_OP=PRELU" + "_X_POS_" + support::cpp11::to_string(post_op->prev_dst_pos()); - build_opts.add_option(eltwise_op); - for(const auto &tensor : post_op->arguments()) - { - const auto height = slot_prefix + "_ELTWISE_ARG" + support::cpp11::to_string(arg_id) + "_HEIGHT=" + support::cpp11::to_string((*tensor)->dimension(1)); - const auto width = slot_prefix + "_ELTWISE_ARG" + support::cpp11::to_string(arg_id) + "_WIDTH=" + support::cpp11::to_string((*tensor)->dimension(0)); - build_opts.add_option(height); - build_opts.add_option(width); - ++arg_id; - } - } - } -} - -void PostOpCLKernelUtils::set_post_ops_cl_kernel_name(std::string &kernel_name, const PostOpList<ITensorInfo *> &post_ops) const -{ - const auto post_op_sequence = get_post_op_sequence(post_ops); - const auto postfix = std::get<0>(_supported_config.at(post_op_sequence)); - kernel_name += postfix; -} -} // namespace experimental - } // namespace arm_compute diff --git a/src/core/CL/CLUtils.h b/src/core/CL/CLUtils.h index e3f12d4b53..f0e79bccfc 100644 --- a/src/core/CL/CLUtils.h +++ b/src/core/CL/CLUtils.h @@ -22,11 +22,10 @@ * SOFTWARE. */ -#ifndef ARM_COMPUTE_CL_CLUTILS_H -#define ARM_COMPUTE_CL_CLUTILS_H +#ifndef ACL_SRC_CORE_CL_CLUTILS_H +#define ACL_SRC_CORE_CL_CLUTILS_H #include "arm_compute/core/CL/OpenCL.h" -#include "arm_compute/core/experimental/IPostOp.h" #include <map> @@ -74,88 +73,6 @@ cl::Image2D create_image2d_from_tensor(const ICLTensor *tensor, CLImage2DType im * @return cl::Image2D object */ cl::Image2D create_image2d_from_buffer(const cl::Context &ctx, const cl::Buffer &buffer, const TensorShape &shape2d, DataType data_type, size_t image_row_pitch, CLImage2DType image_type); +} // namespace arm_compute -namespace experimental -{ -/** @name (EXPERIMENTAL_POST_OPS) - * @{ - */ - -/** Manage validation, building and configurations of PostOp CL kernels */ -class PostOpCLKernelUtils final -{ -public: - /** CL kernel name postfix for post ops */ - using NamePostfix = std::string; - /** CL kernels that supports post ops assign each post op to a 'slot', in accordance with the postfix - * For example, for a kernel with postfix '_act_prelu_eltwiseadd', there are 3 slots - * slot 1: (unary) activation, slot 2: pRelu, slot 3: elementwise addition - * - * Some kernels may allow some slots to be optional, to support multiple combinations of post op sequences. - * In such cases, we need to explicitly set up a mapping between each post op and the slots for that kernel. - * For example, suppose we have 2 kernels with postfixes: _eltwiseadd_prelu, _act_eltwiseadd_act_prelu, where the activations in the - * second kernel are optional. Say we want to support an eltwise addition, followed by a prelu (sequence { eltwiseadd, prelu }). - * Now we can choose which one of the 2 kernels to use, since they both support this post op sequence. - * We can either: - * 1. assign the elementwise to slot 1 and prelu to slot 2 of kernel 1 - * { { Eltwise_Add, PRelu } -> {"_eltwise_act", {1, 2} } } or - * 2. assign the elementwise to slot 2 and prelu to slot 4 of kernel 1 - * { { Eltwise_Add, PRelu } -> {"_act_eltwiseadd_act_prelu", {2, 4} } } - */ - using Slots = std::vector<unsigned int>; - using Config = std::map<PostOpTypeSequence, std::tuple<NamePostfix, Slots>>; - -public: - explicit PostOpCLKernelUtils(const Config &config); - - /** Check if post op argument tensor shapes are compliant - * All post ops must not alter the shape of the original dst tensor (even after broadcasting) - * - * @param[in] dst Dst tensor to apply the post ops to - * @param[in] post_ops Post ops - * - * @return true if shapes are compliant and false otherwise - */ - static bool are_post_op_shapes_compliant(const ITensorInfo *dst, const experimental::PostOpList<ITensorInfo *> &post_ops); - /** Check if the post op sequence is supported in the current configuration - * - * @param[in] post_ops Post ops - * - * @return true if the post op sequence is supported and false otherwise - */ - bool is_post_op_sequence_supported(const PostOpList<ITensorInfo *> &post_ops) const; - /** Helper function to set PostOp related build options - * @note Convention - * 1. Each post op "slot" is prefixed with "P<slot number>", followed by the usual parameters for that post op. - * E.g. If the first slot is an activation, we need to pass 3 definitions in this way: - * -P1_ACTIVATION_TYPE=... -P1_ACTIVATION_A_VAL=... -P1_ACTIVATION_B_VAL=... - * - * 2. For multi-ary post ops, to pass the position of the previous op's dest tensor, - * we append "_X_POS_<pos>" to the post op type. - * E.g. for a single post op add(dst, x), where dst is the result of the main op. - * In this case, the position of the previous op's dest is 0, so we pass - * -P1_ELTWISE_OP=ADD_X_POS_0 - * - * @param[out] built_opts OpenCL kernel build options - * @param[in] post_ops Post ops - * - */ - void set_post_ops_cl_build_options(CLBuildOptions &built_opts, const PostOpList<ITensorInfo *> &post_ops) const; - /** Helper function to set PostOp kernel name - * - * @param[out] kernel_name OpenCL kernel name - * @param[in] post_ops Post ops - * - */ - void set_post_ops_cl_kernel_name(std::string &kernel_name, const PostOpList<ITensorInfo *> &post_ops) const; - -private: - Config _supported_config{}; -}; -/** @} */ // end of group (EXPERIMENTAL_POST_OPS) - -} // namespace experimental - -} // arm_compute - -#endif /* ARM_COMPUTE_CL_CLUTILS_H */ +#endif // ACL_SRC_CORE_CL_CLUTILS_H diff --git a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/fp_post_ops_act_eltwise_op_act.h b/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/fp_post_ops_act_eltwise_op_act.h deleted file mode 100644 index 2c2d60ed13..0000000000 --- a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/fp_post_ops_act_eltwise_op_act.h +++ /dev/null @@ -1,103 +0,0 @@ -/* - * Copyright (c) 2021-2022 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ -#include "common/experimental/gemm_fused_post_ops/fp_mixed_precision_helpers.h" - -/** (EXPERIMENTAL_POST_OPS) Post Op expansions for the post op sequence: - * act (optional): POST_OP1_ACTIVATION_OPTIONAL - * eltwise_op : POST_OP2_ELTWISE_OP - * act (optional): POST_OP3_ACTIVATION_OPTIONAL - */ - -/** Post Op 1: Activation Block (Optional) - * @name POST_OP1_ACTIVATION_OPTIONAL - * Toggled by -DP1_ACTIVATION_TYPE - * params: same as those in @ref MIXED_PRECISION_ACTIVATION_BLOCK - * @{ - */ -#if defined(P1_ACTIVATION_TYPE) && defined(P1_ACTIVATION_A_VAL) && defined(P1_ACTIVATION_B_VAL) -#define POST_OP1_ACTIVATION_OPTIONAL(N, DATA_TYPE, DATA_TYPE_ACCUMULATOR, VEC_SIZE, BASENAME) \ - MIXED_PRECISION_ACTIVATION_BLOCK(N, P1_ACTIVATION_TYPE, DATA_TYPE, VEC_SIZE, BASENAME, P1_ACTIVATION_A_VAL, P1_ACTIVATION_B_VAL, DATA_TYPE_ACCUMULATOR); -#else // defined(P1_ACTIVATION_TYPE) && defined(P1_ACTIVATION_A_VAL) && defined(P1_ACTIVATION_B_VAL) -#define POST_OP1_ACTIVATION_OPTIONAL(N, DATA_TYPE, DATA_TYPE_ACCUMULATOR, VEC_SIZE, BASENAME) // noop -#endif // defined(P1_ACTIVATION_TYPE) && defined(P1_ACTIVATION_A_VAL) && defined(P1_ACTIVATION_B_VAL) -/** @} */ // end of group POST_OP1_ACTIVATION_OPTIONAL - -/** Post Op 2: Eltwise Op Block - * Handles both broadcasting and non-broadcasting cases - * @name POST_OP2_ELTWISE_OP - * - * @param[in] P2_ELTWISE_ARG1_HEIGHT Height (number of rows) of the @ref ELTWISE_OPERAND_NAME tensor - * @param[in] P2_ELTWISE_ARG1_WIDTH Width (number of columns) of the @ref ELTWISE_OPERAND_NAME tensor - * @param[in] OP The elementwise post op - * @param[in] M0 The number of consecutive rows - * @param[in] N0 The number of consecutive columns - * @param[in] BASENAME The basename of the result variables - * @param[in] ELTWISE_OPERAND_NAME The basename of the other operand variables - * @param[in] ELTWISE_OPERAND_ROW The starting row of the other operand variables. Required as different boundary handling strategies are used by different kernels - * E.g. reshaped_only_rhs and native kernels shifts rows (by using COMPUTE_M0_START_ROW) to handle boundary rows, - * whereas reshaped kernels do not shift rows - * @param[in] DATA_TYPE Data type of the result variables - * @param[in] DATA_TYPE_ACCUMULATR Higher-precision accumulator data type in case of mixed-precision op - * @param[in] ZERO Zero vector for z offset - * @param[in] PARTIAL_LOAD_M0 The partial size in y, for partial blocks. Supported: [0, @p M0) - * @param[in] PARTIAL_LOAD_N0 The partial size in x, for partial blocks. Supported: [0, @p N0) - * @param[in] PARTIAL_COND_Y Condition on the y axis to perform the partial load Y. True to use PARTIAL_LOAD_M0 rather than M0. - * @param[in] PARTIAL_COND_X Condition on the x axis to perform the partial load X. True to use PARTIAL_LOAD_N0 rather than N0. - * @{ - */ -#if defined(P2_ELTWISE_ARG1_HEIGHT) && defined(P2_ELTWISE_ARG1_WIDTH) -#if P2_ELTWISE_ARG1_HEIGHT == 1 -#if P2_ELTWISE_ARG1_WIDTH == 1 // Case 1: Broadcasting in both X and Y; op2 arg tile shape[YxX] == [1x1] -#define POST_OP2_ELTWISE_OP(OP, M0, N0, BASENAME, ELTWISE_OPERAND_NAME, ELTWISE_OPERAND_ROW, DATA_TYPE, DATA_TYPE_ACCUMULATOR, ZERO, PARTIAL_LOAD_M0, PARTIAL_LOAD_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \ - __global uchar *ELTWISE_OPERAND_NAME##_addr = ELTWISE_OPERAND_NAME##_ptr + ELTWISE_OPERAND_NAME##_offset_first_element_in_bytes + get_global_id(2) * ELTWISE_OPERAND_NAME##_stride_z; \ - VEC_DATA_TYPE(DATA_TYPE, 1) \ - ELTWISE_OPERAND_NAME##0 = VLOAD(1)(0, (__global DATA_TYPE *)ELTWISE_OPERAND_NAME##_addr); \ - MIXED_PRECISION_ELTWISE_OP_BLOCK_BROADCAST(OP, M0, 1, BASENAME, ELTWISE_OPERAND_NAME, DATA_TYPE_ACCUMULATOR, ELTWISE_OPERAND_NAME##_hp); -#else // P2_ELTWISE_ARG1_WIDTH == 1; Case 2: Broadcasting in only Y; op2 arg tile shape[YxX] == [1xN0] -#define POST_OP2_ELTWISE_OP(OP, M0, N0, BASENAME, ELTWISE_OPERAND_NAME, ELTWISE_OPERAND_ROW, DATA_TYPE, DATA_TYPE_ACCUMULATOR, ZERO, PARTIAL_LOAD_M0, PARTIAL_LOAD_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \ - __global uchar *ELTWISE_OPERAND_NAME##_addr = ELTWISE_OPERAND_NAME##_ptr + ELTWISE_OPERAND_NAME##_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)) + get_global_id(2) * ELTWISE_OPERAND_NAME##_stride_z; \ - LOAD_BLOCK_BOUNDARY_AWARE(1, N0, DATA_TYPE, ELTWISE_OPERAND_NAME, ELTWISE_OPERAND_NAME##_addr, 0, ELTWISE_OPERAND_NAME##_stride_y, ZERO, 1, PARTIAL_LOAD_N0, false, PARTIAL_COND_X); \ - MIXED_PRECISION_ELTWISE_OP_BLOCK_BROADCAST(OP, M0, N0, BASENAME, ELTWISE_OPERAND_NAME, DATA_TYPE_ACCUMULATOR, ELTWISE_OPERAND_NAME##_hp); -#endif // P2_ELTWISE_ARG1_WIDTH == 1 -#else // P2_ELTWISE_ARG1_HEIGHT == 1; Case 3: No broadcasting; op2 arg tile shape[YxX] == [M0xN0] -#define POST_OP2_ELTWISE_OP(OP, M0, N0, BASENAME, ELTWISE_OPERAND_NAME, ELTWISE_OPERAND_ROW, DATA_TYPE, DATA_TYPE_ACCUMULATOR, ZERO, PARTIAL_LOAD_M0, PARTIAL_LOAD_N0, PARTIAL_COND_Y, PARTIAL_COND_X) \ - __global uchar *ELTWISE_OPERAND_NAME##_addr = ELTWISE_OPERAND_NAME##_ptr + ELTWISE_OPERAND_NAME##_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)) + (ELTWISE_OPERAND_ROW * ELTWISE_OPERAND_NAME##_stride_y) + get_global_id(2) * ELTWISE_OPERAND_NAME##_stride_z; \ - LOAD_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, ELTWISE_OPERAND_NAME, ELTWISE_OPERAND_NAME##_addr, 0, ELTWISE_OPERAND_NAME##_stride_y, ZERO, PARTIAL_LOAD_M0, PARTIAL_LOAD_N0, PARTIAL_COND_Y, PARTIAL_COND_X); \ - MIXED_PRECISION_ELTWISE_OP_BLOCK(OP, M0, N0, BASENAME, ELTWISE_OPERAND_NAME, DATA_TYPE_ACCUMULATOR, ELTWISE_OPERAND_NAME##_hp); -#endif // P2_ELTWISE_ARG1_HEIGHT == 1 -#endif // defined(P2_ELTWISE_ARG1_HEIGHT) && defined(P2_ELTWISE_ARG1_WIDTH) -/** @} */ // end of group POST_OP2_ELTWISE_OP -/** Post Op 3: Activation Block (Optional) - * @name POST_OP3_ACTIVATION_OPTIONAL - * Toggled by -DP3_ACTIVATION_TYPE - * params: same as those in @ref MIXED_PRECISION_ACTIVATION_BLOCK - * @{ - */ -#if defined(P3_ACTIVATION_TYPE) && defined(P3_ACTIVATION_A_VAL) && defined(P3_ACTIVATION_B_VAL) -#define POST_OP3_ACTIVATION_OPTIONAL(N, DATA_TYPE, DATA_TYPE_ACCUMULATOR, VEC_SIZE, BASENAME) \ - MIXED_PRECISION_ACTIVATION_BLOCK(N, P3_ACTIVATION_TYPE, DATA_TYPE, VEC_SIZE, BASENAME, P3_ACTIVATION_A_VAL, P3_ACTIVATION_B_VAL, DATA_TYPE_ACCUMULATOR); -#else // defined(P3_ACTIVATION_TYPE) && defined(P3_ACTIVATION_A_VAL) && defined(P3_ACTIVATION_B_VAL) -#define POST_OP3_ACTIVATION_OPTIONAL(N, DATA_TYPE, DATA_TYPE_ACCUMULATOR, VEC_SIZE, BASENAME) // noop -#endif // defined(P3_ACTIVATION_TYPE) && defined(P3_ACTIVATION_A_VAL) && defined(P3_ACTIVATION_B_VAL) -/** @} */ // end of group POST_OP3_ACTIVATION_OPTIONAL diff --git a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_native.cl b/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_native.cl deleted file mode 100644 index 22ae098772..0000000000 --- a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_native.cl +++ /dev/null @@ -1,372 +0,0 @@ -/* - * Copyright (c) 2021-2022 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#include "common/experimental/gemm_fused_post_ops/act_eltwise_op_act/fp_post_ops_act_eltwise_op_act.h" -#include "common/experimental/gemm_fused_post_ops/fp_elementwise_op_helpers.h" -#include "common/experimental/gemm_fused_post_ops/fp_mixed_precision_helpers.h" - -#include "gemm_helpers.h" -#include "repeat.h" - -/** (EXPERIMENTAL_POST_OPS) gemm_mm_native kernel */ -#if defined(M0) && defined(N0) && defined(K0) && defined(DATA_TYPE) && defined(PARTIAL_STORE_M0) && defined(PARTIAL_STORE_N0) -#if defined(P2_ELTWISE_OP) && defined(P2_ELTWISE_ARG1_HEIGHT) && defined(P2_ELTWISE_ARG1_WIDTH) - -#define VFMA(a, b, c) \ - ({ \ - c = fma(a, b, c); \ - }) - -#if M0 == 1 -#define RHS_VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - }) -#elif M0 == 2 // M0 == 2 -#define RHS_VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - }) -#elif M0 == 3 // M0 == 3 -#define RHS_VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##2).s##i), b, (c##2)); \ - }) -#elif M0 == 4 // M0 == 4 -#define RHS_VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##2).s##i), b, (c##2)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##3).s##i), b, (c##3)); \ - }) -#elif M0 == 5 // M0 == 5 -#define RHS_VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##2).s##i), b, (c##2)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##3).s##i), b, (c##3)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##4).s##i), b, (c##4)); \ - }) -#elif M0 == 6 // M0 == 6 -#define RHS_VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##2).s##i), b, (c##2)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##3).s##i), b, (c##3)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##4).s##i), b, (c##4)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##5).s##i), b, (c##5)); \ - }) -#elif M0 == 7 // M0 == 7 -#define RHS_VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##2).s##i), b, (c##2)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##3).s##i), b, (c##3)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##4).s##i), b, (c##4)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##5).s##i), b, (c##5)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##6).s##i), b, (c##6)); \ - }) -#elif M0 == 8 // M0 == 8 -#define RHS_VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##2).s##i), b, (c##2)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##3).s##i), b, (c##3)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##4).s##i), b, (c##4)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##5).s##i), b, (c##5)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##6).s##i), b, (c##6)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##7).s##i), b, (c##7)); \ - }) -#else // M0 not supported -#error "M0 not supported" -#endif // M0 not supported - -#if defined(GEMM_MM_NATIVE_POST_ACT_ELTWISE_OP_ACT) -/** This OpenCL kernel computes the matrix multiplication between 2 matrices plus 3 post ops: - * Post op 1: activation (optional) - * Post op 2: elementwise op - * Post op 3: activation (optional) - * - * @note (Optional) -DP1_ACTIVATION_TYPE, -DP1_ACTIVATION_A_VAL, -DP1_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * @note (Required) -DP2_ELTWISE_OP: The (binary) elementwise post op to perform - * @note (Required) -DP2_ELTWISE_ARG1_HEIGHT: The height (Y dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Required) -DP2_ELTWISE_ARG1_WIDTH: The width (X dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Optional) -DP3_ACTIVATION_TYPE, -DP3_ACTIVATION_A_VAL, -DP3_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * - * All parameters are similarly defined in kernel gemm_mm_native, with these additions: - * - * @param[in] eltwise_operand_ptr Pointer to the eltwise operand matrix. Supported data type: F16/F32 - * @param[in] eltwise_operand_stride_x Stride of the eltwise operand matrix in X dimension (in bytes) - * @param[in] eltwise_operand_step_x eltwise_operand_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_y Stride of the eltwise operand matrix in Y dimension (in bytes) - * @param[in] eltwise_operand_step_y eltwise_operand_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_z Stride of the eltwise operand tensor in Z dimension (in bytes) - */ -__kernel void gemm_mm_native_post_act_eltwise_op_act(IMAGE_DECLARATION(lhs), - IMAGE_DECLARATION(rhs), -#if defined(BETA) - IMAGE_DECLARATION(bias), -#endif // defined(BETA) - IMAGE_DECLARATION(dst), - // Post Op arguments - IMAGE_DECLARATION(eltwise_operand), - uint lhs_stride_z, - uint rhs_stride_z, -#if defined(BETA) - uint bias_stride_z, -#endif //defined(BETA) - uint dst_stride_z, - uint eltwise_operand_stride_z, - const int M, - const int N, - const int K -#if defined(REINTERPRET_INPUT_AS_3D) - , - uint lhs_cross_plane_pad -#endif // REINTERPRET_INPUT_AS_3D -#if defined(REINTERPRET_OUTPUT_AS_3D) - , - uint dst_cross_plane_pad -#endif // REINTERPRET_OUTPUT_AS_3D - ) -{ - // Block size -#define RHS_BLOCK_SIZE ((K0) * (N0)) - - // RHS offset and step X -#define RHS_OFFSET_X (RHS_BLOCK_SIZE) - - uint x = get_global_id(0); - uint y = get_global_id(1); - uint z = get_global_id(2); - -#if defined(DUMMY_WORK_ITEMS) - if((x * N0 >= N) || (y * M0 >= M)) - { - return; - } -#endif // defined(DUMMY_WORK_ITEMS) - - // Compute LHS matrix address - uint lhs_offset = lhs_offset_first_element_in_bytes + COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * (uint)lhs_stride_y; - - // Compute RHS matrix address - uint rhs_offset = rhs_offset_first_element_in_bytes + x * N0 * sizeof(DATA_TYPE); - -#if defined(MATRIX_B_DEPTH) - // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 - rhs_offset += (z % MATRIX_B_DEPTH) * rhs_stride_z; -#else // defined(MATRIX_B_DEPTH) - rhs_offset += z * rhs_stride_z; -#endif // defined(MATRIX_B_DEPTH) - - REPEAT_VAR_INIT_TO_CONST(M0, uint, zlhs, 0); - REPEAT_VAR_INIT_TO_CONST(16, uint, zero, 0); - -#if defined(REINTERPRET_INPUT_AS_3D) - // The plane (zlhs) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zlhs, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), HEIGHT_GEMM3D, DEPTH_GEMM3D, lhs_cross_plane_pad, lhs_stride_y); - - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply lhs_stride_z by DEPTH_GEMM3D - lhs_offset += z * lhs_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_INPUT_AS_3D) - - // Add offset for batched GEMM - lhs_offset += z * lhs_stride_z; - -#endif // defined(REINTERPRET_INPUT_AS_3D) - - // Initialize the accumulators - REPEAT_VAR_INIT_TO_CONST(M0, VEC_DATA_TYPE(DATA_TYPE, N0), c, 0); //VEC_DATA_TYPE(DATA_TYPE, N0) c0=0,c1=0,c2=0,... c(M0-1)=0; - - int i = 0; -#if K0 > 1 - for(; i <= (K - K0); i += K0) - { - // Supported cases (M0, K0): - // 1,2 - 1,3 - 1,4 - 1,8 - 1,16 - // 2,2 - 2,3 - 2,4 - 2,8 - 2,16 - // 3,2 - 3,3 - 3,4 - 3,8 - 3,16 - // 4,2 - 4,3 - 4,4 - 4,8 - 4,16 - // 5,2 - 5,3 - 5,4 - 5,8 - 5,16 - // 6,2 - 6,3 - 6,4 - 6,8 - 6,16 - // 7,2 - 7,3 - 7,4 - 7,8 - 7,16 - // 8,2 - 8,3 - 8,4 - 8,8 - 8,16 - // Load values from LHS matrix - LOAD_BLOCK(M0, K0, DATA_TYPE, a, lhs_ptr, lhs_offset, lhs_stride_y, zlhs); - - // Load values from RHS matrix - LOAD_BLOCK(K0, N0, DATA_TYPE, b, rhs_ptr, rhs_offset, rhs_stride_y, zero); - - RHS_VFMA_M0xN0(0, a, b0, c); - RHS_VFMA_M0xN0(1, a, b1, c); -#if K0 > 2 - RHS_VFMA_M0xN0(2, a, b2, c); -#endif // K0 > 2 -#if K0 > 3 - RHS_VFMA_M0xN0(3, a, b3, c); -#endif // K0 > 3 -#if K0 > 4 - RHS_VFMA_M0xN0(4, a, b4, c); - RHS_VFMA_M0xN0(5, a, b5, c); - RHS_VFMA_M0xN0(6, a, b6, c); - RHS_VFMA_M0xN0(7, a, b7, c); -#endif // K0 > 4 -#if K0 > 8 - RHS_VFMA_M0xN0(8, a, b8, c); - RHS_VFMA_M0xN0(9, a, b9, c); - RHS_VFMA_M0xN0(A, a, bA, c); - RHS_VFMA_M0xN0(B, a, bB, c); - RHS_VFMA_M0xN0(C, a, bC, c); - RHS_VFMA_M0xN0(D, a, bD, c); - RHS_VFMA_M0xN0(E, a, bE, c); - RHS_VFMA_M0xN0(F, a, bF, c); -#endif // K0 > 8 - - lhs_offset += K0 * sizeof(DATA_TYPE); - rhs_offset += K0 * rhs_stride_y; - } -#endif // K0 > 1 - // Left-over accumulations - for(; i < K; ++i) - { - // Load values from LHS matrix - VEC_DATA_TYPE(DATA_TYPE, 2) - a0 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 0 * lhs_stride_y + zlhs0)); -#if M0 > 1 - VEC_DATA_TYPE(DATA_TYPE, 2) - a1 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 1 * lhs_stride_y + zlhs1)); -#endif // M0 > 1 -#if M0 > 2 - VEC_DATA_TYPE(DATA_TYPE, 2) - a2 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 2 * lhs_stride_y + zlhs2)); -#endif // M0 > 2 -#if M0 > 3 - VEC_DATA_TYPE(DATA_TYPE, 2) - a3 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 3 * lhs_stride_y + zlhs3)); -#endif // M0 > 3 -#if M0 > 4 - VEC_DATA_TYPE(DATA_TYPE, 2) - a4 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 4 * lhs_stride_y + zlhs4)); -#endif // M0 > 4 -#if M0 > 5 - VEC_DATA_TYPE(DATA_TYPE, 2) - a5 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 5 * lhs_stride_y + zlhs5)); -#endif // M0 > 5 -#if M0 > 6 - VEC_DATA_TYPE(DATA_TYPE, 2) - a6 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 6 * lhs_stride_y + zlhs6)); -#endif // M0 > 6 -#if M0 > 7 - VEC_DATA_TYPE(DATA_TYPE, 2) - a7 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 7 * lhs_stride_y + zlhs7)); -#endif // M0 > 7 - - VEC_DATA_TYPE(DATA_TYPE, N0) - b = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 0 * rhs_stride_y)); - RHS_VFMA_M0xN0(0, a, b, c); - - lhs_offset += sizeof(DATA_TYPE); - rhs_offset += rhs_stride_y; - } - - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)) + (COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * dst_stride_y); - - REPEAT_VAR_INIT_TO_CONST(M0, uint, zout, 0); - -#if defined(REINTERPRET_OUTPUT_AS_3D) - // The plane (zout) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zout, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), HEIGHT_GEMM3D, DEPTH_GEMM3D, dst_cross_plane_pad, dst_stride_y); - - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply dst_stride_z by DEPTH_GEMM3D - dst_addr += z * dst_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_OUTPUT_AS_3D) - - // Add offset for batched GEMM - dst_addr += z * dst_stride_z; - -#endif // defined(REINTERPRET_OUTPUT_AS_3D) - - // Multiply by the weight of matrix-matrix product and store the result -#if defined(ALPHA) - SCALE_BLOCK(M0, DATA_TYPE, c, ALPHA); -#endif // defined(ALPHA) - - // Add beta*bias -#if defined(BETA) -#if defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)); - - LOAD_BLOCK(1, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero); - -#ifndef UNIT_BETA - SCALE_BLOCK(1, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias[broadcasted] - ADD_BLOCK_BROADCAST(M0, c, bias0); - -#else // defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)) + (COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * bias_stride_y) + z * bias_stride_z; - - LOAD_BLOCK(M0, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero); - -#ifndef UNIT_BETA - SCALE_BLOCK(M0, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias - ADD_BLOCK(M0, c, bias); - -#endif // defined(BROADCAST_BIAS) -#endif // defined(BETA) - - const bool cond_y = y == 0; - const bool cond_x = ((x + 1) * N0 >= N); - - // c = act(c) - POST_OP1_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - // c = c + eltwise_operand (mix-precision, broadcast, boundary aware) - POST_OP2_ELTWISE_OP(P2_ELTWISE_OP, M0, N0, c, eltwise_operand, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), DATA_TYPE, DATA_TYPE_ACCUMULATOR, zero, 1, PARTIAL_STORE_N0, false, cond_x); - // c = act(c) - POST_OP3_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - - // Store output block - STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, c, dst_addr, dst_stride_y, zout, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); -} -#endif // defined(GEMM_MM_NATIVE_POST_ACT_ELTWISE_OP_ACT) -#endif // defined(P2_ELTWISE_OP) && defined(P2_ELTWISE_ARG1_HEIGHT) && defined(P2_ELTWISE_ARG1_WIDTH) -#endif // defined(M0) && defined(N0) && defined(K0) && defined(DATA_TYPE) && defined(PARTIAL_STORE_M0) && defined(PARTIAL_STORE_N0) diff --git a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped.cl b/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped.cl deleted file mode 100644 index 89577e9ebd..0000000000 --- a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped.cl +++ /dev/null @@ -1,1424 +0,0 @@ -/* - * Copyright (c) 2021 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ -#include "fp_post_ops_act_eltwise_op_act.h" -#include "gemm_helpers.h" -#include "repeat.h" - -/** (EXPERIMENTAL_POST_OPS) gemm_mm_reshaped kernel */ - -#if defined(M0) && defined(N0) && defined(K0) && defined(V0) && defined(H0) && defined(DATA_TYPE) && defined(DATA_TYPE_ACCUMULATOR) -#if defined(P2_ELTWISE_OP) && defined(P2_ELTWISE_ARG1_HEIGHT) && defined(P2_ELTWISE_ARG1_WIDTH) - -#if defined(MIXED_PRECISION) -#if K0 == 2 -#define ARM_DOT_K0(a, b, c) \ - ({ \ - c += a.s0 * b.s0; \ - c += a.s1 * b.s1; \ - }) -#elif K0 == 3 // K0 == 3 -#define ARM_DOT_K0(a, b, c) \ - ({ \ - c += a.s0 * b.s0; \ - c += a.s1 * b.s1; \ - c += a.s2 * b.s2; \ - }) -#elif K0 == 4 // K0 == 4 -#define ARM_DOT_K0(a, b, c) \ - ({ \ - c += a.s0 * b.s0; \ - c += a.s1 * b.s1; \ - c += a.s2 * b.s2; \ - c += a.s3 * b.s3; \ - }) -#elif K0 == 8 // K0 == 8 -#define ARM_DOT_K0(a, b, c) \ - ({ \ - c += a.s0 * b.s0; \ - c += a.s1 * b.s1; \ - c += a.s2 * b.s2; \ - c += a.s3 * b.s3; \ - c += a.s4 * b.s4; \ - c += a.s5 * b.s5; \ - c += a.s6 * b.s6; \ - c += a.s7 * b.s7; \ - }) -#elif K0 == 16 // K0 == 16 -#define ARM_DOT_K0(a, b, c) \ - ({ \ - c += a.s0 * b.s0; \ - c += a.s1 * b.s1; \ - c += a.s2 * b.s2; \ - c += a.s3 * b.s3; \ - c += a.s4 * b.s4; \ - c += a.s5 * b.s5; \ - c += a.s6 * b.s6; \ - c += a.s7 * b.s7; \ - c += a.s8 * b.s8; \ - c += a.s9 * b.s9; \ - c += a.sA * b.sA; \ - c += a.sB * b.sB; \ - c += a.sC * b.sC; \ - c += a.sD * b.sD; \ - c += a.sE * b.sE; \ - c += a.sF * b.sF; \ - }) -#else // K0 not supported -#error "K0 value not supported" -#endif // K0 conditions -#else // defined(MIXED_PRECISION) -#if K0 == 2 -#define ARM_DOT_K0(a, b, c) \ - ({ \ - c = fma(a.s0, b.s0, c); \ - c = fma(a.s1, b.s1, c); \ - }) -#elif K0 == 3 // K0 == 3 -#define ARM_DOT_K0(a, b, c) \ - ({ \ - c = fma(a.s0, b.s0, c); \ - c = fma(a.s1, b.s1, c); \ - c = fma(a.s2, b.s2, c); \ - }) -#elif K0 == 4 // K0 == 4 -#define ARM_DOT_K0(a, b, c) \ - ({ \ - c = fma(a.s0, b.s0, c); \ - c = fma(a.s1, b.s1, c); \ - c = fma(a.s2, b.s2, c); \ - c = fma(a.s3, b.s3, c); \ - }) -#elif K0 == 8 // K0 == 8 -#define ARM_DOT_K0(a, b, c) \ - ({ \ - c = fma(a.s0, b.s0, c); \ - c = fma(a.s1, b.s1, c); \ - c = fma(a.s2, b.s2, c); \ - c = fma(a.s3, b.s3, c); \ - c = fma(a.s4, b.s4, c); \ - c = fma(a.s5, b.s5, c); \ - c = fma(a.s6, b.s6, c); \ - c = fma(a.s7, b.s7, c); \ - }) -#elif K0 == 16 // K0 == 16 -#define ARM_DOT_K0(a, b, c) \ - ({ \ - c = fma(a.s0, b.s0, c); \ - c = fma(a.s1, b.s1, c); \ - c = fma(a.s2, b.s2, c); \ - c = fma(a.s3, b.s3, c); \ - c = fma(a.s4, b.s4, c); \ - c = fma(a.s5, b.s5, c); \ - c = fma(a.s6, b.s6, c); \ - c = fma(a.s7, b.s7, c); \ - c = fma(a.s8, b.s8, c); \ - c = fma(a.s9, b.s9, c); \ - c = fma(a.sA, b.sA, c); \ - c = fma(a.sB, b.sB, c); \ - c = fma(a.sC, b.sC, c); \ - c = fma(a.sD, b.sD, c); \ - c = fma(a.sE, b.sE, c); \ - c = fma(a.sF, b.sF, c); \ - }) -#else // K0 not supported -#error "K0 value not supported" -#endif // K0 conditions -#endif // defined(MIXED_PRECISION) - -#if defined(ARM_DOT_K0XN0) -#undef ARM_DOT_K0XN0 -#endif // defined(ARM_DOT_K0XN0) - -#if N0 == 2 -#define ARM_DOT_K0XN0(a, b, c) \ - ({ \ - ARM_DOT_K0((a), (b##0), (c.s0)); \ - ARM_DOT_K0((a), (b##1), (c.s1)); \ - }) -#elif N0 == 3 // N0 == 3 -#define ARM_DOT_K0XN0(a, b, c) \ - ({ \ - ARM_DOT_K0((a), (b##0), (c.s0)); \ - ARM_DOT_K0((a), (b##1), (c.s1)); \ - ARM_DOT_K0((a), (b##2), (c.s2)); \ - }) -#elif N0 == 4 // N0 == 4 -#define ARM_DOT_K0XN0(a, b, c) \ - ({ \ - ARM_DOT_K0((a), (b##0), (c.s0)); \ - ARM_DOT_K0((a), (b##1), (c.s1)); \ - ARM_DOT_K0((a), (b##2), (c.s2)); \ - ARM_DOT_K0((a), (b##3), (c.s3)); \ - }) -#elif N0 == 8 // N0 == 8 -#define ARM_DOT_K0XN0(a, b, c) \ - ({ \ - ARM_DOT_K0((a), (b##0), (c.s0)); \ - ARM_DOT_K0((a), (b##1), (c.s1)); \ - ARM_DOT_K0((a), (b##2), (c.s2)); \ - ARM_DOT_K0((a), (b##3), (c.s3)); \ - ARM_DOT_K0((a), (b##4), (c.s4)); \ - ARM_DOT_K0((a), (b##5), (c.s5)); \ - ARM_DOT_K0((a), (b##6), (c.s6)); \ - ARM_DOT_K0((a), (b##7), (c.s7)); \ - }) -#elif N0 == 16 // N0 == 16 -#define ARM_DOT_K0XN0(a, b, c) \ - ({ \ - ARM_DOT_K0((a), (b##0), (c.s0)); \ - ARM_DOT_K0((a), (b##1), (c.s1)); \ - ARM_DOT_K0((a), (b##2), (c.s2)); \ - ARM_DOT_K0((a), (b##3), (c.s3)); \ - ARM_DOT_K0((a), (b##4), (c.s4)); \ - ARM_DOT_K0((a), (b##5), (c.s5)); \ - ARM_DOT_K0((a), (b##6), (c.s6)); \ - ARM_DOT_K0((a), (b##7), (c.s7)); \ - ARM_DOT_K0((a), (b##8), (c.s8)); \ - ARM_DOT_K0((a), (b##9), (c.s9)); \ - ARM_DOT_K0((a), (b##A), (c.sA)); \ - ARM_DOT_K0((a), (b##B), (c.sB)); \ - ARM_DOT_K0((a), (b##C), (c.sC)); \ - ARM_DOT_K0((a), (b##D), (c.sD)); \ - ARM_DOT_K0((a), (b##E), (c.sE)); \ - ARM_DOT_K0((a), (b##F), (c.sF)); \ - }) -#else // N0 not supported -#error "N0 value not supported" -#endif // N0 conditions - -#if defined(GEMM_MM_RESHAPED_LHS_NT_RHS_T_POST_ACT_ELTWISE_OP_ACT) -/** This OpenCL kernel computes the matrix multiplication between 2 matrices plus 3 post ops: - * Post op 1: activation (optional) - * Post op 2: elementwise op - * Post op 3: activation (optional) - * - * @note (Optional) -DP1_ACTIVATION_TYPE, -DP1_ACTIVATION_A_VAL, -DP1_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * @note (Required) -DP2_ELTWISE_OP: The (binary) elementwise post op to perform - * @note (Required) -DP2_ELTWISE_ARG1_HEIGHT: The height (Y dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Required) -DP2_ELTWISE_ARG1_WIDTH: The width (X dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Optional) -DP3_ACTIVATION_TYPE, -DP3_ACTIVATION_A_VAL, -DP3_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * - * All parameters are similarly defined in kernel gemm_mm_reshaped_lhs_nt_rhs_t, with these additions: - * - * @param[in] eltwise_operand_ptr Pointer to the eltwise operand matrix. Supported data type: F16/F32 - * @param[in] eltwise_operand_stride_x Stride of the eltwise operand matrix in X dimension (in bytes) - * @param[in] eltwise_operand_step_x eltwise_operand_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_y Stride of the eltwise operand matrix in Y dimension (in bytes) - * @param[in] eltwise_operand_step_y eltwise_operand_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_z Stride of the eltwise operand tensor in Z dimension (in bytes) - */ -__kernel void gemm_mm_reshaped_lhs_nt_rhs_t_post_act_eltwise_op_act(IMAGE_DECLARATION(lhs), - IMAGE_DECLARATION(rhs), -#if defined(BETA) - IMAGE_DECLARATION(bias), -#endif // defined(BETA) - IMAGE_DECLARATION(dst), - // Post Op arguments - IMAGE_DECLARATION(eltwise_operand), - uint lhs_stride_z, - uint rhs_stride_z, -#if defined(BETA) - uint bias_stride_z, -#endif //defined(BETA) - uint dst_stride_z, - uint eltwise_operand_stride_z -#if defined(REINTERPRET_OUTPUT_AS_3D) - , - uint dst_cross_plane_pad -#endif // REINTERPRET_OUTPUT_AS_3D - , - const int M, - const int N, - const int K) -{ - // Block size -#define LHS_BLOCK_SIZE ((K0) * (M0)) - -#if defined(LHS_INTERLEAVE) -#define LHS_OFFSET_X (K0) -#define LHS_STEP_X ((K0) * (V0)) -#define LHS_STEP_LOOP (1) -#else // defined(INTERLEAVE) -#define LHS_OFFSET_X (LHS_BLOCK_SIZE) -#define LHS_STEP_X (K0) -#define LHS_STEP_LOOP (V0) -#endif // defined(INTERLEAVE) - - // Block size -#define RHS_BLOCK_SIZE ((K0) * (N0)) - - // RHS offset and step X -#if defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (K0) -#define RHS_STEP_X ((K0) * (H0)) -#define RHS_STEP_LOOP (1) -#else // defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (RHS_BLOCK_SIZE) -#define RHS_STEP_X (K0) -#define RHS_STEP_LOOP (H0) -#endif // defined(RHS_INTERLEAVE) - -#if defined(DUMMY_WORK_ITEMS) - if((get_global_id(0) * N0 >= N) || (get_global_id(1) * M0 >= M)) - { - return; - } -#endif // defined(DUMMY_WORK_ITEMS) - - // Compute LHS matrix address - __global uchar *lhs_addr = lhs_ptr + lhs_offset_first_element_in_bytes + (get_global_id(1) % V0) * (uint)LHS_OFFSET_X * sizeof(DATA_TYPE) + (get_global_id(1) / V0) * (uint)lhs_stride_y + - (get_global_id(2) * lhs_stride_z); - - // Compute RHS matrix address - __global uchar *rhs_addr = rhs_ptr + rhs_offset_first_element_in_bytes + (get_global_id(0) % H0) * (uint)RHS_OFFSET_X * sizeof(DATA_TYPE) + (get_global_id(0) / (uint)H0) * rhs_stride_y; - -#if defined(MATRIX_B_DEPTH) - // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 - rhs_addr += (get_global_id(2) % MATRIX_B_DEPTH) * rhs_stride_z; -#else // defined(MATRIX_B_DEPTH) - rhs_addr += get_global_id(2) * rhs_stride_z; -#endif // defined(MATRIX_B_DEPTH) - - // Initialize the accumulators - REPEAT_VAR_INIT_TO_CONST(M0, VEC_DATA_TYPE(DATA_TYPE_ACCUMULATOR, N0), c, 0); - - REPEAT_VAR_INIT_TO_CONST(M0, uint, zlhs, 0); //uint zlhs0=0,zlhs1=0,zlhs2=0,... zlhs7=0; - REPEAT_VAR_INIT_TO_CONST(16, uint, zero, 0); - - for(int i = 0; i < K; i += K0) - { - // Supported cases (M0, K0): - // 1,2 - 1,3 - 1,4 - 1,8 - 1,16 - // 2,2 - 2,3 - 2,4 - 2,8 - 2,16 - // 3,2 - 3,3 - 3,4 - 3,8 - 3,16 - // 4,2 - 4,3 - 4,4 - 4,8 - 4,16 - // 5,2 - 5,3 - 5,4 - 5,8 - 5,16 - // 6,2 - 6,3 - 6,4 - 6,8 - 6,16 - // 7,2 - 7,3 - 7,4 - 7,8 - 7,16 - // 8,2 - 8,3 - 8,4 - 8,8 - 8,16 - // Load values from LHS matrix - LOAD_BLOCK(M0, K0, DATA_TYPE, a, lhs_addr, 0, LHS_STEP_X * sizeof(DATA_TYPE), zlhs); - - // Load values from RHS matrix - LOAD_BLOCK(N0, K0, DATA_TYPE, b, rhs_addr, 0, RHS_STEP_X * sizeof(DATA_TYPE), zero); - - // Accumulate - ARM_DOT_K0XN0(a0, b, c0); -#if M0 > 1 - ARM_DOT_K0XN0(a1, b, c1); -#endif // M0 > 1 -#if M0 > 2 - ARM_DOT_K0XN0(a2, b, c2); -#endif // M0 > 2 -#if M0 > 3 - ARM_DOT_K0XN0(a3, b, c3); -#endif // M0 > 3 -#if M0 > 4 - ARM_DOT_K0XN0(a4, b, c4); -#endif // M0 > 4 -#if M0 > 5 - ARM_DOT_K0XN0(a5, b, c5); -#endif // M0 > 5 -#if M0 > 6 - ARM_DOT_K0XN0(a6, b, c6); -#endif // M0 > 6 -#if M0 > 7 - ARM_DOT_K0XN0(a7, b, c7); -#endif // M0 > 7 - - lhs_addr += (M0 * LHS_STEP_X * LHS_STEP_LOOP) * sizeof(DATA_TYPE); - rhs_addr += (N0 * RHS_STEP_X * RHS_STEP_LOOP) * sizeof(DATA_TYPE); - } - - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)) + (get_global_id(1) * (uint)M0 * dst_stride_y); - - REPEAT_VAR_INIT_TO_CONST(M0, uint, zout, 0); - - // Boundary conditions: detect if current block is at the "bottom" or "right" boundary - const bool cond_y = ((get_global_id(1) + 1) * M0 >= M); - const bool cond_x = ((get_global_id(0) + 1) * N0 >= N); - -#if defined(REINTERPRET_OUTPUT_AS_3D) - - // The plane (zin) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zout, get_global_id(1) * (uint)M0, HEIGHT_GEMM3D, DEPTH_GEMM3D, dst_cross_plane_pad, dst_stride_y); - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply dst_stride_z by DEPTH_GEMM3D - dst_addr += get_global_id(2) * dst_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_OUTPUT_AS_3D) - - // Add offset for batched GEMM - dst_addr += get_global_id(2) * dst_stride_z; - -#endif // defined(REINTERPRET_OUTPUT_AS_3D) - - // Multiply by the weight of matrix-matrix product and store the result -#if defined(ALPHA) - SCALE_BLOCK(M0, DATA_TYPE, c, ALPHA); -#endif // defined(ALPHA) - - // Add beta*bias -#if defined(BETA) -#if defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)); - - LOAD_BLOCK_BOUNDARY_AWARE(1, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, 1, PARTIAL_STORE_N0, false, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(1, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias[broadcasted] - MIXED_PRECISION_ELTWISE_OP_BLOCK_BROADCAST(ADD, M0, N0, c, bias, DATA_TYPE_ACCUMULATOR, bias_hp); - -#else // defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)) + (get_global_id(1) * (uint)M0 * bias_stride_y) + get_global_id( - 2) * bias_stride_z; - - LOAD_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(M0, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias - MIXED_PRECISION_ELTWISE_OP_BLOCK(ADD, M0, N0, c, bias, DATA_TYPE_ACCUMULATOR, bias_hp); - -#endif // defined(BROADCAST_BIAS) -#endif // defined(BETA) - - // c = act(c) - POST_OP1_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - // c = c + eltwise_operand (mix-precision, broadcast, boundary aware) - POST_OP2_ELTWISE_OP(P2_ELTWISE_OP, M0, N0, c, eltwise_operand, get_global_id(1) * (uint)M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, zero, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - // c = act(c) - POST_OP3_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - - // Store output block - MIXED_PRECISION_STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, c, dst_addr, dst_stride_y, zout, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x, c_lp); - -#undef LHS_BLOCK_SIZE -#undef LHS_OFFSET_X -#undef LHS_STEP_X -#undef RHS_BLOCK_SIZE -#undef RHS_OFFSET_X -#undef RHS_STEP_X -#undef LHS_STEP_LOOP -#undef RHS_STEP_LOOP -} -#endif // defined(GEMM_MM_RESHAPED_LHS_NT_RHS_T_POST_ACT_ELTWISE_OP_ACT) - -#if defined(OPENCL_IMAGE_SUPPORT) && defined(GEMM_MM_RESHAPED_LHS_NT_RHS_T_TEXTURE_POST_ACT_ELTWISE_OP_ACT) -/** This OpenCL kernel computes the matrix multiplication between 2 matrices plus 3 post ops. The RHS matrix is stored in OpenCL image object. - * Post op 1: activation (optional) - * Post op 2: elementwise op - * Post op 3: activation (optional) - * - * @note (Optional) -DP1_ACTIVATION_TYPE, -DP1_ACTIVATION_A_VAL, -DP1_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * @note (Required) -DP2_ELTWISE_OP: The (binary) elementwise post op to perform - * @note (Required) -DP2_ELTWISE_ARG1_HEIGHT: The height (Y dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Required) -DP2_ELTWISE_ARG1_WIDTH: The width (X dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Optional) -DP3_ACTIVATION_TYPE, -DP3_ACTIVATION_A_VAL, -DP3_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * - * All parameters are similarly defined in kernel gemm_mm_reshaped_lhs_nt_rhs_t_texture, with these additions: - * - * @param[in] eltwise_operand_ptr Pointer to the eltwise operand matrix. Supported data type: F16/F32 - * @param[in] eltwise_operand_stride_x Stride of the eltwise operand matrix in X dimension (in bytes) - * @param[in] eltwise_operand_step_x eltwise_operand_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_y Stride of the eltwise operand matrix in Y dimension (in bytes) - * @param[in] eltwise_operand_step_y eltwise_operand_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_z Stride of the eltwise operand tensor in Z dimension (in bytes) - */ -__kernel void gemm_mm_reshaped_lhs_nt_rhs_t_texture_post_act_eltwise_op_act(IMAGE_DECLARATION(lhs), - __read_only image2d_t rhs_img, -#if defined(BETA) - IMAGE_DECLARATION(bias), -#endif // defined(BETA) - IMAGE_DECLARATION(dst), - // Post Op arguments - IMAGE_DECLARATION(eltwise_operand), - uint lhs_stride_z, - uint rhs_stride_z, -#if defined(BETA) - uint bias_stride_z, -#endif //defined(BETA) - uint dst_stride_z, - uint eltwise_operand_stride_z -#if defined(REINTERPRET_OUTPUT_AS_3D) - , - uint dst_cross_plane_pad -#endif // REINTERPRET_OUTPUT_AS_3D - , - const int M, - const int N, - const int K) -{ - // Pixel unit -#define PIXEL_UNIT CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(K0) - - // Block size -#define LHS_BLOCK_SIZE ((K0) * (M0)) - -#if defined(LHS_INTERLEAVE) -#define LHS_OFFSET_X (K0) -#define LHS_STEP_X ((K0) * (V0)) -#define LHS_STEP_LOOP (1) -#else // defined(INTERLEAVE) -#define LHS_OFFSET_X (LHS_BLOCK_SIZE) -#define LHS_STEP_X (K0) -#define LHS_STEP_LOOP (V0) -#endif // defined(INTERLEAVE) - - // Block size -#define RHS_BLOCK_SIZE (PIXEL_UNIT * (N0)) - - // RHS offset and step X -#if defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (PIXEL_UNIT) -#define RHS_STEP_X (PIXEL_UNIT * (H0)) -#define RHS_STEP_LOOP (1) -#else // defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (RHS_BLOCK_SIZE) -#define RHS_STEP_X PIXEL_UNIT -#define RHS_STEP_LOOP (H0) -#endif // defined(RHS_INTERLEAVE) - -#if defined(DUMMY_WORK_ITEMS) - if((get_global_id(0) * N0 >= N) || (get_global_id(1) * M0 >= M)) - { - return; - } -#endif // defined(DUMMY_WORK_ITEMS) - - // Compute LHS matrix address - __global uchar *lhs_addr = lhs_ptr + lhs_offset_first_element_in_bytes + (get_global_id(1) % V0) * (uint)LHS_OFFSET_X * sizeof(DATA_TYPE) + (get_global_id(1) / V0) * (uint)lhs_stride_y + - (get_global_id(2) * lhs_stride_z); - -#if defined(MATRIX_B_DEPTH) - // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 - const uint z_rhs = (get_global_id(2) % MATRIX_B_DEPTH); -#else // defined(MATRIX_B_DEPTH) - const uint z_rhs = get_global_id(2); -#endif // defined(MATRIX_B_DEPTH) - - // Compute RHS matrix coordinates - uint x_rhs = (get_global_id(0) % H0) * (uint)RHS_OFFSET_X; - const uint y_rhs = (get_global_id(0) / (uint)H0) + z_rhs * RHS_HEIGHT; - - // Initialize the accumulators - REPEAT_VAR_INIT_TO_CONST(M0, VEC_DATA_TYPE(DATA_TYPE_ACCUMULATOR, N0), c, 0); - - REPEAT_VAR_INIT_TO_CONST(M0, uint, zlhs, 0); //uint zlhs0=0,zlhs1=0,zlhs2=0,... zlhs7=0; - REPEAT_VAR_INIT_TO_CONST(16, uint, zero, 0); - - for(int i = 0; i < K; i += K0) - { - // Load values from LHS matrix - LOAD_BLOCK(M0, K0, DATA_TYPE, a, lhs_addr, 0, LHS_STEP_X * sizeof(DATA_TYPE), zlhs); - - // Load values from RHS matrix stored in a cl_image - REPEAT_VAR_INIT_TO_CONST(N0, VEC_DATA_TYPE(DATA_TYPE, K0), b, 0); - LOAD_TEXTURE2D(N0, PIXEL_UNIT, DATA_TYPE, b, rhs_img, x_rhs, y_rhs, RHS_STEP_X, 0); - - // Accumulate - ARM_DOT_K0XN0(a0, b, c0); -#if M0 > 1 - ARM_DOT_K0XN0(a1, b, c1); -#endif // M0 > 1 -#if M0 > 2 - ARM_DOT_K0XN0(a2, b, c2); -#endif // M0 > 2 -#if M0 > 3 - ARM_DOT_K0XN0(a3, b, c3); -#endif // M0 > 3 -#if M0 > 4 - ARM_DOT_K0XN0(a4, b, c4); -#endif // M0 > 4 -#if M0 > 5 - ARM_DOT_K0XN0(a5, b, c5); -#endif // M0 > 5 -#if M0 > 6 - ARM_DOT_K0XN0(a6, b, c6); -#endif // M0 > 6 -#if M0 > 7 - ARM_DOT_K0XN0(a7, b, c7); -#endif // M0 > 7 - - lhs_addr += (M0 * LHS_STEP_X * LHS_STEP_LOOP) * sizeof(DATA_TYPE); - - x_rhs += N0 * RHS_STEP_X * RHS_STEP_LOOP; - } - - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)) + (get_global_id(1) * (uint)M0 * dst_stride_y); - - REPEAT_VAR_INIT_TO_CONST(M0, uint, zout, 0); - - // Boundary conditions: detect if current block is at the "bottom" or "right" boundary - const bool cond_y = ((get_global_id(1) + 1) * M0 >= M); - const bool cond_x = ((get_global_id(0) + 1) * N0 >= N); - -#if defined(REINTERPRET_OUTPUT_AS_3D) - - // The plane (zin) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zout, get_global_id(1) * (uint)M0, HEIGHT_GEMM3D, DEPTH_GEMM3D, dst_cross_plane_pad, dst_stride_y); - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply dst_stride_z by DEPTH_GEMM3D - dst_addr += get_global_id(2) * dst_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_OUTPUT_AS_3D) - - // Add offset for batched GEMM - dst_addr += get_global_id(2) * dst_stride_z; - -#endif // defined(REINTERPRET_OUTPUT_AS_3D) - - // Multiply by the weight of matrix-matrix product and store the result -#if defined(ALPHA) - SCALE_BLOCK(M0, DATA_TYPE, c, ALPHA); -#endif // defined(ALPHA) - - // Add beta*bias -#if defined(BETA) -#if defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)); - - LOAD_BLOCK_BOUNDARY_AWARE(1, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, 1, PARTIAL_STORE_N0, false, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(1, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias[broadcasted] - MIXED_PRECISION_ELTWISE_OP_BLOCK_BROADCAST(ADD, M0, N0, c, bias, DATA_TYPE_ACCUMULATOR, bias_hp); - -#else // defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)) + (get_global_id(1) * (uint)M0 * bias_stride_y) + get_global_id( - 2) * bias_stride_z; - - LOAD_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(M0, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias - MIXED_PRECISION_ELTWISE_OP_BLOCK(ADD, M0, N0, c, bias, DATA_TYPE_ACCUMULATOR, bias_hp); - -#endif // defined(BROADCAST_BIAS) -#endif // defined(BETA) - - // c = act(c) - POST_OP1_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - // c = c + eltwise_operand (mix-precision, broadcast, boundary aware) - POST_OP2_ELTWISE_OP(P2_ELTWISE_OP, M0, N0, c, eltwise_operand, get_global_id(1) * (uint)M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, zero, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - // c = act(c) - POST_OP3_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - - // Store output block - MIXED_PRECISION_STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, c, dst_addr, dst_stride_y, zout, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x, c_lp); - -#undef LHS_BLOCK_SIZE -#undef LHS_OFFSET_X -#undef LHS_STEP_X -#undef RHS_BLOCK_SIZE -#undef RHS_OFFSET_X -#undef RHS_STEP_X -#undef PIXEL_UNIT -#undef LHS_STEP_LOOP -#undef RHS_STEP_LOOP -} -#endif // defined(OPENCL_IMAGE_SUPPORT) && defined(GEMM_MM_RESHAPED_LHS_NT_RHS_T_TEXTURE_POST_ACT_ELTWISE_OP_ACT) - -#if defined(LHS_TRANSPOSE) - -#define VTYPE(TYPE, SIZE) VEC_DATA_TYPE(TYPE, SIZE) - -#if defined(MIXED_PRECISION) - -#if(GPU_ARCH == GPU_ARCH_MIDGARD) -#define ARM_VFMA(N0, a, b, c) c += (CONVERT(a, VEC_DATA_TYPE(DATA_TYPE_ACCUMULATOR, N0))) * (CONVERT(b, VEC_DATA_TYPE(DATA_TYPE_ACCUMULATOR, N0))); -#else // GPU_ARCH == GPU_ARCH_MIDGARD -#define ARM_VFMA(N0, a, b, c) c = fma((CONVERT(a, VEC_DATA_TYPE(DATA_TYPE_ACCUMULATOR, N0))), (CONVERT(b, VEC_DATA_TYPE(DATA_TYPE_ACCUMULATOR, N0))), (c)); -#endif // GPU_ARCH == GPU_ARCH_MIDGARD - -#else // defined(MIXED_PRECISION - -#if(GPU_ARCH == GPU_ARCH_MIDGARD) -#define ARM_VFMA(N0, a, b, c) c += (a) * (b); -#else // GPU_ARCH == GPU_ARCH_MIDGARD -#define ARM_VFMA(N0, a, b, c) c = fma((a), (b), (c)); -#endif // GPU_ARCH == GPU_ARCH_MIDGARD - -#endif // defined(MIXED_PRECISION) - -#define ARM_VVM_T_NT_1xN0x1(N0, TYPE, a, b, C) \ - ({ \ - ARM_VFMA(N0, (VTYPE(TYPE, N0))(a), b, (C##0)); \ - }) -#define ARM_VVM_T_NT_2xN0x1(N0, TYPE, a, b, C) \ - ({ \ - ARM_VFMA(N0, (VTYPE(TYPE, N0))(a.s0), b, (C##0)); \ - ARM_VFMA(N0, (VTYPE(TYPE, N0))(a.s1), b, (C##1)); \ - }) -#define ARM_VVM_T_NT_3xN0x1(N0, TYPE, a, b, C) \ - ({ \ - ARM_VVM_T_NT_2xN0x1(N0, TYPE, a, b, C); \ - ARM_VFMA(N0, (VTYPE(TYPE, N0))(a.s2), b, (C##2)); \ - }) -#define ARM_VVM_T_NT_4xN0x1(N0, TYPE, a, b, C) \ - ({ \ - ARM_VVM_T_NT_3xN0x1(N0, TYPE, a, b, C); \ - ARM_VFMA(N0, (VTYPE(TYPE, N0))(a.s3), b, (C##3)); \ - }) -#define ARM_VVM_T_NT_8xN0x1(N0, TYPE, a, b, C) \ - ({ \ - ARM_VVM_T_NT_4xN0x1(N0, TYPE, a, b, C); \ - ARM_VFMA(N0, (VTYPE(TYPE, N0))(a.s4), b, (C##4)); \ - ARM_VFMA(N0, (VTYPE(TYPE, N0))(a.s5), b, (C##5)); \ - ARM_VFMA(N0, (VTYPE(TYPE, N0))(a.s6), b, (C##6)); \ - ARM_VFMA(N0, (VTYPE(TYPE, N0))(a.s7), b, (C##7)); \ - }) - -// Factory macro for the column-vector (transposed) by row-vector (not transposed) multiplication. K0 = 1 -// a is the column-vector (transposed) -// b is the row-vector (not transposed) -// C is the output matrix -// Lower case is a vector (a, b) -// Upper case is a matrix (C) -#define ARM_VVM_T_NT_M0xN0x1(M0, N0, TYPE, a, b, C) ARM_VVM_T_NT_##M0##xN0x1(N0, TYPE, a, b, C) - -#define ARM_MM_T_NT_M0xN0x1(M0, N0, TYPE, A, B, C) \ - ({ \ - ARM_VVM_T_NT_M0xN0x1(M0, N0, TYPE, (A##0), (B##0), C); \ - }) -#define ARM_MM_T_NT_M0xN0x2(M0, N0, TYPE, A, B, C) \ - ({ \ - ARM_MM_T_NT_M0xN0x1(M0, N0, TYPE, A, B, C); \ - ARM_VVM_T_NT_M0xN0x1(M0, N0, TYPE, (A##1), (B##1), C); \ - }) -#define ARM_MM_T_NT_M0xN0x3(M0, N0, TYPE, A, B, C) \ - ({ \ - ARM_MM_T_NT_M0xN0x2(M0, N0, TYPE, A, B, C); \ - ARM_VVM_T_NT_M0xN0x1(M0, N0, TYPE, (A##2), (B##2), C); \ - }) -#define ARM_MM_T_NT_M0xN0x4(M0, N0, TYPE, A, B, C) \ - ({ \ - ARM_MM_T_NT_M0xN0x3(M0, N0, TYPE, A, B, C); \ - ARM_VVM_T_NT_M0xN0x1(M0, N0, TYPE, (A##3), (B##3), C); \ - }) -#define ARM_MM_T_NT_M0xN0x8(M0, N0, TYPE, A, B, C) \ - ({ \ - ARM_MM_T_NT_M0xN0x4(M0, N0, TYPE, A, B, C); \ - ARM_VVM_T_NT_M0xN0x1(M0, N0, TYPE, (A##4), (B##4), C); \ - ARM_VVM_T_NT_M0xN0x1(M0, N0, TYPE, (A##5), (B##5), C); \ - ARM_VVM_T_NT_M0xN0x1(M0, N0, TYPE, (A##6), (B##6), C); \ - ARM_VVM_T_NT_M0xN0x1(M0, N0, TYPE, (A##7), (B##7), C); \ - }) -#define ARM_MM_T_NT_M0xN0x16(M0, N0, TYPE, A, B, C) \ - ({ \ - ARM_MM_T_NT_M0xN0x8(M0, N0, TYPE, A, B, C); \ - ARM_MM_T_NT_M0xN0x1(M0, N0, TYPE, (A##8), (B##8), C); \ - ARM_MM_T_NT_M0xN0x1(M0, N0, TYPE, (A##9), (B##9), C); \ - ARM_MM_T_NT_M0xN0x1(M0, N0, TYPE, (A##A), (B##A), C); \ - ARM_MM_T_NT_M0xN0x1(M0, N0, TYPE, (A##B), (B##B), C); \ - ARM_MM_T_NT_M0xN0x1(M0, N0, TYPE, (A##C), (B##C), C); \ - ARM_MM_T_NT_M0xN0x1(M0, N0, TYPE, (A##D), (B##D), C); \ - ARM_MM_T_NT_M0xN0x1(M0, N0, TYPE, (A##E), (B##E), C); \ - ARM_MM_T_NT_M0xN0x1(M0, N0, TYPE, (A##F), (B##F), C); \ - }) - -// Factory macro for the matrix (transposed) by matrix (not transposed) multiplication. -// The dimensions for this matrix multiplications are defined through M0, N0 and K0 -// The dimensions supported are: -// M0: 1, 2, 3, 4, 8 -// N0: 1, 2, 3, 4, 8, 16 -// K0: 1, 2, 3, 4, 8, 16 -// This macro calls the vector-by-matrix macro K0 times -// A, B and C are matrices -#define ARM_MM_T_NT(M0, N0, K0, TYPE, A, B, C) \ - CONCAT(ARM_MM_T_NT_M0xN0x, K0) \ - (M0, N0, TYPE, A, B, C) - -#if defined(GEMM_MM_RESHAPED_LHS_T_RHS_NT_POST_ACT_ELTWISE_OP_ACT) -/** This OpenCL kernel computes the matrix multiplication between 2 matrices plus 3 post ops: - * Post op 1: activation (optional) - * Post op 2: elementwise op - * Post op 3: activation (optional) - * - * @note (Optional) -DP1_ACTIVATION_TYPE, -DP1_ACTIVATION_A_VAL, -DP1_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * @note (Required) -DP2_ELTWISE_OP: The (binary) elementwise post op to perform - * @note (Required) -DP2_ELTWISE_ARG1_HEIGHT: The height (Y dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Required) -DP2_ELTWISE_ARG1_WIDTH: The width (X dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Optional) -DP3_ACTIVATION_TYPE, -DP3_ACTIVATION_A_VAL, -DP3_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * - * All parameters are similarly defined in kernel gemm_mm_reshaped_lhs_t_rhs_nt, with these additions: - * - * @param[in] eltwise_operand_ptr Pointer to the eltwise operand matrix. Supported data type: F16/F32 - * @param[in] eltwise_operand_stride_x Stride of the eltwise operand matrix in X dimension (in bytes) - * @param[in] eltwise_operand_step_x eltwise_operand_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_y Stride of the eltwise operand matrix in Y dimension (in bytes) - * @param[in] eltwise_operand_step_y eltwise_operand_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_z Stride of the eltwise operand tensor in Z dimension (in bytes) - * @param[in] M Number of rows in LHS matrix not reshaped. - * @param[in] N Number of columns in RHS matrix not reshaped. - * @param[in] K Number of columns in LHS matrix and rows in RHS matrix not reshaped. - */ -__kernel void gemm_mm_reshaped_lhs_t_rhs_nt_post_act_eltwise_op_act(IMAGE_DECLARATION(lhs), - IMAGE_DECLARATION(rhs), -#if defined(BETA) - IMAGE_DECLARATION(bias), -#endif // defined(BETA) - IMAGE_DECLARATION(dst), - // Post Op arguments - IMAGE_DECLARATION(eltwise_operand), - uint lhs_stride_z, - uint rhs_stride_z, -#if defined(BETA) - uint bias_stride_z, -#endif //defined(BETA) - uint dst_stride_z, - uint eltwise_operand_stride_z -#if defined(REINTERPRET_OUTPUT_AS_3D) - , - uint dst_cross_plane_pad -#endif // REINTERPRET_OUTPUT_AS_3D - , - const int M, - const int N, - const int K) -{ - // Block size -#define LHS_BLOCK_SIZE ((K0) * (M0)) - -#if defined(LHS_INTERLEAVE) -#define LHS_OFFSET_X (M0) -#define LHS_STEP_X ((M0) * (V0)) -#define LHS_STEP_LOOP (1) -#else // defined(INTERLEAVE) -#define LHS_OFFSET_X (LHS_BLOCK_SIZE) -#define LHS_STEP_X (M0) -#define LHS_STEP_LOOP (V0) -#endif // defined(INTERLEAVE) - - // Block size -#define RHS_BLOCK_SIZE ((K0) * (N0)) - - // RHS offset and step X -#if defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (N0) -#define RHS_STEP_X ((N0) * (H0)) -#else // defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (RHS_BLOCK_SIZE) -#define RHS_STEP_X (N0) -#endif // defined(RHS_INTERLEAVE) - - const uint x = get_global_id(0); - const uint y = get_global_id(1); - const uint z = get_global_id(2); - - // Boundary conditions: detect if current block is at the "bottom" or "right" boundary - const bool cond_y = ((get_global_id(1) + 1) * M0 >= M); - const bool cond_x = ((get_global_id(0) + 1) * N0 >= N); - -#if defined(DUMMY_WORK_ITEMS) - if((x * N0 >= N) || (y * M0 >= M)) - { - return; - } -#endif // defined(DUMMY_WORK_ITEMS) - - // Compute LHS matrix address - __global uchar *lhs_addr = lhs_ptr + lhs_offset_first_element_in_bytes + (y % V0) * (uint)LHS_OFFSET_X * sizeof(DATA_TYPE) + (y / V0) * (uint)lhs_stride_y + (z * lhs_stride_z); - - // Compute RHS matrix address - __global uchar *rhs_addr = rhs_ptr + rhs_offset_first_element_in_bytes + (x % H0) * (uint)RHS_OFFSET_X * sizeof(DATA_TYPE) + (x / (uint)H0) * rhs_stride_y; - -#if defined(MATRIX_B_DEPTH) - // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 - rhs_addr += (z % MATRIX_B_DEPTH) * rhs_stride_z; -#else // defined(MATRIX_B_DEPTH) - rhs_addr += z * rhs_stride_z; -#endif // defined(MATRIX_B_DEPTH) - - // Initialize the accumulators - REPEAT_VAR_INIT_TO_CONST(M0, VEC_DATA_TYPE(DATA_TYPE_ACCUMULATOR, N0), c, 0); - - REPEAT_VAR_INIT_TO_CONST(M0, uint, zero, 0); - - __global DATA_TYPE *lhs = (__global DATA_TYPE *)(lhs_addr); - __global DATA_TYPE *rhs = (__global DATA_TYPE *)(rhs_addr); - - for(int i = 0; i < K; i += K0) - { - VEC_DATA_TYPE(DATA_TYPE, M0) - a0; - VEC_DATA_TYPE(DATA_TYPE, N0) - b0; - - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; - -#if K0 > 1 - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; -#endif // K0 > 1 - -#if K0 > 2 - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; -#endif // K0 > 2 - -#if K0 > 3 - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; -#endif // K0 > 3 - -#if K0 > 4 - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; -#endif // K0 > 4 - -#if K0 > 8 - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = VLOAD(N0)(0, rhs); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - rhs += RHS_STEP_X; -#endif // K0 > 8 - -#ifndef LHS_INTERLEAVE - lhs += (M0 * K0 * (V0 - 1)); -#endif // LHS_INTERLEAVE - -#ifndef RHS_INTERLEAVE - rhs += (N0 * K0 * (H0 - 1)); -#endif // RHS_INTERLEAVE - } - - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)) + (y * (uint)M0 * dst_stride_y); - - REPEAT_VAR_INIT_TO_CONST(M0, uint, zout, 0); - -#if defined(REINTERPRET_OUTPUT_AS_3D) - - // The plane (zin) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zout, y * (uint)M0, HEIGHT_GEMM3D, DEPTH_GEMM3D, dst_cross_plane_pad, dst_stride_y); - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply dst_stride_z by DEPTH_GEMM3D - dst_addr += z * dst_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_OUTPUT_AS_3D) - - // Add offset for batched GEMM - dst_addr += z * dst_stride_z; - -#endif // defined(REINTERPRET_OUTPUT_AS_3D) - - // Multiply by the weight of matrix-matrix product and store the result -#if defined(ALPHA) - SCALE_BLOCK(M0, DATA_TYPE, c, ALPHA); -#endif // defined(ALPHA) - - // Add beta*bias -#if defined(BETA) -#if defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)); - - LOAD_BLOCK_BOUNDARY_AWARE(1, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, 1, PARTIAL_STORE_N0, false, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(1, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias[broadcasted] - MIXED_PRECISION_ELTWISE_OP_BLOCK_BROADCAST(ADD, M0, N0, c, bias, DATA_TYPE_ACCUMULATOR, bias_hp); - -#else // defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)) + (get_global_id(1) * (uint)M0 * bias_stride_y) + get_global_id( - 2) * bias_stride_z; - - LOAD_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(M0, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias - MIXED_PRECISION_ELTWISE_OP_BLOCK(ADD, M0, N0, c, bias, DATA_TYPE_ACCUMULATOR, bias_hp); - -#endif // defined(BROADCAST_BIAS) -#endif // defined(BETA) - - // c = act(c) - POST_OP1_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - // c = c + eltwise_operand (mix-precision, broadcast, boundary aware) - POST_OP2_ELTWISE_OP(P2_ELTWISE_OP, M0, N0, c, eltwise_operand, get_global_id(1) * (uint)M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, zero, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - // c = act(c) - POST_OP3_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - - // Store output block - MIXED_PRECISION_STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, c, dst_addr, dst_stride_y, zout, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x, c_lp); - -#undef LHS_BLOCK_SIZE -#undef LHS_OFFSET_X -#undef LHS_STEP_X -#undef RHS_BLOCK_SIZE -#undef RHS_OFFSET_X -#undef RHS_STEP_X -} -#endif // defined(GEMM_MM_RESHAPED_LHS_T_RHS_NT_POST_ACT_ELTWISE_OP_ACT) - -#if defined(OPENCL_IMAGE_SUPPORT) && defined(GEMM_MM_RESHAPED_LHS_T_RHS_NT_TEXTURE_POST_ACT_ELTWISE_OP_ACT) -/** This OpenCL kernel computes the matrix multiplication between 2 matrices plus 3 post ops. The RHS matrix is stored in OpenCL image object. - * Post op 1: activation (optional) - * Post op 2: elementwise op - * Post op 3: activation (optional) - * - * @note (Optional) -DP1_ACTIVATION_TYPE, -DP1_ACTIVATION_A_VAL, -DP1_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * @note (Required) -DP2_ELTWISE_OP: The (binary) elementwise post op to perform - * @note (Required) -DP2_ELTWISE_ARG1_HEIGHT: The height (Y dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Required) -DP2_ELTWISE_ARG1_WIDTH: The width (X dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Optional) -DP3_ACTIVATION_TYPE, -DP3_ACTIVATION_A_VAL, -DP3_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * - * All parameters are similarly defined in kernel gemm_mm_reshaped_lhs_t_rhs_nt_texture, with these additions: - * - * @param[in] eltwise_operand_ptr Pointer to the eltwise operand matrix. Supported data type: F16/F32 - * @param[in] eltwise_operand_stride_x Stride of the eltwise operand matrix in X dimension (in bytes) - * @param[in] eltwise_operand_step_x eltwise_operand_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_y Stride of the eltwise operand matrix in Y dimension (in bytes) - * @param[in] eltwise_operand_step_y eltwise_operand_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_z Stride of the eltwise operand tensor in Z dimension (in bytes) - */ -__kernel void gemm_mm_reshaped_lhs_t_rhs_nt_texture_post_act_eltwise_op_act(IMAGE_DECLARATION(lhs), - __read_only image2d_t rhs_img, -#if defined(BETA) - IMAGE_DECLARATION(bias), -#endif // defined(BETA) - IMAGE_DECLARATION(dst), - // Post Op arguments - IMAGE_DECLARATION(eltwise_operand), - uint lhs_stride_z, - uint rhs_stride_z, -#if defined(BETA) - uint bias_stride_z, -#endif //defined(BETA) - uint dst_stride_z, - uint eltwise_operand_stride_z -#if defined(REINTERPRET_OUTPUT_AS_3D) - , - uint dst_cross_plane_pad -#endif // REINTERPRET_OUTPUT_AS_3D - , - const int M, - const int N, - const int K) -{ - // Pixel unit -#define PIXEL_UNIT CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(N0) - - // Block size -#define LHS_BLOCK_SIZE ((K0) * (M0)) - -#if defined(LHS_INTERLEAVE) -#define LHS_OFFSET_X (M0) -#define LHS_STEP_X ((M0) * (V0)) -#define LHS_STEP_LOOP (1) -#else // defined(INTERLEAVE) -#define LHS_OFFSET_X (LHS_BLOCK_SIZE) -#define LHS_STEP_X (M0) -#define LHS_STEP_LOOP (V0) -#endif // defined(INTERLEAVE) - - // Block size -#define RHS_BLOCK_SIZE ((K0) * (PIXEL_UNIT)) - - // RHS offset and step X -#if defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (PIXEL_UNIT) -#define RHS_STEP_X ((PIXEL_UNIT) * (H0)) -#else // defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (RHS_BLOCK_SIZE) -#define RHS_STEP_X (PIXEL_UNIT) -#endif // defined(RHS_INTERLEAVE) - - const uint x = get_global_id(0); - const uint y = get_global_id(1); - const uint z = get_global_id(2); - -#if defined(DUMMY_WORK_ITEMS) - if((x * N0 >= N) || (y * M0 >= M)) - { - return; - } -#endif // defined(DUMMY_WORK_ITEMS) - - // Compute LHS matrix address - __global uchar *lhs_addr = lhs_ptr + lhs_offset_first_element_in_bytes + (y % V0) * (uint)LHS_OFFSET_X * sizeof(DATA_TYPE) + (y / V0) * (uint)lhs_stride_y + (z * lhs_stride_z); - -#if defined(MATRIX_B_DEPTH) - // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 - const uint z_rhs = (z % MATRIX_B_DEPTH); -#else // defined(MATRIX_B_DEPTH) - const uint z_rhs = z; -#endif // defined(MATRIX_B_DEPTH) - - // Compute RHS matrix coordinates - uint x_rhs = (x % H0) * (uint)RHS_OFFSET_X; - const uint y_rhs = (x / (uint)H0) + z_rhs * RHS_HEIGHT; - - // Initialize the accumulators - REPEAT_VAR_INIT_TO_CONST(M0, VEC_DATA_TYPE(DATA_TYPE_ACCUMULATOR, N0), c, 0); - - REPEAT_VAR_INIT_TO_CONST(M0, uint, zero, 0); - - __global DATA_TYPE *lhs = (__global DATA_TYPE *)(lhs_addr); - - for(int i = 0; i < K; i += K0) - { - VEC_DATA_TYPE(DATA_TYPE, M0) - a0; - VEC_DATA_TYPE(DATA_TYPE, N0) - b0; - - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 0 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - -#if K0 > 1 - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 1 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; -#endif // K0 > 1 - -#if K0 > 2 - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 2 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; -#endif // K0 > 2 - -#if K0 > 3 - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 3 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; -#endif // K0 > 3 - -#if K0 > 4 - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 4 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 5 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 6 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 7 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; -#endif // K0 > 4 - -#if K0 > 8 - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 8 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 9 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 10 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 11 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 12 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 13 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 14 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; - - a0 = VLOAD(M0)(0, lhs); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 15 * RHS_STEP_X), (y_rhs)); - - ARM_MM_T_NT(M0, N0, 1, DATA_TYPE, a, b, c); - - lhs += LHS_STEP_X; -#endif // K0 > 8 - -#ifndef LHS_INTERLEAVE - lhs += (M0 * K0 * (V0 - 1)); -#endif // LHS_INTERLEAVE - - x_rhs += K0 * RHS_STEP_X; -#ifndef RHS_INTERLEAVE - x_rhs += (PIXEL_UNIT * K0 * (H0 - 1)); -#endif // RHS_INTERLEAVE - } - - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)) + (y * (uint)M0 * dst_stride_y); - - REPEAT_VAR_INIT_TO_CONST(M0, uint, zout, 0); - - // Boundary conditions: detect if current block is at the "bottom" or "right" boundary - const bool cond_y = ((get_global_id(1) + 1) * M0 >= M); - const bool cond_x = ((get_global_id(0) + 1) * N0 >= N); - -#if defined(REINTERPRET_OUTPUT_AS_3D) - - // The plane (zin) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zout, y * (uint)M0, HEIGHT_GEMM3D, DEPTH_GEMM3D, dst_cross_plane_pad, dst_stride_y); - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply dst_stride_z by DEPTH_GEMM3D - dst_addr += z * dst_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_OUTPUT_AS_3D) - - // Add offset for batched GEMM - dst_addr += z * dst_stride_z; - -#endif // defined(REINTERPRET_OUTPUT_AS_3D) - - // Multiply by the weight of matrix-matrix product and store the result -#if defined(ALPHA) - SCALE_BLOCK(M0, DATA_TYPE, c, ALPHA); -#endif // defined(ALPHA) - - // Add beta*bias -#if defined(BETA) -#if defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)); - - LOAD_BLOCK_BOUNDARY_AWARE(1, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, 1, PARTIAL_STORE_N0, false, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(1, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias[broadcasted] - MIXED_PRECISION_ELTWISE_OP_BLOCK_BROADCAST(ADD, M0, N0, c, bias, DATA_TYPE_ACCUMULATOR, bias_hp); - -#else // defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)) + (y * (uint)M0 * bias_stride_y) + z * bias_stride_z; - - LOAD_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(M0, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - MIXED_PRECISION_ELTWISE_OP_BLOCK(ADD, M0, N0, c, bias, DATA_TYPE_ACCUMULATOR, bias_hp); - -#endif // defined(BROADCAST_BIAS) -#endif // defined(BETA) - - // c = act(c) - POST_OP1_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - // c = c + eltwise_operand (mix-precision, broadcast, boundary aware) - POST_OP2_ELTWISE_OP(P2_ELTWISE_OP, M0, N0, c, eltwise_operand, get_global_id(1) * (uint)M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, zero, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - // c = act(c) - POST_OP3_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - - // Store output block - MIXED_PRECISION_STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, c, dst_addr, dst_stride_y, zout, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x, c_lp); - -#undef LHS_BLOCK_SIZE -#undef LHS_OFFSET_X -#undef LHS_STEP_X -#undef RHS_BLOCK_SIZE -#undef RHS_OFFSET_X -#undef RHS_STEP_X -#undef PIXEL_UNIT -#undef LHS_STEP_LOOP -#undef RHS_STEP_LOOP -} -#endif // defined(OPENCL_IMAGE_SUPPORT) && defined(GEMM_MM_RESHAPED_LHS_T_RHS_NT_TEXTURE_POST_ACT_ELTWISE_OP_ACT) - -#endif // defined(LHS_TRANSPOSE) -#endif // defined(P2_ELTWISE_OP) && defined(P2_ELTWISE_ARG1_HEIGHT) && defined(P2_ELTWISE_ARG1_WIDTH) -#endif // defined(M0) && defined(N0) && defined(K0) && defined(V0) && defined(H0) && defined(DATA_TYPE) && defined(DATA_TYPE_ACCUMULATOR) diff --git a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped_only_rhs.cl b/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped_only_rhs.cl deleted file mode 100644 index 09ddcde043..0000000000 --- a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped_only_rhs.cl +++ /dev/null @@ -1,1399 +0,0 @@ -/* - * Copyright (c) 2021-2022 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ -#include "fp_post_ops_act_eltwise_op_act.h" -#include "gemm_helpers.h" -#include "repeat.h" - -/** (EXPERIMENTAL_POST_OPS) gemm_mm_reshaped_only_rhs kernel */ -#if defined(M0) && defined(N0) && defined(K0) && defined(H0) && defined(DATA_TYPE) -#if defined(P2_ELTWISE_OP) && defined(P2_ELTWISE_ARG1_HEIGHT) && defined(P2_ELTWISE_ARG1_WIDTH) - -#define CONCAT(a, b) a##b - -#define ARM_DOT1(a, b, c) \ - ({ \ - c = fma(a, b, c); \ - }) -#define ARM_DOT2(a, b, c) \ - ({ \ - c = fma(a.s0, b.s0, c); \ - c = fma(a.s1, b.s1, c); \ - }) -#define ARM_DOT3(a, b, c) \ - ({ \ - ARM_DOT2(a, b, c); \ - c = fma((a.s2), (b.s2), c); \ - }) -#define ARM_DOT4(a, b, c) \ - ({ \ - ARM_DOT3(a, b, c); \ - c = fma((a.s3), (b.s3), c); \ - }) -#define ARM_DOT8(a, b, c) \ - ({ \ - ARM_DOT4((a.lo), (b.lo), c); \ - ARM_DOT4((a.hi), (b.hi), c); \ - }) -#define ARM_DOT16(a, b, c) \ - ({ \ - ARM_DOT8((a.lo), (b.lo), c); \ - ARM_DOT8((a.hi), (b.hi), c); \ - }) - -#if N0 == 2 -#define ARM_DOT_K0XN0(k0, a, b, c) \ - ({ \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##0), (c.s0)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##1), (c.s1)); \ - }) -#elif N0 == 3 // N0 == 3 -#define ARM_DOT_K0XN0(k0, a, b, c) \ - ({ \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##0), (c.s0)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##1), (c.s1)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##2), (c.s2)); \ - }) -#elif N0 == 4 // N0 == 4 -#define ARM_DOT_K0XN0(k0, a, b, c) \ - ({ \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##0), (c.s0)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##1), (c.s1)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##2), (c.s2)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##3), (c.s3)); \ - }) -#elif N0 == 8 // N0 == 8 -#define ARM_DOT_K0XN0(k0, a, b, c) \ - ({ \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##0), (c.s0)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##1), (c.s1)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##2), (c.s2)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##3), (c.s3)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##4), (c.s4)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##5), (c.s5)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##6), (c.s6)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##7), (c.s7)); \ - }) -#elif N0 == 16 // N0 == 16 -#define ARM_DOT_K0XN0(k0, a, b, c) \ - ({ \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##0), (c.s0)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##1), (c.s1)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##2), (c.s2)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##3), (c.s3)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##4), (c.s4)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##5), (c.s5)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##6), (c.s6)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##7), (c.s7)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##8), (c.s8)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##9), (c.s9)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##A), (c.sA)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##B), (c.sB)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##C), (c.sC)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##D), (c.sD)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##E), (c.sE)); \ - CONCAT(ARM_DOT, k0) \ - ((a), (b##F), (c.sF)); \ - }) -#else // N0 not supported -#error "N0 value not supported" -#endif // N0 conditions - -#if defined(GEMM_MM_RESHAPED_ONLY_RHS_T_POST_ACT_ELTWISE_OP_ACT) -/** This OpenCL kernel computes the matrix multiplication between 2 matrices plus 3 post ops: - * Post op 1: activation (optional) - * Post op 2: elementwise op - * Post op 3: activation (optional) - * - * @note (Optional) -DP1_ACTIVATION_TYPE, -DP1_ACTIVATION_A_VAL, -DP1_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * @note (Required) -DP2_ELTWISE_OP: The (binary) elementwise post op to perform - * @note (Required) -DP2_ELTWISE_ARG1_HEIGHT: The height (Y dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Required) -DP2_ELTWISE_ARG1_WIDTH: The width (X dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Optional) -DP3_ACTIVATION_TYPE, -DP3_ACTIVATION_A_VAL, -DP3_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * - * All parameters are similarly defined in kernel gemm_mm_reshaped_only_rhs_t, with these additions: - * - * @param[in] eltwise_operand_ptr Pointer to the eltwise operand matrix. Supported data type: F16/F32 - * @param[in] eltwise_operand_stride_x Stride of the eltwise operand matrix in X dimension (in bytes) - * @param[in] eltwise_operand_step_x eltwise_operand_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_y Stride of the eltwise operand matrix in Y dimension (in bytes) - * @param[in] eltwise_operand_step_y eltwise_operand_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_z Stride of the eltwise operand tensor in Z dimension (in bytes) - */ -__kernel void gemm_mm_reshaped_only_rhs_t_post_act_eltwise_op_act(IMAGE_DECLARATION(lhs), - IMAGE_DECLARATION(rhs), -#if defined(BETA) - IMAGE_DECLARATION(bias), -#endif // defined(BETA) - IMAGE_DECLARATION(dst), - // Post Op arguments - IMAGE_DECLARATION(eltwise_operand), - uint lhs_stride_z, - uint rhs_stride_z, -#if defined(BETA) - uint bias_stride_z, -#endif //defined(BETA) - uint dst_stride_z, - uint eltwise_operand_stride_z -#if defined(REINTERPRET_INPUT_AS_3D) - , - uint lhs_cross_plane_pad -#endif // REINTERPRET_INPUT_AS_3D -#if defined(REINTERPRET_OUTPUT_AS_3D) - , - uint dst_cross_plane_pad -#endif // REINTERPRET_OUTPUT_AS_3D - , - const int M, - const int N, - const int K) -{ - // Block size -#define RHS_BLOCK_SIZE ((K0) * (N0)) - - // RHS offset and step X -#if defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (K0) -#define RHS_STEP_X ((K0) * (H0)) -#define RHS_STEP_LOOP (1) -#else // defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (RHS_BLOCK_SIZE) -#define RHS_STEP_X (K0) -#define RHS_STEP_LOOP (H0) -#endif // defined(RHS_INTERLEAVE) - - uint x = get_global_id(0); - uint y = get_global_id(1); - uint z = get_global_id(2); - - const bool cond_y = y == 0; - const bool cond_x = ((x + 1) * N0 >= N); - -#if defined(DUMMY_WORK_ITEMS) - if((x * N0 >= N) || (y * M0 >= M)) - { - return; - } -#endif // defined(DUMMY_WORK_ITEMS) - - // Compute LHS matrix address - uint lhs_offset = lhs_offset_first_element_in_bytes + COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * (uint)lhs_stride_y; - - // Compute RHS reshaped matrix address - uint rhs_offset = rhs_offset_first_element_in_bytes + (x % H0) * (uint)RHS_OFFSET_X * sizeof(DATA_TYPE) + (x / (uint)H0) * rhs_stride_y; - -#if defined(MATRIX_B_DEPTH) - // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 - rhs_offset += (z % MATRIX_B_DEPTH) * rhs_stride_z; -#else // defined(MATRIX_B_DEPTH) - rhs_offset += z * rhs_stride_z; -#endif // defined(MATRIX_B_DEPTH) - - REPEAT_VAR_INIT_TO_CONST(8, uint, zlhs, 0); //uint zlhs0=0,zlhs1=0,zlhs2=0,... zlhs7=0; - REPEAT_VAR_INIT_TO_CONST(16, uint, zero, 0); - -#if defined(REINTERPRET_INPUT_AS_3D) - // The plane (zlhs) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zlhs, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), HEIGHT_GEMM3D, DEPTH_GEMM3D, lhs_cross_plane_pad, lhs_stride_y); - - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply lhs_stride_z by DEPTH_GEMM3D - lhs_offset += z * lhs_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_INPUT_AS_3D) - - // Add offset for batched GEMM - lhs_offset += z * lhs_stride_z; - -#endif // defined(REINTERPRET_INPUT_AS_3D) - - // Initialize the accumulators - REPEAT_VAR_INIT_TO_CONST(M0, VEC_DATA_TYPE(DATA_TYPE, N0), c, 0); //VEC_DATA_TYPE(DATA_TYPE, N0) c0=0,c1=0,c2=0,... c(M0-1)=0; - - int i = 0; - for(; i <= (K - K0); i += K0) - { - // Supported cases (M0, K0): - // 1,2 - 1,3 - 1,4 - 1,8 - 1,16 - // 2,2 - 2,3 - 2,4 - 2,8 - 2,16 - // 3,2 - 3,3 - 3,4 - 3,8 - 3,16 - // 4,2 - 4,3 - 4,4 - 4,8 - 4,16 - // 5,2 - 5,3 - 5,4 - 5,8 - 5,16 - // 6,2 - 6,3 - 6,4 - 6,8 - 6,16 - // 7,2 - 7,3 - 7,4 - 7,8 - 7,16 - // 8,2 - 8,3 - 8,4 - 8,8 - 8,16 - // Load values from LHS matrix - LOAD_BLOCK(M0, K0, DATA_TYPE, a, lhs_ptr, lhs_offset, lhs_stride_y, zlhs); - - // Load values from RHS reshaped matrix - LOAD_BLOCK(N0, K0, DATA_TYPE, b, rhs_ptr, rhs_offset, RHS_STEP_X * sizeof(DATA_TYPE), zero); - - // Accumulate - ARM_DOT_K0XN0(K0, a0, b, c0); -#if M0 > 1 - ARM_DOT_K0XN0(K0, a1, b, c1); -#endif // M0 > 1 -#if M0 > 2 - ARM_DOT_K0XN0(K0, a2, b, c2); -#endif // M0 > 2 -#if M0 > 3 - ARM_DOT_K0XN0(K0, a3, b, c3); -#endif // M0 > 3 -#if M0 > 4 - ARM_DOT_K0XN0(K0, a4, b, c4); -#endif // M0 > 4 -#if M0 > 5 - ARM_DOT_K0XN0(K0, a5, b, c5); -#endif // M0 > 5 -#if M0 > 6 - ARM_DOT_K0XN0(K0, a6, b, c6); -#endif // M0 > 6 -#if M0 > 7 - ARM_DOT_K0XN0(K0, a7, b, c7); -#endif // M0 > 7 - - lhs_offset += K0 * sizeof(DATA_TYPE); - rhs_offset += (N0 * RHS_STEP_X * RHS_STEP_LOOP) * sizeof(DATA_TYPE); - } - - // Left-over accumulations - for(; i < K; ++i) - { - // Load values from LHS matrix - LOAD_BLOCK(M0, 1, DATA_TYPE, a, lhs_ptr, lhs_offset, lhs_stride_y, zlhs); - - // Load values from RHS reshaped matrix - LOAD_BLOCK(N0, 1, DATA_TYPE, b, rhs_ptr, rhs_offset, RHS_STEP_X * sizeof(DATA_TYPE), zero); - - // Accumulate - ARM_DOT_K0XN0(1, a0, b, c0); -#if M0 > 1 - ARM_DOT_K0XN0(1, a1, b, c1); -#endif // M0 > 1 -#if M0 > 2 - ARM_DOT_K0XN0(1, a2, b, c2); -#endif // M0 > 2 -#if M0 > 3 - ARM_DOT_K0XN0(1, a3, b, c3); -#endif // M0 > 3 -#if M0 > 4 - ARM_DOT_K0XN0(1, a4, b, c4); -#endif // M0 > 4 -#if M0 > 5 - ARM_DOT_K0XN0(1, a5, b, c5); -#endif // M0 > 5 -#if M0 > 6 - ARM_DOT_K0XN0(1, a6, b, c6); -#endif // M0 > 6 -#if M0 > 7 - ARM_DOT_K0XN0(1, a7, b, c7); -#endif // M0 > 7 - - lhs_offset += sizeof(DATA_TYPE); - rhs_offset += sizeof(DATA_TYPE); - } - - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)) + (COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * dst_stride_y); - - REPEAT_VAR_INIT_TO_CONST(8, uint, zout, 0); //uint zout0=0,zout1=0,zout2=0,... zout7=0; - -#if defined(REINTERPRET_OUTPUT_AS_3D) - - // The plane (zout) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zout, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), HEIGHT_GEMM3D, DEPTH_GEMM3D, dst_cross_plane_pad, dst_stride_y); - - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply dst_stride_z by DEPTH_GEMM3D - dst_addr += z * dst_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_OUTPUT_AS_3D) - - // Add offset for batched GEMM - dst_addr += z * dst_stride_z; - -#endif // defined(REINTERPRET_OUTPUT_AS_3D) - - // Multiply by the weight of matrix-matrix product and store the result -#if defined(ALPHA) - SCALE_BLOCK(M0, DATA_TYPE, c, ALPHA); -#endif // defined(ALPHA) - - // Add beta*bias -#if defined(BETA) -#if defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)); - - LOAD_BLOCK_BOUNDARY_AWARE(1, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, 1, PARTIAL_STORE_N0, false, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(1, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias[broadcasted] - ADD_BLOCK_BROADCAST(M0, c, bias0); - -#else // defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)) + (COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * bias_stride_y) + z * bias_stride_z; - - LOAD_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(M0, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias - ADD_BLOCK(M0, c, bias); - -#endif // defined(BROADCAST_BIAS) -#endif // defined(BETA) - - // c = act(c) - POST_OP1_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - // c = c + eltwise_operand (mix-precision, broadcast, boundary aware) - POST_OP2_ELTWISE_OP(P2_ELTWISE_OP, M0, N0, c, eltwise_operand, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), DATA_TYPE, DATA_TYPE_ACCUMULATOR, zero, 1, PARTIAL_STORE_N0, false, cond_x); - // c = act(c) - POST_OP3_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - - // Store output block - STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, c, dst_addr, dst_stride_y, zout, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - -#undef RHS_BLOCK_SIZE -#undef RHS_OFFSET_X -#undef RHS_STEP_X -} -#endif // defined(GEMM_MM_RESHAPED_ONLY_RHS_T_POST_ACT_ELTWISE_OP_ACT) - -#if defined(OPENCL_IMAGE_SUPPORT) && defined(GEMM_MM_RESHAPED_ONLY_RHS_T_TEXTURE_POST_ACT_ELTWISE_OP_ACT) -/** This OpenCL kernel computes the matrix multiplication between 2 matrices plus 3 post ops. The RHS matrix is stored in OpenCL image object. - * Post op 1: activation (optional) - * Post op 2: elementwise op - * Post op 3: activation (optional) - * - * @note (Optional) -DP1_ACTIVATION_TYPE, -DP1_ACTIVATION_A_VAL, -DP1_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * @note (Required) -DP2_ELTWISE_OP: The (binary) elementwise post op to perform - * @note (Required) -DP2_ELTWISE_ARG1_HEIGHT: The height (Y dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Required) -DP2_ELTWISE_ARG1_WIDTH: The width (X dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Optional) -DP3_ACTIVATION_TYPE, -DP3_ACTIVATION_A_VAL, -DP3_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * - * All parameters are similarly defined in kernel gemm_mm_reshaped_only_rhs_t_texture, with these additions: - * - * @param[in] eltwise_operand_ptr Pointer to the eltwise operand matrix. Supported data type: F16/F32 - * @param[in] eltwise_operand_stride_x Stride of the eltwise operand matrix in X dimension (in bytes) - * @param[in] eltwise_operand_step_x eltwise_operand_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_y Stride of the eltwise operand matrix in Y dimension (in bytes) - * @param[in] eltwise_operand_step_y eltwise_operand_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_z Stride of the eltwise operand tensor in Z dimension (in bytes) - * @param[in] M Number of rows in LHS matrix not reshaped. - * @param[in] N Number of columns in RHS matrix not reshaped. - * @param[in] K Number of columns in LHS matrix and rows in RHS matrix not reshaped. - */ -__kernel void gemm_mm_reshaped_only_rhs_t_texture_post_act_eltwise_op_act(IMAGE_DECLARATION(lhs), - __read_only image2d_t rhs_img, -#if defined(BETA) - IMAGE_DECLARATION(bias), -#endif // defined(BETA) - IMAGE_DECLARATION(dst), - // Post Op arguments - IMAGE_DECLARATION(eltwise_operand), - uint lhs_stride_z, - uint rhs_stride_z, -#if defined(BETA) - uint bias_stride_z, -#endif //defined(BETA) - uint dst_stride_z, - uint eltwise_operand_stride_z -#if defined(REINTERPRET_INPUT_AS_3D) - , - uint lhs_cross_plane_pad -#endif // REINTERPRET_INPUT_AS_3D -#if defined(REINTERPRET_OUTPUT_AS_3D) - , - uint dst_cross_plane_pad -#endif // REINTERPRET_OUTPUT_AS_3D - , - const int M, - const int N, - const int K) -{ - // Pixel unit -#define PIXEL_UNIT CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(K0) - - const uint LEFTOVER_K = K % K0; - - // Block size -#define RHS_BLOCK_SIZE (PIXEL_UNIT * (N0)) - - // RHS offset and step X -#if defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (PIXEL_UNIT) -#define RHS_STEP_X (PIXEL_UNIT * (H0)) -#define RHS_STEP_LOOP (1) -#else // defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (RHS_BLOCK_SIZE) -#define RHS_STEP_X PIXEL_UNIT -#define RHS_STEP_LOOP (H0) -#endif // defined(RHS_INTERLEAVE) - - uint x = get_global_id(0); - uint y = get_global_id(1); - uint z = get_global_id(2); - - const bool cond_y = y == 0; - const bool cond_x = ((x + 1) * N0 >= N); - -#if defined(DUMMY_WORK_ITEMS) - if((x * N0 >= N) || (y * M0 >= M)) - { - return; - } -#endif // defined(DUMMY_WORK_ITEMS) - - // Compute LHS matrix address - uint lhs_offset = lhs_offset_first_element_in_bytes + COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * (uint)lhs_stride_y; - -#if defined(MATRIX_B_DEPTH) - // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 - const uint z_rhs = (get_global_id(2) % MATRIX_B_DEPTH); -#else // defined(MATRIX_B_DEPTH) - const uint z_rhs = get_global_id(2); -#endif // defined(MATRIX_B_DEPTH) - - // Compute RHS matrix coordinates - uint x_rhs = (get_global_id(0) % H0) * (uint)RHS_OFFSET_X; - const uint y_rhs = (get_global_id(0) / (uint)H0) + z_rhs * RHS_HEIGHT; - - REPEAT_VAR_INIT_TO_CONST(M0, uint, zlhs, 0); - REPEAT_VAR_INIT_TO_CONST(16, uint, zero, 0); - -#if defined(REINTERPRET_INPUT_AS_3D) - // The plane (zlhs) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zlhs, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), HEIGHT_GEMM3D, DEPTH_GEMM3D, lhs_cross_plane_pad, lhs_stride_y); - - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply lhs_stride_z by DEPTH_GEMM3D - lhs_offset += z * lhs_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_INPUT_AS_3D) - - // Add offset for batched GEMM - lhs_offset += z * lhs_stride_z; - -#endif // defined(REINTERPRET_INPUT_AS_3D) - - // Initialize the accumulators - REPEAT_VAR_INIT_TO_CONST(M0, VEC_DATA_TYPE(DATA_TYPE, N0), c, 0); - - int i = 0; - for(; i <= (K - K0); i += K0) - { - // Load values from LHS matrix - LOAD_BLOCK(M0, K0, DATA_TYPE, a, lhs_ptr, lhs_offset, lhs_stride_y, zlhs); - - // Load values from RHS matrix stored in a cl_image - REPEAT_VAR_INIT_TO_CONST(N0, VEC_DATA_TYPE(DATA_TYPE, K0), b, 0); - LOAD_TEXTURE2D(N0, PIXEL_UNIT, DATA_TYPE, b, rhs_img, x_rhs, y_rhs, RHS_STEP_X, 0); - - // Accumulate - ARM_DOT_K0XN0(K0, a0, b, c0); -#if M0 > 1 - ARM_DOT_K0XN0(K0, a1, b, c1); -#endif // M0 > 1 -#if M0 > 2 - ARM_DOT_K0XN0(K0, a2, b, c2); -#endif // M0 > 2 -#if M0 > 3 - ARM_DOT_K0XN0(K0, a3, b, c3); -#endif // M0 > 3 -#if M0 > 4 - ARM_DOT_K0XN0(K0, a4, b, c4); -#endif // M0 > 4 -#if M0 > 5 - ARM_DOT_K0XN0(K0, a5, b, c5); -#endif // M0 > 5 -#if M0 > 6 - ARM_DOT_K0XN0(K0, a6, b, c6); -#endif // M0 > 6 -#if M0 > 7 - ARM_DOT_K0XN0(K0, a7, b, c7); -#endif // M0 > 7 - - lhs_offset += K0 * sizeof(DATA_TYPE); - x_rhs += N0 * RHS_STEP_X * RHS_STEP_LOOP; - } - - if(LEFTOVER_K != 0) - { - // Note: We cannot read out-of-bound elements from the RHS matrix because - // the RHS width is always multiple of K0. This is not be true for the LHS matrix - - union UNION_VEC_TYPE - { - DATA_TYPE s[K0]; - VEC_DATA_TYPE(DATA_TYPE, K0) - v; - }; - - union UNION_VEC_TYPE a0 = {.v = 0 }; -#if M0 > 1 - union UNION_VEC_TYPE a1 = {.v = 0 }; -#endif // M0 > 1 -#if M0 > 2 - union UNION_VEC_TYPE a2 = {.v = 0 }; -#endif // M0 > 2 -#if M0 > 3 - union UNION_VEC_TYPE a3 = {.v = 0 }; -#endif // M0 > 3 -#if M0 > 4 - union UNION_VEC_TYPE a4 = {.v = 0 }; -#endif // M0 > 4 -#if M0 > 5 - union UNION_VEC_TYPE a5 = {.v = 0 }; -#endif // M0 > 5 -#if M0 > 6 - union UNION_VEC_TYPE a6 = {.v = 0 }; -#endif // M0 > 6 -#if M0 > 7 - union UNION_VEC_TYPE a7 = {.v = 0 }; -#endif // M0 > 7 - - REPEAT_VAR_INIT_TO_CONST(N0, VEC_DATA_TYPE(DATA_TYPE, K0), b, 0); - - // Load from RHS matrix - LOAD_TEXTURE2D(N0, PIXEL_UNIT, DATA_TYPE, b, rhs_img, x_rhs, y_rhs, RHS_STEP_X, 0); - - // Load from LHS matrix - for(int k = 0; k < LEFTOVER_K; ++k) - { - a0.s[k] = *(__global DATA_TYPE *)(lhs_ptr + lhs_offset + 0 * lhs_stride_y + zlhs0); -#if M0 > 1 - a1.s[k] = *(__global DATA_TYPE *)(lhs_ptr + lhs_offset + 1 * lhs_stride_y + zlhs1); -#endif // M0 > 1 -#if M0 > 2 - a2.s[k] = *(__global DATA_TYPE *)(lhs_ptr + lhs_offset + 2 * lhs_stride_y + zlhs2); -#endif // M0 > 2 -#if M0 > 3 - a3.s[k] = *(__global DATA_TYPE *)(lhs_ptr + lhs_offset + 3 * lhs_stride_y + zlhs3); -#endif // M0 > 3 -#if M0 > 4 - a4.s[k] = *(__global DATA_TYPE *)(lhs_ptr + lhs_offset + 4 * lhs_stride_y + zlhs4); -#endif // M0 > 4 -#if M0 > 5 - a5.s[k] = *(__global DATA_TYPE *)(lhs_ptr + lhs_offset + 5 * lhs_stride_y + zlhs5); -#endif // M0 > 5 -#if M0 > 6 - a6.s[k] = *(__global DATA_TYPE *)(lhs_ptr + lhs_offset + 6 * lhs_stride_y + zlhs6); -#endif // M0 > 6 -#if M0 > 7 - a7.s[k] = *(__global DATA_TYPE *)(lhs_ptr + lhs_offset + 7 * lhs_stride_y + zlhs7); -#endif // M0 > 7 - - lhs_offset += sizeof(DATA_TYPE); - } - - // Accumulate - ARM_DOT_K0XN0(K0, a0.v, b, c0); -#if M0 > 1 - ARM_DOT_K0XN0(K0, a1.v, b, c1); -#endif // M0 > 1 -#if M0 > 2 - ARM_DOT_K0XN0(K0, a2.v, b, c2); -#endif // M0 > 2 -#if M0 > 3 - ARM_DOT_K0XN0(K0, a3.v, b, c3); -#endif // M0 > 3 -#if M0 > 4 - ARM_DOT_K0XN0(K0, a4.v, b, c4); -#endif // M0 > 4 -#if M0 > 5 - ARM_DOT_K0XN0(K0, a5.v, b, c5); -#endif // M0 > 5 -#if M0 > 6 - ARM_DOT_K0XN0(K0, a6.v, b, c6); -#endif // M0 > 6 -#if M0 > 7 - ARM_DOT_K0XN0(K0, a7.v, b, c7); -#endif // M0 > 7 - } - - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)) + (COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * dst_stride_y); - - REPEAT_VAR_INIT_TO_CONST(M0, uint, zout, 0); //uint zout0=0,zout1=0,zout2=0,... zout7=0; - -#if defined(REINTERPRET_OUTPUT_AS_3D) - - // The plane (zout) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zout, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), HEIGHT_GEMM3D, DEPTH_GEMM3D, dst_cross_plane_pad, dst_stride_y); - - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply dst_stride_z by DEPTH_GEMM3D - dst_addr += z * dst_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_OUTPUT_AS_3D) - - // Add offset for batched GEMM - dst_addr += z * dst_stride_z; - -#endif // defined(REINTERPRET_OUTPUT_AS_3D) - - // Multiply by the weight of matrix-matrix product and store the result -#if defined(ALPHA) - SCALE_BLOCK(M0, DATA_TYPE, c, ALPHA); -#endif // defined(ALPHA) - - // Add beta*bias -#if defined(BETA) -#if defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)); - - LOAD_BLOCK_BOUNDARY_AWARE(1, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, 1, PARTIAL_STORE_N0, false, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(1, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias[broadcasted] - ADD_BLOCK_BROADCAST(M0, c, bias0); - -#else // defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)) + (COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * bias_stride_y) + z * bias_stride_z; - - LOAD_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(M0, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias - ADD_BLOCK(M0, c, bias); - -#endif // defined(BROADCAST_BIAS) -#endif // defined(BETA) - - // c = act(c) - POST_OP1_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - // c = c + eltwise_operand (mix-precision, broadcast, boundary aware) - POST_OP2_ELTWISE_OP(P2_ELTWISE_OP, M0, N0, c, eltwise_operand, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), DATA_TYPE, DATA_TYPE_ACCUMULATOR, zero, 1, PARTIAL_STORE_N0, false, cond_x); - // c = act(c) - POST_OP3_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - - // Store output block - STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, c, dst_addr, dst_stride_y, zout, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - -#undef RHS_BLOCK_SIZE -#undef RHS_OFFSET_X -#undef RHS_STEP_X -#undef PIXEL_UNIT -} -#endif // defined(OPENCL_IMAGE_SUPPORT) && defined(GEMM_MM_RESHAPED_ONLY_RHS_T_TEXTURE_POST_ACT_ELTWISE_OP_ACT) - -#define VFMA(a, b, c) \ - ({ \ - c = fma(a, b, c); \ - }) - -#if M0 == 1 -#define VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - }) -#elif M0 == 2 // M0 == 2 -#define VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - }) -#elif M0 == 3 // M0 == 3 -#define VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##2).s##i), b, (c##2)); \ - }) -#elif M0 == 4 // M0 == 4 -#define VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##2).s##i), b, (c##2)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##3).s##i), b, (c##3)); \ - }) -#elif M0 == 5 // M0 == 5 -#define VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##2).s##i), b, (c##2)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##3).s##i), b, (c##3)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##4).s##i), b, (c##4)); \ - }) -#elif M0 == 6 // M0 == 6 -#define VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##2).s##i), b, (c##2)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##3).s##i), b, (c##3)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##4).s##i), b, (c##4)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##5).s##i), b, (c##5)); \ - }) -#elif M0 == 7 // M0 == 7 -#define VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##2).s##i), b, (c##2)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##3).s##i), b, (c##3)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##4).s##i), b, (c##4)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##5).s##i), b, (c##5)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##6).s##i), b, (c##6)); \ - }) -#elif M0 == 8 // M0 == 8 -#define VFMA_M0xN0(i, a, b, c) \ - ({ \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##0).s##i), b, (c##0)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##1).s##i), b, (c##1)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##2).s##i), b, (c##2)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##3).s##i), b, (c##3)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##4).s##i), b, (c##4)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##5).s##i), b, (c##5)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##6).s##i), b, (c##6)); \ - VFMA((VEC_DATA_TYPE(DATA_TYPE, N0))((a##7).s##i), b, (c##7)); \ - }) -#else // M0 not supported -#error "M0 not supported" -#endif // M0 not supported - -#if defined(GEMM_MM_RESHAPED_ONLY_RHS_NT_POST_ACT_ELTWISE_OP_ACT) -/** This OpenCL kernel computes the matrix multiplication between 2 matrices plus 3 post ops: - * Post op 1: activation (optional) - * Post op 2: elementwise op - * Post op 3: activation (optional) - * - * @note (Optional) -DP1_ACTIVATION_TYPE, -DP1_ACTIVATION_A_VAL, -DP1_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * @note (Required) -DP2_ELTWISE_OP: The (binary) elementwise post op to perform - * @note (Required) -DP2_ELTWISE_ARG1_HEIGHT: The height (Y dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Required) -DP2_ELTWISE_ARG1_WIDTH: The width (X dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Optional) -DP3_ACTIVATION_TYPE, -DP3_ACTIVATION_A_VAL, -DP3_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * - * All parameters are similarly defined in kernel gemm_mm_reshaped_only_rhs_nt, with these additions: - * - * @param[in] eltwise_operand_ptr Pointer to the eltwise operand matrix. Supported data type: F16/F32 - * @param[in] eltwise_operand_stride_x Stride of the eltwise operand matrix in X dimension (in bytes) - * @param[in] eltwise_operand_step_x eltwise_operand_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_y Stride of the eltwise operand matrix in Y dimension (in bytes) - * @param[in] eltwise_operand_step_y eltwise_operand_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_z Stride of the eltwise operand tensor in Z dimension (in bytes) - * @param[in] M Number of rows in LHS matrix not reshaped. - * @param[in] N Number of columns in RHS matrix not reshaped. - * @param[in] K Number of columns in LHS matrix and rows in RHS matrix not reshaped. - */ -__kernel void gemm_mm_reshaped_only_rhs_nt_post_act_eltwise_op_act(IMAGE_DECLARATION(lhs), - IMAGE_DECLARATION(rhs), -#if defined(BETA) - IMAGE_DECLARATION(bias), -#endif // defined(BETA) - IMAGE_DECLARATION(dst), - // Post Op arguments - IMAGE_DECLARATION(eltwise_operand), - uint lhs_stride_z, - uint rhs_stride_z, -#if defined(BETA) - uint bias_stride_z, -#endif //defined(BETA) - uint dst_stride_z, - uint eltwise_operand_stride_z -#if defined(REINTERPRET_INPUT_AS_3D) - , - uint lhs_cross_plane_pad -#endif // REINTERPRET_INPUT_AS_3D -#if defined(REINTERPRET_OUTPUT_AS_3D) - , - uint dst_cross_plane_pad -#endif // REINTERPRET_OUTPUT_AS_3D - , - const int M, - const int N, - const int K) -{ - // Block size -#define RHS_BLOCK_SIZE ((K0) * (N0)) - - // RHS offset and step X -#if defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (N0) -#define RHS_STEP_X ((N0) * (H0)) -#define RHS_STEP_LOOP (1) -#else // defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (RHS_BLOCK_SIZE) -#define RHS_STEP_X (N0) -#define RHS_STEP_LOOP (H0) -#endif // defined(RHS_INTERLEAVE) - - uint x = get_global_id(0); - uint y = get_global_id(1); - uint z = get_global_id(2); - - const bool cond_y = y == 0; - const bool cond_x = ((x + 1) * N0 >= N); - -#if defined(DUMMY_WORK_ITEMS) - if((x * N0 >= N) || (y * M0 >= M)) - { - return; - } -#endif // defined(DUMMY_WORK_ITEMS) - - // Compute LHS matrix address - uint lhs_offset = lhs_offset_first_element_in_bytes + COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * (uint)lhs_stride_y; - - // Compute RHS reshaped matrix address - uint rhs_offset = rhs_offset_first_element_in_bytes + (x % H0) * (uint)RHS_OFFSET_X * sizeof(DATA_TYPE) + (x / (uint)H0) * rhs_stride_y; - -#if defined(MATRIX_B_DEPTH) - // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 - rhs_offset += (z % MATRIX_B_DEPTH) * rhs_stride_z; -#else // defined(MATRIX_B_DEPTH) - rhs_offset += z * rhs_stride_z; -#endif // defined(MATRIX_B_DEPTH) - - REPEAT_VAR_INIT_TO_CONST(8, uint, zin, 0); //uint zin0=0,zin1=0,zin2=0,... zin7=0; - REPEAT_VAR_INIT_TO_CONST(16, uint, zero, 0); //uint zero0=0,zero1=0,zero2=0,... zero7=0; - -#if defined(REINTERPRET_INPUT_AS_3D) - - // The plane (zin) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zin, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), HEIGHT_GEMM3D, DEPTH_GEMM3D, lhs_cross_plane_pad, lhs_stride_y); - - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply lhs_stride_z by DEPTH_GEMM3D - lhs_offset += z * lhs_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_INPUT_AS_3D) - - // Add offset for batched GEMM - lhs_offset += z * lhs_stride_z; - -#endif // defined(REINTERPRET_INPUT_AS_3D) - - // Initialize the accumulators - REPEAT_VAR_INIT_TO_CONST(M0, VEC_DATA_TYPE(DATA_TYPE, N0), c, 0); //VEC_DATA_TYPE(DATA_TYPE, N0) c0=0,c1=0,c2=0,... c(N0-1)=0; - - int i = 0; - for(; i <= (K - K0); i += K0) - { - // Supported cases (M0, K0): - // 1,2 - 1,3 - 1,4 - 1,8 - 1,16 - // 2,2 - 2,3 - 2,4 - 2,8 - 2,16 - // 3,2 - 3,3 - 3,4 - 3,8 - 3,16 - // 4,2 - 4,3 - 4,4 - 4,8 - 4,16 - // 5,2 - 5,3 - 5,4 - 5,8 - 5,16 - // 6,2 - 6,3 - 6,4 - 6,8 - 6,16 - // 7,2 - 7,3 - 7,4 - 7,8 - 7,16 - // 8,2 - 8,3 - 8,4 - 8,8 - 8,16 - // Load values from LHS matrix - LOAD_BLOCK(M0, K0, DATA_TYPE, a, lhs_ptr, lhs_offset, lhs_stride_y, zin); - - VEC_DATA_TYPE(DATA_TYPE, N0) - b0; - - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 0 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(0, a, b0, c); - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 1 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(1, a, b0, c); -#if K0 > 2 - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 2 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(2, a, b0, c); -#endif // K0 > 2 -#if K0 > 3 - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 3 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(3, a, b0, c); -#endif // K0 > 3 -#if K0 > 4 - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 4 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(4, a, b0, c); - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 5 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(5, a, b0, c); - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 6 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(6, a, b0, c); - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 7 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(7, a, b0, c); -#endif // K0 > 4 -#if K0 > 8 - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 8 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(8, a, b0, c); - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 9 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(9, a, b0, c); - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 10 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(A, a, b0, c); - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 11 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(B, a, b0, c); - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 12 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(C, a, b0, c); - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 13 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(D, a, b0, c); - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 14 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(E, a, b0, c); - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 15 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(F, a, b0, c); -#endif // K0 > 8 - - lhs_offset += K0 * sizeof(DATA_TYPE); - rhs_offset += K0 * RHS_STEP_X * RHS_STEP_LOOP * sizeof(DATA_TYPE); - } - - // Left-over accumulations - for(; i < K; ++i) - { - // Load values from LHS matrix - VEC_DATA_TYPE(DATA_TYPE, 2) - a0 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 0 * lhs_stride_y + zin0)); -#if M0 > 1 - VEC_DATA_TYPE(DATA_TYPE, 2) - a1 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 1 * lhs_stride_y + zin1)); -#endif // M0 > 1 -#if M0 > 2 - VEC_DATA_TYPE(DATA_TYPE, 2) - a2 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 2 * lhs_stride_y + zin2)); -#endif // M0 > 2 -#if M0 > 3 - VEC_DATA_TYPE(DATA_TYPE, 2) - a3 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 3 * lhs_stride_y + zin3)); -#endif // M0 > 3 -#if M0 > 4 - VEC_DATA_TYPE(DATA_TYPE, 2) - a4 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 4 * lhs_stride_y + zin4)); -#endif // M0 > 4 -#if M0 > 5 - VEC_DATA_TYPE(DATA_TYPE, 2) - a5 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 5 * lhs_stride_y + zin5)); -#endif // M0 > 5 -#if M0 > 6 - VEC_DATA_TYPE(DATA_TYPE, 2) - a6 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 6 * lhs_stride_y + zin6)); -#endif // M0 > 6 -#if M0 > 7 - VEC_DATA_TYPE(DATA_TYPE, 2) - a7 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 7 * lhs_stride_y + zin7)); -#endif // M0 > 7 - - VEC_DATA_TYPE(DATA_TYPE, N0) - b0; - - b0 = VLOAD(N0)(0, (__global DATA_TYPE *)(rhs_ptr + rhs_offset + 0 * RHS_STEP_X * sizeof(DATA_TYPE))); - VFMA_M0xN0(0, a, b0, c); - - lhs_offset += sizeof(DATA_TYPE); - rhs_offset += RHS_STEP_X * sizeof(DATA_TYPE); - } - - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)) + (COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * dst_stride_y); - - REPEAT_VAR_INIT_TO_CONST(8, uint, zout, 0); //uint zout0=0,zout1=0,zout2=0,... zout7=0; - -#if defined(REINTERPRET_OUTPUT_AS_3D) - // The plane (zout) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zout, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), HEIGHT_GEMM3D, DEPTH_GEMM3D, dst_cross_plane_pad, dst_stride_y); - - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply dst_stride_z by DEPTH_GEMM3D - dst_addr += z * dst_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_OUTPUT_AS_3D) - - // Add offset for batched GEMM - dst_addr += z * dst_stride_z; - -#endif // defined(REINTERPRET_OUTPUT_AS_3D) - - // Multiply by the weight of matrix-matrix product and store the result -#if defined(ALPHA) - SCALE_BLOCK(M0, DATA_TYPE, c, ALPHA); -#endif // defined(ALPHA) - - // Add beta*bias -#if defined(BETA) -#if defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)); - - LOAD_BLOCK_BOUNDARY_AWARE(1, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, 1, PARTIAL_STORE_N0, false, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(1, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias[broadcasted] - ADD_BLOCK_BROADCAST(M0, c, bias0); - -#else // defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)) + (COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * bias_stride_y) + z * bias_stride_z; - - LOAD_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(M0, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias - ADD_BLOCK(M0, c, bias); - -#endif // defined(BROADCAST_BIAS) -#endif // defined(BETA) - - // c = act(c) - POST_OP1_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - // c = c + eltwise_operand (mix-precision, broadcast, boundary aware) - POST_OP2_ELTWISE_OP(P2_ELTWISE_OP, M0, N0, c, eltwise_operand, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), DATA_TYPE, DATA_TYPE_ACCUMULATOR, zero, 1, PARTIAL_STORE_N0, false, cond_x); - // c = act(c) - POST_OP3_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - - // Store output block - STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, c, dst_addr, dst_stride_y, zout, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - -#undef RHS_BLOCK_SIZE -#undef RHS_OFFSET_X -#undef RHS_STEP_X -#undef RHS_STEP_LOOP -} -#endif // defined(GEMM_MM_RESHAPED_ONLY_RHS_NT_POST_ACT_ELTWISE_OP_ACT) - -#if defined(OPENCL_IMAGE_SUPPORT) && defined(GEMM_MM_RESHAPED_ONLY_RHS_NT_TEXTURE_POST_ACT_ELTWISE_OP_ACT) -/** This OpenCL kernel computes the matrix multiplication between 2 matrices plus 3 post ops. The RHS matrix is stored in OpenCL image object. - * Post op 1: activation (optional) - * Post op 2: elementwise op - * Post op 3: activation (optional) - * - * @note (Optional) -DP1_ACTIVATION_TYPE, -DP1_ACTIVATION_A_VAL, -DP1_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * @note (Required) -DP2_ELTWISE_OP: The (binary) elementwise post op to perform - * @note (Required) -DP2_ELTWISE_ARG1_HEIGHT: The height (Y dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Required) -DP2_ELTWISE_ARG1_WIDTH: The width (X dimension) of the eltwise operand matrix of the eltwise post op at slot 2 - * @note (Optional) -DP3_ACTIVATION_TYPE, -DP3_ACTIVATION_A_VAL, -DP3_ACTIVATION_B_VAL: The activation type, alpha and beta values of the activation post op at slot 3 - * - * All parameters are similarly defined in kernel gemm_mm_reshaped_only_rhs_nt_texture, with these additions: - * - * @param[in] eltwise_operand_ptr Pointer to the eltwise operand matrix. Supported data type: F16/F32 - * @param[in] eltwise_operand_stride_x Stride of the eltwise operand matrix in X dimension (in bytes) - * @param[in] eltwise_operand_step_x eltwise_operand_stride_x * number of elements along X processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_y Stride of the eltwise operand matrix in Y dimension (in bytes) - * @param[in] eltwise_operand_step_y eltwise_operand_stride_y * number of elements along Y processed per workitem(in bytes) - * @param[in] eltwise_operand_stride_z Stride of the eltwise operand tensor in Z dimension (in bytes) - * @param[in] M Number of rows in LHS matrix not reshaped. - * @param[in] N Number of columns in RHS matrix not reshaped. - * @param[in] K Number of columns in LHS matrix and rows in RHS matrix not reshaped. - */ -__kernel void gemm_mm_reshaped_only_rhs_nt_texture_post_act_eltwise_op_act(IMAGE_DECLARATION(lhs), - __read_only image2d_t rhs_img, -#if defined(BETA) - IMAGE_DECLARATION(bias), -#endif // defined(BETA) - IMAGE_DECLARATION(dst), - // Post Op arguments - IMAGE_DECLARATION(eltwise_operand), - uint lhs_stride_z, - uint rhs_stride_z, -#if defined(BETA) - uint bias_stride_z, -#endif //defined(BETA) - uint dst_stride_z, - uint eltwise_operand_stride_z -#if defined(REINTERPRET_INPUT_AS_3D) - , - uint lhs_cross_plane_pad -#endif // REINTERPRET_INPUT_AS_3D -#if defined(REINTERPRET_OUTPUT_AS_3D) - , - uint dst_cross_plane_pad -#endif // REINTERPRET_OUTPUT_AS_3D - , - const int M, - const int N, - const int K) -{ - // Pixel unit -#define PIXEL_UNIT CONVERT_VECTOR_SIZE_TO_PIXEL_UNIT(N0) - - // Block size -#define RHS_BLOCK_SIZE ((K0) * (PIXEL_UNIT)) - - // RHS offset and step X -#if defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (PIXEL_UNIT) -#define RHS_STEP_X ((PIXEL_UNIT) * (H0)) -#define RHS_STEP_LOOP (1) -#else // defined(RHS_INTERLEAVE) -#define RHS_OFFSET_X (RHS_BLOCK_SIZE) -#define RHS_STEP_X (PIXEL_UNIT) -#define RHS_STEP_LOOP (H0) -#endif // defined(RHS_INTERLEAVE) - - uint x = get_global_id(0); - uint y = get_global_id(1); - uint z = get_global_id(2); - - const bool cond_y = y == 0; - const bool cond_x = ((x + 1) * N0 >= N); - -#if defined(DUMMY_WORK_ITEMS) - if((x * N0 >= N) || (y * M0 >= M)) - { - return; - } -#endif // defined(DUMMY_WORK_ITEMS) - - // Compute LHS matrix address - uint lhs_offset = lhs_offset_first_element_in_bytes + COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * (uint)lhs_stride_y; - -#if defined(MATRIX_B_DEPTH) - // Do not slide matrix B if the matrix B has 3 dimensions and matrix A more than 3 - const uint z_rhs = (z % MATRIX_B_DEPTH); -#else // defined(MATRIX_B_DEPTH) - const uint z_rhs = z; -#endif // defined(MATRIX_B_DEPTH) - - // Compute RHS matrix coordinates - uint x_rhs = (x % H0) * (uint)RHS_OFFSET_X; - const uint y_rhs = (x / (uint)H0) + z_rhs * RHS_HEIGHT; - - REPEAT_VAR_INIT_TO_CONST(8, uint, zin, 0); - REPEAT_VAR_INIT_TO_CONST(16, uint, zero, 0); - -#if defined(REINTERPRET_INPUT_AS_3D) - - // The plane (zin) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zin, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), HEIGHT_GEMM3D, DEPTH_GEMM3D, lhs_cross_plane_pad, lhs_stride_y); - - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply lhs_stride_z by DEPTH_GEMM3D - lhs_offset += z * lhs_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_INPUT_AS_3D) - - // Add offset for batched GEMM - lhs_offset += z * lhs_stride_z; - -#endif // defined(REINTERPRET_INPUT_AS_3D) - - // Initialize the accumulators - REPEAT_VAR_INIT_TO_CONST(M0, VEC_DATA_TYPE(DATA_TYPE, N0), c, 0); - - int i = 0; - for(; i <= (K - K0); i += K0) - { - // Load values from LHS matrix - LOAD_BLOCK(M0, K0, DATA_TYPE, a, lhs_ptr, lhs_offset, lhs_stride_y, zin); - - VEC_DATA_TYPE(DATA_TYPE, N0) - b0; - - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 0 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(0, a, b0, c); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 1 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(1, a, b0, c); -#if K0 > 2 - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 2 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(2, a, b0, c); -#endif // K0 > 2 -#if K0 > 3 - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 3 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(3, a, b0, c); -#endif // K0 > 3 -#if K0 > 4 - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 4 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(4, a, b0, c); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 5 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(5, a, b0, c); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 6 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(6, a, b0, c); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 7 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(7, a, b0, c); -#endif // K0 > 4 -#if K0 > 8 - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 8 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(8, a, b0, c); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 9 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(9, a, b0, c); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 10 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(A, a, b0, c); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 11 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(B, a, b0, c); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 12 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(C, a, b0, c); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 13 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(D, a, b0, c); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 14 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(E, a, b0, c); - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 15 * RHS_STEP_X), (y_rhs)); - VFMA_M0xN0(F, a, b0, c); -#endif // K0 > 8 - - lhs_offset += K0 * sizeof(DATA_TYPE); - x_rhs += K0 * RHS_STEP_X * RHS_STEP_LOOP; - } - - // Left-over accumulations - for(; i < K; ++i) - { - // Load values from LHS matrix - VEC_DATA_TYPE(DATA_TYPE, 2) - a0 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 0 * lhs_stride_y + zin0)); -#if M0 > 1 - VEC_DATA_TYPE(DATA_TYPE, 2) - a1 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 1 * lhs_stride_y + zin1)); -#endif // M0 > 1 -#if M0 > 2 - VEC_DATA_TYPE(DATA_TYPE, 2) - a2 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 2 * lhs_stride_y + zin2)); -#endif // M0 > 2 -#if M0 > 3 - VEC_DATA_TYPE(DATA_TYPE, 2) - a3 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 3 * lhs_stride_y + zin3)); -#endif // M0 > 3 -#if M0 > 4 - VEC_DATA_TYPE(DATA_TYPE, 2) - a4 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 4 * lhs_stride_y + zin4)); -#endif // M0 > 4 -#if M0 > 5 - VEC_DATA_TYPE(DATA_TYPE, 2) - a5 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 5 * lhs_stride_y + zin5)); -#endif // M0 > 5 -#if M0 > 6 - VEC_DATA_TYPE(DATA_TYPE, 2) - a6 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 6 * lhs_stride_y + zin6)); -#endif // M0 > 6 -#if M0 > 7 - VEC_DATA_TYPE(DATA_TYPE, 2) - a7 = *((__global DATA_TYPE *)(lhs_ptr + lhs_offset + 7 * lhs_stride_y + zin7)); -#endif // M0 > 7 - - VEC_DATA_TYPE(DATA_TYPE, N0) - b0; - b0 = READ_IMAGE2D(DATA_TYPE, PIXEL_UNIT, rhs_img, (x_rhs + 0 * RHS_STEP_X), (y_rhs)); - - VFMA_M0xN0(0, a, b0, c); - - lhs_offset += sizeof(DATA_TYPE); - x_rhs += RHS_STEP_X; - } - - __global uchar *dst_addr = dst_ptr + dst_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)) + (COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * dst_stride_y); - - REPEAT_VAR_INIT_TO_CONST(8, uint, zout, 0); //uint zout0=0,zout1=0,zout2=0,... zout7=0; - -#if defined(REINTERPRET_OUTPUT_AS_3D) - // The plane (zout) is calculated dividing M (y * M0) by HEIGHT_GEMM3D - CALCULATE_Z_OFFSET(M0, uint, zout, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), HEIGHT_GEMM3D, DEPTH_GEMM3D, dst_cross_plane_pad, dst_stride_y); - - // Add offset for batched GEMM. The batches will be in the fourth dimension and for this reason we - // multiply dst_stride_z by DEPTH_GEMM3D - dst_addr += z * dst_stride_z * DEPTH_GEMM3D; - -#else // defined(REINTERPRET_OUTPUT_AS_3D) - - // Add offset for batched GEMM - dst_addr += z * dst_stride_z; - -#endif // defined(REINTERPRET_OUTPUT_AS_3D) - - // Multiply by the weight of matrix-matrix product and store the result -#if defined(ALPHA) - SCALE_BLOCK(M0, DATA_TYPE, c, ALPHA); -#endif // defined(ALPHA) - - // Add beta*bias -#if defined(BETA) -#if defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)); - - LOAD_BLOCK_BOUNDARY_AWARE(1, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, 1, PARTIAL_STORE_N0, false, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(1, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias[broadcasted] - ADD_BLOCK_BROADCAST(M0, c, bias0); - -#else // defined(BROADCAST_BIAS) - __global uchar *bias_addr = bias_ptr + bias_offset_first_element_in_bytes + (x * (uint)N0 * sizeof(DATA_TYPE)) + (COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0) * bias_stride_y) + z * bias_stride_z; - - LOAD_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, bias, bias_addr, 0, bias_stride_y, zero, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - -#ifndef UNIT_BETA - SCALE_BLOCK(M0, DATA_TYPE, bias, BETA); -#endif // UNIT_BIAS - - // c = c + bias - ADD_BLOCK(M0, c, bias); - -#endif // defined(BROADCAST_BIAS) -#endif // defined(BETA) - - // c = act(c) - POST_OP1_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - // c = c + eltwise_operand (mix-precision, broadcast, boundary aware) - POST_OP2_ELTWISE_OP(P2_ELTWISE_OP, M0, N0, c, eltwise_operand, COMPUTE_M0_START_ROW(y, M0, PARTIAL_STORE_M0), DATA_TYPE, DATA_TYPE_ACCUMULATOR, zero, 1, PARTIAL_STORE_N0, false, cond_x); - // c = act(c) - POST_OP3_ACTIVATION_OPTIONAL(M0, DATA_TYPE, DATA_TYPE_ACCUMULATOR, N0, c); - - // Store output block - STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, c, dst_addr, dst_stride_y, zout, PARTIAL_STORE_M0, PARTIAL_STORE_N0, cond_y, cond_x); - -#undef RHS_BLOCK_SIZE -#undef RHS_OFFSET_X -#undef RHS_STEP_X -#undef RHS_STEP_LOOP -} -#endif // defined(OPENCL_IMAGE_SUPPORT) && defined(GEMM_MM_RESHAPED_ONLY_RHS_NT_TEXTURE_POST_ACT_ELTWISE_OP_ACT) -#endif // defined(P2_ELTWISE_OP) && defined(P2_ELTWISE_ARG1_HEIGHT) && defined(P2_ELTWISE_ARG1_WIDTH) -#endif // defined(M0) && defined(N0) && defined(K0) && defined(H0) && defined(DATA_TYPE) diff --git a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/fp_elementwise_op_helpers.h b/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/fp_elementwise_op_helpers.h deleted file mode 100644 index b584251c2a..0000000000 --- a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/fp_elementwise_op_helpers.h +++ /dev/null @@ -1,274 +0,0 @@ -/* - * Copyright (c) 2021 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ -#include "helpers.h" - -/** (EXPERIMENTAL_POST_OPS) Macros for (binary) elementwise operations */ - -/** List of (binary) elementwise operators, accounting for the argument position of argument X - * @note X_Pos denotes the position of argument X. e.g. X_POS_0 means X is in the first place whereas X_POS_1 means X is in the second place - * @name elementwise_post_ops - * @{ - */ -#if defined(N0) && !defined(VEC_SIZE) -#define VEC_SIZE N0 -#endif // defined(N0) && !defined(VEC_SIZE) - -#if defined(VEC_SIZE) && defined(DATA_TYPE) - -#define ADD_X_POS_0(x, y) (x) + (y) -#define SUB_X_POS_0(x, y) (x) - (y) -#define MAX_X_POS_0(x, y) max(x, y) -#define MIN_X_POS_0(x, y) min(x, y) -#define SQUARED_DIFF_X_POS_0(x, y) (x - y) * (x - y) -#define POWER_X_POS_0(x, y) pow(x, y) -#if VEC_SIZE == 1 -#define PRELU_X_POS_0(x, y) (x > 0 ? x : x * y) -#else // VEC_SIZE == 1 - -#if defined(MIXED_PRECISION) -#define PRELU_X_POS_0(x, y) (select(y * x, x, CONVERT((x > (DATA_TYPE_ACCUMULATOR)0), SELECT_VEC_DATA_TYPE(DATA_TYPE_ACCUMULATOR, VEC_SIZE)))) -#else // MIXED_PRECISION -#define PRELU_X_POS_0(x, y) (select(y * x, x, CONVERT((x > (DATA_TYPE)0), SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)))) -#endif // MIXED_PRECISION - -#endif // VEC_SIZE == 1 -#define DIV_X_POS_0(x, y) (x / y) -#define AND_X_POS_0(x, y) (CONVERT((x && y), VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)) & ((VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))1)) -#define OR_X_POS_0(x, y) (CONVERT((x || y), VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)) & ((VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE))1)) - -#define ADD_X_POS_1(x, y) ADD_X_POS_0(x, y) -#define SUB_X_POS_1(x, y) (y) - (x) -#define MAX_X_POS_1(x, y) MAX_X_POS_0(x, y) -#define MIN_X_POS_1(x, y) MIN_X_POS_0(x, y) -#define SQUARED_DIFF_X_POS_1(x, y) SQUARED_DIFF_X_POS_0(x, y) -#define POWER_X_POS_1(x, y) pow(y, x) -#if VEC_SIZE == 1 -#define PRELU_X_POS_1(x, y) (y > 0 ? y : y * x) -#else // VEC_SIZE == 1 - -#if defined(MIXED_PRECISION) -#define PRELU_X_POS_1(x, y) (select(x * y, y, CONVERT((y > (DATA_TYPE_ACCUMULATOR)0), SELECT_VEC_DATA_TYPE(DATA_TYPE_ACCUMULATOR, VEC_SIZE)))) -#else // MIXED_PRECISION -#define PRELU_X_POS_1(x, y) (select(x * y, y, CONVERT((y > (DATA_TYPE)0), SELECT_VEC_DATA_TYPE(DATA_TYPE, VEC_SIZE)))) -#endif // MIXED_PRECISION - -#endif // VEC_SIZE == 1 -#define DIV_X_POS_1(x, y) (y / x) -#define AND_X_POS_1(x, y) AND_X_POS_0(x, y) -#define OR_X_POS_1(x, y) OR_X_POS_0(x, y) - -// By default use the order of the arguments as they are passed in, ie. _X_POS_0 -#define ADD(x, y) ADD_X_POS_0(x, y) -#define SUB(x, y) SUB_X_POS_0(x, y) -#define MAX(x, y) MAX_X_POS_0(x, y) -#define MIN(x, y) MIN_X_POS_0(x, y) -#define SQUARED_DIFF(x, y) SQUARED_DIFF_X_POS_0(x, y) -#define POWER(x, y) POWER_X_POS_0(x, y) -#define PRELU(x, y) PRELU_X_POS_0(x, y) -#define DIV(x, y) DIV_X_POS_0(x, y) -#define AND(x, y) AND_X_POS_0(x, y) -#define OR(x, y) OR_X_POS_0(x, y) - -#endif // defined(VEC_SIZE) && defined(DATA_TYPE) -/** @} */ // end of group elementwise_post_ops - -/** Performs OPERAND1 = OP(OPERAND1, OPERAND2) - * @name ELTWISE_OP_ROW_n - * - * @param[in] OP The elementwise post op - * @param[in, out] OPERAND1 The basename of the destination and operand 1 variables - * @param[in] OPERAND2 The basename of the operand 2 variables - * @{ - */ -#define ELTWISE_OP_ROW_1(OP, OPERAND1, OPERAND2) \ - OPERAND1##0 = OP(OPERAND1##0, OPERAND2##0); - -#define ELTWISE_OP_ROW_2(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_1(OP, OPERAND1, OPERAND2) \ - OPERAND1##1 = OP(OPERAND1##1, OPERAND2##1); - -#define ELTWISE_OP_ROW_3(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_2(OP, OPERAND1, OPERAND2) \ - OPERAND1##2 = OP(OPERAND1##2, OPERAND2##2); - -#define ELTWISE_OP_ROW_4(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_3(OP, OPERAND1, OPERAND2) \ - OPERAND1##3 = OP(OPERAND1##3, OPERAND2##3); - -#define ELTWISE_OP_ROW_5(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_4(OP, OPERAND1, OPERAND2) \ - OPERAND1##4 = OP(OPERAND1##4, OPERAND2##4); - -#define ELTWISE_OP_ROW_6(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_5(OP, OPERAND1, OPERAND2) \ - OPERAND1##5 = OP(OPERAND1##5, OPERAND2##5); - -#define ELTWISE_OP_ROW_7(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_6(OP, OPERAND1, OPERAND2) \ - OPERAND1##6 = OP(OPERAND1##6, OPERAND2##6); - -#define ELTWISE_OP_ROW_8(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_7(OP, OPERAND1, OPERAND2) \ - OPERAND1##7 = OP(OPERAND1##7, OPERAND2##7); - -#define ELTWISE_OP_ROW_9(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_8(OP, OPERAND1, OPERAND2) \ - OPERAND1##8 = OP(OPERAND1##8, OPERAND2##8); - -#define ELTWISE_OP_ROW_10(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_9(OP, OPERAND1, OPERAND2) \ - OPERAND1##9 = OP(OPERAND1##9, OPERAND2##9); - -#define ELTWISE_OP_ROW_11(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_10(OP, OPERAND1, OPERAND2) \ - OPERAND1##A = OP(OPERAND1##A, OPERAND2##A); - -#define ELTWISE_OP_ROW_12(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_11(OP, OPERAND1, OPERAND2) \ - OPERAND1##B = OP(OPERAND1##B, OPERAND2##B); - -#define ELTWISE_OP_ROW_13(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_12(OP, OPERAND1, OPERAND2) \ - OPERAND1##C = OP(OPERAND1##C, OPERAND2##C); - -#define ELTWISE_OP_ROW_14(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_13(OP, OPERAND1, OPERAND2) \ - OPERAND1##D = OP(OPERAND1##D, OPERAND2##D); - -#define ELTWISE_OP_ROW_15(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_14(OP, OPERAND1, OPERAND2) \ - OPERAND1##E = OP(OPERAND1##E, OPERAND2##E); - -#define ELTWISE_OP_ROW_16(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_15(OP, OPERAND1, OPERAND2) \ - OPERAND1##F = OP(OPERAND1##F, OPERAND2##F); - -/** @} */ // end of group ELTWISE_OP_ROW_n - -/** Performs OPERAND1 = OP(OPERAND1, OPERAND2) - * @name ELTWISE_OP_BLOCK - * - * Supported cases are N=1,2,3,...,16 - * - * @param[in] OP The elementwise post op - * @param[in] N The number of vectors in the block - * @param[in] OPERAND1 The basename of the destination and operand 1 variables - * @param[in] OPERAND2 The basename of the operand 2 variables - * @{ - */ -#define ELTWISE_OP_BLOCK_STR(OP, N, OPERAND1, OPERAND2) ELTWISE_OP_ROW_##N(OP, OPERAND1, OPERAND2) -#define ELTWISE_OP_BLOCK(OP, N, OPERAND1, OPERAND2) ELTWISE_OP_BLOCK_STR(OP, N, OPERAND1, OPERAND2) -/** @} */ // end of group ELTWISE_OP_BLOCK - -/** Performs OPERAND1 = OP(OPERAND1, OPERAND2) with broadcasting - * @name ELTWISE_OP_ROW_BROADCAST_n - * - * @param[in] OP The elementwise post op - * @param[in, out] OPERAND1 The basename of the destination and operand 1 variables - * @param[in] OPERAND2 The basename of the broadcast operand 2 variables - * @{ - */ -#define ELTWISE_OP_ROW_BROADCAST_1(OP, OPERAND1, OPERAND2) \ - OPERAND1##0 = OP(OPERAND1##0, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_2(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_1(OP, OPERAND1, OPERAND2) \ - OPERAND1##1 = OP(OPERAND1##1, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_3(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_2(OP, OPERAND1, OPERAND2) \ - OPERAND1##2 = OP(OPERAND1##2, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_4(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_3(OP, OPERAND1, OPERAND2) \ - OPERAND1##3 = OP(OPERAND1##3, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_5(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_4(OP, OPERAND1, OPERAND2) \ - OPERAND1##4 = OP(OPERAND1##4, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_6(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_5(OP, OPERAND1, OPERAND2) \ - OPERAND1##5 = OP(OPERAND1##5, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_7(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_6(OP, OPERAND1, OPERAND2) \ - OPERAND1##6 = OP(OPERAND1##6, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_8(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_7(OP, OPERAND1, OPERAND2) \ - OPERAND1##7 = OP(OPERAND1##7, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_9(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_8(OP, OPERAND1, OPERAND2) \ - OPERAND1##8 = OP(OPERAND1##8, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_10(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_9(OP, OPERAND1, OPERAND2) \ - OPERAND1##9 = OP(OPERAND1##9, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_11(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_10(OP, OPERAND1, OPERAND2) \ - OPERAND1##A = OP(OPERAND1##A, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_12(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_11(OP, OPERAND1, OPERAND2) \ - OPERAND1##B = OP(OPERAND1##B, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_13(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_12(OP, OPERAND1, OPERAND2) \ - OPERAND1##C = OP(OPERAND1##C, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_14(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_13(OP, OPERAND1, OPERAND2) \ - OPERAND1##D = OP(OPERAND1##D, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_15(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_14(OP, OPERAND1, OPERAND2) \ - OPERAND1##E = OP(OPERAND1##E, OPERAND2); - -#define ELTWISE_OP_ROW_BROADCAST_16(OP, OPERAND1, OPERAND2) \ - ELTWISE_OP_ROW_BROADCAST_15(OP, OPERAND1, OPERAND2) \ - OPERAND1##F = OP(OPERAND1##F, OPERAND2); - -/** @} */ // end of group ELTWISE_OP_ROW_BROADCAST_n - -/** Performs OPERAND1 = OP(OPERAND1, OPERAND2) with broadcasting - * @name ELTWISE_OP_BLOCK_BROADCAST - * @note Only support: - * case 1 broadcast in Y dimension : Operand1 [YxX] + Operand2 [1xX]; - * case 2 broadcast in both Y and X dimensions : Operand1 [YxX] + Operand2 [1x1] (scalar); - * Does NOT support broad cast in X dimension: Operand1 [YxX] + Operand2 [Yx1]; - * - * Supported cases are N=1,2,3,...,16 - * - * @param[in] OP The elementwise post op - * @param[in] N The number of vectors in the block - * @param[in] OPERAND1 The basename of the destination and operand 1 variables - * @param[in] OPERAND2 The basename of the operand 2 variables - * @{ - */ -#define ELTWISE_OP_BLOCK_BROADCAST_STR(OP, N, OPERAND1, OPERAND2) ELTWISE_OP_ROW_BROADCAST_##N(OP, OPERAND1, OPERAND2) -#define ELTWISE_OP_BLOCK_BROADCAST(OP, N, OPERAND1, OPERAND2) ELTWISE_OP_BLOCK_BROADCAST_STR(OP, N, OPERAND1, OPERAND2) -/** @} */ // end of group ELTWISE_OP_BLOCK_BROADCAST
\ No newline at end of file diff --git a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/fp_mixed_precision_helpers.h b/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/fp_mixed_precision_helpers.h deleted file mode 100644 index e107f4452d..0000000000 --- a/src/core/CL/cl_kernels/common/experimental/gemm_fused_post_ops/fp_mixed_precision_helpers.h +++ /dev/null @@ -1,113 +0,0 @@ -/* - * Copyright (c) 2021-2022 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ -#include "common/experimental/gemm_fused_post_ops/fp_elementwise_op_helpers.h" -#include "gemm_helpers.h" -#include "load_store_utility.h" - -/** (EXPERIMENTAL_POST_OPS) Convenience macros for automatically handling mixed precision (fp16 and fp32) operations - * -DMIXED_PRECISION toggles mixed precision mode - */ - -/** Mixed-Precision-Aware Activation Block - * @name MIXED_PRECISION_ACTIVATION_BLOCK - * params N ... B_VAL: same as those in @ref ACTIVATION_BLOCK - * - * @param[in] DATA_TYPE_ACCUMULATR Higher-precision accumulator data type in case of mixed-precision op - * @{ - */ -#if defined(MIXED_PRECISION) -#define MIXED_PRECISION_ACTIVATION_BLOCK(N, ACTIVATION_TYPE, DATA_TYPE, VEC_SIZE, BASENAME, A_VAL, B_VAL, DATA_TYPE_ACCUMULATOR) \ - ACTIVATION_BLOCK(N, ACTIVATION_TYPE, DATA_TYPE_ACCUMULATOR, VEC_SIZE, BASENAME, A_VAL, B_VAL); -#else // defined(MIXED_PRECISION) -#define MIXED_PRECISION_ACTIVATION_BLOCK(N, ACTIVATION_TYPE, DATA_TYPE, VEC_SIZE, BASENAME, A_VAL, B_VAL, DATA_TYPE_ACCUMULATOR) \ - ACTIVATION_BLOCK(N, ACTIVATION_TYPE, DATA_TYPE, VEC_SIZE, BASENAME, A_VAL, B_VAL); -#endif // defined(MIXED_PRECISION) -/** @} */ // end of group MIXED_PRECISION_ACTIVATION_BLOCK - -/** Mixed-Precision-Aware Elementwise Op Block - * Performs OPERAND1 = OP(OPERAND1, OPERAND2) - * @name MIXED_PRECISION_ELTWISE_OP_BLOCK - * - * @param[in] OP The elementwise post op - * @param[in] M0 The number of consecutive rows - * @param[in] N0 The number of consecutive columns - * @param[in] OPERAND1 The basename of the first and result operand variables - * @param[in] OPERAND2 The basename of the second operand variables - * @param[in] DATA_TYPE_ACCUMULATR Higher-precision accumulator data type in case of mixed-precision op - * @param[in] CONVERTED_OPERAND2 The basename of the second operand variables converted to higher-precision in case of mixed-precision op - * @{ - */ -#if defined(MIXED_PRECISION) -#define MIXED_PRECISION_ELTWISE_OP_BLOCK(OP, M0, N0, OPERAND1, OPERAND2, DATA_TYPE_ACCUMULATOR, CONVERTED_OPERAND2) \ - CONVERT_BLOCK(M0, N0, DATA_TYPE_ACCUMULATOR, OPERAND2, CONVERTED_OPERAND2); \ - ELTWISE_OP_BLOCK(OP, M0, OPERAND1, CONVERTED_OPERAND2); -#else // defined(MIXED_PRECISION) -#define MIXED_PRECISION_ELTWISE_OP_BLOCK(OP, M0, N0, OPERAND1, OPERAND2, DATA_TYPE_ACCUMULATOR, CONVERTED_OPERAND2) \ - ELTWISE_OP_BLOCK(OP, M0, OPERAND1, OPERAND2); -#endif // defined(MIXED_PRECISION) -/** @} */ // end of group MIXED_PRECISION_ELTWISE_OP_BLOCK - -/** Mixed-Precision-Aware Elementwise Op Broadcast Block - * Performs OPERAND1 = OP(OPERAND1, OPERAND2) - * @name MIXED_PRECISION_ELTWISE_OP_BLOCK_BROADCAST - * @note Only support: - * case 1 broadcast in Y dimension : Operand1 [YxX] + Operand2 [1xX]; this means @p N0 > 1 - * case 2 broadcast in both Y and X dimensions : Operand1 [YxX] + Operand2 [1x1] (scalar) ; this means @p N0 == 1 - * Does NOT support broad cast in X dimension: Operand1 [YxX] + Operand2 [Yx1]; this means @p M0 should never == 1 - * - * @param[in] OP The elementwise post op - * @param[in] M0 The number of consecutive rows, > 1 - * @param[in] N0 The number of consecutive columns, >= 1 - * @param[in] OPERAND1 The basename of the first and result operand variables - * @param[in] OPERAND2 The basename of the second operand variables - * @param[in] DATA_TYPE_ACCUMULATR Higher-precision accumulator data type in case of mixed-precision op - * @param[in] CONVERTED_OPERAND2 The basename of the second operand variables converted to higher-precision in case of mixed-precision op - * @{ - */ -#if defined(MIXED_PRECISION) -#define MIXED_PRECISION_ELTWISE_OP_BLOCK_BROADCAST(OP, M0, N0, OPERAND1, OPERAND2, DATA_TYPE_ACCUMULATOR, CONVERTED_OPERAND2) \ - CONVERT_BLOCK(1, N0, DATA_TYPE_ACCUMULATOR, OPERAND2, CONVERTED_OPERAND2); \ - ELTWISE_OP_BLOCK_BROADCAST(OP, M0, OPERAND1, CONVERTED_OPERAND2##0); -#else // defined(MIXED_PRECISION) -#define MIXED_PRECISION_ELTWISE_OP_BLOCK_BROADCAST(OP, M0, N0, OPERAND1, OPERAND2, DATA_TYPE_ACCUMULATOR, CONVERTED_OPERAND2) \ - ELTWISE_OP_BLOCK_BROADCAST(OP, M0, OPERAND1, OPERAND2##0); -#endif // defined(MIXED_PRECISION) -/** @} */ // end of group MIXED_PRECISION_ELTWISE_OP_BLOCK_BROADCAST - -/** Mixed-Precision-Aware Boundary-Aware Store Block - * @name MIXED_PRECISION_STORE_BLOCK_BOUNDARY_AWARE - * params M0 ... PARTIAL_COND_X, same as those in STORE_BLOCK_BOUNDARY_AWARE - * - * @param[in] BASENAME_LP The name of the low precision variables, converted from BASENAME, in case of mixed-precision op - * @{ - */ -#if defined(MIXED_PRECISION) -#define MIXED_PRECISION_STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X, BASENAME_LP) \ - CONVERT_BLOCK(M0, N0, DATA_TYPE, BASENAME, BASENAME_LP); \ - STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME_LP, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X); -#else // defined(MIXED_PRECISION) -#define MIXED_PRECISION_STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X, BASENAME_LP) \ - STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, BASENAME, PTR, STRIDE_Y, Z, PARTIAL_STORE_M0, PARTIAL_STORE_N0, PARTIAL_COND_Y, PARTIAL_COND_X); -#endif // defined(MIXED_PRECISION) -/** @} */ // end of group MIXED_PRECISION_STORE_BLOCK_BOUNDARY_AWARE
\ No newline at end of file diff --git a/src/core/CL/cl_kernels/common/gemm.cl b/src/core/CL/cl_kernels/common/gemm.cl index a32301d8e3..0c30c0e626 100644 --- a/src/core/CL/cl_kernels/common/gemm.cl +++ b/src/core/CL/cl_kernels/common/gemm.cl @@ -152,7 +152,6 @@ /** This OpenCL kernel computes the matrix multiplication between 2 matrices. * The LHS matrix is NOT reshaped * The RHS is reshaped with @ref CLGEMMReshapeRHSMatrixKernel and the block K0xN0 is transposed - * @note This kernel is duplicated in /experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped_only_rhs.cl * * @note If the first two dimensions of NDRange have been dispatched with "dummy_work_items" support, the option -DDUMMY_WORK_ITEMS must be passed at compile time. * @note The GEMM's dimensions (M,N and K) must be passed at runtime as kernel parameters. @@ -453,7 +452,6 @@ __kernel void gemm_mm_reshaped_only_rhs_t(IMAGE_DECLARATION(lhs), /** This OpenCL kernel computes the matrix multiplication between 2 matrices. The RHS matrix is stored in OpenCL image * The LHS matrix is NOT reshaped * The RHS is reshaped with @ref CLGEMMReshapeRHSMatrixKernel and the block K0xN0 is transposed - * @note This kernel is duplicated in /experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped_only_rhs.cl * * @note -DOPENCL_IMAGE_SUPPORT must be passed at compile time in order to compile this OpenCL kernel * @note If the first two dimensions of NDRange have been dispatched with "dummy_work_items" support, the option -DDUMMY_WORK_ITEMS must be passed at compile time. @@ -887,7 +885,6 @@ __kernel void gemm_mm_reshaped_only_rhs_t_texture(IMAGE_DECLARATION(lhs), /** This OpenCL kernel computes the matrix multiplication between 2 matrices. * The LHS matrix is NOT reshaped * The RHS is reshaped with @ref CLGEMMReshapeRHSMatrixKernel and the block K0xN0 is NOT transposed - * @note This kernel is duplicated in /experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped_only_rhs.cl * * @note If the first two dimensions of NDRange have been dispatched with "dummy_work_items" support, the option -DDUMMY_WORK_ITEMS must be passed at compile time. * @note The GEMM's dimensions (M,N and K) must be passed at runtime as kernel parameters. @@ -1213,7 +1210,6 @@ __kernel void gemm_mm_reshaped_only_rhs_nt(IMAGE_DECLARATION(lhs), /** This OpenCL kernel computes the matrix multiplication between 2 matrices. * The LHS matrix is NOT reshaped * The RHS is reshaped with @ref CLGEMMReshapeRHSMatrixKernel and the block K0xN0 is NOT transposed - * @note This kernel is duplicated in /experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped_only_rhs.cl * * @note -DOPENCL_IMAGE_SUPPORT must be passed at compile time in order to compile this OpenCL kernel * @note If the first two dimensions of NDRange have been dispatched with "dummy_work_items" support, the option -DDUMMY_WORK_ITEMS must be passed at compile time. @@ -1713,7 +1709,6 @@ __kernel void gemm_mm_reshaped_only_rhs_nt_texture(IMAGE_DECLARATION(lhs), /** This OpenCL kernel computes the matrix multiplication between 2 matrices. * The LHS matrix must be reshaped with @ref CLGEMMReshapeLHSMatrixKernel and the M0xK0 must be NOT transposed * The RHS matrix must be reshaped with @ref CLGEMMReshapeRHSMatrixKernel and the K0xN0 must be transposed - * @note This kernel is duplicated in /experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped.cl * * @note The data type must be passed at compile time using -DDATA_TYPE (e.g. -DDATA_TYPE=float) * @note The data type used for the accumulators must be passed at compile time using -DDATA_TYPE_ACCUMULATOR (e.g. -DDATA_TYPE_ACCUMULATOR=float) @@ -1993,7 +1988,6 @@ __kernel void gemm_mm_reshaped_lhs_nt_rhs_t(IMAGE_DECLARATION(lhs), /** This OpenCL kernel computes the matrix multiplication between 2 matrices. The RHS matrix is stored in OpenCL image object. * The LHS matrix must be reshaped with @ref CLGEMMReshapeLHSMatrixKernel and the M0xK0 must be NOT transposed * The RHS matrix must be reshaped with @ref CLGEMMReshapeRHSMatrixKernel and the K0xN0 must be transposed - * @note This kernel is duplicated in /experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped.cl * * @note -DOPENCL_IMAGE_SUPPORT must be passed at compile time in order to compile this OpenCL kernel * @note The data type must be passed at compile time using -DDATA_TYPE (e.g. -DDATA_TYPE=float) @@ -2380,7 +2374,6 @@ __kernel void gemm_mm_reshaped_lhs_nt_rhs_t_texture(IMAGE_DECLARATION(lhs), /** This OpenCL kernel computes the matrix multiplication between 2 matrices. * The LHS matrix must be reshaped with @ref CLGEMMReshapeLHSMatrixKernel and the M0xK0 must be transposed * The RHS matrix must be reshaped with @ref CLGEMMReshapeRHSMatrixKernel and the K0xN0 must be NOT transposed - * @note This kernel is duplicated in /experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped.cl * * @note LHS_TRANSPOSE should be passed at compile time in order to compile this OpenCL kernel (e.g. -DLHS_TRANSPOSE). * @note If the first two dimensions of NDRange have been dispatched with "dummy_work_items" support, the option -DDUMMY_WORK_ITEMS must be passed at compile time. @@ -2767,7 +2760,6 @@ __kernel void gemm_mm_reshaped_lhs_t_rhs_nt(IMAGE_DECLARATION(lhs), /** This OpenCL kernel computes the matrix multiplication between 2 matrices. The RHS matrix is stored in OpenCL image object. * The LHS matrix must be reshaped with @ref CLGEMMReshapeLHSMatrixKernel and the M0xK0 must be transposed * The RHS matrix must be reshaped with @ref CLGEMMReshapeRHSMatrixKernel and the K0xN0 must be NOT transposed - * @note This kernel is duplicated in /experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_reshaped.cl * * @note -DOPENCL_IMAGE_SUPPORT must be passed at compile time in order to compile this OpenCL kernel * @note LHS_TRANSPOSE should be passed at compile time in order to compile this OpenCL kernel (e.g. -DLHS_TRANSPOSE). @@ -3226,7 +3218,6 @@ __kernel void gemm_mm_reshaped_lhs_t_rhs_nt_texture(IMAGE_DECLARATION(lhs), /** This OpenCL kernel computes the matrix multiplication between 2 matrices. * The LHS matrix is NOT reshaped * The RHS matrix is NOT reshaped - * @note This kernel is duplicated in /experimental/gemm_fused_post_ops/act_eltwise_op_act/gemm_mm_native.cl * * @note If the first two dimensions of NDRange have been dispatched with "dummy_work_items" support, the option -DDUMMY_WORK_ITEMS must be passed at compile time. * @note The GEMM's dimensions (M,N and K) must be passed at runtime as kernel parameters. |