diff options
author | Viet-Hoa Do <viet-hoa.do@arm.com> | 2023-08-29 16:01:13 +0100 |
---|---|---|
committer | Viet-Hoa Do <viet-hoa.do@arm.com> | 2023-08-30 12:32:10 +0000 |
commit | d0d8f2e61039826685aa076347eacce526e8c74b (patch) | |
tree | b655ee0c07c8c87ce387c370d20494169ae86d5c /compute_kernel_writer/validation | |
parent | 87706692252b0746e882a9dd34ae64dc60acd767 (diff) | |
download | ComputeLibrary-d0d8f2e61039826685aa076347eacce526e8c74b.tar.gz |
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>
Diffstat (limited to 'compute_kernel_writer/validation')
3 files changed, 153 insertions, 0 deletions
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 |