Move CKW prototype to separate directory
Partially resolves: COMPMID-6283
Signed-off-by: Viet-Hoa Do <viet-hoa.do@arm.com>
Change-Id: I7596e3dc357d6f0b9cbe66534523943a73c26d81
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/9864
Reviewed-by: SiCong Li <sicong.li@arm.com>
Reviewed-by: Jakub Sujak <jakub.sujak@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Benchmark: Arm Jenkins <bsgcomp@arm.com>
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 0000000..b18944e
--- /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 0000000..57a8a40
--- /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 0000000..a2778a9
--- /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 0000000..f4825e1
--- /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 0000000..16c3f6d
--- /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 0000000..00bb60a
--- /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 0000000..2fc5044
--- /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 0000000..2ea65bc
--- /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 0000000..8fba8bb
--- /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 0000000..c071707
--- /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 0000000..bb5d7ce
--- /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