diff options
27 files changed, 1035 insertions, 47 deletions
diff --git a/Android.bp b/Android.bp index b31c9eebd9..85efa54371 100644 --- a/Android.bp +++ b/Android.bp @@ -237,7 +237,6 @@ cc_library_static { "src/core/CL/ICLSimpleKernel.cpp", "src/core/CL/ICLTensor.cpp", "src/core/CL/OpenCL.cpp", - "src/core/CL/ckw/KernelArgumentsHelpers.cpp", "src/core/CL/kernels/CLArgMinMaxLayerKernel.cpp", "src/core/CL/kernels/CLBatchNormalizationLayerKernel.cpp", "src/core/CL/kernels/CLBatchToSpaceLayerKernel.cpp", @@ -625,6 +624,7 @@ cc_library_static { "src/cpu/operators/internal/CpuGemmAssemblyDispatch.cpp", "src/dynamic_fusion/runtime/gpu/cl/ClKernelRuntime.cpp", "src/dynamic_fusion/runtime/gpu/cl/ClWorkloadRuntime.cpp", + "src/dynamic_fusion/runtime/gpu/cl/ckw_driver/GpuCkwKernelArgumentsHelpers.cpp", "src/dynamic_fusion/sketch/attributes/CastAttributes.cpp", "src/dynamic_fusion/sketch/attributes/ClampAttributes.cpp", "src/dynamic_fusion/sketch/attributes/Conv2dAttributes.cpp", @@ -641,6 +641,10 @@ cc_library_static { "src/dynamic_fusion/sketch/gpu/GpuOperatorGroup.cpp", "src/dynamic_fusion/sketch/gpu/GpuWorkloadContext.cpp", "src/dynamic_fusion/sketch/gpu/GpuWorkloadSketch.cpp", + "src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwDriver.cpp", + "src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwVariableTable.cpp", + "src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwElementwiseBinary.cpp", + "src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwStore.cpp", "src/dynamic_fusion/sketch/gpu/components/cl/ClComponentActivation.cpp", "src/dynamic_fusion/sketch/gpu/components/cl/ClComponentCast.cpp", "src/dynamic_fusion/sketch/gpu/components/cl/ClComponentDepthwiseConv2d.cpp", diff --git a/SConscript b/SConscript index c7139552e3..9ddffc4273 100644 --- a/SConscript +++ b/SConscript @@ -127,7 +127,7 @@ def build_library(name, build_env, sources, static=False, libs=[]): cloned_build_env["LINKFLAGS"].remove('-pie') cloned_build_env["LINKFLAGS"].remove('-static-libstdc++') - if env['ckw']: + if env['experimental_dynamic_fusion']: libs.append('libckw.a') if static: @@ -536,11 +536,6 @@ if env['fixed_format_kernels']: if env['experimental_dynamic_fusion']: lib_files += filelist['experimental']['dynamic_fusion'] -# Compute Kernel Writer integration files -if env['ckw']: - if env['opencl']: - lib_files += filelist['experimental']['ckw']['cl'] - # Logging files if env["logging"]: lib_files += filelist['logging'] diff --git a/SConstruct b/SConstruct index 48f17472d3..4994972172 100644 --- a/SConstruct +++ b/SConstruct @@ -102,7 +102,6 @@ vars.AddVariables( BoolVariable("standalone", "Builds the tests as standalone executables, links statically with libgcc, libstdc++ and libarm_compute", False), BoolVariable("opencl", "Enable OpenCL support", True), BoolVariable("neon", "Enable Arm® Neon™ support", False), - BoolVariable("ckw", "Build and link the Compute Kernel Writer subproject", False), BoolVariable("embed_kernels", "Enable if you want the OpenCL kernels to be built in the library's binaries instead of being read from separate '.cl' / '.cs' files. If embed_kernels is set to 0 then the application can set the path to the folder containing the OpenCL kernel files by calling CLKernelLibrary::init(). By default the path is set to './cl_kernels'.", True), BoolVariable("compress_kernels", "Compress embedded OpenCL kernels in library binary using zlib. Useful for reducing the binary size. embed_kernels should be enabled", False), BoolVariable("set_soname", "If enabled the library will contain a SONAME and SHLIBVERSION and some symlinks will automatically be created between the objects. (requires SCons 2.4 or above)", False), @@ -126,7 +125,7 @@ vars.AddVariables( ├── datasets ├── fixtures └── Neon\n""", "", PathVariable.PathAccept), - BoolVariable("experimental_dynamic_fusion", "Build the experimental dynamic fusion files. This option also enables opencl=1 and ckw=1 on which it has a direct dependency.", False), + BoolVariable("experimental_dynamic_fusion", "Build the experimental dynamic fusion files. This option also enables opencl=1 on which it has a direct dependency.", False), BoolVariable("fixed_format_kernels", "Enable fixed format kernels for GEMM", False), BoolVariable("mapfile", "Generate a map file", False), ListVariable("custom_options", "Custom options that can be used to turn on/off features", "none", ["disable_mmla_fp"]), @@ -218,7 +217,6 @@ if env['os'] == 'bare_metal': if env['experimental_dynamic_fusion']: # Dynamic Fusion on GPU has a direct dependency on OpenCL and Compute Kernel Writer env['opencl'] = 1 - env['ckw'] = 1 if env['opencl'] and env['embed_kernels'] and env['compress_kernels'] and env['os'] not in ['android']: print("Compressed kernels are supported only for android builds") @@ -422,7 +420,7 @@ print("CC", env['CC']) print("CXX", env['CXX']) """Build the Compute Kernel Writer subproject""" -if env['ckw']: +if env['experimental_dynamic_fusion']: # Strip ccache prefix from CC and CXX to obtain only the target triple CKW_CC = env['CC'].replace(env['compiler_cache'] + " ", "") CKW_CXX = env['CXX'].replace(env['compiler_cache'] + " ", "") diff --git a/filelist.json b/filelist.json index 0e18b37f5c..b9196cb6b4 100644 --- a/filelist.json +++ b/filelist.json @@ -2336,12 +2336,12 @@ "src/dynamic_fusion/sketch/gpu/template_writer/cl/ClTemplateResize.cpp", "src/dynamic_fusion/sketch/gpu/template_writer/cl/ClTemplateStore.cpp", "src/dynamic_fusion/sketch/gpu/template_writer/cl/ClTemplateWriter.cpp", - "src/dynamic_fusion/sketch/gpu/template_writer/GpuKernelVariableTable.cpp" - ], - "ckw": { - "cl": [ - "src/core/CL/ckw/KernelArgumentsHelpers.cpp" - ] - } + "src/dynamic_fusion/sketch/gpu/template_writer/GpuKernelVariableTable.cpp", + "src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwDriver.cpp", + "src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwElementwiseBinary.cpp", + "src/dynamic_fusion/sketch/gpu/ckw_driver/components/GpuCkwStore.cpp", + "src/dynamic_fusion/sketch/gpu/ckw_driver/GpuCkwVariableTable.cpp", + "src/dynamic_fusion/runtime/gpu/cl/ckw_driver/GpuCkwKernelArgumentsHelpers.cpp" + ] } } 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); |