diff options
Diffstat (limited to 'compute_kernel_writer/prototype/include/ckw')
13 files changed, 376 insertions, 94 deletions
diff --git a/compute_kernel_writer/prototype/include/ckw/Kernel.h b/compute_kernel_writer/prototype/include/ckw/Kernel.h index 57a8a40341..527206feec 100644 --- a/compute_kernel_writer/prototype/include/ckw/Kernel.h +++ b/compute_kernel_writer/prototype/include/ckw/Kernel.h @@ -26,7 +26,7 @@ #define CKW_PROTOTYPE_INCLUDE_CKW_KERNEL_H #include "ckw/OperandBase.h" -#include "ckw/Types.h" +#include "ckw/types/GpuTargetLanguage.h" #include <map> #include <memory> diff --git a/compute_kernel_writer/prototype/include/ckw/KernelWriter.h b/compute_kernel_writer/prototype/include/ckw/KernelWriter.h index 3b1539116a..2bf443cd53 100644 --- a/compute_kernel_writer/prototype/include/ckw/KernelWriter.h +++ b/compute_kernel_writer/prototype/include/ckw/KernelWriter.h @@ -30,6 +30,9 @@ #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> @@ -83,23 +86,23 @@ public: // Tensor and tile declaration // ============================================================================================= - /** Define a tensor argument. + /** Declare 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); + TensorOperand &declare_tensor_argument(const std::string &name, const TensorInfo &info); - /** Define a compile-time constant scalar argument. + /** 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 &create_tile_argument(const char *name, int32_t value); + TileOperand &declare_tile_argument(const std::string &name, int32_t value); /** Declare a new tile. * @@ -111,7 +114,7 @@ public: * @return The @ref TileOperand object. */ template <typename... TArgs> - TileOperand &declare_tile(const char *name, TArgs &&...args) + TileOperand &declare_tile(const std::string &name, TArgs &&...args) { const auto var_name = generate_variable_name(name); auto operand = new TileOperand(var_name, ::std::forward<TArgs>(args)...); @@ -144,29 +147,103 @@ public: // Data processing // ============================================================================================= - /** Write assignment: `<dst> = <src>`. + /** Write assignment: `<dst> = <src>;`. * - * @param[in] dst The destination tile. - * @param[in] src The source tile. + * @param[out] dst The destination tile. + * @param[in] src The source tile. */ - void op_assign(TileOperand &dst, const TileOperand &src); + void op_assign(const TileOperand &dst, const TileOperand &src); - /** Write binary expression: `<dst> = <lhs> <op> <rhs>`. + /** Write the cast: `<dst> = convert_<dst.type><_sat>(<src>);`. * - * @param[in] dst The destination tile. - * @param[in] lhs The LHS operand. - * @param[in] rhs The RHS operand. - * @param[in] op The binary operator. + * @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_binary_expression(TileOperand &dst, const TileOperand &lhs, const TileOperand &rhs, BinaryOp op); + void op_cast_expression(const TileOperand &dst, const TileOperand &src, ConvertPolicy policy); - /** Write function applied to scalar value: `<dst> = <func>(<src>)`. + /** Write the unary expression: `<dst> = <op> <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. + * @param[out] dst The destination tile. + * @param[in] op The unary operator. + * @param[in] src The source tile. */ - void op_scalar_function(TileOperand &dst, const TileOperand &src, ScalarUnaryFunction func); + 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>; <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_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, AssignmentOp update_op, const TileOperand &update_value_name, const std::function<void()> &body); + + /** Write the return statement: `return;` + */ + void op_return(); // ============================================================================================= // Misc @@ -174,8 +251,8 @@ public: /** Set `dst` the global ID of dimension `dim`. * - * @param[in] dst The tile to be written to. - * @param[in] dim The global ID dimension. + * @param[out] dst The tile to be written to. + * @param[in] dim The global ID dimension. */ void op_get_global_id(TileOperand &dst, int32_t dim); @@ -193,7 +270,7 @@ private: * * @return The full variable name. */ - ::std::string generate_variable_name(const char *name) const; + ::std::string generate_variable_name(const std::string &name) const; /** Register the operand to the kernel. * diff --git a/compute_kernel_writer/prototype/include/ckw/OperandBase.h b/compute_kernel_writer/prototype/include/ckw/OperandBase.h index a9e313fc0a..06d9f82756 100644 --- a/compute_kernel_writer/prototype/include/ckw/OperandBase.h +++ b/compute_kernel_writer/prototype/include/ckw/OperandBase.h @@ -25,7 +25,7 @@ #ifndef CKW_PROTOTYPE_INCLUDE_CKW_OPERANDBASE_H #define CKW_PROTOTYPE_INCLUDE_CKW_OPERANDBASE_H -#include "ckw/Types.h" +#include "ckw/types/DataType.h" #include <string> namespace ckw diff --git a/compute_kernel_writer/prototype/include/ckw/TensorInfo.h b/compute_kernel_writer/prototype/include/ckw/TensorInfo.h index 807158896b..8eaa6ae314 100644 --- a/compute_kernel_writer/prototype/include/ckw/TensorInfo.h +++ b/compute_kernel_writer/prototype/include/ckw/TensorInfo.h @@ -25,7 +25,7 @@ #ifndef CKW_PROTOTYPE_INCLUDE_CKW_TENSORINFO_H #define CKW_PROTOTYPE_INCLUDE_CKW_TENSORINFO_H -#include "ckw/Types.h" +#include "ckw/types/DataType.h" #include <array> #include <cstdint> diff --git a/compute_kernel_writer/prototype/include/ckw/TensorOperand.h b/compute_kernel_writer/prototype/include/ckw/TensorOperand.h index 7a663f095b..3a2509e7c8 100644 --- a/compute_kernel_writer/prototype/include/ckw/TensorOperand.h +++ b/compute_kernel_writer/prototype/include/ckw/TensorOperand.h @@ -29,7 +29,7 @@ #include "ckw/TensorInfo.h" #include "ckw/TensorTileSampler.h" #include "ckw/TileOperand.h" -#include "ckw/Types.h" +#include "ckw/types/DataType.h" #include <memory> diff --git a/compute_kernel_writer/prototype/include/ckw/TensorTileSampler.h b/compute_kernel_writer/prototype/include/ckw/TensorTileSampler.h index 2ea65bce9e..e1bf0c52b8 100644 --- a/compute_kernel_writer/prototype/include/ckw/TensorTileSampler.h +++ b/compute_kernel_writer/prototype/include/ckw/TensorTileSampler.h @@ -25,7 +25,7 @@ #ifndef CKW_PROTOTYPE_INCLUDE_CKW_TENSORTILESAMPLER_H #define CKW_PROTOTYPE_INCLUDE_CKW_TENSORTILESAMPLER_H -#include "ckw/Types.h" +#include "ckw/types/TensorSamplerTypes.h" #include <functional> namespace ckw diff --git a/compute_kernel_writer/prototype/include/ckw/TileInfo.h b/compute_kernel_writer/prototype/include/ckw/TileInfo.h index c60880dcd1..de9e47af2b 100644 --- a/compute_kernel_writer/prototype/include/ckw/TileInfo.h +++ b/compute_kernel_writer/prototype/include/ckw/TileInfo.h @@ -25,7 +25,7 @@ #ifndef CKW_PROTOTYPE_INCLUDE_CKW_TILEINFO_H #define CKW_PROTOTYPE_INCLUDE_CKW_TILEINFO_H -#include "ckw/Types.h" +#include "ckw/types/DataType.h" #include <array> #include <cstdint> diff --git a/compute_kernel_writer/prototype/include/ckw/types/ConvertPolicy.h b/compute_kernel_writer/prototype/include/ckw/types/ConvertPolicy.h new file mode 100644 index 0000000000..2a198507eb --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/types/ConvertPolicy.h @@ -0,0 +1,41 @@ +/* + * Copyright (c) 2023 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ + +#ifndef CKW_INCLUDE_CKW_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 new file mode 100644 index 0000000000..3447dd61d6 --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/types/DataType.h @@ -0,0 +1,50 @@ +/* +* Copyright (c) 2023 Arm Limited. +* +* SPDX-License-Identifier: MIT +* +* Permission is hereby granted, free of charge, to any person obtaining a copy +* of this software and associated documentation files (the "Software"), to +* deal in the Software without restriction, including without limitation the +* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or +* sell copies of the Software, and to permit persons to whom the Software is +* furnished to do so, subject to the following conditions: +* +* The above copyright notice and this permission notice shall be included in all +* copies or substantial portions of the Software. +* +* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +* SOFTWARE. +*/ + +#ifndef CKW_INCLUDE_CKW_DATATYPE_H +#define CKW_INCLUDE_CKW_DATATYPE_H + +#include <cstdint> + +namespace ckw +{ + +/** Compute Kernel Writer data types. This data type is used by the code variables and tensor arguments. */ +enum class DataType : int32_t +{ + Unknown = 0x00, + Fp32 = 0x11, + Fp16 = 0x12, + Int32 = 0x21, + Int16 = 0x22, + Int8 = 0x24, + Uint32 = 0x31, + Uint16 = 0x32, + Uint8 = 0x34, + Bool = 0x41 +}; + +} // namespace ckw + +#endif //CKW_INCLUDE_CKW_DATATYPE_H diff --git a/compute_kernel_writer/prototype/include/ckw/types/Functions.h b/compute_kernel_writer/prototype/include/ckw/types/Functions.h new file mode 100644 index 0000000000..68146cb1c8 --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/types/Functions.h @@ -0,0 +1,61 @@ +/* +* Copyright (c) 2023 Arm Limited. +* +* SPDX-License-Identifier: MIT +* +* Permission is hereby granted, free of charge, to any person obtaining a copy +* of this software and associated documentation files (the "Software"), to +* deal in the Software without restriction, including without limitation the +* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or +* sell copies of the Software, and to permit persons to whom the Software is +* furnished to do so, subject to the following conditions: +* +* The above copyright notice and this permission notice shall be included in all +* copies or substantial portions of the Software. +* +* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +* SOFTWARE. +*/ + +#ifndef CKW_INCLUDE_CKW_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, + IsGreaterEqual = 0x0005, + Log = 0x0006, + Round = 0x0007, + + // Misc + SizeOf = 0x0008, +}; + +enum class BinaryFunction : int32_t +{ + Min = 0x0000, + Max = 0x0001, +}; + +enum class TernaryFunction : int32_t +{ + Select = 0x0000, +}; + +} // 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 new file mode 100644 index 0000000000..6c08617949 --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/types/GpuTargetLanguage.h @@ -0,0 +1,41 @@ +/* + * Copyright (c) 2023 Arm Limited. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + */ + +#ifndef CKW_INCLUDE_CKW_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 new file mode 100644 index 0000000000..78027f1ed5 --- /dev/null +++ b/compute_kernel_writer/prototype/include/ckw/types/Operators.h @@ -0,0 +1,74 @@ +/* +* 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, // ! +}; + +/* 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, // || +}; + +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.h b/compute_kernel_writer/prototype/include/ckw/types/TensorSamplerTypes.h index bb5d7ce077..836bd13c95 100644 --- a/compute_kernel_writer/prototype/include/ckw/Types.h +++ b/compute_kernel_writer/prototype/include/ckw/types/TensorSamplerTypes.h @@ -22,76 +22,14 @@ * SOFTWARE. */ -#ifndef CKW_PROTOTYPE_INCLUDE_CKW_TYPES_H -#define CKW_PROTOTYPE_INCLUDE_CKW_TYPES_H +#ifndef CKW_INCLUDE_CKW_TENSORSAMPLERTYPES_H +#define CKW_INCLUDE_CKW_TENSORSAMPLERTYPES_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, @@ -137,4 +75,4 @@ enum class TensorSamplerAddressModeZ : int32_t } // namespace ckw -#endif // CKW_PROTOTYPE_INCLUDE_CKW_TYPES_H +#endif //CKW_INCLUDE_CKW_TENSORSAMPLERTYPES_H |