diff options
Diffstat (limited to 'compute_kernel_writer/src')
41 files changed, 5564 insertions, 0 deletions
diff --git a/compute_kernel_writer/src/Error.cpp b/compute_kernel_writer/src/Error.cpp new file mode 100644 index 0000000000..e1e4bffcec --- /dev/null +++ b/compute_kernel_writer/src/Error.cpp @@ -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. + */ + +#include "ckw/Error.h" + +#include <string> + +namespace ckw +{ +std::string +create_error_msg(const std::string &file, const std::string &func, const std::string &line, const std::string &msg) +{ + std::string err; + err += "[COMPUTE_KERNEL_WRITER][ERROR]:"; + err += " " + file + ":" + line; + err += " " + func; + err += " " + msg; + return err; +} +} // namespace ckw diff --git a/compute_kernel_writer/src/Helpers.cpp b/compute_kernel_writer/src/Helpers.cpp new file mode 100644 index 0000000000..82d4c4e917 --- /dev/null +++ b/compute_kernel_writer/src/Helpers.cpp @@ -0,0 +1,63 @@ +/* + * 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. + */ + +#include "src/Helpers.h" + +#include "ckw/Error.h" + +namespace ckw +{ +std::string dec_to_hex_as_string(int32_t dec) +{ + switch (dec) + { + case 0: + case 1: + case 2: + case 3: + case 4: + case 5: + case 6: + case 7: + case 8: + case 9: + return std::to_string(dec); + case 10: + return "A"; + case 11: + return "B"; + case 12: + return "C"; + case 13: + return "D"; + case 14: + return "E"; + case 15: + return "F"; + default: + COMPUTE_KERNEL_WRITER_ERROR_ON_MSG("Unsupported decimal number"); + return ""; + } +} +} // namespace ckw diff --git a/compute_kernel_writer/src/Helpers.h b/compute_kernel_writer/src/Helpers.h new file mode 100644 index 0000000000..16c06d60e7 --- /dev/null +++ b/compute_kernel_writer/src/Helpers.h @@ -0,0 +1,56 @@ +/* + * 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 COMPUTE_KERNEL_WRITER_SRC_HELPERS_H +#define COMPUTE_KERNEL_WRITER_SRC_HELPERS_H + +#include <cstdint> +#include <string> + +/** Generic helper functions */ +namespace ckw +{ +/** Helper function to convert a decimal number passed as int32_t variable to hexadecimal number as string + * + * @param[in] dec Decimal number. It must be >= 0 and < 16 + * + * @return the OpenCL datatype as a string + */ +std::string dec_to_hex_as_string(int32_t dec); + +/** Helper function to clamp a value between min_val and max_val + * + * @param[in] val Value to clamp + * @param[in] min_val Lower value + * @param[in] max_val Upper value + * + * @return the clamped value + */ +template <typename T> +T clamp(const T &val, const T &min_val, const T &max_val) +{ + return std::max(min_val, std::min(val, max_val)); +} +} // namespace ckw +#endif /* COMPUTE_KERNEL_WRITER_SRC_HELPERS_H */ diff --git a/compute_kernel_writer/src/ITensor.h b/compute_kernel_writer/src/ITensor.h new file mode 100644 index 0000000000..4c1c56fd35 --- /dev/null +++ b/compute_kernel_writer/src/ITensor.h @@ -0,0 +1,46 @@ +/* + * 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_SRC_ITENSOR_H +#define CKW_SRC_ITENSOR_H + +#include "src/ITensorArgument.h" + +namespace ckw +{ + +/** The generic class for all tensor objects in CKW. + * + * Tensors in CKW are always kernel arguments consisting of: + * - Essential information such as name, tensor info, etc. + * - Tensor storage access: allowing load/store operation to perform. + * - Tensor component access: allowing interaction with tensor information such as shape, strides, etc. in the form of tile objects. + */ +class ITensor : public ITensorArgument, public ITensorStorageAccess, public ITensorComponentAccess +{ +}; + +} // namespace ckw + +#endif // CKW_SRC_ITENSOR_H
\ No newline at end of file diff --git a/compute_kernel_writer/src/ITensorArgument.h b/compute_kernel_writer/src/ITensorArgument.h new file mode 100644 index 0000000000..ece45a4dc4 --- /dev/null +++ b/compute_kernel_writer/src/ITensorArgument.h @@ -0,0 +1,135 @@ +/* + * 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_SRC_ITENSORARGUMENT_H +#define CKW_SRC_ITENSORARGUMENT_H + +#include "ckw/TensorInfo.h" +#include "ckw/types/TensorComponentType.h" +#include "ckw/types/TensorStorageType.h" + +#include "src/ITile.h" + +#include <string> +#include <vector> + +namespace ckw +{ + +class ITensorComponent; + +/** Tensor storage variable */ +struct TensorStorageVariable +{ + std::string val{""}; /** Tensor storage as a string */ + TensorStorageType type{TensorStorageType::Unknown}; /** Tensor storage type */ +}; + +/** Tensor argument base class. + * A tensor is a multidimensional array used to store data. To access an element (or multiple elements) from a tensor, + * the following information are required: + * -# The data memory object. For example, the pointer to the array + * -# The tensor components, such as the size of each tensor dimension, or the number of elements in bytes contained in each dimension (also known as the "stride") + */ +class ITensorArgument +{ +public: + virtual ~ITensorArgument() = default; + /** Method to get the name of the tensor argument. + * + * @return the name of the tensor argument + */ + std::string name() const + { + return _basename; + } + + /** Method to get the tensor info + * + * @return the @ref TensorInfo + */ + TensorInfo &info() + { + return _info; + } + + /** Method to get the tensor info + * + * @return the @ref TensorInfo + */ + const TensorInfo &info() const + { + return _info; + } + +protected: + TensorInfo _info{}; // Tensor info + std::string _basename{""}; // Tensor name +}; + +/** Tensor component argument base class */ +class ITensorComponentAccess +{ +public: + virtual ~ITensorComponentAccess() = default; + /** Method to get the tensor component variable as a tile. + * + * @param[in] x The tensor component to query + * + * @return the tensor component variable as a @ref ITile. + */ + virtual ITile &component(TensorComponentType x) = 0; + /** Method to get all tensor components needed to access the data in the tensor + * + * The tensor components returned by this method must be all passed as kernel argument + * + * @return a vector containing all the tensor components as pointers to @ref ITensorComponent objects. + */ + virtual std::vector<const ITensorComponent *> components() const = 0; +}; + +/** Tensor storage argument base class */ +class ITensorStorageAccess +{ +public: + virtual ~ITensorStorageAccess() = default; + /** Method to get the tensor storage as a string + * + * @param[in] x The tensor storage to query + * + * @return the tensor storage as a @ref TensorStorageVariable + */ + virtual TensorStorageVariable &storage(TensorStorageType x) = 0; + /** Method to get all tensor storages needed to access the data in the tensor + * + * The tensor storages returned by this method must be all passed as kernel argument + * + * @return a vector containing all the tensor storages as @ref TensorStorageVariable objects + */ + virtual std::vector<TensorStorageVariable> storages() const = 0; +}; + +} // namespace ckw + +#endif // CKW_SRC_ITENSORARGUMENT_H diff --git a/compute_kernel_writer/src/ITensorComponent.h b/compute_kernel_writer/src/ITensorComponent.h new file mode 100644 index 0000000000..f9c9d8fd81 --- /dev/null +++ b/compute_kernel_writer/src/ITensorComponent.h @@ -0,0 +1,54 @@ +/* + * 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_SRC_ITENSORCOMPONENT_H +#define CKW_SRC_ITENSORCOMPONENT_H + +#include "ckw/types/TensorComponentType.h" + +#include "src/ITile.h" + +namespace ckw +{ + +/** A tensor component provides access to tensor information such as shape, strides, etc. in the form of @ref ITile objects. */ +class ITensorComponent +{ +public: + /** Destructor. */ + virtual ~ITensorComponent() = default; + + /** Get the tile variable for the component. */ + virtual ITile &tile() = 0; + + /** Get the const tile variable for the component. */ + virtual const ITile &tile() const = 0; + + /** Get the component type. */ + virtual TensorComponentType component_type() const = 0; +}; + +} // namespace ckw + +#endif // CKW_SRC_ITENSORCOMPONENT_H diff --git a/compute_kernel_writer/src/ITile.cpp b/compute_kernel_writer/src/ITile.cpp new file mode 100644 index 0000000000..eeb7816068 --- /dev/null +++ b/compute_kernel_writer/src/ITile.cpp @@ -0,0 +1,35 @@ +/* + * 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. + */ + +#include "src/ITile.h" + +namespace ckw +{ + +bool ITile::is_scalar() const +{ + return info().width() == 1 && info().height() == 1; +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/ITile.h b/compute_kernel_writer/src/ITile.h new file mode 100644 index 0000000000..8eaac5ac12 --- /dev/null +++ b/compute_kernel_writer/src/ITile.h @@ -0,0 +1,141 @@ +/* + * 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_SRC_ITILE_H +#define CKW_SRC_ITILE_H + +#include "ckw/TileInfo.h" + +#include <string> +#include <vector> + +namespace ckw +{ +/** Compute Kernel Writer tile container. It contains the variables stored in the tile as a string */ +using TileContainer = std::vector<std::vector<std::string>>; + +/** Tile descriptor which reports the underlying datatype and vector length */ +struct TileVariableDescriptor +{ + DataType dt{DataType::Unknown}; /** Data type */ + int32_t len{1}; /** Number of elements in a single variable. For example, 1 for scalar */ +}; + +/** Tile variable */ +struct TileVariable +{ + std::string str{""}; /** Tile variable as a string */ + TileVariableDescriptor desc{}; /** Tile value descriptor which reports the datatype and vector length */ +}; + +/** Interface to provide support for scalar access for a Tile. + */ +class IScalarAccess +{ +public: + virtual ~IScalarAccess() = default; + + /** Method to get the scalar variable from a tile as a string + * @param[in] row Tile row. If out-of-bound, the row is clamped to the nearest valid edge + * @param[in] col Tile column. If out-of-bound, the column is clamped to the nearest valid edge + * + * @return the @ref TileVariable + */ + virtual TileVariable scalar(int32_t row, int32_t col) const = 0; +}; + +/** Interface to provide support for vector access for a tile. + */ +class IVectorAccess +{ +public: + virtual ~IVectorAccess() = default; + + /** Method to get the vector variable from a tile. + * The user can query the list of supported vector lengths through the supported_vector_lengths() method. + * + * @param[in] row Tile row. If out-of-bound, the row is clamped to the nearest valid edge + * + * @return the vector variable as a @ref TileVariable + */ + virtual TileVariable vector(int32_t row) const = 0; + + /** Method to get a sub-vector variable. The length of the sub-vector must be supported by the derived IVectorAccess class + * + * @param[in] row Tile row. If out-of-bound, the row is clamped to the nearest valid edge + * @param[in] col_start Tile starting column to get the sub-vector. If out-of-bound, the derived IVectorAccess class may throw an assert. + * @param[in] width The width of the sub-vector. The width must be supported by the derived IVectorAccess class and the last element must be in-bound. + * + * @return the vector variable as a @ref TileVariable + */ + virtual TileVariable vector(int32_t row, int32_t col_start, int32_t width) const = 0; + + /** Method to get the supported vector length. + * + * @return a vector containing the supported vector lengths + */ + virtual std::vector<int32_t> supported_vector_lengths() const = 0; +}; + +/** Tile base class. + * A Tile is a collection of variables (either program variables or constants) used to express a 2D data. + */ +class ITile : public IScalarAccess +{ +public: + virtual ~ITile() = default; + + /** Method to get all TileVariable objects + * + * @return a vector containing all @ref TileVariable objects + */ + virtual std::vector<TileVariable> all() const = 0; + + /** Method to get the name of the tile. + * + * @return the name of the tile + */ + virtual const std::string &name() const = 0; + + /** Method to get the tile info + * + * @return the @ref TileInfo + */ + virtual const TileInfo &info() const = 0; + + /** Method to know whether the tile is assignable or not. + * For example, a constant tile is not assignable. + * + * @return true if the tile is assignable + */ + virtual bool is_assignable() const = 0; + + /** Get whether the tile is scalar, i.e. the width and height are both 1. + * + * @return true if the tile is scalar. + */ + bool is_scalar() const; +}; +} // namespace ckw + +#endif // CKW_SRC_ITILE_H diff --git a/compute_kernel_writer/src/Kernel.cpp b/compute_kernel_writer/src/Kernel.cpp new file mode 100644 index 0000000000..12389b3816 --- /dev/null +++ b/compute_kernel_writer/src/Kernel.cpp @@ -0,0 +1,54 @@ +/* + * 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. + */ + +#include "ckw/Kernel.h" + +#include "ckw/types/TargetLanguage.h" + +namespace ckw +{ + +Kernel::~Kernel() = default; + +Kernel::Kernel(TargetLanguage language, const std::vector<KernelArgument> &arguments, const std::string &source_code) + : _language(language), _arguments(arguments), _source_code(source_code) +{ +} + +TargetLanguage Kernel::target_language() const +{ + return _language; +} + +const std::vector<KernelArgument> &Kernel::arguments() const +{ + return _arguments; +} + +const std::string &Kernel::source_code() const +{ + return _source_code; +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/KernelArgument.cpp b/compute_kernel_writer/src/KernelArgument.cpp new file mode 100644 index 0000000000..a640d36507 --- /dev/null +++ b/compute_kernel_writer/src/KernelArgument.cpp @@ -0,0 +1,68 @@ +/* + * 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. + */ + +#include "ckw/KernelArgument.h" + +#include "ckw/Error.h" + +namespace ckw +{ + +KernelArgument::KernelArgument(int32_t tensor_id, TensorStorageType storage_type) + : _type(Type::TensorStorage), _id(tensor_id) +{ + _sub_id.tensor_storage_type = storage_type; +} + +KernelArgument::KernelArgument(int32_t tensor_id, TensorComponentType component_type) + : _type(Type::TensorComponent), _id(tensor_id) +{ + _sub_id.tensor_component_type = component_type; +} + +KernelArgument::Type KernelArgument::type() const +{ + return _type; +} + +int32_t KernelArgument::id() const +{ + return _id; +} + +TensorStorageType KernelArgument::tensor_storage_type() const +{ + CKW_ASSERT(_type == Type::TensorStorage); + + return _sub_id.tensor_storage_type; +} + +TensorComponentType KernelArgument::tensor_component_type() const +{ + CKW_ASSERT(_type == Type::TensorComponent); + + return _sub_id.tensor_component_type; +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/KernelWriter.cpp b/compute_kernel_writer/src/KernelWriter.cpp new file mode 100644 index 0000000000..92a36746ce --- /dev/null +++ b/compute_kernel_writer/src/KernelWriter.cpp @@ -0,0 +1,124 @@ +/* + * 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. + */ + +#include "ckw/KernelWriter.h" + +#include "ckw/Error.h" +#include "ckw/TileOperand.h" +#include "ckw/types/TargetArchitecture.h" +#include "ckw/types/TargetLanguage.h" + +#include "src/cl/CLKernelWriter.h" +#include "src/cl/CLTensorArgument.h" +#include "src/cl/CLTile.h" +#include "src/TileView.h" + +#include <tuple> + +namespace ckw +{ + +KernelWriter::~KernelWriter() = default; + +std::unique_ptr<KernelWriter> KernelWriter::create_instance(TargetArchitecture architecture, TargetLanguage language) +{ + CKW_UNUSED(architecture); + switch (language) + { + case TargetLanguage::OpenCL: + // Currently this is the oldest and the only supported GPU architecture. + CKW_ASSERT(architecture == TargetArchitecture::GpuArmMaliValhall); + return std::make_unique<CLKernelWriter>(); + + default: + CKW_THROW_MSG("Language not supported!"); + } +} + +int32_t KernelWriter::new_id_space() +{ + _id_space = ++_last_created_id_space; + + return _id_space; +} + +int32_t KernelWriter::id_space() const +{ + return _id_space; +} + +KernelWriter &KernelWriter::id_space(int32_t value) +{ + CKW_ASSERT(value <= _last_created_id_space); + + _id_space = value; + + return *this; +} + +void KernelWriter::write_body(const std::function<void()> &body) +{ + const auto curr_id_space = id_space(); + new_id_space(); + body(); + id_space(curr_id_space); +} + +std::string KernelWriter::generate_full_name(const std::string &name) const +{ + return "G" + std::to_string(id_space()) + "__" + name; +} + +TileOperand KernelWriter::create_tile_operand(ITile &tile) +{ + return TileOperand(tile); +} + +std::tuple<ITile &, TileArea> KernelWriter::get_tile(const TileOperand &operand) +{ + return {*operand._tile, {operand._row_start, operand._row_end, operand._col_start, operand._col_end}}; +} + +TensorOperand KernelWriter::create_tensor_operand(ITensor &tensor) +{ + return TensorOperand(tensor); +} + +ITensor &KernelWriter::get_tensor(const TensorOperand &operand) +{ + CKW_ASSERT(operand._tensor != nullptr); + return *operand._tensor; +} + +const std::vector<std::vector<std::string>> &KernelWriter::get_values(const ConstantData &data) +{ + return data.values(); +} + +DataType KernelWriter::get_data_type(const ConstantData &data) +{ + return data.data_type(); +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/Tensor3dMapper.cpp b/compute_kernel_writer/src/Tensor3dMapper.cpp new file mode 100644 index 0000000000..acef6412a4 --- /dev/null +++ b/compute_kernel_writer/src/Tensor3dMapper.cpp @@ -0,0 +1,155 @@ +/* + * 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. + */ + +#include "Tensor3dMapper.h" + +#include "ckw/Error.h" +#include "ckw/types/TensorSamplerTypes.h" + +#include "src/ITensor.h" +#include "src/ITile.h" + +namespace ckw +{ +Tensor3dMapper::Tensor3dMapper(ITensor *tensor, TensorSamplerFormat format) : _tensor(tensor), _format(format) +{ +} + +TileVariable Tensor3dMapper::dim_x() const +{ + switch (_format) + { + case TensorSamplerFormat::Dim0_Dim1xDim2_1: + case TensorSamplerFormat::Dim0_Dim1_Dim2: + return _tensor->component(TensorComponentType::Dim0).scalar(0, 0); + default: + CKW_THROW_MSG("Unsupported tensor format"); + return _tensor->component(TensorComponentType::Unknown).scalar(0, 0); + } +} + +TileVariable Tensor3dMapper::dim_y() const +{ + switch (_format) + { + case TensorSamplerFormat::Dim0_Dim1xDim2_1: + return _tensor->component(TensorComponentType::Dim1xDim2).scalar(0, 0); + case TensorSamplerFormat::Dim0_Dim1_Dim2: + return _tensor->component(TensorComponentType::Dim1).scalar(0, 0); + default: + CKW_THROW_MSG("Unsupported tensor format"); + return _tensor->component(TensorComponentType::Unknown).scalar(0, 0); + } +} + +TileVariable Tensor3dMapper::dim_z() const +{ + TileVariable dim_one; + + switch (_format) + { + case TensorSamplerFormat::Dim0_Dim1xDim2_1: + dim_one = _tensor->component(TensorComponentType::Dim3).scalar(0, 0); + dim_one.str = "1"; + return dim_one; + case TensorSamplerFormat::Dim0_Dim1_Dim2: + return _tensor->component(TensorComponentType::Dim2).scalar(0, 0); + default: + CKW_THROW_MSG("Unsupported tensor format"); + return _tensor->component(TensorComponentType::Unknown).scalar(0, 0); + } +} + +TileVariable Tensor3dMapper::dim_batch() const +{ + TileVariable dim_one; + + switch (_format) + { + case TensorSamplerFormat::Dim0_Dim1xDim2_1: + case TensorSamplerFormat::Dim0_Dim1_Dim2: + return _tensor->component(TensorComponentType::Dim3).scalar(0, 0); + default: + CKW_THROW_MSG("Unsupported tensor format"); + return _tensor->component(TensorComponentType::Unknown).scalar(0, 0); + } +} + +TileVariable Tensor3dMapper::stride_x() const +{ + switch (_format) + { + case TensorSamplerFormat::Dim0_Dim1xDim2_1: + case TensorSamplerFormat::Dim0_Dim1_Dim2: + return _tensor->component(TensorComponentType::Stride0).scalar(0, 0); + default: + CKW_THROW_MSG("Unsupported tensor format"); + return _tensor->component(TensorComponentType::Unknown).scalar(0, 0); + } +} + +TileVariable Tensor3dMapper::stride_y() const +{ + switch (_format) + { + case TensorSamplerFormat::Dim0_Dim1xDim2_1: + case TensorSamplerFormat::Dim0_Dim1_Dim2: + return _tensor->component(TensorComponentType::Stride1).scalar(0, 0); + default: + CKW_THROW_MSG("Unsupported tensor format"); + return _tensor->component(TensorComponentType::Unknown).scalar(0, 0); + } +} + +TileVariable Tensor3dMapper::stride_z() const +{ + TileVariable stride_zero; + + switch (_format) + { + case TensorSamplerFormat::Dim0_Dim1xDim2_1: + stride_zero = _tensor->component(TensorComponentType::Stride3).scalar(0, 0); + stride_zero.str = "0"; + return stride_zero; + case TensorSamplerFormat::Dim0_Dim1_Dim2: + return _tensor->component(TensorComponentType::Stride2).scalar(0, 0); + default: + CKW_THROW_MSG("Unsupported tensor format"); + return _tensor->component(TensorComponentType::Unknown).scalar(0, 0); + } +} + +TileVariable Tensor3dMapper::stride_batch() const +{ + switch (_format) + { + case TensorSamplerFormat::Dim0_Dim1xDim2_1: + case TensorSamplerFormat::Dim0_Dim1_Dim2: + return _tensor->component(TensorComponentType::Stride3).scalar(0, 0); + default: + CKW_THROW_MSG("Unsupported tensor format"); + return _tensor->component(TensorComponentType::Unknown).scalar(0, 0); + } +} +} // namespace ckw diff --git a/compute_kernel_writer/src/Tensor3dMapper.h b/compute_kernel_writer/src/Tensor3dMapper.h new file mode 100644 index 0000000000..e94b595193 --- /dev/null +++ b/compute_kernel_writer/src/Tensor3dMapper.h @@ -0,0 +1,82 @@ +/* + * 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_SRC_TENSOR3DMAPPER_H +#define CKW_SRC_TENSOR3DMAPPER_H + +#include <string> + +namespace ckw +{ +// Forward declarations +class ITensor; +enum class TensorSamplerFormat; +struct TileVariable; + +/** This internal-only class is responsible to map an Nd tensor spatial dimensions to a 3d tensor spatial dimensions with the + * help of TensorSamplerFormat. + * Attention: The batch is not considered as a spatial dimension and it is treated as an offset + * + * The aim of the dimensionality reduction is primarily to reduce + * the address calculation to: + * x + y * stride_y + z * stride_z + offset, where offset is determined by the batch (for example, b * stride_batch). + * + */ +class Tensor3dMapper +{ +public: + /** Constructor */ + Tensor3dMapper(ITensor *tensor, TensorSamplerFormat format); + + /** Get dimension x as string */ + TileVariable dim_x() const; + + /** Get dimension y as string */ + TileVariable dim_y() const; + + /** Get dimension z as string */ + TileVariable dim_z() const; + + /** Get batch dimension as string */ + TileVariable dim_batch() const; + + /** Get stride for dimension x as string */ + TileVariable stride_x() const; + + /** Get stride for dimension y as string */ + TileVariable stride_y() const; + + /** Get stride for dimension z as string */ + TileVariable stride_z() const; + + /** Get stride for batch dimension as string */ + TileVariable stride_batch() const; + +private: + ITensor *_tensor; + TensorSamplerFormat _format; +}; +} // namespace ckw + +#endif /* CKW_SRC_TENSOR3DMAPPER_H */ diff --git a/compute_kernel_writer/src/TensorInfo.cpp b/compute_kernel_writer/src/TensorInfo.cpp new file mode 100644 index 0000000000..561c126469 --- /dev/null +++ b/compute_kernel_writer/src/TensorInfo.cpp @@ -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. + */ + +#include "ckw/TensorInfo.h" + +namespace ckw +{ +TensorInfo::TensorInfo(DataType dt, const TensorShape &shape, TensorDataLayout dl, int32_t id) + : _shape(shape), _dt(dt), _dl(dl), _id(id) +{ +} + +TensorInfo &TensorInfo::shape(const TensorShape &shape) +{ + _shape = shape; + return *this; +} + +TensorShape TensorInfo::shape() const +{ + return _shape; +} + +TensorInfo &TensorInfo::data_type(DataType dt) +{ + _dt = dt; + return *this; +} + +DataType TensorInfo::data_type() const +{ + return _dt; +} + +TensorInfo &TensorInfo::data_layout(TensorDataLayout dl) +{ + _dl = dl; + return *this; +} + +TensorDataLayout TensorInfo::data_layout() const +{ + return _dl; +} + +TensorInfo &TensorInfo::id(int32_t id) +{ + _id = id; + return *this; +} + +int32_t TensorInfo::id() const +{ + return _id; +} +} // namespace ckw diff --git a/compute_kernel_writer/src/TensorOperand.cpp b/compute_kernel_writer/src/TensorOperand.cpp new file mode 100644 index 0000000000..94997537d8 --- /dev/null +++ b/compute_kernel_writer/src/TensorOperand.cpp @@ -0,0 +1,135 @@ +/* + * 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. + */ +#include "ckw/TensorOperand.h" + +#include "ckw/Error.h" + +#include "src/ITensor.h" + +namespace ckw +{ + +TensorOperand::TensorOperand() : _tensor(nullptr) +{ +} + +TensorOperand::TensorOperand(ITensor &tensor) : _tensor(&tensor) +{ +} + +bool TensorOperand::is_valid() const +{ + return _tensor != nullptr; +} + +const TensorInfo &TensorOperand::info() const +{ + CKW_ASSERT(is_valid() == true); + return _tensor->info(); +} + +TileOperand TensorOperand::stride0() +{ + CKW_ASSERT(is_valid() == true); + return TileOperand(_tensor->component(TensorComponentType::Stride0)); +} + +TileOperand TensorOperand::stride1() +{ + CKW_ASSERT(is_valid() == true); + return TileOperand(_tensor->component(TensorComponentType::Stride1)); +} + +TileOperand TensorOperand::stride2() +{ + CKW_ASSERT(is_valid() == true); + return TileOperand(_tensor->component(TensorComponentType::Stride2)); +} + +TileOperand TensorOperand::stride3() +{ + CKW_ASSERT(is_valid() == true); + return TileOperand(_tensor->component(TensorComponentType::Stride3)); +} + +TileOperand TensorOperand::stride4() +{ + CKW_ASSERT(is_valid() == true); + return TileOperand(_tensor->component(TensorComponentType::Stride4)); +} + +TileOperand TensorOperand::dim0() +{ + return TileOperand(_tensor->component(TensorComponentType::Dim0)); +} + +TileOperand TensorOperand::dim1() +{ + CKW_ASSERT(is_valid() == true); + return TileOperand(_tensor->component(TensorComponentType::Dim1)); +} + +TileOperand TensorOperand::dim2() +{ + CKW_ASSERT(is_valid() == true); + return TileOperand(_tensor->component(TensorComponentType::Dim2)); +} + +TileOperand TensorOperand::dim3() +{ + CKW_ASSERT(is_valid() == true); + return TileOperand(_tensor->component(TensorComponentType::Dim3)); +} + +TileOperand TensorOperand::dim4() +{ + CKW_ASSERT(is_valid() == true); + return TileOperand(_tensor->component(TensorComponentType::Dim4)); +} + +TileOperand TensorOperand::dim1_dim2() +{ + CKW_ASSERT(is_valid() == true); + return TileOperand(_tensor->component(TensorComponentType::Dim1xDim2)); +} + +TileOperand TensorOperand::dim1_dim2_dim3() +{ + CKW_ASSERT(is_valid() == true); + return TileOperand(_tensor->component(TensorComponentType::Dim1xDim2xDim3)); +} + +TileOperand TensorOperand::dim2_dim3() +{ + CKW_ASSERT(is_valid() == true); + return TileOperand(_tensor->component(TensorComponentType::Dim2xDim3)); +} + +TileOperand TensorOperand::offset_first_element_in_bytes() +{ + CKW_ASSERT(is_valid() == true); + return TileOperand(_tensor->component(TensorComponentType::OffsetFirstElement)); +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/TensorSampler.cpp b/compute_kernel_writer/src/TensorSampler.cpp new file mode 100644 index 0000000000..e81c5f9d66 --- /dev/null +++ b/compute_kernel_writer/src/TensorSampler.cpp @@ -0,0 +1,108 @@ +/* + * 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. + */ + +#include "ckw/TensorSampler.h" + +namespace ckw +{ + +TensorSampler::TensorSampler() = default; + +TensorSampler::TensorSampler(TensorStorageType storage, + TensorSamplerFormat format, + TensorSamplerAddressModeX address_mode_x, + TensorSamplerAddressModeY address_mode_y, + TensorSamplerAddressModeZ address_mode_z) + : _storage(storage), + _format(format), + _address_mode_x(address_mode_x), + _address_mode_y(address_mode_y), + _address_mode_z(address_mode_z) +{ +} + +TensorStorageType TensorSampler::storage() const +{ + return _storage; +} + +TensorSampler &TensorSampler::storage(TensorStorageType storage) +{ + _storage = storage; + return *this; +} + +/** Get the format of the tensor. */ +TensorSamplerFormat TensorSampler::format() const +{ + return _format; +} + +/** Set the format of the tensor. */ +TensorSampler &TensorSampler::format(TensorSamplerFormat format) +{ + _format = format; + return *this; +} + +/** Get the address mode of the x dimension. */ +TensorSamplerAddressModeX TensorSampler::address_mode_x() const +{ + return _address_mode_x; +} + +/** Set the address mode of the x-dimension. */ +TensorSampler &TensorSampler::address_mode_x(TensorSamplerAddressModeX address_mode_x) +{ + _address_mode_x = address_mode_x; + return *this; +} + +/** Get the address mode of the y dimension. */ +TensorSamplerAddressModeY TensorSampler::address_mode_y() const +{ + return _address_mode_y; +} + +/** Set the address mode of the y dimension. */ +TensorSampler &TensorSampler::address_mode_y(TensorSamplerAddressModeY address_mode_y) +{ + _address_mode_y = address_mode_y; + return *this; +} + +/** Get the address mode of the z dimension. */ +TensorSamplerAddressModeZ TensorSampler::address_mode_z() const +{ + return _address_mode_z; +} + +/** Set the address mode of the z dimension. */ +TensorSampler &TensorSampler::address_mode_z(TensorSamplerAddressModeZ address_mode_z) +{ + _address_mode_z = address_mode_z; + return *this; +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/TensorUtils.cpp b/compute_kernel_writer/src/TensorUtils.cpp new file mode 100644 index 0000000000..17fc9547ae --- /dev/null +++ b/compute_kernel_writer/src/TensorUtils.cpp @@ -0,0 +1,116 @@ +/* + * 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. + */ + +#include "src/TensorUtils.h" + +#include "ckw/Error.h" +#include "ckw/TensorInfo.h" +#include "ckw/types/TensorComponentType.h" + +namespace ckw +{ +TensorComponentType get_tensor_dimension(TensorDataLayout layout, TensorDataLayoutComponent component) +{ + switch (layout) + { + case TensorDataLayout::Nhwc: + switch (component) + { + case TensorDataLayoutComponent::C: + return TensorComponentType::Dim0; + case TensorDataLayoutComponent::W: + return TensorComponentType::Dim1; + case TensorDataLayoutComponent::H: + return TensorComponentType::Dim2; + case TensorDataLayoutComponent::N: + return TensorComponentType::Dim3; + default: + COMPUTE_KERNEL_WRITER_ERROR_ON_MSG("Unsupported tensor component for NHWC"); + return TensorComponentType::Unknown; + } + case TensorDataLayout::Ndhwc: + switch (component) + { + case TensorDataLayoutComponent::C: + return TensorComponentType::Dim0; + case TensorDataLayoutComponent::W: + return TensorComponentType::Dim1; + case TensorDataLayoutComponent::H: + return TensorComponentType::Dim2; + case TensorDataLayoutComponent::D: + return TensorComponentType::Dim3; + case TensorDataLayoutComponent::N: + return TensorComponentType::Dim4; + default: + COMPUTE_KERNEL_WRITER_ERROR_ON_MSG("Unsupported tensor component for NDHWC"); + return TensorComponentType::Unknown; + } + default: + COMPUTE_KERNEL_WRITER_ERROR_ON_MSG("Unsupported tensor data layout"); + return TensorComponentType::Unknown; + } +} + +TensorComponentType get_tensor_stride(TensorDataLayout layout, TensorDataLayoutComponent component) +{ + switch (layout) + { + case TensorDataLayout::Nhwc: + switch (component) + { + case TensorDataLayoutComponent::C: + return TensorComponentType::Stride0; + case TensorDataLayoutComponent::W: + return TensorComponentType::Stride1; + case TensorDataLayoutComponent::H: + return TensorComponentType::Stride2; + case TensorDataLayoutComponent::N: + return TensorComponentType::Stride3; + default: + COMPUTE_KERNEL_WRITER_ERROR_ON_MSG("Unsupported tensor component for NHWC"); + return TensorComponentType::Unknown; + } + case TensorDataLayout::Ndhwc: + switch (component) + { + case TensorDataLayoutComponent::C: + return TensorComponentType::Stride0; + case TensorDataLayoutComponent::W: + return TensorComponentType::Stride1; + case TensorDataLayoutComponent::H: + return TensorComponentType::Stride2; + case TensorDataLayoutComponent::D: + return TensorComponentType::Stride3; + case TensorDataLayoutComponent::N: + return TensorComponentType::Stride4; + default: + COMPUTE_KERNEL_WRITER_ERROR_ON_MSG("Unsupported tensor component for NDHWC"); + return TensorComponentType::Unknown; + } + default: + COMPUTE_KERNEL_WRITER_ERROR_ON_MSG("Unsupported tensor data layout"); + return TensorComponentType::Unknown; + } +} +} // namespace ckw diff --git a/compute_kernel_writer/src/TensorUtils.h b/compute_kernel_writer/src/TensorUtils.h new file mode 100644 index 0000000000..bb0af5c0b9 --- /dev/null +++ b/compute_kernel_writer/src/TensorUtils.h @@ -0,0 +1,57 @@ +/* + * 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_SRC_TENSORUTILS_H +#define CKW_SRC_TENSORUTILS_H + +#include <cstdint> + +/** Tensor specific utility functions */ +namespace ckw +{ +// Forward declarations +enum class TensorDataLayout; +enum class TensorDataLayoutComponent; +enum class TensorComponentType : uint32_t; + +/** Get tensor dimension from a given data layout and data layout component + * + * @param[in] layout Layout of the tensor + * @param[in] component Data layout component + * + * @return the @ref TensorComponent + */ +TensorComponentType get_tensor_dimension(TensorDataLayout layout, TensorDataLayoutComponent component); + +/** Get tensor stride from a given data layout and data layout component + * + * @param[in] layout Layout of the tensor + * @param[in] component Data layout component + * + * @return the @ref TensorComponent + */ +TensorComponentType get_tensor_stride(TensorDataLayout layout, TensorDataLayoutComponent component); +} // namespace ckw + +#endif // CKW_SRC_TENSORUTILS_H diff --git a/compute_kernel_writer/src/TileInfo.cpp b/compute_kernel_writer/src/TileInfo.cpp new file mode 100644 index 0000000000..273266eedc --- /dev/null +++ b/compute_kernel_writer/src/TileInfo.cpp @@ -0,0 +1,73 @@ +/* + * 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. + */ + +#include "ckw/TileInfo.h" + +namespace ckw +{ +TileInfo::TileInfo(DataType dt) : _dt(dt), _shape({{1, 1}}) +{ +} + +TileInfo::TileInfo(DataType dt, int32_t w) : _dt(dt), _shape({{w, 1}}) +{ +} + +TileInfo::TileInfo(DataType dt, int32_t h, int32_t w) : _dt(dt), _shape({{w, h}}) +{ +} + +TileInfo &TileInfo::width(int32_t w) +{ + _shape[kTileWidthIdx] = w; + return *this; +} + +int32_t TileInfo::width() const +{ + return _shape[kTileWidthIdx]; +} + +TileInfo &TileInfo::height(int32_t h) +{ + _shape[kTileHeightIdx] = h; + return *this; +} + +int32_t TileInfo::height() const +{ + return _shape[kTileHeightIdx]; +} + +TileInfo &TileInfo::data_type(DataType dt) +{ + _dt = dt; + return *this; +} + +DataType TileInfo::data_type() const +{ + return _dt; +} +} // namespace ckw diff --git a/compute_kernel_writer/src/TileOperand.cpp b/compute_kernel_writer/src/TileOperand.cpp new file mode 100644 index 0000000000..8ced6cfe3f --- /dev/null +++ b/compute_kernel_writer/src/TileOperand.cpp @@ -0,0 +1,89 @@ +/* + * 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. + */ + +#include "ckw/TileOperand.h" + +#include "ckw/Error.h" + +#include "src/ITile.h" + +namespace ckw +{ + +TileOperand::TileOperand() : _tile(nullptr), _row_start(0), _row_end(0), _col_start(0), _col_end(0) +{ +} + +TileOperand::TileOperand(ITile &tile) + : _tile(&tile), _row_start(0), _row_end(tile.info().height()), _col_start(0), _col_end(tile.info().width()) +{ +} + +TileOperand::TileOperand( + const TileOperand &operand, int32_t row_start, int32_t row_end, int32_t col_start, int32_t col_end) + : _tile(operand._tile), _row_start(row_start), _row_end(row_end), _col_start(col_start), _col_end(col_end) +{ + CKW_ASSERT(row_start >= 0 && row_start < _tile->info().height()); + CKW_ASSERT(row_end > row_start && row_end <= _tile->info().height()); + CKW_ASSERT(col_start >= 0 && col_start < _tile->info().width()); + CKW_ASSERT(col_end > col_start && col_end <= _tile->info().width()); +} + +bool TileOperand::is_valid() const +{ + return _tile != nullptr; +} + +const TileInfo &TileOperand::tile_info() const +{ + return _tile->info(); +} + +TileOperand TileOperand::tile(int32_t row_start, int32_t row_end, int32_t col_start, int32_t col_end) const +{ + CKW_ASSERT(row_start >= 0 && _row_start + row_start < _row_end); + CKW_ASSERT(row_end > row_start && _row_start + row_end <= _row_end); + CKW_ASSERT(col_start >= 0 && _col_start + col_start < _col_end); + CKW_ASSERT(col_end > col_start && _col_start + col_end <= _col_end); + + return TileOperand(*this, _row_start + row_start, _row_start + row_end, _col_start + col_start, + _col_start + col_end); +} + +TileOperand TileOperand::row(int32_t row) const +{ + CKW_ASSERT(row >= 0 && _row_start + row < _row_end); + + return tile(_row_start + row, _row_start + row + 1, _col_start, _col_end); +} + +TileOperand TileOperand::scalar(int32_t row, int32_t col) const +{ + CKW_ASSERT(row >= 0 && _row_start + row < _row_end); + CKW_ASSERT(col >= 0 && _col_start + col < _col_end); + + return tile(_row_start + row, _row_start + row + 1, _col_start + col, _col_start + col + 1); +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/TileView.cpp b/compute_kernel_writer/src/TileView.cpp new file mode 100644 index 0000000000..ea803f92f4 --- /dev/null +++ b/compute_kernel_writer/src/TileView.cpp @@ -0,0 +1,57 @@ +/* + * 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. + */ + +#include "src/TileView.h" + +#include <cstdint> + +namespace ckw +{ + +TileArea::TileArea(int32_t row_start, int32_t row_end, int32_t col_start, int32_t col_end) + : _row_start(row_start), _row_end(row_end), _col_start(col_start), _col_end(col_end) +{ +} + +int32_t TileArea::row_start() const +{ + return _row_start; +} + +int32_t TileArea::row_end() const +{ + return _row_end; +} + +int32_t TileArea::col_start() const +{ + return _col_start; +} + +int32_t TileArea::col_end() const +{ + return _col_end; +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/TileView.h b/compute_kernel_writer/src/TileView.h new file mode 100644 index 0000000000..42854ac823 --- /dev/null +++ b/compute_kernel_writer/src/TileView.h @@ -0,0 +1,209 @@ +/* + * 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_SRC_TILEVIEW_H +#define CKW_SRC_TILEVIEW_H + +#include "ckw/Error.h" +#include "ckw/types/DataType.h" + +#include "src/ITile.h" + +#include <cstdint> + +namespace ckw +{ + +/** A rectangular active area of a tile. */ +class TileArea +{ +public: + /** Create a new tile rectangular active area. + * + * The range of rows and columns is defined by pairs of start and end indices, inclusive lower and exclusive upper. + * In other word, any row and column indices satisfied the following conditions will be part of the active area: + * + * row_start <= row_index < row_end + * col_start <= col_index < col_end + * + * @param[in] row_start The start index of the row range. + * @param[in] row_end The end index of the row range. + * @param[in] col_start The start index of the column range. + * @param[in] col_end The end index of the column range. + */ + TileArea(int32_t row_start, int32_t row_end, int32_t col_start, int32_t col_end); + + /** Get the start row index. */ + int32_t row_start() const; + + /** Get the end row (exclusive) index. */ + int32_t row_end() const; + + /** Get the start column index. */ + int32_t col_start() const; + + /** Get the end column (exclusive) index. */ + int32_t col_end() const; + +private: + int32_t _row_start; + int32_t _row_end; + int32_t _col_start; + int32_t _col_end; +}; + +/** A rectangular view of a tile. */ +template <typename T> +class TileView +{ +public: + /** Default constructor */ + TileView() : _tile(nullptr), _area(0, 0, 0, 0) + { + } + /** Create a tile view that refers to the whole tile. + * + * @param[in] tile The tile object. + */ + TileView(const T &tile) : _tile(&tile), _area(0, tile.info().height(), 0, tile.info().width()) + { + } + + /** Create a new rectangular view of the given tile. + * + * @param[in] tile The tile object. + * @param[in] area The rectangular active area. + */ + TileView(const T &tile, const TileArea &area) : _tile(&tile), _area(area) + { + } + + /** Get the tile object. + * + * The caller must guarantee that the tile view refers to the whole tile. + */ + const T &full_tile() const + { + CKW_ASSERT(is_full_tile()); + + return *_tile; + } + + /** Get the data type of the tile. */ + DataType data_type() const + { + return _tile->info().data_type(); + } + + /** Get the start row index. */ + int32_t row_start() const + { + return _area.row_start(); + } + + /** Get the end row index. */ + int32_t row_end() const + { + return _area.row_end(); + } + + /** Get the start column index. */ + int32_t col_start() const + { + return _area.col_start(); + } + + /** Get the end column index. */ + int32_t col_end() const + { + return _area.col_end(); + } + + /** Get the height of the tile view. */ + int32_t height() const + { + return _area.row_end() - _area.row_start(); + } + + /** Get the width of the tile view. */ + int32_t width() const + { + return _area.col_end() - _area.col_start(); + } + + /** See @ref IVectorAccess::vector. */ + TileVariable vector(int32_t row) const + { + return _tile->vector(row_start() + row, col_start(), width()); + } + + /** See @ref IScalarAccess::scalar. */ + TileVariable scalar(int32_t row, int32_t col) const + { + return _tile->scalar(row_start() + row, col_start() + col); + } + + /** Get the name of the tile. */ + const std::string &name() const + { + return _tile->name(); + } + + /** Get whether the tile view is a scalar element. */ + bool is_scalar() const + { + return height() == 1 && width() == 1; + } + + /** Get whether the tile view refers to the whole tile. */ + bool is_full_tile() const + { + return row_start() == 0 && row_end() == _tile->info().height() && col_start() == 0 && + col_end() == _tile->info().width(); + } + + /** Set the rectangular active area. + * + * @param[in] area The rectangular active area. + */ + TileView &area(const TileArea &area) + { + _area = area; + return *this; + } + + /** Get the tile area */ + TileArea area() const + { + return _area; + } + +private: + const T *_tile; + TileArea _area; +}; + +} // namespace ckw + +#endif // CKW_SRC_TILEVIEW_H diff --git a/compute_kernel_writer/src/cl/CLHelpers.cpp b/compute_kernel_writer/src/cl/CLHelpers.cpp new file mode 100644 index 0000000000..252c5cdfcb --- /dev/null +++ b/compute_kernel_writer/src/cl/CLHelpers.cpp @@ -0,0 +1,353 @@ +/* + * 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. + */ + +#include "src/cl/CLHelpers.h" + +#include "ckw/Error.h" +#include "ckw/types/DataType.h" +#include "ckw/types/Operators.h" +#include "ckw/types/TensorStorageType.h" + +#include "src/types/DataTypeHelpers.h" + +namespace ckw +{ +bool cl_validate_vector_length(int32_t len) +{ + bool valid_vector_length = true; + if (len < 1 || len > 16 || (len > 4 && len < 8) || (len > 8 && len < 16)) + { + valid_vector_length = false; + } + return valid_vector_length; +} + +std::string cl_get_variable_datatype_as_string(DataType dt, int32_t len) +{ + if (cl_validate_vector_length(len) == false) + { + CKW_THROW_MSG("Unsupported vector length"); + return ""; + } + + std::string res; + switch (dt) + { + case DataType::Fp32: + res += "float"; + break; + case DataType::Fp16: + res += "half"; + break; + case DataType::Int8: + res += "char"; + break; + case DataType::Uint8: + res += "uchar"; + break; + case DataType::Uint16: + res += "ushort"; + break; + case DataType::Int16: + res += "short"; + break; + case DataType::Uint32: + res += "uint"; + break; + case DataType::Int32: + res += "int"; + break; + case DataType::Bool: + res += "bool"; + break; + default: + CKW_THROW_MSG("Unsupported datatype"); + return ""; + } + + if (len > 1) + { + res += std::to_string(len); + } + + return res; +} + +int32_t cl_round_up_to_nearest_valid_vector_width(int32_t width) +{ + switch (width) + { + case 1: + return 1; + case 2: + return 2; + case 3: + return 3; + case 4: + return 4; + case 5: + case 6: + case 7: + case 8: + return 8; + case 9: + case 10: + case 11: + case 12: + case 13: + case 14: + case 15: + case 16: + return 16; + default: + CKW_THROW_MSG("Unsupported width to convert to OpenCL vector"); + return 0; + } +} + +std::string cl_get_variable_storagetype_as_string(TensorStorageType storage) +{ + std::string res; + switch (storage) + { + case TensorStorageType::BufferUint8Ptr: + res += "__global uchar*"; + break; + case TensorStorageType::Texture2dReadOnly: + res += "__read_only image2d_t"; + break; + case TensorStorageType::Texture2dWriteOnly: + res += "__write_only image2d_t"; + break; + default: + CKW_THROW_MSG("Unsupported storage type"); + } + + return res; +} + +std::string cl_get_assignment_op_as_string(AssignmentOp op) +{ + switch (op) + { + case AssignmentOp::Increment: + return "+="; + + case AssignmentOp::Decrement: + return "-="; + + default: + CKW_THROW_MSG("Unsupported assignment operator!"); + } +} + +std::tuple<bool, std::string> cl_get_unary_op(UnaryOp op) +{ + switch (op) + { + case UnaryOp::LogicalNot: + return {false, "!"}; + + case UnaryOp::BitwiseNot: + return {false, "~"}; + + case UnaryOp::Exp: + return {true, "exp"}; + + case UnaryOp::Tanh: + return {true, "tanh"}; + + case UnaryOp::Sqrt: + return {true, "sqrt"}; + + case UnaryOp::Erf: + return {true, "erf"}; + + case UnaryOp::Fabs: + return {true, "fabs"}; + + case UnaryOp::Log: + return {true, "log"}; + + case UnaryOp::Round: + return {true, "round"}; + + case UnaryOp::Floor: + return {true, "floor"}; + + default: + CKW_THROW_MSG("Unsupported unary operation!"); + } +} + +std::tuple<bool, std::string> cl_get_binary_op(BinaryOp op, DataType data_type) +{ + const auto is_float = is_data_type_float(data_type); + + switch (op) + { + case BinaryOp::Add: + return {false, "+"}; + + case BinaryOp::Sub: + return {false, "-"}; + + case BinaryOp::Mul: + return {false, "*"}; + + case BinaryOp::Div: + return {false, "/"}; + + case BinaryOp::Mod: + return {false, "%"}; + + case BinaryOp::Equal: + return {false, "=="}; + + case BinaryOp::Less: + return {false, "<"}; + + case BinaryOp::LessEqual: + return {false, "<="}; + + case BinaryOp::Greater: + return {false, ">"}; + + case BinaryOp::GreaterEqual: + return {false, ">="}; + + case BinaryOp::LogicalAnd: + return {false, "&&"}; + + case BinaryOp::LogicalOr: + return {false, "||"}; + + case BinaryOp::BitwiseXOR: + return {false, "^"}; + + case BinaryOp::Min: + return {true, is_float ? "fmin" : "min"}; + + case BinaryOp::Max: + return {true, is_float ? "fmax" : "max"}; + + default: + CKW_THROW_MSG("Unsupported binary operator/function!"); + } +} + +std::tuple<bool, std::string> cl_get_ternary_op(TernaryOp op) +{ + switch (op) + { + case TernaryOp::Select: + return {true, "select"}; + + case TernaryOp::Clamp: + return {true, "clamp"}; + + default: + CKW_THROW_MSG("Unsupported ternary function!"); + } +} + +std::string cl_data_type_rounded_up_to_valid_vector_width(DataType dt, int32_t width) +{ + std::string data_type; + const int32_t w = cl_round_up_to_nearest_valid_vector_width(width); + data_type += cl_get_variable_datatype_as_string(dt, 1); + if (w != 1) + { + data_type += std::to_string(w); + } + return data_type; +} + +std::vector<int32_t> cl_decompose_vector_width(int32_t vector_width) +{ + std::vector<int32_t> x; + + switch (vector_width) + { + case 0: + break; + case 1: + case 2: + case 3: + case 4: + case 8: + case 16: + x.push_back(vector_width); + break; + case 5: + x.push_back(4); + x.push_back(1); + break; + case 6: + x.push_back(4); + x.push_back(2); + break; + case 7: + x.push_back(4); + x.push_back(3); + break; + case 9: + x.push_back(8); + x.push_back(1); + break; + case 10: + x.push_back(8); + x.push_back(2); + break; + case 11: + x.push_back(8); + x.push_back(3); + break; + case 12: + x.push_back(8); + x.push_back(4); + break; + case 13: + x.push_back(8); + x.push_back(4); + x.push_back(1); + break; + case 14: + x.push_back(8); + x.push_back(4); + x.push_back(2); + break; + case 15: + x.push_back(8); + x.push_back(4); + x.push_back(3); + break; + + default: + CKW_THROW_MSG("Vector width is too large"); + } + return x; +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/cl/CLHelpers.h b/compute_kernel_writer/src/cl/CLHelpers.h new file mode 100644 index 0000000000..370ffc700c --- /dev/null +++ b/compute_kernel_writer/src/cl/CLHelpers.h @@ -0,0 +1,138 @@ +/* + * 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_SRC_CL_CLHELPERS_H +#define CKW_SRC_CL_CLHELPERS_H + +#include "ckw/types/Operators.h" + +#include <cstdint> +#include <string> +#include <tuple> +#include <vector> + +/** OpenCL specific helper functions */ +namespace ckw +{ +// Forward declarations +enum class DataType; +enum class TensorStorageType : uint32_t; + +/** Helper function to validate the vector length of OpenCL vector data types + * + * @param[in] len Vector length + * + * @return true if the vector lenght is valid. It returns false, otherwise. + */ +bool cl_validate_vector_length(int32_t len); + +/** Helper function to return the OpenCL datatype as a string from a @ref DataType and vector length as int32_t variable + * + * @param[in] dt Datatype + * @param[in] len Vector length + * + * @return the OpenCL datatype as a string + */ +std::string cl_get_variable_datatype_as_string(DataType dt, int32_t len); + +/** Return the assignment operator in OpenCL language. + * + * @param[in] op The assignment operator. + * + * @return The operator in OpenCL language as a string. + */ +std::string cl_get_assignment_op_as_string(AssignmentOp op); + +/** Return the information about the unary operation. + * + * The result contains: + * - is_func: true if it's a function and false if it's an unary operator in OpenCL language. + * - str: the function name or the operator in OpenCL language. + * + * @param[in] op The unary operator. + * + * @return The information about the unary operation. + */ +std::tuple<bool, std::string> cl_get_unary_op(UnaryOp op); + +/** Return the information about the binary operation. + * + * The result contains: + * - is_func: true if it's a function and false if it's an binary operator in OpenCL language. + * - str: the function name or the operator in OpenCL language. + * + * @param[in] op The binary operator. + * @param[in] data_type The input data type. + * + * @return The information about the binary operation. + */ +std::tuple<bool, std::string> cl_get_binary_op(BinaryOp op, DataType data_type); + +/** Return the information about the ternary operation. + * + * The result contains: + * - is_func: true if it's a function and false if it's a ternary operator in OpenCL language. + * - str: the function name or the operator in OpenCL language. + * + * @param[in] op The ternary operator. + * + * @return The information about the ternary operation. + */ +std::tuple<bool, std::string> cl_get_ternary_op(TernaryOp op); + +/** Helper function to return the OpenCL vector size that accommodate the the desired width + * + * @param[in] width The desired width + * + * @return the OpenCL vector size +*/ +int32_t cl_round_up_to_nearest_valid_vector_width(int32_t width); + +/** Helper function to return the OpenCL storage type as a string from a @ref TensorStorage + * + * @param[in] storage Storage type + * + * @return the OpenCL storage type as a string + */ +std::string cl_get_variable_storagetype_as_string(TensorStorageType storage); + +/** Helper function to decompose a vector width into a summation of valid OpenCL vector widths. + * + * @param[in] vector_width Vector width to be decomposed + * + * @return a vector of OpenCL vector widths + */ +std::vector<int32_t> cl_decompose_vector_width(int32_t vector_width); + +/** Helper function to get OpenCL data type from the data type enum and width + * It'll round up the given vector width to the nearest valid OpenCL vector width. + * + * @param[in] dt data type enum + * @param[in] width vector width + * + * @return a string representation of the data type + */ +std::string cl_data_type_rounded_up_to_valid_vector_width(DataType dt, int32_t width); +} // namespace ckw + +#endif /* CKW_SRC_CL_CLHELPERS_H */ diff --git a/compute_kernel_writer/src/cl/CLKernelWriter.cpp b/compute_kernel_writer/src/cl/CLKernelWriter.cpp new file mode 100644 index 0000000000..91512bde23 --- /dev/null +++ b/compute_kernel_writer/src/cl/CLKernelWriter.cpp @@ -0,0 +1,833 @@ +/* + * Copyright (c) 2023-2024 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. + */ + +#include "src/cl/CLKernelWriter.h" + +#include "ckw/Error.h" +#include "ckw/Kernel.h" +#include "ckw/TensorSampler.h" +#include "ckw/TileOperand.h" +#include "ckw/types/DataType.h" +#include "ckw/types/MemoryOperation.h" +#include "ckw/types/TargetLanguage.h" + +#include "src/cl/CLHelpers.h" +#include "src/cl/CLTensorArgument.h" +#include "src/cl/CLTile.h" +#include "src/cl/helpers/CLMemoryOpBufferHelper.h" +#include "src/cl/helpers/CLMemoryOpImage2dHelper.h" +#include "src/cl/helpers/ICLMemoryOpHelper.h" +#include "src/ITensorComponent.h" +#include "src/TileView.h" +#include "src/types/DataTypeHelpers.h" + +#include <algorithm> +#include <cstdint> +#include <tuple> +#include <vector> + +namespace +{ +std::string generate_cl_extensions() +{ + std::string ext = R"( +#if defined(cl_khr_fp16) +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +#endif // defined(cl_khr_fp16) + +#if defined(cl_arm_printf) +#pragma OPENCL EXTENSION cl_arm_printf : enable +#endif // defined(cl_arm_printf); + +#define inf (INFINITY) +)"; + return ext; +} +} // namespace + +namespace ckw +{ + +CLKernelWriter::CLKernelWriter() = default; +CLKernelWriter::~CLKernelWriter() = default; + +std::unique_ptr<Kernel> CLKernelWriter::emit_kernel(const std::string &name) +{ + std::string code; + code += generate_cl_extensions(); + code += "__kernel void "; + code += name; + code += "\n(\n"; + + // Create the list of arguments. + std::vector<KernelArgument> arguments; + + for (const auto &tensor : _tensors) + { + const auto tensor_id = tensor->info().id(); + + const auto storages = tensor->storages(); + const auto components = tensor->components(); + + for (const auto &storage : storages) + { + code += cl_get_variable_storagetype_as_string(storage.type); + code += " "; + code += storage.val; + code += ",\n"; + + arguments.emplace_back(tensor_id, storage.type); + } + + for (const auto &component : components) + { + const auto &tile = component->tile(); + const auto &tile_info = tile.info(); + + CKW_ASSERT(tile.is_scalar()); + + code += cl_get_variable_datatype_as_string(tile_info.data_type(), 1); + code += " "; + code += tile.name(); + code += ",\n"; + + arguments.emplace_back(tensor_id, component->component_type()); + } + } + + if (code.size() >= 2 && code[code.size() - 2] == ',' && code[code.size() - 1] == '\n') + { + // Remove the last comma in the argument list. + code.pop_back(); + code[code.size() - 1] = '\n'; + } + + code += ")\n{\n"; + + code += _body_source_code; + + code += "}\n"; + + return std::make_unique<Kernel>(TargetLanguage::OpenCL, arguments, code); +} + +void CLKernelWriter::op_assign(const TileOperand &dst, const TileOperand &src) +{ + const auto dst_view = to_cl_tile_view(dst); + const auto src_view = to_cl_tile_view(src); + + const auto dst_w = dst_view.width(); + const auto dst_h = dst_view.height(); + const auto src_w = src_view.width(); + + const auto data_type_str = cl_get_variable_datatype_as_string(dst_view.data_type(), dst_w); + + const auto broadcast_src_x = dst_w != 1 && src_w == 1; + const std::string src_prefix = broadcast_src_x ? "(" + data_type_str + ")" : ""; + + CKW_ASSERT_MSG(src_view.data_type() == dst_view.data_type(), "Source and destination type must match."); + CKW_ASSERT_MSG(src_view.height() == dst_h || src_view.height() == 1, + "Tile height must match or source is broadcasting in y dimension."); + CKW_ASSERT_MSG(src_w == dst_w || src_w == 1, "Tile width must match or source is broadcasting in x dimension."); + + // Broadcasting on y dimension is automatic (see CLTile::vector). + for (int32_t y = 0; y < dst_h; ++y) + { + append_code(dst_view.vector(y).str, " = ", src_prefix, src_view.vector(y).str, ";\n"); + } +} + +void CLKernelWriter::op_cast(const TileOperand &dst, const TileOperand &src, ConvertPolicy policy) +{ + const auto dst_view = to_cl_tile_view(dst); + const auto src_view = to_cl_tile_view(src); + + const auto dst_w = dst_view.width(); + const auto dst_h = dst_view.height(); + const auto src_w = src_view.width(); + + const auto dst_type = dst_view.data_type(); + + const auto convert_type_str = cl_get_variable_datatype_as_string(dst_type, src_w); + const auto dst_type_str = cl_get_variable_datatype_as_string(dst_type, dst_w); + + const std::string sat = policy == ConvertPolicy::Saturate ? "_sat" : ""; + + CKW_ASSERT_IF(policy == ConvertPolicy::Saturate, !is_data_type_float(dst_type)); + + const auto broadcast_x = dst_w != 1 && src_w == 1; + const std::string prefix = broadcast_x ? "(" + dst_type_str + ")" : ""; + + CKW_ASSERT_MSG(src_view.height() == dst_h || src_view.height() == 1, + "Tile height must match or source is broadcasting in y dimension."); + CKW_ASSERT_MSG(src_w == dst_w || src_w == 1, "Tile width must match or source is broadcasting in x dimension."); + + // Broadcasting on y dimension is automatic (see CLTile::vector). + if (src_view.data_type() == dst_view.data_type()) + { + for (int32_t y = 0; y < dst_h; ++y) + { + append_code(dst_view.vector(y).str, " = ", src_view.vector(y).str, ";\n"); + } + } + else + { + for (int32_t y = 0; y < dst_h; ++y) + { + append_code(dst_view.vector(y).str, " = ", prefix, "convert_", convert_type_str, sat, "(", + src_view.vector(y).str, ");\n"); + } + } +} + +void CLKernelWriter::op_unary(const TileOperand &dst, UnaryOp op, const TileOperand &src) +{ + const auto dst_view = to_cl_tile_view(dst); + const auto src_view = to_cl_tile_view(src); + + const auto dst_w = dst_view.width(); + const auto dst_h = dst_view.height(); + const auto src_w = src_view.width(); + + const auto data_type_str = cl_get_variable_datatype_as_string(dst_view.data_type(), dst_w); + const auto broadcast_src_x = dst_w != 1 && src_w == 1; + + const std::string src_prefix = broadcast_src_x ? "(" + data_type_str + ")" : ""; + + const auto op_info = cl_get_unary_op(op); + const auto op_is_func = std::get<0>(op_info); + const auto &op_name = std::get<1>(op_info); + const auto op_prefix = op_is_func ? op_name + "(" : op_name; + const auto op_suffix = op_is_func ? ")" : ""; + + CKW_ASSERT_MSG(src_view.data_type() == dst_view.data_type(), "Source and destination type must match."); + CKW_ASSERT_MSG(src_view.height() == dst_h || src_view.height() == 1, + "Tile height must match or source is broadcasting in y dimension."); + CKW_ASSERT_MSG(src_w == dst_w || src_w == 1, "Tile width must match or source is broadcasting in x dimension."); + + // Broadcasting on y dimension is automatic (see CLTile::vector). + for (int32_t y = 0; y < dst_h; ++y) + { + append_code(dst_view.vector(y).str, " = ", src_prefix, op_prefix, src_view.vector(y).str, op_suffix, ";\n"); + } +} + +void CLKernelWriter::op_binary(const TileOperand &dst, BinaryOp op, const TileOperand &first, const TileOperand &second) +{ + const auto dst_view = to_cl_tile_view(dst); + const auto lhs_view = to_cl_tile_view(first); + const auto rhs_view = to_cl_tile_view(second); + + const auto dst_w = dst_view.width(); + const auto dst_h = dst_view.height(); + const auto lhs_w = lhs_view.width(); + const auto rhs_w = rhs_view.width(); + + const auto data_type = lhs_view.data_type(); + + CKW_ASSERT_MSG(lhs_view.data_type() == rhs_view.data_type(), "LHS and RHS type must match."); + + if (op == BinaryOp::MatMul_Nt_T) + { + CKW_ASSERT_MSG(lhs_view.height() == dst_h, "LHS tile height must match the DST tile height"); + CKW_ASSERT_MSG(rhs_view.height() == dst_w, "RHS tile height must match the DST tile width"); + CKW_ASSERT_MSG(lhs_view.width() == rhs_view.width(), "LHS tile width must match the LHS tile width"); + + CKW_ASSERT(is_data_type_float(data_type)); + + for (int32_t y = 0; y < dst_h; ++y) + { + for (int32_t x = 0; x < dst_w; ++x) + { + for (int32_t k = 0; k < lhs_w; ++k) + { + append_code(dst_view.scalar(y, x).str, " = fma(", lhs_view.scalar(y, k).str, ", ", + rhs_view.scalar(x, k).str, ", ", dst_view.scalar(y, x).str, ");\n"); + } + } + } + } + else + { + CKW_ASSERT_MSG(lhs_view.height() == dst_h || lhs_view.height() == 1, + "LHS tile height must match or source is broadcasting in y dimension."); + CKW_ASSERT_MSG(rhs_view.height() == dst_h || rhs_view.height() == 1, + "RHS tile height must match or source is broadcasting in y dimension."); + + CKW_ASSERT_MSG(lhs_w == dst_w || lhs_w == 1, + "LHS tile width must match destination or LHS is broadcasting in x dimension."); + CKW_ASSERT_MSG(rhs_w == dst_w || rhs_w == 1, + "RHS tile width must match destination or RHS is broadcasting in x dimension."); + + const auto op_info = cl_get_binary_op(op, data_type); + const auto op_is_func = std::get<0>(op_info); + const auto &op_name = std::get<1>(op_info); + + const auto data_type_str = cl_get_variable_datatype_as_string(data_type, dst_w); + + const auto broadcast_lhs_x = dst_w != 1 && lhs_w == 1; + const auto broadcast_rhs_x = dst_w != 1 && rhs_w == 1; + + const std::string lhs_prefix = broadcast_lhs_x ? "(" + data_type_str + ")" : ""; + const std::string rhs_prefix = broadcast_rhs_x ? "(" + data_type_str + ")" : ""; + + const std::string op_prefix = op_is_func ? " = " + op_name + "(" : " = "; + const std::string op_separator = op_is_func ? ", " : " " + op_name + " "; + const std::string op_suffix = op_is_func ? ");\n" : ";\n"; + + // Broadcasting on y dimension is automatic (see CLTile::vector). + for (int32_t y = 0; y < dst_h; ++y) + { + append_code(dst_view.vector(y).str, op_prefix, lhs_prefix, lhs_view.vector(y).str, op_separator, rhs_prefix, + rhs_view.vector(y).str, op_suffix); + } + } +} + +void CLKernelWriter::op_ternary( + const TileOperand &dst, TernaryOp op, const TileOperand &first, const TileOperand &second, const TileOperand &third) +{ + const auto dst_view = to_cl_tile_view(dst); + const auto first_view = to_cl_tile_view(first); + const auto second_view = to_cl_tile_view(second); + const auto third_view = to_cl_tile_view(third); + + const auto dst_w = dst_view.width(); + const auto dst_h = dst_view.height(); + const auto first_w = first_view.width(); + const auto second_w = second_view.width(); + const auto third_w = third_view.width(); + + const auto data_type = dst_view.data_type(); + const auto data_type_str = cl_get_variable_datatype_as_string(data_type, dst_w); + + const auto op_info = cl_get_ternary_op(op); + const auto op_is_func = std::get<0>(op_info); + const auto &op_name = std::get<1>(op_info); + + const auto broadcast_first_x = dst_w != 1 && first_w == 1; + const auto broadcast_second_x = dst_w != 1 && second_w == 1; + const auto broadcast_third_x = dst_w != 1 && third_w == 1; + + const std::string first_prefix = broadcast_first_x ? "(" + data_type_str + ")" : ""; + const std::string second_prefix = broadcast_second_x ? "(" + data_type_str + ")" : ""; + const std::string third_prefix = broadcast_third_x ? "(" + data_type_str + ")" : ""; + + CKW_UNUSED(op_is_func); + CKW_ASSERT_MSG(op_is_func, "The only supported ternary operator is function."); + CKW_ASSERT_MSG(second_view.data_type() == dst_view.data_type(), "2nd source and destination type must match."); + CKW_ASSERT_MSG(third_view.data_type() == dst_view.data_type(), "3rd source and destination type must match."); + + CKW_ASSERT_MSG(first_view.height() == dst_h || first_view.height() == 1, + "1st tile height must match or source is broadcasting in y dimension."); + CKW_ASSERT_MSG(second_view.height() == dst_h || second_view.height() == 1, + "2nd tile height must match or source is broadcasting in y dimension."); + CKW_ASSERT_MSG(third_view.height() == dst_h || third_view.height() == 1, + "3rd tile height must match or source is broadcasting in y dimension."); + + CKW_ASSERT_MSG(first_w == dst_w || first_w == 1, + "1st tile width must match or source is broadcasting in x dimension."); + CKW_ASSERT_MSG(second_w == dst_w || second_w == 1, + "2nd tile width must match or source is broadcasting in x dimension."); + CKW_ASSERT_MSG(third_w == dst_w || third_w == 1, + "3rd tile width must match or source is broadcasting in x dimension."); + + // Broadcasting on y dimension is automatic (see CLTile::vector). + for (int32_t y = 0; y < dst_h; ++y) + { + append_code(dst_view.vector(y).str, " = ", op_name, "(", first_prefix, first_view.vector(y).str, ", ", + second_prefix, second_view.vector(y).str, ", ", third_prefix, third_view.vector(y).str, ");\n"); + } +} + +void CLKernelWriter::op_if_generic( + const TileOperand &lhs, BinaryOp op, const TileOperand &rhs, const std::function<void()> &body, bool is_else_if) +{ + const auto lhs_view = to_cl_tile_view(lhs); + const auto rhs_view = to_cl_tile_view(rhs); + + const auto op_name = std::get<1>(cl_get_binary_op(op, lhs_view.data_type())); + CKW_ASSERT(op == BinaryOp::Less || op == BinaryOp::LessEqual || op == BinaryOp::Equal || + op == BinaryOp::GreaterEqual || op == BinaryOp::Greater); + + CKW_ASSERT(lhs_view.is_scalar()); + CKW_ASSERT(rhs_view.is_scalar()); + + if (is_else_if) + { + append_code("else "); + } + + append_code("if (", lhs_view.scalar(0, 0).str, " ", op_name, " ", rhs_view.scalar(0, 0).str, ")\n{\n"); + write_body(body); + append_code("}\n"); +} + +void CLKernelWriter::op_if(const TileOperand &lhs, + BinaryOp op, + const TileOperand &rhs, + const std::function<void()> &body) +{ + op_if_generic(lhs, op, rhs, body, false /* is_else_if */); +} + +void CLKernelWriter::op_else_if(const TileOperand &lhs, + BinaryOp op, + const TileOperand &rhs, + const std::function<void()> &body) +{ + op_if_generic(lhs, op, rhs, body, true /* is_else_if */); +} + +void CLKernelWriter::op_else(const std::function<void()> &body) +{ + append_code("else\n{\n"); + write_body(body); + append_code("}\n"); +} + +void CLKernelWriter::op_for_loop(const TileOperand &var, + BinaryOp cond_op, + const TileOperand &cond_value, + const TileOperand &update_var, + AssignmentOp update_op, + const TileOperand &update_value, + const std::function<void()> &body) +{ + const auto var_view = to_cl_tile_view(var); + const auto cond_value_view = to_cl_tile_view(cond_value); + const auto update_var_view = to_cl_tile_view(update_var); + const auto update_value_view = to_cl_tile_view(update_value); + + CKW_ASSERT(var_view.is_scalar()); + CKW_ASSERT(cond_value_view.is_scalar()); + CKW_ASSERT(update_var_view.is_scalar()); + CKW_ASSERT(update_value_view.is_scalar()); + + CKW_ASSERT(var_view.data_type() == cond_value_view.data_type()); + CKW_ASSERT(update_var_view.data_type() == update_value_view.data_type()); + + const auto cond_op_name = std::get<1>(cl_get_binary_op(cond_op, var_view.data_type())); + CKW_ASSERT(cond_op == BinaryOp::Less || cond_op == BinaryOp::LessEqual || cond_op == BinaryOp::Equal || + cond_op == BinaryOp::GreaterEqual || cond_op == BinaryOp::Greater); + + append_code("for (; ", var_view.scalar(0, 0).str, " ", cond_op_name, " ", cond_value_view.scalar(0, 0).str, "; ", + update_var_view.scalar(0, 0).str, " ", cl_get_assignment_op_as_string(update_op), " ", + update_value_view.scalar(0, 0).str, ")\n{\n"); + write_body(body); + append_code("}\n"); +} + +void CLKernelWriter::op_return() +{ + append_code("return;\n"); +} + +void CLKernelWriter::op_get_global_id(const TileOperand &dst, int32_t dim) +{ + const auto tile_view = to_cl_tile_view(dst); + + CKW_ASSERT(tile_view.is_scalar()); + CKW_ASSERT(tile_view.data_type() == DataType::Int32 || tile_view.data_type() == DataType::Uint32); + + CKW_ASSERT(dim >= 0 && dim <= 2); + + append_code(tile_view.scalar(0, 0).str, " = get_global_id(", std::to_string(dim), ");\n"); +} + +void CLKernelWriter::op_print(const std::string &prefix, const std::vector<TileOperand> &operands) +{ + std::string format_code; + std::string args_code; + + for (auto &op : operands) + { + const auto tile_view = to_cl_tile_view(op); + + const auto name = tile_view.name(); + const auto width = tile_view.width(); + const auto height = tile_view.height(); + const auto data_type = tile_view.data_type(); + + // Construct the format specifier to print out one row of the tile. + std::string row_format("%"); + + if (width > 1) + { + row_format += "v" + std::to_string(width); + } + + switch (data_type) + { + case DataType::Fp32: + row_format += "hlg"; + break; + case DataType::Fp16: + row_format += "hg"; + break; + case DataType::Int32: + case DataType::Bool: + row_format += (width > 1) ? "hli" : "i"; + break; + case DataType::Int16: + row_format += "hi"; + break; + case DataType::Int8: + row_format += "hhi"; + break; + case DataType::Uint32: + row_format += (width > 1) ? "hlu" : "u"; + break; + case DataType::Uint16: + row_format += "hu"; + break; + case DataType::Uint8: + row_format += "hhu"; + break; + default: + CKW_THROW_MSG("Unsupported data type!"); + } + + if (width > 1) + { + row_format = "[" + row_format + "]"; + } + + // Construct the format specifier for the printf statement. + format_code += name + " = "; + + if (height == 1) + { + format_code += row_format; + } + else + { + format_code += "[" + row_format; + for (int32_t row = 1; row < height; ++row) + { + format_code += ", " + row_format; + } + format_code += "]"; + } + + format_code += "\\n"; + + // Construct the variable arguments for the printf statement. + for (int32_t row = 0; row < height; ++row) + { + args_code += ", " + tile_view.vector(row).str; + } + } + + append_code("printf(\"", prefix, "\\n", format_code, "\"", args_code, ");\n"); +} + +void CLKernelWriter::op_comment(const std::string &text) +{ +#ifdef COMPUTE_KERNEL_WRITER_DEBUG_ENABLED + + CKW_ASSERT(text.find("\n") == text.npos); + CKW_ASSERT(text.find("\r") == text.npos); + + append_code("// ", text, "\n"); + +#else // COMPUTE_KERNEL_WRITER_DEBUG_ENABLED + + CKW_UNUSED(text); + +#endif // COMPUTE_KERNEL_WRITER_DEBUG_ENABLED +} + +const std::string &CLKernelWriter::body_source_code() const +{ + return _body_source_code; +} + +TensorOperand CLKernelWriter::declare_tensor_argument(const std::string &name, const TensorInfo &info) +{ + const auto fullname = generate_full_name(name); + + auto tensor = std::make_unique<CLTensorArgument>(fullname, info, false /* return_dims_by_value */); + const auto operand = create_tensor_operand(*tensor); + + _tensors.insert(std::move(tensor)); + + return operand; +} + +TileOperand CLKernelWriter::declare_tile(const std::string &name, const TileInfo &tile_info) +{ + const std::string fullname = generate_full_name(name); + + const int32_t height = tile_info.height(); + const int32_t width = tile_info.width(); + const DataType data_type = tile_info.data_type(); + + CKW_ASSERT_MSG(std::find_if(_tiles.begin(), _tiles.end(), + [=](const std::unique_ptr<CLTile> &e) + { return e->name() == fullname; }) == _tiles.end(), + "There is already a tile with name: " + fullname); + + auto tile = std::make_unique<CLTile>(fullname, tile_info); + + for (int32_t row = 0; row < height; ++row) + { + const std::string cl_type = cl_get_variable_datatype_as_string(data_type, width); + append_code(cl_type, " ", tile->vector(row).str, ";\n"); + } + + const auto operand = create_tile_operand(*tile); + + _tiles.insert(std::move(tile)); + + return operand; +} + +TileOperand CLKernelWriter::declare_constant_tile(const ConstantData &data) +{ + auto tile = std::make_unique<CLTile>(get_values(data), get_data_type(data)); + const TileOperand operand = create_tile_operand(*tile); + _constant_tiles.insert(std::move(tile)); + + return operand; +} + +void CLKernelWriter::op_write_raw_code(const std::string &raw_code) +{ + append_code(raw_code); +} + +TileView<CLTile> CLKernelWriter::to_cl_tile_view(const TileOperand &operand) const +{ + const auto tile_and_area = get_tile(operand); + ITile &tile = std::get<0>(tile_and_area); + const TileArea area = std::get<1>(tile_and_area); + +#ifdef COMPUTE_KERNEL_WRITER_ASSERTS_ENABLED + // Check if the tile is a CLTile created by this kernel writer. + + { + bool found = false; + + for (const auto &t : _tiles) + { + if (&tile == t.get()) + { + found = true; + break; + } + } + + for (const auto &t : _constant_tiles) + { + if (&tile == t.get()) + { + found = true; + break; + } + } + + if (!found) + { + for (const auto &t : _tensors) + { + const auto components = t->components(); + + for (const auto component : components) + { + if (&tile == &component->tile()) + { + found = true; + break; + } + } + + if (found) + { + break; + } + } + } + + CKW_ASSERT_MSG(found, "The tile is not found!"); + } +#endif // COMPUTE_KERNEL_WRITER_ASSERTS_ENABLED + + return {static_cast<CLTile &>(tile), area}; +} + +void CLKernelWriter::op_load(const TileOperand &tile_op, + const TensorOperand &tensor_op, + TensorSampler &sampler, + const TileOperand &x, + const TileOperand &y, + const TileOperand &z, + const TileOperand &batch) +{ + const CLTile dilation_x({{"1"}}, DataType::Int32); + const CLTile dilation_y({{"1"}}, DataType::Int32); + + op_load_store(MemoryOperation::Load, tile_op, tensor_op, sampler, x, y, z, batch, dilation_x, dilation_y, + false /* indirect buffer */); +} + +void CLKernelWriter::op_load_dilated(const TileOperand &tile_op, + const TensorOperand &tensor_op, + TensorSampler &sampler, + const TileOperand &x, + const TileOperand &y, + const TileOperand &z, + const TileOperand &batch, + const TileOperand &dilation_x, + const TileOperand &dilation_y) +{ + const auto dil_x_view = to_cl_tile_view(dilation_x); + const auto dil_y_view = to_cl_tile_view(dilation_y); + + op_load_store(MemoryOperation::Load, tile_op, tensor_op, sampler, x, y, z, batch, dil_x_view, dil_y_view, + false /* indirect buffer */); +} + +void CLKernelWriter::op_store(const TensorOperand &tensor_op, + const TileOperand &tile_op, + TensorSampler &sampler, + const TileOperand &x, + const TileOperand &y, + const TileOperand &z, + const TileOperand &batch) +{ + const CLTile dilation_x({{"1"}}, DataType::Int32); + const CLTile dilation_y({{"1"}}, DataType::Int32); + + op_load_store(MemoryOperation::Store, tile_op, tensor_op, sampler, x, y, z, batch, dilation_x, dilation_y, + false /* indirect buffer */); +} + +void CLKernelWriter::op_store_dilated(const TensorOperand &tensor_op, + const TileOperand &tile_op, + TensorSampler &sampler, + const TileOperand &x, + const TileOperand &y, + const TileOperand &z, + const TileOperand &batch, + const TileOperand &dilation_x, + const TileOperand &dilation_y) +{ + const auto dil_x_view = to_cl_tile_view(dilation_x); + const auto dil_y_view = to_cl_tile_view(dilation_y); + + op_load_store(MemoryOperation::Store, tile_op, tensor_op, sampler, x, y, z, batch, dil_x_view, dil_y_view, + false /* indirect buffer */); +} + +void CLKernelWriter::op_load_indirect(const TileOperand &tile_op, + const TensorOperand &tensor_op, + TensorSampler &sampler, + const TileOperand &x, + const TileOperand &y, + const TileOperand &z, + const TileOperand &batch) +{ + const CLTile dilation_x({{"1"}}, DataType::Int32); + const CLTile dilation_y({{"1"}}, DataType::Int32); + + op_load_store(MemoryOperation::Load, tile_op, tensor_op, sampler, x, y, z, batch, dilation_x, dilation_y, + true /* indirect buffer */); +} + +void CLKernelWriter::op_load_store(MemoryOperation op, + const TileOperand &tile_op, + const TensorOperand &tensor_op, + TensorSampler &sampler, + const TileOperand &x, + const TileOperand &y, + const TileOperand &z, + const TileOperand &batch, + const TileView<CLTile> &dilation_x, + const TileView<CLTile> &dilation_y, + bool indirect_buffer) +{ + CKW_UNUSED(dilation_x); + CKW_ASSERT(dilation_x.is_scalar()); + CKW_ASSERT(dilation_y.is_scalar()); + CKW_ASSERT(dilation_x.scalar(0, 0).str == "((int)(1))"); // Dilation in x dimension is not implemented yet + + if (indirect_buffer) + { + CKW_ASSERT(dilation_y.scalar(0, 0).str == "((int)(1))" && dilation_x.scalar(0, 0).str == "((int)(1))"); + } + + ITensor &tensor = get_tensor(tensor_op); + + const auto tile = to_cl_tile_view(tile_op); + const auto x_tile = to_cl_tile_view(x).full_tile(); + const auto y_tile = to_cl_tile_view(y).full_tile(); + const auto z_tile = to_cl_tile_view(z).full_tile(); + const auto batch_tile = to_cl_tile_view(batch).full_tile(); + + std::unique_ptr<ICLMemoryOpHelper> helper; + switch (sampler.storage()) + { + case TensorStorageType::BufferUint8Ptr: + helper = std::make_unique<CLMemoryOpBufferHelper>(this, &tensor, &sampler, op, tile); + break; + case TensorStorageType::Texture2dReadOnly: + case TensorStorageType::Texture2dWriteOnly: + helper = std::make_unique<CLMemoryOpImage2dHelper>(this, &tensor, &sampler, op, tile); + break; + default: + CKW_THROW_MSG("Unsupported tensor storage"); + } + + CKW_ASSERT(x_tile.is_scalar()); + CKW_ASSERT(z_tile.is_scalar()); + CKW_ASSERT_IF(indirect_buffer, y_tile.info().width() == 1); + CKW_ASSERT_IF(!indirect_buffer, y_tile.is_scalar()); + CKW_ASSERT(batch_tile.is_scalar()); + + helper->initialize(&x_tile, &z_tile, &batch_tile); + + for (int row = 0; row < tile.height(); ++row) + { + if (!indirect_buffer) + { + std::string coord_y = y_tile.scalar(0, 0).str + " + " + std::to_string(row); + + if (dilation_y.scalar(0, 0).str != "((int)(1))") + { + coord_y += " * " + dilation_y.scalar(0, 0).str; + } + + helper->write_row(row, coord_y); + } + else + { + helper->write_row(row, y_tile.scalar(row, 0).str); + } + } + + helper->finalize(); +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/cl/CLKernelWriter.h b/compute_kernel_writer/src/cl/CLKernelWriter.h new file mode 100644 index 0000000000..6485bae512 --- /dev/null +++ b/compute_kernel_writer/src/cl/CLKernelWriter.h @@ -0,0 +1,261 @@ +/* + * 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_SRC_CL_CLKERNELWRITER_H +#define CKW_SRC_CL_CLKERNELWRITER_H + +#include "ckw/KernelWriter.h" + +#include "src/TileView.h" + +#include <memory> +#include <set> +#include <string> +#include <utility> + +namespace ckw +{ + +// Forward Declarations +class CLTile; +class CLTensorArgument; +class ConstantData; +class TensorOperand; +class TensorSampler; +class TileOperand; + +enum class DataType; +enum class MemoryOperation; + +/** OpenCL kernel writer. */ +class CLKernelWriter : public KernelWriter +{ +public: + // ============================================================================================= + // Construtors and destructor + // ============================================================================================= + + /** Initialize a new instance of @ref CLKernelWriter class. */ + CLKernelWriter(); + + /** Destructor */ + ~CLKernelWriter(); + + // ============================================================================================= + // Data processing + // ============================================================================================= + + void op_assign(const TileOperand &dst, const TileOperand &src) override; + + void op_cast(const TileOperand &dst, const TileOperand &src, ConvertPolicy policy) override; + + void op_unary(const TileOperand &dst, UnaryOp op, const TileOperand &src) override; + + void op_binary(const TileOperand &dst, BinaryOp op, const TileOperand &first, const TileOperand &second) override; + + void op_ternary(const TileOperand &dst, + TernaryOp op, + const TileOperand &first, + const TileOperand &second, + const TileOperand &third) override; + + // ============================================================================================= + // Flow control + // ============================================================================================= + + void op_if(const TileOperand &lhs, BinaryOp op, const TileOperand &rhs, const std::function<void()> &body) override; + + void + op_else_if(const TileOperand &lhs, BinaryOp op, const TileOperand &rhs, const std::function<void()> &body) override; + + void op_else(const std::function<void()> &body) override; + + void op_for_loop(const TileOperand &var, + BinaryOp cond_op, + const TileOperand &cond_value, + const TileOperand &update_var, + AssignmentOp update_op, + const TileOperand &update_value, + const std::function<void()> &body) override; + + void op_return() override; + + // ============================================================================================= + // Misc + // ============================================================================================= + + void op_get_global_id(const TileOperand &dst, int32_t dim) override; + + void op_comment(const std::string &text) override; + + void op_write_raw_code(const std::string &raw_code) override; + + void op_print(const std::string &prefix, const std::vector<TileOperand> &operands) override; + + // ============================================================================================= + // Code generation + // ============================================================================================= + + std::unique_ptr<Kernel> emit_kernel(const std::string &name) override; + + // ============================================================================================= + // Tensor and tile declaration + // ============================================================================================= + + TensorOperand declare_tensor_argument(const std::string &name, const TensorInfo &info) override; + + /** Declare a tile given name and tile information + * + * Similar to @ref KernelWriter::declare_tile() + */ + TileOperand declare_tile(const std::string &name, const TileInfo &tile_info) override; + + /** Declare a constant tile given a @ref:ConstantData object + * + * Similar to @ref KernelWriter::declare_constant_tile() + */ + TileOperand declare_constant_tile(const ConstantData &data) override; + + // ============================================================================================= + // Memory Operations + // ============================================================================================= + + void op_load(const TileOperand &tile_op, + const TensorOperand &tensor_op, + TensorSampler &sampler, + const TileOperand &x, + const TileOperand &y, + const TileOperand &z, + const TileOperand &batch) override; + + void op_load_dilated(const TileOperand &tile_op, + const TensorOperand &tensor_op, + TensorSampler &sampler, + const TileOperand &x, + const TileOperand &y, + const TileOperand &z, + const TileOperand &batch, + const TileOperand &dilation_x, + const TileOperand &dilation_y) override; + + void op_store(const TensorOperand &tensor_op, + const TileOperand &tile_op, + TensorSampler &sampler, + const TileOperand &x, + const TileOperand &y, + const TileOperand &z, + const TileOperand &batch) override; + + void op_store_dilated(const TensorOperand &tensor_op, + const TileOperand &tile_op, + TensorSampler &sampler, + const TileOperand &x, + const TileOperand &y, + const TileOperand &z, + const TileOperand &batch, + const TileOperand &dilation_x, + const TileOperand &dilation_y) override; + + void op_load_indirect(const TileOperand &tile_op, + const TensorOperand &tensor_op, + TensorSampler &sampler, + const TileOperand &x, + const TileOperand &y, + const TileOperand &z, + const TileOperand &batch) override; + +protected: + /** Return a tile view containing a reference to @ref CLTile object and the active area. + * + * This function performs appropriate check before doing type casting. + */ + TileView<CLTile> to_cl_tile_view(const TileOperand &operand) const; + + /** Append the specified code to the kernel body source code. */ + template <typename T, typename... TArgs> + void append_code(T &&code, TArgs &&...args) + { + append_code(std::forward<T>(code)); + append_code(std::forward<TArgs>(args)...); + } + + /** Append the specified code to the kernel body source code. */ + template <typename T> + void append_code(T &&code) + { + _body_source_code += std::forward<T>(code); + } + + /** Get the current kernel body source code. */ + const std::string &body_source_code() const; + + // For helper functions +private: + /** Helper method to consolidate all load/store logic in this class */ + void op_load_store(MemoryOperation op, + const TileOperand &tile_op, + const TensorOperand &tensor_op, + TensorSampler &sampler, + const TileOperand &x, + const TileOperand &y, + const TileOperand &z, + const TileOperand &batch, + const TileView<CLTile> &dilation_x, + const TileView<CLTile> &dilation_y, + bool indirect_buffer); + + /** This function is the generic function to write both `if` and `else if` blocks. + * + * It is used for both @ref CLKernelWriter::op_if and @ref CLKernelWriter::op_else_if. + * + * @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 function that writes the body of the else-if block. + * @param[in] is_else_if True if this is an `else if` block, otherwise this is an `if` block. + */ + void op_if_generic(const TileOperand &lhs, + BinaryOp op, + const TileOperand &rhs, + const std::function<void()> &body, + bool is_else_if); + + // For attributes +private: + /** This string contains the kernel body source code, not the full CL source code. + * The full source code will only be generated when the user calls @ref KernelWriter::emit_kernel. + * + * In order to add code to this, use @ref CLKernelWriter::append_code. + * Do not attempt to concatenate and alter this string directly. + */ + std::string _body_source_code{}; + + std::set<std::unique_ptr<CLTensorArgument>> _tensors{}; + std::set<std::unique_ptr<CLTile>> _tiles{}; + std::set<std::unique_ptr<CLTile>> _constant_tiles{}; +}; + +} // namespace ckw + +#endif // CKW_SRC_CL_CLKERNELWRITER_H diff --git a/compute_kernel_writer/src/cl/CLTensorArgument.cpp b/compute_kernel_writer/src/cl/CLTensorArgument.cpp new file mode 100644 index 0000000000..e53de2830d --- /dev/null +++ b/compute_kernel_writer/src/cl/CLTensorArgument.cpp @@ -0,0 +1,207 @@ +/* + * 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. + */ + +#include "src/cl/CLTensorArgument.h" + +#include "ckw/Error.h" + +#include "src/cl/CLHelpers.h" +#include "src/cl/CLTensorComponent.h" +#include "src/ITensorArgument.h" +#include "src/ITensorComponent.h" +#include "src/types/TensorComponentType.h" + +#include <algorithm> +#include <vector> + +namespace ckw +{ +CLTensorArgument::CLTensorArgument(const std::string &name, const TensorInfo &info, bool return_dims_by_value) +{ + _return_dims_by_value = return_dims_by_value; + _basename = name; + _info = info; +} + +CLTensorArgument::~CLTensorArgument() = default; + +CLTensorComponent &CLTensorArgument::cl_component(TensorComponentType x) +{ + // Return the component if it has already been created. + { + const auto it = + std::find_if(_components_used.begin(), _components_used.end(), + [=](const std::unique_ptr<CLTensorComponent> &item) { return item->component_type() == x; }); + + if (it != _components_used.end()) + { + return **it; + } + } + + if (_return_dims_by_value) + { + uint32_t component_type = static_cast<uint32_t>(x); + + const bool is_dimension = (component_type & static_cast<uint32_t>(TensorComponentBitmask::Dimension)) != 0; + const bool is_folded_dimensions = + (component_type & static_cast<uint32_t>(TensorComponentBitmask::FoldedDimensions)) != 0; + + constexpr auto bitmask_all = static_cast<uint32_t>(TensorComponentIndexBitmask::All); + constexpr auto bitmask_index_0 = static_cast<uint32_t>(TensorComponentIndexBitmask::Index0); +#ifdef COMPUTE_KERNEL_WRITER_ASSERTS_ENABLED + constexpr auto bitmask_index_1 = static_cast<uint32_t>(TensorComponentIndexBitmask::Index1); + constexpr auto bitmask_index_2 = static_cast<uint32_t>(TensorComponentIndexBitmask::Index2); + constexpr auto bitmask_index_3 = static_cast<uint32_t>(TensorComponentIndexBitmask::Index3); +#endif // COMPUTE_KERNEL_WRITER_ASSERTS_ENABLED + + // Make sure that the encoding of component type hasn't changed and each nibble is 4 bits apart. + CKW_ASSERT(bitmask_all == (bitmask_index_0 | bitmask_index_1 | bitmask_index_2 | bitmask_index_3)); + CKW_ASSERT(bitmask_index_0 == bitmask_index_1 >> 4); + CKW_ASSERT(bitmask_index_1 == bitmask_index_2 >> 4); + CKW_ASSERT(bitmask_index_2 == bitmask_index_3 >> 4); + + // If we have a dimension or folded dimensions, we can return the corresponding value if it is not dynamic (not equal to -1) + if (is_dimension == true || is_folded_dimensions == true) + { + component_type = component_type & bitmask_all; + + int32_t idx = 1; + for (int32_t i = 0; i < tensor_component_index_max_count; ++i) + { + uint32_t dim_idx = component_type & bitmask_index_0; + + if (dim_idx == 0) + { + // Stop at the first nibble containing 0 + break; + } + + // Subtract - 1. Please refer to the TensorComponentIndexBitmask documentation + dim_idx -= 1; + + // Get the dimension value + const int32_t dim_val = _info.shape()[dim_idx]; + + if (dim_val == kDynamicTensorDimensionValue) + { + // We cannot return the dimension by value if it is dynamic. + // Therefore, force the idx variable to kDynamicTensorDimensionValue and break the loop. + idx = kDynamicTensorDimensionValue; + break; + } + + idx *= dim_val; + + // Go to the next nibble + component_type >>= 4; + } + + if (idx != kDynamicTensorDimensionValue) + { + _components_used.emplace_back(std::make_unique<CLTensorComponent>(*this, x, idx)); + + return *_components_used.back(); + } + } + } + + _components_used.emplace_back(std::make_unique<CLTensorComponent>(*this, x)); + + return *_components_used.back(); +} + +ITile &CLTensorArgument::component(TensorComponentType x) +{ + return cl_component(x); +} + +TensorStorageVariable &CLTensorArgument::storage(TensorStorageType x) +{ + // Return the storage if it has already been created. + { + const auto it = std::find_if(_storages_used.begin(), _storages_used.end(), + [=](const TensorStorageVariable &item) { return item.type == x; }); + + if (it != _storages_used.end()) + { + return *it; + } + } + + TensorStorageVariable t; + t.val = create_storage_name(x); + t.type = x; + + _storages_used.emplace_back(t); + + return _storages_used.back(); +} + +std::string CLTensorArgument::create_storage_name(TensorStorageType x) const +{ + std::string var_name = _basename; + + switch (x) + { + case TensorStorageType::BufferUint8Ptr: + var_name += "_ptr"; + break; + case TensorStorageType::Texture2dReadOnly: + case TensorStorageType::Texture2dWriteOnly: + var_name += "_img2d"; + break; + default: + CKW_ASSERT_FAILED_MSG("Unsupported tensor storage"); + return ""; + } + + return var_name; +} + +std::vector<TensorStorageVariable> CLTensorArgument::storages() const +{ + std::vector<TensorStorageVariable> storages; + storages.reserve(_storages_used.size()); + + std::copy(_storages_used.begin(), _storages_used.end(), std::back_inserter(storages)); + + return storages; +} + +std::vector<const ITensorComponent *> CLTensorArgument::components() const +{ + std::vector<const ITensorComponent *> components; + + for (const auto &component : _components_used) + { + if (component->is_assignable()) + { + components.push_back(component.get()); + } + } + + return components; +} +} // namespace ckw diff --git a/compute_kernel_writer/src/cl/CLTensorArgument.h b/compute_kernel_writer/src/cl/CLTensorArgument.h new file mode 100644 index 0000000000..a79cf340bb --- /dev/null +++ b/compute_kernel_writer/src/cl/CLTensorArgument.h @@ -0,0 +1,89 @@ +/* + * 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_SRC_CL_CLTENSORARGUMENT_H +#define CKW_SRC_CL_CLTENSORARGUMENT_H + +#include "ckw/types/TensorComponentType.h" +#include "ckw/types/TensorStorageType.h" + +#include "src/cl/CLTensorComponent.h" +#include "src/ITensor.h" + +#include <memory> +#include <string> +#include <vector> + +namespace ckw +{ +// Forward declarations +class TensorInfo; + +class ITensorComponent; + +/** OpenCL specific tensor argument + * Internally, the object keeps track of the components and storages used to minimize the number + * of kernel arguments required. Therefore, if we create this object but we do not access any components + * or storages, the storages() and components() method will return an empty list. +*/ +class CLTensorArgument : public ITensor +{ +public: + /** Constructor + * + * @param[in] name Tensor name + * @param[in] info Tensor info + * @param[in] return_dims_by_value Flag to return the dimensions by value whenever it is possible. + * True, if the dimensions should be returned as value instead as variable. + */ + CLTensorArgument(const std::string &name, const TensorInfo &info, bool return_dims_by_value); + + /** Destructor. */ + ~CLTensorArgument(); + + /** Get a tensor component of the given type. + * + * This function is for internal use as it returns a reference to @ref CLTensorComponent object. + * It provides rich functionalities and doesn't require unnecessary casting + * unlike @ref CLTensorComponent::component which is for the public API and only returns + * a reference to a generic @ref ITile object. + */ + CLTensorComponent &cl_component(TensorComponentType component_type); + + // Inherited method overridden + TensorStorageVariable &storage(TensorStorageType x) override; + ITile &component(TensorComponentType x) override; + std::vector<TensorStorageVariable> storages() const override; + std::vector<const ITensorComponent *> components() const override; + +private: + std::string create_storage_name(TensorStorageType x) const; + + bool _return_dims_by_value{false}; + std::vector<TensorStorageVariable> _storages_used{}; + std::vector<std::unique_ptr<CLTensorComponent>> _components_used{}; +}; + +} // namespace ckw + +#endif // CKW_SRC_CL_CLTENSORARGUMENT_H diff --git a/compute_kernel_writer/src/cl/CLTensorComponent.cpp b/compute_kernel_writer/src/cl/CLTensorComponent.cpp new file mode 100644 index 0000000000..dbe2036768 --- /dev/null +++ b/compute_kernel_writer/src/cl/CLTensorComponent.cpp @@ -0,0 +1,126 @@ +/* + * 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. + */ + +#include "src/cl/CLTensorComponent.h" + +#include "ckw/Error.h" +#include "ckw/types/TensorComponentType.h" + +#include "src/cl/CLTensorArgument.h" +#include "src/cl/CLTile.h" + +namespace ckw +{ + +namespace +{ + +std::string create_component_name(const std::string &name, TensorComponentType x) +{ + std::string var_name(name); + + switch (x) + { + case TensorComponentType::OffsetFirstElement: + var_name += "_offset_first_element"; + break; + case TensorComponentType::Stride0: + var_name += "_stride0"; + break; + case TensorComponentType::Stride1: + var_name += "_stride1"; + break; + case TensorComponentType::Stride2: + var_name += "_stride2"; + break; + case TensorComponentType::Stride3: + var_name += "_stride3"; + break; + case TensorComponentType::Stride4: + var_name += "_stride4"; + break; + case TensorComponentType::Dim0: + var_name += "_dim0"; + break; + case TensorComponentType::Dim1: + var_name += "_dim1"; + break; + case TensorComponentType::Dim2: + var_name += "_dim2"; + break; + case TensorComponentType::Dim3: + var_name += "_dim3"; + break; + case TensorComponentType::Dim4: + var_name += "_dim4"; + break; + case TensorComponentType::Dim1xDim2: + var_name += "_dim1xdim2"; + break; + case TensorComponentType::Dim2xDim3: + var_name += "_dim2xdim3"; + break; + case TensorComponentType::Dim1xDim2xDim3: + var_name += "_dim1xdim2xdim3"; + break; + default: + CKW_THROW_MSG("Unsupported tensor component"); + return ""; + } + + return var_name; +} + +} // namespace + +CLTensorComponent::CLTensorComponent(const CLTensorArgument &tensor, TensorComponentType component_type) + : CLTile(create_component_name(tensor.name(), component_type), TileInfo(DataType::Int32)), + _component_type(component_type) +{ +} + +CLTensorComponent::CLTensorComponent(const CLTensorArgument &tensor, TensorComponentType component_type, int32_t value) + : CLTile({{std::to_string(value)}}, DataType::Int32), _component_type(component_type) +{ + CKW_UNUSED(tensor); +} + +CLTensorComponent::~CLTensorComponent() = default; + +ITile &CLTensorComponent::tile() +{ + return *this; +} + +const ITile &CLTensorComponent::tile() const +{ + return *this; +} + +TensorComponentType CLTensorComponent::component_type() const +{ + return _component_type; +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/cl/CLTensorComponent.h b/compute_kernel_writer/src/cl/CLTensorComponent.h new file mode 100644 index 0000000000..731597ebbf --- /dev/null +++ b/compute_kernel_writer/src/cl/CLTensorComponent.h @@ -0,0 +1,81 @@ +/* + * 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_SRC_CL_CLTENSORCOMPONENT_H +#define CKW_SRC_CL_CLTENSORCOMPONENT_H + +#include "ckw/types/TensorComponentType.h" + +#include "src/cl/CLTile.h" +#include "src/ITensorComponent.h" + +namespace ckw +{ + +class CLTensorArgument; + +/** A tensor component object that can be used as a tile. + * + * The tensor component is created by @ref CLTensorArgument object when it is used + * either by the user or internally by a kernel writer operation. + * It allows the user to perform operation on tensor component just like any other tile. + * + * Because of the nature of tensor component, it's always a scalar tile of 32-bit integer. + * + * To find the list of all tensor components, see @ref TensorComponentType. + */ +class CLTensorComponent : public CLTile, public ITensorComponent +{ +public: + /** Initialize a new instance of @ref CLTensorComponent class for dynamic component. + * + * @param[in] tensor The tensor to which this component belongs. + * @param[in] component_type The tensor component type. + */ + CLTensorComponent(const CLTensorArgument &tensor, TensorComponentType component_type); + + /** Initialize a new instance of @ref CLTensorComponent class for compile-time constant component. + * + * @param[in] tensor The tensor to which this component belongs. + * @param[in] component_type The tensor component type. + * @param[in] value The value of the component. + */ + CLTensorComponent(const CLTensorArgument &tensor, TensorComponentType component_type, int32_t value); + + /** Destructor. */ + virtual ~CLTensorComponent(); + + ITile &tile() override; + + const ITile &tile() const override; + + TensorComponentType component_type() const override; + +private: + TensorComponentType _component_type{TensorComponentType::Unknown}; +}; + +} // namespace ckw + +#endif // CKW_SRC_CL_CLTENSORCOMPONENT_H diff --git a/compute_kernel_writer/src/cl/CLTile.cpp b/compute_kernel_writer/src/cl/CLTile.cpp new file mode 100644 index 0000000000..f6e271e813 --- /dev/null +++ b/compute_kernel_writer/src/cl/CLTile.cpp @@ -0,0 +1,234 @@ +/* + * 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. + */ +#include "src/cl/CLTile.h" + +#include "ckw/Error.h" +#include "ckw/TileInfo.h" + +#include "src/cl/CLHelpers.h" +#include "src/Helpers.h" + +#include <algorithm> +#include <vector> + +namespace ckw +{ +CLTile::CLTile(const std::string &name, const TileInfo &info) : _is_constant(false) +{ + validate_tile_info(info); + + _basename = name; + _info = info; +} + +CLTile::CLTile(const TileContainer &vals, DataType dt) : _is_constant(true) +{ + const int32_t w = vals[0].size(); + const int32_t h = vals.size(); + + _info.width(w); + _info.height(h); + _info.data_type(dt); + + validate_tile_info(_info); + + _vals = TileContainer(h, std::vector<std::string>(w)); + + for (int32_t y = 0; y < h; ++y) + { + for (int32_t x = 0; x < w; ++x) + { + _vals[y][x] = vals[y][x]; + } + } +} + +const std::string &CLTile::name() const +{ + return _basename; +} + +const TileInfo &CLTile::info() const +{ + return _info; +} + +TileVariable CLTile::scalar(int32_t row, int32_t col) const +{ + // Clamp to nearest valid edge + col = clamp(col, static_cast<int32_t>(0), _info.width() - 1); + row = clamp(row, static_cast<int32_t>(0), _info.height() - 1); + + if (_is_constant) + { + // We can use the vector method to retrieve the scalar variable stored in the constant tile + return vector(row, col, 1); + } + else + { + TileVariable t; + t.str = create_var_name(row); + t.desc.dt = _info.data_type(); + t.desc.len = 1; + + // This check is required because if the width has only one element, we cannot use .s0 + if (_info.width() != 1) + { + // Automatic broadcasting + t.str += ".s" + dec_to_hex_as_string(col); + } + + return t; + } +} + +TileVariable CLTile::vector(int32_t row) const +{ + // Clamp to nearest valid edge + row = clamp(row, static_cast<int32_t>(0), _info.height() - 1); + + if (_is_constant) + { + return vector(row, 0, _info.width()); + } + else + { + TileVariable t; + t.str = create_var_name(row); + t.desc.dt = _info.data_type(); + t.desc.len = _info.width(); + return t; + } +} + +TileVariable CLTile::vector(int32_t row, int32_t col_start, int32_t width) const +{ + CKW_ASSERT(col_start >= 0 && col_start < _info.width()); + CKW_ASSERT(col_start + width <= _info.width()); + + // Validate the new vector length + cl_validate_vector_length(width); + + // Clamp to nearest valid edge + row = clamp(row, static_cast<int32_t>(0), _info.height() - 1); + + TileVariable t; + t.desc.dt = _info.data_type(); + t.desc.len = width; + + if (_is_constant) + { + // The vector has the following form: ((data_typeN)(val0, val1,..., ValN-1)) + t.str = "((" + cl_get_variable_datatype_as_string(t.desc.dt, width) + ")"; + t.str += "("; + + int32_t col = col_start; + for (; col < width - 1; ++col) + { + t.str += _vals[row][col]; + t.str += ", "; + } + t.str += _vals[row][col]; + t.str += "))"; + } + else + { + t.str = create_var_name(row); + + if (_info.width() != 1 && _info.width() != width) + { + t.str += ".s"; + for (int i = 0; i < width; ++i) + { + t.str += dec_to_hex_as_string(col_start + i); + } + } + } + + return t; +} + +std::vector<TileVariable> CLTile::all() const +{ + std::vector<TileVariable> vars; + + if (_is_constant) + { + for (int32_t y = 0; y < _info.height(); ++y) + { + for (int32_t x = 0; x < _info.width(); ++x) + { + // We can use the vector method to retrieve all the scalar variables stored in the constant tile + TileVariable t = vector(y, x, 1); + vars.push_back(t); + } + } + } + else + { + for (int32_t y = 0; y < _info.height(); ++y) + { + TileVariable t; + t.str = create_var_name(y); + t.desc.dt = _info.data_type(); + t.desc.len = _info.width(); + vars.push_back(t); + } + } + + return vars; +} + +bool CLTile::is_assignable() const +{ + return !_is_constant; +} + +std::string CLTile::create_var_name(int32_t row) const +{ + std::string var_name = _basename; + + // If a scalar variable, we do not append the row index + if (_info.height() > 1) + { + var_name += "__"; + var_name += std::to_string(row); + } + + return var_name; +} + +std::vector<int32_t> CLTile::supported_vector_lengths() const +{ + return std::vector<int32_t>{1, 2, 3, 4, 8, 16}; +} + +void CLTile::validate_tile_info(const TileInfo &info) const +{ + CKW_UNUSED(info); + CKW_ASSERT_MSG(cl_validate_vector_length(info.width()), "Unsupported TileInfo width"); + CKW_ASSERT_MSG(info.data_type() != DataType::Unknown, "DataType::Unknown is not supported"); +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/cl/CLTile.h b/compute_kernel_writer/src/cl/CLTile.h new file mode 100644 index 0000000000..498cf51034 --- /dev/null +++ b/compute_kernel_writer/src/cl/CLTile.h @@ -0,0 +1,86 @@ +/* + * 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 COMPUTE_KERNEL_WRITER_SRC_CL_CLTILE_H +#define COMPUTE_KERNEL_WRITER_SRC_CL_CLTILE_H + +#include "src/ITile.h" + +#include <string> + +namespace ckw +{ +// Forward declarations +class TileInfo; + +/** OpenCL specific tile */ +class CLTile : public ITile, public IVectorAccess +{ +public: + /** Initialize a new instance of @ref CLTile class for variable tile. + * + * @param[in] name Tile name + * @param[in] info Tile info + */ + CLTile(const std::string &name, const TileInfo &info); + + /** Initialize a new instane of @ref CLTile class for compile-time constant tile. + * + * @note A constant tile does not need a name since this object does not return variable's name but rather + * values stored as string type + * + * @param[in] vals The tile container with the constant values as std::string + * @param[in] dt Datatype of the values stored in the tile container + */ + CLTile(const TileContainer &vals, DataType dt); + + // Inherited method overridden + const std::string &name() const override; + + const TileInfo &info() const override; + + TileVariable scalar(int32_t row, int32_t col) const override; + + TileVariable vector(int32_t row) const override; + + TileVariable vector(int32_t row, int32_t col_start, int32_t width) const override; + + std::vector<TileVariable> all() const override; + + bool is_assignable() const override; + + std::vector<int32_t> supported_vector_lengths() const override; + +private: + void validate_tile_info(const TileInfo &info) const; + + std::string create_var_name(int32_t row) const; + + TileInfo _info{DataType::Unknown}; + std::string _basename{""}; + bool _is_constant{false}; + TileContainer _vals{}; +}; +} // namespace ckw + +#endif /* COMPUTE_KERNEL_WRITER_SRC_CL_CLTILE_H */ diff --git a/compute_kernel_writer/src/cl/helpers/CLMemoryOpBufferHelper.cpp b/compute_kernel_writer/src/cl/helpers/CLMemoryOpBufferHelper.cpp new file mode 100644 index 0000000000..7d16f35fbe --- /dev/null +++ b/compute_kernel_writer/src/cl/helpers/CLMemoryOpBufferHelper.cpp @@ -0,0 +1,353 @@ +/* + * 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. + */ +#include "src/cl/helpers/CLMemoryOpBufferHelper.h" + +#include "ckw/Error.h" +#include "ckw/TensorSampler.h" +#include "ckw/types/MemoryOperation.h" +#include "ckw/types/TensorStorageType.h" + +#include "src/cl/CLHelpers.h" +#include "src/cl/CLKernelWriter.h" +#include "src/cl/CLTensorArgument.h" +#include "src/cl/CLTile.h" +#include "src/ITensor.h" +#include "src/Tensor3dMapper.h" +#include "src/TileView.h" + +namespace ckw +{ +bool CLMemoryOpBufferHelper::validate(const CLKernelWriter *writer, + const ITensor *tensor, + const TensorSampler *sampler, + const Tensor3dMapper *mapper, + MemoryOperation op, + const TileView<CLTile> &dst) +{ + CKW_UNUSED(writer, tensor, mapper, op, dst); + + if (sampler->storage() != TensorStorageType::BufferUint8Ptr) + { + return false; + } + return true; +} + +/** Initialization and Finalizing Logic + * + * The meanings of if/elses in different dimensions and how they're constructed: + * - x: partial load/store + * - y: no load/store operation + * - z: no load/store operation + * if(x) + * { + * if(z) + * { + * if(y) + * { + * // full load/store width + * } + * else + * { + * // no load/store + * } + * } + * else + * { + * // no load/store + * } + * } + * else + * { + * if(z) + * { + * if(y) + * { + * // partial load/store width + * } + * else + * { + * // no load/store + * } + * } + * else + * { + * // no load/store + * } + * } + * + * In general, initialize() writes if conditions, and finalize() writes else conditions. + * The outermost block is x, then z and then y. This is why, if/else's covering for y are initialized + * at each row write. In some addressing modes, such as None, no if/else conditions are written. + */ +void CLMemoryOpBufferHelper::initialize(const CLTile *x, const CLTile *z, const CLTile *b) +{ + CKW_ASSERT(validate(_writer, _tensor, _sampler, _mapper.get(), _op, _dst)); + + _coord_x = x->scalar(0, 0).str; + _coord_z = z->scalar(0, 0).str; + _coord_b = b->scalar(0, 0).str; + _coord_orig_z = _coord_z; + + out_of_bound_initialize_x(_coord_x); + out_of_bound_initialize_z(_coord_z); +} + +void CLMemoryOpBufferHelper::write_row(int32_t row_id, const std::string &coord_y) +{ + // The only check required is on Y. + out_of_bound_initialize_y(coord_y); + + const std::string dst = _dst.vector(row_id).str; + const std::string address = to_buffer_address(_coord_x, coord_y, _coord_z, _coord_b); + const std::string ls_buf = to_statement(_op, _ls_width_full, dst, address); + + _writer->op_write_raw_code(ls_buf); + _writer->op_write_raw_code(";\n"); + + out_of_bound_finalize_y(dst); + + // The left over load/store will be written in the finalize stage + if (_ls_width_part.size() != 0) + { + int32_t col_start = 0; + const TileArea original_area = _dst.area(); + + for (int32_t partial_width : _ls_width_part) + { + // Set the active area + const TileArea area(original_area.row_start(), original_area.row_end(), col_start, + col_start + partial_width); + _dst.area(area); + + const std::string dst = _dst.vector(row_id).str; + const std::string coord_x = _coord_x + " + " + std::to_string(col_start); + const std::string address = to_buffer_address(coord_x, coord_y, _coord_z, _coord_b); + const std::string statement = to_statement(_op, partial_width, dst, address); + _leftovers_x.emplace_back(dst, coord_y, statement); + + col_start += partial_width; + } + // Restore the original area + _dst.area(original_area); + } +} + +void CLMemoryOpBufferHelper::finalize() +{ + out_of_bound_finalize_z(); + out_of_bound_finalize_x(); +} + +void CLMemoryOpBufferHelper::out_of_bound_initialize_x(const std::string &coord) +{ + if (_sampler->address_mode_x() == TensorSamplerAddressModeX::OverlappingMin) + { + TensorInfo tensor_info = _tensor->info(); + TensorShape shape = tensor_info.shape(); + + _ls_width_part = cl_decompose_vector_width(shape[0] % _ls_width_full); + if (_ls_width_part.size() != 0) + { + _writer->op_write_raw_code("if(" + coord + " > 0)\n{\n"); + } + } +} + +void CLMemoryOpBufferHelper::out_of_bound_finalize_x() +{ + if (_sampler->address_mode_x() == TensorSamplerAddressModeX::OverlappingMin) + { + if (_ls_width_part.size() != 0) + { + _writer->op_write_raw_code("}\nelse\n{\n"); + + out_of_bound_initialize_z(_coord_orig_z); + for (LeftoverDescriptor leftover_desc : _leftovers_x) + { + out_of_bound_initialize_y(leftover_desc.coord); + _writer->op_write_raw_code(leftover_desc.statement); + _writer->op_write_raw_code(";\n"); + out_of_bound_finalize_y(leftover_desc.dst); + } + out_of_bound_finalize_z(); + _writer->op_write_raw_code("}\n"); + } + } +} + +void CLMemoryOpBufferHelper::out_of_bound_initialize_y(const std::string &coord) +{ + std::string max = ""; + + const TensorSamplerAddressModeY address_mode_y = _sampler->address_mode_y(); + + switch (address_mode_y) + { + case TensorSamplerAddressModeY::ClampToBorderMaxOnly: + // Not to be moved outside the case because it marks the relevant tensor component as used even if we dont't use the variable + max = _mapper->dim_y().str; + _writer->op_write_raw_code("if(" + coord + " < " + max + ")\n{\n"); + break; + case TensorSamplerAddressModeY::SkipLessThanZero: + _writer->op_write_raw_code("if(" + coord + " >= 0)\n{\n"); + break; + case TensorSamplerAddressModeY::None: + break; + default: + CKW_THROW_MSG("Unsupported address mode for Y dimension"); + } +} + +void CLMemoryOpBufferHelper::out_of_bound_finalize_y(const std::string &dst) +{ + const TensorSamplerAddressModeY address_mode_y = _sampler->address_mode_y(); + + switch (address_mode_y) + { + case TensorSamplerAddressModeY::ClampToBorderMaxOnly: + _writer->op_write_raw_code("}\nelse\n{\n"); + _writer->op_write_raw_code(dst); + _writer->op_write_raw_code(" = 0.0f;\n}\n"); + break; + case TensorSamplerAddressModeY::SkipLessThanZero: + _writer->op_write_raw_code("}\n"); + break; + case TensorSamplerAddressModeY::None: + break; + default: + CKW_THROW_MSG("Unsupported address mode for Y dimension"); + } +} + +void CLMemoryOpBufferHelper::out_of_bound_initialize_z(const std::string &coord) +{ + CKW_UNUSED(coord); + + const TensorSamplerAddressModeZ address_mode_z = _sampler->address_mode_z(); + switch (address_mode_z) + { + case TensorSamplerAddressModeZ::None: + break; + default: + CKW_THROW_MSG("Unsupported address mode for Z dimension"); + } +} + +void CLMemoryOpBufferHelper::out_of_bound_finalize_z() +{ + const TensorSamplerAddressModeZ address_mode_z = _sampler->address_mode_z(); + + switch (address_mode_z) + { + case TensorSamplerAddressModeZ::None: + break; + default: + CKW_THROW_MSG("Unsupported address mode for Z dimension"); + } +} + +std::string CLMemoryOpBufferHelper::to_statement(MemoryOperation op, + int32_t vector_width, + const std::string &data, + const std::string &address) const +{ + switch (op) + { + case MemoryOperation::Load: + if (vector_width != 1) + { + return data + " = vload" + std::to_string(vector_width) + "(0, " + address + ")"; + } + else + { + return data + " = *(" + address + ")"; + } + break; + case MemoryOperation::Store: + if (vector_width != 1) + { + return "vstore" + std::to_string(vector_width) + "(" + data + ", 0, " + address + ")"; + } + else + { + return "*(" + address + ") = " + data; + } + break; + default: + CKW_THROW_MSG("Unsupported MemoryOperation"); + } + + return ""; +} + +std::string CLMemoryOpBufferHelper::to_buffer_address(const std::string &x, + const std::string &y, + const std::string &z, + const std::string &b) const +{ + TensorStorageType tensor_storage = _sampler->storage(); + CKW_ASSERT(tensor_storage == TensorStorageType::BufferUint8Ptr); + + const std::string ptr_buf = _tensor->storage(tensor_storage).val; + const std::string dst_type = cl_data_type_rounded_up_to_valid_vector_width(_dst.data_type(), 1); + + std::string address; + address += "(__global "; + address += dst_type; + address += "*)("; + address += ptr_buf; + if (x != "0" && (_mapper->dim_x().str != "1")) + { + address += " + ("; + address += x + ") * sizeof(" + dst_type + ")"; + } + if (y != "0") + { + const std::string stride_y = _mapper->stride_y().str; + address += " + ("; + address += y + ")"; + address += " * "; + address += stride_y; + } + if (z != "0" && (_mapper->dim_z().str != "1")) + { + const std::string stride_z = _mapper->stride_z().str; + address += " + ("; + address += z + ")"; + address += " * "; + address += stride_z; + } + if (b != "0" && (_mapper->dim_batch().str != "1")) + { + const std::string stride_b = _mapper->stride_batch().str; + address += " + ("; + address += b + ")"; + address += " * "; + address += stride_b; + } + address += ")"; + return address; +} +} // namespace ckw diff --git a/compute_kernel_writer/src/cl/helpers/CLMemoryOpBufferHelper.h b/compute_kernel_writer/src/cl/helpers/CLMemoryOpBufferHelper.h new file mode 100644 index 0000000000..a6b3272f32 --- /dev/null +++ b/compute_kernel_writer/src/cl/helpers/CLMemoryOpBufferHelper.h @@ -0,0 +1,108 @@ +/* + * 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_SRC_CL_HELPERS_CLMEMORYOPBUFFERHELPER_H +#define CKW_SRC_CL_HELPERS_CLMEMORYOPBUFFERHELPER_H + +#include "src/cl/helpers/ICLMemoryOpHelper.h" + +#include <cstdint> +#include <string> +#include <vector> + +namespace ckw +{ + +// Forward Declarations +class CLKernelWriter; +class CLTile; +template <class CLTile> +class TileView; +enum class MemoryOperation; + +/** Helper class to write memory operations (like load/store) in OpenCL + */ +class CLMemoryOpBufferHelper : public ICLMemoryOpHelper +{ +public: + /** Constructor similar to @ref ICLMemoryOpHelper() */ + CLMemoryOpBufferHelper(CLKernelWriter *writer, + ITensor *tensor, + TensorSampler *sampler, + MemoryOperation op, + const TileView<CLTile> &dst) + : ICLMemoryOpHelper(writer, tensor, sampler, op, dst) + { + } + + /** Copy constructor */ + CLMemoryOpBufferHelper(const CLMemoryOpBufferHelper &) = delete; + + /** Assignment operator overload */ + CLMemoryOpBufferHelper &operator=(const CLMemoryOpBufferHelper &) = delete; + + // Methods overridden + void initialize(const CLTile *x, const CLTile *z, const CLTile *b) override; + void write_row(int32_t row_id, const std::string &coord_y) override; + void finalize() override; + +private: + struct LeftoverDescriptor + { + LeftoverDescriptor(const std::string &dst, const std::string &coord, const std::string &statement) + : dst(dst), coord(coord), statement(statement) + { + } + + std::string dst{}; // Describes the destination tile or part of it + std::string coord{}; // Describes the coordinate to be used in boundary checks + std::string statement{}; // Describes the memory operation statement + }; + + std::vector<int32_t> _ls_width_part{}; + std::vector<LeftoverDescriptor> _leftovers_x{}; + std::string _coord_orig_z{}; + + static bool validate(const CLKernelWriter *writer, + const ITensor *tensor, + const TensorSampler *sampler, + const Tensor3dMapper *mapper, + MemoryOperation op, + const TileView<CLTile> &dst); + + void out_of_bound_initialize_x(const std::string &coord); + void out_of_bound_finalize_x(); + void out_of_bound_initialize_y(const std::string &coord); + void out_of_bound_finalize_y(const std::string &dst); + void out_of_bound_initialize_z(const std::string &coord); + void out_of_bound_finalize_z(); + + std::string + to_statement(MemoryOperation op, int32_t vector_width, const std::string &data, const std::string &address) const; + std::string + to_buffer_address(const std::string &x, const std::string &y, const std::string &z, const std::string &b) const; +}; +} // namespace ckw + +#endif // CKW_SRC_CL_HELPERS_CLMEMORYOPBUFFERHELPER_H diff --git a/compute_kernel_writer/src/cl/helpers/CLMemoryOpImage2dHelper.cpp b/compute_kernel_writer/src/cl/helpers/CLMemoryOpImage2dHelper.cpp new file mode 100644 index 0000000000..f392cd89cc --- /dev/null +++ b/compute_kernel_writer/src/cl/helpers/CLMemoryOpImage2dHelper.cpp @@ -0,0 +1,213 @@ +/* + * 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. + */ +#include "src/cl/helpers/CLMemoryOpImage2dHelper.h" + +#include "ckw/Error.h" +#include "ckw/TensorSampler.h" +#include "ckw/types/MemoryOperation.h" +#include "ckw/types/TensorStorageType.h" + +#include "src/cl/CLKernelWriter.h" +#include "src/cl/CLTensorArgument.h" +#include "src/cl/CLTile.h" +#include "src/ITensor.h" +#include "src/Tensor3dMapper.h" +#include "src/TileView.h" + +namespace ckw +{ +void CLMemoryOpImage2dHelper::initialize(const CLTile *x, const CLTile *z, const CLTile *b) +{ + _coord_x = x->scalar(0, 0).str; + _coord_z = z->scalar(0, 0).str; + _coord_b = b->scalar(0, 0).str; +} + +void CLMemoryOpImage2dHelper::write_row(int32_t row_id, const std::string &coord_y) +{ + // The only check required is on Y. + out_of_bound_initialize_y(coord_y); + + const std::string dst = _dst.vector(row_id).str; + const std::string sampler = to_ls_image2d_sampler(); + const std::string coord = to_ls_image2d_address(_coord_x, coord_y, _coord_z, _coord_b); + const std::string ls_buf = to_ls_image2d(_op, _ls_width_full, dst, sampler, coord); + + _writer->op_write_raw_code(ls_buf + ";\n"); + + out_of_bound_finalize_y(); +} + +void CLMemoryOpImage2dHelper::finalize() +{ +} + +bool CLMemoryOpImage2dHelper::validate(const CLKernelWriter *writer, + const ITensor *tensor, + const TensorSampler *sampler, + const Tensor3dMapper *mapper, + MemoryOperation op, + const TileView<CLTile> &dst) +{ + CKW_UNUSED(writer, tensor, mapper); + + if (dst.width() != 4) + { + return false; + } + if (sampler->address_mode_x() != TensorSamplerAddressModeX::None) + { + return false; + } + if (sampler->address_mode_z() != TensorSamplerAddressModeZ::None) + { + return false; + } + if (sampler->storage() != TensorStorageType::Texture2dReadOnly && op == MemoryOperation::Load) + { + return false; + } + if (sampler->storage() != TensorStorageType::Texture2dWriteOnly && op == MemoryOperation::Store) + { + return false; + } + if ((dst.data_type() != DataType::Fp32) && (dst.data_type() != DataType::Fp16)) + { + return false; + } + return true; +} + +void CLMemoryOpImage2dHelper::out_of_bound_initialize_y(const std::string &coord) +{ + CKW_UNUSED(coord); + + const TensorSamplerAddressModeY address_mode_y = _sampler->address_mode_y(); + switch (address_mode_y) + { + case TensorSamplerAddressModeY::SkipLessThanZero: + _writer->op_write_raw_code("if(" + coord + " >= 0)\n{\n"); + break; + case TensorSamplerAddressModeY::ClampToBorderMaxOnly: + case TensorSamplerAddressModeY::None: + break; + default: + CKW_THROW_MSG("Unsupported address mode for Y dimension"); + } +} + +void CLMemoryOpImage2dHelper::out_of_bound_finalize_y() +{ + const TensorSamplerAddressModeY address_mode_y = _sampler->address_mode_y(); + switch (address_mode_y) + { + case TensorSamplerAddressModeY::SkipLessThanZero: + _writer->op_write_raw_code("}\n"); + break; + case TensorSamplerAddressModeY::ClampToBorderMaxOnly: + case TensorSamplerAddressModeY::None: + break; + default: + CKW_THROW_MSG("Unsupported address mode for Y dimension"); + } +} + +std::string CLMemoryOpImage2dHelper::to_ls_image2d(MemoryOperation op, + int32_t vector_width, + const std::string &data, + const std::string &sampler, + const std::string &address) const +{ + CKW_UNUSED(vector_width); + CKW_ASSERT_MSG(_dst.data_type() == DataType::Fp32 || _dst.data_type() == DataType::Fp16, + "Image2d only supports floating-point data type"); + + const TensorStorageType tensor_storage = _sampler->storage(); + const std::string image2d_obj = _tensor->storage(tensor_storage).val; + const std::string post_fix = _dst.data_type() == DataType::Fp32 ? "f" : "h"; + + switch (op) + { + case MemoryOperation::Load: + return data + " = read_image" + post_fix + "(" + image2d_obj + ", " + sampler + ", " + address + ")"; + break; + case MemoryOperation::Store: + return "write_image" + post_fix + "(" + image2d_obj + ", " + address + ", " + data + ")"; + default: + CKW_THROW_MSG("Unsupported MemoryOperation"); + } +} + +std::string CLMemoryOpImage2dHelper::to_ls_image2d_sampler() const +{ + const auto address_mode_y = _sampler->address_mode_y(); + + switch (address_mode_y) + { + case TensorSamplerAddressModeY::None: + return "CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST"; + case TensorSamplerAddressModeY::SkipLessThanZero: + case TensorSamplerAddressModeY::ClampToBorderMaxOnly: + return "CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST"; + default: + CKW_THROW_MSG("Unsupported address_mode_coord"); + } +} + +std::string CLMemoryOpImage2dHelper::to_ls_image2d_address(const std::string &x, + const std::string &y, + const std::string &z, + const std::string &b) const +{ + std::string coord_x = "(" + x + ") >> 2"; + std::string coord_y = "("; + + if (y != "0") + { + coord_y += y; + } + if (z != "0" && (_mapper->dim_z().str != "1")) + { + const std::string dim = _mapper->dim_y().str; + coord_y += " + ("; + coord_y += z + ")"; + coord_y += " * "; + coord_y += dim; + } + if (b != "0" && (_mapper->dim_batch().str != "1")) + { + const std::string dim0 = _mapper->dim_y().str; + const std::string dim1 = _mapper->dim_z().str; + coord_y += " + ("; + coord_y += b + ")"; + coord_y += " * "; + coord_y += dim0; + coord_y += " * "; + coord_y += dim1; + } + coord_y += ")"; + return "(int2)(" + coord_x + ", " + coord_y + ")"; +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/cl/helpers/CLMemoryOpImage2dHelper.h b/compute_kernel_writer/src/cl/helpers/CLMemoryOpImage2dHelper.h new file mode 100644 index 0000000000..6c42c132d9 --- /dev/null +++ b/compute_kernel_writer/src/cl/helpers/CLMemoryOpImage2dHelper.h @@ -0,0 +1,89 @@ +/* + * 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_SRC_CL_HELPERS_CLMEMORYOPIMAGE2DHELPER_H +#define CKW_SRC_CL_HELPERS_CLMEMORYOPIMAGE2DHELPER_H + +#include "src/cl/helpers/ICLMemoryOpHelper.h" + +#include <string> + +namespace ckw +{ + +// Forward Declarations +class CLKernelWriter; +class CLTile; +template <class CLTile> +class TileView; +enum class MemoryOperation; + +/** Helper class to write memory operations (like load/store) in OpenCL for Image2d type */ +class CLMemoryOpImage2dHelper : public ICLMemoryOpHelper +{ +public: + /** Constructor similar to @ref ICLMemoryOpHelper() */ + CLMemoryOpImage2dHelper(CLKernelWriter *writer, + ITensor *tensor, + TensorSampler *sampler, + MemoryOperation op, + const TileView<CLTile> &dst) + : ICLMemoryOpHelper(writer, tensor, sampler, op, dst) + { + } + + /** Copy constructor */ + CLMemoryOpImage2dHelper(const CLMemoryOpImage2dHelper &) = delete; + + /** Assignment operator overload */ + CLMemoryOpImage2dHelper &operator=(const CLMemoryOpImage2dHelper &) = delete; + + // Methods overridden + void initialize(const CLTile *x, const CLTile *z, const CLTile *b) override; + void write_row(int32_t row_id, const std::string &coord_y) override; + void finalize() override; + +private: + static bool validate(const CLKernelWriter *writer, + const ITensor *tensor, + const TensorSampler *sampler, + const Tensor3dMapper *mapper, + MemoryOperation op, + const TileView<CLTile> &dst); + + void out_of_bound_initialize_y(const std::string &coord); + void out_of_bound_finalize_y(); + + std::string to_ls_image2d(MemoryOperation op, + int32_t vector_width, + const std::string &data, + const std::string &sampler, + const std::string &address) const; + std::string to_ls_image2d_sampler() const; + std::string + to_ls_image2d_address(const std::string &x, const std::string &y, const std::string &z, const std::string &b) const; +}; +} // namespace ckw + +#endif // CKW_SRC_CL_HELPERS_CLMEMORYOPIMAGE2DHELPER_H diff --git a/compute_kernel_writer/src/cl/helpers/ICLMemoryOpHelper.h b/compute_kernel_writer/src/cl/helpers/ICLMemoryOpHelper.h new file mode 100644 index 0000000000..a5b679ac03 --- /dev/null +++ b/compute_kernel_writer/src/cl/helpers/ICLMemoryOpHelper.h @@ -0,0 +1,121 @@ +/* + * 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_SRC_CL_HELPERS_ICLMEMORYOPHELPER_H +#define CKW_SRC_CL_HELPERS_ICLMEMORYOPHELPER_H + +#include "ckw/TensorSampler.h" + +#include "src/Tensor3dMapper.h" +#include "src/TileView.h" + +#include <cstdint> +#include <memory> +#include <string> + +namespace ckw +{ + +// Forward Declarations +class CLTile; +class CLKernelWriter; +class ITensor; +class TensorSampler; +enum class MemoryOperation; + +/** Base class OpenCL memory operation helper classes + * that helps writing code for memory operations like load/store. + */ +class ICLMemoryOpHelper +{ +public: + /** Constructor + * + * @param[in] writer @ref ckw::CLKernelWriter object to write the code + * @param[in] tensor @ref ckw::ITensor object to perform the memory operation on + * @param[in] sampler @ref ckw::TensorSampler object that tells how to sample a tensor + * @param[in] op The memory operation to be done (e.g. Load/Store) + * @param[in] dst The tile to perform the memory operation on + */ + ICLMemoryOpHelper(CLKernelWriter *writer, + ITensor *tensor, + TensorSampler *sampler, + MemoryOperation op, + const TileView<CLTile> &dst) + : _writer(writer), _tensor(tensor), _sampler(sampler), _op(op), _dst(dst) + { + _mapper = std::make_unique<Tensor3dMapper>(tensor, sampler->format()); + _ls_width_full = _dst.width(); + } + + /** Copy constructor */ + ICLMemoryOpHelper(const ICLMemoryOpHelper &) = delete; + + /** Assignment operator overload */ + ICLMemoryOpHelper &operator=(const ICLMemoryOpHelper &) = delete; + + /** Destructor */ + virtual ~ICLMemoryOpHelper() = default; + + /** Initialization method that takes a 3D tensor's x, z dimensions and + * the batch offset as a tile object, and initializes the code inside + * the writer object. + * + * @param[in] x tile object that describes the x-coordinate of the tensor involved + * @param[in] z tile object that describes the z-coordinate of the tensor involved + * @param[in] b tile object that describes the batch offset of the tensor involved + */ + virtual void initialize(const CLTile *x, const CLTile *z, const CLTile *b) = 0; + + /** Method that writes the actual code to the writer that performs the mentioned memory + * operation on the tile initialized. It writes the code for a specific row given in the + * arguments. + * + * @param[in] row_id row id + * @param[in] coord_y y-coordinate as string + */ + virtual void write_row(int32_t row_id, const std::string &coord_y) = 0; + + /** Method that finalizes the code in the writer object. This part is usually for taking + * care of finalizing anything that's been initialized inside @ref IMemoryHelper::initialize() + * such as matching compound statements, checking certain boundary conditions etc. No inputs + * and/or outputs, only the writer object is affected. + */ + virtual void finalize() = 0; + +protected: + CLKernelWriter *_writer{nullptr}; + ITensor *_tensor{nullptr}; + TensorSampler *_sampler{nullptr}; + MemoryOperation _op; + std::unique_ptr<Tensor3dMapper> _mapper{nullptr}; + TileView<CLTile> _dst{}; + int32_t _ls_width_full{0}; + std::string _coord_x{}; + std::string _coord_z{}; + std::string _coord_b{}; +}; +} // namespace ckw + +#endif // CKW_SRC_CL_HELPERS_ICLMEMORYOPHELPER_H diff --git a/compute_kernel_writer/src/types/ConstantData.cpp b/compute_kernel_writer/src/types/ConstantData.cpp new file mode 100644 index 0000000000..6d15eab407 --- /dev/null +++ b/compute_kernel_writer/src/types/ConstantData.cpp @@ -0,0 +1,141 @@ +/* + * 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. + */ + +#include "ckw/types/ConstantData.h" + +#include <limits> + +namespace ckw +{ +namespace +{ +template <typename T> +typename std::enable_if<std::is_same<T, float>::value, std::string>::type to_str(T value) +{ + std::stringstream ss; + ss << std::scientific << std::setprecision(std::numeric_limits<T>::max_digits10) << value; + return ss.str(); +} + +template <typename T> +typename std::enable_if<!std::is_same<T, float>::value && !std::is_same<T, bool>::value, std::string>::type +to_str(T value) +{ + return std::to_string(value); +} + +template <typename T> +typename std::enable_if<std::is_same<T, bool>::value, std::string>::type to_str(T value) +{ + return std::to_string((int)value); +} +} // namespace + +template <typename T> +ConstantData::ConstantData(std::initializer_list<std::initializer_list<T>> values, DataType data_type) + : _data_type(data_type) +{ + CKW_ASSERT(validate<T>(data_type)); + CKW_ASSERT(values.size() > 0); + + for (auto value_arr : values) + { + // Each row must have the same number of elements + CKW_ASSERT(value_arr.size() == (*values.begin()).size()); + + StringVector vec; + std::transform(value_arr.begin(), value_arr.end(), std::back_inserter(vec), [](T val) { return to_str(val); }); + + _values.push_back(std::move(vec)); + } +} + +template <typename T> +ConstantData::ConstantData(const std::vector<std::vector<T>> &values, DataType data_type) : _data_type(data_type) +{ + CKW_ASSERT(validate<T>(data_type)); + CKW_ASSERT(values.size() > 0); + + for (auto value_arr : values) + { + // Each row must have the same number of elements + CKW_ASSERT(value_arr.size() == (*values.begin()).size()); + + StringVector vec; + std::transform(value_arr.begin(), value_arr.end(), std::back_inserter(vec), [](T val) { return to_str(val); }); + + _values.push_back(std::move(vec)); + } +} + +template <typename T> +bool ConstantData::validate(DataType data_type) +{ + switch (data_type) + { + case DataType::Fp32: + case DataType::Fp16: + return std::is_same<T, float>::value; + case DataType::Bool: + return std::is_same<T, bool>::value; + case DataType::Int32: + case DataType::Int16: + case DataType::Int8: + return std::is_same<T, int32_t>::value; + case DataType::Uint32: + case DataType::Uint16: + case DataType::Uint8: + return std::is_same<T, uint32_t>::value; + default: + CKW_THROW_MSG("Unknown data type!"); + break; + } +} + +// Necessary instantiations for compiler to recognize +template ConstantData::ConstantData(std::initializer_list<std::initializer_list<int32_t>>, DataType); +template ConstantData::ConstantData(std::initializer_list<std::initializer_list<uint32_t>>, DataType); +template ConstantData::ConstantData(std::initializer_list<std::initializer_list<bool>>, DataType); +template ConstantData::ConstantData(std::initializer_list<std::initializer_list<float>>, DataType); +template ConstantData::ConstantData(const std::vector<std::vector<int32_t>> &, DataType); +template ConstantData::ConstantData(const std::vector<std::vector<uint32_t>> &, DataType); +template ConstantData::ConstantData(const std::vector<std::vector<bool>> &, DataType); +template ConstantData::ConstantData(const std::vector<std::vector<float>> &, DataType); + +template bool ConstantData::validate<int32_t>(DataType); +template bool ConstantData::validate<uint32_t>(DataType); +template bool ConstantData::validate<bool>(DataType); +template bool ConstantData::validate<float>(DataType); + +const std::vector<std::vector<std::string>> &ConstantData::values() const +{ + return _values; +} + +DataType ConstantData::data_type() const +{ + return _data_type; +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/types/DataTypeHelpers.cpp b/compute_kernel_writer/src/types/DataTypeHelpers.cpp new file mode 100644 index 0000000000..7f0c33fb72 --- /dev/null +++ b/compute_kernel_writer/src/types/DataTypeHelpers.cpp @@ -0,0 +1,35 @@ +/* +* 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. +*/ + +#include "src/types/DataTypeHelpers.h" + +namespace ckw +{ + +bool is_data_type_float(DataType data_type) +{ + return (data_type == DataType::Fp32 || data_type == DataType::Fp16); +} + +} // namespace ckw diff --git a/compute_kernel_writer/src/types/DataTypeHelpers.h b/compute_kernel_writer/src/types/DataTypeHelpers.h new file mode 100644 index 0000000000..b6ec6ccd19 --- /dev/null +++ b/compute_kernel_writer/src/types/DataTypeHelpers.h @@ -0,0 +1,43 @@ +/* +* 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_SRC_TYPES_DATATYPEHELPERS_H +#define CKW_SRC_TYPES_DATATYPEHELPERS_H + +#include "ckw/types/DataType.h" + +namespace ckw +{ + +/** Return a value indicating whether the data type is floating-point. + * + * @param[in] data_type The data type to check. + * + * @return Whether the data type is floating-point. + */ +bool is_data_type_float(DataType data_type); + +} // namespace ckw + +#endif // CKW_SRC_TYPES_DATATYPEHELPERS_H diff --git a/compute_kernel_writer/src/types/TensorComponentType.h b/compute_kernel_writer/src/types/TensorComponentType.h new file mode 100644 index 0000000000..03f4f4f5c8 --- /dev/null +++ b/compute_kernel_writer/src/types/TensorComponentType.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_SRC_TYPES_TENSORCOMPONENTTYPE_H +#define CKW_SRC_TYPES_TENSORCOMPONENTTYPE_H + +#include <cstdint> + +namespace ckw +{ + +/** 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 TensorComponent + Stride = 0x02000000, // For example, stride0 in TensorComponent + Dimension = 0x04000000, // For example, Dim0 in TensorComponent + FoldedDimensions = 0x08000000, // For example, Dim0xDim1 in TensorComponent +}; + +/** Mask to retrieve the component index (for example, 1 for stride1, 2 for stride2, or 1 and 2 for Dim1xDim2). + * + * The 4 least significant half-bytes (nibbles) of the @ref TensorComponent are used to retrieve the specific component index. + * TensorComponent = | i7 | i6 | i5 | i4 | i3 | i2 | i1 | i0 |, where i7,...i0 are the nibbles + * of the TensorComponent hexadecimal number. i0, i1, i2 and i3 are reserved to the component index. + * + * In particular: + * + * -# i0: reserved to the first folded dimension component index + * -# i1: reserved to the second folded dimension component index + * -# i2: reserved to the third folded dimension component index + * -# i3: reserved to the fourth folded dimension component index + * + * Therefore, if there are no folded dimensions (dimensions and strides), only i0 is used. + * Instead, if there are two folded dimensions, only i0 and i1 are used. + * + * The component index is stored with the corresponding hexadecimal number + 1, + * hence the component index 0 is represented as 1, while the component index 3 is represented as 4. + */ +enum class TensorComponentIndexBitmask : uint32_t +{ + All = 0x0000ffff, // All nibbles reserved to the tensor component index + Index0 = 0x0000000f, // Folded dimension 0 + Index1 = 0x000000f0, // Folded dimension 1 + Index2 = 0x00000f00, // Folded dimension 2 + Index3 = 0x0000f000 // Folded dimension 3 +}; + +/** The maximum number of folded dimensions. */ +constexpr int tensor_component_index_max_count = 4; + +} // namespace ckw + +#endif // CKW_SRC_TYPES_TENSORCOMPONENTTYPE_H |