| /* |
| * 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/TensorSampler.h" |
| #include "ckw/TileInfo.h" |
| #include "ckw/TileOperand.h" |
| #include "ckw/types/ConstantData.h" |
| #include "ckw/types/ConvertPolicy.h" |
| #include "ckw/types/DataType.h" |
| #include "ckw/types/Operators.h" |
| #include "ckw/types/TargetArchitecture.h" |
| #include "ckw/types/TargetLanguage.h" |
| #include "ckw/types/TensorComponentType.h" |
| #include "ckw/types/TensorDataLayout.h" |
| #include "ckw/types/TensorSamplerTypes.h" |
| #include "ckw/types/TensorStorageType.h" |
| |
| #include <functional> |
| #include <memory> |
| #include <string> |
| #include <tuple> |
| |
| namespace ckw |
| { |
| |
| /** Forward Declarations */ |
| class TileArea; |
| |
| /** A kernel writer. |
| * |
| * This class is used to construct a new kernel by defining arguments, declaring variable and writing code. |
| * |
| * Use @ref KernelWriter::create_instance method to create the kernel writer for the specific target architecture and language. |
| * |
| * After having finished constructing the kernel, call @ref KernelWriter::emit_kernel to get the kernel object. |
| */ |
| class KernelWriter |
| { |
| public: |
| // ============================================================================================= |
| // Construtors and destructor |
| // ============================================================================================= |
| |
| /** Initialize a new instance of @ref KernelWriter class for the specific architecture and language. |
| * |
| * Supported target architectures and languages: |
| * |
| * Architecture | Languages | |
| * ------------------------------|------------------------------| |
| * GpuArmMaliValhall | OpenCL | |
| * |
| * @param[in] architecture The architecture on which the kernel is executed. |
| * @param[in] language The language to write the kernel. |
| */ |
| static std::unique_ptr<KernelWriter> create_instance(TargetArchitecture architecture, TargetLanguage language); |
| |
| /** Destructor */ |
| virtual ~KernelWriter(); |
| |
| // ============================================================================================= |
| // Data processing |
| // ============================================================================================= |
| |
| /** Write assignment statement: `<dst> = <src>;`. |
| * |
| * @param[in] dst The destination tile. |
| * @param[in] src The source tile. |
| */ |
| virtual void op_assign(const TileOperand &dst, const TileOperand &src) = 0; |
| |
| /** Write the cast statement: `<dst> = convert_<dst.type><policy>(<src>);`. |
| * |
| * @param[in] dst The destination tile. |
| * @param[in] src The source tile. |
| * @param[in] policy The policy governing the behavior of the cast. |
| */ |
| virtual void op_cast(const TileOperand &dst, const TileOperand &src, ConvertPolicy policy) = 0; |
| |
| /** Write the unary expression statement: `<dst> = <op> <src>;`. |
| * |
| * @param[in] dst The destination tile. |
| * @param[in] op The unary operator. |
| * @param[in] src The source tile. |
| */ |
| virtual void op_unary(const TileOperand &dst, UnaryOp op, const TileOperand &src) = 0; |
| |
| /** Write the binary expression statement: `<dst> = <op>(<first>, <second>);`. |
| * |
| * @param[in] dst The destination tile. |
| * @param[in] op The binary operator. |
| * @param[in] first The first source tile. |
| * @param[in] second The second source tile. |
| */ |
| virtual void |
| op_binary(const TileOperand &dst, BinaryOp op, const TileOperand &first, const TileOperand &second) = 0; |
| |
| /** Write ternary expression statement: `<dst> = <op>(<first>, <second>, <third>);`. |
| * |
| * @param[in] dst The destination tile. |
| * @param[in] op The ternary operator. |
| * @param[in] first The first source tile. |
| * @param[in] second The second source tile. |
| * @param[in] third The third source tile. |
| */ |
| virtual void op_ternary(const TileOperand &dst, |
| TernaryOp op, |
| const TileOperand &first, |
| const TileOperand &second, |
| const TileOperand &third) = 0; |
| |
| // ============================================================================================= |
| // Flow control |
| // ============================================================================================= |
| |
| /** Write if block: `if(<lhs> <op> <rhs>) { <body> }`. |
| * |
| * @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 function that writes the body of the if block. |
| */ |
| virtual void |
| op_if(const TileOperand &lhs, BinaryOp op, const TileOperand &rhs, const std::function<void()> &body) = 0; |
| |
| /** Write else-if block: `else if(<lhs> <op> <rhs>) { <body> }`. |
| * |
| * @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 function that writes the body of the else-if block. |
| */ |
| virtual void |
| op_else_if(const TileOperand &lhs, BinaryOp op, const TileOperand &rhs, const std::function<void()> &body) = 0; |
| |
| /** Write an else block: `else { <body> }`. |
| * |
| * @param[in] body The function that writes the body of the else block. |
| */ |
| virtual void op_else(const std::function<void()> &body) = 0; |
| |
| /** Write for-loop block: `for(; <var> <cond_op> <cond_value>; <update_var> <update_op> <update_value>) { body }`. |
| * |
| * @param[in] var The scalar tile used in loop condition. |
| * @param[in] cond_op The relational binary operator used in loop condition. |
| * @param[in] cond_value The value which the variable is compared against. |
| * @param[in] update_var The scalar tile which is updated each iteration. |
| * @param[in] update_op The assignment operator used for updating the update value. |
| * @param[in] update_value The value which is updated at every iteration. |
| * @param[in] body The function that writes the body of the for-loop block. |
| */ |
| virtual void op_for_loop(const TileOperand &var, |
| BinaryOp cond_op, |
| const TileOperand &cond_value, |
| const TileOperand &update_var, |
| AssignmentOp update_op, |
| const TileOperand &update_value, |
| const std::function<void()> &body) = 0; |
| |
| /** Write the return statement. */ |
| virtual void op_return() = 0; |
| |
| // ============================================================================================= |
| // Misc |
| // ============================================================================================= |
| |
| /** Write the statement to get the global ID of the specified dimension. |
| * |
| * @param[in] dst The tile to write the global ID into. |
| * @param[in] dim The dimension. |
| */ |
| virtual void op_get_global_id(const TileOperand &dst, int32_t dim) = 0; |
| |
| /** Write the line comment in debug build. |
| * |
| * This function does not take effect on release build. |
| * |
| * The comment must only contain one line (i.e. no newline character is allowed). |
| * |
| * @param[in] text The comment to be written. |
| */ |
| virtual void op_comment(const std::string &text) = 0; |
| |
| /** Write the statement to print out the value of all the specified tiles. |
| * |
| * The printing statement is constructed so that the prefix and each of the operand are printed in separate lines. |
| * The format for each operand varies depending on whether it is a 2D tile, a vector or a scalar value. |
| * |
| * Example output of the printing statement when it is executed: |
| * |
| * prefix |
| * scalar_name = scalar_value |
| * vector_name = [vector_value_0, vector_value_1, vector_value_2] |
| * tile_name = [[tile_value_00, tile_value_01], [tile_value_10, tile_value_11]] |
| * |
| * @param[in] prefix The first string to be printed out before the list of operands. |
| * @param[in] operands The list of tiles to be included in the printing statement. |
| */ |
| virtual void op_print(const std::string &prefix, const std::vector<TileOperand> &operands) = 0; |
| |
| /** Write the given raw code to kernel source code |
| * It's used to address the cases where the user needs to |
| * explicitly add a code where it's not (yet) supported by |
| * the kernel writer utility calls. |
| * |
| * @param[in] raw_code raw code to write as string |
| */ |
| virtual void op_write_raw_code(const std::string &raw_code) = 0; |
| |
| // ============================================================================================= |
| // Code generation |
| // ============================================================================================= |
| |
| /** Emit the kernel object. |
| * |
| * @param[in] name The name of the kernel object to be generated. |
| */ |
| virtual std::unique_ptr<Kernel> emit_kernel(const std::string &name) = 0; |
| |
| // ============================================================================================= |
| // Tensor and tile declaration |
| // ============================================================================================= |
| |
| /** Declare a tensor argument. |
| * |
| * @param[in] name The name of the tensor. |
| * @param[in] info The tensor info. |
| * |
| * @return The @ref TensorOperand object. |
| */ |
| virtual TensorOperand declare_tensor_argument(const std::string &name, const TensorInfo &info) = 0; |
| |
| /** Declare a tile given its name and tile info |
| * |
| * @param[in] name Name of the tile |
| * @param[in] tile_info Shape and data type of the tile |
| * |
| * @return The created tile operand |
| */ |
| virtual TileOperand declare_tile(const std::string &name, const TileInfo &tile_info) = 0; |
| |
| /** Declare a constant tile given a @ref:ConstantData object |
| * |
| * @param[in] data a @ref ckw::ConstantData object that has the values and the |
| * underlying data type of the constant tile |
| * |
| * @return The created constant tile operand |
| */ |
| virtual TileOperand declare_constant_tile(const ConstantData &data) = 0; |
| |
| /** Load the data from the tensor memory to the tile using the sampling information. |
| * |
| * @param[in] tile_op The tile to be loaded. |
| * @param[in] tensor_op The tensor to be read. |
| * @param[in] sampler The tensor sampling information. |
| * @param[in] x x-coordinate |
| * @param[in] y y-coordinate |
| * @param[in] z z-coordinate |
| * @param[in] batch batch |
| */ |
| virtual void op_load(const TileOperand &tile_op, |
| const TensorOperand &tensor_op, |
| TensorSampler &sampler, |
| const TileOperand &x, |
| const TileOperand &y, |
| const TileOperand &z, |
| const TileOperand &batch) = 0; |
| |
| /** Load the data from the tensor memory to the tile in a dilated way using the sampling information. |
| * |
| * Similar to @ref KernelWriter::op_load() and |
| * |
| * @param[in] dilation_x Dilation while reading in x-dimension |
| * @param[in] dilation_y Dilation while reading in y-dimension |
| */ |
| virtual void op_load_dilated(const TileOperand &tile_op, |
| const TensorOperand &tensor_op, |
| TensorSampler &sampler, |
| const TileOperand &x, |
| const TileOperand &y, |
| const TileOperand &z, |
| const TileOperand &batch, |
| const TileOperand &dilation_x, |
| const TileOperand &dilation_y) = 0; |
| |
| /** Store the data to the tensor memory from the tile using the sampling information. |
| * |
| * Similar to @ref KernelWriter::op_load() |
| */ |
| virtual void op_store(const TensorOperand &tensor_op, |
| const TileOperand &tile_op, |
| TensorSampler &sampler, |
| const TileOperand &x, |
| const TileOperand &y, |
| const TileOperand &z, |
| const TileOperand &batch) = 0; |
| |
| /** Store the data to the tensor memory from the tile in a dilated way using the sampling information. |
| * |
| * Similar to @ref KernelWriter::op_load_dilated() |
| */ |
| virtual void op_store_dilated(const TensorOperand &tensor_op, |
| const TileOperand &tile_op, |
| TensorSampler &sampler, |
| const TileOperand &x, |
| const TileOperand &y, |
| const TileOperand &z, |
| const TileOperand &batch, |
| const TileOperand &dilation_x, |
| const TileOperand &dilation_y) = 0; |
| |
| /** Load the data from the tensor memory to the tile using the indirect buffer approach and respecting the sampling information. |
| * |
| * @param[in] tile_op The tile to be loaded. |
| * @param[in] tensor_op The tensor to be read. |
| * @param[in] sampler The tensor sampling information. |
| * @param[in] x x-coordinate |
| * @param[in] y y-coordinate |
| * @param[in] z z-coordinate |
| * @param[in] batch batch |
| */ |
| virtual void op_load_indirect(const TileOperand &tile_op, |
| const TensorOperand &tensor_op, |
| TensorSampler &sampler, |
| const TileOperand &x, |
| const TileOperand &y, |
| const TileOperand &z, |
| const TileOperand &batch_op) = 0; |
| |
| // ============================================================================================= |
| // ID space management |
| // ============================================================================================= |
| |
| /** Create the new unique ID space and return the value. |
| * |
| * This function changes the ID space to a new number which hasn't been used since the creation |
| * of this kernel writer object. |
| * |
| * @return The new ID space value. |
| */ |
| int32_t new_id_space(); |
| |
| /** Get the current ID space. */ |
| int32_t id_space() const; |
| |
| protected: |
| /** Set the current ID space. |
| * |
| * @param[in] value The ID space to be used. |
| */ |
| KernelWriter &id_space(int32_t value); |
| |
| /** Write the body code using the specified function. |
| * |
| * This function makes sure that a new ID space is created before and then is used solely |
| * by the specified body writing function. |
| * The ID space will not be reused after that. |
| * |
| * @param[in] body The function that writes the body code. |
| */ |
| void write_body(const std::function<void()> &body); |
| |
| protected: |
| /** Generate full variable name by prefixing it with id space */ |
| std::string generate_full_name(const std::string &name) const; |
| |
| /** Create a new tile operand referring to the specified tile object. */ |
| static TileOperand create_tile_operand(ITile &tile); |
| |
| /** Get the reference to the tile object and the active area from the tile operand. */ |
| static std::tuple<ITile &, TileArea> get_tile(const TileOperand &operand); |
| |
| /** Create a new tensor operand from a tensor object. */ |
| static TensorOperand create_tensor_operand(ITensor &tensor); |
| |
| /** Get the reference to tensor object from the tensor operand. */ |
| static ITensor &get_tensor(const TensorOperand &operand); |
| |
| /** Get the values of a constant data object. */ |
| static const std::vector<std::vector<std::string>> &get_values(const ConstantData &data); |
| |
| /** Get the data type of a constant data object. */ |
| static DataType get_data_type(const ConstantData &data); |
| |
| private: |
| int32_t _id_space{0}; |
| int32_t _last_created_id_space{0}; |
| }; |
| |
| } // namespace ckw |
| |
| #endif // CKW_INCLUDE_CKW_KERNELWRITER_H |