From cd11d020bb4d046c546b99c3fbef657085348f07 Mon Sep 17 00:00:00 2001 From: adstraw Date: Tue, 19 Apr 2022 14:57:45 -0700 Subject: [PATCH 01/20] link gtest to tvm runtime --- CMakeLists.txt | 7 +++++++ src/runtime/hexagon/hexagon_device_api_v2.cc | 11 +++++++++++ 2 files changed, 18 insertions(+) diff --git a/CMakeLists.txt b/CMakeLists.txt index 151173ac5759..9a512522bef4 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -598,6 +598,13 @@ endif() target_link_libraries(tvm PRIVATE ${TVM_LINKER_LIBS} ${TVM_RUNTIME_LINKER_LIBS}) target_link_libraries(tvm_runtime PRIVATE ${TVM_RUNTIME_LINKER_LIBS}) +include(FetchContent) +FetchContent_Declare(googletest SOURCE_DIR "/local/mnt/workspace/Qualcomm/Hexagon_SDK/4.5.0.3/utils/googletest/gtest") +set(gtest_force_shared_crt ON CACHE BOOL "" FORCE) +FetchContent_MakeAvailable(googletest) +include(GoogleTest) +target_link_libraries(tvm_runtime PUBLIC gtest) + # Set flags for clang include(cmake/modules/ClangFlags.cmake) set(CRC16_INCLUDE_PATH "3rdparty/libcrc/include") diff --git a/src/runtime/hexagon/hexagon_device_api_v2.cc b/src/runtime/hexagon/hexagon_device_api_v2.cc index 8da66ad1d0b8..2e5682793214 100644 --- a/src/runtime/hexagon/hexagon_device_api_v2.cc +++ b/src/runtime/hexagon/hexagon_device_api_v2.cc @@ -35,6 +35,8 @@ #include "hexagon_buffer.h" #include "hexagon_common.h" +#include "gtest/gtest.h" + namespace tvm { namespace runtime { namespace hexagon { @@ -245,6 +247,15 @@ TVM_REGISTER_GLOBAL("device_api.hexagon.v2").set_body([](TVMArgs args, TVMRetVal *rv = static_cast(ptr); }); +TEST(HexagonDeviceApi, my_first_test) { + ASSERT_EQ(42, 42); +} + +TVM_REGISTER_GLOBAL("device_api.hexagon.run_all_tests").set_body([](TVMArgs args, TVMRetValue* rv) { + *rv = RUN_ALL_TESTS(); +}); + + } // namespace hexagon } // namespace runtime } // namespace tvm From 31fc623ed553ba6d893b70903f86a0d9a81556c6 Mon Sep 17 00:00:00 2001 From: adstraw Date: Thu, 21 Apr 2022 10:22:19 -0700 Subject: [PATCH 02/20] first test running! --- CMakeLists.txt | 4 +- src/runtime/hexagon/hexagon_device_api_v2.cc | 2 +- .../test_hexagon/test_cache_read_write.py | 37 +++---------------- 3 files changed, 9 insertions(+), 34 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 9a512522bef4..481d95634c11 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -602,8 +602,9 @@ include(FetchContent) FetchContent_Declare(googletest SOURCE_DIR "/local/mnt/workspace/Qualcomm/Hexagon_SDK/4.5.0.3/utils/googletest/gtest") set(gtest_force_shared_crt ON CACHE BOOL "" FORCE) FetchContent_MakeAvailable(googletest) -include(GoogleTest) target_link_libraries(tvm_runtime PUBLIC gtest) +target_link_libraries(tvm PUBLIC gtest) # is this necessary if we move the tests out of device api? +include_directories("/local/mnt/workspace/Qualcomm/Hexagon_SDK/4.5.0.3/utils/googletest/gtest/include") # Set flags for clang include(cmake/modules/ClangFlags.cmake) @@ -656,6 +657,7 @@ add_custom_target(runtime DEPENDS tvm_runtime) # Installation rules install(TARGETS tvm EXPORT ${PROJECT_NAME}Targets DESTINATION lib${LIB_SUFFIX}) install(TARGETS tvm_runtime EXPORT ${PROJECT_NAME}Targets DESTINATION lib${LIB_SUFFIX}) +install(TARGETS gtest EXPORT ${PROJECT_NAME}Targets DESTINATION lib${LIB_SUFFIX}) if (INSTALL_DEV) install( diff --git a/src/runtime/hexagon/hexagon_device_api_v2.cc b/src/runtime/hexagon/hexagon_device_api_v2.cc index 2e5682793214..6640b3647d59 100644 --- a/src/runtime/hexagon/hexagon_device_api_v2.cc +++ b/src/runtime/hexagon/hexagon_device_api_v2.cc @@ -251,7 +251,7 @@ TEST(HexagonDeviceApi, my_first_test) { ASSERT_EQ(42, 42); } -TVM_REGISTER_GLOBAL("device_api.hexagon.run_all_tests").set_body([](TVMArgs args, TVMRetValue* rv) { +TVM_REGISTER_GLOBAL("hexagon.run_all_tests").set_body([](TVMArgs args, TVMRetValue* rv) { *rv = RUN_ALL_TESTS(); }); diff --git a/tests/python/contrib/test_hexagon/test_cache_read_write.py b/tests/python/contrib/test_hexagon/test_cache_read_write.py index 8f9453187169..a67a14687d1c 100644 --- a/tests/python/contrib/test_hexagon/test_cache_read_write.py +++ b/tests/python/contrib/test_hexagon/test_cache_read_write.py @@ -144,35 +144,8 @@ def layout_transform_2d(n): @requires_hexagon_toolchain def test_cache_read_write_2d(hexagon_session): - size = 128 - outer_shape = (size,) - factor = 16 - inner_shape = (factor,) - dtype = "int8" - - x = te.placeholder(shape=outer_shape, dtype=dtype, name="x") - y = te.placeholder(shape=outer_shape, dtype=dtype, name="y") - z = te.compute(outer_shape, lambda i: x[i] + y[i], name="z") - s = te.create_schedule(z.op) - - x_vtcm = s.cache_read(x, "global.vtcm", [z]) - y_vtcm = s.cache_read(y, "global.vtcm", [z]) - z_vtcm = s.cache_write(z, "global.vtcm") - - layout_x_vtcm = s[x_vtcm].transform_layout(layout_transform_2d) - layout_y_vtcm = s[y_vtcm].transform_layout(layout_transform_2d) - layout_z_vtcm = s[z_vtcm].transform_layout(layout_transform_2d) - - mem_copy_read = intrin_mem_copy(inner_shape, dtype, "global.vtcm", "global") - s[x_vtcm].tensorize(layout_x_vtcm[1], mem_copy_read) - s[y_vtcm].tensorize(layout_y_vtcm[1], mem_copy_read) - - # The loop schedule over `z` is not modified when calling `transform_layout` - # on `z_vtcm` above therefore we must call `split` to modify the loop schedule - # over `z` to match the layout of `z_vtcm` such that we can accurately write - # `z_vtcm` back to `z` using memory copy intrinsic - zouter, zinner = s[z].split(z.op.axis[0], factor=factor) - mem_copy_write = intrin_mem_copy(inner_shape, dtype, "global", "global.vtcm") - s[z].tensorize(zinner, mem_copy_write) - - verify(hexagon_session, s, x, y, z, size) + func = hexagon_session._rpc.get_function("hexagon.run_all_tests") + x = func() + print("PRETTY SURE I JUST RAN MY TEST, PRINTING THE RETURN VALUE") + print(x) + np.testing.assert_equal(x, 0) From 79c21105b9eb939ebdf2d34416a0067562dc1f5e Mon Sep 17 00:00:00 2001 From: adstraw Date: Thu, 21 Apr 2022 10:44:09 -0700 Subject: [PATCH 03/20] HexagonBuffer tests running in sim --- .../runtime/hexagon/hexagon_buffer_tests.cc | 4 ++++ src/runtime/hexagon/hexagon_device_api_v2.cc | 2 +- 2 files changed, 5 insertions(+), 1 deletion(-) rename tests/cpp/runtime/hexagon_buffer.cc => src/runtime/hexagon/hexagon_buffer_tests.cc (98%) diff --git a/tests/cpp/runtime/hexagon_buffer.cc b/src/runtime/hexagon/hexagon_buffer_tests.cc similarity index 98% rename from tests/cpp/runtime/hexagon_buffer.cc rename to src/runtime/hexagon/hexagon_buffer_tests.cc index 715d9b1b695d..ee76ad0296e7 100644 --- a/tests/cpp/runtime/hexagon_buffer.cc +++ b/src/runtime/hexagon/hexagon_buffer_tests.cc @@ -18,7 +18,11 @@ */ #include +<<<<<<< HEAD:tests/cpp/runtime/hexagon_buffer.cc #include +======= +#include "hexagon_buffer.h" +>>>>>>> 499d5ee4d (HexagonBuffer tests running in sim):src/runtime/hexagon/hexagon/hexagon_buffer_tests.cc #include using namespace tvm::runtime; diff --git a/src/runtime/hexagon/hexagon_device_api_v2.cc b/src/runtime/hexagon/hexagon_device_api_v2.cc index 6640b3647d59..3f9fc608f89a 100644 --- a/src/runtime/hexagon/hexagon_device_api_v2.cc +++ b/src/runtime/hexagon/hexagon_device_api_v2.cc @@ -248,7 +248,7 @@ TVM_REGISTER_GLOBAL("device_api.hexagon.v2").set_body([](TVMArgs args, TVMRetVal }); TEST(HexagonDeviceApi, my_first_test) { - ASSERT_EQ(42, 42); + ASSERT_EQ(42, 41); } TVM_REGISTER_GLOBAL("hexagon.run_all_tests").set_body([](TVMArgs args, TVMRetValue* rv) { From 8a2ff6d779148d0f6e6946b299f9b22673a186d8 Mon Sep 17 00:00:00 2001 From: adstraw Date: Thu, 21 Apr 2022 11:38:19 -0700 Subject: [PATCH 04/20] move to new tests directory --- CMakeLists.txt | 1 - src/runtime/hexagon/hexagon_buffer_tests.cc | 5 + src/runtime/hexagon/hexagon_device_api_v2.cc | 11 - .../hexagon/tests/hexagon_buffer_tests.cc | 473 ++++++++++++++++++ src/runtime/hexagon/tests/run_all_tests.cc | 15 + 5 files changed, 493 insertions(+), 12 deletions(-) create mode 100644 src/runtime/hexagon/tests/hexagon_buffer_tests.cc create mode 100644 src/runtime/hexagon/tests/run_all_tests.cc diff --git a/CMakeLists.txt b/CMakeLists.txt index 481d95634c11..e3967fb61b1f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -603,7 +603,6 @@ FetchContent_Declare(googletest SOURCE_DIR "/local/mnt/workspace/Qualcomm/Hexago set(gtest_force_shared_crt ON CACHE BOOL "" FORCE) FetchContent_MakeAvailable(googletest) target_link_libraries(tvm_runtime PUBLIC gtest) -target_link_libraries(tvm PUBLIC gtest) # is this necessary if we move the tests out of device api? include_directories("/local/mnt/workspace/Qualcomm/Hexagon_SDK/4.5.0.3/utils/googletest/gtest/include") # Set flags for clang diff --git a/src/runtime/hexagon/hexagon_buffer_tests.cc b/src/runtime/hexagon/hexagon_buffer_tests.cc index ee76ad0296e7..4078ed9d7e37 100644 --- a/src/runtime/hexagon/hexagon_buffer_tests.cc +++ b/src/runtime/hexagon/hexagon_buffer_tests.cc @@ -17,12 +17,17 @@ * under the License. */ +#include "../hexagon_buffer.h" + #include +<<<<<<<< HEAD:src/runtime/hexagon/hexagon_buffer_tests.cc <<<<<<< HEAD:tests/cpp/runtime/hexagon_buffer.cc #include ======= #include "hexagon_buffer.h" >>>>>>> 499d5ee4d (HexagonBuffer tests running in sim):src/runtime/hexagon/hexagon/hexagon_buffer_tests.cc +======== +>>>>>>>> 2300f1904 (move to new tests directory):src/runtime/hexagon/hexagon/tests/hexagon_buffer_tests.cc #include using namespace tvm::runtime; diff --git a/src/runtime/hexagon/hexagon_device_api_v2.cc b/src/runtime/hexagon/hexagon_device_api_v2.cc index 3f9fc608f89a..8da66ad1d0b8 100644 --- a/src/runtime/hexagon/hexagon_device_api_v2.cc +++ b/src/runtime/hexagon/hexagon_device_api_v2.cc @@ -35,8 +35,6 @@ #include "hexagon_buffer.h" #include "hexagon_common.h" -#include "gtest/gtest.h" - namespace tvm { namespace runtime { namespace hexagon { @@ -247,15 +245,6 @@ TVM_REGISTER_GLOBAL("device_api.hexagon.v2").set_body([](TVMArgs args, TVMRetVal *rv = static_cast(ptr); }); -TEST(HexagonDeviceApi, my_first_test) { - ASSERT_EQ(42, 41); -} - -TVM_REGISTER_GLOBAL("hexagon.run_all_tests").set_body([](TVMArgs args, TVMRetValue* rv) { - *rv = RUN_ALL_TESTS(); -}); - - } // namespace hexagon } // namespace runtime } // namespace tvm diff --git a/src/runtime/hexagon/tests/hexagon_buffer_tests.cc b/src/runtime/hexagon/tests/hexagon_buffer_tests.cc new file mode 100644 index 000000000000..4078ed9d7e37 --- /dev/null +++ b/src/runtime/hexagon/tests/hexagon_buffer_tests.cc @@ -0,0 +1,473 @@ +/* + * Licensed to the Apache Software Foundation (ASF) under one + * or more contributor license agreements. See the NOTICE file + * distributed with this work for additional information + * regarding copyright ownership. The ASF licenses this file + * to you under the Apache License, Version 2.0 (the + * "License"); you may not use this file except in compliance + * with the License. You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, + * software distributed under the License is distributed on an + * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY + * KIND, either express or implied. See the License for the + * specific language governing permissions and limitations + * under the License. + */ + +#include "../hexagon_buffer.h" + +#include +<<<<<<<< HEAD:src/runtime/hexagon/hexagon_buffer_tests.cc +<<<<<<< HEAD:tests/cpp/runtime/hexagon_buffer.cc +#include +======= +#include "hexagon_buffer.h" +>>>>>>> 499d5ee4d (HexagonBuffer tests running in sim):src/runtime/hexagon/hexagon/hexagon_buffer_tests.cc +======== +>>>>>>>> 2300f1904 (move to new tests directory):src/runtime/hexagon/hexagon/tests/hexagon_buffer_tests.cc +#include + +using namespace tvm::runtime; +using namespace tvm::runtime::hexagon; + +TEST(HexagonBuffer, default_scope) { + Optional scope; + HexagonBuffer hb(8 /* nbytes */, 8 /* alignment */, scope); + EXPECT_EQ(hb.GetStorageScope(), HexagonBuffer::StorageScope::kDDR); +} + +TEST(HexagonBuffer, ddr_scope) { + Optional scope("global"); + HexagonBuffer hb(8 /* nbytes */, 8 /* alignment */, scope); + EXPECT_EQ(hb.GetStorageScope(), HexagonBuffer::StorageScope::kDDR); +} + +TEST(HexagonBuffer, vtcm_scope) { + Optional scope("global.vtcm"); + HexagonBuffer hb(8 /* nbytes */, 8 /* alignment */, scope); + EXPECT_EQ(hb.GetStorageScope(), HexagonBuffer::StorageScope::kVTCM); +} + +TEST(HexagonBuffer, invalid_scope) { + Optional scope("invalid"); + EXPECT_THROW(HexagonBuffer hb(8 /* nbytes */, 8 /* alignment */, scope), InternalError); +} + +TEST(HexagonBuffer, micro_copies_corresponding_regions) { + auto ptr = [](auto val) { return reinterpret_cast(val); }; + + std::vector src_ptr{ptr(0), ptr(16)}; + BufferSet src(src_ptr.data(), src_ptr.size(), 16); + + std::vector dest_ptr{ptr(64), ptr(80)}; + BufferSet dest(dest_ptr.data(), dest_ptr.size(), 16); + + auto micro_copies = BufferSet::MemoryCopies(dest, src, 32); + EXPECT_EQ(micro_copies.size(), 2); + for (size_t i = 0; i < micro_copies.size(); i++) { + EXPECT_EQ(micro_copies[i].src, ptr(16 * i)); + EXPECT_EQ(micro_copies[i].dest, ptr(64 + 16 * i)); + EXPECT_EQ(micro_copies[i].num_bytes, 16); + } +} + +TEST(HexagonBuffer, micro_copies_src_bigger) { + auto ptr = [](auto val) { return reinterpret_cast(val); }; + + std::vector src_ptr{ptr(0), ptr(16)}; + BufferSet src(src_ptr.data(), src_ptr.size(), 16); + + std::vector dest_ptr{ptr(64), ptr(72), ptr(80), ptr(88)}; + BufferSet dest(dest_ptr.data(), dest_ptr.size(), 8); + + auto micro_copies = BufferSet::MemoryCopies(dest, src, 32); + EXPECT_EQ(micro_copies.size(), 4); + for (size_t i = 0; i < micro_copies.size(); i++) { + EXPECT_EQ(micro_copies[i].src, ptr(8 * i)); + EXPECT_EQ(micro_copies[i].dest, ptr(64 + 8 * i)); + EXPECT_EQ(micro_copies[i].num_bytes, 8); + } +} + +TEST(HexagonBuffer, micro_copies_dest_bigger) { + auto ptr = [](auto val) { return reinterpret_cast(val); }; + + std::vector src_ptr{ptr(0), ptr(8), ptr(16), ptr(24)}; + BufferSet src(src_ptr.data(), src_ptr.size(), 8); + + std::vector dest_ptr{ptr(64), ptr(80)}; + BufferSet dest(dest_ptr.data(), dest_ptr.size(), 16); + + auto micro_copies = BufferSet::MemoryCopies(dest, src, 32); + EXPECT_EQ(micro_copies.size(), 4); + for (size_t i = 0; i < micro_copies.size(); i++) { + EXPECT_EQ(micro_copies[i].src, ptr(8 * i)); + EXPECT_EQ(micro_copies[i].dest, ptr(64 + 8 * i)); + EXPECT_EQ(micro_copies[i].num_bytes, 8); + } +} + +TEST(HexagonBuffer, micro_copies_src_overlaps_dest_region) { + auto ptr = [](auto val) { return reinterpret_cast(val); }; + + std::vector src_ptr{ptr(0), ptr(16)}; + BufferSet src(src_ptr.data(), src_ptr.size(), 16); + + std::vector dest_ptr{ptr(64), ptr(76)}; + BufferSet dest(dest_ptr.data(), dest_ptr.size(), 12); + + auto micro_copies = BufferSet::MemoryCopies(dest, src, 24); + EXPECT_EQ(micro_copies.size(), 3); + + // First region of source, first region of dest + EXPECT_EQ(micro_copies[0].src, ptr(0)); + EXPECT_EQ(micro_copies[0].dest, ptr(64)); + EXPECT_EQ(micro_copies[0].num_bytes, 12); + + // First region of source, second region of dest + EXPECT_EQ(micro_copies[1].src, ptr(12)); + EXPECT_EQ(micro_copies[1].dest, ptr(76)); + EXPECT_EQ(micro_copies[1].num_bytes, 4); + + // Second region of source, second region of dest + EXPECT_EQ(micro_copies[2].src, ptr(16)); + EXPECT_EQ(micro_copies[2].dest, ptr(80)); + EXPECT_EQ(micro_copies[2].num_bytes, 8); +} + +TEST(HexagonBuffer, micro_copies_dest_overlaps_src_region) { + auto ptr = [](auto val) { return reinterpret_cast(val); }; + + std::vector src_ptr{ptr(0), ptr(12)}; + BufferSet src(src_ptr.data(), src_ptr.size(), 12); + + std::vector dest_ptr{ptr(64), ptr(80)}; + BufferSet dest(dest_ptr.data(), dest_ptr.size(), 16); + + auto micro_copies = BufferSet::MemoryCopies(dest, src, 24); + EXPECT_EQ(micro_copies.size(), 3); + + // First region of source, first region of dest + EXPECT_EQ(micro_copies[0].src, ptr(0)); + EXPECT_EQ(micro_copies[0].dest, ptr(64)); + EXPECT_EQ(micro_copies[0].num_bytes, 12); + + // Second region of source, first region of dest + EXPECT_EQ(micro_copies[1].src, ptr(12)); + EXPECT_EQ(micro_copies[1].dest, ptr(76)); + EXPECT_EQ(micro_copies[1].num_bytes, 4); + + // Second region of source, second region of dest + EXPECT_EQ(micro_copies[2].src, ptr(16)); + EXPECT_EQ(micro_copies[2].dest, ptr(80)); + EXPECT_EQ(micro_copies[2].num_bytes, 8); +} + +TEST(HexagonBuffer, micro_copies_discontiguous_regions) { + auto ptr = [](auto val) { return reinterpret_cast(val); }; + + // Stride of 16, but only first 11 bytes in each region belong to + // this buffer. + std::vector src_ptr{ptr(0), ptr(16)}; + BufferSet src(src_ptr.data(), src_ptr.size(), 11); + + std::vector dest_ptr{ptr(64), ptr(80)}; + BufferSet dest(dest_ptr.data(), dest_ptr.size(), 13); + + auto micro_copies = BufferSet::MemoryCopies(dest, src, 16); + EXPECT_EQ(micro_copies.size(), 3); + + // First region of source, first region of dest + EXPECT_EQ(micro_copies[0].src, ptr(0)); + EXPECT_EQ(micro_copies[0].dest, ptr(64)); + EXPECT_EQ(micro_copies[0].num_bytes, 11); + + // Second region of source, first region of dest + EXPECT_EQ(micro_copies[1].src, ptr(16)); + EXPECT_EQ(micro_copies[1].dest, ptr(75)); + EXPECT_EQ(micro_copies[1].num_bytes, 2); + + // Second region of source, second region of dest + EXPECT_EQ(micro_copies[2].src, ptr(18)); + EXPECT_EQ(micro_copies[2].dest, ptr(80)); + EXPECT_EQ(micro_copies[2].num_bytes, 3); +} + +TEST(HexagonBuffer, micro_copies_invalid_size) { + auto ptr = [](auto val) { return reinterpret_cast(val); }; + + std::vector src_ptr{ptr(0), ptr(16)}; + std::vector dest_ptr{ptr(64), ptr(80)}; + + { + BufferSet src(src_ptr.data(), 1, 16); + BufferSet dest(dest_ptr.data(), 2, 16); + EXPECT_THROW(BufferSet::MemoryCopies(dest, src, 24), InternalError); + } + + { + BufferSet src(src_ptr.data(), 2, 16); + BufferSet dest(dest_ptr.data(), 1, 16); + EXPECT_THROW(BufferSet::MemoryCopies(dest, src, 24), InternalError); + } +} + +TEST(HexagonBuffer, macro_copies_adjacent_corresponding_regions_merged) { + auto ptr = [](auto val) { return reinterpret_cast(val); }; + + std::vector src_ptr{ptr(0), ptr(16)}; + BufferSet src(src_ptr.data(), src_ptr.size(), 16); + + std::vector dest_ptr{ptr(64), ptr(80)}; + BufferSet dest(dest_ptr.data(), dest_ptr.size(), 16); + + auto micro_copies = BufferSet::MemoryCopies(dest, src, 32); + auto macro_copies = MemoryCopy::MergeAdjacent(std::move(micro_copies)); + + ASSERT_EQ(macro_copies.size(), 1); + EXPECT_EQ(macro_copies[0].src, ptr(0)); + EXPECT_EQ(macro_copies[0].dest, ptr(64)); + EXPECT_EQ(macro_copies[0].num_bytes, 32); +} + +TEST(HexagonBuffer, macro_copies_discontiguous_regions_not_merged) { + auto ptr = [](auto val) { return reinterpret_cast(val); }; + + std::vector src_ptr{ptr(0), ptr(16)}; + BufferSet src(src_ptr.data(), src_ptr.size(), 12); + + std::vector dest_ptr{ptr(64), ptr(80)}; + BufferSet dest(dest_ptr.data(), dest_ptr.size(), 12); + + auto micro_copies = BufferSet::MemoryCopies(dest, src, 24); + auto macro_copies = MemoryCopy::MergeAdjacent(std::move(micro_copies)); + + ASSERT_EQ(macro_copies.size(), 2); + + EXPECT_EQ(macro_copies[0].src, ptr(0)); + EXPECT_EQ(macro_copies[0].dest, ptr(64)); + EXPECT_EQ(macro_copies[0].num_bytes, 12); + + EXPECT_EQ(macro_copies[1].src, ptr(16)); + EXPECT_EQ(macro_copies[1].dest, ptr(80)); + EXPECT_EQ(macro_copies[1].num_bytes, 12); +} + +TEST(HexagonBuffer, macro_copies_overlapping_regions_merged) { + auto ptr = [](auto val) { return reinterpret_cast(val); }; + + std::vector src_ptr{ptr(0), ptr(12)}; + BufferSet src(src_ptr.data(), src_ptr.size(), 12); + + std::vector dest_ptr{ptr(64), ptr(80)}; + BufferSet dest(dest_ptr.data(), dest_ptr.size(), 16); + + auto micro_copies = BufferSet::MemoryCopies(dest, src, 24); + auto macro_copies = MemoryCopy::MergeAdjacent(std::move(micro_copies)); + + ASSERT_EQ(macro_copies.size(), 1); + EXPECT_EQ(macro_copies[0].src, ptr(0)); + EXPECT_EQ(macro_copies[0].dest, ptr(64)); + EXPECT_EQ(macro_copies[0].num_bytes, 24); +} + +TEST(HexagonBuffer, copy_from) { + Optional scope("global"); + HexagonBuffer hb(8 /* nbytes */, 8 /* alignment */, scope); + + std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; + hb.CopyFrom(data.data(), data.size()); + + uint8_t* ptr = static_cast(hb.GetPointer()); + for (size_t i = 0; i < data.size(); ++i) { + EXPECT_EQ(ptr[i], data[i]); + } +} + +TEST(HexagonBuffer, copy_from_invalid_size) { + Optional scope("global"); + std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; + + // HexagonBuffer too small + HexagonBuffer toosmall(4 /* nbytes */, 8 /* alignment */, scope); + EXPECT_THROW(toosmall.CopyFrom(data.data(), data.size()), InternalError); +} + +TEST(HexagonBuffer, copy_from_smaller_size) { + Optional scope("global"); + std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; + + // HexagonBuffer is big + HexagonBuffer big(16 /* nbytes */, 16 /* alignment */, scope); + EXPECT_NO_THROW(big.CopyFrom(data.data(), data.size())); +} + +TEST(HexagonBuffer, nd) { + Optional def; + HexagonBuffer hb_default(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, def); + EXPECT_EQ(hb_default.GetStorageScope(), HexagonBuffer::StorageScope::kDDR); + + Optional global("global"); + HexagonBuffer hb_global(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, global); + EXPECT_EQ(hb_global.GetStorageScope(), HexagonBuffer::StorageScope::kDDR); + + Optional vtcm("global.vtcm"); + HexagonBuffer hb_vtcm(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, vtcm); + EXPECT_EQ(hb_vtcm.GetStorageScope(), HexagonBuffer::StorageScope::kVTCM); + + Optional invalid("invalid"); + EXPECT_THROW(HexagonBuffer hb_invalid(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, invalid), + InternalError); +} + +TEST(HexagonBuffer, nd_copy_from) { + Optional scope("global"); + HexagonBuffer hb(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, scope); + + std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; + hb.CopyFrom(data.data(), data.size()); + + uint8_t** ptr = static_cast(hb.GetPointer()); + EXPECT_EQ(ptr[0][0], data[0]); + EXPECT_EQ(ptr[0][1], data[1]); + EXPECT_EQ(ptr[0][2], data[2]); + EXPECT_EQ(ptr[0][3], data[3]); + EXPECT_EQ(ptr[1][0], data[4]); + EXPECT_EQ(ptr[1][1], data[5]); + EXPECT_EQ(ptr[1][2], data[6]); + EXPECT_EQ(ptr[1][3], data[7]); +} + +TEST(HexagonBuffer, 1d_copy_from_1d) { + Optional global("global"); + HexagonBuffer from(8 /* nbytes */, 8 /* alignment */, global); + + Optional vtcm("global.vtcm"); + HexagonBuffer to(8 /* nbytes */, 8 /* alignment */, vtcm); + + std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; + from.CopyFrom(data.data(), data.size()); + to.CopyFrom(from, 8); + + uint8_t* ptr = static_cast(to.GetPointer()); + for (size_t i = 0; i < data.size(); ++i) { + EXPECT_EQ(ptr[i], data[i]); + } +} + +TEST(HexagonBuffer, 2d_copy_from_1d) { + Optional vtcm("global.vtcm"); + HexagonBuffer hb1d(8 /* nbytes */, 8 /* alignment */, vtcm); + + Optional global("global"); + HexagonBuffer hb2d(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, global); + + std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; + hb1d.CopyFrom(data.data(), data.size()); + hb2d.CopyFrom(hb1d, 8); + + uint8_t** ptr = static_cast(hb2d.GetPointer()); + EXPECT_EQ(ptr[0][0], data[0]); + EXPECT_EQ(ptr[0][1], data[1]); + EXPECT_EQ(ptr[0][2], data[2]); + EXPECT_EQ(ptr[0][3], data[3]); + EXPECT_EQ(ptr[1][0], data[4]); + EXPECT_EQ(ptr[1][1], data[5]); + EXPECT_EQ(ptr[1][2], data[6]); + EXPECT_EQ(ptr[1][3], data[7]); +} + +TEST(HexagonBuffer, 1d_copy_from_2d) { + Optional vtcm("global.vtcm"); + HexagonBuffer hb2d(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, vtcm); + + Optional global("global.vtcm"); + HexagonBuffer hb1d(8 /* nbytes */, 8 /* alignment */, global); + + std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; + hb2d.CopyFrom(data.data(), data.size()); + hb1d.CopyFrom(hb2d, 8); + + uint8_t* ptr = static_cast(hb1d.GetPointer()); + for (size_t i = 0; i < data.size(); ++i) { + EXPECT_EQ(ptr[i], data[i]); + } +} + +TEST(HexagonBuffer, nd_copy_from_nd_invalid_size) { + Optional scope("global"); + HexagonBuffer hb1d(8 /* nbytes */, 8 /* alignment */, scope); + HexagonBuffer hb2d(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, scope); + + HexagonBuffer toosbig1d(16 /* nbytes */, 16 /* alignment */, scope); + EXPECT_THROW(hb1d.CopyFrom(toosbig1d, 16), InternalError); + EXPECT_THROW(hb2d.CopyFrom(toosbig1d, 16), InternalError); + + HexagonBuffer toobig2d(2 /* ndim */, 16 /* nbytes */, 16 /* alignment */, scope); + EXPECT_THROW(hb1d.CopyFrom(toobig2d, 32), InternalError); + EXPECT_THROW(hb2d.CopyFrom(toobig2d, 32), InternalError); +} + +TEST(HexagonBuffer, nd_copy_from_nd_smaller_size) { + Optional scope("global"); + HexagonBuffer hb1d(8 /* nbytes */, 8 /* alignment */, scope); + HexagonBuffer hb2d(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, scope); + + HexagonBuffer small1d(4 /* nbytes */, 8 /* alignment */, scope); + EXPECT_NO_THROW(hb1d.CopyFrom(small1d, 4)); + EXPECT_NO_THROW(hb2d.CopyFrom(small1d, 4)); + + HexagonBuffer small2d(2 /* ndim */, 2 /* nbytes */, 8 /* alignment */, scope); + EXPECT_NO_THROW(hb1d.CopyFrom(small2d, 4)); + EXPECT_NO_THROW(hb2d.CopyFrom(small2d, 4)); +} + +TEST(HexagonBuffer, md_copy_from_nd) { + Optional scope("global"); + HexagonBuffer hb3d(3 /* ndim */, 4 /* nbytes */, 8 /* alignment */, scope); + HexagonBuffer hb4d(4 /* ndim */, 3 /* nbytes */, 8 /* alignment */, scope); + + std::vector data{0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11}; + + hb3d.CopyFrom(data.data(), data.size()); + hb4d.CopyFrom(hb3d, data.size()); + + uint8_t** hb3d_ptr = static_cast(hb3d.GetPointer()); + uint8_t** hb4d_ptr = static_cast(hb4d.GetPointer()); + for (size_t i = 0; i < 12; i++) { + EXPECT_EQ(hb3d_ptr[i / 4][i % 4], hb4d_ptr[i / 3][i % 3]); + } +} + +TEST(HexagonBuffer, copy_to) { + Optional scope("global"); + HexagonBuffer hb(8 /* nbytes */, 8 /* alignment */, scope); + + std::vector data_in{0, 1, 2, 3, 4, 5, 6, 7}; + hb.CopyFrom(data_in.data(), data_in.size()); + + std::vector data_out{7, 6, 5, 4, 3, 2, 1, 0}; + hb.CopyTo(data_out.data(), data_out.size()); + + for (size_t i = 0; i < data_in.size(); ++i) { + EXPECT_EQ(data_in[i], data_out[i]); + } +} + +TEST(HexagonBuffer, nd_copy_to) { + Optional scope("global"); + HexagonBuffer hb(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, scope); + + std::vector data_in{0, 1, 2, 3, 4, 5, 6, 7}; + hb.CopyFrom(data_in.data(), data_in.size()); + + std::vector data_out{7, 6, 5, 4, 3, 2, 1, 0}; + hb.CopyTo(data_out.data(), data_out.size()); + + for (size_t i = 0; i < data_in.size(); ++i) { + EXPECT_EQ(data_in[i], data_out[i]); + } +} diff --git a/src/runtime/hexagon/tests/run_all_tests.cc b/src/runtime/hexagon/tests/run_all_tests.cc new file mode 100644 index 000000000000..61f3dc1d4909 --- /dev/null +++ b/src/runtime/hexagon/tests/run_all_tests.cc @@ -0,0 +1,15 @@ +#include +#include + +#include "gtest/gtest.h" + +namespace tvm { +namespace runtime { +namespace hexagon { + + +TVM_REGISTER_GLOBAL("hexagon.run_all_tests").set_body([](TVMArgs args, TVMRetValue* rv) { + *rv = RUN_ALL_TESTS(); +}); + +}}} \ No newline at end of file From a85c333c3fb34580d412edf3627e9faafd08b48a Mon Sep 17 00:00:00 2001 From: adstraw Date: Thu, 21 Apr 2022 14:49:14 -0700 Subject: [PATCH 05/20] use USE_HEXAGON_SDK --- CMakeLists.txt | 5 +++-- apps/hexagon_api/CMakeLists.txt | 1 + 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index e3967fb61b1f..7871f1ab3a5d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -599,11 +599,12 @@ target_link_libraries(tvm PRIVATE ${TVM_LINKER_LIBS} ${TVM_RUNTIME_LINKER_LIBS}) target_link_libraries(tvm_runtime PRIVATE ${TVM_RUNTIME_LINKER_LIBS}) include(FetchContent) -FetchContent_Declare(googletest SOURCE_DIR "/local/mnt/workspace/Qualcomm/Hexagon_SDK/4.5.0.3/utils/googletest/gtest") +FetchContent_Declare(googletest SOURCE_DIR "${USE_HEXAGON_SDK}/utils/googletest/gtest") set(gtest_force_shared_crt ON CACHE BOOL "" FORCE) FetchContent_MakeAvailable(googletest) target_link_libraries(tvm_runtime PUBLIC gtest) -include_directories("/local/mnt/workspace/Qualcomm/Hexagon_SDK/4.5.0.3/utils/googletest/gtest/include") +target_link_libraries(tvm PUBLIC gtest) # workaround +include_directories("${USE_HEXAGON_SDK}/utils/googletest/gtest/include") # Set flags for clang include(cmake/modules/ClangFlags.cmake) diff --git a/apps/hexagon_api/CMakeLists.txt b/apps/hexagon_api/CMakeLists.txt index 40f070513e3d..8d83cebbc540 100644 --- a/apps/hexagon_api/CMakeLists.txt +++ b/apps/hexagon_api/CMakeLists.txt @@ -39,6 +39,7 @@ ExternalProject_Add(x86_tvm_runtime_rpc "-DUSE_HEXAGON_RPC=ON" "-DBUILD_STATIC_RUNTIME=ON" "-DCMAKE_BUILD_TYPE=${CMAKE_BUILD_TYPE}" + "-DUSE_HEXAGON_SDK=${USE_HEXAGON_SDK}" INSTALL_COMMAND "" BUILD_ALWAYS ON ) From b4db712d507296cfb69ca11b6ccfef5963e4d343 Mon Sep 17 00:00:00 2001 From: adstraw Date: Thu, 21 Apr 2022 14:58:03 -0700 Subject: [PATCH 06/20] add python frontend for Hexagon unit tests --- .../test_hexagon/test_cache_read_write.py | 37 ++++++++++++++++--- .../python/contrib/test_hexagon/unit_tests.py | 11 ++++++ 2 files changed, 43 insertions(+), 5 deletions(-) create mode 100644 tests/python/contrib/test_hexagon/unit_tests.py diff --git a/tests/python/contrib/test_hexagon/test_cache_read_write.py b/tests/python/contrib/test_hexagon/test_cache_read_write.py index a67a14687d1c..8f9453187169 100644 --- a/tests/python/contrib/test_hexagon/test_cache_read_write.py +++ b/tests/python/contrib/test_hexagon/test_cache_read_write.py @@ -144,8 +144,35 @@ def layout_transform_2d(n): @requires_hexagon_toolchain def test_cache_read_write_2d(hexagon_session): - func = hexagon_session._rpc.get_function("hexagon.run_all_tests") - x = func() - print("PRETTY SURE I JUST RAN MY TEST, PRINTING THE RETURN VALUE") - print(x) - np.testing.assert_equal(x, 0) + size = 128 + outer_shape = (size,) + factor = 16 + inner_shape = (factor,) + dtype = "int8" + + x = te.placeholder(shape=outer_shape, dtype=dtype, name="x") + y = te.placeholder(shape=outer_shape, dtype=dtype, name="y") + z = te.compute(outer_shape, lambda i: x[i] + y[i], name="z") + s = te.create_schedule(z.op) + + x_vtcm = s.cache_read(x, "global.vtcm", [z]) + y_vtcm = s.cache_read(y, "global.vtcm", [z]) + z_vtcm = s.cache_write(z, "global.vtcm") + + layout_x_vtcm = s[x_vtcm].transform_layout(layout_transform_2d) + layout_y_vtcm = s[y_vtcm].transform_layout(layout_transform_2d) + layout_z_vtcm = s[z_vtcm].transform_layout(layout_transform_2d) + + mem_copy_read = intrin_mem_copy(inner_shape, dtype, "global.vtcm", "global") + s[x_vtcm].tensorize(layout_x_vtcm[1], mem_copy_read) + s[y_vtcm].tensorize(layout_y_vtcm[1], mem_copy_read) + + # The loop schedule over `z` is not modified when calling `transform_layout` + # on `z_vtcm` above therefore we must call `split` to modify the loop schedule + # over `z` to match the layout of `z_vtcm` such that we can accurately write + # `z_vtcm` back to `z` using memory copy intrinsic + zouter, zinner = s[z].split(z.op.axis[0], factor=factor) + mem_copy_write = intrin_mem_copy(inner_shape, dtype, "global", "global.vtcm") + s[z].tensorize(zinner, mem_copy_write) + + verify(hexagon_session, s, x, y, z, size) diff --git a/tests/python/contrib/test_hexagon/unit_tests.py b/tests/python/contrib/test_hexagon/unit_tests.py new file mode 100644 index 000000000000..435c40a75ea6 --- /dev/null +++ b/tests/python/contrib/test_hexagon/unit_tests.py @@ -0,0 +1,11 @@ +import pytest +import numpy as np +from tvm.contrib.hexagon.build import HexagonLauncher +#import tvm.contrib.hexagon as hexagon +from .conftest import requires_hexagon_toolchain + +@requires_hexagon_toolchain +def test_cache_read_write_2d(hexagon_session): + func = hexagon_session._rpc.get_function("hexagon.run_all_tests") + x = func() + np.testing.assert_equal(x, 0) From 8ec34af823c7b2a21cf00cc87066ef2a89de5d4d Mon Sep 17 00:00:00 2001 From: adstraw Date: Fri, 22 Apr 2022 08:55:50 -0700 Subject: [PATCH 07/20] clean up after rebase --- cmake/modules/Hexagon.cmake | 1 + src/runtime/hexagon/hexagon_buffer_tests.cc | 473 ------------------ .../hexagon/tests/hexagon_buffer_tests.cc | 11 +- 3 files changed, 2 insertions(+), 483 deletions(-) delete mode 100644 src/runtime/hexagon/hexagon_buffer_tests.cc diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index 3b0ff7dfeae3..8a6a0c3acf65 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -123,6 +123,7 @@ endfunction() # Common sources for TVM runtime with Hexagon support file_glob_append(RUNTIME_HEXAGON_SRCS "${TVMRT_SOURCE_DIR}/hexagon/*.cc" + "${TVMRT_SOURCE_DIR}/hexagon/tests/*.cc" ) diff --git a/src/runtime/hexagon/hexagon_buffer_tests.cc b/src/runtime/hexagon/hexagon_buffer_tests.cc deleted file mode 100644 index 4078ed9d7e37..000000000000 --- a/src/runtime/hexagon/hexagon_buffer_tests.cc +++ /dev/null @@ -1,473 +0,0 @@ -/* - * Licensed to the Apache Software Foundation (ASF) under one - * or more contributor license agreements. See the NOTICE file - * distributed with this work for additional information - * regarding copyright ownership. The ASF licenses this file - * to you under the Apache License, Version 2.0 (the - * "License"); you may not use this file except in compliance - * with the License. You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, - * software distributed under the License is distributed on an - * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY - * KIND, either express or implied. See the License for the - * specific language governing permissions and limitations - * under the License. - */ - -#include "../hexagon_buffer.h" - -#include -<<<<<<<< HEAD:src/runtime/hexagon/hexagon_buffer_tests.cc -<<<<<<< HEAD:tests/cpp/runtime/hexagon_buffer.cc -#include -======= -#include "hexagon_buffer.h" ->>>>>>> 499d5ee4d (HexagonBuffer tests running in sim):src/runtime/hexagon/hexagon/hexagon_buffer_tests.cc -======== ->>>>>>>> 2300f1904 (move to new tests directory):src/runtime/hexagon/hexagon/tests/hexagon_buffer_tests.cc -#include - -using namespace tvm::runtime; -using namespace tvm::runtime::hexagon; - -TEST(HexagonBuffer, default_scope) { - Optional scope; - HexagonBuffer hb(8 /* nbytes */, 8 /* alignment */, scope); - EXPECT_EQ(hb.GetStorageScope(), HexagonBuffer::StorageScope::kDDR); -} - -TEST(HexagonBuffer, ddr_scope) { - Optional scope("global"); - HexagonBuffer hb(8 /* nbytes */, 8 /* alignment */, scope); - EXPECT_EQ(hb.GetStorageScope(), HexagonBuffer::StorageScope::kDDR); -} - -TEST(HexagonBuffer, vtcm_scope) { - Optional scope("global.vtcm"); - HexagonBuffer hb(8 /* nbytes */, 8 /* alignment */, scope); - EXPECT_EQ(hb.GetStorageScope(), HexagonBuffer::StorageScope::kVTCM); -} - -TEST(HexagonBuffer, invalid_scope) { - Optional scope("invalid"); - EXPECT_THROW(HexagonBuffer hb(8 /* nbytes */, 8 /* alignment */, scope), InternalError); -} - -TEST(HexagonBuffer, micro_copies_corresponding_regions) { - auto ptr = [](auto val) { return reinterpret_cast(val); }; - - std::vector src_ptr{ptr(0), ptr(16)}; - BufferSet src(src_ptr.data(), src_ptr.size(), 16); - - std::vector dest_ptr{ptr(64), ptr(80)}; - BufferSet dest(dest_ptr.data(), dest_ptr.size(), 16); - - auto micro_copies = BufferSet::MemoryCopies(dest, src, 32); - EXPECT_EQ(micro_copies.size(), 2); - for (size_t i = 0; i < micro_copies.size(); i++) { - EXPECT_EQ(micro_copies[i].src, ptr(16 * i)); - EXPECT_EQ(micro_copies[i].dest, ptr(64 + 16 * i)); - EXPECT_EQ(micro_copies[i].num_bytes, 16); - } -} - -TEST(HexagonBuffer, micro_copies_src_bigger) { - auto ptr = [](auto val) { return reinterpret_cast(val); }; - - std::vector src_ptr{ptr(0), ptr(16)}; - BufferSet src(src_ptr.data(), src_ptr.size(), 16); - - std::vector dest_ptr{ptr(64), ptr(72), ptr(80), ptr(88)}; - BufferSet dest(dest_ptr.data(), dest_ptr.size(), 8); - - auto micro_copies = BufferSet::MemoryCopies(dest, src, 32); - EXPECT_EQ(micro_copies.size(), 4); - for (size_t i = 0; i < micro_copies.size(); i++) { - EXPECT_EQ(micro_copies[i].src, ptr(8 * i)); - EXPECT_EQ(micro_copies[i].dest, ptr(64 + 8 * i)); - EXPECT_EQ(micro_copies[i].num_bytes, 8); - } -} - -TEST(HexagonBuffer, micro_copies_dest_bigger) { - auto ptr = [](auto val) { return reinterpret_cast(val); }; - - std::vector src_ptr{ptr(0), ptr(8), ptr(16), ptr(24)}; - BufferSet src(src_ptr.data(), src_ptr.size(), 8); - - std::vector dest_ptr{ptr(64), ptr(80)}; - BufferSet dest(dest_ptr.data(), dest_ptr.size(), 16); - - auto micro_copies = BufferSet::MemoryCopies(dest, src, 32); - EXPECT_EQ(micro_copies.size(), 4); - for (size_t i = 0; i < micro_copies.size(); i++) { - EXPECT_EQ(micro_copies[i].src, ptr(8 * i)); - EXPECT_EQ(micro_copies[i].dest, ptr(64 + 8 * i)); - EXPECT_EQ(micro_copies[i].num_bytes, 8); - } -} - -TEST(HexagonBuffer, micro_copies_src_overlaps_dest_region) { - auto ptr = [](auto val) { return reinterpret_cast(val); }; - - std::vector src_ptr{ptr(0), ptr(16)}; - BufferSet src(src_ptr.data(), src_ptr.size(), 16); - - std::vector dest_ptr{ptr(64), ptr(76)}; - BufferSet dest(dest_ptr.data(), dest_ptr.size(), 12); - - auto micro_copies = BufferSet::MemoryCopies(dest, src, 24); - EXPECT_EQ(micro_copies.size(), 3); - - // First region of source, first region of dest - EXPECT_EQ(micro_copies[0].src, ptr(0)); - EXPECT_EQ(micro_copies[0].dest, ptr(64)); - EXPECT_EQ(micro_copies[0].num_bytes, 12); - - // First region of source, second region of dest - EXPECT_EQ(micro_copies[1].src, ptr(12)); - EXPECT_EQ(micro_copies[1].dest, ptr(76)); - EXPECT_EQ(micro_copies[1].num_bytes, 4); - - // Second region of source, second region of dest - EXPECT_EQ(micro_copies[2].src, ptr(16)); - EXPECT_EQ(micro_copies[2].dest, ptr(80)); - EXPECT_EQ(micro_copies[2].num_bytes, 8); -} - -TEST(HexagonBuffer, micro_copies_dest_overlaps_src_region) { - auto ptr = [](auto val) { return reinterpret_cast(val); }; - - std::vector src_ptr{ptr(0), ptr(12)}; - BufferSet src(src_ptr.data(), src_ptr.size(), 12); - - std::vector dest_ptr{ptr(64), ptr(80)}; - BufferSet dest(dest_ptr.data(), dest_ptr.size(), 16); - - auto micro_copies = BufferSet::MemoryCopies(dest, src, 24); - EXPECT_EQ(micro_copies.size(), 3); - - // First region of source, first region of dest - EXPECT_EQ(micro_copies[0].src, ptr(0)); - EXPECT_EQ(micro_copies[0].dest, ptr(64)); - EXPECT_EQ(micro_copies[0].num_bytes, 12); - - // Second region of source, first region of dest - EXPECT_EQ(micro_copies[1].src, ptr(12)); - EXPECT_EQ(micro_copies[1].dest, ptr(76)); - EXPECT_EQ(micro_copies[1].num_bytes, 4); - - // Second region of source, second region of dest - EXPECT_EQ(micro_copies[2].src, ptr(16)); - EXPECT_EQ(micro_copies[2].dest, ptr(80)); - EXPECT_EQ(micro_copies[2].num_bytes, 8); -} - -TEST(HexagonBuffer, micro_copies_discontiguous_regions) { - auto ptr = [](auto val) { return reinterpret_cast(val); }; - - // Stride of 16, but only first 11 bytes in each region belong to - // this buffer. - std::vector src_ptr{ptr(0), ptr(16)}; - BufferSet src(src_ptr.data(), src_ptr.size(), 11); - - std::vector dest_ptr{ptr(64), ptr(80)}; - BufferSet dest(dest_ptr.data(), dest_ptr.size(), 13); - - auto micro_copies = BufferSet::MemoryCopies(dest, src, 16); - EXPECT_EQ(micro_copies.size(), 3); - - // First region of source, first region of dest - EXPECT_EQ(micro_copies[0].src, ptr(0)); - EXPECT_EQ(micro_copies[0].dest, ptr(64)); - EXPECT_EQ(micro_copies[0].num_bytes, 11); - - // Second region of source, first region of dest - EXPECT_EQ(micro_copies[1].src, ptr(16)); - EXPECT_EQ(micro_copies[1].dest, ptr(75)); - EXPECT_EQ(micro_copies[1].num_bytes, 2); - - // Second region of source, second region of dest - EXPECT_EQ(micro_copies[2].src, ptr(18)); - EXPECT_EQ(micro_copies[2].dest, ptr(80)); - EXPECT_EQ(micro_copies[2].num_bytes, 3); -} - -TEST(HexagonBuffer, micro_copies_invalid_size) { - auto ptr = [](auto val) { return reinterpret_cast(val); }; - - std::vector src_ptr{ptr(0), ptr(16)}; - std::vector dest_ptr{ptr(64), ptr(80)}; - - { - BufferSet src(src_ptr.data(), 1, 16); - BufferSet dest(dest_ptr.data(), 2, 16); - EXPECT_THROW(BufferSet::MemoryCopies(dest, src, 24), InternalError); - } - - { - BufferSet src(src_ptr.data(), 2, 16); - BufferSet dest(dest_ptr.data(), 1, 16); - EXPECT_THROW(BufferSet::MemoryCopies(dest, src, 24), InternalError); - } -} - -TEST(HexagonBuffer, macro_copies_adjacent_corresponding_regions_merged) { - auto ptr = [](auto val) { return reinterpret_cast(val); }; - - std::vector src_ptr{ptr(0), ptr(16)}; - BufferSet src(src_ptr.data(), src_ptr.size(), 16); - - std::vector dest_ptr{ptr(64), ptr(80)}; - BufferSet dest(dest_ptr.data(), dest_ptr.size(), 16); - - auto micro_copies = BufferSet::MemoryCopies(dest, src, 32); - auto macro_copies = MemoryCopy::MergeAdjacent(std::move(micro_copies)); - - ASSERT_EQ(macro_copies.size(), 1); - EXPECT_EQ(macro_copies[0].src, ptr(0)); - EXPECT_EQ(macro_copies[0].dest, ptr(64)); - EXPECT_EQ(macro_copies[0].num_bytes, 32); -} - -TEST(HexagonBuffer, macro_copies_discontiguous_regions_not_merged) { - auto ptr = [](auto val) { return reinterpret_cast(val); }; - - std::vector src_ptr{ptr(0), ptr(16)}; - BufferSet src(src_ptr.data(), src_ptr.size(), 12); - - std::vector dest_ptr{ptr(64), ptr(80)}; - BufferSet dest(dest_ptr.data(), dest_ptr.size(), 12); - - auto micro_copies = BufferSet::MemoryCopies(dest, src, 24); - auto macro_copies = MemoryCopy::MergeAdjacent(std::move(micro_copies)); - - ASSERT_EQ(macro_copies.size(), 2); - - EXPECT_EQ(macro_copies[0].src, ptr(0)); - EXPECT_EQ(macro_copies[0].dest, ptr(64)); - EXPECT_EQ(macro_copies[0].num_bytes, 12); - - EXPECT_EQ(macro_copies[1].src, ptr(16)); - EXPECT_EQ(macro_copies[1].dest, ptr(80)); - EXPECT_EQ(macro_copies[1].num_bytes, 12); -} - -TEST(HexagonBuffer, macro_copies_overlapping_regions_merged) { - auto ptr = [](auto val) { return reinterpret_cast(val); }; - - std::vector src_ptr{ptr(0), ptr(12)}; - BufferSet src(src_ptr.data(), src_ptr.size(), 12); - - std::vector dest_ptr{ptr(64), ptr(80)}; - BufferSet dest(dest_ptr.data(), dest_ptr.size(), 16); - - auto micro_copies = BufferSet::MemoryCopies(dest, src, 24); - auto macro_copies = MemoryCopy::MergeAdjacent(std::move(micro_copies)); - - ASSERT_EQ(macro_copies.size(), 1); - EXPECT_EQ(macro_copies[0].src, ptr(0)); - EXPECT_EQ(macro_copies[0].dest, ptr(64)); - EXPECT_EQ(macro_copies[0].num_bytes, 24); -} - -TEST(HexagonBuffer, copy_from) { - Optional scope("global"); - HexagonBuffer hb(8 /* nbytes */, 8 /* alignment */, scope); - - std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; - hb.CopyFrom(data.data(), data.size()); - - uint8_t* ptr = static_cast(hb.GetPointer()); - for (size_t i = 0; i < data.size(); ++i) { - EXPECT_EQ(ptr[i], data[i]); - } -} - -TEST(HexagonBuffer, copy_from_invalid_size) { - Optional scope("global"); - std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; - - // HexagonBuffer too small - HexagonBuffer toosmall(4 /* nbytes */, 8 /* alignment */, scope); - EXPECT_THROW(toosmall.CopyFrom(data.data(), data.size()), InternalError); -} - -TEST(HexagonBuffer, copy_from_smaller_size) { - Optional scope("global"); - std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; - - // HexagonBuffer is big - HexagonBuffer big(16 /* nbytes */, 16 /* alignment */, scope); - EXPECT_NO_THROW(big.CopyFrom(data.data(), data.size())); -} - -TEST(HexagonBuffer, nd) { - Optional def; - HexagonBuffer hb_default(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, def); - EXPECT_EQ(hb_default.GetStorageScope(), HexagonBuffer::StorageScope::kDDR); - - Optional global("global"); - HexagonBuffer hb_global(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, global); - EXPECT_EQ(hb_global.GetStorageScope(), HexagonBuffer::StorageScope::kDDR); - - Optional vtcm("global.vtcm"); - HexagonBuffer hb_vtcm(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, vtcm); - EXPECT_EQ(hb_vtcm.GetStorageScope(), HexagonBuffer::StorageScope::kVTCM); - - Optional invalid("invalid"); - EXPECT_THROW(HexagonBuffer hb_invalid(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, invalid), - InternalError); -} - -TEST(HexagonBuffer, nd_copy_from) { - Optional scope("global"); - HexagonBuffer hb(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, scope); - - std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; - hb.CopyFrom(data.data(), data.size()); - - uint8_t** ptr = static_cast(hb.GetPointer()); - EXPECT_EQ(ptr[0][0], data[0]); - EXPECT_EQ(ptr[0][1], data[1]); - EXPECT_EQ(ptr[0][2], data[2]); - EXPECT_EQ(ptr[0][3], data[3]); - EXPECT_EQ(ptr[1][0], data[4]); - EXPECT_EQ(ptr[1][1], data[5]); - EXPECT_EQ(ptr[1][2], data[6]); - EXPECT_EQ(ptr[1][3], data[7]); -} - -TEST(HexagonBuffer, 1d_copy_from_1d) { - Optional global("global"); - HexagonBuffer from(8 /* nbytes */, 8 /* alignment */, global); - - Optional vtcm("global.vtcm"); - HexagonBuffer to(8 /* nbytes */, 8 /* alignment */, vtcm); - - std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; - from.CopyFrom(data.data(), data.size()); - to.CopyFrom(from, 8); - - uint8_t* ptr = static_cast(to.GetPointer()); - for (size_t i = 0; i < data.size(); ++i) { - EXPECT_EQ(ptr[i], data[i]); - } -} - -TEST(HexagonBuffer, 2d_copy_from_1d) { - Optional vtcm("global.vtcm"); - HexagonBuffer hb1d(8 /* nbytes */, 8 /* alignment */, vtcm); - - Optional global("global"); - HexagonBuffer hb2d(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, global); - - std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; - hb1d.CopyFrom(data.data(), data.size()); - hb2d.CopyFrom(hb1d, 8); - - uint8_t** ptr = static_cast(hb2d.GetPointer()); - EXPECT_EQ(ptr[0][0], data[0]); - EXPECT_EQ(ptr[0][1], data[1]); - EXPECT_EQ(ptr[0][2], data[2]); - EXPECT_EQ(ptr[0][3], data[3]); - EXPECT_EQ(ptr[1][0], data[4]); - EXPECT_EQ(ptr[1][1], data[5]); - EXPECT_EQ(ptr[1][2], data[6]); - EXPECT_EQ(ptr[1][3], data[7]); -} - -TEST(HexagonBuffer, 1d_copy_from_2d) { - Optional vtcm("global.vtcm"); - HexagonBuffer hb2d(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, vtcm); - - Optional global("global.vtcm"); - HexagonBuffer hb1d(8 /* nbytes */, 8 /* alignment */, global); - - std::vector data{0, 1, 2, 3, 4, 5, 6, 7}; - hb2d.CopyFrom(data.data(), data.size()); - hb1d.CopyFrom(hb2d, 8); - - uint8_t* ptr = static_cast(hb1d.GetPointer()); - for (size_t i = 0; i < data.size(); ++i) { - EXPECT_EQ(ptr[i], data[i]); - } -} - -TEST(HexagonBuffer, nd_copy_from_nd_invalid_size) { - Optional scope("global"); - HexagonBuffer hb1d(8 /* nbytes */, 8 /* alignment */, scope); - HexagonBuffer hb2d(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, scope); - - HexagonBuffer toosbig1d(16 /* nbytes */, 16 /* alignment */, scope); - EXPECT_THROW(hb1d.CopyFrom(toosbig1d, 16), InternalError); - EXPECT_THROW(hb2d.CopyFrom(toosbig1d, 16), InternalError); - - HexagonBuffer toobig2d(2 /* ndim */, 16 /* nbytes */, 16 /* alignment */, scope); - EXPECT_THROW(hb1d.CopyFrom(toobig2d, 32), InternalError); - EXPECT_THROW(hb2d.CopyFrom(toobig2d, 32), InternalError); -} - -TEST(HexagonBuffer, nd_copy_from_nd_smaller_size) { - Optional scope("global"); - HexagonBuffer hb1d(8 /* nbytes */, 8 /* alignment */, scope); - HexagonBuffer hb2d(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, scope); - - HexagonBuffer small1d(4 /* nbytes */, 8 /* alignment */, scope); - EXPECT_NO_THROW(hb1d.CopyFrom(small1d, 4)); - EXPECT_NO_THROW(hb2d.CopyFrom(small1d, 4)); - - HexagonBuffer small2d(2 /* ndim */, 2 /* nbytes */, 8 /* alignment */, scope); - EXPECT_NO_THROW(hb1d.CopyFrom(small2d, 4)); - EXPECT_NO_THROW(hb2d.CopyFrom(small2d, 4)); -} - -TEST(HexagonBuffer, md_copy_from_nd) { - Optional scope("global"); - HexagonBuffer hb3d(3 /* ndim */, 4 /* nbytes */, 8 /* alignment */, scope); - HexagonBuffer hb4d(4 /* ndim */, 3 /* nbytes */, 8 /* alignment */, scope); - - std::vector data{0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11}; - - hb3d.CopyFrom(data.data(), data.size()); - hb4d.CopyFrom(hb3d, data.size()); - - uint8_t** hb3d_ptr = static_cast(hb3d.GetPointer()); - uint8_t** hb4d_ptr = static_cast(hb4d.GetPointer()); - for (size_t i = 0; i < 12; i++) { - EXPECT_EQ(hb3d_ptr[i / 4][i % 4], hb4d_ptr[i / 3][i % 3]); - } -} - -TEST(HexagonBuffer, copy_to) { - Optional scope("global"); - HexagonBuffer hb(8 /* nbytes */, 8 /* alignment */, scope); - - std::vector data_in{0, 1, 2, 3, 4, 5, 6, 7}; - hb.CopyFrom(data_in.data(), data_in.size()); - - std::vector data_out{7, 6, 5, 4, 3, 2, 1, 0}; - hb.CopyTo(data_out.data(), data_out.size()); - - for (size_t i = 0; i < data_in.size(); ++i) { - EXPECT_EQ(data_in[i], data_out[i]); - } -} - -TEST(HexagonBuffer, nd_copy_to) { - Optional scope("global"); - HexagonBuffer hb(2 /* ndim */, 4 /* nbytes */, 8 /* alignment */, scope); - - std::vector data_in{0, 1, 2, 3, 4, 5, 6, 7}; - hb.CopyFrom(data_in.data(), data_in.size()); - - std::vector data_out{7, 6, 5, 4, 3, 2, 1, 0}; - hb.CopyTo(data_out.data(), data_out.size()); - - for (size_t i = 0; i < data_in.size(); ++i) { - EXPECT_EQ(data_in[i], data_out[i]); - } -} diff --git a/src/runtime/hexagon/tests/hexagon_buffer_tests.cc b/src/runtime/hexagon/tests/hexagon_buffer_tests.cc index 4078ed9d7e37..b3ff808b6135 100644 --- a/src/runtime/hexagon/tests/hexagon_buffer_tests.cc +++ b/src/runtime/hexagon/tests/hexagon_buffer_tests.cc @@ -17,17 +17,8 @@ * under the License. */ -#include "../hexagon_buffer.h" - #include -<<<<<<<< HEAD:src/runtime/hexagon/hexagon_buffer_tests.cc -<<<<<<< HEAD:tests/cpp/runtime/hexagon_buffer.cc -#include -======= -#include "hexagon_buffer.h" ->>>>>>> 499d5ee4d (HexagonBuffer tests running in sim):src/runtime/hexagon/hexagon/hexagon_buffer_tests.cc -======== ->>>>>>>> 2300f1904 (move to new tests directory):src/runtime/hexagon/hexagon/tests/hexagon_buffer_tests.cc +#include "../hexagon_buffer.h" #include using namespace tvm::runtime; From 4eb97fde68fc6401b88024f3152e9e7901dca451 Mon Sep 17 00:00:00 2001 From: adstraw Date: Fri, 22 Apr 2022 11:11:49 -0700 Subject: [PATCH 08/20] isolate cmake changes to Hexagon --- CMakeLists.txt | 21 ++++++++------ apps/hexagon_api/CMakeLists.txt | 1 - cmake/modules/Hexagon.cmake | 3 -- .../hexagon/tests/hexagon_buffer_tests.cc | 7 ++++- src/runtime/hexagon/tests/run_all_tests.cc | 28 +++++++++++++++++-- 5 files changed, 44 insertions(+), 16 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 7871f1ab3a5d..c19374adef84 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -598,13 +598,15 @@ endif() target_link_libraries(tvm PRIVATE ${TVM_LINKER_LIBS} ${TVM_RUNTIME_LINKER_LIBS}) target_link_libraries(tvm_runtime PRIVATE ${TVM_RUNTIME_LINKER_LIBS}) -include(FetchContent) -FetchContent_Declare(googletest SOURCE_DIR "${USE_HEXAGON_SDK}/utils/googletest/gtest") -set(gtest_force_shared_crt ON CACHE BOOL "" FORCE) -FetchContent_MakeAvailable(googletest) -target_link_libraries(tvm_runtime PUBLIC gtest) -target_link_libraries(tvm PUBLIC gtest) # workaround -include_directories("${USE_HEXAGON_SDK}/utils/googletest/gtest/include") +if(BUILD_FOR_HEXAGON) + include(FetchContent) + FetchContent_Declare(googletest SOURCE_DIR "${USE_HEXAGON_SDK}/utils/googletest/gtest") + set(gtest_force_shared_crt ON CACHE BOOL "" FORCE) + FetchContent_MakeAvailable(googletest) + target_link_libraries(tvm_runtime PUBLIC gtest) + target_link_libraries(tvm PUBLIC gtest) # workaround + include_directories("${USE_HEXAGON_SDK}/utils/googletest/gtest/include") +endif() # Set flags for clang include(cmake/modules/ClangFlags.cmake) @@ -642,7 +644,6 @@ if(GTEST_FOUND) tvm_file_glob(GLOB_RECURSE TEST_SRCS tests/cpp/*.cc) add_executable(cpptest ${TEST_SRCS}) # include runtime files for unit testing - target_include_directories(cpptest PUBLIC "src/runtime") target_link_libraries(cpptest PRIVATE ${TVM_TEST_LIBRARY_NAME} GTest::GTest GTest::Main GTest::gmock pthread dl) set_target_properties(cpptest PROPERTIES EXCLUDE_FROM_ALL 1) set_target_properties(cpptest PROPERTIES EXCLUDE_FROM_DEFAULT_BUILD 1) @@ -657,7 +658,9 @@ add_custom_target(runtime DEPENDS tvm_runtime) # Installation rules install(TARGETS tvm EXPORT ${PROJECT_NAME}Targets DESTINATION lib${LIB_SUFFIX}) install(TARGETS tvm_runtime EXPORT ${PROJECT_NAME}Targets DESTINATION lib${LIB_SUFFIX}) -install(TARGETS gtest EXPORT ${PROJECT_NAME}Targets DESTINATION lib${LIB_SUFFIX}) +if(BUILD_FOR_HEXAGON) + install(TARGETS gtest EXPORT ${PROJECT_NAME}Targets DESTINATION lib${LIB_SUFFIX}) +endif() if (INSTALL_DEV) install( diff --git a/apps/hexagon_api/CMakeLists.txt b/apps/hexagon_api/CMakeLists.txt index 8d83cebbc540..40f070513e3d 100644 --- a/apps/hexagon_api/CMakeLists.txt +++ b/apps/hexagon_api/CMakeLists.txt @@ -39,7 +39,6 @@ ExternalProject_Add(x86_tvm_runtime_rpc "-DUSE_HEXAGON_RPC=ON" "-DBUILD_STATIC_RUNTIME=ON" "-DCMAKE_BUILD_TYPE=${CMAKE_BUILD_TYPE}" - "-DUSE_HEXAGON_SDK=${USE_HEXAGON_SDK}" INSTALL_COMMAND "" BUILD_ALWAYS ON ) diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index 8a6a0c3acf65..e788985fb73a 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -84,9 +84,6 @@ if(NOT USE_HEXAGON) if(BUILD_FOR_HOST) list(APPEND COMPILER_SRCS src/target/opt/build_hexagon_off.cc) endif() - list(APPEND RUNTIME_SRCS src/runtime/hexagon/hexagon_buffer.cc) - list(APPEND RUNTIME_SRCS src/runtime/hexagon/hexagon_common.cc) - list(APPEND RUNTIME_SRCS src/runtime/hexagon/hexagon_user_dma.cc) return() endif() diff --git a/src/runtime/hexagon/tests/hexagon_buffer_tests.cc b/src/runtime/hexagon/tests/hexagon_buffer_tests.cc index b3ff808b6135..4319bd2814b4 100644 --- a/src/runtime/hexagon/tests/hexagon_buffer_tests.cc +++ b/src/runtime/hexagon/tests/hexagon_buffer_tests.cc @@ -17,10 +17,13 @@ * under the License. */ +#if defined(__hexagon__) + #include -#include "../hexagon_buffer.h" #include +#include "../hexagon_buffer.h" + using namespace tvm::runtime; using namespace tvm::runtime::hexagon; @@ -462,3 +465,5 @@ TEST(HexagonBuffer, nd_copy_to) { EXPECT_EQ(data_in[i], data_out[i]); } } + +#endif diff --git a/src/runtime/hexagon/tests/run_all_tests.cc b/src/runtime/hexagon/tests/run_all_tests.cc index 61f3dc1d4909..be8625258c5d 100644 --- a/src/runtime/hexagon/tests/run_all_tests.cc +++ b/src/runtime/hexagon/tests/run_all_tests.cc @@ -1,3 +1,24 @@ +/* + * Licensed to the Apache Software Foundation (ASF) under one + * or more contributor license agreements. See the NOTICE file + * distributed with this work for additional information + * regarding copyright ownership. The ASF licenses this file + * to you under the Apache License, Version 2.0 (the + * "License"); you may not use this file except in compliance + * with the License. You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, + * software distributed under the License is distributed on an + * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY + * KIND, either express or implied. See the License for the + * specific language governing permissions and limitations + * under the License. + */ + +#if defined(__hexagon__) + #include #include @@ -7,9 +28,12 @@ namespace tvm { namespace runtime { namespace hexagon { - TVM_REGISTER_GLOBAL("hexagon.run_all_tests").set_body([](TVMArgs args, TVMRetValue* rv) { *rv = RUN_ALL_TESTS(); }); -}}} \ No newline at end of file +} +} // namespace runtime +} // namespace tvm + +#endif From 262dc73e53469f069d3794bfa4a6fe2d023e861f Mon Sep 17 00:00:00 2001 From: adstraw Date: Fri, 22 Apr 2022 08:57:24 -0700 Subject: [PATCH 09/20] add gtest init with arguments --- src/runtime/hexagon/tests/run_all_tests.cc | 29 +++++++++++++++++-- .../python/contrib/test_hexagon/unit_tests.py | 13 +++++++-- 2 files changed, 37 insertions(+), 5 deletions(-) diff --git a/src/runtime/hexagon/tests/run_all_tests.cc b/src/runtime/hexagon/tests/run_all_tests.cc index be8625258c5d..466132fd375e 100644 --- a/src/runtime/hexagon/tests/run_all_tests.cc +++ b/src/runtime/hexagon/tests/run_all_tests.cc @@ -19,20 +19,45 @@ #if defined(__hexagon__) +#include #include #include -#include "gtest/gtest.h" +#include +#include + +#include "../src/support/utils.h" namespace tvm { namespace runtime { namespace hexagon { TVM_REGISTER_GLOBAL("hexagon.run_all_tests").set_body([](TVMArgs args, TVMRetValue* rv) { + // gtest args are passed into this packed func as a singular string + // split gtest args using delimiter and build argument vector + std::vector parsed_args = tvm::support::Split(args[0], ' '); + std::vector argv; + + // add executable name + argv.push_back(const_cast("hexagon_run_all_tests")); + + // add parsed arguments + for (int i = 0; i < parsed_args.size(); ++i) { + argv.push_back(const_cast(parsed_args[i].data())); + } + + // end of parsed arguments + argv.push_back(nullptr); + + // set argument count + int argc = argv.size() - 1; + + // initialize gtest with arguments and run + ::testing::InitGoogleTest(&argc, argv.data()); *rv = RUN_ALL_TESTS(); }); -} +} // namespace hexagon } // namespace runtime } // namespace tvm diff --git a/tests/python/contrib/test_hexagon/unit_tests.py b/tests/python/contrib/test_hexagon/unit_tests.py index 435c40a75ea6..7795799d70b0 100644 --- a/tests/python/contrib/test_hexagon/unit_tests.py +++ b/tests/python/contrib/test_hexagon/unit_tests.py @@ -1,11 +1,18 @@ import pytest import numpy as np from tvm.contrib.hexagon.build import HexagonLauncher -#import tvm.contrib.hexagon as hexagon from .conftest import requires_hexagon_toolchain + @requires_hexagon_toolchain def test_cache_read_write_2d(hexagon_session): + # arguments to pass to gtest + # e.g. + # 1) to run all tests use + # gtest_args = "" + # 2) to run all tests with "foo" in their name twice use: + # gtest_args = "--gtest_repeat=2 --gtest_filter=*foo*" + gtest_args = "--gtest_repeat=2 --gtest_filter=*micro*" func = hexagon_session._rpc.get_function("hexagon.run_all_tests") - x = func() - np.testing.assert_equal(x, 0) + result = func(gtest_args) + np.testing.assert_equal(result, 0) From b6b0c4bc4330de31eb286c7d35873e6f1c1bff18 Mon Sep 17 00:00:00 2001 From: adstraw Date: Wed, 27 Apr 2022 10:20:09 -0700 Subject: [PATCH 10/20] add hexagon sources only if building for Hexagon; remove workaround --- CMakeLists.txt | 1 - cmake/modules/Hexagon.cmake | 14 +++++++------- tests/python/contrib/test_hexagon/unit_tests.py | 4 ++-- 3 files changed, 9 insertions(+), 10 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index c19374adef84..3cf4f99f6fd6 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -604,7 +604,6 @@ if(BUILD_FOR_HEXAGON) set(gtest_force_shared_crt ON CACHE BOOL "" FORCE) FetchContent_MakeAvailable(googletest) target_link_libraries(tvm_runtime PUBLIC gtest) - target_link_libraries(tvm PUBLIC gtest) # workaround include_directories("${USE_HEXAGON_SDK}/utils/googletest/gtest/include") endif() diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index e788985fb73a..6eb996e6d2f5 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -116,13 +116,13 @@ function(add_hexagon_wrapper_paths) link_directories("${HEXAGON_TOOLCHAIN}/lib/iss") endfunction() - -# Common sources for TVM runtime with Hexagon support -file_glob_append(RUNTIME_HEXAGON_SRCS - "${TVMRT_SOURCE_DIR}/hexagon/*.cc" - "${TVMRT_SOURCE_DIR}/hexagon/tests/*.cc" -) - +if(BUILD_FOR_HEXAGON OR BUILD_FOR_ANDROID OR USE_HEXAGON_RPC) + # Common sources for TVM runtime with Hexagon support + file_glob_append(RUNTIME_HEXAGON_SRCS + "${TVMRT_SOURCE_DIR}/hexagon/*.cc" + "${TVMRT_SOURCE_DIR}/hexagon/tests/*.cc" + ) +endif() if(BUILD_FOR_HEXAGON) get_hexagon_sdk_property("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}" diff --git a/tests/python/contrib/test_hexagon/unit_tests.py b/tests/python/contrib/test_hexagon/unit_tests.py index 7795799d70b0..535f1cfb6ac1 100644 --- a/tests/python/contrib/test_hexagon/unit_tests.py +++ b/tests/python/contrib/test_hexagon/unit_tests.py @@ -8,11 +8,11 @@ def test_cache_read_write_2d(hexagon_session): # arguments to pass to gtest # e.g. - # 1) to run all tests use + # 1) to run all tests use: # gtest_args = "" # 2) to run all tests with "foo" in their name twice use: # gtest_args = "--gtest_repeat=2 --gtest_filter=*foo*" - gtest_args = "--gtest_repeat=2 --gtest_filter=*micro*" + gtest_args = "" func = hexagon_session._rpc.get_function("hexagon.run_all_tests") result = func(gtest_args) np.testing.assert_equal(result, 0) From 273a2ddaf90e8f3b88b1eadc39090fb33df04822 Mon Sep 17 00:00:00 2001 From: adstraw Date: Wed, 27 Apr 2022 10:36:25 -0700 Subject: [PATCH 11/20] format & lint --- tests/python/contrib/test_hexagon/unit_tests.py | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/tests/python/contrib/test_hexagon/unit_tests.py b/tests/python/contrib/test_hexagon/unit_tests.py index 535f1cfb6ac1..d84a035698f8 100644 --- a/tests/python/contrib/test_hexagon/unit_tests.py +++ b/tests/python/contrib/test_hexagon/unit_tests.py @@ -1,3 +1,20 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you under the Apache License, Version 2.0 (the +# "License"); you may not use this file except in compliance +# with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, +# software distributed under the License is distributed on an +# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY +# KIND, either express or implied. See the License for the +# specific language governing permissions and limitations +# under the License. + import pytest import numpy as np from tvm.contrib.hexagon.build import HexagonLauncher From 89167680144f5cc8cc56aea536df5bc17a1864de Mon Sep 17 00:00:00 2001 From: adstraw Date: Wed, 27 Apr 2022 14:15:41 -0700 Subject: [PATCH 12/20] fix Hexagon build error --- cmake/modules/Hexagon.cmake | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index 6eb996e6d2f5..32bd7e4bafcb 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -116,12 +116,16 @@ function(add_hexagon_wrapper_paths) link_directories("${HEXAGON_TOOLCHAIN}/lib/iss") endfunction() -if(BUILD_FOR_HEXAGON OR BUILD_FOR_ANDROID OR USE_HEXAGON_RPC) +if(BUILD_FOR_HEXAGON OR USE_HEXAGON_RPC) # Common sources for TVM runtime with Hexagon support file_glob_append(RUNTIME_HEXAGON_SRCS "${TVMRT_SOURCE_DIR}/hexagon/*.cc" "${TVMRT_SOURCE_DIR}/hexagon/tests/*.cc" ) +else() + file_glob_append(RUNTIME_HEXAGON_SRCS + "${TVMRT_SOURCE_DIR}/hexagon/hexagon_module.cc" + ) endif() if(BUILD_FOR_HEXAGON) From 39c2275b63654d45be8f1be09b91a8e0b5666b0c Mon Sep 17 00:00:00 2001 From: adstraw Date: Thu, 28 Apr 2022 09:49:03 -0700 Subject: [PATCH 13/20] remove x86 implementation and win32 code --- cmake/modules/Hexagon.cmake | 4 ++- src/runtime/hexagon/hexagon_buffer.cc | 25 ++++--------------- src/runtime/hexagon/hexagon_user_dma.cc | 13 +--------- .../hexagon/tests/hexagon_buffer_tests.cc | 4 --- 4 files changed, 9 insertions(+), 37 deletions(-) diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index 32bd7e4bafcb..e7c3bb43fc1b 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -120,7 +120,6 @@ if(BUILD_FOR_HEXAGON OR USE_HEXAGON_RPC) # Common sources for TVM runtime with Hexagon support file_glob_append(RUNTIME_HEXAGON_SRCS "${TVMRT_SOURCE_DIR}/hexagon/*.cc" - "${TVMRT_SOURCE_DIR}/hexagon/tests/*.cc" ) else() file_glob_append(RUNTIME_HEXAGON_SRCS @@ -129,6 +128,9 @@ else() endif() if(BUILD_FOR_HEXAGON) + file_glob_append(RUNTIME_HEXAGON_SRCS + "${TVMRT_SOURCE_DIR}/hexagon/tests/*.cc" + ) get_hexagon_sdk_property("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}" SDK_INCLUDE SDK_INCLUDE_DIRS QURT_INCLUDE QURT_INCLUDE_DIRS diff --git a/src/runtime/hexagon/hexagon_buffer.cc b/src/runtime/hexagon/hexagon_buffer.cc index cfe2b528bb9f..909d37481147 100644 --- a/src/runtime/hexagon/hexagon_buffer.cc +++ b/src/runtime/hexagon/hexagon_buffer.cc @@ -52,26 +52,15 @@ struct Allocation { struct DDRAllocation : public Allocation { DDRAllocation(size_t nbytes, size_t alignment) : Allocation(nbytes, alignment) { -#ifdef _WIN32 - data_ = _aligned_malloc(nbytes, alignment); - CHECK(data_ != nullptr); -#else int ret = posix_memalign(&data_, alignment, nbytes); CHECK_EQ(ret, 0); -#endif - } - ~DDRAllocation() { -#ifdef _WIN32 - _aligned_free(data_); -#else - free(data_); -#endif } + ~DDRAllocation() { free(data_); } }; -#if defined(__hexagon__) struct VTCMAllocation : public Allocation { VTCMAllocation(size_t nbytes, size_t alignment) : Allocation(nbytes, alignment) { +#if defined(__hexagon__) compute_res_attr_t res_info; HEXAGON_SAFE_CALL(HAP_compute_res_attr_init(&res_info)); @@ -94,20 +83,16 @@ struct VTCMAllocation : public Allocation { LOG(ERROR) << "ERROR: Unable to acquire requeisted resource."; return; } - // LOG(INFO) << "VTCMAllocation() - Context ID: " << context_id_ << ", VTCM ptr: " << data_; +#endif } ~VTCMAllocation() { - // LOG(INFO) << "~VTCMAllocation() - Context ID: " << context_id_ << ", VTCM ptr: " << data_; +#if defined(__hexagon__) HEXAGON_SAFE_CALL(HAP_compute_res_release(context_id_)); data_ = nullptr; +#endif } unsigned int context_id_{0}; }; -#else -struct VTCMAllocation : public DDRAllocation { - VTCMAllocation(size_t nbytes, size_t alignment) : DDRAllocation(nbytes, alignment) {} -}; -#endif template std::unique_ptr Allocator(size_t nbytes, size_t alignment); diff --git a/src/runtime/hexagon/hexagon_user_dma.cc b/src/runtime/hexagon/hexagon_user_dma.cc index 6e286ae8b3f4..9bf7a9f6c1d4 100644 --- a/src/runtime/hexagon/hexagon_user_dma.cc +++ b/src/runtime/hexagon/hexagon_user_dma.cc @@ -68,14 +68,10 @@ int hexagon_user_dma_1d_sync_helper(void* dst, void* src, uint32_t length) { void* dma_desc = nullptr; -#ifdef _WIN32 - dma_desc = _aligned_malloc(DMA_DESC_2D_SIZE, DMA_DESC_2D_SIZE); -#else int ret = posix_memalign(&dma_desc, DMA_DESC_2D_SIZE, DMA_DESC_2D_SIZE); if (ret) { return DMA_FAILURE; } -#endif if (!dma_desc) { return DMA_FAILURE; @@ -98,20 +94,13 @@ int hexagon_user_dma_1d_sync_helper(void* dst, void* src, uint32_t length) { unsigned int status = dmwait() & DM0_STATUS_MASK; unsigned int done = dma_desc_get_done(dma_desc); -#ifdef _WIN32 - _aligned_free(dma_desc); -#else free(dma_desc); -#endif if (status == DM0_STATUS_IDLE && done == DESC_DONE_COMPLETE) { return DMA_SUCCESS; } - return DMA_FAILURE; -#else - memcpy(dst, src, length); - return DMA_SUCCESS; #endif + return DMA_FAILURE; } int hexagon_user_dma_1d_sync(void* dst, void* src, uint32_t length) { diff --git a/src/runtime/hexagon/tests/hexagon_buffer_tests.cc b/src/runtime/hexagon/tests/hexagon_buffer_tests.cc index 4319bd2814b4..6735e7dc2a17 100644 --- a/src/runtime/hexagon/tests/hexagon_buffer_tests.cc +++ b/src/runtime/hexagon/tests/hexagon_buffer_tests.cc @@ -17,8 +17,6 @@ * under the License. */ -#if defined(__hexagon__) - #include #include @@ -465,5 +463,3 @@ TEST(HexagonBuffer, nd_copy_to) { EXPECT_EQ(data_in[i], data_out[i]); } } - -#endif From d1f17841afb33853f7b89fa05e5e444ce4e2e98f Mon Sep 17 00:00:00 2001 From: adstraw Date: Thu, 28 Apr 2022 17:07:18 -0700 Subject: [PATCH 14/20] check if hexagon gtest path exists before linking --- CMakeLists.txt | 8 ++++---- apps/hexagon_api/CMakeLists.txt | 1 + cmake/modules/Hexagon.cmake | 8 +++++--- src/runtime/hexagon/tests/run_all_tests.cc | 4 ---- 4 files changed, 10 insertions(+), 11 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 3cf4f99f6fd6..50e39c8992c5 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -598,13 +598,13 @@ endif() target_link_libraries(tvm PRIVATE ${TVM_LINKER_LIBS} ${TVM_RUNTIME_LINKER_LIBS}) target_link_libraries(tvm_runtime PRIVATE ${TVM_RUNTIME_LINKER_LIBS}) -if(BUILD_FOR_HEXAGON) +if(BUILD_FOR_HEXAGON AND EXISTS ${USE_HEXAGON_GTEST}) include(FetchContent) - FetchContent_Declare(googletest SOURCE_DIR "${USE_HEXAGON_SDK}/utils/googletest/gtest") + FetchContent_Declare(googletest SOURCE_DIR "${USE_HEXAGON_GTEST}") set(gtest_force_shared_crt ON CACHE BOOL "" FORCE) FetchContent_MakeAvailable(googletest) target_link_libraries(tvm_runtime PUBLIC gtest) - include_directories("${USE_HEXAGON_SDK}/utils/googletest/gtest/include") + include_directories("${USE_HEXAGON_GTEST}/include") endif() # Set flags for clang @@ -657,7 +657,7 @@ add_custom_target(runtime DEPENDS tvm_runtime) # Installation rules install(TARGETS tvm EXPORT ${PROJECT_NAME}Targets DESTINATION lib${LIB_SUFFIX}) install(TARGETS tvm_runtime EXPORT ${PROJECT_NAME}Targets DESTINATION lib${LIB_SUFFIX}) -if(BUILD_FOR_HEXAGON) +if(BUILD_FOR_HEXAGON AND EXISTS ${USE_HEXAGON_GTEST}) install(TARGETS gtest EXPORT ${PROJECT_NAME}Targets DESTINATION lib${LIB_SUFFIX}) endif() diff --git a/apps/hexagon_api/CMakeLists.txt b/apps/hexagon_api/CMakeLists.txt index 40f070513e3d..7d4f13443d86 100644 --- a/apps/hexagon_api/CMakeLists.txt +++ b/apps/hexagon_api/CMakeLists.txt @@ -109,6 +109,7 @@ ExternalProject_Add(hexagon_tvm_runtime_rpc "-DCMAKE_BUILD_TYPE=${CMAKE_BUILD_TYPE}" "-DUSE_ALTERNATIVE_LINKER=OFF" "-DUSE_CUSTOM_LOGGING=ON" + "-DUSE_HEXAGON_GTEST=${USE_HEXAGON_SDK}/utils/googletest/gtest" INSTALL_COMMAND "" BUILD_ALWAYS ON ) diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index e7c3bb43fc1b..9627610e7e9b 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -128,9 +128,11 @@ else() endif() if(BUILD_FOR_HEXAGON) - file_glob_append(RUNTIME_HEXAGON_SRCS - "${TVMRT_SOURCE_DIR}/hexagon/tests/*.cc" - ) + if(EXISTS ${USE_HEXAGON_GTEST}) + file_glob_append(RUNTIME_HEXAGON_SRCS + "${TVMRT_SOURCE_DIR}/hexagon/tests/*.cc" + ) + endif() get_hexagon_sdk_property("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}" SDK_INCLUDE SDK_INCLUDE_DIRS QURT_INCLUDE QURT_INCLUDE_DIRS diff --git a/src/runtime/hexagon/tests/run_all_tests.cc b/src/runtime/hexagon/tests/run_all_tests.cc index 466132fd375e..166d89b63566 100644 --- a/src/runtime/hexagon/tests/run_all_tests.cc +++ b/src/runtime/hexagon/tests/run_all_tests.cc @@ -17,8 +17,6 @@ * under the License. */ -#if defined(__hexagon__) - #include #include #include @@ -60,5 +58,3 @@ TVM_REGISTER_GLOBAL("hexagon.run_all_tests").set_body([](TVMArgs args, TVMRetVal } // namespace hexagon } // namespace runtime } // namespace tvm - -#endif From 32e6dd2f2fc6f2e534b189ce3f9adb6e08ab7d63 Mon Sep 17 00:00:00 2001 From: adstraw Date: Fri, 29 Apr 2022 09:48:58 -0700 Subject: [PATCH 15/20] make USE_HEXAGON_GTEST an optional cmake param --- CMakeLists.txt | 4 ++-- apps/hexagon_api/CMakeLists.txt | 10 +++++++++- cmake/modules/Hexagon.cmake | 2 +- tests/python/contrib/test_hexagon/unit_tests.py | 9 +++++++-- 4 files changed, 19 insertions(+), 6 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 50e39c8992c5..4371dbf7e559 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -598,7 +598,7 @@ endif() target_link_libraries(tvm PRIVATE ${TVM_LINKER_LIBS} ${TVM_RUNTIME_LINKER_LIBS}) target_link_libraries(tvm_runtime PRIVATE ${TVM_RUNTIME_LINKER_LIBS}) -if(BUILD_FOR_HEXAGON AND EXISTS ${USE_HEXAGON_GTEST}) +if(BUILD_FOR_HEXAGON AND DEFINED USE_HEXAGON_GTEST AND EXISTS ${USE_HEXAGON_GTEST}) include(FetchContent) FetchContent_Declare(googletest SOURCE_DIR "${USE_HEXAGON_GTEST}") set(gtest_force_shared_crt ON CACHE BOOL "" FORCE) @@ -657,7 +657,7 @@ add_custom_target(runtime DEPENDS tvm_runtime) # Installation rules install(TARGETS tvm EXPORT ${PROJECT_NAME}Targets DESTINATION lib${LIB_SUFFIX}) install(TARGETS tvm_runtime EXPORT ${PROJECT_NAME}Targets DESTINATION lib${LIB_SUFFIX}) -if(BUILD_FOR_HEXAGON AND EXISTS ${USE_HEXAGON_GTEST}) +if(BUILD_FOR_HEXAGON AND DEFINED USE_HEXAGON_GTEST AND EXISTS ${USE_HEXAGON_GTEST}) install(TARGETS gtest EXPORT ${PROJECT_NAME}Targets DESTINATION lib${LIB_SUFFIX}) endif() diff --git a/apps/hexagon_api/CMakeLists.txt b/apps/hexagon_api/CMakeLists.txt index 7d4f13443d86..1cb7931523cb 100644 --- a/apps/hexagon_api/CMakeLists.txt +++ b/apps/hexagon_api/CMakeLists.txt @@ -13,6 +13,7 @@ include(ExternalProject) # USE_HEXAGON_TOOLCHAIN (Path to Hexagon toolchain ending with "Tools") # Optional variable: # USE_OUTPUT_BINARY_DIR (Path to copy the output binaries to) +# USE_HEXAGON_GTEST (Path to Hexagon gtest) set(TVM_SOURCE_DIR "${CMAKE_SOURCE_DIR}/../..") @@ -23,6 +24,11 @@ else() endif() file(MAKE_DIRECTORY ${HEXAGON_API_BINARY_DIR}) +if(DEFINED USE_HEXAGON_GTEST AND NOT EXISTS $USE_HEXAGON_GTEST) + message(WARNING "The specified Hexagon gtest path USE_HEXAGON_GTEST = ${USE_HEXAGON_GTEST} does not exist. Disabling Hexagon gtest support.") + unset(USE_HEXAGON_GTEST) +endif() + # Build X86 binaries: # - tvm_rpc_x86 @@ -109,7 +115,9 @@ ExternalProject_Add(hexagon_tvm_runtime_rpc "-DCMAKE_BUILD_TYPE=${CMAKE_BUILD_TYPE}" "-DUSE_ALTERNATIVE_LINKER=OFF" "-DUSE_CUSTOM_LOGGING=ON" - "-DUSE_HEXAGON_GTEST=${USE_HEXAGON_SDK}/utils/googletest/gtest" + if(DEFINED USE_HEXAGON_GTEST) + "-DUSE_HEXAGON_GTEST=${USE_HEXAGON_GTEST}" + endif() INSTALL_COMMAND "" BUILD_ALWAYS ON ) diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index 9627610e7e9b..d8cf26bfdb70 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -128,7 +128,7 @@ else() endif() if(BUILD_FOR_HEXAGON) - if(EXISTS ${USE_HEXAGON_GTEST}) + if(DEFINED USE_HEXAGON_GTEST AND EXISTS ${USE_HEXAGON_GTEST}) file_glob_append(RUNTIME_HEXAGON_SRCS "${TVMRT_SOURCE_DIR}/hexagon/tests/*.cc" ) diff --git a/tests/python/contrib/test_hexagon/unit_tests.py b/tests/python/contrib/test_hexagon/unit_tests.py index d84a035698f8..f4e539a76c35 100644 --- a/tests/python/contrib/test_hexagon/unit_tests.py +++ b/tests/python/contrib/test_hexagon/unit_tests.py @@ -30,6 +30,11 @@ def test_cache_read_write_2d(hexagon_session): # 2) to run all tests with "foo" in their name twice use: # gtest_args = "--gtest_repeat=2 --gtest_filter=*foo*" gtest_args = "" - func = hexagon_session._rpc.get_function("hexagon.run_all_tests") - result = func(gtest_args) + try: + func = hexagon_session._rpc.get_function("hexagon.run_all_tests") + result = func(gtest_args) + except: + print("WARNING: Skipping Hexagon unit tests because they have not been enabled") + result = 0 + np.testing.assert_equal(result, 0) From f12efd7e19091f01033b7ca1cf17954620f088a8 Mon Sep 17 00:00:00 2001 From: adstraw Date: Fri, 29 Apr 2022 10:09:29 -0700 Subject: [PATCH 16/20] turn on Hexagon gtest in Hexagon CI --- tests/scripts/task_build_hexagon_api.sh | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tests/scripts/task_build_hexagon_api.sh b/tests/scripts/task_build_hexagon_api.sh index 89b7545f4d89..08096706982f 100755 --- a/tests/scripts/task_build_hexagon_api.sh +++ b/tests/scripts/task_build_hexagon_api.sh @@ -33,6 +33,7 @@ cmake -DANDROID_ABI=arm64-v8a \ -DUSE_HEXAGON_ARCH=v68 \ -DUSE_HEXAGON_SDK="${HEXAGON_SDK_PATH}" \ -DUSE_HEXAGON_TOOLCHAIN="${HEXAGON_TOOLCHAIN}" \ - -DUSE_OUTPUT_BINARY_DIR="${output_binary_directory}" .. + -DUSE_OUTPUT_BINARY_DIR="${output_binary_directory}" \ + -DUSE_HEXAGON_GTEST="${HEXAGON_SDK_PATH}/utils/googletest/gtest" .. make -j$(nproc) From fa19d826f1c286e2c9ccc8770ff00cd71ddf49a1 Mon Sep 17 00:00:00 2001 From: adstraw Date: Fri, 29 Apr 2022 14:27:06 -0700 Subject: [PATCH 17/20] Hexagon unit tests should fail if run without proper gtest linkage --- tests/python/contrib/test_hexagon/unit_tests.py | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/tests/python/contrib/test_hexagon/unit_tests.py b/tests/python/contrib/test_hexagon/unit_tests.py index f4e539a76c35..d340cba5b150 100644 --- a/tests/python/contrib/test_hexagon/unit_tests.py +++ b/tests/python/contrib/test_hexagon/unit_tests.py @@ -34,7 +34,9 @@ def test_cache_read_write_2d(hexagon_session): func = hexagon_session._rpc.get_function("hexagon.run_all_tests") result = func(gtest_args) except: - print("WARNING: Skipping Hexagon unit tests because they have not been enabled") - result = 0 + print( + "This test requires the USE_HEXAGON_GTEST cmake flag to be specified with a path to a Hexagon gtest version normally located at /path/to/hexagon/sdk/utils/googletest/gtest" + ) + result = 1 np.testing.assert_equal(result, 0) From e6222e8a83f5119b5a84e4efd7acc57f7c047465 Mon Sep 17 00:00:00 2001 From: adstraw Date: Mon, 2 May 2022 16:16:27 -0700 Subject: [PATCH 18/20] add tvm option; move Hexagon tests to test/cpp-runtime/hexagon --- CMakeLists.txt | 1 + apps/hexagon_api/CMakeLists.txt | 12 ++++++++---- cmake/modules/Hexagon.cmake | 2 +- .../cpp-runtime/hexagon}/hexagon_buffer_tests.cc | 2 +- .../cpp-runtime/hexagon}/run_all_tests.cc | 0 5 files changed, 11 insertions(+), 6 deletions(-) rename {src/runtime/hexagon/tests => tests/cpp-runtime/hexagon}/hexagon_buffer_tests.cc (99%) rename {src/runtime/hexagon/tests => tests/cpp-runtime/hexagon}/run_all_tests.cc (100%) diff --git a/CMakeLists.txt b/CMakeLists.txt index 4371dbf7e559..873d794871b7 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -43,6 +43,7 @@ tvm_option(ROCM_PATH "The path to rocm" /opt/rocm) tvm_option(USE_HEXAGON "Build with Hexagon support" OFF) tvm_option(USE_HEXAGON_SDK "Path to the Hexagon SDK root (required for Hexagon support)" /path/to/sdk) tvm_option(USE_HEXAGON_RPC "Enable Hexagon RPC using minRPC implementation over Android." OFF) +tvm_option(USE_HEXAGON_GTEST "Path to Hexagon specific gtest version for runtime cpp tests." /path/to/hexagon/gtest) tvm_option(USE_RPC "Build with RPC" ON) tvm_option(USE_THREADS "Build with thread support" ON) tvm_option(USE_LLVM "Build with LLVM, can be set to specific llvm-config path" OFF) diff --git a/apps/hexagon_api/CMakeLists.txt b/apps/hexagon_api/CMakeLists.txt index 1cb7931523cb..1f9e982970c3 100644 --- a/apps/hexagon_api/CMakeLists.txt +++ b/apps/hexagon_api/CMakeLists.txt @@ -13,7 +13,7 @@ include(ExternalProject) # USE_HEXAGON_TOOLCHAIN (Path to Hexagon toolchain ending with "Tools") # Optional variable: # USE_OUTPUT_BINARY_DIR (Path to copy the output binaries to) -# USE_HEXAGON_GTEST (Path to Hexagon gtest) +# USE_HEXAGON_GTEST (Path to Hexagon specific gtest version) set(TVM_SOURCE_DIR "${CMAKE_SOURCE_DIR}/../..") @@ -24,9 +24,13 @@ else() endif() file(MAKE_DIRECTORY ${HEXAGON_API_BINARY_DIR}) -if(DEFINED USE_HEXAGON_GTEST AND NOT EXISTS $USE_HEXAGON_GTEST) - message(WARNING "The specified Hexagon gtest path USE_HEXAGON_GTEST = ${USE_HEXAGON_GTEST} does not exist. Disabling Hexagon gtest support.") - unset(USE_HEXAGON_GTEST) +if(DEFINED USE_HEXAGON_GTEST) + if(EXISTS ${USE_HEXAGON_GTEST}) + message(STATUS "Found Hexagon gtest at ${USE_HEXAGON_GTEST}") + else() + message(WARNING "Could not find Hexagon gtest at ${USE_HEXAGON_GTEST}. Disabling Hexagon gtest support.") + unset(USE_HEXAGON_GTEST) + endif() endif() # Build X86 binaries: diff --git a/cmake/modules/Hexagon.cmake b/cmake/modules/Hexagon.cmake index d8cf26bfdb70..d45311a87fec 100644 --- a/cmake/modules/Hexagon.cmake +++ b/cmake/modules/Hexagon.cmake @@ -130,7 +130,7 @@ endif() if(BUILD_FOR_HEXAGON) if(DEFINED USE_HEXAGON_GTEST AND EXISTS ${USE_HEXAGON_GTEST}) file_glob_append(RUNTIME_HEXAGON_SRCS - "${TVMRT_SOURCE_DIR}/hexagon/tests/*.cc" + "${CMAKE_SOURCE_DIR}/tests/cpp-runtime/hexagon/*.cc" ) endif() get_hexagon_sdk_property("${USE_HEXAGON_SDK}" "${USE_HEXAGON_ARCH}" diff --git a/src/runtime/hexagon/tests/hexagon_buffer_tests.cc b/tests/cpp-runtime/hexagon/hexagon_buffer_tests.cc similarity index 99% rename from src/runtime/hexagon/tests/hexagon_buffer_tests.cc rename to tests/cpp-runtime/hexagon/hexagon_buffer_tests.cc index 6735e7dc2a17..803e67785413 100644 --- a/src/runtime/hexagon/tests/hexagon_buffer_tests.cc +++ b/tests/cpp-runtime/hexagon/hexagon_buffer_tests.cc @@ -20,7 +20,7 @@ #include #include -#include "../hexagon_buffer.h" +#include "../src/runtime/hexagon/hexagon_buffer.h" using namespace tvm::runtime; using namespace tvm::runtime::hexagon; diff --git a/src/runtime/hexagon/tests/run_all_tests.cc b/tests/cpp-runtime/hexagon/run_all_tests.cc similarity index 100% rename from src/runtime/hexagon/tests/run_all_tests.cc rename to tests/cpp-runtime/hexagon/run_all_tests.cc From 816f3bec24c54f7ae7f62bdf2a38f6b03d3668f4 Mon Sep 17 00:00:00 2001 From: adstraw Date: Mon, 2 May 2022 16:34:29 -0700 Subject: [PATCH 19/20] add libinfo --- cmake/modules/LibInfo.cmake | 1 + src/support/libinfo.cc | 5 +++++ 2 files changed, 6 insertions(+) diff --git a/cmake/modules/LibInfo.cmake b/cmake/modules/LibInfo.cmake index eefa7036a0ff..76ddbede8ac0 100644 --- a/cmake/modules/LibInfo.cmake +++ b/cmake/modules/LibInfo.cmake @@ -74,6 +74,7 @@ function(add_lib_info src_file) TVM_INFO_USE_HEXAGON="${USE_HEXAGON}" TVM_INFO_USE_HEXAGON_RPC="${USE_HEXAGON_RPC}" TVM_INFO_USE_HEXAGON_SDK="${USE_HEXAGON_SDK}" + TVM_INFO_USE_HEXAGON_GTEST="${USE_HEXAGON_GTEST}" TVM_INFO_USE_IOS_RPC="${USE_IOS_RPC}" TVM_INFO_USE_KHRONOS_SPIRV="${USE_KHRONOS_SPIRV}" TVM_INFO_USE_LIBBACKTRACE="${USE_LIBBACKTRACE}" diff --git a/src/support/libinfo.cc b/src/support/libinfo.cc index c6cf916ae8a2..e6f322885e3a 100644 --- a/src/support/libinfo.cc +++ b/src/support/libinfo.cc @@ -67,6 +67,10 @@ #define TVM_INFO_USE_HEXAGON_SDK "NOT-FOUND" #endif +#ifndef TVM_INFO_USE_HEXAGON_GTEST +#define TVM_INFO_USE_HEXAGON_GTEST "NOT-FOUND" +#endif + #ifndef TVM_INFO_USE_RPC #define TVM_INFO_USE_RPC "NOT-FOUND" #endif @@ -267,6 +271,7 @@ TVM_DLL Map GetLibInfo() { {"USE_HEXAGON", TVM_INFO_USE_HEXAGON}, {"USE_HEXAGON_RPC", TVM_INFO_USE_HEXAGON_RPC}, {"USE_HEXAGON_SDK", TVM_INFO_USE_HEXAGON_SDK}, + {"USE_HEXAGON_GTEST", TVM_INFO_USE_HEXAGON_GTEST}, {"USE_IOS_RPC", TVM_INFO_USE_IOS_RPC}, {"USE_KHRONOS_SPIRV", TVM_INFO_USE_KHRONOS_SPIRV}, {"USE_LIBBACKTRACE", TVM_INFO_USE_LIBBACKTRACE}, From 3946e23aa27e8894099cb8ecb5d789aeb7b241e7 Mon Sep 17 00:00:00 2001 From: adstraw Date: Tue, 3 May 2022 08:28:56 -0700 Subject: [PATCH 20/20] trigger ci