diff options
Diffstat (limited to 'compute_kernel_writer/prototype/examples')
8 files changed, 0 insertions, 758 deletions
diff --git a/compute_kernel_writer/prototype/examples/add_exp_store.cpp b/compute_kernel_writer/prototype/examples/add_exp_store.cpp deleted file mode 100644 index 2b640ca01b..0000000000 --- a/compute_kernel_writer/prototype/examples/add_exp_store.cpp +++ /dev/null @@ -1,206 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#include "ckw/Error.h" -#include "ckw/KernelArgument.h" -#include "ckw/KernelWriter.h" -#include "ckw/TensorOperand.h" -#include "ckw/TensorTileSampler.h" -#include "ckw/TileOperand.h" - -#include "common/ExampleComponentArgument.h" -#include "common/ExampleKernelWriter.h" -#include "common/ExampleScopedKernelWriter.h" -#include <iostream> -#include <vector> - -using namespace ckw; - -TensorTileSampler create_simple_sampler(ExampleScopedKernelWriter writer) -{ - TensorTileSampler sampler; - - constexpr int32_t m0 = 4; - constexpr int32_t n0 = 4; - - auto &gid_0 = writer->declare_tile("gid_0", DataType::Int32); - auto &gid_1 = writer->declare_tile("gid_1", DataType::Int32); - auto &gid_2 = writer->declare_tile("gid_2", DataType::Int32); - - auto &const_0 = writer->declare_tile("0", 0); - - writer->op_get_global_id(gid_0, 0); - writer->op_get_global_id(gid_1, 1); - writer->op_get_global_id(gid_2, 2); - - sampler.x(gid_0); - sampler.y(gid_1); - sampler.z(const_0); - sampler.b(gid_2); - - sampler.width(n0); - sampler.height(m0); - - sampler.format(TensorSamplerFormat::C_WH_1); - sampler.address_mode_x(TensorSamplerAddressModeX::None); - sampler.address_mode_y(TensorSamplerAddressModeY::ClampToBorder); - sampler.address_mode_z(TensorSamplerAddressModeZ::Skip); - - return sampler; -} - -void op_binary_elementwise(ExampleScopedKernelWriter writer, std::vector<ExampleComponentArgument *> operands) -{ - auto lhs = operands.at(0); - auto rhs = operands.at(1); - auto dst = operands.at(2); - - // Load the LHS and RHS tile and prepare the tensor sampler. - if (!lhs->has_tile() && !rhs->has_tile()) - { - const auto sampler = create_simple_sampler(writer); - - writer->op_load_once(lhs, sampler); - writer->op_load_once(rhs, sampler); - } - else if (lhs->has_tile()) - { - const auto &sampler = lhs->tile_sampler(); - writer->op_load_once(rhs, sampler); - } - else - { - const auto &sampler = rhs->tile_sampler(); - writer->op_load_once(lhs, sampler); - } - - auto &lhs_tile = lhs->tile(); - auto &rhs_tile = rhs->tile(); - const auto &sampler = lhs->tile_sampler(); - - // Prepare the output tile. - if (!dst->has_tile()) - { - auto &tile = writer->declare_tile("dst_tile", lhs_tile.tile_info()); - dst->init_virtual_tensor(tile, sampler); - } - - auto &dst_tile = dst->tile(); - - // Perform the operation. - writer->op_binary_expression(dst_tile, lhs_tile, BinaryOp::Add, rhs_tile); -} - -void op_exp(ExampleScopedKernelWriter writer, std::vector<ExampleComponentArgument *> operands) -{ - auto src = operands.at(0); - auto dst = operands.at(1); - - // Load the source tile and prepare the sampler. - if (!src->has_tile()) - { - const auto sampler = create_simple_sampler(writer); - writer->op_load_once(src, sampler); - } - - auto &src_tile = src->tile(); - const auto &sampler = src->tile_sampler(); - - // Prepare the output tile. - if (!dst->has_tile()) - { - auto &tile = writer->declare_tile("dst_tile", src_tile.tile_info()); - dst->init_virtual_tensor(tile, sampler); - } - - auto &dst_tile = dst->tile(); - - // Perform the operation. - writer->op_unary_elementwise_function(dst_tile, UnaryFunction::Exp, src_tile); -} - -void op_store(ExampleScopedKernelWriter writer, std::vector<ExampleComponentArgument *> operands) -{ - auto src = operands.at(0); - auto dst = operands.at(1); - - auto &src_tile = src->tile(); - const auto &sampler = src->tile_sampler(); - auto &dst_tensor = dst->tensor(); - - writer->op_store(dst_tensor, src_tile, sampler); -} - -int main() -{ - Kernel kernel("example", GpuTargetLanguage::OpenCL); - ExampleKernelWriter root_writer(kernel); - - ExampleScopedKernelWriter writer(&root_writer); - - const TensorInfo src0_info(DataType::Fp32, TensorShape({3, 10, 20, 1, 1}), TensorDataLayout::Nhwc, 0); - const TensorInfo src1_info(DataType::Fp32, TensorShape({3, 10, 20, 1, 1}), TensorDataLayout::Nhwc, 1); - const TensorInfo dst_info(DataType::Fp32, TensorShape({3, 10, 20, 1, 1}), TensorDataLayout::Nhwc, 2); - - ExampleComponentArgument src0( - writer->declare_tensor_argument("src0", src0_info, TensorStorageType::BufferUint8Ptr)); - ExampleComponentArgument src1( - writer->declare_tensor_argument("src1", src1_info, TensorStorageType::BufferUint8Ptr)); - ExampleComponentArgument dst(writer->declare_tensor_argument("dst", dst_info, TensorStorageType::BufferUint8Ptr)); - - ExampleComponentArgument ans; - - op_binary_elementwise(writer, {&src0, &src1, &ans}); - op_exp(writer, {&ans, &ans}); - op_store(writer, {&ans, &dst}); - - const auto arguments = kernel.arguments(); - - std::cout << "\n====================\nArguments:\n====================\n"; - - for (auto &arg : arguments) - { - switch (arg.type()) - { - case ckw::KernelArgument::Type::TensorStorage: - std::cout << "* Tensor storage: ID = " << arg.id() << ", type = " << std::hex << "0x" - << static_cast<uint32_t>(arg.tensor_storage_type()) << std::dec << "\n"; - break; - - case ckw::KernelArgument::Type::TensorComponent: - std::cout << "* Tensor component: ID = " << arg.id() << ", type = " << std::hex << "0x" - << static_cast<uint32_t>(arg.tensor_component_type()) << std::dec << "\n"; - break; - - default: - CKW_ASSERT(false); - } - } - - std::cout << "\n====================\nCode:\n====================\n"; - const auto code = root_writer.generate_code(); - std::cout << code; - - return 0; -} diff --git a/compute_kernel_writer/prototype/examples/common/ExampleComponentArgument.cpp b/compute_kernel_writer/prototype/examples/common/ExampleComponentArgument.cpp deleted file mode 100644 index 55223dae0e..0000000000 --- a/compute_kernel_writer/prototype/examples/common/ExampleComponentArgument.cpp +++ /dev/null @@ -1,98 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#include "ExampleComponentArgument.h" - -#include "ckw/Error.h" - -ExampleComponentArgument::ExampleComponentArgument() -{ -} - -ExampleComponentArgument::ExampleComponentArgument(ckw::TensorOperand &tensor) : _tensor(&tensor) -{ -} - -ExampleComponentArgument &ExampleComponentArgument::init_virtual_tensor(ckw::TileOperand &tile, - const ckw::TensorTileSampler &tile_sampler) -{ - CKW_ASSERT(_tile == nullptr); - - _tile = &tile; - _tile_sampler = tile_sampler; - - return *this; -} - -bool ExampleComponentArgument::has_tensor() const -{ - return _tensor != nullptr; -} - -ckw::TensorOperand &ExampleComponentArgument::tensor() -{ - CKW_ASSERT(_tensor != nullptr); - - return *_tensor; -} - -const ckw::TensorOperand &ExampleComponentArgument::tensor() const -{ - CKW_ASSERT(_tensor != nullptr); - - return *_tensor; -} - -bool ExampleComponentArgument::has_tile() const -{ - return _tile != nullptr; -} - -ckw::TileOperand &ExampleComponentArgument::tile() -{ - CKW_ASSERT(_tile != nullptr); - - return *_tile; -} - -const ckw::TileOperand &ExampleComponentArgument::tile() const -{ - CKW_ASSERT(_tile != nullptr); - - return *_tile; -} - -ckw::TensorTileSampler &ExampleComponentArgument::tile_sampler() -{ - CKW_ASSERT(_tile != nullptr); - - return _tile_sampler; -} - -const ckw::TensorTileSampler &ExampleComponentArgument::tile_sampler() const -{ - CKW_ASSERT(_tile != nullptr); - - return _tile_sampler; -} diff --git a/compute_kernel_writer/prototype/examples/common/ExampleComponentArgument.h b/compute_kernel_writer/prototype/examples/common/ExampleComponentArgument.h deleted file mode 100644 index 0e029b1157..0000000000 --- a/compute_kernel_writer/prototype/examples/common/ExampleComponentArgument.h +++ /dev/null @@ -1,112 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_PROTOTYPE_EXAMPLES_COMMON_EXAMPLECOMPONENTARGUMENT_H -#define CKW_PROTOTYPE_EXAMPLES_COMMON_EXAMPLECOMPONENTARGUMENT_H - -#include "ckw/TensorTileSampler.h" - -namespace ckw -{ -class TensorOperand; - -class TileOperand; -} // namespace ckw - -/** The argument of a dynamic fusion component which can be either user tensor or virtual tensor. */ -class ExampleComponentArgument -{ -public: - /** Initialize a new instance of @ref ExampleComponentArgument class for empty virtual tensor. */ - ExampleComponentArgument(); - - /** Initialize a new instance of @ref ExampleComponentArgument class for user tensor. - * - * @param[in] tensor The user tensor. - */ - explicit ExampleComponentArgument(ckw::TensorOperand &tensor); - - /** Set virtual tensor information (tile, sampler) for the argument. - * - * If the component is a user tensor, it can be treated as virtual tensor as well - * and won't be loaded again using @ref ExampleKernelWriter::op_load_once method. - * - * @param[in] tile The tile that has been loaded. - * @param[in] sampler The tensor sampling information that has been used to load the tile. - */ - ExampleComponentArgument &init_virtual_tensor(ckw::TileOperand &tile, const ckw::TensorTileSampler &sampler); - - /** Get whether the argument is a user tensor. */ - bool has_tensor() const; - - /** Get the tensor operand. - * - * If the tensor is not available, throw an error. - */ - ckw::TensorOperand &tensor(); - - /** Get the tensor operand. - * - * If the tensor is not available, throw an error. - */ - const ckw::TensorOperand &tensor() const; - - /** Get whether the argument contains a tile. - * - * The argument can be either a user tensor that has been loaded, - * or a virtual tensor (i.e. a tile with tensor sampling information). - */ - bool has_tile() const; - - /** Get the tile operand. - * - * If the tile is not available, throw an error. - */ - ckw::TileOperand &tile(); - - /** Get the tile operand. - * - * If the tile is not available, throw an error. - */ - const ckw::TileOperand &tile() const; - - /** Get the tensor sampling information for the tile. - * - * If the tile is not available, throw an error. - */ - ckw::TensorTileSampler &tile_sampler(); - - /** Get the tensor sampling information for the tile. - * - * If the tile is not available, throw an error. - */ - const ckw::TensorTileSampler &tile_sampler() const; - -private: - ckw::TensorOperand *_tensor{nullptr}; - ckw::TileOperand *_tile{nullptr}; - ckw::TensorTileSampler _tile_sampler{}; -}; - -#endif // CKW_PROTOTYPE_EXAMPLES_COMMON_EXAMPLECOMPONENTARGUMENT_H diff --git a/compute_kernel_writer/prototype/examples/common/ExampleKernelWriter.cpp b/compute_kernel_writer/prototype/examples/common/ExampleKernelWriter.cpp deleted file mode 100644 index 1734ce8823..0000000000 --- a/compute_kernel_writer/prototype/examples/common/ExampleKernelWriter.cpp +++ /dev/null @@ -1,52 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#include "ExampleKernelWriter.h" - -#include "ckw/Error.h" -#include "ckw/TileInfo.h" - -#include "ExampleComponentArgument.h" - -ExampleKernelWriter::ExampleKernelWriter(ckw::Kernel &kernel) : KernelWriter(kernel) -{ -} - -void ExampleKernelWriter::op_load_once(ExampleComponentArgument *tensor_or_tile, const ckw::TensorTileSampler &sampler) -{ - if (!tensor_or_tile->has_tile()) - { - CKW_ASSERT(tensor_or_tile->has_tensor()); - - auto &tensor = tensor_or_tile->tensor(); - - const auto tile_name = tensor.name() + "_tile"; - auto &tile = - declare_tile(tile_name.c_str(), ckw::TileInfo(tensor.data_type(), sampler.height(), sampler.width())); - - op_load(tile, tensor, sampler); - - tensor_or_tile->init_virtual_tensor(tile, sampler); - } -} diff --git a/compute_kernel_writer/prototype/examples/common/ExampleKernelWriter.h b/compute_kernel_writer/prototype/examples/common/ExampleKernelWriter.h deleted file mode 100644 index 1528c3d933..0000000000 --- a/compute_kernel_writer/prototype/examples/common/ExampleKernelWriter.h +++ /dev/null @@ -1,56 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_PROTOTYPE_EXAMPLES_COMMON_EXAMPLEKERNELWRITER_H -#define CKW_PROTOTYPE_EXAMPLES_COMMON_EXAMPLEKERNELWRITER_H - -#include "ckw/KernelWriter.h" -#include "ckw/TensorTileSampler.h" - -class ExampleComponentArgument; - -namespace ckw -{ -class Kernel; -} // namespace ckw - -/** Extended implementation of kernel writer for dynamic fusion. */ -class ExampleKernelWriter : public ckw::KernelWriter -{ -public: - /** Initialize a new instance of @ref ExampleKernelWriter class. - * - * @param[in] kernel The kernel to be generated. - */ - explicit ExampleKernelWriter(ckw::Kernel &kernel); - - /** Load the user tensor to the tile in the same component argument if it hasn't been loaded. - * - * @param[in] tensor_or_tile The component argument that is either a user tensor or a virtual tensor. - * @param[in] sampler The tensor sampling information to load the tile. - */ - void op_load_once(ExampleComponentArgument *tensor_or_tile, const ckw::TensorTileSampler &sampler); -}; - -#endif // CKW_PROTOTYPE_EXAMPLES_COMMON_EXAMPLEKERNELWRITER_H diff --git a/compute_kernel_writer/prototype/examples/common/ExampleScopedKernelWriter.cpp b/compute_kernel_writer/prototype/examples/common/ExampleScopedKernelWriter.cpp deleted file mode 100644 index 784d5ffb96..0000000000 --- a/compute_kernel_writer/prototype/examples/common/ExampleScopedKernelWriter.cpp +++ /dev/null @@ -1,59 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#include "ExampleScopedKernelWriter.h" - -#include "ExampleKernelWriter.h" - -ExampleScopedKernelWriter::ExampleScopedKernelWriter(ExampleKernelWriter *writer) - : _writer(writer), _parent_id_space(writer->id_space()) -{ - _writer->next_id_space(); -} - -ExampleScopedKernelWriter::ExampleScopedKernelWriter(const ExampleScopedKernelWriter &other) - : _writer(other._writer), _parent_id_space(other._writer->id_space()) -{ - _writer->next_id_space(); -} - -ExampleKernelWriter *ExampleScopedKernelWriter::operator->() -{ - return _writer; -} - -const ExampleKernelWriter *ExampleScopedKernelWriter::operator->() const -{ - return _writer; -} - -ExampleKernelWriter *ExampleScopedKernelWriter::writer() -{ - return _writer; -} - -const ExampleKernelWriter *ExampleScopedKernelWriter::writer() const -{ - return _writer; -} diff --git a/compute_kernel_writer/prototype/examples/common/ExampleScopedKernelWriter.h b/compute_kernel_writer/prototype/examples/common/ExampleScopedKernelWriter.h deleted file mode 100644 index 4655b1897e..0000000000 --- a/compute_kernel_writer/prototype/examples/common/ExampleScopedKernelWriter.h +++ /dev/null @@ -1,62 +0,0 @@ -/* - * Copyright (c) 2023 Arm Limited. - * - * SPDX-License-Identifier: MIT - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal in the Software without restriction, including without limitation the - * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or - * sell copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - */ - -#ifndef CKW_PROTOTYPE_EXAMPLES_COMMON_EXAMPLESCOPEDKERNELWRITER_H -#define CKW_PROTOTYPE_EXAMPLES_COMMON_EXAMPLESCOPEDKERNELWRITER_H - -#include <cstdint> - -class ExampleKernelWriter; - -/** Helper to automatically manage kernel writer ID space. */ -class ExampleScopedKernelWriter -{ -public: - /** Initialize a new instance of @ref ExampleScopedKernelWriter class. */ - explicit ExampleScopedKernelWriter(ExampleKernelWriter *writer); - - /** Create a new scope from the specified scoped kernel writer. */ - ExampleScopedKernelWriter(const ExampleScopedKernelWriter &other); - - /** Assignment is disallowed. */ - ExampleScopedKernelWriter &operator=(const ExampleScopedKernelWriter &) = delete; - - /** Access the underlying kernel writer. */ - ExampleKernelWriter *operator->(); - - /** Access the underlying kernel writer. */ - const ExampleKernelWriter *operator->() const; - - /** Get the kernel writer. */ - ExampleKernelWriter *writer(); - - /** Get the kernel writer. */ - const ExampleKernelWriter *writer() const; - -private: - ExampleKernelWriter *_writer; - int32_t _parent_id_space; -}; - -#endif // CKW_PROTOTYPE_EXAMPLES_COMMON_EXAMPLESCOPEDKERNELWRITER_H diff --git a/compute_kernel_writer/prototype/examples/writer_helper.cpp b/compute_kernel_writer/prototype/examples/writer_helper.cpp deleted file mode 100644 index 8623afbf50..0000000000 --- a/compute_kernel_writer/prototype/examples/writer_helper.cpp +++ /dev/null @@ -1,113 +0,0 @@ -/* -* Copyright (c) 2023 Arm Limited. -* -* SPDX-License-Identifier: MIT -* -* Permission is hereby granted, free of charge, to any person obtaining a copy -* of this software and associated documentation files (the "Software"), to -* deal in the Software without restriction, including without limitation the -* rights to use, copy, modify, merge, publish, distribute, sublicense, and/or -* sell copies of the Software, and to permit persons to whom the Software is -* furnished to do so, subject to the following conditions: -* -* The above copyright notice and this permission notice shall be included in all -* copies or substantial portions of the Software. -* -* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -* SOFTWARE. -*/ - -#include "ckw/KernelWriter.h" -#include "ckw/TensorTileSampler.h" - -#include "../include/ckw/KernelWriterHelper.h" -#include <iostream> - -using namespace ckw; - -TensorTileSampler create_simple_sampler(KernelWriter &writer) -{ - TensorTileSampler sampler; - - constexpr int32_t m0 = 1; - constexpr int32_t n0 = 1; - - auto &gid_0 = writer.declare_tile("gid_0", DataType::Int32); - auto &gid_1 = writer.declare_tile("gid_1", DataType::Int32); - auto &gid_2 = writer.declare_tile("gid_2", DataType::Int32); - - auto &const_0 = writer.declare_tile("0", 0); - - writer.op_get_global_id(gid_0, 0); - writer.op_get_global_id(gid_1, 1); - writer.op_get_global_id(gid_2, 2); - - sampler.x(gid_0); - sampler.y(gid_1); - sampler.z(gid_2); - sampler.b(const_0); - - sampler.width(n0); - sampler.height(m0); - - sampler.format(TensorSamplerFormat::C_WH_1); - sampler.address_mode_x(TensorSamplerAddressModeX::None); - sampler.address_mode_y(TensorSamplerAddressModeY::ClampToBorder); - sampler.address_mode_z(TensorSamplerAddressModeZ::Skip); - - return sampler; -} - -int main() -{ - Kernel kernel("test", GpuTargetLanguage::OpenCL); - KernelWriterHelper<KernelWriter> writer(kernel); - - const TensorInfo src_info(DataType::Fp32, TensorShape({1, 1, 1, 1, 1}), TensorDataLayout::Nhwc, 0); - const TensorInfo dst_info(DataType::Fp32, TensorShape({1, 1, 1, 1, 1}), TensorDataLayout::Nhwc, 1); - - auto &src_tensor = writer.declare_tensor_argument("src", src_info); - auto &dst_tensor = writer.declare_tensor_argument("dst", dst_info); - - const auto sampler = create_simple_sampler(writer); - - auto &src = writer.declare_tile("src_tile", TileInfo(src_tensor.data_type(), sampler.height(), sampler.width())); - auto &other = - writer.declare_tile("other_tile", TileInfo(src_tensor.data_type(), sampler.height(), sampler.width())); - auto &dst = writer.declare_tile("dst_tile", TileInfo(src_tensor.data_type(), sampler.height(), sampler.width())); - - writer.op_load(src, src_tensor, sampler); - writer.op_load(other, src_tensor, sampler); - writer.op_load(dst, dst_tensor, sampler); - - auto test = dst ^ src ^ other; - auto other_test = logical_and(dst, src, other); - writer.op_assign(dst, logical_and(dst, src, other)); - writer.op_assign(dst, test); - writer.op_assign(dst, other_test); - writer.op_assign(dst, operator^(operator^(dst, src), other)); - - writer.op_if(exp(src) == dst, [&] { writer.op_binary_expression(dst, src, BinaryOp::Add, src); }) - .op_else_if(exp(src) > dst, [&] { writer.op_binary_expression(dst, src, BinaryOp::Add, src); }) - .op_else([&] { writer.op_assign(dst, src); }); - - writer.op_assign(dst, src + src * src); - writer.op_assign(dst, src * max(src, dst) + src); - writer.op_assign(dst, src * select(src, dst, src) + src); - - writer.op_assign(dst, src ^ dst); - writer.op_assign(dst, ~src); - - writer.op_for_loop(dst < src, dst += src, [&] { writer.op_assign(dst, src + dst); }); - - writer.op_assign(dst += src); - writer.op_assign(dst += exp(src)); - - std::cout << "======== KERNEL ========" << std::endl; - std::cout << writer.generate_code() << std::endl; -} |