/* * 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_PROTOTYPE_INCLUDE_CKW_KERNELWRITER_H #define CKW_PROTOTYPE_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 "ckw/types/ConvertPolicy.h" #include "ckw/types/Functions.h" #include "ckw/types/Operators.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 // ============================================================================================= /** Declare a tensor argument. * * @param[in] name The name of the tensor. * @param[in] info The tensor info. * @param[in] storage_type The tensor storage type. * * @return The @ref TensorOperand object. */ TensorOperand &declare_tensor_argument(const std::string &name, const TensorInfo &info, TensorStorageType storage_type = TensorStorageType::BufferUint8Ptr); /** Declare 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 &declare_tile_argument(const std::string &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 std::string &name, TArgs &&...args) { const auto var_name = generate_variable_name(name); auto operand = std::make_unique(var_name, ::std::forward(args)...); return declare_tile_operand(std::move(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. * @param[in] dilation_y Dilation in the Y dimension. */ void op_load(TileOperand &tile, const TensorOperand &tensor, const TensorTileSampler &sampler, const TileOperand &dilation_y = TileOperand("dil_y", 1)); /** Load the data from the tensor memory to the tile using the indirect buffer approach and respective of 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_indirect(TileOperand &tile, const TensorOperand &tensor, const TensorTileSampler &sampler); /** Construct an indirection buffer in @p tile containing the precalculated addresses of elements in the source tensor. * * @param[out] tile The tile to be loaded. * @param[in] tensor The tensor the be read. * @param[in] sampler The tensor sampling information. * @param[in] x The X coordinate. * @param[in] y The Y coordinate. * @param[in] x_off Offset in the X dimension. * @param[in] y_off Offset in the Y dimension. */ void util_get_indirect_buffer(TileOperand &tile, const TensorOperand &tensor, const TensorTileSampler &sampler, const TileOperand &x, const TileOperand &y, const TileOperand &x_off, const TileOperand &y_off); /** 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[out] dst The destination tile. * @param[in] src The source tile. */ void op_assign(const TileOperand &dst, const TileOperand &src); /** Write the cast: ` = convert_<_sat>();`. * * @param[out] dst The destination tile. * @param[in] src The source tile. * @param[in] policy The policy governing the behavior of the cast. */ void op_cast_expression(const TileOperand &dst, const TileOperand &src, ConvertPolicy policy); /** Write the unary expression: ` = `. * * @param[out] dst The destination tile. * @param[in] op The unary operator. * @param[in] src The source tile. */ void op_unary_expression(const TileOperand &dst, UnaryOp op, const TileOperand &src); /** Write binary expression: ` = ;`. * * @param[out] dst The destination tile. * @param[in] lhs The LHS tile. * @param[in] op The binary operator. * @param[in] rhs The RHS tile. */ void op_binary_expression(const TileOperand &dst, const TileOperand &lhs, BinaryOp op, const TileOperand &rhs); /** Write function applied to scalar value: ` = ();`. * * @param[out] dst The destination tile. * @param[in] func The function to be applied to the source tile. * @param[in] src The source tile. */ void op_unary_elementwise_function(const TileOperand &dst, UnaryFunction func, const TileOperand &src); /** Write function applied to scalar value: ` = (, );`. * * @param[out] dst The destination tile. * @param[in] func The function to be applied to the source tiles. * @param[in] first The first argument tile. * @param[in] second The second argument tile. */ void op_binary_elementwise_function(const TileOperand &dst, BinaryFunction func, const TileOperand &first, const TileOperand &second); /** Write function applied to scalar value: ` = (, , );`. * * @param[out] dst The destination tile. * @param[in] func The function to be applied to the source tiles. * @param[in] first The first argument tile. * @param[in] second The second argument tile. * @param[in] third The third argument tile. */ void op_ternary_elementwise_function(const TileOperand &dst, TernaryFunction func, const TileOperand &first, const TileOperand &second, const TileOperand &third); /** Write if-statement: `if( ) { }`. * * @param[in] lhs The LHS tile of the condition. * @param[in] op The relational binary operator. * @param[in] rhs The RHS tile of the condition. * @param[in] body The body of the if-statement. */ void op_if(const TileOperand &lhs, BinaryOp op, const TileOperand &rhs, const std::function &body); /** Write else-if-statement: `else if( ) { }`. * * @param[in] lhs The LHS tile of the condition. * @param[in] op The relational binary operator. * @param[in] rhs The RHS tile of the condition. * @param[in] body The body of the else-if-statement. */ void op_else_if(const TileOperand &lhs, BinaryOp op, const TileOperand &rhs, const std::function &body); /** Write an else-statement: `else { }`. * * @param[in] body The body of the else-statement. */ void op_else(const std::function &body); /** Write for-loops: `for(; ; ) { body }`. * * @param[in] var_name The name of the variable used in condition. * @param[in] cond_op The relational binary operator used in condition. * @param[in] cond_value_name The value which the variable is compared against. * @param[in] update_var_name The name of the variable which is updated. * @param[in] update_op The assignment operator used for updating the update value. * @param[in, out] update_value The value which is updated at every iteration. * @param[in] body The body of the for-loop. */ void op_for_loop(const TileOperand &var_name, BinaryOp cond_op, const TileOperand &cond_value_name, const TileOperand &update_var_name, AssignmentOp update_op, const TileOperand &update_value_name, const std::function &body); /** Write the return statement: `return;` */ void op_return(); // ============================================================================================= // Misc // ============================================================================================= /** Set `dst` the global ID of dimension `dim`. * * @param[out] dst The tile to be written to. * @param[in] dim The global ID dimension. */ void op_get_global_id(const 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 std::string &name) const; /** Declare the tile operand. * * @param[in] operand The tile operand to be declared. */ TileOperand &declare_tile_operand(std::unique_ptr operand); 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_PROTOTYPE_INCLUDE_CKW_KERNELWRITER_H