diff options
Diffstat (limited to 'compute_kernel_writer/prototype/include/ckw')
18 files changed, 0 insertions, 3217 deletions
diff --git a/compute_kernel_writer/prototype/include/ckw/Error.h b/compute_kernel_writer/prototype/include/ckw/Error.h deleted file mode 100644 index aab713c817..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/Error.h +++ /dev/null @@ -1,78 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_PROTOTYPE_INCLUDE_CKW_ERROR_H -#define CKW_PROTOTYPE_INCLUDE_CKW_ERROR_H - -#include <stdexcept> -#include <string> - -namespace ckw -{ - -/** If the condition is not met, throw an std::runtime_error with the specified message. - * - * @param[in] cond The condition that is expected to be true. - * @param[in] msg The error message when the condition is not met. - */ -#define CKW_ASSERT_MSG(cond, msg) \ - do \ - { \ - if (!(cond)) \ - { \ - throw ::std::runtime_error(msg); \ - } \ - } while (false) - -/** If the condition is not met, throw an std::runtime_error. - * - * @param[in] cond The condition that is expected to be true. - */ -#define CKW_ASSERT(cond) CKW_ASSERT_MSG(cond, #cond) - -/** If the precondition is met but the consequence is not met, throw an std::runtime_error. - * - * @param[in] precond The condition if is met requires the consequence must also be met. - * @param[in] cond The condition that is expected to be true if the precondition is true. - */ -#define CKW_ASSERT_IF(precond, cond) CKW_ASSERT_MSG(!(precond) || ((precond) && (cond)), #precond " |-> " #cond) - -/** Mark the variables as unused. - * - * @param[in] ... Variables which are unused. - */ -#define CKW_UNUSED(...) ::ckw::ignore_unused(__VA_ARGS__) // NOLINT - -/** Mark the variables as unused. - * - * @param[in] ... Variables which are unused. - */ -template <typename... T> -inline void ignore_unused(T &&...) -{ -} - -} // namespace ckw - -#endif // CKW_INCLUDE_CKW_ERROR_H diff --git a/compute_kernel_writer/prototype/include/ckw/Kernel.h b/compute_kernel_writer/prototype/include/ckw/Kernel.h deleted file mode 100644 index ba31a29ba7..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/Kernel.h +++ /dev/null @@ -1,102 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_PROTOTYPE_INCLUDE_CKW_KERNEL_H -#define CKW_PROTOTYPE_INCLUDE_CKW_KERNEL_H - -#include "ckw/KernelArgument.h" -#include "ckw/OperandBase.h" -#include "ckw/types/GpuTargetLanguage.h" - -#include <map> -#include <memory> -#include <string> -#include <vector> - -namespace ckw -{ - -class TileOperand; - -namespace prototype -{ -class GpuKernelWriterDataHolder; -} // namespace prototype - -/** The target for kernel writer to write into. */ -class Kernel -{ -public: - /** Constructor - * - * @param[in] language The programming language to write the kernel. - */ - Kernel(GpuTargetLanguage language); - /** Constructor - * - * @param[in] name The name of the kernel function. - * @param[in] language The programming language to write the kernel. - */ - Kernel(const char *name, GpuTargetLanguage language); - - /** Destructor */ - ~Kernel(); - - /** Get the name of the kernel function. */ - const std::string &name() const; - - /** Set the name of the kernel function. - * - * @param[in] name The name of the kernel function. - */ - void name(const std::string &name); - - /** Get the list of kernel arguments. */ - ::std::vector<KernelArgument> arguments() const; - - /** (Internal use only) Register the tile operand. - * - * @param operand The tile operand to be registered. - */ - TileOperand ®ister_operand(::std::unique_ptr<TileOperand> operand); - - /** (Internal use only) Register the tensor operand. - * - * @param operand The tensor operand to be registered. - */ - TensorOperand ®ister_operand(::std::unique_ptr<TensorOperand> operand); - - /** (Internal use only) Get the implementation data. */ - prototype::GpuKernelWriterDataHolder *impl(); - -private: - ::std::string _name; - ::std::unique_ptr<prototype::GpuKernelWriterDataHolder> _kernel; - ::std::map<::std::string, ::std::unique_ptr<OperandBase>> _operands; - ::std::map<int32_t, TensorOperand *> _tensor_id_operands; -}; - -} // namespace ckw - -#endif // CKW_PROTOTYPE_INCLUDE_CKW_KERNEL_H diff --git a/compute_kernel_writer/prototype/include/ckw/KernelArgument.h b/compute_kernel_writer/prototype/include/ckw/KernelArgument.h deleted file mode 100644 index 3384a20aef..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/KernelArgument.h +++ /dev/null @@ -1,107 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_PROTOTYPE_INCLUDE_CKW_KERNELARGUMENT_H -#define CKW_PROTOTYPE_INCLUDE_CKW_KERNELARGUMENT_H - -#include "ckw/TensorInfo.h" - -#include <cstdint> - -namespace ckw -{ - -class TensorOperand; -class TensorComponentOperand; - -/** 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::TensorStorage to see the list of supported storage type. - */ - TensorStorage, - - /** The argument that provides extra information about the tensor. - * - * See @ref ckw::TensorComponent to see the list of supported component. - */ - TensorComponent, - }; - - /** Initialize a new instance of kernel argument class for a tensor storage argument. - * - * @param[in] tensor The tensor whose storage is exposed to kernel arguments. - */ - KernelArgument(TensorOperand &tensor); - - /** Initialize a new instance of kernel argument class for a tensor component argument. - * - * @param[in] tensor_component The tensor component to be exposed to kernel arguments. - */ - KernelArgument(TensorComponentOperand &tensor_component); - - /** 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_PROTOTYPE_INCLUDE_CKW_KERNELARGUMENT_H diff --git a/compute_kernel_writer/prototype/include/ckw/KernelWriter.h b/compute_kernel_writer/prototype/include/ckw/KernelWriter.h deleted file mode 100644 index f9e0066f91..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/KernelWriter.h +++ /dev/null @@ -1,338 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_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 <memory> - -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 <typename... TArgs> - TileOperand &declare_tile(const std::string &name, TArgs &&...args) - { - const auto var_name = generate_variable_name(name); - auto operand = std::make_unique<TileOperand>(var_name, ::std::forward<TArgs>(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: `<dst> = <src>;`. - * - * @param[out] dst The destination tile. - * @param[in] src The source tile. - */ - void op_assign(const TileOperand &dst, const TileOperand &src); - - /** Write the cast: `<dst> = convert_<dst.type><_sat>(<src>);`. - * - * @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: `<dst> = <op> <src>`. - * - * @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: `<dst> = <lhs> <op> <rhs>;`. - * - * @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: `<dst> = <func>(<src>);`. - * - * @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: `<dst> = <func>(<first>, <second>);`. - * - * @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: `<dst> = <func>(<first>, <second>, <third>);`. - * - * @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(<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 body of the if-statement. - */ - void op_if(const TileOperand &lhs, BinaryOp op, const TileOperand &rhs, const std::function<void()> &body); - - /** Write else-if-statement: `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 body of the else-if-statement. - */ - void op_else_if(const TileOperand &lhs, BinaryOp op, const TileOperand &rhs, const std::function<void()> &body); - - /** Write an else-statement: `else { <body> }`. - * - * @param[in] body The body of the else-statement. - */ - void op_else(const std::function<void()> &body); - - /** Write for-loops: `for(; <var> <cond_op> <cond_value>; <var> <update_op> <update_value>) { 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<void()> &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<TileOperand> operand); - -private: - Kernel *_kernel; - ::std::unique_ptr<prototype::GpuKernelWriterAttribute> _impl_attr; - ::std::unique_ptr<prototype::IGpuKernelWriter> _impl; - - int32_t _id_space{0}; - int32_t _max_id_space{0}; -}; - -} // namespace ckw - -#endif // CKW_PROTOTYPE_INCLUDE_CKW_KERNELWRITER_H diff --git a/compute_kernel_writer/prototype/include/ckw/KernelWriterHelper.h b/compute_kernel_writer/prototype/include/ckw/KernelWriterHelper.h deleted file mode 100644 index 3ba079bbc2..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/KernelWriterHelper.h +++ /dev/null @@ -1,1286 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_INCLUDE_CKW_KERNELWRITERHELPER_H -#define CKW_INCLUDE_CKW_KERNELWRITERHELPER_H - -#include "ckw/KernelWriter.h" -#include "ckw/TensorOperand.h" -#include "ckw/TileOperand.h" - -#include <iostream> -#include <type_traits> - -/* - * By including this header file you will be able to supplement the default - * Compute Kernel Writer API with additional syntax to help ease the use of CKW. - * - * To use the KernelWriterHelper you need to wrap your instance of KernelWriter - * (or any class deriving from KernelWriter): - * KernelWriterHelper<KernelWriter> writer; - * The resulting writer object comprises the original KernelWriter - * functionality (drop-in replacement), but extends the syntax as follows. - * - * Common functions/operators have natural syntax: - * 1. Unary expressions: - * writer.op_assign(dst, !src); // Logical NOT - * writer.op_assign(dst, ~src); // Bitwise NOT - * - * 2. Binary expressions: - * writer.op_assign(dst, lhs + rhs); // Addition - * writer.op_assign(dst, lhs - rhs); // Subtraction - * writer.op_assign(dst, lhs * rhs); // Multiplication - * writer.op_assign(dst, lhs / rhs); // Division - * writer.op_assign(dst, lhs % rhs); // Modulo - * writer.op_assign(dst, lhs == rhs); // Equality - * writer.op_assign(dst, lhs < rhs); // Less-than - * writer.op_assign(dst, lhs <= rhs); // Less-than-or-equal - * writer.op_assign(dst, lhs > rhs); // Greater-than - * writer.op_assign(dst, lhs >= rhs); // Greater-than-or-equal - * writer.op_assign(dst, lhs ^ rhs); // Bitwise XOR - * writer.op_assign(dst, logical_and(lhs, rhs)); // Logical AND - * writer.op_assign(dst, logical_or(lhs, rhs)); // Logical OR - * - * 3. Unary elementwise functions: - * writer.op_assign(dst, exp(src)); // Exponent - * writer.op_assign(dst, tanh(src)); // Hyperbolic tangent - * writer.op_assign(dst, sqrt(src)); // Square root - * writer.op_assign(dst, erf(src)); // Error function - * writer.op_assign(dst, fabs(src)); // Absolute of floating-point number - * writer.op_assign(dst, log(src)); // Natural logarithm - * writer.op_assign(dst, round(src)); // Round - * writer.op_assign(dst, sizeOf(src)); // sizeof - * - * 4. Binary elementwise functions: - * writer.op_assign(dst, max(first, second)); // Max - * writer.op_assign(dst, min(first, second)); // Min - * - * 5. Ternary elementwise functions: - * writer.op_assign(dst, select(first, second, third)); // Select - * - * NOTE: All the above examples support nesting, so you could write - * something like: writer.op_assign(dst, src * (log(arg) + sqrt(abs(arg))); - * - * - * 6. If-statements. The preceding syntax also allows easier writing of if-statements: - * writer.op_if(<cond>, <body>); - * - * For example: - * writer.op_if(exp(first_arg) == dst, [&]{ - * //... - * }).op_else_if(exp(first_arg) > dst, [&]{ - * //... - * }).op_else([&] { - * //... - * }); - * - * 7. For-loops. A similar syntax exists for for-loops: - * writer.op_for_loop(<cond>, <updater>, <body>); - * - * For example: - * writer.op_for_loop(index < limit, index += step, [&]{ - * //... - * }); - * - * NOTE: There are limitations on the for-loop <cond> and <updater> parameters. - * In neither the <cond> (Binary expression) or <updater> (Increment/Decrement) - * is it allowed to use nesting. For example, `(index + other) < limit` and - * `index < round(limit)` are invalid <cond> parameters. This is because the - * semantics of for-loops rely on the condition being evaluated at every iteration, - * but as temporary variables might be defined for nested expressions the semantics - * cannot be guaranteed. - */ - -namespace ckw -{ - -// ================================================== -// Type traits -// ================================================== - -/** Specifies if the type can be used as an operand for functions (e.g. max), operations (e.g. *), or assignments. */ -template <typename T> -struct can_be_operand : ::std::false_type -{ -}; - -/** Specifies if the type can be assigned/written to. */ -template <typename T> -struct can_be_assigned : ::std::false_type -{ -}; - -template <> -struct can_be_operand<TileOperand &> : ::std::true_type -{ -}; - -template <> -struct can_be_assigned<TileOperand &> : ::std::true_type -{ -}; - -// ================================================== -// Assignment -// ================================================== - -/** AST node for assignments. - * - * Note that \p TRight must be an operand, and \p TLeft must be assignable. - * - * @tparam TLeft The type of the destination of the assignment. - * @tparam TRight The type of the source assigned to the destination. - */ -template <typename TLeft, - typename TRight, - typename = ::std::enable_if<can_be_operand<TRight>::value && can_be_assigned<TLeft>::value>> -struct Assignment -{ - TLeft lhs; - TRight rhs; - AssignmentOp opcode; -}; - -/** Represents the expression: `\p lhs += \p rhs`. - * - * @tparam TLeft The type of the LHS of the assignment. - * @tparam TRight The type of the RHS of the assignment. - * @param[in] lhs The LHS of the assignment. - * @param[in] rhs The RHS of the assignment. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline Assignment<TLeft, TRight> operator+=(TLeft &&lhs, TRight &&rhs) -{ - return Assignment<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), AssignmentOp::Increment}; -} - -/** Represents the expression: `\p lhs -= \p rhs`. - * - * @tparam TLeft The type of the LHS of the assignment. - * @tparam TRight The type of the RHS of the assignment. - * @param[in] lhs The LHS of the assignment. - * @param[in] rhs The RHS of the assignment. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline Assignment<TLeft, TRight> operator-=(TLeft &&lhs, TRight &&rhs) -{ - return Assignment<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), AssignmentOp::Decrement}; -} - -// ================================================== -// Unary expression -// ================================================== - -/** AST node for unary expressions. - * - * Note that \p TSrc must be an operand. - * - * @tparam TSrc The type of the argument to the expression. - */ -template <typename TSrc, typename = ::std::enable_if<can_be_operand<TSrc>::value>> -struct UnaryExpression -{ - TSrc src; - UnaryOp opcode; -}; - -template <typename TLeft> -struct can_be_operand<UnaryExpression<TLeft>> : ::std::true_type -{ -}; - -/** Represents the expression: `!\p src`. - * - * @tparam TSrc The type of the argument. - * @param[in] src The argument. - * @return The resulting AST node. - */ -template <typename TSrc> -inline UnaryExpression<TSrc> operator!(TSrc &&src) -{ - return UnaryExpression<TSrc>{std::forward<TSrc>(src), UnaryOp::LogicalNot}; -} - -/** Represents the expression: `~\p src`. - * - * @tparam TSrc The type of the argument. - * @param[in] src The argument. - * @return The resulting AST node. - */ -template <typename TSrc> -inline UnaryExpression<TSrc> operator~(TSrc &&src) -{ - return UnaryExpression<TSrc>{std::forward<TSrc>(src), UnaryOp::BitwiseNot}; -} - -// ================================================== -// Binary expressions -// ================================================== - -/** AST node for binary expressions. - * - * Note that both \p TLeft and \p TRight must be operands. - * - * @tparam TLeft The type of the left argument of the expression. - * @tparam TRight The type of the right argument of the expression. - */ -template <typename TLeft, - typename TRight, - typename = ::std::enable_if_t<can_be_operand<TLeft>::value && can_be_operand<TRight>::value>> -struct BinaryExpression -{ - TLeft lhs; - TRight rhs; - BinaryOp opcode; -}; - -template <typename TLeft, typename TRight> -struct can_be_operand<BinaryExpression<TLeft, TRight>> : ::std::true_type -{ -}; - -/** Represents the expression: `\p lhs + \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline BinaryExpression<TLeft, TRight> operator+(TLeft &&lhs, TRight &&rhs) -{ - return BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::Add}; -} - -/** Represents the expression: `\p lhs - \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline BinaryExpression<TLeft, TRight> operator-(TLeft &&lhs, TRight &&rhs) -{ - return BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::Sub}; -} - -/** Represents the expression: `\p lhs * \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline BinaryExpression<TLeft, TRight> operator*(TLeft &&lhs, TRight &&rhs) -{ - return BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::Mul}; -} - -/** Represents the expression: `\p lhs / \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline BinaryExpression<TLeft, TRight> operator/(TLeft &&lhs, TRight &&rhs) -{ - return BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::Div}; -} - -/** Represents the expression: `\p lhs % \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline BinaryExpression<TLeft, TRight> operator%(TLeft &&lhs, TRight &&rhs) -{ - return BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::Mod}; -} - -/** Represents the expression: `\p lhs == \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline BinaryExpression<TLeft, TRight> operator==(TLeft &&lhs, TRight &&rhs) -{ - return BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::Equal}; -} - -/** Represents the expression: `\p lhs < \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline BinaryExpression<TLeft, TRight> operator<(TLeft &&lhs, TRight &&rhs) -{ - return BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::Less}; -} - -/** Represents the expression: `\p lhs <= \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline BinaryExpression<TLeft, TRight> operator<=(TLeft &&lhs, TRight &&rhs) -{ - return BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::LessEqual}; -} - -/** Represents the expression: `\p lhs > \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline BinaryExpression<TLeft, TRight> operator>(TLeft &&lhs, TRight &&rhs) -{ - return BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::Greater}; -} - -/** Represents the expression: `\p lhs >= \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline BinaryExpression<TLeft, TRight> operator>=(TLeft &&lhs, TRight &&rhs) -{ - return BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::GreaterEqual}; -} - -/** Represents the expression: `\p lhs ^ \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline BinaryExpression<TLeft, TRight> operator^(TLeft &&lhs, TRight &&rhs) -{ - return BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::BitwiseXOR}; -} - -/** Represents the expression: `\p lhs && \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline BinaryExpression<TLeft, TRight> logical_and(TLeft &&lhs, TRight &&rhs) -{ - return BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::LogicalAnd}; -} - -/** Represents the expression: `\p lhs && \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight, typename... TOps> -inline BinaryExpression<BinaryExpression<TLeft, TRight>, TOps...> logical_and(TLeft &&lhs, TRight &&rhs, TOps &&...ops) -{ - return logical_and( - BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::LogicalAnd}, - std::forward<TOps>(ops)...); -} - -/** Represents the expression: `\p lhs || \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight> -inline BinaryExpression<TLeft, TRight> logical_or(TLeft &&lhs, TRight &&rhs) -{ - return BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::LogicalOr}; -} - -/** Represents the expression: `\p lhs || \p rhs`. - * - * @tparam TLeft The type of the LHS of the expression. - * @tparam TRight The type of the RHS of the expression. - * @param[in] lhs The LHS of the expression. - * @param[in] rhs The RHS of the expression. - * @return The resulting AST node. - */ -template <typename TLeft, typename TRight, typename... TOps> -inline BinaryExpression<BinaryExpression<TLeft, TRight>, TOps...> logical_or(TLeft &&lhs, TRight &&rhs, TOps &&...ops) -{ - return logical_or( - BinaryExpression<TLeft, TRight>{std::forward<TLeft>(lhs), std::forward<TRight>(rhs), BinaryOp::LogicalOr}, - std::forward<TOps>(ops)...); -} - -// ================================================== -// Unary elementwise functions -// ================================================== - -/** AST node for unary elementwise functions. - * - * Note that \p TSrc must be an operand. - * - * @tparam TSrc The type of the argument to the function. - */ -template <typename TSrc, typename = ::std::enable_if<can_be_operand<TSrc>::value>> -struct UnaryElementwiseFunction -{ - TSrc src; - UnaryFunction opcode; -}; - -template <typename TLeft> -struct can_be_operand<UnaryElementwiseFunction<TLeft>> : ::std::true_type -{ -}; - -/** Represents the expression: `exp(\p src)`. - * - * @tparam TSrc The type of the argument. - * @param[in] src The argument. - * @return The resulting AST node. - */ -template <typename TSrc> -UnaryElementwiseFunction<TSrc> exp(TSrc &&src) -{ - return UnaryElementwiseFunction<TSrc>{std::forward<TSrc>(src), UnaryFunction::Exp}; -} - -/** Represents the expression: `tanh(\p src)`. - * - * @tparam TSrc The type of the argument. - * @param[in] src The argument. - * @return The resulting AST node. - */ -template <typename TSrc> -UnaryElementwiseFunction<TSrc> tanh(TSrc &&src) -{ - return UnaryElementwiseFunction<TSrc>{std::forward<TSrc>(src), UnaryFunction::Tanh}; -} - -/** Represents the expression: `sqrt(\p src)`. - * - * @tparam TSrc The type of the argument. - * @param[in] src The argument. - * @return The resulting AST node. - */ -template <typename TSrc> -UnaryElementwiseFunction<TSrc> sqrt(TSrc &&src) -{ - return UnaryElementwiseFunction<TSrc>{std::forward<TSrc>(src), UnaryFunction::Sqrt}; -} - -/** Represents the expression: `erf(\p src)`. - * - * @tparam TSrc The type of the argument. - * @param[in] src The argument. - * @return The resulting AST node. - */ -template <typename TSrc> -UnaryElementwiseFunction<TSrc> erf(TSrc &&src) -{ - return UnaryElementwiseFunction<TSrc>{std::forward<TSrc>(src), UnaryFunction::Erf}; -} - -/** Represents the expression: `fabs(\p src)`. - * - * @tparam TSrc The type of the argument. - * @param[in] src The argument. - * @return The resulting AST node. - */ -template <typename TSrc> -UnaryElementwiseFunction<TSrc> fabs(TSrc &&src) -{ - return UnaryElementwiseFunction<TSrc>{std::forward<TSrc>(src), UnaryFunction::Fabs}; -} - -/** Represents the expression: `log(\p src)`. - * - * @tparam TSrc The type of the argument. - * @param[in] src The argument. - * @return The resulting AST node. - */ -template <typename TSrc> -UnaryElementwiseFunction<TSrc> log(TSrc &&src) -{ - return UnaryElementwiseFunction<TSrc>{std::forward<TSrc>(src), UnaryFunction::Log}; -} - -/** Represents the expression: `round(\p src)`. - * - * @tparam TSrc The type of the argument. - * @param[in] src The argument. - * @return The resulting AST node. - */ -template <typename TSrc> -UnaryElementwiseFunction<TSrc> round(TSrc &&src) -{ - return UnaryElementwiseFunction<TSrc>{std::forward<TSrc>(src), UnaryFunction::Round}; -} - -/** Represents the expression: `sizeof(\p src)`. - * - * @tparam TSrc The type of the argument. - * @param[in] src The argument. - * @return The resulting AST node. - */ -template <typename TSrc> -UnaryElementwiseFunction<TSrc> sizeOf(TSrc &&src) -{ - return UnaryElementwiseFunction<TSrc>{std::forward<TSrc>(src), UnaryFunction::SizeOf}; -} - -// ================================================== -// Binary elementwise functions -// ================================================== - -/** AST node for binary elementwise functions. - * - * Note that both \p TFirst and \p TSecond must be operands. - * - * @tparam TFirst The type of the left argument of the function. - * @tparam TSecond The type of the right argument of the function. - */ -template <typename TFirst, - typename TSecond, - typename = ::std::enable_if<can_be_operand<TFirst>::value && can_be_operand<TSecond>::value>> -struct BinaryElementwiseFunction -{ - TFirst first; - TSecond second; - BinaryFunction opcode; -}; - -template <typename TFirst, typename TSecond> -struct can_be_operand<BinaryElementwiseFunction<TFirst, TSecond>> : ::std::true_type -{ -}; - -/** Represents the function call: `max(\p first, \p second)`. - * - * @tparam TFirst The type of the first argument. - * @tparam TSecond The type of the second argument. - * @param[in] first The first argument. - * @param[in] second The second argument. - * @return The resulting AST node. - */ -template <typename TFirst, typename TSecond> -BinaryElementwiseFunction<TFirst, TSecond> max(TFirst &&first, TSecond &&second) -{ - return BinaryElementwiseFunction<TFirst, TSecond>{std::forward<TFirst>(first), std::forward<TSecond>(second), - BinaryFunction::Max}; -} - -/** Represents the function call: `min(\p first, \p second)`. - * - * @tparam TFirst The type of the first argument. - * @tparam TSecond The type of the second argument. - * @param[in] first The first argument. - * @param[in] second The second argument. - * @return The resulting AST node. - */ -template <typename TFirst, typename TSecond> -BinaryElementwiseFunction<TFirst, TSecond> min(TFirst &&first, TSecond &&second) -{ - return BinaryElementwiseFunction<TFirst, TSecond>{std::forward<TFirst>(first), std::forward<TSecond>(second), - BinaryFunction::Min}; -} - -// ================================================== -// Ternary elementwise functions -// ================================================== - -/** AST node for ternary elementwise functions. - * - * Note that \p TFirst, \p TSecond, and \p TThird all must be operands. - * - * @tparam TFirst The type of the first argument to the function. - * @tparam TSecond The type of the second argument to the function. - * @tparam TThird The type of the third argument to the function. - */ -template <typename TFirst, - typename TSecond, - typename TThird, - typename = ::std::enable_if<can_be_operand<TFirst>::value && can_be_operand<TSecond>::value && - can_be_operand<TThird>::value>> -struct TernaryElementwiseFunction -{ - TFirst first; - TSecond second; - TThird third; - TernaryFunction opcode; -}; - -template <typename TFirst, typename TSecond, typename TThird> -struct can_be_operand<TernaryElementwiseFunction<TFirst, TSecond, TThird>> : ::std::true_type -{ -}; - -/** Represents the function call: `select(\p first, \p second, \p third)`. - * - * @tparam TFirst The type of the first argument. - * @tparam TSecond The type of the second argument. - * @tparam TThird The type of the third argument. - * @param[in] first The first argument. - * @param[in] second The second argument. - * @param[in] third The third argument. - * @return The resulting AST node. - */ -template <typename TFirst, typename TSecond, typename TThird> -TernaryElementwiseFunction<TFirst, TSecond, TThird> select(TFirst &&first, TSecond &&second, TThird &&third) -{ - return TernaryElementwiseFunction<TFirst, TSecond, TThird>{std::forward<TFirst>(first), - std::forward<TSecond>(second), - std::forward<TThird>(third), TernaryFunction::Select}; -} - -/** Helper class used to extend a KernelWriter with additional functionality - * in order to make writing easier. - * - * This extension automatically handles creation of temporary variables, and - * allows nested function calls and operations. - * - * @tparam TWriter The type of KernelWriter to be overloaded. This must inherit from KernelWriter. - */ -template <class TWriter, typename = std::enable_if<std::is_base_of<KernelWriter, TWriter>::value>> -class KernelWriterHelper : public TWriter -{ -public: - using TWriter::TWriter; - - // ================================================== - // If-statements - // ================================================== - - // Un-hide original implementation, in case the original implementation is required. - using TWriter::op_if; - - /** Represents the if-statement: `if(\p cond) { \p body }`. - * - * The BinaryExpression is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] cond The BinaryExpression representing the condition. - * @param[in] body The body of the if-statement. - */ - KernelWriterHelper<TWriter> &op_if(const BinaryExpression<TileOperand &, TileOperand &> &cond, - const std::function<void()> &body) - { - TWriter::op_if(cond.lhs, cond.opcode, cond.rhs, body); - return *this; - } - - /** Represents the if-statement: `if(\p cond) { \p body }`. - * - * The BinaryExpression is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] cond The BinaryExpression representing the condition. - * @param[in] body The body of the if-statement. - */ - template <typename TRight> - KernelWriterHelper<TWriter> &op_if(const BinaryExpression<TileOperand &, TRight> &cond, - const std::function<void()> &body) - { - auto &tmp1 = declare_temp_tile(cond.lhs.tile_info()); - op_assign(tmp1, cond.rhs); - TWriter::op_if(cond.lhs, cond.opcode, tmp1, body); - return *this; - } - - /** Represents the if-statement: `if(\p cond) { \p body }`. - * - * The BinaryExpression is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] cond The BinaryExpression representing the condition. - * @param[in] body The body of the if-statement. - */ - template <typename TLeft> - KernelWriterHelper<TWriter> &op_if(const BinaryExpression<TLeft, TileOperand &> &cond, - const std::function<void()> &body) - { - auto &tmp1 = declare_temp_tile(cond.rhs.tile_info()); - op_assign(tmp1, cond.lhs); - TWriter::op_if(tmp1, cond.opcode, cond.rhs, body); - return *this; - } - - // Un-hide original implementation, in case the original implementation is required. - using TWriter::op_else_if; - - /** Represents the else-if-statement: `else if(\p cond) { \p body }`. - * - * The BinaryExpression is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] cond The BinaryExpression representing the condition. - * @param[in] body The body of the else-if-statement. - */ - KernelWriterHelper<TWriter> &op_else_if(const BinaryExpression<TileOperand &, TileOperand &> &cond, - const std::function<void()> &body) - { - TWriter::op_else_if(cond.lhs, cond.opcode, cond.rhs, body); - return *this; - } - - /** Represents the else-if-statement: `else if(\p cond) { \p body }`. - * - * The BinaryExpression is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] cond The BinaryExpression representing the condition. - * @param[in] body The body of the else-if-statement. - */ - template <typename TRight> - KernelWriterHelper<TWriter> &op_else_if(const BinaryExpression<TileOperand &, TRight> &cond, - const std::function<void()> &body) - { - auto &tmp1 = declare_temp_tile(cond.lhs.tile_info()); - op_assign(tmp1, cond.rhs); - TWriter::op_else_if(cond.lhs, cond.opcode, tmp1, body); - return *this; - } - - /** Represents the else-if-statement: `else if(\p cond) { \p body }`. - * - * The BinaryExpression is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] cond The BinaryExpression representing the condition. - * @param[in] body The body of the else-if-statement. - */ - template <typename TLeft> - KernelWriterHelper<TWriter> &op_else_if(const BinaryExpression<TLeft, TileOperand &> &cond, - const std::function<void()> &body) - { - auto &tmp1 = declare_temp_tile(cond.rhs.tile_info()); - op_assign(tmp1, cond.lhs); - TWriter::op_else_if(tmp1, cond.opcode, cond.rhs, body); - return *this; - } - - // ================================================== - // For-loops - // ================================================== - - // Un-hide original implementation, in case the original implementation is required. - using TWriter::op_for_loop; - - /** Represents the for-loop: `for(;\p cond; \p updater) { \p body }`. - * - * The BinaryExpression for the condition and the Assignment - * for the updater are unpacked and their components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] cond The BinaryExpression representing the condition. - * @param[in] updater The Assignment representing the updater. - * @param[in] body The body of the for-loop. - */ - void op_for_loop(const BinaryExpression<TileOperand &, TileOperand &> &cond, - const Assignment<TileOperand &, TileOperand &> &updater, - const std::function<void()> &body) - { - TWriter::op_for_loop(cond.lhs, cond.opcode, cond.rhs, updater.lhs, updater.opcode, updater.rhs, body); - } - - // ================================================== - // Unary expressions - // ================================================== - - // Un-hide original implementation, in case the original implementation is required. - using TWriter::op_assign; - - /** Represents the assignment: `\p dst = \p exp`. - * - * The UnaryExpression is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The UnaryExpression representing the expression to be evaluated and assigned. - */ - void op_assign(const TileOperand &dst, const UnaryExpression<TileOperand &> &exp) - { - TWriter::op_unary_expression(dst, exp.opcode, exp.src); - } - - // ================================================== - // Binary expressions - // ================================================== - - /** Represents the assignment: `\p dst = \p exp`. - * - * The BinaryExpression is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The BinaryExpression representing the expression to be evaluated and assigned. - */ - void op_assign(const TileOperand &dst, const BinaryExpression<TileOperand &, TileOperand &> &exp) - { - TWriter::op_binary_expression(dst, exp.lhs, exp.opcode, exp.rhs); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The BinaryExpression is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The BinaryExpression representing the expression to be evaluated and assigned. - */ - template <typename TRight> - void op_assign(const TileOperand &dst, const BinaryExpression<TileOperand &, TRight> &exp) - { - std::cout << "Beginning assignment!" << std::endl; - auto &tmp1 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.rhs); - TWriter::op_binary_expression(dst, exp.lhs, exp.opcode, tmp1); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The BinaryExpression is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The BinaryExpression representing the expression to be evaluated and assigned. - */ - template <typename TLeft> - void op_assign(const TileOperand &dst, const BinaryExpression<TLeft, TileOperand &> &exp) - { - std::cout << "Beginning assignment!" << std::endl; - auto &tmp1 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.lhs); - TWriter::op_binary_expression(dst, tmp1, exp.opcode, exp.rhs); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The BinaryExpression is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The BinaryExpression representing the expression to be evaluated and assigned. - */ - template <typename TLeft, typename TRight> - void op_assign(const TileOperand &dst, const BinaryExpression<TLeft, TRight> &exp) - { - auto &tmp1 = declare_temp_tile(dst.tile_info()); - auto &tmp2 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.lhs); - op_assign(tmp2, exp.rhs); - TWriter::op_binary_expression(dst, tmp1, exp.opcode, tmp2); - } - - // ================================================== - // Unary elementwise functions - // ================================================== - - /** Represents the assignment: `\p dst = \p exp`. - * - * The UnaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The UnaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - void op_assign(const TileOperand &dst, const UnaryElementwiseFunction<TileOperand &> &exp) - { - TWriter::op_unary_elementwise_function(dst, exp.opcode, exp.src); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The UnaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The UnaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - template <typename TArg> - void op_assign(const TileOperand &dst, const UnaryElementwiseFunction<TArg> &exp) - { - auto &tmp1 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.lhs); - TWriter::op_unary_elementwise_function(dst, exp.opcode, tmp1); - } - - // ================================================== - // Binary elementwise functions - // ================================================== - - /** Represents the assignment: `\p dst = \p exp`. - * - * The BinaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The BinaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - void op_assign(const TileOperand &dst, const BinaryElementwiseFunction<TileOperand &, TileOperand &> &exp) - { - TWriter::op_binary_elementwise_function(dst, exp.opcode, exp.first, exp.second); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The BinaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The BinaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - template <typename TRight> - void op_assign(const TileOperand &dst, const BinaryElementwiseFunction<TileOperand &, TRight> &exp) - { - auto &tmp1 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.second); - TWriter::op_binary_elementwise_function(dst, exp.opcode, exp.first, tmp1); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The BinaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The BinaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - template <typename TLeft> - void op_assign(const TileOperand &dst, const BinaryElementwiseFunction<TLeft, TileOperand &> &exp) - { - auto &tmp1 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.first); - TWriter::op_binary_elementwise_function(dst, exp.opcode, tmp1, exp.second); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The BinaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The BinaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - template <typename TLeft, typename TRight> - void op_assign(const TileOperand &dst, const BinaryElementwiseFunction<TLeft, TRight> &exp) - { - auto &tmp1 = declare_temp_tile(dst.tile_info()); - auto &tmp2 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.first); - op_assign(tmp2, exp.second); - TWriter::op_binary_elementwise_function(dst, exp.opcode, tmp1, tmp2); - } - - // ================================================== - // Ternary elementwise functions - // ================================================== - - /** Represents the assignment: `\p dst = \p exp`. - * - * The TernaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The TernaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - void op_assign(const TileOperand &dst, - const TernaryElementwiseFunction<TileOperand &, TileOperand &, TileOperand &> &exp) - { - TWriter::op_ternary_elementwise_function(dst, exp.opcode, exp.first, exp.second, exp.third); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The TernaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The TernaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - template <typename TFirst> - void op_assign(const TileOperand &dst, const TernaryElementwiseFunction<TFirst, TileOperand &, TileOperand &> &exp) - { - auto &tmp1 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.first); - TWriter::op_ternary_elementwise_function(dst, exp.opcode, tmp1, exp.second, exp.third); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The TernaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The TernaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - template <typename TSecond> - void op_assign(const TileOperand &dst, const TernaryElementwiseFunction<TileOperand &, TSecond, TileOperand &> &exp) - { - auto &tmp1 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.second); - TWriter::op_ternary_elementwise_function(dst, exp.opcode, exp.first, tmp1, exp.third); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The TernaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The TernaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - template <typename TThird> - void op_assign(const TileOperand &dst, const TernaryElementwiseFunction<TileOperand &, TileOperand &, TThird> &exp) - { - auto &tmp1 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.third); - TWriter::op_ternary_elementwise_function(dst, exp.opcode, exp.first, exp.second, tmp1); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The TernaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The TernaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - template <typename TFirst, typename TSecond> - void op_assign(const TileOperand &dst, const TernaryElementwiseFunction<TFirst, TSecond, TileOperand &> &exp) - { - auto &tmp1 = declare_temp_tile(dst.tile_info()); - auto &tmp2 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.first); - op_assign(tmp2, exp.second); - TWriter::op_ternary_elementwise_function(dst, exp.opcode, tmp1, tmp2, exp.third); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The TernaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The TernaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - template <typename TFirst, typename TThird> - void op_assign(const TileOperand &dst, const TernaryElementwiseFunction<TFirst, TileOperand &, TThird> &exp) - { - auto &tmp1 = declare_temp_tile(dst.tile_info()); - auto &tmp2 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.first); - op_assign(tmp2, exp.third); - TWriter::op_ternary_elementwise_function(dst, exp.opcode, tmp1, exp.second, tmp2); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The TernaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The TernaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - template <typename TSecond, typename TThird> - void op_assign(const TileOperand &dst, const TernaryElementwiseFunction<TileOperand &, TSecond, TThird> &exp) - { - auto &tmp1 = declare_temp_tile(dst.tile_info()); - auto &tmp2 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.second); - op_assign(tmp2, exp.third); - TWriter::op_ternary_elementwise_function(dst, exp.opcode, exp.first, tmp1, tmp2); - } - - /** Represents the assignment: `\p dst = \p exp`. - * - * The TernaryElementwiseFunction is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] dst The tile which is assigned to. - * @param[in] exp The TernaryElementwiseFunction representing the expression to be evaluated and assigned. - */ - template <typename TFirst, typename TSecond, typename TThird> - void op_assign(const TileOperand &dst, const TernaryElementwiseFunction<TFirst, TSecond, TThird> &exp) - { - auto &tmp1 = declare_temp_tile(dst.tile_info(), dst.tile_info(), dst.tile_info()); - auto &tmp2 = declare_temp_tile(dst.tile_info()); - auto &tmp3 = declare_temp_tile(dst.tile_info()); - op_assign(tmp1, exp.first); - op_assign(tmp2, exp.second); - op_assign(tmp3, exp.third); - TWriter::op_ternary_elementwise_function(dst, exp.opcode, tmp1, tmp2, tmp3); - } - - // ================================================== - // Assignments - // ================================================== - - /** Represents the assignment: `\p lhs += \p rhs` or `\p lhs -= \p rhs`. - * - * The Assignment is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @param[in] exp The Assignment representing the expression to be evaluated. - */ - void op_assign(const Assignment<TileOperand &, TileOperand &> &exp) - { - if (exp.opcode == AssignmentOp::Increment) - { - TWriter::op_binary_expression(exp.lhs, exp.lhs, BinaryOp::Add, exp.rhs); - } - else if (exp.opcode == AssignmentOp::Decrement) - { - TWriter::op_binary_expression(exp.lhs, exp.lhs, BinaryOp::Sub, exp.rhs); - } - } - - /** Represents the assignment: `\p lhs += \p rhs` or `\p lhs -= \p rhs`. - * - * The Assignment is unpacked and its components are forwarded to - * the underlying KernelWriter's implementation. - * - * @tparam TRight The type of the RHS of the assignment. - * @param[in] exp The Assignment representing the expression to be evaluated. - */ - template <typename TRight> - void op_assign(const Assignment<TileOperand &, TRight> &exp) - { - auto &tmp1 = declare_temp_tile(exp.lhs.tile_info()); - op_assign(tmp1, exp.rhs); - op_assign(Assignment<TileOperand &, TileOperand &>{exp.lhs, tmp1, exp.opcode}); - } - -private: - unsigned int temp_var_counter = 0; - - /** Return the current counter value, then increment it. - * - * @return The current counter value. - */ - int next_ctr() - { - return temp_var_counter++; - } - - /** Gets the next temporary variable counter value, - * and returns a suitable temporary variable name. - * - * @return A temporary variable name. - */ - std::string next_tmp_var_name() - { - return "tmp_" + std::to_string(next_ctr()); - } - - /** Returns the argument. - * - * Used for recursion with the variadic function version of this function. - * - * @param[in] arg The TileInfo to return. - * @return The \p arg. - */ - TileInfo get_largest_size(const TileInfo &arg) - { - return arg; - } - - /** Returns a TileInfo object where the size in each dimension (width, height) is the largest - * of either TileInfo argument in the corresponding dimension. - * - * @tparam TOps Must be of TileInfo type. - * @param[in] first A TileInfo object. - * @param[in] second A TileInfo object. - * @param[in] ops A number of TileInfo objects. - * @return A TileInfo object which represents the largest shape in each dimension across the arguments. - */ - template <typename... TOps, typename = ::std::enable_if_t<std::is_same<TOps..., TileInfo>::value>> - TileInfo get_largest_size(const TileInfo &first, const TileInfo &second, const TOps &...ops) - { - TileInfo largest = {first.data_type(), std::max(first.width(), second.width()), - std::max(first.height(), second.height())}; - return get_largest_size(largest, ops...); - } - - /** Helper function to define a suitable TileOperand with appropriate TileInfo - * such that broadcasting is taken into account, based on the arguments provided. - * - * @tparam TArgs Must be of TileInfo type. - * @param[in] args A number of TileInfo which determine the shape of the TileOperand to declare. - * @return A newly created TileOperand. - */ - template <typename... TArgs, typename = ::std::enable_if_t<std::is_same<TArgs..., TileInfo>::value>> - TileOperand &declare_temp_tile(const TArgs &...args) - { - return TWriter::declare_tile(next_tmp_var_name().c_str(), get_largest_size(args...)); - } -}; - -} // namespace ckw - -#endif // CKW_INCLUDE_CKW_KERNELWRITERHELPER_H diff --git a/compute_kernel_writer/prototype/include/ckw/OperandBase.h b/compute_kernel_writer/prototype/include/ckw/OperandBase.h deleted file mode 100644 index 9842127339..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/OperandBase.h +++ /dev/null @@ -1,78 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_PROTOTYPE_INCLUDE_CKW_OPERANDBASE_H -#define CKW_PROTOTYPE_INCLUDE_CKW_OPERANDBASE_H - -#include "ckw/types/DataType.h" - -#include <string> - -namespace ckw -{ -namespace prototype -{ -class IGpuKernelWriter; - -class Operand; -} // namespace prototype - -/** The base class for all operands. */ -class OperandBase -{ -public: - /** Constructor - * - * @param[in] name The name of the operand. - */ - explicit OperandBase(const ::std::string &name); - - /** Destructor */ - virtual ~OperandBase(); - - /** (Internal use only) Create the implementation operand. - * - * @param[in] writer The implementation kernel writer. - */ - virtual prototype::Operand create_impl_operand(prototype::IGpuKernelWriter *writer) const = 0; - - /** Get the name of the operand. */ - const ::std::string &name() const; - - /** Set the name of the operand. */ - OperandBase &name(const ::std::string &name); - - /** Get the data type of the operand. */ - virtual DataType data_type() const = 0; - - /** Get whether the operand is compile-time constant. */ - virtual bool is_constant() const = 0; - -private: - ::std::string _name; -}; - -} // namespace ckw - -#endif // CKW_PROTOTYPE_INCLUDE_CKW_OPERANDBASE_H diff --git a/compute_kernel_writer/prototype/include/ckw/ScalarValue.h b/compute_kernel_writer/prototype/include/ckw/ScalarValue.h deleted file mode 100644 index 2a9c42acc8..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/ScalarValue.h +++ /dev/null @@ -1,137 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_PROTOTYPE_INCLUDE_CKW_SCALARVALUE_H -#define CKW_PROTOTYPE_INCLUDE_CKW_SCALARVALUE_H - -#include "ckw/Error.h" - -#include <cstdint> - -namespace ckw -{ - -/** The scalar value known at compile-time. */ -class ScalarValue -{ -public: - /** Initialize a new instance of @ref ScalarValue class with integer value 0. */ - ScalarValue() - { - _type = Type::INT; - _value.i64 = 0; - } - - /** Initialize a new instance of @ref ScalarValue class with the specified value. */ - template <typename T> - ScalarValue(T value) - { - set(value); - } - - /** Set the value. */ - template <typename T> - void set(T value) - { - CKW_ASSERT(::std::is_integral<T>::value || ::std::is_floating_point<T>::value); - CKW_ASSERT(sizeof(T) <= 8); - - _size = sizeof(T); - - if (::std::is_integral<T>::value) - { - if (::std::is_signed<T>::value) - { - _type = Type::INT; - _value.i64 = value; - } - else - { - _type = Type::UINT; - _value.u64 = value; - } - } - else - { - _type = Type::FLOAT; - _value.f64 = value; - } - } - - /** Get the value. - * - * The caller must make sure that what has been stored in the object must fit - * the output data type without data corruption or loss of accuracy. - */ - template <typename T> - T get() const - { - CKW_ASSERT(::std::is_integral<T>::value || ::std::is_floating_point<T>::value); - CKW_ASSERT(sizeof(T) >= _size); - - if (::std::is_integral<T>::value) - { - if (::std::is_signed<T>::value) - { - CKW_ASSERT(_type == Type::INT || _type == Type::UINT); - CKW_ASSERT_IF(_type == Type::UINT, sizeof(T) > _size); - - return _value.i64; - } - else - { - CKW_ASSERT(_type == Type::INT); - - return _value.u64; - } - } - else - { - return _value.f64; - } - } - -private: - union Value - { - int64_t i64; - uint64_t u64; - double f64; - }; - - enum class Type : int32_t - { - UINT, - INT, - FLOAT, - }; - - Value _value{}; - Type _type{}; - uint32_t _size{}; -}; - -} // namespace ckw - -#endif // CKW_PROTOTYPE_INCLUDE_CKW_SCALARVALUE_H diff --git a/compute_kernel_writer/prototype/include/ckw/TensorInfo.h b/compute_kernel_writer/prototype/include/ckw/TensorInfo.h deleted file mode 100644 index 24da7dc8ab..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/TensorInfo.h +++ /dev/null @@ -1,153 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_PROTOTYPE_INCLUDE_CKW_TENSORINFO_H -#define CKW_PROTOTYPE_INCLUDE_CKW_TENSORINFO_H - -#include "ckw/types/DataType.h" - -#include <array> -#include <cstdint> - -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, -}; - -/** Compute Kernel Writer tensor component bitmask. The bitmask can be used to retrieve - * the info from @ref TensorComponent. - */ -enum class TensorComponentBitmask : uint32_t -{ - OffsetFirstElement = 0x01000000, // For example, OffsetFirstElement in @ref TensorComponent - Stride = 0x02000000, // For example, stride0 in @ref TensorComponent - Dimension = 0x04000000, // For example, Dim0 in @ref TensorComponent - FoldedDimensions = 0x08000000, // For example, Dim0xDim1 in @ref TensorComponent -}; - -/** 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 data type is represented as an 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 = 0x02000010, - Stride2 = 0x02000100, - Stride3 = 0x02001000, - Stride4 = 0x02010000, - Dim0 = 0x04000001, - Dim1 = 0x04000010, - Dim2 = 0x04000100, - Dim3 = 0x04001000, - Dim4 = 0x04010000, - Dim1xDim2 = 0x08000110, - Dim2xDim3 = 0x08001100, - Dim1xDim2xDim3 = 0x08001110 -}; - -/** 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, -}; - -/** Compute Kernel Writer tensor shape - * Negative dimensions can be interpreted as dynamic dimensions by the Compute Kernel Writer - */ -using TensorShape = std::array<int32_t, 5>; - -/** Compute Kernel Writer tensor info */ -class TensorInfo -{ -public: - /** 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 bound user tensor. 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 /* CKW_PROTOTYPE_INCLUDE_CKW_TENSORINFO_H */ diff --git a/compute_kernel_writer/prototype/include/ckw/TensorOperand.h b/compute_kernel_writer/prototype/include/ckw/TensorOperand.h deleted file mode 100644 index c221b449fa..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/TensorOperand.h +++ /dev/null @@ -1,196 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_PROTOTYPE_INCLUDE_CKW_TENSOROPERAND_H -#define CKW_PROTOTYPE_INCLUDE_CKW_TENSOROPERAND_H - -#include "ckw/OperandBase.h" -#include "ckw/TensorInfo.h" -#include "ckw/TensorTileSampler.h" -#include "ckw/TileOperand.h" -#include "ckw/types/DataType.h" - -#include <memory> - -namespace ckw -{ - -class TensorComponentOperand; - -// ================================================================================================= -// TensorOperand -// ================================================================================================= - -/** Tensor operand */ -class TensorOperand : public OperandBase -{ -public: - /** Initialize a new instance of @ref TensorOperand class. - * - * @param[in] name The name of the tensor. - * @param[in] info The tensor info. - * @param[in] storage_type The tensor storage type. - */ - TensorOperand(const ::std::string &name, const TensorInfo &info, TensorStorageType storage_type); - - /** No copy constructor. */ - TensorOperand(const TensorOperand &other) = delete; - - /** No copy assignment. */ - TensorOperand &operator=(const TensorOperand &other) = delete; - - /** (Internal use only) Create the implementation operand. - * - * @param[in] writer The implementation kernel writer. - */ - virtual prototype::Operand create_impl_operand(prototype::IGpuKernelWriter *writer) const override; - - /** Get the tensor info. */ - const TensorInfo &info() const; - - /** Get the tensor info. */ - TensorInfo &info(); - - /** Get the tensor storage type. */ - TensorStorageType storage_type() const; - - /** Get the data type. */ - virtual DataType data_type() const override; - - /** Get whether the tensor is compile-time constant. */ - virtual bool is_constant() const override; - - /** Get the default tile attached to the tensor. */ - const TileOperand &tile() const; - - /** Get the default tile attached to the tensor. */ - TileOperand &tile(); - - /** Set the default tile attached to the tensor. */ - TensorOperand &tile(TileOperand &tile); - - /** Get the tensor sampler of the default tile. */ - const TensorTileSampler &tile_sampler() const; - - /** Get the tensor sampler of the default tile. */ - TensorTileSampler &tile_sampler(); - - /** Set the tensor sampler of the default tile. */ - TensorOperand &tile_sampler(const TensorTileSampler &value); - - /** Get the operand that contains the stride in y dimension of the tensor. */ - TensorComponentOperand &stride1(); - - /** Get the operand that contains the stride in z dimension of the tensor. */ - TensorComponentOperand &stride2(); - - /** Get the operand that contains the stride in w dimension of the tensor. */ - TensorComponentOperand &stride3(); - - /** Get the operand that contains the stride in w dimension of the tensor. */ - TensorComponentOperand &stride4(); - - /** Get the operand that contains the size of dimension 0 of the tensor. */ - TensorComponentOperand &dim0(); - - /** Get the operand that contains the size of dimension 1 of the tensor. */ - TensorComponentOperand &dim1(); - - /** Get the operand that contains the size of dimension 2 of the tensor. */ - TensorComponentOperand &dim2(); - - /** Get the operand that contains the size of dimension 3 of the tensor. */ - TensorComponentOperand &dim3(); - - /** Get the operand that contains the size of dimension 4 of the tensor. */ - TensorComponentOperand &dim4(); - - /** Get the operand that contains the size of dimensions 1 and 2 collapsed. */ - TensorComponentOperand &dim1_dim2(); - - /** Get the operand that contains the size of dimensions 1, 2 and 3 collapsed. */ - TensorComponentOperand &dim1_dim2_dim3(); - - /** Get the operand that contains the offset in bytes to the first element. */ - TensorComponentOperand &offset_first_element_in_bytes(); - -private: - TensorInfo _info; - TensorStorageType _storage_type; - - TileOperand *_tile{nullptr}; - TensorTileSampler _tile_sampler{}; - - ::std::unique_ptr<TensorComponentOperand> _stride1{nullptr}; - ::std::unique_ptr<TensorComponentOperand> _stride2{nullptr}; - ::std::unique_ptr<TensorComponentOperand> _stride3{nullptr}; - ::std::unique_ptr<TensorComponentOperand> _stride4{nullptr}; - ::std::unique_ptr<TensorComponentOperand> _dim0{nullptr}; - ::std::unique_ptr<TensorComponentOperand> _dim1{nullptr}; - ::std::unique_ptr<TensorComponentOperand> _dim2{nullptr}; - ::std::unique_ptr<TensorComponentOperand> _dim3{nullptr}; - ::std::unique_ptr<TensorComponentOperand> _dim4{nullptr}; - ::std::unique_ptr<TensorComponentOperand> _dim1_dim2{nullptr}; - ::std::unique_ptr<TensorComponentOperand> _dim1_dim2_dim3{nullptr}; - ::std::unique_ptr<TensorComponentOperand> _offset_first_element_in_bytes{nullptr}; -}; - -// ================================================================================================= -// TensorComponentOperand -// ================================================================================================= - -/** Tile operand that contains tensor information. */ -class TensorComponentOperand : public TileOperand -{ -public: - /** Initialize a new instance of @ref TensorComponentOperand class. - * - * @param[in] tensor The tensor operand. - * @param[in] component The tensor info component. - */ - TensorComponentOperand(TensorOperand &tensor, TensorComponentType component); - - /** Get the tensor operand. */ - TensorOperand &tensor(); - - /** Get the tensor operand. */ - const TensorOperand &tensor() const; - - /** Get the tensor component. */ - TensorComponentType component_type() const; - - /** (Internal use only) Create the implementation operand. - * - * @param[in] writer The implementation kernel writer. - */ - virtual prototype::Operand create_impl_operand(prototype::IGpuKernelWriter *writer) const override; - -private: - TensorOperand &_tensor; - TensorComponentType _component; -}; - -} // namespace ckw - -#endif // CKW_PROTOTYPE_INCLUDE_CKW_TENSOROPERAND_H diff --git a/compute_kernel_writer/prototype/include/ckw/TensorTileSampler.h b/compute_kernel_writer/prototype/include/ckw/TensorTileSampler.h deleted file mode 100644 index 606dec3535..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/TensorTileSampler.h +++ /dev/null @@ -1,169 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_PROTOTYPE_INCLUDE_CKW_TENSORTILESAMPLER_H -#define CKW_PROTOTYPE_INCLUDE_CKW_TENSORTILESAMPLER_H - -#include "ckw/types/TensorSamplerTypes.h" - -#include <functional> - -namespace ckw -{ - -class TileOperand; - -/** Tensor sampler - * - * It contains information about how the result tile should be stored to tensor memory. - * It can also be used to dictate how the subsequent operators fetch the input tensor. - */ -class TensorTileSampler -{ -public: - /** Initialize a new instance of @ref TensorSampler class. */ - TensorTileSampler(); - - /** Initialize a new instance of @ref TensorSampler class. - * - * @param[in] x The coordinate in the x dimension. - * @param[in] y The coordinate in the y dimension. - * @param[in] z The coordinate in the z dimension. - * @param[in] b The coordinate in the batch dimension. - * @param[in] format The tensor data format. - * @param[in] address_mode_x The address mode of the x dimension. - * @param[in] address_mode_y The address mode of the y dimension. - * @param[in] address_mode_z The address mode of the z dimension. - */ - TensorTileSampler(TileOperand &x, - TileOperand &y, - TileOperand &z, - TileOperand &b, - TensorSamplerFormat format, - TensorSamplerAddressModeX address_mode_x, - TensorSamplerAddressModeY address_mode_y, - TensorSamplerAddressModeZ address_mode_z); - - /** Initialize a new instance of @ref TensorSampler class. - * - * @param[in] x The coordinate in the x dimension. - * @param[in] y The coordinate in the y dimension. - * @param[in] z The coordinate in the z dimension. - * @param[in] b The coordinate in the batch dimension. - * @param[in] height The height of the tile. - * @param[in] width The width of the tile. - * @param[in] format The tensor data format. - * @param[in] address_mode_x The address mode of the x dimension. - * @param[in] address_mode_y The address mode of the y dimension. - * @param[in] address_mode_z The address mode of the z dimension. - */ - TensorTileSampler(TileOperand &x, - TileOperand &y, - TileOperand &z, - TileOperand &b, - int32_t height, - int32_t width, - TensorSamplerFormat format, - TensorSamplerAddressModeX address_mode_x, - TensorSamplerAddressModeY address_mode_y, - TensorSamplerAddressModeZ address_mode_z); - - /** Get the coordinate in the x dimension. */ - const TileOperand &x() const; - - /** Set the coordinate in the x dimension. */ - TensorTileSampler &x(TileOperand &x); - - /** Get the coordinate in the y dimension. */ - const TileOperand &y() const; - - /** Set the coordinate in the y dimension. */ - TensorTileSampler &y(TileOperand &y); - - /** Get the coordinate in the z dimension. */ - const TileOperand &z() const; - - /** Set the coordinate in the z dimension. */ - TensorTileSampler &z(TileOperand &z); - - /** Get the coordinate in the batch dimension. */ - const TileOperand &b() const; - - /** Set the coordinate in the batch dimension. */ - TensorTileSampler &b(TileOperand &b); - - /** Get the width of the tile. */ - int32_t width() const; - - /** Set the width of the tile. */ - TensorTileSampler &width(int32_t width); - - /** Get the height of the tile. */ - int32_t height() const; - - /** Set the height of the tile. */ - TensorTileSampler &height(int32_t height); - - /** Get the format of the tensor. */ - TensorSamplerFormat format() const; - - /** Set the format of the tensor. */ - TensorTileSampler &format(TensorSamplerFormat format); - - /** Get the address mode of the x dimension. */ - TensorSamplerAddressModeX address_mode_x() const; - - /** Set the address mode of the x-dimension. */ - TensorTileSampler &address_mode_x(TensorSamplerAddressModeX address_mode_x); - - /** Get the address mode of the y dimension. */ - TensorSamplerAddressModeY address_mode_y() const; - - /** Set the address mode of the y dimension. */ - TensorTileSampler &address_mode_y(TensorSamplerAddressModeY address_mode_y); - - /** Get the address mode of the z dimension. */ - TensorSamplerAddressModeZ address_mode_z() const; - - /** Set the address mode of the z dimension. */ - TensorTileSampler &address_mode_z(TensorSamplerAddressModeZ address_mode_z); - -private: - TileOperand *_x{nullptr}; - TileOperand *_y{nullptr}; - TileOperand *_z{nullptr}; - TileOperand *_b{nullptr}; - - int32_t _height{0}; - int32_t _width{0}; - - TensorSamplerFormat _format{TensorSamplerFormat::Unknown}; - TensorSamplerAddressModeX _address_mode_x{TensorSamplerAddressModeX::Unknown}; - TensorSamplerAddressModeY _address_mode_y{TensorSamplerAddressModeY::Unknown}; - TensorSamplerAddressModeZ _address_mode_z{TensorSamplerAddressModeZ::Unknown}; -}; - -} // namespace ckw - -#endif // CKW_PROTOTYPE_INCLUDE_CKW_TENSORTILESAMPLER_H diff --git a/compute_kernel_writer/prototype/include/ckw/TileInfo.h b/compute_kernel_writer/prototype/include/ckw/TileInfo.h deleted file mode 100644 index e0d064169e..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/TileInfo.h +++ /dev/null @@ -1,92 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_PROTOTYPE_INCLUDE_CKW_TILEINFO_H -#define CKW_PROTOTYPE_INCLUDE_CKW_TILEINFO_H - -#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_H */ diff --git a/compute_kernel_writer/prototype/include/ckw/TileOperand.h b/compute_kernel_writer/prototype/include/ckw/TileOperand.h deleted file mode 100644 index 24ee373a24..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/TileOperand.h +++ /dev/null @@ -1,127 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_PROTOTYPE_INCLUDE_CKW_TILEOPERAND_H -#define CKW_PROTOTYPE_INCLUDE_CKW_TILEOPERAND_H - -#include "ckw/Error.h" -#include "ckw/OperandBase.h" -#include "ckw/ScalarValue.h" -#include "ckw/TileInfo.h" - -#include <vector> - -namespace ckw -{ - -class Kernel; - -using TileContainer = std::vector<std::vector<std::string>>; - -/** Tile operand which can be either scalar, vector or 2D tile. */ -class TileOperand : public OperandBase -{ -public: - /** Initialize a new instance of @ref TileOperand class with the tile information. - * - * @param[in] name The name of the tile. - * @param[in] tile_info The tile info. - */ - TileOperand(const ::std::string &name, const TileInfo &tile_info); - - /** Initialize a new instance of @ref TileOperand for scalar variable. - * - * @param[in] name The name of the tile. - * @param[in] data_type The data type of the tile. - */ - TileOperand(const ::std::string &name, DataType data_type); - - /** Initialize a new instance of @ref TileOperand for compile-time constant scalar variable. - * - * @param[in] name The name of the tile. - * @param[in] value The value of the tile. - */ - TileOperand(const ::std::string &name, int32_t value); - - /** Initialize a new instance of @ref TileOperand for compile-time constant scalar variable. - * - * @param[in] name The name of the tile. - * @param[in] value The value of the tile. - */ - TileOperand(const ::std::string &name, float value); - - /** Initialize a new instance of @ref TileOperand for compile-time constant variable. - * - * @param[in] name The name of the tile. - * @param[in] value The value of the tile. - */ - TileOperand(const ::std::string &name, const ::std::vector<std::vector<std::string>> &value, DataType dt); - - /** Prohibit copy of tile operand. */ - TileOperand(const TileOperand &) = delete; - - /** Prohibit copy of tile operand. */ - TileOperand &operator=(const TileOperand &) = delete; - - /** (Internal use only) Create the implementation operand. - * - * @param[in] writer The implementation kernel writer. - */ - virtual prototype::Operand create_impl_operand(prototype::IGpuKernelWriter *writer) const override; - - /** Get the tile info. */ - const TileInfo &tile_info() const; - - /** Get the data type of the tile. */ - virtual DataType data_type() const override; - - /** Get whether the tile is compile-time constant. */ - virtual bool is_constant() const override; - - /** Get whether the tile is a scalar value. */ - bool is_scalar() const; - - /** Get the scalar value of the tile. - * - * The tile must have the shape of 1, 1 (i.e. scalar). - * - * @return Scalar value as a string. - */ - std::string scalar_value() const; - - /** Get the values of the tile. - * - * @return 2D container of values. - */ - const TileContainer &value() const; - -private: - TileInfo _info; - TileContainer _value{}; - bool _constant; -}; - -} // namespace ckw - -#endif // CKW_PROTOTYPE_INCLUDE_CKW_TILEOPERAND_H diff --git a/compute_kernel_writer/prototype/include/ckw/types/ConvertPolicy.h b/compute_kernel_writer/prototype/include/ckw/types/ConvertPolicy.h deleted file mode 100644 index 2a198507eb..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/types/ConvertPolicy.h +++ /dev/null @@ -1,41 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_INCLUDE_CKW_CONVERTPOLICY_H -#define CKW_INCLUDE_CKW_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_CONVERTPOLICY_H diff --git a/compute_kernel_writer/prototype/include/ckw/types/DataType.h b/compute_kernel_writer/prototype/include/ckw/types/DataType.h deleted file mode 100644 index 3447dd61d6..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/types/DataType.h +++ /dev/null @@ -1,50 +0,0 @@ -/* -* Copyright (c) 2023 Arm Limited. -* -* SPDX-License-Identifier: MIT -* -* Permission is hereby granted, free of charge, to any person obtaining a copy -* of this software and associated documentation files (the "Software"), to -* deal in the Software without restriction, including without limitation the -* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or -* sell copies of the Software, and to permit persons to whom the Software is -* furnished to do so, subject to the following conditions: -* -* The above copyright notice and this permission notice shall be included in all -* copies or substantial portions of the Software. -* -* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -* SOFTWARE. -*/ - -#ifndef CKW_INCLUDE_CKW_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/prototype/include/ckw/types/Functions.h b/compute_kernel_writer/prototype/include/ckw/types/Functions.h deleted file mode 100644 index c6afaa0ac8..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/types/Functions.h +++ /dev/null @@ -1,62 +0,0 @@ -/* -* Copyright (c) 2023 Arm Limited. -* -* SPDX-License-Identifier: MIT -* -* Permission is hereby granted, free of charge, to any person obtaining a copy -* of this software and associated documentation files (the "Software"), to -* deal in the Software without restriction, including without limitation the -* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or -* sell copies of the Software, and to permit persons to whom the Software is -* furnished to do so, subject to the following conditions: -* -* The above copyright notice and this permission notice shall be included in all -* copies or substantial portions of the Software. -* -* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -* SOFTWARE. -*/ - -#ifndef CKW_INCLUDE_CKW_FUNCTIONS_H -#define CKW_INCLUDE_CKW_FUNCTIONS_H - -#include <cstdint> - -namespace ckw -{ - -enum class UnaryFunction : int32_t -{ - Exp = 0x0000, - Tanh = 0x0001, - Sqrt = 0x0002, - Erf = 0x0003, - Fabs = 0x0004, - Log = 0x0006, - Round = 0x0007, - Floor = 0x0008, - - // Misc - SizeOf = 0x0009, -}; - -enum class BinaryFunction : int32_t -{ - Min = 0x0000, - Max = 0x0001, -}; - -enum class TernaryFunction : int32_t -{ - Select = 0x0000, - Clamp = 0x0001, -}; - -} // namespace ckw - -#endif //CKW_INCLUDE_CKW_FUNCTIONS_H diff --git a/compute_kernel_writer/prototype/include/ckw/types/GpuTargetLanguage.h b/compute_kernel_writer/prototype/include/ckw/types/GpuTargetLanguage.h deleted file mode 100644 index 6c08617949..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/types/GpuTargetLanguage.h +++ /dev/null @@ -1,41 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_INCLUDE_CKW_GPUTARGETLANGUAGE_H -#define CKW_INCLUDE_CKW_GPUTARGETLANGUAGE_H - -#include <cstdint> - -namespace ckw -{ - -enum class GpuTargetLanguage : int32_t -{ - Unknown, - OpenCL -}; - -} // namespace ckw - -#endif //CKW_INCLUDE_CKW_GPUTARGETLANGUAGE_H diff --git a/compute_kernel_writer/prototype/include/ckw/types/Operators.h b/compute_kernel_writer/prototype/include/ckw/types/Operators.h deleted file mode 100644 index b560996837..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/types/Operators.h +++ /dev/null @@ -1,78 +0,0 @@ -/* -* Copyright (c) 2023 Arm Limited. -* -* SPDX-License-Identifier: MIT -* -* Permission is hereby granted, free of charge, to any person obtaining a copy -* of this software and associated documentation files (the "Software"), to -* deal in the Software without restriction, including without limitation the -* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or -* sell copies of the Software, and to permit persons to whom the Software is -* furnished to do so, subject to the following conditions: -* -* The above copyright notice and this permission notice shall be included in all -* copies or substantial portions of the Software. -* -* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -* SOFTWARE. -*/ - -#ifndef CKW_INCLUDE_CKW_OPERATORS_H -#define CKW_INCLUDE_CKW_OPERATORS_H - -#include <cstdint> - -namespace ckw -{ - -enum class UnaryOp : int32_t -{ - LogicalNot = 0x0000, // ! - BitwiseNot = 0x0001, // ~ - Negate = 0x0002, // - -}; - -/* Binary operations -*/ -enum class BinaryOp : int32_t -{ - // Elementwise - Add = 0x0000, // + - Sub = 0x0001, // - - Mul = 0x0002, // * - Div = 0x0003, // / - Mod = 0x0004, // % - // Relational - Equal = 0x1000, // == - Less = 0x1001, // < - LessEqual = 0x1002, // <= - Greater = 0x1003, // > - GreaterEqual = 0x1004, // >= - // Algebra - MatMul_Nt_Nt = 0x2000, // X - MatMul_Nt_T = 0x2001, // X - MatMul_T_Nt = 0x2002, // X - MatMul_T_T = 0x2003, // X - Dot = 0x2004, // . - // Logical - LogicalAnd = 0x3000, // && - LogicalOr = 0x3001, // || - // Bitwise - BitwiseXOR = 0x4000, // ^ -}; - -enum class AssignmentOp : int32_t -{ - // Unary - Increment = 0x0000, // += - Decrement = 0x0001, // -= -}; - -} // namespace ckw - -#endif //CKW_INCLUDE_CKW_OPERATORS_H diff --git a/compute_kernel_writer/prototype/include/ckw/types/TensorSamplerTypes.h b/compute_kernel_writer/prototype/include/ckw/types/TensorSamplerTypes.h deleted file mode 100644 index 63405a0764..0000000000 --- a/compute_kernel_writer/prototype/include/ckw/types/TensorSamplerTypes.h +++ /dev/null @@ -1,82 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_INCLUDE_CKW_TENSORSAMPLERTYPES_H -#define CKW_INCLUDE_CKW_TENSORSAMPLERTYPES_H - -#include <cstdint> - -namespace ckw -{ - -enum class TensorSamplerFormat : int32_t -{ - Unknown = 0, - C_WH_1 = 1, - C_W_H = 2 -}; - -enum class TensorSamplerAddressModeX : int32_t -{ - Unknown = 0, - None = 1, // The user guarantees that the X coordinate is always in-bound - OverlappingMin = - 2 // (FIXED shapes only) Reduce the load/store length when x == 0 (MIN). The load length will be width % original length - // Leftover elements can be handled using overlapping. This involves processing some of the elements in the array twice. -}; - -enum class TensorSamplerAddressModeY : int32_t -{ - Unknown = 0, - None = 1, // The user guarantees that the Y coordinate is always in-bound - OverlappingMin = - 2, // (FIXED shapes only) Reduce the load/store length when x == 0 (MIN). The load length will be width % original length - Skip = 3, // Skip the read/write - SkipMinEdgeOnly = - 4, // Skip greater than or equal to max only. The user guarantees that the Y coordinate is always >= 0 - SkipMaxEdgeOnly = 5, // Skip less than 0 only - ClampToNearest = 6, // Clamp the coordinate to nearest edge (0 or max value allowed on Y) - ClampToMinEdgeOnly = 7, // Clamp the negative coordinate to 0 only. Therefore, we expect Y to be always < MAX - ClampToMaxEdgeOnly = 8, // Clamp the coordinate to the max value allowed on Y only. We expect Y to be always >= 0 - ClampToBorder = 9, // Clamp to border which always has 0 value - ClampToBorderMinEdgeOnly = 10, - ClampToBorderMaxEdgeOnly = 11 -}; - -enum class TensorSamplerAddressModeZ : int32_t -{ - Unknown = 0, - None = 1, // The user guarantees that the Y coordinate is always in-bound - Skip = 3, // Skip the read/write - SkipMinEdgeOnly = - 4, // Skip greater than or equal to max only. The user guarantees that the Y coordinate is always >= 0 - SkipMaxEdgeOnly = 5, // Skip less than 0 only - ClampToNearest = 6, // Clamp the coordinate to nearest edge (0 or max value allowed on Y) - ClampToMinEdgeOnly = 7, // Clamp the negative coordinate to 0 only. Therefore, we expect Y to be always < MAX - ClampToMaxEdgeOnly = 8, // Clamp the coordinate to the max value allowed on Y only. We expect Y to be always >= 0 -}; - -} // namespace ckw - -#endif //CKW_INCLUDE_CKW_TENSORSAMPLERTYPES_H |