diff options
Diffstat (limited to 'compute_kernel_writer/prototype/include')
11 files changed, 1409 insertions, 0 deletions
diff --git a/compute_kernel_writer/prototype/include/ckw/Error.h b/compute_kernel_writer/prototype/include/ckw/Error.h new file mode 100644 index 0000000000..b18944eac5 --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/Error.h @@ -0,0 +1,79 @@ +/* + * 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 new file mode 100644 index 0000000000..57a8a40341 --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/Kernel.h @@ -0,0 +1,77 @@ +/* + * 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/OperandBase.h" +#include "ckw/Types.h" + +#include <map> +#include <memory> +#include <string> + +namespace ckw +{ + +namespace prototype +{ +class GpuKernelWriterDataHolder; +} // namespace prototype + +/** The target for kernel writer to write into. */ +class Kernel +{ +public: + /** Constructor + * + * @param[in] name The name of the kernel function. + * @param[in] language The programming language to write the kernel. + */ + Kernel(const char *name, GpuTargetLanguage language); + + /** Destructor */ + ~Kernel(); + + /** Get the name of the kernel function. */ + const std::string &name() const; + + /** (Internal use only) Get the map from operand name to the operand declared in this kernel. */ + const ::std::map<::std::string, ::std::unique_ptr<OperandBase>> &operands() const; + + /** (Internal use only) Get the map from operand name to the operand declared in this kernel. */ + ::std::map<::std::string, ::std::unique_ptr<OperandBase>> &operands(); + + /** (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; +}; + +} // namespace ckw + +#endif // CKW_PROTOTYPE_INCLUDE_CKW_KERNEL_H diff --git a/compute_kernel_writer/prototype/include/ckw/KernelWriter.h b/compute_kernel_writer/prototype/include/ckw/KernelWriter.h new file mode 100644 index 0000000000..a2778a9485 --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/KernelWriter.h @@ -0,0 +1,217 @@ +/* + * 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 <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 + // ============================================================================================= + + /** Define a tensor argument. + * + * @param[in] name The name of the tensor. + * @param[in] info The tensor info. + * + * @return The @ref TensorOperand object. + */ + TensorOperand &create_tensor_argument(const char *name, const TensorInfo &info); + + /** Define a compile-time constant scalar argument. + * + * @param[in] name The name of the tile. + * @param[in] value The value of the tile. + * + * @return The @ref TileOperand object. + */ + TileOperand &create_tile_argument(const char *name, int32_t value); + + /** Declare a new tile. + * + * The name of the tile must be unique in the current ID space. + * + * @param[in] name The name of the tile. + * @param[in] ... The necessary arguments to create a new @ref TileOperand. + * + * @return The @ref TileOperand object. + */ + template <typename... TArgs> + TileOperand &declare_tile(const char *name, TArgs &&...args) + { + const auto var_name = generate_variable_name(name); + auto operand = new TileOperand(var_name, ::std::forward<TArgs>(args)...); + register_operand(operand, true); + + return *operand; + } + + // ============================================================================================= + // Load and store + // ============================================================================================= + + /** Load the data from the tensor memory to the tile using the sampling information. + * + * @param[out] tile The tile to be loaded. + * @param[in] tensor The tensor to be read. + * @param[in] sampler The tensor sampling information. + */ + void op_load(TileOperand &tile, TensorOperand &tensor, const TensorTileSampler &sampler); + + /** Store the tile to the tensor using the specified sampling information. + * + * @param[out] dst The tensor that the tile is written to. + * @param[in] src The tile to be stored. + * @param[in] sampler The tensor sampling information. + */ + void op_store(TensorOperand &tensor, const TileOperand &tile, const TensorTileSampler &sampler); + + // ============================================================================================= + // Data processing + // ============================================================================================= + + /** Write assignment: `<dst> = <src>`. + * + * @param[in] dst The destination tile. + * @param[in] src The source tile. + */ + void op_assign(TileOperand &dst, const TileOperand &src); + + /** Write binary expression: `<dst> = <lhs> <op> <rhs>`. + * + * @param[in] dst The destination tile. + * @param[in] lhs The LHS operand. + * @param[in] rhs The RHS operand. + * @param[in] op The binary operator. + */ + void op_binary_expression(TileOperand &dst, const TileOperand &lhs, const TileOperand &rhs, BinaryOp op); + + /** Write function applied to scalar value: `<dst> = <func>(<src>)`. + * + * @param[in] dst The destination tile. + * @param[in] src The source tile. + * @param[in] func The function to be applied to the source tile. + */ + void op_scalar_function(TileOperand &dst, const TileOperand &src, ScalarUnaryFunction func); + + // ============================================================================================= + // Misc + // ============================================================================================= + + /** Set `dst` the global ID of dimension `dim`. + * + * @param[in] dst The tile to be written to. + * @param[in] dim The global ID dimension. + */ + void op_get_global_id(TileOperand &dst, int32_t dim); + + // ============================================================================================= + // Code generation + // ============================================================================================= + + /** Generate the source code of the kernel. */ + ::std::string generate_code(); + +private: + /** Generate the full variable name based on the original name and the ID space. + * + * @param[in] name The name of the variable. + * + * @return The full variable name. + */ + ::std::string generate_variable_name(const char *name) const; + + /** Register the operand to the kernel. + * + * The operand is uniquely owned by the kernel afterward. + * + * @param[in] operand The operand to be registered. + * @param[in] declaring Whether the tile declaration is generated. + */ + void register_operand(OperandBase *operand, bool declaring); + +private: + Kernel *_kernel; + ::std::unique_ptr<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/OperandBase.h b/compute_kernel_writer/prototype/include/ckw/OperandBase.h new file mode 100644 index 0000000000..f4825e16a7 --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/OperandBase.h @@ -0,0 +1,76 @@ +/* + * Copyright (c) 2023 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ + +#ifndef CKW_PROTOTYPE_INCLUDE_CKW_OPERANDBASE_H +#define CKW_PROTOTYPE_INCLUDE_CKW_OPERANDBASE_H + +#include "ckw/Types.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 new file mode 100644 index 0000000000..16c3f6d441 --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/ScalarValue.h @@ -0,0 +1,137 @@ +/* + * 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 new file mode 100644 index 0000000000..00bb60a444 --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/TensorInfo.h @@ -0,0 +1,145 @@ +/* + * 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.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 TensorComponent : 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 TensorStorage : 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 kw + +#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 new file mode 100644 index 0000000000..2fc5044d1c --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/TensorOperand.h @@ -0,0 +1,181 @@ +/* + * 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.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. + */ + TensorOperand(const ::std::string &name, const TensorInfo &info); + + /** No copy constructor. */ + TensorOperand(const TensorOperand &other) = delete; + + /** No copy assignment. */ + TensorOperand &operator=(const TensorOperand &other) = delete; + + /** (Internal use only) Create the implementation operand. + * + * @param[in] writer The implementation kernel writer. + */ + virtual prototype::Operand create_impl_operand(prototype::IGpuKernelWriter *writer) const override; + + /** Get the tensor info. */ + const TensorInfo &info() const; + + /** Get the tensor info. */ + TensorInfo &info(); + + /** Get the data type. */ + virtual DataType data_type() const override; + + /** Get whether the tensor is compile-time constant. */ + virtual bool is_constant() const override; + + /** Get the default tile attached to the tensor. */ + const TileOperand &tile() const; + + /** Get the default tile attached to the tensor. */ + TileOperand &tile(); + + /** Set the default tile attached to the tensor. */ + TensorOperand &tile(TileOperand &tile); + + /** Get the tensor sampler of the default tile. */ + const TensorTileSampler &tile_sampler() const; + + /** Get the tensor sampler of the default tile. */ + TensorTileSampler &tile_sampler(); + + /** Set the tensor sampler of the default tile. */ + TensorOperand &tile_sampler(const TensorTileSampler &value); + + /** Get the operand that contains the stride in y dimension of the tensor. */ + TileOperand &stride1(); + + /** Get the operand that contains the stride in z dimension of the tensor. */ + TileOperand &stride2(); + + /** Get the operand that contains the stride in w dimension of the tensor. */ + TileOperand &stride3(); + + /** Get the operand that contains the stride in w dimension of the tensor. */ + TileOperand &stride4(); + + /** Get the operand that contains the size of dimension 0 of the tensor. */ + TileOperand &dim0(); + + /** Get the operand that contains the size of dimension 1 of the tensor. */ + TileOperand &dim1(); + + /** Get the operand that contains the size of dimension 2 of the tensor. */ + TileOperand &dim2(); + + /** Get the operand that contains the size of dimension 3 of the tensor. */ + TileOperand &dim3(); + + /** Get the operand that contains the size of dimension 4 of the tensor. */ + TileOperand &dim4(); + + /** Get the operand that contains the size of dimensions 1 and 2 collapsed. */ + TileOperand &dim1_dim2(); + + /** Get the operand that contains the size of dimensions 1, 2 and 3 collapsed. */ + TileOperand &dim1_dim2_dim3(); + + /** Get the operand that contains the offset in bytes to the first element. */ + TileOperand &offset_first_element_in_bytes(); + +private: + TensorInfo _info; + + TileOperand *_tile{ nullptr }; + TensorTileSampler _tile_sampler{}; + + ::std::unique_ptr<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] name The name of the operand. + * @param[in] component The tensor info component. + */ + TensorComponentOperand(const ::std::string &name, TensorComponent component); + + /** (Internal use only) Create the implementation operand. + * + * @param[in] writer The implementation kernel writer. + */ + virtual prototype::Operand create_impl_operand(prototype::IGpuKernelWriter *writer) const override; + +private: + TensorComponent _component; +}; + +} // namespace ckw + +#endif // CKW_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 new file mode 100644 index 0000000000..2ea65bce9e --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/TensorTileSampler.h @@ -0,0 +1,163 @@ +/* + * 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.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 new file mode 100644 index 0000000000..8fba8bb827 --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/TileInfo.h @@ -0,0 +1,84 @@ +/* + * Copyright (c) 2023 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ + +#ifndef CKW_PROTOTYPE_INCLUDE_CKW_TILEINFO_H +#define CKW_PROTOTYPE_INCLUDE_CKW_TILEINFO_H + +#include "ckw/Types.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 new file mode 100644 index 0000000000..c071707a45 --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/TileOperand.h @@ -0,0 +1,110 @@ +/* + * 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; + +/** Tile operand which can be either scalar, vector or 2D tile. */ +class TileOperand : public OperandBase +{ +public: + /** Initialize a new instance of @ref TileOperand class with the tile information. + * + * @param[in] name The name of the tile. + * @param[in] tile_info The tile info. + */ + TileOperand(const ::std::string &name, const TileInfo &tile_info); + + /** Initialize a new instance of @ref TileOperand for scalar variable. + * + * @param[in] name The name of the tile. + * @param[in] data_type The data type of the tile. + */ + TileOperand(const ::std::string &name, DataType data_type); + + /** Initialize a new instance of @ref TileOperand for compile-time constant scalar variable. + * + * @param[in] name The name of the tile. + * @param[in] value The value of the tile. + */ + TileOperand(const ::std::string &name, int32_t value); + + /** Initialize a new instance of @ref TileOperand for compile-time constant scalar variable. + * + * @param[in] name The name of the tile. + * @param[in] value The value of the tile. + */ + TileOperand(const ::std::string &name, float value); + + /** Prohibit copy of tile operand. */ + TileOperand(const TileOperand &) = delete; + + /** Prohibit copy of tile operand. */ + TileOperand &operator=(const TileOperand &) = delete; + + /** (Internal use only) Create the implementation operand. + * + * @param[in] writer The implementation kernel writer. + */ + virtual prototype::Operand create_impl_operand(prototype::IGpuKernelWriter *writer) const override; + + /** Get the tile info. */ + const TileInfo &tile_info() const; + + /** Get the data type of the tile. */ + virtual DataType data_type() const override; + + /** Get whether the tile is compile-time constant. */ + virtual bool is_constant() const override; + + /** Get whether the tile is a scalar value. */ + bool is_scalar() const; + + /** Get the scalar value of the tile. + * + * The tile must have the shape of 1, 1 (i.e. scalar). + */ + ScalarValue scalar_value() const; + +private: + TileInfo _info; + ScalarValue _value{}; + bool _constant; +}; + +} // namespace ckw + +#endif // CKW_PROTOTYPE_INCLUDE_CKW_TILEOPERAND_H diff --git a/compute_kernel_writer/prototype/include/ckw/Types.h b/compute_kernel_writer/prototype/include/ckw/Types.h new file mode 100644 index 0000000000..bb5d7ce077 --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/Types.h @@ -0,0 +1,140 @@ +/* + * 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_TYPES_H +#define CKW_PROTOTYPE_INCLUDE_CKW_TYPES_H + +#include <array> +#include <cstdint> + +namespace ckw +{ + +/** Compute Kernel Writer data types. This data type is used by the code variables and tensor arguments. */ +enum class DataType +{ + Unknown = 0x00, + Fp32 = 0x11, + Fp16 = 0x12, + Int32 = 0x21, + Int16 = 0x22, + Int8 = 0x24, + Uint32 = 0x31, + Uint16 = 0x32, + Uint8 = 0x34, + Bool = 0x41 +}; + +enum class GpuTargetLanguage +{ + Unknown, + OpenCL +}; + +/* Binary operations +*/ +enum class BinaryOp : int32_t +{ + // Elementwise + Add = 0x0000, // + + Sub = 0x0001, // - + Mul = 0x0002, // * + Div = 0x0003, // / + Mod = 0x0004, // % + // Relational + Equal = 0x1000, // == + Less = 0x1001, // < + LessEqual = 0x1002, // <= + Greater = 0x1003, // > + GreaterEqual = 0x1004, // >= + // Algebra + MatMul_Nt_Nt = 0x2000, // X + MatMul_Nt_T = 0x2001, // X + MatMul_T_Nt = 0x2002, // X + MatMul_T_T = 0x2003, // X + Dot = 0x2004, // . + // Logical + LogicalAnd = 0x3000, // && + LogicalOr = 0x3001, // || + LogicalNot = 0x3002 // ! +}; + +enum class AssignmentOp : int32_t +{ + // Unary + Increment = 0x0000, // += + Decrement = 0x0001, // -= +}; + +enum class ScalarUnaryFunction : int32_t +{ + Exp, +}; + +enum class TensorSamplerFormat : int32_t +{ + Unknown = 0, + C_WH_1 = 1, + C_W_H = 2 +}; + +enum class TensorSamplerAddressModeX : int32_t +{ + Unknown = 0, + None = 1, // The user guarantees that the X coordinate is always in-bound + OverlappingMin = 2 // (FIXED shapes only) Reduce the load/store length when x == 0 (MIN). The load length will be width % original length + // Leftover elements can be handled using overlapping. This involves processing some of the elements in the array twice. +}; + +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_PROTOTYPE_INCLUDE_CKW_TYPES_H |