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>

6 files changed, 281 insertions, 0 deletions

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<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 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<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 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<TileOperand> &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<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 @@ 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<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 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<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