Port operations to CKW prototype
Resolves: COMPMID-6334
Signed-off-by: Nikolaj Jensen <nikolaj.jensen@arm.com>
Change-Id: I500d30f09daec4087eb3e7aecd1de77dc8fd53b4
Signed-off-by: Nikolaj Jensen <nikolaj.jensen@arm.com>
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/9828
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Reviewed-by: Gunes Bayir <gunes.bayir@arm.com>
Reviewed-by: Jakub Sujak <jakub.sujak@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Benchmark: Arm Jenkins <bsgcomp@arm.com>
diff --git a/compute_kernel_writer/prototype/include/ckw/Kernel.h b/compute_kernel_writer/prototype/include/ckw/Kernel.h
index 57a8a40..527206f 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 3b15391..2bf443c 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 @@
// 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 @@
* @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 @@
// 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 @@
/** 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 @@
*
* @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 a9e313f..06d9f82 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 8071588..8eaa6ae 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 7a663f0..3a2509e 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 2ea65bc..e1bf0c5 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 c60880d..de9e47a 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 0000000..2a19850
--- /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 0000000..3447dd6
--- /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 0000000..68146cb
--- /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 0000000..6c08617
--- /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 0000000..78027f1
--- /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
similarity index 72%
rename from compute_kernel_writer/prototype/include/ckw/Types.h
rename to compute_kernel_writer/prototype/include/ckw/types/TensorSamplerTypes.h
index bb5d7ce..836bd13 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 @@
} // namespace ckw
-#endif // CKW_PROTOTYPE_INCLUDE_CKW_TYPES_H
+#endif //CKW_INCLUDE_CKW_TENSORSAMPLERTYPES_H