Add get_global_id and printf for CKW
Resolves: COMPMID-6387
Signed-off-by: Viet-Hoa Do <viet-hoa.do@arm.com>
Change-Id: I5bedb2fdb658a6eb5f1d5053b3840ca81cf75d03
Reviewed-on: https://review.mlplatform.org/c/ml/ComputeLibrary/+/10214
Reviewed-by: Gunes Bayir <gunes.bayir@arm.com>
Comments-Addressed: Arm Jenkins <bsgcomp@arm.com>
Tested-by: Arm Jenkins <bsgcomp@arm.com>
Benchmark: Arm Jenkins <bsgcomp@arm.com>
diff --git a/compute_kernel_writer/include/ckw/KernelWriter.h b/compute_kernel_writer/include/ckw/KernelWriter.h
index 23237ac..0c8f3de 100644
--- a/compute_kernel_writer/include/ckw/KernelWriter.h
+++ b/compute_kernel_writer/include/ckw/KernelWriter.h
@@ -175,6 +175,13 @@
// 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 @@
*/
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<TileOperand> &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 4284388..a946b98 100644
--- a/compute_kernel_writer/src/cl/CLKernelWriter.cpp
+++ b/compute_kernel_writer/src/cl/CLKernelWriter.cpp
@@ -385,6 +385,106 @@
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<TileOperand> &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 9458ced..c494847 100644
--- a/compute_kernel_writer/src/cl/CLKernelWriter.h
+++ b/compute_kernel_writer/src/cl/CLKernelWriter.h
@@ -95,10 +95,14 @@
// 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
// =============================================================================================
diff --git a/compute_kernel_writer/validation/Validation.cpp b/compute_kernel_writer/validation/Validation.cpp
index c8d0f6b..06af610 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 @@
const auto test32 = std::make_unique<CLKernelWriterIfTest>();
const auto test33 = std::make_unique<CLKernelWriterForTest>();
const auto test34 = std::make_unique<CLKernelWriterReturnTest>();
+ const auto test35 = std::make_unique<CLKernelWriterGetGlobalIdTest>();
+ const auto test36 = std::make_unique<CLKernelWriterPrintTest>();
tests.push_back(test3.get());
tests.push_back(test4.get());
@@ -130,6 +134,8 @@
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 0000000..fa34b3f
--- /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<CLKernelWriter> 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 0000000..6229dfb
--- /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<CLKernelWriter> 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