diff options
Diffstat (limited to 'src')
23 files changed, 1020 insertions, 29 deletions
diff --git a/src/dynamic_fusion/runtime/gpu/cl/ClKernelRuntime.cpp b/src/dynamic_fusion/runtime/gpu/cl/ClKernelRuntime.cpp index b3ec39362c..6a57482bb2 100644 --- a/src/dynamic_fusion/runtime/gpu/cl/ClKernelRuntime.cpp +++ b/src/dynamic_fusion/runtime/gpu/cl/ClKernelRuntime.cpp @@ -59,6 +59,8 @@ void ClKernelRuntime::configure(const ClCompileContext &compile_ctx, const GpuKe inline void ClKernelRuntime::add_tensor_argument(unsigned int &idx, const GpuKernelArgumentInfo &arg, const ICLTensor *tensor, const Window &arg_slice, std::vector<cl::Image2D> &cl_images) { + ARM_COMPUTE_ERROR_ON_NULLPTR(tensor); + switch(arg.type) { case GpuKernelArgumentInfo::Type::Scalar: @@ -142,6 +144,18 @@ inline void ClKernelRuntime::add_tensor_argument(unsigned int &idx, const GpuKer add_4d_tensor_nhwc_argument(idx, tensor); break; } + case GpuKernelArgumentInfo::Type::Tensor_Special_0: + { + const ITensorInfo *info = tensor->info(); + const Strides &strides = info->strides_in_bytes(); + + _kernel.setArg(idx++, tensor->cl_buffer()); + const size_t dim1xdim2 = info->tensor_shape()[1] * info->tensor_shape()[2]; + _kernel.setArg<cl_int>(idx++, static_cast<int32_t>(dim1xdim2)); + const size_t stride1 = strides[1]; + _kernel.setArg<cl_int>(idx++, static_cast<int32_t>(stride1)); + break; + } default: { ARM_COMPUTE_ERROR("Unsupported"); diff --git a/src/core/CL/ckw/KernelArgumentsHelpers.cpp b/src/dynamic_fusion/runtime/gpu/cl/ckw_driver/GpuCkwKernelArgumentsHelpers.cpp index 3fbdc46a3a..40934cf6e9 100644 --- a/src/core/CL/ckw/KernelArgumentsHelpers.cpp +++ b/src/dynamic_fusion/runtime/gpu/cl/ckw_driver/GpuCkwKernelArgumentsHelpers.cpp @@ -22,11 +22,10 @@ * SOFTWARE. */ -#include "KernelArgumentsHelpers.h" +#include "GpuCkwKernelArgumentsHelpers.h" namespace arm_compute { - void cl_add_tensor_component_argument(cl::Kernel &kernel, unsigned int &idx, ICLTensor *tensor, ckw::TensorComponent component) { ARM_COMPUTE_ERROR_ON(tensor == nullptr); diff --git a/src/core/CL/ckw/KernelArgumentsHelpers.h b/src/dynamic_fusion/runtime/gpu/cl/ckw_driver/GpuCkwKernelArgumentsHelpers.h index b681636c26..206ad2e401 100644 --- a/src/core/CL/ckw/KernelArgumentsHelpers.h +++ b/src/dynamic_fusion/runtime/gpu/cl/ckw_driver/GpuCkwKernelArgumentsHelpers.h @@ -22,8 +22,8 @@ * SOFTWARE. */ -#ifndef ARM_COMPUTE_SRC_CKW_KERNELARGUMENTSHELPERS_H -#define ARM_COMPUTE_SRC_CKW_KERNELARGUMENTSHELPERS_H +#ifndef ACL_SRC_DYNAMIC_FUSION_RUNTIME_GPU_CL_CKW_DRIVER_GPUCKWKERNELARGUMENTSHELPERS +#define ACL_SRC_DYNAMIC_FUSION_RUNTIME_GPU_CL_CKW_DRIVER_GPUCKWKERNELARGUMENTSHELPERS #include "arm_compute/core/CL/ICLTensor.h" @@ -31,7 +31,6 @@ namespace arm_compute { - /** Select a Compute Kernel Writer tensor component from a tensor and add to the kernel's arguments at the specified index idx. * * @param[in,out] kernel OpenCL kernel to configure with the provided argument. @@ -59,4 +58,4 @@ void cl_add_texture_argument(cl::Kernel &kernel, unsigned int &idx, const cl::Im } // namespace arm_compute -#endif //ARM_COMPUTE_SRC_CKW_KERNELARGUMENTSHELPERS_H +#endif /* ACL_SRC_DYNAMIC_FUSION_RUNTIME_GPU_CL_CKW_DRIVER_GPUCKWKERNELARGUMENTSHELPERS */ diff --git a/src/dynamic_fusion/sketch/gpu/GpuKernelArgument.h b/src/dynamic_fusion/sketch/gpu/GpuKernelArgument.h index eb36e91d48..302d4c8562 100644 --- a/src/dynamic_fusion/sketch/gpu/GpuKernelArgument.h +++ b/src/dynamic_fusion/sketch/gpu/GpuKernelArgument.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022 Arm Limited. + * Copyright (c) 2022-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -53,7 +53,9 @@ struct GpuKernelArgumentInfo Tensor_3D, Tensor_4D, Tensor_4D_t_Buffer, - Tensor_4D_t_Image + Tensor_4D_t_Image, + + Tensor_Special_0, }; /** Default constructor */ GpuKernelArgumentInfo() = default; diff --git a/src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwDriver.cpp b/src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwDriver.cpp new file mode 100644 index 0000000000..7c8ec8777d --- /dev/null +++ b/src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwDriver.cpp @@ -0,0 +1,115 @@ +/* + * Copyright (c) 2023 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 "src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwDriver.h" + +#include "src/dynamic_fusion/sketch/gpu/ckw_driver/IGpuCkwComponentDriver.h" +#include "src/dynamic_fusion/sketch/gpu/components/IGpuKernelComponent.h" + +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Window.h" +#include "src/common/utils/Log.h" +#include "src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwVariableTable.h" + +#include "acl/AclKernelWriter.h" +#include "acl/AclScopedKernelWriter.h" + +using namespace ckw; +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +GpuCkwDriver::GpuCkwDriver(const GpuKernelComponentGroup &components) + : _components{ components } +{ +} + +std::string GpuCkwDriver::get_name() +{ + ARM_COMPUTE_LOG_PARAMS(std::string("[V1] TODO")); + return "todo_get_name"; +} + +std::string GpuCkwDriver::get_code() +{ + ARM_COMPUTE_LOG_PARAMS(std::string("[V1] TODO")); + ckw::Kernel kernel(get_name().c_str(), GpuTargetLanguage::OpenCL); + AclKernelWriter root_writer(kernel); + AclScopedKernelWriter writer(&root_writer); + GpuCkwVariableTable vtable{}; + + // Global Kernel Writer Driver code + + // The following is just an incomplete example of using the kernel writer + + // Iterate over component specific Ckw Driver; generate component code and concatenate them + for(auto &comp : _components) + { + auto ckw_driver = comp->ckw_component_driver(); + ARM_COMPUTE_ERROR_ON(ckw_driver == nullptr); + ckw_driver->write_component_code(_components, vtable, writer); + } + + std::string code = root_writer.generate_code(); + + return code; +} + +CLBuildOptions GpuCkwDriver::get_build_options() +{ + ARM_COMPUTE_LOG_PARAMS(std::string("[V1] TO REMOVE")); + return CLBuildOptions{}; +} + +std::string GpuCkwDriver::get_config_id() +{ + ARM_COMPUTE_LOG_PARAMS(std::string("[V1] TODO")); + return ""; +} + +Window GpuCkwDriver::get_window() const +{ + const auto root_comp = _components.get_root_component(); + ARM_COMPUTE_ERROR_ON_MSG(root_comp == nullptr, "No root component found"); + return root_comp->ckw_component_driver()->get_window(); +} + +std::map<ITensorInfo::Id, GpuKernelArgument> GpuCkwDriver::get_tensors() +{ + ARM_COMPUTE_LOG_PARAMS(std::string("[V1] TODO")); + // Assemble GpuKernelArguments + std::map<ITensorInfo::Id, GpuKernelArgument> tensors; + for(const auto t : _components.get_argument_tensors()) + { + tensors.emplace( + t->id(), + GpuKernelArgument{ *t, { GpuKernelArgumentInfo::Type::Tensor_Special_0 } }); + } + return tensors; +} + +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute diff --git a/src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwDriver.h b/src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwDriver.h new file mode 100644 index 0000000000..2084b72098 --- /dev/null +++ b/src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwDriver.h @@ -0,0 +1,80 @@ +/* + * Copyright (c) 2023 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. + */ +#ifndef ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_GPUCKWDRIVER +#define ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_GPUCKWDRIVER + +#include "src/dynamic_fusion/sketch/gpu/GpuKernelArgument.h" +#include "src/dynamic_fusion/sketch/gpu/GpuKernelComponentGroup.h" +#include "src/dynamic_fusion/sketch/gpu/IGpuKernelWriter.h" + +#include <map> +#include <string> + +namespace arm_compute +{ +/** Forward declarations */ +class Window; +class CLBuildOptions; + +namespace experimental +{ +namespace dynamic_fusion +{ +/** Use Kernel Writer to write kernel code + * Used by dynamic_fusion module + */ +class GpuCkwDriver : public IGpuKernelWriter +{ +public: + /** Default constructor */ + GpuCkwDriver() = default; + /** Constructor + * + * @param[in] components Kernel component group from which the kernel will be generated + */ + GpuCkwDriver(const GpuKernelComponentGroup &components); + /** Destructor */ + ~GpuCkwDriver() override = default; + /** Generate kernel name */ + std::string get_name() override; + /** Generate kernel code */ + std::string get_code() override; + /** Generate build options */ + CLBuildOptions get_build_options() override; + /** Generate config id string of the entire kernel. This is used for tuning */ + std::string get_config_id() override; + /** Generate execution window */ + Window get_window() const override; + /** Get the kernel argument lists of the kernel*/ + std::map<ITensorInfo::Id, GpuKernelArgument> get_tensors() override; + +private: + GpuKernelComponentGroup _components{}; +}; + +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute + +#endif /* ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_GPUCKWDRIVER */ diff --git a/src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwVariableTable.cpp b/src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwVariableTable.cpp new file mode 100644 index 0000000000..85aed282d1 --- /dev/null +++ b/src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwVariableTable.cpp @@ -0,0 +1,69 @@ +/* + * Copyright (c) 2023 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 "src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwVariableTable.h" + +#include "acl/AclKernelWriter.h" +#include "acl/AclScopedKernelWriter.h" +#include "src/dynamic_fusion/sketch/gpu/GpuKernelComponentGroup.h" +#include "src/dynamic_fusion/sketch/gpu/ckw_driver/components/utils/TypeConverter.h" +#include <sstream> + +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +AclComponentArgument *GpuCkwVariableTable::declare_variable(const GpuKernelComponentGroup &comp_group, AclScopedKernelWriter &writer, const ITensorInfo *tensor, const std::string &alias) +{ + ARM_COMPUTE_ERROR_ON_MSG(!tensor->has_valid_id(), "Tensor info with valid id expected"); + + // Do not re-declare if the variable associated with the tensor has already been declared + auto it = _vars.find(tensor->id()); + + if(it != _vars.end()) + { + return &it->second; + } + if(comp_group.is_intermediate_tensor(tensor)) + { + // Create a virtual tensor variable + AclComponentArgument var; + auto &&inserted = _vars.emplace(tensor->id(), var); + return &(inserted.first->second); + } + else + { + // Create a user tensor variable + std::stringstream ss; + ss << alias << "_t" << abs(tensor->id()); + const auto uniq_name = ss.str(); + AclComponentArgument var{ writer->create_tensor_argument(uniq_name.c_str(), to_ckw(*tensor)) }; + auto &&inserted = _vars.emplace(tensor->id(), var); + return &(inserted.first->second); + } +} +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute diff --git a/src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwVariableTable.h b/src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwVariableTable.h new file mode 100644 index 0000000000..170fda451b --- /dev/null +++ b/src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwVariableTable.h @@ -0,0 +1,69 @@ +/* + * Copyright (c) 2023 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. + */ +#ifndef ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_GPUCKWVARIABLETABLE +#define ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_GPUCKWVARIABLETABLE + +#include "acl/AclComponentArgument.h" +#include "arm_compute/core/ITensorInfo.h" + +#include <map> + +class AclScopedKernelWriter; + +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +class GpuKernelComponentGroup; + +/** A table of all the variables used in the kernel. + * + * It determines whether we create an virtual tensor var or a user tensor var + * It avoids duplicating variables for the same tensors (Tensors with the same id) + * Each kernel has exactly one variable table. + */ +class GpuCkwVariableTable +{ +public: + /** Declare a kernel component variable(argument) for the corresponding tensor info. + * + * @param[in] comp_group Component group the tensor belongs to + * @param[in] writer Compute Kernel Writer + * @param[in] tensor Tensor info with which the new variable is associated + * @param[in] alias Alias for the variable. Will be used as part of the variable name + * + * @return AclComponentArgument* + */ + AclComponentArgument *declare_variable(const GpuKernelComponentGroup &comp_group, AclScopedKernelWriter &writer, const ITensorInfo *tensor, const std::string &alias = "unnamed"); + +private: + std::map<ITensorInfo::Id, AclComponentArgument> _vars{}; +}; + +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute +#endif /* ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_GPUCKWVARIABLETABLE */ diff --git a/src/dynamic_fusion/sketch/gpu/ckw_driver/IGpuCkwComponentDriver.h b/src/dynamic_fusion/sketch/gpu/ckw_driver/IGpuCkwComponentDriver.h new file mode 100644 index 0000000000..15402bc330 --- /dev/null +++ b/src/dynamic_fusion/sketch/gpu/ckw_driver/IGpuCkwComponentDriver.h @@ -0,0 +1,118 @@ +/* + * Copyright (c) 2023 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. + */ +#ifndef ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_IGPUCKWCOMPONENTDRIVER +#define ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_IGPUCKWCOMPONENTDRIVER + +#include "arm_compute/core/Window.h" +#include "src/dynamic_fusion/sketch/ArgumentPack.h" +#include "src/dynamic_fusion/sketch/gpu/components/Types.h" + +class AclScopedKernelWriter; + +namespace arm_compute +{ +class ITensorInfo; +namespace experimental +{ +namespace dynamic_fusion +{ +/** Forward declaration */ +class GpuKernelComponentGroup; +class GpuCkwVariableTable; + +/** An interface used by @ref GpuCkwDriver to write source code for a kernel component + * + * There are 3 main architecture layers for using Compute Kernel Writer (Ckw) inside ACL's dynamic fusion module + * From top level to bottom level: + * | Layer | Library + * =========================== + * | dynamic_fusion | acl + * | ckw_driver | acl + * | ckw | ckw + * + * ckw_driver is a glue layer that directs how fused code is produced using the ckw library + * + * There are two main groups within ckw_driver: + * - @ref GpuCkwDriver is a global driver that coordinates how the final fused code along with all the info necessary + * for run time execution is produced using ckw + * - Various classes implementing @ref IGpuCkwComponentDriver is a component driver that directs ckw to generate kernel component code (e.g. activation, store etc.) + * + * The overall flow goes like this: + * In dynamic_fusion module, @ref GpuLogicalKernel instantiates a @ref GpuCkwDriver from a @ref GpuKernelComponentGroup + * The logical kernel then uses the global driver's various interfaces to generate the code info. + * In particular, the @ref GpuCkwDriver::get_code() interface will call into each @ref IGpuCkwComponentDriver::write_component_code() + */ +class IGpuCkwComponentDriver +{ +public: + using ComponentGroup = GpuKernelComponentGroup; + +public: + /** Constructor + * + * @param[in] id Component id + * @param[in] tensors Tensor arguments to the components + */ + IGpuCkwComponentDriver(ComponentId id, const ArgumentPack<ITensorInfo> &tensors) + : _id{ id }, _tensors{ tensors } + { + } + /** Destructor */ + virtual ~IGpuCkwComponentDriver() + { + } + /** Generate kernel component code + * + * @param[in] comp_group Component group of which the component is a part of + * @param[in, out] vtable Table of variables declared by each component + * @param[in, out] writer CKW writer that writes code scoped to this kernel component. + * + * @note @p writer can only be passed via value since the new scope is created in the copy constructor + */ + virtual void write_component_code(const ComponentGroup &comp_group, GpuCkwVariableTable &vtable, AclScopedKernelWriter writer) const = 0; + /** Get tensor arguments */ + ArgumentPack<ITensorInfo> tensors() const + { + return _tensors; + } + /** Generate the execution window for the component */ + virtual Window get_window() const + { + return Window{}; + } + /** Get component id */ + ComponentId id() const + { + return _id; + } + +private: + ComponentId _id{ -1 }; + ArgumentPack<ITensorInfo> _tensors{}; +}; +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute + +#endif /* ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_IGPUCKWCOMPONENTDRIVER */ diff --git a/src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwElementwiseBinary.cpp b/src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwElementwiseBinary.cpp new file mode 100644 index 0000000000..9895bbeb77 --- /dev/null +++ b/src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwElementwiseBinary.cpp @@ -0,0 +1,143 @@ +/* + * Copyright (c) 2023 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 "GpuCkwElementwiseBinary.h" + +#include "acl/AclKernelWriter.h" +#include "acl/AclScopedKernelWriter.h" +#include "arm_compute/core/Error.h" +#include "arm_compute/core/Validate.h" +#include "ckw/TensorTileSampler.h" +#include "ckw/Types.h" +#include "src/core/helpers/WindowHelpers.h" +#include "src/dynamic_fusion/sketch/gpu/GpuKernelComponentGroup.h" +#include "src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwVariableTable.h" +#include "src/dynamic_fusion/sketch/gpu/ckw_driver/components/utils/WriterHelper.h" +#include <string> + +using namespace ckw; +namespace arm_compute +{ +namespace experimental +{ +namespace +{ +/** Create a simple sampler from tile of dimension [m0, n0] + */ +inline TensorTileSampler create_simple_sampler(AclScopedKernelWriter &writer, int32_t m0, int32_t n0) +{ + TensorTileSampler sampler; + + auto &gid_0 = writer->declare_tile("gid_0", ckw::DataType::Int32); + auto &gid_1 = writer->declare_tile("gid_1", ckw::DataType::Int32); + auto &gid_2 = writer->declare_tile("gid_2", ckw::DataType::Int32); + + auto &const_0 = writer->declare_tile("0", 0); + + writer->op_get_global_id(gid_0, 0); + writer->op_get_global_id(gid_1, 1); + writer->op_get_global_id(gid_2, 2); + + sampler.x(gid_0); + sampler.y(gid_1); + sampler.z(const_0); // 3rd dimension collapsed with 2nd dimension + sampler.b(gid_2); + + sampler.width(n0); + sampler.height(m0); + + sampler.format(TensorSamplerFormat::C_WH_1); // 3rd dimension collapsed with 2nd dimension + sampler.address_mode_x(TensorSamplerAddressModeX::None); + sampler.address_mode_y(TensorSamplerAddressModeY::ClampToBorder); + sampler.address_mode_z(TensorSamplerAddressModeZ::Skip); // Dimensions higher than 3 not supported yet + + return sampler; +} +} // namespace + +namespace dynamic_fusion +{ +GpuCkwElementwiseBinary::GpuCkwElementwiseBinary(ComponentId id, + const ArgumentPack<ITensorInfo> &tensors, + const Attributes &attributes) + : IGpuCkwComponentDriver{ id, tensors }, + _lhs{}, + _rhs{}, + _dst{}, + _attributes{ attributes } +{ + _lhs = this->tensors().get_const_tensor(TensorType::ACL_SRC_0); + _rhs = this->tensors().get_const_tensor(TensorType::ACL_SRC_1); + _dst = this->tensors().get_const_tensor(TensorType::ACL_DST_0); + ARM_COMPUTE_ERROR_ON_NULLPTR(_lhs, _rhs, _dst); +} + +void GpuCkwElementwiseBinary::write_component_code(const ComponentGroup &comp_group, GpuCkwVariableTable &vtable, AclScopedKernelWriter writer) const +{ + const auto root_window = comp_group.get_root_component()->ckw_component_driver()->get_window(); + const unsigned int n0 = root_window.x().step(); + const unsigned int m0 = root_window.y().step(); + + AclComponentArgument *lhs = vtable.declare_variable(comp_group, writer, _lhs, "lhs"); + AclComponentArgument *rhs = vtable.declare_variable(comp_group, writer, _rhs, "rhs"); + AclComponentArgument *dst = vtable.declare_variable(comp_group, writer, _dst, "dst"); + + // Load the LHS and RHS tiles and prepare the tensor sampler. + load_lhs_rhs_tiles_and_prepare_sampler(writer, lhs, rhs, m0, n0, create_simple_sampler); + + auto &lhs_tile = lhs->tile(); + auto &rhs_tile = rhs->tile(); + const auto &sampler = lhs->tile_sampler(); + + // Prepare the output tile. + if(!dst->has_tile()) + { + auto &tile = writer->declare_tile("dst_tile", lhs_tile.tile_info()); + dst->init_virtual_tensor(tile, sampler); + } + + auto &dst_tile = dst->tile(); + + // Perform the operation. + writer->op_binary_expression(dst_tile, lhs_tile, rhs_tile, BinaryOp::Add); +} + +Window GpuCkwElementwiseBinary::get_window() const +{ + ARM_COMPUTE_ERROR_ON_MSG(_dst->tensor_shape().total_size() == 0U, "Destination tensor is not initialized"); + + TensorShape output_shape = _dst->tensor_shape(); + // Collapse Dim 1 (W) and Dim 2 (H) together, leave Dim 0 (C) unchanged + // This is in line with the collapsing convention used by operators like Conv2d + output_shape.collapse(2U, 1U); + // constexpr unsigned int vector_size_byte_opencl = 16; + // const unsigned int num_elems_processed_per_iteration = adjust_vec_size(vector_size_byte_opencl / _dst->element_size(), _dst->dimension(0)); + const unsigned int num_elems_processed_per_iteration = 1U; // Hard-coded for now + Window win = calculate_max_window(output_shape, Steps(num_elems_processed_per_iteration)); + + return win; +} + +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute diff --git a/src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwElementwiseBinary.h b/src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwElementwiseBinary.h new file mode 100644 index 0000000000..1a79754d1d --- /dev/null +++ b/src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwElementwiseBinary.h @@ -0,0 +1,69 @@ +/* + * Copyright (c) 2023 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. + */ +#ifndef ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_COMPONENTS_GPUCKWELEMENTWISEBINARY +#define ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_COMPONENTS_GPUCKWELEMENTWISEBINARY + +#include "src/core/common/Macros.h" +#include "src/dynamic_fusion/sketch/gpu/ckw_driver/IGpuCkwComponentDriver.h" +#include "src/dynamic_fusion/sketch/gpu/components/cl/ClComponentElementwiseBinary.h" + +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +class GpuCkwElementwiseBinary : public IGpuCkwComponentDriver +{ +public: + using Attributes = ClComponentElementwiseBinary::Attributes; + /** Constructor + * + * For supported configurations please refer to @ref ClComponentElementwiseBinary::validate() + * + * @param[in] id Component id + * @param[in] tensors Tensor arguments to the component + * @param[in] attributes Component attributes + */ + GpuCkwElementwiseBinary(ComponentId id, + const ArgumentPack<ITensorInfo> &tensors, + const Attributes &attributes); + ARM_COMPUTE_DISALLOW_COPY_ALLOW_MOVE(GpuCkwElementwiseBinary); + /** Destructor */ + ~GpuCkwElementwiseBinary() override = default; + // Inherited methods overriden: + virtual void write_component_code(const ComponentGroup &comp_group, GpuCkwVariableTable &vtable, AclScopedKernelWriter writer) const override; + Window get_window() const override; + +private: + const ITensorInfo *_lhs; + const ITensorInfo *_rhs; + const ITensorInfo *_dst; + Attributes _attributes; +}; +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute + +#endif /* ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_COMPONENTS_GPUCKWELEMENTWISEBINARY */ diff --git a/src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwStore.cpp b/src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwStore.cpp new file mode 100644 index 0000000000..1a1dfc135a --- /dev/null +++ b/src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwStore.cpp @@ -0,0 +1,57 @@ +/* + * Copyright (c) 2023 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 "GpuCkwStore.h" + +#include "arm_compute/core/Error.h" +#include "compute_kernel_writer/include/acl/AclKernelWriter.h" +#include "compute_kernel_writer/include/acl/AclScopedKernelWriter.h" +#include "src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwVariableTable.h" +#include <string> + +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +GpuCkwStore::GpuCkwStore(ComponentId id, const ArgumentPack<ITensorInfo> &tensors) + : IGpuCkwComponentDriver{ id, tensors }, _src{}, _dst{} +{ + _src = this->tensors().get_const_tensor(TensorType::ACL_SRC_0); + _dst = this->tensors().get_const_tensor(TensorType::ACL_DST_0); +} +void GpuCkwStore::write_component_code(const ComponentGroup &comp_group, GpuCkwVariableTable &vtable, AclScopedKernelWriter writer) const +{ + auto src = vtable.declare_variable(comp_group, writer, _src, "src"); + auto dst = vtable.declare_variable(comp_group, writer, _dst, "dst"); + + auto &src_tile = src->tile(); + const auto &sampler = src->tile_sampler(); + auto &dst_tensor = dst->tensor(); + + writer->op_store(dst_tensor, src_tile, sampler); +} +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute diff --git a/src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwStore.h b/src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwStore.h new file mode 100644 index 0000000000..45cc43fe62 --- /dev/null +++ b/src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwStore.h @@ -0,0 +1,61 @@ +/* + * Copyright (c) 2023 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. + */ +#ifndef ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_COMPONENTS_GPUCKWSTORE +#define ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_COMPONENTS_GPUCKWSTORE + +#include "src/core/common/Macros.h" +#include "src/dynamic_fusion/sketch/gpu/ckw_driver/IGpuCkwComponentDriver.h" + +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +/** An interface used by @ref ClTemplateWriter to write source code for a kernel component + */ +class GpuCkwStore : public IGpuCkwComponentDriver +{ +public: + /** Constructor + * + * @param[in] id Component id + * @param[in] tensors Tensor arguments to the component + */ + GpuCkwStore(ComponentId id, const ArgumentPack<ITensorInfo> &tensors); + ARM_COMPUTE_DISALLOW_COPY_ALLOW_MOVE(GpuCkwStore); + /** Destructor */ + ~GpuCkwStore() override = default; + // Inherited methods overriden: + virtual void write_component_code(const ComponentGroup &comp_group, GpuCkwVariableTable &vtable, AclScopedKernelWriter writer) const override; + +private: + const ITensorInfo *_src; + const ITensorInfo *_dst; +}; +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute + +#endif /* ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_COMPONENTS_GPUCKWSTORE */ diff --git a/src/dynamic_fusion/sketch/gpu/ckw_driver/components/utils/TypeConverter.h b/src/dynamic_fusion/sketch/gpu/ckw_driver/components/utils/TypeConverter.h new file mode 100644 index 0000000000..2531fb7379 --- /dev/null +++ b/src/dynamic_fusion/sketch/gpu/ckw_driver/components/utils/TypeConverter.h @@ -0,0 +1,102 @@ +/* + * Copyright (c) 2023 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. + */ +#ifndef ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_COMPONENTS_UTILS_TYPECONVERTER +#define ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_COMPONENTS_UTILS_TYPECONVERTER + +#include "arm_compute/core/ITensorInfo.h" +#include "arm_compute/core/TensorShape.h" +#include "arm_compute/core/Types.h" +#include "ckw/TensorInfo.h" + +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +inline ckw::DataType to_ckw(DataType dt) +{ + switch(dt) + { + case DataType::F32: + return ckw::DataType::Fp32; + case DataType::F16: + return ckw::DataType::Fp16; + case DataType::S32: + return ckw::DataType::Int32; + case DataType::S16: + return ckw::DataType::Int16; + case DataType::S8: + return ckw::DataType::Int8; + case DataType::U32: + return ckw::DataType::Uint32; + case DataType::U16: + return ckw::DataType::Uint16; + case DataType::U8: + return ckw::DataType::Uint8; + default: + return ckw::DataType::Unknown; + } +} + +inline ckw::TensorShape to_ckw(const TensorShape &shape) +{ + ARM_COMPUTE_ERROR_ON(shape.num_max_dimensions < std::tuple_size<ckw::TensorShape> {}); + ARM_COMPUTE_ERROR_ON(std::tuple_size<ckw::TensorShape> {} != 5); + /// NOTE: Overflow danger. Use size_t? + return ckw::TensorShape + { + static_cast<int32_t>(shape[0]), + static_cast<int32_t>(shape[1]), + static_cast<int32_t>(shape[2]), + static_cast<int32_t>(shape[3]), + static_cast<int32_t>(shape[4]) + }; +} +inline ckw::TensorDataLayout to_ckw(DataLayout dl) +{ + switch(dl) + { + case DataLayout::NHWC: + return ckw::TensorDataLayout::Nhwc; + case DataLayout::NDHWC: + return ckw::TensorDataLayout::Ndhwc; + default: + return ckw::TensorDataLayout::Unknown; + } +} +inline ckw::TensorInfo to_ckw(const ITensorInfo &tensor_info) +{ + return ckw::TensorInfo + { + to_ckw(tensor_info.data_type()), + to_ckw(tensor_info.tensor_shape()), + to_ckw(tensor_info.data_layout()), + tensor_info.id() + }; +} +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute +#endif /* ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_COMPONENTS_UTILS_TYPECONVERTER */ diff --git a/src/dynamic_fusion/sketch/gpu/ckw_driver/components/utils/WriterHelper.h b/src/dynamic_fusion/sketch/gpu/ckw_driver/components/utils/WriterHelper.h new file mode 100644 index 0000000000..d94ebd5ce9 --- /dev/null +++ b/src/dynamic_fusion/sketch/gpu/ckw_driver/components/utils/WriterHelper.h @@ -0,0 +1,68 @@ +/* + * Copyright (c) 2023 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. + */ +#ifndef ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_COMPONENTS_UTILS_WRITERHELPER +#define ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_COMPONENTS_UTILS_WRITERHELPER + +#include "acl/AclComponentArgument.h" +#include "acl/AclScopedKernelWriter.h" +#include "ckw/TensorTileSampler.h" + +#include <functional> + +using namespace ckw; +namespace arm_compute +{ +namespace experimental +{ +namespace dynamic_fusion +{ +using SamplerCreator = std::function<TensorTileSampler(AclScopedKernelWriter &, int32_t /* m0 */, int32_t /* n0 */)>; + +/** Load lhs and rhs tiles of dimension [m0, n0] only when not loaded and prepare the sampler + */ +inline void load_lhs_rhs_tiles_and_prepare_sampler(AclScopedKernelWriter &writer, AclComponentArgument *lhs, AclComponentArgument *rhs, int32_t m0, int32_t n0, SamplerCreator create_sampler) +{ + if(!lhs->has_tile() && !rhs->has_tile()) + { + const auto sampler = create_sampler(writer, m0, n0); + + writer->op_load_once(lhs, sampler); + writer->op_load_once(rhs, sampler); + } + else if(lhs->has_tile()) + { + const auto &sampler = lhs->tile_sampler(); + writer->op_load_once(rhs, sampler); + } + else + { + const auto &sampler = rhs->tile_sampler(); + writer->op_load_once(lhs, sampler); + } +} + +} // namespace dynamic_fusion +} // namespace experimental +} // namespace arm_compute +#endif /* ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_CKW_DRIVER_COMPONENTS_UTILS_WRITERHELPER */ diff --git a/src/dynamic_fusion/sketch/gpu/components/IGpuKernelComponent.h b/src/dynamic_fusion/sketch/gpu/components/IGpuKernelComponent.h index 8bb19155a2..d600956b4f 100644 --- a/src/dynamic_fusion/sketch/gpu/components/IGpuKernelComponent.h +++ b/src/dynamic_fusion/sketch/gpu/components/IGpuKernelComponent.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022 Arm Limited. + * Copyright (c) 2022-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -21,8 +21,8 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. */ -#ifndef SRC_DYNAMIC_FUSION_SKETCH_GPU_COMPONENTS_IGPUKERNELCOMPONENT -#define SRC_DYNAMIC_FUSION_SKETCH_GPU_COMPONENTS_IGPUKERNELCOMPONENT +#ifndef ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_COMPONENTS_IGPUKERNELCOMPONENT +#define ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_COMPONENTS_IGPUKERNELCOMPONENT #include "Types.h" @@ -60,6 +60,7 @@ inline bool operator==(const KernelProperties &config0, const KernelProperties & /** Forward declaration */ class IGpuTemplateComponentWriter; +class IGpuCkwComponentDriver; /** An abstract interface of a component. It enables manipulation by the component graph for purposes like fusion */ @@ -105,6 +106,11 @@ public: } /** Get template writer for the component */ virtual const IGpuTemplateComponentWriter *template_writer() const = 0; + /** Get compute kernel writer driver for the component */ + virtual const IGpuCkwComponentDriver *ckw_component_driver() const + { + return nullptr; + } /** Get component type */ virtual GpuComponentType type() const = 0; @@ -116,4 +122,4 @@ private: } // namespace dynamic_fusion } // namespace experimental } // namespace arm_compute -#endif /* SRC_DYNAMIC_FUSION_SKETCH_GPU_COMPONENTS_IGPUKERNELCOMPONENT */ +#endif /* ACL_SRC_DYNAMIC_FUSION_SKETCH_GPU_COMPONENTS_IGPUKERNELCOMPONENT */ diff --git a/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentActivation.cpp b/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentActivation.cpp index 6eaa45c25d..3e8d256a08 100644 --- a/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentActivation.cpp +++ b/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentActivation.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022 Arm Limited. + * Copyright (c) 2022-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -32,7 +32,6 @@ namespace experimental { namespace dynamic_fusion { - Status ClComponentActivation::validate(const Properties &properties, const ArgumentPack<ITensorInfo> &tensors, const Attributes &attributes) @@ -70,6 +69,10 @@ ClComponentActivation::ClComponentActivation(ComponentId { } +ClComponentActivation::~ClComponentActivation() +{ +} + const IGpuTemplateComponentWriter *ClComponentActivation::template_writer() const { return _component_writer.get(); diff --git a/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentActivation.h b/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentActivation.h index d1b849ec73..d5013acddf 100644 --- a/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentActivation.h +++ b/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentActivation.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022 Arm Limited. + * Copyright (c) 2022-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -88,7 +88,7 @@ public: const Attributes &attributes); /** Destructor */ - ~ClComponentActivation() override = default; + ~ClComponentActivation() override; /** Prevent instances of this class from being copy constructed */ ClComponentActivation(const ClComponentActivation &component) = delete; diff --git a/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentElementwiseBinary.cpp b/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentElementwiseBinary.cpp index b21c7c382f..52739e23c0 100644 --- a/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentElementwiseBinary.cpp +++ b/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentElementwiseBinary.cpp @@ -25,6 +25,7 @@ #include "arm_compute/core/Validate.h" #include "src/core/CL/CLValidate.h" +#include "src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwElementwiseBinary.h" #include "src/dynamic_fusion/sketch/gpu/template_writer/cl/ClTemplateElementwiseBinary.h" namespace arm_compute @@ -111,7 +112,8 @@ ClComponentElementwiseBinary::ClComponentElementwiseBinary( const ArgumentPack<ITensorInfo> &tensors, const Attributes &attributes) : IGpuKernelComponent{ id, properties, tensors }, - _component_writer{ std::make_unique<ClTemplateElementwiseBinary>(id, tensors, attributes) } + _component_writer{ std::make_unique<ClTemplateElementwiseBinary>(id, tensors, attributes) }, + _ckw_driver{ std::make_unique<GpuCkwElementwiseBinary>(id, tensors, attributes) } { } ClComponentElementwiseBinary::~ClComponentElementwiseBinary() @@ -121,6 +123,11 @@ const IGpuTemplateComponentWriter *ClComponentElementwiseBinary::template_writer { return _component_writer.get(); } + +const IGpuCkwComponentDriver *ClComponentElementwiseBinary::ckw_component_driver() const +{ + return _ckw_driver.get(); +} } // namespace dynamic_fusion } // namespace experimental } // namespace arm_compute diff --git a/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentElementwiseBinary.h b/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentElementwiseBinary.h index 02e61019f4..a56dd8b37d 100644 --- a/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentElementwiseBinary.h +++ b/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentElementwiseBinary.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022 Arm Limited. + * Copyright (c) 2022-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -24,7 +24,6 @@ #ifndef SRC_DYNAMIC_FUSION_SKETCH_GPU_COMPONENTS_CL_CLCOMPONENTELEMENTWISEBINARY #define SRC_DYNAMIC_FUSION_SKETCH_GPU_COMPONENTS_CL_CLCOMPONENTELEMENTWISEBINARY -#include "arm_compute/core/Error.h" #include "src/dynamic_fusion/sketch/gpu/components/IGpuKernelComponent.h" #include "src/dynamic_fusion/sketch/gpu/operators/internal/GpuElementwiseBinaryCommon.h" @@ -42,6 +41,7 @@ class ArgumentPack; /** Forward declaration */ class ClTemplateElementwiseBinary; +class GpuCkwElementwiseBinary; class ClComponentElementwiseBinary final : public IGpuKernelComponent { @@ -103,6 +103,8 @@ public: ClComponentElementwiseBinary &operator=(ClComponentElementwiseBinary &&component) = default; /** Get template writer for the component */ const IGpuTemplateComponentWriter *template_writer() const override; + + const IGpuCkwComponentDriver *ckw_component_driver() const override; /** Get component type */ GpuComponentType type() const override { @@ -111,6 +113,7 @@ public: private: std::unique_ptr<ClTemplateElementwiseBinary> _component_writer; + std::unique_ptr<GpuCkwElementwiseBinary> _ckw_driver; }; } // namespace dynamic_fusion } // namespace experimental diff --git a/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentStore.cpp b/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentStore.cpp index f49f397ec1..a3283b1866 100644 --- a/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentStore.cpp +++ b/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentStore.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022 Arm Limited. + * Copyright (c) 2022-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -24,6 +24,7 @@ #include "ClComponentStore.h" #include "src/dynamic_fusion/sketch/ArgumentPack.h" +#include "src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwStore.h" #include "src/dynamic_fusion/sketch/gpu/template_writer/cl/ClTemplateStore.h" #include <memory> @@ -42,7 +43,7 @@ Status ClComponentStore::validate( return Status{}; } ClComponentStore::ClComponentStore(ComponentId id, const Properties &properties, const ArgumentPack<ITensorInfo> &tensors) - : IGpuKernelComponent{ id, properties, tensors }, _component_writer{ std::make_unique<ClTemplateStore>(id, tensors) } + : IGpuKernelComponent{ id, properties, tensors }, _component_writer{ std::make_unique<ClTemplateStore>(id, tensors) }, _ckw_driver{ std::make_unique<GpuCkwStore>(id, tensors) } { } ClComponentStore::~ClComponentStore() @@ -52,6 +53,10 @@ const IGpuTemplateComponentWriter *ClComponentStore::template_writer() const { return _component_writer.get(); } +const IGpuCkwComponentDriver *ClComponentStore::ckw_component_driver() const +{ + return _ckw_driver.get(); +} } // namespace dynamic_fusion } // namespace experimental } // namespace arm_compute diff --git a/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentStore.h b/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentStore.h index bf8c9f031e..f168ccb97e 100644 --- a/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentStore.h +++ b/src/dynamic_fusion/sketch/gpu/components/cl/ClComponentStore.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022 Arm Limited. + * Copyright (c) 2022-2023 Arm Limited. * * SPDX-License-Identifier: MIT * @@ -24,7 +24,6 @@ #ifndef SRC_DYNAMIC_FUSION_SKETCH_GPU_COMPONENTS_CL_CLCOMPONENTSTORE #define SRC_DYNAMIC_FUSION_SKETCH_GPU_COMPONENTS_CL_CLCOMPONENTSTORE -#include "arm_compute/core/Error.h" #include "src/dynamic_fusion/sketch/gpu/components/IGpuKernelComponent.h" #include "src/dynamic_fusion/sketch/gpu/template_writer/cl/ClTemplateStore.h" #include <memory> @@ -40,6 +39,7 @@ namespace dynamic_fusion /** Forward declaration */ template <typename T> class ArgumentPack; +class GpuCkwStore; class ClComponentStore final : public IGpuKernelComponent { @@ -87,6 +87,8 @@ public: ClComponentStore &operator=(ClComponentStore &&component) = default; /** Get template writer for the component */ const IGpuTemplateComponentWriter *template_writer() const override; + + const IGpuCkwComponentDriver *ckw_component_driver() const override; /** Get component type */ GpuComponentType type() const override { @@ -95,6 +97,7 @@ public: private: std::unique_ptr<ClTemplateStore> _component_writer; + std::unique_ptr<GpuCkwStore> _ckw_driver; }; } // namespace dynamic_fusion } // namespace experimental diff --git a/src/dynamic_fusion/sketch/gpu/operators/GpuSigmoid.cpp b/src/dynamic_fusion/sketch/gpu/operators/GpuSigmoid.cpp index 75b726bfda..87439c4e3d 100644 --- a/src/dynamic_fusion/sketch/gpu/operators/GpuSigmoid.cpp +++ b/src/dynamic_fusion/sketch/gpu/operators/GpuSigmoid.cpp @@ -25,12 +25,11 @@ #include "arm_compute/dynamic_fusion/sketch/gpu/operators/GpuSigmoid.h" #include "arm_compute/core/experimental/Types.h" +#include "src/common/utils/Log.h" +#include "src/core/helpers/AutoConfiguration.h" #include "src/dynamic_fusion/sketch/ArgumentPack.h" #include "src/dynamic_fusion/sketch/gpu/GpuWorkloadSketchImpl.h" #include "src/dynamic_fusion/sketch/gpu/components/cl/ClComponentActivation.h" -#include "src/dynamic_fusion/sketch/gpu/template_writer/cl/ClTemplateActivation.h" -#include "src/core/helpers/AutoConfiguration.h" -#include "src/common/utils/Log.h" namespace arm_compute { @@ -81,7 +80,7 @@ constexpr GpuOperatorType operator_type = GpuOperatorType::Simple; } // namespace Status GpuSigmoid::is_supported_op(const GpuWorkloadContext &context, - const ITensorInfo *src) + const ITensorInfo *src) { return is_supported_op_helper(context, src, nullptr); } @@ -112,8 +111,8 @@ Status GpuSigmoid::validate_op(const GpuWorkloadSketch &sketch, return is_supported_op_helper(*sketch.gpu_context(), src, &dst_info_to_validate); } -ITensorInfo *GpuSigmoid::create_op(GpuWorkloadSketch &sketch, - ITensorInfo *src) +ITensorInfo *GpuSigmoid::create_op(GpuWorkloadSketch &sketch, + ITensorInfo *src) { ARM_COMPUTE_ERROR_ON_NULLPTR(src); ARM_COMPUTE_LOG_PARAMS(src); |