From d0d8f2e61039826685aa076347eacce526e8c74b Mon Sep 17 00:00:00 2001 From: Viet-Hoa Do Date: Tue, 29 Aug 2023 16:01:13 +0100 Subject: Add get_global_id and printf for CKW Resolves: COMPMID-6387 Signed-off-by: Viet-Hoa Do Change-Id: I5bedb2fdb658a6eb5f1d5053b3840ca81cf75d03 Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/10214 Reviewed-by: Gunes Bayir Comments-Addressed: Arm Jenkins Tested-by: Arm Jenkins Benchmark: Arm Jenkins --- compute_kernel_writer/include/ckw/KernelWriter.h | 24 +++++ compute_kernel_writer/src/cl/CLKernelWriter.cpp | 100 +++++++++++++++++++++ compute_kernel_writer/src/cl/CLKernelWriter.h | 4 + compute_kernel_writer/validation/Validation.cpp | 6 ++ .../tests/CLKernelWriterGetGlobalIdTest.h | 72 +++++++++++++++ .../validation/tests/CLKernelWriterPrintTest.h | 75 ++++++++++++++++ 6 files changed, 281 insertions(+) create mode 100644 compute_kernel_writer/validation/tests/CLKernelWriterGetGlobalIdTest.h create mode 100644 compute_kernel_writer/validation/tests/CLKernelWriterPrintTest.h (limited to 'compute_kernel_writer') diff --git a/compute_kernel_writer/include/ckw/KernelWriter.h b/compute_kernel_writer/include/ckw/KernelWriter.h index 23237ace28..0c8f3de0a1 100644 --- a/compute_kernel_writer/include/ckw/KernelWriter.h +++ b/compute_kernel_writer/include/ckw/KernelWriter.h @@ -175,6 +175,13 @@ public: // Misc // ============================================================================================= + /** Write the statement to get the global ID of the specified dimension. + * + * @param[in] dst The tile to write the global ID into. + * @param[in] dim The dimension. + */ + virtual void op_get_global_id(const TileOperand &dst, int32_t dim) = 0; + /** Write the line comment in debug build. * * This function does not take effect on release build. @@ -185,6 +192,23 @@ public: */ virtual void op_comment(const std::string &text) = 0; + /** Write the statement to print out the value of all the specified tiles. + * + * The printing statement is constructed so that the prefix and each of the operand are printed in separate lines. + * The format for each operand varies depending on whether it is a 2D tile, a vector or a scalar value. + * + * Example output of the printing statement when it is executed: + * + * prefix + * scalar_name = scalar_value + * vector_name = [vector_value_0, vector_value_1, vector_value_2] + * tile_name = [[tile_value_00, tile_value_01], [tile_value_10, tile_value_11]] + * + * @param[in] prefix The first string to be printed out before the list of operands. + * @param[in] operands The list of tiles to be included in the printing statement. + */ + virtual void op_print(const std::string &prefix, const std::vector &operands) = 0; + /** Write the given raw code to kernel source code * It's used to address the cases where the user needs to * explicitly add a code where it's not (yet) supported by diff --git a/compute_kernel_writer/src/cl/CLKernelWriter.cpp b/compute_kernel_writer/src/cl/CLKernelWriter.cpp index 4284388c0b..a946b989d7 100644 --- a/compute_kernel_writer/src/cl/CLKernelWriter.cpp +++ b/compute_kernel_writer/src/cl/CLKernelWriter.cpp @@ -385,6 +385,106 @@ void CLKernelWriter::op_return() append_code("return;\n"); } +void CLKernelWriter::op_get_global_id(const TileOperand &dst, int32_t dim) +{ + const auto &tile = to_cl_tile(dst); + + CKW_ASSERT(tile.is_scalar()); + CKW_ASSERT(tile.info().data_type() == DataType::Int32 || tile.info().data_type() == DataType::Uint32); + + CKW_ASSERT(dim >= 0 && dim <= 2); + + append_code(tile.scalar(0, 0).str, " = get_global_id(", std::to_string(dim), ");\n"); +} + +void CLKernelWriter::op_print(const std::string &prefix, const std::vector &operands) +{ + std::string format_code; + std::string args_code; + + for(auto &op : operands) + { + const auto &tile = to_cl_tile(op); + const auto &info = tile.info(); + + const auto &name = tile.name(); + const auto width = info.width(); + const auto height = info.height(); + const auto data_type = info.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.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 diff --git a/compute_kernel_writer/src/cl/CLKernelWriter.h b/compute_kernel_writer/src/cl/CLKernelWriter.h index 9458ced916..c494847944 100644 --- a/compute_kernel_writer/src/cl/CLKernelWriter.h +++ b/compute_kernel_writer/src/cl/CLKernelWriter.h @@ -95,10 +95,14 @@ public: // 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 &operands) override; + // ============================================================================================= // Code generation // ============================================================================================= diff --git a/compute_kernel_writer/validation/Validation.cpp b/compute_kernel_writer/validation/Validation.cpp index c8d0f6b45d..06af610456 100644 --- a/compute_kernel_writer/validation/Validation.cpp +++ b/compute_kernel_writer/validation/Validation.cpp @@ -31,8 +31,10 @@ #include "validation/tests/CLKernelWriterDeclareTensorTest.h" #include "validation/tests/CLKernelWriterDeclareTileTest.h" #include "validation/tests/CLKernelWriterForTest.h" +#include "validation/tests/CLKernelWriterGetGlobalIdTest.h" #include "validation/tests/CLKernelWriterIfTest.h" #include "validation/tests/CLKernelWriterOpLoadStoreTest.h" +#include "validation/tests/CLKernelWriterPrintTest.h" #include "validation/tests/CLKernelWriterReturnTest.h" #include "validation/tests/CLKernelWriterTernaryOpTest.h" #include "validation/tests/CLKernelWriterUnaryExpressionTest.h" @@ -95,6 +97,8 @@ int32_t main() const auto test32 = std::make_unique(); const auto test33 = std::make_unique(); const auto test34 = std::make_unique(); + const auto test35 = std::make_unique(); + const auto test36 = std::make_unique(); tests.push_back(test3.get()); tests.push_back(test4.get()); @@ -130,6 +134,8 @@ int32_t main() tests.push_back(test32.get()); tests.push_back(test33.get()); tests.push_back(test34.get()); + tests.push_back(test35.get()); + tests.push_back(test36.get()); #endif /* COMPUTE_KERNEL_WRITER_OPENCL_ENABLED */ bool all_test_passed = true; diff --git a/compute_kernel_writer/validation/tests/CLKernelWriterGetGlobalIdTest.h b/compute_kernel_writer/validation/tests/CLKernelWriterGetGlobalIdTest.h new file mode 100644 index 0000000000..fa34b3f5df --- /dev/null +++ b/compute_kernel_writer/validation/tests/CLKernelWriterGetGlobalIdTest.h @@ -0,0 +1,72 @@ +/* + * 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_VALIDATION_TESTS_CLKERNELWRITERGETGLOBALIDTEST_H +#define CKW_VALIDATION_TESTS_CLKERNELWRITERGETGLOBALIDTEST_H + +#include "ckw/TileInfo.h" +#include "src/cl/CLKernelWriter.h" +#include "validation/tests/common/Common.h" +#include "validation/tests/common/KernelWriterInterceptor.h" + +namespace ckw +{ + +class CLKernelWriterGetGlobalIdTest : public ITest +{ +public: + CLKernelWriterGetGlobalIdTest() + { + } + + bool run() override + { + bool all_tests_passed = true; + + KernelWriterInterceptor writer; + + auto gid = writer.declare_tile("gid", TileInfo(DataType::Int32)); + + writer.start_capture_code(); + + writer.op_get_global_id(gid, 0); + writer.op_get_global_id(gid, 1); + writer.op_get_global_id(gid, 2); + + constexpr auto expected_code = "G0__gid = get_global_id(0);\nG0__gid = get_global_id(1);\nG0__gid = get_global_id(2);\n"; + + VALIDATE_TEST(writer.check_added_code(expected_code), all_tests_passed, 0); + + return all_tests_passed; + } + + std::string name() override + { + return "CLKernelWriterGetGlobalIdTest"; + } +}; + +} // namespace ckw + +#endif // CKW_VALIDATION_TESTS_CLKERNELWRITERGETGLOBALIDTEST_H diff --git a/compute_kernel_writer/validation/tests/CLKernelWriterPrintTest.h b/compute_kernel_writer/validation/tests/CLKernelWriterPrintTest.h new file mode 100644 index 0000000000..6229dfb8c0 --- /dev/null +++ b/compute_kernel_writer/validation/tests/CLKernelWriterPrintTest.h @@ -0,0 +1,75 @@ +/* + * 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_VALIDATION_TESTS_CLKERNELWRITERPRINT_H +#define CKW_VALIDATION_TESTS_CLKERNELWRITERPRINT_H + +#include "ckw/TileInfo.h" +#include "src/cl/CLKernelWriter.h" +#include "validation/tests/common/Common.h" +#include "validation/tests/common/KernelWriterInterceptor.h" + +namespace ckw +{ + +class CLKernelWriterPrintTest : public ITest +{ +public: + CLKernelWriterPrintTest() + { + } + + bool run() override + { + bool all_tests_passed = true; + + KernelWriterInterceptor writer; + + const auto tile2x3f16 = writer.declare_tile("tile2x3f16", TileInfo(DataType::Fp16, 2, 3)); + const auto tile1x2i32 = writer.declare_tile("tile1x2i32", TileInfo(DataType::Int32, 1, 2)); + const auto tile2x1s32 = writer.declare_tile("tile2x1s32", TileInfo(DataType::Int32, 2, 1)); + const auto tile1x1u32 = writer.declare_tile("tile1x1u32", TileInfo(DataType::Uint32, 1, 1)); + + writer.start_capture_code(); + + writer.op_print("debug_log", { tile2x3f16, tile1x2i32, tile2x1s32, tile1x1u32 }); + + constexpr auto expected_code = + "printf(\"debug_log\\nG0__tile2x3f16 = [[%v3hg], [%v3hg]]\\nG0__tile1x2i32 = [%v2hli]\\nG0__tile2x1s32 = [%i, %i]\\nG0__tile1x1u32 = %u\\n\", " + "G0__tile2x3f16__0, G0__tile2x3f16__1, G0__tile1x2i32, G0__tile2x1s32__0, G0__tile2x1s32__1, G0__tile1x1u32);\n"; + + VALIDATE_TEST(writer.check_added_code(expected_code), all_tests_passed, 0); + + return all_tests_passed; + } + + std::string name() override + { + return "CLKernelWriterPrintTest"; + } +}; + +} // namespace ckw + +#endif // CKW_VALIDATION_TESTS_CLKERNELWRITERPRINT_H -- cgit v1.2.1