diff options
Diffstat (limited to 'src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components')
7 files changed, 758 insertions, 0 deletions
diff --git a/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClElementwiseAddKernelComponent.cpp b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClElementwiseAddKernelComponent.cpp new file mode 100644 index 0000000000..a44b5faee2 --- /dev/null +++ b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClElementwiseAddKernelComponent.cpp @@ -0,0 +1,92 @@ +/* + * Copyright (c) 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. + */ +#if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION) + +#include "src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClElementwiseAddKernelComponent.h" + +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +ComponentType ClElementwiseAddKernelComponent::get_component_type() const +{ + return ComponentType::Simple; +} + +std::set<std::string> ClElementwiseAddKernelComponent::get_headers_list() const +{ + return std::set<std::string> { "gemm_helpers.h", "repeat.h" }; +} + +std::string ClElementwiseAddKernelComponent::get_component_code() const +{ + std::string code; + return R"_( + //------------------ START KERNEL {{meta_kernel_id}} ELTWISE_ADD --------------------- + // IN_0(Accumulator) {{acc}} + // IN_1(Addend) {{addend}} + + // c = addend + c (mix-precision, broadcast, boundary aware) + { + __global uchar *addend_addr = {{addend}}_ptr + {{addend}}_offset_first_element_in_bytes + (get_global_id(0) * (uint)N0 * sizeof(DATA_TYPE)) + (COMPUTE_M0_START_ROW(g_y, M0, PARTIAL_STORE_M0) * {{addend}}_stride_y) + get_global_id(2) * {{addend}}_stride_z; \ + LOAD_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, addend, addend_addr, 0, {{addend}}_stride_y, g_zero, PARTIAL_LOAD_M0, PARTIAL_LOAD_N0, PARTIAL_COND_Y, PARTIAL_COND_X); \ + MIXED_PRECISION_ELTWISE_OP_BLOCK(ADD_X_POS_0, M0, N0, {{acc}}, addend, DATA_TYPE_ACCUMULATOR, addend_hp); + } + //------------------ END KERNEL {{meta_kernel_id}} ELTWISE_ADD --------------------- + +)_"; +} +ClElementwiseAddKernelComponent::TagLUT ClElementwiseAddKernelComponent::allocate_vars(SharedVarTable &vtable) const +{ + // Determine which argument is the accumulator + Link accumulator; + Link addend; + if(_lhs.group == SharedVarGroup::Automatic) + { + accumulator = _lhs; + addend = _rhs; + } + else if(_rhs.group == SharedVarGroup::Automatic) + { + accumulator = _rhs; + addend = _lhs; + } + else + { + ARM_COMPUTE_ERROR("Invalid elementwise component linking"); + } + return { + { "meta_kernel_id", id() }, + { "acc", vtable.add(accumulator, ClKernelArgRuntimeDescriptor(accumulator.arg_id, TensorArgType::Image_3D), "add_acc") }, + { "addend", vtable.add(addend, ClKernelArgRuntimeDescriptor(addend.arg_id, TensorArgType::Image_3D), "add_addend") }, + // {"dst", vtable.add(_dst, ClKernelArgRuntimeDescriptor(_dst.arg_id, TensorArgType::Image_3D), "dst")}, // dst is needed for the root version and/or non-inplace version should we need one + }; +} +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute + +#endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
\ No newline at end of file diff --git a/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClElementwiseAddKernelComponent.h b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClElementwiseAddKernelComponent.h new file mode 100644 index 0000000000..c0de4ac9b8 --- /dev/null +++ b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClElementwiseAddKernelComponent.h @@ -0,0 +1,71 @@ +/* + * Copyright (c) 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. + */ +#if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION) + +#ifndef ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_COMPONENTS_CLELEMENTWISEADDKERNELCOMPONENT_H +#define ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_COMPONENTS_CLELEMENTWISEADDKERNELCOMPONENT_H + +#include "src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/Common.h" + +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +class ClElementwiseAddKernelComponent : public IClKernelComponent +{ +public: + ClElementwiseAddKernelComponent(const Link &lhs, const Link &rhs, const Link &dst) + : _lhs{ lhs }, _rhs{ rhs }, _dst{ dst } + { + } + ComponentType get_component_type() const override; + std::set<std::string> get_headers_list() const override; + std::string get_component_code() const override; + + virtual std::vector<Link> get_links() const override + { + return { _lhs, _rhs, _dst }; + } + + virtual TagLUT allocate_vars(SharedVarTable &vtable) const override; + + virtual std::string name() const override + { + return "eltwise_add_" + std::to_string(id()); + } + +private: + Link _lhs{}; + Link _rhs{}; + Link _dst{}; +}; + +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute +#endif // ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_COMPONENTS_CLELEMENTWISEADDKERNELCOMPONENT_H + +#endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
\ No newline at end of file diff --git a/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClGemmNativeKernelComponent.cpp b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClGemmNativeKernelComponent.cpp new file mode 100644 index 0000000000..1521973d55 --- /dev/null +++ b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClGemmNativeKernelComponent.cpp @@ -0,0 +1,341 @@ +/* + * Copyright (c) 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. + */ +#if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION) + +#include "src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClGemmNativeKernelComponent.h" + +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +ComponentType ClGemmNativeKernelComponent::get_component_type() const +{ + return ComponentType::Complex; +} + +std::set<std::string> ClGemmNativeKernelComponent::get_headers_list() const +{ + return std::set<std::string> { "./common/experimental/gemm_fused_post_ops/act_eltwise_op_act/fp_post_ops_act_eltwise_op_act.h", "gemm_helpers.h", "repeat.h" }; +} + +std::string ClGemmNativeKernelComponent::get_additional_macros() const +{ + return R"_( +#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 +)_"; +} + +std::string ClGemmNativeKernelComponent::get_component_code() const +{ + std::string code = R"_( + //------------------ START KERNEL {{meta_kernel_id}} --------------------- + // IN_0(lhs) {{lhs}} + // IN_1(rhs) {{rhs}} + )_"; + + if(!_bias.is_empty()) + { + code += R"_( + // IN_2(bias) {{bias}} + )_"; + } + + code += R"_( + // OUT(dst, accum) {{dst}} + + // Initialize the accumulators + REPEAT_VAR_INIT_TO_CONST(M0, VEC_DATA_TYPE(DATA_TYPE, N0), {{dst}}, 0); //VEC_DATA_TYPE(DATA_TYPE, N0) c0=0,c1=0,c2=0,... c(M0-1)=0; + { +#if defined(DUMMY_WORK_ITEMS) + if((g_x * N0 >= N) || (g_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(g_y, M0, PARTIAL_STORE_M0) * (uint){{lhs}}_stride_y; + + // Compute RHS matrix address + uint rhs_offset = {{rhs}}_offset_first_element_in_bytes + g_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 += (g_z % MATRIX_B_DEPTH) * {{rhs}}_stride_z; +#else // defined(MATRIX_B_DEPTH) + rhs_offset += g_z * {{rhs}}_stride_z; +#endif // defined(MATRIX_B_DEPTH) + + REPEAT_VAR_INIT_TO_CONST(M0, uint, zlhs, 0); + +#if defined(REINTERPRET_INPUT_AS_3D) + // The plane (zlhs) is calculated dividing M (g_y * M0) by HEIGHT_GEMM3D + CALCULATE_Z_OFFSET(M0, uint, zlhs, COMPUTE_M0_START_ROW(g_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 += g_z * {{lhs}}_stride_z * DEPTH_GEMM3D; + +#else // defined(REINTERPRET_INPUT_AS_3D) + + // Add offset for batched GEMM + lhs_offset += g_z * {{lhs}}_stride_z; + +#endif // defined(REINTERPRET_INPUT_AS_3D) + + 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, g_zero); + + RHS_VFMA_M0xN0(0, a, b0, {{dst}}); + RHS_VFMA_M0xN0(1, a, b1, {{dst}}); +#if K0 > 2 + RHS_VFMA_M0xN0(2, a, b2, {{dst}}); +#endif // K0 > 2 +#if K0 > 3 + RHS_VFMA_M0xN0(3, a, b3, {{dst}}); +#endif // K0 > 3 +#if K0 > 4 + RHS_VFMA_M0xN0(4, a, b4, {{dst}}); + RHS_VFMA_M0xN0(5, a, b5, {{dst}}); + RHS_VFMA_M0xN0(6, a, b6, {{dst}}); + RHS_VFMA_M0xN0(7, a, b7, {{dst}}); +#endif // K0 > 4 +#if K0 > 8 + RHS_VFMA_M0xN0(8, a, b8, {{dst}}); + RHS_VFMA_M0xN0(9, a, b9, {{dst}}); + RHS_VFMA_M0xN0(A, a, bA, {{dst}}); + RHS_VFMA_M0xN0(B, a, bB, {{dst}}); + RHS_VFMA_M0xN0(C, a, bC, {{dst}}); + RHS_VFMA_M0xN0(D, a, bD, {{dst}}); + RHS_VFMA_M0xN0(E, a, bE, {{dst}}); + RHS_VFMA_M0xN0(F, a, bF, {{dst}}); +#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, {{dst}}); + + lhs_offset += sizeof(DATA_TYPE); + rhs_offset += {{rhs}}_stride_y; + } + + // Multiply by the weight of matrix-matrix product and store the result +#if defined(ALPHA) + SCALE_BLOCK(M0, DATA_TYPE, {{dst}}, ALPHA); +#endif // defined(ALPHA) + )_"; + + if(!_bias.is_empty()) + { + code += R"_( + // Add beta*bias +#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, g_zero); + +#ifndef UNIT_BETA + SCALE_BLOCK(1, DATA_TYPE, bias, BETA); +#endif // UNIT_BIAS + + // c = c + bias[broadcasted] + ADD_BLOCK_BROADCAST(M0, {{dst}}, bias0); + +#else // defined(BROADCAST_BIAS) + __global uchar *bias_addr = {{bias}}_ptr + {{bias}}_offset_first_element_in_bytes + (g_x * (uint)N0 * sizeof(DATA_TYPE)) + (COMPUTE_M0_START_ROW(g_y, M0, + PARTIAL_STORE_M0) + * {{bias}}_stride_y) + + g_z * {{bias}}_stride_z; + + LOAD_BLOCK(M0, N0, DATA_TYPE, bias, bias_addr, 0, {{bias}}_stride_y, g_zero); + +#ifndef UNIT_BETA + SCALE_BLOCK(M0, DATA_TYPE, bias, BETA); +#endif // UNIT_BIAS + + // c = c + bias + ADD_BLOCK(M0, {{dst}}, bias); + +#endif // defined(BROADCAST_BIAS) + )_"; + } + + code += R"_( + } + //------------------ END KERNEL {{meta_kernel_id}} --------------------- + )_"; + return code.c_str(); +} + +ClGemmNativeKernelComponent::TagLUT ClGemmNativeKernelComponent::allocate_vars(SharedVarTable &vtable) const +{ + TagLUT lut{}; + + lut["meta_kernel_id"] = id(); + lut["lhs"] = vtable.add(_lhs, ClKernelArgRuntimeDescriptor(_lhs.arg_id, TensorArgType::Image_3D), "lhs"); + lut["rhs"] = vtable.add(_rhs, ClKernelArgRuntimeDescriptor(_rhs.arg_id, TensorArgType::Image_3D), "rhs"); + if(!_bias.is_empty()) // optional bias + { + lut["bias"] = vtable.add(_bias, ClKernelArgRuntimeDescriptor(_bias.arg_id, TensorArgType::Image_3D), "bias"); + } + lut["dst"] = vtable.add(_dst, ClKernelArgRuntimeDescriptor(_dst.arg_id, TensorArgType::Image_3D), "dst"); + return lut; +} +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute + +#endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
\ No newline at end of file diff --git a/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClGemmNativeKernelComponent.h b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClGemmNativeKernelComponent.h new file mode 100644 index 0000000000..38f007c07c --- /dev/null +++ b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClGemmNativeKernelComponent.h @@ -0,0 +1,74 @@ +/* + * Copyright (c) 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. + */ +#if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION) + +#ifndef ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_COMPONENTS_CLGEMMNATIVEKERNELCOMPONENT_H +#define ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_COMPONENTS_CLGEMMNATIVEKERNELCOMPONENT_H + +#include "src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/Common.h" + +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +class ClGemmNativeKernelComponent : public IClKernelComponent +{ +public: + ClGemmNativeKernelComponent(const Link &lhs, const Link &rhs, const Link &dst, const Link &bias = Link{}) + : _lhs{ lhs }, _rhs{ rhs }, _bias{ bias }, _dst{ dst } + { + } + ComponentType get_component_type() const override; + std::set<std::string> get_headers_list() const override; + std::string get_additional_macros() const override; + std::string get_component_code() const override; + ClKernelArgList get_args(); + + virtual std::vector<Link> get_links() const override + { + return { _lhs, _rhs, _bias, _dst }; + } + + virtual TagLUT allocate_vars(SharedVarTable &vtable) const override; + + virtual std::string name() const override + { + return "gemm_mm_native_" + std::to_string(id()); + } + +private: + Link _lhs{}; + Link _rhs{}; + Link _bias{}; + Link _dst{}; +}; + +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute +#endif // ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_COMPONENTS_CLGEMMNATIVEKERNELCOMPONENT_H + +#endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
\ No newline at end of file diff --git a/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClKernelComponents.h b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClKernelComponents.h new file mode 100644 index 0000000000..b751ce237f --- /dev/null +++ b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClKernelComponents.h @@ -0,0 +1,35 @@ +/* + * Copyright (c) 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. + */ +#if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION) + +#ifndef ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_CLKERNELCOMPONENTS_H +#define ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_CLKERNELCOMPONENTS_H + +#include "src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClElementwiseAddKernelComponent.h" +#include "src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClGemmNativeKernelComponent.h" +#include "src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClStoreKernelComponents.h" + +#endif //ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_CLKERNELCOMPONENTS_H + +#endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
\ No newline at end of file diff --git a/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClStoreKernelComponents.cpp b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClStoreKernelComponents.cpp new file mode 100644 index 0000000000..430fafb89f --- /dev/null +++ b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClStoreKernelComponents.cpp @@ -0,0 +1,76 @@ +/* + * Copyright (c) 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. + */ +#if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION) + +#include "src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClStoreKernelComponents.h" + +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +ComponentType ClStoreBlockBoundaryAwareKernelComponent::get_component_type() const +{ + return ComponentType::Store; +} + +std::string ClStoreBlockBoundaryAwareKernelComponent::get_component_code() const +{ + return R"_( + //------------------ START KERNEL {{meta_kernel_id}} STORE --------------------- + + __global uchar *dst_addr = {{dst}}_ptr + {{dst}}_offset_first_element_in_bytes + (g_x * (uint)N0 * sizeof(DATA_TYPE)) + (COMPUTE_M0_START_ROW(g_y, M0, PARTIAL_STORE_M0) * {{dst}}_stride_y); + +#if defined(REINTERPRET_OUTPUT_AS_3D) + // 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 += g_z * {{dst}}_stride_z * DEPTH_GEMM3D; + +#else // defined(REINTERPRET_OUTPUT_AS_3D) + + // Add offset for batched GEMM + dst_addr += g_z * {{dst}}_stride_z; + +#endif // defined(REINTERPRET_OUTPUT_AS_3D) + + STORE_BLOCK_BOUNDARY_AWARE(M0, N0, DATA_TYPE, {{src}}, dst_addr, {{dst}}_stride_y, g_zout, PARTIAL_STORE_M0, PARTIAL_STORE_N0, g_cond_y, g_cond_x); + + //------------------ END KERNEL {{meta_kernel_id}} STORE --------------------- + +)_"; +} +ClStoreBlockBoundaryAwareKernelComponent::TagLUT ClStoreBlockBoundaryAwareKernelComponent::allocate_vars(SharedVarTable &vtable) const +{ + return { + { "meta_kernel_id", id() }, + { "src", vtable.add(_src, ClKernelArgRuntimeDescriptor(_src.arg_id, TensorArgType::Image_3D), "src") }, + { "dst", vtable.add(_dst, ClKernelArgRuntimeDescriptor(_dst.arg_id, TensorArgType::Image_3D), "dst") }, + }; +} +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute + +#endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
\ No newline at end of file diff --git a/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClStoreKernelComponents.h b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClStoreKernelComponents.h new file mode 100644 index 0000000000..f0d01d30a9 --- /dev/null +++ b/src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/components/ClStoreKernelComponents.h @@ -0,0 +1,69 @@ +/* + * Copyright (c) 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. + */ +#if defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION) + +#ifndef ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_COMPONENTS_CLSTOREKERNELCOMPONENTS_H +#define ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_COMPONENTS_CLSTOREKERNELCOMPONENTS_H + +#include "src/core/experimental/dynamic_fusion/ClKernelBuildingImpl/Common.h" + +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +class ClStoreBlockBoundaryAwareKernelComponent : public IClKernelComponent +{ +public: + ClStoreBlockBoundaryAwareKernelComponent(const Link &src, const Link &dst) + : _src{ src }, _dst{ dst } + { + } + ComponentType get_component_type() const override; + std::string get_component_code() const override; + + virtual std::vector<Link> get_links() const override + { + return { _src, _dst }; + } + + virtual TagLUT allocate_vars(SharedVarTable &vtable) const override; + + virtual std::string name() const override + { + return ""; + } + +private: + Link _src{}; + Link _dst{}; +}; + +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute +#endif // ARM_COMPUTE_EXPERIMENTAL_DYNAMICFUSION_IMPL_COMPONENTS_CLSTOREKERNELCOMPONENTS_H + +#endif // defined(ENABLE_EXPERIMENTAL_DYNAMIC_FUSION)
\ No newline at end of file |