diff options
Diffstat (limited to 'compute_kernel_writer/include/ckw')
20 files changed, 1880 insertions, 0 deletions
diff --git a/compute_kernel_writer/include/ckw/Error.h b/compute_kernel_writer/include/ckw/Error.h new file mode 100644 index 0000000000..6b80778957 --- /dev/null +++ b/compute_kernel_writer/include/ckw/Error.h @@ -0,0 +1,131 @@ +/* + * 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_ERROR_H +#define CKW_INCLUDE_CKW_ERROR_H + +#include <stdexcept> +#include <string> + +namespace ckw +{ +/** Creates the error message + * + * @param[in] file File in which the error occurred. + * @param[in] func Function in which the error occurred. + * @param[in] line Line in which the error occurred. + * @param[in] msg Message to display before abandoning. + * + * @return status containing the error + */ +std::string +create_error_msg(const std::string &file, const std::string &func, const std::string &line, const std::string &msg); + +/** Print the given message then throw an std::runtime_error. + * + * @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); \ + std::runtime_error(create_error_msg(arg0, arg1, arg2, arg3)); \ + } while (false) + +/** 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 <typename... T> +inline void ignore_unused(T &&...) +{ +} + +/** Throw an std::runtime_error with the specified message. + * + * @param[in] msg The error message. + */ +#define CKW_THROW_MSG(msg) \ + do \ + { \ + const std::string file(__FILE__); \ + const std::string func(__func__); \ + const std::string line(std::to_string(__LINE__)); \ + const std::string message(msg); \ + \ + throw std::runtime_error(ckw::create_error_msg(file, func, line, message)); \ + } while (false) + +#ifdef COMPUTE_KERNEL_WRITER_ASSERTS_ENABLED + +/** If the condition is not met, throw an std::runtime_error with the specified message if assertion is enabled. + * + * @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)) \ + { \ + CKW_THROW_MSG(msg); \ + } \ + } while (false) + +#else // COMPUTE_KERNEL_WRITER_ASSERTS_ENABLED + +#define CKW_ASSERT_MSG(cond, msg) + +#endif // COMPUTE_KERNEL_WRITER_ASSERTS_ENABLED + +/** If the condition is not met, throw an std::runtime_error if assertion is enabled. + * + * @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 condition is not met, throw an std::runtime_error if assertion is enabled. + * + * @param[in] precond The precondition that triggers the check. + * @param[in] cond The condition that is expected to be true if precondition is true. + */ +#define CKW_ASSERT_IF(precond, cond) CKW_ASSERT(!(precond) || (cond)) + +/** Throw an std::runtime_error with the specified message if assertion is enabled. + * + * @param[in] msg The error message when the condition is not met. + */ +#define CKW_ASSERT_FAILED_MSG(msg) CKW_ASSERT_MSG(false, msg) + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_ERROR_H diff --git a/compute_kernel_writer/include/ckw/Kernel.h b/compute_kernel_writer/include/ckw/Kernel.h new file mode 100644 index 0000000000..f9b7bbb82e --- /dev/null +++ b/compute_kernel_writer/include/ckw/Kernel.h @@ -0,0 +1,76 @@ +/* + * 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/KernelArgument.h" + +#include <string> +#include <vector> + +namespace ckw +{ + +// Forward Declerations +class TileInfo; +class TileOperand; + +enum class TargetLanguage; + +/** The kernel that has been emitted by the kernel writer. + * + * It contains all the necessary information to compile and execute the kernel. + */ +class Kernel +{ +public: + virtual ~Kernel(); + + /** Initialize a new instance of @ref Kernel class with all emitted kernel information. + * + * @param[in] language The target language of the kernel. + * @param[in] arguments The list of kernel arguments. + * @param[in] source_code The source code of the kernel. + */ + Kernel(TargetLanguage language, const std::vector<KernelArgument> &arguments, const std::string &source_code); + + /** Get the target language. */ + TargetLanguage target_language() const; + + /** Get the list of arguments. */ + const std::vector<KernelArgument> &arguments() const; + + /** Get the source code. */ + const std::string &source_code() const; + +private: + TargetLanguage _language; + std::vector<KernelArgument> _arguments; + std::string _source_code; +}; + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_KERNEL_H diff --git a/compute_kernel_writer/include/ckw/KernelArgument.h b/compute_kernel_writer/include/ckw/KernelArgument.h new file mode 100644 index 0000000000..7e9bcbf1ee --- /dev/null +++ b/compute_kernel_writer/include/ckw/KernelArgument.h @@ -0,0 +1,99 @@ +/* + * 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_KERNELARGUMENT_H +#define CKW_INCLUDE_CKW_KERNELARGUMENT_H + +#include "ckw/types/TensorComponentType.h" +#include "ckw/types/TensorStorageType.h" + +#include <cstdint> + +namespace ckw +{ + +/** A kernel argument which can be either a tensor storage or a tensor component. */ +class KernelArgument +{ +public: + /** The type of kernel argument. */ + enum class Type : int32_t + { + /** The argument that provides the read and/or write access to the tensor data. + * + * See @ref ckw::TensorStorageType to see the list of supported storage type. + */ + TensorStorage, + + /** The argument that provides extra information about the tensor. + * + * See @ref ckw::TensorComponentType to see the list of supported component. + */ + TensorComponent, + }; + + /** Initialize a new instance of kernel argument class for a tensor storage argument. */ + KernelArgument(int32_t tensor_id, TensorStorageType storage_type); + + /** Initialize a new instance of kernel argument class for a tensor component argument. */ + KernelArgument(int32_t tensor_id, TensorComponentType component_type); + + /** Get the type of kernel argument. */ + Type type() const; + + /** Get the argument ID. + * + * This method can be used to get the tensor info ID of both tensor storage and tensor component arguments. + */ + int32_t id() const; + + /** Get the type of tensor storage. + * + * This method can only be used for tensor storage argument. + */ + TensorStorageType tensor_storage_type() const; + + /** Get the tensor component type. + * + * This method can only be used for tensor component argument. + */ + TensorComponentType tensor_component_type() const; + +private: + Type _type; + int32_t _id; + + union SubId + { + int32_t unknown; + TensorStorageType tensor_storage_type; + TensorComponentType tensor_component_type; + }; + + SubId _sub_id{0}; +}; + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_KERNELARGUMENT_H diff --git a/compute_kernel_writer/include/ckw/KernelWriter.h b/compute_kernel_writer/include/ckw/KernelWriter.h new file mode 100644 index 0000000000..da41b940d7 --- /dev/null +++ b/compute_kernel_writer/include/ckw/KernelWriter.h @@ -0,0 +1,418 @@ +/* + * 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 diff --git a/compute_kernel_writer/include/ckw/TensorInfo.h b/compute_kernel_writer/include/ckw/TensorInfo.h new file mode 100644 index 0000000000..5c87cb5b12 --- /dev/null +++ b/compute_kernel_writer/include/ckw/TensorInfo.h @@ -0,0 +1,96 @@ +/* + * 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 COMPUTE_KERNEL_WRITER_INCLUDE_CKW_TENSORINFO_H +#define COMPUTE_KERNEL_WRITER_INCLUDE_CKW_TENSORINFO_H + +#include "ckw/types/DataType.h" +#include "ckw/types/TensorDataLayout.h" + +#include <array> +#include <cstdint> + +namespace ckw +{ + +/** Compute Kernel Writer tensor shape + * The value -1 for the tensor dimension is reserved to dynamic dimensions. + */ +using TensorShape = std::array<int32_t, 5>; + +/** Tensor dimension value reserved to dynamic dimensions */ +constexpr int32_t kDynamicTensorDimensionValue = -1; + +/** Compute Kernel Writer tensor info */ +class TensorInfo +{ +public: + /** Default constructor */ + TensorInfo() = default; + /** Constructor + * + * @param[in] dt Tensor data type + * @param[in] shape Tensor shape + * @param[in] dl Tensor data layout + * @param[in] id Tensor id. The id is used to keep track of the user tensor binded. Through the id, + * the user can know what tensor has been used by the Compute Kernel Writer. + * Possible id values: + * - greater than or equal to 0: bind a user specific tensors + * - less than 0: bind a virtual tensor (tile) + */ + TensorInfo(DataType dt, const TensorShape &shape, TensorDataLayout dl, int32_t id); + + /** Set shape */ + TensorInfo &shape(const TensorShape &shape); + + /** Get shape */ + TensorShape shape() const; + + /** Set data type */ + TensorInfo &data_type(DataType dt); + + /** Get data type */ + DataType data_type() const; + + /** Set data layout */ + TensorInfo &data_layout(TensorDataLayout dl); + + /** Get data layout */ + TensorDataLayout data_layout() const; + + /** Set id */ + TensorInfo &id(int32_t id); + + /** Get layout */ + int32_t id() const; + +private: + TensorShape _shape{{0}}; + DataType _dt{DataType::Unknown}; + TensorDataLayout _dl{TensorDataLayout::Unknown}; + int32_t _id{-1}; +}; +} // namespace ckw + +#endif /* COMPUTE_KERNEL_WRITER_INCLUDE_CKW_TENSORINFO_H */ diff --git a/compute_kernel_writer/include/ckw/TensorOperand.h b/compute_kernel_writer/include/ckw/TensorOperand.h new file mode 100644 index 0000000000..a3e53d1314 --- /dev/null +++ b/compute_kernel_writer/include/ckw/TensorOperand.h @@ -0,0 +1,109 @@ +/* + * 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/TileOperand.h" + +namespace ckw +{ + +class ITensor; +class TensorInfo; + +/** A tensor operand provides access to the tensor info, tensor storages for load/store operations + * and tensor components (e.g. shape, strides, etc.) in the form of @ref TileOperand objects. + */ +class TensorOperand +{ +public: + // _tensor field is completely hidden from the public API to avoid any misuse. + // Only kernel writer class interacts with tensor operand hence we allow it to access this field. + friend class KernelWriter; + + /** Create an empty tensor operand. + * + * The new tensor operand doesn't refer to any tensor therefore it is not useable. + */ + TensorOperand(); + + /** Check if the tensor operand contains a tensor and therefore useable. */ + bool is_valid() const; + + /** Get the tensor info. */ + const TensorInfo &info() const; + + /** Get the operand that contains the stride in dimension 0 of the tensor. */ + TileOperand stride0(); + + /** Get the operand that contains the stride in dimension 1 of the tensor. */ + TileOperand stride1(); + + /** Get the operand that contains the stride in dimension 2 of the tensor. */ + TileOperand stride2(); + + /** Get the operand that contains the stride in dimension 3 of the tensor. */ + TileOperand stride3(); + + /** Get the operand that contains the stride in dimension 4 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 size of dimensions 2 and 3 collapsed. */ + TileOperand dim2_dim3(); + + /** Get the operand that contains the offset in bytes to the first element. */ + TileOperand offset_first_element_in_bytes(); + +private: + /** Initialize a new instance of @ref TensorOperand class for a tensor. */ + TensorOperand(ITensor &tensor); + + ITensor *_tensor; +}; + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_TENSOROPERAND_H diff --git a/compute_kernel_writer/include/ckw/TensorSampler.h b/compute_kernel_writer/include/ckw/TensorSampler.h new file mode 100644 index 0000000000..117e8de2cf --- /dev/null +++ b/compute_kernel_writer/include/ckw/TensorSampler.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 CKW_INCLUDE_CKW_TENSORSAMPLER_H +#define CKW_INCLUDE_CKW_TENSORSAMPLER_H + +#include "ckw/types/TensorSamplerTypes.h" +#include "ckw/types/TensorStorageType.h" + +namespace ckw +{ + +/** Tensor sampler + * + * It contains information about how the tensor is sampled. It can be used to + * tell how a tile should be stored to tensor memory, and how a tensor should be + * sampled to get the values stored in a tile. Where to sample the tensor is + * defined with the coordinates respecting the addressing modes, storage type and + * the tensor format defined in this class. + */ +class TensorSampler +{ +public: + /** Initialize a new instance of @ref TensorSampler class. */ + TensorSampler(); + + /** Initialize a new instance of @ref TensorSampler class. + * + * @param[in] storage Tensor storage to load/store the tensor from/to + * @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. + */ + TensorSampler(TensorStorageType storage, + TensorSamplerFormat format, + TensorSamplerAddressModeX address_mode_x, + TensorSamplerAddressModeY address_mode_y, + TensorSamplerAddressModeZ address_mode_z); + + /** Get the storage for the tensor */ + TensorStorageType storage() const; + + /** Set the storage for the tensor */ + TensorSampler &storage(TensorStorageType storage); + + /** Get the format of the tensor. */ + TensorSamplerFormat format() const; + + /** Set the format of the tensor. */ + TensorSampler &format(TensorSamplerFormat format); + + /** Get the address mode of the x dimension. */ + TensorSamplerAddressModeX address_mode_x() const; + + /** Set the address mode of the x dimension. */ + TensorSampler &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. */ + TensorSampler &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. */ + TensorSampler &address_mode_z(TensorSamplerAddressModeZ address_mode_z); + +private: + TensorStorageType _storage{TensorStorageType::BufferUint8Ptr}; + 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_PROTOTYPE_INCLUDE_CKW_TENSORSAMPLER_H diff --git a/compute_kernel_writer/include/ckw/TileInfo.h b/compute_kernel_writer/include/ckw/TileInfo.h new file mode 100644 index 0000000000..678bb7aaf6 --- /dev/null +++ b/compute_kernel_writer/include/ckw/TileInfo.h @@ -0,0 +1,92 @@ +/* + * 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 COMPUTE_KERNEL_WRITER_INCLUDE_CKW_TILEINFO +#define COMPUTE_KERNEL_WRITER_INCLUDE_CKW_TILEINFO + +#include "ckw/types/DataType.h" + +#include <array> +#include <cstdint> + +namespace ckw +{ +// Constants to access the tile width and height in the TileShape +constexpr int32_t kTileWidthIdx = 0; +constexpr int32_t kTileHeightIdx = 1; + +/** Compute Kernel Writer tile shape. It is used to define the shape of the tile */ +using TileShape = std::array<int32_t, 2>; + +/** Compute Kernel Writer tile info */ +class TileInfo +{ +public: + /** Constructor used to initialize a scalar variable with a given data type + * + * @param[in] dt Tile data type + */ + TileInfo(DataType dt); + + /** Constructor used to initialize a vector with a given data type and vector length. + * + * @param[in] dt Tile data type + * @param[in] w Tile width (or vector length) + */ + TileInfo(DataType dt, int32_t w); + + /** Constructor used to initialize a tile with a given data type and tile sizes. + * + * @param[in] dt Tile data type + * @param[in] h Tile height + * @param[in] w Tile width + */ + TileInfo(DataType dt, int32_t h, int32_t w); + + /** Set width */ + TileInfo &width(int32_t w); + + /** Get width */ + int32_t width() const; + + /** Set height */ + TileInfo &height(int32_t h); + + /** Get height */ + int32_t height() const; + + /** Set data type */ + TileInfo &data_type(DataType dt); + + /** Get data type */ + DataType data_type() const; + +private: + DataType _dt{DataType::Unknown}; + TileShape _shape{}; +}; + +} // namespace ckw + +#endif /* COMPUTE_KERNEL_WRITER_INCLUDE_CKW_TILEINFO */ diff --git a/compute_kernel_writer/include/ckw/TileOperand.h b/compute_kernel_writer/include/ckw/TileOperand.h new file mode 100644 index 0000000000..556d589bc0 --- /dev/null +++ b/compute_kernel_writer/include/ckw/TileOperand.h @@ -0,0 +1,112 @@ +/* + * 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 <cstdint> + +namespace ckw +{ + +class KernelWriter; +class TensorOperand; +class ITile; +class TileInfo; + +/** A tile operand refers to a tile object that can be used for kernel writing. */ +class TileOperand +{ +public: + // The constructor and _tile field is completely hidden from the public API to avoid any misuse. + // Only kernel writer and tensor operand classes create and interact with tile operand hence we allow them to access this field. + friend class KernelWriter; + friend class TensorOperand; + + /** Create an empty tile operand. + * + * The new tile operand doesn't refer to any tile therefore it is not useable. + */ + TileOperand(); + + /** Check if the tile operand contains a tile and therefore useable. */ + bool is_valid() const; + + /** Get the tile info. */ + const TileInfo &tile_info() const; + + /** Get a row vector of the current tile operand. + * + * @param[in] row The index of the row to be accessed in the current tile operand. + * + * @return A new tile operand referring to a row of the current tile operand. + */ + TileOperand row(int32_t row) const; + + /** Get a scalar element of the current tile operand. + * + * @param[in] row The index of the row to be accessed in the current tile operand. + * @param[in] col The index of the column to be accessed in the current tile operand. + * + * @return A new tile operand referring to a scalar element of the current tile operand. + */ + TileOperand scalar(int32_t row, int32_t col) const; + +private: + // These are hidden from the public API to avoid any misuse. + + /** Initialize a new instance of @ref TileOperand class for the given tile. */ + TileOperand(ITile &tile); + + /** Initialize a new instance of @ref TileOperand class that is the sub-tile of the given tile. */ + TileOperand(const TileOperand &operand, int32_t row_start, int32_t row_end, int32_t col_start, int32_t col_end); + + /** Get a sub-tile of the current tile operand. + * + * The range of rows and columns is defined by pairs of start and end indices, inclusive lower and exclusive upper. + * In other words, any row and column indices satisfying the following conditions will be part of the sub-tile: + * + * row_start <= row_index < row_end + * col_start <= col_index < col_end + * + * @param[in] row_start The start index of the row range. + * @param[in] row_end The end index of the row range. + * @param[in] col_start The start index of the column range. + * @param[in] col_end The end index of the column range. + * + * @return A new tile operand refering to the same tile but with the new active area. + */ + TileOperand tile(int32_t row_start, int32_t row_end, int32_t col_start, int32_t col_end) const; + + ITile *_tile; + + int32_t _row_start; + int32_t _row_end; + int32_t _col_start; + int32_t _col_end; +}; + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_TILEOPERAND_H diff --git a/compute_kernel_writer/include/ckw/types/ConstantData.h b/compute_kernel_writer/include/ckw/types/ConstantData.h new file mode 100644 index 0000000000..ea95049c9e --- /dev/null +++ b/compute_kernel_writer/include/ckw/types/ConstantData.h @@ -0,0 +1,93 @@ +/* + * 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_TYPES_CONSTANTDATA_H +#define CKW_INCLUDE_CKW_TYPES_CONSTANTDATA_H + +#include "ckw/Error.h" +#include "ckw/types/DataType.h" + +#include <algorithm> +#include <cstdint> +#include <initializer_list> +#include <iomanip> +#include <iterator> +#include <sstream> +#include <string> +#include <type_traits> +#include <vector> + +namespace ckw +{ +// Forward Declarations +class KernelWriter; + +class ConstantData +{ + using String = std::string; + using StringVector = std::vector<String>; + +public: + /** Templated constructor */ + template <typename T> + ConstantData(std::initializer_list<std::initializer_list<T>> values, DataType data_type); + + /** Templated constructor */ + template <typename T> + ConstantData(const std::vector<std::vector<T>> &values, DataType data_type); + +private: + /** Validate the given data type and the template type + * + * @param[in] data_type data type + * + * @return true if user provided data type and the template type are conformant + */ + template <typename T> + bool validate(DataType data_type); + + /** Get the constant data as a 2d vector of string values + * + * @return a 2d vector of strings that has the string-converted values + */ + const std::vector<StringVector> &values() const; + + /** Get the underlying data type of the constant values + * + * @return a @ref ckw::DataType object that represents the underlying data type + */ + DataType data_type() const; + + // Friends + friend class KernelWriter; + +private: + // Data members + std::vector<StringVector> _values{}; + DataType _data_type{}; +}; + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_TYPES_CONSTANTDATA_H diff --git a/compute_kernel_writer/include/ckw/types/ConvertPolicy.h b/compute_kernel_writer/include/ckw/types/ConvertPolicy.h new file mode 100644 index 0000000000..43a37ff118 --- /dev/null +++ b/compute_kernel_writer/include/ckw/types/ConvertPolicy.h @@ -0,0 +1,41 @@ +/* + * 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_TYPES_CONVERTPOLICY_H +#define CKW_INCLUDE_CKW_TYPES_CONVERTPOLICY_H + +#include <cstdint> + +namespace ckw +{ + +enum class ConvertPolicy : int32_t +{ + None = 0, // No policy specified. + Saturate = 1, // Saturated. +}; + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_TYPES_CONVERTPOLICY_H diff --git a/compute_kernel_writer/include/ckw/types/DataType.h b/compute_kernel_writer/include/ckw/types/DataType.h new file mode 100644 index 0000000000..3447dd61d6 --- /dev/null +++ b/compute_kernel_writer/include/ckw/types/DataType.h @@ -0,0 +1,50 @@ +/* +* 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_DATATYPE_H +#define CKW_INCLUDE_CKW_DATATYPE_H + +#include <cstdint> + +namespace ckw +{ + +/** Compute Kernel Writer data types. This data type is used by the code variables and tensor arguments. */ +enum class DataType : int32_t +{ + Unknown = 0x00, + Fp32 = 0x11, + Fp16 = 0x12, + Int32 = 0x21, + Int16 = 0x22, + Int8 = 0x24, + Uint32 = 0x31, + Uint16 = 0x32, + Uint8 = 0x34, + Bool = 0x41 +}; + +} // namespace ckw + +#endif //CKW_INCLUDE_CKW_DATATYPE_H diff --git a/compute_kernel_writer/include/ckw/types/MemoryOperation.h b/compute_kernel_writer/include/ckw/types/MemoryOperation.h new file mode 100644 index 0000000000..f93f60c51a --- /dev/null +++ b/compute_kernel_writer/include/ckw/types/MemoryOperation.h @@ -0,0 +1,37 @@ +/* + * 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_TYPES_MEMORYOPERATION +#define CKW_INCLUDE_CKW_TYPES_MEMORYOPERATION + +namespace ckw +{ +enum class MemoryOperation +{ + Load = 1, + Store = 2 +}; +} // namespace ckw + +#endif /* CKW_INCLUDE_CKW_TYPES_MEMORYOPERATION */ diff --git a/compute_kernel_writer/include/ckw/types/Operators.h b/compute_kernel_writer/include/ckw/types/Operators.h new file mode 100644 index 0000000000..77b0519422 --- /dev/null +++ b/compute_kernel_writer/include/ckw/types/Operators.h @@ -0,0 +1,101 @@ +/* +* 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_TYPES_OPERATORS_H +#define CKW_INCLUDE_CKW_TYPES_OPERATORS_H + +#include <cstdint> + +namespace ckw +{ + +/** Unary operators and functions. */ +enum class UnaryOp : int32_t +{ + LogicalNot = 0x0000, // ! + BitwiseNot = 0x0001, // ~ + + Exp = 0x0010, + Tanh = 0x0011, + Sqrt = 0x0012, + Erf = 0x0013, + Fabs = 0x0014, + Log = 0x0015, + Round = 0x0016, + Floor = 0x0017, +}; + +/** Assignment operators. */ +enum class AssignmentOp : int32_t +{ + Increment = 0x0000, // += + Decrement = 0x0001, // -= +}; + +/** Binary operators. */ +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, // || + + // Bitwise + BitwiseXOR = 0x4000, // ^ + + // Functions + Min = 0x8000, + Max = 0x8001, +}; + +/** Ternary operators. */ +enum class TernaryOp : int32_t +{ + Select = 0x0000, + Clamp = 0x0001, +}; + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_TYPES_OPERATORS_H diff --git a/compute_kernel_writer/include/ckw/types/TargetArchitecture.h b/compute_kernel_writer/include/ckw/types/TargetArchitecture.h new file mode 100644 index 0000000000..25662a01f0 --- /dev/null +++ b/compute_kernel_writer/include/ckw/types/TargetArchitecture.h @@ -0,0 +1,40 @@ +/* +* 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_TYPES_TARGETARCHITECTURE_H +#define CKW_INCLUDE_CKW_TYPES_TARGETARCHITECTURE_H + +namespace ckw +{ + +/** Target platform architecture. */ +enum class TargetArchitecture +{ + Unknown, + GpuArmMaliValhall, +}; + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_TYPES_TARGETARCHITECTURE_H diff --git a/compute_kernel_writer/include/ckw/types/TargetLanguage.h b/compute_kernel_writer/include/ckw/types/TargetLanguage.h new file mode 100644 index 0000000000..1f507573dd --- /dev/null +++ b/compute_kernel_writer/include/ckw/types/TargetLanguage.h @@ -0,0 +1,40 @@ +/* +* 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_TYPES_TARGETLANGUAGE_H +#define CKW_INCLUDE_CKW_TYPES_TARGETLANGUAGE_H + +namespace ckw +{ + +/** Target language. */ +enum class TargetLanguage +{ + Unknown, + OpenCL +}; + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_TYPES_TARGETLANGUAGE_H diff --git a/compute_kernel_writer/include/ckw/types/TensorComponentType.h b/compute_kernel_writer/include/ckw/types/TensorComponentType.h new file mode 100644 index 0000000000..7a5031d8c0 --- /dev/null +++ b/compute_kernel_writer/include/ckw/types/TensorComponentType.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 CKW_INCLUDE_CKW_TYPES_TENSORCOMPONENTTYPE_H +#define CKW_INCLUDE_CKW_TYPES_TENSORCOMPONENTTYPE_H + +#include <cstdint> + +namespace ckw +{ + +/** Compute Kernel Writer tensor component. + * + * The tensor components are used to access specific backend-agnostic tensor arguments, + * such as the tensor dimensions and tensor strides. + * The tensor component is represented as an unsigned integer. The value of the integer value + * is assigned to retrieve the information through the @ref TensorComponentBitmask. + */ +enum class TensorComponentType : uint32_t +{ + Unknown = 0x00000000, + OffsetFirstElement = 0x01000000, + Stride0 = 0x02000001, + Stride1 = 0x02000002, + Stride2 = 0x02000003, + Stride3 = 0x02000004, + Stride4 = 0x02000005, + Dim0 = 0x04000001, + Dim1 = 0x04000002, + Dim2 = 0x04000003, + Dim3 = 0x04000004, + Dim4 = 0x04000005, + Dim1xDim2 = 0x08000032, + Dim2xDim3 = 0x08000043, + Dim1xDim2xDim3 = 0x08000432 +}; + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_TYPES_TENSORCOMPONENTTYPE_H diff --git a/compute_kernel_writer/include/ckw/types/TensorDataLayout.h b/compute_kernel_writer/include/ckw/types/TensorDataLayout.h new file mode 100644 index 0000000000..532b299910 --- /dev/null +++ b/compute_kernel_writer/include/ckw/types/TensorDataLayout.h @@ -0,0 +1,52 @@ +/* + * 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_TYPES_TENSORDATALAYOUT_H +#define CKW_INCLUDE_CKW_TYPES_TENSORDATALAYOUT_H + +namespace ckw +{ + +/** Compute Kernel Writer tensor data layout (or memory format) */ +enum class TensorDataLayout +{ + Unknown, + Nhwc, + Ndhwc +}; + +/** Compute Kernel Writer tensor data layout component */ +enum class TensorDataLayoutComponent +{ + Unknown, + N, + D, + H, + W, + C, +}; + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_TYPES_TENSORDATALAYOUT_H diff --git a/compute_kernel_writer/include/ckw/types/TensorSamplerTypes.h b/compute_kernel_writer/include/ckw/types/TensorSamplerTypes.h new file mode 100644 index 0000000000..512d0b4501 --- /dev/null +++ b/compute_kernel_writer/include/ckw/types/TensorSamplerTypes.h @@ -0,0 +1,84 @@ +/* + * 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_TYPES_TENSORSAMPLERTYPES_H +#define CKW_INCLUDE_CKW_TYPES_TENSORSAMPLERTYPES_H + +#include <cstdint> + +namespace ckw +{ + +// This enum class defines how the dimensions of a 3d tensor is mapped into x,y and z coordianates. +enum class TensorSamplerFormat : int32_t +{ + Unknown = 0, + Dim0_Dim1xDim2_1 = 1, // Original dimensions 1 and 2 are collapsed onto y-axis + Dim0_Dim1_Dim2 = 2 // Original dimensions stays as they're defined. No collapsing. +}; + +/** Tensor sampler address mode enum class for X dimension + * + * The following address modes are available in total: + * Unknown + * None : The user guarantees that the coordinate is always in-bound + * OverlappingMin : (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. + * ClampToBorderMaxOnly : Clamp to max value allowed in the corresponding dimension, and construct an if/else guard to prevent out of bound access, + * e.g. if( y < size-of-dimension-y ){ <do the operation> } + * SkipLessThanZero : Skip loading/storing if the index is less than 0 + * + * Individual dimensions choose which adddress mode to implement in their respective enum classes. + */ +enum class TensorSamplerAddressModeX : int32_t +{ + Unknown = 0, + None = 1, + OverlappingMin = 2 +}; + +/** + * Similar to @ref TensorSamplerAddressModeX + */ +enum class TensorSamplerAddressModeY : int32_t +{ + Unknown = 0, + None = 1, + OverlappingMin = 2, + ClampToBorderMaxOnly = 3, + SkipLessThanZero = 4 +}; + +/** + * Similar to @ref TensorSamplerAddressModeX + */ +enum class TensorSamplerAddressModeZ : int32_t +{ + Unknown = 0, + None = 1, +}; + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_TYPES_TENSORSAMPLERTYPES_H diff --git a/compute_kernel_writer/include/ckw/types/TensorStorageType.h b/compute_kernel_writer/include/ckw/types/TensorStorageType.h new file mode 100644 index 0000000000..5a2f17d520 --- /dev/null +++ b/compute_kernel_writer/include/ckw/types/TensorStorageType.h @@ -0,0 +1,46 @@ +/* + * 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_TYPES_TENSORSTORAGETYPE_H +#define CKW_INCLUDE_CKW_TYPES_TENSORSTORAGETYPE_H + +#include <cstdint> + +namespace ckw +{ + +/** Compute Kernel Writer tensor storage. + * The tensor storage represents the type of tensor memory object. + */ +enum class TensorStorageType : uint32_t +{ + Unknown = 0x00000000, + BufferUint8Ptr = 0x01000000, + Texture2dReadOnly = 0x02000001, + Texture2dWriteOnly = 0x02000010, +}; + +} // namespace ckw + +#endif // CKW_INCLUDE_CKW_TYPES_TENSORSTORAGETYPE_H |