From ce3c48c7af02555f81c0f5e7ef2677916cecef34 Mon Sep 17 00:00:00 2001 From: Viet-Hoa Do Date: Mon, 3 Jul 2023 13:44:43 +0100 Subject: Move CKW prototype to separate directory Partially resolves: COMPMID-6283 Signed-off-by: Viet-Hoa Do Change-Id: I7596e3dc357d6f0b9cbe66534523943a73c26d81 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/9864 Reviewed-by: SiCong Li Reviewed-by: Jakub Sujak Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins Benchmark: Arm Jenkins --- compute_kernel_writer/include/ckw/Error.h | 66 ++----- compute_kernel_writer/include/ckw/Kernel.h | 77 -------- compute_kernel_writer/include/ckw/KernelWriter.h | 217 --------------------- compute_kernel_writer/include/ckw/OperandBase.h | 76 -------- compute_kernel_writer/include/ckw/ScalarValue.h | 137 ------------- compute_kernel_writer/include/ckw/TensorOperand.h | 181 ----------------- .../include/ckw/TensorTileSampler.h | 163 ---------------- compute_kernel_writer/include/ckw/TileOperand.h | 110 ----------- compute_kernel_writer/include/ckw/Types.h | 119 ++--------- 9 files changed, 23 insertions(+), 1123 deletions(-) delete mode 100644 compute_kernel_writer/include/ckw/Kernel.h delete mode 100644 compute_kernel_writer/include/ckw/KernelWriter.h delete mode 100644 compute_kernel_writer/include/ckw/OperandBase.h delete mode 100644 compute_kernel_writer/include/ckw/ScalarValue.h delete mode 100644 compute_kernel_writer/include/ckw/TensorOperand.h delete mode 100644 compute_kernel_writer/include/ckw/TensorTileSampler.h delete mode 100644 compute_kernel_writer/include/ckw/TileOperand.h (limited to 'compute_kernel_writer/include/ckw') diff --git a/compute_kernel_writer/include/ckw/Error.h b/compute_kernel_writer/include/ckw/Error.h index 8c4853722b..996893823e 100644 --- a/compute_kernel_writer/include/ckw/Error.h +++ b/compute_kernel_writer/include/ckw/Error.h @@ -21,12 +21,11 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. */ +#ifndef COMPUTE_KERNEL_WRITER_INCLUDE_CKW_ERROR_H +#define COMPUTE_KERNEL_WRITER_INCLUDE_CKW_ERROR_H -#ifndef CKW_INCLUDE_CKW_ERROR_H -#define CKW_INCLUDE_CKW_ERROR_H - -#include #include +#include namespace ckw { @@ -45,59 +44,16 @@ std::string create_error_msg(const std::string &file, const std::string &func, c * * @param[in] msg Message to display. */ -#define COMPUTE_KERNEL_WRITER_ERROR_ON_MSG(msg) \ - do \ - { \ - const std::string arg0(__FILE__); \ - const std::string arg1(__func__); \ - const std::string arg2(std::to_string(__LINE__)); \ - const std::string arg3(msg); \ +#define COMPUTE_KERNEL_WRITER_ERROR_ON_MSG(msg) \ + do \ + { \ + const std::string arg0(__FILE__); \ + const std::string arg1(__func__); \ + const std::string arg2(std::to_string(__LINE__)); \ + const std::string arg3(msg); \ std::runtime_error(create_error_msg(arg0, arg1, arg2, arg3)); \ } while(false) -/** If the condition is not met, throw an std::runtime_error with the specified message. - * - * @param[in] cond The condition that is expected to be true. - * @param[in] msg The error message when the condition is not met. - */ -#define CKW_ASSERT_MSG(cond, msg) \ - do \ - { \ - if(!(cond)) \ - { \ - throw ::std::runtime_error(msg); \ - } \ - } while(false) - -/** If the condition is not met, throw an std::runtime_error. - * - * @param[in] cond The condition that is expected to be true. - */ -#define CKW_ASSERT(cond) CKW_ASSERT_MSG(cond, #cond) - -/** If the precondition is met but the consequence is not met, throw an std::runtime_error. - * - * @param[in] precond The condition if is met requires the consequence must also be met. - * @param[in] cond The condition that is expected to be true if the precondition is true. - */ -#define CKW_ASSERT_IF(precond, cond) \ - CKW_ASSERT_MSG(!(precond) || ((precond) && (cond)), #precond " |-> " #cond) - -/** Mark the variables as unused. - * - * @param[in] ... Variables which are unused. - */ -#define CKW_UNUSED(...) ::ckw::ignore_unused(__VA_ARGS__) // NOLINT - -/** Mark the variables as unused. - * - * @param[in] ... Variables which are unused. - */ -template -inline void ignore_unused(T &&...) -{ -} - } // namespace ckw -#endif // CKW_INCLUDE_CKW_ERROR_H +#endif /* COMPUTE_KERNEL_WRITER_INCLUDE_CKW_ERROR_H */ diff --git a/compute_kernel_writer/include/ckw/Kernel.h b/compute_kernel_writer/include/ckw/Kernel.h deleted file mode 100644 index cbc7700c22..0000000000 --- a/compute_kernel_writer/include/ckw/Kernel.h +++ /dev/null @@ -1,77 +0,0 @@ -/* - * 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 CKW_INCLUDE_CKW_KERNEL_H -#define CKW_INCLUDE_CKW_KERNEL_H - -#include "ckw/OperandBase.h" -#include "ckw/Types.h" - -#include -#include -#include - -namespace ckw -{ - -namespace prototype -{ -class GpuKernelWriterDataHolder; -} // namespace prototype - -/** The target for kernel writer to write into. */ -class Kernel -{ -public: - /** Constructor - * - * @param[in] name The name of the kernel function. - * @param[in] language The programming language to write the kernel. - */ - Kernel(const char *name, GpuTargetLanguage language); - - /** Destructor */ - ~Kernel(); - - /** Get the name of the kernel function. */ - const std::string &name() const; - - /** (Internal use only) Get the map from operand name to the operand declared in this kernel. */ - const ::std::map<::std::string, ::std::unique_ptr> &operands() const; - - /** (Internal use only) Get the map from operand name to the operand declared in this kernel. */ - ::std::map<::std::string, ::std::unique_ptr> &operands(); - - /** (Internal use only) Get the implementation data. */ - prototype::GpuKernelWriterDataHolder *impl(); - -private: - ::std::string _name; - ::std::unique_ptr _kernel; - ::std::map<::std::string, ::std::unique_ptr> _operands; -}; - -} // namespace ckw - -#endif // CKW_INCLUDE_CKW_KERNEL_H diff --git a/compute_kernel_writer/include/ckw/KernelWriter.h b/compute_kernel_writer/include/ckw/KernelWriter.h deleted file mode 100644 index 5dce62a14c..0000000000 --- a/compute_kernel_writer/include/ckw/KernelWriter.h +++ /dev/null @@ -1,217 +0,0 @@ -/* - * 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 CKW_INCLUDE_CKW_KERNELWRITER_H -#define CKW_INCLUDE_CKW_KERNELWRITER_H - -#include "ckw/Kernel.h" -#include "ckw/TensorInfo.h" -#include "ckw/TensorOperand.h" -#include "ckw/TileInfo.h" -#include "ckw/TileOperand.h" - -#include - -namespace ckw -{ - -namespace prototype -{ -struct GpuKernelWriterAttribute; -class IGpuKernelWriter; -} // namespace prototype - -/** Kernel writer. */ -class KernelWriter -{ -public: - // ============================================================================================= - // Constructors and destructor - // ============================================================================================= - - /** Initialize a new instance of kernel writer. - * - * @param[in] kernel The kernel to be written to. - */ - explicit KernelWriter(Kernel &kernel); - - /** Destructor */ - ~KernelWriter(); - - /** No copy constructor. */ - KernelWriter(const KernelWriter &) = delete; - - /** No copy assignment. */ - KernelWriter &operator=(const KernelWriter &) = delete; - - // ============================================================================================= - // Scope management - // ============================================================================================= - - /** Get the current ID space. */ - int32_t id_space() const; - - /** Set the current ID space. */ - KernelWriter &id_space(int32_t id_space); - - /** Switch to and return a new ID space. */ - int32_t next_id_space(); - - // ============================================================================================= - // Tensor and tile declaration - // ============================================================================================= - - /** Define a tensor argument. - * - * @param[in] name The name of the tensor. - * @param[in] info The tensor info. - * - * @return The @ref TensorOperand object. - */ - TensorOperand &create_tensor_argument(const char *name, const TensorInfo &info); - - /** Define a compile-time constant scalar argument. - * - * @param[in] name The name of the tile. - * @param[in] value The value of the tile. - * - * @return The @ref TileOperand object. - */ - TileOperand &create_tile_argument(const char *name, int32_t value); - - /** Declare a new tile. - * - * The name of the tile must be unique in the current ID space. - * - * @param[in] name The name of the tile. - * @param[in] ... The necessary arguments to create a new @ref TileOperand. - * - * @return The @ref TileOperand object. - */ - template - TileOperand &declare_tile(const char *name, TArgs &&...args) - { - const auto var_name = generate_variable_name(name); - auto operand = new TileOperand(var_name, ::std::forward(args)...); - register_operand(operand, true); - - return *operand; - } - - // ============================================================================================= - // Load and store - // ============================================================================================= - - /** Load the data from the tensor memory to the tile using the sampling information. - * - * @param[out] tile The tile to be loaded. - * @param[in] tensor The tensor to be read. - * @param[in] sampler The tensor sampling information. - */ - void op_load(TileOperand &tile, TensorOperand &tensor, const TensorTileSampler &sampler); - - /** Store the tile to the tensor using the specified sampling information. - * - * @param[out] dst The tensor that the tile is written to. - * @param[in] src The tile to be stored. - * @param[in] sampler The tensor sampling information. - */ - void op_store(TensorOperand &tensor, const TileOperand &tile, const TensorTileSampler &sampler); - - // ============================================================================================= - // Data processing - // ============================================================================================= - - /** Write assignment: ` = `. - * - * @param[in] dst The destination tile. - * @param[in] src The source tile. - */ - void op_assign(TileOperand &dst, const TileOperand &src); - - /** Write binary expression: ` = `. - * - * @param[in] dst The destination tile. - * @param[in] lhs The LHS operand. - * @param[in] rhs The RHS operand. - * @param[in] op The binary operator. - */ - void op_binary_expression(TileOperand &dst, const TileOperand &lhs, const TileOperand &rhs, BinaryOp op); - - /** Write function applied to scalar value: ` = ()`. - * - * @param[in] dst The destination tile. - * @param[in] src The source tile. - * @param[in] func The function to be applied to the source tile. - */ - void op_scalar_function(TileOperand &dst, const TileOperand &src, ScalarUnaryFunction func); - - // ============================================================================================= - // Misc - // ============================================================================================= - - /** Set `dst` the global ID of dimension `dim`. - * - * @param[in] dst The tile to be written to. - * @param[in] dim The global ID dimension. - */ - void op_get_global_id(TileOperand &dst, int32_t dim); - - // ============================================================================================= - // Code generation - // ============================================================================================= - - /** Generate the source code of the kernel. */ - ::std::string generate_code(); - -private: - /** Generate the full variable name based on the original name and the ID space. - * - * @param[in] name The name of the variable. - * - * @return The full variable name. - */ - ::std::string generate_variable_name(const char *name) const; - - /** Register the operand to the kernel. - * - * The operand is uniquely owned by the kernel afterward. - * - * @param[in] operand The operand to be registered. - * @param[in] declaring Whether the tile declaration is generated. - */ - void register_operand(OperandBase *operand, bool declaring); - -private: - Kernel *_kernel; - ::std::unique_ptr _impl_attr; - ::std::unique_ptr _impl; - - int32_t _id_space{ 0 }; - int32_t _max_id_space{ 0 }; -}; - -} // namespace ckw - -#endif // CKW_INCLUDE_CKW_KERNELWRITER_H diff --git a/compute_kernel_writer/include/ckw/OperandBase.h b/compute_kernel_writer/include/ckw/OperandBase.h deleted file mode 100644 index 0ea5030968..0000000000 --- a/compute_kernel_writer/include/ckw/OperandBase.h +++ /dev/null @@ -1,76 +0,0 @@ -/* - * 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 CKW_INCLUDE_CKW_OPERANDBASE_H -#define CKW_INCLUDE_CKW_OPERANDBASE_H - -#include "ckw/Types.h" -#include - -namespace ckw -{ -namespace prototype -{ -class IGpuKernelWriter; -class Operand; -} // namespace prototype - -/** The base class for all operands. */ -class OperandBase -{ -public: - /** Constructor - * - * @param[in] name The name of the operand. - */ - explicit OperandBase(const ::std::string &name); - - /** Destructor */ - virtual ~OperandBase(); - - /** (Internal use only) Create the implementation operand. - * - * @param[in] writer The implementation kernel writer. - */ - virtual prototype::Operand create_impl_operand(prototype::IGpuKernelWriter *writer) const = 0; - - /** Get the name of the operand. */ - const ::std::string &name() const; - - /** Set the name of the operand. */ - OperandBase &name(const ::std::string &name); - - /** Get the data type of the operand. */ - virtual DataType data_type() const = 0; - - /** Get whether the operand is compile-time constant. */ - virtual bool is_constant() const = 0; - -private: - ::std::string _name; -}; - -} // namespace ckw - -#endif // CKW_INCLUDE_CKW_OPERANDBASE_H diff --git a/compute_kernel_writer/include/ckw/ScalarValue.h b/compute_kernel_writer/include/ckw/ScalarValue.h deleted file mode 100644 index cf017d435f..0000000000 --- a/compute_kernel_writer/include/ckw/ScalarValue.h +++ /dev/null @@ -1,137 +0,0 @@ -/* - * 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 CKW_INCLUDE_CKW_SCALARVALUE_H -#define CKW_INCLUDE_CKW_SCALARVALUE_H - -#include "ckw/Error.h" - -#include - -namespace ckw -{ - -/** The scalar value known at compile-time. */ -class ScalarValue -{ -public: - /** Initialize a new instance of @ref ScalarValue class with integer value 0. */ - ScalarValue() - { - _type = Type::INT; - _value.i64 = 0; - } - - /** Initialize a new instance of @ref ScalarValue class with the specified value. */ - template - ScalarValue(T value) - { - set(value); - } - - /** Set the value. */ - template - void set(T value) - { - CKW_ASSERT(::std::is_integral::value || ::std::is_floating_point::value); - CKW_ASSERT(sizeof(T) <= 8); - - _size = sizeof(T); - - if(::std::is_integral::value) - { - if(::std::is_signed::value) - { - _type = Type::INT; - _value.i64 = value; - } - else - { - _type = Type::UINT; - _value.u64 = value; - } - } - else - { - _type = Type::FLOAT; - _value.f64 = value; - } - } - - /** Get the value. - * - * The caller must make sure that what has been stored in the object must fit - * the output data type without data corruption or loss of accuracy. - */ - template - T get() const - { - CKW_ASSERT(::std::is_integral::value || ::std::is_floating_point::value); - CKW_ASSERT(sizeof(T) >= _size); - - if(::std::is_integral::value) - { - if(::std::is_signed::value) - { - CKW_ASSERT(_type == Type::INT || _type == Type::UINT); - CKW_ASSERT_IF(_type == Type::UINT, sizeof(T) > _size); - - return _value.i64; - } - else - { - CKW_ASSERT(_type == Type::INT); - - return _value.u64; - } - } - else - { - return _value.f64; - } - } - -private: - union Value - { - int64_t i64; - uint64_t u64; - double f64; - }; - - enum class Type : int32_t - { - UINT, - INT, - FLOAT, - }; - - Value _value{}; - Type _type{}; - uint32_t _size{}; -}; - -} // namespace ckw - -#endif // CKW_INCLUDE_CKW_SCALARVALUE_H diff --git a/compute_kernel_writer/include/ckw/TensorOperand.h b/compute_kernel_writer/include/ckw/TensorOperand.h deleted file mode 100644 index 130ab596fb..0000000000 --- a/compute_kernel_writer/include/ckw/TensorOperand.h +++ /dev/null @@ -1,181 +0,0 @@ -/* - * 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 CKW_INCLUDE_CKW_TENSOROPERAND_H -#define CKW_INCLUDE_CKW_TENSOROPERAND_H - -#include "ckw/OperandBase.h" -#include "ckw/TensorInfo.h" -#include "ckw/TensorTileSampler.h" -#include "ckw/TileOperand.h" -#include "ckw/Types.h" - -#include - -namespace ckw -{ - -class TensorComponentOperand; - -// ================================================================================================= -// TensorOperand -// ================================================================================================= - -/** Tensor operand */ -class TensorOperand : public OperandBase -{ -public: - /** Initialize a new instance of @ref TensorOperand class. - * - * @param[in] name The name of the tensor. - * @param[in] info The tensor info. - */ - TensorOperand(const ::std::string &name, const TensorInfo &info); - - /** No copy constructor. */ - TensorOperand(const TensorOperand &other) = delete; - - /** No copy assignment. */ - TensorOperand &operator=(const TensorOperand &other) = delete; - - /** (Internal use only) Create the implementation operand. - * - * @param[in] writer The implementation kernel writer. - */ - virtual prototype::Operand create_impl_operand(prototype::IGpuKernelWriter *writer) const override; - - /** Get the tensor info. */ - const TensorInfo &info() const; - - /** Get the tensor info. */ - TensorInfo &info(); - - /** Get the data type. */ - virtual DataType data_type() const override; - - /** Get whether the tensor is compile-time constant. */ - virtual bool is_constant() const override; - - /** Get the default tile attached to the tensor. */ - const TileOperand &tile() const; - - /** Get the default tile attached to the tensor. */ - TileOperand &tile(); - - /** Set the default tile attached to the tensor. */ - TensorOperand &tile(TileOperand &tile); - - /** Get the tensor sampler of the default tile. */ - const TensorTileSampler &tile_sampler() const; - - /** Get the tensor sampler of the default tile. */ - TensorTileSampler &tile_sampler(); - - /** Set the tensor sampler of the default tile. */ - TensorOperand &tile_sampler(const TensorTileSampler &value); - - /** Get the operand that contains the stride in y dimension of the tensor. */ - TileOperand &stride1(); - - /** Get the operand that contains the stride in z dimension of the tensor. */ - TileOperand &stride2(); - - /** Get the operand that contains the stride in w dimension of the tensor. */ - TileOperand &stride3(); - - /** Get the operand that contains the stride in w dimension of the tensor. */ - TileOperand &stride4(); - - /** Get the operand that contains the size of dimension 0 of the tensor. */ - TileOperand &dim0(); - - /** Get the operand that contains the size of dimension 1 of the tensor. */ - TileOperand &dim1(); - - /** Get the operand that contains the size of dimension 2 of the tensor. */ - TileOperand &dim2(); - - /** Get the operand that contains the size of dimension 3 of the tensor. */ - TileOperand &dim3(); - - /** Get the operand that contains the size of dimension 4 of the tensor. */ - TileOperand &dim4(); - - /** Get the operand that contains the size of dimensions 1 and 2 collapsed. */ - TileOperand &dim1_dim2(); - - /** Get the operand that contains the size of dimensions 1, 2 and 3 collapsed. */ - TileOperand &dim1_dim2_dim3(); - - /** Get the operand that contains the offset in bytes to the first element. */ - TileOperand &offset_first_element_in_bytes(); - -private: - TensorInfo _info; - - TileOperand *_tile{ nullptr }; - TensorTileSampler _tile_sampler{}; - - ::std::unique_ptr _stride1{ nullptr }; - ::std::unique_ptr _stride2{ nullptr }; - ::std::unique_ptr _stride3{ nullptr }; - ::std::unique_ptr _stride4{ nullptr }; - ::std::unique_ptr _dim0{ nullptr }; - ::std::unique_ptr _dim1{ nullptr }; - ::std::unique_ptr _dim2{ nullptr }; - ::std::unique_ptr _dim3{ nullptr }; - ::std::unique_ptr _dim4{ nullptr }; - ::std::unique_ptr _dim1_dim2{ nullptr }; - ::std::unique_ptr _dim1_dim2_dim3{ nullptr }; - ::std::unique_ptr _offset_first_element_in_bytes{ nullptr }; -}; - -// ================================================================================================= -// TensorComponentOperand -// ================================================================================================= - -/** Tile operand that contains tensor information. */ -class TensorComponentOperand : public TileOperand -{ -public: - /** Initialize a new instance of @ref TensorComponentOperand class. - * - * @param[in] name The name of the operand. - * @param[in] component The tensor info component. - */ - TensorComponentOperand(const ::std::string &name, TensorComponent component); - - /** (Internal use only) Create the implementation operand. - * - * @param[in] writer The implementation kernel writer. - */ - virtual prototype::Operand create_impl_operand(prototype::IGpuKernelWriter *writer) const override; - -private: - TensorComponent _component; -}; - -} // namespace ckw - -#endif // CKW_INCLUDE_CKW_TENSOROPERAND_H diff --git a/compute_kernel_writer/include/ckw/TensorTileSampler.h b/compute_kernel_writer/include/ckw/TensorTileSampler.h deleted file mode 100644 index 5ef7bca647..0000000000 --- a/compute_kernel_writer/include/ckw/TensorTileSampler.h +++ /dev/null @@ -1,163 +0,0 @@ -/* - * 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 CKW_INCLUDE_CKW_TENSORTILESAMPLER_H -#define CKW_INCLUDE_CKW_TENSORTILESAMPLER_H - -#include "ckw/Types.h" -#include - -namespace ckw -{ - -class TileOperand; - -/** Tensor sampler - * - * It contains information about how the result tile should be stored to tensor memory. - * It can also be used to dictate how the subsequent operators fetch the input tensor. - */ -class TensorTileSampler -{ -public: - /** Initialize a new instance of @ref TensorSampler class. */ - TensorTileSampler(); - - /** Initialize a new instance of @ref TensorSampler class. - * - * @param[in] x The coordinate in the x dimension. - * @param[in] y The coordinate in the y dimension. - * @param[in] z The coordinate in the z dimension. - * @param[in] b The coordinate in the batch dimension. - * @param[in] format The tensor data format. - * @param[in] address_mode_x The address mode of the x dimension. - * @param[in] address_mode_y The address mode of the y dimension. - * @param[in] address_mode_z The address mode of the z dimension. - */ - TensorTileSampler( - TileOperand &x, TileOperand &y, TileOperand &z, TileOperand &b, - TensorSamplerFormat format, - TensorSamplerAddressModeX address_mode_x, - TensorSamplerAddressModeY address_mode_y, - TensorSamplerAddressModeZ address_mode_z); - - /** Initialize a new instance of @ref TensorSampler class. - * - * @param[in] x The coordinate in the x dimension. - * @param[in] y The coordinate in the y dimension. - * @param[in] z The coordinate in the z dimension. - * @param[in] b The coordinate in the batch dimension. - * @param[in] height The height of the tile. - * @param[in] width The width of the tile. - * @param[in] format The tensor data format. - * @param[in] address_mode_x The address mode of the x dimension. - * @param[in] address_mode_y The address mode of the y dimension. - * @param[in] address_mode_z The address mode of the z dimension. - */ - TensorTileSampler( - TileOperand &x, TileOperand &y, TileOperand &z, TileOperand &b, - int32_t height, int32_t width, - TensorSamplerFormat format, - TensorSamplerAddressModeX address_mode_x, - TensorSamplerAddressModeY address_mode_y, - TensorSamplerAddressModeZ address_mode_z); - - /** Get the coordinate in the x dimension. */ - const TileOperand &x() const; - - /** Set the coordinate in the x dimension. */ - TensorTileSampler &x(TileOperand &x); - - /** Get the coordinate in the y dimension. */ - const TileOperand &y() const; - - /** Set the coordinate in the y dimension. */ - TensorTileSampler &y(TileOperand &y); - - /** Get the coordinate in the z dimension. */ - const TileOperand &z() const; - - /** Set the coordinate in the z dimension. */ - TensorTileSampler &z(TileOperand &z); - - /** Get the coordinate in the batch dimension. */ - const TileOperand &b() const; - - /** Set the coordinate in the batch dimension. */ - TensorTileSampler &b(TileOperand &b); - - /** Get the width of the tile. */ - int32_t width() const; - - /** Set the width of the tile. */ - TensorTileSampler &width(int32_t width); - - /** Get the height of the tile. */ - int32_t height() const; - - /** Set the height of the tile. */ - TensorTileSampler &height(int32_t height); - - /** Get the format of the tensor. */ - TensorSamplerFormat format() const; - - /** Set the format of the tensor. */ - TensorTileSampler &format(TensorSamplerFormat format); - - /** Get the address mode of the x dimension. */ - TensorSamplerAddressModeX address_mode_x() const; - - /** Set the address mode of the x-dimension. */ - TensorTileSampler &address_mode_x(TensorSamplerAddressModeX address_mode_x); - - /** Get the address mode of the y dimension. */ - TensorSamplerAddressModeY address_mode_y() const; - - /** Set the address mode of the y dimension. */ - TensorTileSampler &address_mode_y(TensorSamplerAddressModeY address_mode_y); - - /** Get the address mode of the z dimension. */ - TensorSamplerAddressModeZ address_mode_z() const; - - /** Set the address mode of the z dimension. */ - TensorTileSampler &address_mode_z(TensorSamplerAddressModeZ address_mode_z); - -private: - TileOperand *_x{ nullptr }; - TileOperand *_y{ nullptr }; - TileOperand *_z{ nullptr }; - TileOperand *_b{ nullptr }; - - int32_t _height{ 0 }; - int32_t _width{ 0 }; - - TensorSamplerFormat _format{ TensorSamplerFormat::Unknown }; - TensorSamplerAddressModeX _address_mode_x{ TensorSamplerAddressModeX::Unknown }; - TensorSamplerAddressModeY _address_mode_y{ TensorSamplerAddressModeY::Unknown }; - TensorSamplerAddressModeZ _address_mode_z{ TensorSamplerAddressModeZ::Unknown }; -}; - -} // namespace ckw - -#endif // CKW_INCLUDE_CKW_TENSORTILESAMPLER_H diff --git a/compute_kernel_writer/include/ckw/TileOperand.h b/compute_kernel_writer/include/ckw/TileOperand.h deleted file mode 100644 index 1eee18589f..0000000000 --- a/compute_kernel_writer/include/ckw/TileOperand.h +++ /dev/null @@ -1,110 +0,0 @@ -/* - * 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 CKW_INCLUDE_CKW_TILEOPERAND_H -#define CKW_INCLUDE_CKW_TILEOPERAND_H - -#include "ckw/Error.h" -#include "ckw/OperandBase.h" -#include "ckw/ScalarValue.h" -#include "ckw/TileInfo.h" - -#include - -namespace ckw -{ - -class Kernel; - -/** Tile operand which can be either scalar, vector or 2D tile. */ -class TileOperand : public OperandBase -{ -public: - /** Initialize a new instance of @ref TileOperand class with the tile information. - * - * @param[in] name The name of the tile. - * @param[in] tile_info The tile info. - */ - TileOperand(const ::std::string &name, const TileInfo &tile_info); - - /** Initialize a new instance of @ref TileOperand for scalar variable. - * - * @param[in] name The name of the tile. - * @param[in] data_type The data type of the tile. - */ - TileOperand(const ::std::string &name, DataType data_type); - - /** Initialize a new instance of @ref TileOperand for compile-time constant scalar variable. - * - * @param[in] name The name of the tile. - * @param[in] value The value of the tile. - */ - TileOperand(const ::std::string &name, int32_t value); - - /** Initialize a new instance of @ref TileOperand for compile-time constant scalar variable. - * - * @param[in] name The name of the tile. - * @param[in] value The value of the tile. - */ - TileOperand(const ::std::string &name, float value); - - /** Prohibit copy of tile operand. */ - TileOperand(const TileOperand &) = delete; - - /** Prohibit copy of tile operand. */ - TileOperand &operator=(const TileOperand &) = delete; - - /** (Internal use only) Create the implementation operand. - * - * @param[in] writer The implementation kernel writer. - */ - virtual prototype::Operand create_impl_operand(prototype::IGpuKernelWriter *writer) const override; - - /** Get the tile info. */ - const TileInfo &tile_info() const; - - /** Get the data type of the tile. */ - virtual DataType data_type() const override; - - /** Get whether the tile is compile-time constant. */ - virtual bool is_constant() const override; - - /** Get whether the tile is a scalar value. */ - bool is_scalar() const; - - /** Get the scalar value of the tile. - * - * The tile must have the shape of 1, 1 (i.e. scalar). - */ - ScalarValue scalar_value() const; - -private: - TileInfo _info; - ScalarValue _value{}; - bool _constant; -}; - -} // namespace ckw - -#endif // CKW_INCLUDE_CKW_TILEOPERAND_H diff --git a/compute_kernel_writer/include/ckw/Types.h b/compute_kernel_writer/include/ckw/Types.h index 5516718e54..c9f80b65e0 100644 --- a/compute_kernel_writer/include/ckw/Types.h +++ b/compute_kernel_writer/include/ckw/Types.h @@ -21,120 +21,25 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE * SOFTWARE. */ - -#ifndef CKW_INCLUDE_CKW_TYPES_H -#define CKW_INCLUDE_CKW_TYPES_H - -#include -#include +#ifndef COMPUTE_KERNEL_WRITER_INCLUDE_CKW_TYPES_H +#define COMPUTE_KERNEL_WRITER_INCLUDE_CKW_TYPES_H namespace ckw { - /** Compute Kernel Writer data types. This data type is used by the code variables and tensor arguments. */ enum class DataType -{ - Unknown = 0x00, - Fp32 = 0x11, - Fp16 = 0x12, - Int32 = 0x21, - Int16 = 0x22, - Int8 = 0x24, - Uint32 = 0x31, - Uint16 = 0x32, - Uint8 = 0x34, - Bool = 0x41 -}; - -enum class GpuTargetLanguage { Unknown, - OpenCL -}; - -/* Binary operations -*/ -enum class BinaryOp : int32_t -{ - // Elementwise - Add = 0x0000, // + - Sub = 0x0001, // - - Mul = 0x0002, // * - Div = 0x0003, // / - Mod = 0x0004, // % - // Relational - Equal = 0x1000, // == - Less = 0x1001, // < - LessEqual = 0x1002, // <= - Greater = 0x1003, // > - GreaterEqual = 0x1004, // >= - // Algebra - MatMul_Nt_Nt = 0x2000, // X - MatMul_Nt_T = 0x2001, // X - MatMul_T_Nt = 0x2002, // X - MatMul_T_T = 0x2003, // X - Dot = 0x2004, // . - // Logical - LogicalAnd = 0x3000, // && - LogicalOr = 0x3001, // || - LogicalNot = 0x3002 // ! -}; - -enum class AssignmentOp : int32_t -{ - // Unary - Increment = 0x0000, // += - Decrement = 0x0001, // -= -}; - -enum class ScalarUnaryFunction : int32_t -{ - Exp, -}; - -enum class TensorSamplerFormat : int32_t -{ - Unknown = 0, - C_WH_1 = 1, - C_W_H = 2 -}; - -enum class TensorSamplerAddressModeX : int32_t -{ - Unknown = 0, - None = 1, // The user guarantees that the X coordinate is always in-bound - OverlappingMin = 2 // (FIXED shapes only) Reduce the load/store length when x == 0 (MIN). The load length will be width % original length - // Leftover elements can be handled using overlapping. This involves processing some of the elements in the array twice. + Fp32, + Fp16, + Int32, + Int16, + Int8, + Uint32, + Uint16, + Uint8, + Bool }; - -enum class TensorSamplerAddressModeY : int32_t -{ - Unknown = 0, - None = 1, // The user guarantees that the Y coordinate is always in-bound - OverlappingMin = 2, // (FIXED shapes only) Reduce the load/store length when x == 0 (MIN). The load length will be width % original length - Skip = 3, // Skip the read/write - SkipMinEdgeOnly = 4, // Skip greater than or equal to max only. The user guarantees that the Y coordinate is always >= 0 - SkipMaxEdgeOnly = 5, // Skip less than 0 only - ClampToNearest = 6, // Clamp the coordinate to nearest edge (0 or max value allowed on Y) - ClampToMinEdgeOnly = 7, // Clamp the negative coordinate to 0 only. Therefore, we expect Y to be always < MAX - ClampToMaxEdgeOnly = 8, // Clamp the coordinate to the max value allowed on Y only. We expect Y to be always >= 0 - ClampToBorder = 9, // Clamp to border which always has 0 value - ClampToBorderMinEdgeOnly = 10, - ClampToBorderMaxEdgeOnly = 11 -}; - -enum class TensorSamplerAddressModeZ : int32_t -{ - Unknown = 0, - None = 1, // The user guarantees that the Y coordinate is always in-bound - Skip = 3, // Skip the read/write - SkipMinEdgeOnly = 4, // Skip greater than or equal to max only. The user guarantees that the Y coordinate is always >= 0 - SkipMaxEdgeOnly = 5, // Skip less than 0 only - ClampToNearest = 6, // Clamp the coordinate to nearest edge (0 or max value allowed on Y) - ClampToMinEdgeOnly = 7, // Clamp the negative coordinate to 0 only. Therefore, we expect Y to be always < MAX - ClampToMaxEdgeOnly = 8, // Clamp the coordinate to the max value allowed on Y only. We expect Y to be always >= 0 -}; - } // namespace ckw -#endif // CKW_INCLUDE_CKW_TYPES_H +#endif /* COMPUTE_KERNEL_WRITER_INCLUDE_CKW_TYPES_H */ -- cgit v1.2.1