From 5ff480265a110ea1f2ce24491e082f52348b0f92 Mon Sep 17 00:00:00 2001
From: Nikolaj Jensen <nikolaj.jensen@arm.com>
Date: Tue, 27 Jun 2023 14:13:24 +0100
Subject: 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>
---
 .../prototype/include/ckw/Kernel.h                 |   2 +-
 .../prototype/include/ckw/KernelWriter.h           | 123 ++++++++++++++----
 .../prototype/include/ckw/OperandBase.h            |   2 +-
 .../prototype/include/ckw/TensorInfo.h             |   2 +-
 .../prototype/include/ckw/TensorOperand.h          |   2 +-
 .../prototype/include/ckw/TensorTileSampler.h      |   2 +-
 .../prototype/include/ckw/TileInfo.h               |   2 +-
 .../prototype/include/ckw/Types.h                  | 140 ---------------------
 .../prototype/include/ckw/types/ConvertPolicy.h    |  41 ++++++
 .../prototype/include/ckw/types/DataType.h         |  50 ++++++++
 .../prototype/include/ckw/types/Functions.h        |  61 +++++++++
 .../include/ckw/types/GpuTargetLanguage.h          |  41 ++++++
 .../prototype/include/ckw/types/Operators.h        |  74 +++++++++++
 .../include/ckw/types/TensorSamplerTypes.h         |  78 ++++++++++++
 14 files changed, 451 insertions(+), 169 deletions(-)
 delete mode 100644 compute_kernel_writer/prototype/include/ckw/Types.h
 create mode 100644 compute_kernel_writer/prototype/include/ckw/types/ConvertPolicy.h
 create mode 100644 compute_kernel_writer/prototype/include/ckw/types/DataType.h
 create mode 100644 compute_kernel_writer/prototype/include/ckw/types/Functions.h
 create mode 100644 compute_kernel_writer/prototype/include/ckw/types/GpuTargetLanguage.h
 create mode 100644 compute_kernel_writer/prototype/include/ckw/types/Operators.h
 create mode 100644 compute_kernel_writer/prototype/include/ckw/types/TensorSamplerTypes.h

(limited to 'compute_kernel_writer/prototype/include')

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.h b/compute_kernel_writer/prototype/include/ckw/Types.h
deleted file mode 100644
index bb5d7ce077..0000000000
--- a/compute_kernel_writer/prototype/include/ckw/Types.h
+++ /dev/null
@@ -1,140 +0,0 @@
-/*
- * Copyright (c) 2023 Arm Limited.
- *
- * SPDX-License-Identifier: MIT
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and associated documentation files (the "Software"), to
- * deal in the Software without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- */
-
-#ifndef CKW_PROTOTYPE_INCLUDE_CKW_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
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/TensorSamplerTypes.h b/compute_kernel_writer/prototype/include/ckw/types/TensorSamplerTypes.h
new file mode 100644
index 0000000000..836bd13c95
--- /dev/null
+++ b/compute_kernel_writer/prototype/include/ckw/types/TensorSamplerTypes.h
@@ -0,0 +1,78 @@
+/*
+ * Copyright (c) 2023 Arm Limited.
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ */
+
+#ifndef CKW_INCLUDE_CKW_TENSORSAMPLERTYPES_H
+#define CKW_INCLUDE_CKW_TENSORSAMPLERTYPES_H
+
+#include <cstdint>
+
+namespace ckw
+{
+
+enum class TensorSamplerFormat : int32_t
+{
+    Unknown = 0,
+    C_WH_1  = 1,
+    C_W_H   = 2
+};
+
+enum class TensorSamplerAddressModeX : int32_t
+{
+    Unknown        = 0,
+    None           = 1, // The user guarantees that the X coordinate is always in-bound
+    OverlappingMin = 2  // (FIXED shapes only) Reduce the load/store length when x == 0 (MIN). The load length will be width % original length
+                        // Leftover elements can be handled using overlapping. This involves processing some of the elements in the array twice.
+};
+
+enum class TensorSamplerAddressModeY : int32_t
+{
+    Unknown                  = 0,
+    None                     = 1, // The user guarantees that the Y coordinate is always in-bound
+    OverlappingMin           = 2, // (FIXED shapes only) Reduce the load/store length when x == 0 (MIN). The load length will be width % original length
+    Skip                     = 3, // Skip the read/write
+    SkipMinEdgeOnly          = 4, // Skip greater than or equal to max only. The user guarantees that the Y coordinate is always >= 0
+    SkipMaxEdgeOnly          = 5, // Skip less than 0 only
+    ClampToNearest           = 6, // Clamp the coordinate to nearest edge (0 or max value allowed on Y)
+    ClampToMinEdgeOnly       = 7, // Clamp the negative coordinate to 0 only. Therefore, we expect Y to be always < MAX
+    ClampToMaxEdgeOnly       = 8, // Clamp the coordinate to the max value allowed on Y only. We expect Y to be always >= 0
+    ClampToBorder            = 9, // Clamp to border which always has 0 value
+    ClampToBorderMinEdgeOnly = 10,
+    ClampToBorderMaxEdgeOnly = 11
+};
+
+enum class TensorSamplerAddressModeZ : int32_t
+{
+    Unknown            = 0,
+    None               = 1, // The user guarantees that the Y coordinate is always in-bound
+    Skip               = 3, // Skip the read/write
+    SkipMinEdgeOnly    = 4, // Skip greater than or equal to max only. The user guarantees that the Y coordinate is always >= 0
+    SkipMaxEdgeOnly    = 5, // Skip less than 0 only
+    ClampToNearest     = 6, // Clamp the coordinate to nearest edge (0 or max value allowed on Y)
+    ClampToMinEdgeOnly = 7, // Clamp the negative coordinate to 0 only. Therefore, we expect Y to be always < MAX
+    ClampToMaxEdgeOnly = 8, // Clamp the coordinate to the max value allowed on Y only. We expect Y to be always >= 0
+};
+
+} // namespace ckw
+
+#endif //CKW_INCLUDE_CKW_TENSORSAMPLERTYPES_H
-- 
cgit v1.2.1